[ir] Add explicit Return instructions.

This Cl adds a `ret` instruction into the IR. The `FunctionTerminator`
block has been removed.

Bug: tint:1718
Change-Id: Ie5fcdbfa8983b4c960773494b0c58793bd9ef503
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/134461
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>

src/tint/BUILD.gn

@@ -1236,8 +1236,6 @@ if (tint_build_ir) {
       "ir/function.h",
       "ir/function_param.cc",
       "ir/function_param.h",
-      "ir/function_terminator.cc",
-      "ir/function_terminator.h",
       "ir/if.cc",
       "ir/if.h",
       "ir/instruction.cc",
@@ -1248,6 +1246,8 @@ if (tint_build_ir) {
       "ir/loop.h",
       "ir/module.cc",
       "ir/module.h",
+      "ir/return.cc",
+      "ir/return.h",
       "ir/root_terminator.cc",
       "ir/root_terminator.h",
       "ir/store.cc",

src/tint/CMakeLists.txt

@@ -746,8 +746,6 @@ if(${TINT_BUILD_IR})
     ir/function.h
     ir/function_param.cc
     ir/function_param.h
-    ir/function_terminator.cc
-    ir/function_terminator.h
     ir/if.cc
     ir/if.h
     ir/instruction.cc
@@ -758,6 +756,8 @@ if(${TINT_BUILD_IR})
     ir/loop.h
     ir/module.cc
     ir/module.h
+    ir/return.cc
+    ir/return.h
     ir/root_terminator.cc
     ir/root_terminator.h
     ir/store.cc

src/tint/ir/branch.cc

@@ -22,14 +22,19 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Branch);
 
 namespace tint::ir {
 
-Branch::Branch(Block* to, utils::VectorRef<Value*> args) : to_(to), args_(std::move(args)) {
+Branch::Branch(utils::VectorRef<Value*> args) : args_(std::move(args)) {
-    TINT_ASSERT(IR, to_);
-    to_->AddInboundBranch(this);
     for (auto* arg : args) {
         arg->AddUsage(this);
     }
 }
 
+Branch::Branch(Block* to, utils::VectorRef<Value*> args) : Branch(args) {
+    to_ = to;
+
+    TINT_ASSERT(IR, to_);
+    to_->AddInboundBranch(this);
+}
+
 Branch::~Branch() = default;
 
 }  // namespace tint::ir

src/tint/ir/branch.h

@@ -26,7 +26,7 @@ class Block;
 
 namespace tint::ir {
 
-/// A branch instruction. A branch is a walk terminating jump.
+/// A branch instruction.
 class Branch : public utils::Castable<Branch, Instruction> {
   public:
     /// Constructor
@@ -41,8 +41,13 @@ class Branch : public utils::Castable<Branch, Instruction> {
     /// @returns the branch arguments
     utils::VectorRef<Value*> Args() const { return args_; }
 
+  protected:
+    /// Constructor
+    /// @param args the branch arguments
+    explicit Branch(utils::VectorRef<Value*> args);
+
   private:
-    Block* to_;
+    Block* to_ = nullptr;
     utils::Vector<Value*, 2> args_;
 };
 

src/tint/ir/builder.cc

@@ -41,10 +41,6 @@ RootTerminator* Builder::CreateRootTerminator() {
     return ir.blocks.Create<RootTerminator>();
 }
 
-FunctionTerminator* Builder::CreateFunctionTerminator() {
-    return ir.blocks.Create<FunctionTerminator>();
-}
-
 Function* Builder::CreateFunction(std::string_view name,
                                   const type::Type* return_type,
                                   Function::PipelineStage stage,
@@ -53,7 +49,6 @@ Function* Builder::CreateFunction(std::string_view name,
 
     auto* ir_func = ir.values.Create<Function>(return_type, stage, wg_size);
     ir_func->SetStartTarget(CreateBlock());
-    ir_func->SetEndTarget(CreateFunctionTerminator());
     ir.SetName(ir_func, name);
     return ir_func;
 }
@@ -214,6 +209,10 @@ ir::Branch* Builder::Branch(Block* to, utils::VectorRef<Value*> args) {
     return ir.values.Create<ir::Branch>(to, args);
 }
 
+ir::Return* Builder::Return(Function* func, utils::VectorRef<Value*> args) {
+    return ir.values.Create<ir::Return>(func, args);
+}
+
 ir::BlockParam* Builder::BlockParam(const type::Type* type) {
     return ir.values.Create<ir::BlockParam>(type);
 }

src/tint/ir/builder.h

@@ -28,11 +28,11 @@
 #include "src/tint/ir/discard.h"
 #include "src/tint/ir/function.h"
 #include "src/tint/ir/function_param.h"
-#include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/if.h"
 #include "src/tint/ir/load.h"
 #include "src/tint/ir/loop.h"
 #include "src/tint/ir/module.h"
+#include "src/tint/ir/return.h"
 #include "src/tint/ir/root_terminator.h"
 #include "src/tint/ir/store.h"
 #include "src/tint/ir/switch.h"
@@ -65,9 +65,6 @@ class Builder {
     /// @returns a new root terminator flow node
     RootTerminator* CreateRootTerminator();
 
-    /// @returns a new function terminator flow node
-    FunctionTerminator* CreateFunctionTerminator();
-
     /// Creates a function flow node
     /// @param name the function name
     /// @param return_type the function return type
@@ -333,6 +330,12 @@ class Builder {
     /// @returns the instruction
     ir::Var* Declare(const type::Type* type);
 
+    /// Creates a return instruction
+    /// @param func the function being returned
+    /// @param args the return arguments
+    /// @returns the instruction
+    ir::Return* Return(Function* func, utils::VectorRef<Value*> args = {});
+
     /// Creates a branch declaration
     /// @param to the node being branched too
     /// @param args the branch arguments

src/tint/ir/debug.cc

@@ -18,9 +18,9 @@
 #include <unordered_set>
 
 #include "src/tint/ir/block.h"
-#include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/if.h"
 #include "src/tint/ir/loop.h"
+#include "src/tint/ir/return.h"
 #include "src/tint/ir/switch.h"
 #include "src/tint/switch.h"
 #include "src/tint/utils/string_stream.h"
@@ -54,11 +54,7 @@ std::string Debug::AsDotGraph(const Module* mod) {
         }
         visited.insert(blk);
 
-        tint::Switch(
+        tint::Switch(blk,  //
-            blk,
-            [&](const ir::FunctionTerminator*) {
-                // Already done
-            },
                      [&](const ir::Block* b) {
                          if (block_to_name.count(b) == 0) {
                              out << name_for(b) << R"( [label="block"])" << std::endl;
@@ -71,7 +67,12 @@ std::string Debug::AsDotGraph(const Module* mod) {
                          }
 
                          out << std::endl;
+
+                         if (b->Branch()->Is<ir::Return>()) {
+                             return;
+                         } else {
                              Graph(b->Branch()->To());
+                         }
                      });
     };
 
@@ -81,7 +82,6 @@ std::string Debug::AsDotGraph(const Module* mod) {
         out << "subgraph cluster_" << mod->NameOf(func).Name() << " {" << std::endl;
         out << R"(label=")" << mod->NameOf(func).Name() << R"(")" << std::endl;
         out << name_for(func->StartTarget()) << R"( [label="start"])" << std::endl;
-        out << name_for(func->EndTarget()) << R"( [label="end"])" << std::endl;
         Graph(func->StartTarget());
         out << "}" << std::endl;
     }

src/tint/ir/disassembler.cc

@@ -25,10 +25,10 @@
 #include "src/tint/ir/construct.h"
 #include "src/tint/ir/convert.h"
 #include "src/tint/ir/discard.h"
-#include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/if.h"
 #include "src/tint/ir/load.h"
 #include "src/tint/ir/loop.h"
+#include "src/tint/ir/return.h"
 #include "src/tint/ir/root_terminator.h"
 #include "src/tint/ir/store.h"
 #include "src/tint/ir/switch.h"
@@ -108,11 +108,6 @@ void Disassembler::Walk(const Block* blk) {
 
     tint::Switch(
         blk,
-        [&](const ir::FunctionTerminator* t) {
-            TINT_ASSERT(IR, in_function_);
-            Indent() << "%b" << IdOf(t) << " = func_terminator" << std::endl;
-            in_function_ = false;
-        },
         [&](const ir::RootTerminator* t) {
             TINT_ASSERT(IR, !in_function_);
             Indent() << "%b" << IdOf(t) << " = root_terminator" << std::endl << std::endl;
@@ -142,7 +137,7 @@ void Disassembler::Walk(const Block* blk) {
             }
             Indent() << "}" << std::endl;
 
-            if (!b->Branch()->To()->Is<FunctionTerminator>()) {
+            if (!b->Branch()->Is<ir::Return>()) {
                 out_ << std::endl;
             }
         });
@@ -186,7 +181,6 @@ void Disassembler::EmitFunction(const Function* func) {
     {
         ScopedIndent si(indent_size_);
         Walk(func->StartTarget());
-        Walk(func->EndTarget());
     }
     Indent() << "}" << std::endl;
 }
@@ -415,12 +409,14 @@ void Disassembler::EmitSwitch(const Switch* s) {
 
 void Disassembler::EmitBranch(const Branch* b) {
     std::string suffix = "";
+    if (b->Is<ir::Return>()) {
+        out_ << "ret";
+    } else {
         out_ << "br %b" << IdOf(b->To());
-    if (b->To()->Is<FunctionTerminator>()) {
+        if (b->To()->Is<RootTerminator>()) {
-        suffix = "return";
-    } else if (b->To()->Is<RootTerminator>()) {
             suffix = "root_end";
         }
+    }
 
     if (!b->Args().IsEmpty()) {
         out_ << " ";

src/tint/ir/from_program.cc

@@ -166,6 +166,21 @@ class Impl {
         diagnostics_.add_error(tint::diag::System::IR, err, s);
     }
 
+    void SetBranch(Branch* br) {
+        TINT_ASSERT(IR, current_flow_block_);
+        TINT_ASSERT(IR, !current_flow_block_->HasBranchTarget());
+
+        current_flow_block_->Instructions().Push(br);
+        current_flow_block_ = nullptr;
+    }
+
+    void SetBranchIfNeeded(Branch* br) {
+        if (!current_flow_block_ || current_flow_block_->HasBranchTarget()) {
+            return;
+        }
+        SetBranch(br);
+    }
+
     void BranchTo(Block* node, utils::VectorRef<Value*> args = {}) {
         TINT_ASSERT(IR, current_flow_block_);
         TINT_ASSERT(IR, !current_flow_block_->HasBranchTarget());
@@ -337,7 +352,7 @@ class Impl {
 
             // If the branch target has already been set then a `return` was called. Only set in
             // the case where `return` wasn't called.
-            BranchToIfNeeded(current_function_->EndTarget());
+            SetBranchIfNeeded(builder_.Return(current_function_));
         }
 
         TINT_ASSERT(IR, control_stack_.IsEmpty());
@@ -716,8 +731,7 @@ class Impl {
             }
             ret_value.Push(ret.Get());
         }
-
+        SetBranch(builder_.Return(current_function_, std::move(ret_value)));
-        BranchTo(current_function_->EndTarget(), std::move(ret_value));
     }
 
     void EmitBreak(const ast::BreakStatement*) {

src/tint/ir/from_program_binary_test.cc

@@ -36,17 +36,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Add) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:u32 = add %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -72,9 +70,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Increment) {
     %3:u32 = load %v1
     %4:u32 = add %3, 1u
     store %v1, %4
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -100,9 +97,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAdd) {
     %3:u32 = load %v1
     %4:u32 = add %3, 1u
     store %v1, %4
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -117,17 +113,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Subtract) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:u32 = sub %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -153,9 +147,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Decrement) {
     %3:i32 = load %v1
     %4:i32 = sub %3, 1i
     store %v1, %4
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -181,9 +174,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundSubtract) {
     %3:u32 = load %v1
     %4:u32 = sub %3, 1u
     store %v1, %4
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -198,17 +190,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Multiply) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:u32 = mul %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -234,9 +224,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundMultiply) {
     %3:u32 = load %v1
     %4:u32 = mul %3, 1u
     store %v1, %4
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -251,17 +240,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Div) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:u32 = div %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -287,9 +274,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundDiv) {
     %3:u32 = load %v1
     %4:u32 = div %3, 1u
     store %v1, %4
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -304,17 +290,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Modulo) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:u32 = mod %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -340,9 +324,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundModulo) {
     %3:u32 = load %v1
     %4:u32 = mod %3, 1u
     store %v1, %4
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -357,17 +340,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_And) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:u32 = and %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -393,9 +374,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAnd) {
     %3:bool = load %v1
     %4:bool = and %3, false
     store %v1, %4
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -410,17 +390,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Or) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:u32 = or %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -446,9 +424,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundOr) {
     %3:bool = load %v1
     %4:bool = or %3, false
     store %v1, %4
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -463,17 +440,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Xor) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:u32 = xor %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -499,9 +474,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundXor) {
     %3:u32 = load %v1
     %4:u32 = xor %3, 1u
     store %v1, %4
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -516,40 +490,39 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():bool -> %b1 {
   %b1 = block {
-    br %b2 true  # return
+    ret true
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:bool = call %my_func
-    if %3 [t: %b4, f: %b5, m: %b6]
+    if %3 [t: %b3, f: %b4, m: %b5]
       # True block
+      %b3 = block {
+        br %b5 false
+      }
+
+      # False block
       %b4 = block {
-        br %b6 false
+        br %b5 %3
-      }
-
-      # False block
-      %b5 = block {
-        br %b6 %3
       }
 
     # Merge block
-    %b6 = block (%4:bool) {
+    %b5 = block (%4:bool) {
-      if %4:bool [t: %b7, f: %b8, m: %b9]
+      if %4:bool [t: %b6, f: %b7, m: %b8]
         # True block
-        %b7 = block {
+        %b6 = block {
-          br %b9
+          br %b8
         }
 
         # False block
-        %b8 = block {
+        %b7 = block {
-          br %b9
+          br %b8
         }
 
       # Merge block
-      %b9 = block {
+      %b8 = block {
-        br %b10  # return
+        ret
       }
 
     }
@@ -557,7 +530,6 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) {
 
   }
 
-  %b10 = func_terminator
 }
 )");
 }
@@ -572,40 +544,39 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():bool -> %b1 {
   %b1 = block {
-    br %b2 true  # return
+    ret true
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:bool = call %my_func
-    if %3 [t: %b4, f: %b5, m: %b6]
+    if %3 [t: %b3, f: %b4, m: %b5]
       # True block
+      %b3 = block {
+        br %b5 %3
+      }
+
+      # False block
       %b4 = block {
-        br %b6 %3
+        br %b5 true
-      }
-
-      # False block
-      %b5 = block {
-        br %b6 true
       }
 
     # Merge block
-    %b6 = block (%4:bool) {
+    %b5 = block (%4:bool) {
-      if %4:bool [t: %b7, f: %b8, m: %b9]
+      if %4:bool [t: %b6, f: %b7, m: %b8]
         # True block
-        %b7 = block {
+        %b6 = block {
-          br %b9
+          br %b8
         }
 
         # False block
-        %b8 = block {
+        %b7 = block {
-          br %b9
+          br %b8
         }
 
       # Merge block
-      %b9 = block {
+      %b8 = block {
-        br %b10  # return
+        ret
       }
 
     }
@@ -613,7 +584,6 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) {
 
   }
 
-  %b10 = func_terminator
 }
 )");
 }
@@ -628,17 +598,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Equal) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:bool = eq %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -653,17 +621,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_NotEqual) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:bool = neq %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -678,17 +644,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LessThan) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:bool = lt %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -703,17 +667,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_GreaterThan) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:bool = gt %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -728,17 +690,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LessThanEqual) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:bool = lte %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -753,17 +713,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_GreaterThanEqual) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:bool = gte %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -778,17 +736,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_ShiftLeft) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:u32 = shiftl %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -814,9 +770,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftLeft) {
     %3:u32 = load %v1
     %4:u32 = shiftl %3, 1u
     store %v1, %4
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -831,17 +786,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_ShiftRight) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 0u  # return
+    ret 0u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:u32 = shiftr %3, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -867,9 +820,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftRight) {
     %3:u32 = load %v1
     %4:u32 = shiftr %3, 1u
     store %v1, %4
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -886,38 +838,36 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():f32 -> %b1 {
   %b1 = block {
-    br %b2 0.0f  # return
+    ret 0.0f
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:f32 = call %my_func
     %4:bool = lt %3, 2.0f
-    if %4 [t: %b4, f: %b5, m: %b6]
+    if %4 [t: %b3, f: %b4, m: %b5]
       # True block
-      %b4 = block {
+      %b3 = block {
         %5:f32 = call %my_func
         %6:f32 = call %my_func
         %7:f32 = mul 2.29999995231628417969f, %6
         %8:f32 = div %5, %7
         %9:bool = gt 2.5f, %8
-        br %b6 %9
+        br %b5 %9
       }
 
       # False block
-      %b5 = block {
+      %b4 = block {
-        br %b6 %4
+        br %b5 %4
       }
 
     # Merge block
-    %b6 = block (%tint_symbol:bool) {
+    %b5 = block (%tint_symbol:bool) {
-      br %b7  # return
+      ret
     }
 
   }
 
-  %b7 = func_terminator
 }
 )");
 }
@@ -933,16 +883,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound_WithConstEval) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func(%p:bool):bool -> %b1 {
   %b1 = block {
-    br %b2 true  # return
+    ret true
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %tint_symbol:bool = call %my_func, false
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }

src/tint/ir/from_program_builtin_test.cc

@@ -46,9 +46,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Builtin) {
   %b3 = block {
     %3:f32 = load %i
     %tint_symbol:f32 = asin %3
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }

src/tint/ir/from_program_call_test.cc

@@ -37,17 +37,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Bitcast) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():f32 -> %b1 {
   %b1 = block {
-    br %b2 0.0f  # return
+    ret 0.0f
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:f32 = call %my_func
     %tint_symbol:f32 = bitcast %3
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -65,9 +63,8 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Discard) {
     EXPECT_EQ(Disassemble(m.Get()), R"(%test_function = func():void [@fragment] -> %b1 {
   %b1 = block {
     discard
-    br %b2  # return
+    ret
   }
-  %b2 = func_terminator
 }
 )");
 }
@@ -82,16 +79,14 @@ TEST_F(IR_BuilderImplTest, EmitStatement_UserFunction) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func(%p:f32):void -> %b1 {
   %b1 = block {
-    br %b2  # return
+    ret
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %4:void = call %my_func, 6.0f
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -116,9 +111,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) {
   %b3 = block {
     %3:i32 = load %i
     %tint_symbol:f32 = convert i32, %3
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -161,9 +155,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) {
   %b3 = block {
     %3:f32 = load %i
     %tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %3
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }

src/tint/ir/from_program_function_test.cc

@@ -36,9 +36,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Vertex) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():vec4<f32> [@vertex ra: @position] -> %b1 {
   %b1 = block {
-    br %b2 vec4<f32> 0.0f  # return
+    ret vec4<f32> 0.0f
   }
-  %b2 = func_terminator
 }
 )");
 }
@@ -52,9 +51,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Fragment) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():void [@fragment] -> %b1 {
   %b1 = block {
-    br %b2  # return
+    ret
   }
-  %b2 = func_terminator
 }
 )");
 }
