diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 344607d825..bff6c37886 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -1230,12 +1230,16 @@ if (tint_build_ir) {
       "ir/continue.h",
       "ir/convert.cc",
       "ir/convert.h",
-      "ir/debug.cc",
-      "ir/debug.h",
       "ir/disassembler.cc",
       "ir/disassembler.h",
       "ir/discard.cc",
       "ir/discard.h",
+      "ir/exit_if.cc",
+      "ir/exit_if.h",
+      "ir/exit_loop.cc",
+      "ir/exit_loop.h",
+      "ir/exit_switch.cc",
+      "ir/exit_switch.h",
       "ir/function.cc",
       "ir/function.h",
       "ir/function_param.cc",
@@ -1252,8 +1256,6 @@ if (tint_build_ir) {
       "ir/module.h",
       "ir/return.cc",
       "ir/return.h",
-      "ir/root_terminator.cc",
-      "ir/root_terminator.h",
       "ir/store.cc",
       "ir/store.h",
       "ir/switch.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index ff69240998..e5e529b10b 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -738,12 +738,16 @@ if(${TINT_BUILD_IR})
     ir/continue.h
     ir/convert.cc
     ir/convert.h
-    ir/debug.cc
-    ir/debug.h
     ir/disassembler.cc
     ir/disassembler.h
     ir/discard.cc
     ir/discard.h
+    ir/exit_if.cc
+    ir/exit_if.h
+    ir/exit_loop.cc
+    ir/exit_loop.h
+    ir/exit_switch.cc
+    ir/exit_switch.h
     ir/from_program.cc
     ir/from_program.h
     ir/function.cc
@@ -762,8 +766,6 @@ if(${TINT_BUILD_IR})
     ir/module.h
     ir/return.cc
     ir/return.h
-    ir/root_terminator.cc
-    ir/root_terminator.h
     ir/store.cc
     ir/store.h
     ir/switch.cc
diff --git a/src/tint/cmd/main.cc b/src/tint/cmd/main.cc
index bfdc41fd18..5d82479a7e 100644
--- a/src/tint/cmd/main.cc
+++ b/src/tint/cmd/main.cc
@@ -49,7 +49,6 @@
 #include "tint/tint.h"
 
 #if TINT_BUILD_IR
-#include "src/tint/ir/debug.h"         // nogncheck
 #include "src/tint/ir/disassembler.h"  // nogncheck
 #include "src/tint/ir/from_program.h"  // nogncheck
 #include "src/tint/ir/module.h"        // nogncheck
@@ -110,7 +109,6 @@ struct Options {
 
 #if TINT_BUILD_IR
     bool dump_ir = false;
-    bool dump_ir_graph = false;
     bool use_ir = false;
 #endif  // TINT_BUILD_IR
 
@@ -374,8 +372,6 @@ bool ParseArgs(const std::vector<std::string>& args, Options* opts) {
 #if TINT_BUILD_IR
         } else if (arg == "--dump-ir") {
             opts->dump_ir = true;
-        } else if (arg == "--dump-ir-graph") {
-            opts->dump_ir_graph = true;
         } else if (arg == "--use-ir") {
             opts->use_ir = true;
 #endif  // TINT_BUILD_IR
@@ -1072,7 +1068,7 @@ int main(int argc, const char** argv) {
 #endif  // TINT_BUILD_SYNTAX_TREE_WRITER
 
 #if TINT_BUILD_IR
-    if (options.dump_ir || options.dump_ir_graph) {
+    if (options.dump_ir) {
         auto result = tint::ir::FromProgram(program.get());
         if (!result) {
             std::cerr << "Failed to build IR from program: " << result.Failure() << std::endl;
@@ -1082,10 +1078,6 @@ int main(int argc, const char** argv) {
                 tint::ir::Disassembler d(mod);
                 std::cout << d.Disassemble() << std::endl;
             }
-            if (options.dump_ir_graph) {
-                auto graph = tint::ir::Debug::AsDotGraph(&mod);
-                WriteFile("tint.dot", "w", graph);
-            }
         }
     }
 #endif  // TINT_BUILD_IR
diff --git a/src/tint/ir/block.h b/src/tint/ir/block.h
index 597ff2cbb8..d40002d0b1 100644
--- a/src/tint/ir/block.h
+++ b/src/tint/ir/block.h
@@ -46,18 +46,6 @@ class Block : public utils::Castable<Block> {
         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 Block* 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
     void SetInstructions(utils::VectorRef<const Instruction*> instructions) {
diff --git a/src/tint/ir/branch.cc b/src/tint/ir/branch.cc
index 0918962b5c..191831fe00 100644
--- a/src/tint/ir/branch.cc
+++ b/src/tint/ir/branch.cc
@@ -28,13 +28,6 @@ Branch::Branch(utils::VectorRef<Value*> args) : args_(std::move(args)) {
     }
 }
 
-Branch::Branch(Block* to, utils::VectorRef<Value*> args) : Branch(args) {
-    to_ = to;
-
-    TINT_ASSERT(IR, to_);
-    to_->AddInboundBranch(this);
-}
-
 Branch::~Branch() = default;
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/branch.h b/src/tint/ir/branch.h
index aa3d1a0b4a..5c926e9fba 100644
--- a/src/tint/ir/branch.h
+++ b/src/tint/ir/branch.h
@@ -29,15 +29,8 @@ namespace tint::ir {
 /// A branch instruction.
 class Branch : public utils::Castable<Branch, Instruction> {
   public:
-    /// Constructor
-    /// @param to the block to branch too
-    /// @param args the branch arguments
-    explicit Branch(Block* to, utils::VectorRef<Value*> args = {});
     ~Branch() override;
 
-    /// @returns the block being branched too.
-    const Block* To() const { return to_; }
-
     /// @returns the branch arguments
     utils::VectorRef<Value*> Args() const { return args_; }
 
@@ -47,7 +40,6 @@ class Branch : public utils::Castable<Branch, Instruction> {
     explicit Branch(utils::VectorRef<Value*> args);
 
   private:
-    Block* to_ = nullptr;
     utils::Vector<Value*, 2> args_;
 };
 
diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc
index 4bdb53d2bf..1daf4583ae 100644
--- a/src/tint/ir/builder.cc
+++ b/src/tint/ir/builder.cc
@@ -37,10 +37,6 @@ Block* Builder::CreateBlock() {
     return ir.blocks.Create<Block>();
 }
 
-RootTerminator* Builder::CreateRootTerminator() {
-    return ir.blocks.Create<RootTerminator>();
-}
-
 Function* Builder::CreateFunction(std::string_view name,
                                   const type::Type* return_type,
                                   Function::PipelineStage stage,
@@ -205,10 +201,6 @@ ir::Var* Builder::Declare(const type::Type* type) {
     return ir.values.Create<ir::Var>(type);
 }
 
-ir::Branch* Builder::Branch(Block* to, utils::VectorRef<Value*> args) {
-    return ir.values.Create<ir::Branch>(to, args);
-}
-
 ir::Return* Builder::Return(Function* func, utils::VectorRef<Value*> args) {
     return ir.values.Create<ir::Return>(func, args);
 }
@@ -220,6 +212,17 @@ ir::BreakIf* Builder::BreakIf(Value* condition, Loop* loop) {
 ir::Continue* Builder::Continue(Loop* loop) {
     return ir.values.Create<ir::Continue>(loop);
 }
+ir::ExitSwitch* Builder::ExitSwitch(Switch* sw) {
+    return ir.values.Create<ir::ExitSwitch>(sw);
+}
+
+ir::ExitLoop* Builder::ExitLoop(Loop* loop) {
+    return ir.values.Create<ir::ExitLoop>(loop);
+}
+
+ir::ExitIf* Builder::ExitIf(If* i, utils::VectorRef<Value*> args) {
+    return ir.values.Create<ir::ExitIf>(i, args);
+}
 
 ir::BlockParam* Builder::BlockParam(const type::Type* type) {
     return ir.values.Create<ir::BlockParam>(type);
diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h
index 9a2fe18098..8a4c0213ed 100644
--- a/src/tint/ir/builder.h
+++ b/src/tint/ir/builder.h
@@ -28,6 +28,9 @@
 #include "src/tint/ir/continue.h"
 #include "src/tint/ir/convert.h"
 #include "src/tint/ir/discard.h"
+#include "src/tint/ir/exit_if.h"
+#include "src/tint/ir/exit_loop.h"
+#include "src/tint/ir/exit_switch.h"
 #include "src/tint/ir/function.h"
 #include "src/tint/ir/function_param.h"
 #include "src/tint/ir/if.h"
@@ -35,7 +38,6 @@
 #include "src/tint/ir/loop.h"
 #include "src/tint/ir/module.h"
 #include "src/tint/ir/return.h"
-#include "src/tint/ir/root_terminator.h"
 #include "src/tint/ir/store.h"
 #include "src/tint/ir/switch.h"
 #include "src/tint/ir/unary.h"
@@ -64,9 +66,6 @@ class Builder {
     /// @returns a new block flow node
     Block* CreateBlock();
 
-    /// @returns a new root terminator flow node
-    RootTerminator* CreateRootTerminator();
-
     /// Creates a function flow node
     /// @param name the function name
     /// @param return_type the function return type
@@ -349,11 +348,21 @@ class Builder {
     /// @returns the instruction
     ir::Continue* Continue(Loop* loop);
 
-    /// Creates a branch declaration
-    /// @param to the node being branched too
+    /// Creates an exit switch instruction
+    /// @param sw the switch being exited
+    /// @returns the instruction
+    ir::ExitSwitch* ExitSwitch(Switch* sw);
+
+    /// Creates an exit loop instruction
+    /// @param loop the loop being exited
+    /// @returns the instruction
+    ir::ExitLoop* ExitLoop(Loop* loop);
+
+    /// Creates an exit if instruction
+    /// @param i the if being exited
     /// @param args the branch arguments
     /// @returns the instruction
-    ir::Branch* Branch(Block* to, utils::VectorRef<Value*> args = {});
+    ir::ExitIf* ExitIf(If* i, utils::VectorRef<Value*> args = {});
 
     /// Creates a new `BlockParam`
     /// @param type the parameter type
diff --git a/src/tint/ir/debug.cc b/src/tint/ir/debug.cc
deleted file mode 100644
index c7196984fd..0000000000
--- a/src/tint/ir/debug.cc
+++ /dev/null
@@ -1,95 +0,0 @@
-// Copyright 2022 The Tint Authors.
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "src/tint/ir/debug.h"
-
-#include <unordered_map>
-#include <unordered_set>
-
-#include "src/tint/ir/block.h"
-#include "src/tint/ir/continue.h"
-#include "src/tint/ir/if.h"
-#include "src/tint/ir/loop.h"
-#include "src/tint/ir/return.h"
-#include "src/tint/ir/switch.h"
-#include "src/tint/switch.h"
-#include "src/tint/utils/string_stream.h"
-
-namespace tint::ir {
-
-// static
-std::string Debug::AsDotGraph(const Module* mod) {
-    size_t block_count = 0;
-
-    std::unordered_set<const Block*> visited;
-    std::unordered_set<const Block*> merge_blocks;
-    std::unordered_map<const Block*, std::string> block_to_name;
-    utils::StringStream out;
-
-    auto name_for = [&](const Block* blk) -> std::string {
-        if (block_to_name.count(blk) > 0) {
-            return block_to_name[blk];
-        }
-
-        std::string name = "blk_" + std::to_string(block_count);
-        block_count += 1;
-
-        block_to_name[blk] = name;
-        return name;
-    };
-
-    std::function<void(const Block*)> Graph = [&](const Block* blk) {
-        if (visited.count(blk) > 0) {
-            return;
-        }
-        visited.insert(blk);
-
-        tint::Switch(blk,  //
-                     [&](const ir::Block* b) {
-                         if (block_to_name.count(b) == 0) {
-                             out << name_for(b) << R"( [label="block"])" << std::endl;
-                         }
-                         out << name_for(b) << " -> " << name_for(b->Branch()->To());
-
-                         // Dashed lines to merge blocks
-                         if (merge_blocks.count(b->Branch()->To()) != 0) {
-                             out << " [style=dashed]";
-                         }
-
-                         out << std::endl;
-
-                         if (b->Branch()->Is<ir::Return>()) {
-                             return;
-                         } else if (auto* cont = b->Branch()->As<ir::Continue>()) {
-                             Graph(cont->Loop()->Continuing());
-                         } else {
-                             Graph(b->Branch()->To());
-                         }
-                     });
-    };
-
-    out << "digraph G {" << std::endl;
-    for (const auto* func : mod->functions) {
-        // Cluster each function to label and draw a box around it.
-        out << "subgraph cluster_" << mod->NameOf(func).Name() << " {" << std::endl;
-        out << R"(label=")" << mod->NameOf(func).Name() << R"(")" << std::endl;
-        out << name_for(func->StartTarget()) << R"( [label="start"])" << std::endl;
-        Graph(func->StartTarget());
-        out << "}" << std::endl;
-    }
-    out << "}";
-    return out.str();
-}
-
-}  // namespace tint::ir
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index 6097726ee6..d39438aec3 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -27,11 +27,13 @@
 #include "src/tint/ir/continue.h"
 #include "src/tint/ir/convert.h"
 #include "src/tint/ir/discard.h"
+#include "src/tint/ir/exit_if.h"
+#include "src/tint/ir/exit_loop.h"
+#include "src/tint/ir/exit_switch.h"
 #include "src/tint/ir/if.h"
 #include "src/tint/ir/load.h"
 #include "src/tint/ir/loop.h"
 #include "src/tint/ir/return.h"
-#include "src/tint/ir/root_terminator.h"
 #include "src/tint/ir/store.h"
 #include "src/tint/ir/switch.h"
 #include "src/tint/ir/user_call.h"
@@ -70,7 +72,6 @@ void Disassembler::EmitBlockInstructions(const Block* b) {
     for (const auto* inst : b->Instructions()) {
         Indent();
         EmitInstruction(inst);
-        out_ << std::endl;
     }
 }
 
@@ -92,8 +93,8 @@ std::string_view Disassembler::IdOf(const Value* value) {
 std::string Disassembler::Disassemble() {
     if (mod_.root_block) {
         Indent() << "# Root block" << std::endl;
-        Walk(mod_.root_block);
-        Walk(mod_.root_block->Branch()->To());
+        WalkInternal(mod_.root_block);
+        out_ << std::endl;
     }
 
     for (auto* func : mod_.functions) {
@@ -108,41 +109,33 @@ void Disassembler::Walk(const Block* blk) {
     }
     visited_.Add(blk);
 
-    tint::Switch(
-        blk,
-        [&](const ir::RootTerminator* t) {
-            TINT_ASSERT(IR, !in_function_);
-            Indent() << "%b" << IdOf(t) << " = root_terminator" << std::endl << std::endl;
-        },
-        [&](const ir::Block* b) {
-            // If this block is dead, nothing to do
-            if (!b->HasBranchTarget()) {
-                return;
-            }
+    // If this block is dead, nothing to do
+    if (!blk->HasBranchTarget()) {
+        return;
+    }
 
-            Indent() << "%b" << IdOf(b) << " = block";
-            if (!b->Params().IsEmpty()) {
-                out_ << " (";
-                for (auto* p : b->Params()) {
-                    if (p != b->Params().Front()) {
-                        out_ << ", ";
-                    }
-                    EmitValue(p);
-                }
-                out_ << ")";
-            }
+    WalkInternal(blk);
+}
 
-            out_ << " {" << std::endl;
-            {
-                ScopedIndent si(indent_size_);
-                EmitBlockInstructions(b);
+void Disassembler::WalkInternal(const Block* blk) {
+    Indent() << "%b" << IdOf(blk) << " = block";
+    if (!blk->Params().IsEmpty()) {
+        out_ << " (";
+        for (auto* p : blk->Params()) {
+            if (p != blk->Params().Front()) {
+                out_ << ", ";
             }
-            Indent() << "}" << std::endl;
+            EmitValue(p);
+        }
+        out_ << ")";
+    }
 
-            if (!b->Branch()->Is<ir::Return>()) {
-                out_ << std::endl;
-            }
-        });
+    out_ << " {" << std::endl;
+    {
+        ScopedIndent si(indent_size_);
+        EmitBlockInstructions(blk);
+    }
+    Indent() << "}" << std::endl;
 }
 
 void Disassembler::EmitFunction(const Function* func) {
@@ -258,33 +251,39 @@ void Disassembler::EmitInstruction(const Instruction* inst) {
             EmitValueWithType(b);
             out_ << " = bitcast ";
             EmitArgs(b);
+            out_ << std::endl;
         },
-        [&](const ir::Discard*) { out_ << "discard"; },
+        [&](const ir::Discard*) { out_ << "discard" << std::endl; },
         [&](const ir::Builtin* b) {
             EmitValueWithType(b);
             out_ << " = " << builtin::str(b->Func()) << " ";
             EmitArgs(b);
+            out_ << std::endl;
         },
         [&](const ir::Construct* c) {
             EmitValueWithType(c);
             out_ << " = construct ";
             EmitArgs(c);
+            out_ << std::endl;
         },
         [&](const ir::Convert* c) {
             EmitValueWithType(c);
             out_ << " = convert " << c->FromType()->FriendlyName() << ", ";
             EmitArgs(c);
+            out_ << std::endl;
         },
         [&](const ir::Load* l) {
             EmitValueWithType(l);
             out_ << " = load ";
             EmitValue(l->From());
+            out_ << std::endl;
         },
         [&](const ir::Store* s) {
             out_ << "store ";
             EmitValue(s->To());
             out_ << ", ";
             EmitValue(s->From());
+            out_ << std::endl;
         },
         [&](const ir::UserCall* uc) {
             EmitValueWithType(uc);
@@ -293,6 +292,7 @@ void Disassembler::EmitInstruction(const Instruction* inst) {
                 out_ << ", ";
             }
             EmitArgs(uc);
+            out_ << std::endl;
         },
         [&](const ir::Var* v) {
             EmitValueWithType(v);
@@ -301,6 +301,7 @@ void Disassembler::EmitInstruction(const Instruction* inst) {
                 out_ << ", ";
                 EmitValue(v->Initializer());
             }
+            out_ << std::endl;
         },
         [&](const ir::Branch* b) { EmitBranch(b); },
         [&](Default) { out_ << "Unknown instruction: " << inst->TypeInfo().name; });
@@ -332,15 +333,18 @@ void Disassembler::EmitIf(const If* i) {
         ScopedIndent si(indent_size_);
         Indent() << "# True block" << std::endl;
         Walk(i->True());
+        out_ << std::endl;
     }
     if (has_false) {
         ScopedIndent si(indent_size_);
         Indent() << "# False block" << std::endl;
         Walk(i->False());
+        out_ << std::endl;
     }
     if (i->Merge()->HasBranchTarget()) {
         Indent() << "# Merge block" << std::endl;
         Walk(i->Merge());
+        out_ << std::endl;
     }
 }
 
@@ -358,16 +362,19 @@ void Disassembler::EmitLoop(const Loop* l) {
     {
         ScopedIndent si(indent_size_);
         Walk(l->Start());
+        out_ << std::endl;
     }
 
     if (l->Continuing()->HasBranchTarget()) {
         ScopedIndent si(indent_size_);
         Indent() << "# Continuing block" << std::endl;
         Walk(l->Continuing());
+        out_ << std::endl;
     }
     if (l->Merge()->HasBranchTarget()) {
         Indent() << "# Merge block" << std::endl;
         Walk(l->Merge());
+        out_ << std::endl;
     }
 }
 
@@ -402,29 +409,31 @@ void Disassembler::EmitSwitch(const Switch* s) {
         ScopedIndent si(indent_size_);
         Indent() << "# Case block" << std::endl;
         Walk(c.Start());
+        out_ << std::endl;
     }
     if (s->Merge()->HasBranchTarget()) {
         Indent() << "# Merge block" << std::endl;
         Walk(s->Merge());
+        out_ << std::endl;
     }
 }
 
 void Disassembler::EmitBranch(const Branch* b) {
-    std::string suffix = "";
-    if (b->Is<ir::Return>()) {
-        out_ << "ret";
-    } else if (auto* cont = b->As<ir::Continue>()) {
-        out_ << "continue %b" << IdOf(cont->Loop()->Continuing());
-    } else if (auto* bi = b->As<ir::BreakIf>()) {
-        out_ << "break_if ";
-        EmitValue(bi->Condition());
-        out_ << " %b" << IdOf(bi->Loop()->Start());
-    } else {
-        out_ << "br %b" << IdOf(b->To());
-        if (b->To()->Is<RootTerminator>()) {
-            suffix = "root_end";
-        }
-    }
+    tint::Switch(
+        b,  //
+        [&](const ir::Return*) { out_ << "ret"; },
+        [&](const ir::Continue* cont) {
+            out_ << "continue %b" << IdOf(cont->Loop()->Continuing());
+        },
+        [&](const ir::ExitIf* ei) { out_ << "exit_if %b" << IdOf(ei->If()->Merge()); },
+        [&](const ir::ExitSwitch* es) { out_ << "exit_switch %b" << IdOf(es->Switch()->Merge()); },
+        [&](const ir::ExitLoop* el) { out_ << "exit_loop %b" << IdOf(el->Loop()->Merge()); },
+        [&](const ir::BreakIf* bi) {
+            out_ << "break_if ";
+            EmitValue(bi->Condition());
+            out_ << " %b" << IdOf(bi->Loop()->Start());
+        },
+        [&](Default) { out_ << "Unknown branch " << b->TypeInfo().name; });
 
     if (!b->Args().IsEmpty()) {
         out_ << " ";
@@ -435,9 +444,7 @@ void Disassembler::EmitBranch(const Branch* b) {
             EmitValue(v);
         }
     }
-    if (!suffix.empty()) {
-        out_ << "  # " << suffix;
-    }
+    out_ << std::endl;
 }
 
 void Disassembler::EmitArgs(const Call* call) {
@@ -508,6 +515,7 @@ void Disassembler::EmitBinary(const Binary* b) {
     EmitValue(b->LHS());
     out_ << ", ";
     EmitValue(b->RHS());
+    out_ << std::endl;
 }
 
 void Disassembler::EmitUnary(const Unary* u) {
@@ -523,6 +531,7 @@ void Disassembler::EmitUnary(const Unary* u) {
     }
     out_ << " ";
     EmitValue(u->Val());
+    out_ << std::endl;
 }
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/disassembler.h b/src/tint/ir/disassembler.h
index f17103145d..06f5b08278 100644
--- a/src/tint/ir/disassembler.h
+++ b/src/tint/ir/disassembler.h
@@ -57,6 +57,7 @@ class Disassembler {
     std::string_view IdOf(const Value* node);
 
     void Walk(const Block* blk);
+    void WalkInternal(const Block* blk);
     void EmitFunction(const Function* func);
     void EmitInstruction(const Instruction* inst);
     void EmitValueWithType(const Value* val);
diff --git a/src/tint/ir/root_terminator.cc b/src/tint/ir/exit_if.cc
similarity index 63%
rename from src/tint/ir/root_terminator.cc
rename to src/tint/ir/exit_if.cc
index bfccf46a21..8b7de7f738 100644
--- a/src/tint/ir/root_terminator.cc
+++ b/src/tint/ir/exit_if.cc
@@ -1,4 +1,4 @@
-// Copyright 2022 The Tint Authors.
+// Copyright 2023 The Tint Authors.
 //
 // Licensed under the Apache License, Version 2.0 (the "License");
 // you may not use this file except in compliance with the License.
@@ -12,14 +12,20 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#include "src/tint/ir/root_terminator.h"
+#include "src/tint/ir/exit_if.h"
 
-TINT_INSTANTIATE_TYPEINFO(tint::ir::RootTerminator);
+#include "src/tint/ir/if.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::ExitIf);
 
 namespace tint::ir {
 
-RootTerminator::RootTerminator() : Base() {}
+ExitIf::ExitIf(ir::If* i, utils::VectorRef<Value*> args) : Base(args), if_(i) {
+    TINT_ASSERT(IR, if_);
+    if_->AddUsage(this);
+    if_->Merge()->AddInboundBranch(this);
+}
 
-RootTerminator::~RootTerminator() = default;
+ExitIf::~ExitIf() = default;
 
 }  // namespace tint::ir
diff --git a/src/tint/ir/exit_if.h b/src/tint/ir/exit_if.h
new file mode 100644
index 0000000000..9ba1421501
--- /dev/null
+++ b/src/tint/ir/exit_if.h
@@ -0,0 +1,46 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_TINT_IR_EXIT_IF_H_
+#define SRC_TINT_IR_EXIT_IF_H_
+
+#include "src/tint/ir/branch.h"
+#include "src/tint/utils/castable.h"
+
+// Forward declarations
+namespace tint::ir {
+class If;
+}  // namespace tint::ir
+
+namespace tint::ir {
+
+/// A exit if instruction.
+class ExitIf : public utils::Castable<ExitIf, Branch> {
+  public:
+    /// Constructor
+    /// @param i the if being exited
+    /// @param args the branch arguments
+    explicit ExitIf(ir::If* i, utils::VectorRef<Value*> args = {});
+    ~ExitIf() override;
+
+    /// @returns the if being exited
+    const ir::If* If() const { return if_; }
+
+  private:
+    ir::If* if_ = nullptr;
+};
+
+}  // namespace tint::ir
+
+#endif  // SRC_TINT_IR_EXIT_IF_H_
diff --git a/src/tint/ir/root_terminator.h b/src/tint/ir/exit_loop.cc
similarity index 58%
rename from src/tint/ir/root_terminator.h
rename to src/tint/ir/exit_loop.cc
index 4a52b3290f..5fe3910e93 100644
--- a/src/tint/ir/root_terminator.h
+++ b/src/tint/ir/exit_loop.cc
@@ -1,4 +1,4 @@
-// Copyright 2022 The Tint Authors.
+// Copyright 2023 The Tint Authors.
 //
 // Licensed under the Apache License, Version 2.0 (the "License");
 // you may not use this file except in compliance with the License.
@@ -12,21 +12,20 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#ifndef SRC_TINT_IR_ROOT_TERMINATOR_H_
-#define SRC_TINT_IR_ROOT_TERMINATOR_H_
+#include "src/tint/ir/exit_loop.h"
 
-#include "src/tint/ir/block.h"
+#include "src/tint/ir/loop.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::ExitLoop);
 
 namespace tint::ir {
 
-/// Block used as the end of a root block. There are no instructions in this block.
-class RootTerminator : public utils::Castable<RootTerminator, Block> {
-  public:
-    /// Constructor
-    RootTerminator();
-    ~RootTerminator() override;
-};
+ExitLoop::ExitLoop(ir::Loop* loop) : Base(utils::Empty), loop_(loop) {
+    TINT_ASSERT(IR, loop_);
+    loop_->AddUsage(this);
+    loop_->Merge()->AddInboundBranch(this);
+}
+
+ExitLoop::~ExitLoop() = default;
 
 }  // namespace tint::ir
-
-#endif  // SRC_TINT_IR_ROOT_TERMINATOR_H_
diff --git a/src/tint/ir/exit_loop.h b/src/tint/ir/exit_loop.h
new file mode 100644
index 0000000000..1df1119bc9
--- /dev/null
+++ b/src/tint/ir/exit_loop.h
@@ -0,0 +1,45 @@
+// 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_EXIT_LOOP_H_
+#define SRC_TINT_IR_EXIT_LOOP_H_
+
+#include "src/tint/ir/branch.h"
+#include "src/tint/utils/castable.h"
+
+// Forward declarations
+namespace tint::ir {
+class Loop;
+}  // namespace tint::ir
+
+namespace tint::ir {
+
+/// A exit loop instruction.
+class ExitLoop : public utils::Castable<ExitLoop, Branch> {
+  public:
+    /// Constructor
+    /// @param loop the loop being exited
+    explicit ExitLoop(ir::Loop* loop);
+    ~ExitLoop() override;
+
+    /// @returns the loop being exited
+    const ir::Loop* Loop() const { return loop_; }
+
+  private:
+    ir::Loop* loop_ = nullptr;
+};
+
+}  // namespace tint::ir
+
+#endif  // SRC_TINT_IR_EXIT_LOOP_H_
diff --git a/src/tint/ir/debug.h b/src/tint/ir/exit_switch.cc
similarity index 50%
rename from src/tint/ir/debug.h
rename to src/tint/ir/exit_switch.cc
index 2363776b8d..ba6a178c05 100644
--- a/src/tint/ir/debug.h
+++ b/src/tint/ir/exit_switch.cc
@@ -1,4 +1,4 @@
-// Copyright 2022 The Tint Authors.
+// Copyright 2023 The Tint Authors.
 //
 // Licensed under the Apache License, Version 2.0 (the "License");
 // you may not use this file except in compliance with the License.
@@ -12,29 +12,20 @@
 // See the License for the specific language governing permissions and
 // limitations under the License.
 
-#ifndef SRC_TINT_IR_DEBUG_H_
-#define SRC_TINT_IR_DEBUG_H_
+#include "src/tint/ir/exit_switch.h"
 
-#include <string>
+#include "src/tint/ir/switch.h"
 
-#include "src/tint/ir/module.h"
+TINT_INSTANTIATE_TYPEINFO(tint::ir::ExitSwitch);
 
 namespace tint::ir {
 
-/// Helper class to debug IR.
-class Debug {
-  public:
-    /// Returns the module as a dot graph
-    /// @param mod the module to emit
-    /// @returns the dot graph for the given module
-    static std::string AsDotGraph(const Module* mod);
+ExitSwitch::ExitSwitch(ir::Switch* sw) : Base(utils::Empty), switch_(sw) {
+    TINT_ASSERT(IR, switch_);
+    switch_->AddUsage(this);
+    switch_->Merge()->AddInboundBranch(this);
+}
 
-    /// Returns the module as a string
-    /// @param mod the module to emit
-    /// @returns the string representation of the module
-    static std::string AsString(const Module* mod);
-};
+ExitSwitch::~ExitSwitch() = default;
 
 }  // namespace tint::ir
-
-#endif  // SRC_TINT_IR_DEBUG_H_
diff --git a/src/tint/ir/exit_switch.h b/src/tint/ir/exit_switch.h
new file mode 100644
index 0000000000..6b406fe5c2
--- /dev/null
+++ b/src/tint/ir/exit_switch.h
@@ -0,0 +1,45 @@
+// 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_EXIT_SWITCH_H_
+#define SRC_TINT_IR_EXIT_SWITCH_H_
+
+#include "src/tint/ir/branch.h"
+#include "src/tint/utils/castable.h"
+
+// Forward declarations
+namespace tint::ir {
+class Switch;
+}  // namespace tint::ir
+
+namespace tint::ir {
+
+/// A exit switch instruction.
+class ExitSwitch : public utils::Castable<ExitSwitch, Branch> {
+  public:
+    /// Constructor
+    /// @param sw the switch being exited
+    explicit ExitSwitch(ir::Switch* sw);
+    ~ExitSwitch() override;
+
+    /// @returns the switch being exited
+    const ir::Switch* Switch() const { return switch_; }
+
+  private:
+    ir::Switch* switch_ = nullptr;
+};
+
+}  // namespace tint::ir
+
+#endif  // SRC_TINT_IR_EXIT_SWITCH_H_
diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc
index a4d95e0c29..67c9f3f53c 100644
--- a/src/tint/ir/from_program.cc
+++ b/src/tint/ir/from_program.cc
@@ -62,6 +62,9 @@
 #include "src/tint/ast/while_statement.h"
 #include "src/tint/ir/block_param.h"
 #include "src/tint/ir/builder.h"
+#include "src/tint/ir/exit_if.h"
+#include "src/tint/ir/exit_loop.h"
+#include "src/tint/ir/exit_switch.h"
 #include "src/tint/ir/function.h"
 #include "src/tint/ir/if.h"
 #include "src/tint/ir/loop.h"
@@ -98,11 +101,8 @@ namespace {
 
 using ResultType = utils::Result<Module, diag::List>;
 
-// 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 Block* b, uint32_t count) {
-    return b->InboundBranches().Length() > count;
+bool IsConnected(const Block* b) {
+    return b->InboundBranches().Length() > 0;
 }
 
 /// Impl is the private-implementation of FromProgram().
@@ -176,21 +176,6 @@ class Impl {
         current_flow_block_ = nullptr;
     }
 
-    void BranchTo(Block* node, utils::VectorRef<Value*> args = {}) {
-        TINT_ASSERT(IR, current_flow_block_);
-        TINT_ASSERT(IR, !current_flow_block_->HasBranchTarget());
-
-        current_flow_block_->Instructions().Push(builder_.Branch(node, args));
-        current_flow_block_ = nullptr;
-    }
-
-    void BranchToIfNeeded(Block* node) {
-        if (!NeedBranch()) {
-            return;
-        }
-        BranchTo(node);
-    }
-
     Branch* FindEnclosingControl(ControlFlags flags) {
         for (auto it = control_stack_.rbegin(); it != control_stack_.rend(); ++it) {
             if ((*it)->Is<Loop>()) {
@@ -238,11 +223,6 @@ 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_));
         }
@@ -541,7 +521,9 @@ class Impl {
             EmitBlock(stmt->body);
 
             // If the true branch did not execute control flow, then go to the Merge().target
-            BranchToIfNeeded(if_inst->Merge());
+            if (NeedBranch()) {
+                SetBranch(builder_.ExitIf(if_inst));
+            }
 
             current_flow_block_ = if_inst->False();
             if (stmt->else_statement) {
@@ -549,14 +531,16 @@ class Impl {
             }
 
             // If the false branch did not execute control flow, then go to the Merge().target
-            BranchToIfNeeded(if_inst->Merge());
+            if (NeedBranch()) {
+                SetBranch(builder_.ExitIf(if_inst));
+            }
         }
         current_flow_block_ = nullptr;
 
         // If both branches went somewhere, then they both returned, continued or broke. So,
         // there is no need for the if merge-block and there is nothing to branch to the merge
         // block anyway.
-        if (IsConnected(if_inst->Merge(), 1)) {
+        if (IsConnected(if_inst->Merge())) {
             current_flow_block_ = if_inst->Merge();
         }
     }
@@ -580,7 +564,7 @@ class Impl {
                 SetBranch(builder_.Continue(loop_inst));
             }
 
-            if (IsConnected(loop_inst->Continuing(), 0)) {
+            if (IsConnected(loop_inst->Continuing())) {
                 // Note, even if there is no continuing block, we may have branched into the
                 // continue so we have to set the current block and then emit the branch if needed
                 // below otherwise empty continuing blocks will fail to branch back to the start
@@ -600,7 +584,7 @@ class Impl {
         // target branches, eventually, to the merge, but nothing branched to the
         // Continuing() block.
         current_flow_block_ = loop_inst->Merge();
-        if (!IsConnected(loop_inst->Merge(), 0)) {
+        if (!IsConnected(loop_inst->Merge())) {
             current_flow_block_ = nullptr;
         }
     }
@@ -626,10 +610,14 @@ class Impl {
 
             // Create an `if (cond) {} else {break;}` control flow
             auto* if_inst = builder_.CreateIf(reg.Get());
-            if_inst->True()->Instructions().Push(builder_.Branch(if_inst->Merge()));
-            if_inst->False()->Instructions().Push(builder_.Branch(loop_inst->Merge()));
             current_flow_block_->Instructions().Push(if_inst);
 
+            current_flow_block_ = if_inst->True();
+            SetBranch(builder_.ExitIf(if_inst));
+
+            current_flow_block_ = if_inst->False();
+            SetBranch(builder_.ExitLoop(loop_inst));
+
             current_flow_block_ = if_inst->Merge();
             EmitBlock(stmt->body);
 
@@ -669,10 +657,14 @@ class Impl {
 
                 // Create an `if (cond) {} else {break;}` control flow
                 auto* if_inst = builder_.CreateIf(reg.Get());
-                if_inst->True()->Instructions().Push(builder_.Branch(if_inst->Merge()));
-                if_inst->False()->Instructions().Push(builder_.Branch(loop_inst->Merge()));
                 current_flow_block_->Instructions().Push(if_inst);
 
+                current_flow_block_ = if_inst->True();
+                SetBranch(builder_.ExitIf(if_inst));
+
+                current_flow_block_ = if_inst->False();
+                SetBranch(builder_.ExitLoop(loop_inst));
+
                 current_flow_block_ = if_inst->Merge();
             }
 
@@ -719,12 +711,14 @@ class Impl {
                 current_flow_block_ = builder_.CreateCase(switch_inst, selectors);
                 EmitBlock(c->Body()->Declaration());
 
-                BranchToIfNeeded(switch_inst->Merge());
+                if (NeedBranch()) {
+                    SetBranch(builder_.ExitSwitch(switch_inst));
+                }
             }
         }
         current_flow_block_ = nullptr;
 
-        if (IsConnected(switch_inst->Merge(), 1)) {
+        if (IsConnected(switch_inst->Merge())) {
             current_flow_block_ = switch_inst->Merge();
         }
     }
@@ -746,9 +740,9 @@ class Impl {
         TINT_ASSERT(IR, current_control);
 
         if (auto* c = current_control->As<Loop>()) {
-            BranchTo(c->Merge());
+            SetBranch(builder_.ExitLoop(c));
         } else if (auto* s = current_control->As<Switch>()) {
-            BranchTo(s->Merge());
+            SetBranch(builder_.ExitSwitch(s));
         } else {
             TINT_UNREACHABLE(IR, diagnostics_);
         }
@@ -964,14 +958,14 @@ class Impl {
                 // If the lhs is false, then that is the result we want to pass to the merge
                 // block as our argument
                 current_flow_block_ = if_inst->False();
-                BranchTo(if_inst->Merge(), std::move(alt_args));
+                SetBranch(builder_.ExitIf(if_inst, std::move(alt_args)));
 
                 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 as our argument
                 current_flow_block_ = if_inst->True();
-                BranchTo(if_inst->Merge(), std::move(alt_args));
+                SetBranch(builder_.ExitIf(if_inst, std::move(alt_args)));
 
                 current_flow_block_ = if_inst->False();
             }
@@ -983,7 +977,7 @@ class Impl {
             utils::Vector<Value*, 1> args;
             args.Push(rhs.Get());
 
-            BranchTo(if_inst->Merge(), std::move(args));
+            SetBranch(builder_.ExitIf(if_inst, std::move(args)));
         }
         current_flow_block_ = if_inst->Merge();
 
diff --git a/src/tint/ir/from_program_binary_test.cc b/src/tint/ir/from_program_binary_test.cc
index 1a81f56598..8cc56047e4 100644
--- a/src/tint/ir/from_program_binary_test.cc
+++ b/src/tint/ir/from_program_binary_test.cc
@@ -60,13 +60,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Increment) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %v1:ptr<private, u32, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     %3:u32 = load %v1
     %4:u32 = add %3, 1u
     store %v1, %4
@@ -87,13 +84,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAdd) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %v1:ptr<private, u32, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     %3:u32 = load %v1
     %4:u32 = add %3, 1u
     store %v1, %4
@@ -137,13 +131,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Decrement) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %v1:ptr<private, i32, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     %3:i32 = load %v1
     %4:i32 = sub %3, 1i
     store %v1, %4
@@ -164,13 +155,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundSubtract) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %v1:ptr<private, u32, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     %3:u32 = load %v1
     %4:u32 = sub %3, 1u
     store %v1, %4
@@ -214,13 +202,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundMultiply) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %v1:ptr<private, u32, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     %3:u32 = load %v1
     %4:u32 = mul %3, 1u
     store %v1, %4
@@ -264,13 +249,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundDiv) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %v1:ptr<private, u32, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     %3:u32 = load %v1
     %4:u32 = div %3, 1u
     store %v1, %4
@@ -314,13 +296,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundModulo) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %v1:ptr<private, u32, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     %3:u32 = load %v1
     %4:u32 = mod %3, 1u
     store %v1, %4
@@ -364,13 +343,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAnd) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %v1:ptr<private, bool, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     %3:bool = load %v1
     %4:bool = and %3, false
     store %v1, %4
@@ -414,13 +390,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundOr) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %v1:ptr<private, bool, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     %3:bool = load %v1
     %4:bool = or %3, false
     store %v1, %4
@@ -464,13 +437,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundXor) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %v1:ptr<private, u32, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     %3:u32 = load %v1
     %4:u32 = xor %3, 1u
     store %v1, %4
@@ -499,12 +469,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) {
     if %3 [t: %b3, f: %b4, m: %b5]
       # True block
       %b3 = block {
-        br %b5 false
+        exit_if %b5 false
       }
 
       # False block
       %b4 = block {
-        br %b5 %3
+        exit_if %b5 %3
       }
 
     # Merge block
@@ -512,12 +482,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) {
       if %4:bool [t: %b6, f: %b7, m: %b8]
         # True block
         %b6 = block {
-          br %b8
+          exit_if %b8
         }
 
         # False block
         %b7 = block {
-          br %b8
+          exit_if %b8
         }
 
       # Merge block
@@ -527,9 +497,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) {
 
     }
 
-
   }
-
 }
 )");
 }
@@ -553,12 +521,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) {
     if %3 [t: %b3, f: %b4, m: %b5]
       # True block
       %b3 = block {
-        br %b5 %3
+        exit_if %b5 %3
       }
 
       # False block
       %b4 = block {
-        br %b5 true
+        exit_if %b5 true
       }
 
     # Merge block
@@ -566,12 +534,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) {
       if %4:bool [t: %b6, f: %b7, m: %b8]
         # True block
         %b6 = block {
-          br %b8
+          exit_if %b8
         }
 
         # False block
         %b7 = block {
-          br %b8
+          exit_if %b8
         }
 
       # Merge block
@@ -581,9 +549,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) {
 
     }
 
-
   }
-
 }
 )");
 }
@@ -760,13 +726,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftLeft) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %v1:ptr<private, u32, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     %3:u32 = load %v1
     %4:u32 = shiftl %3, 1u
     store %v1, %4
@@ -810,13 +773,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftRight) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %v1:ptr<private, u32, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     %3:u32 = load %v1
     %4:u32 = shiftr %3, 1u
     store %v1, %4
@@ -853,12 +813,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound) {
         %7:f32 = mul 2.29999995231628417969f, %6
         %8:f32 = div %5, %7
         %9:bool = gt 2.5f, %8
-        br %b5 %9
+        exit_if %b5 %9
       }
 
       # False block
       %b4 = block {
-        br %b5 %4
+        exit_if %b5 %4
       }
 
     # Merge block
