diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index adf0e48..344607d 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -1214,6 +1214,8 @@
       "ir/block_param.h",
       "ir/branch.cc",
       "ir/branch.h",
+      "ir/break_if.cc",
+      "ir/break_if.h",
       "ir/builder.cc",
       "ir/builder.h",
       "ir/builtin.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 39bc28b..ff69240 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -722,6 +722,8 @@
     ir/block_param.h
     ir/branch.cc
     ir/branch.h
+    ir/break_if.cc
+    ir/break_if.h
     ir/builder.cc
     ir/builder.h
     ir/builtin.cc
diff --git a/src/tint/ir/break_if.cc b/src/tint/ir/break_if.cc
new file mode 100644
index 0000000..f19fb79
--- /dev/null
+++ b/src/tint/ir/break_if.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/break_if.h"
+
+#include "src/tint/ir/loop.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::BreakIf);
+
+namespace tint::ir {
+
+BreakIf::BreakIf(Value* condition, ir::Loop* loop)
+    : Base(utils::Empty), condition_(condition), loop_(loop) {
+    TINT_ASSERT(IR, condition_);
+    TINT_ASSERT(IR, loop_);
+    condition_->AddUsage(this);
+    loop_->AddUsage(this);
+    loop_->Start()->AddInboundBranch(this);
+    loop_->Merge()->AddInboundBranch(this);
+}
+
+BreakIf::~BreakIf() = default;
+
+}  // namespace tint::ir
diff --git a/src/tint/ir/break_if.h b/src/tint/ir/break_if.h
new file mode 100644
index 0000000..47fd4e8
--- /dev/null
+++ b/src/tint/ir/break_if.h
@@ -0,0 +1,51 @@
+// 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_BREAK_IF_H_
+#define SRC_TINT_IR_BREAK_IF_H_
+
+#include "src/tint/ir/branch.h"
+#include "src/tint/ir/value.h"
+#include "src/tint/utils/castable.h"
+
+// Forward declarations
+namespace tint::ir {
+class Loop;
+}  // namespace tint::ir
+
+namespace tint::ir {
+
+/// A break-if iteration instruction.
+class BreakIf : public utils::Castable<BreakIf, Branch> {
+  public:
+    /// Constructor
+    /// @param condition the break condition
+    /// @param loop the loop containing the break-if
+    BreakIf(Value* condition, ir::Loop* loop);
+    ~BreakIf() override;
+
+    /// @returns the break condition
+    const Value* Condition() const { return condition_; }
+
+    /// @returns the loop containing the break-if
+    const ir::Loop* Loop() const { return loop_; }
+
+  private:
+    Value* condition_ = nullptr;
+    ir::Loop* loop_ = nullptr;
+};
+
+}  // namespace tint::ir
+
+#endif  // SRC_TINT_IR_BREAK_IF_H_
diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc
index 7c94648..4bdb53d 100644
--- a/src/tint/ir/builder.cc
+++ b/src/tint/ir/builder.cc
@@ -213,6 +213,10 @@
     return ir.values.Create<ir::Return>(func, args);
 }
 
+ir::BreakIf* Builder::BreakIf(Value* condition, Loop* loop) {
+    return ir.values.Create<ir::BreakIf>(condition, loop);
+}
+
 ir::Continue* Builder::Continue(Loop* loop) {
     return ir.values.Create<ir::Continue>(loop);
 }
diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h
index a63515e..9a2fe18 100644
--- a/src/tint/ir/builder.h
+++ b/src/tint/ir/builder.h
@@ -21,6 +21,7 @@
 #include "src/tint/ir/binary.h"
 #include "src/tint/ir/bitcast.h"
 #include "src/tint/ir/block_param.h"
+#include "src/tint/ir/break_if.h"
 #include "src/tint/ir/builtin.h"
 #include "src/tint/ir/constant.h"
 #include "src/tint/ir/construct.h"
@@ -337,6 +338,12 @@
     /// @returns the instruction
     ir::Return* Return(Function* func, utils::VectorRef<Value*> args = {});
 