@@ -69,9 +67,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Compute) {
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%test = func():void [@compute @workgroup_size(8, 4, 2)] -> %b1 {
   %b1 = block {
-    br %b2  # return
+    ret
   }
-  %b2 = func_terminator
 }
 )");
 }
@@ -85,9 +82,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Return) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():vec3<f32> -> %b1 {
   %b1 = block {
-    br %b2 vec3<f32> 0.0f  # return
+    ret vec3<f32> 0.0f
   }
-  %b2 = func_terminator
 }
 )");
 }
@@ -102,9 +98,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPosition) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():vec4<f32> [@vertex ra: @position] -> %b1 {
   %b1 = block {
-    br %b2 vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f  # return
+    ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
   }
-  %b2 = func_terminator
 }
 )");
 }
@@ -120,9 +115,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPositionInvariant) {
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%test = func():vec4<f32> [@vertex ra: @position @invariant] -> %b1 {
   %b1 = block {
-    br %b2 vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f  # return
+    ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
   }
-  %b2 = func_terminator
 }
 )");
 }
@@ -137,9 +131,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_ReturnLocation) {
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%test = func():vec4<f32> [@fragment ra: @location(1)] -> %b1 {
   %b1 = block {
-    br %b2 vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f  # return
+    ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
   }
-  %b2 = func_terminator
 }
 )");
 }