@@ -867,7 +827,6 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound) {
     }
 
   }
-
 }
 )");
 }
diff --git a/src/tint/ir/from_program_builtin_test.cc b/src/tint/ir/from_program_builtin_test.cc
index 3993fbbbf8..acdbc8a5ed 100644
--- a/src/tint/ir/from_program_builtin_test.cc
+++ b/src/tint/ir/from_program_builtin_test.cc
@@ -37,13 +37,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Builtin) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %i:ptr<private, f32, read_write> = var, 1.0f
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     %3:f32 = load %i
     %tint_symbol:f32 = asin %3
     ret
diff --git a/src/tint/ir/from_program_call_test.cc b/src/tint/ir/from_program_call_test.cc
index 43a9709be4..155e42c61c 100644
--- a/src/tint/ir/from_program_call_test.cc
+++ b/src/tint/ir/from_program_call_test.cc
@@ -102,13 +102,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %i:ptr<private, i32, read_write> = var, 1i
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     %3:i32 = load %i
     %tint_symbol:f32 = convert i32, %3
     ret
@@ -127,11 +124,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_ConstructEmpty) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %i:ptr<private, vec3<f32>, read_write> = var, vec3<f32> 0.0f
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
 )");
 }
 
@@ -146,13 +140,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %i:ptr<private, f32, read_write> = var, 1.0f
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     %3:f32 = load %i
     %tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %3
     ret
