[ir] Dissolve the flow graph

The ir::Value objects each have a list of instructions in which their
used. These lists allow us to determine all the places the value is
used. Currently this is unable to track the usage of a value in an
`if` or `switch` condition. It is also unable to track the usage of a
value as a branch argument.

In order to facilitate this tracking, the flow graph has been resolved.
Branches are moved to branch instructions (and jump instructions). A
jump is walk continue branch. A branch is a walk terminating branch. The
`if`, `switch` and `loop` flow nodes are moved to instructions as well.

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

src/tint/BUILD.gn

@@ -1210,6 +1210,8 @@ if (tint_build_ir) {
       "ir/block.h",
       "ir/block_param.cc",
       "ir/block_param.h",
+      "ir/branch.cc",
+      "ir/branch.h",
       "ir/builder.cc",
       "ir/builder.h",
       "ir/builtin.cc",
@@ -1240,6 +1242,8 @@ if (tint_build_ir) {
       "ir/if.h",
       "ir/instruction.cc",
       "ir/instruction.h",
+      "ir/jump.cc",
+      "ir/jump.h",
       "ir/load.cc",
       "ir/load.h",
       "ir/loop.cc",

src/tint/CMakeLists.txt

@@ -718,6 +718,8 @@ if(${TINT_BUILD_IR})
     ir/block.h
     ir/block_param.cc
     ir/block_param.h
+    ir/branch.cc
+    ir/branch.h
     ir/builder.cc
     ir/builder.h
     ir/builtin.cc
@@ -750,6 +752,8 @@ if(${TINT_BUILD_IR})
     ir/if.h
     ir/instruction.cc
     ir/instruction.h
+    ir/jump.cc
+    ir/jump.h
     ir/load.cc
     ir/load.h
     ir/loop.cc

src/tint/ir/block.cc

@@ -22,13 +22,4 @@ Block::Block() : Base() {}
 
 Block::~Block() = default;
 
-void Block::BranchTo(FlowNode* to, utils::VectorRef<Value*> args) {
-    TINT_ASSERT(IR, to);
-    branch_.target = to;
-    branch_.args = args;
-    if (to) {
-        to->AddInboundBranch(this);
-    }
-}
-
 }  // namespace tint::ir

src/tint/ir/block.h

@@ -34,16 +34,30 @@ class Block : public utils::Castable<Block, FlowNode> {
     Block();
     ~Block() override;
 
-    /// Sets the blocks branch target to the given node.
-    /// @param to the node to branch too
-    /// @param args the branch arguments
-    void BranchTo(FlowNode* to, utils::VectorRef<Value*> args = {});
-
     /// @returns true if this is block has a branch target set
-    bool HasBranchTarget() const override { return branch_.target != nullptr; }
+    bool HasBranchTarget() const override {
+        return !instructions_.IsEmpty() && instructions_.Back()->Is<ir::Branch>();
+    }
 
-    /// @return the node this block branches too.
+    /// @return the node this block branches to or nullptr if the block doesn't branch
-    const ir::Branch& Branch() const { return branch_; }
+    const ir::Branch* Branch() const {
+        if (!HasBranchTarget()) {
+            return nullptr;
+        }
+        return instructions_.Back()->As<ir::Branch>();
+    }
+
+    /// @param target the block to see if we trampoline too
+    /// @returns if this block just branches to the provided target.
+    bool IsTrampoline(const FlowNode* target) const {
+        if (instructions_.Length() != 1) {
+            return false;
+        }
+        if (auto* inst = instructions_.Front()->As<ir::Branch>()) {
+            return inst->To() == target;
+        }
+        return false;
+    }
 
     /// Sets the instructions in the block
     /// @param instructions the instructions to set
@@ -59,14 +73,12 @@ class Block : public utils::Castable<Block, FlowNode> {
     /// Sets the params to the block
     /// @param params the params for the block
     void SetParams(utils::VectorRef<const BlockParam*> params) { params_ = std::move(params); }
+    /// @return the parameters passed into the block
+    utils::VectorRef<const BlockParam*> Params() const { return params_; }
     /// @returns the params to the block
     utils::Vector<const BlockParam*, 0>& Params() { return params_; }
 
-    /// @return the parameters passed into the block
-    utils::VectorRef<const BlockParam*> Params() const { return params_; }
-
   private:
-    ir::Branch branch_ = {};
     utils::Vector<const Instruction*, 16> instructions_;
     utils::Vector<const BlockParam*, 0> params_;
 };

src/tint/ir/branch.cc

@@ -0,0 +1,35 @@
+// 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.
+
+#include "src/tint/ir/branch.h"
+
+#include <utility>
+
+#include "src/tint/ir/flow_node.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::Branch);
+
+namespace tint::ir {
+
+Branch::Branch(FlowNode* to, utils::VectorRef<Value*> args) : to_(to), args_(std::move(args)) {
+    TINT_ASSERT(IR, to_);
+    to_->AddInboundBranch(this);
+    for (auto* arg : args) {
+        arg->AddUsage(this);
+    }
+}
+
+Branch::~Branch() = default;
+
+}  // namespace tint::ir

src/tint/ir/branch.h

@@ -15,20 +15,35 @@
 #ifndef SRC_TINT_IR_BRANCH_H_
 #define SRC_TINT_IR_BRANCH_H_
 
-#include "src/tint/ir/flow_node.h"
+#include "src/tint/ir/instruction.h"
 #include "src/tint/ir/value.h"
+#include "src/tint/utils/castable.h"
+
+// Forward declarations
+namespace tint::ir {
+class FlowNode;
+}  // namespace tint::ir
 
 namespace tint::ir {
 
-/// A information on a branch to another block
+/// A branch instruction. A branch is a walk terminating jump.
-struct Branch {
+class Branch : public utils::Castable<Branch, Instruction> {
-    /// The block being branched too.
+  public:
-    FlowNode* target = nullptr;
+    /// Constructor
+    /// @param to the block to branch too
+    /// @param args the branch arguments
+    explicit Branch(FlowNode* to, utils::VectorRef<Value*> args = {});
+    ~Branch() override;
 
-    /// The arguments provided for that branch. These arguments could be the
+    /// @returns the block being branched too.
-    /// return value in the case of a branch to the function terminator, or they could
+    const FlowNode* To() const { return to_; }
-    /// be the basic block arguments passed into the block.
+
-    utils::Vector<Value*, 2> args;
+    /// @returns the branch arguments
+    utils::VectorRef<Value*> Args() const { return args_; }
+
+  private:
+    FlowNode* to_;
+    utils::Vector<Value*, 2> args_;
 };
 
 }  // namespace tint::ir

src/tint/ir/builder.cc

@@ -29,10 +29,6 @@ Builder::~Builder() = default;
 ir::Block* Builder::CreateRootBlockIfNeeded() {
     if (!ir.root_block) {
         ir.root_block = CreateBlock();
-
-        // Everything in the module scope must have been const-eval's, so everything will go into a
-        // single block. So, we can create the root terminator for the root-block now.
-        ir.root_block->BranchTo(CreateRootTerminator());
     }
     return ir.root_block;
 }
@@ -59,50 +55,26 @@ Function* Builder::CreateFunction(Symbol name,
     ir_func->SetStartTarget(CreateBlock());
     ir_func->SetEndTarget(CreateFunctionTerminator());
 
-    // Function is always branching into the Start().target
-    ir_func->StartTarget()->AddInboundBranch(ir_func);
-
     return ir_func;
 }
 
 If* Builder::CreateIf(Value* condition) {
     TINT_ASSERT(IR, condition);
-
+    return ir.values.Create<If>(condition, CreateBlock(), CreateBlock(), CreateBlock());
-    auto* ir_if = ir.flow_nodes.Create<If>(condition);
-    ir_if->True().target = CreateBlock();
-    ir_if->False().target = CreateBlock();
-    ir_if->Merge().target = CreateBlock();
-
-    // An if always branches to both the true and false block.
-    ir_if->True().target->AddInboundBranch(ir_if);
-    ir_if->False().target->AddInboundBranch(ir_if);
-
-    return ir_if;
 }
 
 Loop* Builder::CreateLoop() {
-    auto* ir_loop = ir.flow_nodes.Create<Loop>();
+    return ir.values.Create<Loop>(CreateBlock(), CreateBlock(), CreateBlock());
-    ir_loop->Start().target = CreateBlock();
-    ir_loop->Continuing().target = CreateBlock();
-    ir_loop->Merge().target = CreateBlock();
-
-    // A loop always branches to the start block.
-    ir_loop->Start().target->AddInboundBranch(ir_loop);
-
-    return ir_loop;
 }
 
 Switch* Builder::CreateSwitch(Value* condition) {
-    auto* ir_switch = ir.flow_nodes.Create<Switch>(condition);
+    return ir.values.Create<Switch>(condition, CreateBlock());
-    ir_switch->Merge().target = CreateBlock();
-    return ir_switch;
 }
 
 Block* Builder::CreateCase(Switch* s, utils::VectorRef<Switch::CaseSelector> selectors) {
-    s->Cases().Push(Switch::Case{selectors, {CreateBlock(), utils::Empty}});
+    s->Cases().Push(Switch::Case{std::move(selectors), CreateBlock()});
 
-    Block* b = s->Cases().Back().Start().target->As<Block>();
+    Block* b = s->Cases().Back().Start();
-    // Switch branches into the case block
     b->AddInboundBranch(s);
     return b;
 }
@@ -238,6 +210,14 @@ ir::Var* Builder::Declare(const type::Type* type) {
     return ir.values.Create<ir::Var>(type);
 }
 
+ir::Branch* Builder::Branch(FlowNode* to, utils::VectorRef<Value*> args) {
+    return ir.values.Create<ir::Branch>(to, args);
+}
+
+ir::Jump* Builder::Jump(FlowNode* to, utils::VectorRef<Value*> args) {
+    return ir.values.Create<ir::Jump>(to, args);
+}
+
 ir::BlockParam* Builder::BlockParam(const type::Type* type) {
     return ir.values.Create<ir::BlockParam>(type);
 }

src/tint/ir/builder.h

@@ -30,6 +30,7 @@
 #include "src/tint/ir/function_param.h"
 #include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/if.h"
+#include "src/tint/ir/jump.h"
 #include "src/tint/ir/load.h"
 #include "src/tint/ir/loop.h"
 #include "src/tint/ir/module.h"
@@ -351,6 +352,18 @@ class Builder {
     /// @returns the instruction
     ir::Var* Declare(const type::Type* type);
 
+    /// Creates a branch declaration
+    /// @param to the node being branched too
+    /// @param args the branch arguments
+    /// @returns the instruction
+    ir::Branch* Branch(FlowNode* to, utils::VectorRef<Value*> args = {});
+
+    /// Creates a jump declaration
+    /// @param to the node being branched too
+    /// @param args the branch arguments
+    /// @returns the instruction
+    ir::Jump* Jump(FlowNode* to, utils::VectorRef<Value*> args = {});
+
     /// Creates a new `BlockParam`
     /// @param type the parameter type
     /// @returns the value

src/tint/ir/debug.cc

@@ -60,81 +60,15 @@ std::string Debug::AsDotGraph(const Module* mod) {
                 if (node_to_name.count(b) == 0) {
                     out << name_for(b) << R"( [label="block"])" << std::endl;
                 }
-                out << name_for(b) << " -> " << name_for(b->Branch().target);
+                out << name_for(b) << " -> " << name_for(b->Branch()->To());
 
                 // Dashed lines to merge blocks
-                if (merge_nodes.count(b->Branch().target) != 0) {
+                if (merge_nodes.count(b->Branch()->To()) != 0) {
                     out << " [style=dashed]";
                 }
 
                 out << std::endl;
-                Graph(b->Branch().target);
+                Graph(b->Branch()->To());
-            },
-            [&](const ir::Switch* s) {
-                out << name_for(s) << R"( [label="switch"])" << std::endl;
-                out << name_for(s->Merge().target) << R"( [label="switch merge"])" << std::endl;
-                merge_nodes.insert(s->Merge().target);
-
-                size_t i = 0;
-                for (const auto& c : s->Cases()) {
-                    out << name_for(c.Start().target)
-                        << R"( [label="case )" + std::to_string(i++) + R"("])" << std::endl;
-                }
-                out << name_for(s) << " -> {";
-                for (const auto& c : s->Cases()) {
-                    if (&c != &(s->Cases().Front())) {
-                        out << ", ";
-                    }
-                    out << name_for(c.Start().target);
-                }
-                out << "}" << std::endl;
-
-                for (const auto& c : s->Cases()) {
-                    Graph(c.Start().target);
-                }
-                Graph(s->Merge().target);
-            },
-            [&](const ir::If* i) {
-                out << name_for(i) << R"( [label="if"])" << std::endl;
-                out << name_for(i->True().target) << R"( [label="true"])" << std::endl;
-                out << name_for(i->False().target) << R"( [label="false"])" << std::endl;
-                out << name_for(i->Merge().target) << R"( [label="if merge"])" << std::endl;
-                merge_nodes.insert(i->Merge().target);
-
-                out << name_for(i) << " -> {";
-                out << name_for(i->True().target) << ", " << name_for(i->False().target);
-                out << "}" << std::endl;
-
-                // Subgraph if true/false branches so they draw on the same line
-                out << "subgraph sub_" << name_for(i) << " {" << std::endl;
-                out << R"(rank="same")" << std::endl;
-                out << name_for(i->True().target) << std::endl;
-                out << name_for(i->False().target) << std::endl;
-                out << "}" << std::endl;
-
-                Graph(i->True().target);
-                Graph(i->False().target);
-                Graph(i->Merge().target);
-            },
-            [&](const ir::Loop* l) {
-                out << name_for(l) << R"( [label="loop"])" << std::endl;
-                out << name_for(l->Start().target) << R"( [label="start"])" << std::endl;
-                out << name_for(l->Continuing().target) << R"( [label="continuing"])" << std::endl;
-                out << name_for(l->Merge().target) << R"( [label="loop merge"])" << std::endl;
-                merge_nodes.insert(l->Merge().target);
-
-                // Subgraph the continuing and merge so they get drawn on the same line
-                out << "subgraph sub_" << name_for(l) << " {" << std::endl;
-                out << R"(rank="same")" << std::endl;
-                out << name_for(l->Continuing().target) << std::endl;
-                out << name_for(l->Merge().target) << std::endl;
-                out << "}" << std::endl;
-
-                out << name_for(l) << " -> " << name_for(l->Start().target) << std::endl;
-
-                Graph(l->Start().target);
-                Graph(l->Continuing().target);
-                Graph(l->Merge().target);
             },
             [&](const ir::FunctionTerminator*) {
                 // Already done

src/tint/ir/disassembler.cc

@@ -27,6 +27,7 @@
 #include "src/tint/ir/discard.h"
 #include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/if.h"
+#include "src/tint/ir/jump.h"
 #include "src/tint/ir/load.h"
 #include "src/tint/ir/loop.h"
 #include "src/tint/ir/root_terminator.h"
@@ -41,22 +42,6 @@
 namespace tint::ir {
 namespace {
 
-class ScopedStopNode {
-    static constexpr size_t N = 32;
-
-  public:
-    ScopedStopNode(utils::Hashset<const FlowNode*, N>& stop_nodes, const FlowNode* node)
-        : stop_nodes_(stop_nodes), node_(node) {
-        stop_nodes_.Add(node_);
-    }
-
-    ~ScopedStopNode() { stop_nodes_.Remove(node_); }
-
-  private:
-    utils::Hashset<const FlowNode*, N>& stop_nodes_;
-    const FlowNode* node_;
-};
-
 class ScopedIndent {
   public:
     explicit ScopedIndent(uint32_t& indent) : indent_(indent) { indent_ += 2; }
@@ -103,261 +88,115 @@ std::string_view Disassembler::IdOf(const Value* value) {
     });
 }
 
-void Disassembler::Walk(const FlowNode* node) {
-    if (visited_.Contains(node) || stop_nodes_.Contains(node)) {
-        return;
-    }
-    visited_.Add(node);
-
-    tint::Switch(
-        node,
-        [&](const ir::Function* f) {
-            TINT_SCOPED_ASSIGNMENT(in_function_, true);
-
-            Indent() << "%fn" << IdOf(f) << " = func " << f->Name().Name() << "(";
-            for (auto* p : f->Params()) {
-                if (p != f->Params().Front()) {
-                    out_ << ", ";
-                }
-                out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
-            }
-            out_ << "):" << f->ReturnType()->FriendlyName();
-
-            if (f->Stage() != Function::PipelineStage::kUndefined) {
-                out_ << " [@" << f->Stage();
-
-                if (f->WorkgroupSize()) {
-                    auto arr = f->WorkgroupSize().value();
-                    out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2]
-                         << ")";
-                }
-
-                if (!f->ReturnAttributes().IsEmpty()) {
-                    out_ << " ra:";
-
-                    for (auto attr : f->ReturnAttributes()) {
-                        out_ << " @" << attr;
-                        if (attr == Function::ReturnAttribute::kLocation) {
-                            out_ << "(" << f->ReturnLocation().value() << ")";
-                        }
-                    }
-                }
-
-                out_ << "]";
-            }
-            out_ << " {" << std::endl;
-
-            {
-                ScopedIndent func_indent(indent_size_);
-                ScopedStopNode scope(stop_nodes_, f->EndTarget());
-                Walk(f->StartTarget());
-            }
-            out_ << "} ";
-            Walk(f->EndTarget());
-        },
-        [&](const ir::Block* b) {
-            // If this block is dead, nothing to do
-            if (!b->HasBranchTarget()) {
-                return;
-            }
-
-            Indent() << "%fn" << IdOf(b) << " = block";
-            if (!b->Params().IsEmpty()) {
-                out_ << " (";
-                for (const auto* p : b->Params()) {
-                    if (p != b->Params().Front()) {
-                        out_ << ", ";
-                    }
-                    EmitValue(p);
-                }
-                out_ << ")";
-            }
-
-            out_ << " {" << std::endl;
-            {
-                ScopedIndent si(indent_size_);
-                EmitBlockInstructions(b);
-            }
-            Indent() << "}";
-
-            std::string suffix = "";
-            if (b->Branch().target->Is<FunctionTerminator>()) {
-                out_ << " -> %func_end";
-                suffix = "return";
-            } else if (b->Branch().target->Is<RootTerminator>()) {
-                // Nothing to do
-            } else {
-                out_ << " -> "
-                     << "%fn" << IdOf(b->Branch().target);
-                suffix = "branch";
-            }
-            if (!b->Branch().args.IsEmpty()) {
-                out_ << " ";
-                for (const auto* v : b->Branch().args) {
-                    if (v != b->Branch().args.Front()) {
-                        out_ << ", ";
-                    }
-                    EmitValue(v);
-                }
-            }
-            if (!suffix.empty()) {
-                out_ << " # " << suffix;
-            }
-            out_ << std::endl;
-
-            if (!b->Branch().target->Is<FunctionTerminator>()) {
-                out_ << std::endl;
-            }
-
-            Walk(b->Branch().target);
-        },
-        [&](const ir::Switch* s) {
-            Indent() << "%fn" << IdOf(s) << " = switch ";
-            EmitValue(s->Condition());
-            out_ << " [";
-            for (const auto& c : s->Cases()) {
-                if (&c != &s->Cases().Front()) {
-                    out_ << ", ";
-                }
-                out_ << "c: (";
-                for (const auto& selector : c.selectors) {
-                    if (&selector != &c.selectors.Front()) {
-                        out_ << " ";
-                    }
-
-                    if (selector.IsDefault()) {
-                        out_ << "default";
-                    } else {
-                        EmitValue(selector.val);
-                    }
-                }
-                out_ << ", %fn" << IdOf(c.Start().target) << ")";
-            }
-            if (s->Merge().target->IsConnected()) {
-                out_ << ", m: %fn" << IdOf(s->Merge().target);
-            }
-            out_ << "]" << std::endl;
-
-            {
-                ScopedIndent switch_indent(indent_size_);
-                ScopedStopNode scope(stop_nodes_, s->Merge().target);
-                for (const auto& c : s->Cases()) {
-                    Indent() << "# case ";
-                    for (const auto& selector : c.selectors) {
-                        if (&selector != &c.selectors.Front()) {
-                            out_ << " ";
-                        }
-
-                        if (selector.IsDefault()) {
-                            out_ << "default";
-                        } else {
-                            EmitValue(selector.val);
-                        }
-                    }
-                    out_ << std::endl;
-                    Walk(c.Start().target);
-                }
-            }
-
-            if (s->Merge().target->IsConnected()) {
-                Indent() << "# switch merge" << std::endl;
-                Walk(s->Merge().target);
-            }
-        },
-        [&](const ir::If* i) {
-            Indent() << "%fn" << IdOf(i) << " = if ";
-            EmitValue(i->Condition());
-
-            bool has_true = i->True().target->HasBranchTarget();
-            bool has_false = i->False().target->HasBranchTarget();
-
-            out_ << " [";
-            if (has_true) {
-                out_ << "t: %fn" << IdOf(i->True().target);
-            }
-            if (has_false) {
-                if (has_true) {
-                    out_ << ", ";
-                }
-                out_ << "f: %fn" << IdOf(i->False().target);
-            }
-            if (i->Merge().target->IsConnected()) {
-                out_ << ", m: %fn" << IdOf(i->Merge().target);
-            }
-            out_ << "]" << std::endl;
-
-            {
-                ScopedIndent if_indent(indent_size_);
-                ScopedStopNode scope(stop_nodes_, i->Merge().target);
-
-                if (has_true) {
-                    Indent() << "# true branch" << std::endl;
-                    Walk(i->True().target);
-                }
-
-                if (has_false) {
-                    Indent() << "# false branch" << std::endl;
-                    Walk(i->False().target);
-                }
-            }
-
-            if (i->Merge().target->IsConnected()) {
-                Indent() << "# if merge" << std::endl;
-                Walk(i->Merge().target);
-            }
-        },
-        [&](const ir::Loop* l) {
-            Indent() << "%fn" << IdOf(l) << " = loop [s: %fn" << IdOf(l->Start().target);
-
-            if (l->Continuing().target->IsConnected()) {
-                out_ << ", c: %fn" << IdOf(l->Continuing().target);
-            }
-            if (l->Merge().target->IsConnected()) {
-                out_ << ", m: %fn" << IdOf(l->Merge().target);
-            }
-            out_ << "]" << std::endl;
-
-            {
-                ScopedStopNode loop_scope(stop_nodes_, l->Merge().target);
-                ScopedIndent loop_indent(indent_size_);
-                {
-                    ScopedStopNode inner_scope(stop_nodes_, l->Continuing().target);
-                    Indent() << "# loop start" << std::endl;
-                    Walk(l->Start().target);
-                }
-
-                if (l->Continuing().target->IsConnected()) {
-                    Indent() << "# loop continuing" << std::endl;
-                    Walk(l->Continuing().target);
-                }
-            }
-
-            if (l->Merge().target->IsConnected()) {
-                Indent() << "# loop merge" << std::endl;
-                Walk(l->Merge().target);
-            }
-        },
-        [&](const ir::FunctionTerminator*) {
-            TINT_ASSERT(IR, in_function_);
-            Indent() << "%func_end" << std::endl << std::endl;
-        },
-        [&](const ir::RootTerminator*) {
-            TINT_ASSERT(IR, !in_function_);
-            out_ << std::endl;
-        });
-}
-
 std::string Disassembler::Disassemble() {
     if (mod_.root_block) {
-        Walk(mod_.root_block);
+        walk_list_.push_back(mod_.root_block);
+        Walk();
+        TINT_ASSERT(IR, walk_list_.empty());
     }
 
-    for (const auto* func : mod_.functions) {
+    for (auto* func : mod_.functions) {
-        Walk(func);
+        walk_list_.push_back(func);
+        Walk();
+        TINT_ASSERT(IR, walk_list_.empty());
     }
     return out_.str();
 }
 
+void Disassembler::Walk() {
+    utils::Hashset<const FlowNode*, 32> visited_;
+
+    while (!walk_list_.empty()) {
+        const FlowNode* node = walk_list_.front();
+        walk_list_.pop_front();
+
+        if (visited_.Contains(node)) {
+            continue;
+        }
+        visited_.Add(node);
+
+        tint::Switch(
+            node,
+            [&](const ir::Function* f) {
+                in_function_ = true;
+
+                Indent() << "%fn" << IdOf(f) << " = func " << f->Name().Name() << "(";
+                for (auto* p : f->Params()) {
+                    if (p != f->Params().Front()) {
+                        out_ << ", ";
+                    }
+                    out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
+                }
+                out_ << "):" << f->ReturnType()->FriendlyName();
+
+                if (f->Stage() != Function::PipelineStage::kUndefined) {
+                    out_ << " [@" << f->Stage();
+
+                    if (f->WorkgroupSize()) {
+                        auto arr = f->WorkgroupSize().value();
+                        out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2]
+                             << ")";
+                    }
+
+                    if (!f->ReturnAttributes().IsEmpty()) {
+                        out_ << " ra:";
+
+                        for (auto attr : f->ReturnAttributes()) {
+                            out_ << " @" << attr;
+                            if (attr == Function::ReturnAttribute::kLocation) {
+                                out_ << "(" << f->ReturnLocation().value() << ")";
+                            }
+                        }
+                    }
+
+                    out_ << "]";
+                }
+                out_ << " -> %fn" << IdOf(f->StartTarget()) << std::endl;
+                walk_list_.push_back(f->StartTarget());
+            },
+            [&](const ir::Block* b) {
+                // If this block is dead, nothing to do
+                if (!b->HasBranchTarget()) {
+                    return;
+                }
+
+                Indent() << "%fn" << IdOf(b) << " = block";
+                if (!b->Params().IsEmpty()) {
+                    out_ << " (";
+                    for (auto* p : b->Params()) {
+                        if (p != b->Params().Front()) {
+                            out_ << ", ";
+                        }
+                        EmitValue(p);
+                    }
+                    out_ << ")";
+                }
+
+                out_ << " {" << std::endl;
+                {
+                    ScopedIndent si(indent_size_);
+                    EmitBlockInstructions(b);
+                }
+                Indent() << "}" << std::endl;
+
+                if (!b->Branch()->To()->Is<FunctionTerminator>()) {
+                    out_ << std::endl;
+                }
+
+                walk_list_.push_back(b->Branch()->To());
+            },
+            [&](const ir::FunctionTerminator* t) {
+                TINT_ASSERT(IR, in_function_);
+                Indent() << "%fn" << IdOf(t) << " = func_terminator" << std::endl << std::endl;
+                in_function_ = false;
+            },
+            [&](const ir::RootTerminator* t) {
+                TINT_ASSERT(IR, !in_function_);
+                Indent() << "%fn" << IdOf(t) << " = root_terminator" << std::endl << std::endl;
+            });
+    }
+}
+
 void Disassembler::EmitValueWithType(const Value* val) {
     EmitValue(val);
     if (auto* i = val->As<ir::Instruction>(); i->Type() != nullptr) {
@@ -419,8 +258,12 @@ void Disassembler::EmitValue(const Value* val) {
 
 void Disassembler::EmitInstruction(const Instruction* inst) {
     tint::Switch(
-        inst,  //
+        inst,                                         //
-        [&](const ir::Binary* b) { EmitBinary(b); }, [&](const ir::Unary* u) { EmitUnary(u); },
+        [&](const ir::Switch* s) { EmitSwitch(s); },  //
+        [&](const ir::If* i) { EmitIf(i); },          //
+        [&](const ir::Loop* l) { EmitLoop(l); },      //
+        [&](const ir::Binary* b) { EmitBinary(b); },  //
+        [&](const ir::Unary* u) { EmitUnary(u); },
         [&](const ir::Bitcast* b) {
             EmitValueWithType(b);
             out_ << " = bitcast ";
@@ -468,7 +311,131 @@ void Disassembler::EmitInstruction(const Instruction* inst) {
                 out_ << ", ";
                 EmitValue(v->Initializer());
             }
-        });
+        },
+        [&](const ir::Branch* b) { EmitBranch(b); },
+        [&](Default) { out_ << "Unknown instruction: " << inst->TypeInfo().name; });
+}
+
+void Disassembler::EmitIf(const If* i) {
+    out_ << "if ";
+    EmitValue(i->Condition());
+
+    bool has_true = i->True()->HasBranchTarget();
+    bool has_false = i->False()->HasBranchTarget();
+
+    out_ << " [";
+    if (has_true) {
+        out_ << "t: %fn" << IdOf(i->True());
+    }
+    if (has_false) {
+        if (has_true) {
+            out_ << ", ";
+        }
+        out_ << "f: %fn" << IdOf(i->False());
+    }
+    if (i->Merge()->IsConnected()) {
+        out_ << ", m: %fn" << IdOf(i->Merge());
+    }
+    out_ << "]";
+
+    if (has_true) {
+        walk_list_.push_back(i->True());
+    }
+    if (has_false) {
+        walk_list_.push_back(i->False());
+    }
+    if (i->Merge()->IsConnected()) {
+        walk_list_.push_back(i->Merge());
+    }
+}
+
+void Disassembler::EmitLoop(const Loop* l) {
+    out_ << "loop [s: %fn" << IdOf(l->Start());
+
+    if (l->Continuing()->IsConnected()) {
+        out_ << ", c: %fn" << IdOf(l->Continuing());
+    }
+    if (l->Merge()->IsConnected()) {
+        out_ << ", m: %fn" << IdOf(l->Merge());
+    }
+    out_ << "]";
+
+    { walk_list_.push_back(l->Start()); }
+
+    if (l->Continuing()->IsConnected()) {
+        walk_list_.push_back(l->Continuing());
+    }
+    if (l->Merge()->IsConnected()) {
+        walk_list_.push_back(l->Merge());
+    }
+}
+
+void Disassembler::EmitSwitch(const Switch* s) {
+    out_ << "switch ";
+    EmitValue(s->Condition());
+    out_ << " [";
+    for (const auto& c : s->Cases()) {
+        if (&c != &s->Cases().Front()) {
+            out_ << ", ";
+        }
+        out_ << "c: (";
+        for (const auto& selector : c.selectors) {
+            if (&selector != &c.selectors.Front()) {
+                out_ << " ";
+            }
+
+            if (selector.IsDefault()) {
+                out_ << "default";
+            } else {
+                EmitValue(selector.val);
+            }
+        }
+        out_ << ", %fn" << IdOf(c.Start()) << ")";
+    }
+    if (s->Merge()->IsConnected()) {
+        out_ << ", m: %fn" << IdOf(s->Merge());
+    }
+    out_ << "]";
+
+    for (auto& c : s->Cases()) {
+        walk_list_.push_back(c.Start());
+    }
+    if (s->Merge()->IsConnected()) {
+        walk_list_.push_back(s->Merge());
+    }
+}
+
+void Disassembler::EmitBranch(const Branch* b) {
+    if (b->Is<Jump>()) {
+        out_ << "jmp ";
+
+        // Stuff the thing we're jumping too into the front of the walk list so it will be emitted
+        // next.
+        walk_list_.push_front(b->To());
+    } else {
+        out_ << "br ";
+    }
+
+    std::string suffix = "";
+    out_ << "%fn" << IdOf(b->To());
+    if (b->To()->Is<FunctionTerminator>()) {
+        suffix = "return";
+    } else if (b->To()->Is<RootTerminator>()) {
+        suffix = "root_end";
+    }
+
+    if (!b->Args().IsEmpty()) {
+        out_ << " ";
+        for (auto* v : b->Args()) {
+            if (v != b->Args().Front()) {
+                out_ << ", ";
+            }
+            EmitValue(v);
+        }
+    }
+    if (!suffix.empty()) {
+        out_ << "  # " << suffix;
+    }
 }
 
 void Disassembler::EmitArgs(const Call* call) {

src/tint/ir/disassembler.h

@@ -15,15 +15,18 @@
 #ifndef SRC_TINT_IR_DISASSEMBLER_H_
 #define SRC_TINT_IR_DISASSEMBLER_H_
 
+#include <deque>
 #include <string>
 
 #include "src/tint/ir/binary.h"
 #include "src/tint/ir/call.h"
 #include "src/tint/ir/flow_node.h"
+#include "src/tint/ir/if.h"
+#include "src/tint/ir/loop.h"
 #include "src/tint/ir/module.h"
+#include "src/tint/ir/switch.h"
 #include "src/tint/ir/unary.h"
 #include "src/tint/utils/hashmap.h"
-#include "src/tint/utils/hashset.h"
 #include "src/tint/utils/string_stream.h"
 
 namespace tint::ir {
@@ -53,18 +56,21 @@ class Disassembler {
     size_t IdOf(const FlowNode* node);
     std::string_view IdOf(const Value* node);
 
-    void Walk(const FlowNode* node);
+    void Walk();
     void EmitInstruction(const Instruction* inst);
     void EmitValueWithType(const Value* val);
     void EmitValue(const Value* val);
     void EmitArgs(const Call* call);
     void EmitBinary(const Binary* b);
     void EmitUnary(const Unary* b);
+    void EmitBranch(const Branch* b);
+    void EmitSwitch(const Switch* s);
+    void EmitLoop(const Loop* l);
+    void EmitIf(const If* i);
 
     const Module& mod_;
     utils::StringStream out_;
-    utils::Hashset<const FlowNode*, 32> visited_;
+    std::deque<const FlowNode*> walk_list_;
-    utils::Hashset<const FlowNode*, 32> stop_nodes_;
     utils::Hashmap<const FlowNode*, size_t, 32> flow_node_ids_;
     utils::Hashmap<const Value*, std::string, 32> value_ids_;
     uint32_t indent_size_ = 0;

src/tint/ir/flow_node.h

@@ -18,6 +18,11 @@
 #include "src/tint/utils/castable.h"
 #include "src/tint/utils/vector.h"
 
+// Forward Declarations
+namespace tint::ir {
+class Branch;
+}  // namespace tint::ir
+
 namespace tint::ir {
 
 /// Base class for flow nodes
@@ -26,17 +31,17 @@ class FlowNode : public utils::Castable<FlowNode> {
     ~FlowNode() override;
 
     /// @returns true if this node has inbound branches and branches out
-    bool IsConnected() const { return HasBranchTarget() && !inbound_branches_.IsEmpty(); }
+    bool IsConnected() const { return HasBranchTarget(); }
 
     /// @returns true if the node has a branch target
     virtual bool HasBranchTarget() const { return false; }
 
     /// @returns the inbound branch list for the flow node
-    utils::VectorRef<FlowNode*> InboundBranches() const { return inbound_branches_; }
+    utils::VectorRef<Branch*> InboundBranches() const { return inbound_branches_; }
 
     /// Adds the given node to the inbound branches
     /// @param node the node to add
-    void AddInboundBranch(FlowNode* node) { inbound_branches_.Push(node); }
+    void AddInboundBranch(Branch* node) { inbound_branches_.Push(node); }
 
   protected:
     /// Constructor
@@ -48,7 +53,7 @@ class FlowNode : public utils::Castable<FlowNode> {
     ///   - Node is a start node
     ///   - Node is a merge target outside control flow (e.g. an if that returns in both branches)
     ///   - Node is a continue target outside control flow (e.g. a loop that returns)
-    utils::Vector<FlowNode*, 2> inbound_branches_;
+    utils::Vector<Branch*, 2> inbound_branches_;
 };
 
 }  // namespace tint::ir

src/tint/ir/from_program.cc

@@ -98,19 +98,15 @@ namespace {
 
 using ResultType = utils::Result<Module, diag::List>;
 
-bool IsConnected(const FlowNode* b) {
+// For an `if` and `switch` block, the merge has a registered incoming branch instruction of the
+// `if` and `switch. So, to determine if the merge is connected to any of the branches that happend
+// in the `if` or `switch` we need a `count` value that is larger then 1.
+bool IsConnected(const FlowNode* b, uint32_t count) {
     // Function is always connected as it's the start.
     if (b->Is<ir::Function>()) {
         return true;
     }
-
+    return b->InboundBranches().Length() > count;
-    for (auto* parent : b->InboundBranches()) {
-        if (IsConnected(parent)) {
-            return true;
-        }
-    }
-    // Getting here means all the incoming branches are disconnected.
-    return false;
 }
 
 /// Impl is the private-implementation of FromProgram().
@@ -145,8 +141,8 @@ class Impl {
         /* dst */ {&builder_.ir.constants_arena},
     };
 
-    /// The stack of flow control blocks.
+    /// The stack of control blocks.
-    utils::Vector<FlowNode*, 8> flow_stack_;
+    utils::Vector<Branch*, 8> control_stack_;
 
     /// The current flow block for expressions.
     Block* current_flow_block_ = nullptr;
@@ -160,15 +156,11 @@ class Impl {
     /// The diagnostic that have been raised.
     diag::List diagnostics_;
 
-    /// Map from ast nodes to flow nodes, used to retrieve the flow node for a given AST node.
+    class ControlStackScope {
-    /// Used for testing purposes.
-    std::unordered_map<const ast::Node*, const FlowNode*> ast_to_flow_;
-
-    class FlowStackScope {
       public:
-        FlowStackScope(Impl* impl, FlowNode* node) : impl_(impl) { impl_->flow_stack_.Push(node); }
+        ControlStackScope(Impl* impl, Branch* b) : impl_(impl) { impl_->control_stack_.Push(b); }
 
-        ~FlowStackScope() { impl_->flow_stack_.Pop(); }
+        ~ControlStackScope() { impl_->control_stack_.Pop(); }
 
       private:
         Impl* impl_;
@@ -178,11 +170,25 @@ class Impl {
         diagnostics_.add_error(tint::diag::System::IR, err, s);
     }
 
+    void JumpTo(FlowNode* node, utils::VectorRef<Value*> args = {}) {
+        TINT_ASSERT(IR, current_flow_block_);
+        TINT_ASSERT(IR, !current_flow_block_->HasBranchTarget());
+
+        current_flow_block_->Instructions().Push(builder_.Jump(node, args));
+        current_flow_block_ = nullptr;
+    }
+    void JumpToIfNeeded(FlowNode* node) {
+        if (!current_flow_block_ || current_flow_block_->HasBranchTarget()) {
+            return;
+        }
+        JumpTo(node);
+    }
+
     void BranchTo(FlowNode* node, utils::VectorRef<Value*> args = {}) {
         TINT_ASSERT(IR, current_flow_block_);
         TINT_ASSERT(IR, !current_flow_block_->HasBranchTarget());
 
-        current_flow_block_->BranchTo(node, args);
+        current_flow_block_->Instructions().Push(builder_.Branch(node, args));
         current_flow_block_ = nullptr;
     }
 
@@ -193,8 +199,8 @@ class Impl {
         BranchTo(node);
     }
 
-    FlowNode* FindEnclosingControl(ControlFlags flags) {
+    Branch* FindEnclosingControl(ControlFlags flags) {
-        for (auto it = flow_stack_.rbegin(); it != flow_stack_.rend(); ++it) {
+        for (auto it = control_stack_.rbegin(); it != control_stack_.rend(); ++it) {
             if ((*it)->Is<Loop>()) {
                 return *it;
             }
@@ -244,6 +250,11 @@ class Impl {
                 });
         }
 
+        // Add the root terminator if needed
+        if (mod.root_block) {
+            mod.root_block->Instructions().Push(builder_.Branch(builder_.CreateRootTerminator()));
+        }
+
         if (diagnostics_.contains_errors()) {
             return ResultType(std::move(diagnostics_));
         }
@@ -253,7 +264,7 @@ class Impl {
 
     void EmitFunction(const ast::Function* ast_func) {
         // The flow stack should have been emptied when the previous function finished building.
-        TINT_ASSERT(IR, flow_stack_.IsEmpty());
+        TINT_ASSERT(IR, control_stack_.IsEmpty());
 
         const auto* sem = program_->Sem().Get(ast_func);
 
@@ -262,8 +273,6 @@ class Impl {
         current_function_ = ir_func;
         builder_.ir.functions.Push(ir_func);
 
-        ast_to_flow_[ast_func] = ir_func;
-
         if (ast_func->IsEntryPoint()) {
             switch (ast_func->PipelineStage()) {
                 case ast::PipelineStage::kVertex:
@@ -343,17 +352,15 @@ class Impl {
         ir_func->SetParams(params);
 
         {
-            FlowStackScope scope(this, ir_func);
-
             current_flow_block_ = ir_func->StartTarget();
             EmitBlock(ast_func->body);
 
-            // If the branch target has already been set then a `return` was called. Only set in the
+            // If the branch target has already been set then a `return` was called. Only set in
-            // case where `return` wasn't called.
+            // the case where `return` wasn't called.
-            BranchToIfNeeded(current_function_->EndTarget());
+            JumpToIfNeeded(current_function_->EndTarget());
         }
 
-        TINT_ASSERT(IR, flow_stack_.IsEmpty());
+        TINT_ASSERT(IR, control_stack_.IsEmpty());
         current_flow_block_ = nullptr;
         current_function_ = nullptr;
     }
@@ -362,8 +369,8 @@ class Impl {
         for (auto* s : stmts) {
             EmitStatement(s);
 
-            // If the current flow block has a branch target then the rest of the statements in this
+            // If the current flow block has a branch target then the rest of the statements in
-            // block are dead code. Skip them.
+            // this block are dead code. Skip them.
             if (!current_flow_block_ || current_flow_block_->HasBranchTarget()) {
                 break;
             }
@@ -399,11 +406,11 @@ class Impl {
     }
 
     void EmitAssignment(const ast::AssignmentStatement* stmt) {
-        // If assigning to a phony, just generate the RHS and we're done. Note that, because this
+        // If assigning to a phony, just generate the RHS and we're done. Note that, because
-        // isn't used, a subsequent transform could remove it due to it being dead code. This could
+        // this isn't used, a subsequent transform could remove it due to it being dead code.
-        // then change the interface for the program (i.e. a global var no longer used). If that
+        // This could then change the interface for the program (i.e. a global var no longer
-        // happens we have to either fix this to store to a phony value, or make sure we pull the
+        // used). If that happens we have to either fix this to store to a phony value, or make
-        // interface before doing the dead code elimination.
+        // sure we pull the interface before doing the dead code elimination.
         if (stmt->lhs->Is<ast::PhonyExpression>()) {
             (void)EmitExpression(stmt->rhs);
             return;
@@ -523,8 +530,8 @@ class Impl {
         TINT_DEFER(scopes_.Pop());
 
         // Note, this doesn't need to emit a Block as the current block flow node should be
-        // sufficient as the blocks all get flattened. Each flow control node will inject the basic
+        // sufficient as the blocks all get flattened. Each flow control node will inject the
-        // blocks it requires.
+        // basic blocks it requires.
         EmitStatements(block->statements);
     }
 
@@ -534,50 +541,43 @@ class Impl {
         if (!reg) {
             return;
         }
-        auto* if_node = builder_.CreateIf(reg.Get());
+        auto* if_inst = builder_.CreateIf(reg.Get());
-
+        current_flow_block_->Instructions().Push(if_inst);
-        BranchTo(if_node);
-
-        ast_to_flow_[stmt] = if_node;
 
         {
-            FlowStackScope scope(this, if_node);
+            ControlStackScope scope(this, if_inst);
 
-            current_flow_block_ = if_node->True().target->As<Block>();
+            current_flow_block_ = if_inst->True();
             EmitBlock(stmt->body);
 
             // If the true branch did not execute control flow, then go to the Merge().target
-            BranchToIfNeeded(if_node->Merge().target);
+            BranchToIfNeeded(if_inst->Merge());
 
-            current_flow_block_ = if_node->False().target->As<Block>();
+            current_flow_block_ = if_inst->False();
             if (stmt->else_statement) {
                 EmitStatement(stmt->else_statement);
             }
 
             // If the false branch did not execute control flow, then go to the Merge().target
-            BranchToIfNeeded(if_node->Merge().target);
+            BranchToIfNeeded(if_inst->Merge());
         }
         current_flow_block_ = nullptr;
 
-        // If both branches went somewhere, then they both returned, continued or broke. So, there
+        // If both branches went somewhere, then they both returned, continued or broke. So,
-        // is no need for the if merge-block and there is nothing to branch to the merge block
+        // there is no need for the if merge-block and there is nothing to branch to the merge
-        // anyway.
+        // block anyway.
-        if (IsConnected(if_node->Merge().target)) {
+        if (IsConnected(if_inst->Merge(), 1)) {
-            current_flow_block_ = if_node->Merge().target->As<Block>();
+            current_flow_block_ = if_inst->Merge();
         }
     }
 
     void EmitLoop(const ast::LoopStatement* stmt) {
-        auto* loop_node = builder_.CreateLoop();
+        auto* loop_inst = builder_.CreateLoop();
-
+        current_flow_block_->Instructions().Push(loop_inst);
-        BranchTo(loop_node);
-
-        ast_to_flow_[stmt] = loop_node;
 
         {
-            FlowStackScope scope(this, loop_node);
+            ControlStackScope scope(this, loop_inst);
-
+            current_flow_block_ = loop_inst->Start();
-            current_flow_block_ = loop_node->Start().target->As<Block>();
 
             // The loop doesn't use EmitBlock because it needs the scope stack to not get popped
             // until after the continuing block.
@@ -585,41 +585,39 @@ class Impl {
             TINT_DEFER(scopes_.Pop());
             EmitStatements(stmt->body->statements);
 
-            // The current block didn't `break`, `return` or `continue`, go to the continuing block.
+            // The current block didn't `break`, `return` or `continue`, go to the continuing
-            BranchToIfNeeded(loop_node->Continuing().target);
+            // block.
+            JumpToIfNeeded(loop_inst->Continuing());
 
-            current_flow_block_ = loop_node->Continuing().target->As<Block>();
+            current_flow_block_ = loop_inst->Continuing();
             if (stmt->continuing) {
                 EmitBlock(stmt->continuing);
             }
 
             // Branch back to the start node if the continue target didn't branch out already
-            BranchToIfNeeded(loop_node->Start().target);
+            BranchToIfNeeded(loop_inst->Start());
         }
 
         // The loop merge can get disconnected if the loop returns directly, or the continuing
         // target branches, eventually, to the merge, but nothing branched to the
         // Continuing().target.
-        current_flow_block_ = loop_node->Merge().target->As<Block>();
+        current_flow_block_ = loop_inst->Merge();
-        if (!IsConnected(loop_node->Merge().target)) {
+        if (!IsConnected(loop_inst->Merge(), 0)) {
             current_flow_block_ = nullptr;
         }
     }
 
     void EmitWhile(const ast::WhileStatement* stmt) {
-        auto* loop_node = builder_.CreateLoop();
+        auto* loop_inst = builder_.CreateLoop();
+        current_flow_block_->Instructions().Push(loop_inst);
+
         // Continue is always empty, just go back to the start
-        TINT_ASSERT(IR, loop_node->Continuing().target->Is<Block>());
+        loop_inst->Continuing()->Instructions().Push(builder_.Branch(loop_inst->Start()));
-        loop_node->Continuing().target->As<Block>()->BranchTo(loop_node->Start().target);
-
-        BranchTo(loop_node);
-
-        ast_to_flow_[stmt] = loop_node;
 
         {
-            FlowStackScope scope(this, loop_node);
+            ControlStackScope scope(this, loop_inst);
 
-            current_flow_block_ = loop_node->Start().target->As<Block>();
+            current_flow_block_ = loop_inst->Start();
 
             // Emit the while condition into the Start().target of the loop
             auto reg = EmitExpression(stmt->condition);
@@ -628,25 +626,26 @@ class Impl {
             }
 
             // Create an `if (cond) {} else {break;}` control flow
-            auto* if_node = builder_.CreateIf(reg.Get());
+            auto* if_inst = builder_.CreateIf(reg.Get());
-            if_node->True().target->As<Block>()->BranchTo(if_node->Merge().target);
+            if_inst->True()->Instructions().Push(builder_.Branch(if_inst->Merge()));
-            if_node->False().target->As<Block>()->BranchTo(loop_node->Merge().target);
+            if_inst->False()->Instructions().Push(builder_.Branch(loop_inst->Merge()));
+            current_flow_block_->Instructions().Push(if_inst);
 
-            BranchTo(if_node);
+            current_flow_block_ = if_inst->Merge();
-
-            current_flow_block_ = if_node->Merge().target->As<Block>();
             EmitBlock(stmt->body);
 
-            BranchToIfNeeded(loop_node->Continuing().target);
+            JumpToIfNeeded(loop_inst->Continuing());
         }
         // The while loop always has a path to the Merge().target as the break statement comes
         // before anything inside the loop.
-        current_flow_block_ = loop_node->Merge().target->As<Block>();
+        current_flow_block_ = loop_inst->Merge();
     }
 
     void EmitForLoop(const ast::ForLoopStatement* stmt) {
-        auto* loop_node = builder_.CreateLoop();
+        auto* loop_inst = builder_.CreateLoop();
-        loop_node->Continuing().target->As<Block>()->BranchTo(loop_node->Start().target);
+        current_flow_block_->Instructions().Push(loop_inst);
+
+        loop_inst->Continuing()->Instructions().Push(builder_.Branch(loop_inst->Start()));
 
         // Make sure the initializer ends up in a contained scope
         scopes_.Push();
@@ -657,14 +656,10 @@ class Impl {
             EmitStatement(stmt->initializer);
         }
 
-        BranchTo(loop_node);
-
-        ast_to_flow_[stmt] = loop_node;
-
         {
-            FlowStackScope scope(this, loop_node);
+            ControlStackScope scope(this, loop_inst);
 
-            current_flow_block_ = loop_node->Start().target->As<Block>();
+            current_flow_block_ = loop_inst->Start();
 
             if (stmt->condition) {
                 // Emit the condition into the target target of the loop
@@ -674,26 +669,26 @@ class Impl {
                 }
 
                 // Create an `if (cond) {} else {break;}` control flow
-                auto* if_node = builder_.CreateIf(reg.Get());
+                auto* if_inst = builder_.CreateIf(reg.Get());
-                if_node->True().target->As<Block>()->BranchTo(if_node->Merge().target);
+                if_inst->True()->Instructions().Push(builder_.Branch(if_inst->Merge()));
-                if_node->False().target->As<Block>()->BranchTo(loop_node->Merge().target);
+                if_inst->False()->Instructions().Push(builder_.Branch(loop_inst->Merge()));
+                current_flow_block_->Instructions().Push(if_inst);
 
-                BranchTo(if_node);
+                current_flow_block_ = if_inst->Merge();
-                current_flow_block_ = if_node->Merge().target->As<Block>();
             }
 
             EmitBlock(stmt->body);
-            BranchToIfNeeded(loop_node->Continuing().target);
+            JumpToIfNeeded(loop_inst->Continuing());
 
             if (stmt->continuing) {
-                current_flow_block_ = loop_node->Continuing().target->As<Block>();
+                current_flow_block_ = loop_inst->Continuing();
                 EmitStatement(stmt->continuing);
             }
         }
 
         // The while loop always has a path to the Merge().target as the break statement comes
         // before anything inside the loop.
-        current_flow_block_ = loop_node->Merge().target->As<Block>();
+        current_flow_block_ = loop_inst->Merge();
     }
 
     void EmitSwitch(const ast::SwitchStatement* stmt) {
@@ -702,14 +697,11 @@ class Impl {
         if (!reg) {
             return;
         }
-        auto* switch_node = builder_.CreateSwitch(reg.Get());
+        auto* switch_inst = builder_.CreateSwitch(reg.Get());
-
+        current_flow_block_->Instructions().Push(switch_inst);
-        BranchTo(switch_node);
-
-        ast_to_flow_[stmt] = switch_node;
 
         {
-            FlowStackScope scope(this, switch_node);
+            ControlStackScope scope(this, switch_inst);
 
             const auto* sem = program_->Sem().Get(stmt);
             for (const auto* c : sem->Cases()) {
@@ -722,16 +714,16 @@ class Impl {
                     }
                 }
 
-                current_flow_block_ = builder_.CreateCase(switch_node, selectors);
+                current_flow_block_ = builder_.CreateCase(switch_inst, selectors);
                 EmitBlock(c->Body()->Declaration());
 
-                BranchToIfNeeded(switch_node->Merge().target);
+                BranchToIfNeeded(switch_inst->Merge());
             }
         }
         current_flow_block_ = nullptr;
 
-        if (IsConnected(switch_node->Merge().target)) {
+        if (IsConnected(switch_inst->Merge(), 1)) {
-            current_flow_block_ = switch_node->Merge().target->As<Block>();
+            current_flow_block_ = switch_inst->Merge();
         }
     }
 
@@ -753,9 +745,9 @@ class Impl {
         TINT_ASSERT(IR, current_control);
 
         if (auto* c = current_control->As<Loop>()) {
-            BranchTo(c->Merge().target);
+            BranchTo(c->Merge());
         } else if (auto* s = current_control->As<Switch>()) {
-            BranchTo(s->Merge().target);
+            BranchTo(s->Merge());
         } else {
             TINT_UNREACHABLE(IR, diagnostics_);
         }
@@ -766,14 +758,14 @@ class Impl {
         TINT_ASSERT(IR, current_control);
 
         if (auto* c = current_control->As<Loop>()) {
-            BranchTo(c->Continuing().target);
+            BranchTo(c->Continuing());
         } else {
             TINT_UNREACHABLE(IR, diagnostics_);
         }
     }
 
-    // Discard is being treated as an instruction. The semantics in WGSL is demote_to_helper, so the
+    // Discard is being treated as an instruction. The semantics in WGSL is demote_to_helper, so
-    // code has to continue as before it just predicates writes. If WGSL grows some kind of
+    // the code has to continue as before it just predicates writes. If WGSL grows some kind of
     // terminating discard that would probably make sense as a FlowNode but would then require
     // figuring out the multi-level exit that is triggered.
     void EmitDiscard(const ast::DiscardStatement*) {
@@ -787,11 +779,8 @@ class Impl {
         if (!reg) {
             return;
         }
-        auto* if_node = builder_.CreateIf(reg.Get());
+        auto* if_inst = builder_.CreateIf(reg.Get());
-
+        current_flow_block_->Instructions().Push(if_inst);
-        BranchTo(if_node);
-
-        ast_to_flow_[stmt] = if_node;
 
         auto* current_control = FindEnclosingControl(ControlFlags::kExcludeSwitch);
         TINT_ASSERT(IR, current_control);
@@ -799,17 +788,17 @@ class Impl {
 
         auto* loop = current_control->As<Loop>();
 
-        current_flow_block_ = if_node->True().target->As<Block>();
+        current_flow_block_ = if_inst->True();
-        BranchTo(loop->Merge().target);
+        BranchTo(loop->Merge());
 
-        current_flow_block_ = if_node->False().target->As<Block>();
+        current_flow_block_ = if_inst->False();
-        BranchTo(if_node->Merge().target);
+        BranchTo(if_inst->Merge());
 
-        current_flow_block_ = if_node->Merge().target->As<Block>();
+        current_flow_block_ = if_inst->Merge();
 
-        // The `break-if` has to be the last item in the continuing block. The false branch of the
+        // The `break-if` has to be the last item in the continuing block. The false branch of
-        // `break-if` will always take us back to the start of the loop.
+        // the `break-if` will always take us back to the start of the loop.
-        BranchTo(loop->Start().target);
+        BranchTo(loop->Start());
     }
 
     utils::Result<Value*> EmitExpression(const ast::Expression* expr) {
@@ -845,8 +834,8 @@ class Impl {
             // TODO(dsinclair): Implement
             // },
             [&](const ast::UnaryOpExpression* u) { return EmitUnary(u); },
-            // Note, ast::PhonyExpression is explicitly not handled here as it should never get into
+            // Note, ast::PhonyExpression is explicitly not handled here as it should never get
-            // this method. The assignment statement should have filtered it out already.
+            // into this method. The assignment statement should have filtered it out already.
             [&](Default) {
                 add_error(expr->source,
                           "unknown expression type: " + std::string(expr->TypeInfo().name));
@@ -891,8 +880,8 @@ class Impl {
                 builder_.ir.SetName(val, v->name->symbol.Name());
             },
             [&](const ast::Let* l) {
-                // A `let` doesn't exist as a standalone item in the IR, it's just the result of the
+                // A `let` doesn't exist as a standalone item in the IR, it's just the result of
-                // initializer.
+                // the initializer.
                 auto init = EmitExpression(l->initializer);
                 if (!init) {
                     return;
@@ -911,12 +900,12 @@ class Impl {
             },
             [&](const ast::Const*) {
                 // Skip. This should be handled by const-eval already, so the const will be a
-                // `constant::` value at the usage sites. Can just ignore the `const` variable as it
+                // `constant::` value at the usage sites. Can just ignore the `const` variable
-                // should never be used.
+                // as it should never be used.
                 //
-                // TODO(dsinclair): Probably want to store the const variable somewhere and then in
+                // TODO(dsinclair): Probably want to store the const variable somewhere and then
-                // identifier expression log an error if we ever see a const identifier. Add this
+                // in identifier expression log an error if we ever see a const identifier. Add
-                // when identifiers and variables are supported.
+                // this when identifiers and variables are supported.
             },
             [&](Default) {
                 add_error(var->source, "unknown variable: " + std::string(var->TypeInfo().name));
@@ -953,8 +942,8 @@ class Impl {
         return inst;
     }
 
-    // A short-circut needs special treatment. The short-circuit is decomposed into the relevant if
+    // A short-circut needs special treatment. The short-circuit is decomposed into the relevant
-    // statements and declarations.
+    // if statements and declarations.
     utils::Result<Value*> EmitShortCircuit(const ast::BinaryExpression* expr) {
         switch (expr->op) {
             case ast::BinaryOp::kLogicalAnd:
@@ -972,15 +961,15 @@ class Impl {
             return utils::Failure;
         }
 
-        auto* if_node = builder_.CreateIf(lhs.Get());
+        auto* if_inst = builder_.CreateIf(lhs.Get());
-        BranchTo(if_node);
+        current_flow_block_->Instructions().Push(if_inst);
 
         auto* result = builder_.BlockParam(builder_.ir.types.Get<type::Bool>());
-        if_node->Merge().target->As<Block>()->SetParams(utils::Vector{result});
+        if_inst->Merge()->SetParams(utils::Vector{result});
 
         utils::Result<Value*> rhs;
         {
-            FlowStackScope scope(this, if_node);
+            ControlStackScope scope(this, if_inst);
 
             utils::Vector<Value*, 1> alt_args;
             alt_args.Push(lhs.Get());
@@ -988,19 +977,19 @@ class Impl {
             // If this is an `&&` then we only evaluate the RHS expression in the true block.
             // If this is an `||` then we only evaluate the RHS expression in the false block.
             if (expr->op == ast::BinaryOp::kLogicalAnd) {
-                // If the lhs is false, then that is the result we want to pass to the merge block
+                // If the lhs is false, then that is the result we want to pass to the merge
-                // as our argument
+                // block as our argument
-                current_flow_block_ = if_node->False().target->As<Block>();
+                current_flow_block_ = if_inst->False();
-                BranchTo(if_node->Merge().target, std::move(alt_args));
+                BranchTo(if_inst->Merge(), std::move(alt_args));
 
-                current_flow_block_ = if_node->True().target->As<Block>();
+                current_flow_block_ = if_inst->True();
             } else {
-                // If the lhs is true, then that is the result we want to pass to the merge block
+                // If the lhs is true, then that is the result we want to pass to the merge
-                // as our argument
+                // block as our argument
-                current_flow_block_ = if_node->True().target->As<Block>();
+                current_flow_block_ = if_inst->True();
-                BranchTo(if_node->Merge().target, std::move(alt_args));
+                BranchTo(if_inst->Merge(), std::move(alt_args));
 
-                current_flow_block_ = if_node->False().target->As<Block>();
+                current_flow_block_ = if_inst->False();
             }
 
             rhs = EmitExpression(expr->rhs);
@@ -1010,9 +999,9 @@ class Impl {
             utils::Vector<Value*, 1> args;
             args.Push(rhs.Get());
 
-            BranchTo(if_node->Merge().target, std::move(args));
+            BranchTo(if_inst->Merge(), std::move(args));
         }
-        current_flow_block_ = if_node->Merge().target->As<Block>();
+        current_flow_block_ = if_inst->Merge();
 
         return result;
     }
@@ -1191,67 +1180,6 @@ class Impl {
         }
         return builder_.Constant(cv);
     }
-
-    //    void EmitAttributes(utils::VectorRef<const ast::Attribute*> attrs) {
-    //        for (auto* attr : attrs) {
-    //            EmitAttribute(attr);
-    //        }
-    //    }
-    //
-    //    void EmitAttribute(const ast::Attribute* attr) {
-    //        tint::Switch(  //
-    //            attr,
-    //            [&](const ast::WorkgroupAttribute* wg) {
-    //              // TODO(dsinclair): Implement
-    //            },
-    //            [&](const ast::StageAttribute* s) {
-    //              // TODO(dsinclair): Implement
-    //            },
-    //            [&](const ast::BindingAttribute* b) {
-    //              // TODO(dsinclair): Implement
-    //            },
-    //            [&](const ast::GroupAttribute* g) {
-    //              // TODO(dsinclair): Implement
-    //            },
-    //            [&](const ast::LocationAttribute* l) {
-    //              // TODO(dsinclair): Implement
-    //            },
-    //            [&](const ast::BuiltinAttribute* b) {
-    //              // TODO(dsinclair): Implement
-    //            },
-    //            [&](const ast::InterpolateAttribute* i) {
-    //              // TODO(dsinclair): Implement
-    //            },
-    //            [&](const ast::InvariantAttribute* i) {
-    //              // TODO(dsinclair): Implement
-    //            },
-    //            [&](const ast::MustUseAttribute* i) {
-    //              // TODO(dsinclair): Implement
-    //            },
-    //            [&](const ast::IdAttribute*) {
-    //                add_error(attr->source,
-    //                          "found an `Id` attribute. The SubstituteOverrides transform "
-    //                          "must be run before converting to IR");
-    //            },
-    //            [&](const ast::StructMemberSizeAttribute*) {
-    //                TINT_ICE(IR, diagnostics_)
-    //                    << "StructMemberSizeAttribute encountered during IR conversion";
-    //            },
-    //            [&](const ast::StructMemberAlignAttribute*) {
-    //                TINT_ICE(IR, diagnostics_)
-    //                    << "StructMemberAlignAttribute encountered during IR conversion";
-    //            },
-    //            [&](const ast::StrideAttribute* s) {
-    //              // TODO(dsinclair): Implement
-    //            },
-    //            [&](const ast::InternalAttribute *i) {
-    //              // TODO(dsinclair): Implement
-    //            },
-    //            [&](Default) {
-    //                add_error(attr->source, "unknown attribute: " +
-    //                std::string(attr->TypeInfo().name));
-    //            });
-    //    }
 };
 
 }  // namespace

src/tint/ir/from_program_call_test.cc

@@ -35,17 +35,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Bitcast) {
     auto m = Build();
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():f32 {
+    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():f32 -> %fn2
-  %fn2 = block {
+%fn2 = block {
-  } -> %func_end 0.0f # return
+  br %fn3 0.0f  # return
-} %func_end
+}
+%fn3 = func_terminator
 
-%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
+%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5
-  %fn4 = block {
+%fn5 = block {
-    %1:f32 = call my_func
+  %1:f32 = call my_func
-    %tint_symbol:f32 = bitcast %1
+  %tint_symbol:f32 = bitcast %1
-  } -> %func_end # return
+  jmp %fn6  # return
-} %func_end
+}
+%fn6 = func_terminator
 
 )");
 }
@@ -60,11 +62,12 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Discard) {
     auto m = Build();
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func test_function():void [@fragment] {
+    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func test_function():void [@fragment] -> %fn2
-  %fn2 = block {
+%fn2 = block {
-    discard
+  discard
-  } -> %func_end # return
+  jmp %fn3  # return
-} %func_end
+}
+%fn3 = func_terminator
 
 )");
 }
@@ -77,16 +80,18 @@ TEST_F(IR_BuilderImplTest, EmitStatement_UserFunction) {
     auto m = Build();
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func(%p:f32):void {
+    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func(%p:f32):void -> %fn2
-  %fn2 = block {
+%fn2 = block {
-  } -> %func_end # return
+  jmp %fn3  # return
-} %func_end
+}
+%fn3 = func_terminator
 
-%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
+%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5
-  %fn4 = block {
+%fn5 = block {
-    %2:void = call my_func, 6.0f
+  %2:void = call my_func, 6.0f
-  } -> %func_end # return
+  jmp %fn6  # return
-} %func_end
+}
+%fn6 = func_terminator
 
 )");
 }
@@ -101,15 +106,18 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
   %i:ptr<private, i32, read_write> = var, 1i
+  br %fn2  # root_end
 }
 
+%fn2 = root_terminator
 
-%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
+%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4
-  %fn3 = block {
+%fn4 = block {
-    %2:i32 = load %i
+  %2:i32 = load %i
-    %tint_symbol:f32 = convert i32, %2
+  %tint_symbol:f32 = convert i32, %2
-  } -> %func_end # return
+  jmp %fn5  # return
-} %func_end
+}
+%fn5 = func_terminator
 
 )");
 }
@@ -123,8 +131,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_ConstructEmpty) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
   %i:ptr<private, vec3<f32>, read_write> = var, vec3<f32> 0.0f
+  br %fn2  # root_end
 }
 
+%fn2 = root_terminator
 
 )");
 }
@@ -139,15 +149,18 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
   %i:ptr<private, f32, read_write> = var, 1.0f
+  br %fn2  # root_end
 }
 
+%fn2 = root_terminator
 
-%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
+%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4
-  %fn3 = block {
+%fn4 = block {
-    %2:f32 = load %i
+  %2:f32 = load %i
-    %tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %2
+  %tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %2
-  } -> %func_end # return
+  jmp %fn5  # return
-} %func_end
+}
+%fn5 = func_terminator
 
 )");
 }

src/tint/ir/from_program_materialize_test.cc

@@ -34,10 +34,11 @@ TEST_F(IR_BuilderImplTest, EmitExpression_MaterializedCall) {
     auto m = Build();
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func test_function():f32 {
+    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func test_function():f32 -> %fn2
-  %fn2 = block {
+%fn2 = block {
-  } -> %func_end 2.0f # return
+  br %fn3 2.0f  # return
-} %func_end
+}
+%fn3 = func_terminator
 
 )");
 }

src/tint/ir/from_program_store_test.cc

@@ -37,14 +37,17 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Assign) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
   %a:ptr<private, u32, read_write> = var
+  br %fn2  # root_end
 }
 
+%fn2 = root_terminator
 
-%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
+%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4
-  %fn3 = block {
+%fn4 = block {
-    store %a, 4u
+  store %a, 4u
-  } -> %func_end # return
+  jmp %fn5  # return
-} %func_end
+}
+%fn5 = func_terminator
 
 )");
 }

src/tint/ir/from_program_unary_test.cc

@@ -34,17 +34,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Not) {
     auto m = Build();
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():bool {
+    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():bool -> %fn2
-  %fn2 = block {
+%fn2 = block {
-  } -> %func_end false # return
+  br %fn3 false  # return
-} %func_end
+}
+%fn3 = func_terminator
 
-%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
+%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5
-  %fn4 = block {
+%fn5 = block {
-    %1:bool = call my_func
+  %1:bool = call my_func
-    %tint_symbol:bool = eq %1, false
+  %tint_symbol:bool = eq %1, false
-  } -> %func_end # return
+  jmp %fn6  # return
-} %func_end
+}
+%fn6 = func_terminator
 
 )");
 }
@@ -57,17 +59,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Complement) {
     auto m = Build();
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 {
+    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2
-  %fn2 = block {
+%fn2 = block {
-  } -> %func_end 1u # return
+  br %fn3 1u  # return
-} %func_end
+}
+%fn3 = func_terminator
 
-%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
+%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5
-  %fn4 = block {
+%fn5 = block {
-    %1:u32 = call my_func
+  %1:u32 = call my_func
-    %tint_symbol:u32 = complement %1
+  %tint_symbol:u32 = complement %1
-  } -> %func_end # return
+  jmp %fn6  # return
-} %func_end
+}
+%fn6 = func_terminator
 
 )");
 }
@@ -80,17 +84,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Negation) {
     auto m = Build();
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
-    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():i32 {
+    EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():i32 -> %fn2
-  %fn2 = block {
+%fn2 = block {
-  } -> %func_end 1i # return
+  br %fn3 1i  # return
-} %func_end
+}
+%fn3 = func_terminator
 
-%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
+%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5
-  %fn4 = block {
+%fn5 = block {
-    %1:i32 = call my_func
+  %1:i32 = call my_func
-    %tint_symbol:i32 = negation %1
+  %tint_symbol:i32 = negation %1
-  } -> %func_end # return
+  jmp %fn6  # return
-} %func_end
+}
+%fn6 = func_terminator
 
 )");
 }
@@ -106,13 +112,16 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_AddressOf) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
   %v2:ptr<private, i32, read_write> = var
+  br %fn2  # root_end
 }
 
+%fn2 = root_terminator
 
-%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
+%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4
-  %fn3 = block {
+%fn4 = block {
-  } -> %func_end # return
+  jmp %fn5  # return
-} %func_end
+}
+%fn5 = func_terminator
 
 )");
 }
@@ -130,14 +139,17 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Indirection) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
   %v3:ptr<private, i32, read_write> = var