+    /// Creates a loop break-if instruction
+    /// @param condition the break condition
+    /// @param loop the loop being iterated
+    /// @returns the instruction
+    ir::BreakIf* BreakIf(Value* condition, Loop* loop);
+
     /// Creates a continue instruction
     /// @param loop the loop being continued
     /// @returns the instruction
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index ae89902..6097726 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -21,6 +21,7 @@
 #include "src/tint/ir/binary.h"
 #include "src/tint/ir/bitcast.h"
 #include "src/tint/ir/block.h"
+#include "src/tint/ir/break_if.h"
 #include "src/tint/ir/builtin.h"
 #include "src/tint/ir/construct.h"
 #include "src/tint/ir/continue.h"
@@ -414,6 +415,10 @@
         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>()) {
diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc
index 444950e..a4d95e0 100644
--- a/src/tint/ir/from_program.cc
+++ b/src/tint/ir/from_program.cc
@@ -590,7 +590,9 @@
                     EmitBlock(stmt->continuing);
                 }
                 // Branch back to the start node if the continue target didn't branch out already
-                BranchToIfNeeded(loop_inst->Start());
+                if (NeedBranch()) {
+                    SetBranch(builder_.BreakIf(builder_.Constant(false), loop_inst));
+                }
             }
         }
 
@@ -608,7 +610,8 @@
         current_flow_block_->Instructions().Push(loop_inst);
 
         // Continue is always empty, just go back to the start
-        loop_inst->Continuing()->Instructions().Push(builder_.Branch(loop_inst->Start()));
+        current_flow_block_ = loop_inst->Continuing();
+        SetBranch(builder_.BreakIf(builder_.Constant(false), loop_inst));
 
         {
             ControlStackScope scope(this, loop_inst);
@@ -681,7 +684,7 @@
             if (stmt->continuing) {
                 current_flow_block_ = loop_inst->Continuing();
                 EmitStatement(stmt->continuing);
-                loop_inst->Continuing()->Instructions().Push(builder_.Branch(loop_inst->Start()));
+                SetBranch(builder_.BreakIf(builder_.Constant(false), loop_inst));
             }
         }
 