@@ -154,9 +147,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_ReturnFragDepth) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():f32 [@fragment ra: @frag_depth] -> %b1 {
   %b1 = block {
-    br %b2 1.0f  # return
+    ret 1.0f
   }
-  %b2 = func_terminator
 }
 )");
 }
@@ -171,9 +163,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_ReturnSampleMask) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():u32 [@fragment ra: @sample_mask] -> %b1 {
   %b1 = block {
-    br %b2 1u  # return
+    ret 1u
   }
-  %b2 = func_terminator
 }
 )");
 }

src/tint/ir/from_program_materialize_test.cc

@@ -36,9 +36,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_MaterializedCall) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%test_function = func():f32 -> %b1 {
   %b1 = block {
-    br %b2 2.0f  # return
+    ret 2.0f
   }
-  %b2 = func_terminator
 }
 )");
 }

src/tint/ir/from_program_store_test.cc

@@ -46,9 +46,8 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Assign) {
 %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
   %b3 = block {
     store %a, 4u
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }

src/tint/ir/from_program_test.cc

@@ -19,7 +19,6 @@
 #include "src/tint/ast/int_literal_expression.h"
 #include "src/tint/constant/scalar.h"
 #include "src/tint/ir/block.h"
-#include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/if.h"
 #include "src/tint/ir/loop.h"
 #include "src/tint/ir/switch.h"
@@ -64,17 +63,13 @@ TEST_F(IR_BuilderImplTest, Func) {
 
     auto* f = m->functions[0];
     ASSERT_NE(f->StartTarget(), nullptr);
-    ASSERT_NE(f->EndTarget(), nullptr);
-
-    EXPECT_EQ(1u, f->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined);
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():void -> %b1 {
   %b1 = block {
-    br %b2  # return
+    ret
   }
-  %b2 = func_terminator
 }
 )");
 }
@@ -89,17 +84,13 @@ TEST_F(IR_BuilderImplTest, Func_WithParam) {
 
     auto* f = m->functions[0];
     ASSERT_NE(f->StartTarget(), nullptr);
-    ASSERT_NE(f->EndTarget(), nullptr);
-
-    EXPECT_EQ(1u, f->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined);
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%f = func(%a:u32):u32 -> %b1 {
   %b1 = block {
-    br %b2 %a  # return
+    ret %a
   }
-  %b2 = func_terminator
 }
 )");
 }
