[tint][msl]: Use a volatile bool to control loop preservation

To preserve loops, compare the condition with the contents
of a volatile boolean variable that is always set to 'true'.

Do this instead of using __asm__(""); that has proved faulty
in certain macOS environments.

Bug: tint:2125
Change-Id: Ia700db8d8612006eb614c30e6c647ca30032d120
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/170720
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
index 63dcf81..58af4e0 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -2122,8 +2122,7 @@
     };
 
     TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
-    Line() << "while (true) {";
-    EmitLoopPreserver();
+    EmitUnconditionalLoopHeader();
     {
         ScopedIndent si(this);
         if (!EmitStatements(stmt->body->statements)) {
@@ -2193,8 +2192,7 @@
         };
 
         TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
-        Line() << "while (true) {";
-        EmitLoopPreserver();
+        EmitUnconditionalLoopHeader();
         IncrementIndent();
         TINT_DEFER({
             DecrementIndent();
@@ -2227,7 +2225,8 @@
                     out << "; ";
                 }
 
-                out << cond_buf.str() << "; ";
+                EmitLoopCondition(out, cond_buf.str());
+                out << "; ";
 
                 if (!cont_buf.lines.empty()) {
                     out << tint::TrimSuffix(cont_buf.lines[0].content, ";");
@@ -2235,7 +2234,6 @@
             }
             out << " {";
         }
-        EmitLoopPreserver();
         {
             auto emit_continuing = [] { return true; };
             TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
@@ -2268,8 +2266,7 @@
     // as a regular while in MSL. Instead we need to generate a `while(true)` loop.
     bool emit_as_loop = cond_pre.lines.size() > 0;
     if (emit_as_loop) {
-        Line() << "while (true) {";
-        EmitLoopPreserver();
+        EmitUnconditionalLoopHeader();
         IncrementIndent();
         TINT_DEFER({
             DecrementIndent();
@@ -2288,11 +2285,10 @@
             out << "while";
             {
                 ScopedParen sp(out);
-                out << cond_buf.str();
+                EmitLoopCondition(out, cond_buf.str());
             }
             out << " {";
         }
-        EmitLoopPreserver();
         if (!EmitStatementsWithIndent(stmt->body->statements)) {
             return false;
         }
@@ -3031,32 +3027,25 @@
     return true;
 }
 
-void ASTPrinter::EmitLoopPreserver() {
-    IncrementIndent();
-    // This statement prevents the MSL compiler from erasing a loop during
-    // optimizations.  In the AIR dialiect of LLVM IR, WGSL loops should compile
-    // to a loop that contains an 'asm' call with a 'sideeffect' annotation.
-    //
-    // For example, compile a WGSL file with a trivial while(1) loop to 'a.metal',
-    // then compile that to AIR (LLVM IR dialect):
-    //
-    //    xcrun metal a.metal -S -o -
-    //
-    // The loop in the AIR should look something like this:
-    //
-    //    1: ...
-    //      br label %2
-    //
-    //    2:                                      ; preds = %1, %2
-    //      tail call void asm sideeffect "", ""() #1, !srcloc !27
-    //      br label %2, !llvm.loop !28
-    //
-    // It is important that the 'sideeffect' annotation exist. That tells the
-    // optimizer that the instruction has side effects invisible to the
-    // optimizer, and therefore the loop should not be eliminated.
-    Line() << R"(__asm__("");)";
+std::string_view ASTPrinter::LoopPreservingVar() {
+    if (loop_preserving_var_.empty()) {
+        loop_preserving_var_ = UniqueIdentifier("tint_preserve_loop");
+        Line(&helpers_) << "constant static volatile bool " << loop_preserving_var_ << " = true;";
+        Line(&helpers_);
+    }
+    return loop_preserving_var_;
+}
 
-    DecrementIndent();
+void ASTPrinter::EmitLoopCondition(StringStream& out, const std::string& cond) {
+    if (cond.empty()) {
+        out << LoopPreservingVar();
+    } else {
+        out << "(" << cond << ") == " << LoopPreservingVar();
+    }
+}
+
+void ASTPrinter::EmitUnconditionalLoopHeader() {
+    Line() << "while (" << LoopPreservingVar() << ") {";
 }
 
 template <typename F>
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.h b/src/tint/lang/msl/writer/ast_printer/ast_printer.h
index af251db..a13f907 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.h
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.h
@@ -378,9 +378,18 @@
                               const ast::CallExpression* expr,
                               const sem::BuiltinFn* builtin);
 
-    /// Emits a code sequence that preserves a loop during
-    /// optimizations even if the loop is infinite.
-    void EmitLoopPreserver();
+    /// Lazily emits the global-scope declaration of the helper static volatile boolean variable
+    /// used to fool the MSL compiler into thinking loops might exit.
+    /// @return the name of the variable.
+    std::string_view LoopPreservingVar();
+
+    /// Emits a loop condition, but santized so that the MSL compiler can't infer that the loop
+    /// never exits.
+    void EmitLoopCondition(StringStream& out, const std::string& cond);
+
+    /// Emits the header of an unconditional loop, but use the loop-preserving condition to fool the
+    /// MSL compiler into thinking that the loop might exit.
+    void EmitUnconditionalLoopHeader();
 
     /// Handles generating a builtin name
     /// @param builtin the semantic info for the builtin
@@ -437,6 +446,9 @@
 
     std::function<bool()> emit_continuing_;
 
+    // Name of the variable used to ensure the MSL compiler thinks a loop will terminate.
+    std::string loop_preserving_var_;
+
     /// Name of atomicCompareExchangeWeak() helper for the given pointer storage
     /// class and struct return type
     using ACEWKeyType =
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
index 0477e48..b0a48d8 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
@@ -226,13 +226,14 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_3 {
   tint_array<float2x2, 4> m;
 };
 
 void comp_main_inner(uint local_invocation_index, threadgroup tint_array<float2x2, 4>* const tint_symbol) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i] = float2x2(float2(0.0f), float2(0.0f));
   }
diff --git a/src/tint/lang/msl/writer/ast_printer/continue_test.cc b/src/tint/lang/msl/writer/ast_printer/continue_test.cc
index a79ddc3..9b4dd02 100644
--- a/src/tint/lang/msl/writer/ast_printer/continue_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/continue_test.cc
@@ -39,16 +39,23 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
+    ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
+    EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
 
-    ASSERT_TRUE(gen.EmitStatement(loop)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while (true) {
-    __asm__("");
+using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
+kernel void test_function() {
+  while (tint_preserve_loop) {
     if (false) {
       break;
     }
     continue;
   }
+  return;
+}
+
 )");
 }
 
diff --git a/src/tint/lang/msl/writer/ast_printer/loop_test.cc b/src/tint/lang/msl/writer/ast_printer/loop_test.cc
index 450592b..73e7cdc 100644
--- a/src/tint/lang/msl/writer/ast_printer/loop_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/loop_test.cc
@@ -46,13 +46,20 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
+    ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
+    EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
 
-    ASSERT_TRUE(gen.EmitStatement(l)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while (true) {
-    __asm__("");
+using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
+fragment void F() {
+  while (tint_preserve_loop) {
     break;
   }
+  return;
+}
+
 )");
 }
 
@@ -67,16 +74,26 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
+    ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
+    EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
 
-    ASSERT_TRUE(gen.EmitStatement(l)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while (true) {
-    __asm__("");
+using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
+void a_statement() {
+}
+
+fragment void F() {
+  while (tint_preserve_loop) {
     break;
     {
       a_statement();
     }
   }
+  return;
+}
+
 )");
 }
 
@@ -91,17 +108,27 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
+    ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
+    EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
 
-    ASSERT_TRUE(gen.EmitStatement(l)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while (true) {
-    __asm__("");
+using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
+void a_statement() {
+}
+
+fragment void F() {
+  while (tint_preserve_loop) {
     break;
     {
       a_statement();
       if (true) { break; }
     }
   }
+  return;
+}
+
 )");
 }
 
@@ -125,23 +152,19 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(outer)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while (true) {
-    __asm__("");
-    while (true) {
-      __asm__("");
-      break;
-      {
-        a_statement();
-      }
-    }
+    EXPECT_EQ(gen.Result(), R"(while (tint_preserve_loop) {
+  while (tint_preserve_loop) {
+    break;
     {
-      lhs = rhs;
-      if (true) { break; }
+      a_statement();
     }
   }
+  {
+    lhs = rhs;
+    if (true) { break; }
+  }
+}
 )");
 }
 
@@ -167,18 +190,15 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(outer)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while (true) {
-    __asm__("");
-    float lhs = 2.5f;
-    float other = 0.0f;
-    break;
-    {
-      lhs = rhs;
-    }
+    EXPECT_EQ(gen.Result(), R"(while (tint_preserve_loop) {
+  float lhs = 2.5f;
+  float other = 0.0f;
+  break;
+  {
+    lhs = rhs;
   }
+}
 )");
 }
 
@@ -193,13 +213,10 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  for(; ; ) {
-    __asm__("");
-    return;
-  }
+    EXPECT_EQ(gen.Result(), R"(for(; tint_preserve_loop; ) {
+  return;
+}
 )");
 }
 
@@ -214,13 +231,10 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  for(int i = 0; ; ) {
-    __asm__("");
-    return;
-  }
+    EXPECT_EQ(gen.Result(), R"(for(int i = 0; tint_preserve_loop; ) {
+  return;
+}
 )");
 }
 
@@ -243,19 +257,16 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(loop)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  {
-    {
-      f(1);
-      f(2);
-    }
-    for(; ; ) {
-      __asm__("");
-      return;
-    }
+    EXPECT_EQ(gen.Result(), R"({
+  {
+    f(1);
+    f(2);
   }
+  for(; tint_preserve_loop; ) {
+    return;
+  }
+}
 )");
 }
 
@@ -270,13 +281,10 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  for(; true; ) {
-    __asm__("");
-    return;
-  }
+    EXPECT_EQ(gen.Result(), R"(for(; (true) == tint_preserve_loop; ) {
+  return;
+}
 )");
 }
 
@@ -292,14 +300,12 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(),
-              R"(  for(; ; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
-    return;
-  }
+    EXPECT_EQ(
+        gen.Result(),
+        R"(for(; tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  return;
+}
 )");
 }
 
@@ -322,17 +328,14 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(loop)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while (true) {
-    __asm__("");
-    return;
-    {
-      f(1);
-      f(2);
-    }
+    EXPECT_EQ(gen.Result(), R"(while (tint_preserve_loop) {
+  return;
+  {
+    f(1);
+    f(2);
   }
+}
 )");
 }
 
@@ -349,14 +352,12 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(),
-              R"(  for(int i = 0; true; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
-    a_statement();
-  }
+    EXPECT_EQ(
+        gen.Result(),
+        R"(for(int i = 0; (true) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  a_statement();
+}
 )");
 }
 
@@ -380,24 +381,21 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(loop)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  {
+    EXPECT_EQ(gen.Result(), R"({
+  {
+    f(1);
+    f(2);
+  }
+  while (tint_preserve_loop) {
+    if (!(true)) { break; }
+    return;
     {
-      f(1);
-      f(2);
-    }
-    while (true) {
-      __asm__("");
-      if (!(true)) { break; }
-      return;
-      {
-        f(3);
-        f(4);
-      }
+      f(3);
+      f(4);
     }
   }
+}
 )");
 }
 
@@ -411,13 +409,10 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while(true) {
-    __asm__("");
-    return;
-  }
+    EXPECT_EQ(gen.Result(), R"(while((true) == tint_preserve_loop) {
+  return;
+}
 )");
 }
 
@@ -431,13 +426,10 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while(true) {
-    __asm__("");
-    continue;
-  }
+    EXPECT_EQ(gen.Result(), R"(while((true) == tint_preserve_loop) {
+  continue;
+}
 )");
 }
 
@@ -454,13 +446,10 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while((t && false)) {
-    __asm__("");
-    return;
-  }
+    EXPECT_EQ(gen.Result(), R"(while(((t && false)) == tint_preserve_loop) {
+  return;
+}
 )");
 }
 
