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
   }