diff --git a/src/tint/ir/from_program_store_test.cc b/src/tint/ir/from_program_store_test.cc
index 2d3f3d730c..38efea07dc 100644
--- a/src/tint/ir/from_program_store_test.cc
+++ b/src/tint/ir/from_program_store_test.cc
@@ -38,13 +38,10 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Assign) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %a:ptr<private, u32, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     store %a, 4u
     ret
   }
diff --git a/src/tint/ir/from_program_test.cc b/src/tint/ir/from_program_test.cc
index e0f161d914..f2b9c19aaf 100644
--- a/src/tint/ir/from_program_test.cc
+++ b/src/tint/ir/from_program_test.cc
@@ -141,7 +141,7 @@ TEST_F(IR_BuilderImplTest, IfStatement) {
 
     EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
-    EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length());
+    EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -149,12 +149,12 @@ TEST_F(IR_BuilderImplTest, IfStatement) {
     if true [t: %b2, f: %b3, m: %b4]
       # True block
       %b2 = block {
-        br %b4
+        exit_if %b4
       }
 
       # False block
       %b3 = block {
-        br %b4
+        exit_if %b4
       }
 
     # Merge block
@@ -163,7 +163,6 @@ TEST_F(IR_BuilderImplTest, IfStatement) {
     }
 
   }