@@ -115,17 +106,13 @@ TEST_F(IR_BuilderImplTest, Func_WithMultipleParam) {
 
     auto* f = m->functions[0];
     ASSERT_NE(f->StartTarget(), nullptr);
-    ASSERT_NE(f->EndTarget(), nullptr);
-
-    EXPECT_EQ(1u, f->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined);
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%f = func(%a:u32, %b:i32, %c:bool):void -> %b1 {
   %b1 = block {
-    br %b2  # return
+    ret
   }
-  %b2 = func_terminator
 }
 )");
 }
@@ -151,12 +138,10 @@ TEST_F(IR_BuilderImplTest, IfStatement) {
     auto* flow = FindSingleValue<ir::If>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
     EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -174,12 +159,11 @@ TEST_F(IR_BuilderImplTest, IfStatement) {
 
     # Merge block
     %b4 = block {
-      br %b5  # return
+      ret
     }
 
   }
 
-  %b5 = func_terminator
 }
 )");
 }
@@ -195,12 +179,10 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) {
     auto* flow = FindSingleValue<ir::If>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
     EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -208,7 +190,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) {
     if true [t: %b2, f: %b3, m: %b4]
       # True block
       %b2 = block {
-        br %b5  # return
+        ret
       }
       # False block
       %b3 = block {
@@ -217,12 +199,11 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) {
 
     # Merge block
     %b4 = block {
-      br %b5  # return
+      ret
     }
 
   }
 
-  %b5 = func_terminator
 }
 )");
 }
@@ -238,12 +219,10 @@ TEST_F(IR_BuilderImplTest, IfStatement_FalseReturns) {
     auto* flow = FindSingleValue<ir::If>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
     EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -256,16 +235,15 @@ TEST_F(IR_BuilderImplTest, IfStatement_FalseReturns) {
 
       # False block
       %b3 = block {
-        br %b5  # return
+        ret
       }
     # Merge block
     %b4 = block {
-      br %b5  # return
+      ret
     }
 
   }
 
-  %b5 = func_terminator
 }
 )");
 }
@@ -281,12 +259,10 @@ TEST_F(IR_BuilderImplTest, IfStatement_BothReturn) {
     auto* flow = FindSingleValue<ir::If>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -294,16 +270,15 @@ TEST_F(IR_BuilderImplTest, IfStatement_BothReturn) {
     if true [t: %b2, f: %b3]
       # True block
       %b2 = block {
-        br %b4  # return
+        ret
       }
       # False block
       %b3 = block {
-        br %b4  # return
+        ret
       }
 
   }
 
-  %b4 = func_terminator
 }
 )");
 }
@@ -354,12 +329,11 @@ TEST_F(IR_BuilderImplTest, IfStatement_JumpChainToMerge) {
 
     # Merge block
     %b4 = block {
-      br %b8  # return
+      ret
     }
 
   }
 
-  %b8 = func_terminator
 }
 )");
 }
@@ -375,12 +349,10 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) {
     auto* flow = FindSingleValue<ir::Loop>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     EXPECT_EQ(2u, flow->Start()->InboundBranches().Length());
     EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -397,12 +369,11 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) {
 
     # Merge block
     %b4 = block {
-      br %b5  # return
+      ret
     }
 
   }
 
-  %b5 = func_terminator
 }
 )");
 }
@@ -421,7 +392,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
     auto* if_flow = FindSingleValue<ir::If>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
     EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
@@ -429,7 +399,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
     EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
     EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -462,12 +431,11 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
 
     # Merge block
     %b4 = block {
-      br %b8  # return
+      ret
     }
 
   }
 
-  %b8 = func_terminator
 }
 )");
 }
@@ -485,7 +453,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) {
     auto* break_if_flow = FindSingleValue<ir::If>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
     EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
@@ -493,7 +460,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) {
     EXPECT_EQ(1u, break_if_flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, break_if_flow->False()->InboundBranches().Length());
     EXPECT_EQ(2u, break_if_flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -526,12 +492,11 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) {
 
     # Merge block
     %b4 = block {
-      br %b8  # return
+      ret
     }
 
   }
 
-  %b8 = func_terminator
 }
 )");
 }
@@ -577,12 +542,11 @@ TEST_F(IR_BuilderImplTest, Loop_Continuing_Body_Scope) {
 
     # Merge block
     %b4 = block {
-      br %b8  # return
+      ret
     }
 
   }
 
-  %b8 = func_terminator
 }
 )");
 }
@@ -600,7 +564,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
     auto* if_flow = FindSingleValue<ir::If>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
     EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
@@ -608,7 +571,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
     EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
     EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -618,7 +580,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
         if true [t: %b4, f: %b5, m: %b6]
           # True block
           %b4 = block {
-            br %b7  # return
+            ret
           }
           # False block
           %b5 = block {
@@ -641,7 +603,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
 
   }
 
-  %b7 = func_terminator
 }
 )");
 }
@@ -657,19 +618,17 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn) {
     auto* loop_flow = FindSingleValue<ir::Loop>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
     EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
   %b1 = block {
     loop [s: %b2, c: %b3]
       %b2 = block {
-        br %b4  # return
+        ret
       }
       # Continuing block
       %b3 = block {
@@ -679,7 +638,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn) {
 
   }
 
-  %b4 = func_terminator
 }
 )");
 }
@@ -704,35 +662,33 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) {
     auto* loop_flow = FindSingleValue<ir::Loop>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
     EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(3u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
   %b1 = block {
     loop [s: %b2, c: %b3, m: %b4]
       %b2 = block {
-        br %b5  # return
+        ret
       }
       # Continuing block
       %b3 = block {
-        if true [t: %b6, f: %b7, m: %b8]
+        if true [t: %b5, f: %b6, m: %b7]
           # True block
-          %b6 = block {
+          %b5 = block {
             br %b4
           }
 
           # False block
-          %b7 = block {
+          %b6 = block {
-            br %b8
+            br %b7
           }
 
         # Merge block
-        %b8 = block {
+        %b7 = block {
           br %b2
         }
 
@@ -741,19 +697,19 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) {
 
     # Merge block
     %b4 = block {
-      if true [t: %b9, f: %b10, m: %b11]
+      if true [t: %b8, f: %b9, m: %b10]
         # True block
-        %b9 = block {
+        %b8 = block {
-          br %b5  # return
+          ret
         }
         # False block
-        %b10 = block {
+        %b9 = block {
-          br %b11
+          br %b10
         }
 
       # Merge block
-      %b11 = block {
+      %b10 = block {
-        br %b5  # return
+        ret
       }
 
     }
@@ -761,7 +717,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) {
 
   }
 
-  %b5 = func_terminator
 }
 )");
 }
@@ -779,7 +734,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) {
     auto* if_flow = FindSingleValue<ir::If>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
     EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
@@ -787,7 +741,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) {
     EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -815,12 +768,11 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) {
 
     # Merge block
     %b4 = block {
-      br %b7  # return
+      ret
     }
 
   }
 
-  %b7 = func_terminator
 }
 )");
 }
@@ -966,12 +918,11 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
 
     # Merge block
     %b4 = block {
-      br %b26  # return
+      ret
     }
 
   }
 
-  %b26 = func_terminator
 }
 )");
 }