+  br %fn2  # root_end
 }
 
+%fn2 = root_terminator
 
-%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
+%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4
-  %fn3 = block {
+%fn4 = block {
-    store %v3, 42i
+  store %v3, 42i
-  } -> %func_end # return
+  jmp %fn5  # return
-} %func_end
+}
+%fn5 = func_terminator
 
 )");
 }

src/tint/ir/from_program_var_test.cc

@@ -34,8 +34,10 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_NoInit) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
   %a:ptr<private, u32, read_write> = var
+  br %fn2  # root_end
 }
 
+%fn2 = root_terminator
 
 )");
 }
@@ -49,8 +51,10 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_Init) {
 
     EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
   %a:ptr<private, u32, read_write> = var, 2u
+  br %fn2  # root_end
 }
 
+%fn2 = root_terminator
 
 )");
 }
@@ -63,11 +67,12 @@ TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()),
-              R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
+              R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2
-  %fn2 = block {
+%fn2 = block {
-    %a:ptr<function, u32, read_write> = var
+  %a:ptr<function, u32, read_write> = var
-  } -> %func_end # return
+  jmp %fn3  # return
-} %func_end
+}
+%fn3 = func_terminator
 
 )");
 }
@@ -81,11 +86,12 @@ TEST_F(IR_BuilderImplTest, Emit_Var_Init) {
     ASSERT_TRUE(m) << (!m ? m.Failure() : "");
 
     EXPECT_EQ(Disassemble(m.Get()),
-              R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
+              R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2
-  %fn2 = block {
+%fn2 = block {
-    %a:ptr<function, u32, read_write> = var, 2u
+  %a:ptr<function, u32, read_write> = var, 2u
-  } -> %func_end # return
+  jmp %fn3  # return
-} %func_end
+}
+%fn3 = func_terminator
 
 )");
 }

