[ir] Add function return information.

This Cl adds information into the IR on the function return type and any
associated attributes.

Bug: tint:1915
Change-Id: I74cbf2613b4ae575e33a61d04d30b515df6ba796
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/131300
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>

src/tint/ir/builder_impl.cc

@@ -39,6 +39,7 @@
 #include "src/tint/ast/identifier_expression.h"
 #include "src/tint/ast/if_statement.h"
 #include "src/tint/ast/int_literal_expression.h"
+#include "src/tint/ast/invariant_attribute.h"
 #include "src/tint/ast/let.h"
 #include "src/tint/ast/literal_expression.h"
 #include "src/tint/ast/loop_statement.h"
@@ -211,6 +212,7 @@ void BuilderImpl::EmitFunction(const ast::Function* ast_func) {
 
     ast_to_flow_[ast_func] = ir_func;
 
+    const auto* sem = program_->Sem().Get(ast_func);
     if (ast_func->IsEntryPoint()) {
         builder.ir.entry_points.Push(ir_func);
 
@@ -224,7 +226,6 @@ void BuilderImpl::EmitFunction(const ast::Function* ast_func) {
             case ast::PipelineStage::kCompute: {
                 ir_func->pipeline_stage = Function::PipelineStage::kCompute;
 
-                const auto* sem = program_->Sem().Get(ast_func);
                 auto wg_size = sem->WorkgroupSize();
 
                 uint32_t x = wg_size[0].value();
@@ -246,7 +247,49 @@ void BuilderImpl::EmitFunction(const ast::Function* ast_func) {
                 return;
             }
         }
+
+        for (auto* attr : ast_func->return_type_attributes) {
+            tint::Switch(
+                attr,  //
+                [&](const ast::LocationAttribute*) {
+                    ir_func->return_attributes.Push(Function::ReturnAttribute::kLocation);
+                },
+                [&](const ast::InvariantAttribute*) {
+                    ir_func->return_attributes.Push(Function::ReturnAttribute::kInvariant);
+                },
+                [&](const ast::BuiltinAttribute* b) {
+                    if (auto* ident_sem =
+                            program_->Sem()
+                                .Get(b)
+                                ->As<sem::BuiltinEnumExpression<builtin::BuiltinValue>>()) {
+                        switch (ident_sem->Value()) {
+                            case builtin::BuiltinValue::kPosition:
+                                ir_func->return_attributes.Push(
+                                    Function::ReturnAttribute::kPosition);
+                                break;
+                            case builtin::BuiltinValue::kFragDepth:
+                                ir_func->return_attributes.Push(
+                                    Function::ReturnAttribute::kFragDepth);
+                                break;
+                            case builtin::BuiltinValue::kSampleMask:
+                                ir_func->return_attributes.Push(
+                                    Function::ReturnAttribute::kSampleMask);
+                                break;
+                            default:
+                                TINT_ICE(IR, diagnostics_)
+                                    << "Unknown builtin value in return attributes "
+                                    << ident_sem->Value();
+                                return;
+                        }
+                    } else {
+                        TINT_ICE(IR, diagnostics_) << "Builtin attribute sem invalid";
+                        return;
+                    }
+                });
+        }
     }
+    ir_func->return_type = sem->ReturnType()->Clone(clone_ctx_.type_ctx);
+    ir_func->return_location = sem->ReturnLocation();
 
     {
         FlowStackScope scope(this, ir_func);
@@ -400,9 +443,9 @@ void BuilderImpl::EmitBlock(const ast::BlockStatement* block) {
     scopes_.Push();
     TINT_DEFER(scopes_.Pop());
 
-    // Note, this doesn't need to emit a Block as the current block flow node should be
+    // Note, this doesn't need to emit a Block as the current block flow node should be sufficient
-    // sufficient as the blocks all get flattened. Each flow control node will inject the basic
+    // as the blocks all get flattened. Each flow control node will inject the basic blocks it
-    // blocks it requires.
+    // requires.
     EmitStatements(block->statements);
 }
 
@@ -439,9 +482,8 @@ void BuilderImpl::EmitIf(const ast::IfStatement* stmt) {
     }
     current_flow_block = nullptr;
 
-    // If both branches went somewhere, then they both returned, continued or broke. So,
+    // If both branches went somewhere, then they both returned, continued or broke. So, there is no
-    // there is no need for the if merge-block and there is nothing to branch to the merge
+    // need for the if merge-block and there is nothing to branch to the merge block anyway.
-    // block anyway.
     if (IsConnected(if_node->merge.target)) {
         current_flow_block = if_node->merge.target->As<Block>();
     }
@@ -472,8 +514,8 @@ void BuilderImpl::EmitLoop(const ast::LoopStatement* stmt) {
         BranchToIfNeeded(loop_node->start.target);
     }
 
-    // The loop merge can get disconnected if the loop returns directly, or the continuing
+    // The loop merge can get disconnected if the loop returns directly, or the continuing target
-    // target branches, eventually, to the merge, but nothing branched to the continuing target.
+    // branches, eventually, to the merge, but nothing branched to the continuing target.
     current_flow_block = loop_node->merge.target->As<Block>();
     if (!IsConnected(loop_node->merge.target)) {
         current_flow_block = nullptr;
@@ -661,9 +703,9 @@ void BuilderImpl::EmitContinue(const ast::ContinueStatement*) {
 }
 
 // Discard is being treated as an instruction. The semantics in WGSL is demote_to_helper, so the
-// code has to continue as before it just predicates writes. If WGSL grows some kind of
+// code has to continue as before it just predicates writes. If WGSL grows some kind of terminating
-// terminating discard that would probably make sense as a FlowNode but would then require
+// discard that would probably make sense as a FlowNode but would then require figuring out the
-// figuring out the multi-level exit that is triggered.
+// multi-level exit that is triggered.
 void BuilderImpl::EmitDiscard(const ast::DiscardStatement*) {
     auto* inst = builder.Discard();
     current_flow_block->instructions.Push(inst);
@@ -783,8 +825,8 @@ void BuilderImpl::EmitVariable(const ast::Variable* var) {
             // should never be used.
             //
             // TODO(dsinclair): Probably want to store the const variable somewhere and then in
-            // identifier expression log an error if we ever see a const identifier. Add this
+            // identifier expression log an error if we ever see a const identifier. Add this when
-            // when identifiers and variables are supported.
+            // identifiers and variables are supported.
         },
         [&](Default) {
             add_error(var->source, "unknown variable: " + std::string(var->TypeInfo().name));

src/tint/ir/builder_impl_test.cc

@@ -42,7 +42,7 @@ TEST_F(IR_BuilderImplTest, Func) {
     EXPECT_EQ(1u, f->start_target->inbound_branches.Length());
     EXPECT_EQ(1u, f->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func f
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func f(void)
   %fn1 = block
   ret
 func_end
@@ -88,7 +88,7 @@ TEST_F(IR_BuilderImplTest, IfStatement) {
     EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
     EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -136,7 +136,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) {
     EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
     EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -183,7 +183,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_FalseReturns) {
     EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
     EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -230,7 +230,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_BothReturn) {
     EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
     EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -273,7 +273,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_JumpChainToMerge) {
     ASSERT_NE(loop_flow->continuing.target, nullptr);
     ASSERT_NE(loop_flow->merge.target, nullptr);
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -330,7 +330,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) {
     EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
     EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -388,7 +388,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
     EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
     EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -463,7 +463,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) {
     EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
     EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -538,7 +538,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
     EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
     EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -595,7 +595,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn) {
     EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
     EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -657,7 +657,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) {
     // This is 1 because only the loop branch happens. The subsequent if return is dead code.
     EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -711,7 +711,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) {
     EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
     EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -857,7 +857,7 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
     EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
     EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -1001,7 +1001,7 @@ TEST_F(IR_BuilderImplTest, While) {
     EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
     EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -1071,7 +1071,7 @@ TEST_F(IR_BuilderImplTest, While_Return) {
     EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
     EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -1178,7 +1178,7 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) {
     EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
     EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -1237,7 +1237,7 @@ TEST_F(IR_BuilderImplTest, Switch) {
     EXPECT_EQ(3u, flow->merge.target->inbound_branches.Length());
     EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -1301,7 +1301,7 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) {
     EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
     EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -1345,7 +1345,7 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) {
     EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
     EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -1398,7 +1398,7 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) {
     // This is 1 because the if is dead-code eliminated and the return doesn't happen.
     EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -1457,7 +1457,7 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) {
     EXPECT_EQ(0u, flow->merge.target->inbound_branches.Length());
     EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   branch %fn2
 
@@ -1596,7 +1596,7 @@ TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) {
     ASSERT_TRUE(r) << Error();
     auto m = r.Move();
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   %1(ref<function, u32, read_write>) = var function read_write
   ret
@@ -1614,7 +1614,7 @@ TEST_F(IR_BuilderImplTest, Emit_Var_Init) {
     ASSERT_TRUE(r) << Error();
     auto m = r.Move();
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function [@compute @workgroup_size(1, 1, 1)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn1 = block
   %1(ref<function, u32, read_write>) = var function read_write
   store %1(ref<function, u32, read_write>), 2u
@@ -1777,12 +1777,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) {
     ASSERT_TRUE(r) << Error();
     auto m = r.Move();
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func(bool)
   %fn1 = block
   ret true
 func_end
 
-%fn2 = func test_function [@compute @workgroup_size(1, 1, 1)]
+%fn2 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn3 = block
   %1(bool) = call my_func
   %2(bool) = var function read_write
@@ -1812,12 +1812,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) {
     ASSERT_TRUE(r) << Error();
     auto m = r.Move();
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func(bool)
   %fn1 = block
   ret true
 func_end
 
-%fn2 = func test_function [@compute @workgroup_size(1, 1, 1)]
+%fn2 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn3 = block
   %1(bool) = call my_func
   %2(bool) = var function read_write
@@ -1994,12 +1994,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound) {
     ASSERT_TRUE(r) << Error();
     auto m = r.Move();
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func(f32)
   %fn1 = block
   ret 0.0f
 func_end
 
-%fn2 = func test_function [@compute @workgroup_size(1, 1, 1)]
+%fn2 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn3 = block
   %1(f32) = call my_func
   %2(bool) = lt %1(f32), 2.0f
@@ -2036,12 +2036,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound_WithConstEval) {
     ASSERT_TRUE(r) << Error();
     auto m = r.Move();
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func(bool)
   %fn1 = block
   ret true
 func_end
 
-%fn2 = func test_function [@compute @workgroup_size(1, 1, 1)]
+%fn2 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn3 = block
   %1(bool) = call my_func, false
   ret
@@ -2137,7 +2137,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) {
 store %1(ref<private, f32, read_write>), 1.0f
 ret
 
-%fn1 = func test_function [@compute @workgroup_size(1, 1, 1)]
+%fn1 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn2 = block
   %2(vec3<f32>) = construct 2.0f, 3.0f, %1(ref<private, f32, read_write>)
   ret
@@ -2161,7 +2161,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) {
 store %1(ref<private, i32, read_write>), 1i
 ret
 
-%fn1 = func test_function [@compute @workgroup_size(1, 1, 1)]
+%fn1 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn2 = block
   %2(f32) = convert i32, %1(ref<private, i32, read_write>)
   ret
@@ -2179,7 +2179,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_MaterializedCall) {
     ASSERT_TRUE(r) << Error();
     auto m = r.Move();
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(f32)
   %fn1 = block
   ret 2.0f
 func_end
@@ -2201,7 +2201,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Builtin) {
 store %1(ref<private, f32, read_write>), 1.0f
 ret
 
-%fn1 = func test_function [@compute @workgroup_size(1, 1, 1)]
+%fn1 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn2 = block
   %2(f32) = asin %1(ref<private, f32, read_write>)
   ret
@@ -2219,7 +2219,7 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Vertex) {
     ASSERT_TRUE(r) << Error();
     auto m = r.Move();
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test [@vertex]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(vec4<f32>) [@vertex ra: @position]
   %fn1 = block
   ret vec4<f32> 0.0f
 func_end
@@ -2235,7 +2235,7 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Fragment) {
     ASSERT_TRUE(r) << Error();
     auto m = r.Move();
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test [@fragment]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(void) [@fragment]
   %fn1 = block
   ret
 func_end
@@ -2251,7 +2251,7 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Compute) {
     ASSERT_TRUE(r) << Error();
     auto m = r.Move();
 
-    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test [@compute @workgroup_size(8, 4, 2)]
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(void) [@compute @workgroup_size(8, 4, 2)]
   %fn1 = block
   ret
 func_end
@@ -2259,5 +2259,105 @@ func_end
 )");
 }
 
+TEST_F(IR_BuilderImplTest, EmitFunction_Return) {
+    Func("test", utils::Empty, ty.vec3<f32>(), utils::Vector{Return(vec3<f32>(0_f, 0_f, 0_f))},
+         utils::Empty);
+
+    auto r = Build();
+    ASSERT_TRUE(r) << Error();
+    auto m = r.Move();
+
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(vec3<f32>)
+  %fn1 = block
+  ret vec3<f32> 0.0f
+func_end
+
+)");
+}
+
+TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPosition) {
+    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::kVertex)},
+         utils::Vector{Builtin(builtin::BuiltinValue::kPosition)});
+
+    auto r = Build();
+    ASSERT_TRUE(r) << Error();
+    auto m = r.Move();
+
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(vec4<f32>) [@vertex ra: @position]
+  %fn1 = block
+  ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
+func_end
+
+)");
+}
+
+TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPositionInvariant) {
+    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::kVertex)},
+         utils::Vector{Builtin(builtin::BuiltinValue::kPosition), Invariant()});
+
+    auto r = Build();
+    ASSERT_TRUE(r) << Error();
+    auto m = r.Move();
+
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(vec4<f32>) [@vertex ra: @position @invariant]
+  %fn1 = block
+  ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
+func_end
+
+)");
+}
+
+TEST_F(IR_BuilderImplTest, EmitFunction_ReturnLocation) {
+    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)});
+
+    auto r = Build();
+    ASSERT_TRUE(r) << Error();
+    auto m = r.Move();
+
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(vec4<f32>) [@fragment ra: @location(1)]
+  %fn1 = block
+  ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
+func_end
+
+)");
+}
+
+TEST_F(IR_BuilderImplTest, EmitFunction_ReturnFragDepth) {
+    Func("test", utils::Empty, ty.f32(), utils::Vector{Return(1_f)},
+         utils::Vector{Stage(ast::PipelineStage::kFragment)},
+         utils::Vector{Builtin(builtin::BuiltinValue::kFragDepth)});
+
+    auto r = Build();
+    ASSERT_TRUE(r) << Error();
+    auto m = r.Move();
+
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(f32) [@fragment ra: @frag_depth]
+  %fn1 = block
+  ret 1.0f
+func_end
+
+)");
+}
+
+TEST_F(IR_BuilderImplTest, EmitFunction_ReturnSampleMask) {
+    Func("test", utils::Empty, ty.u32(), utils::Vector{Return(1_u)},
+         utils::Vector{Stage(ast::PipelineStage::kFragment)},
+         utils::Vector{Builtin(builtin::BuiltinValue::kSampleMask)});
+
+    auto r = Build();
+    ASSERT_TRUE(r) << Error();
+    auto m = r.Move();
+
+    EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(u32) [@fragment ra: @sample_mask]
+  %fn1 = block
+  ret 1u
+func_end
+
+)");
+}
+
 }  // namespace
 }  // namespace tint::ir