-
 }
 )");
 }
@@ -182,7 +181,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) {
 
     EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
-    EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -192,9 +191,10 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) {
       %b2 = block {
         ret
       }
+
       # False block
       %b3 = block {
-        br %b4
+        exit_if %b4
       }
 
     # Merge block
@@ -203,7 +203,6 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) {
     }
 
   }
-
 }
 )");
 }
@@ -222,7 +221,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_FalseReturns) {
 
     EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
-    EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -230,20 +229,20 @@ TEST_F(IR_BuilderImplTest, IfStatement_FalseReturns) {
     if true [t: %b2, f: %b3, m: %b4]
       # True block
       %b2 = block {
-        br %b4
+        exit_if %b4
       }
 
       # False block
       %b3 = block {
         ret
       }
+
     # Merge block
     %b4 = block {
       ret
     }
 
   }
-
 }
 )");
 }
@@ -262,7 +261,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_BothReturn) {
 
     EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
-    EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
+    EXPECT_EQ(0u, flow->Merge()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -272,13 +271,13 @@ TEST_F(IR_BuilderImplTest, IfStatement_BothReturn) {
       %b2 = block {
         ret
       }
+
       # False block
       %b3 = block {
         ret
       }
 
   }
-
 }
 )");
 }
@@ -306,20 +305,19 @@ TEST_F(IR_BuilderImplTest, IfStatement_JumpChainToMerge) {
       %b2 = block {
         loop [s: %b5, m: %b6]
           %b5 = block {
-            br %b6
+            exit_loop %b6
           }
 
         # Merge block
         %b6 = block {
-          br %b4
+          exit_if %b4
         }
 
-
       }
 
       # False block
       %b3 = block {
-        br %b4
+        exit_if %b4
       }
 
     # Merge block
@@ -328,7 +326,6 @@ TEST_F(IR_BuilderImplTest, IfStatement_JumpChainToMerge) {
     }
 
   }
