[ir] Change function return attributes to not be list.
This CL updates the function return attributes to store individual
members instead of a list of attributes. This matches the
FunctionParam and Var usage for attributes.
Bug: tint:1915
Change-Id: I99e8a7a055afbe841af7c149078f547fd0608344
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/134840
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index 3e7b50e..d8e472a 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -143,7 +143,20 @@
out_ << "@binding_point(" << p.group << ", " << p.binding << ")";
}
-void Disassembler::EmitParamAttributes(FunctionParam* p) {
+void Disassembler::EmitLocation(Location loc) {
+ out_ << "@location(" << loc.value << ")";
+ if (loc.interpolation.has_value()) {
+ out_ << ", @interpolate(";
+ out_ << loc.interpolation->type;
+ if (loc.interpolation->sampling != builtin::InterpolationSampling::kUndefined) {
+ out_ << ", ";
+ out_ << loc.interpolation->sampling;
+ }
+ out_ << ")";
+ }
+}
+
+void Disassembler::EmitParamAttributes(const FunctionParam* p) {
if (!p->Invariant() && !p->Location().has_value() && !p->BindingPoint().has_value() &&
!p->Builtin().has_value()) {
return;
@@ -164,17 +177,7 @@
need_comma = true;
}
if (p->Location().has_value()) {
- out_ << "@location(" << p->Location()->value << ")";
- if (p->Location()->interpolation.has_value()) {
- out_ << ", @interpolate(";
- out_ << p->Location()->interpolation->type;
- if (p->Location()->interpolation->sampling !=
- builtin::InterpolationSampling::kUndefined) {
- out_ << ", ";
- out_ << p->Location()->interpolation->sampling;
- }
- out_ << ")";
- }
+ EmitLocation(p->Location().value());
need_comma = true;
}
if (p->BindingPoint().has_value()) {
@@ -190,11 +193,54 @@
out_ << "]";
}
+void Disassembler::EmitReturnAttributes(const Function* func) {
+ if (!func->ReturnInvariant() && !func->ReturnLocation().has_value() &&
+ !func->ReturnBuiltin().has_value()) {
+ return;
+ }
+
+ out_ << " [";
+
+ bool need_comma = false;
+ auto comma = [&]() {
+ if (need_comma) {
+ out_ << ", ";
+ }
+ };
+ if (func->ReturnInvariant()) {
+ comma();
+ out_ << "@invariant";
+ need_comma = true;
+ }
+ if (func->ReturnLocation().has_value()) {
+ comma();
+ EmitLocation(func->ReturnLocation().value());
+ need_comma = true;
+ }
+ if (func->ReturnBuiltin().has_value()) {
+ comma();
+ out_ << "@" << func->ReturnBuiltin().value();
+ need_comma = true;
+ }
+ out_ << "]";
+}
+
void Disassembler::EmitFunction(const Function* func) {
in_function_ = true;
- Indent() << "%" << IdOf(func) << " = func(";
- for (auto* p : func->Params()) {
+ Indent() << "%" << IdOf(func) << " =";
+
+ if (func->Stage() != Function::PipelineStage::kUndefined) {
+ out_ << " @" << func->Stage();
+ }
+ if (func->WorkgroupSize()) {
+ auto arr = func->WorkgroupSize().value();
+ out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2] << ")";
+ }
+
+ out_ << " func(";
+
+ for (const auto* p : func->Params()) {
if (p != func->Params().Front()) {
out_ << ", ";
}
@@ -204,27 +250,8 @@
}
out_ << "):" << func->ReturnType()->FriendlyName();
- if (func->Stage() != Function::PipelineStage::kUndefined) {
- out_ << " [@" << func->Stage();
+ EmitReturnAttributes(func);
- if (func->WorkgroupSize()) {
- auto arr = func->WorkgroupSize().value();
- out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2] << ")";
- }
-
- if (!func->ReturnAttributes().IsEmpty()) {
- out_ << " ra:";
-
- for (auto attr : func->ReturnAttributes()) {
- out_ << " @" << attr;
- if (attr == Function::ReturnAttribute::kLocation) {
- out_ << "(" << func->ReturnLocation().value() << ")";
- }
- }
- }
-
- out_ << "]";
- }
out_ << " -> %b" << IdOf(func->StartTarget()) << " {" << std::endl;
{
diff --git a/src/tint/ir/disassembler.h b/src/tint/ir/disassembler.h
index 99d783b..9194942 100644
--- a/src/tint/ir/disassembler.h
+++ b/src/tint/ir/disassembler.h
@@ -59,8 +59,10 @@
void Walk(const Block* blk);
void WalkInternal(const Block* blk);
void EmitFunction(const Function* func);
- void EmitParamAttributes(FunctionParam* p);
+ void EmitParamAttributes(const FunctionParam* p);
+ void EmitReturnAttributes(const Function* func);
void EmitBindingPoint(BindingPoint p);
+ void EmitLocation(Location loc);
void EmitInstruction(const Instruction* inst);
void EmitValueWithType(const Value* val);
void EmitValue(const Value* val);
diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc
index 6e917e1..19b6199 100644
--- a/src/tint/ir/from_program.cc
+++ b/src/tint/ir/from_program.cc
@@ -231,6 +231,24 @@
return ResultType{std::move(mod)};
}
+ builtin::Interpolation ExtractInterpolation(const ast::InterpolateAttribute* interp) {
+ auto type = program_->Sem()
+ .Get(interp->type)
+ ->As<sem::BuiltinEnumExpression<builtin::InterpolationType>>();
+ builtin::InterpolationType interpolation_type = type->Value();
+
+ builtin::InterpolationSampling interpolation_sampling =
+ builtin::InterpolationSampling::kUndefined;
+ if (interp->sampling) {
+ auto sampling = program_->Sem()
+ .Get(interp->sampling)
+ ->As<sem::BuiltinEnumExpression<builtin::InterpolationSampling>>();
+ interpolation_sampling = sampling->Value();
+ }
+
+ return builtin::Interpolation{interpolation_type, interpolation_sampling};
+ }
+
void EmitFunction(const ast::Function* ast_func) {
// The flow stack should have been emptied when the previous function finished building.
TINT_ASSERT(IR, control_stack_.IsEmpty());
@@ -266,16 +284,16 @@
}
}
- utils::Vector<Function::ReturnAttribute, 1> return_attributes;
+ // Note, interpolated is only valid when paired with Location, so it will only be set
+ // when the location is set.
+ std::optional<builtin::Interpolation> interpolation;
for (auto* attr : ast_func->return_type_attributes) {
tint::Switch(
attr, //
- [&](const ast::LocationAttribute*) {
- return_attributes.Push(Function::ReturnAttribute::kLocation);
+ [&](const ast::InterpolateAttribute* interp) {
+ interpolation = ExtractInterpolation(interp);
},
- [&](const ast::InvariantAttribute*) {
- return_attributes.Push(Function::ReturnAttribute::kInvariant);
- },
+ [&](const ast::InvariantAttribute*) { ir_func->SetReturnInvariant(true); },
[&](const ast::BuiltinAttribute* b) {
if (auto* ident_sem =
program_->Sem()
@@ -283,13 +301,13 @@
->As<sem::BuiltinEnumExpression<builtin::BuiltinValue>>()) {
switch (ident_sem->Value()) {
case builtin::BuiltinValue::kPosition:
- return_attributes.Push(Function::ReturnAttribute::kPosition);
+ ir_func->SetReturnBuiltin(Function::ReturnBuiltin::kPosition);
break;
case builtin::BuiltinValue::kFragDepth:
- return_attributes.Push(Function::ReturnAttribute::kFragDepth);
+ ir_func->SetReturnBuiltin(Function::ReturnBuiltin::kFragDepth);
break;
case builtin::BuiltinValue::kSampleMask:
- return_attributes.Push(Function::ReturnAttribute::kSampleMask);
+ ir_func->SetReturnBuiltin(Function::ReturnBuiltin::kSampleMask);
break;
default:
TINT_ICE(IR, diagnostics_)
@@ -303,9 +321,10 @@
}
});
}
- ir_func->SetReturnAttributes(return_attributes);
+ if (sem->ReturnLocation().has_value()) {
+ ir_func->SetReturnLocation(sem->ReturnLocation().value(), interpolation);
+ }
}
- ir_func->SetReturnLocation(sem->ReturnLocation());
scopes_.Push();
TINT_DEFER(scopes_.Pop());
@@ -323,24 +342,7 @@
tint::Switch(
attr, //
[&](const ast::InterpolateAttribute* interp) {
- auto type =
- program_->Sem()
- .Get(interp->type)
- ->As<sem::BuiltinEnumExpression<builtin::InterpolationType>>();
- builtin::InterpolationType interpolation_type = type->Value();
-
- builtin::InterpolationSampling interpolation_sampling =
- builtin::InterpolationSampling::kUndefined;
- if (interp->sampling) {
- auto sampling = program_->Sem()
- .Get(interp->sampling)
- ->As<sem::BuiltinEnumExpression<
- builtin::InterpolationSampling>>();
- interpolation_sampling = sampling->Value();
- }
-
- interpolation =
- builtin::Interpolation{interpolation_type, interpolation_sampling};
+ interpolation = ExtractInterpolation(interp);
},
[&](const ast::InvariantAttribute*) { param->SetInvariant(true); },
[&](const ast::BuiltinAttribute* b) {
diff --git a/src/tint/ir/from_program_binary_test.cc b/src/tint/ir/from_program_binary_test.cc
index 8cc5604..5f7bcb2 100644
--- a/src/tint/ir/from_program_binary_test.cc
+++ b/src/tint/ir/from_program_binary_test.cc
@@ -39,7 +39,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = add %3, 4u
@@ -62,7 +62,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = add %3, 1u
@@ -86,7 +86,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = add %3, 1u
@@ -110,7 +110,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = sub %3, 4u
@@ -133,7 +133,7 @@
%v1:ptr<private, i32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:i32 = load %v1
%4:i32 = sub %3, 1i
@@ -157,7 +157,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = sub %3, 1u
@@ -181,7 +181,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = mul %3, 4u
@@ -204,7 +204,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = mul %3, 1u
@@ -228,7 +228,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = div %3, 4u
@@ -251,7 +251,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = div %3, 1u
@@ -275,7 +275,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = mod %3, 4u
@@ -298,7 +298,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = mod %3, 1u
@@ -322,7 +322,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = and %3, 4u
@@ -345,7 +345,7 @@
%v1:ptr<private, bool, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:bool = load %v1
%4:bool = and %3, false
@@ -369,7 +369,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = or %3, 4u
@@ -392,7 +392,7 @@
%v1:ptr<private, bool, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:bool = load %v1
%4:bool = or %3, false
@@ -416,7 +416,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = xor %3, 4u
@@ -439,7 +439,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = xor %3, 1u
@@ -463,7 +463,7 @@
ret true
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:bool = call %my_func
if %3 [t: %b3, f: %b4, m: %b5]
@@ -515,7 +515,7 @@
ret true
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:bool = call %my_func
if %3 [t: %b3, f: %b4, m: %b5]
@@ -567,7 +567,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:bool = eq %3, 4u
@@ -590,7 +590,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:bool = neq %3, 4u
@@ -613,7 +613,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:bool = lt %3, 4u
@@ -636,7 +636,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:bool = gt %3, 4u
@@ -659,7 +659,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:bool = lte %3, 4u
@@ -682,7 +682,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:bool = gte %3, 4u
@@ -705,7 +705,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = shiftl %3, 4u
@@ -728,7 +728,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = shiftl %3, 1u
@@ -752,7 +752,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = shiftr %3, 4u
@@ -775,7 +775,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = shiftr %3, 1u
@@ -801,7 +801,7 @@
ret 0.0f
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:f32 = call %my_func
%4:bool = lt %3, 2.0f
@@ -845,7 +845,7 @@
ret true
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%tint_symbol:bool = call %my_func, false
ret
diff --git a/src/tint/ir/from_program_builtin_test.cc b/src/tint/ir/from_program_builtin_test.cc
index acdbc8a..4d21b48 100644
--- a/src/tint/ir/from_program_builtin_test.cc
+++ b/src/tint/ir/from_program_builtin_test.cc
@@ -39,7 +39,7 @@
%i:ptr<private, f32, read_write> = var, 1.0f
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:f32 = load %i
%tint_symbol:f32 = asin %3
diff --git a/src/tint/ir/from_program_call_test.cc b/src/tint/ir/from_program_call_test.cc
index 155e42c..d305855 100644
--- a/src/tint/ir/from_program_call_test.cc
+++ b/src/tint/ir/from_program_call_test.cc
@@ -40,7 +40,7 @@
ret 0.0f
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:f32 = call %my_func
%tint_symbol:f32 = bitcast %3
@@ -60,7 +60,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(%test_function = func():void [@fragment] -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%test_function = @fragment func():void -> %b1 {
%b1 = block {
discard
ret
@@ -82,7 +82,7 @@
ret
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%4:void = call %my_func, 6.0f
ret
@@ -104,7 +104,7 @@
%i:ptr<private, i32, read_write> = var, 1i
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:i32 = load %i
%tint_symbol:f32 = convert i32, %3
@@ -142,7 +142,7 @@
%i:ptr<private, f32, read_write> = var, 1.0f
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:f32 = load %i
%tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %3
diff --git a/src/tint/ir/from_program_function_test.cc b/src/tint/ir/from_program_function_test.cc
index 1be388a..57e78a3 100644
--- a/src/tint/ir/from_program_function_test.cc
+++ b/src/tint/ir/from_program_function_test.cc
@@ -34,7 +34,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():vec4<f32> [@vertex ra: @position] -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%test = @vertex func():vec4<f32> [@position] -> %b1 {
%b1 = block {
ret vec4<f32> 0.0f
}
@@ -49,7 +49,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():void [@fragment] -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%test = @fragment func():void -> %b1 {
%b1 = block {
ret
}
@@ -65,7 +65,7 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test = func():void [@compute @workgroup_size(8, 4, 2)] -> %b1 {
+ R"(%test = @compute @workgroup_size(8, 4, 2) func():void -> %b1 {
%b1 = block {
ret
}
@@ -96,7 +96,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():vec4<f32> [@vertex ra: @position] -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%test = @vertex func():vec4<f32> [@position] -> %b1 {
%b1 = block {
ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
}
@@ -113,7 +113,7 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test = func():vec4<f32> [@vertex ra: @position @invariant] -> %b1 {
+ R"(%test = @vertex func():vec4<f32> [@invariant, @position] -> %b1 {
%b1 = block {
ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
}
@@ -129,7 +129,26 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test = func():vec4<f32> [@fragment ra: @location(1)] -> %b1 {
+ R"(%test = @fragment func():vec4<f32> [@location(1)] -> %b1 {
+ %b1 = block {
+ ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
+ }
+}
+)");
+}
+
+TEST_F(IR_BuilderImplTest, EmitFunction_ReturnLocation_Interpolate) {
+ Func("test", utils::Empty, ty.vec4<f32>(), utils::Vector{Return(vec4<f32>(1_f, 2_f, 3_f, 4_f))},
+ utils::Vector{Stage(ast::PipelineStage::kFragment)},
+ utils::Vector{Location(1_i), Interpolate(builtin::InterpolationType::kLinear,
+ builtin::InterpolationSampling::kCentroid)});
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ EXPECT_EQ(
+ Disassemble(m.Get()),
+ R"(%test = @fragment func():vec4<f32> [@location(1), @interpolate(linear, centroid)] -> %b1 {
%b1 = block {
ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
}
@@ -145,7 +164,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():f32 [@fragment ra: @frag_depth] -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%test = @fragment func():f32 [@frag_depth] -> %b1 {
%b1 = block {
ret 1.0f
}
@@ -161,7 +180,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():u32 [@fragment ra: @sample_mask] -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%test = @fragment func():u32 [@sample_mask] -> %b1 {
%b1 = block {
ret 1u
}
diff --git a/src/tint/ir/from_program_store_test.cc b/src/tint/ir/from_program_store_test.cc
index 38efea0..20ea290 100644
--- a/src/tint/ir/from_program_store_test.cc
+++ b/src/tint/ir/from_program_store_test.cc
@@ -40,7 +40,7 @@
%a:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
store %a, 4u
ret
diff --git a/src/tint/ir/from_program_test.cc b/src/tint/ir/from_program_test.cc
index c8dbd39..8f5348b 100644
--- a/src/tint/ir/from_program_test.cc
+++ b/src/tint/ir/from_program_test.cc
@@ -144,7 +144,7 @@
EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
if true [t: %b2, f: %b3, m: %b4]
# True block
@@ -184,7 +184,7 @@
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
if true [t: %b2, f: %b3, m: %b4]
# True block
@@ -224,7 +224,7 @@
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
if true [t: %b2, f: %b3, m: %b4]
# True block
@@ -264,7 +264,7 @@
EXPECT_EQ(0u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
if true [t: %b2, f: %b3]
# True block
@@ -298,7 +298,7 @@
ASSERT_NE(loop_flow, nullptr);
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
if true [t: %b2, f: %b3, m: %b4]
# True block
@@ -347,7 +347,7 @@
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, m: %b3]
%b2 = block {
@@ -387,7 +387,7 @@
EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, c: %b3, m: %b4]
%b2 = block {
@@ -442,7 +442,7 @@
EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, c: %b3, m: %b4]
%b2 = block {
@@ -475,7 +475,7 @@
auto m = res.Move();
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, c: %b3, m: %b4]
%b2 = block {
@@ -519,7 +519,7 @@
EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, c: %b3]
%b2 = block {
@@ -568,7 +568,7 @@
EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2]
%b2 = block {
@@ -606,7 +606,7 @@
EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2]
%b2 = block {
@@ -640,7 +640,7 @@
EXPECT_EQ(0u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, m: %b3]
%b2 = block {
@@ -685,7 +685,7 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, c: %b3, m: %b4]
%b2 = block {
@@ -813,7 +813,7 @@
EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, c: %b3, m: %b4]
%b2 = block {
@@ -874,7 +874,7 @@
EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, c: %b3, m: %b4]
%b2 = block {
@@ -967,7 +967,7 @@
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, m: %b3]
%b2 = block {
@@ -1021,7 +1021,7 @@
EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
switch 1i [c: (0i, %b2), c: (1i, %b3), c: (default, %b4), m: %b5]
# Case block
@@ -1082,7 +1082,7 @@
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
switch 1i [c: (0i 1i default, %b2), m: %b3]
# Case block
@@ -1121,7 +1121,7 @@
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
switch 1i [c: (default, %b2), m: %b3]
# Case block
@@ -1169,7 +1169,7 @@
// This is 1 because the if is dead-code eliminated and the return doesn't happen.
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
switch 1i [c: (0i, %b2), c: (default, %b3), m: %b4]
# Case block
@@ -1224,7 +1224,7 @@
EXPECT_EQ(0u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
switch 1i [c: (0i, %b2), c: (default, %b3)]
# Case block
@@ -1255,7 +1255,7 @@
ret 1i
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:i32 = call %b
ret
@@ -1276,7 +1276,7 @@
EXPECT_EQ(
Disassemble(m.Get()),
- R"(%f = func(%a:vec4<f32> [@invariant, @position]):vec4<f32> [@fragment ra: @location(1)] -> %b1 {
+ R"(%f = @fragment func(%a:vec4<f32> [@invariant, @position]):vec4<f32> [@location(1)] -> %b1 {
%b1 = block {
ret %a
}
@@ -1293,7 +1293,7 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%f = func(%a:f32 [@location(2)]):f32 [@fragment ra: @location(1)] -> %b1 {
+ R"(%f = @fragment func(%a:f32 [@location(2)]):f32 [@location(1)] -> %b1 {
%b1 = block {
ret %a
}
@@ -1315,7 +1315,7 @@
EXPECT_EQ(
Disassemble(m.Get()),
- R"(%f = func(%a:f32 [@location(2), @interpolate(linear, centroid)]):f32 [@fragment ra: @location(1)] -> %b1 {
+ R"(%f = @fragment func(%a:f32 [@location(2), @interpolate(linear, centroid)]):f32 [@location(1)] -> %b1 {
%b1 = block {
ret %a
}
@@ -1336,7 +1336,7 @@
EXPECT_EQ(
Disassemble(m.Get()),
- R"(%f = func(%a:f32 [@location(2), @interpolate(flat)]):f32 [@fragment ra: @location(1)] -> %b1 {
+ R"(%f = @fragment func(%a:f32 [@location(2), @interpolate(flat)]):f32 [@location(1)] -> %b1 {
%b1 = block {
ret %a
}
diff --git a/src/tint/ir/from_program_unary_test.cc b/src/tint/ir/from_program_unary_test.cc
index bb58c02..b83e5d5 100644
--- a/src/tint/ir/from_program_unary_test.cc
+++ b/src/tint/ir/from_program_unary_test.cc
@@ -39,7 +39,7 @@
ret false
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:bool = call %my_func
%tint_symbol:bool = eq %3, false
@@ -62,7 +62,7 @@
ret 1u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = complement %3
@@ -85,7 +85,7 @@
ret 1i
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:i32 = call %my_func
%tint_symbol:i32 = negation %3
@@ -109,7 +109,7 @@
%v2:ptr<private, i32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
ret
}
@@ -133,7 +133,7 @@
%v3:ptr<private, i32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
store %v3, 42i
ret
diff --git a/src/tint/ir/from_program_var_test.cc b/src/tint/ir/from_program_var_test.cc
index 6d1d7a9..91b8b8f 100644
--- a/src/tint/ir/from_program_var_test.cc
+++ b/src/tint/ir/from_program_var_test.cc
@@ -78,7 +78,7 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:ptr<function, u32, read_write> = var
ret
@@ -96,7 +96,7 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:ptr<function, u32, read_write> = var, 2u
ret
@@ -114,7 +114,7 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:ptr<function, u32, read_write> = var
%3:u32 = load %a
diff --git a/src/tint/ir/function.cc b/src/tint/ir/function.cc
index ef09065..3769d4f 100644
--- a/src/tint/ir/function.cc
+++ b/src/tint/ir/function.cc
@@ -21,7 +21,9 @@
Function::Function(const type::Type* rt,
PipelineStage stage,
std::optional<std::array<uint32_t, 3>> wg_size)
- : Base(), return_type_(rt), pipeline_stage_(stage), workgroup_size_(wg_size) {}
+ : Base(), pipeline_stage_(stage), workgroup_size_(wg_size) {
+ return_.type = rt;
+}
Function::~Function() = default;
@@ -39,18 +41,14 @@
return out << "<unknown>";
}
-utils::StringStream& operator<<(utils::StringStream& out, Function::ReturnAttribute value) {
+utils::StringStream& operator<<(utils::StringStream& out, enum Function::ReturnBuiltin value) {
switch (value) {
- case Function::ReturnAttribute::kLocation:
- return out << "location";
- case Function::ReturnAttribute::kFragDepth:
+ case Function::ReturnBuiltin::kFragDepth:
return out << "frag_depth";
- case Function::ReturnAttribute::kSampleMask:
+ case Function::ReturnBuiltin::kSampleMask:
return out << "sample_mask";
- case Function::ReturnAttribute::kPosition:
+ case Function::ReturnBuiltin::kPosition:
return out << "position";
- case Function::ReturnAttribute::kInvariant:
- return out << "invariant";
}
return out << "<unknown>";
}
diff --git a/src/tint/ir/function.h b/src/tint/ir/function.h
index dbee570..fa5c547 100644
--- a/src/tint/ir/function.h
+++ b/src/tint/ir/function.h
@@ -20,6 +20,7 @@
#include <utility>
#include "src/tint/ir/function_param.h"
+#include "src/tint/ir/location.h"
#include "src/tint/ir/value.h"
#include "src/tint/type/type.h"
@@ -46,18 +47,14 @@
kVertex,
};
- /// Attributes attached to return types
- enum class ReturnAttribute {
- /// Location attribute
- kLocation,
+ /// Builtin attached to return types
+ enum class ReturnBuiltin {
/// Builtin Position attribute
kPosition,
/// Builtin FragDepth attribute
kFragDepth,
/// Builtin SampleMask
kSampleMask,
- /// Invariant attribute
- kInvariant,
};
/// Constructor
@@ -86,26 +83,35 @@
std::optional<std::array<uint32_t, 3>> WorkgroupSize() const { return workgroup_size_; }
/// @returns the return type for the function
- const type::Type* ReturnType() const { return return_type_; }
+ const type::Type* ReturnType() const { return return_.type; }
/// Sets the return attributes
- /// @param attrs the attributes to set
- void SetReturnAttributes(utils::VectorRef<ReturnAttribute> attrs) {
- return_attributes_ = std::move(attrs);
+ /// @param builtin the builtin to set
+ void SetReturnBuiltin(ReturnBuiltin builtin) {
+ TINT_ASSERT(IR, !return_.builtin.has_value());
+ return_.builtin = builtin;
}
- /// @returns the return attributes
- utils::VectorRef<ReturnAttribute> ReturnAttributes() const { return return_attributes_; }
+ /// @returns the return builtin attribute
+ std::optional<enum ReturnBuiltin> ReturnBuiltin() const { return return_.builtin; }
/// Sets the return location
/// @param loc the location to set
- void SetReturnLocation(std::optional<uint32_t> loc) { return_location_ = loc; }
+ /// @param interp the interpolation
+ void SetReturnLocation(uint32_t loc, std::optional<builtin::Interpolation> interp) {
+ return_.location = {loc, interp};
+ }
/// @returns the return location
- std::optional<uint32_t> ReturnLocation() const { return return_location_; }
+ std::optional<Location> ReturnLocation() const { return return_.location; }
+
+ /// Sets the return as invariant
+ /// @param val the invariant value to set
+ void SetReturnInvariant(bool val) { return_.invariant = val; }
+ /// @returns the return invariant value
+ bool ReturnInvariant() const { return return_.invariant; }
/// Sets the function parameters
/// @param params the function paramters
void SetParams(utils::VectorRef<FunctionParam*> params) { params_ = std::move(params); }
-
/// @returns the function parameters
utils::VectorRef<FunctionParam*> Params() const { return params_; }
@@ -116,20 +122,22 @@
Block* StartTarget() const { return start_target_; }
private:
- const type::Type* return_type_;
PipelineStage pipeline_stage_;
std::optional<std::array<uint32_t, 3>> workgroup_size_;
- utils::Vector<ReturnAttribute, 1> return_attributes_;
- std::optional<uint32_t> return_location_;
+ struct {
+ const type::Type* type = nullptr;
+ std::optional<enum ReturnBuiltin> builtin;
+ std::optional<Location> location;
+ bool invariant = false;
+ } return_;
utils::Vector<FunctionParam*, 1> params_;
-
Block* start_target_ = nullptr;
};
utils::StringStream& operator<<(utils::StringStream& out, Function::PipelineStage value);
-utils::StringStream& operator<<(utils::StringStream& out, Function::ReturnAttribute value);
+utils::StringStream& operator<<(utils::StringStream& out, enum Function::ReturnBuiltin value);
} // namespace tint::ir
diff --git a/src/tint/ir/transform/add_empty_entry_point_test.cc b/src/tint/ir/transform/add_empty_entry_point_test.cc
index ca918f2..1da2e7d 100644
--- a/src/tint/ir/transform/add_empty_entry_point_test.cc
+++ b/src/tint/ir/transform/add_empty_entry_point_test.cc
@@ -25,7 +25,7 @@
TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) {
auto* expect = R"(
-%unused_entry_point = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+%unused_entry_point = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
ret
}
@@ -43,7 +43,7 @@
mod.functions.Push(ep);
auto* expect = R"(
-%main = func():void [@fragment] -> %b1 {
+%main = @fragment func():void -> %b1 {
%b1 = block {
ret
}