diff --git a/test/tint/array/assign_to_function_var.wgsl.expected.msl b/test/tint/array/assign_to_function_var.wgsl.expected.msl
index 6f9b5bb..794e774 100644
--- a/test/tint/array/assign_to_function_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_function_var.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   tint_array<int4, 4> src_private;
 };
@@ -54,8 +56,7 @@
 }
 
 void tint_symbol_inner(uint local_invocation_index, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_8, const constant S* const tint_symbol_9, device S* const tint_symbol_10) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol_8))[i] = int4(0);
   }
diff --git a/test/tint/array/assign_to_private_var.wgsl.expected.msl b/test/tint/array/assign_to_private_var.wgsl.expected.msl
index 4046b9e..bf83fb8 100644
--- a/test/tint/array/assign_to_private_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_private_var.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   tint_array<int4, 4> src_private;
   tint_array<int4, 4> dst;
@@ -54,8 +56,7 @@
 }
 
 void tint_symbol_inner(uint local_invocation_index, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_8, const constant S* const tint_symbol_9, device S* const tint_symbol_10) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol_8))[i] = int4(0);
   }
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.msl b/test/tint/array/assign_to_storage_var.wgsl.expected.msl
index 94abf80..accc11a 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   tint_array<int4, 4> src_private;
 };
@@ -56,8 +58,7 @@
 }
 
 void tint_symbol_inner(uint local_invocation_index, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_10, device S* const tint_symbol_11, const constant S* const tint_symbol_12, device S* const tint_symbol_13, device S_nested* const tint_symbol_14) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol_10))[i] = int4(0);
   }
diff --git a/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl b/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl
index d644a06..3b0ae7e 100644
--- a/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl
+++ b/test/tint/array/assign_to_workgroup_var.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   tint_array<int4, 4> src_private;
 };
@@ -52,14 +54,12 @@
 }
 
 void tint_symbol_inner(uint local_invocation_index, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<int4, 4>* const tint_symbol_10, threadgroup tint_array<int4, 4>* const tint_symbol_11, threadgroup tint_array<tint_array<tint_array<int, 2>, 3>, 4>* const tint_symbol_12, const constant S* const tint_symbol_13, device S* const tint_symbol_14) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol_10))[i] = int4(0);
     (*(tint_symbol_11))[i] = int4(0);
   }
-  for(uint idx_1 = local_invocation_index; (idx_1 < 24u); idx_1 = (idx_1 + 1u)) {
-    __asm__("");
+  for(uint idx_1 = local_invocation_index; ((idx_1 < 24u)) == tint_preserve_loop; idx_1 = (idx_1 + 1u)) {
     uint const i_1 = (idx_1 / 6u);
     uint const i_2 = ((idx_1 % 6u) / 2u);
     uint const i_3 = (idx_1 % 2u);
diff --git a/test/tint/array/strides.spvasm.expected.msl b/test/tint/array/strides.spvasm.expected.msl
index 38ff24a..f67348f 100644
--- a/test/tint/array/strides.spvasm.expected.msl
+++ b/test/tint/array/strides.spvasm.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct strided_arr {
   /* 0x0000 */ float el;
   /* 0x0004 */ tint_array<int8_t, 4> tint_pad;
@@ -33,15 +35,13 @@
 }
 
 void assign_and_preserve_padding_3(device tint_array<strided_arr, 2>* const dest, tint_array<strided_arr, 2> value) {
-  for(uint i = 0u; (i < 2u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 2u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_4(&((*(dest))[i]), value[i]);
   }
 }
 
 void assign_and_preserve_padding_2(device tint_array<tint_array<strided_arr, 2>, 3>* const dest, tint_array<tint_array<strided_arr, 2>, 3> value) {
-  for(uint i = 0u; (i < 3u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 3u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_3(&((*(dest))[i]), value[i]);
   }
 }
@@ -51,8 +51,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<strided_arr_1, 4>* const dest, tint_array<strided_arr_1, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl
index 31962aa..fc97070 100644
--- a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl
+++ b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
   /* 0x000c */ tint_array<int8_t, 4> tint_pad;
@@ -102,8 +104,7 @@
 }
 
 void assign_and_preserve_padding_3(device tint_array<tint_packed_vec3_f32_array_element, 2>* const dest, tint_array<float3, 2> value) {
-  for(uint i = 0u; (i < 2u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 2u)) == tint_preserve_loop; i = (i + 1u)) {
     (*(dest))[i].elements = packed_float3(value[i]);
   }
 }
diff --git a/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.msl b/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.msl
index a984832..204d507 100644
--- a/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.msl
+++ b/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
   /* 0x000c */ tint_array<int8_t, 4> tint_pad;
@@ -157,8 +159,7 @@
 }
 
 void assign_and_preserve_padding_6(device tint_array<tint_packed_vec3_f32_array_element, 2>* const dest, tint_array<float3, 2> value) {
-  for(uint i = 0u; (i < 2u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 2u)) == tint_preserve_loop; i = (i + 1u)) {
     (*(dest))[i].elements = packed_float3(value[i]);
   }
 }
diff --git a/test/tint/buffer/storage/static_index/write.wgsl.expected.msl b/test/tint/buffer/storage/static_index/write.wgsl.expected.msl
index 93c261e..705ddf0 100644
--- a/test/tint/buffer/storage/static_index/write.wgsl.expected.msl
+++ b/test/tint/buffer/storage/static_index/write.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
   /* 0x000c */ tint_array<int8_t, 4> tint_pad;
@@ -104,8 +106,7 @@
 }
 
 void assign_and_preserve_padding_3(device tint_array<tint_packed_vec3_f32_array_element, 2>* const dest, tint_array<float3, 2> value) {
-  for(uint i = 0u; (i < 2u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 2u)) == tint_preserve_loop; i = (i + 1u)) {
     (*(dest))[i].elements = packed_float3(value[i]);
   }
 }
diff --git a/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.msl b/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.msl
index 6141100..291a4c5 100644
--- a/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.msl
+++ b/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
   /* 0x000c */ tint_array<int8_t, 4> tint_pad;
@@ -161,8 +163,7 @@
 }
 
 void assign_and_preserve_padding_6(device tint_array<tint_packed_vec3_f32_array_element, 2>* const dest, tint_array<float3, 2> value) {
-  for(uint i = 0u; (i < 2u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 2u)) == tint_preserve_loop; i = (i + 1u)) {
     (*(dest))[i].elements = packed_float3(value[i]);
   }
 }