-
 }
 )");
 }
@@ -345,7 +342,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) {
 
     ASSERT_EQ(1u, m.functions.Length());
 
-    EXPECT_EQ(1u, flow->Start()->InboundBranches().Length());
+    EXPECT_EQ(0u, flow->Start()->InboundBranches().Length());
     EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
 
@@ -354,7 +351,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) {
   %b1 = block {
     loop [s: %b2, m: %b3]
       %b2 = block {
-        br %b3
+        exit_loop %b3
       }
 
     # Merge block
@@ -363,7 +360,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) {
     }
 
   }
-
 }
 )");
 }
@@ -383,12 +379,12 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
 
     ASSERT_EQ(1u, m.functions.Length());
 
-    EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
     EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(2u, loop_flow->Merge()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
-    EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -398,12 +394,12 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
         if true [t: %b5, f: %b6, m: %b7]
           # True block
           %b5 = block {
-            br %b4
+            exit_loop %b4
           }
 
           # False block
           %b6 = block {
-            br %b7
+            exit_if %b7
           }
 
         # Merge block
@@ -411,7 +407,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
           continue %b3
         }
 
-
       }
 
       # Continuing block
@@ -425,7 +420,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
     }
 
   }
-
 }
 )");
 }
@@ -443,7 +437,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) {
 
     ASSERT_EQ(1u, m.functions.Length());
 
-    EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
     EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length());
 
