[tint][ir][val] Improve checks for shader IO annotations
Fixes: 379652398
Change-Id: I110c31169fd8f116b4e658309a505e57341f44fb
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/215814
Commit-Queue: Ryan Harrison <rharrison@chromium.org>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Auto-Submit: Ryan Harrison <rharrison@chromium.org>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/core/ir/transform/demote_to_helper_test.cc b/src/tint/lang/core/ir/transform/demote_to_helper_test.cc
index 861de81..4e0d1c6 100644
--- a/src/tint/lang/core/ir/transform/demote_to_helper_test.cc
+++ b/src/tint/lang/core/ir/transform/demote_to_helper_test.cc
@@ -538,7 +538,11 @@
auto* front_facing = b.FunctionParam("front_facing", ty.bool_());
front_facing->SetBuiltin(BuiltinValue::kFrontFacing);
+
auto* coord = b.FunctionParam("coord", ty.vec2<i32>());
+ IOAttributes coord_attr;
+ coord_attr.location = 0;
+ coord->SetAttributes(coord_attr);
auto* ep = b.Function("ep", ty.f32(), Function::PipelineStage::kFragment);
ep->SetParams({front_facing, coord});
ep->SetReturnLocation(0_u);
@@ -559,7 +563,7 @@
%texture:ptr<handle, texture_storage_2d<r32float, write>, read> = var @binding_point(0, 0)
}
-%ep = @fragment func(%front_facing:bool [@front_facing], %coord:vec2<i32>):f32 [@location(0)] {
+%ep = @fragment func(%front_facing:bool [@front_facing], %coord:vec2<i32> [@location(0)]):f32 [@location(0)] {
$B2: {
if %front_facing [t: $B3] { # if_1
$B3: { # true
@@ -581,7 +585,7 @@
%continue_execution:ptr<private, bool, read_write> = var, true
}
-%ep = @fragment func(%front_facing:bool [@front_facing], %coord:vec2<i32>):f32 [@location(0)] {
+%ep = @fragment func(%front_facing:bool [@front_facing], %coord:vec2<i32> [@location(0)]):f32 [@location(0)] {
$B2: {
if %front_facing [t: $B3] { # if_1
$B3: { # true
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index 5bafb6e..0744f3f 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -140,16 +140,6 @@
return false;
}
-/// @returns true if @p attr contains both a location and builtin decoration
-bool HasLocationAndBuiltin(const tint::core::IOAttributes& attr) {
- return attr.builtin.has_value() && attr.location.has_value();
-}
-
-/// @returns true if @p attr contains one of location or builtin decoration
-bool HasLocationOrBuiltin(const tint::core::IOAttributes& attr) {
- return attr.builtin.has_value() || attr.location.has_value();
-}
-
/// @return true if @param attr does not have invariant decoration or if it also has position
/// decoration
bool InvariantOnlyIfAlsoPosition(const tint::core::IOAttributes& attr) {
@@ -187,7 +177,7 @@
/// If the type that the attributes are attached to is a struct, the check is run over the members,
/// otherwise it run on the attributes directly.
///
-/// @param msg_anchor what to associate errors with, i.e. the 'foo' of AddError(foo)
+/// @param msg_anchor what to associate errors with, e.g. the 'foo' of AddError(foo)
/// @param ty_attr the directly attached attributes
/// @param ty the type of the thing that the attributes are attached to
/// @param is_not_struct_impl has the signature 'void(const MSG_ANCHOR*, const IOAttributes&)' and
@@ -237,7 +227,7 @@
/// If the type that the attributes are attached to is a struct, the check is run over the members,
/// otherwise it run on the attributes directly.
///
-/// @param msg_anchor what to associate errors with, i.e. the 'foo' of AddError(foo)
+/// @param msg_anchor what to associate errors with, e.g. the 'foo' of AddError(foo)
/// @param ty_attr the directly attached attributes
/// @param ty the type of the thing that the attributes are attached to
/// @param is_not_struct_impl has the signature 'void(const MSG_ANCHOR*, const IOAttributes&, const
@@ -650,6 +640,86 @@
return Success;
}
+// Annotations that can be associated with a value that are used for shader IO, e.g. binding_points,
+// @location, being in workgroup address space, etc.
+enum class IOAnnotation : uint8_t {
+ /// @group + @binding
+ kBindingPoint,
+ /// @location
+ kLocation,
+ /// @builtin(...)
+ kBuiltin,
+ /// Pointer to Workgroup address space
+ kWorkgroup,
+ /// @color
+ kColor,
+};
+
+/// @returns text describing the annotation for error logging
+std::string ToString(IOAnnotation value) {
+ switch (value) {
+ case IOAnnotation::kBindingPoint:
+ return "@group + @binding";
+ case IOAnnotation::kLocation:
+ return "@location";
+ case IOAnnotation::kBuiltin:
+ return "built-in";
+ case IOAnnotation::kWorkgroup:
+ return "<workgroup>";
+ case IOAnnotation::kColor:
+ return "@color";
+ }
+ TINT_ICE() << "Unknown enum passed to ToString(IOAnnotation)";
+}
+
+/// @returns a human-readable string of all the entries in a set of IOAnnotations
+std::string ToString(const EnumSet<IOAnnotation>& values) {
+ std::stringstream result;
+ result << "[ ";
+ bool first = true;
+ for (auto v : values) {
+ if (!first) {
+ result << ", ";
+ }
+ first = false;
+ result << ToString(v);
+ }
+ result << " ]";
+ return result.str();
+}
+
+/// Adds appropriate entries to annotations, based on what values are present in attributes
+/// @param annotations the set to updated
+/// @param attr the attributes to be examined
+/// @returns Success if none of the values being added where already present, otherwise returns the
+/// first non-unique value as a Failure
+Result<SuccessType, IOAnnotation> AddIOAnnotationsFromIOAttributes(
+ EnumSet<IOAnnotation>& annotations,
+ const IOAttributes& attr) {
+ if (attr.location.has_value()) {
+ if (annotations.Contains(IOAnnotation::kLocation)) {
+ return IOAnnotation::kLocation;
+ }
+ annotations.Add(IOAnnotation::kLocation);
+ }
+
+ if (attr.builtin.has_value()) {
+ if (annotations.Contains(IOAnnotation::kBuiltin)) {
+ return IOAnnotation::kBuiltin;
+ }
+ annotations.Add(IOAnnotation::kBuiltin);
+ }
+
+ if (attr.color.has_value()) {
+ if (annotations.Contains(IOAnnotation::kColor)) {
+ return IOAnnotation::kColor;
+ }
+ annotations.Add(IOAnnotation::kColor);
+ }
+
+ return Success;
+}
+
/// The core IR validator.
class Validator {
public:
@@ -935,29 +1005,6 @@
};
}
- /// @returns a function that validates that location and builtin attributes are not present at
- /// the same time
- /// @param err error message to log when check fails
- template <typename MSG_ANCHOR>
- auto CheckDoesNotHaveBothLocationAndBuiltinFunc(const std::string& err) {
- return [this, err](const MSG_ANCHOR* msg_anchor, const IOAttributes& attr) {
- if (HasLocationAndBuiltin(attr)) {
- AddError(msg_anchor) << err;
- }
- };
- }
-
- /// @returns a function that validates that either a location or builtin attribute are present
- /// @param err error message to log when check fails
- template <typename MSG_ANCHOR>
- auto CheckHasLocationOrBuiltinFunc(const std::string& err) {
- return [this, err](const MSG_ANCHOR* msg_anchor, const IOAttributes& attr) {
- if (!HasLocationOrBuiltin(attr)) {
- AddError(msg_anchor) << err;
- }
- };
- }
-
/// @returns a function that validates that type is bool iff decorated with
/// @builtin(front_facing)
/// @param err error message to log when check fails
@@ -1006,6 +1053,22 @@
AddressSpace address_space,
const std::string& target_str = "variable");
+ /// Validates shader IO annotations for entry point input/output
+ /// Note: Call is required to ensure that the value being validated is associated with an entry
+ /// point function
+ /// @param ty type of the value under test
+ /// @param binding_point the binding information associated with the value
+ /// @param attr IO attributes associated with the values
+ /// @param target_str string to insert in error message describing what has a binding_point,
+ /// something like 'input param' or 'return value'
+ /// @returns Success if one, and only one, shader IO is present, otherwise a Failure with the
+ /// error reason is returned
+ Result<SuccessType, std::string> ValidateShaderIOAnnotations(
+ const core::type::Type* ty,
+ const std::optional<struct BindingPoint>& binding_point,
+ const core::IOAttributes& attr,
+ const std::string& target_str);
+
/// Validates the given let
/// @param l the let to validate
void CheckLet(const Let* l);
@@ -1874,15 +1937,6 @@
"invariant can only decorate a param member iff it is also "
"decorated with position"));
- if (func->Stage() != Function::PipelineStage::kUndefined) {
- CheckFunctionParamAttributes(
- param,
- CheckDoesNotHaveBothLocationAndBuiltinFunc<FunctionParam>(
- "a builtin and location cannot be both declared for a param"),
- CheckDoesNotHaveBothLocationAndBuiltinFunc<FunctionParam>(
- "a builtin and location cannot be both declared for a struct member"));
- }
-
if (func->Stage() == Function::PipelineStage::kFragment) {
CheckFunctionParamAttributesAndType(
param,
@@ -1910,9 +1964,19 @@
}
if (func->Stage() != Function::PipelineStage::kUndefined) {
- auto result = ValidateBindingPoint(param->BindingPoint(), address_space, "input param");
- if (result != Success) {
- AddError(param) << result.Failure();
+ {
+ auto result = ValidateShaderIOAnnotations(param->Type(), param->BindingPoint(),
+ param->Attributes(), "input param");
+ if (result != Success) {
+ AddError(param) << result.Failure();
+ }
+ }
+ {
+ auto result =
+ ValidateBindingPoint(param->BindingPoint(), address_space, "input param");
+ if (result != Success) {
+ AddError(param) << result.Failure();
+ }
}
} else {
if (param->BindingPoint().has_value()) {
@@ -1937,16 +2001,6 @@
"invariant can only decorate outputs iff they are also position builtins"),
CheckInvariantFunc<Function>(
"invariant can only decorate output members iff they are also position builtins"));
-
- if (func->Stage() != Function::PipelineStage::kUndefined) {
- CheckFunctionReturnAttributes(
- func,
- CheckDoesNotHaveBothLocationAndBuiltinFunc<Function>(
- "a builtin and location cannot be both declared for a function return"),
- CheckDoesNotHaveBothLocationAndBuiltinFunc<Function>(
- "a builtin and location cannot be both declared for a struct member"));
- }
-
// void needs to be filtered out, since it isn't constructible, but used in the IR when no
// return is specified.
if (DAWN_UNLIKELY(!func->ReturnType()->Is<core::type::Void>() &&
@@ -1968,20 +2022,13 @@
}
}
- if (func->Stage() == Function::PipelineStage::kFragment) {
- if (!func->ReturnType()->Is<core::type::Void>()) {
- CheckFunctionReturnAttributes(
- func,
- CheckHasLocationOrBuiltinFunc<Function>(
- "a non-void return for an entry point must have a "
- "builtin or location decoration"),
- CheckHasLocationOrBuiltinFunc<Function>(
- "members of struct used for returns of entry points must have a builtin or "
- "location decoration"));
- }
- }
-
if (func->Stage() != Function::PipelineStage::kUndefined) {
+ auto result = ValidateShaderIOAnnotations(func->ReturnType(), std::nullopt,
+ func->ReturnAttributes(), "return values");
+ if (result != Success) {
+ AddError(func) << result.Failure();
+ }
+
CheckFunctionReturnAttributesAndType(
func, CheckFrontFacingIfBoolFunc<Function>("entry point returns can not be bool"),
CheckFrontFacingIfBoolFunc<Function>("entry point return members can not be bool"));
@@ -2338,20 +2385,10 @@
if (var->Block() == mod_.root_block) {
if (mv->AddressSpace() == AddressSpace::kIn || mv->AddressSpace() == AddressSpace::kOut) {
- if (HasLocationAndBuiltin(var->Attributes())) {
- AddError(var)
- << "a builtin and location cannot be both declared for a module scope var";
- return;
- }
-
- if (auto* s = var->Result(0)->Type()->UnwrapPtrOrRef()->As<core::type::Struct>()) {
- for (auto* mem : s->Members()) {
- if (HasLocationAndBuiltin(mem->Attributes())) {
- AddError(var) << "a builtin and location cannot be both declared for a "
- "module scope var struct member";
- return;
- }
- }
+ auto result = ValidateShaderIOAnnotations(var->Result(0)->Type(), var->BindingPoint(),
+ var->Attributes(), "module scope variable");
+ if (result != Success) {
+ AddError(var) << result.Failure();
}
}
}
@@ -2384,6 +2421,70 @@
return Success;
}
+Result<SuccessType, std::string> Validator::ValidateShaderIOAnnotations(
+ const core::type::Type* ty,
+ const std::optional<struct BindingPoint>& binding_point,
+ const core::IOAttributes& attr,
+ const std::string& target_str) {
+ EnumSet<IOAnnotation> annotations;
+ // Since there is no entries in the set at this point, this should never fail.
+ TINT_ASSERT(AddIOAnnotationsFromIOAttributes(annotations, attr) == Success);
+ if (binding_point.has_value()) {
+ annotations.Add(IOAnnotation::kBindingPoint);
+ }
+ if (auto* mv = ty->As<core::type::MemoryView>()) {
+ if (mv->AddressSpace() == AddressSpace::kWorkgroup) {
+ annotations.Add(IOAnnotation::kWorkgroup);
+ }
+ }
+
+ // void being annotated should never occur
+ TINT_ASSERT(!ty->Is<core::type::Void>() || annotations.Empty());
+ if (ty->Is<core::type::Void>()) {
+ return Success;
+ }
+
+ if (auto* ty_struct = ty->UnwrapPtrOrRef()->As<core::type::Struct>()) {
+ for (const auto* mem : ty_struct->Members()) {
+ EnumSet<IOAnnotation> mem_annotations = annotations;
+ auto add_result = AddIOAnnotationsFromIOAttributes(mem_annotations, mem->Attributes());
+ if (add_result != Success) {
+ return target_str +
+ " struct member has same IO annotation, as top-level struct, '" +
+ ToString(add_result.Failure()) + "'";
+ }
+
+ if (capabilities_.Contains(Capability::kAllowPointersInStructures)) {
+ if (auto* mv = mem->Type()->As<core::type::MemoryView>()) {
+ if (mv->AddressSpace() == AddressSpace::kWorkgroup) {
+ mem_annotations.Add(IOAnnotation::kWorkgroup);
+ }
+ }
+ }
+
+ if (mem_annotations.Empty()) {
+ return target_str +
+ " struct members must have at least one IO annotation, e.g. a binding "
+ "point, a location, etc";
+ }
+
+ if (mem_annotations.Size() > 1) {
+ return target_str + " struct member has more than one IO annotation, " +
+ ToString(mem_annotations);
+ }
+ }
+ } else {
+ if (annotations.Empty()) {
+ return target_str +
+ " must have at least one IO annotation, e.g. a binding point, a location, etc";
+ }
+ if (annotations.Size() > 1) {
+ return target_str + " has more than one IO annotation, " + ToString(annotations);
+ }
+ }
+ return Success;
+}
+
void Validator::CheckLet(const Let* l) {
if (!CheckResultsAndOperands(l, Let::kNumResults, Let::kNumOperands)) {
return;
diff --git a/src/tint/lang/core/ir/validator_test.cc b/src/tint/lang/core/ir/validator_test.cc
index 774ad8c..f1187cd 100644
--- a/src/tint/lang/core/ir/validator_test.cc
+++ b/src/tint/lang/core/ir/validator_test.cc
@@ -66,9 +66,7 @@
/// Builds and returns a basic 'vertex' entry point function, named @p name
Function* VertexEntryPoint(const std::string& name = "f") {
auto* f = b.Function(name, ty.vec4<f32>(), Function::PipelineStage::kVertex);
- IOAttributes attr;
- attr.builtin = BuiltinValue::kPosition;
- f->SetReturnAttributes(attr);
+ f->SetReturnBuiltin(BuiltinValue::kPosition);
return f;
}
@@ -78,10 +76,8 @@
const std::string& name,
BuiltinValue builtin,
const core::type::Type* type) {
- IOAttributes attr;
- attr.builtin = builtin;
auto* p = b.FunctionParam(name, type);
- p->SetAttributes(attr);
+ p->SetBuiltin(builtin);
func->AppendParam(p);
}
@@ -438,14 +434,12 @@
)");
}
-TEST_F(IR_ValidatorTest, Function_Param_BothLocationAndBuiltin) {
+TEST_F(IR_ValidatorTest, Function_Param_MultipleIOAnnotations) {
auto* f = FragmentEntryPoint("my_func");
auto* p = b.FunctionParam("my_param", ty.vec4<f32>());
- IOAttributes attr;
- attr.builtin = BuiltinValue::kPosition;
- attr.location = 0;
- p->SetAttributes(attr);
+ p->SetBuiltin(BuiltinValue::kPosition);
+ p->SetLocation(0);
f->SetParams({p});
b.Append(f->Block(), [&] { b.Return(f); });
@@ -453,7 +447,7 @@
auto res = ir::Validate(mod);
ASSERT_NE(res, Success);
EXPECT_EQ(res.Failure().reason.Str(),
- R"(:1:27 error: a builtin and location cannot be both declared for a param
+ R"(:1:27 error: input param has more than one IO annotation, [ @location, built-in ]
%my_func = @fragment func(%my_param:vec4<f32> [@location(0), @position]):void {
^^^^^^^^^^^^^^^^^^^
@@ -466,12 +460,12 @@
)");
}
-TEST_F(IR_ValidatorTest, Function_Param_Struct_BothLocationAndBuiltin) {
+TEST_F(IR_ValidatorTest, Function_Param_Struct_MultipleIOAnnotations) {
auto* f = FragmentEntryPoint("my_func");
IOAttributes attr;
attr.builtin = BuiltinValue::kPosition;
- attr.location = 0;
+ attr.color = 0;
auto* str_ty =
ty.Struct(mod.symbols.New("MyStruct"), {
{mod.symbols.New("a"), ty.vec4<f32>(), attr},
@@ -483,14 +477,15 @@
auto res = ir::Validate(mod);
ASSERT_NE(res, Success);
- EXPECT_EQ(res.Failure().reason.Str(),
- R"(:5:27 error: a builtin and location cannot be both declared for a struct member
+ EXPECT_EQ(
+ res.Failure().reason.Str(),
+ R"(:5:27 error: input param struct member has more than one IO annotation, [ built-in, @color ]
%my_func = @fragment func(%my_param:MyStruct):void {
^^^^^^^^^^^^^^^^^^
note: # Disassembly
MyStruct = struct @align(16) {
- a:vec4<f32> @offset(0), @location(0), @builtin(position)
+ a:vec4<f32> @offset(0), @color(0), @builtin(position)
}
%my_func = @fragment func(%my_param:MyStruct):void {
@@ -501,6 +496,158 @@
)");
}
+TEST_F(IR_ValidatorTest, Function_Param_MissingIOAnnotations) {
+ auto* f = FragmentEntryPoint("my_func");
+
+ auto* p = b.FunctionParam("my_param", ty.vec4<f32>());
+ f->SetParams({p});
+
+ b.Append(f->Block(), [&] { b.Return(f); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(
+ res.Failure().reason.Str(),
+ R"(:1:27 error: input param must have at least one IO annotation, e.g. a binding point, a location, etc
+%my_func = @fragment func(%my_param:vec4<f32>):void {
+ ^^^^^^^^^^^^^^^^^^^
+
+note: # Disassembly
+%my_func = @fragment func(%my_param:vec4<f32>):void {
+ $B1: {
+ ret
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Function_Param_Struct_MissingIOAnnotations) {
+ auto* f = ComputeEntryPoint("my_func");
+
+ auto* str_ty =
+ ty.Struct(mod.symbols.New("MyStruct"), {
+ {mod.symbols.New("a"), ty.vec4<f32>(), {}},
+ });
+ auto* p = b.FunctionParam("my_param", str_ty);
+ f->SetParams({p});
+
+ b.Append(f->Block(), [&] { b.Return(f); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(
+ res.Failure().reason.Str(),
+ R"(:5:54 error: input param struct members must have at least one IO annotation, e.g. a binding point, a location, etc
+%my_func = @compute @workgroup_size(1u, 1u, 1u) func(%my_param:MyStruct):void {
+ ^^^^^^^^^^^^^^^^^^
+
+note: # Disassembly
+MyStruct = struct @align(16) {
+ a:vec4<f32> @offset(0)
+}
+
+%my_func = @compute @workgroup_size(1u, 1u, 1u) func(%my_param:MyStruct):void {
+ $B1: {
+ ret
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Function_Param_Struct_DuplicateAnnotations) {
+ auto* f = ComputeEntryPoint("my_func");
+ IOAttributes attr;
+ attr.location = 0;
+ auto* str_ty =
+ ty.Struct(mod.symbols.New("MyStruct"), {
+ {mod.symbols.New("a"), ty.vec4<f32>(), attr},
+ });
+ auto* p = b.FunctionParam("my_param", str_ty);
+ p->SetLocation(0);
+ f->SetParams({p});
+
+ b.Append(f->Block(), [&] { b.Return(f); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(
+ res.Failure().reason.Str(),
+ R"(:5:54 error: input param struct member has same IO annotation, as top-level struct, '@location'
+%my_func = @compute @workgroup_size(1u, 1u, 1u) func(%my_param:MyStruct [@location(0)]):void {
+ ^^^^^^^^^^^^^^^^^^
+
+note: # Disassembly
+MyStruct = struct @align(16) {
+ a:vec4<f32> @offset(0), @location(0)
+}
+
+%my_func = @compute @workgroup_size(1u, 1u, 1u) func(%my_param:MyStruct [@location(0)]):void {
+ $B1: {
+ ret
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Function_Param_WorkgroupPlusOtherIOAnnotation) {
+ auto* f = ComputeEntryPoint("my_func");
+ auto* p = b.FunctionParam("my_param", ty.ptr<workgroup, i32>());
+ p->SetLocation(0);
+ f->SetParams({p});
+
+ b.Append(f->Block(), [&] { b.Return(f); });
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(
+ res.Failure().reason.Str(),
+ R"(:1:54 error: input param has more than one IO annotation, [ @location, <workgroup> ]
+%my_func = @compute @workgroup_size(1u, 1u, 1u) func(%my_param:ptr<workgroup, i32, read_write> [@location(0)]):void {
+ ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+note: # Disassembly
+%my_func = @compute @workgroup_size(1u, 1u, 1u) func(%my_param:ptr<workgroup, i32, read_write> [@location(0)]):void {
+ $B1: {
+ ret
+ }
+}
+)");
+}
+
+TEST_F(IR_ValidatorTest, Function_Param_Struct_WorkgroupPlusOtherIOAnnotations) {
+ auto* f = ComputeEntryPoint("my_func");
+ IOAttributes attr;
+ attr.location = 0;
+ auto* str_ty = ty.Struct(mod.symbols.New("MyStruct"),
+ {
+ {mod.symbols.New("a"), ty.ptr<workgroup, i32>(), attr},
+ });
+ auto* p = b.FunctionParam("my_param", str_ty);
+ f->SetParams({p});
+
+ b.Append(f->Block(), [&] { b.Return(f); });
+
+ auto res = ir::Validate(mod, Capabilities{Capability::kAllowPointersInStructures});
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(
+ res.Failure().reason.Str(),
+ R"(:5:54 error: input param struct member has more than one IO annotation, [ @location, <workgroup> ]
+%my_func = @compute @workgroup_size(1u, 1u, 1u) func(%my_param:MyStruct):void {
+ ^^^^^^^^^^^^^^^^^^
+
+note: # Disassembly
+MyStruct = struct @align(1) {
+ a:ptr<workgroup, i32, read_write> @offset(0), @location(0)
+}
+
+%my_func = @compute @workgroup_size(1u, 1u, 1u) func(%my_param:MyStruct):void {
+ $B1: {
+ ret
+ }
+}
+)");
+}
+
TEST_F(IR_ValidatorTest, Function_ParameterWithConstructibleType) {
auto* f = b.Function("my_func", ty.void_());
auto* p = b.FunctionParam("my_param", ty.u32());
@@ -568,10 +715,8 @@
auto* f = b.Function("my_func", ty.void_(), Function::PipelineStage::kFragment);
auto* p = b.FunctionParam("my_param", ty.vec4<f32>());
- IOAttributes attr;
- attr.builtin = BuiltinValue::kPosition;
- attr.invariant = true;
- p->SetAttributes(attr);
+ p->SetInvariant(true);
+ p->SetBuiltin(BuiltinValue::kPosition);
f->SetParams({p});
b.Append(f->Block(), [&] { b.Return(f); });
@@ -583,9 +728,7 @@
TEST_F(IR_ValidatorTest, Function_Param_InvariantWithoutPosition) {
auto* f = b.Function("my_func", ty.void_());
auto* p = b.FunctionParam("my_param", ty.vec4<f32>());
- IOAttributes attr;
- attr.invariant = true;
- p->SetAttributes(attr);
+ p->SetInvariant(true);
f->SetParams({p});
b.Append(f->Block(), [&] { b.Return(f); });
@@ -686,19 +829,16 @@
)");
}
-TEST_F(IR_ValidatorTest, Function_Return_BothLocationAndBuiltin) {
+TEST_F(IR_ValidatorTest, Function_Return_MultipleIOAnnotations) {
auto* f = VertexEntryPoint("my_func");
- IOAttributes attr;
- attr.builtin = BuiltinValue::kPosition;
- attr.location = 0;
- f->SetReturnAttributes(attr);
+ f->SetReturnLocation(0);
b.Append(f->Block(), [&] { b.Unreachable(); });
auto res = ir::Validate(mod);
ASSERT_NE(res, Success);
EXPECT_EQ(res.Failure().reason.Str(),
- R"(:1:1 error: a builtin and location cannot be both declared for a function return
+ R"(:1:1 error: return values has more than one IO annotation, [ @location, built-in ]
%my_func = @vertex func():vec4<f32> [@location(0), @position] {
^^^^^^^^
@@ -711,7 +851,7 @@
)");
}
-TEST_F(IR_ValidatorTest, Function_Return_Struct_BothLocationAndBuiltin) {
+TEST_F(IR_ValidatorTest, Function_Return_Struct_MultipleIOAnnotations) {
IOAttributes attr;
attr.builtin = BuiltinValue::kPosition;
attr.location = 0;
@@ -724,8 +864,9 @@
auto res = ir::Validate(mod);
ASSERT_NE(res, Success);
- EXPECT_EQ(res.Failure().reason.Str(),
- R"(:5:1 error: a builtin and location cannot be both declared for a struct member
+ EXPECT_EQ(
+ res.Failure().reason.Str(),
+ R"(:5:1 error: return values struct member has more than one IO annotation, [ @location, built-in ]
%my_func = @vertex func():MyStruct {
^^^^^^^^
@@ -742,7 +883,7 @@
)");
}
-TEST_F(IR_ValidatorTest, Function_Return_NonVoid_MissingLocationAndBuiltin) {
+TEST_F(IR_ValidatorTest, Function_Return_NonVoid_MissingIOAnnotations) {
auto* f = b.Function("my_func", ty.f32(), Function::PipelineStage::kFragment);
b.Append(f->Block(), [&] { b.Unreachable(); });
@@ -751,7 +892,7 @@
ASSERT_NE(res, Success);
EXPECT_EQ(
res.Failure().reason.Str(),
- R"(:1:1 error: a non-void return for an entry point must have a builtin or location decoration
+ R"(:1:1 error: return values must have at least one IO annotation, e.g. a binding point, a location, etc
%my_func = @fragment func():f32 {
^^^^^^^^
@@ -764,7 +905,7 @@
)");
}
-TEST_F(IR_ValidatorTest, Function_Return_NonVoid_Struct_MissingLocationAndBuiltin) {
+TEST_F(IR_ValidatorTest, Function_Return_NonVoid_Struct_MissingIOAnnotations) {
auto* str_ty = ty.Struct(mod.symbols.New("MyStruct"), {
{mod.symbols.New("a"), ty.f32(), {}},
});
@@ -776,7 +917,7 @@
ASSERT_NE(res, Success);
EXPECT_EQ(
res.Failure().reason.Str(),
- R"(:5:1 error: members of struct used for returns of entry points must have a builtin or location decoration
+ R"(:5:1 error: return values struct members must have at least one IO annotation, e.g. a binding point, a location, etc
%my_func = @fragment func():MyStruct {
^^^^^^^^
@@ -794,12 +935,9 @@
}
TEST_F(IR_ValidatorTest, Function_Return_InvariantWithPosition) {
- IOAttributes attr;
- attr.builtin = BuiltinValue::kPosition;
- attr.invariant = true;
-
auto* f = b.Function("my_func", ty.vec4<f32>(), Function::PipelineStage::kVertex);
- f->SetReturnAttributes(attr);
+ f->SetReturnBuiltin(BuiltinValue::kPosition);
+ f->SetReturnInvariant(true);
b.Append(f->Block(), [&] { b.Unreachable(); });
@@ -808,11 +946,8 @@
}
TEST_F(IR_ValidatorTest, Function_Return_InvariantWithoutPosition) {
- IOAttributes attr;
- attr.invariant = true;
-
auto* f = b.Function("my_func", ty.vec4<f32>());
- f->SetReturnAttributes(attr);
+ f->SetReturnInvariant(true);
b.Append(f->Block(), [&] { b.Unreachable(); });
@@ -973,6 +1108,7 @@
TEST_F(IR_ValidatorTest, Function_Compute_NonVoidReturn) {
auto* f = b.Function("my_func", ty.f32(), core::ir::Function::PipelineStage::kCompute);
f->SetWorkgroupSize(b.Constant(1_u), b.Constant(1_u), b.Constant(1_u));
+ f->SetReturnLocation(0);
b.Append(f->Block(), [&] { b.Unreachable(); });
@@ -980,11 +1116,11 @@
ASSERT_NE(res, Success);
EXPECT_EQ(res.Failure().reason.Str(),
R"(:1:1 error: compute entry point must not have a return type
-%my_func = @compute @workgroup_size(1u, 1u, 1u) func():f32 {
+%my_func = @compute @workgroup_size(1u, 1u, 1u) func():f32 [@location(0)] {
^^^^^^^^
note: # Disassembly
-%my_func = @compute @workgroup_size(1u, 1u, 1u) func():f32 {
+%my_func = @compute @workgroup_size(1u, 1u, 1u) func():f32 [@location(0)] {
$B1: {
unreachable
}
@@ -1229,17 +1365,19 @@
TEST_F(IR_ValidatorTest, Function_Vertex_MissingPosition) {
auto* f = b.Function("my_func", ty.vec4<f32>(), Function::PipelineStage::kVertex);
+ f->SetReturnLocation(0);
+
b.Append(f->Block(), [&] { b.Unreachable(); });
auto res = ir::Validate(mod);
ASSERT_NE(res, Success);
EXPECT_EQ(res.Failure().reason.Str(),
R"(:1:1 error: position must be declared for vertex entry point output
-%my_func = @vertex func():vec4<f32> {
+%my_func = @vertex func():vec4<f32> [@location(0)] {
^^^^^^^^
note: # Disassembly
-%my_func = @vertex func():vec4<f32> {
+%my_func = @vertex func():vec4<f32> [@location(0)] {
$B1: {
unreachable
}
@@ -1249,18 +1387,20 @@
TEST_F(IR_ValidatorTest, Function_NonFragment_BoolInput) {
auto* f = VertexEntryPoint();
- f->AppendParam(b.FunctionParam("invalid", ty.bool_()));
+ auto* p = b.FunctionParam("invalid", ty.bool_());
+ p->SetLocation(0);
+ f->AppendParam(p);
b.Append(f->Block(), [&] { b.Unreachable(); });
auto res = ir::Validate(mod);
ASSERT_NE(res, Success);
EXPECT_EQ(res.Failure().reason.Str(),
R"(:1:19 error: entry point params can only be a bool for fragment shaders
-%f = @vertex func(%invalid:bool):vec4<f32> [@position] {
+%f = @vertex func(%invalid:bool [@location(0)]):vec4<f32> [@position] {
^^^^^^^^^^^^^
note: # Disassembly
-%f = @vertex func(%invalid:bool):vec4<f32> [@position] {
+%f = @vertex func(%invalid:bool [@location(0)]):vec4<f32> [@position] {
$B1: {
unreachable
}
@@ -1270,7 +1410,9 @@
TEST_F(IR_ValidatorTest, Function_NonFragment_BoolOutput) {
auto* f = VertexEntryPoint();
- AddReturn(f, "invalid", ty.bool_());
+ IOAttributes attr;
+ attr.location = 0;
+ AddReturn(f, "invalid", ty.bool_(), attr);
b.Append(f->Block(), [&] { b.Unreachable(); });
auto res = ir::Validate(mod);
@@ -1283,7 +1425,7 @@
note: # Disassembly
OutputStruct = struct @align(16) {
pos:vec4<f32> @offset(0), @builtin(position)
- invalid:bool @offset(16)
+ invalid:bool @offset(16), @location(0)
}
%f = @vertex func():OutputStruct {
@@ -1296,7 +1438,9 @@
TEST_F(IR_ValidatorTest, Function_Fragment_BoolInputWithoutFrontFacing) {
auto* f = FragmentEntryPoint();
- f->AppendParam(b.FunctionParam("invalid", ty.bool_()));
+ auto* p = b.FunctionParam("invalid", ty.bool_());
+ p->SetLocation(0);
+ f->AppendParam(p);
b.Append(f->Block(), [&] { b.Unreachable(); });
auto res = ir::Validate(mod);
@@ -1304,11 +1448,11 @@
EXPECT_EQ(
res.Failure().reason.Str(),
R"(:1:21 error: fragment entry point params can only be a bool if decorated with @builtin(front_facing)
-%f = @fragment func(%invalid:bool):void {
+%f = @fragment func(%invalid:bool [@location(0)]):void {
^^^^^^^^^^^^^
note: # Disassembly
-%f = @fragment func(%invalid:bool):void {
+%f = @fragment func(%invalid:bool [@location(0)]):void {
$B1: {
unreachable
}
@@ -1343,6 +1487,9 @@
auto* f = ComputeEntryPoint();
auto* v = b.Var(ty.ptr(AddressSpace::kOut, ty.bool_(), core::Access::kReadWrite));
+ IOAttributes attr;
+ attr.location = 0;
+ v->SetAttributes(attr);
mod.root_block->Append(v);
b.Append(f->Block(), [&] {
@@ -1361,7 +1508,7 @@
note: # Disassembly
$B1: { # root
- %1:ptr<__out, bool, read_write> = var
+ %1:ptr<__out, bool, read_write> = var @location(0)
}
%f = @compute @workgroup_size(1u, 1u, 1u) func():void {
@@ -1377,6 +1524,9 @@
auto* f = FragmentEntryPoint();
auto* invalid = b.Var("invalid", AddressSpace::kIn, ty.bool_());
+ IOAttributes attr;
+ attr.location = 0;
+ invalid->SetAttributes(attr);
mod.root_block->Append(invalid);
b.Append(f->Block(), [&] {
@@ -1396,7 +1546,7 @@
note: # Disassembly
$B1: { # root
- %invalid:ptr<__in, bool, read> = var
+ %invalid:ptr<__in, bool, read> = var @location(0)
}
%f = @fragment func():void {
@@ -5200,7 +5350,7 @@
)");
}
-TEST_F(IR_ValidatorTest, Var_IOBothLocationAndBuiltin) {
+TEST_F(IR_ValidatorTest, Var_MultipleIOAnnotations) {
auto* v = b.Var<AddressSpace::kIn, vec4<f32>>();
IOAttributes attr;
attr.builtin = BuiltinValue::kPosition;
@@ -5212,7 +5362,7 @@
ASSERT_NE(res, Success);
EXPECT_EQ(
res.Failure().reason.Str(),
- R"(:2:35 error: var: a builtin and location cannot be both declared for a module scope var
+ R"(:2:35 error: var: module scope variable has more than one IO annotation, [ @location, built-in ]
%1:ptr<__in, vec4<f32>, read> = var @location(0) @builtin(position)
^^^
@@ -5228,10 +5378,10 @@
)");
}
-TEST_F(IR_ValidatorTest, Var_Struct_IOBothLocationAndBuiltin) {
+TEST_F(IR_ValidatorTest, Var_Struct_MultipleIOAnnotations) {
IOAttributes attr;
attr.builtin = BuiltinValue::kPosition;
- attr.location = 0;
+ attr.color = 0;
auto* str_ty =
ty.Struct(mod.symbols.New("MyStruct"), {
@@ -5244,7 +5394,7 @@
ASSERT_NE(res, Success);
EXPECT_EQ(
res.Failure().reason.Str(),
- R"(:6:41 error: var: a builtin and location cannot be both declared for a module scope var struct member
+ R"(:6:41 error: var: module scope variable struct member has more than one IO annotation, [ built-in, @color ]
%1:ptr<__out, MyStruct, read_write> = var
^^^
@@ -5254,7 +5404,62 @@
note: # Disassembly
MyStruct = struct @align(4) {
- a:f32 @offset(0), @location(0), @builtin(position)
+ a:f32 @offset(0), @color(0), @builtin(position)
+}
+
+$B1: { # root
+ %1:ptr<__out, MyStruct, read_write> = var
+}
+
+)");
+}
+
+TEST_F(IR_ValidatorTest, Var_MissingIOAnnotations) {
+ auto* v = b.Var<AddressSpace::kIn, vec4<f32>>();
+ mod.root_block->Append(v);
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(
+ res.Failure().reason.Str(),
+ R"(:2:35 error: var: module scope variable must have at least one IO annotation, e.g. a binding point, a location, etc
+ %1:ptr<__in, vec4<f32>, read> = var
+ ^^^
+
+:1:1 note: in block
+$B1: { # root
+^^^
+
+note: # Disassembly
+$B1: { # root
+ %1:ptr<__in, vec4<f32>, read> = var
+}
+
+)");
+}
+
+TEST_F(IR_ValidatorTest, Var_Struct_MissingIOAnnotations) {
+ auto* str_ty = ty.Struct(mod.symbols.New("MyStruct"), {
+ {mod.symbols.New("a"), ty.f32(), {}},
+ });
+ auto* v = b.Var(ty.ptr(AddressSpace::kOut, str_ty, read_write));
+ mod.root_block->Append(v);
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_EQ(
+ res.Failure().reason.Str(),
+ R"(:6:41 error: var: module scope variable struct members must have at least one IO annotation, e.g. a binding point, a location, etc
+ %1:ptr<__out, MyStruct, read_write> = var
+ ^^^
+
+:5:1 note: in block
+$B1: { # root
+^^^
+
+note: # Disassembly
+MyStruct = struct @align(4) {
+ a:f32 @offset(0)
}
$B1: { # root
diff --git a/src/tint/lang/hlsl/writer/raise/decompose_storage_access_test.cc b/src/tint/lang/hlsl/writer/raise/decompose_storage_access_test.cc
index 1d2db45..b4f1e84 100644
--- a/src/tint/lang/hlsl/writer/raise/decompose_storage_access_test.cc
+++ b/src/tint/lang/hlsl/writer/raise/decompose_storage_access_test.cc
@@ -1386,7 +1386,10 @@
b.ir.root_block->Append(var);
auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ core::IOAttributes index_attr;
+ index_attr.location = 0;
auto index = b.FunctionParam(ty.u32());
+ index->SetAttributes(index_attr);
func->SetParams({index});
b.Append(func->Block(), [&] {
auto* access = b.Access(ty.ptr<storage>(ty.atomic<i32>()), var, 0_u, index, 1_u, index);
@@ -1408,7 +1411,7 @@
%v:ptr<storage, S2, read_write> = var @binding_point(0, 0)
}
-%foo = @fragment func(%3:u32):void {
+%foo = @fragment func(%3:u32 [@location(0)]):void {
$B2: {
%4:ptr<storage, atomic<i32>, read_write> = access %v, 0u, %3, 1u, %3
%5:void = atomicStore %4, 123i
@@ -1432,7 +1435,7 @@
%v:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
}
-%foo = @fragment func(%3:u32):void {
+%foo = @fragment func(%3:u32 [@location(0)]):void {
$B2: {
%4:u32 = mul %3, 32u
%5:u32 = mul %3, 4u
@@ -1576,7 +1579,10 @@
b.ir.root_block->Append(var);
auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ core::IOAttributes index_attr;
+ index_attr.location = 0;
auto index = b.FunctionParam(ty.u32());
+ index->SetAttributes(index_attr);
func->SetParams({index});
b.Append(func->Block(), [&] {
auto* access = b.Access(ty.ptr<storage>(ty.atomic<i32>()), var, 0_u, index, 1_u, index);
@@ -1598,7 +1604,7 @@
%v:ptr<storage, S2, read_write> = var @binding_point(0, 0)
}
-%foo = @fragment func(%3:u32):void {
+%foo = @fragment func(%3:u32 [@location(0)]):void {
$B2: {
%4:ptr<storage, atomic<i32>, read_write> = access %v, 0u, %3, 1u, %3
%5:i32 = atomicLoad %4
@@ -1623,7 +1629,7 @@
%v:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
}
-%foo = @fragment func(%3:u32):void {
+%foo = @fragment func(%3:u32 [@location(0)]):void {
$B2: {
%4:u32 = mul %3, 32u
%5:u32 = mul %3, 4u
@@ -1773,7 +1779,10 @@
b.ir.root_block->Append(var);
auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ core::IOAttributes index_attr;
+ index_attr.location = 0;
auto index = b.FunctionParam(ty.u32());
+ index->SetAttributes(index_attr);
func->SetParams({index});
b.Append(func->Block(), [&] {
auto* access = b.Access(ty.ptr<storage>(ty.atomic<i32>()), var, 0_u, index, 1_u, index);
@@ -1795,7 +1804,7 @@
%v:ptr<storage, S2, read_write> = var @binding_point(0, 0)
}
-%foo = @fragment func(%3:u32):void {
+%foo = @fragment func(%3:u32 [@location(0)]):void {
$B2: {
%4:ptr<storage, atomic<i32>, read_write> = access %v, 0u, %3, 1u, %3
%5:i32 = atomicSub %4, 123i
@@ -1820,7 +1829,7 @@
%v:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
}
-%foo = @fragment func(%3:u32):void {
+%foo = @fragment func(%3:u32 [@location(0)]):void {
$B2: {
%4:u32 = mul %3, 32u
%5:u32 = mul %3, 4u
@@ -1985,7 +1994,10 @@
b.ir.root_block->Append(var);
auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ core::IOAttributes index_attr;
+ index_attr.location = 0;
auto index = b.FunctionParam(ty.u32());
+ index->SetAttributes(index_attr);
func->SetParams({index});
b.Append(func->Block(), [&] {
auto* access = b.Access(ty.ptr<storage>(ty.atomic<i32>()), var, 0_u, index, 1_u, index);
@@ -2013,7 +2025,7 @@
%v:ptr<storage, S2, read_write> = var @binding_point(0, 0)
}
-%foo = @fragment func(%3:u32):void {
+%foo = @fragment func(%3:u32 [@location(0)]):void {
$B2: {
%4:ptr<storage, atomic<i32>, read_write> = access %v, 0u, %3, 1u, %3
%5:__atomic_compare_exchange_result_i32 = atomicCompareExchangeWeak %4, 123i, 345i
@@ -2043,7 +2055,7 @@
%v:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
}
-%foo = @fragment func(%3:u32):void {
+%foo = @fragment func(%3:u32 [@location(0)]):void {
$B2: {
%4:u32 = mul %3, 32u
%5:u32 = mul %3, 4u
@@ -2274,6 +2286,7 @@
auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
auto index = b.FunctionParam(ty.u32());
+ index->SetLocation(0);
func->SetParams({index});
b.Append(func->Block(), [&] {
auto* access = b.Access(ty.ptr<storage>(ty.atomic<u32>()), var, 0_u, index, 0_u, index);
@@ -2294,7 +2307,7 @@
%v:ptr<storage, S2, read_write> = var @binding_point(0, 0)
}
-%foo = @fragment func(%3:u32):void {
+%foo = @fragment func(%3:u32 [@location(0)]):void {
$B2: {
%4:ptr<storage, atomic<u32>, read_write> = access %v, 0u, %3, 0u, %3
%5:u32 = )" +
@@ -2319,7 +2332,7 @@
%v:hlsl.byte_address_buffer<read_write> = var @binding_point(0, 0)
}
-%foo = @fragment func(%3:u32):void {
+%foo = @fragment func(%3:u32 [@location(0)]):void {
$B2: {
%4:u32 = mul %3, 12u
%5:u32 = mul %3, 4u
diff --git a/src/tint/lang/msl/writer/raise/simd_ballot_test.cc b/src/tint/lang/msl/writer/raise/simd_ballot_test.cc
index 3e8e1f7..c026411 100644
--- a/src/tint/lang/msl/writer/raise/simd_ballot_test.cc
+++ b/src/tint/lang/msl/writer/raise/simd_ballot_test.cc
@@ -44,6 +44,9 @@
TEST_F(MslWriter_SimdBallotTest, SimdBallot_WithUserDeclaredSubgroupSize) {
auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
auto* subgroup_size = b.FunctionParam("user_subgroup_size", ty.u32());
+ core::IOAttributes attr;
+ attr.location = 0;
+ subgroup_size->SetAttributes(attr);
func->SetParams({subgroup_size});
b.Append(func->Block(), [&] { //
b.Call<vec4<u32>>(core::BuiltinFn::kSubgroupBallot, true);
@@ -51,7 +54,7 @@
});
auto* src = R"(
-%foo = @fragment func(%user_subgroup_size:u32):void {
+%foo = @fragment func(%user_subgroup_size:u32 [@location(0)]):void {
$B1: {
%3:vec4<u32> = subgroupBallot true
ret
@@ -65,7 +68,7 @@
%tint_subgroup_size_mask:ptr<private, vec2<u32>, read_write> = var
}
-%foo = @fragment func(%user_subgroup_size:u32, %tint_subgroup_size:u32 [@subgroup_size]):void {
+%foo = @fragment func(%user_subgroup_size:u32 [@location(0)], %tint_subgroup_size:u32 [@subgroup_size]):void {
$B2: {
%5:bool = gt %tint_subgroup_size, 32u
%6:u32 = sub 32u, %tint_subgroup_size
@@ -160,6 +163,9 @@
auto* ep1 = b.Function("ep1", ty.void_(), core::ir::Function::PipelineStage::kFragment);
auto* subgroup_size = b.FunctionParam("user_subgroup_size", ty.u32());
+ core::IOAttributes attr;
+ attr.location = 0;
+ subgroup_size->SetAttributes(attr);
ep1->SetParams({subgroup_size});
b.Append(ep1->Block(), [&] { //
b.Call<vec4<u32>>(foo, true);
@@ -179,7 +185,7 @@
ret %3
}
}
-%ep1 = @fragment func(%user_subgroup_size:u32):void {
+%ep1 = @fragment func(%user_subgroup_size:u32 [@location(0)]):void {
$B2: {
%6:vec4<u32> = call %foo, true
ret
@@ -205,7 +211,7 @@
ret %4
}
}
-%ep1 = @fragment func(%user_subgroup_size:u32, %tint_subgroup_size:u32 [@subgroup_size]):void {
+%ep1 = @fragment func(%user_subgroup_size:u32 [@location(0)], %tint_subgroup_size:u32 [@subgroup_size]):void {
$B3: {
%9:bool = gt %tint_subgroup_size, 32u
%10:u32 = sub 32u, %tint_subgroup_size
diff --git a/src/tint/lang/spirv/writer/raise/merge_return_test.cc b/src/tint/lang/spirv/writer/raise/merge_return_test.cc
index 860f1b9..3c6446c 100644
--- a/src/tint/lang/spirv/writer/raise/merge_return_test.cc
+++ b/src/tint/lang/spirv/writer/raise/merge_return_test.cc
@@ -213,6 +213,9 @@
TEST_F(SpirvWriter_MergeReturnTest, NoModify_EntryPoint_IfElse_OneSideReturns) {
auto* cond = b.FunctionParam(ty.u32());
+ core::IOAttributes attr;
+ attr.location = 0;
+ cond->SetAttributes(attr);
auto* func = b.ComputeFunction("entrypointfunction", 2_u, 3_u, 4_u);
func->SetParams({cond});
b.Append(func->Block(), [&] {
@@ -224,7 +227,7 @@
});
auto* src = R"(
-%entrypointfunction = @compute @workgroup_size(2u, 3u, 4u) func(%2:u32):void {
+%entrypointfunction = @compute @workgroup_size(2u, 3u, 4u) func(%2:u32 [@location(0)]):void {
$B1: {
%3:bool = eq %2, 0u
if %3 [t: $B2, f: $B3] { # if_1
diff --git a/src/tint/lang/spirv/writer/raise/shader_io_test.cc b/src/tint/lang/spirv/writer/raise/shader_io_test.cc
index 4d09417..417b406 100644
--- a/src/tint/lang/spirv/writer/raise/shader_io_test.cc
+++ b/src/tint/lang/spirv/writer/raise/shader_io_test.cc
@@ -1528,7 +1528,7 @@
auto* in1 = b.FunctionParam("in1", ty.f16());
auto* in2 = b.FunctionParam("in2", ty.vec4<f16>());
in1->SetLocation(1);
- in1->SetLocation(2);
+ in2->SetLocation(2);
auto* func = b.Function("main", outputs, core::ir::Function::PipelineStage::kFragment);
func->SetParams({in1, in2});
b.Append(func->Block(), [&] { //
@@ -1541,7 +1541,7 @@
out2:vec4<f16> @offset(8), @location(2)
}
-%main = @fragment func(%in1:f16 [@location(2)], %in2:vec4<f16>):Outputs {
+%main = @fragment func(%in1:f16 [@location(1)], %in2:vec4<f16> [@location(2)]):Outputs {
$B1: {
%4:Outputs = construct %in1, %in2
ret %4
@@ -1557,8 +1557,8 @@
}
$B1: { # root
- %main_loc2_Input:ptr<__in, f16, read> = var @location(2)
- %main_Input:ptr<__in, vec4<f16>, read> = var
+ %main_loc1_Input:ptr<__in, f16, read> = var @location(1)
+ %main_loc2_Input:ptr<__in, vec4<f16>, read> = var @location(2)
%main_loc1_Output:ptr<__out, f16, write> = var @location(1)
%main_loc2_Output:ptr<__out, vec4<f16>, write> = var @location(2)
}
@@ -1571,8 +1571,8 @@
}
%main = @fragment func():void {
$B3: {
- %10:f16 = load %main_loc2_Input
- %11:vec4<f16> = load %main_Input
+ %10:f16 = load %main_loc1_Input
+ %11:vec4<f16> = load %main_loc2_Input
%12:Outputs = call %main_inner, %10, %11
%13:f16 = access %12, 0u
store %main_loc1_Output, %13
@@ -1622,7 +1622,7 @@
auto* in1 = b.FunctionParam("in1", ty.f16());
auto* in2 = b.FunctionParam("in2", ty.vec4<f16>());
in1->SetLocation(1);
- in1->SetLocation(2);
+ in2->SetLocation(2);
auto* func = b.Function("main", outputs, core::ir::Function::PipelineStage::kFragment);
func->SetParams({in1, in2});
b.Append(func->Block(), [&] { //
@@ -1635,7 +1635,7 @@
out2:vec4<f16> @offset(8), @location(2)
}
-%main = @fragment func(%in1:f16 [@location(2)], %in2:vec4<f16>):Outputs {
+%main = @fragment func(%in1:f16 [@location(1)], %in2:vec4<f16> [@location(2)]):Outputs {
$B1: {
%4:Outputs = construct %in1, %in2
ret %4
@@ -1651,8 +1651,8 @@
}
$B1: { # root
- %main_loc2_Input:ptr<__in, f32, read> = var @location(2)
- %main_Input:ptr<__in, vec4<f32>, read> = var
+ %main_loc1_Input:ptr<__in, f32, read> = var @location(1)
+ %main_loc2_Input:ptr<__in, vec4<f32>, read> = var @location(2)
%main_loc1_Output:ptr<__out, f32, write> = var @location(1)
%main_loc2_Output:ptr<__out, vec4<f32>, write> = var @location(2)
}
@@ -1665,9 +1665,9 @@
}
%main = @fragment func():void {
$B3: {
- %10:f32 = load %main_loc2_Input
+ %10:f32 = load %main_loc1_Input
%11:f16 = convert %10
- %12:vec4<f32> = load %main_Input
+ %12:vec4<f32> = load %main_loc2_Input
%13:vec4<f16> = convert %12
%14:Outputs = call %main_inner, %11, %13
%15:f16 = access %14, 0u