src/tint/ir/if.cc

@@ -18,7 +18,16 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::If);
 
 namespace tint::ir {
 
-If::If(Value* cond) : Base(), condition_(cond) {}
+If::If(Value* cond, Block* t, Block* f, Block* m)
+    : Base(m), condition_(cond), true_(t), false_(f), merge_(m) {
+    TINT_ASSERT(IR, true_);
+    TINT_ASSERT(IR, false_);
+    TINT_ASSERT(IR, merge_);
+
+    condition_->AddUsage(this);
+    true_->AddInboundBranch(this);
+    false_->AddInboundBranch(this);
+}
 
 If::~If() = default;
 

src/tint/ir/if.h

@@ -15,8 +15,8 @@
 #ifndef SRC_TINT_IR_IF_H_
 #define SRC_TINT_IR_IF_H_
 
+#include "src/tint/ir/block.h"
 #include "src/tint/ir/branch.h"
-#include "src/tint/ir/flow_node.h"
 #include "src/tint/ir/value.h"
 
 // Forward declarations
@@ -26,37 +26,42 @@ class Block;
 
 namespace tint::ir {
 
-/// A flow node representing an if statement.
+/// An if instruction
-class If : public utils::Castable<If, FlowNode> {
+class If : public utils::Castable<If, Branch> {
   public:
     /// Constructor
     /// @param cond the if condition
-    explicit If(Value* cond);
+    /// @param t the true block
+    /// @param f the false block
+    /// @param m the merge block
+    explicit If(Value* cond, Block* t, Block* f, Block* m);
     ~If() override;
 
     /// @returns the if condition
     const Value* Condition() const { return condition_; }
+    /// @returns the if condition
+    Value* Condition() { return condition_; }
 
     /// @returns the true branch block
-    const Branch& True() const { return true_; }
+    const Block* True() const { return true_; }
     /// @returns the true branch block
-    Branch& True() { return true_; }
+    Block* True() { return true_; }
 
     /// @returns the false branch block
-    const Branch& False() const { return false_; }
+    const Block* False() const { return false_; }
     /// @returns the false branch block
-    Branch& False() { return false_; }
+    Block* False() { return false_; }
 
     /// @returns the merge branch block
-    const Branch& Merge() const { return merge_; }
+    const Block* Merge() const { return merge_; }
     /// @returns the merge branch block
-    Branch& Merge() { return merge_; }
+    Block* Merge() { return merge_; }
 
   private:
-    Branch true_ = {};
+    Value* condition_ = nullptr;
-    Branch false_ = {};
+    Block* true_ = nullptr;
-    Branch merge_ = {};
+    Block* false_ = nullptr;
-    Value* condition_;
+    Block* merge_ = nullptr;
 };
 
 }  // namespace tint::ir

src/tint/ir/jump.cc

@@ -0,0 +1,25 @@
+// 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.
+
+#include "src/tint/ir/jump.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::Jump);
+
+namespace tint::ir {
+
+Jump::Jump(FlowNode* to, utils::VectorRef<Value*> args) : Base(to, args) {}
+
+Jump::~Jump() = default;
+
+}  // namespace tint::ir

src/tint/ir/jump.h

@@ -0,0 +1,37 @@
+// 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_JUMP_H_
+#define SRC_TINT_IR_JUMP_H_
+
+#include "src/tint/ir/block.h"
+#include "src/tint/ir/branch.h"
+#include "src/tint/ir/value.h"
+#include "src/tint/utils/castable.h"
+
+namespace tint::ir {
+
+/// A jump instruction. A jump is walk continuing.
+class Jump : public utils::Castable<Jump, Branch> {
+  public:
+    /// Constructor
+    /// @param to the block to branch too
+    /// @param args the branch arguments
+    explicit Jump(FlowNode* to, utils::VectorRef<Value*> args = {});
+    ~Jump() override;
+};
+
+}  // namespace tint::ir
+
+#endif  // SRC_TINT_IR_JUMP_H_

src/tint/ir/loop.cc

@@ -18,7 +18,11 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Loop);
 
 namespace tint::ir {
 
-Loop::Loop() : Base() {}
+Loop::Loop(Block* s, Block* c, Block* m) : Base(s), start_(s), continuing_(c), merge_(m) {
+    TINT_ASSERT(IR, start_);
+    TINT_ASSERT(IR, continuing_);
+    TINT_ASSERT(IR, merge_);
+}
 
 Loop::~Loop() = default;
 

src/tint/ir/loop.h

@@ -17,36 +17,38 @@
 
 #include "src/tint/ir/block.h"
 #include "src/tint/ir/branch.h"
-#include "src/tint/ir/flow_node.h"
 
 namespace tint::ir {
 
 /// Flow node describing a loop.
-class Loop : public utils::Castable<Loop, FlowNode> {
+class Loop : public utils::Castable<Loop, Branch> {
   public:
     /// Constructor
-    Loop();
+    /// @param s the start block
+    /// @param c the continuing block
+    /// @param m the merge block
+    Loop(Block* s, Block* c, Block* m);
     ~Loop() override;
 
     /// @returns the switch start branch
-    const Branch& Start() const { return start_; }
+    const Block* Start() const { return start_; }
     /// @returns the switch start branch
-    Branch& Start() { return start_; }
+    Block* Start() { return start_; }
 
     /// @returns the switch continuing branch
-    const Branch& Continuing() const { return continuing_; }
+    const Block* Continuing() const { return continuing_; }
     /// @returns the switch continuing branch
-    Branch& Continuing() { return continuing_; }
+    Block* Continuing() { return continuing_; }
 
     /// @returns the switch merge branch
-    const Branch& Merge() const { return merge_; }
+    const Block* Merge() const { return merge_; }
     /// @returns the switch merge branch
-    Branch& Merge() { return merge_; }
+    Block* Merge() { return merge_; }
 
   private:
-    Branch start_ = {};
+    Block* start_ = nullptr;
-    Branch continuing_ = {};
+    Block* continuing_ = nullptr;
-    Branch merge_ = {};
+    Block* merge_ = nullptr;
 };
 
 }  // namespace tint::ir

src/tint/ir/switch.cc

@@ -18,7 +18,11 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Switch);
 
 namespace tint::ir {
 
-Switch::Switch(Value* cond) : Base(), condition_(cond) {}
+Switch::Switch(Value* cond, Block* m) : Base(m), condition_(cond), merge_(m) {
+    TINT_ASSERT(IR, condition_);
+    TINT_ASSERT(IR, merge_);
+    condition_->AddUsage(this);
+}
 
 Switch::~Switch() = default;
 

src/tint/ir/switch.h

@@ -18,13 +18,12 @@
 #include "src/tint/ir/block.h"
 #include "src/tint/ir/branch.h"
 #include "src/tint/ir/constant.h"
-#include "src/tint/ir/flow_node.h"
 #include "src/tint/ir/value.h"
 
 namespace tint::ir {
 
 /// Flow node representing a switch statement
-class Switch : public utils::Castable<Switch, FlowNode> {
+class Switch : public utils::Castable<Switch, Branch> {
   public:
     /// A case selector
     struct CaseSelector {
@@ -40,23 +39,24 @@ class Switch : public utils::Castable<Switch, FlowNode> {
         /// The case selector for this node
         utils::Vector<CaseSelector, 4> selectors;
         /// The start block for the case block.
-        Branch start = {};
+        Block* start = nullptr;
 
         /// @returns the case start target
-        const Branch& Start() const { return start; }
+        const Block* Start() const { return start; }
         /// @returns the case start target
-        Branch& Start() { return start; }
+        Block* Start() { return start; }
     };
 
     /// Constructor
     /// @param cond the condition
-    explicit Switch(Value* cond);
+    /// @param m the merge block
+    explicit Switch(Value* cond, Block* m);
     ~Switch() override;
 
     /// @returns the switch merge branch
-    const Branch& Merge() const { return merge_; }
+    const Block* Merge() const { return merge_; }
     /// @returns the switch merge branch
-    Branch& Merge() { return merge_; }
+    Block* Merge() { return merge_; }
 
     /// @returns the switch cases
     utils::VectorRef<Case> Cases() const { return cases_; }
@@ -65,11 +65,13 @@ class Switch : public utils::Castable<Switch, FlowNode> {
 
     /// @returns the condition
     const Value* Condition() const { return condition_; }
+    /// @returns the condition
+    Value* Condition() { return condition_; }
 
   private:
-    Branch merge_ = {};
+    Value* condition_ = nullptr;
+    Block* merge_ = nullptr;
     utils::Vector<Case, 4> cases_;
-    Value* condition_;
 };
 
 }  // namespace tint::ir

src/tint/ir/to_program.cc

@@ -23,6 +23,7 @@
 #include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/if.h"
 #include "src/tint/ir/instruction.h"
+#include "src/tint/ir/jump.h"
 #include "src/tint/ir/load.h"
 #include "src/tint/ir/module.h"
 #include "src/tint/ir/store.h"
@@ -108,25 +109,26 @@ class State {
                       std::move(ret_attrs));
     }
 
-    const ast::BlockStatement* FlowNodeGraph(ir::FlowNode* start_node,
+    const ast::BlockStatement* FlowNodeGraph(const ir::Block* start_node) {
-                                             ir::FlowNode* stop_at = nullptr) {
         // TODO(crbug.com/tint/1902): Check if the block is dead
         utils::Vector<const ast::Statement*,
                       decltype(ast::BlockStatement::statements)::static_length>
             stmts;
 
-        ir::Branch root_branch{start_node, {}};
+        const ir::FlowNode* block = start_node;
-        const ir::Branch* branch = &root_branch;
 
         // TODO(crbug.com/tint/1902): Handle block arguments.
 
-        while (branch->target != stop_at) {
+        while (block) {
-            enum Status { kContinue, kStop, kError };
+            TINT_ASSERT(IR, block->HasBranchTarget());
-            Status status = tint::Switch(
-                branch->target,
 
-                [&](const ir::Block* block) {
+            enum Status { kContinue, kStop, kError };
-                    for (const auto* inst : block->Instructions()) {
+
+            Status status = tint::Switch(
+                block,
+
+                [&](const ir::Block* blk) {
+                    for (auto* inst : blk->Instructions()) {
                         auto stmt = Stmt(inst);
                         if (TINT_UNLIKELY(!stmt)) {
                             return kError;
@@ -135,43 +137,27 @@ class State {
                             stmts.Push(s);
                         }
                     }
-                    branch = &block->Branch();
+                    if (blk->Branch()->Is<Jump>() && blk->Branch()->To()->Is<Block>()) {
-                    return kContinue;
+                        block = blk->Branch()->To()->As<Block>();
-                },
+                        return kContinue;
-
+                    } else if (auto* if_ = blk->Branch()->As<ir::If>()) {
-                [&](const ir::If* if_) {
+                        if (if_->Merge()->HasBranchTarget()) {
-                    auto* stmt = If(if_);
+                            block = if_->Merge();
-                    if (TINT_UNLIKELY(!stmt)) {
+                            return kContinue;
-                        return kError;
+                        }
-                    }
+                    } else if (auto* switch_ = blk->Branch()->As<ir::Switch>()) {
-                    stmts.Push(stmt);
+                        if (switch_->Merge()->HasBranchTarget()) {
-                    branch = &if_->Merge();
+                            block = switch_->Merge();
-                    return branch->target->InboundBranches().IsEmpty() ? kStop : kContinue;
+                            return kContinue;
-                },
+                        }
-
-                [&](const ir::Switch* switch_) {
-                    auto* stmt = Switch(switch_);
-                    if (TINT_UNLIKELY(!stmt)) {
-                        return kError;
-                    }
-                    stmts.Push(stmt);
-                    branch = &switch_->Merge();
-                    return branch->target->InboundBranches().IsEmpty() ? kStop : kContinue;
-                },
-
-                [&](const ir::FunctionTerminator*) {
-                    auto res = FunctionTerminator(branch);
-                    if (TINT_UNLIKELY(!res)) {
-                        return kError;
-                    }
-                    if (auto* stmt = res.Get()) {
-                        stmts.Push(stmt);
                     }
                     return kStop;
                 },
 
+                [&](const ir::FunctionTerminator*) { return kStop; },
+
                 [&](Default) {
-                    UNHANDLED_CASE(branch->target);
+                    UNHANDLED_CASE(block);
                     return kError;
                 });
 
@@ -188,26 +174,24 @@ class State {
 
     const ast::IfStatement* If(const ir::If* i) {
         SCOPED_NESTING();
-
         auto* cond = Expr(i->Condition());
-        auto* t = FlowNodeGraph(i->True().target, i->Merge().target);
+        auto* t = FlowNodeGraph(i->True());
         if (TINT_UNLIKELY(!t)) {
             return nullptr;
         }
 
-        if (!IsEmpty(i->False().target, i->Merge().target)) {
+        if (!IsEmpty(i->False(), i->Merge())) {
-            // If the else target is an if flow node with the same Merge().target as this if, then
+            // If the else target is an `if` which has a merge target that just bounces to the outer
-            // emit an 'else if' instead of a block statement for the else.
+            // if merge target then emit an 'else if' instead of a block statement for the else.
-            if (auto* else_if = As<ir::If>(NextNonEmptyNode(i->False().target));
+            if (auto* inst = i->False()->Instructions().Front()->As<ir::If>();
-                else_if &&
+                inst && inst->Merge()->IsTrampoline(i->Merge())) {
-                NextNonEmptyNode(i->Merge().target) == NextNonEmptyNode(else_if->Merge().target)) {
+                auto* f = If(inst);
-                auto* f = If(else_if);
                 if (!f) {
                     return nullptr;
                 }
                 return b.If(cond, t, b.Else(f));
             } else {
-                auto* f = FlowNodeGraph(i->False().target, i->Merge().target);
+                auto* f = FlowNodeGraph(i->False());
                 if (!f) {
                     return nullptr;
                 }
@@ -226,11 +210,11 @@ class State {
             return nullptr;
         }
 
-        auto cases = utils::Transform<1>(
+        auto cases = utils::Transform<2>(
             s->Cases(),  //
-            [&](const ir::Switch::Case& c) -> const tint::ast::CaseStatement* {
+            [&](const ir::Switch::Case c) -> const tint::ast::CaseStatement* {
                 SCOPED_NESTING();
-                auto* body = FlowNodeGraph(c.start.target, s->Merge().target);
+                auto* body = FlowNodeGraph(c.start);
                 if (!body) {
                     return nullptr;
                 }
@@ -261,26 +245,27 @@ class State {
     }
 
     utils::Result<const ast::ReturnStatement*> FunctionTerminator(const ir::Branch* branch) {
-        if (branch->args.IsEmpty()) {
+        if (branch->Args().IsEmpty()) {
             // Branch to function terminator has no arguments.
-            // If this block is nested withing some control flow, then we must emit a
+            // If this block is nested withing some control flow, then we must
-            // 'return' statement, otherwise we've just naturally reached the end of the
+            // emit a 'return' statement, otherwise we've just naturally reached
-            // function where the 'return' is redundant.
+            // the end of the function where the 'return' is redundant.
             if (nesting_depth_ > 1) {
                 return b.Return();
             }
             return nullptr;
         }
 
-        // Branch to function terminator has arguments - this is the return value.
+        // Branch to function terminator has arguments - this is the return
-        if (branch->args.Length() != 1) {
+        // value.
-            TINT_ICE(IR, b.Diagnostics())
+        if (branch->Args().Length() != 1) {
-                << "expected 1 value for function terminator (return value), got "
+            TINT_ICE(IR, b.Diagnostics()) << "expected 1 value for function "
-                << branch->args.Length();
+                                             "terminator (return value), got "
+                                          << branch->Args().Length();
             return utils::Failure;
         }
 
-        auto* val = Expr(branch->args.Front());
+        auto* val = Expr(branch->Args().Front());
         if (TINT_UNLIKELY(!val)) {
             return utils::Failure;
         }
@@ -289,36 +274,16 @@ class State {
     }
 
     /// @return true if there are no instructions between @p node and and @p stop_at
-    bool IsEmpty(const ir::FlowNode* node, const ir::FlowNode* stop_at) {
+    bool IsEmpty(const ir::Block* node, const ir::FlowNode* stop_at) {
-        while (node != stop_at) {
+        if (node->Instructions().IsEmpty()) {
-            if (auto* block = node->As<ir::Block>()) {
+            return true;
-                if (!block->Instructions().IsEmpty()) {
-                    return false;
-                }
-                node = block->Branch().target;
-            } else {
-                return false;
-            }
         }
-        return true;
+        if (auto* br = node->Instructions().Front()->As<Branch>()) {
-    }
+            return br->To() == stop_at;
-
-    /// @return the next flow node that isn't an empty block
-    const ir::FlowNode* NextNonEmptyNode(const ir::FlowNode* node) {
-        while (node) {
-            if (auto* block = node->As<ir::Block>()) {
-                for (const auto* inst : block->Instructions()) {
-                    // Load instructions will be inlined, so ignore them.
-                    if (!inst->Is<ir::Load>()) {
-                        return node;
-                    }
-                }
-                node = block->Branch().target;
-            } else {
-                return node;
-            }
         }
-        return nullptr;
+        // TODO(dsinclair): This should possibly walk over Jump instructions that
+        // just jump to empty blocks if we want to be comprehensive.
+        return false;
     }
 
     utils::Result<const ast::Statement*> Stmt(const ir::Instruction* inst) {
@@ -328,6 +293,14 @@ class State {
             [&](const ir::Var* i) { return Var(i); },        //
             [&](const ir::Load*) { return nullptr; },
             [&](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) {
+                if (branch->To()->Is<ir::FunctionTerminator>()) {
+                    return utils::Result<const ast::Statement*>{FunctionTerminator(branch)};
+                }
+                return utils::Result<const ast::Statement*>{nullptr};
+            },
             [&](Default) {
                 UNHANDLED_CASE(inst);
                 return utils::Failure;

src/tint/ir/to_program_roundtrip_test.cc

@@ -229,10 +229,9 @@ fn c() {
 
 fn f() {
   var cond_a : bool = true;
-  var cond_b : bool = true;
   if (cond_a) {
     a();
-  } else if (cond_b) {
+  } else if (false) {
     b();
   }
   c();

src/tint/ir/transform/add_empty_entry_point.cc

@@ -38,7 +38,7 @@ void AddEmptyEntryPoint::Run(ir::Module* ir, const DataMap&, DataMap&) const {
     auto* ep =
         builder.CreateFunction(ir->symbols.New("unused_entry_point"), ir->types.Get<type::Void>(),
                                Function::PipelineStage::kCompute, std::array{1u, 1u, 1u});
-    ep->StartTarget()->BranchTo(ep->EndTarget());
+    ep->StartTarget()->SetInstructions(utils::Vector{builder.Branch(ep->EndTarget())});
     ir->functions.Push(ep);
 }
 

src/tint/ir/transform/add_empty_entry_point_test.cc

@@ -25,10 +25,11 @@ using IR_AddEmptyEntryPointTest = TransformTest;
 
 TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) {
     auto* expect = R"(
-%fn1 = func unused_entry_point():void [@compute @workgroup_size(1, 1, 1)] {
+%fn1 = func unused_entry_point():void [@compute @workgroup_size(1, 1, 1)] -> %fn2
-  %fn2 = block {
+%fn2 = block {
-  } -> %func_end # return
+  br %fn3  # return
-} %func_end
+}
+%fn3 = func_terminator
 
 )";
 
@@ -40,14 +41,15 @@ TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) {
 TEST_F(IR_AddEmptyEntryPointTest, ExistingEntryPoint) {
     auto* ep = b.CreateFunction(mod.symbols.New("main"), mod.types.Get<type::Void>(),
                                 Function::PipelineStage::kFragment);
-    ep->StartTarget()->BranchTo(ep->EndTarget());
+    ep->StartTarget()->SetInstructions(utils::Vector{b.Branch(ep->EndTarget())});
     mod.functions.Push(ep);
 
     auto* expect = R"(
-%fn1 = func main():void [@fragment] {
+%fn1 = func main():void [@fragment] -> %fn2
-  %fn2 = block {
+%fn2 = block {
-  } -> %func_end # return
+  br %fn3  # return
-} %func_end
+}
+%fn3 = func_terminator
 
 )";
 

src/tint/ir/unary.h

@@ -41,6 +41,8 @@ class Unary : public utils::Castable<Unary, Instruction> {
 
     /// @returns the value for the instruction
     const Value* Val() const { return val_; }
+    /// @returns the value for the instruction
+    Value* Val() { return val_; }
 
     /// @returns the kind of unary instruction
     enum Kind Kind() const { return kind_; }

src/tint/ir/unary_test.cc

@@ -26,7 +26,7 @@ using IR_InstructionTest = TestHelper;
 TEST_F(IR_InstructionTest, CreateComplement) {
     Module mod;
     Builder b{mod};
-    const auto* inst = b.Complement(b.ir.types.Get<type::I32>(), b.Constant(4_i));
+    auto* inst = b.Complement(b.ir.types.Get<type::I32>(), b.Constant(4_i));
 
     ASSERT_TRUE(inst->Is<Unary>());
     EXPECT_EQ(inst->Kind(), Unary::Kind::kComplement);
@@ -40,7 +40,7 @@ TEST_F(IR_InstructionTest, CreateComplement) {
 TEST_F(IR_InstructionTest, CreateNegation) {
     Module mod;
     Builder b{mod};
-    const auto* inst = b.Negation(b.ir.types.Get<type::I32>(), b.Constant(4_i));
+    auto* inst = b.Negation(b.ir.types.Get<type::I32>(), b.Constant(4_i));
 
     ASSERT_TRUE(inst->Is<Unary>());
     EXPECT_EQ(inst->Kind(), Unary::Kind::kNegation);
@@ -54,7 +54,7 @@ TEST_F(IR_InstructionTest, CreateNegation) {
 TEST_F(IR_InstructionTest, Unary_Usage) {
     Module mod;
     Builder b{mod};
-    const auto* inst = b.Negation(b.ir.types.Get<type::I32>(), b.Constant(4_i));
+    auto* inst = b.Negation(b.ir.types.Get<type::I32>(), b.Constant(4_i));
 
     EXPECT_EQ(inst->Kind(), Unary::Kind::kNegation);
 

src/tint/transform/manager_test.cc

@@ -52,7 +52,7 @@ class IR_AddFunction final : public ir::transform::Transform {
         ir::Builder builder(*mod);
         auto* func =
             builder.CreateFunction(mod->symbols.New("ir_func"), mod->types.Get<type::Void>());
-        func->StartTarget()->BranchTo(func->EndTarget());
+        func->StartTarget()->SetInstructions(utils::Vector{builder.Branch(func->EndTarget())});
         mod->functions.Push(func);
     }
 };
@@ -70,7 +70,7 @@ ir::Module MakeIR() {
     ir::Builder builder(mod);
     auto* func =
         builder.CreateFunction(builder.ir.symbols.New("main"), builder.ir.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
+    func->StartTarget()->SetInstructions(utils::Vector{builder.Branch(func->EndTarget())});
     builder.ir.functions.Push(func);
     return mod;
 }

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

@@ -292,8 +292,15 @@ void GeneratorImplIr::EmitBlock(const ir::Block* block) {
         current_function_.push_inst(spv::Op::OpLabel, {Label(block)});
     }
 
+    // If there are no instructions in the block, it's a dead end, so we shouldn't be able to get
+    // here to begin with.
+    if (block->Instructions().IsEmpty()) {
+        current_function_.push_inst(spv::Op::OpUnreachable, {});
+        return;
+    }
+
     // Emit the instructions.
-    for (const auto* inst : block->Instructions()) {
+    for (auto* inst : block->Instructions()) {
         auto result = Switch(
             inst,  //
             [&](const ir::Binary* b) { return EmitBinary(b); },
@@ -303,6 +310,14 @@ void GeneratorImplIr::EmitBlock(const ir::Block* block) {
                 return 0u;
             },
             [&](const ir::Var* v) { return EmitVar(v); },
+            [&](const ir::If* i) {
+                EmitIf(i);
+                return 0u;
+            },
+            [&](const ir::Branch* b) {
+                EmitBranch(b);
+                return 0u;
+            },
             [&](Default) {
                 TINT_ICE(Writer, diagnostics_)
                     << "unimplemented instruction: " << inst->TypeInfo().name;
@@ -310,46 +325,42 @@ void GeneratorImplIr::EmitBlock(const ir::Block* block) {
             });
         instructions_.Add(inst, result);
     }
+}
 
-    // Handle the branch at the end of the block.
+void GeneratorImplIr::EmitBranch(const ir::Branch* b) {
     Switch(
-        block->Branch().target,
+        b->To(),
-        [&](const ir::Block* b) { current_function_.push_inst(spv::Op::OpBranch, {Label(b)}); },
+        [&](const ir::Block* blk) { current_function_.push_inst(spv::Op::OpBranch, {Label(blk)}); },
-        [&](const ir::If* i) { EmitIf(i); },
         [&](const ir::FunctionTerminator*) {
             // TODO(jrprice): Handle the return value, which will be a branch argument.
-            if (!block->Branch().args.IsEmpty()) {
+            if (!b->Args().IsEmpty()) {
                 TINT_ICE(Writer, diagnostics_) << "unimplemented return value";
             }
             current_function_.push_inst(spv::Op::OpReturn, {});
         },
         [&](Default) {
-            if (!block->Branch().target) {
+            // A block may not have an outward branch (e.g. an unreachable merge
-                // A block may not have an outward branch (e.g. an unreachable merge block).
+            // block).
-                current_function_.push_inst(spv::Op::OpUnreachable, {});
+            current_function_.push_inst(spv::Op::OpUnreachable, {});
-            } else {
-                TINT_ICE(Writer, diagnostics_)
-                    << "unimplemented branch target: " << block->Branch().target->TypeInfo().name;
-            }
         });
 }
 
 void GeneratorImplIr::EmitIf(const ir::If* i) {
-    auto* merge_block = i->Merge().target->As<ir::Block>();
+    auto* merge_block = i->Merge();
-    auto* true_block = i->True().target->As<ir::Block>();
+    auto* true_block = i->True();
-    auto* false_block = i->False().target->As<ir::Block>();
+    auto* false_block = i->False();
 
     // Generate labels for the blocks. We emit the true or false block if it:
-    // 1. contains instructions, or
+    // 1. contains instructions other then the branch, or
-    // 2. branches somewhere other then the Merge().target.
+    // 2. branches somewhere other then the Merge().
     // Otherwise we skip them and branch straight to the merge block.
     uint32_t merge_label = Label(merge_block);
     uint32_t true_label = merge_label;
     uint32_t false_label = merge_label;
-    if (!true_block->Instructions().IsEmpty() || true_block->Branch().target != merge_block) {
+    if (true_block->Instructions().Length() > 1 || true_block->Branch()->To() != merge_block) {
         true_label = Label(true_block);
     }
-    if (!false_block->Instructions().IsEmpty() || false_block->Branch().target != merge_block) {
+    if (false_block->Instructions().Length() > 1 || false_block->Branch()->To() != merge_block) {
         false_label = Label(false_block);
     }
 

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

@@ -30,6 +30,7 @@
 namespace tint::ir {
 class Binary;
 class Block;
+class Branch;
 class If;
 class Function;
 class Load;
@@ -121,6 +122,10 @@ class GeneratorImplIr {
     /// @returns the result ID of the instruction
     uint32_t EmitVar(const ir::Var* var);
 
+    /// Emit a branch instruction.
+    /// @param b the branch instruction to emit
+    void EmitBranch(const ir::Branch* b);
+
   private:
     /// Get the result ID of the constant `constant`, emitting its instruction if necessary.
     /// @param constant the constant to get the ID for

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

@@ -21,10 +21,9 @@ namespace {
 
 TEST_F(SpvGeneratorImplTest, Binary_Add_I32) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
-
     func->StartTarget()->SetInstructions(
-        utils::Vector{b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(2_i))});
+        utils::Vector{b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(2_i)),
+                      b.Branch(func->EndTarget())});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -43,10 +42,9 @@ OpFunctionEnd
 
 TEST_F(SpvGeneratorImplTest, Binary_Add_U32) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
-
     func->StartTarget()->SetInstructions(
-        utils::Vector{b.Add(mod.types.Get<type::U32>(), b.Constant(1_u), b.Constant(2_u))});
+        utils::Vector{b.Add(mod.types.Get<type::U32>(), b.Constant(1_u), b.Constant(2_u)),
+                      b.Branch(func->EndTarget())});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -65,10 +63,9 @@ OpFunctionEnd
 
 TEST_F(SpvGeneratorImplTest, Binary_Add_F32) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
-
     func->StartTarget()->SetInstructions(
-        utils::Vector{b.Add(mod.types.Get<type::F32>(), b.Constant(1_f), b.Constant(2_f))});
+        utils::Vector{b.Add(mod.types.Get<type::F32>(), b.Constant(1_f), b.Constant(2_f)),
+                      b.Branch(func->EndTarget())});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -87,10 +84,9 @@ OpFunctionEnd
 
 TEST_F(SpvGeneratorImplTest, Binary_Sub_I32) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
-
     func->StartTarget()->SetInstructions(
-        utils::Vector{b.Subtract(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(2_i))});
+        utils::Vector{b.Subtract(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(2_i)),
+                      b.Branch(func->EndTarget())});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -109,10 +105,9 @@ OpFunctionEnd
 
 TEST_F(SpvGeneratorImplTest, Binary_Sub_U32) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
-
     func->StartTarget()->SetInstructions(
-        utils::Vector{b.Subtract(mod.types.Get<type::U32>(), b.Constant(1_u), b.Constant(2_u))});
+        utils::Vector{b.Subtract(mod.types.Get<type::U32>(), b.Constant(1_u), b.Constant(2_u)),
+                      b.Branch(func->EndTarget())});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -131,10 +126,9 @@ OpFunctionEnd
 
 TEST_F(SpvGeneratorImplTest, Binary_Sub_F32) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
-
     func->StartTarget()->SetInstructions(
-        utils::Vector{b.Subtract(mod.types.Get<type::F32>(), b.Constant(1_f), b.Constant(2_f))});
+        utils::Vector{b.Subtract(mod.types.Get<type::F32>(), b.Constant(1_f), b.Constant(2_f)),
+                      b.Branch(func->EndTarget())});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -153,8 +147,6 @@ OpFunctionEnd
 
 TEST_F(SpvGeneratorImplTest, Binary_Sub_Vec2i) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
-
     auto* lhs = mod.constants_arena.Create<constant::Composite>(
         mod.types.Get<type::Vector>(mod.types.Get<type::I32>(), 2u),
         utils::Vector{b.Constant(42_i)->Value(), b.Constant(-1_i)->Value()}, false, false);
@@ -163,7 +155,8 @@ TEST_F(SpvGeneratorImplTest, Binary_Sub_Vec2i) {
         utils::Vector{b.Constant(0_i)->Value(), b.Constant(-43_i)->Value()}, false, false);
     func->StartTarget()->SetInstructions(
         utils::Vector{b.Subtract(mod.types.Get<type::Vector>(mod.types.Get<type::I32>(), 2u),
-                                 b.Constant(lhs), b.Constant(rhs))});
+                                 b.Constant(lhs), b.Constant(rhs)),
+                      b.Branch(func->EndTarget())});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -187,8 +180,6 @@ OpFunctionEnd
 
 TEST_F(SpvGeneratorImplTest, Binary_Sub_Vec4f) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
-
     auto* lhs = mod.constants_arena.Create<constant::Composite>(
         mod.types.Get<type::Vector>(mod.types.Get<type::F32>(), 4u),
         utils::Vector{b.Constant(42_f)->Value(), b.Constant(-1_f)->Value(),
@@ -201,7 +192,8 @@ TEST_F(SpvGeneratorImplTest, Binary_Sub_Vec4f) {
         false, false);
     func->StartTarget()->SetInstructions(
         utils::Vector{b.Subtract(mod.types.Get<type::Vector>(mod.types.Get<type::F32>(), 4u),
-                                 b.Constant(lhs), b.Constant(rhs))});
+                                 b.Constant(lhs), b.Constant(rhs)),
+                      b.Branch(func->EndTarget())});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -227,10 +219,9 @@ OpFunctionEnd
 
 TEST_F(SpvGeneratorImplTest, Binary_Chain) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
-
     auto* a = b.Subtract(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(2_i));
-    func->StartTarget()->SetInstructions(utils::Vector{a, b.Add(mod.types.Get<type::I32>(), a, a)});
+    func->StartTarget()->SetInstructions(
+        utils::Vector{a, b.Add(mod.types.Get<type::I32>(), a, a), b.Branch(func->EndTarget())});
 
     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(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
+    func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
 
     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(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
+    func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
 
     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(mod.symbols.Register("main"), mod.types.Get<type::Void>(),
                                   ir::Function::PipelineStage::kCompute, {{32, 4, 1}});
-    func->StartTarget()->BranchTo(func->EndTarget());
+    func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
 
     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(mod.symbols.Register("main"), mod.types.Get<type::Void>(),
                                   ir::Function::PipelineStage::kFragment);
-    func->StartTarget()->BranchTo(func->EndTarget());
+    func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
 
     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(mod.symbols.Register("main"), mod.types.Get<type::Void>(),
                                   ir::Function::PipelineStage::kVertex);
-    func->StartTarget()->BranchTo(func->EndTarget());
+    func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
 
     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(mod.symbols.Register("main1"), mod.types.Get<type::Void>(),
                                 ir::Function::PipelineStage::kCompute, {{32, 4, 1}});
-    f1->StartTarget()->BranchTo(f1->EndTarget());
+    f1->StartTarget()->SetInstructions(utils::Vector{b.Branch(f1->EndTarget())});
 
     auto* f2 = b.CreateFunction(mod.symbols.Register("main2"), mod.types.Get<type::Void>(),
                                 ir::Function::PipelineStage::kCompute, {{8, 2, 16}});
-    f2->StartTarget()->BranchTo(f2->EndTarget());
+    f2->StartTarget()->SetInstructions(utils::Vector{b.Branch(f2->EndTarget())});
 
     auto* f3 = b.CreateFunction(mod.symbols.Register("main3"), mod.types.Get<type::Void>(),
                                 ir::Function::PipelineStage::kFragment);
-    f3->StartTarget()->BranchTo(f3->EndTarget());
+    f3->StartTarget()->SetInstructions(utils::Vector{b.Branch(f3->EndTarget())});
 
     generator_.EmitFunction(f1);
     generator_.EmitFunction(f2);

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

@@ -23,11 +23,11 @@ TEST_F(SpvGeneratorImplTest, If_TrueEmpty_FalseEmpty) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
 
     auto* i = b.CreateIf(b.Constant(true));
-    i->True().target->As<ir::Block>()->BranchTo(i->Merge().target);
+    i->True()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
-    i->False().target->As<ir::Block>()->BranchTo(i->Merge().target);
+    i->False()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
-    i->Merge().target->As<ir::Block>()->BranchTo(func->EndTarget());
+    i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
 
-    func->StartTarget()->BranchTo(i);
+    func->StartTarget()->SetInstructions(utils::Vector{i});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -49,15 +49,14 @@ TEST_F(SpvGeneratorImplTest, If_FalseEmpty) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
 
     auto* i = b.CreateIf(b.Constant(true));
-    i->False().target->As<ir::Block>()->BranchTo(i->Merge().target);
+    i->False()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
-    i->Merge().target->As<ir::Block>()->BranchTo(func->EndTarget());
+    i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
 
-    auto* true_block = i->True().target->As<ir::Block>();
+    auto* true_block = i->True();
-    true_block->SetInstructions(
+    true_block->SetInstructions(utils::Vector{
-        utils::Vector{b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(1_i))});
+        b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(1_i)), b.Branch(i->Merge())});
-    true_block->BranchTo(i->Merge().target);
 
-    func->StartTarget()->BranchTo(i);
+    func->StartTarget()->SetInstructions(utils::Vector{i});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -84,15 +83,14 @@ TEST_F(SpvGeneratorImplTest, If_TrueEmpty) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
 
     auto* i = b.CreateIf(b.Constant(true));
-    i->True().target->As<ir::Block>()->BranchTo(i->Merge().target);
+    i->True()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
-    i->Merge().target->As<ir::Block>()->BranchTo(func->EndTarget());
+    i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
 
-    auto* false_block = i->False().target->As<ir::Block>();
+    auto* false_block = i->False();
-    false_block->SetInstructions(
+    false_block->SetInstructions(utils::Vector{
-        utils::Vector{b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(1_i))});
+        b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(1_i)), b.Branch(i->Merge())});
-    false_block->BranchTo(i->Merge().target);
 
-    func->StartTarget()->BranchTo(i);
+    func->StartTarget()->SetInstructions(utils::Vector{i});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -119,11 +117,10 @@ TEST_F(SpvGeneratorImplTest, If_BothBranchesReturn) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
 
     auto* i = b.CreateIf(b.Constant(true));
-    i->True().target->As<ir::Block>()->BranchTo(func->EndTarget());
+    i->True()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
-    i->False().target->As<ir::Block>()->BranchTo(func->EndTarget());
+    i->False()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
-    i->Merge().target->As<ir::Block>()->BranchTo(nullptr);
 
-    func->StartTarget()->BranchTo(i);
+    func->StartTarget()->SetInstructions(utils::Vector{i});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"

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

@@ -22,11 +22,10 @@ namespace {
 
 TEST_F(SpvGeneratorImplTest, FunctionVar_NoInit) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
 
     auto* ty = mod.types.Get<type::Pointer>(
         mod.types.Get<type::I32>(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
-    func->StartTarget()->SetInstructions(utils::Vector{b.Declare(ty)});
+    func->StartTarget()->SetInstructions(utils::Vector{b.Declare(ty), b.Branch(func->EndTarget())});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -44,14 +43,13 @@ OpFunctionEnd
 
 TEST_F(SpvGeneratorImplTest, FunctionVar_WithInit) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
 
     auto* ty = mod.types.Get<type::Pointer>(
         mod.types.Get<type::I32>(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
     auto* v = b.Declare(ty);
     v->SetInitializer(b.Constant(42_i));
 
-    func->StartTarget()->SetInstructions(utils::Vector{v});
+    func->StartTarget()->SetInstructions(utils::Vector{v, b.Branch(func->EndTarget())});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -71,12 +69,11 @@ OpFunctionEnd
 
 TEST_F(SpvGeneratorImplTest, FunctionVar_Name) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
 
     auto* ty = mod.types.Get<type::Pointer>(
         mod.types.Get<type::I32>(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
     auto* v = b.Declare(ty);
-    func->StartTarget()->SetInstructions(utils::Vector{v});
+    func->StartTarget()->SetInstructions(utils::Vector{v, b.Branch(func->EndTarget())});
     mod.SetName(v, "myvar");
 
     generator_.EmitFunction(func);
@@ -96,7 +93,6 @@ OpFunctionEnd
 
 TEST_F(SpvGeneratorImplTest, FunctionVar_DeclInsideBlock) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
 
     auto* ty = mod.types.Get<type::Pointer>(
         mod.types.Get<type::I32>(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
@@ -104,14 +100,11 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_DeclInsideBlock) {
     v->SetInitializer(b.Constant(42_i));
 
     auto* i = b.CreateIf(b.Constant(true));
-    i->False().target->As<ir::Block>()->BranchTo(func->EndTarget());
+    i->True()->SetInstructions(utils::Vector{v, b.Branch(i->Merge())});
-    i->Merge().target->As<ir::Block>()->BranchTo(func->EndTarget());
+    i->False()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
+    i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())});
 
-    auto* true_block = i->True().target->As<ir::Block>();
+    func->StartTarget()->SetInstructions(utils::Vector{i});
-    true_block->SetInstructions(utils::Vector{v});
-    true_block->BranchTo(i->Merge().target);
-
-    func->StartTarget()->BranchTo(i);
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -140,13 +133,12 @@ OpFunctionEnd
 
 TEST_F(SpvGeneratorImplTest, FunctionVar_Load) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
 
     auto* store_ty = mod.types.Get<type::I32>();
     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)});
+    func->StartTarget()->SetInstructions(utils::Vector{v, b.Load(v), b.Branch(func->EndTarget())});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -165,12 +157,12 @@ OpFunctionEnd
 
 TEST_F(SpvGeneratorImplTest, FunctionVar_Store) {
     auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
-    func->StartTarget()->BranchTo(func->EndTarget());
 
     auto* ty = mod.types.Get<type::Pointer>(
         mod.types.Get<type::I32>(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
     auto* v = b.Declare(ty);
-    func->StartTarget()->SetInstructions(utils::Vector{v, b.Store(v, b.Constant(42_i))});
+    func->StartTarget()->SetInstructions(
+        utils::Vector{v, b.Store(v, b.Constant(42_i)), b.Branch(func->EndTarget())});
 
     generator_.EmitFunction(func);
     EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"