diff --git a/src/tint/lang/msl/writer/loop_test.cc b/src/tint/lang/msl/writer/loop_test.cc
index 330dd2c..f9e0d17 100644
--- a/src/tint/lang/msl/writer/loop_test.cc
+++ b/src/tint/lang/msl/writer/loop_test.cc
@@ -47,12 +47,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 a() {
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       break;
     }
   }
@@ -74,12 +74,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 a() {
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       {
         if (true) { break; }
       }
@@ -108,12 +108,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 a() {
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       bool v = true;
       {
         if (v) { break; }
@@ -144,13 +144,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 a() {
   {
     bool v = true;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       {
         if (v) { break; }
       }
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index 401d3f4..aa84321 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -219,6 +219,12 @@
         return array_template_name_;
     }
 
+    /// 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.
     /// @returns a call to the TINT_ISOLATE_UB macro, creating that macro on first call
     std::string IsolateUB() {
         if (isolate_ub_macro_name_.empty()) {
@@ -226,11 +232,10 @@
             isolate_ub_macro_name_ = UniqueIdentifier("TINT_ISOLATE_UB");
             Line();
             Line() << "#define " << isolate_ub_macro_name_ << "(VOLATILE_NAME) \\";
-            Line() << "  volatile bool VOLATILE_NAME = true; \\";
-            Line() << "  if (VOLATILE_NAME)";
+            Line() << "  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}";
         }
         StringStream ss;
-        ss << isolate_ub_macro_name_ << "(" << UniqueIdentifier("tint_volatile_true") << ")";
+        ss << isolate_ub_macro_name_ << "(" << UniqueIdentifier("tint_volatile_false") << ")";
         return ss.str();
     }
 
@@ -668,9 +673,10 @@
             ScopedIndent init(current_buffer_);
             EmitBlock(l->Initializer());
 
-            Line() << IsolateUB() << " while(true) {";
+            Line() << "while(true) {";
             {
                 ScopedIndent si(current_buffer_);
+                Line() << IsolateUB();
                 EmitBlock(l->Body());
             }
             Line() << "}";
diff --git a/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl
index 14a5111..fc71736 100644
--- a/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl
@@ -25,8 +25,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_2 {
   tint_array<int4, 4> tint_symbol_1;
@@ -63,7 +62,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl
index 09ba79a..27c0213 100644
--- a/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl
@@ -27,8 +27,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_2 {
   tint_array<int4, 4> tint_symbol_1;
@@ -63,7 +62,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl
index eeaddb1..184604a 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl
@@ -31,8 +31,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_2 {
   tint_array<int4, 4> tint_symbol_1;
@@ -67,7 +66,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl
index 9b28283..033a240 100644
--- a/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl
@@ -27,8 +27,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_4 {
   tint_array<int4, 4> tint_symbol_1;
@@ -65,7 +64,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
@@ -81,7 +81,8 @@
   {
     uint v_2 = 0u;
     v_2 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1)
       uint const v_3 = v_2;
       if ((v_3 >= 24u)) {
         break;
diff --git a/test/tint/array/strides.spvasm.expected.ir.msl b/test/tint/array/strides.spvasm.expected.ir.msl
index 431a06a..b355790 100644
--- a/test/tint/array/strides.spvasm.expected.ir.msl
+++ b/test/tint/array/strides.spvasm.expected.ir.msl
@@ -19,8 +19,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_1 {
   /* 0x0000 */ tint_array<tint_array<strided_arr, 2>, 3> el;
@@ -43,7 +42,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 2u)) {
         break;
@@ -61,7 +61,8 @@
   {
     uint v_2 = 0u;
     v_2 = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1)
       uint const v_3 = v_2;
       if ((v_3 >= 3u)) {
         break;
@@ -83,7 +84,8 @@
   {
     uint v_4 = 0u;
     v_4 = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_2)
       uint const v_5 = v_4;
       if ((v_5 >= 4u)) {
         break;
diff --git a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.ir.msl b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.ir.msl
index 351609c..d798fea 100644
--- a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.ir.msl
@@ -19,8 +19,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 Inner_packed_vec3 {
   /* 0x0000 */ float scalar_f32;
@@ -65,7 +64,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 2u)) {
         break;
diff --git a/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.ir.msl b/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.ir.msl
index 378d780..40572dd 100644
--- a/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.ir.msl
@@ -19,8 +19,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 packed_1;
@@ -88,7 +87,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 2u)) {
         break;
diff --git a/test/tint/buffer/storage/static_index/write.wgsl.expected.ir.msl b/test/tint/buffer/storage/static_index/write.wgsl.expected.ir.msl
index df814f3..5fee146 100644
--- a/test/tint/buffer/storage/static_index/write.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/static_index/write.wgsl.expected.ir.msl
@@ -19,8 +19,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 Inner {
   /* 0x0000 */ int scalar_i32;
@@ -69,7 +68,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 2u)) {
         break;
diff --git a/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.ir.msl b/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.ir.msl
index ea90738..f6d86f4 100644
--- a/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.ir.msl
@@ -21,8 +21,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 packed;
@@ -100,7 +99,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
@@ -118,7 +118,8 @@
   {
     uint v_2 = 0u;
     v_2 = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1)
       uint const v_3 = v_2;
       if ((v_3 >= 2u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
index ba3dfa8..4903a1f 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -19,8 +19,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 {
   tint_array<float2x2, 4> tint_symbol;
@@ -30,7 +29,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.ir.msl
index dec692e..0ef6ce0 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.ir.msl
@@ -19,8 +19,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_module_vars_struct {
   const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* u;
@@ -36,7 +35,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
index 9940bdc0..9ba26c5b 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -25,8 +25,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 {
   tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4> tint_symbol;
@@ -62,7 +61,8 @@
   {
     uint v_11 = 0u;
     v_11 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_12 = v_11;
       if ((v_12 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.ir.msl
index 0a23ec0..010faf1 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.ir.msl
@@ -19,8 +19,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_module_vars_struct {
   const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* u;
@@ -36,7 +35,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
index 121862f..6eede57 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -24,8 +24,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 {
   tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4> tint_symbol;
@@ -61,7 +60,8 @@
   {
     uint v_11 = 0u;
     v_11 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_12 = v_11;
       if ((v_12 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
index 3a70542..430020c 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -19,8 +19,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 {
   tint_array<half2x4, 4> tint_symbol;
@@ -30,7 +29,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
index e3cde16..61fc3d3 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -19,8 +19,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 {
   tint_array<float2x4, 4> tint_symbol;
@@ -30,7 +29,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.ir.msl
index 33d13f1..8c5e9a6 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.ir.msl
@@ -19,8 +19,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_module_vars_struct {
   const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* u;
@@ -37,7 +36,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
index 026c275..b3d3a51 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -24,8 +24,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 {
   tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4> tint_symbol;
@@ -69,7 +68,8 @@
   {
     uint v_15 = 0u;
     v_15 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_16 = v_15;
       if ((v_16 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
index aee2d2e..1a6b5dc 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -19,8 +19,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 {
   tint_array<float3x4, 4> tint_symbol;
@@ -30,7 +29,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
index b4321d6..c99d751 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -19,8 +19,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 {
   tint_array<half4x2, 4> tint_symbol;
@@ -30,7 +29,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
index ab5e665..e185b2a 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -19,8 +19,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 {
   tint_array<float4x2, 4> tint_symbol;
@@ -30,7 +29,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.ir.msl
index b00d8c5..76face9 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.ir.msl
@@ -19,8 +19,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_module_vars_struct {
   const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* u;
@@ -38,7 +37,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
index 52c9321..e2ab200 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -24,8 +24,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 {
   tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4> tint_symbol;
@@ -77,7 +76,8 @@
   {
     uint v_19 = 0u;
     v_19 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_20 = v_19;
       if ((v_20 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.ir.msl
index a6ad8aa..125ce31 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.ir.msl
@@ -19,8 +19,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_module_vars_struct {
   const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* u;
@@ -38,7 +37,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
index 1d30101..13068e7 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -24,8 +24,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 {
   tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4> tint_symbol;
@@ -77,7 +76,8 @@
   {
     uint v_19 = 0u;
     v_19 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_20 = v_19;
       if ((v_20 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
index 4f0e022..3505775 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -19,8 +19,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 {
   tint_array<half4x4, 4> tint_symbol;
@@ -30,7 +29,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
index 35ccec5..440abbd 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -19,8 +19,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 {
   tint_array<float4x4, 4> tint_symbol;
@@ -30,7 +29,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.ir.msl
index 89d8544..8b67dc7 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.ir.msl
@@ -22,8 +22,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_module_vars_struct {
   const constant tint_array<S, 4>* u;
@@ -40,7 +39,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
index 201ecd7..02b3af7 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -27,8 +27,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 {
   tint_array<S, 4> tint_symbol;
@@ -38,7 +37,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.ir.msl
index f51c619..7ebea3d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.ir.msl
@@ -23,8 +23,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_module_vars_struct {
   const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
index 81a2c8d..1ebc4a5 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -28,8 +28,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 {
   tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.ir.msl
index 2fa25c3..a4fd815 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.ir.msl
@@ -34,8 +34,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_module_vars_struct {
   const constant tint_array<S_packed_vec3, 4>* u;
@@ -65,7 +64,8 @@
   {
     uint v_4 = 0u;
     v_4 = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_5 = v_4;
       if ((v_5 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
index 8ba6fb0..e3ef41a 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -39,8 +39,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 {
   tint_array<S_packed_vec3, 4> tint_symbol;
@@ -79,7 +78,8 @@
   {
     uint v_7 = 0u;
     v_7 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_8 = v_7;
       if ((v_8 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.ir.msl
index f080673..1ea3778 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.ir.msl
@@ -34,8 +34,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_module_vars_struct {
   const constant tint_array<S_packed_vec3, 4>* u;
@@ -65,7 +64,8 @@
   {
     uint v_4 = 0u;
     v_4 = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_5 = v_4;
       if ((v_5 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
index ccddfbf..e8fa2b3 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -39,8 +39,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 {
   tint_array<S_packed_vec3, 4> tint_symbol;
@@ -79,7 +78,8 @@
   {
     uint v_7 = 0u;
     v_7 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_8 = v_7;
       if ((v_8 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.ir.msl
index 913cd10..851d6d2 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.ir.msl
@@ -23,8 +23,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_module_vars_struct {
   const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
index ad9574a..9b35ae4 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -28,8 +28,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 {
   tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.ir.msl
index 5d712dc..0c3eb5b 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.ir.msl
@@ -23,8 +23,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_module_vars_struct {
   const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
index 7d3170c..2d3491e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -28,8 +28,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 {
   tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.ir.msl
index e70ceb1..bf52720 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.ir.msl
@@ -22,8 +22,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_module_vars_struct {
   const constant tint_array<S, 4>* u;
@@ -40,7 +39,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
index 35f763e..44ff961 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -27,8 +27,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 {
   tint_array<S, 4> tint_symbol;
@@ -38,7 +37,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.ir.msl
index e425614..bba164c 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.ir.msl
@@ -23,8 +23,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_module_vars_struct {
   const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
index b845d9e..d3191c8 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -28,8 +28,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 {
   tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.ir.msl
index 07623e3..3c4752f 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.ir.msl
@@ -34,8 +34,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_module_vars_struct {
   const constant tint_array<S_packed_vec3, 4>* u;
@@ -67,7 +66,8 @@
   {
     uint v_5 = 0u;
     v_5 = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_6 = v_5;
       if ((v_6 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
index 1316725..bc96eb1 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -39,8 +39,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 {
   tint_array<S_packed_vec3, 4> tint_symbol;
@@ -81,7 +80,8 @@
   {
     uint v_8 = 0u;
     v_8 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_9 = v_8;
       if ((v_9 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.ir.msl
index 1d18969d..083aa88 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.ir.msl
@@ -33,8 +33,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_module_vars_struct {
   const constant tint_array<S_packed_vec3, 4>* u;
@@ -66,7 +65,8 @@
   {
     uint v_5 = 0u;
     v_5 = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_6 = v_5;
       if ((v_6 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
index c4cc786..70ffec1 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -38,8 +38,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 {
   tint_array<S_packed_vec3, 4> tint_symbol;
@@ -80,7 +79,8 @@
   {
     uint v_8 = 0u;
     v_8 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_9 = v_8;
       if ((v_9 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.ir.msl
index 6b56138..31faeaa 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.ir.msl
@@ -23,8 +23,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_module_vars_struct {
   const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
index 0b5d8a2..8153372 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -28,8 +28,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 {
   tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.ir.msl
index 8ff7e82..fb2085f 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.ir.msl
@@ -22,8 +22,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_module_vars_struct {
   const constant tint_array<S, 4>* u;
@@ -40,7 +39,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
index 6c637af..acad8ac 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -27,8 +27,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 {
   tint_array<S, 4> tint_symbol;
@@ -38,7 +37,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.ir.msl
index 139cb9e..2355cd3 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.ir.msl
@@ -22,8 +22,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_module_vars_struct {
   const constant tint_array<S, 4>* u;
@@ -40,7 +39,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
index 6f8d3af..9510bf9 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -27,8 +27,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 {
   tint_array<S, 4> tint_symbol;
@@ -38,7 +37,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.ir.msl
index 0a827db..ca17dbd 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.ir.msl
@@ -23,8 +23,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_module_vars_struct {
   const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
index 9a43ba5..04e94dc 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -28,8 +28,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 {
   tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.ir.msl
index c71d8eb..8dd5d90 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.ir.msl
@@ -34,8 +34,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_module_vars_struct {
   const constant tint_array<S_packed_vec3, 4>* u;
@@ -69,7 +68,8 @@
   {
     uint v_6 = 0u;
     v_6 = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_7 = v_6;
       if ((v_7 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
index 0c76569..98652ab 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -39,8 +39,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 {
   tint_array<S_packed_vec3, 4> tint_symbol;
@@ -83,7 +82,8 @@
   {
     uint v_9 = 0u;
     v_9 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_10 = v_9;
       if ((v_10 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.ir.msl
index d9bf522..79faedd 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.ir.msl
@@ -34,8 +34,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_module_vars_struct {
   const constant tint_array<S_packed_vec3, 4>* u;
@@ -69,7 +68,8 @@
   {
     uint v_6 = 0u;
     v_6 = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_7 = v_6;
       if ((v_7 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
index eb89dec..e1ca55e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -39,8 +39,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 {
   tint_array<S_packed_vec3, 4> tint_symbol;
@@ -83,7 +82,8 @@
   {
     uint v_9 = 0u;
     v_9 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_10 = v_9;
       if ((v_10 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.ir.msl
index 198b3a0..494cafd 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.ir.msl
@@ -23,8 +23,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_module_vars_struct {
   const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
index ea18e24..a5d3929 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -28,8 +28,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 {
   tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.ir.msl
index 473a54d..8bc1c3a 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.ir.msl
@@ -23,8 +23,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_module_vars_struct {
   const constant tint_array<S, 4>* u;
@@ -41,7 +40,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
index 08984e5..9e2e17a 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -28,8 +28,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 {
   tint_array<S, 4> tint_symbol;
@@ -39,7 +38,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl b/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl
index fe4107f..e8c2e82 100644
--- a/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl
@@ -2,13 +2,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 d() {
   int j = 0;
   {
-    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/chromium/1449538.wgsl.expected.ir.msl b/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl
index 3e5d3f6..b6f1da3 100644
--- a/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl
@@ -2,13 +2,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 f() {
   {
     int i0520 = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if (false) {
       } else {
         break;
@@ -20,7 +20,8 @@
   }
   {
     int i62 = 0;
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1)
       if (false) {
       } else {
         break;
@@ -32,7 +33,8 @@
   }
   {
     int i0520 = 0;
-    TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_2)
       if (false) {
       } else {
         break;
@@ -44,7 +46,8 @@
   }
   {
     int i62 = 0;
-    TINT_ISOLATE_UB(tint_volatile_true_3) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_3)
       if (false) {
       } else {
         break;
@@ -56,7 +59,8 @@
   }
   {
     int i62 = 0;
-    TINT_ISOLATE_UB(tint_volatile_true_4) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_4)
       if (false) {
       } else {
         break;
@@ -68,7 +72,8 @@
   }
   {
     int i60 = 0;
-    TINT_ISOLATE_UB(tint_volatile_true_5) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_5)
       if (false) {
       } else {
         break;
@@ -80,7 +85,8 @@
   }
   {
     int i62 = 0;
-    TINT_ISOLATE_UB(tint_volatile_true_6) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_6)
       if (false) {
       } else {
         break;
@@ -92,7 +98,8 @@
   }
   {
     int i60 = 0;
-    TINT_ISOLATE_UB(tint_volatile_true_7) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_7)
       if (false) {
       } else {
         break;
diff --git a/test/tint/bug/chromium/344265982.wgsl.expected.ir.msl b/test/tint/bug/chromium/344265982.wgsl.expected.ir.msl
index 4617798..f1b46a3 100644
--- a/test/tint/bug/chromium/344265982.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/344265982.wgsl.expected.ir.msl
@@ -14,8 +14,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_module_vars_struct {
   device tint_array<int, 4>* tint_symbol;
@@ -24,7 +23,8 @@
 void foo(device tint_array<int, 4>* const arg) {
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 4)) {
       } else {
         break;
diff --git a/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.ir.msl b/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.ir.msl
index 3a5398f..b8b2bc7 100644
--- a/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.ir.msl
@@ -32,8 +32,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 {
   S tint_symbol;
@@ -43,7 +42,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 64u)) {
         break;
diff --git a/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.ir.msl b/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.ir.msl
index d725486..c5c0ef0 100644
--- a/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.ir.msl
@@ -32,8 +32,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 {
   S tint_symbol;
@@ -43,7 +42,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 64u)) {
         break;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl
index 69b47a5..9ad691d 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl
@@ -9,13 +9,13 @@
 };
 
 #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(tint_module_vars_struct tint_module_vars) {
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 2)) {
       } else {
         break;
@@ -40,7 +40,8 @@
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v2f=(&v2f), .v3i=(&v3i), .v4u=(&v4u), .v2b=(&v2b)};
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1)
       if ((i < 2)) {
       } else {
         break;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl
index 4eb9cf4..0e0353d 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl
@@ -9,8 +9,7 @@
 };
 
 #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(tint_module_vars_struct tint_module_vars) {
   int i = 0;
@@ -28,7 +27,8 @@
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v2f=(&v2f), .v3i=(&v3i), .v4u=(&v4u), .v2b=(&v2b)};
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 2)) {
       } else {
         break;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl
index 34a3cd7..0c1bd3b 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl
@@ -2,8 +2,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;
@@ -20,7 +19,8 @@
   bool4 v4b = false;
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 2)) {
       } else {
         break;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl
index 3484c4e..44bc789 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl
@@ -2,8 +2,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;
@@ -16,7 +15,8 @@
   bool2 v2b_2 = false;
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 2)) {
       } else {
         break;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl
index 70f3124..e33ec1c 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl
@@ -2,8 +2,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;
@@ -20,7 +19,8 @@
   bool4 v4b = false;
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 2)) {
       } else {
         break;
diff --git a/test/tint/bug/tint/1064.wgsl.expected.ir.msl b/test/tint/bug/tint/1064.wgsl.expected.ir.msl
index e728d67..d7f982f 100644
--- a/test/tint/bug/tint/1064.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1064.wgsl.expected.ir.msl
@@ -2,12 +2,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;}
 
 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.ir.msl b/test/tint/bug/tint/1081.wgsl.expected.ir.msl
index 8bf645b..f562caa 100644
--- a/test/tint/bug/tint/1081.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1081.wgsl.expected.ir.msl
@@ -6,8 +6,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_outputs {
   int tint_symbol_1 [[color(2)]];
@@ -27,7 +26,8 @@
 int tint_symbol_inner(int3 x, tint_module_vars_struct tint_module_vars) {
   int y = x[0u];
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       int const r = f(y, tint_module_vars);
       if ((r == 0)) {
         break;
diff --git a/test/tint/bug/tint/1121.wgsl.expected.ir.msl b/test/tint/bug/tint/1121.wgsl.expected.ir.msl
index b0ca2dd..9c99a38 100644
--- a/test/tint/bug/tint/1121.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1121.wgsl.expected.ir.msl
@@ -57,8 +57,7 @@
 };
 
 #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_inner(uint3 GlobalInvocationID, tint_module_vars_struct tint_module_vars) {
   uint index = GlobalInvocationID[0u];
@@ -91,14 +90,16 @@
   int const TILE_COUNT_Y = 2;
   {
     int y = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((y < TILE_COUNT_Y)) {
       } else {
         break;
       }
       {
         int x = 0;
-        TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+        while(true) {
+          TINT_ISOLATE_UB(tint_volatile_false_1)
           if ((x < TILE_COUNT_X)) {
           } else {
             break;
@@ -118,7 +119,8 @@
           float dp = 0.0f;
           {
             uint i = 0u;
-            TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
+            while(true) {
+              TINT_ISOLATE_UB(tint_volatile_false_2)
               if ((i < 6u)) {
               } else {
                 break;
diff --git a/test/tint/bug/tint/1321.wgsl.expected.ir.msl b/test/tint/bug/tint/1321.wgsl.expected.ir.msl
index 537c53a..58ac15d 100644
--- a/test/tint/bug/tint/1321.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1321.wgsl.expected.ir.msl
@@ -14,8 +14,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;
@@ -25,7 +24,8 @@
   tint_array<float, 4> arr = tint_array<float, 4>{};
   {
     thread float* const a = (&arr[foo()]);
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       float const x = (*a);
       break;
     }
diff --git a/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl b/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl
index 9444e09..4d10141 100644
--- a/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl
@@ -2,12 +2,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() {
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if (true) {
       } else {
         break;
diff --git a/test/tint/bug/tint/1538.wgsl.expected.ir.msl b/test/tint/bug/tint/1538.wgsl.expected.ir.msl
index 2c5da95..102080b 100644
--- a/test/tint/bug/tint/1538.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1538.wgsl.expected.ir.msl
@@ -2,8 +2,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 {
@@ -27,7 +26,8 @@
 
 int f() {
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       g();
       break;
     }
@@ -39,7 +39,8 @@
 kernel void tint_symbol(device tint_array<uint, 1>* buf [[buffer(0)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.buf=buf};
   {
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1)
       if (((*tint_module_vars.buf)[0] == 0u)) {
         break;
       }
diff --git a/test/tint/bug/tint/1557.wgsl.expected.ir.msl b/test/tint/bug/tint/1557.wgsl.expected.ir.msl
index ad98ec0..adc3222 100644
--- a/test/tint/bug/tint/1557.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1557.wgsl.expected.ir.msl
@@ -2,8 +2,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_module_vars_struct {
   const constant int* u;
@@ -16,7 +15,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.ir.msl b/test/tint/bug/tint/1604.wgsl.expected.ir.msl
index 80c92e3..afba81b 100644
--- a/test/tint/bug/tint/1604.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1604.wgsl.expected.ir.msl
@@ -6,8 +6,7 @@
 };
 
 #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* x [[buffer(0)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.x=x};
@@ -15,7 +14,8 @@
     case 0:
     {
       {
-        TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+        while(true) {
+          TINT_ISOLATE_UB(tint_volatile_false)
           return;
         }
       }
diff --git a/test/tint/bug/tint/1605.wgsl.expected.ir.msl b/test/tint/bug/tint/1605.wgsl.expected.ir.msl
index 11318e0f..15bb159 100644
--- a/test/tint/bug/tint/1605.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1605.wgsl.expected.ir.msl
@@ -6,20 +6,21 @@
 };
 
 #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(tint_module_vars_struct tint_module_vars) {
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < (*tint_module_vars.b))) {
       } else {
         break;
       }
       {
         int j = -1;
-        TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+        while(true) {
+          TINT_ISOLATE_UB(tint_volatile_false_1)
           if ((j == 1)) {
           } else {
             break;
diff --git a/test/tint/bug/tint/1764.wgsl.expected.ir.msl b/test/tint/bug/tint/1764.wgsl.expected.ir.msl
index 49a49c8..f23a45c 100644
--- a/test/tint/bug/tint/1764.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1764.wgsl.expected.ir.msl
@@ -18,8 +18,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_2 {
   tint_array<int, 246> tint_symbol_1;
@@ -29,7 +28,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 246u)) {
         break;
diff --git a/test/tint/bug/tint/2010.spvasm.expected.ir.msl b/test/tint/bug/tint/2010.spvasm.expected.ir.msl
index 3e1fb27..51fd847 100644
--- a/test/tint/bug/tint/2010.spvasm.expected.ir.msl
+++ b/test/tint/bug/tint/2010.spvasm.expected.ir.msl
@@ -47,8 +47,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<S, 4096> tint_symbol_1;
@@ -66,7 +65,8 @@
   uint const x_52 = (*tint_module_vars.x_3)[0u];
   x_54 = 0u;
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint x_55 = 0u;
       x_58 = (*tint_module_vars.x_6).field0.field0;
       if ((x_54 < x_58)) {
@@ -100,7 +100,8 @@
   x_85 = x_76.xyxy;
   x_88 = 1u;
   {
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1)
       float4 x_111 = 0.0f;
       float4 x_86 = 0.0f;
       uint x_89 = 0u;
@@ -157,7 +158,8 @@
   {
     uint v_3 = 0u;
     v_3 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_2)
       uint const v_4 = v_3;
       if ((v_4 >= 4096u)) {
         break;
diff --git a/test/tint/bug/tint/2039.wgsl.expected.ir.msl b/test/tint/bug/tint/2039.wgsl.expected.ir.msl
index 8401406a..11378a4 100644
--- a/test/tint/bug/tint/2039.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2039.wgsl.expected.ir.msl
@@ -2,13 +2,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;
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       bool tint_continue = false;
       switch(2) {
         case 1:
diff --git a/test/tint/bug/tint/2059.wgsl.expected.ir.msl b/test/tint/bug/tint/2059.wgsl.expected.ir.msl
index 0d038a4..7fdf29c 100644
--- a/test/tint/bug/tint/2059.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2059.wgsl.expected.ir.msl
@@ -19,8 +19,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 S2_packed_vec3 {
   /* 0x0000 */ tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 1> m;
@@ -75,7 +74,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 1u)) {
         break;
@@ -97,7 +97,8 @@
   {
     uint v_2 = 0u;
     v_2 = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1)
       uint const v_3 = v_2;
       if ((v_3 >= 1u)) {
         break;
@@ -119,7 +120,8 @@
   {
     uint v_4 = 0u;
     v_4 = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_2)
       uint const v_5 = v_4;
       if ((v_5 >= 1u)) {
         break;
@@ -146,7 +148,8 @@
   float3x3 m = float3x3(0.0f);
   {
     uint c = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true_3) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_3)
       if ((c < 3u)) {
       } else {
         break;
diff --git a/test/tint/bug/tint/2201.wgsl.expected.ir.msl b/test/tint/bug/tint/2201.wgsl.expected.ir.msl
index aae30ea..2cc5558 100644
--- a/test/tint/bug/tint/2201.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2201.wgsl.expected.ir.msl
@@ -6,12 +6,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() {
   {
-    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.ir.msl b/test/tint/bug/tint/2202.wgsl.expected.ir.msl
index f4bd798..b7f4c947 100644
--- a/test/tint/bug/tint/2202.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2202.wgsl.expected.ir.msl
@@ -6,14 +6,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() {
   {
-    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)
           return;
         }
       }
diff --git a/test/tint/bug/tint/221.wgsl.expected.ir.msl b/test/tint/bug/tint/221.wgsl.expected.ir.msl
index 5fb7153..0025dd0 100644
--- a/test/tint/bug/tint/221.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/221.wgsl.expected.ir.msl
@@ -23,8 +23,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 uint tint_mod_u32(uint lhs, uint rhs) {
   uint const v = select(rhs, 1u, (rhs == 0u));
@@ -35,7 +34,8 @@
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.b=b};
   uint i = 0u;
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i >= (*tint_module_vars.b).count)) {
         break;
       }
diff --git a/test/tint/bug/tint/349291130.wgsl.expected.ir.msl b/test/tint/bug/tint/349291130.wgsl.expected.ir.msl
index be80066..57661c0 100644
--- a/test/tint/bug/tint/349291130.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/349291130.wgsl.expected.ir.msl
@@ -6,14 +6,14 @@
 };
 
 #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 [[texture(0)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.tint_symbol=tint_symbol};
   {
     uint level = tint_module_vars.tint_symbol.get_num_mip_levels();
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((level > 0u)) {
       } else {
         break;
diff --git a/test/tint/bug/tint/354627692.wgsl.expected.ir.msl b/test/tint/bug/tint/354627692.wgsl.expected.ir.msl
index 568f3b8..fb5f664 100644
--- a/test/tint/bug/tint/354627692.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/354627692.wgsl.expected.ir.msl
@@ -6,17 +6,18 @@
 };
 
 #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 [[buffer(0)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.tint_symbol=tint_symbol};
   int i = (*tint_module_vars.tint_symbol);
   {
-    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.ir.msl b/test/tint/bug/tint/366037039.wgsl.expected.ir.msl
index 3da7d69..85977d4 100644
--- a/test/tint/bug/tint/366037039.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/366037039.wgsl.expected.ir.msl
@@ -30,8 +30,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_module_vars_struct {
   const constant S_packed_vec3* ubuffer;
@@ -56,7 +55,8 @@
   {
     uint v = 0u;
     v = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/bug/tint/534.wgsl.expected.ir.msl b/test/tint/bug/tint/534.wgsl.expected.ir.msl
index bf985e2..270c65a 100644
--- a/test/tint/bug/tint/534.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/534.wgsl.expected.ir.msl
@@ -32,8 +32,7 @@
 };
 
 #define TINT_ISOLATE_UB(VOLATILE_NAME) \
-  volatile bool VOLATILE_NAME = true; \
-  if (VOLATILE_NAME)
+  {volatile bool VOLATILE_NAME = false; if (VOLATILE_NAME) break;}
 
 uint ConvertToFp16FloatValue(float fp32) {
   return 1u;
@@ -58,7 +57,8 @@
   uint4 dstColorBits = tint_v4f32_to_v4u32(dstColor);
   {
     uint i = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < (*tint_module_vars.uniforms).channelCount)) {
       } else {
         break;
diff --git a/test/tint/bug/tint/744.wgsl.expected.ir.msl b/test/tint/bug/tint/744.wgsl.expected.ir.msl
index f86490f..1a98eae 100644
--- a/test/tint/bug/tint/744.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/744.wgsl.expected.ir.msl
@@ -31,8 +31,7 @@
 };
 
 #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_inner(uint3 global_id, tint_module_vars_struct tint_module_vars) {
   uint2 const resultCell = uint2(global_id[1u], global_id[0u]);
@@ -41,7 +40,8 @@
   uint result = 0u;
   {
     uint i = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < dimInner)) {
       } else {
         break;
diff --git a/test/tint/bug/tint/757.wgsl.expected.ir.msl b/test/tint/bug/tint/757.wgsl.expected.ir.msl
index c4b51b5..0acd061 100644
--- a/test/tint/bug/tint/757.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/757.wgsl.expected.ir.msl
@@ -28,8 +28,7 @@
 };
 
 #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_inner(uint3 GlobalInvocationID, tint_module_vars_struct tint_module_vars) {
   uint flatIndex = (((4u * GlobalInvocationID[2u]) + (2u * GlobalInvocationID[1u])) + GlobalInvocationID[0u]);
@@ -37,7 +36,8 @@
   float4 texel = tint_module_vars.myTexture.read(uint2(int2(GlobalInvocationID.xy)), 0, 0);
   {
     uint i = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 1u)) {
       } else {
         break;
diff --git a/test/tint/bug/tint/914.wgsl.expected.ir.msl b/test/tint/bug/tint/914.wgsl.expected.ir.msl
index ebedfcc..ac94aea 100644
--- a/test/tint/bug/tint/914.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/914.wgsl.expected.ir.msl
@@ -33,8 +33,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_3 {
   tint_array<tint_array<float, 64>, 64> tint_symbol_1;
@@ -90,7 +89,8 @@
   {
     uint v_3 = 0u;
     v_3 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_4 = v_3;
       if ((v_4 >= 4096u)) {
         break;
@@ -114,7 +114,8 @@
   tint_array<float, 4> BCached = {};
   {
     uint index = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1)
       if ((index < 16u)) {
       } else {
         break;
@@ -132,21 +133,24 @@
   uint const tileRowB = (local_id[1u] * RowPerThreadB);
   {
     uint t = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_2)
       if ((t < numTiles)) {
       } else {
         break;
       }
       {
         uint innerRow = 0u;
-        TINT_ISOLATE_UB(tint_volatile_true_3) while(true) {
+        while(true) {
+          TINT_ISOLATE_UB(tint_volatile_false_3)
           if ((innerRow < 4u)) {
           } else {
             break;
           }
           {
             uint innerCol = 0u;
-            TINT_ISOLATE_UB(tint_volatile_true_4) while(true) {
+            while(true) {
+              TINT_ISOLATE_UB(tint_volatile_false_4)
               if ((innerCol < ColPerThreadA)) {
               } else {
                 break;
@@ -168,14 +172,16 @@
       }
       {
         uint innerRow = 0u;
-        TINT_ISOLATE_UB(tint_volatile_true_5) while(true) {
+        while(true) {
+          TINT_ISOLATE_UB(tint_volatile_false_5)
           if ((innerRow < RowPerThreadB)) {
           } else {
             break;
           }
           {
             uint innerCol = 0u;
-            TINT_ISOLATE_UB(tint_volatile_true_6) while(true) {
+            while(true) {
+              TINT_ISOLATE_UB(tint_volatile_false_6)
               if ((innerCol < 4u)) {
               } else {
                 break;
@@ -199,14 +205,16 @@
       threadgroup_barrier(mem_flags::mem_threadgroup);
       {
         uint k = 0u;
-        TINT_ISOLATE_UB(tint_volatile_true_7) while(true) {
+        while(true) {
+          TINT_ISOLATE_UB(tint_volatile_false_7)
           if ((k < 64u)) {
           } else {
             break;
           }
           {
             uint inner = 0u;
-            TINT_ISOLATE_UB(tint_volatile_true_8) while(true) {
+            while(true) {
+              TINT_ISOLATE_UB(tint_volatile_false_8)
               if ((inner < 4u)) {
               } else {
                 break;
@@ -220,7 +228,8 @@
           }
           {
             uint innerRow = 0u;
-            TINT_ISOLATE_UB(tint_volatile_true_9) while(true) {
+            while(true) {
+              TINT_ISOLATE_UB(tint_volatile_false_9)
               if ((innerRow < 4u)) {
               } else {
                 break;
@@ -228,7 +237,8 @@
               ACached = (*tint_module_vars.mm_Asub)[(tileRow + innerRow)][k];
               {
                 uint innerCol = 0u;
-                TINT_ISOLATE_UB(tint_volatile_true_10) while(true) {
+                while(true) {
+                  TINT_ISOLATE_UB(tint_volatile_false_10)
                   if ((innerCol < 4u)) {
                   } else {
                     break;
@@ -262,14 +272,16 @@
   }
   {
     uint innerRow = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true_11) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_11)
       if ((innerRow < 4u)) {
       } else {
         break;
       }
       {
         uint innerCol = 0u;
-        TINT_ISOLATE_UB(tint_volatile_true_12) while(true) {
+        while(true) {
+          TINT_ISOLATE_UB(tint_volatile_false_12)
           if ((innerCol < 4u)) {
           } else {
             break;
diff --git a/test/tint/bug/tint/942.wgsl.expected.ir.msl b/test/tint/bug/tint/942.wgsl.expected.ir.msl
index d053cc5..998db89 100644
--- a/test/tint/bug/tint/942.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/942.wgsl.expected.ir.msl
@@ -36,8 +36,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_2 {
   tint_array<tint_array<tint_packed_vec3_f32_array_element, 256>, 4> tint_symbol_1;
@@ -51,7 +50,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 1024u)) {
         break;
@@ -72,14 +72,16 @@
   uint2 const baseIndex = (v_4 - uint2(filterOffset, 0u));
   {
     uint r = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1)
       if ((r < 4u)) {
       } else {
         break;
       }
       {
         uint c = 0u;
-        TINT_ISOLATE_UB(tint_volatile_true_2) while(true) {
+        while(true) {
+          TINT_ISOLATE_UB(tint_volatile_false_2)
           if ((c < 4u)) {
           } else {
             break;
@@ -107,14 +109,16 @@
   threadgroup_barrier(mem_flags::mem_threadgroup);
   {
     uint r = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true_3) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_3)
       if ((r < 4u)) {
       } else {
         break;
       }
       {
         uint c = 0u;
-        TINT_ISOLATE_UB(tint_volatile_true_4) while(true) {
+        while(true) {
+          TINT_ISOLATE_UB(tint_volatile_false_4)
           if ((c < 4u)) {
           } else {
             break;
@@ -140,7 +144,8 @@
             float3 acc = float3(0.0f);
             {
               uint f = 0u;
-              TINT_ISOLATE_UB(tint_volatile_true_5) while(true) {
+              while(true) {
+                TINT_ISOLATE_UB(tint_volatile_false_5)
                 if ((f < (*tint_module_vars.params).filterDim)) {
                 } else {
                   break;
diff --git a/test/tint/bug/tint/948.wgsl.expected.ir.msl b/test/tint/bug/tint/948.wgsl.expected.ir.msl
index 5a500d2..d738f09 100644
--- a/test/tint/bug/tint/948.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/948.wgsl.expected.ir.msl
@@ -49,8 +49,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 main_out {
   float4 glFragColor_1;
@@ -122,7 +121,8 @@
   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 {
@@ -163,7 +163,8 @@
         (*tint_module_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/990.wgsl.expected.ir.msl b/test/tint/bug/tint/990.wgsl.expected.ir.msl
index b2ba140..b225169 100644
--- a/test/tint/bug/tint/990.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/990.wgsl.expected.ir.msl
@@ -2,14 +2,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 f() {
   int i = 0;
   {
     thread int* const p = (&i);
-    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/builtins/atomicStore/array/aliased_arrays.spvasm.expected.ir.msl b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.ir.msl
index 57f860a..e14cb85 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.ir.msl
@@ -19,8 +19,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 {
   tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol;
@@ -39,7 +38,8 @@
   uint idx = 0u;
   idx = local_invocation_index_2;
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if (!((idx < 6u))) {
         break;
       }
@@ -68,7 +68,8 @@
   {
     uint v_3 = 0u;
     v_3 = local_invocation_index_1_param;
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1)
       uint const v_4 = v_3;
       if ((v_4 >= 6u)) {
         break;
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.ir.msl b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.ir.msl
index ce7ebc2..85326e1 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.ir.msl
@@ -18,8 +18,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 {
   tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol;
@@ -29,7 +28,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 6u)) {
         break;
diff --git a/test/tint/builtins/atomicStore/array/array.spvasm.expected.ir.msl b/test/tint/builtins/atomicStore/array/array.spvasm.expected.ir.msl
index 8a9a376..cc5f715 100644
--- a/test/tint/builtins/atomicStore/array/array.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/array.spvasm.expected.ir.msl
@@ -19,8 +19,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 {
   tint_array<atomic_uint, 4> tint_symbol;
@@ -30,7 +29,8 @@
   uint idx = 0u;
   idx = local_invocation_index_2;
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if (!((idx < 4u))) {
         break;
       }
@@ -55,7 +55,8 @@
   {
     uint v = 0u;
     v = local_invocation_index_1_param;
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/builtins/atomicStore/array/array.wgsl.expected.ir.msl b/test/tint/builtins/atomicStore/array/array.wgsl.expected.ir.msl
index 62f1481..eede400 100644
--- a/test/tint/builtins/atomicStore/array/array.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.ir.msl
@@ -18,8 +18,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 {
   tint_array<atomic_uint, 4> tint_symbol;
@@ -29,7 +28,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.ir.msl b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.ir.msl
index 57f860a..e14cb85 100644
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.ir.msl
@@ -19,8 +19,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 {
   tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol;
@@ -39,7 +38,8 @@
   uint idx = 0u;
   idx = local_invocation_index_2;
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if (!((idx < 6u))) {
         break;
       }
@@ -68,7 +68,8 @@
   {
     uint v_3 = 0u;
     v_3 = local_invocation_index_1_param;
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1)
       uint const v_4 = v_3;
       if ((v_4 >= 6u)) {
         break;
diff --git a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.ir.msl b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.ir.msl
index ce7ebc2..85326e1 100644
--- a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.ir.msl
@@ -18,8 +18,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 {
   tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol;
@@ -29,7 +28,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 6u)) {
         break;
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.ir.msl b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.ir.msl
index cefd49b..ca989a3 100644
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.ir.msl
@@ -25,8 +25,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 {
   tint_array<S_atomic, 10> tint_symbol;
@@ -36,7 +35,8 @@
   uint idx = 0u;
   idx = local_invocation_index_2;
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if (!((idx < 10u))) {
         break;
       }
@@ -63,7 +63,8 @@
   {
     uint v = 0u;
     v = local_invocation_index_1_param;
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1)
       uint const v_1 = v;
       if ((v_1 >= 10u)) {
         break;
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.ir.msl b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.ir.msl
index a3173ad..6140181 100644
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.ir.msl
@@ -24,8 +24,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 {
   tint_array<S, 10> tint_symbol;
@@ -35,7 +34,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 10u)) {
         break;
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.ir.msl b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.ir.msl
index 2f3f91d..1ef5578 100644
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.ir.msl
@@ -25,8 +25,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 {
   S_atomic tint_symbol;
@@ -38,7 +37,8 @@
   (*tint_module_vars.wg).y = 0u;
   idx = local_invocation_index_2;
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if (!((idx < 10u))) {
         break;
       }
@@ -67,7 +67,8 @@
   {
     uint v = 0u;
     v = local_invocation_index_1_param;
-    TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false_1)
       uint const v_1 = v;
       if ((v_1 >= 10u)) {
         break;
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.ir.msl b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.ir.msl
index 0e4dbc6..e123238 100644
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.ir.msl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.ir.msl
@@ -24,8 +24,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 {
   S tint_symbol;
@@ -39,7 +38,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 10u)) {
         break;
diff --git a/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.ir.msl b/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.ir.msl
index 84b3346..c0a5876 100644
--- a/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.ir.msl
+++ b/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.ir.msl
@@ -6,13 +6,13 @@
 };
 
 #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(tint_module_vars_struct tint_module_vars) {
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 3)) {
       } else {
         break;
diff --git a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl
index efd369f..df06f9b 100644
--- a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl
@@ -7,13 +7,13 @@
 };
 
 #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(tint_module_vars_struct tint_module_vars) {
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       int const v = i;
       threadgroup_barrier(mem_flags::mem_threadgroup);
       int const v_1 = (*tint_module_vars.a);
diff --git a/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.ir.msl
index ce5135c..49b79da 100644
--- a/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.ir.msl
@@ -14,8 +14,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_inputs {
   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 v_1 = false;
       if ((x > v[0u])) {
         v_1 = (dfdx(1.0f) > 0.0f);
diff --git a/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl
index 3d98bff..7d349b6 100644
--- a/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl
@@ -19,8 +19,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_inputs {
   float x [[user(locn0)]];
@@ -29,7 +28,8 @@
 void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
   float4 v = float4(0.0f);
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((x > v[0u])) {
       } else {
         break;
diff --git a/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl
index eefa0a6..de7b62f 100644
--- a/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl
@@ -14,8 +14,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_inputs {
   float x [[user(locn0)]];
@@ -23,7 +22,8 @@
 
 void tint_symbol_inner(float x) {
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       dfdx(1.0f);
       {
         if ((x > 0.0f)) { break; }
diff --git a/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.ir.msl
index 5aca6c4..18c03d7 100644
--- a/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.ir.msl
@@ -14,8 +14,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_inputs {
   float x [[user(locn0)]];
@@ -23,7 +22,8 @@
 
 void tint_symbol_inner(float x) {
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       dfdx(1.0f);
       {
         if ((x > 0.0f)) { break; }
diff --git a/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.ir.msl
index d05523a..efa9d91 100644
--- a/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.ir.msl
@@ -14,8 +14,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_inputs {
   float x [[user(locn0)]];
@@ -23,7 +22,8 @@
 
 void tint_symbol_inner(float x) {
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       {
         dfdx(1.0f);
         if ((x > 0.0f)) { break; }
diff --git a/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.ir.msl
index 0edc31c..4c7af50 100644
--- a/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.ir.msl
@@ -14,8 +14,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_inputs {
   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 v_1 = false;
       if ((x > 0.0f)) {
         v_1 = (dfdx(1.0f) > 0.0f);
diff --git a/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl
index 8e828b2..93cc724 100644
--- a/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl
@@ -19,8 +19,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_inputs {
   float x [[user(locn0)]];
@@ -29,7 +28,8 @@
 void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
   float4 v = float4(0.0f);
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((x > v[0u])) {
       } else {
         break;
diff --git a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.ir.msl b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.ir.msl
index b966a93..6a9c791 100644
--- a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.ir.msl
+++ b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.ir.msl
@@ -19,8 +19,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 SSBO {
   /* 0x0000 */ tint_array<strided_arr, 2> m;
@@ -47,7 +46,8 @@
   {
     uint v_1 = 0u;
     v_1 = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_2 = v_1;
       if ((v_2 >= 2u)) {
         break;
diff --git a/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl b/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl
index ec8684d..8c5fb99 100644
--- a/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl
@@ -2,13 +2,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;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 4)) {
       } else {
         break;
diff --git a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl
index 5755d1f..d9fbf58 100644
--- a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl
@@ -2,13 +2,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;
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       bool tint_continue = false;
       switch(i) {
         case 0:
diff --git a/test/tint/loops/loop.wgsl.expected.ir.msl b/test/tint/loops/loop.wgsl.expected.ir.msl
index 5062fd6..d5e5fa2 100644
--- a/test/tint/loops/loop.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop.wgsl.expected.ir.msl
@@ -2,13 +2,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;}
 
 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.ir.msl b/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl
index 6348b86..b32bbacb 100644
--- a/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl
@@ -2,13 +2,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;}
 
 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.ir.msl b/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl
index 6d7a2b3..e21dbe4 100644
--- a/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl
@@ -2,13 +2,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;}
 
 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.ir.msl b/test/tint/loops/multiple_continues.wgsl.expected.ir.msl
index 8ef35c5..a0c6a80 100644
--- a/test/tint/loops/multiple_continues.wgsl.expected.ir.msl
+++ b/test/tint/loops/multiple_continues.wgsl.expected.ir.msl
@@ -2,13 +2,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;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 2)) {
       } else {
         break;
diff --git a/test/tint/loops/multiple_switch.wgsl.expected.ir.msl b/test/tint/loops/multiple_switch.wgsl.expected.ir.msl
index 2a0046b..5edf8d5 100644
--- a/test/tint/loops/multiple_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/multiple_switch.wgsl.expected.ir.msl
@@ -2,14 +2,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() {
   int i = 0;
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 2)) {
       } else {
         break;
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl
index 1eff924..3090c93 100644
--- a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl
@@ -2,20 +2,21 @@
 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;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 2)) {
       } else {
         break;
       }
       {
         int j = 0;
-        TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+        while(true) {
+          TINT_ISOLATE_UB(tint_volatile_false_1)
           if ((j < 2)) {
           } else {
             break;
diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl
index 08e2cec..eb3d97c 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl
@@ -2,13 +2,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;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 2)) {
       } else {
         break;
@@ -19,7 +19,8 @@
         {
           {
             int j = 0;
-            TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+            while(true) {
+              TINT_ISOLATE_UB(tint_volatile_false_1)
               if ((j < 2)) {
               } else {
                 break;
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl
index 14e9c8b..6bc5dec 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl
@@ -2,14 +2,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() {
   int k = 0;
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 2)) {
       } else {
         break;
@@ -20,7 +20,8 @@
         {
           {
             int j = 0;
-            TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+            while(true) {
+              TINT_ISOLATE_UB(tint_volatile_false_1)
               if ((j < 2)) {
               } else {
                 break;
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl
index 6ed9ba5..208823c 100644
--- a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl
@@ -2,14 +2,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() {
   int j = 0;
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 2)) {
       } else {
         break;
diff --git a/test/tint/loops/nested_loops.wgsl.expected.ir.msl b/test/tint/loops/nested_loops.wgsl.expected.ir.msl
index fe13807..daa4a15 100644
--- a/test/tint/loops/nested_loops.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loops.wgsl.expected.ir.msl
@@ -2,20 +2,21 @@
 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.ir.msl b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl
index 1e711ce..8d2609f 100644
--- a/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl
@@ -2,19 +2,20 @@
 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.ir.msl b/test/tint/loops/single_continue.wgsl.expected.ir.msl
index f830ab8..5321463 100644
--- a/test/tint/loops/single_continue.wgsl.expected.ir.msl
+++ b/test/tint/loops/single_continue.wgsl.expected.ir.msl
@@ -2,13 +2,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;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 2)) {
       } else {
         break;
diff --git a/test/tint/loops/while.wgsl.expected.ir.msl b/test/tint/loops/while.wgsl.expected.ir.msl
index a2e7eb1..62623af 100644
--- a/test/tint/loops/while.wgsl.expected.ir.msl
+++ b/test/tint/loops/while.wgsl.expected.ir.msl
@@ -2,13 +2,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;}
 
 int f() {
   int i = 0;
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 4)) {
       } else {
         break;
diff --git a/test/tint/loops/while_with_continue.wgsl.expected.ir.msl b/test/tint/loops/while_with_continue.wgsl.expected.ir.msl
index a2e7eb1..62623af 100644
--- a/test/tint/loops/while_with_continue.wgsl.expected.ir.msl
+++ b/test/tint/loops/while_with_continue.wgsl.expected.ir.msl
@@ -2,13 +2,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;}
 
 int f() {
   int i = 0;
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 4)) {
       } else {
         break;
diff --git a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.ir.msl
index 88d037f..fbcdcaa 100644
--- a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.ir.msl
@@ -22,8 +22,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_2 {
   str tint_symbol_1;
@@ -37,7 +36,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.ir.msl
index 6410c42..a50727f 100644
--- a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.ir.msl
@@ -22,8 +22,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_2 {
   tint_array<str, 4> tint_symbol_1;
@@ -37,7 +36,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.ir.msl
index 801cdb0..e3a3a94 100644
--- a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.ir.msl
@@ -22,8 +22,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_2 {
   str tint_symbol_1;
@@ -37,7 +36,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.ir.msl
index 458af39..cb2fb04 100644
--- a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.ir.msl
@@ -22,8 +22,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_2 {
   tint_array<str, 4> tint_symbol_1;
@@ -37,7 +36,8 @@
   {
     uint v = 0u;
     v = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_1 = v;
       if ((v_1 >= 4u)) {
         break;
diff --git a/test/tint/samples/compute_boids.wgsl.expected.ir.msl b/test/tint/samples/compute_boids.wgsl.expected.ir.msl
index 394da95..5144d1b 100644
--- a/test/tint/samples/compute_boids.wgsl.expected.ir.msl
+++ b/test/tint/samples/compute_boids.wgsl.expected.ir.msl
@@ -39,8 +39,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 vert_main_outputs {
   float4 tint_symbol [[position]];
@@ -85,7 +84,8 @@
   float2 vel = 0.0f;
   {
     uint i = 0u;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 5u)) {
       } else {
         break;
diff --git a/test/tint/shadowing/loop.wgsl.expected.ir.msl b/test/tint/shadowing/loop.wgsl.expected.ir.msl
index 32f654a..bf61e09 100644
--- a/test/tint/shadowing/loop.wgsl.expected.ir.msl
+++ b/test/tint/shadowing/loop.wgsl.expected.ir.msl
@@ -18,14 +18,14 @@
 };
 
 #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>* output [[buffer(0)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.output=output};
   int i = 0;
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       int x = (*tint_module_vars.output)[i];
       {
         int x = (*tint_module_vars.output)[x];
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.ir.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.ir.msl
index c4065b5..a5ef9e6 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.ir.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.ir.msl
@@ -30,8 +30,7 @@
 };
 
 #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 Uniforms* uniforms [[buffer(0)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.uniforms=uniforms};
@@ -39,7 +38,8 @@
   OuterS s1 = {};
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 4)) {
       } else {
         break;
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.ir.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.ir.msl
index 5c483ed..140855f 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.ir.msl
@@ -30,8 +30,7 @@
 };
 
 #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 Uniforms* uniforms [[buffer(0)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.uniforms=uniforms};
@@ -39,7 +38,8 @@
   OuterS s1 = {};
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 4)) {
       } else {
         break;
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.ir.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.ir.msl
index 1226caa..b2a5dc7 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.ir.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.ir.msl
@@ -30,8 +30,7 @@
 };
 
 #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 Uniforms* uniforms [[buffer(0)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.uniforms=uniforms};
@@ -40,7 +39,8 @@
   int i = 0;
   {
     s1.a1[(*tint_module_vars.uniforms).i] = v;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 4)) {
       } else {
         break;
diff --git a/test/tint/statements/compound_assign/for_loop.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/for_loop.wgsl.expected.ir.msl
index c84338f..f3c5980 100644
--- a/test/tint/statements/compound_assign/for_loop.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/for_loop.wgsl.expected.ir.msl
@@ -25,8 +25,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 idx1(tint_module_vars_struct tint_module_vars) {
   (*tint_module_vars.i) = ((*tint_module_vars.i) + 1u);
@@ -48,7 +47,8 @@
   {
     thread float* const v_1 = (&a[idx1(tint_module_vars)]);
     (*v_1) = ((*v_1) * 2.0f);
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((a[idx2(tint_module_vars)] < 10.0f)) {
       } else {
         break;
diff --git a/test/tint/statements/decrement/complex.wgsl.expected.ir.msl b/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
index d932ca4..e9505f5 100644
--- a/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
@@ -23,8 +23,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 idx1(tint_module_vars_struct tint_module_vars) {
   (*tint_module_vars.v) = ((*tint_module_vars.v) - 1u);
@@ -62,7 +61,8 @@
     device int4* const v_2 = (&(*tint_module_vars.tint_symbol)[v_1].a[idx2(tint_module_vars)]);
     int const v_3 = idx3(tint_module_vars);
     (*v_2)[v_3] = as_type<int>((as_type<uint>((*v_2)[v_3]) - as_type<uint>(1)));
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if (((*tint_module_vars.v) < 10u)) {
       } else {
         break;
diff --git a/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.ir.msl b/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.ir.msl
index a809d8b..df04b49 100644
--- a/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.ir.msl
@@ -6,12 +6,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 tint_symbol(tint_module_vars_struct tint_module_vars) {
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if (((*tint_module_vars.i) < 10u)) {
       } else {
         break;
diff --git a/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.ir.msl b/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.ir.msl
index aecb1e3..6bb40ad 100644
--- a/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.ir.msl
@@ -6,13 +6,13 @@
 };
 
 #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(tint_module_vars_struct tint_module_vars) {
   {
     (*tint_module_vars.i) = ((*tint_module_vars.i) - 1u);
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if (((*tint_module_vars.i) < 10u)) {
       } else {
         break;
diff --git a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
index 39ed593..7e4bc07 100644
--- a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
@@ -9,8 +9,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 foo_outputs {
   int tint_symbol [[color(0)]];
@@ -32,7 +31,8 @@
   int result = tint_f32_to_i32(tint_module_vars.t.sample(tint_module_vars.s, coord)[0u]);
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 10)) {
       } else {
         break;
diff --git a/test/tint/statements/discard/loop_discard_return.wgsl.expected.ir.msl b/test/tint/statements/discard/loop_discard_return.wgsl.expected.ir.msl
index 44045f7..752536e 100644
--- a/test/tint/statements/discard/loop_discard_return.wgsl.expected.ir.msl
+++ b/test/tint/statements/discard/loop_discard_return.wgsl.expected.ir.msl
@@ -6,12 +6,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 f(tint_module_vars_struct tint_module_vars) {
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       (*tint_module_vars.continue_execution) = false;
       return;
     }
diff --git a/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl b/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl
index 7303f09..1c9efef 100644
--- a/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl
+++ b/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl
@@ -8,8 +8,7 @@
 };
 
 #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(device int* non_uniform_global [[buffer(0)]], device float* output [[buffer(1)]]) {
   thread bool continue_execution = true;
@@ -24,7 +23,8 @@
   if (((*tint_module_vars.output) < 0.0f)) {
     int i = 0;
     {
-      TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+      while(true) {
+        TINT_ISOLATE_UB(tint_volatile_false)
         float const v_1 = (*tint_module_vars.output);
         if ((v_1 > float(i))) {
           float const v_2 = float(i);
diff --git a/test/tint/statements/for/basic.wgsl.expected.ir.msl b/test/tint/statements/for/basic.wgsl.expected.ir.msl
index fb42656..97465a8 100644
--- a/test/tint/statements/for/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/basic.wgsl.expected.ir.msl
@@ -2,8 +2,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;}
 
 void some_loop_body() {
 }
@@ -11,7 +10,8 @@
 void f() {
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 5)) {
       } else {
         break;
diff --git a/test/tint/statements/for/complex.wgsl.expected.ir.msl b/test/tint/statements/for/complex.wgsl.expected.ir.msl
index 16f4e4e..4dd60ee 100644
--- a/test/tint/statements/for/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/complex.wgsl.expected.ir.msl
@@ -2,8 +2,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;}
 
 void some_loop_body() {
 }
@@ -12,7 +11,8 @@
   int j = 0;
   {
     int i = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       bool v = false;
       if ((i < 5)) {
         v = (j < 10);
diff --git a/test/tint/statements/for/condition/array_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/condition/array_ctor.wgsl.expected.ir.msl
index 3204d4a..02a71f0 100644
--- a/test/tint/statements/for/condition/array_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/condition/array_ctor.wgsl.expected.ir.msl
@@ -2,13 +2,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 f() {
   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/statements/for/condition/basic.wgsl.expected.ir.msl b/test/tint/statements/for/condition/basic.wgsl.expected.ir.msl
index 52e6a2f..fae57d4 100644
--- a/test/tint/statements/for/condition/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/condition/basic.wgsl.expected.ir.msl
@@ -2,13 +2,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 f() {
   int i = 0;
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if ((i < 4)) {
       } else {
         break;
diff --git a/test/tint/statements/for/condition/struct_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/condition/struct_ctor.wgsl.expected.ir.msl
index 3204d4a..02a71f0 100644
--- a/test/tint/statements/for/condition/struct_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/condition/struct_ctor.wgsl.expected.ir.msl
@@ -2,13 +2,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 f() {
   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/statements/for/continuing/array_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl
index 9fe44a7..76a659f 100644
--- a/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl
@@ -2,13 +2,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 f() {
   int i = 0;
   {
-    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/statements/for/continuing/basic.wgsl.expected.ir.msl b/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl
index 9fe44a7..76a659f 100644
--- a/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl
@@ -2,13 +2,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 f() {
   int i = 0;
   {
-    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/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl
index c5a4b75..ae63ee6 100644
--- a/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl
@@ -2,13 +2,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 f() {
   {
     int i = 0;
-    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/statements/for/empty.wgsl.expected.ir.msl b/test/tint/statements/for/empty.wgsl.expected.ir.msl
index 986f6cc..7b3f2ce 100644
--- a/test/tint/statements/for/empty.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/empty.wgsl.expected.ir.msl
@@ -2,12 +2,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() {
   {
-    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/statements/for/initializer/array_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/initializer/array_ctor.wgsl.expected.ir.msl
index 6be977a..a1a7b9a 100644
--- a/test/tint/statements/for/initializer/array_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/initializer/array_ctor.wgsl.expected.ir.msl
@@ -2,13 +2,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 f() {
   {
     int i = 1;
-    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/statements/for/initializer/basic.wgsl.expected.ir.msl b/test/tint/statements/for/initializer/basic.wgsl.expected.ir.msl
index b86deae..c3e6673 100644
--- a/test/tint/statements/for/initializer/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/initializer/basic.wgsl.expected.ir.msl
@@ -2,13 +2,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 f() {
   {
     int i = 0;
-    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/statements/for/initializer/struct_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.ir.msl
index 6be977a..a1a7b9a 100644
--- a/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.ir.msl
@@ -2,13 +2,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 f() {
   {
     int i = 1;
-    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/statements/for/scoping.wgsl.expected.ir.msl b/test/tint/statements/for/scoping.wgsl.expected.ir.msl
index efaba43..5b08034 100644
--- a/test/tint/statements/for/scoping.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/scoping.wgsl.expected.ir.msl
@@ -2,13 +2,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 f() {
   {
     int must_not_collide = 0;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       break;
     }
   }
diff --git a/test/tint/statements/increment/complex.wgsl.expected.ir.msl b/test/tint/statements/increment/complex.wgsl.expected.ir.msl
index 7f62d2b..b4cc61f 100644
--- a/test/tint/statements/increment/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/complex.wgsl.expected.ir.msl
@@ -23,8 +23,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 idx1(tint_module_vars_struct tint_module_vars) {
   (*tint_module_vars.v) = ((*tint_module_vars.v) + 1u);
@@ -62,7 +61,8 @@
     device int4* const v_2 = (&(*tint_module_vars.tint_symbol)[v_1].a[idx2(tint_module_vars)]);
     int const v_3 = idx3(tint_module_vars);
     (*v_2)[v_3] = as_type<int>((as_type<uint>((*v_2)[v_3]) + as_type<uint>(1)));
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if (((*tint_module_vars.v) < 10u)) {
       } else {
         break;
diff --git a/test/tint/statements/increment/for_loop_continuing.wgsl.expected.ir.msl b/test/tint/statements/increment/for_loop_continuing.wgsl.expected.ir.msl
index 8c17b12..1aa5603 100644
--- a/test/tint/statements/increment/for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/for_loop_continuing.wgsl.expected.ir.msl
@@ -6,12 +6,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 tint_symbol(tint_module_vars_struct tint_module_vars) {
   {
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if (((*tint_module_vars.i) < 10u)) {
       } else {
         break;
diff --git a/test/tint/statements/increment/for_loop_initializer.wgsl.expected.ir.msl b/test/tint/statements/increment/for_loop_initializer.wgsl.expected.ir.msl
index 0a1e473..550762b 100644
--- a/test/tint/statements/increment/for_loop_initializer.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/for_loop_initializer.wgsl.expected.ir.msl
@@ -6,13 +6,13 @@
 };
 
 #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(tint_module_vars_struct tint_module_vars) {
   {
     (*tint_module_vars.i) = ((*tint_module_vars.i) + 1u);
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       if (((*tint_module_vars.i) < 10u)) {
       } else {
         break;
diff --git a/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.ir.msl b/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.ir.msl
index fdd9f42..4d50db2 100644
--- a/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.ir.msl
@@ -18,8 +18,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_2 {
   tint_array<tint_array<int, 3>, 2> tint_symbol_1;
@@ -29,7 +28,8 @@
   {
     uint v_1 = 0u;
     v_1 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_2 = v_1;
       if ((v_2 >= 6u)) {
         break;
diff --git a/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.ir.msl b/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.ir.msl
index a400836..1fc2b21 100644
--- a/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.ir.msl
@@ -18,8 +18,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_2 {
   tint_array<int, 3> tint_symbol_1;
@@ -29,7 +28,8 @@
   {
     uint v_1 = 0u;
     v_1 = tint_local_index;
-    TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    while(true) {
+      TINT_ISOLATE_UB(tint_volatile_false)
       uint const v_2 = v_1;
       if ((v_2 >= 3u)) {
         break;