@@ -991,9 +942,7 @@ TEST_F(IR_BuilderImplTest, While) {
     auto* if_flow = flow->Start()->Branch()->As<ir::If>();
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
-    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
     EXPECT_EQ(2u, flow->Start()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
@@ -1032,12 +981,11 @@ TEST_F(IR_BuilderImplTest, While) {
 
     # Merge block
     %b4 = block {
-      br %b8  # return
+      ret
     }
 
   }
 
-  %b8 = func_terminator
 }
 )");
 }
@@ -1057,9 +1005,7 @@ TEST_F(IR_BuilderImplTest, While_Return) {
     auto* if_flow = flow->Start()->Branch()->As<ir::If>();
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
-    EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
     EXPECT_EQ(2u, flow->Start()->InboundBranches().Length());
     EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
@@ -1085,7 +1031,7 @@ TEST_F(IR_BuilderImplTest, While_Return) {
 
         # Merge block
         %b7 = block {
-          br %b8  # return
+          ret
         }
 
       }
@@ -1097,12 +1043,11 @@ TEST_F(IR_BuilderImplTest, While_Return) {
 
     # Merge block
     %b4 = block {
-      br %b8  # return
+      ret
     }
 
   }
 
-  %b8 = func_terminator
 }
 )");
 }
@@ -1135,9 +1080,7 @@ TEST_F(IR_BuilderImplTest, DISABLED_For) {
     auto* if_flow = flow->Start()->Branch()->As<ir::If>();
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
-    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
     EXPECT_EQ(2u, flow->Start()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
@@ -1159,12 +1102,10 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) {
     auto* flow = FindSingleValue<ir::Loop>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     EXPECT_EQ(2u, flow->Start()->InboundBranches().Length());
     EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1181,12 +1122,11 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) {
 
     # Merge block
     %b4 = block {
-      br %b5  # return
+      ret
     }
 
   }
 
-  %b5 = func_terminator
 }
 )");
 }
@@ -1205,7 +1145,6 @@ TEST_F(IR_BuilderImplTest, Switch) {
     auto* flow = FindSingleValue<ir::Switch>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     auto cases = flow->Cases();
     ASSERT_EQ(3u, cases.Length());
@@ -1227,7 +1166,6 @@ TEST_F(IR_BuilderImplTest, Switch) {
     EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length());
     EXPECT_EQ(1u, cases[2].Start()->InboundBranches().Length());
     EXPECT_EQ(4u, flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1250,12 +1188,11 @@ TEST_F(IR_BuilderImplTest, Switch) {
 
     # Merge block
     %b5 = block {
-      br %b6  # return
+      ret
     }
 
   }
 
-  %b6 = func_terminator
 }
 )");
 }
@@ -1275,7 +1212,6 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) {
     auto* flow = FindSingleValue<ir::Switch>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     auto cases = flow->Cases();
     ASSERT_EQ(1u, cases.Length());
@@ -1292,7 +1228,6 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) {
 
     EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
     EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1305,12 +1240,11 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) {
 
     # Merge block
     %b3 = block {
-      br %b4  # return
+      ret
     }
 
   }
 
-  %b4 = func_terminator
 }
 )");
 }
@@ -1326,7 +1260,6 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) {
     auto* flow = FindSingleValue<ir::Switch>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     auto cases = flow->Cases();
     ASSERT_EQ(1u, cases.Length());
@@ -1335,7 +1268,6 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) {
 
     EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
     EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1348,12 +1280,11 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) {
 
     # Merge block
     %b3 = block {
-      br %b4  # return
+      ret
     }
 
   }
 
-  %b4 = func_terminator
 }
 )");
 }
@@ -1371,7 +1302,6 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) {
     auto* flow = FindSingleValue<ir::Switch>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     auto cases = flow->Cases();
     ASSERT_EQ(2u, cases.Length());
@@ -1387,7 +1317,6 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) {
     EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length());
     EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length());
     // This is 1 because the if is dead-code eliminated and the return doesn't happen.
-    EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1405,12 +1334,11 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) {
 
     # Merge block
     %b4 = block {
-      br %b5  # return
+      ret
     }
 
   }
 
-  %b5 = func_terminator
 }
 )");
 }
@@ -1431,7 +1359,6 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) {
     auto* flow = FindSingleValue<ir::Switch>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
-    auto* func = m.functions[0];
 
     auto cases = flow->Cases();
     ASSERT_EQ(2u, cases.Length());
@@ -1446,7 +1373,6 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) {
     EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
     EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1454,16 +1380,15 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) {
     switch 1i [c: (0i, %b2), c: (default, %b3)]
       # Case block
       %b2 = block {
-        br %b4  # return
+        ret
       }
       # Case block
       %b3 = block {
-        br %b4  # return
+        ret
       }
 
   }
 
-  %b4 = func_terminator
 }
 )");
 }
@@ -1478,16 +1403,14 @@ TEST_F(IR_BuilderImplTest, Emit_Phony) {
     EXPECT_EQ(Disassemble(m.Get()),
               R"(%b = func():i32 -> %b1 {
   %b1 = block {
-    br %b2 1i  # return
+    ret 1i
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:i32 = call %b
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }

src/tint/ir/from_program_unary_test.cc

@@ -36,17 +36,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Not) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():bool -> %b1 {
   %b1 = block {
-    br %b2 false  # return
+    ret false
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:bool = call %my_func
     %tint_symbol:bool = eq %3, false
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -61,17 +59,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Complement) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 {
   %b1 = block {
-    br %b2 1u  # return
+    ret 1u
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:u32 = call %my_func
     %tint_symbol:u32 = complement %3
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -86,17 +82,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Negation) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():i32 -> %b1 {
   %b1 = block {
-    br %b2 1i  # return
+    ret 1i
   }
-  %b2 = func_terminator
 }
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
-  %b3 = block {
+  %b2 = block {
     %3:i32 = call %my_func
     %tint_symbol:i32 = negation %3
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -120,9 +114,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_AddressOf) {
 
 %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
   %b3 = block {
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }
@@ -149,9 +142,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Indirection) {
 %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
   %b3 = block {
     store %v3, 42i
-    br %b4  # return
+    ret
   }
-  %b4 = func_terminator
 }
 )");
 }

src/tint/ir/from_program_var_test.cc

@@ -72,9 +72,8 @@ TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) {
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
   %b1 = block {
     %a:ptr<function, u32, read_write> = var
-    br %b2  # return
+    ret
   }
-  %b2 = func_terminator
 }
 )");
 }
@@ -91,9 +90,8 @@ TEST_F(IR_BuilderImplTest, Emit_Var_Init) {
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
   %b1 = block {
     %a:ptr<function, u32, read_write> = var, 2u
-    br %b2  # return
+    ret
   }
-  %b2 = func_terminator
 }
 )");
 }

src/tint/ir/function.h