@@ -466,7 +460,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) {
     }
 
   }
-
 }
 )");
 }
@@ -500,7 +493,6 @@ TEST_F(IR_BuilderImplTest, Loop_Continuing_Body_Scope) {
     }
 
   }
-
 }
 )");
 }
@@ -519,12 +511,12 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
 
     ASSERT_EQ(1u, m.functions.Length());
 
-    EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
+    EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
     EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
-    EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -536,9 +528,10 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
           %b5 = block {
             ret
           }
+
           # False block
           %b6 = block {
-            br %b7
+            exit_if %b7
           }
 
         # Merge block
@@ -546,7 +539,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
           continue %b3
         }
 
-
       }
 
       # Continuing block
@@ -560,7 +552,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
     }
 
   }
-
 }
 )");
 }
@@ -577,7 +568,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn) {
 
     ASSERT_EQ(1u, m.functions.Length());
 
-    EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
+    EXPECT_EQ(0u, loop_flow->Start()->InboundBranches().Length());
     EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length());
 
@@ -590,7 +581,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn) {
       }
 
   }
-
 }
 )");
 }
@@ -616,7 +606,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) {
 
     ASSERT_EQ(1u, m.functions.Length());
 
-    EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
+    EXPECT_EQ(0u, loop_flow->Start()->InboundBranches().Length());
     EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length());
 
@@ -629,7 +619,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) {
       }
 
   }
-
 }
 )");
 }
@@ -648,12 +637,12 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) {
 
     ASSERT_EQ(1u, m.functions.Length());
 
-    EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
+    EXPECT_EQ(0u, loop_flow->Start()->InboundBranches().Length());
     EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(2u, loop_flow->Merge()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
-    EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
+    EXPECT_EQ(0u, if_flow->Merge()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -663,15 +652,14 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) {
         if true [t: %b4, f: %b5]
           # True block
           %b4 = block {
-            br %b3
+            exit_loop %b3
           }
 
           # False block
           %b5 = block {
-            br %b3
+            exit_loop %b3
           }
 
-
       }
 
     # Merge block
@@ -680,7 +668,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) {
     }
 
   }
-
 }
 )");
 }
@@ -712,12 +699,12 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
             if true [t: %b8, f: %b9, m: %b10]
               # True block
               %b8 = block {
-                br %b7
+                exit_loop %b7
               }
 
               # False block
               %b9 = block {
-                br %b10
+                exit_if %b10
               }
 
             # Merge block
@@ -730,7 +717,7 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
 
                 # False block
                 %b12 = block {
-                  br %b13
+                  exit_if %b13
                 }
 
               # Merge block
@@ -738,17 +725,15 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
                 continue %b6
               }
 
-
             }
 
-
           }
 
           # Continuing block
           %b6 = block {
             loop [s: %b14, m: %b15]
               %b14 = block {
-                br %b15
+                exit_loop %b15
               }
 
             # Merge block
@@ -768,10 +753,8 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
                 break_if false %b5
               }
 
-
             }
 
-
           }
 
         # Merge block
@@ -779,12 +762,12 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
           if true [t: %b19, f: %b20, m: %b21]
             # True block
             %b19 = block {
-              br %b4
+              exit_loop %b4
             }
 
             # False block
             %b20 = block {
-              br %b21
+              exit_if %b21
             }
 
           # Merge block
@@ -792,10 +775,8 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
             continue %b3
           }
 
-
         }
 
-
       }
 
       # Continuing block
@@ -809,7 +790,6 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
     }
 
   }
-
 }
 )");
 }
@@ -830,12 +810,12 @@ TEST_F(IR_BuilderImplTest, While) {
 
     ASSERT_EQ(1u, m.functions.Length());
 
-    EXPECT_EQ(2u, flow->Start()->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Start()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
-    EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -845,12 +825,12 @@ TEST_F(IR_BuilderImplTest, While) {
         if false [t: %b5, f: %b6, m: %b7]
           # True block
           %b5 = block {
-            br %b7
+            exit_if %b7
           }
 
           # False block
           %b6 = block {
-            br %b4
+            exit_loop %b4
           }
 
         # Merge block
@@ -858,7 +838,6 @@ TEST_F(IR_BuilderImplTest, While) {
           continue %b3
         }
 
-
       }
 
       # Continuing block
@@ -872,7 +851,6 @@ TEST_F(IR_BuilderImplTest, While) {
     }
 
   }
-
 }
 )");
 }
