diff --git a/src/dawn/tests/end2end/InfiniteLoopTests.cpp b/src/dawn/tests/end2end/InfiniteLoopTests.cpp
index e0c321c..21172f4 100644
--- a/src/dawn/tests/end2end/InfiniteLoopTests.cpp
+++ b/src/dawn/tests/end2end/InfiniteLoopTests.cpp
@@ -269,7 +269,10 @@
 };
 
 TEST_P(InfiniteLoopTests, LoopDeletedThenBoundedWrite) {
-    DAWN_SKIP_TEST_IF_BASE(!IsMetal(), "platform-specific", "only test on Metal");
+    DAWN_SKIP_TEST_IF_BASE(!IsMetal(), "infinite-loops", "only test on Metal");
+    DAWN_SKIP_TEST_IF_BASE(
+        IsMetal(), "infinite-loops",
+        "Metal loops run forever: TODO(crbug.com/371840056) rewrite as death tests with watchdog?");
 
     const uint32_t sentinelValue = 777;
     std::string shader = Shader(sentinelValue);
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
index e1d16b9..5d61ad4 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -2324,9 +2324,10 @@
     };
 
     TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
-    Line() << IsolateUB() << " while(true) {";
+    Line() << "while(true) {";
     {
         ScopedIndent si(this);
+        Line() << IsolateUB();
         if (!EmitStatements(stmt->body->statements)) {
             return false;
         }
@@ -2394,12 +2395,13 @@
         };
 
         TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
-        Line() << IsolateUB() << " while(true) {";
+        Line() << "while(true) {";
         IncrementIndent();
         TINT_DEFER({
             DecrementIndent();
             Line() << "}";
         });
+        Line() << IsolateUB();
 
         if (stmt->condition) {
             current_buffer_->Append(cond_pre);
@@ -2417,7 +2419,7 @@
         // For-loop can be generated.
         {
             auto out = Line();
-            out << IsolateUB() << " for";
+            out << "for";
             {
                 ScopedParen sp(out);
 
@@ -2436,6 +2438,9 @@
             out << " {";
         }
         {
+            IncrementIndent();
+            Line() << IsolateUB();
+            DecrementIndent();
             auto emit_continuing = [] { return true; };
             TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
             if (!EmitStatementsWithIndent(stmt->body->statements)) {
@@ -2467,8 +2472,9 @@
     // as a regular while in MSL. Instead we need to generate a `while(true)` loop.
     bool emit_as_loop = cond_pre.lines.size() > 0;
     if (emit_as_loop) {
-        Line() << IsolateUB() << " while(true) {";
+        Line() << "while(true) {";
         IncrementIndent();
+        Line() << IsolateUB();
         TINT_DEFER({
             DecrementIndent();
             Line() << "}";
@@ -2481,7 +2487,10 @@
         }
     } else {
         // While can be generated.
-        Line() << IsolateUB() << " while(" << cond_buf.str() << ") {";
+        Line() << "while(" << cond_buf.str() << ") {";
+        IncrementIndent();
+        Line() << IsolateUB();
+        DecrementIndent();
         if (!EmitStatementsWithIndent(stmt->body->statements)) {
             return false;
         }
@@ -3240,12 +3249,11 @@
     if (isolate_ub_macro_name_.empty()) {
         isolate_ub_macro_name_ = UniqueIdentifier("TINT_ISOLATE_UB");
         Line(&helpers_) << "#define " << isolate_ub_macro_name_ << "(VOLATILE_NAME) \\";
-        Line(&helpers_) << "  volatile bool VOLATILE_NAME = true; \\";
-        Line(&helpers_) << "  if (VOLATILE_NAME)";
+        Line(&helpers_) << "  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}";
         Line(&helpers_);
     }
     StringStream ss;
-    ss << isolate_ub_macro_name_ << "(" << UniqueIdentifier("tint_volatile_true") << ")";
+    ss << isolate_ub_macro_name_ << "(" << UniqueIdentifier("tint_volatile_false") << ");";
     return ss.str();
 }
 
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.h b/src/tint/lang/msl/writer/ast_printer/ast_printer.h
index 8004d61..2f7542b 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.h
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.h
@@ -373,8 +373,13 @@
                               const ast::CallExpression* expr,
                               const sem::BuiltinFn* builtin);
 
-    /// Lazilly generates the TINT_ISOLATE_UB macro, used to prevent UB statements from affecting
-    /// later logic.
+    /// Lazily generates the TINT_ISOLATE_UB macro, and returns a call to
+    /// the macro, passing in a unique identifier. The call tricks the MSL
+    /// compiler into thinking it might execute a `break`, but otherwise
+    /// has no effect in the generated code.
+    ///
+    /// Invoke this inside the body of a loop to prevent the MSL compiler
+    /// from inferring the loop never terminates.
     /// @return the MSL to call the TINT_ISOLATE_UB macro.
     std::string IsolateUB();
 
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
index 4f60358..bfc485a 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
@@ -231,15 +231,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_4 {
   tint_array<float2x2, 4> m;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<float2x2, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol))[i] = float2x2(float2(0.0f), float2(0.0f));
   }
diff --git a/src/tint/lang/msl/writer/ast_printer/continue_test.cc b/src/tint/lang/msl/writer/ast_printer/continue_test.cc
index ad68153..266daaf 100644
--- a/src/tint/lang/msl/writer/ast_printer/continue_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/continue_test.cc
@@ -45,11 +45,11 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void test_function() {
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     if (false) {
       break;
     }
diff --git a/src/tint/lang/msl/writer/ast_printer/loop_test.cc b/src/tint/lang/msl/writer/ast_printer/loop_test.cc
index 3d1dd64..f2ac3b1 100644
--- a/src/tint/lang/msl/writer/ast_printer/loop_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/loop_test.cc
@@ -52,11 +52,11 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 fragment void F() {
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     break;
   }
   return;
@@ -82,14 +82,14 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void a_statement() {
 }
 
 fragment void F() {
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     break;
     {
       a_statement();
@@ -118,14 +118,14 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void a_statement() {
 }
 
 fragment void F() {
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     break;
     {
       a_statement();
@@ -159,8 +159,10 @@
     ASTPrinter& gen = Build();
 
     ASSERT_TRUE(gen.EmitStatement(outer)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB(tint_volatile_true) while(true) {
-  TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    EXPECT_EQ(gen.Result(), R"(while(true) {
+  TINT_ISOLATE_UB(tint_volatile_false);
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
     break;
     {
       a_statement();
@@ -197,7 +199,8 @@
     ASTPrinter& gen = Build();
 
     ASSERT_TRUE(gen.EmitStatement(outer)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    EXPECT_EQ(gen.Result(), R"(while(true) {
+  TINT_ISOLATE_UB(tint_volatile_false);
   float lhs = 2.5f;
   float other = 0.0f;
   break;
@@ -220,7 +223,8 @@
     ASTPrinter& gen = Build();
 
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB(tint_volatile_true) for(; ; ) {
+    EXPECT_EQ(gen.Result(), R"(for(; ; ) {
+  TINT_ISOLATE_UB(tint_volatile_false);
   return;
 }
 )");
@@ -238,7 +242,8 @@
     ASTPrinter& gen = Build();
 
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; ; ) {
+    EXPECT_EQ(gen.Result(), R"(for(int i = 0; ; ) {
+  TINT_ISOLATE_UB(tint_volatile_false);
   return;
 }
 )");
@@ -269,7 +274,8 @@
     f(1);
     f(2);
   }
-  TINT_ISOLATE_UB(tint_volatile_true) for(; ; ) {
+  for(; ; ) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     return;
   }
 }
@@ -288,7 +294,8 @@
     ASTPrinter& gen = Build();
 
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB(tint_volatile_true) for(; true; ) {
+    EXPECT_EQ(gen.Result(), R"(for(; true; ) {
+  TINT_ISOLATE_UB(tint_volatile_false);
   return;
 }
 )");
@@ -307,9 +314,9 @@
     ASTPrinter& gen = Build();
 
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(
-        gen.Result(),
-        R"(TINT_ISOLATE_UB(tint_volatile_true) for(; ; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    EXPECT_EQ(gen.Result(),
+              R"(for(; ; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  TINT_ISOLATE_UB(tint_volatile_false);
   return;
 }
 )");
@@ -335,7 +342,8 @@
     ASTPrinter& gen = Build();
 
     ASSERT_TRUE(gen.EmitStatement(loop)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    EXPECT_EQ(gen.Result(), R"(while(true) {
+  TINT_ISOLATE_UB(tint_volatile_false);
   return;
   {
     f(1);
@@ -359,9 +367,9 @@
     ASTPrinter& gen = Build();
 
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(
-        gen.Result(),
-        R"(TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; true; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    EXPECT_EQ(gen.Result(),
+              R"(for(int i = 0; true; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  TINT_ISOLATE_UB(tint_volatile_false);
   a_statement();
 }
 )");
@@ -393,7 +401,8 @@
     f(1);
     f(2);
   }
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     if (!(true)) { break; }
     return;
     {
@@ -416,7 +425,8 @@
     ASTPrinter& gen = Build();
 
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    EXPECT_EQ(gen.Result(), R"(while(true) {
+  TINT_ISOLATE_UB(tint_volatile_false);
   return;
 }
 )");
@@ -433,7 +443,8 @@
     ASTPrinter& gen = Build();
 
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    EXPECT_EQ(gen.Result(), R"(while(true) {
+  TINT_ISOLATE_UB(tint_volatile_false);
   continue;
 }
 )");
@@ -453,7 +464,8 @@
     ASTPrinter& gen = Build();
 
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB(tint_volatile_true) while((t && false)) {
+    EXPECT_EQ(gen.Result(), R"(while((t && false)) {
+  TINT_ISOLATE_UB(tint_volatile_false);
   return;
 }
 )");
diff --git a/test/tint/array/assign_to_function_var.wgsl.expected.msl b/test/tint/array/assign_to_function_var.wgsl.expected.msl
index 6b23fac..872caa8 100644
--- a/test/tint/array/assign_to_function_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_function_var.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   tint_array<int4, 4> src_private;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<int4, 4>* const tint_symbol_5) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol_5))[i] = int4(0);
   }
diff --git a/test/tint/array/assign_to_private_var.wgsl.expected.msl b/test/tint/array/assign_to_private_var.wgsl.expected.msl
index fb4dbcd..fecfa2b 100644
--- a/test/tint/array/assign_to_private_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_private_var.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   tint_array<int4, 4> src_private;
@@ -25,7 +24,8 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<int4, 4>* const tint_symbol_5) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol_5))[i] = int4(0);
   }
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.msl b/test/tint/array/assign_to_storage_var.wgsl.expected.msl
index f477833..7ce6ea8 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   tint_array<int4, 4> src_private;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<int4, 4>* const tint_symbol_5) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol_5))[i] = int4(0);
   }