@@ -117,12 +117,6 @@ class Function : public utils::Castable<Function, Value> {
     /// @returns the function start target
     Block* StartTarget() const { return start_target_; }
 
-    /// Sets the end target for the function
-    /// @param target the end target
-    void SetEndTarget(FunctionTerminator* target) { end_target_ = target; }
-    /// @returns the function end target
-    FunctionTerminator* EndTarget() const { return end_target_; }
-
   private:
     const type::Type* return_type_;
     PipelineStage pipeline_stage_;
@@ -134,7 +128,6 @@ class Function : public utils::Castable<Function, Value> {
     utils::Vector<FunctionParam*, 1> params_;
 
     Block* start_target_ = nullptr;
-    FunctionTerminator* end_target_ = nullptr;
 };
 
 utils::StringStream& operator<<(utils::StringStream& out, Function::PipelineStage value);

src/tint/ir/function_terminator.h

@@ -1,33 +0,0 @@
-// Copyright 2022 The Tint Authors.
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef SRC_TINT_IR_FUNCTION_TERMINATOR_H_
-#define SRC_TINT_IR_FUNCTION_TERMINATOR_H_
-
-#include "src/tint/ir/block.h"
-
-namespace tint::ir {
-
-/// Block used as the end of a function. Must only be used as the `end_target` in a function. There
-/// are no instructions in this block.
-class FunctionTerminator : public utils::Castable<FunctionTerminator, Block> {
-  public:
-    /// Constructor
-    FunctionTerminator();
-    ~FunctionTerminator() override;
-};
-
-}  // namespace tint::ir
-
-#endif  // SRC_TINT_IR_FUNCTION_TERMINATOR_H_

src/tint/ir/return.cc

@@ -1,4 +1,4 @@
-// Copyright 2022 The Tint Authors.
+// Copyright 2023 The Tint Authors.
 //
 // Licensed under the Apache License, Version 2.0 (the "License");
 // you may not use this file except in compliance with the License.
@@ -12,14 +12,19 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "src/tint/ir/function_terminator.h"
+#include "src/tint/ir/return.h"
 
-TINT_INSTANTIATE_TYPEINFO(tint::ir::FunctionTerminator);
+#include "src/tint/ir/function.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::Return);
 
 namespace tint::ir {
 
-FunctionTerminator::FunctionTerminator() : Base() {}
+Return::Return(Function* func, utils::VectorRef<Value*> args) : Base(args), func_(func) {
+    TINT_ASSERT(IR, func_);
+    func_->AddUsage(this);
+}
 
-FunctionTerminator::~FunctionTerminator() = default;
+Return::~Return() = default;
 
 }  // namespace tint::ir

src/tint/ir/return.h

@@ -0,0 +1,46 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_TINT_IR_RETURN_H_
+#define SRC_TINT_IR_RETURN_H_
+
+#include "src/tint/ir/branch.h"
+#include "src/tint/utils/castable.h"
+
+// Forward declarations
+namespace tint::ir {
+class Function;
+}  // namespace tint::ir
+
+namespace tint::ir {
+
+/// A return instruction.
+class Return : public utils::Castable<Return, Branch> {
+  public:
+    /// Constructor
+    /// @param func the function being returned
+    /// @param args the branch arguments
+    explicit Return(Function* func, utils::VectorRef<Value*> args = {});
+    ~Return() override;
+
+    /// @returns the function being returned
+    const Function* Func() const { return func_; }
+
+  private:
+    Function* func_ = nullptr;
+};
+
+}  // namespace tint::ir
+
+#endif  // SRC_TINT_IR_RETURN_H_

src/tint/ir/to_program.cc

@@ -20,11 +20,11 @@
 #include "src/tint/ir/block.h"
 #include "src/tint/ir/call.h"
 #include "src/tint/ir/constant.h"
-#include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/if.h"
 #include "src/tint/ir/instruction.h"
 #include "src/tint/ir/load.h"
 #include "src/tint/ir/module.h"
+#include "src/tint/ir/return.h"
 #include "src/tint/ir/store.h"
 #include "src/tint/ir/switch.h"
 #include "src/tint/ir/user_call.h"
@@ -126,8 +126,6 @@ class State {
             Status status = tint::Switch(
                 block,
 
-                [&](const ir::FunctionTerminator*) { return kStop; },
-
                 [&](const ir::Block* blk) {
                     for (auto* inst : blk->Instructions()) {
                         auto stmt = Stmt(inst);
@@ -240,9 +238,9 @@ class State {
         return b.Switch(cond, std::move(cases));
     }
 
-    utils::Result<const ast::ReturnStatement*> FunctionTerminator(const ir::Branch* branch) {
+    utils::Result<const ast::ReturnStatement*> Return(const ir::Return* ret) {
-        if (branch->Args().IsEmpty()) {
+        if (ret->Args().IsEmpty()) {
-            // Branch to function terminator has no arguments.
+            // Return has no arguments.
             // If this block is nested withing some control flow, then we must
             // emit a 'return' statement, otherwise we've just naturally reached
             // the end of the function where the 'return' is redundant.
@@ -252,16 +250,14 @@ class State {
             return nullptr;
         }
 
-        // Branch to function terminator has arguments - this is the return
+        // Return has arguments - this is the return value.
-        // value.
+        if (ret->Args().Length() != 1) {
-        if (branch->Args().Length() != 1) {
+            TINT_ICE(IR, b.Diagnostics())
-            TINT_ICE(IR, b.Diagnostics()) << "expected 1 value for function "
+                << "expected 1 value for return, got " << ret->Args().Length();
-                                             "terminator (return value), got "
-                                          << branch->Args().Length();
             return utils::Failure;
         }
 
-        auto* val = Expr(branch->Args().Front());
+        auto* val = Expr(ret->Args().Front());
         if (TINT_UNLIKELY(!val)) {
             return utils::Failure;
         }
@@ -275,10 +271,8 @@ class State {
             return true;
         }
         if (auto* br = node->Instructions().Front()->As<Branch>()) {
-            return br->To() == stop_at;
+            return !br->Is<ir::Return>() && br->To() == stop_at;
         }
-        // TODO(dsinclair): This should possibly walk over Jump instructions that
-        // just jump to empty blocks if we want to be comprehensive.
         return false;
     }
 
@@ -291,12 +285,9 @@ class State {
             [&](const ir::Store* i) { return Store(i); },  //
             [&](const ir::If* if_) { return If(if_); },
             [&](const ir::Switch* switch_) { return Switch(switch_); },
-            [&](const ir::Branch* branch) {
+            [&](const ir::Return* ret) { return Return(ret); },
-                if (branch->To()->Is<ir::FunctionTerminator>()) {
+            // TODO(dsinclair): Remove when branch is only a parent ...
-                    return utils::Result<const ast::Statement*>{FunctionTerminator(branch)};
+            [&](const ir::Branch*) { return utils::Result<const ast::Statement*>{nullptr}; },
-                }
-                return utils::Result<const ast::Statement*>{nullptr};
-            },
             [&](Default) {
                 UNHANDLED_CASE(inst);
                 return utils::Failure;

src/tint/ir/transform/add_empty_entry_point.cc

@@ -37,7 +37,7 @@ void AddEmptyEntryPoint::Run(ir::Module* ir, const DataMap&, DataMap&) const {
     ir::Builder builder(*ir);
     auto* ep = builder.CreateFunction("unused_entry_point", ir->Types().void_(),
                                       Function::PipelineStage::kCompute, std::array{1u, 1u, 1u});
-    ep->StartTarget()->SetInstructions(utils::Vector{builder.Branch(ep->EndTarget())});
+    ep->StartTarget()->SetInstructions(utils::Vector{builder.Return(ep)});
     ir->functions.Push(ep);
 }
 

src/tint/ir/transform/add_empty_entry_point_test.cc

@@ -27,9 +27,8 @@ TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) {
     auto* expect = R"(
 %unused_entry_point = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
   %b1 = block {
-    br %b2  # return
+    ret
   }
-  %b2 = func_terminator
 }
 )";
 
@@ -40,15 +39,14 @@ TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) {
 
 TEST_F(IR_AddEmptyEntryPointTest, ExistingEntryPoint) {
     auto* ep = b.CreateFunction("main", mod.Types().void_(), Function::PipelineStage::kFragment);
-    ep->StartTarget()->SetInstructions(utils::Vector{b.Branch(ep->EndTarget())});
+    ep->StartTarget()->SetInstructions(utils::Vector{b.Return(ep)});
     mod.functions.Push(ep);
 
     auto* expect = R"(
 %main = func():void [@fragment] -> %b1 {
   %b1 = block {
-    br %b2  # return
+    ret
   }
-  %b2 = func_terminator
 }
 )";
 

src/tint/transform/manager_test.cc

@@ -51,7 +51,7 @@ class IR_AddFunction final : public ir::transform::Transform {
     void Run(ir::Module* mod, const DataMap&, DataMap&) const override {
         ir::Builder builder(*mod);
         auto* func = builder.CreateFunction("ir_func", mod->Types().Get<type::Void>());
-        func->StartTarget()->SetInstructions(utils::Vector{builder.Branch(func->EndTarget())});
+        func->StartTarget()->SetInstructions(utils::Vector{builder.Return(func)});
         mod->functions.Push(func);
     }
 };
@@ -68,7 +68,7 @@ ir::Module MakeIR() {
     ir::Module mod;
     ir::Builder builder(mod);
     auto* func = builder.CreateFunction("main", mod.Types().Get<type::Void>());
-    func->StartTarget()->SetInstructions(utils::Vector{builder.Branch(func->EndTarget())});
+    func->StartTarget()->SetInstructions(utils::Vector{builder.Return(func)});
     builder.ir.functions.Push(func);
     return mod;
 }

src/tint/writer/spirv/ir/generator_impl_ir.cc

@@ -19,10 +19,10 @@
 #include "spirv/unified1/spirv.h"
 #include "src/tint/ir/binary.h"
 #include "src/tint/ir/block.h"
-#include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/if.h"
 #include "src/tint/ir/load.h"
 #include "src/tint/ir/module.h"
+#include "src/tint/ir/return.h"
 #include "src/tint/ir/store.h"
 #include "src/tint/ir/transform/add_empty_entry_point.h"
 #include "src/tint/ir/user_call.h"
@@ -354,9 +354,7 @@ void GeneratorImplIr::EmitBlock(const ir::Block* block) {
 }
 
 void GeneratorImplIr::EmitBranch(const ir::Branch* b) {
-    Switch(
+    if (b->Is<ir::Return>()) {
-        b->To(),
-        [&](const ir::FunctionTerminator*) {
         if (!b->Args().IsEmpty()) {
             TINT_ASSERT(Writer, b->Args().Length() == 1u);
             OperandList operands;
@@ -365,7 +363,11 @@ void GeneratorImplIr::EmitBranch(const ir::Branch* b) {
         } else {
             current_function_.push_inst(spv::Op::OpReturn, {});
         }
-        },
+        return;
+    }
+
+    Switch(
+        b->To(),
         [&](const ir::Block* blk) { current_function_.push_inst(spv::Op::OpBranch, {Label(blk)}); },
         [&](Default) {
             // A block may not have an outward branch (e.g. an unreachable merge

src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc

@@ -125,7 +125,7 @@ TEST_P(Arithmetic, Scalar) {
     func->StartTarget()->SetInstructions(
         utils::Vector{b.CreateBinary(params.kind, MakeScalarType(params.type),
                                      MakeScalarValue(params.type), MakeScalarValue(params.type)),
-                      b.Branch(func->EndTarget())});
+                      b.Return(func)});
 
     generator_.EmitFunction(func);
     EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst));
@@ -138,7 +138,7 @@ TEST_P(Arithmetic, Vector) {
         utils::Vector{b.CreateBinary(params.kind, MakeVectorType(params.type),
                                      MakeVectorValue(params.type), MakeVectorValue(params.type)),
 
-                      b.Branch(func->EndTarget())});
+                      b.Return(func)});
 
     generator_.EmitFunction(func);
     EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst));
@@ -172,7 +172,7 @@ TEST_P(Bitwise, Scalar) {
     func->StartTarget()->SetInstructions(
         utils::Vector{b.CreateBinary(params.kind, MakeScalarType(params.type),
                                      MakeScalarValue(params.type), MakeScalarValue(params.type)),
-                      b.Branch(func->EndTarget())});
+                      b.Return(func)});
 
     generator_.EmitFunction(func);
     EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst));
@@ -185,7 +185,7 @@ TEST_P(Bitwise, Vector) {
         utils::Vector{b.CreateBinary(params.kind, MakeVectorType(params.type),
                                      MakeVectorValue(params.type), MakeVectorValue(params.type)),
 
-                      b.Branch(func->EndTarget())});
+                      b.Return(func)});
 
     generator_.EmitFunction(func);
     EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst));