@@ -893,12 +871,12 @@ TEST_F(IR_BuilderImplTest, While_Return) {
 
     ASSERT_EQ(1u, m.functions.Length());
 
-    EXPECT_EQ(2u, flow->Start()->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Start()->InboundBranches().Length());
     EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
-    EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -908,12 +886,12 @@ TEST_F(IR_BuilderImplTest, While_Return) {
         if true [t: %b5, f: %b6, m: %b7]
           # True block
           %b5 = block {
-            br %b7
+            exit_if %b7
           }
 
           # False block
           %b6 = block {
-            br %b4
+            exit_loop %b4
           }
 
         # Merge block
@@ -934,7 +912,6 @@ TEST_F(IR_BuilderImplTest, While_Return) {
     }
 
   }
-
 }
 )");
 }
@@ -973,7 +950,7 @@ TEST_F(IR_BuilderImplTest, DISABLED_For) {
     EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
     EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
-    EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length());
+    EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m), R"()");
 }
@@ -990,7 +967,7 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) {
 
     ASSERT_EQ(1u, m.functions.Length());
 
-    EXPECT_EQ(1u, flow->Start()->InboundBranches().Length());
+    EXPECT_EQ(0u, flow->Start()->InboundBranches().Length());
     EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
 
@@ -999,7 +976,7 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) {
   %b1 = block {
     loop [s: %b2, m: %b3]
       %b2 = block {
-        br %b3
+        exit_loop %b3
       }
 
     # Merge block
@@ -1008,7 +985,6 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) {
     }
 
   }
-
 }
 )");
 }
@@ -1047,7 +1023,7 @@ TEST_F(IR_BuilderImplTest, Switch) {
     EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
     EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length());
     EXPECT_EQ(1u, cases[2].Start()->InboundBranches().Length());
-    EXPECT_EQ(4u, flow->Merge()->InboundBranches().Length());
+    EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1055,17 +1031,17 @@ TEST_F(IR_BuilderImplTest, Switch) {
     switch 1i [c: (0i, %b2), c: (1i, %b3), c: (default, %b4), m: %b5]
       # Case block
       %b2 = block {
-        br %b5
+        exit_switch %b5
       }
 
       # Case block
       %b3 = block {
-        br %b5
+        exit_switch %b5
       }
 
       # Case block
       %b4 = block {
-        br %b5
+        exit_switch %b5
       }
 
     # Merge block
@@ -1074,7 +1050,6 @@ TEST_F(IR_BuilderImplTest, Switch) {
     }
 
   }
-
 }
 )");
 }
@@ -1109,7 +1084,7 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) {
     EXPECT_TRUE(cases[0].selectors[2].IsDefault());
 
     EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
-    EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1117,7 +1092,7 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) {
     switch 1i [c: (0i 1i default, %b2), m: %b3]
       # Case block
       %b2 = block {
-        br %b3
+        exit_switch %b3
       }
 
     # Merge block
@@ -1126,7 +1101,6 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) {
     }
 
   }
-
 }
 )");
 }
@@ -1149,7 +1123,7 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) {
     EXPECT_TRUE(cases[0].selectors[0].IsDefault());
 
     EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
-    EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
+    EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1157,7 +1131,7 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) {
     switch 1i [c: (default, %b2), m: %b3]
       # Case block
       %b2 = block {
-        br %b3
+        exit_switch %b3
       }
 
     # Merge block
@@ -1166,7 +1140,6 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) {
     }
 
   }
-
 }
 )");
 }
@@ -1197,7 +1170,7 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) {
 
     EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
     EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length());
-    EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length());
+    EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
     // This is 1 because the if is dead-code eliminated and the return doesn't happen.
 
     EXPECT_EQ(Disassemble(m),
@@ -1206,12 +1179,12 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) {
     switch 1i [c: (0i, %b2), c: (default, %b3), m: %b4]
       # Case block
       %b2 = block {
-        br %b4
+        exit_switch %b4
       }
 
       # Case block
       %b3 = block {
-        br %b4
+        exit_switch %b4
       }
 
     # Merge block
@@ -1220,7 +1193,6 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) {
     }
 
   }
-
 }
 )");
 }
@@ -1254,7 +1226,7 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) {
 
     EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
     EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length());
-    EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
+    EXPECT_EQ(0u, flow->Merge()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1264,13 +1236,13 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) {
       %b2 = block {
         ret
       }
+
       # Case block
       %b3 = block {
         ret
       }
 
   }
-
 }
 )");
 }
diff --git a/src/tint/ir/from_program_unary_test.cc b/src/tint/ir/from_program_unary_test.cc
index 774be8cb51..bb58c02fd5 100644
--- a/src/tint/ir/from_program_unary_test.cc
+++ b/src/tint/ir/from_program_unary_test.cc
@@ -107,13 +107,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_AddressOf) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %v2:ptr<private, i32, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     ret
   }
 }
@@ -134,13 +131,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Indirection) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %v3:ptr<private, i32, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
-  %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+  %b2 = block {
     store %v3, 42i
     ret
   }
diff --git a/src/tint/ir/from_program_var_test.cc b/src/tint/ir/from_program_var_test.cc
index 29377e4f33..6cecf47d1b 100644
--- a/src/tint/ir/from_program_var_test.cc
+++ b/src/tint/ir/from_program_var_test.cc
@@ -35,11 +35,8 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_NoInit) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %a:ptr<private, u32, read_write> = var
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
 )");
 }
 
@@ -53,11 +50,8 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_Init) {
     EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
 %b1 = block {
   %a:ptr<private, u32, read_write> = var, 2u
-  br %b2  # root_end
 }
 
-%b2 = root_terminator
-
 )");
 }
 