diff --git a/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl b/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl
index 5973bc4..63276cb 100644
--- a/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl
@@ -15,20 +15,21 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   tint_array<int4, 4> src_private;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<int4, 4>* const tint_symbol_5, threadgroup tint_array<int4, 4>* const tint_symbol_6, threadgroup tint_array<tint_array<tint_array<int, 2>, 3>, 4>* const tint_symbol_7) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol_5))[i] = int4(0);
     (*(tint_symbol_6))[i] = int4(0);
   }
-  TINT_ISOLATE_UB(tint_volatile_true_1) for(uint idx_1 = local_idx; (idx_1 < 24u); idx_1 = (idx_1 + 1u)) {
+  for(uint idx_1 = local_idx; (idx_1 < 24u); idx_1 = (idx_1 + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
     uint const i_1 = (idx_1 / 6u);
     uint const i_2 = ((idx_1 % 6u) / 2u);
     uint const i_3 = (idx_1 % 2u);
diff --git a/test/tint/array/strides.spvasm.expected.msl b/test/tint/array/strides.spvasm.expected.msl
index 331afc2..4c70c23 100644
--- a/test/tint/array/strides.spvasm.expected.msl
+++ b/test/tint/array/strides.spvasm.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct strided_arr {
   /* 0x0000 */ float el;
@@ -37,13 +36,15 @@
 }
 
 void assign_and_preserve_padding_3(device tint_array<strided_arr, 2>* const dest, tint_array<strided_arr, 2> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 2u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 2u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_4(&((*(dest))[i]), value[i]);
   }
 }
 
 void assign_and_preserve_padding_2(device tint_array<tint_array<strided_arr, 2>, 3>* const dest, tint_array<tint_array<strided_arr, 2>, 3> value) {
-  TINT_ISOLATE_UB(tint_volatile_true_1) for(uint i = 0u; (i < 3u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 3u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
     assign_and_preserve_padding_3(&((*(dest))[i]), value[i]);
   }
 }
@@ -53,7 +54,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<strided_arr_1, 4>* const dest, tint_array<strided_arr_1, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true_2) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false_2);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl
index beaa591..37da0cf 100644
--- a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl
+++ b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
@@ -106,7 +105,8 @@
 }
 
 void assign_and_preserve_padding_3(device tint_array<tint_packed_vec3_f32_array_element, 2>* const dest, tint_array<float3, 2> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 2u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 2u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     (*(dest))[i].elements = packed_float3(value[i]);
   }
 }
diff --git a/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.msl b/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.msl
index ba02219..6ca902a 100644
--- a/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.msl
+++ b/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
@@ -161,7 +160,8 @@
 }
 
 void assign_and_preserve_padding_6(device tint_array<tint_packed_vec3_f32_array_element, 2>* const dest, tint_array<float3, 2> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 2u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 2u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     (*(dest))[i].elements = packed_float3(value[i]);
   }
 }
diff --git a/test/tint/buffer/storage/static_index/write.wgsl.expected.msl b/test/tint/buffer/storage/static_index/write.wgsl.expected.msl
index 69afba3..dddbdcc 100644
--- a/test/tint/buffer/storage/static_index/write.wgsl.expected.msl
+++ b/test/tint/buffer/storage/static_index/write.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
@@ -108,7 +107,8 @@
 }
 
 void assign_and_preserve_padding_3(device tint_array<tint_packed_vec3_f32_array_element, 2>* const dest, tint_array<float3, 2> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 2u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 2u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     (*(dest))[i].elements = packed_float3(value[i]);
   }
 }
diff --git a/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.msl b/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.msl
index ef66275..52e6d56 100644
--- a/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.msl
+++ b/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
@@ -165,7 +164,8 @@
 }
 
 void assign_and_preserve_padding_6(device tint_array<tint_packed_vec3_f32_array_element, 2>* const dest, tint_array<float3, 2> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 2u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 2u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     (*(dest))[i].elements = packed_float3(value[i]);
   }
 }