@@ -174,8 +175,7 @@
 }
 
 void assign_and_preserve_padding_8(device tint_array<Inner, 4>* const dest, tint_array<Inner, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_7(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.msl
index f431584..be62ea6 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.msl
@@ -14,13 +14,14 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_5 {
   tint_array<float2x2, 4> w;
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<float2x2, 4>* const tint_symbol, const constant tint_array<float2x2, 4>* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i] = float2x2(float2(0.0f), float2(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl
index 87a9625..14f6a5b 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f16_array_element {
   /* 0x0000 */ packed_half3 elements;
   /* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@@ -35,8 +37,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* const dest, tint_array<half2x3, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.msl
index 3543956..26033d9 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_7 {
   tint_array<half2x3, 4> w;
 };
@@ -34,8 +36,7 @@
 }
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<half2x3, 4>* const tint_symbol, const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 2>, 4>* const tint_symbol_1, device half* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i] = half2x3(half3(0.0h), half3(0.0h));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl
index 9bfa849..89464a7 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
   /* 0x000c */ tint_array<int8_t, 4> tint_pad;
@@ -35,8 +37,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* const dest, tint_array<float2x3, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.msl
index 0b49f03..fce938a 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_5 {
   tint_array<float2x3, 4> w;
 };
@@ -34,8 +36,7 @@
 }
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<float2x3, 4>* const tint_symbol, const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 2>, 4>* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i] = float2x3(float3(0.0f), float3(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.msl
index ec6ff10..e9ffa1c 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.msl
@@ -14,13 +14,14 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_5 {
   tint_array<half2x4, 4> w;
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<half2x4, 4>* const tint_symbol, const constant tint_array<half2x4, 4>* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i] = half2x4(half4(0.0h), half4(0.0h));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.msl
index ab18a13..1c75e49 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.msl
@@ -14,13 +14,14 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_5 {
   tint_array<float2x4, 4> w;
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<float2x4, 4>* const tint_symbol, const constant tint_array<float2x4, 4>* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i] = float2x4(float4(0.0f), float4(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl
index 2efe350..6699025 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
   /* 0x000c */ tint_array<int8_t, 4> tint_pad;
@@ -36,8 +38,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* const dest, tint_array<float3x3, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.msl
index 4eae64c..13d0e9f 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_5 {
   tint_array<float3x3, 4> w;
 };
@@ -34,8 +36,7 @@
 }
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<float3x3, 4>* const tint_symbol, const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 4>* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i] = float3x3(float3(0.0f), float3(0.0f), float3(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.msl
index 3a2f692..e1bf4de 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.msl
@@ -14,13 +14,14 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_5 {
   tint_array<float3x4, 4> w;
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<float3x4, 4>* const tint_symbol, const constant tint_array<float3x4, 4>* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i] = float3x4(float4(0.0f), float4(0.0f), float4(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.msl
index 527f902..9d0cb69 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.msl
@@ -14,13 +14,14 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_5 {
   tint_array<half4x2, 4> w;
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<half4x2, 4>* const tint_symbol, const constant tint_array<half4x2, 4>* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i] = half4x2(half2(0.0h), half2(0.0h), half2(0.0h), half2(0.0h));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.msl
index e523c29..34a353f 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.msl
@@ -14,13 +14,14 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_5 {
   tint_array<float4x2, 4> w;
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<float4x2, 4>* const tint_symbol, const constant tint_array<float4x2, 4>* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i] = float4x2(float2(0.0f), float2(0.0f), float2(0.0f), float2(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl
index 7ae0b35..b9ae76b 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f16_array_element {
   /* 0x0000 */ packed_half3 elements;
   /* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@@ -37,8 +39,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* const dest, tint_array<half4x3, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.msl
index ade6756..e4983aa 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_5 {
   tint_array<half4x3, 4> w;
 };
@@ -34,8 +36,7 @@
 }
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<half4x3, 4>* const tint_symbol, const constant tint_array<tint_array<tint_packed_vec3_f16_array_element, 4>, 4>* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i] = half4x3(half3(0.0h), half3(0.0h), half3(0.0h), half3(0.0h));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl
index 9b88b8f..a13be34 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
   /* 0x000c */ tint_array<int8_t, 4> tint_pad;
@@ -37,8 +39,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* const dest, tint_array<float4x3, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.msl
index d06150d..d6d36ab 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_5 {
   tint_array<float4x3, 4> w;
 };
@@ -34,8 +36,7 @@
 }
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<float4x3, 4>* const tint_symbol, const constant tint_array<tint_array<tint_packed_vec3_f32_array_element, 4>, 4>* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i] = float4x3(float3(0.0f), float3(0.0f), float3(0.0f), float3(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.msl
index 79d94e5..3b9b765 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.msl
@@ -14,13 +14,14 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_5 {
   tint_array<half4x4, 4> w;
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<half4x4, 4>* const tint_symbol, const constant tint_array<half4x4, 4>* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i] = half4x4(half4(0.0h), half4(0.0h), half4(0.0h), half4(0.0h));
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.msl
index 78fac42..248ef96 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.msl
@@ -14,13 +14,14 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_5 {
   tint_array<float4x4, 4> w;
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<float4x4, 4>* const tint_symbol, const constant tint_array<float4x4, 4>* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i] = float4x4(float4(0.0f), float4(0.0f), float4(0.0f), float4(0.0f));
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl
index 5d62b61..57455e1 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ half2x2 m;
@@ -29,8 +31,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.msl
index fb0567d..dce4eac 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ half2x2 m;
@@ -27,8 +29,7 @@
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl
index e63293d..d4f859f 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 4> tint_pad;
@@ -30,8 +32,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.msl
index d2a1def..5575d59 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 4> tint_pad;
@@ -28,8 +30,7 @@
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl
index f4659e2..8866890 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f16_array_element {
   /* 0x0000 */ packed_half3 elements;
   /* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@@ -64,8 +66,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.msl
index 3da4e4d..a282ff6 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   int before;
   half2x3 m;
@@ -57,8 +59,7 @@
 }
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S_tint_packed_vec3, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl
index 73a8802..273e906 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
   /* 0x000c */ tint_array<int8_t, 4> tint_pad;
@@ -64,8 +66,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.msl
index bf35778..e1b2d71 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   int before;
   float2x3 m;
@@ -57,8 +59,7 @@
 }
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S_tint_packed_vec3, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl
index 28a9ca2..934101a 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 4> tint_pad;
@@ -30,8 +32,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.msl
index 7607f1a..70dfb67 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 4> tint_pad;
@@ -28,8 +30,7 @@
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl
index 7aac521..10d60c6 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 12> tint_pad;
@@ -30,8 +32,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.msl
index e93e10a..93a2642 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 12> tint_pad;
@@ -28,8 +30,7 @@
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl
index 312ef3a..8f4151b 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ half3x2 m;
@@ -29,8 +31,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.msl
index 62fba42..a83fab0 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ half3x2 m;
@@ -27,8 +29,7 @@
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl
index cc8345c..a15a0bd 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 4> tint_pad;
@@ -30,8 +32,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.msl
index fd7fd71..786d13e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 4> tint_pad;
@@ -28,8 +30,7 @@
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl
index e40b77b..da1c72f 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f16_array_element {
   /* 0x0000 */ packed_half3 elements;
   /* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@@ -65,8 +67,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.msl
index 4a381a8..53c0ed3 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   int before;
   half3x3 m;
@@ -57,8 +59,7 @@
 }
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S_tint_packed_vec3, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl
index 92efeb1..18cf65e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
   /* 0x000c */ tint_array<int8_t, 4> tint_pad;
@@ -64,8 +66,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.msl
index 9f4531c..7c3c470 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   int before;
   float3x3 m;
@@ -56,8 +58,7 @@
 }
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S_tint_packed_vec3, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl
index ce88e3e..29eae07 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 4> tint_pad;
@@ -30,8 +32,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.msl
index c9e815c..0ab8876 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 4> tint_pad;
@@ -28,8 +30,7 @@
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl
index 663a39a..4e9f71d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 12> tint_pad;
@@ -29,8 +31,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.msl
index 4ebd2e0..ae7c285 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 12> tint_pad;
@@ -27,8 +29,7 @@
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl
index 4838775..500386d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ half4x2 m;
@@ -29,8 +31,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.msl
index ad88575..767775c 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ half4x2 m;
@@ -27,8 +29,7 @@
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl
index e71864b..cd0013d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 4> tint_pad;
@@ -30,8 +32,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.msl
index 37b7b89..ce195f0 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 4> tint_pad;
@@ -28,8 +30,7 @@
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl
index 2cecf05..4fc1cf1 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f16_array_element {
   /* 0x0000 */ packed_half3 elements;
   /* 0x0006 */ tint_array<int8_t, 2> tint_pad;
@@ -66,8 +68,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.msl
index ce4b9c8..c53df85 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   int before;
   half4x3 m;
@@ -57,8 +59,7 @@
 }
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S_tint_packed_vec3, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl
index d59b971..f09b05a 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
   /* 0x000c */ tint_array<int8_t, 4> tint_pad;
@@ -66,8 +68,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S_tint_packed_vec3, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.msl
index 60449fe..857b0bb 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   int before;
   float4x3 m;
@@ -57,8 +59,7 @@
 }
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S_tint_packed_vec3, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl
index 8223731..121388d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 4> tint_pad;
@@ -30,8 +32,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.msl
index d43d742..8a4308c 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 4> tint_pad;
@@ -28,8 +30,7 @@
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl
index 54e8326..a97b159 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 12> tint_pad;
@@ -30,8 +32,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<S, 4>* const dest, tint_array<S, 4> value) {
-  for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 4u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.msl
index efb661e..92d0210 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   /* 0x0000 */ int before;
   /* 0x0004 */ tint_array<int8_t, 12> tint_pad;
@@ -28,8 +30,7 @@
 };
 
 void f_inner(uint local_invocation_index, threadgroup tint_array<S, 4>* const tint_symbol_1, const constant tint_array<S, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     S const tint_symbol = S{};
     (*(tint_symbol_1))[i] = tint_symbol;
diff --git a/test/tint/bug/chromium/1403752.wgsl.expected.msl b/test/tint/bug/chromium/1403752.wgsl.expected.msl
index 4ae1c69..22558cd 100644
--- a/test/tint/bug/chromium/1403752.wgsl.expected.msl
+++ b/test/tint/bug/chromium/1403752.wgsl.expected.msl
@@ -1,10 +1,12 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void d() {
   int j = 0;
-  for(; false; ) {
-    __asm__("");
+  for(; (false) == tint_preserve_loop; ) {
   }
 }
 
diff --git a/test/tint/bug/chromium/1449538.wgsl.expected.msl b/test/tint/bug/chromium/1449538.wgsl.expected.msl
index e80c545..ed8df03 100644
--- a/test/tint/bug/chromium/1449538.wgsl.expected.msl
+++ b/test/tint/bug/chromium/1449538.wgsl.expected.msl
@@ -1,30 +1,25 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void f() {
-  for(int i0520 = 0; false; ) {
-    __asm__("");
+  for(int i0520 = 0; (false) == tint_preserve_loop; ) {
   }
-  for(int i62 = 0; false; ) {
-    __asm__("");
+  for(int i62 = 0; (false) == tint_preserve_loop; ) {
   }
-  for(int i0520 = 0; false; ) {
-    __asm__("");
+  for(int i0520 = 0; (false) == tint_preserve_loop; ) {
   }
-  for(int i62 = 0; false; ) {
-    __asm__("");
+  for(int i62 = 0; (false) == tint_preserve_loop; ) {
   }
-  for(int i62 = 0; false; ) {
-    __asm__("");
+  for(int i62 = 0; (false) == tint_preserve_loop; ) {
   }
-  for(int i60 = 0; false; ) {
-    __asm__("");
+  for(int i60 = 0; (false) == tint_preserve_loop; ) {
   }
-  for(int i62 = 0; false; ) {
-    __asm__("");
+  for(int i62 = 0; (false) == tint_preserve_loop; ) {
   }
-  for(int i60 = 0; false; ) {
-    __asm__("");
+  for(int i60 = 0; (false) == tint_preserve_loop; ) {
   }
 }
 
diff --git a/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl b/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl
index 28c36c5..c934bb2 100644
--- a/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct UBO {
   /* 0x0000 */ int dynamic_idx;
 };
@@ -27,8 +29,7 @@
 };
 
 void f_inner(uint local_invocation_index, threadgroup S* const tint_symbol, device Result* const tint_symbol_1, const constant UBO* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 64u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 64u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol)).data[i] = 0;
   }
diff --git a/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl b/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl
index 8c88e24..1c3c26d 100644
--- a/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct UBO {
   /* 0x0000 */ int dynamic_idx;
 };
@@ -27,8 +29,7 @@
 };
 
 void f_inner(uint local_invocation_index, threadgroup S* const tint_symbol, const constant UBO* const tint_symbol_1, device Result* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 64u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 64u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol)).data[i] = 0;
   }
diff --git a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl
index 3987e5d..e64a323 100644
--- a/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl
+++ b/test/tint/bug/fxc/gradient_in_varying_loop/1112.wgsl.expected.msl
@@ -1,6 +1,9 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_2 {
   float2 vUV [[user(locn0)]];
 };
@@ -12,8 +15,7 @@
 float4 tint_symbol_inner(float2 vUV, texture2d<float, access::sample> tint_symbol_4, sampler tint_symbol_5) {
   float3 const random = tint_symbol_4.sample(tint_symbol_5, vUV).rgb;
   int i = 0;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     if ((i < 1)) {
     } else {
       break;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
index e81bc5f..89f7f51 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
@@ -1,6 +1,9 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   float2 v2f;
   int3 v3i;
@@ -9,8 +12,7 @@
 };
 
 void foo(thread tint_private_vars_struct* const tint_private_vars) {
-  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     (*(tint_private_vars)).v2f[i] = 1.0f;
     (*(tint_private_vars)).v3i[i] = 1;
     (*(tint_private_vars)).v4u[i] = 1u;
@@ -20,8 +22,7 @@
 
 kernel void tint_symbol() {
   thread tint_private_vars_struct tint_private_vars = {};
-  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     foo(&(tint_private_vars));
   }
   return;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
index 9d73779..046f8c4 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
@@ -1,6 +1,9 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   float2 v2f;
   int3 v3i;
@@ -18,8 +21,7 @@
 
 kernel void tint_symbol() {
   thread tint_private_vars_struct tint_private_vars = {};
-  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     foo(&(tint_private_vars));
   }
   return;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
index 664550c..c62bf40 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
@@ -1,6 +1,9 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 kernel void tint_symbol() {
   float2 v2f = 0.0f;
   float3 v3f = 0.0f;
@@ -14,8 +17,7 @@
   bool2 v2b = false;
   bool3 v3b = false;
   bool4 v4b = false;
-  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     v2f[i] = 1.0f;
     v3f[i] = 1.0f;
     v4f[i] = 1.0f;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
index 2d918b9..847e8f6 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
@@ -1,6 +1,9 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 kernel void tint_symbol() {
   float2 v2f = 0.0f;
   float2 v2f_2 = 0.0f;
@@ -10,8 +13,7 @@
   uint4 v4u_2 = 0u;
   bool2 v2b = false;
   bool2 v2b_2 = false;
-  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     v2f[i] = 1.0f;
     v3i[i] = 1;
     v4u[i] = 1u;
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
index 5425ca7..91ac2be 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
@@ -1,6 +1,9 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 kernel void tint_symbol() {
   float2 v2f = 0.0f;
   float3 v3f = 0.0f;
@@ -14,8 +17,7 @@
   bool2 v2b = false;
   bool3 v3b = false;
   bool4 v4b = false;
-  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     v2f[i] = 1.0f;
     v2i[i] = 1;
     v2u[i] = 1u;
diff --git a/test/tint/bug/tint/1064.wgsl.expected.msl b/test/tint/bug/tint/1064.wgsl.expected.msl
index 454d69c..7b81e57 100644
--- a/test/tint/bug/tint/1064.wgsl.expected.msl
+++ b/test/tint/bug/tint/1064.wgsl.expected.msl
@@ -1,9 +1,11 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 fragment void tint_symbol() {
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     if (false) {
     } else {
       break;
diff --git a/test/tint/bug/tint/1081.wgsl.expected.msl b/test/tint/bug/tint/1081.wgsl.expected.msl
index c1181ce..8b120a8 100644
--- a/test/tint/bug/tint/1081.wgsl.expected.msl
+++ b/test/tint/bug/tint/1081.wgsl.expected.msl
@@ -1,6 +1,9 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   bool tint_discarded;
 };
@@ -22,8 +25,7 @@
 
 int tint_symbol_inner(int3 x, thread tint_private_vars_struct* const tint_private_vars) {
   int y = x[0];
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     int const r = f(y, tint_private_vars);
     if ((r == 0)) {
       break;
diff --git a/test/tint/bug/tint/1121.wgsl.expected.msl b/test/tint/bug/tint/1121.wgsl.expected.msl
index f0db93a..2abb7d5 100644
--- a/test/tint/bug/tint/1121.wgsl.expected.msl
+++ b/test/tint/bug/tint/1121.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct LightData_tint_packed_vec3 {
   /* 0x0000 */ float4 position;
   /* 0x0010 */ packed_float3 color;
@@ -84,10 +86,8 @@
   int const TILE_SIZE = 16;
   int const TILE_COUNT_X = 2;
   int const TILE_COUNT_Y = 2;
-  for(int y = 0; (y < TILE_COUNT_Y); y = as_type<int>((as_type<uint>(y) + as_type<uint>(1)))) {
-    __asm__("");
-    for(int x = 0; (x < TILE_COUNT_X); x = as_type<int>((as_type<uint>(x) + as_type<uint>(1)))) {
-      __asm__("");
+  for(int y = 0; ((y < TILE_COUNT_Y)) == tint_preserve_loop; y = as_type<int>((as_type<uint>(y) + as_type<uint>(1)))) {
+    for(int x = 0; ((x < TILE_COUNT_X)) == tint_preserve_loop; x = as_type<int>((as_type<uint>(x) + as_type<uint>(1)))) {
       int2 tilePixel0Idx = int2(as_type<int>((as_type<uint>(x) * as_type<uint>(TILE_SIZE))), as_type<int>((as_type<uint>(y) * as_type<uint>(TILE_SIZE))));
       float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / (*(tint_symbol_3)).fullScreenSize.xy) - float2(1.0f));
       float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(tilePixel0Idx) + as_type<uint2>(int2(TILE_SIZE)))))) / (*(tint_symbol_3)).fullScreenSize.xy) - float2(1.0f));
@@ -98,8 +98,7 @@
       frustumPlanes[2] = float4(0.0f, 1.0f, (-(viewFloorCoord[1]) / viewNear), 0.0f);
       frustumPlanes[3] = float4(0.0f, -1.0f, (viewCeilCoord[1] / viewNear), 0.0f);
       float dp = 0.0f;
-      for(uint i = 0u; (i < 6u); i = (i + 1u)) {
-        __asm__("");
+      for(uint i = 0u; ((i < 6u)) == tint_preserve_loop; i = (i + 1u)) {
         float4 p = 0.0f;
         if ((frustumPlanes[i][0] > 0.0f)) {
           p[0] = boxMax[0];
diff --git a/test/tint/bug/tint/1321.wgsl.expected.msl b/test/tint/bug/tint/1321.wgsl.expected.msl
index 7d3878a..99f3002 100644
--- a/test/tint/bug/tint/1321.wgsl.expected.msl
+++ b/test/tint/bug/tint/1321.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 int foo() {
   return 1;
 }
@@ -23,8 +25,7 @@
   {
     int const tint_symbol_1 = foo();
     int const a_save = tint_symbol_1;
-    while (true) {
-      __asm__("");
+    while (tint_preserve_loop) {
       {
         float const x = arr[a_save];
         break;
diff --git a/test/tint/bug/tint/1474-a.wgsl.expected.msl b/test/tint/bug/tint/1474-a.wgsl.expected.msl
index 4cb74c7..02c5b90 100644
--- a/test/tint/bug/tint/1474-a.wgsl.expected.msl
+++ b/test/tint/bug/tint/1474-a.wgsl.expected.msl
@@ -1,9 +1,11 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 kernel void tint_symbol() {
-  while(true) {
-    __asm__("");
+  while((true) == tint_preserve_loop) {
     if (true) {
       break;
     } else {
diff --git a/test/tint/bug/tint/1538.wgsl.expected.msl b/test/tint/bug/tint/1538.wgsl.expected.msl
index 6e5e6e1..ea91f6c 100644
--- a/test/tint/bug/tint/1538.wgsl.expected.msl
+++ b/test/tint/bug/tint/1538.wgsl.expected.msl
@@ -2,6 +2,8 @@
 
 using namespace metal;
 
+constant static volatile bool tint_preserve_loop = true;
+
 template<typename T, size_t N>
 struct tint_array {
     const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -19,8 +21,7 @@
 }
 
 int f() {
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     g();
     break;
   }
@@ -29,8 +30,7 @@
 }
 
 kernel void tint_symbol(device tint_array<uint, 1>* tint_symbol_1 [[buffer(0)]]) {
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     if (((*(tint_symbol_1))[0] == 0u)) {
       break;
     }
diff --git a/test/tint/bug/tint/1557.wgsl.expected.msl b/test/tint/bug/tint/1557.wgsl.expected.msl
index e61acec..9952e88 100644
--- a/test/tint/bug/tint/1557.wgsl.expected.msl
+++ b/test/tint/bug/tint/1557.wgsl.expected.msl
@@ -1,14 +1,16 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 int f() {
   return 0;
 }
 
 void g() {
   int j = 0;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     if ((j >= 1)) {
       break;
     }
diff --git a/test/tint/bug/tint/1604.wgsl.expected.msl b/test/tint/bug/tint/1604.wgsl.expected.msl
index d2f041d..2dfebe4 100644
--- a/test/tint/bug/tint/1604.wgsl.expected.msl
+++ b/test/tint/bug/tint/1604.wgsl.expected.msl
@@ -1,11 +1,13 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 kernel void tint_symbol(const constant int* tint_symbol_1 [[buffer(0)]]) {
   switch(*(tint_symbol_1)) {
     case 0: {
-      while (true) {
-        __asm__("");
+      while (tint_preserve_loop) {
         return;
       }
       break;
diff --git a/test/tint/bug/tint/1605.wgsl.expected.msl b/test/tint/bug/tint/1605.wgsl.expected.msl
index e4a6e0b..78953c3 100644
--- a/test/tint/bug/tint/1605.wgsl.expected.msl
+++ b/test/tint/bug/tint/1605.wgsl.expected.msl
@@ -1,11 +1,12 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 bool func_3(const constant int* const tint_symbol_1) {
-  for(int i = 0; (i < *(tint_symbol_1)); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
-    for(int j = -1; (j == 1); j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)))) {
-      __asm__("");
+  for(int i = 0; ((i < *(tint_symbol_1))) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    for(int j = -1; ((j == 1)) == tint_preserve_loop; j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)))) {
       return false;
     }
   }
diff --git a/test/tint/bug/tint/1764.wgsl.expected.msl b/test/tint/bug/tint/1764.wgsl.expected.msl
index b1cc428..f4dcf7f 100644
--- a/test/tint/bug/tint/1764.wgsl.expected.msl
+++ b/test/tint/bug/tint/1764.wgsl.expected.msl
@@ -14,9 +14,10 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 void tint_symbol_inner(uint local_invocation_index, threadgroup tint_array<int, 246>* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 246u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 246u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol_1))[i] = 0;
   }
diff --git a/test/tint/bug/tint/2010.spvasm.expected.msl b/test/tint/bug/tint/2010.spvasm.expected.msl
index f46cd3a..061aa4e 100644
--- a/test/tint/bug/tint/2010.spvasm.expected.msl
+++ b/test/tint/bug/tint/2010.spvasm.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   uint3 x_3;
 };
@@ -46,8 +48,7 @@
   uint x_88 = 0u;
   uint const x_52 = (*(tint_private_vars)).x_3[0];
   x_54 = 0u;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     uint x_55 = 0u;
     x_58 = (*(tint_symbol_3)).field0.field0;
     if ((x_54 < x_58)) {
@@ -79,8 +80,7 @@
   }
   x_85 = x_76.xyxy;
   x_88 = 1u;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     float4 x_111 = 0.0f;
     float4 x_86 = 0.0f;
     uint x_89 = 0u;
@@ -130,8 +130,7 @@
     atomic_store_explicit(tint_symbol_13, 0u, memory_order_relaxed);
     atomic_store_explicit(tint_symbol_14, 0u, memory_order_relaxed);
   }
-  for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 32u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4096u)) == tint_preserve_loop; idx = (idx + 32u)) {
     uint const i = idx;
     S const tint_symbol_1 = S{};
     (*(tint_symbol_15))[i] = tint_symbol_1;
diff --git a/test/tint/bug/tint/2059.wgsl.expected.msl b/test/tint/bug/tint/2059.wgsl.expected.msl
index 9bc2e44..097729c 100644
--- a/test/tint/bug/tint/2059.wgsl.expected.msl
+++ b/test/tint/bug/tint/2059.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_packed_vec3_f32_array_element {
   /* 0x0000 */ packed_float3 elements;
   /* 0x000c */ tint_array<int8_t, 4> tint_pad;
@@ -62,8 +64,7 @@
 }
 
 void assign_and_preserve_padding_3(device tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 1>* const dest, tint_array<float3x3, 1> value) {
-  for(uint i = 0u; (i < 1u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 1u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding(&((*(dest))[i]), value[i]);
   }
 }
@@ -77,8 +78,7 @@
 }
 
 void assign_and_preserve_padding_6(device tint_array<S_tint_packed_vec3, 1>* const dest, tint_array<S, 1> value) {
-  for(uint i = 0u; (i < 1u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 1u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
@@ -88,16 +88,14 @@
 }
 
 void assign_and_preserve_padding_7(device tint_array<S2_tint_packed_vec3, 1>* const dest, tint_array<S2, 1> value) {
-  for(uint i = 0u; (i < 1u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 1u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_2(&((*(dest))[i]), value[i]);
   }
 }
 
 kernel void tint_symbol(device tint_array<tint_packed_vec3_f32_array_element, 3>* tint_symbol_8 [[buffer(0)]], device S_tint_packed_vec3* tint_symbol_9 [[buffer(1)]], device S2_tint_packed_vec3* tint_symbol_10 [[buffer(2)]], device S3_tint_packed_vec3* tint_symbol_11 [[buffer(3)]], device S4_tint_packed_vec3* tint_symbol_12 [[buffer(4)]], device tint_array<tint_array<tint_packed_vec3_f32_array_element, 3>, 1>* tint_symbol_13 [[buffer(5)]], device tint_array<S_tint_packed_vec3, 1>* tint_symbol_14 [[buffer(6)]], device tint_array<S2_tint_packed_vec3, 1>* tint_symbol_15 [[buffer(7)]]) {
   float3x3 m = float3x3(0.0f);
-  for(uint c = 0u; (c < 3u); c = (c + 1u)) {
-    __asm__("");
+  for(uint c = 0u; ((c < 3u)) == tint_preserve_loop; c = (c + 1u)) {
     m[c] = float3(float(((c * 3u) + 1u)), float(((c * 3u) + 2u)), float(((c * 3u) + 3u)));
   }
   {
diff --git a/test/tint/bug/tint/221.wgsl.expected.msl b/test/tint/bug/tint/221.wgsl.expected.msl
index 83beacb..c14e039 100644
--- a/test/tint/bug/tint/221.wgsl.expected.msl
+++ b/test/tint/bug/tint/221.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct Buf {
   /* 0x0000 */ uint count;
   /* 0x0004 */ tint_array<uint, 50> data;
@@ -25,8 +27,7 @@
 
 kernel void tint_symbol(device Buf* tint_symbol_1 [[buffer(0)]]) {
   uint i = 0u;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     if ((i >= (*(tint_symbol_1)).count)) {
       break;
     }
diff --git a/test/tint/bug/tint/534.wgsl.expected.msl b/test/tint/bug/tint/534.wgsl.expected.msl
index 2b51953..fd3dbbc 100644
--- a/test/tint/bug/tint/534.wgsl.expected.msl
+++ b/test/tint/bug/tint/534.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 uint4 tint_ftou(float4 v) {
   return select(uint4(4294967295u), select(uint4(v), uint4(0u), (v < float4(0.0f))), (v < float4(4294967040.0f)));
 }
@@ -45,8 +47,7 @@
   bool success = true;
   uint4 srcColorBits = 0u;
   uint4 dstColorBits = tint_ftou(dstColor);
-  for(uint i = 0u; (i < (*(tint_symbol_3)).channelCount); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < (*(tint_symbol_3)).channelCount)) == tint_preserve_loop; i = (i + 1u)) {
     uint const tint_symbol_1 = i;
     srcColorBits[tint_symbol_1] = ConvertToFp16FloatValue(srcColor[i]);
     success = (success && (srcColorBits[i] == dstColorBits[i]));
diff --git a/test/tint/bug/tint/744.wgsl.expected.msl b/test/tint/bug/tint/744.wgsl.expected.msl
index d539045..85ae3c4 100644
--- a/test/tint/bug/tint/744.wgsl.expected.msl
+++ b/test/tint/bug/tint/744.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct Uniforms {
   /* 0x0000 */ uint2 aShape;
   /* 0x0008 */ uint2 bShape;
@@ -29,8 +31,7 @@
   uint const dimInner = (*(tint_symbol_1)).aShape[1];
   uint const dimOutter = (*(tint_symbol_1)).outShape[1];
   uint result = 0u;
-  for(uint i = 0u; (i < dimInner); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < dimInner)) == tint_preserve_loop; i = (i + 1u)) {
     uint const a = (i + (resultCell[0] * dimInner));
     uint const b = (resultCell[1] + (i * dimOutter));
     result = (result + ((*(tint_symbol_2)).numbers[a] * (*(tint_symbol_3)).numbers[b]));
diff --git a/test/tint/bug/tint/757.wgsl.expected.msl b/test/tint/bug/tint/757.wgsl.expected.msl
index e6671f0..11910c8 100644
--- a/test/tint/bug/tint/757.wgsl.expected.msl
+++ b/test/tint/bug/tint/757.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct Constants {
   int level;
 };
@@ -26,8 +28,7 @@
   uint flatIndex = (((4u * GlobalInvocationID[2]) + (2u * GlobalInvocationID[1])) + GlobalInvocationID[0]);
   flatIndex = (flatIndex * 1u);
   float4 texel = tint_symbol_1.read(uint2(int2(GlobalInvocationID.xy)), 0, 0);
-  for(uint i = 0u; (i < 1u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 1u)) == tint_preserve_loop; i = (i + 1u)) {
     (*(tint_symbol_2)).values[(flatIndex + i)] = texel[0];
   }
 }
diff --git a/test/tint/bug/tint/914.wgsl.expected.msl b/test/tint/bug/tint/914.wgsl.expected.msl
index adb0758..98d18dd 100644
--- a/test/tint/bug/tint/914.wgsl.expected.msl
+++ b/test/tint/bug/tint/914.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct Uniforms {
   /* 0x0000 */ uint dimAOuter;
   /* 0x0004 */ uint dimInner;
@@ -52,8 +54,7 @@
 }
 
 void tint_symbol_inner(uint3 local_id, uint3 global_id, uint local_invocation_index, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_11, threadgroup tint_array<tint_array<float, 64>, 64>* const tint_symbol_12, const constant Uniforms* const tint_symbol_13, const device Matrix* const tint_symbol_14, const device Matrix* const tint_symbol_15, device Matrix* const tint_symbol_16) {
-  for(uint idx = local_invocation_index; (idx < 4096u); idx = (idx + 256u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4096u)) == tint_preserve_loop; idx = (idx + 256u)) {
     uint const i = (idx / 64u);
     uint const i_1 = (idx % 64u);
     (*(tint_symbol_11))[i][i_1] = 0.0f;
@@ -68,20 +69,16 @@
   tint_array<float, 16> acc = {};
   float ACached = 0.0f;
   tint_array<float, 4> BCached = {};
-  for(uint index = 0u; (index < 16u); index = (index + 1u)) {
-    __asm__("");
+  for(uint index = 0u; ((index < 16u)) == tint_preserve_loop; index = (index + 1u)) {
     acc[index] = 0.0f;
   }
   uint const ColPerThreadA = 4u;
   uint const tileColA = (local_id[0] * ColPerThreadA);
   uint const RowPerThreadB = 4u;
   uint const tileRowB = (local_id[1] * RowPerThreadB);
-  for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
-    __asm__("");
-    for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
-      __asm__("");
-      for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
-        __asm__("");
+  for(uint t = 0u; ((t < numTiles)) == tint_preserve_loop; t = (t + 1u)) {
+    for(uint innerRow = 0u; ((innerRow < 4u)) == tint_preserve_loop; innerRow = (innerRow + 1u)) {
+      for(uint innerCol = 0u; ((innerCol < ColPerThreadA)) == tint_preserve_loop; innerCol = (innerCol + 1u)) {
         uint const inputRow = (tileRow + innerRow);
         uint const inputCol = (tileColA + innerCol);
         uint const tint_symbol_1 = inputRow;
@@ -89,10 +86,8 @@
         (*(tint_symbol_11))[tint_symbol_1][tint_symbol_2] = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol), tint_symbol_13, tint_symbol_14);
       }
     }
-    for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
-      __asm__("");
-      for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
-        __asm__("");
+    for(uint innerRow = 0u; ((innerRow < RowPerThreadB)) == tint_preserve_loop; innerRow = (innerRow + 1u)) {
+      for(uint innerCol = 0u; ((innerCol < 4u)) == tint_preserve_loop; innerCol = (innerCol + 1u)) {
         uint const inputRow = (tileRowB + innerRow);
         uint const inputCol = (tileCol + innerCol);
         uint const tint_symbol_3 = innerCol;
@@ -101,17 +96,13 @@
       }
     }
     threadgroup_barrier(mem_flags::mem_threadgroup);
-    for(uint k = 0u; (k < 64u); k = (k + 1u)) {
-      __asm__("");
-      for(uint inner = 0u; (inner < 4u); inner = (inner + 1u)) {
-        __asm__("");
+    for(uint k = 0u; ((k < 64u)) == tint_preserve_loop; k = (k + 1u)) {
+      for(uint inner = 0u; ((inner < 4u)) == tint_preserve_loop; inner = (inner + 1u)) {
         BCached[inner] = (*(tint_symbol_12))[k][(tileCol + inner)];
       }
-      for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
-        __asm__("");
+      for(uint innerRow = 0u; ((innerRow < 4u)) == tint_preserve_loop; innerRow = (innerRow + 1u)) {
         ACached = (*(tint_symbol_11))[(tileRow + innerRow)][k];
-        for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
-          __asm__("");
+        for(uint innerCol = 0u; ((innerCol < 4u)) == tint_preserve_loop; innerCol = (innerCol + 1u)) {
           uint const index = ((innerRow * 4u) + innerCol);
           acc[index] = (acc[index] + (ACached * BCached[innerCol]));
         }
@@ -119,10 +110,8 @@
     }
     threadgroup_barrier(mem_flags::mem_threadgroup);
   }
-  for(uint innerRow = 0u; (innerRow < 4u); innerRow = (innerRow + 1u)) {
-    __asm__("");
-    for(uint innerCol = 0u; (innerCol < 4u); innerCol = (innerCol + 1u)) {
-      __asm__("");
+  for(uint innerRow = 0u; ((innerRow < 4u)) == tint_preserve_loop; innerRow = (innerRow + 1u)) {
+    for(uint innerCol = 0u; ((innerCol < 4u)) == tint_preserve_loop; innerCol = (innerCol + 1u)) {
       uint const index = ((innerRow * 4u) + innerCol);
       mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index], tint_symbol_13, tint_symbol_16);
     }
diff --git a/test/tint/bug/tint/942.wgsl.expected.msl b/test/tint/bug/tint/942.wgsl.expected.msl
index 81b81f1..ea207c3 100644
--- a/test/tint/bug/tint/942.wgsl.expected.msl
+++ b/test/tint/bug/tint/942.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct Params {
   /* 0x0000 */ uint filterDim;
   /* 0x0004 */ uint blockDim;
@@ -28,8 +30,7 @@
 }
 
 void tint_symbol_inner(uint3 WorkGroupID, uint3 LocalInvocationID, uint local_invocation_index, threadgroup tint_array<tint_array<float3, 256>, 4>* const tint_symbol_1, const constant Params* const tint_symbol_2, texture2d<float, access::sample> tint_symbol_3, const constant Flip* const tint_symbol_4, sampler tint_symbol_5, texture2d<float, access::write> tint_symbol_6) {
-  for(uint idx = local_invocation_index; (idx < 1024u); idx = (idx + 64u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 1024u)) == tint_preserve_loop; idx = (idx + 64u)) {
     uint const i_1 = (idx / 256u);
     uint const i_2 = (idx % 256u);
     (*(tint_symbol_1))[i_1][i_2] = float3(0.0f);
@@ -38,10 +39,8 @@
   uint const filterOffset = tint_div(((*(tint_symbol_2)).filterDim - 1u), 2u);
   uint2 const dims = uint2(tint_symbol_3.get_width(0), tint_symbol_3.get_height(0));
   uint2 const baseIndex = (((WorkGroupID.xy * uint2((*(tint_symbol_2)).blockDim, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u))) - uint2(filterOffset, 0u));
-  for(uint r = 0u; (r < 4u); r = (r + 1u)) {
-    __asm__("");
-    for(uint c = 0u; (c < 4u); c = (c + 1u)) {
-      __asm__("");
+  for(uint r = 0u; ((r < 4u)) == tint_preserve_loop; r = (r + 1u)) {
+    for(uint c = 0u; ((c < 4u)) == tint_preserve_loop; c = (c + 1u)) {
       uint2 loadIndex = (baseIndex + uint2(c, r));
       if (((*(tint_symbol_4)).value != 0u)) {
         loadIndex = loadIndex.yx;
@@ -50,10 +49,8 @@
     }
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  for(uint r = 0u; (r < 4u); r = (r + 1u)) {
-    __asm__("");
-    for(uint c = 0u; (c < 4u); c = (c + 1u)) {
-      __asm__("");
+  for(uint r = 0u; ((r < 4u)) == tint_preserve_loop; r = (r + 1u)) {
+    for(uint c = 0u; ((c < 4u)) == tint_preserve_loop; c = (c + 1u)) {
       uint2 writeIndex = (baseIndex + uint2(c, r));
       if (((*(tint_symbol_4)).value != 0u)) {
         writeIndex = writeIndex.yx;
@@ -61,8 +58,7 @@
       uint const center = ((4u * LocalInvocationID[0]) + c);
       if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
         float3 acc = float3(0.0f);
-        for(uint f = 0u; (f < (*(tint_symbol_2)).filterDim); f = (f + 1u)) {
-          __asm__("");
+        for(uint f = 0u; ((f < (*(tint_symbol_2)).filterDim)) == tint_preserve_loop; f = (f + 1u)) {
           uint i = ((center + f) - filterOffset);
           acc = (acc + ((1.0f / float((*(tint_symbol_2)).filterDim)) * (*(tint_symbol_1))[r][i]));
         }
diff --git a/test/tint/bug/tint/948.wgsl.expected.msl b/test/tint/bug/tint/948.wgsl.expected.msl
index 57ef1a9..510b303 100644
--- a/test/tint/bug/tint/948.wgsl.expected.msl
+++ b/test/tint/bug/tint/948.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   float2 tUV;
   float mt;
@@ -98,8 +100,7 @@
   float2 const x_111 = (*(tint_symbol_8)).stageSize;
   stageUnits = (float2(1.0f) / x_111);
   i = 0;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     int const x_122 = i;
     if ((x_122 < 2)) {
     } else {
@@ -135,8 +136,7 @@
       float const x_184 = animationData[2];
       (*(tint_private_vars)).mt = fmod((x_181 * x_184), 1.0f);
       f = 0.0f;
-      while (true) {
-        __asm__("");
+      while (tint_preserve_loop) {
         float const x_193 = f;
         if ((x_193 < 8.0f)) {
         } else {
diff --git a/test/tint/bug/tint/949.wgsl.expected.msl b/test/tint/bug/tint/949.wgsl.expected.msl
index 5c21f24..b936ff4 100644
--- a/test/tint/bug/tint/949.wgsl.expected.msl
+++ b/test/tint/bug/tint/949.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   float u_Float;
   float3 u_Color;
@@ -330,8 +332,7 @@
   lastSampledHeight = 1.0f;
   currSampledHeight = 1.0f;
   i = 0;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     int const x_388 = i;
     if ((x_388 < 15)) {
     } else {
diff --git a/test/tint/bug/tint/990.wgsl.expected.msl b/test/tint/bug/tint/990.wgsl.expected.msl
index ff2feee..7d1836d1 100644
--- a/test/tint/bug/tint/990.wgsl.expected.msl
+++ b/test/tint/bug/tint/990.wgsl.expected.msl
@@ -1,10 +1,12 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void f() {
   int i = 0;
-  for(; false; ) {
-    __asm__("");
+  for(; (false) == tint_preserve_loop; ) {
   }
 }
 
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
index 6407f10..d612d9e 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   uint local_invocation_index_1;
 };
@@ -29,8 +31,7 @@
 void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
   uint idx = 0u;
   idx = local_invocation_index_2;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     if (!((idx < 6u))) {
       break;
     }
@@ -54,8 +55,7 @@
 }
 
 void compute_main_inner_1(uint local_invocation_index_1_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_2) {
-  for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
-    __asm__("");
+  for(uint idx_1 = local_invocation_index_1_param; ((idx_1 < 6u)) == tint_preserve_loop; idx_1 = (idx_1 + 1u)) {
     uint const i = (idx_1 / 2u);
     uint const i_1 = (idx_1 % 2u);
     uint const i_2 = (idx_1 % 1u);
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl
index 7deaf57..7cf8ff3 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl
@@ -14,9 +14,10 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
-  for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 6u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = (idx / 2u);
     uint const i_1 = (idx % 2u);
     uint const i_2 = (idx % 1u);
diff --git a/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl b/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
index 856968a..7177f3f 100644
--- a/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   uint local_invocation_index_1;
 };
@@ -21,8 +23,7 @@
 void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<atomic_uint, 4>* const tint_symbol) {
   uint idx = 0u;
   idx = local_invocation_index_2;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     if (!((idx < 4u))) {
       break;
     }
@@ -44,8 +45,7 @@
 }
 
 void compute_main_inner_1(uint local_invocation_index_1_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<atomic_uint, 4>* const tint_symbol_2) {
-  for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 4u); idx_1 = (idx_1 + 1u)) {
-    __asm__("");
+  for(uint idx_1 = local_invocation_index_1_param; ((idx_1 < 4u)) == tint_preserve_loop; idx_1 = (idx_1 + 1u)) {
     uint const i = idx_1;
     atomic_store_explicit(&((*(tint_symbol_2))[i]), 0u, memory_order_relaxed);
   }
diff --git a/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl b/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl
index d3036aa..a5fb3f5 100644
--- a/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl
@@ -14,9 +14,10 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 void compute_main_inner(uint local_invocation_index, threadgroup tint_array<atomic_uint, 4>* const tint_symbol) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     atomic_store_explicit(&((*(tint_symbol))[i]), 0u, memory_order_relaxed);
   }
diff --git a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
index 6407f10..d612d9e 100644
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   uint local_invocation_index_1;
 };
@@ -29,8 +31,7 @@
 void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
   uint idx = 0u;
   idx = local_invocation_index_2;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     if (!((idx < 6u))) {
       break;
     }
@@ -54,8 +55,7 @@
 }
 
 void compute_main_inner_1(uint local_invocation_index_1_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_2) {
-  for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
-    __asm__("");
+  for(uint idx_1 = local_invocation_index_1_param; ((idx_1 < 6u)) == tint_preserve_loop; idx_1 = (idx_1 + 1u)) {
     uint const i = (idx_1 / 2u);
     uint const i_1 = (idx_1 % 2u);
     uint const i_2 = (idx_1 % 1u);
diff --git a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl
index 7deaf57..7cf8ff3 100644
--- a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl
@@ -14,9 +14,10 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
-  for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 6u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = (idx / 2u);
     uint const i_1 = (idx % 2u);
     uint const i_2 = (idx % 1u);
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
index 81f7e0e..29ecbd8 100644
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   uint local_invocation_index_1;
 };
@@ -33,8 +35,7 @@
 void compute_main_inner(uint local_invocation_index_2, threadgroup tint_array<S_atomic, 10>* const tint_symbol) {
   uint idx = 0u;
   idx = local_invocation_index_2;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     if (!((idx < 10u))) {
       break;
     }
@@ -58,8 +59,7 @@
 }
 
 void compute_main_inner_1(uint local_invocation_index_1_param, thread tint_private_vars_struct* const tint_private_vars, threadgroup tint_array<S_atomic, 10>* const tint_symbol_2) {
-  for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
-    __asm__("");
+  for(uint idx_1 = local_invocation_index_1_param; ((idx_1 < 10u)) == tint_preserve_loop; idx_1 = (idx_1 + 1u)) {
     uint const i = idx_1;
     (*(tint_symbol_2))[i].x = 0;
     atomic_store_explicit(&((*(tint_symbol_2))[i].a), 0u, memory_order_relaxed);
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl
index 6d38f69..f46810f 100644
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   int x;
   atomic_uint a;
@@ -21,8 +23,7 @@
 };
 
 void compute_main_inner(uint local_invocation_index, threadgroup tint_array<S, 10>* const tint_symbol) {
-  for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 10u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i].x = 0;
     atomic_store_explicit(&((*(tint_symbol))[i].a), 0u, memory_order_relaxed);
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
index 9f335ca..e543c39 100644
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   uint local_invocation_index_1;
 };
@@ -35,8 +37,7 @@
   (*(tint_symbol)).x = 0;
   (*(tint_symbol)).y = 0u;
   idx = local_invocation_index_2;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     if (!((idx < 10u))) {
       break;
     }
@@ -62,8 +63,7 @@
     (*(tint_symbol_2)).x = 0;
     (*(tint_symbol_2)).y = 0u;
   }
-  for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
-    __asm__("");
+  for(uint idx_1 = local_invocation_index_1_param; ((idx_1 < 10u)) == tint_preserve_loop; idx_1 = (idx_1 + 1u)) {
     uint const i = idx_1;
     atomic_store_explicit(&((*(tint_symbol_2)).a[i]), 0u, memory_order_relaxed);
   }
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl
index edaee4cb..d4e7b34 100644
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   int x;
   tint_array<atomic_uint, 10> a;
@@ -25,8 +27,7 @@
     (*(tint_symbol)).x = 0;
     (*(tint_symbol)).y = 0u;
   }
-  for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 10u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     atomic_store_explicit(&((*(tint_symbol)).a[i]), 0u, memory_order_relaxed);
   }
diff --git a/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.msl b/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.msl
index 82b5b7d..1998f18 100644
--- a/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.msl
+++ b/test/tint/builtins/textureStore/loop_continuing_read_write_texture.wgsl.expected.msl
@@ -1,11 +1,13 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void foo(texture2d<int, access::read_write> tint_symbol) {
   {
     int i = 0;
-    while (true) {
-      __asm__("");
+    while (tint_preserve_loop) {
       if (!((i < 3))) {
         break;
       }
diff --git a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.msl b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.msl
index 5819f07..3fa0ee1 100644
--- a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.msl
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.msl
@@ -1,6 +1,9 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 int tint_workgroupUniformLoad(threadgroup int* const p) {
   threadgroup_barrier(mem_flags::mem_threadgroup);
   int const result = *(p);
@@ -11,8 +14,7 @@
 void foo(threadgroup int* const tint_symbol_4, threadgroup int* const tint_symbol_5) {
   {
     int i = 0;
-    while (true) {
-      __asm__("");
+    while (tint_preserve_loop) {
       int const tint_symbol = i;
       int const tint_symbol_1 = tint_workgroupUniformLoad(tint_symbol_4);
       if (!((tint_symbol < tint_symbol_1))) {
diff --git a/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.msl
index f722f40..4c582c9 100644
--- a/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.msl
@@ -13,14 +13,16 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_2 {
   float x [[user(locn0)]];
 };
 
 void tint_symbol_inner(float x) {
   float4 v = float4(0.0f);
-  for(; ((x > v[0]) && (dfdx(1.0f) > 0.0f)); ) {
-    __asm__("");
+  for(; (((x > v[0]) && (dfdx(1.0f) > 0.0f))) == tint_preserve_loop; ) {
   }
 }
 
diff --git a/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.msl
index 918a01b..63ce7b4 100644
--- a/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.msl
@@ -13,14 +13,16 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_2 {
   float x [[user(locn0)]];
 };
 
 void tint_symbol_inner(float x, texture2d<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
   float4 v = float4(0.0f);
-  for(; (x > v[0]); ) {
-    __asm__("");
+  for(; ((x > v[0])) == tint_preserve_loop; ) {
     v = tint_symbol_3.sample(tint_symbol_4, float2(0.0f));
   }
 }
diff --git a/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.msl
index 86fc7d2..7fb929f 100644
--- a/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.msl
@@ -13,13 +13,15 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_2 {
   float x [[user(locn0)]];
 };
 
 void tint_symbol_inner(float x) {
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     {
       if ((x > 0.0f)) { break; }
     }
diff --git a/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.msl
index 063fabb..7dc46dd 100644
--- a/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.msl
@@ -13,13 +13,15 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_2 {
   float x [[user(locn0)]];
 };
 
 void tint_symbol_inner(float x) {
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     {
       if ((x > 0.0f)) { break; }
     }
diff --git a/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.msl
index 741fbad..38bded8 100644
--- a/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.msl
@@ -13,13 +13,15 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_2 {
   float x [[user(locn0)]];
 };
 
 void tint_symbol_inner(float x) {
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     {
       if ((x > 0.0f)) { break; }
     }
diff --git a/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.msl
index 7858cd2..6a4adfd 100644
--- a/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.msl
@@ -13,14 +13,16 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_2 {
   float x [[user(locn0)]];
 };
 
 void tint_symbol_inner(float x) {
   float4 v = float4(0.0f);
-  while(((x > 0.0f) && (dfdx(1.0f) > 0.0f))) {
-    __asm__("");
+  while((((x > 0.0f) && (dfdx(1.0f) > 0.0f))) == tint_preserve_loop) {
   }
 }
 
diff --git a/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.msl b/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.msl
index 4eb60f8..9f81385 100644
--- a/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.msl
+++ b/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.msl
@@ -13,14 +13,16 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_2 {
   float x [[user(locn0)]];
 };
 
 void tint_symbol_inner(float x, texture2d<float, access::sample> tint_symbol_3, sampler tint_symbol_4) {
   float4 v = float4(0.0f);
-  while((x > v[0])) {
-    __asm__("");
+  while(((x > v[0])) == tint_preserve_loop) {
     v = tint_symbol_3.sample(tint_symbol_4, float2(0.0f));
   }
 }
diff --git a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.msl b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.msl
index 03ea0b0..666903f 100644
--- a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.msl
+++ b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct strided_arr {
   /* 0x0000 */ float2 el;
   /* 0x0008 */ tint_array<int8_t, 8> tint_pad;
@@ -39,8 +41,7 @@
 }
 
 void assign_and_preserve_padding(device tint_array<strided_arr, 2>* const dest, tint_array<strided_arr, 2> value) {
-  for(uint i = 0u; (i < 2u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 2u)) == tint_preserve_loop; i = (i + 1u)) {
     assign_and_preserve_padding_1(&((*(dest))[i]), value[i]);
   }
 }
diff --git a/test/tint/loops/continue_in_switch.wgsl.expected.msl b/test/tint/loops/continue_in_switch.wgsl.expected.msl
index 5d5f9c6..db2c859 100644
--- a/test/tint/loops/continue_in_switch.wgsl.expected.msl
+++ b/test/tint/loops/continue_in_switch.wgsl.expected.msl
@@ -1,9 +1,11 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 kernel void f() {
-  for(int i = 0; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(int i = 0; ((i < 4)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     switch(i) {
       case 0: {
         continue;
diff --git a/test/tint/loops/loop.wgsl.expected.msl b/test/tint/loops/loop.wgsl.expected.msl
index 6eae6c8..096ffe4 100644
--- a/test/tint/loops/loop.wgsl.expected.msl
+++ b/test/tint/loops/loop.wgsl.expected.msl
@@ -1,10 +1,12 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 int f() {
   int i = 0;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
     if ((i > 4)) {
       return i;
diff --git a/test/tint/loops/loop_with_break_if.wgsl.expected.msl b/test/tint/loops/loop_with_break_if.wgsl.expected.msl
index 7dc8ba3..865b88d 100644
--- a/test/tint/loops/loop_with_break_if.wgsl.expected.msl
+++ b/test/tint/loops/loop_with_break_if.wgsl.expected.msl
@@ -1,10 +1,12 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 int f() {
   int i = 0;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     if ((i > 4)) {
       return i;
     }
diff --git a/test/tint/loops/loop_with_continuing.wgsl.expected.msl b/test/tint/loops/loop_with_continuing.wgsl.expected.msl
index 9a4f4d0..cdb7a82 100644
--- a/test/tint/loops/loop_with_continuing.wgsl.expected.msl
+++ b/test/tint/loops/loop_with_continuing.wgsl.expected.msl
@@ -1,10 +1,12 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 int f() {
   int i = 0;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     if ((i > 4)) {
       return i;
     }
diff --git a/test/tint/loops/multiple_continues.wgsl.expected.msl b/test/tint/loops/multiple_continues.wgsl.expected.msl
index ed9d214..53d5235 100644
--- a/test/tint/loops/multiple_continues.wgsl.expected.msl
+++ b/test/tint/loops/multiple_continues.wgsl.expected.msl
@@ -1,9 +1,11 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 kernel void tint_symbol() {
-  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     switch(i) {
       case 0: {
         continue;
diff --git a/test/tint/loops/multiple_switch.wgsl.expected.msl b/test/tint/loops/multiple_switch.wgsl.expected.msl
index adc1ac8..2bb537d 100644
--- a/test/tint/loops/multiple_switch.wgsl.expected.msl
+++ b/test/tint/loops/multiple_switch.wgsl.expected.msl
@@ -1,10 +1,12 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 kernel void tint_symbol() {
   int i = 0;
-  for(int i_1 = 0; (i_1 < 2); i_1 = as_type<int>((as_type<uint>(i_1) + as_type<uint>(1)))) {
-    __asm__("");
+  for(int i_1 = 0; ((i_1 < 2)) == tint_preserve_loop; i_1 = as_type<int>((as_type<uint>(i_1) + as_type<uint>(1)))) {
     switch(i_1) {
       case 0: {
         continue;
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
index 6f58d1b..dc0686f 100644
--- a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
@@ -1,11 +1,12 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 kernel void tint_symbol() {
-  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
-    __asm__("");
-    for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
-      __asm__("");
+  for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+    for(int j = 0; ((j < 2)) == tint_preserve_loop; j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
       switch(i) {
         case 0: {
           continue;
diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
index ce28e7d..0dab2a6 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
@@ -1,13 +1,14 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 kernel void tint_symbol() {
-  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
-    __asm__("");
+  for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
     switch(i) {
       case 0: {
-        for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
-          __asm__("");
+        for(int j = 0; ((j < 2)) == tint_preserve_loop; j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
           switch(j) {
             case 0: {
               continue;
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl
index 2690672..23a482e 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl
@@ -1,14 +1,15 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 kernel void tint_symbol() {
   int k = 0;
-  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
-    __asm__("");
+  for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
     switch(i) {
       case 0: {
-        for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
-          __asm__("");
+        for(int j = 0; ((j < 2)) == tint_preserve_loop; j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
           switch(j) {
             case 0: {
               continue;
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
index 8402967..5885440 100644
--- a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
@@ -1,10 +1,12 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 kernel void tint_symbol() {
   int j = 0;
-  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
-    __asm__("");
+  for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
     switch(i) {
       case 0: {
         switch(j) {
diff --git a/test/tint/loops/nested_loops.wgsl.expected.msl b/test/tint/loops/nested_loops.wgsl.expected.msl
index a84996f..2fd8895 100644
--- a/test/tint/loops/nested_loops.wgsl.expected.msl
+++ b/test/tint/loops/nested_loops.wgsl.expected.msl
@@ -1,17 +1,18 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 int f() {
   int i = 0;
   int j = 0;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
     if ((i > 4)) {
       return 1;
     }
-    while (true) {
-      __asm__("");
+    while (tint_preserve_loop) {
       j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
       if ((j > 4)) {
         return 2;
diff --git a/test/tint/loops/nested_loops_with_continuing.wgsl.expected.msl b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.msl
index 17720f6..e0b172b 100644
--- a/test/tint/loops/nested_loops_with_continuing.wgsl.expected.msl
+++ b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.msl
@@ -1,16 +1,17 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 int f() {
   int i = 0;
   int j = 0;
-  while (true) {
-    __asm__("");
+  while (tint_preserve_loop) {
     if ((i > 4)) {
       return 1;
     }
-    while (true) {
-      __asm__("");
+    while (tint_preserve_loop) {
       if ((j > 4)) {
         return 2;
       }
diff --git a/test/tint/loops/single_continue.wgsl.expected.msl b/test/tint/loops/single_continue.wgsl.expected.msl
index 467b2e5..2e85db7 100644
--- a/test/tint/loops/single_continue.wgsl.expected.msl
+++ b/test/tint/loops/single_continue.wgsl.expected.msl
@@ -1,9 +1,11 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 kernel void tint_symbol() {
-  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(int i = 0; ((i < 2)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     switch(i) {
       case 0: {
         continue;
diff --git a/test/tint/loops/while.wgsl.expected.msl b/test/tint/loops/while.wgsl.expected.msl
index dd2c7ec..e2655d3 100644
--- a/test/tint/loops/while.wgsl.expected.msl
+++ b/test/tint/loops/while.wgsl.expected.msl
@@ -1,10 +1,12 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 int f() {
   int i = 0;
-  while((i < 4)) {
-    __asm__("");
+  while(((i < 4)) == tint_preserve_loop) {
     i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
   }
   return i;
diff --git a/test/tint/loops/while_with_continue.wgsl.expected.msl b/test/tint/loops/while_with_continue.wgsl.expected.msl
index 6c9936f..68d388a 100644
--- a/test/tint/loops/while_with_continue.wgsl.expected.msl
+++ b/test/tint/loops/while_with_continue.wgsl.expected.msl
@@ -1,10 +1,12 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 int f() {
   int i = 0;
-  while((i < 4)) {
-    __asm__("");
+  while(((i < 4)) == tint_preserve_loop) {
     i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
     continue;
   }
diff --git a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.msl b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.msl
index 8b496c6..7541aea 100644
--- a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct str {
   tint_array<int, 4> arr;
 };
@@ -23,8 +25,7 @@
 }
 
 void tint_symbol_inner(uint local_invocation_index, threadgroup str* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol_1)).arr[i] = 0;
   }
diff --git a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.msl b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.msl
index b424dc8..28efaee 100644
--- a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct str {
   int i;
 };
@@ -23,8 +25,7 @@
 }
 
 void tint_symbol_inner(uint local_invocation_index, threadgroup tint_array<str, 4>* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i_1 = idx;
     str const tint_symbol_1 = str{};
     (*(tint_symbol_2))[i_1] = tint_symbol_1;
diff --git a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.msl b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.msl
index b63f3c0..fd889f8 100644
--- a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct str {
   tint_array<int, 4> arr;
 };
@@ -24,8 +26,7 @@
 }
 
 void tint_symbol_inner(uint local_invocation_index, threadgroup str* const tint_symbol_2) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol_2)).arr[i] = 0;
   }
diff --git a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.msl b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.msl
index 01ad03d..1ff952b 100644
--- a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct str {
   int i;
 };
@@ -24,8 +26,7 @@
 }
 
 void tint_symbol_inner(uint local_invocation_index, threadgroup tint_array<str, 4>* const tint_symbol_3) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 4u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i_1 = idx;
     str const tint_symbol_1 = str{};
     (*(tint_symbol_3))[i_1] = tint_symbol_1;
diff --git a/test/tint/samples/compute_boids.wgsl.expected.msl b/test/tint/samples/compute_boids.wgsl.expected.msl
index e6b5e22..b7fe17a 100644
--- a/test/tint/samples/compute_boids.wgsl.expected.msl
+++ b/test/tint/samples/compute_boids.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_symbol_1 {
   float2 a_particlePos [[attribute(0)]];
   float2 a_particleVel [[attribute(1)]];
@@ -85,8 +87,7 @@
   int cVelCount = 0;
   float2 pos = 0.0f;
   float2 vel = 0.0f;
-  for(uint i = 0u; (i < 5u); i = (i + 1u)) {
-    __asm__("");
+  for(uint i = 0u; ((i < 5u)) == tint_preserve_loop; i = (i + 1u)) {
     if ((i == index)) {
       continue;
     }
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl
index 95e334c..64296fb 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct Uniforms {
   /* 0x0000 */ uint i;
 };
@@ -29,8 +31,7 @@
 kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
   InnerS v = {};
   OuterS s1 = {};
-  for(int i = 0; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(int i = 0; ((i < 4)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     s1.a1[(*(tint_symbol_1)).i] = v;
   }
   return;
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl
index 33d1be5..74b50fc 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct Uniforms {
   /* 0x0000 */ uint i;
 };
@@ -29,8 +31,7 @@
 kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
   InnerS v = {};
   OuterS s1 = {};
-  for(int i = 0; (i < 4); s1.a1[(*(tint_symbol_1)).i] = v) {
-    __asm__("");
+  for(int i = 0; ((i < 4)) == tint_preserve_loop; s1.a1[(*(tint_symbol_1)).i] = v) {
     i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
   }
   return;
diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl
index 163f692..76e6c8d 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct Uniforms {
   /* 0x0000 */ uint i;
 };
@@ -30,8 +32,7 @@
   InnerS v = {};
   OuterS s1 = {};
   int i = 0;
-  for(s1.a1[(*(tint_symbol_1)).i] = v; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(s1.a1[(*(tint_symbol_1)).i] = v; ((i < 4)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
   }
   return;
 }
diff --git a/test/tint/statements/compound_assign/for_loop.wgsl.expected.msl b/test/tint/statements/compound_assign/for_loop.wgsl.expected.msl
index f9d89b4..99a18b6 100644
--- a/test/tint/statements/compound_assign/for_loop.wgsl.expected.msl
+++ b/test/tint/statements/compound_assign/for_loop.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   uint i;
 };
@@ -45,8 +47,7 @@
     int const tint_symbol_2 = idx1(tint_private_vars);
     int const tint_symbol_save = tint_symbol_2;
     a[tint_symbol_save] = (a[tint_symbol_save] * 2.0f);
-    while (true) {
-      __asm__("");
+    while (tint_preserve_loop) {
       int const tint_symbol_3 = idx2(tint_private_vars);
       if (!((a[tint_symbol_3] < 10.0f))) {
         break;
diff --git a/test/tint/statements/decrement/complex.wgsl.expected.msl b/test/tint/statements/decrement/complex.wgsl.expected.msl
index de484b1..7efbaf0 100644
--- a/test/tint/statements/decrement/complex.wgsl.expected.msl
+++ b/test/tint/statements/decrement/complex.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   uint v;
 };
@@ -60,8 +62,7 @@
     int const tint_symbol_2_save_1 = tint_symbol_7;
     int const tint_symbol_3 = idx3(tint_private_vars);
     (*(tint_symbol_10))[tint_symbol_2_save].a[tint_symbol_2_save_1][tint_symbol_3] = as_type<int>((as_type<uint>((*(tint_symbol_10))[tint_symbol_2_save].a[tint_symbol_2_save_1][tint_symbol_3]) - as_type<uint>(1)));
-    while (true) {
-      __asm__("");
+    while (tint_preserve_loop) {
       if (!(((*(tint_private_vars)).v < 10u))) {
         break;
       }
diff --git a/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.msl b/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.msl
index faeb478..3c3e4da 100644
--- a/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.msl
+++ b/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.msl
@@ -1,9 +1,11 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void tint_symbol(device uint* const tint_symbol_1) {
-  for(; (*(tint_symbol_1) < 10u); *(tint_symbol_1) = (*(tint_symbol_1) - 1u)) {
-    __asm__("");
+  for(; ((*(tint_symbol_1) < 10u)) == tint_preserve_loop; *(tint_symbol_1) = (*(tint_symbol_1) - 1u)) {
   }
 }
 
diff --git a/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.msl b/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.msl
index b700dca..89002b5 100644
--- a/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.msl
+++ b/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.msl
@@ -1,9 +1,11 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void tint_symbol(device uint* const tint_symbol_1) {
-  for(*(tint_symbol_1) = (*(tint_symbol_1) - 1u); (*(tint_symbol_1) < 10u); ) {
-    __asm__("");
+  for(*(tint_symbol_1) = (*(tint_symbol_1) - 1u); ((*(tint_symbol_1) < 10u)) == tint_preserve_loop; ) {
   }
 }
 
diff --git a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.msl b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.msl
index 4931e68..fd1c50f 100644
--- a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.msl
+++ b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.msl
@@ -1,6 +1,9 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   bool tint_discarded;
 };
@@ -25,8 +28,7 @@
   int result = tint_ftoi(tint_symbol_4.sample(tint_symbol_5, coord)[0]);
   {
     int i = 0;
-    while (true) {
-      __asm__("");
+    while (tint_preserve_loop) {
       if (!((i < 10))) {
         break;
       }
diff --git a/test/tint/statements/discard/multiple_returns.wgsl.expected.msl b/test/tint/statements/discard/multiple_returns.wgsl.expected.msl
index 037a078..75f3dd5 100644
--- a/test/tint/statements/discard/multiple_returns.wgsl.expected.msl
+++ b/test/tint/statements/discard/multiple_returns.wgsl.expected.msl
@@ -1,6 +1,9 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   bool tint_discarded;
 };
@@ -16,8 +19,7 @@
   }
   if ((*(tint_symbol_2) < 0.0f)) {
     int i = 0;
-    while (true) {
-      __asm__("");
+    while (tint_preserve_loop) {
       if ((*(tint_symbol_2) > float(i))) {
         if (!(tint_private_vars.tint_discarded)) {
           *(tint_symbol_2) = float(i);
diff --git a/test/tint/statements/for/basic.wgsl.expected.msl b/test/tint/statements/for/basic.wgsl.expected.msl
index d4ee785..453a4d3 100644
--- a/test/tint/statements/for/basic.wgsl.expected.msl
+++ b/test/tint/statements/for/basic.wgsl.expected.msl
@@ -1,12 +1,14 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void some_loop_body() {
 }
 
 void f() {
-  for(int i = 0; (i < 5); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(int i = 0; ((i < 5)) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     some_loop_body();
   }
 }
diff --git a/test/tint/statements/for/complex.wgsl.expected.msl b/test/tint/statements/for/complex.wgsl.expected.msl
index 297f77e..cda792d 100644
--- a/test/tint/statements/for/complex.wgsl.expected.msl
+++ b/test/tint/statements/for/complex.wgsl.expected.msl
@@ -1,13 +1,15 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void some_loop_body() {
 }
 
 void f() {
   int j = 0;
-  for(int i = 0; ((i < 5) && (j < 10)); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(int i = 0; (((i < 5) && (j < 10))) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     some_loop_body();
     j = as_type<int>((as_type<uint>(i) * as_type<uint>(30)));
   }
diff --git a/test/tint/statements/for/condition/array_ctor.wgsl.expected.msl b/test/tint/statements/for/condition/array_ctor.wgsl.expected.msl
index 0bfd498..c2e0634 100644
--- a/test/tint/statements/for/condition/array_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/condition/array_ctor.wgsl.expected.msl
@@ -1,10 +1,12 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void f() {
   int i = 0;
-  for(; (i < 1); ) {
-    __asm__("");
+  for(; ((i < 1)) == tint_preserve_loop; ) {
   }
 }
 
diff --git a/test/tint/statements/for/condition/basic.wgsl.expected.msl b/test/tint/statements/for/condition/basic.wgsl.expected.msl
index c29c578..7da63a4 100644
--- a/test/tint/statements/for/condition/basic.wgsl.expected.msl
+++ b/test/tint/statements/for/condition/basic.wgsl.expected.msl
@@ -1,10 +1,12 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void f() {
   int i = 0;
-  for(; (i < 4); ) {
-    __asm__("");
+  for(; ((i < 4)) == tint_preserve_loop; ) {
   }
 }
 
diff --git a/test/tint/statements/for/condition/struct_ctor.wgsl.expected.msl b/test/tint/statements/for/condition/struct_ctor.wgsl.expected.msl
index 4b4bd74..ca21e5a 100644
--- a/test/tint/statements/for/condition/struct_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/condition/struct_ctor.wgsl.expected.msl
@@ -1,14 +1,16 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   int i;
 };
 
 void f() {
   int i = 0;
-  for(; (i < 1); ) {
-    __asm__("");
+  for(; ((i < 1)) == tint_preserve_loop; ) {
   }
 }
 
diff --git a/test/tint/statements/for/continuing/array_ctor.wgsl.expected.msl b/test/tint/statements/for/continuing/array_ctor.wgsl.expected.msl
index 6c7e914..2da7062 100644
--- a/test/tint/statements/for/continuing/array_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/continuing/array_ctor.wgsl.expected.msl
@@ -1,10 +1,12 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void f() {
   int i = 0;
-  for(; false; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(; (false) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
   }
 }
 
diff --git a/test/tint/statements/for/continuing/basic.wgsl.expected.msl b/test/tint/statements/for/continuing/basic.wgsl.expected.msl
index 6c7e914..2da7062 100644
--- a/test/tint/statements/for/continuing/basic.wgsl.expected.msl
+++ b/test/tint/statements/for/continuing/basic.wgsl.expected.msl
@@ -1,10 +1,12 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void f() {
   int i = 0;
-  for(; false; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(; (false) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
   }
 }
 
diff --git a/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.msl b/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.msl
index 45ae3b6..95ed7eb 100644
--- a/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.msl
@@ -1,13 +1,15 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   int i;
 };
 
 void f() {
-  for(int i = 0; false; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
+  for(int i = 0; (false) == tint_preserve_loop; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
   }
 }
 
diff --git a/test/tint/statements/for/empty.wgsl.expected.msl b/test/tint/statements/for/empty.wgsl.expected.msl
index 0032ddd..d05f6d5 100644
--- a/test/tint/statements/for/empty.wgsl.expected.msl
+++ b/test/tint/statements/for/empty.wgsl.expected.msl
@@ -1,9 +1,11 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void f() {
-  for(; false; ) {
-    __asm__("");
+  for(; (false) == tint_preserve_loop; ) {
   }
 }
 
diff --git a/test/tint/statements/for/initializer/array_ctor.wgsl.expected.msl b/test/tint/statements/for/initializer/array_ctor.wgsl.expected.msl
index 2baee65..6e66415 100644
--- a/test/tint/statements/for/initializer/array_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/initializer/array_ctor.wgsl.expected.msl
@@ -1,9 +1,11 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void f() {
-  for(int i = 1; false; ) {
-    __asm__("");
+  for(int i = 1; (false) == tint_preserve_loop; ) {
   }
 }
 
diff --git a/test/tint/statements/for/initializer/basic.wgsl.expected.msl b/test/tint/statements/for/initializer/basic.wgsl.expected.msl
index 505277d..c50d83b 100644
--- a/test/tint/statements/for/initializer/basic.wgsl.expected.msl
+++ b/test/tint/statements/for/initializer/basic.wgsl.expected.msl
@@ -1,9 +1,11 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void f() {
-  for(int i = 0; false; ) {
-    __asm__("");
+  for(int i = 0; (false) == tint_preserve_loop; ) {
   }
 }
 
diff --git a/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.msl b/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.msl
index 6718a1f..60dbd8a 100644
--- a/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.msl
+++ b/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.msl
@@ -1,13 +1,15 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 struct S {
   int i;
 };
 
 void f() {
-  for(int i = 1; false; ) {
-    __asm__("");
+  for(int i = 1; (false) == tint_preserve_loop; ) {
   }
 }
 
diff --git a/test/tint/statements/for/scoping.wgsl.expected.msl b/test/tint/statements/for/scoping.wgsl.expected.msl
index 4a7d54e..0206ba1 100644
--- a/test/tint/statements/for/scoping.wgsl.expected.msl
+++ b/test/tint/statements/for/scoping.wgsl.expected.msl
@@ -1,9 +1,11 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void f() {
-  for(int must_not_collide = 0; ; ) {
-    __asm__("");
+  for(int must_not_collide = 0; tint_preserve_loop; ) {
     break;
   }
   int must_not_collide = 0;
diff --git a/test/tint/statements/increment/complex.wgsl.expected.msl b/test/tint/statements/increment/complex.wgsl.expected.msl
index ec39dfd..509a125 100644
--- a/test/tint/statements/increment/complex.wgsl.expected.msl
+++ b/test/tint/statements/increment/complex.wgsl.expected.msl
@@ -14,6 +14,8 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 struct tint_private_vars_struct {
   uint v;
 };
@@ -60,8 +62,7 @@
     int const tint_symbol_2_save_1 = tint_symbol_7;
     int const tint_symbol_3 = idx3(tint_private_vars);
     (*(tint_symbol_10))[tint_symbol_2_save].a[tint_symbol_2_save_1][tint_symbol_3] = as_type<int>((as_type<uint>((*(tint_symbol_10))[tint_symbol_2_save].a[tint_symbol_2_save_1][tint_symbol_3]) + as_type<uint>(1)));
-    while (true) {
-      __asm__("");
+    while (tint_preserve_loop) {
       if (!(((*(tint_private_vars)).v < 10u))) {
         break;
       }
diff --git a/test/tint/statements/increment/for_loop_continuing.wgsl.expected.msl b/test/tint/statements/increment/for_loop_continuing.wgsl.expected.msl
index 239618b..8092545 100644
--- a/test/tint/statements/increment/for_loop_continuing.wgsl.expected.msl
+++ b/test/tint/statements/increment/for_loop_continuing.wgsl.expected.msl
@@ -1,9 +1,11 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void tint_symbol(device uint* const tint_symbol_1) {
-  for(; (*(tint_symbol_1) < 10u); *(tint_symbol_1) = (*(tint_symbol_1) + 1u)) {
-    __asm__("");
+  for(; ((*(tint_symbol_1) < 10u)) == tint_preserve_loop; *(tint_symbol_1) = (*(tint_symbol_1) + 1u)) {
   }
 }
 
diff --git a/test/tint/statements/increment/for_loop_initializer.wgsl.expected.msl b/test/tint/statements/increment/for_loop_initializer.wgsl.expected.msl
index 37b9419..2cc9f36 100644
--- a/test/tint/statements/increment/for_loop_initializer.wgsl.expected.msl
+++ b/test/tint/statements/increment/for_loop_initializer.wgsl.expected.msl
@@ -1,9 +1,11 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+constant static volatile bool tint_preserve_loop = true;
+
 void tint_symbol(device uint* const tint_symbol_1) {
-  for(*(tint_symbol_1) = (*(tint_symbol_1) + 1u); (*(tint_symbol_1) < 10u); ) {
-    __asm__("");
+  for(*(tint_symbol_1) = (*(tint_symbol_1) + 1u); ((*(tint_symbol_1) < 10u)) == tint_preserve_loop; ) {
   }
 }
 
diff --git a/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.msl b/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.msl
index cc02af2..a578747 100644
--- a/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.msl
+++ b/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.msl
@@ -14,9 +14,10 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 void tint_symbol_inner(uint local_invocation_index, threadgroup tint_array<tint_array<int, 3>, 2>* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 6u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = (idx / 3u);
     uint const i_1 = (idx % 3u);
     (*(tint_symbol_1))[i][i_1] = 0;
diff --git a/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.msl b/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.msl
index 00dd95f..513d0a3 100644
--- a/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.msl
+++ b/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.msl
@@ -14,9 +14,10 @@
     T elements[N];
 };
 
+constant static volatile bool tint_preserve_loop = true;
+
 void tint_symbol_inner(uint local_invocation_index, threadgroup tint_array<int, 3>* const tint_symbol_1) {
-  for(uint idx = local_invocation_index; (idx < 3u); idx = (idx + 1u)) {
-    __asm__("");
+  for(uint idx = local_invocation_index; ((idx < 3u)) == tint_preserve_loop; idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol_1))[i] = 0;
   }