diff --git a/src/tint/ir/if.cc b/src/tint/ir/if.cc
index a89a51f765..2069c6d010 100644
--- a/src/tint/ir/if.cc
+++ b/src/tint/ir/if.cc
@@ -19,7 +19,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::If);
 namespace tint::ir {
 
 If::If(Value* cond, Block* t, Block* f, Block* m)
-    : Base(m), condition_(cond), true_(t), false_(f), merge_(m) {
+    : Base(utils::Empty), condition_(cond), true_(t), false_(f), merge_(m) {
     TINT_ASSERT(IR, true_);
     TINT_ASSERT(IR, false_);
     TINT_ASSERT(IR, merge_);
diff --git a/src/tint/ir/loop.cc b/src/tint/ir/loop.cc
index 0bbb710e54..fe34283f3b 100644
--- a/src/tint/ir/loop.cc
+++ b/src/tint/ir/loop.cc
@@ -18,7 +18,8 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Loop);
 
 namespace tint::ir {
 
-Loop::Loop(Block* s, Block* c, Block* m) : Base(s), start_(s), continuing_(c), merge_(m) {
+Loop::Loop(Block* s, Block* c, Block* m)
+    : Base(utils::Empty), start_(s), continuing_(c), merge_(m) {
     TINT_ASSERT(IR, start_);
     TINT_ASSERT(IR, continuing_);
     TINT_ASSERT(IR, merge_);
diff --git a/src/tint/ir/switch.cc b/src/tint/ir/switch.cc
index a28666a6b1..003feb20c9 100644
--- a/src/tint/ir/switch.cc
+++ b/src/tint/ir/switch.cc
@@ -18,7 +18,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Switch);
 
 namespace tint::ir {
 
-Switch::Switch(Value* cond, Block* m) : Base(m), condition_(cond), merge_(m) {
+Switch::Switch(Value* cond, Block* m) : Base(utils::Empty), condition_(cond), merge_(m) {
     TINT_ASSERT(IR, condition_);
     TINT_ASSERT(IR, merge_);
     condition_->AddUsage(this);
diff --git a/src/tint/ir/to_program.cc b/src/tint/ir/to_program.cc
index af4a957a8e..08436f3780 100644
--- a/src/tint/ir/to_program.cc
+++ b/src/tint/ir/to_program.cc
@@ -20,6 +20,7 @@
 #include "src/tint/ir/block.h"
 #include "src/tint/ir/call.h"
 #include "src/tint/ir/constant.h"
+#include "src/tint/ir/exit_if.h"
 #include "src/tint/ir/if.h"
 #include "src/tint/ir/instruction.h"
 #include "src/tint/ir/load.h"
@@ -121,46 +122,27 @@ class State {
         while (block) {
             TINT_ASSERT(IR, block->HasBranchTarget());
 
-            enum Status { kContinue, kStop, kError };
-
-            Status status = tint::Switch(
-                block,
-
-                [&](const ir::Block* blk) {
-                    for (auto* inst : blk->Instructions()) {
-                        auto stmt = Stmt(inst);
-                        if (TINT_UNLIKELY(!stmt)) {
-                            return kError;
-                        }
-                        if (auto* s = stmt.Get()) {
-                            stmts.Push(s);
-                        }
-                    }
-                    if (auto* if_ = blk->Branch()->As<ir::If>()) {
-                        if (if_->Merge()->HasBranchTarget()) {
-                            block = if_->Merge();
-                            return kContinue;
-                        }
-                    } else if (auto* switch_ = blk->Branch()->As<ir::Switch>()) {
-                        if (switch_->Merge()->HasBranchTarget()) {
-                            block = switch_->Merge();
-                            return kContinue;
-                        }
-                    }
-                    return kStop;
-                },
-
-                [&](Default) {
-                    UNHANDLED_CASE(block);
-                    return kError;
-                });
-
-            if (TINT_UNLIKELY(status == kError)) {
-                return nullptr;
+            for (auto* inst : block->Instructions()) {
+                auto stmt = Stmt(inst);
+                if (TINT_UNLIKELY(!stmt)) {
+                    return nullptr;
+                }
+                if (auto* s = stmt.Get()) {
+                    stmts.Push(s);
+                }
             }
-            if (status == kStop) {
-                break;
+            if (auto* if_ = block->Branch()->As<ir::If>()) {
+                if (if_->Merge()->HasBranchTarget()) {
+                    block = if_->Merge();
+                    continue;
+                }
+            } else if (auto* switch_ = block->Branch()->As<ir::Switch>()) {
+                if (switch_->Merge()->HasBranchTarget()) {
+                    block = switch_->Merge();
+                    continue;
+                }
             }
+            break;
         }
 
         return b.Block(std::move(stmts));
@@ -174,16 +156,20 @@ class State {
             return nullptr;
         }
 
-        if (!IsEmpty(i->False(), i->Merge())) {
+        auto* false_blk = i->False();
+        if (false_blk->Instructions().Length() > 1 ||
+            (false_blk->Instructions().Length() == 1 && false_blk->HasBranchTarget() &&
+             !false_blk->Branch()->Is<ir::ExitIf>())) {
             // If the else target is an `if` which has a merge target that just bounces to the outer
             // if merge target then emit an 'else if' instead of a block statement for the else.
-            if (auto* inst = i->False()->Instructions().Front()->As<ir::If>();
-                inst && inst->Merge()->IsTrampoline(i->Merge())) {
-                auto* f = If(inst);
-                if (!f) {
-                    return nullptr;
+            if (auto* inst = i->False()->Instructions().Front()->As<ir::If>()) {
+                if (auto* br = inst->Merge()->Branch()->As<ir::ExitIf>(); br && br->If() == i) {
+                    auto* f = If(inst);
+                    if (!f) {
+                        return nullptr;
+                    }
+                    return b.If(cond, t, b.Else(f));
                 }
-                return b.If(cond, t, b.Else(f));
             } else {
                 auto* f = BlockGraph(i->False());
                 if (!f) {
@@ -192,7 +178,6 @@ class State {
                 return b.If(cond, t, b.Else(f));
             }
         }
-
         return b.If(cond, t);
     }
 
@@ -265,17 +250,6 @@ class State {
         return b.Return(val);
     }
 
-    /// @return true if there are no instructions between @p node and and @p stop_at
-    bool IsEmpty(const ir::Block* node, const ir::Block* stop_at) {
-        if (node->Instructions().IsEmpty()) {
-            return true;
-        }
-        if (auto* br = node->Instructions().Front()->As<Branch>()) {
-            return !br->Is<ir::Return>() && br->To() == stop_at;
-        }
-        return false;
-    }
-
     utils::Result<const ast::Statement*> Stmt(const ir::Instruction* inst) {
         return tint::Switch<utils::Result<const ast::Statement*>>(
             inst,                                            //
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.cc b/src/tint/writer/spirv/ir/generator_impl_ir.cc
index 29c89e6619..7da87c1305 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir.cc
@@ -19,6 +19,7 @@
 #include "spirv/unified1/spirv.h"
 #include "src/tint/ir/binary.h"
 #include "src/tint/ir/block.h"
+#include "src/tint/ir/exit_if.h"
 #include "src/tint/ir/if.h"
 #include "src/tint/ir/load.h"
 #include "src/tint/ir/module.h"
@@ -354,25 +355,24 @@ void GeneratorImplIr::EmitBlock(const ir::Block* block) {
 }
 
 void GeneratorImplIr::EmitBranch(const ir::Branch* b) {
-    if (b->Is<ir::Return>()) {
-        if (!b->Args().IsEmpty()) {
-            TINT_ASSERT(Writer, b->Args().Length() == 1u);
-            OperandList operands;
-            operands.push_back(Value(b->Args()[0]));
-            current_function_.push_inst(spv::Op::OpReturnValue, operands);
-        } else {
-            current_function_.push_inst(spv::Op::OpReturn, {});
-        }
-        return;
-    }
-
-    Switch(
-        b->To(),
-        [&](const ir::Block* blk) { current_function_.push_inst(spv::Op::OpBranch, {Label(blk)}); },
+    tint::Switch(  //
+        b,         //
+        [&](const ir::Return*) {
+            if (!b->Args().IsEmpty()) {
+                TINT_ASSERT(Writer, b->Args().Length() == 1u);
+                OperandList operands;
+                operands.push_back(Value(b->Args()[0]));
+                current_function_.push_inst(spv::Op::OpReturnValue, operands);
+            } else {
+                current_function_.push_inst(spv::Op::OpReturn, {});
+            }
+            return;
+        },
+        [&](const ir::ExitIf* if_) {
+            current_function_.push_inst(spv::Op::OpBranch, {Label(if_->If()->Merge())});
+        },
         [&](Default) {
-            // A block may not have an outward branch (e.g. an unreachable merge
-            // block).
-            current_function_.push_inst(spv::Op::OpUnreachable, {});
+            TINT_ICE(Writer, diagnostics_) << "unimplemented branch: " << b->TypeInfo().name;
         });
 }
 
@@ -388,10 +388,12 @@ void GeneratorImplIr::EmitIf(const ir::If* i) {
     uint32_t merge_label = Label(merge_block);
     uint32_t true_label = merge_label;
     uint32_t false_label = merge_label;
-    if (true_block->Instructions().Length() > 1 || true_block->Branch()->To() != merge_block) {
+    if (true_block->Instructions().Length() > 1 ||
+        (true_block->HasBranchTarget() && !true_block->Branch()->Is<ir::ExitIf>())) {
         true_label = Label(true_block);
     }
-    if (false_block->Instructions().Length() > 1 || false_block->Branch()->To() != merge_block) {
+    if (false_block->Instructions().Length() > 1 ||
+        (false_block->HasBranchTarget() && !false_block->Branch()->Is<ir::ExitIf>())) {
         false_label = Label(false_block);
     }
 
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc
index b526c91c32..f43a09a9f6 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc
@@ -23,8 +23,8 @@ TEST_F(SpvGeneratorImplTest, If_TrueEmpty_FalseEmpty) {
     auto* func = b.CreateFunction("foo", mod.Types().void_());
 
     auto* i = b.CreateIf(b.Constant(true));
-    i->True()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
-    i->False()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
+    i->True()->SetInstructions(utils::Vector{b.ExitIf(i)});
+    i->False()->SetInstructions(utils::Vector{b.ExitIf(i)});
     i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
 
     func->StartTarget()->SetInstructions(utils::Vector{i});
@@ -49,12 +49,12 @@ TEST_F(SpvGeneratorImplTest, If_FalseEmpty) {
     auto* func = b.CreateFunction("foo", mod.Types().void_());
 
     auto* i = b.CreateIf(b.Constant(true));
-    i->False()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
+    i->False()->SetInstructions(utils::Vector{b.ExitIf(i)});
     i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
 
     auto* true_block = i->True();
-    true_block->SetInstructions(utils::Vector{
-        b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.Branch(i->Merge())});
+    true_block->SetInstructions(
+        utils::Vector{b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.ExitIf(i)});
 
     func->StartTarget()->SetInstructions(utils::Vector{i});
 
@@ -83,12 +83,12 @@ TEST_F(SpvGeneratorImplTest, If_TrueEmpty) {
     auto* func = b.CreateFunction("foo", mod.Types().void_());
 
     auto* i = b.CreateIf(b.Constant(true));
-    i->True()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
+    i->True()->SetInstructions(utils::Vector{b.ExitIf(i)});
     i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
 
     auto* false_block = i->False();
-    false_block->SetInstructions(utils::Vector{
-        b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.Branch(i->Merge())});
+    false_block->SetInstructions(
+        utils::Vector{b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.ExitIf(i)});
 
     func->StartTarget()->SetInstructions(utils::Vector{i});
 
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc
index d66df98885..574a0b0336 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc
@@ -100,7 +100,7 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_DeclInsideBlock) {
     v->SetInitializer(b.Constant(42_i));
 
     auto* i = b.CreateIf(b.Constant(true));
-    i->True()->SetInstructions(utils::Vector{v, b.Branch(i->Merge())});
+    i->True()->SetInstructions(utils::Vector{v, b.ExitIf(i)});
     i->False()->SetInstructions(utils::Vector{b.Return(func)});
     i->Merge()->SetInstructions(utils::Vector{b.Return(func)});