@@ -772,31 +775,14 @@
     }
 
     void EmitBreakIf(const ast::BreakIfStatement* stmt) {
+        auto* current_control = FindEnclosingControl(ControlFlags::kExcludeSwitch);
+
         // Emit the break-if condition into the end of the preceding block
-        auto reg = EmitExpression(stmt->condition);
-        if (!reg) {
+        auto cond = EmitExpression(stmt->condition);
+        if (!cond) {
             return;
         }
-        auto* if_inst = builder_.CreateIf(reg.Get());
-        current_flow_block_->Instructions().Push(if_inst);
-
-        auto* current_control = FindEnclosingControl(ControlFlags::kExcludeSwitch);
-        TINT_ASSERT(IR, current_control);
-        TINT_ASSERT(IR, current_control->Is<Loop>());
-
-        auto* loop = current_control->As<Loop>();
-
-        current_flow_block_ = if_inst->True();
-        BranchTo(loop->Merge());
-
-        current_flow_block_ = if_inst->False();
-        BranchTo(if_inst->Merge());
-
-        current_flow_block_ = if_inst->Merge();
-
-        // The `break-if` has to be the last item in the continuing block. The false branch of
-        // the `break-if` will always take us back to the start of the loop.
-        BranchTo(loop->Start());
+        SetBranch(builder_.BreakIf(cond.Get(), current_control->As<ir::Loop>()));
     }
 
     utils::Result<Value*> EmitExpression(const ast::Expression* expr) {
diff --git a/src/tint/ir/from_program_test.cc b/src/tint/ir/from_program_test.cc
index 0b23ec5..e0f161d 100644
--- a/src/tint/ir/from_program_test.cc
+++ b/src/tint/ir/from_program_test.cc
@@ -385,7 +385,7 @@
 
     EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
     EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
-    EXPECT_EQ(1u, loop_flow->Merge()->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());
@@ -416,7 +416,7 @@
 
       # Continuing block
       %b3 = block {
-        br %b2
+        break_if false %b2
       }
 
     # Merge block
@@ -440,16 +440,12 @@
 
     auto m = res.Move();
     auto* loop_flow = FindSingleValue<ir::Loop>(m);
-    auto* break_if_flow = FindSingleValue<ir::If>(m);
 
     ASSERT_EQ(1u, m.functions.Length());
 
     EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
     EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
     EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length());
-    EXPECT_EQ(1u, break_if_flow->True()->InboundBranches().Length());
-    EXPECT_EQ(1u, break_if_flow->False()->InboundBranches().Length());
-    EXPECT_EQ(2u, break_if_flow->Merge()->InboundBranches().Length());
 
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -461,23 +457,7 @@
 
       # Continuing block
       %b3 = block {
-        if true [t: %b5, f: %b6, m: %b7]
-          # True block
-          %b5 = block {
-            br %b4
-          }
-
-          # False block
-          %b6 = block {
-            br %b7
-          }
-
-        # Merge block
-        %b7 = block {
-          br %b2
-        }
-
-
+        break_if true %b2
       }
 
     # Merge block
@@ -511,23 +491,7 @@
 
       # Continuing block
       %b3 = block {
-        if true [t: %b5, f: %b6, m: %b7]
-          # True block
-          %b5 = block {
-            br %b4
-          }
-
-          # False block
-          %b6 = block {
-            br %b7
-          }
-
-        # Merge block
-        %b7 = block {
-          br %b2
-        }
-
-
+        break_if true %b2
       }
 
     # Merge block
@@ -557,7 +521,7 @@
 
     EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
     EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
-    EXPECT_EQ(0u, loop_flow->Merge()->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());
@@ -565,20 +529,20 @@
     EXPECT_EQ(Disassemble(m),
               R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
   %b1 = block {
-    loop [s: %b2, c: %b3]
+    loop [s: %b2, c: %b3, m: %b4]
       %b2 = block {
-        if true [t: %b4, f: %b5, m: %b6]
+        if true [t: %b5, f: %b6, m: %b7]
           # True block
-          %b4 = block {
+          %b5 = block {
             ret
           }
           # False block
-          %b5 = block {
-            br %b6
+          %b6 = block {
+            br %b7
           }
 
         # Merge block
-        %b6 = block {
+        %b7 = block {
           continue %b3
         }
 
@@ -587,9 +551,13 @@
 
       # Continuing block
       %b3 = block {
-        br %b2
+        break_if false %b2
       }
 
+    # Merge block
+    %b4 = block {
+      ret
+    }
 
   }
 
@@ -792,28 +760,12 @@
 
                 # Continuing block
                 %b17 = block {
-                  if true [t: %b19, f: %b20, m: %b21]
-                    # True block
-                    %b19 = block {
-                      br %b18
-                    }
-
-                    # False block
-                    %b20 = block {
-                      br %b21
-                    }
-
-                  # Merge block
-                  %b21 = block {
-                    br %b16
-                  }
-
-
+                  break_if true %b16
                 }
 
               # Merge block
               %b18 = block {
-                br %b5
+                break_if false %b5
               }
 
 
@@ -824,19 +776,19 @@
 
         # Merge block
         %b7 = block {
-          if true [t: %b22, f: %b23, m: %b24]
+          if true [t: %b19, f: %b20, m: %b21]
             # True block
-            %b22 = block {
+            %b19 = block {
               br %b4
             }
 
             # False block
-            %b23 = block {
-              br %b24
+            %b20 = block {
+              br %b21
             }
 
           # Merge block
-          %b24 = block {
+          %b21 = block {
             continue %b3
           }
 
@@ -848,7 +800,7 @@
 
       # Continuing block
       %b3 = block {
-        br %b2
+        break_if false %b2
       }
 
     # Merge block
@@ -880,7 +832,7 @@
 
     EXPECT_EQ(2u, flow->Start()->InboundBranches().Length());
     EXPECT_EQ(1u, flow->Continuing()->InboundBranches().Length());
-    EXPECT_EQ(1u, flow->Merge()->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());
@@ -911,7 +863,7 @@
 
       # Continuing block
       %b3 = block {
-        br %b2
+        break_if false %b2
       }
 
     # Merge block
@@ -943,7 +895,7 @@
 
     EXPECT_EQ(2u, flow->Start()->InboundBranches().Length());
     EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
-    EXPECT_EQ(1u, flow->Merge()->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());
@@ -973,7 +925,7 @@
 
       # Continuing block
       %b3 = block {
-        br %b2
+        break_if false %b2
       }
 
     # Merge block