@@ -211,7 +211,7 @@ TEST_P(Comparison, Scalar) {
     func->StartTarget()->SetInstructions(
         utils::Vector{b.CreateBinary(params.kind, mod.Types().bool_(), MakeScalarValue(params.type),
                                      MakeScalarValue(params.type)),
-                      b.Branch(func->EndTarget())});
+                      b.Return(func)});
 
     generator_.EmitFunction(func);
     EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst));
@@ -224,7 +224,7 @@ TEST_P(Comparison, Vector) {
         utils::Vector{b.CreateBinary(params.kind, mod.Types().vec2(mod.Types().bool_()),
                                      MakeVectorValue(params.type), MakeVectorValue(params.type)),
 
-                      b.Branch(func->EndTarget())});
+                      b.Return(func)});
 
     generator_.EmitFunction(func);
     EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst));
@@ -279,7 +279,7 @@ TEST_F(SpvGeneratorImplTest, Binary_Chain) {
     auto* func = b.CreateFunction("foo", mod.Types().void_());
     auto* a = b.Subtract(mod.Types().i32(), b.Constant(1_i), b.Constant(2_i));
     func->StartTarget()->SetInstructions(
-        utils::Vector{a, b.Add(mod.Types().i32(), a, a), b.Branch(func->EndTarget())});
+        utils::Vector{a, b.Add(mod.Types().i32(), a, a), b.Return(func)});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"

src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc

@@ -19,7 +19,7 @@ namespace {
 
 TEST_F(SpvGeneratorImplTest, Function_Empty) {
     auto* func = b.CreateFunction("foo", mod.Types().void_());
-    func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
+    func->StartTarget()->SetInstructions(utils::Vector{b.Return(func)});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -35,7 +35,7 @@ OpFunctionEnd
 // Test that we do not emit the same function type more than once.
 TEST_F(SpvGeneratorImplTest, Function_DeduplicateType) {
     auto* func = b.CreateFunction("foo", mod.Types().void_());
-    func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
+    func->StartTarget()->SetInstructions(utils::Vector{b.Return(func)});
 
     generator_.EmitFunction(func);
     generator_.EmitFunction(func);
@@ -48,7 +48,7 @@ TEST_F(SpvGeneratorImplTest, Function_DeduplicateType) {
 TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Compute) {
     auto* func = b.CreateFunction("main", mod.Types().void_(),
                                   ir::Function::PipelineStage::kCompute, {{32, 4, 1}});
-    func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
+    func->StartTarget()->SetInstructions(utils::Vector{b.Return(func)});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint GLCompute %1 "main"
@@ -66,7 +66,7 @@ OpFunctionEnd
 TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Fragment) {
     auto* func =
         b.CreateFunction("main", mod.Types().void_(), ir::Function::PipelineStage::kFragment);
-    func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
+    func->StartTarget()->SetInstructions(utils::Vector{b.Return(func)});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint Fragment %1 "main"
@@ -84,7 +84,7 @@ OpFunctionEnd
 TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Vertex) {
     auto* func =
         b.CreateFunction("main", mod.Types().void_(), ir::Function::PipelineStage::kVertex);
-    func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
+    func->StartTarget()->SetInstructions(utils::Vector{b.Return(func)});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint Vertex %1 "main"
@@ -101,15 +101,15 @@ OpFunctionEnd
 TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Multiple) {
     auto* f1 = b.CreateFunction("main1", mod.Types().void_(), ir::Function::PipelineStage::kCompute,
                                 {{32, 4, 1}});
-    f1->StartTarget()->SetInstructions(utils::Vector{b.Branch(f1->EndTarget())});
+    f1->StartTarget()->SetInstructions(utils::Vector{b.Return(f1)});
 
     auto* f2 = b.CreateFunction("main2", mod.Types().void_(), ir::Function::PipelineStage::kCompute,
                                 {{8, 2, 16}});
-    f2->StartTarget()->SetInstructions(utils::Vector{b.Branch(f2->EndTarget())});
+    f2->StartTarget()->SetInstructions(utils::Vector{b.Return(f2)});
 
     auto* f3 =
         b.CreateFunction("main3", mod.Types().void_(), ir::Function::PipelineStage::kFragment);
-    f3->StartTarget()->SetInstructions(utils::Vector{b.Branch(f3->EndTarget())});
+    f3->StartTarget()->SetInstructions(utils::Vector{b.Return(f3)});
 
     generator_.EmitFunction(f1);
     generator_.EmitFunction(f2);
@@ -143,7 +143,7 @@ OpFunctionEnd
 TEST_F(SpvGeneratorImplTest, Function_ReturnValue) {
     auto* func = b.CreateFunction("foo", mod.Types().i32());
     func->StartTarget()->SetInstructions(
-        utils::Vector{b.Branch(func->EndTarget(), utils::Vector{b.Constant(i32(42))})});
+        utils::Vector{b.Return(func, utils::Vector{b.Constant(i32(42))})});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -165,7 +165,7 @@ TEST_F(SpvGeneratorImplTest, Function_Parameters) {
     auto* func = b.CreateFunction("foo", i32);
     func->SetParams(utils::Vector{x, y});
     func->StartTarget()->SetInstructions(
-        utils::Vector{result, b.Branch(func->EndTarget(), utils::Vector{result})});
+        utils::Vector{result, b.Return(func, utils::Vector{result})});
     mod.SetName(x, "x");
     mod.SetName(y, "y");
 
@@ -193,12 +193,12 @@ TEST_F(SpvGeneratorImplTest, Function_Call) {
     auto* foo = b.CreateFunction("foo", i32_ty);
     foo->SetParams(utils::Vector{x, y});
     foo->StartTarget()->SetInstructions(
-        utils::Vector{result, b.Branch(foo->EndTarget(), utils::Vector{result})});
+        utils::Vector{result, b.Return(foo, utils::Vector{result})});
 
     auto* bar = b.CreateFunction("bar", mod.Types().void_());
     bar->StartTarget()->SetInstructions(utils::Vector{
         b.UserCall(i32_ty, foo, utils::Vector{b.Constant(i32(2)), b.Constant(i32(3))}),
-        b.Branch(bar->EndTarget())});
+        b.Return(bar)});
 
     generator_.EmitFunction(foo);
     generator_.EmitFunction(bar);
@@ -227,11 +227,11 @@ OpFunctionEnd
 
 TEST_F(SpvGeneratorImplTest, Function_Call_Void) {
     auto* foo = b.CreateFunction("foo", mod.Types().void_());
-    foo->StartTarget()->SetInstructions(utils::Vector{b.Branch(foo->EndTarget())});
+    foo->StartTarget()->SetInstructions(utils::Vector{b.Return(foo)});
 
     auto* bar = b.CreateFunction("bar", mod.Types().void_());
-    bar->StartTarget()->SetInstructions(utils::Vector{
+    bar->StartTarget()->SetInstructions(
-        b.UserCall(mod.Types().void_(), foo, utils::Empty), b.Branch(bar->EndTarget())});
+        utils::Vector{b.UserCall(mod.Types().void_(), foo, utils::Empty), b.Return(bar)});
 
     generator_.EmitFunction(foo);
     generator_.EmitFunction(bar);

src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc

@@ -25,7 +25,7 @@ TEST_F(SpvGeneratorImplTest, If_TrueEmpty_FalseEmpty) {
     auto* i = b.CreateIf(b.Constant(true));
     i->True()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
     i->False()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
-    i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
+    i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
 
     func->StartTarget()->SetInstructions(utils::Vector{i});
 
@@ -50,7 +50,7 @@ TEST_F(SpvGeneratorImplTest, If_FalseEmpty) {
 
     auto* i = b.CreateIf(b.Constant(true));
     i->False()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
-    i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
+    i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
 
     auto* true_block = i->True();
     true_block->SetInstructions(utils::Vector{
@@ -84,7 +84,7 @@ TEST_F(SpvGeneratorImplTest, If_TrueEmpty) {
 
     auto* i = b.CreateIf(b.Constant(true));
     i->True()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
-    i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
+    i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
 
     auto* false_block = i->False();
     false_block->SetInstructions(utils::Vector{
@@ -117,8 +117,8 @@ TEST_F(SpvGeneratorImplTest, If_BothBranchesReturn) {
     auto* func = b.CreateFunction("foo", mod.Types().void_());
 
     auto* i = b.CreateIf(b.Constant(true));
-    i->True()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
+    i->True()->SetInstructions(utils::Vector{b.Return(func)});
-    i->False()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
+    i->False()->SetInstructions(utils::Vector{b.Return(func)});
 
     func->StartTarget()->SetInstructions(utils::Vector{i});
 

src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc

@@ -25,7 +25,7 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_NoInit) {
 
     auto* ty = mod.Types().Get<type::Pointer>(mod.Types().i32(), builtin::AddressSpace::kFunction,
                                               builtin::Access::kReadWrite);
-    func->StartTarget()->SetInstructions(utils::Vector{b.Declare(ty), b.Branch(func->EndTarget())});
+    func->StartTarget()->SetInstructions(utils::Vector{b.Declare(ty), b.Return(func)});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -49,7 +49,7 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_WithInit) {
     auto* v = b.Declare(ty);
     v->SetInitializer(b.Constant(42_i));
 
-    func->StartTarget()->SetInstructions(utils::Vector{v, b.Branch(func->EndTarget())});
+    func->StartTarget()->SetInstructions(utils::Vector{v, b.Return(func)});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -73,7 +73,7 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_Name) {
     auto* ty = mod.Types().Get<type::Pointer>(mod.Types().i32(), builtin::AddressSpace::kFunction,
                                               builtin::Access::kReadWrite);
     auto* v = b.Declare(ty);
-    func->StartTarget()->SetInstructions(utils::Vector{v, b.Branch(func->EndTarget())});
+    func->StartTarget()->SetInstructions(utils::Vector{v, b.Return(func)});
     mod.SetName(v, "myvar");
 
     generator_.EmitFunction(func);
@@ -101,8 +101,8 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_DeclInsideBlock) {
 
     auto* i = b.CreateIf(b.Constant(true));
     i->True()->SetInstructions(utils::Vector{v, b.Branch(i->Merge())});
-    i->False()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
+    i->False()->SetInstructions(utils::Vector{b.Return(func)});
-    i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
+    i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
 
     func->StartTarget()->SetInstructions(utils::Vector{i});
 
@@ -138,7 +138,7 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_Load) {
     auto* ty = mod.Types().Get<type::Pointer>(store_ty, builtin::AddressSpace::kFunction,
                                               builtin::Access::kReadWrite);
     auto* v = b.Declare(ty);
-    func->StartTarget()->SetInstructions(utils::Vector{v, b.Load(v), b.Branch(func->EndTarget())});
+    func->StartTarget()->SetInstructions(utils::Vector{v, b.Load(v), b.Return(func)});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -162,7 +162,7 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_Store) {
                                               builtin::Access::kReadWrite);
     auto* v = b.Declare(ty);
     func->StartTarget()->SetInstructions(
-        utils::Vector{v, b.Store(v, b.Constant(42_i)), b.Branch(func->EndTarget())});
+        utils::Vector{v, b.Store(v, b.Constant(42_i)), b.Return(func)});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"