@@ -177,7 +177,8 @@
 }
 
 void assign_and_preserve_padding_8(device tint_array<Inner, 4>* const dest, tint_array<Inner, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true_1) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
     assign_and_preserve_padding_7(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.msl
index 24e6469..972347e 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_6 {
   tint_array<float2x2, 4> w;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<float2x2, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol))[i] = float2x2(float2(0.0f), float2(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl
index 064fbc7..629f1b7 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f16_array_element {
   /* 0x0000 */ packed_half3 elements;
@@ -39,7 +38,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* const dest, tint_array<half2x3, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.msl
index c83e23b..4db31c8 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_8 {
   tint_array<half2x3, 4> w;
@@ -38,7 +37,8 @@
 }
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<half2x3, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol))[i] = half2x3(half3(0.0h), half3(0.0h));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl
index caa5898..b3a4701 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
@@ -39,7 +38,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* const dest, tint_array<float2x3, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.msl
index a4a3284..0220335 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_6 {
   tint_array<float2x3, 4> w;
@@ -38,7 +37,8 @@
 }
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<float2x3, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol))[i] = float2x3(float3(0.0f), float3(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.msl
index 448ea42..ea42e796 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_6 {
   tint_array<half2x4, 4> w;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<half2x4, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol))[i] = half2x4(half4(0.0h), half4(0.0h));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.msl
index 2fa5365..2b487fa 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_6 {
   tint_array<float2x4, 4> w;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<float2x4, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol))[i] = float2x4(float4(0.0f), float4(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl
index 774ab90..af77aeb 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
@@ -40,7 +39,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* const dest, tint_array<float3x3, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.msl
index ad83201..fc315e0 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_6 {
   tint_array<float3x3, 4> w;
@@ -38,7 +37,8 @@
 }
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<float3x3, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol))[i] = float3x3(float3(0.0f), float3(0.0f), float3(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.msl
index 2a269f7..c629971 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_6 {
   tint_array<float3x4, 4> w;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<float3x4, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol))[i] = float3x4(float4(0.0f), float4(0.0f), float4(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.msl
index fda58c9..309bd54 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_6 {
   tint_array<half4x2, 4> w;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<half4x2, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol))[i] = half4x2(half2(0.0h), half2(0.0h), half2(0.0h), half2(0.0h));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.msl
index 89d931f..c30d5be 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_6 {
   tint_array<float4x2, 4> w;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<float4x2, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol))[i] = float4x2(float2(0.0f), float2(0.0f), float2(0.0f), float2(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl
index 4193d09..fc75a24 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f16_array_element {
   /* 0x0000 */ packed_half3 elements;
@@ -41,7 +40,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* const dest, tint_array<half4x3, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.msl
index 159bd76..d1b8c98 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_6 {
   tint_array<half4x3, 4> w;
@@ -38,7 +37,8 @@
 }
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<half4x3, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol))[i] = half4x3(half3(0.0h), half3(0.0h), half3(0.0h), half3(0.0h));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl
index 3aa3e7d..8898c0b 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
@@ -41,7 +40,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* const dest, tint_array<float4x3, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.msl
index 2f82577..3d83a69 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_6 {
   tint_array<float4x3, 4> w;
@@ -38,7 +37,8 @@
 }
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<float4x3, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol))[i] = float4x3(float3(0.0f), float3(0.0f), float3(0.0f), float3(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.msl
index 7ed4a40..dfbd2b7 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_6 {
   tint_array<half4x4, 4> w;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<half4x4, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol))[i] = half4x4(half4(0.0h), half4(0.0h), half4(0.0h), half4(0.0h));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.msl
index 1a07b1d..0e7c57c 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_6 {
   tint_array<float4x4, 4> w;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<float4x4, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol))[i] = float4x4(float4(0.0f), float4(0.0f), float4(0.0f), float4(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl
index cb8a8d8..b80e4f5 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -33,7 +32,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.msl
index 43d8d35..2da2d6f 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -31,7 +30,8 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl
index 0084bbe..3594727 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -34,7 +33,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.msl
index 62f3d27..2f17b98 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -32,7 +31,8 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl
index 2d8a663..1e2dc8e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f16_array_element {
   /* 0x0000 */ packed_half3 elements;
@@ -68,7 +67,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.msl
index 32a680e..07f4037 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   int before;
@@ -61,7 +60,8 @@
 }
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl
index 6eac46f..81fad65 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
@@ -68,7 +67,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.msl
index 0c87556..7268dcd 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   int before;
@@ -61,7 +60,8 @@
 }
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl
index ccc6b0a..e05f119 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -34,7 +33,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.msl
index 758eb27..a342e45 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -32,7 +31,8 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl
index ed35a67..8eefeb9 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -34,7 +33,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.msl
index 81b6ebf..cb2a119 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -32,7 +31,8 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl
index 805b281..8d07ecd 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -33,7 +32,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.msl
index 23b0e59..1176487 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -31,7 +30,8 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl
index 005e620..918b124 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -34,7 +33,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.msl
index d74aaff..6023e07 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -32,7 +31,8 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl
index 655bc2d..6ac1f82 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f16_array_element {
   /* 0x0000 */ packed_half3 elements;
@@ -69,7 +68,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.msl
index f982bed..5ce7f2a 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   int before;
@@ -61,7 +60,8 @@
 }
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl
index 07d8848..6fd37f6 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
@@ -68,7 +67,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.msl
index f86d50e..92e880e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   int before;
@@ -60,7 +59,8 @@
 }
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl
index 71ee1ce..25e0dcf 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -34,7 +33,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.msl
index 052885c..e06b7b1 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -32,7 +31,8 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl
index 4ca161f..e92a747 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -33,7 +32,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.msl
index e99d6c9..3425faa 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -31,7 +30,8 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl
index 9350b1f..8826480 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -33,7 +32,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.msl
index c5e0a04..3c5a01a 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -31,7 +30,8 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl
index f463cad..2b4e5f3 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -34,7 +33,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.msl
index 36d1cca..a1a2e9d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -32,7 +31,8 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl
index a32e29f..9e9a5ca 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f16_array_element {
   /* 0x0000 */ packed_half3 elements;
@@ -70,7 +69,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.msl
index f4055dd..c1071f3 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   int before;
@@ -61,7 +60,8 @@
 }
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl
index eebace9..53ad59b 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
@@ -70,7 +69,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.msl
index daf7a4e..04664b4 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   int before;
@@ -61,7 +60,8 @@
 }
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl
index ed271f9..83a0dda 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -34,7 +33,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.msl
index 23ecd0d..21fda59 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -32,7 +31,8 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl
index e62b530..582e0f2 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -34,7 +33,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.msl
index c37dd00..8e1b9e3 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   /* 0x0000 */ int before;
@@ -32,7 +31,8 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/bug/chromium/1403752.wgsl.expected.msl b/test/tint/bug/chromium/1403752.wgsl.expected.msl
index 4b4bc84..bbfdc8b 100644
--- a/test/tint/bug/chromium/1403752.wgsl.expected.msl
+++ b/test/tint/bug/chromium/1403752.wgsl.expected.msl
@@ -3,12 +3,12 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void d() {
   int j = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) for(; false; ) {
+  for(; false; ) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/bug/chromium/1449538.wgsl.expected.msl b/test/tint/bug/chromium/1449538.wgsl.expected.msl
index f317e66..be2c460 100644
--- a/test/tint/bug/chromium/1449538.wgsl.expected.msl
+++ b/test/tint/bug/chromium/1449538.wgsl.expected.msl
@@ -3,25 +3,32 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void f() {
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i0520 = 0; false; ) {
+  for(int i0520 = 0; false; ) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
-  TINT_ISOLATE_UB(tint_volatile_true_1) for(int i62 = 0; false; ) {
+  for(int i62 = 0; false; ) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
   }
-  TINT_ISOLATE_UB(tint_volatile_true_2) for(int i0520 = 0; false; ) {
+  for(int i0520 = 0; false; ) {
+    TINT_ISOLATE_UB(tint_volatile_false_2);
   }
-  TINT_ISOLATE_UB(tint_volatile_true_3) for(int i62 = 0; false; ) {
+  for(int i62 = 0; false; ) {
+    TINT_ISOLATE_UB(tint_volatile_false_3);
   }
-  TINT_ISOLATE_UB(tint_volatile_true_4) for(int i62 = 0; false; ) {
+  for(int i62 = 0; false; ) {
+    TINT_ISOLATE_UB(tint_volatile_false_4);
   }
-  TINT_ISOLATE_UB(tint_volatile_true_5) for(int i60 = 0; false; ) {
+  for(int i60 = 0; false; ) {
+    TINT_ISOLATE_UB(tint_volatile_false_5);
   }
-  TINT_ISOLATE_UB(tint_volatile_true_6) for(int i62 = 0; false; ) {
+  for(int i62 = 0; false; ) {
+    TINT_ISOLATE_UB(tint_volatile_false_6);
   }
-  TINT_ISOLATE_UB(tint_volatile_true_7) for(int i60 = 0; false; ) {
+  for(int i60 = 0; false; ) {
+    TINT_ISOLATE_UB(tint_volatile_false_7);
   }
 }
 
diff --git a/test/tint/bug/chromium/344265982.wgsl.expected.msl b/test/tint/bug/chromium/344265982.wgsl.expected.msl
index b9966d0..144c609 100644
--- a/test/tint/bug/chromium/344265982.wgsl.expected.msl
+++ b/test/tint/bug/chromium/344265982.wgsl.expected.msl
@@ -15,12 +15,12 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void foo(device tint_array<int, 4>* const arg) {
   bool tint_continue = false;
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     tint_continue = false;
     switch((*(arg))[i]) {
       case 1: {
diff --git a/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl b/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl
index a52486c..796ab73 100644
--- a/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   tint_array<int, 64> data;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup S* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 64u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 64u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol)).data[i] = 0;
   }
diff --git a/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl b/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl
index 5b651d5..46fa1cd 100644
--- a/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   tint_array<int, 64> data;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup S* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 64u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 64u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol)).data[i] = 0;
   }
diff --git a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl
index 7efaf69..674321f 100644
--- a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl
+++ b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl
@@ -3,8 +3,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_3 {
   float2 vUV [[user(locn0)]];
@@ -18,7 +17,8 @@
   float4 const tint_symbol_1 = tint_symbol_5.sample(tint_symbol_6, vUV);
   float3 const random = tint_symbol_1.rgb;
   int i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     if ((i < 1)) {
     } else {
       break;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
index 54555c1..4670e28 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
@@ -3,8 +3,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   float2 v2f;
@@ -14,7 +13,8 @@
 };
 
 void foo(thread tint_private_vars_struct* const tint_private_vars) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     (*(tint_private_vars)).v2f[i] = 1.0f;
     (*(tint_private_vars)).v3i[i] = 1;
     (*(tint_private_vars)).v4u[i] = 1u;
@@ -24,7 +24,8 @@
 
 kernel void tint_symbol() {
   thread tint_private_vars_struct tint_private_vars = {};
-  TINT_ISOLATE_UB(tint_volatile_true_1) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
     foo(&(tint_private_vars));
   }
   return;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
index 0be7944..fe9c834 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
@@ -3,8 +3,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   float2 v2f;
@@ -23,7 +22,8 @@
 
 kernel void tint_symbol() {
   thread tint_private_vars_struct tint_private_vars = {};
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     foo(&(tint_private_vars));
   }
   return;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
index 2304b9f..d76e9ef 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
@@ -3,8 +3,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol() {
   float2 v2f = 0.0f;
@@ -19,7 +18,8 @@
   bool2 v2b = false;
   bool3 v3b = false;
   bool4 v4b = false;
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     v2f[i] = 1.0f;
     v3f[i] = 1.0f;
     v4f[i] = 1.0f;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
index 478a818..8cfeedb 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
@@ -3,8 +3,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol() {
   float2 v2f = 0.0f;
@@ -15,7 +14,8 @@
   uint4 v4u_2 = 0u;
   bool2 v2b = false;
   bool2 v2b_2 = false;
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     v2f[i] = 1.0f;
     v3i[i] = 1;
     v4u[i] = 1u;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
index 3faf230..705ba69 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
@@ -3,8 +3,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol() {
   float2 v2f = 0.0f;
@@ -19,7 +18,8 @@
   bool2 v2b = false;
   bool3 v3b = false;
   bool4 v4b = false;
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     v2f[i] = 1.0f;
     v2i[i] = 1;
     v2u[i] = 1u;
diff --git a/test/tint/bug/tint/1064.wgsl.expected.msl b/test/tint/bug/tint/1064.wgsl.expected.msl
index 4ea4a7c..b0bf1a3 100644
--- a/test/tint/bug/tint/1064.wgsl.expected.msl
+++ b/test/tint/bug/tint/1064.wgsl.expected.msl
@@ -3,11 +3,11 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 fragment void tint_symbol() {
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     if (false) {
     } else {
       break;
diff --git a/test/tint/bug/tint/1081.wgsl.expected.msl b/test/tint/bug/tint/1081.wgsl.expected.msl
index 64f9887..d955e00 100644
--- a/test/tint/bug/tint/1081.wgsl.expected.msl
+++ b/test/tint/bug/tint/1081.wgsl.expected.msl
@@ -3,8 +3,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   bool tint_discarded;
@@ -27,7 +26,8 @@
 
 int tint_symbol_inner(int3 x, thread tint_private_vars_struct* const tint_private_vars) {
   int y = x[0];
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     int const r = f(y, tint_private_vars);
     if ((r == 0)) {
       break;
diff --git a/test/tint/bug/tint/1121.wgsl.expected.msl b/test/tint/bug/tint/1121.wgsl.expected.msl
index bf996f9..5d77cc0 100644
--- a/test/tint/bug/tint/1121.wgsl.expected.msl
+++ b/test/tint/bug/tint/1121.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct LightData_tint_packed_vec3 {
   /* 0x0000 */ float4 position;
@@ -88,8 +87,10 @@
   int const TILE_SIZE = 16;
   int const TILE_COUNT_X = 2;
   int const TILE_COUNT_Y = 2;
-  TINT_ISOLATE_UB(tint_volatile_true) for(int y = 0; (y < TILE_COUNT_Y); y = as_type<int>((as_type<uint>(y) + as_type<uint>(1)))) {
-    TINT_ISOLATE_UB(tint_volatile_true_1) for(int x = 0; (x < TILE_COUNT_X); x = as_type<int>((as_type<uint>(x) + as_type<uint>(1)))) {
+  for(int y = 0; (y < TILE_COUNT_Y); y = as_type<int>((as_type<uint>(y) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
+    for(int x = 0; (x < TILE_COUNT_X); x = as_type<int>((as_type<uint>(x) + as_type<uint>(1)))) {
+      TINT_ISOLATE_UB(tint_volatile_false_1);
       int2 tilePixel0Idx = int2(as_type<int>((as_type<uint>(x) * as_type<uint>(TILE_SIZE))), as_type<int>((as_type<uint>(y) * as_type<uint>(TILE_SIZE))));
       float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / (*(tint_symbol_3)).fullScreenSize.xy) - float2(1.0f));
       float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(tilePixel0Idx) + as_type<uint2>(int2(TILE_SIZE)))))) / (*(tint_symbol_3)).fullScreenSize.xy) - float2(1.0f));
@@ -100,7 +101,8 @@
       frustumPlanes[2] = float4(0.0f, 1.0f, (-(viewFloorCoord[1]) / viewNear), 0.0f);
       frustumPlanes[3] = float4(0.0f, -1.0f, (viewCeilCoord[1] / viewNear), 0.0f);
       float dp = 0.0f;
-      TINT_ISOLATE_UB(tint_volatile_true_2) for(uint i = 0u; (i < 6u); i = (i + 1u)) {
+      for(uint i = 0u; (i < 6u); i = (i + 1u)) {
+        TINT_ISOLATE_UB(tint_volatile_false_2);
         float4 p = 0.0f;
         if ((frustumPlanes[i][0] > 0.0f)) {
           p[0] = boxMax[0];
diff --git a/test/tint/bug/tint/1321.wgsl.expected.msl b/test/tint/bug/tint/1321.wgsl.expected.msl
index 8dfe5a5..7e00141 100644
--- a/test/tint/bug/tint/1321.wgsl.expected.msl
+++ b/test/tint/bug/tint/1321.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 int foo() {
   return 1;
@@ -27,7 +26,8 @@
   {
     int const tint_symbol_1 = foo();
     int const a_save = tint_symbol_1;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false);
       {
         float const x = arr[a_save];
         break;
diff --git a/test/tint/bug/tint/1474-a.wgsl.expected.msl b/test/tint/bug/tint/1474-a.wgsl.expected.msl
index 4d8d438..fab12b4 100644
--- a/test/tint/bug/tint/1474-a.wgsl.expected.msl
+++ b/test/tint/bug/tint/1474-a.wgsl.expected.msl
@@ -3,11 +3,11 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol() {
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     if (true) {
       break;
     } else {
diff --git a/test/tint/bug/tint/1538.wgsl.expected.msl b/test/tint/bug/tint/1538.wgsl.expected.msl
index 87ec4d5..05c5d2e 100644
--- a/test/tint/bug/tint/1538.wgsl.expected.msl
+++ b/test/tint/bug/tint/1538.wgsl.expected.msl
@@ -3,8 +3,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 template<typename T, size_t N>
 struct tint_array {
@@ -23,7 +22,8 @@
 }
 
 int f() {
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     g();
     break;
   }
@@ -32,7 +32,8 @@
 }
 
 kernel void tint_symbol(device tint_array<uint, 1>* tint_symbol_1 [[buffer(0)]]) {
-  TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
     if (((*(tint_symbol_1))[0] == 0u)) {
       break;
     }
diff --git a/test/tint/bug/tint/1557.wgsl.expected.msl b/test/tint/bug/tint/1557.wgsl.expected.msl
index c4a412f..7bc61d7 100644
--- a/test/tint/bug/tint/1557.wgsl.expected.msl
+++ b/test/tint/bug/tint/1557.wgsl.expected.msl
@@ -3,8 +3,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 int f() {
   return 0;
@@ -12,7 +11,8 @@
 
 void g() {
   int j = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     if ((j >= 1)) {
       break;
     }
diff --git a/test/tint/bug/tint/1604.wgsl.expected.msl b/test/tint/bug/tint/1604.wgsl.expected.msl
index beb2f4d..50eb1f2 100644
--- a/test/tint/bug/tint/1604.wgsl.expected.msl
+++ b/test/tint/bug/tint/1604.wgsl.expected.msl
@@ -3,13 +3,13 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol(const constant int* tint_symbol_1 [[buffer(0)]]) {
   switch(*(tint_symbol_1)) {
     case 0: {
-      TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+      while(true) {
+        TINT_ISOLATE_UB(tint_volatile_false);
         return;
       }
       break;
diff --git a/test/tint/bug/tint/1605.wgsl.expected.msl b/test/tint/bug/tint/1605.wgsl.expected.msl
index 86a658c..65c5357 100644
--- a/test/tint/bug/tint/1605.wgsl.expected.msl
+++ b/test/tint/bug/tint/1605.wgsl.expected.msl
@@ -3,12 +3,13 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 bool func_3(const constant int* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < *(tint_symbol_1)); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    TINT_ISOLATE_UB(tint_volatile_true_1) for(int j = -1; (j == 1); j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)))) {
+  for(int i = 0; (i < *(tint_symbol_1)); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
+    for(int j = -1; (j == 1); j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)))) {
+      TINT_ISOLATE_UB(tint_volatile_false_1);
       return false;
     }
   }
diff --git a/test/tint/bug/tint/1764.wgsl.expected.msl b/test/tint/bug/tint/1764.wgsl.expected.msl
index 58d5684..0205a7e 100644
--- a/test/tint/bug/tint/1764.wgsl.expected.msl
+++ b/test/tint/bug/tint/1764.wgsl.expected.msl
@@ -15,11 +15,11 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<int, 246>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 246u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 246u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol_1))[i] = 0;
   }
diff --git a/test/tint/bug/tint/2010.spvasm.expected.msl b/test/tint/bug/tint/2010.spvasm.expected.msl
index d3c2676..34ac79e 100644
--- a/test/tint/bug/tint/2010.spvasm.expected.msl
+++ b/test/tint/bug/tint/2010.spvasm.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   uint3 x_3;
@@ -34,7 +33,8 @@
     atomic_store_explicit(tint_symbol_5, 0u, memory_order_relaxed);
     atomic_store_explicit(tint_symbol_6, 0u, memory_order_relaxed);
   }
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4096u); idx = (idx + 32u)) {
+  for(uint idx = local_idx; (idx < 4096u); idx = (idx + 32u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     S const tint_symbol_1 = S{};
     (*(tint_symbol_7))[i] = tint_symbol_1;
@@ -65,7 +65,8 @@
   uint x_88 = 0u;
   uint const x_52 = (*(tint_private_vars)).x_3[0];
   x_54 = 0u;
-  TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
     uint x_55 = 0u;
     x_58 = (*(tint_symbol_8)).field0.field0;
     if ((x_54 < x_58)) {
@@ -97,7 +98,8 @@
   }
   x_85 = x_76.xyxy;
   x_88 = 1u;
-  TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false_2);
     float4 x_111 = 0.0f;
     float4 x_86 = 0.0f;
     uint x_89 = 0u;
diff --git a/test/tint/bug/tint/2039.wgsl.expected.msl b/test/tint/bug/tint/2039.wgsl.expected.msl
index c756acb..28d2df8 100644
--- a/test/tint/bug/tint/2039.wgsl.expected.msl
+++ b/test/tint/bug/tint/2039.wgsl.expected.msl
@@ -3,13 +3,13 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol() {
   uint out = 0u;
   bool tint_continue = false;
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     tint_continue = false;
     switch(2) {
       case 1: {
diff --git a/test/tint/bug/tint/2059.wgsl.expected.msl b/test/tint/bug/tint/2059.wgsl.expected.msl
index 5588ec6..bd2706e 100644
--- a/test/tint/bug/tint/2059.wgsl.expected.msl
+++ b/test/tint/bug/tint/2059.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
@@ -66,7 +65,8 @@
 }
 
 void assign_and_preserve_padding_3(device tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 1>* const dest, tint_array<float3x3, 1> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 1u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 1u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding(&((*(dest))[i]), value[i]);
   }
 }
@@ -80,7 +80,8 @@
 }
 
 void assign_and_preserve_padding_6(device tint_array<S_tint_packed_vec3, 1>* const dest, tint_array<S, 1> value) {
-  TINT_ISOLATE_UB(tint_volatile_true_1) for(uint i = 0u; (i < 1u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 1u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
@@ -90,14 +91,16 @@
 }
 
 void assign_and_preserve_padding_7(device tint_array<S2_tint_packed_vec3, 1>* const dest, tint_array<S2, 1> value) {
-  TINT_ISOLATE_UB(tint_volatile_true_2) for(uint i = 0u; (i < 1u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 1u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false_2);
     assign_and_preserve_padding_2(&((*(dest))[i]), value[i]);
   }
 }
 
 kernel void tint_symbol(device tint_array<tint_packed_vec3_f32_array_element, 3>* tint_symbol_8 [[buffer(0)]], device S_tint_packed_vec3* tint_symbol_9 [[buffer(1)]], device S2_tint_packed_vec3* tint_symbol_10 [[buffer(2)]], device S3_tint_packed_vec3* tint_symbol_11 [[buffer(3)]], device S4_tint_packed_vec3* tint_symbol_12 [[buffer(4)]], device tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 1>* tint_symbol_13 [[buffer(5)]], device tint_array<S_tint_packed_vec3, 1>* tint_symbol_14 [[buffer(6)]], device tint_array<S2_tint_packed_vec3, 1>* tint_symbol_15 [[buffer(7)]]) {
   float3x3 m = float3x3(0.0f);
-  TINT_ISOLATE_UB(tint_volatile_true_3) for(uint c = 0u; (c < 3u); c = (c + 1u)) {
+  for(uint c = 0u; (c < 3u); c = (c + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false_3);
     m[c] = float3(float(((c * 3u) + 1u)), float(((c * 3u) + 2u)), float(((c * 3u) + 3u)));
   }
   {
diff --git a/test/tint/bug/tint/2201.wgsl.expected.msl b/test/tint/bug/tint/2201.wgsl.expected.msl
index 5852fb6..2aba54c 100644
--- a/test/tint/bug/tint/2201.wgsl.expected.msl
+++ b/test/tint/bug/tint/2201.wgsl.expected.msl
@@ -7,11 +7,11 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol() {
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     if (true) {
       break;
     } else {
diff --git a/test/tint/bug/tint/2202.wgsl.expected.msl b/test/tint/bug/tint/2202.wgsl.expected.msl
index 2d601f3..767e10c 100644
--- a/test/tint/bug/tint/2202.wgsl.expected.msl
+++ b/test/tint/bug/tint/2202.wgsl.expected.msl
@@ -7,12 +7,13 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol() {
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1);
       return;
     }
     bool const _e9 = true;
diff --git a/test/tint/bug/tint/221.wgsl.expected.msl b/test/tint/bug/tint/221.wgsl.expected.msl
index c418ea4..5d5f951 100644
--- a/test/tint/bug/tint/221.wgsl.expected.msl
+++ b/test/tint/bug/tint/221.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct Buf {
   /* 0x0000 */ uint count;
@@ -29,7 +28,8 @@
 
 kernel void tint_symbol(device Buf* tint_symbol_1 [[buffer(0)]]) {
   uint i = 0u;
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     if ((i >= (*(tint_symbol_1)).count)) {
       break;
     }
diff --git a/test/tint/bug/tint/349291130.wgsl.expected.msl b/test/tint/bug/tint/349291130.wgsl.expected.msl
index b3390c3..83fff02 100644
--- a/test/tint/bug/tint/349291130.wgsl.expected.msl
+++ b/test/tint/bug/tint/349291130.wgsl.expected.msl
@@ -3,12 +3,12 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void e(texture2d<float, access::sample> tint_symbol_1 [[texture(0)]]) {
   {
-    TINT_ISOLATE_UB(tint_volatile_true) for(uint level = tint_symbol_1.get_num_mip_levels(); (level > 0u); ) {
+    for(uint level = tint_symbol_1.get_num_mip_levels(); (level > 0u); ) {
+      TINT_ISOLATE_UB(tint_volatile_false);
     }
   }
   return;
diff --git a/test/tint/bug/tint/354627692.wgsl.expected.msl b/test/tint/bug/tint/354627692.wgsl.expected.msl
index 977f57a..87e2356 100644
--- a/test/tint/bug/tint/354627692.wgsl.expected.msl
+++ b/test/tint/bug/tint/354627692.wgsl.expected.msl
@@ -3,14 +3,15 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol_1(device int* tint_symbol_2 [[buffer(0)]]) {
   int i = *(tint_symbol_2);
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     {
-      TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+      while(true) {
+        TINT_ISOLATE_UB(tint_volatile_false_1);
         if ((i > 5)) {
           i = as_type<int>((as_type<uint>(i) * as_type<uint>(2)));
           break;
diff --git a/test/tint/bug/tint/366037039.wgsl.expected.msl b/test/tint/bug/tint/366037039.wgsl.expected.msl
index bf0f065..d819678 100644
--- a/test/tint/bug/tint/366037039.wgsl.expected.msl
+++ b/test/tint/bug/tint/366037039.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_packed_vec3_u32_array_element {
   /* 0x0000 */ packed_uint3 elements;
@@ -49,7 +48,8 @@
 }
 
 void assign_and_preserve_padding_1(device tint_array<tint_packed_vec3_u32_array_element, 4>* const dest, tint_array<uint3, 4> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     (*(dest))[i].elements = packed_uint3(value[i]);
   }
 }
diff --git a/test/tint/bug/tint/534.wgsl.expected.msl b/test/tint/bug/tint/534.wgsl.expected.msl
index 3fe0797..415cbb0 100644
--- a/test/tint/bug/tint/534.wgsl.expected.msl
+++ b/test/tint/bug/tint/534.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 uint4 tint_ftou(float4 v) {
   return select(uint4(4294967295u), select(uint4(v), uint4(0u), (v < float4(0.0f))), (v <= float4(4294967040.0f)));
@@ -49,7 +48,8 @@
   bool success = true;
   uint4 srcColorBits = 0u;
   uint4 dstColorBits = tint_ftou(dstColor);
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < (*(tint_symbol_3)).channelCount); i = (i + 1u)) {
+  for(uint i = 0u; (i < (*(tint_symbol_3)).channelCount); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const tint_symbol_1 = i;
     srcColorBits[tint_symbol_1] = ConvertToFp16FloatValue(srcColor[i]);
     success = (success && (srcColorBits[i] == dstColorBits[i]));
diff --git a/test/tint/bug/tint/744.wgsl.expected.msl b/test/tint/bug/tint/744.wgsl.expected.msl
index ccb621f..295e92f 100644
--- a/test/tint/bug/tint/744.wgsl.expected.msl
+++ b/test/tint/bug/tint/744.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct Uniforms {
   /* 0x0000 */ uint2 aShape;
@@ -33,7 +32,8 @@
   uint const dimInner = (*(tint_symbol_1)).aShape[1];
   uint const dimOutter = (*(tint_symbol_1)).outShape[1];
   uint result = 0u;
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < dimInner); i = (i + 1u)) {
+  for(uint i = 0u; (i < dimInner); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const a = (i + (resultCell[0] * dimInner));
     uint const b = (resultCell[1] + (i * dimOutter));
     result = (result + ((*(tint_symbol_2)).numbers[a] * (*(tint_symbol_3)).numbers[b]));
diff --git a/test/tint/bug/tint/757.wgsl.expected.msl b/test/tint/bug/tint/757.wgsl.expected.msl
index b59e7b3..ab84934 100644
--- a/test/tint/bug/tint/757.wgsl.expected.msl
+++ b/test/tint/bug/tint/757.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct Constants {
   int level;
@@ -30,7 +29,8 @@
   uint flatIndex = (((4u * GlobalInvocationID[2]) + (2u * GlobalInvocationID[1])) + GlobalInvocationID[0]);
   flatIndex = (flatIndex * 1u);
   float4 texel = tint_symbol_1.read(uint2(int2(GlobalInvocationID.xy)), 0, 0);
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 1u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 1u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     (*(tint_symbol_2)).values[(flatIndex + i)] = texel[0];
   }
 }
diff --git a/test/tint/bug/tint/914.wgsl.expected.msl b/test/tint/bug/tint/914.wgsl.expected.msl
index 6f1bb66..1eaf522 100644
--- a/test/tint/bug/tint/914.wgsl.expected.msl
+++ b/test/tint/bug/tint/914.wgsl.expected.msl
@@ -15,11 +15,11 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_5, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_6) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4096u); idx = (idx + 256u)) {
+  for(uint idx = local_idx; (idx < 4096u); idx = (idx + 256u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = (idx / 64u);
     uint const i_1 = (idx % 64u);
     (*(tint_symbol_5))[i][i_1] = 0.0f;
@@ -75,16 +75,20 @@
   tint_array<float, 16> acc = {};
   float ACached = 0.0f;
   tint_array<float, 4> BCached = {};
-  TINT_ISOLATE_UB(tint_volatile_true_1) for(uint index = 0u; (index < 16u); index = (index + 1u)) {
+  for(uint index = 0u; (index < 16u); index = (index + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
     acc[index] = 0.0f;
   }
   uint const ColPerThreadA = 4u;
   uint const tileColA = (local_id[0] * ColPerThreadA);
   uint const RowPerThreadB = 4u;
   uint const tileRowB = (local_id[1] * RowPerThreadB);
-  TINT_ISOLATE_UB(tint_volatile_true_2) for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
-    TINT_ISOLATE_UB(tint_volatile_true_3) for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
-      TINT_ISOLATE_UB(tint_volatile_true_4) for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
+  for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false_2);
+    for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
+      TINT_ISOLATE_UB(tint_volatile_false_3);
+      for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
+        TINT_ISOLATE_UB(tint_volatile_false_4);
         uint const inputRow = (tileRow + innerRow);
         uint const inputCol = (tileColA + innerCol);
         uint const tint_symbol_1 = inputRow;
@@ -92,8 +96,10 @@
         (*(tint_symbol_13))[tint_symbol_1][tint_symbol_2] = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol), tint_symbol_15, tint_symbol_16);
       }
     }
-    TINT_ISOLATE_UB(tint_volatile_true_5) for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
-      TINT_ISOLATE_UB(tint_volatile_true_6) for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
+    for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
+      TINT_ISOLATE_UB(tint_volatile_false_5);
+      for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
+        TINT_ISOLATE_UB(tint_volatile_false_6);
         uint const inputRow = (tileRowB + innerRow);
         uint const inputCol = (tileCol + innerCol);
         uint const tint_symbol_3 = innerCol;
@@ -102,13 +108,17 @@
       }
     }
     threadgroup_barrier(mem_flags::mem_threadgroup);
-    TINT_ISOLATE_UB(tint_volatile_true_7) for(uint k = 0u; (k < 64u); k = (k + 1u)) {
-      TINT_ISOLATE_UB(tint_volatile_true_8) for(uint inner = 0u; (inner < 4u); inner = (inner + 1u)) {
+    for(uint k = 0u; (k < 64u); k = (k + 1u)) {
+      TINT_ISOLATE_UB(tint_volatile_false_7);
+      for(uint inner = 0u; (inner < 4u); inner = (inner + 1u)) {
+        TINT_ISOLATE_UB(tint_volatile_false_8);
         BCached[inner] = (*(tint_symbol_14))[k][(tileCol + inner)];
       }
-      TINT_ISOLATE_UB(tint_volatile_true_9) for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
+      for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
+        TINT_ISOLATE_UB(tint_volatile_false_9);
         ACached = (*(tint_symbol_13))[(tileRow + innerRow)][k];
-        TINT_ISOLATE_UB(tint_volatile_true_10) for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
+        for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
+          TINT_ISOLATE_UB(tint_volatile_false_10);
           uint const index = ((innerRow * 4u) + innerCol);
           acc[index] = (acc[index] + (ACached * BCached[innerCol]));
         }
@@ -116,8 +126,10 @@
     }
     threadgroup_barrier(mem_flags::mem_threadgroup);
   }
-  TINT_ISOLATE_UB(tint_volatile_true_11) for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
-    TINT_ISOLATE_UB(tint_volatile_true_12) for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
+  for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false_11);
+    for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
+      TINT_ISOLATE_UB(tint_volatile_false_12);
       uint const index = ((innerRow * 4u) + innerCol);
       mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index], tint_symbol_15, tint_symbol_18);
     }
diff --git a/test/tint/bug/tint/942.wgsl.expected.msl b/test/tint/bug/tint/942.wgsl.expected.msl
index fa62439..fb02b8a 100644
--- a/test/tint/bug/tint/942.wgsl.expected.msl
+++ b/test/tint/bug/tint/942.wgsl.expected.msl
@@ -15,11 +15,11 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<tint_array<float3, 256>, 4>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 1024u); idx = (idx + 64u)) {
+  for(uint idx = local_idx; (idx < 1024u); idx = (idx + 64u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i_1 = (idx / 256u);
     uint const i_2 = (idx % 256u);
     (*(tint_symbol_1))[i_1][i_2] = float3(0.0f);
@@ -45,8 +45,10 @@
   uint const filterOffset = tint_div(((*(tint_symbol_3)).filterDim - 1u), 2u);
   uint2 const dims = uint2(tint_symbol_4.get_width(0), tint_symbol_4.get_height(0));
   uint2 const baseIndex = (((WorkGroupID.xy * uint2((*(tint_symbol_3)).blockDim, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u))) - uint2(filterOffset, 0u));
-  TINT_ISOLATE_UB(tint_volatile_true_1) for(uint r = 0u; (r < 4u); r = (r + 1u)) {
-    TINT_ISOLATE_UB(tint_volatile_true_2) for(uint c = 0u; (c < 4u); c = (c + 1u)) {
+  for(uint r = 0u; (r < 4u); r = (r + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
+    for(uint c = 0u; (c < 4u); c = (c + 1u)) {
+      TINT_ISOLATE_UB(tint_volatile_false_2);
       uint2 loadIndex = (baseIndex + uint2(c, r));
       if (((*(tint_symbol_5)).value != 0u)) {
         loadIndex = loadIndex.yx;
@@ -55,8 +57,10 @@
     }
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  TINT_ISOLATE_UB(tint_volatile_true_3) for(uint r = 0u; (r < 4u); r = (r + 1u)) {
-    TINT_ISOLATE_UB(tint_volatile_true_4) for(uint c = 0u; (c < 4u); c = (c + 1u)) {
+  for(uint r = 0u; (r < 4u); r = (r + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false_3);
+    for(uint c = 0u; (c < 4u); c = (c + 1u)) {
+      TINT_ISOLATE_UB(tint_volatile_false_4);
       uint2 writeIndex = (baseIndex + uint2(c, r));
       if (((*(tint_symbol_5)).value != 0u)) {
         writeIndex = writeIndex.yx;
@@ -64,7 +68,8 @@
       uint const center = ((4u * LocalInvocationID[0]) + c);
       if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
         float3 acc = float3(0.0f);
-        TINT_ISOLATE_UB(tint_volatile_true_5) for(uint f = 0u; (f < (*(tint_symbol_3)).filterDim); f = (f + 1u)) {
+        for(uint f = 0u; (f < (*(tint_symbol_3)).filterDim); f = (f + 1u)) {
+          TINT_ISOLATE_UB(tint_volatile_false_5);
           uint i = ((center + f) - filterOffset);
           acc = (acc + ((1.0f / float((*(tint_symbol_3)).filterDim)) * (*(tint_symbol_2))[r][i]));
         }
diff --git a/test/tint/bug/tint/948.wgsl.expected.msl b/test/tint/bug/tint/948.wgsl.expected.msl
index cc95fb8..f4fafbe 100644
--- a/test/tint/bug/tint/948.wgsl.expected.msl
+++ b/test/tint/bug/tint/948.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   float2 tUV;
@@ -102,7 +101,8 @@
   float2 const x_111 = (*(tint_symbol_8)).stageSize;
   stageUnits = (float2(1.0f) / x_111);
   i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     int const x_122 = i;
     if ((x_122 < 2)) {
     } else {
@@ -138,7 +138,8 @@
       float const x_184 = animationData[2];
       (*(tint_private_vars)).mt = fmod((x_181 * x_184), 1.0f);
       f = 0.0f;
-      TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+      while(true) {
+        TINT_ISOLATE_UB(tint_volatile_false_1);
         float const x_193 = f;
         if ((x_193 < 8.0f)) {
         } else {
diff --git a/test/tint/bug/tint/949.wgsl.expected.msl b/test/tint/bug/tint/949.wgsl.expected.msl
index 2ed231b..08f3822 100644
--- a/test/tint/bug/tint/949.wgsl.expected.msl
+++ b/test/tint/bug/tint/949.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   float u_Float;
@@ -334,7 +333,8 @@
   lastSampledHeight = 1.0f;
   currSampledHeight = 1.0f;
   i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     int const x_388 = i;
     if ((x_388 < 15)) {
     } else {
diff --git a/test/tint/bug/tint/990.wgsl.expected.msl b/test/tint/bug/tint/990.wgsl.expected.msl
index 09a7dbc..e63c043 100644
--- a/test/tint/bug/tint/990.wgsl.expected.msl
+++ b/test/tint/bug/tint/990.wgsl.expected.msl
@@ -3,12 +3,12 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void f() {
   int i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) for(; false; ) {
+  for(; false; ) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
index c6c2363..f10bcec 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   uint local_invocation_index_1;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx_1 = local_idx; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
+  for(uint idx_1 = local_idx; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = (idx_1 / 2u);
     uint const i_1 = (idx_1 % 2u);
     uint const i_2 = (idx_1 % 1u);
@@ -43,7 +43,8 @@
 void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_1) {
   uint idx = 0u;
   idx = local_invocation_index_2;
-  TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
     if (!((idx < 6u))) {
       break;
     }
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl
index cd424bf..cdc5187 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl
@@ -15,11 +15,11 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 6u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 6u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = (idx / 2u);
     uint const i_1 = (idx % 2u);
     uint const i_2 = (idx % 1u);
diff --git a/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl b/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
index b2c82f3..f2a057b 100644
--- a/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   uint local_invocation_index_1;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<atomic_uint, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx_1 = local_idx; (idx_1 < 4u); idx_1 = (idx_1 + 1u)) {
+  for(uint idx_1 = local_idx; (idx_1 < 4u); idx_1 = (idx_1 + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx_1;
     atomic_store_explicit(&((*(tint_symbol))[i]), 0u, memory_order_relaxed);
   }
@@ -33,7 +33,8 @@
 void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<atomic_uint, 4>* const tint_symbol_1) {
   uint idx = 0u;
   idx = local_invocation_index_2;
-  TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
     if (!((idx < 4u))) {
       break;
     }
diff --git a/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl b/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl
index edf13ac..61cb568 100644
--- a/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl
@@ -15,11 +15,11 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<atomic_uint, 4>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     atomic_store_explicit(&((*(tint_symbol))[i]), 0u, memory_order_relaxed);
   }
diff --git a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
index c6c2363..f10bcec 100644
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   uint local_invocation_index_1;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx_1 = local_idx; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
+  for(uint idx_1 = local_idx; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = (idx_1 / 2u);
     uint const i_1 = (idx_1 % 2u);
     uint const i_2 = (idx_1 % 1u);
@@ -43,7 +43,8 @@
 void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_1) {
   uint idx = 0u;
   idx = local_invocation_index_2;
-  TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
     if (!((idx < 6u))) {
       break;
     }
diff --git a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl
index cd424bf..cdc5187 100644
--- a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl
@@ -15,11 +15,11 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 6u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 6u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = (idx / 2u);
     uint const i_1 = (idx % 2u);
     uint const i_2 = (idx % 1u);
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
index cff2fea..c87552c 100644
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   uint local_invocation_index_1;
@@ -29,7 +28,8 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S_atomic, 10>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx_1 = local_idx; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
+  for(uint idx_1 = local_idx; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx_1;
     (*(tint_symbol))[i].x = 0;
     atomic_store_explicit(&((*(tint_symbol))[i].a), 0u, memory_order_relaxed);
@@ -47,7 +47,8 @@
 void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<S_atomic, 10>* const tint_symbol_1) {
   uint idx = 0u;
   idx = local_invocation_index_2;
-  TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
     if (!((idx < 10u))) {
       break;
     }
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl
index b90f6d6..4dfe73a 100644
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   int x;
@@ -25,7 +24,8 @@
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<S, 10>* const tint_symbol) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 10u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 10u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol))[i].x = 0;
     atomic_store_explicit(&((*(tint_symbol))[i].a), 0u, memory_order_relaxed);
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
index d5e8aca..ccce777 100644
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   uint local_invocation_index_1;
@@ -33,7 +32,8 @@
     (*(tint_symbol)).x = 0;
     (*(tint_symbol)).y = 0u;
   }
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx_1 = local_idx; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
+  for(uint idx_1 = local_idx; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx_1;
     atomic_store_explicit(&((*(tint_symbol)).a[i]), 0u, memory_order_relaxed);
   }
@@ -51,7 +51,8 @@
   (*(tint_symbol_1)).x = 0;
   (*(tint_symbol_1)).y = 0u;
   idx = local_invocation_index_2;
-  TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false_1);
     if (!((idx < 10u))) {
       break;
     }
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl
index a8acff4..f11766f 100644
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   int x;
@@ -29,7 +28,8 @@
     (*(tint_symbol)).x = 0;
     (*(tint_symbol)).y = 0u;
   }
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 10u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 10u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     atomic_store_explicit(&((*(tint_symbol)).a[i]), 0u, memory_order_relaxed);
   }
diff --git a/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.msl b/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.msl
index 9d7fb7d..9fe0a5f 100644
--- a/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.msl
+++ b/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.msl
@@ -3,13 +3,13 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void foo(texture2d<int, access::read_write> tint_symbol) {
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false);
       if (!((i < 3))) {
         break;
       }
diff --git a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.msl b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.msl
index 5a52532..3ae8c42 100644
--- a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.msl
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.msl
@@ -3,8 +3,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 int tint_workgroupUniformLoad(threadgroup int* const p) {
   threadgroup_barrier(mem_flags::mem_threadgroup);
@@ -16,7 +15,8 @@
 void foo(threadgroup int* const tint_symbol_4, threadgroup int* const tint_symbol_5) {
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false);
       int const tint_symbol = i;
       int const tint_symbol_1 = tint_workgroupUniformLoad(tint_symbol_4);
       if (!((tint_symbol < tint_symbol_1))) {
diff --git a/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.msl
index 6354472..7c3a9fa 100644
--- a/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.msl
@@ -15,8 +15,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_4 {
   float x [[user(locn0)]];
@@ -24,7 +23,8 @@
 
 void tint_symbol_inner(float x) {
   float4 v = float4(0.0f);
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     bool tint_symbol_1 = (x > v[0]);
     if (tint_symbol_1) {
       float const tint_symbol_2 = dfdx(1.0f);
diff --git a/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.msl
index 68189b2..119f37e 100644
--- a/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.msl
@@ -15,8 +15,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_2 {
   float x [[user(locn0)]];
@@ -24,7 +23,8 @@
 
 void tint_symbol_inner(float x, texture2d<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
   float4 v = float4(0.0f);
-  TINT_ISOLATE_UB(tint_volatile_true) for(; (x > v[0]); ) {
+  for(; (x > v[0]); ) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     v = tint_symbol_3.sample(tint_symbol_4, float2(0.0f));
   }
 }
diff --git a/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.msl
index 98e918f..6cf6c95 100644
--- a/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.msl
@@ -15,15 +15,15 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_2 {
   float x [[user(locn0)]];
 };
 
 void tint_symbol_inner(float x) {
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     float const tint_phony = dfdx(1.0f);
     {
       if ((x > 0.0f)) { break; }
diff --git a/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.msl
index dde39cd..cf466f9 100644
--- a/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.msl
@@ -15,15 +15,15 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_2 {
   float x [[user(locn0)]];
 };
 
 void tint_symbol_inner(float x) {
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     float const tint_phony = dfdx(1.0f);
     {
       if ((x > 0.0f)) { break; }
diff --git a/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.msl
index 3a6f0c6..06c50d7 100644
--- a/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.msl
@@ -15,15 +15,15 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_2 {
   float x [[user(locn0)]];
 };
 
 void tint_symbol_inner(float x) {
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     {
       float const tint_phony = dfdx(1.0f);
       if ((x > 0.0f)) { break; }
diff --git a/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.msl
index 67d0ba6..96b069d 100644
--- a/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.msl
@@ -15,8 +15,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_4 {
   float x [[user(locn0)]];
@@ -24,7 +23,8 @@
 
 void tint_symbol_inner(float x) {
   float4 v = float4(0.0f);
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     bool tint_symbol_1 = (x > 0.0f);
     if (tint_symbol_1) {
       float const tint_symbol_2 = dfdx(1.0f);
diff --git a/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.msl
index b2e7ce2..9551cf0 100644
--- a/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.msl
@@ -15,8 +15,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_2 {
   float x [[user(locn0)]];
@@ -24,7 +23,8 @@
 
 void tint_symbol_inner(float x, texture2d<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
   float4 v = float4(0.0f);
-  TINT_ISOLATE_UB(tint_volatile_true) while((x > v[0])) {
+  while((x > v[0])) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     v = tint_symbol_3.sample(tint_symbol_4, float2(0.0f));
   }
 }
diff --git a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.msl b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.msl
index 74e58e2..ea779b1 100644
--- a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.msl
+++ b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct strided_arr {
   /* 0x0000 */ float2 el;
@@ -43,7 +42,8 @@
 }
 
 void assign_and_preserve_padding(device tint_array<strided_arr, 2>* const dest, tint_array<strided_arr, 2> value) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 2u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 2u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/loops/continue_in_switch.wgsl.expected.msl b/test/tint/loops/continue_in_switch.wgsl.expected.msl
index e7349b0c..93cfb99 100644
--- a/test/tint/loops/continue_in_switch.wgsl.expected.msl
+++ b/test/tint/loops/continue_in_switch.wgsl.expected.msl
@@ -3,12 +3,12 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void f() {
   bool tint_continue = false;
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     tint_continue = false;
     switch(i) {
       case 0: {
diff --git a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.msl b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.msl
index e0ef3e1..e944287 100644
--- a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.msl
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.msl
@@ -3,13 +3,13 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void f() {
   int i = 0;
   bool tint_continue = false;
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     tint_continue = false;
     switch(i) {
       case 0: {
diff --git a/test/tint/loops/loop.wgsl.expected.msl b/test/tint/loops/loop.wgsl.expected.msl
index 3ab3d8e..e641dc5 100644
--- a/test/tint/loops/loop.wgsl.expected.msl
+++ b/test/tint/loops/loop.wgsl.expected.msl
@@ -3,12 +3,12 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 int f() {
   int i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
     if ((i > 4)) {
       return i;
diff --git a/test/tint/loops/loop_with_break_if.wgsl.expected.msl b/test/tint/loops/loop_with_break_if.wgsl.expected.msl
index 08f6c09..e97f8c6 100644
--- a/test/tint/loops/loop_with_break_if.wgsl.expected.msl
+++ b/test/tint/loops/loop_with_break_if.wgsl.expected.msl
@@ -3,12 +3,12 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 int f() {
   int i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     if ((i > 4)) {
       return i;
     }
diff --git a/test/tint/loops/loop_with_continuing.wgsl.expected.msl b/test/tint/loops/loop_with_continuing.wgsl.expected.msl
index 2f3b467..78da287 100644
--- a/test/tint/loops/loop_with_continuing.wgsl.expected.msl
+++ b/test/tint/loops/loop_with_continuing.wgsl.expected.msl
@@ -3,12 +3,12 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 int f() {
   int i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     if ((i > 4)) {
       return i;
     }
diff --git a/test/tint/loops/multiple_continues.wgsl.expected.msl b/test/tint/loops/multiple_continues.wgsl.expected.msl
index 76e84d7..73cc270 100644
--- a/test/tint/loops/multiple_continues.wgsl.expected.msl
+++ b/test/tint/loops/multiple_continues.wgsl.expected.msl
@@ -3,12 +3,12 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol() {
   bool tint_continue = false;
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     tint_continue = false;
     switch(i) {
       case 0: {
diff --git a/test/tint/loops/multiple_switch.wgsl.expected.msl b/test/tint/loops/multiple_switch.wgsl.expected.msl
index 7728952..8dd8fe4 100644
--- a/test/tint/loops/multiple_switch.wgsl.expected.msl
+++ b/test/tint/loops/multiple_switch.wgsl.expected.msl
@@ -3,13 +3,13 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol() {
   int i = 0;
   bool tint_continue = false;
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i_1 = 0; (i_1 < 2); i_1 = as_type<int>((as_type<uint>(i_1) + as_type<uint>(1)))) {
+  for(int i_1 = 0; (i_1 < 2); i_1 = as_type<int>((as_type<uint>(i_1) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     tint_continue = false;
     switch(i_1) {
       case 0: {
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
index eb32ea2..5b0f057 100644
--- a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
@@ -3,13 +3,14 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol() {
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     bool tint_continue = false;
-    TINT_ISOLATE_UB(tint_volatile_true_1) for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+    for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+      TINT_ISOLATE_UB(tint_volatile_false_1);
       tint_continue = false;
       switch(i) {
         case 0: {
diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
index 9ac0463..0455c39 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
@@ -3,17 +3,18 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol() {
   bool tint_continue_1 = false;
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     tint_continue_1 = false;
     switch(i) {
       case 0: {
         bool tint_continue = false;
-        TINT_ISOLATE_UB(tint_volatile_true_1) for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+        for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+          TINT_ISOLATE_UB(tint_volatile_false_1);
           tint_continue = false;
           switch(j) {
             case 0: {
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl
index c3b1e93..3cfa334 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl
@@ -3,18 +3,19 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol() {
   int k = 0;
   bool tint_continue_1 = false;
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     tint_continue_1 = false;
     switch(i) {
       case 0: {
         bool tint_continue = false;
-        TINT_ISOLATE_UB(tint_volatile_true_1) for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+        for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+          TINT_ISOLATE_UB(tint_volatile_false_1);
           tint_continue = false;
           switch(j) {
             case 0: {
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
index e122be4..8564b09 100644
--- a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
@@ -3,13 +3,13 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol() {
   int j = 0;
   bool tint_continue = false;
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     tint_continue = false;
     switch(i) {
       case 0: {
diff --git a/test/tint/loops/nested_loops.wgsl.expected.msl b/test/tint/loops/nested_loops.wgsl.expected.msl
index 8f176cb..d3f2bcf 100644
--- a/test/tint/loops/nested_loops.wgsl.expected.msl
+++ b/test/tint/loops/nested_loops.wgsl.expected.msl
@@ -3,18 +3,19 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 int f() {
   int i = 0;
   int j = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
     if ((i > 4)) {
       return 1;
     }
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1);
       j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
       if ((j > 4)) {
         return 2;
diff --git a/test/tint/loops/nested_loops_with_continuing.wgsl.expected.msl b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.msl
index 25ea8bc..d0155f3 100644
--- a/test/tint/loops/nested_loops_with_continuing.wgsl.expected.msl
+++ b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.msl
@@ -3,17 +3,18 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 int f() {
   int i = 0;
   int j = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     if ((i > 4)) {
       return 1;
     }
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1);
       if ((j > 4)) {
         return 2;
       }
diff --git a/test/tint/loops/single_continue.wgsl.expected.msl b/test/tint/loops/single_continue.wgsl.expected.msl
index 113c92c..d095344 100644
--- a/test/tint/loops/single_continue.wgsl.expected.msl
+++ b/test/tint/loops/single_continue.wgsl.expected.msl
@@ -3,12 +3,12 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void tint_symbol() {
   bool tint_continue = false;
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     tint_continue = false;
     switch(i) {
       case 0: {
diff --git a/test/tint/loops/while.wgsl.expected.msl b/test/tint/loops/while.wgsl.expected.msl
index d044fb0..1d8717d 100644
--- a/test/tint/loops/while.wgsl.expected.msl
+++ b/test/tint/loops/while.wgsl.expected.msl
@@ -3,12 +3,12 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 int f() {
   int i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) while((i < 4)) {
+  while((i < 4)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
   }
   return i;
diff --git a/test/tint/loops/while_with_continue.wgsl.expected.msl b/test/tint/loops/while_with_continue.wgsl.expected.msl
index 35233d4..6827095 100644
--- a/test/tint/loops/while_with_continue.wgsl.expected.msl
+++ b/test/tint/loops/while_with_continue.wgsl.expected.msl
@@ -3,12 +3,12 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 int f() {
   int i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) while((i < 4)) {
+  while((i < 4)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
     continue;
   }
diff --git a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.msl b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.msl
index ce4c545..479e0f5 100644
--- a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct str {
   tint_array<int, 4> arr;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup str* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol_1)).arr[i] = 0;
   }
diff --git a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.msl b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.msl
index bdaac08..c8a00fd 100644
--- a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct str {
   int i;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<str, 4>* const tint_symbol_2) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i_1 = idx;
     str const tint_symbol_1 = str{};
     (*(tint_symbol_2))[i_1] = tint_symbol_1;
diff --git a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.msl b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.msl
index ce43c2a..669dc4c 100644
--- a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct str {
   tint_array<int, 4> arr;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup str* const tint_symbol_2) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol_2)).arr[i] = 0;
   }
diff --git a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.msl b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.msl
index efaed5d..5d52a1f 100644
--- a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.msl
@@ -15,15 +15,15 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct str {
   int i;
 };
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<str, 4>* const tint_symbol_3) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 4u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i_1 = idx;
     str const tint_symbol_1 = str{};
     (*(tint_symbol_3))[i_1] = tint_symbol_1;
diff --git a/test/tint/samples/compute_boids.wgsl.expected.msl b/test/tint/samples/compute_boids.wgsl.expected.msl
index e64ec48..9a51ae7 100644
--- a/test/tint/samples/compute_boids.wgsl.expected.msl
+++ b/test/tint/samples/compute_boids.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_symbol_1 {
   float2 a_particlePos [[attribute(0)]];
@@ -89,7 +88,8 @@
   int cVelCount = 0;
   float2 pos = 0.0f;
   float2 vel = 0.0f;
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint i = 0u; (i < 5u); i = (i + 1u)) {
+  for(uint i = 0u; (i < 5u); i = (i + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     if ((i == index)) {
       continue;
     }
diff --git a/test/tint/shadowing/loop.wgsl.expected.msl b/test/tint/shadowing/loop.wgsl.expected.msl
index 5a225bd..cec9976 100644
--- a/test/tint/shadowing/loop.wgsl.expected.msl
+++ b/test/tint/shadowing/loop.wgsl.expected.msl
@@ -15,12 +15,12 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 kernel void foo(device tint_array<int, 10>* tint_symbol [[buffer(0)]]) {
   int i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     int x = (*(tint_symbol))[i];
     {
       int x_1 = (*(tint_symbol))[x];
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl
index 091fcb0..a554a58 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct Uniforms {
   /* 0x0000 */ uint i;
@@ -33,7 +32,8 @@
 kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
   InnerS v = {};
   OuterS s1 = {};
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     s1.a1[(*(tint_symbol_1)).i] = v;
   }
   return;
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl
index 884c030..7d5ad6c 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct Uniforms {
   /* 0x0000 */ uint i;
@@ -33,7 +32,8 @@
 kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
   InnerS v = {};
   OuterS s1 = {};
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 4); s1.a1[(*(tint_symbol_1)).i] = v) {
+  for(int i = 0; (i < 4); s1.a1[(*(tint_symbol_1)).i] = v) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
   }
   return;
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl
index 4431501..454c8e7 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct Uniforms {
   /* 0x0000 */ uint i;
@@ -34,7 +33,8 @@
   InnerS v = {};
   OuterS s1 = {};
   int i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) for(s1.a1[(*(tint_symbol_1)).i] = v; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(s1.a1[(*(tint_symbol_1)).i] = v; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
   return;
 }
diff --git a/test/tint/statements/compound_assign/for_loop.wgsl.expected.msl b/test/tint/statements/compound_assign/for_loop.wgsl.expected.msl
index 190496d..b19b1d6 100644
--- a/test/tint/statements/compound_assign/for_loop.wgsl.expected.msl
+++ b/test/tint/statements/compound_assign/for_loop.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   uint i;
@@ -49,7 +48,8 @@
     int const tint_symbol_2 = idx1(tint_private_vars);
     int const tint_symbol_save = tint_symbol_2;
     a[tint_symbol_save] = (a[tint_symbol_save] * 2.0f);
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false);
       int const tint_symbol_3 = idx2(tint_private_vars);
       if (!((a[tint_symbol_3] < 10.0f))) {
         break;
diff --git a/test/tint/statements/decrement/complex.wgsl.expected.msl b/test/tint/statements/decrement/complex.wgsl.expected.msl
index e31aa32..93c8fd2 100644
--- a/test/tint/statements/decrement/complex.wgsl.expected.msl
+++ b/test/tint/statements/decrement/complex.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   uint v;
@@ -64,7 +63,8 @@
     int const tint_symbol_2_save_1 = tint_symbol_7;
     int const tint_symbol_3 = idx3(tint_private_vars);
     (*(tint_symbol_10))[tint_symbol_2_save].a[tint_symbol_2_save_1][tint_symbol_3] = as_type<int>((as_type<uint>((*(tint_symbol_10))[tint_symbol_2_save].a[tint_symbol_2_save_1][tint_symbol_3]) - as_type<uint>(1)));
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false);
       if (!(((*(tint_private_vars)).v < 10u))) {
         break;
       }
diff --git a/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.msl b/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.msl
index 46ba1a3..2f44b37 100644
--- a/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.msl
+++ b/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.msl
@@ -3,11 +3,11 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void tint_symbol(device uint* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(; (*(tint_symbol_1) < 10u); *(tint_symbol_1) = (*(tint_symbol_1) - 1u)) {
+  for(; (*(tint_symbol_1) < 10u); *(tint_symbol_1) = (*(tint_symbol_1) - 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.msl b/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.msl
index 53fcca7..783853d 100644
--- a/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.msl
+++ b/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.msl
@@ -3,11 +3,11 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void tint_symbol(device uint* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(*(tint_symbol_1) = (*(tint_symbol_1) - 1u); (*(tint_symbol_1) < 10u); ) {
+  for(*(tint_symbol_1) = (*(tint_symbol_1) - 1u); (*(tint_symbol_1) < 10u); ) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.msl b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.msl
index 5c54c35..f2295b4 100644
--- a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.msl
+++ b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.msl
@@ -3,8 +3,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   bool tint_discarded;
@@ -31,7 +30,8 @@
   int result = tint_ftoi(tint_symbol[0]);
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false);
       if (!((i < 10))) {
         break;
       }
diff --git a/test/tint/statements/discard/loop_discard_return.wgsl.expected.msl b/test/tint/statements/discard/loop_discard_return.wgsl.expected.msl
index 8779140..b522ea0 100644
--- a/test/tint/statements/discard/loop_discard_return.wgsl.expected.msl
+++ b/test/tint/statements/discard/loop_discard_return.wgsl.expected.msl
@@ -3,11 +3,11 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void f() {
-  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+  while(true) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     discard_fragment();
     return;
   }
diff --git a/test/tint/statements/discard/multiple_returns.wgsl.expected.msl b/test/tint/statements/discard/multiple_returns.wgsl.expected.msl
index 6373d1b..669ce6c 100644
--- a/test/tint/statements/discard/multiple_returns.wgsl.expected.msl
+++ b/test/tint/statements/discard/multiple_returns.wgsl.expected.msl
@@ -3,8 +3,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   bool tint_discarded;
@@ -22,7 +21,8 @@
   }
   if ((*(tint_symbol_3) < 0.0f)) {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false);
       if ((*(tint_symbol_3) > float(i))) {
         if (!(tint_private_vars.tint_discarded)) {
           *(tint_symbol_3) = float(i);
diff --git a/test/tint/statements/for/basic.wgsl.expected.msl b/test/tint/statements/for/basic.wgsl.expected.msl
index d43d02c..3fb7d2b 100644
--- a/test/tint/statements/for/basic.wgsl.expected.msl
+++ b/test/tint/statements/for/basic.wgsl.expected.msl
@@ -3,14 +3,14 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void some_loop_body() {
 }
 
 void f() {
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 5); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 5); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     some_loop_body();
   }
 }
diff --git a/test/tint/statements/for/complex.wgsl.expected.msl b/test/tint/statements/for/complex.wgsl.expected.msl
index a33d5d0..b21df71 100644
--- a/test/tint/statements/for/complex.wgsl.expected.msl
+++ b/test/tint/statements/for/complex.wgsl.expected.msl
@@ -3,15 +3,15 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void some_loop_body() {
 }
 
 void f() {
   int j = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; ((i < 5) && (j < 10)); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; ((i < 5) && (j < 10)); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     some_loop_body();
     j = as_type<int>((as_type<uint>(i) * as_type<uint>(30)));
   }
diff --git a/test/tint/statements/for/condition/array_ctor.wgsl.expected.msl b/test/tint/statements/for/condition/array_ctor.wgsl.expected.msl
index 1ff0ed4..2246141 100644
--- a/test/tint/statements/for/condition/array_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/condition/array_ctor.wgsl.expected.msl
@@ -3,12 +3,12 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void f() {
   int i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) for(; (i < 1); ) {
+  for(; (i < 1); ) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/statements/for/condition/basic.wgsl.expected.msl b/test/tint/statements/for/condition/basic.wgsl.expected.msl
index 960d032..5cb6019 100644
--- a/test/tint/statements/for/condition/basic.wgsl.expected.msl
+++ b/test/tint/statements/for/condition/basic.wgsl.expected.msl
@@ -3,12 +3,12 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void f() {
   int i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) for(; (i < 4); ) {
+  for(; (i < 4); ) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/statements/for/condition/struct_ctor.wgsl.expected.msl b/test/tint/statements/for/condition/struct_ctor.wgsl.expected.msl
index 848e21b..3e5e131 100644
--- a/test/tint/statements/for/condition/struct_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/condition/struct_ctor.wgsl.expected.msl
@@ -3,8 +3,7 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   int i;
@@ -12,7 +11,8 @@
 
 void f() {
   int i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) for(; (i < 1); ) {
+  for(; (i < 1); ) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/statements/for/continuing/array_ctor.wgsl.expected.msl b/test/tint/statements/for/continuing/array_ctor.wgsl.expected.msl
index f40f0a7..f26ddd1 100644
--- a/test/tint/statements/for/continuing/array_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/continuing/array_ctor.wgsl.expected.msl
@@ -3,12 +3,12 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void f() {
   int i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) for(; false; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(; false; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/statements/for/continuing/basic.wgsl.expected.msl b/test/tint/statements/for/continuing/basic.wgsl.expected.msl
index f40f0a7..f26ddd1 100644
--- a/test/tint/statements/for/continuing/basic.wgsl.expected.msl
+++ b/test/tint/statements/for/continuing/basic.wgsl.expected.msl
@@ -3,12 +3,12 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void f() {
   int i = 0;
-  TINT_ISOLATE_UB(tint_volatile_true) for(; false; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(; false; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.msl b/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.msl
index 549cf31..b6436c8 100644
--- a/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.msl
@@ -3,15 +3,15 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   int i;
 };
 
 void f() {
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; false; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; false; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/statements/for/empty.wgsl.expected.msl b/test/tint/statements/for/empty.wgsl.expected.msl
index 3637121..fb1eb0f 100644
--- a/test/tint/statements/for/empty.wgsl.expected.msl
+++ b/test/tint/statements/for/empty.wgsl.expected.msl
@@ -3,11 +3,11 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void f() {
-  TINT_ISOLATE_UB(tint_volatile_true) for(; false; ) {
+  for(; false; ) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/statements/for/initializer/array_ctor.wgsl.expected.msl b/test/tint/statements/for/initializer/array_ctor.wgsl.expected.msl
index 5caf7e9..722361f 100644
--- a/test/tint/statements/for/initializer/array_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/initializer/array_ctor.wgsl.expected.msl
@@ -3,11 +3,11 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void f() {
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 1; false; ) {
+  for(int i = 1; false; ) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/statements/for/initializer/basic.wgsl.expected.msl b/test/tint/statements/for/initializer/basic.wgsl.expected.msl
index d8c6ba5..25e4731 100644
--- a/test/tint/statements/for/initializer/basic.wgsl.expected.msl
+++ b/test/tint/statements/for/initializer/basic.wgsl.expected.msl
@@ -3,11 +3,11 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void f() {
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; false; ) {
+  for(int i = 0; false; ) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.msl b/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.msl
index 3afd169..a3917b9 100644
--- a/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.msl
@@ -3,15 +3,15 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct S {
   int i;
 };
 
 void f() {
-  TINT_ISOLATE_UB(tint_volatile_true) for(int i = 1; false; ) {
+  for(int i = 1; false; ) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/statements/for/scoping.wgsl.expected.msl b/test/tint/statements/for/scoping.wgsl.expected.msl
index 7ac803c..d4637af 100644
--- a/test/tint/statements/for/scoping.wgsl.expected.msl
+++ b/test/tint/statements/for/scoping.wgsl.expected.msl
@@ -3,11 +3,11 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void f() {
-  TINT_ISOLATE_UB(tint_volatile_true) for(int must_not_collide = 0; ; ) {
+  for(int must_not_collide = 0; ; ) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     break;
   }
   int must_not_collide = 0;
diff --git a/test/tint/statements/increment/complex.wgsl.expected.msl b/test/tint/statements/increment/complex.wgsl.expected.msl
index 8ed663a..55f9461 100644
--- a/test/tint/statements/increment/complex.wgsl.expected.msl
+++ b/test/tint/statements/increment/complex.wgsl.expected.msl
@@ -15,8 +15,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 struct tint_private_vars_struct {
   uint v;
@@ -64,7 +63,8 @@
     int const tint_symbol_2_save_1 = tint_symbol_7;
     int const tint_symbol_3 = idx3(tint_private_vars);
     (*(tint_symbol_10))[tint_symbol_2_save].a[tint_symbol_2_save_1][tint_symbol_3] = as_type<int>((as_type<uint>((*(tint_symbol_10))[tint_symbol_2_save].a[tint_symbol_2_save_1][tint_symbol_3]) + as_type<uint>(1)));
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false);
       if (!(((*(tint_private_vars)).v < 10u))) {
         break;
       }
diff --git a/test/tint/statements/increment/for_loop_continuing.wgsl.expected.msl b/test/tint/statements/increment/for_loop_continuing.wgsl.expected.msl
index c3f779c..a25d81a 100644
--- a/test/tint/statements/increment/for_loop_continuing.wgsl.expected.msl
+++ b/test/tint/statements/increment/for_loop_continuing.wgsl.expected.msl
@@ -3,11 +3,11 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void tint_symbol(device uint* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(; (*(tint_symbol_1) < 10u); *(tint_symbol_1) = (*(tint_symbol_1) + 1u)) {
+  for(; (*(tint_symbol_1) < 10u); *(tint_symbol_1) = (*(tint_symbol_1) + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/statements/increment/for_loop_initializer.wgsl.expected.msl b/test/tint/statements/increment/for_loop_initializer.wgsl.expected.msl
index 53257b5..91602a6 100644
--- a/test/tint/statements/increment/for_loop_initializer.wgsl.expected.msl
+++ b/test/tint/statements/increment/for_loop_initializer.wgsl.expected.msl
@@ -3,11 +3,11 @@
 using namespace metal;
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void tint_symbol(device uint* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(*(tint_symbol_1) = (*(tint_symbol_1) + 1u); (*(tint_symbol_1) < 10u); ) {
+  for(*(tint_symbol_1) = (*(tint_symbol_1) + 1u); (*(tint_symbol_1) < 10u); ) {
+    TINT_ISOLATE_UB(tint_volatile_false);
   }
 }
 
diff --git a/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.msl b/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.msl
index b806c1b..b32f26b 100644
--- a/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.msl
+++ b/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.msl
@@ -15,11 +15,11 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<tint_array<int, 3>, 2>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 6u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 6u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = (idx / 3u);
     uint const i_1 = (idx % 3u);
     (*(tint_symbol_1))[i][i_1] = 0;
diff --git a/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.msl b/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.msl
index 71c1d2d..f18d653 100644
--- a/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.msl
+++ b/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.msl
@@ -15,11 +15,11 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 void tint_zero_workgroup_memory(uint local_idx, threadgroup tint_array<int, 3>* const tint_symbol_1) {
-  TINT_ISOLATE_UB(tint_volatile_true) for(uint idx = local_idx; (idx < 3u); idx = (idx + 1u)) {
+  for(uint idx = local_idx; (idx < 3u); idx = (idx + 1u)) {
+    TINT_ISOLATE_UB(tint_volatile_false);
     uint const i = idx;
     (*(tint_symbol_1))[i] = 0;
   }