src/tint/ir/disassembler.cc

@@ -20,6 +20,7 @@
 #include "src/tint/ir/switch.h"
 #include "src/tint/ir/terminator.h"
 #include "src/tint/switch.h"
+#include "src/tint/type/type.h"
 #include "src/tint/utils/scoped_assignment.h"
 
 namespace tint::ir {
@@ -92,7 +93,9 @@ void Disassembler::Walk(const FlowNode* node) {
         [&](const ir::Function* f) {
             TINT_SCOPED_ASSIGNMENT(in_function_, true);
 
-            Indent() << "%fn" << GetIdForNode(f) << " = func " << f->name.Name();
+            Indent() << "%fn" << GetIdForNode(f) << " = func " << f->name.Name() << "("
+                     << f->return_type->FriendlyName() << ")";
+
             if (f->pipeline_stage != Function::PipelineStage::kUndefined) {
                 out_ << " [@" << f->pipeline_stage;
 
@@ -101,6 +104,18 @@ void Disassembler::Walk(const FlowNode* node) {
                     out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2]
                          << ")";
                 }
+
+                if (!f->return_attributes.IsEmpty()) {
+                    out_ << " ra:";
+
+                    for (auto attr : f->return_attributes) {
+                        out_ << " @" << attr;
+                        if (attr == Function::ReturnAttribute::kLocation) {
+                            out_ << "(" << f->return_location.value() << ")";
+                        }
+                    }
+                }
+
                 out_ << "]";
             }
             out_ << std::endl;

src/tint/ir/function.cc

@@ -36,4 +36,22 @@ utils::StringStream& operator<<(utils::StringStream& out, Function::PipelineStag
     return out << "<unknown>";
 }
 
+utils::StringStream& operator<<(utils::StringStream& out, Function::ReturnAttribute value) {
+    switch (value) {
+        case Function::ReturnAttribute::kLocation:
+            return out << "location";
+        case Function::ReturnAttribute::kFragDepth:
+            return out << "frag_depth";
+        case Function::ReturnAttribute::kSampleMask:
+            return out << "sample_mask";
+        case Function::ReturnAttribute::kPosition:
+            return out << "position";
+        case Function::ReturnAttribute::kInvariant:
+            return out << "invariant";
+        default:
+            break;
+    }
+    return out << "<unknown>";
+}
+
 }  // namespace tint::ir

src/tint/ir/function.h

@@ -19,6 +19,7 @@
 
 #include "src/tint/ir/flow_node.h"
 #include "src/tint/symbol.h"
+#include "src/tint/type/type.h"
 
 // Forward declarations
 namespace tint::ir {
@@ -43,6 +44,22 @@ class Function : public utils::Castable<Function, FlowNode> {
         kVertex,
     };
 
+    /// Attributes attached to return types
+    enum class ReturnAttribute {
+        /// No return attribute
+        kNone,
+        /// Location attribute
+        kLocation,
+        /// Builtin Position attribute
+        kPosition,
+        /// Builtin FragDepth attribute
+        kFragDepth,
+        /// Builtin SampleMask
+        kSampleMask,
+        /// Invariant attribute
+        kInvariant,
+    };
+
     /// Constructor
     Function();
     ~Function() override;
@@ -56,6 +73,13 @@ class Function : public utils::Castable<Function, FlowNode> {
     /// If this is a `compute` entry point, holds the workgroup size information
     std::optional<std::array<uint32_t, 3>> workgroup_size;
 
+    /// The function return type
+    const type::Type* return_type = nullptr;
+    /// The function return attributes if any
+    utils::Vector<ReturnAttribute, 1> return_attributes;
+    /// If the return attribute is `kLocation` this stores the location value.
+    std::optional<uint32_t> return_location;
+
     /// The start target is the first block in a function.
     Block* start_target = nullptr;
     /// The end target is the end of the function. It is used as the branch target if a return is
@@ -64,6 +88,7 @@ class Function : public utils::Castable<Function, FlowNode> {
 };
 
 utils::StringStream& operator<<(utils::StringStream& out, Function::PipelineStage value);
+utils::StringStream& operator<<(utils::StringStream& out, Function::ReturnAttribute value);
 
 }  // namespace tint::ir