diff --git a/src/reader/wgsl/parser_impl.cc b/src/reader/wgsl/parser_impl.cc
index d80e8c7..a3b9b23 100644
--- a/src/reader/wgsl/parser_impl.cc
+++ b/src/reader/wgsl/parser_impl.cc
@@ -2071,7 +2071,7 @@
 
 // for_statement
 //   : FOR PAREN_LEFT for_header PAREN_RIGHT BRACE_LEFT statements BRACE_RIGHT
-Maybe<ast::Statement*> ParserImpl::for_stmt() {
+Maybe<ast::ForLoopStatement*> ParserImpl::for_stmt() {
   Source source;
   if (!match(Token::Type::kFor, &source))
     return Failure::kNoMatch;
@@ -2086,45 +2086,9 @@
   if (stmts.errored)
     return Failure::kErrored;
 
-  // The for statement is a syntactic sugar on top of the loop statement.
-  // We create corresponding nodes in ast with the exact same behaviour
-  // as we would expect from the loop statement.
-  if (header->condition != nullptr) {
-    // !condition
-    auto* not_condition = create<ast::UnaryOpExpression>(
-        header->condition->source(), ast::UnaryOp::kNot, header->condition);
-    // { break; }
-    auto* break_stmt = create<ast::BreakStatement>(not_condition->source());
-    auto* break_body =
-        create<ast::BlockStatement>(not_condition->source(), ast::StatementList{
-                                                                 break_stmt,
-                                                             });
-    // if (!condition) { break; }
-    auto* break_if_not_condition =
-        create<ast::IfStatement>(not_condition->source(), not_condition,
-                                 break_body, ast::ElseStatementList{});
-    stmts.value.insert(stmts.value.begin(), break_if_not_condition);
-  }
-
-  ast::BlockStatement* continuing_body = nullptr;
-  if (header->continuing != nullptr) {
-    continuing_body = create<ast::BlockStatement>(header->continuing->source(),
-                                                  ast::StatementList{
-                                                      header->continuing,
-                                                  });
-  }
-
-  auto* body = create<ast::BlockStatement>(source, stmts.value);
-  auto* loop = create<ast::LoopStatement>(source, body, continuing_body);
-
-  if (header->initializer != nullptr) {
-    return create<ast::BlockStatement>(source, ast::StatementList{
-                                                   header->initializer,
-                                                   loop,
-                                               });
-  }
-
-  return loop;
+  return create<ast::ForLoopStatement>(
+      source, header->initializer, header->condition, header->continuing,
+      create<ast::BlockStatement>(stmts.value));
 }
 
 // func_call_stmt
diff --git a/src/reader/wgsl/parser_impl.h b/src/reader/wgsl/parser_impl.h
index 232d952..78192aa 100644
--- a/src/reader/wgsl/parser_impl.h
+++ b/src/reader/wgsl/parser_impl.h
@@ -554,7 +554,7 @@
   Expect<std::unique_ptr<ForHeader>> expect_for_header();
   /// Parses a `for_stmt` grammar element
   /// @returns the parsed for loop or nullptr
-  Maybe<ast::Statement*> for_stmt();
+  Maybe<ast::ForLoopStatement*> for_stmt();
   /// Parses a `continuing_stmt` grammar element
   /// @returns the parsed statements
   Maybe<ast::BlockStatement*> continuing_stmt();
diff --git a/src/reader/wgsl/parser_impl_for_stmt_test.cc b/src/reader/wgsl/parser_impl_for_stmt_test.cc
index 5b72897..1aac60e 100644
--- a/src/reader/wgsl/parser_impl_for_stmt_test.cc
+++ b/src/reader/wgsl/parser_impl_for_stmt_test.cc
@@ -14,141 +14,154 @@
 
 #include "src/reader/wgsl/parser_impl_test_helper.h"
 
+#include "src/ast/discard_statement.h"
+
 namespace tint {
 namespace reader {
 namespace wgsl {
 namespace {
 
-class ForStmtTest : public ParserImplTest {
- public:
-  void TestForLoop(std::string loop_str, std::string for_str) {
-    auto p_loop = parser(loop_str);
-    auto e_loop = p_loop->expect_statements();
-    EXPECT_FALSE(e_loop.errored);
-    EXPECT_FALSE(p_loop->has_error()) << p_loop->error();
-
-    auto p_for = parser(for_str);
-    auto e_for = p_for->expect_statements();
-    EXPECT_FALSE(e_for.errored);
-    EXPECT_FALSE(p_for->has_error()) << p_for->error();
-
-    std::string loop = ast::BlockStatement({}, {}, e_loop.value).str(Sem());
-    std::string for_ = ast::BlockStatement({}, {}, e_for.value).str(Sem());
-    EXPECT_EQ(loop, for_);
-  }
-};
+using ForStmtTest = ParserImplTest;
 
 // Test an empty for loop.
 TEST_F(ForStmtTest, Empty) {
-  std::string for_str = "for (;;) { }";
-  std::string loop_str = "loop { }";
-
-  TestForLoop(loop_str, for_str);
+  auto p = parser("for (;;) { }");
+  auto fl = p->for_stmt();
+  EXPECT_FALSE(p->has_error()) << p->error();
+  EXPECT_FALSE(fl.errored);
+  ASSERT_TRUE(fl.matched);
+  EXPECT_EQ(fl->initializer(), nullptr);
+  EXPECT_EQ(fl->condition(), nullptr);
+  EXPECT_EQ(fl->continuing(), nullptr);
+  EXPECT_TRUE(fl->body()->empty());
 }
 
 // Test a for loop with non-empty body.
 TEST_F(ForStmtTest, Body) {
-  std::string for_str = "for (;;) { discard; }";
-  std::string loop_str = "loop { discard; }";
-
-  TestForLoop(loop_str, for_str);
+  auto p = parser("for (;;) { discard; }");
+  auto fl = p->for_stmt();
+  EXPECT_FALSE(p->has_error()) << p->error();
+  EXPECT_FALSE(fl.errored);
+  ASSERT_TRUE(fl.matched);
+  EXPECT_EQ(fl->initializer(), nullptr);
+  EXPECT_EQ(fl->condition(), nullptr);
+  EXPECT_EQ(fl->continuing(), nullptr);
+  ASSERT_EQ(fl->body()->size(), 1u);
+  EXPECT_TRUE(fl->body()->statements()[0]->Is<ast::DiscardStatement>());
 }
 
 // Test a for loop declaring a variable in the initializer statement.
 TEST_F(ForStmtTest, InitializerStatementDecl) {
-  std::string for_str = "for (var i: i32 ;;) { }";
-  std::string loop_str = "{ var i: i32; loop { } }";
-
-  TestForLoop(loop_str, for_str);
+  auto p = parser("for (var i: i32 ;;) { }");
+  auto fl = p->for_stmt();
+  EXPECT_FALSE(p->has_error()) << p->error();
+  EXPECT_FALSE(fl.errored);
+  ASSERT_TRUE(fl.matched);
+  ASSERT_TRUE(Is<ast::VariableDeclStatement>(fl->initializer()));
+  auto* var = fl->initializer()->As<ast::VariableDeclStatement>()->variable();
+  EXPECT_FALSE(var->is_const());
+  EXPECT_EQ(var->constructor(), nullptr);
+  EXPECT_EQ(fl->condition(), nullptr);
+  EXPECT_EQ(fl->continuing(), nullptr);
+  EXPECT_TRUE(fl->body()->empty());
 }
 
 // Test a for loop declaring and initializing a variable in the initializer
 // statement.
 TEST_F(ForStmtTest, InitializerStatementDeclEqual) {
-  std::string for_str = "for (var i: i32 = 0 ;;) { }";
-  std::string loop_str = "{ var i: i32 = 0; loop { } }";
-
-  TestForLoop(loop_str, for_str);
+  auto p = parser("for (var i: i32 = 0 ;;) { }");
+  auto fl = p->for_stmt();
+  EXPECT_FALSE(p->has_error()) << p->error();
+  EXPECT_FALSE(fl.errored);
+  ASSERT_TRUE(fl.matched);
+  ASSERT_TRUE(Is<ast::VariableDeclStatement>(fl->initializer()));
+  auto* var = fl->initializer()->As<ast::VariableDeclStatement>()->variable();
+  EXPECT_FALSE(var->is_const());
+  EXPECT_NE(var->constructor(), nullptr);
+  EXPECT_EQ(fl->condition(), nullptr);
+  EXPECT_EQ(fl->continuing(), nullptr);
+  EXPECT_TRUE(fl->body()->empty());
 }
 
 // Test a for loop declaring a const variable in the initializer statement.
 TEST_F(ForStmtTest, InitializerStatementConstDecl) {
-  std::string for_str = "for (let i: i32 = 0 ;;) { }";
-  std::string loop_str = "{ let i: i32 = 0; loop { } }";
-
-  TestForLoop(loop_str, for_str);
+  auto p = parser("for (let i: i32 = 0 ;;) { }");
+  auto fl = p->for_stmt();
+  EXPECT_FALSE(p->has_error()) << p->error();
+  EXPECT_FALSE(fl.errored);
+  ASSERT_TRUE(fl.matched);
+  ASSERT_TRUE(Is<ast::VariableDeclStatement>(fl->initializer()));
+  auto* var = fl->initializer()->As<ast::VariableDeclStatement>()->variable();
+  EXPECT_TRUE(var->is_const());
+  EXPECT_NE(var->constructor(), nullptr);
+  EXPECT_EQ(fl->condition(), nullptr);
+  EXPECT_EQ(fl->continuing(), nullptr);
+  EXPECT_TRUE(fl->body()->empty());
 }
 
 // Test a for loop assigning a variable in the initializer statement.
 TEST_F(ForStmtTest, InitializerStatementAssignment) {
-  std::string for_str = "var i: i32; for (i = 0 ;;) { }";
-  std::string loop_str = "var i: i32; { i = 0; loop { } }";
-
-  TestForLoop(loop_str, for_str);
+  auto p = parser("for (i = 0 ;;) { }");
+  auto fl = p->for_stmt();
+  EXPECT_FALSE(p->has_error()) << p->error();
+  EXPECT_FALSE(fl.errored);
+  ASSERT_TRUE(fl.matched);
+  EXPECT_TRUE(Is<ast::AssignmentStatement>(fl->initializer()));
+  EXPECT_EQ(fl->condition(), nullptr);
+  EXPECT_EQ(fl->continuing(), nullptr);
+  EXPECT_TRUE(fl->body()->empty());
 }
 
 // Test a for loop calling a function in the initializer statement.
 TEST_F(ForStmtTest, InitializerStatementFuncCall) {
-  std::string for_str = "for (a(b,c) ;;) { }";
-  std::string loop_str = "{ a(b,c); loop { } }";
-
-  TestForLoop(loop_str, for_str);
+  auto p = parser("for (a(b,c) ;;) { }");
+  auto fl = p->for_stmt();
+  EXPECT_FALSE(p->has_error()) << p->error();
+  EXPECT_FALSE(fl.errored);
+  ASSERT_TRUE(fl.matched);
+  EXPECT_TRUE(Is<ast::CallStatement>(fl->initializer()));
+  EXPECT_EQ(fl->condition(), nullptr);
+  EXPECT_EQ(fl->continuing(), nullptr);
+  EXPECT_TRUE(fl->body()->empty());
 }
 
 // Test a for loop with a break condition
 TEST_F(ForStmtTest, BreakCondition) {
-  std::string for_str = "for (; 0 == 1;) { }";
-  std::string loop_str = "loop { if (!(0 == 1)) { break; } }";
-
-  TestForLoop(loop_str, for_str);
+  auto p = parser("for (; 0 == 1;) { }");
+  auto fl = p->for_stmt();
+  EXPECT_FALSE(p->has_error()) << p->error();
+  EXPECT_FALSE(fl.errored);
+  ASSERT_TRUE(fl.matched);
+  EXPECT_EQ(fl->initializer(), nullptr);
+  EXPECT_TRUE(Is<ast::BinaryExpression>(fl->condition()));
+  EXPECT_EQ(fl->continuing(), nullptr);
+  EXPECT_TRUE(fl->body()->empty());
 }
 
 // Test a for loop assigning a variable in the continuing statement.
 TEST_F(ForStmtTest, ContinuingAssignment) {
-  std::string for_str = "var x: i32; for (;; x = 2) { }";
-  std::string loop_str = "var x: i32; loop { continuing { x = 2; }}";
-
-  TestForLoop(loop_str, for_str);
+  auto p = parser("for (;; x = 2) { }");
+  auto fl = p->for_stmt();
+  EXPECT_FALSE(p->has_error()) << p->error();
+  EXPECT_FALSE(fl.errored);
+  ASSERT_TRUE(fl.matched);
+  EXPECT_EQ(fl->initializer(), nullptr);
+  EXPECT_EQ(fl->condition(), nullptr);
+  EXPECT_TRUE(Is<ast::AssignmentStatement>(fl->continuing()));
+  EXPECT_TRUE(fl->body()->empty());
 }
 
 // Test a for loop calling a function in the continuing statement.
 TEST_F(ForStmtTest, ContinuingFuncCall) {
-  std::string for_str = "for (;; a(b,c)) { }";
-  std::string loop_str = "loop { continuing { a(b,c); }}";
-
-  TestForLoop(loop_str, for_str);
-}
-
-// Test a for loop with all statements non-empty.
-TEST_F(ForStmtTest, All) {
-  std::string for_str =
-      R"(for(var i : i32 = 0; i < 4; i = i + 1) {
-       if (a == 0) {
-        continue;
-       }
-       a = a + 2;
-     })";
-
-  std::string loop_str =
-      R"({ // Introduce new scope for loop variable i
-      var i : i32 = 0;
-      loop {
-        if (!(i < 4)) {
-          break;
-        }
-
-        if (a == 0) {
-          continue;
-        }
-        a = a + 2;
-
-        continuing {
-          i = i + 1;
-        }
-      }
-    };)";
-
-  TestForLoop(loop_str, for_str);
+  auto p = parser("for (;; a(b,c)) { }");
+  auto fl = p->for_stmt();
+  EXPECT_FALSE(p->has_error()) << p->error();
+  EXPECT_FALSE(fl.errored);
+  ASSERT_TRUE(fl.matched);
+  EXPECT_EQ(fl->initializer(), nullptr);
+  EXPECT_EQ(fl->condition(), nullptr);
+  EXPECT_TRUE(Is<ast::CallStatement>(fl->continuing()));
+  EXPECT_TRUE(fl->body()->empty());
 }
 
 class ForStmtErrorTest : public ParserImplTest {
diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc
index 65c0ab0..881d380 100644
--- a/src/resolver/resolver.cc
+++ b/src/resolver/resolver.cc
@@ -62,6 +62,7 @@
 #include "src/sem/storage_texture_type.h"
 #include "src/sem/struct.h"
 #include "src/sem/variable.h"
+#include "src/utils/defer.h"
 #include "src/utils/get_or_create.h"
 #include "src/utils/math.h"
 #include "src/utils/scoped_assignment.h"
@@ -1915,6 +1916,10 @@
       stmt->body(), current_statement_);
   builder_->Sem().Add(stmt->body(), sem_block_body);
   TINT_SCOPED_ASSIGNMENT(current_statement_, sem_block_body);
+  TINT_SCOPED_ASSIGNMENT(current_block_, sem_block_body);
+
+  variable_stack_.push_scope();
+  TINT_DEFER(variable_stack_.pop_scope());
 
   if (auto* initializer = stmt->initializer()) {
     Mark(initializer);
@@ -4011,9 +4016,8 @@
   TINT_SCOPED_ASSIGNMENT(current_block_,
                          const_cast<sem::BlockStatement*>(sem_block));
   variable_stack_.push_scope();
-  bool result = callback();
-  variable_stack_.pop_scope();
-  return result;
+  TINT_DEFER(variable_stack_.pop_scope());
+  return callback();
 }
 
 std::string Resolver::VectorPretty(uint32_t size,
diff --git a/src/transform/spirv.cc b/src/transform/spirv.cc
index 1d5a460..0d84459 100644
--- a/src/transform/spirv.cc
+++ b/src/transform/spirv.cc
@@ -29,6 +29,7 @@
 #include "src/sem/variable.h"
 #include "src/transform/external_texture_transform.h"
 #include "src/transform/fold_constants.h"
+#include "src/transform/for_loop_to_loop.h"
 #include "src/transform/inline_pointer_lets.h"
 #include "src/transform/manager.h"
 #include "src/transform/simplify.h"
@@ -50,6 +51,7 @@
   manager.Add<Simplify>();           // Required for arrayLength()
   manager.Add<FoldConstants>();
   manager.Add<ExternalTextureTransform>();
+  manager.Add<ForLoopToLoop>();  // Must come after ZeroInitWorkgroupMemory
   auto transformedInput = manager.Run(in, data);
 
   auto* cfg = data.Get<Config>();
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index c93caa7..7ab9ea5 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -43,6 +43,7 @@
 #include "src/sem/variable.h"
 #include "src/transform/calculate_array_length.h"
 #include "src/transform/hlsl.h"
+#include "src/utils/defer.h"
 #include "src/utils/get_or_create.h"
 #include "src/utils/scoped_assignment.h"
 #include "src/writer/append_vector.h"
@@ -2574,6 +2575,15 @@
 }
 
 bool GeneratorImpl::EmitForLoop(ast::ForLoopStatement* stmt) {
+  // Nest a for loop with a new block. In HLSL the initializer scope is not
+  // nested by the for-loop, so we may get variable redefinitions.
+  line() << "{";
+  increment_indent();
+  TINT_DEFER({
+    decrement_indent();
+    line() << "}";
+  });
+
   TextBuffer init_buf;
   if (auto* init = stmt->initializer()) {
     TINT_SCOPED_ASSIGNMENT(current_buffer_, &init_buf);
@@ -2581,16 +2591,6 @@
       return false;
     }
   }
-  bool multi_stmt_init = init_buf.lines.size() > 1;
-  // For-loop has multi-statement initializer.
-  // This cannot be emitted with a regular for loop, so instead nest the loop in
-  // a new block scope prefixed with these initializer statements.
-  if (multi_stmt_init) {
-    line() << "{";
-    increment_indent();
-    current_buffer_->Append(init_buf);
-    init_buf.lines.clear();  // Don't emit the initializer again in the 'for'
-  }
 
   TextBuffer cond_pre;
   std::stringstream cond_buf;
@@ -2609,10 +2609,20 @@
     }
   }
 
-  if (cond_pre.lines.size() > 0 || cont_buf.lines.size() > 1) {
-    // For-loop has multi-statement conditional and / or continuing.
-    // This cannot be emitted with a regular for loop, so instead generate a
-    // `while(true)` loop.
+  // If the for-loop has a multi-statement conditional and / or continuing, then
+  // we cannot emit this as a regular for-loop in HLSL. Instead we need to
+  // generate a `while(true)` loop.
+  bool emit_as_loop = cond_pre.lines.size() > 0 || cont_buf.lines.size() > 1;
+
+  // If the for-loop has multi-statement initializer, or is going to be emitted
+  // as a `while(true)` loop, then declare the initializer statement(s) before
+  // the loop.
+  if (init_buf.lines.size() > 1 || (stmt->initializer() && emit_as_loop)) {
+    current_buffer_->Append(init_buf);
+    init_buf.lines.clear();  // Don't emit the initializer again in the 'for'
+  }
+
+  if (emit_as_loop) {
     auto emit_continuing = [&]() {
       current_buffer_->Append(cont_buf);
       return true;
@@ -2620,23 +2630,24 @@
 
     TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
     line() << "while (true) {";
-    {
-      ScopedIndent si(this);
+    increment_indent();
+    TINT_DEFER({
+      decrement_indent();
+      line() << "}";
+    });
 
-      if (stmt->condition()) {
-        current_buffer_->Append(cond_pre);
-        line() << "if (!(" << cond_buf.str() << ")) { break; }";
-      }
-
-      if (!EmitStatements(stmt->body()->statements())) {
-        return false;
-      }
-
-      if (!emit_continuing()) {
-        return false;
-      }
+    if (stmt->condition()) {
+      current_buffer_->Append(cond_pre);
+      line() << "if (!(" << cond_buf.str() << ")) { break; }";
     }
-    line() << "}";
+
+    if (!EmitStatements(stmt->body()->statements())) {
+      return false;
+    }
+
+    if (!emit_continuing()) {
+      return false;
+    }
   } else {
     // For-loop can be generated.
     {
@@ -2666,12 +2677,6 @@
         return false;
       }
     }
-
-    line() << "}";
-  }
-
-  if (multi_stmt_init) {
-    decrement_indent();
     line() << "}";
   }
 
diff --git a/src/writer/hlsl/generator_impl_loop_test.cc b/src/writer/hlsl/generator_impl_loop_test.cc
index 0811a6e..b9d006f 100644
--- a/src/writer/hlsl/generator_impl_loop_test.cc
+++ b/src/writer/hlsl/generator_impl_loop_test.cc
@@ -161,8 +161,10 @@
   gen.increment_indent();
 
   ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(  for(; ; ) {
-    return;
+  EXPECT_EQ(gen.result(), R"(  {
+    for(; ; ) {
+      return;
+    }
   }
 )");
 }
@@ -180,8 +182,10 @@
   gen.increment_indent();
 
   ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(  for(int i = 0; ; ) {
-    return;
+  EXPECT_EQ(gen.result(), R"(  {
+    for(int i = 0; ; ) {
+      return;
+    }
   }
 )");
 }
@@ -227,8 +231,10 @@
   gen.increment_indent();
 
   ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(  for(; true; ) {
-    return;
+  EXPECT_EQ(gen.result(), R"(  {
+    for(; true; ) {
+      return;
+    }
   }
 )");
 }
@@ -248,13 +254,15 @@
   gen.increment_indent();
 
   ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(  while (true) {
-    bool tint_tmp = true;
-    if (tint_tmp) {
-      tint_tmp = false;
+  EXPECT_EQ(gen.result(), R"(  {
+    while (true) {
+      bool tint_tmp = true;
+      if (tint_tmp) {
+        tint_tmp = false;
+      }
+      if (!((tint_tmp))) { break; }
+      return;
     }
-    if (!((tint_tmp))) { break; }
-    return;
   }
 )");
 }
@@ -273,8 +281,10 @@
   gen.increment_indent();
 
   ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(  for(; ; i = (i + 1)) {
-    return;
+  EXPECT_EQ(gen.result(), R"(  {
+    for(; ; i = (i + 1)) {
+      return;
+    }
   }
 )");
 }
@@ -295,13 +305,15 @@
   gen.increment_indent();
 
   ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(  while (true) {
-    return;
-    bool tint_tmp = true;
-    if (tint_tmp) {
-      tint_tmp = false;
+  EXPECT_EQ(gen.result(), R"(  {
+    while (true) {
+      return;
+      bool tint_tmp = true;
+      if (tint_tmp) {
+        tint_tmp = false;
+      }
+      i = (tint_tmp);
     }
-    i = (tint_tmp);
   }
 )");
 }
@@ -320,8 +332,10 @@
   gen.increment_indent();
 
   ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(  for(int i = 0; true; i = (i + 1)) {
-    return;
+  EXPECT_EQ(gen.result(), R"(  {
+    for(int i = 0; true; i = (i + 1)) {
+      return;
+    }
   }
 )");
 }
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index 42d5a36..809ca10 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -54,6 +54,7 @@
 #include "src/sem/vector_type.h"
 #include "src/sem/void_type.h"
 #include "src/transform/msl.h"
+#include "src/utils/defer.h"
 #include "src/utils/scoped_assignment.h"
 #include "src/writer/float_to_string.h"
 
@@ -1493,6 +1494,116 @@
   return true;
 }
 
+bool GeneratorImpl::EmitForLoop(ast::ForLoopStatement* stmt) {
+  TextBuffer init_buf;
+  if (auto* init = stmt->initializer()) {
+    TINT_SCOPED_ASSIGNMENT(current_buffer_, &init_buf);
+    if (!EmitStatement(init)) {
+      return false;
+    }
+  }
+
+  TextBuffer cond_pre;
+  std::stringstream cond_buf;
+  if (auto* cond = stmt->condition()) {
+    TINT_SCOPED_ASSIGNMENT(current_buffer_, &cond_pre);
+    if (!EmitExpression(cond_buf, cond)) {
+      return false;
+    }
+  }
+
+  TextBuffer cont_buf;
+  if (auto* cont = stmt->continuing()) {
+    TINT_SCOPED_ASSIGNMENT(current_buffer_, &cont_buf);
+    if (!EmitStatement(cont)) {
+      return false;
+    }
+  }
+
+  // If the for-loop has a multi-statement conditional and / or continuing, then
+  // we cannot emit this as a regular for-loop in MSL. Instead we need to
+  // generate a `while(true)` loop.
+  bool emit_as_loop = cond_pre.lines.size() > 0 || cont_buf.lines.size() > 1;
+
+  // If the for-loop has multi-statement initializer, or is going to be emitted
+  // as a `while(true)` loop, then declare the initializer statement(s) before
+  // the loop in a new block.
+  bool nest_in_block =
+      init_buf.lines.size() > 1 || (stmt->initializer() && emit_as_loop);
+  if (nest_in_block) {
+    line() << "{";
+    increment_indent();
+    current_buffer_->Append(init_buf);
+    init_buf.lines.clear();  // Don't emit the initializer again in the 'for'
+  }
+  TINT_DEFER({
+    if (nest_in_block) {
+      decrement_indent();
+      line() << "}";
+    }
+  });
+
+  if (emit_as_loop) {
+    auto emit_continuing = [&]() {
+      current_buffer_->Append(cont_buf);
+      return true;
+    };
+
+    TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
+    line() << "while (true) {";
+    increment_indent();
+    TINT_DEFER({
+      decrement_indent();
+      line() << "}";
+    });
+
+    if (stmt->condition()) {
+      current_buffer_->Append(cond_pre);
+      line() << "if (!(" << cond_buf.str() << ")) { break; }";
+    }
+
+    if (!EmitStatements(stmt->body()->statements())) {
+      return false;
+    }
+
+    if (!emit_continuing()) {
+      return false;
+    }
+  } else {
+    // For-loop can be generated.
+    {
+      auto out = line();
+      out << "for";
+      {
+        ScopedParen sp(out);
+
+        if (!init_buf.lines.empty()) {
+          out << init_buf.lines[0].content << " ";
+        } else {
+          out << "; ";
+        }
+
+        out << cond_buf.str() << "; ";
+
+        if (!cont_buf.lines.empty()) {
+          out << TrimSuffix(cont_buf.lines[0].content, ";");
+        }
+      }
+      out << " {";
+    }
+    {
+      auto emit_continuing = [] { return true; };
+      TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
+      if (!EmitStatementsWithIndent(stmt->body()->statements())) {
+        return false;
+      }
+    }
+    line() << "}";
+  }
+
+  return true;
+}
+
 bool GeneratorImpl::EmitDiscard(ast::DiscardStatement*) {
   // TODO(dsinclair): Verify this is correct when the discard semantics are
   // defined for WGSL (https://github.com/gpuweb/gpuweb/issues/361)
@@ -1635,6 +1746,9 @@
   if (auto* l = stmt->As<ast::LoopStatement>()) {
     return EmitLoop(l);
   }
+  if (auto* l = stmt->As<ast::ForLoopStatement>()) {
+    return EmitForLoop(l);
+  }
   if (auto* r = stmt->As<ast::ReturnStatement>()) {
     return EmitReturn(r);
   }
diff --git a/src/writer/msl/generator_impl.h b/src/writer/msl/generator_impl.h
index 0f5bf1a..1495178 100644
--- a/src/writer/msl/generator_impl.h
+++ b/src/writer/msl/generator_impl.h
@@ -170,6 +170,10 @@
   /// @param stmt the statement to emit
   /// @returns true if the statement was emitted
   bool EmitLoop(ast::LoopStatement* stmt);
+  /// Handles a for loop statement
+  /// @param stmt the statement to emit
+  /// @returns true if the statement was emitted
+  bool EmitForLoop(ast::ForLoopStatement* stmt);
   /// Handles a member accessor expression
   /// @param out the output of the expression stream
   /// @param expr the member accessor expression
diff --git a/src/writer/msl/generator_impl_loop_test.cc b/src/writer/msl/generator_impl_loop_test.cc
index 0d18299..a6ae411 100644
--- a/src/writer/msl/generator_impl_loop_test.cc
+++ b/src/writer/msl/generator_impl_loop_test.cc
@@ -141,6 +141,227 @@
 )");
 }
 
+TEST_F(MslGeneratorImplTest, Emit_ForLoop) {
+  // for(; ; ) {
+  //   return;
+  // }
+
+  auto* f = For(nullptr, nullptr, nullptr, Block(Return()));
+  WrapInFunction(f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  for(; ; ) {
+    return;
+  }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithSimpleInit) {
+  // for(var i : i32; ; ) {
+  //   return;
+  // }
+
+  auto* f = For(Decl(Var("i", ty.i32())), nullptr, nullptr, Block(Return()));
+  WrapInFunction(f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  for(int i = 0; ; ) {
+    return;
+  }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithMultiStmtInit) {
+  // var<workgroup> a : atomic<i32>;
+  // for(var b = atomicCompareExchangeWeak(&a, 1, 2); ; ) {
+  //   return;
+  // }
+  Global("a", ty.atomic<i32>(), ast::StorageClass::kWorkgroup);
+  auto* multi_stmt = Call("atomicCompareExchangeWeak", AddressOf("a"), 1, 2);
+  auto* f = For(Decl(Var("b", nullptr, multi_stmt)), nullptr, nullptr,
+                Block(Return()));
+  WrapInFunction(f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  {
+    int prev_value = 1;
+    bool matched = atomic_compare_exchange_weak_explicit(&(a), &prev_value, 2, memory_order_relaxed, memory_order_relaxed);
+    int2 b = int2(prev_value, matched);
+    for(; ; ) {
+      return;
+    }
+  }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithSimpleCond) {
+  // for(; true; ) {
+  //   return;
+  // }
+
+  auto* f = For(nullptr, true, nullptr, Block(Return()));
+  WrapInFunction(f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  for(; true; ) {
+    return;
+  }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithMultiStmtCond) {
+  // var<workgroup> a : atomic<i32>;
+  // for(; atomicCompareExchangeWeak(&a, 1, 2).x == 0; ) {
+  //   return;
+  // }
+
+  Global("a", ty.atomic<i32>(), ast::StorageClass::kWorkgroup);
+  auto* multi_stmt = create<ast::BinaryExpression>(
+      ast::BinaryOp::kEqual,
+      MemberAccessor(Call("atomicCompareExchangeWeak", AddressOf("a"), 1, 2),
+                     "x"),
+      Expr(0));
+  auto* f = For(nullptr, multi_stmt, nullptr, Block(Return()));
+  WrapInFunction(f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  while (true) {
+    int prev_value = 1;
+    bool matched = atomic_compare_exchange_weak_explicit(&(a), &prev_value, 2, memory_order_relaxed, memory_order_relaxed);
+    if (!((int2(prev_value, matched).x == 0))) { break; }
+    return;
+  }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithSimpleCont) {
+  // for(; ; i = i + 1) {
+  //   return;
+  // }
+
+  auto* v = Decl(Var("i", ty.i32()));
+  auto* f = For(nullptr, nullptr, Assign("i", Add("i", 1)), Block(Return()));
+  WrapInFunction(v, f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  for(; ; i = (i + 1)) {
+    return;
+  }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithMultiStmtCont) {
+  // var<workgroup> a : atomic<i32>;
+  // for(; ; ignore(atomicCompareExchangeWeak(&a, 1, 2))) {
+  //   return;
+  // }
+
+  Global("a", ty.atomic<i32>(), ast::StorageClass::kWorkgroup);
+  auto* multi_stmt =
+      Ignore(Call("atomicCompareExchangeWeak", AddressOf("a"), 1, 2));
+  auto* f = For(nullptr, nullptr, multi_stmt, Block(Return()));
+  WrapInFunction(f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  while (true) {
+    return;
+    int prev_value = 1;
+    bool matched = atomic_compare_exchange_weak_explicit(&(a), &prev_value, 2, memory_order_relaxed, memory_order_relaxed);
+    (void) int2(prev_value, matched);
+  }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithSimpleInitCondCont) {
+  // for(var i : i32; true; i = i + 1) {
+  //   return;
+  // }
+
+  auto* f = For(Decl(Var("i", ty.i32())), true, Assign("i", Add("i", 1)),
+                Block(Return()));
+  WrapInFunction(f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  for(int i = 0; true; i = (i + 1)) {
+    return;
+  }
+)");
+}
+
+TEST_F(MslGeneratorImplTest, Emit_ForLoopWithMultiStmtInitCondCont) {
+  // var<workgroup> a : atomic<i32>;
+  // for(var b = atomicCompareExchangeWeak(&a, 1, 2);
+  //     atomicCompareExchangeWeak(&a, 1, 2).x == 0;
+  //     ignore(atomicCompareExchangeWeak(&a, 1, 2))) {
+  //   return;
+  // }
+  Global("a", ty.atomic<i32>(), ast::StorageClass::kWorkgroup);
+  auto* multi_stmt_a = Call("atomicCompareExchangeWeak", AddressOf("a"), 1, 2);
+  auto* multi_stmt_b = create<ast::BinaryExpression>(
+      ast::BinaryOp::kEqual,
+      MemberAccessor(Call("atomicCompareExchangeWeak", AddressOf("a"), 1, 2),
+                     "x"),
+      Expr(0));
+  auto* multi_stmt_c =
+      Ignore(Call("atomicCompareExchangeWeak", AddressOf("a"), 1, 2));
+  auto* f = For(Decl(Var("b", nullptr, multi_stmt_a)), multi_stmt_b,
+                multi_stmt_c, Block(Return()));
+  WrapInFunction(f);
+
+  GeneratorImpl& gen = Build();
+
+  gen.increment_indent();
+
+  ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
+  EXPECT_EQ(gen.result(), R"(  {
+    int prev_value = 1;
+    bool matched = atomic_compare_exchange_weak_explicit(&(a), &prev_value, 2, memory_order_relaxed, memory_order_relaxed);
+    int2 b = int2(prev_value, matched);
+    while (true) {
+      int prev_value_1 = 1;
+      bool matched_1 = atomic_compare_exchange_weak_explicit(&(a), &prev_value_1, 2, memory_order_relaxed, memory_order_relaxed);
+      if (!((int2(prev_value_1, matched_1).x == 0))) { break; }
+      return;
+      int prev_value_2 = 1;
+      bool matched_2 = atomic_compare_exchange_weak_explicit(&(a), &prev_value_2, 2, memory_order_relaxed, memory_order_relaxed);
+      (void) int2(prev_value_2, matched_2);
+    }
+  }
+)");
+}
+
 }  // namespace
 }  // namespace msl
 }  // namespace writer
diff --git a/test/bug/tint/534.wgsl.expected.hlsl b/test/bug/tint/534.wgsl.expected.hlsl
index 2c9ea16..e7a735c 100644
--- a/test/bug/tint/534.wgsl.expected.hlsl
+++ b/test/bug/tint/534.wgsl.expected.hlsl
@@ -43,7 +43,7 @@
     uint i = 0u;
     while (true) {
       const uint scalar_offset_1 = (12u) / 4;
-      if (!(!(!((i < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]))))) { break; }
+      if (!((i < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]))) { break; }
       Set_uint4(srcColorBits, i, ConvertToFp16FloatValue(srcColor[i]));
       bool tint_tmp_1 = success;
       if (tint_tmp_1) {
diff --git a/test/bug/tint/534.wgsl.expected.msl b/test/bug/tint/534.wgsl.expected.msl
index 429c94b..356f44d 100644
--- a/test/bug/tint/534.wgsl.expected.msl
+++ b/test/bug/tint/534.wgsl.expected.msl
@@ -27,18 +27,9 @@
   bool success = true;
   uint4 srcColorBits = 0u;
   uint4 dstColorBits = uint4(dstColor);
-  {
-    uint i = 0u;
-    while (true) {
-      if (!((i < uniforms.channelCount))) {
-        break;
-      }
-      srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
-      success = (success && (srcColorBits[i] == dstColorBits[i]));
-      {
-        i = (i + 1u);
-      }
-    }
+  for(uint i = 0u; (i < uniforms.channelCount); i = (i + 1u)) {
+    srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
+    success = (success && (srcColorBits[i] == dstColorBits[i]));
   }
   uint outputIndex = ((GlobalInvocationID.y * uint(size.x)) + GlobalInvocationID.x);
   if (success) {
diff --git a/test/bug/tint/534.wgsl.expected.wgsl b/test/bug/tint/534.wgsl.expected.wgsl
index 91a330b..1fd409c 100644
--- a/test/bug/tint/534.wgsl.expected.wgsl
+++ b/test/bug/tint/534.wgsl.expected.wgsl
@@ -36,19 +36,9 @@
   var success : bool = true;
   var srcColorBits : vec4<u32>;
   var dstColorBits : vec4<u32> = vec4<u32>(dstColor);
-  {
-    var i : u32 = 0u;
-    loop {
-      if (!((i < uniforms.channelCount))) {
-        break;
-      }
-      srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
-      success = (success && (srcColorBits[i] == dstColorBits[i]));
-
-      continuing {
-        i = (i + 1u);
-      }
-    }
+  for(var i : u32 = 0u; (i < uniforms.channelCount); i = (i + 1u)) {
+    srcColorBits[i] = ConvertToFp16FloatValue(srcColor[i]);
+    success = (success && (srcColorBits[i] == dstColorBits[i]));
   }
   var outputIndex : u32 = ((GlobalInvocationID.y * u32(size.x)) + GlobalInvocationID.x);
   if (success) {
diff --git a/test/bug/tint/744.wgsl.expected.hlsl b/test/bug/tint/744.wgsl.expected.hlsl
index 0314864..47f47a7 100644
--- a/test/bug/tint/744.wgsl.expected.hlsl
+++ b/test/bug/tint/744.wgsl.expected.hlsl
@@ -19,8 +19,7 @@
   const uint dimOutter = uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4];
   uint result = 0u;
   {
-    uint i = 0u;
-    for(; !(!((i < dimInner))); i = (i + 1u)) {
+    for(uint i = 0u; (i < dimInner); i = (i + 1u)) {
       const uint a = (i + (resultCell.x * dimInner));
       const uint b = (resultCell.y + (i * dimOutter));
       result = (result + (firstMatrix.Load((4u * a)) * secondMatrix.Load((4u * b))));
diff --git a/test/bug/tint/744.wgsl.expected.msl b/test/bug/tint/744.wgsl.expected.msl
index 7855d18..16ebd85 100644
--- a/test/bug/tint/744.wgsl.expected.msl
+++ b/test/bug/tint/744.wgsl.expected.msl
@@ -15,19 +15,10 @@
   uint const dimInner = uniforms.aShape.y;
   uint const dimOutter = uniforms.outShape.y;
   uint result = 0u;
-  {
-    uint i = 0u;
-    while (true) {
-      if (!((i < dimInner))) {
-        break;
-      }
-      uint const a = (i + (resultCell.x * dimInner));
-      uint const b = (resultCell.y + (i * dimOutter));
-      result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
-      {
-        i = (i + 1u);
-      }
-    }
+  for(uint i = 0u; (i < dimInner); i = (i + 1u)) {
+    uint const a = (i + (resultCell.x * dimInner));
+    uint const b = (resultCell.y + (i * dimOutter));
+    result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
   }
   uint const index = (resultCell.y + (resultCell.x * dimOutter));
   resultMatrix.numbers[index] = result;
diff --git a/test/bug/tint/744.wgsl.expected.wgsl b/test/bug/tint/744.wgsl.expected.wgsl
index d96d510..ad6c01e 100644
--- a/test/bug/tint/744.wgsl.expected.wgsl
+++ b/test/bug/tint/744.wgsl.expected.wgsl
@@ -24,20 +24,10 @@
   let dimInner : u32 = uniforms.aShape.y;
   let dimOutter : u32 = uniforms.outShape.y;
   var result : u32 = 0u;
-  {
-    var i : u32 = 0u;
-    loop {
-      if (!((i < dimInner))) {
-        break;
-      }
-      let a : u32 = (i + (resultCell.x * dimInner));
-      let b : u32 = (resultCell.y + (i * dimOutter));
-      result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
-
-      continuing {
-        i = (i + 1u);
-      }
-    }
+  for(var i : u32 = 0u; (i < dimInner); i = (i + 1u)) {
+    let a : u32 = (i + (resultCell.x * dimInner));
+    let b : u32 = (resultCell.y + (i * dimOutter));
+    result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b]));
   }
   let index : u32 = (resultCell.y + (resultCell.x * dimOutter));
   resultMatrix.numbers[index] = result;
diff --git a/test/bug/tint/757.wgsl.expected.hlsl b/test/bug/tint/757.wgsl.expected.hlsl
index 72ed5bd..da3073c 100644
--- a/test/bug/tint/757.wgsl.expected.hlsl
+++ b/test/bug/tint/757.wgsl.expected.hlsl
@@ -16,8 +16,7 @@
   flatIndex = (flatIndex * 1u);
   float4 texel = myTexture.Load(int4(GlobalInvocationID.xy, 0, 0));
   {
-    uint i = 0u;
-    for(; !(!((i < 1u))); i = (i + 1u)) {
+    for(uint i = 0u; (i < 1u); i = (i + 1u)) {
       result.Store((4u * (flatIndex + i)), asuint(texel.r));
     }
   }
diff --git a/test/bug/tint/757.wgsl.expected.msl b/test/bug/tint/757.wgsl.expected.msl
index 03e2067..8699607 100644
--- a/test/bug/tint/757.wgsl.expected.msl
+++ b/test/bug/tint/757.wgsl.expected.msl
@@ -12,17 +12,8 @@
   uint flatIndex = ((((2u * 2u) * GlobalInvocationID.z) + (2u * GlobalInvocationID.y)) + GlobalInvocationID.x);
   flatIndex = (flatIndex * 1u);
   float4 texel = tint_symbol_2.read(uint2(int2(GlobalInvocationID.xy)), 0, 0);
-  {
-    uint i = 0u;
-    while (true) {
-      if (!((i < 1u))) {
-        break;
-      }
-      result.values[(flatIndex + i)] = texel.r;
-      {
-        i = (i + 1u);
-      }
-    }
+  for(uint i = 0u; (i < 1u); i = (i + 1u)) {
+    result.values[(flatIndex + i)] = texel.r;
   }
   return;
 }
diff --git a/test/bug/tint/757.wgsl.expected.wgsl b/test/bug/tint/757.wgsl.expected.wgsl
index c7c0d8c..9690af5 100644
--- a/test/bug/tint/757.wgsl.expected.wgsl
+++ b/test/bug/tint/757.wgsl.expected.wgsl
@@ -19,17 +19,7 @@
   var flatIndex : u32 = ((((2u * 2u) * GlobalInvocationID.z) + (2u * GlobalInvocationID.y)) + GlobalInvocationID.x);
   flatIndex = (flatIndex * 1u);
   var texel : vec4<f32> = textureLoad(myTexture, vec2<i32>(GlobalInvocationID.xy), 0, 0);
-  {
-    var i : u32 = 0u;
-    loop {
-      if (!((i < 1u))) {
-        break;
-      }
-      result.values[(flatIndex + i)] = texel.r;
-
-      continuing {
-        i = (i + 1u);
-      }
-    }
+  for(var i : u32 = 0u; (i < 1u); i = (i + 1u)) {
+    result.values[(flatIndex + i)] = texel.r;
   }
 }
diff --git a/test/bug/tint/914.wgsl.expected.hlsl b/test/bug/tint/914.wgsl.expected.hlsl
index 3de9408..5e0ebb4 100644
--- a/test/bug/tint/914.wgsl.expected.hlsl
+++ b/test/bug/tint/914.wgsl.expected.hlsl
@@ -69,14 +69,22 @@
   const uint3 global_id = tint_symbol.global_id;
   const uint local_invocation_index = tint_symbol.local_invocation_index;
   if ((local_invocation_index == 0u)) {
-    for(int i = 0; (i < 64); i = (i + 1)) {
-      for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
-        mm_Asub[i][i_1] = 0.0f;
+    {
+      for(int i = 0; (i < 64); i = (i + 1)) {
+        {
+          for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
+            mm_Asub[i][i_1] = 0.0f;
+          }
+        }
       }
     }
-    for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
-      for(int i_3 = 0; (i_3 < 64); i_3 = (i_3 + 1)) {
-        mm_Bsub[i_2][i_3] = 0.0f;
+    {
+      for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
+        {
+          for(int i_3 = 0; (i_3 < 64); i_3 = (i_3 + 1)) {
+            mm_Bsub[i_2][i_3] = 0.0f;
+          }
+        }
       }
     }
   }
@@ -91,8 +99,7 @@
   float ACached = 0.0f;
   float BCached[4] = (float[4])0;
   {
-    uint index = 0u;
-    for(; !(!((index < (RowPerThread * ColPerThread)))); index = (index + 1u)) {
+    for(uint index = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
       acc[index] = 0.0f;
     }
   }
@@ -101,14 +108,11 @@
   const uint RowPerThreadB = (TileInner / 16u);
   const uint tileRowB = (local_id.y * RowPerThreadB);
   {
-    uint t = 0u;
-    for(; !(!((t < numTiles))); t = (t + 1u)) {
+    for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
       {
-        uint innerRow = 0u;
-        for(; !(!((innerRow < RowPerThread))); innerRow = (innerRow + 1u)) {
+        for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
           {
-            uint innerCol = 0u;
-            for(; !(!((innerCol < ColPerThreadA))); innerCol = (innerCol + 1u)) {
+            for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
               const uint inputRow = (tileRow + innerRow);
               const uint inputCol = (tileColA + innerCol);
               mm_Asub[inputRow][inputCol] = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
@@ -117,11 +121,9 @@
         }
       }
       {
-        uint innerRow = 0u;
-        for(; !(!((innerRow < RowPerThreadB))); innerRow = (innerRow + 1u)) {
+        for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
           {
-            uint innerCol = 0u;
-            for(; !(!((innerCol < ColPerThread))); innerCol = (innerCol + 1u)) {
+            for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
               const uint inputRow = (tileRowB + innerRow);
               const uint inputCol = (tileCol + innerCol);
               mm_Bsub[innerCol][inputCol] = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
@@ -131,21 +133,17 @@
       }
       GroupMemoryBarrierWithGroupSync();
       {
-        uint k = 0u;
-        for(; !(!((k < TileInner))); k = (k + 1u)) {
+        for(uint k = 0u; (k < TileInner); k = (k + 1u)) {
           {
-            uint inner = 0u;
-            for(; !(!((inner < ColPerThread))); inner = (inner + 1u)) {
+            for(uint inner = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
               BCached[inner] = mm_Bsub[k][(tileCol + inner)];
             }
           }
           {
-            uint innerRow = 0u;
-            for(; !(!((innerRow < RowPerThread))); innerRow = (innerRow + 1u)) {
+            for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
               ACached = mm_Asub[(tileRow + innerRow)][k];
               {
-                uint innerCol = 0u;
-                for(; !(!((innerCol < ColPerThread))); innerCol = (innerCol + 1u)) {
+                for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
                   const uint index = ((innerRow * ColPerThread) + innerCol);
                   acc[index] = (acc[index] + (ACached * BCached[innerCol]));
                 }
@@ -158,11 +156,9 @@
     }
   }
   {
-    uint innerRow = 0u;
-    for(; !(!((innerRow < RowPerThread))); innerRow = (innerRow + 1u)) {
+    for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
       {
-        uint innerCol = 0u;
-        for(; !(!((innerCol < ColPerThread))); innerCol = (innerCol + 1u)) {
+        for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
           const uint index = ((innerRow * ColPerThread) + innerCol);
           mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index]);
         }
diff --git a/test/bug/tint/914.wgsl.expected.msl b/test/bug/tint/914.wgsl.expected.msl
index d9a4b82..eb27d80 100644
--- a/test/bug/tint/914.wgsl.expected.msl
+++ b/test/bug/tint/914.wgsl.expected.msl
@@ -68,155 +68,47 @@
   tint_array_wrapper_2 acc = {};
   float ACached = 0.0f;
   tint_array_wrapper_3 BCached = {};
-  {
-    uint index = 0u;
-    while (true) {
-      if (!((index < (RowPerThread * ColPerThread)))) {
-        break;
-      }
-      acc.arr[index] = 0.0f;
-      {
-        index = (index + 1u);
-      }
-    }
+  for(uint index = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
+    acc.arr[index] = 0.0f;
   }
   uint const ColPerThreadA = (TileInner / 16u);
   uint const tileColA = (local_id.x * ColPerThreadA);
   uint const RowPerThreadB = (TileInner / 16u);
   uint const tileRowB = (local_id.y * RowPerThreadB);
-  {
-    uint t = 0u;
-    while (true) {
-      if (!((t < numTiles))) {
-        break;
-      }
-      {
-        uint innerRow = 0u;
-        while (true) {
-          if (!((innerRow < RowPerThread))) {
-            break;
-          }
-          {
-            uint innerCol = 0u;
-            while (true) {
-              if (!((innerCol < ColPerThreadA))) {
-                break;
-              }
-              uint const inputRow = (tileRow + innerRow);
-              uint const inputCol = (tileColA + innerCol);
-              tint_symbol_4.arr[inputRow].arr[inputCol] = mm_readA(uniforms, firstMatrix, (globalRow + innerRow), ((t * TileInner) + inputCol));
-              {
-                innerCol = (innerCol + 1u);
-              }
-            }
-          }
-          {
-            innerRow = (innerRow + 1u);
-          }
-        }
-      }
-      {
-        uint innerRow = 0u;
-        while (true) {
-          if (!((innerRow < RowPerThreadB))) {
-            break;
-          }
-          {
-            uint innerCol = 0u;
-            while (true) {
-              if (!((innerCol < ColPerThread))) {
-                break;
-              }
-              uint const inputRow = (tileRowB + innerRow);
-              uint const inputCol = (tileCol + innerCol);
-              tint_symbol_5.arr[innerCol].arr[inputCol] = mm_readB(uniforms, secondMatrix, ((t * TileInner) + inputRow), (globalCol + innerCol));
-              {
-                innerCol = (innerCol + 1u);
-              }
-            }
-          }
-          {
-            innerRow = (innerRow + 1u);
-          }
-        }
-      }
-      threadgroup_barrier(mem_flags::mem_threadgroup);
-      {
-        uint k = 0u;
-        while (true) {
-          if (!((k < TileInner))) {
-            break;
-          }
-          {
-            uint inner = 0u;
-            while (true) {
-              if (!((inner < ColPerThread))) {
-                break;
-              }
-              BCached.arr[inner] = tint_symbol_5.arr[k].arr[(tileCol + inner)];
-              {
-                inner = (inner + 1u);
-              }
-            }
-          }
-          {
-            uint innerRow = 0u;
-            while (true) {
-              if (!((innerRow < RowPerThread))) {
-                break;
-              }
-              ACached = tint_symbol_4.arr[(tileRow + innerRow)].arr[k];
-              {
-                uint innerCol = 0u;
-                while (true) {
-                  if (!((innerCol < ColPerThread))) {
-                    break;
-                  }
-                  uint const index = ((innerRow * ColPerThread) + innerCol);
-                  acc.arr[index] = (acc.arr[index] + (ACached * BCached.arr[innerCol]));
-                  {
-                    innerCol = (innerCol + 1u);
-                  }
-                }
-              }
-              {
-                innerRow = (innerRow + 1u);
-              }
-            }
-          }
-          {
-            k = (k + 1u);
-          }
-        }
-      }
-      threadgroup_barrier(mem_flags::mem_threadgroup);
-      {
-        t = (t + 1u);
+  for(uint t = 0u; (t < numTiles); t = (t + 1u)) {
+    for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+      for(uint innerCol = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
+        uint const inputRow = (tileRow + innerRow);
+        uint const inputCol = (tileColA + innerCol);
+        tint_symbol_4.arr[inputRow].arr[inputCol] = mm_readA(uniforms, firstMatrix, (globalRow + innerRow), ((t * TileInner) + inputCol));
       }
     }
-  }
-  {
-    uint innerRow = 0u;
-    while (true) {
-      if (!((innerRow < RowPerThread))) {
-        break;
+    for(uint innerRow = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
+      for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
+        uint const inputRow = (tileRowB + innerRow);
+        uint const inputCol = (tileCol + innerCol);
+        tint_symbol_5.arr[innerCol].arr[inputCol] = mm_readB(uniforms, secondMatrix, ((t * TileInner) + inputRow), (globalCol + innerCol));
       }
-      {
-        uint innerCol = 0u;
-        while (true) {
-          if (!((innerCol < ColPerThread))) {
-            break;
-          }
+    }
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+    for(uint k = 0u; (k < TileInner); k = (k + 1u)) {
+      for(uint inner = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
+        BCached.arr[inner] = tint_symbol_5.arr[k].arr[(tileCol + inner)];
+      }
+      for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+        ACached = tint_symbol_4.arr[(tileRow + innerRow)].arr[k];
+        for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
           uint const index = ((innerRow * ColPerThread) + innerCol);
-          mm_write(uniforms, resultMatrix, (globalRow + innerRow), (globalCol + innerCol), acc.arr[index]);
-          {
-            innerCol = (innerCol + 1u);
-          }
+          acc.arr[index] = (acc.arr[index] + (ACached * BCached.arr[innerCol]));
         }
       }
-      {
-        innerRow = (innerRow + 1u);
-      }
+    }
+    threadgroup_barrier(mem_flags::mem_threadgroup);
+  }
+  for(uint innerRow = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+    for(uint innerCol = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
+      uint const index = ((innerRow * ColPerThread) + innerCol);
+      mm_write(uniforms, resultMatrix, (globalRow + innerRow), (globalCol + innerCol), acc.arr[index]);
     }
   }
   return;
diff --git a/test/bug/tint/914.wgsl.expected.wgsl b/test/bug/tint/914.wgsl.expected.wgsl
index ed45fab..17ca3f3 100644
--- a/test/bug/tint/914.wgsl.expected.wgsl
+++ b/test/bug/tint/914.wgsl.expected.wgsl
@@ -65,167 +65,47 @@
   var acc : array<f32, 16>;
   var ACached : f32;
   var BCached : array<f32, 4>;
-  {
-    var index : u32 = 0u;
-    loop {
-      if (!((index < (RowPerThread * ColPerThread)))) {
-        break;
-      }
-      acc[index] = 0.0;
-
-      continuing {
-        index = (index + 1u);
-      }
-    }
+  for(var index : u32 = 0u; (index < (RowPerThread * ColPerThread)); index = (index + 1u)) {
+    acc[index] = 0.0;
   }
   let ColPerThreadA : u32 = (TileInner / 16u);
   let tileColA : u32 = (local_id.x * ColPerThreadA);
   let RowPerThreadB : u32 = (TileInner / 16u);
   let tileRowB : u32 = (local_id.y * RowPerThreadB);
-  {
-    var t : u32 = 0u;
-    loop {
-      if (!((t < numTiles))) {
-        break;
-      }
-      {
-        var innerRow : u32 = 0u;
-        loop {
-          if (!((innerRow < RowPerThread))) {
-            break;
-          }
-          {
-            var innerCol : u32 = 0u;
-            loop {
-              if (!((innerCol < ColPerThreadA))) {
-                break;
-              }
-              let inputRow : u32 = (tileRow + innerRow);
-              let inputCol : u32 = (tileColA + innerCol);
-              mm_Asub[inputRow][inputCol] = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
-
-              continuing {
-                innerCol = (innerCol + 1u);
-              }
-            }
-          }
-
-          continuing {
-            innerRow = (innerRow + 1u);
-          }
-        }
-      }
-      {
-        var innerRow : u32 = 0u;
-        loop {
-          if (!((innerRow < RowPerThreadB))) {
-            break;
-          }
-          {
-            var innerCol : u32 = 0u;
-            loop {
-              if (!((innerCol < ColPerThread))) {
-                break;
-              }
-              let inputRow : u32 = (tileRowB + innerRow);
-              let inputCol : u32 = (tileCol + innerCol);
-              mm_Bsub[innerCol][inputCol] = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
-
-              continuing {
-                innerCol = (innerCol + 1u);
-              }
-            }
-          }
-
-          continuing {
-            innerRow = (innerRow + 1u);
-          }
-        }
-      }
-      workgroupBarrier();
-      {
-        var k : u32 = 0u;
-        loop {
-          if (!((k < TileInner))) {
-            break;
-          }
-          {
-            var inner : u32 = 0u;
-            loop {
-              if (!((inner < ColPerThread))) {
-                break;
-              }
-              BCached[inner] = mm_Bsub[k][(tileCol + inner)];
-
-              continuing {
-                inner = (inner + 1u);
-              }
-            }
-          }
-          {
-            var innerRow : u32 = 0u;
-            loop {
-              if (!((innerRow < RowPerThread))) {
-                break;
-              }
-              ACached = mm_Asub[(tileRow + innerRow)][k];
-              {
-                var innerCol : u32 = 0u;
-                loop {
-                  if (!((innerCol < ColPerThread))) {
-                    break;
-                  }
-                  let index : u32 = ((innerRow * ColPerThread) + innerCol);
-                  acc[index] = (acc[index] + (ACached * BCached[innerCol]));
-
-                  continuing {
-                    innerCol = (innerCol + 1u);
-                  }
-                }
-              }
-
-              continuing {
-                innerRow = (innerRow + 1u);
-              }
-            }
-          }
-
-          continuing {
-            k = (k + 1u);
-          }
-        }
-      }
-      workgroupBarrier();
-
-      continuing {
-        t = (t + 1u);
+  for(var t : u32 = 0u; (t < numTiles); t = (t + 1u)) {
+    for(var innerRow : u32 = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+      for(var innerCol : u32 = 0u; (innerCol < ColPerThreadA); innerCol = (innerCol + 1u)) {
+        let inputRow : u32 = (tileRow + innerRow);
+        let inputCol : u32 = (tileColA + innerCol);
+        mm_Asub[inputRow][inputCol] = mm_readA((globalRow + innerRow), ((t * TileInner) + inputCol));
       }
     }
-  }
-  {
-    var innerRow : u32 = 0u;
-    loop {
-      if (!((innerRow < RowPerThread))) {
-        break;
+    for(var innerRow : u32 = 0u; (innerRow < RowPerThreadB); innerRow = (innerRow + 1u)) {
+      for(var innerCol : u32 = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
+        let inputRow : u32 = (tileRowB + innerRow);
+        let inputCol : u32 = (tileCol + innerCol);
+        mm_Bsub[innerCol][inputCol] = mm_readB(((t * TileInner) + inputRow), (globalCol + innerCol));
       }
-      {
-        var innerCol : u32 = 0u;
-        loop {
-          if (!((innerCol < ColPerThread))) {
-            break;
-          }
+    }
+    workgroupBarrier();
+    for(var k : u32 = 0u; (k < TileInner); k = (k + 1u)) {
+      for(var inner : u32 = 0u; (inner < ColPerThread); inner = (inner + 1u)) {
+        BCached[inner] = mm_Bsub[k][(tileCol + inner)];
+      }
+      for(var innerRow : u32 = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+        ACached = mm_Asub[(tileRow + innerRow)][k];
+        for(var innerCol : u32 = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
           let index : u32 = ((innerRow * ColPerThread) + innerCol);
-          mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index]);
-
-          continuing {
-            innerCol = (innerCol + 1u);
-          }
+          acc[index] = (acc[index] + (ACached * BCached[innerCol]));
         }
       }
-
-      continuing {
-        innerRow = (innerRow + 1u);
-      }
+    }
+    workgroupBarrier();
+  }
+  for(var innerRow : u32 = 0u; (innerRow < RowPerThread); innerRow = (innerRow + 1u)) {
+    for(var innerCol : u32 = 0u; (innerCol < ColPerThread); innerCol = (innerCol + 1u)) {
+      let index : u32 = ((innerRow * ColPerThread) + innerCol);
+      mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index]);
     }
   }
 }
diff --git a/test/bug/tint/942.wgsl.expected.hlsl b/test/bug/tint/942.wgsl.expected.hlsl
index ff2739f..695e8ae 100644
--- a/test/bug/tint/942.wgsl.expected.hlsl
+++ b/test/bug/tint/942.wgsl.expected.hlsl
@@ -22,9 +22,13 @@
   const uint3 LocalInvocationID = tint_symbol.LocalInvocationID;
   const uint local_invocation_index = tint_symbol.local_invocation_index;
   if ((local_invocation_index == 0u)) {
-    for(int i_1 = 0; (i_1 < 4); i_1 = (i_1 + 1)) {
-      for(int i_2 = 0; (i_2 < 256); i_2 = (i_2 + 1)) {
-        tile[i_1][i_2] = float3(0.0f, 0.0f, 0.0f);
+    {
+      for(int i_1 = 0; (i_1 < 4); i_1 = (i_1 + 1)) {
+        {
+          for(int i_2 = 0; (i_2 < 256); i_2 = (i_2 + 1)) {
+            tile[i_1][i_2] = float3(0.0f, 0.0f, 0.0f);
+          }
+        }
       }
     }
   }
@@ -37,11 +41,9 @@
   const uint scalar_offset_1 = (4u) / 4;
   const int2 baseIndex = (int2(((WorkGroupID.xy * uint2(params[scalar_offset_1 / 4][scalar_offset_1 % 4], 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
   {
-    uint r = 0u;
-    for(; !(!((r < 4u))); r = (r + 1u)) {
+    for(uint r = 0u; (r < 4u); r = (r + 1u)) {
       {
-        uint c = 0u;
-        for(; !(!((c < 4u))); c = (c + 1u)) {
+        for(uint c = 0u; (c < 4u); c = (c + 1u)) {
           int2 loadIndex = (baseIndex + int2(int(c), int(r)));
           const uint scalar_offset_2 = (0u) / 4;
           if ((flip[scalar_offset_2 / 4][scalar_offset_2 % 4] != 0u)) {
@@ -54,11 +56,9 @@
   }
   GroupMemoryBarrierWithGroupSync();
   {
-    uint r = 0u;
-    for(; !(!((r < 4u))); r = (r + 1u)) {
+    for(uint r = 0u; (r < 4u); r = (r + 1u)) {
       {
-        uint c = 0u;
-        for(; !(!((c < 4u))); c = (c + 1u)) {
+        for(uint c = 0u; (c < 4u); c = (c + 1u)) {
           int2 writeIndex = (baseIndex + int2(int(c), int(r)));
           const uint scalar_offset_3 = (0u) / 4;
           if ((flip[scalar_offset_3 / 4][scalar_offset_3 % 4] != 0u)) {
@@ -79,7 +79,7 @@
               uint f = 0u;
               while (true) {
                 const uint scalar_offset_4 = (0u) / 4;
-                if (!(!(!((f < params[scalar_offset_4 / 4][scalar_offset_4 % 4]))))) { break; }
+                if (!((f < params[scalar_offset_4 / 4][scalar_offset_4 % 4]))) { break; }
                 uint i = ((center + f) - filterOffset);
                 const uint scalar_offset_5 = (0u) / 4;
                 acc = (acc + ((1.0f / float(params[scalar_offset_5 / 4][scalar_offset_5 % 4])) * tile[r][i]));
diff --git a/test/bug/tint/942.wgsl.expected.msl b/test/bug/tint/942.wgsl.expected.msl
index 16e440d..cecfa13 100644
--- a/test/bug/tint/942.wgsl.expected.msl
+++ b/test/bug/tint/942.wgsl.expected.msl
@@ -25,75 +25,30 @@
   uint const filterOffset = ((params.filterDim - 1u) / 2u);
   int2 const dims = int2(tint_symbol_4.get_width(0), tint_symbol_4.get_height(0));
   int2 const baseIndex = (int2(((WorkGroupID.xy * uint2(params.blockDim, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
-  {
-    uint r = 0u;
-    while (true) {
-      if (!((r < 4u))) {
-        break;
+  for(uint r = 0u; (r < 4u); r = (r + 1u)) {
+    for(uint c = 0u; (c < 4u); c = (c + 1u)) {
+      int2 loadIndex = (baseIndex + int2(int(c), int(r)));
+      if ((flip.value != 0u)) {
+        loadIndex = loadIndex.yx;
       }
-      {
-        uint c = 0u;
-        while (true) {
-          if (!((c < 4u))) {
-            break;
-          }
-          int2 loadIndex = (baseIndex + int2(int(c), int(r)));
-          if ((flip.value != 0u)) {
-            loadIndex = loadIndex.yx;
-          }
-          tint_symbol_3.arr[r].arr[((4u * LocalInvocationID.x) + c)] = tint_symbol_4.sample(tint_symbol_5, ((float2(loadIndex) + float2(0.25f, 0.25f)) / float2(dims)), level(0.0f)).rgb;
-          {
-            c = (c + 1u);
-          }
-        }
-      }
-      {
-        r = (r + 1u);
-      }
+      tint_symbol_3.arr[r].arr[((4u * LocalInvocationID.x) + c)] = tint_symbol_4.sample(tint_symbol_5, ((float2(loadIndex) + float2(0.25f, 0.25f)) / float2(dims)), level(0.0f)).rgb;
     }
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  {
-    uint r = 0u;
-    while (true) {
-      if (!((r < 4u))) {
-        break;
+  for(uint r = 0u; (r < 4u); r = (r + 1u)) {
+    for(uint c = 0u; (c < 4u); c = (c + 1u)) {
+      int2 writeIndex = (baseIndex + int2(int(c), int(r)));
+      if ((flip.value != 0u)) {
+        writeIndex = writeIndex.yx;
       }
-      {
-        uint c = 0u;
-        while (true) {
-          if (!((c < 4u))) {
-            break;
-          }
-          int2 writeIndex = (baseIndex + int2(int(c), int(r)));
-          if ((flip.value != 0u)) {
-            writeIndex = writeIndex.yx;
-          }
-          uint const center = ((4u * LocalInvocationID.x) + c);
-          if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
-            float3 acc = float3(0.0f, 0.0f, 0.0f);
-            {
-              uint f = 0u;
-              while (true) {
-                if (!((f < params.filterDim))) {
-                  break;
-                }
-                uint i = ((center + f) - filterOffset);
-                acc = (acc + ((1.0f / float(params.filterDim)) * tint_symbol_3.arr[r].arr[i]));
-                {
-                  f = (f + 1u);
-                }
-              }
-            }
-            tint_symbol_6.write(float4(acc, 1.0f), uint2(writeIndex));
-          }
-          {
-            c = (c + 1u);
-          }
+      uint const center = ((4u * LocalInvocationID.x) + c);
+      if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
+        float3 acc = float3(0.0f, 0.0f, 0.0f);
+        for(uint f = 0u; (f < params.filterDim); f = (f + 1u)) {
+          uint i = ((center + f) - filterOffset);
+          acc = (acc + ((1.0f / float(params.filterDim)) * tint_symbol_3.arr[r].arr[i]));
         }
-      }
-      {
-        r = (r + 1u);
+        tint_symbol_6.write(float4(acc, 1.0f), uint2(writeIndex));
       }
     }
   }
diff --git a/test/bug/tint/942.wgsl.expected.wgsl b/test/bug/tint/942.wgsl.expected.wgsl
index d261497..b936d4e 100644
--- a/test/bug/tint/942.wgsl.expected.wgsl
+++ b/test/bug/tint/942.wgsl.expected.wgsl
@@ -26,80 +26,30 @@
   let filterOffset : u32 = ((params.filterDim - 1u) / 2u);
   let dims : vec2<i32> = textureDimensions(inputTex, 0);
   let baseIndex = (vec2<i32>(((WorkGroupID.xy * vec2<u32>(params.blockDim, 4u)) + (LocalInvocationID.xy * vec2<u32>(4u, 1u)))) - vec2<i32>(i32(filterOffset), 0));
-  {
-    var r : u32 = 0u;
-    loop {
-      if (!((r < 4u))) {
-        break;
+  for(var r : u32 = 0u; (r < 4u); r = (r + 1u)) {
+    for(var c : u32 = 0u; (c < 4u); c = (c + 1u)) {
+      var loadIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
+      if ((flip.value != 0u)) {
+        loadIndex = loadIndex.yx;
       }
-      {
-        var c : u32 = 0u;
-        loop {
-          if (!((c < 4u))) {
-            break;
-          }
-          var loadIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
-          if ((flip.value != 0u)) {
-            loadIndex = loadIndex.yx;
-          }
-          tile[r][((4u * LocalInvocationID.x) + c)] = textureSampleLevel(inputTex, samp, ((vec2<f32>(loadIndex) + vec2<f32>(0.25, 0.25)) / vec2<f32>(dims)), 0.0).rgb;
-
-          continuing {
-            c = (c + 1u);
-          }
-        }
-      }
-
-      continuing {
-        r = (r + 1u);
-      }
+      tile[r][((4u * LocalInvocationID.x) + c)] = textureSampleLevel(inputTex, samp, ((vec2<f32>(loadIndex) + vec2<f32>(0.25, 0.25)) / vec2<f32>(dims)), 0.0).rgb;
     }
   }
   workgroupBarrier();
-  {
-    var r : u32 = 0u;
-    loop {
-      if (!((r < 4u))) {
-        break;
+  for(var r : u32 = 0u; (r < 4u); r = (r + 1u)) {
+    for(var c : u32 = 0u; (c < 4u); c = (c + 1u)) {
+      var writeIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
+      if ((flip.value != 0u)) {
+        writeIndex = writeIndex.yx;
       }
-      {
-        var c : u32 = 0u;
-        loop {
-          if (!((c < 4u))) {
-            break;
-          }
-          var writeIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
-          if ((flip.value != 0u)) {
-            writeIndex = writeIndex.yx;
-          }
-          let center : u32 = ((4u * LocalInvocationID.x) + c);
-          if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
-            var acc : vec3<f32> = vec3<f32>(0.0, 0.0, 0.0);
-            {
-              var f : u32 = 0u;
-              loop {
-                if (!((f < params.filterDim))) {
-                  break;
-                }
-                var i : u32 = ((center + f) - filterOffset);
-                acc = (acc + ((1.0 / f32(params.filterDim)) * tile[r][i]));
-
-                continuing {
-                  f = (f + 1u);
-                }
-              }
-            }
-            textureStore(outputTex, writeIndex, vec4<f32>(acc, 1.0));
-          }
-
-          continuing {
-            c = (c + 1u);
-          }
+      let center : u32 = ((4u * LocalInvocationID.x) + c);
+      if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
+        var acc : vec3<f32> = vec3<f32>(0.0, 0.0, 0.0);
+        for(var f : u32 = 0u; (f < params.filterDim); f = (f + 1u)) {
+          var i : u32 = ((center + f) - filterOffset);
+          acc = (acc + ((1.0 / f32(params.filterDim)) * tile[r][i]));
         }
-      }
-
-      continuing {
-        r = (r + 1u);
+        textureStore(outputTex, writeIndex, vec4<f32>(acc, 1.0));
       }
     }
   }
diff --git a/test/bug/tint/943.spvasm.expected.hlsl b/test/bug/tint/943.spvasm.expected.hlsl
index 47f46cd..a317e46 100644
--- a/test/bug/tint/943.spvasm.expected.hlsl
+++ b/test/bug/tint/943.spvasm.expected.hlsl
@@ -191,10 +191,14 @@
   const int x_152 = dimInner;
   numTiles = (((x_152 - 1) / 64) + 1);
   innerRow = 0;
-  for(; (innerRow < 1); innerRow = (innerRow + 1)) {
-    innerCol = 0;
-    for(; (innerCol < 1); innerCol = (innerCol + 1)) {
-      acc[innerRow][innerCol] = 0.0f;
+  {
+    for(; (innerRow < 1); innerRow = (innerRow + 1)) {
+      innerCol = 0;
+      {
+        for(; (innerCol < 1); innerCol = (innerCol + 1)) {
+          acc[innerRow][innerCol] = 0.0f;
+        }
+      }
     }
   }
   const uint x_187 = gl_LocalInvocationID.x;
@@ -202,100 +206,120 @@
   const uint x_192 = gl_LocalInvocationID.y;
   tileRowB = (asint(x_192) * 1);
   t = 0;
-  for(; (t < numTiles); t = (t + 1)) {
-    innerRow_1 = 0;
-    for(; (innerRow_1 < 1); innerRow_1 = (innerRow_1 + 1)) {
-      innerCol_1 = 0;
-      for(; (innerCol_1 < 64); innerCol_1 = (innerCol_1 + 1)) {
-        inputRow = (tileRow + innerRow_1);
-        inputCol = (tileColA + innerCol_1);
-        const int x_233 = inputRow;
-        const int x_234 = inputCol;
-        const int x_238 = t;
-        const int x_240 = inputCol;
-        param_3 = (globalRow + innerRow_1);
-        param_4 = ((x_238 * 64) + x_240);
-        const float x_244 = mm_readA_i1_i1_(param_3, param_4);
-        mm_Asub[x_233][x_234] = x_244;
-      }
-    }
-    innerRow_2 = 0;
-    for(; (innerRow_2 < 1); innerRow_2 = (innerRow_2 + 1)) {
-      innerCol_2 = 0;
-      for(; (innerCol_2 < 1); innerCol_2 = (innerCol_2 + 1)) {
-        inputRow_1 = (tileRowB + innerRow_2);
-        inputCol_1 = (tileCol + innerCol_2);
-        const int x_278 = inputRow_1;
-        const int x_279 = inputCol_1;
-        const int x_284 = globalCol;
-        const int x_285 = innerCol_2;
-        param_5 = ((t * 64) + inputRow_1);
-        param_6 = (x_284 + x_285);
-        const float x_289 = mm_readB_i1_i1_(param_5, param_6);
-        mm_Bsub[x_278][x_279] = x_289;
-      }
-    }
-    GroupMemoryBarrierWithGroupSync();
-    k = 0;
-    for(; (k < 64); k = (k + 1)) {
-      inner = 0;
-      for(; (inner < 1); inner = (inner + 1)) {
-        const int x_314 = inner;
-        const float x_320 = mm_Bsub[k][(tileCol + inner)];
-        BCached[x_314] = x_320;
-      }
-      innerRow_3 = 0;
-      for(; (innerRow_3 < 1); innerRow_3 = (innerRow_3 + 1)) {
-        const float x_338 = mm_Asub[(tileRow + innerRow_3)][k];
-        ACached = x_338;
-        innerCol_3 = 0;
-        for(; (innerCol_3 < 1); innerCol_3 = (innerCol_3 + 1)) {
-          const int x_347 = innerRow_3;
-          const int x_348 = innerCol_3;
-          const float x_349 = ACached;
-          const float x_352 = BCached[innerCol_3];
-          const float x_355 = acc[x_347][x_348];
-          acc[x_347][x_348] = (x_355 + (x_349 * x_352));
+  {
+    for(; (t < numTiles); t = (t + 1)) {
+      innerRow_1 = 0;
+      {
+        for(; (innerRow_1 < 1); innerRow_1 = (innerRow_1 + 1)) {
+          innerCol_1 = 0;
+          {
+            for(; (innerCol_1 < 64); innerCol_1 = (innerCol_1 + 1)) {
+              inputRow = (tileRow + innerRow_1);
+              inputCol = (tileColA + innerCol_1);
+              const int x_233 = inputRow;
+              const int x_234 = inputCol;
+              const int x_238 = t;
+              const int x_240 = inputCol;
+              param_3 = (globalRow + innerRow_1);
+              param_4 = ((x_238 * 64) + x_240);
+              const float x_244 = mm_readA_i1_i1_(param_3, param_4);
+              mm_Asub[x_233][x_234] = x_244;
+            }
+          }
         }
       }
+      innerRow_2 = 0;
+      {
+        for(; (innerRow_2 < 1); innerRow_2 = (innerRow_2 + 1)) {
+          innerCol_2 = 0;
+          {
+            for(; (innerCol_2 < 1); innerCol_2 = (innerCol_2 + 1)) {
+              inputRow_1 = (tileRowB + innerRow_2);
+              inputCol_1 = (tileCol + innerCol_2);
+              const int x_278 = inputRow_1;
+              const int x_279 = inputCol_1;
+              const int x_284 = globalCol;
+              const int x_285 = innerCol_2;
+              param_5 = ((t * 64) + inputRow_1);
+              param_6 = (x_284 + x_285);
+              const float x_289 = mm_readB_i1_i1_(param_5, param_6);
+              mm_Bsub[x_278][x_279] = x_289;
+            }
+          }
+        }
+      }
+      GroupMemoryBarrierWithGroupSync();
+      k = 0;
+      {
+        for(; (k < 64); k = (k + 1)) {
+          inner = 0;
+          {
+            for(; (inner < 1); inner = (inner + 1)) {
+              const int x_314 = inner;
+              const float x_320 = mm_Bsub[k][(tileCol + inner)];
+              BCached[x_314] = x_320;
+            }
+          }
+          innerRow_3 = 0;
+          {
+            for(; (innerRow_3 < 1); innerRow_3 = (innerRow_3 + 1)) {
+              const float x_338 = mm_Asub[(tileRow + innerRow_3)][k];
+              ACached = x_338;
+              innerCol_3 = 0;
+              {
+                for(; (innerCol_3 < 1); innerCol_3 = (innerCol_3 + 1)) {
+                  const int x_347 = innerRow_3;
+                  const int x_348 = innerCol_3;
+                  const float x_349 = ACached;
+                  const float x_352 = BCached[innerCol_3];
+                  const float x_355 = acc[x_347][x_348];
+                  acc[x_347][x_348] = (x_355 + (x_349 * x_352));
+                }
+              }
+            }
+          }
+        }
+      }
+      GroupMemoryBarrierWithGroupSync();
     }
-    GroupMemoryBarrierWithGroupSync();
   }
   innerRow_4 = 0;
-  for(; (innerRow_4 < 1); innerRow_4 = (innerRow_4 + 1)) {
-    innerCol_4 = 0;
-    while (true) {
-      bool x_393 = false;
-      bool x_394_phi = false;
-      if ((innerCol_4 < 1)) {
-      } else {
-        break;
-      }
-      const int x_382 = globalCol;
-      const int x_383 = innerCol_4;
-      const int x_385 = dimBOuter;
-      const bool x_386 = ((x_382 + x_383) < x_385);
-      x_394_phi = x_386;
-      if (x_386) {
-        const int x_389 = globalRow;
-        const int x_390 = innerRow_4;
-        const int x_392 = dimAOuter;
-        x_393 = ((x_389 + x_390) < x_392);
-        x_394_phi = x_393;
-      }
-      if (x_394_phi) {
-        const int x_400 = globalCol;
-        const int x_401 = innerCol_4;
-        const int x_403 = innerRow_4;
-        const int x_404 = innerCol_4;
-        param_7 = (globalRow + innerRow_4);
-        param_8 = (x_400 + x_401);
-        const float x_409 = acc[x_403][x_404];
-        param_9 = x_409;
-        mm_write_i1_i1_f1_(param_7, param_8, param_9);
-      }
-      {
-        innerCol_4 = (innerCol_4 + 1);
+  {
+    for(; (innerRow_4 < 1); innerRow_4 = (innerRow_4 + 1)) {
+      innerCol_4 = 0;
+      while (true) {
+        bool x_393 = false;
+        bool x_394_phi = false;
+        if ((innerCol_4 < 1)) {
+        } else {
+          break;
+        }
+        const int x_382 = globalCol;
+        const int x_383 = innerCol_4;
+        const int x_385 = dimBOuter;
+        const bool x_386 = ((x_382 + x_383) < x_385);
+        x_394_phi = x_386;
+        if (x_386) {
+          const int x_389 = globalRow;
+          const int x_390 = innerRow_4;
+          const int x_392 = dimAOuter;
+          x_393 = ((x_389 + x_390) < x_392);
+          x_394_phi = x_393;
+        }
+        if (x_394_phi) {
+          const int x_400 = globalCol;
+          const int x_401 = innerCol_4;
+          const int x_403 = innerRow_4;
+          const int x_404 = innerCol_4;
+          param_7 = (globalRow + innerRow_4);
+          param_8 = (x_400 + x_401);
+          const float x_409 = acc[x_403][x_404];
+          param_9 = x_409;
+          mm_write_i1_i1_f1_(param_7, param_8, param_9);
+        }
+        {
+          innerCol_4 = (innerCol_4 + 1);
+        }
       }
     }
   }
@@ -336,14 +360,20 @@
   const uint3 gl_GlobalInvocationID_param = tint_symbol.gl_GlobalInvocationID_param;
   const uint local_invocation_index = tint_symbol.local_invocation_index;
   if ((local_invocation_index == 0u)) {
-    for(int i = 0; (i < 64); i = (i + 1)) {
-      for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
-        mm_Asub[i][i_1] = 0.0f;
+    {
+      for(int i = 0; (i < 64); i = (i + 1)) {
+        {
+          for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
+            mm_Asub[i][i_1] = 0.0f;
+          }
+        }
       }
     }
-    for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
-      const float tint_symbol_6[1] = (float[1])0;
-      mm_Bsub[i_2] = tint_symbol_6;
+    {
+      for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
+        const float tint_symbol_6[1] = (float[1])0;
+        mm_Bsub[i_2] = tint_symbol_6;
+      }
     }
   }
   GroupMemoryBarrierWithGroupSync();
diff --git a/test/bug/tint/948.wgsl.expected.hlsl b/test/bug/tint/948.wgsl.expected.hlsl
index 2454c3b..d6f3a02 100644
--- a/test/bug/tint/948.wgsl.expected.hlsl
+++ b/test/bug/tint/948.wgsl.expected.hlsl
@@ -67,89 +67,93 @@
   const float2 x_111 = asfloat(((scalar_offset_3 & 2) ? ubo_load_1.zw : ubo_load_1.xy));
   stageUnits = (float2(1.0f, 1.0f) / x_111);
   i = 0;
-  for(; (i < 2); i = (i + 1)) {
-    switch(i) {
-      case 1: {
-        const float2 x_150 = tileID;
-        const uint scalar_offset_4 = (88u) / 4;
-        uint4 ubo_load_2 = x_20[scalar_offset_4 / 4];
-        const float2 x_154 = asfloat(((scalar_offset_4 & 2) ? ubo_load_2.zw : ubo_load_2.xy));
-        const float4 x_156 = tileMapsTexture1.SampleBias(tileMapsSampler, ((x_150 + float2(0.5f, 0.5f)) / x_154), 0.0f);
-        frameID_1 = x_156.x;
-        break;
-      }
-      case 0: {
-        const float2 x_136 = tileID;
-        const uint scalar_offset_5 = (88u) / 4;
-        uint4 ubo_load_3 = x_20[scalar_offset_5 / 4];
-        const float2 x_140 = asfloat(((scalar_offset_5 & 2) ? ubo_load_3.zw : ubo_load_3.xy));
-        const float4 x_142 = tileMapsTexture0.SampleBias(tileMapsSampler, ((x_136 + float2(0.5f, 0.5f)) / x_140), 0.0f);
-        frameID_1 = x_142.x;
-        break;
-      }
-      default: {
-        break;
-      }
-    }
-    const float x_166 = frameID_1;
-    const uint scalar_offset_6 = (108u) / 4;
-    const float x_169 = asfloat(x_20[scalar_offset_6 / 4][scalar_offset_6 % 4]);
-    const float4 x_172 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_166 + 0.5f) / x_169), 0.0f), 0.0f);
-    animationData = x_172;
-    const float x_174 = animationData.y;
-    if ((x_174 > 0.0f)) {
-      const uint scalar_offset_7 = (0u) / 4;
-      const float x_181 = asfloat(x_20[scalar_offset_7 / 4][scalar_offset_7 % 4]);
-      const float x_184 = animationData.z;
-      mt = ((x_181 * x_184) % 1.0f);
-      f = 0.0f;
-      for(; (f < 8.0f); f = (f + 1.0f)) {
-        const float x_197 = animationData.y;
-        if ((x_197 > mt)) {
-          const float x_203 = animationData.x;
-          frameID_1 = x_203;
+  {
+    for(; (i < 2); i = (i + 1)) {
+      switch(i) {
+        case 1: {
+          const float2 x_150 = tileID;
+          const uint scalar_offset_4 = (88u) / 4;
+          uint4 ubo_load_2 = x_20[scalar_offset_4 / 4];
+          const float2 x_154 = asfloat(((scalar_offset_4 & 2) ? ubo_load_2.zw : ubo_load_2.xy));
+          const float4 x_156 = tileMapsTexture1.SampleBias(tileMapsSampler, ((x_150 + float2(0.5f, 0.5f)) / x_154), 0.0f);
+          frameID_1 = x_156.x;
           break;
         }
-        const float x_208 = frameID_1;
-        const uint scalar_offset_8 = (108u) / 4;
-        const float x_211 = asfloat(x_20[scalar_offset_8 / 4][scalar_offset_8 % 4]);
-        const float4 x_217 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_208 + 0.5f) / x_211), (0.125f * f)), 0.0f);
-        animationData = x_217;
+        case 0: {
+          const float2 x_136 = tileID;
+          const uint scalar_offset_5 = (88u) / 4;
+          uint4 ubo_load_3 = x_20[scalar_offset_5 / 4];
+          const float2 x_140 = asfloat(((scalar_offset_5 & 2) ? ubo_load_3.zw : ubo_load_3.xy));
+          const float4 x_142 = tileMapsTexture0.SampleBias(tileMapsSampler, ((x_136 + float2(0.5f, 0.5f)) / x_140), 0.0f);
+          frameID_1 = x_142.x;
+          break;
+        }
+        default: {
+          break;
+        }
       }
-    }
-    param = (frameID_1 + 0.5f);
-    const float4x4 x_225 = getFrameData_f1_(param);
-    frameData = x_225;
-    const float4 x_228 = frameData[0];
-    const uint scalar_offset_9 = (96u) / 4;
-    uint4 ubo_load_4 = x_20[scalar_offset_9 / 4];
-    const float2 x_231 = asfloat(((scalar_offset_9 & 2) ? ubo_load_4.zw : ubo_load_4.xy));
-    frameSize = (float2(x_228.w, x_228.z) / x_231);
-    const float4 x_235 = frameData[0];
-    offset_1 = (float2(x_235.x, x_235.y) * sheetUnits);
-    const float4 x_241 = frameData[2];
-    const float4 x_244 = frameData[0];
-    ratio = (float2(x_241.x, x_241.y) / float2(x_244.w, x_244.z));
-    const float x_248 = frameData[2].z;
-    if ((x_248 == 1.0f)) {
-      const float2 x_252 = tileUV;
-      tileUV = float2(x_252.y, x_252.x);
-    }
-    if ((i == 0)) {
-      const float4 x_268 = spriteSheetTexture.Sample(spriteSheetSampler, ((tileUV * frameSize) + offset_1));
-      color = x_268;
-    } else {
-      const float4 x_279 = spriteSheetTexture.Sample(spriteSheetSampler, ((tileUV * frameSize) + offset_1));
-      nc = x_279;
-      const float x_283 = color.w;
-      const float x_285 = nc.w;
-      alpha = min((x_283 + x_285), 1.0f);
-      const float4 x_290 = color;
-      const float4 x_292 = nc;
-      const float x_295 = nc.w;
-      mixed = lerp(float3(x_290.x, x_290.y, x_290.z), float3(x_292.x, x_292.y, x_292.z), float3(x_295, x_295, x_295));
-      const float3 x_298 = mixed;
-      color = float4(x_298.x, x_298.y, x_298.z, alpha);
+      const float x_166 = frameID_1;
+      const uint scalar_offset_6 = (108u) / 4;
+      const float x_169 = asfloat(x_20[scalar_offset_6 / 4][scalar_offset_6 % 4]);
+      const float4 x_172 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_166 + 0.5f) / x_169), 0.0f), 0.0f);
+      animationData = x_172;
+      const float x_174 = animationData.y;
+      if ((x_174 > 0.0f)) {
+        const uint scalar_offset_7 = (0u) / 4;
+        const float x_181 = asfloat(x_20[scalar_offset_7 / 4][scalar_offset_7 % 4]);
+        const float x_184 = animationData.z;
+        mt = ((x_181 * x_184) % 1.0f);
+        f = 0.0f;
+        {
+          for(; (f < 8.0f); f = (f + 1.0f)) {
+            const float x_197 = animationData.y;
+            if ((x_197 > mt)) {
+              const float x_203 = animationData.x;
+              frameID_1 = x_203;
+              break;
+            }
+            const float x_208 = frameID_1;
+            const uint scalar_offset_8 = (108u) / 4;
+            const float x_211 = asfloat(x_20[scalar_offset_8 / 4][scalar_offset_8 % 4]);
+            const float4 x_217 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_208 + 0.5f) / x_211), (0.125f * f)), 0.0f);
+            animationData = x_217;
+          }
+        }
+      }
+      param = (frameID_1 + 0.5f);
+      const float4x4 x_225 = getFrameData_f1_(param);
+      frameData = x_225;
+      const float4 x_228 = frameData[0];
+      const uint scalar_offset_9 = (96u) / 4;
+      uint4 ubo_load_4 = x_20[scalar_offset_9 / 4];
+      const float2 x_231 = asfloat(((scalar_offset_9 & 2) ? ubo_load_4.zw : ubo_load_4.xy));
+      frameSize = (float2(x_228.w, x_228.z) / x_231);
+      const float4 x_235 = frameData[0];
+      offset_1 = (float2(x_235.x, x_235.y) * sheetUnits);
+      const float4 x_241 = frameData[2];
+      const float4 x_244 = frameData[0];
+      ratio = (float2(x_241.x, x_241.y) / float2(x_244.w, x_244.z));
+      const float x_248 = frameData[2].z;
+      if ((x_248 == 1.0f)) {
+        const float2 x_252 = tileUV;
+        tileUV = float2(x_252.y, x_252.x);
+      }
+      if ((i == 0)) {
+        const float4 x_268 = spriteSheetTexture.Sample(spriteSheetSampler, ((tileUV * frameSize) + offset_1));
+        color = x_268;
+      } else {
+        const float4 x_279 = spriteSheetTexture.Sample(spriteSheetSampler, ((tileUV * frameSize) + offset_1));
+        nc = x_279;
+        const float x_283 = color.w;
+        const float x_285 = nc.w;
+        alpha = min((x_283 + x_285), 1.0f);
+        const float4 x_290 = color;
+        const float4 x_292 = nc;
+        const float x_295 = nc.w;
+        mixed = lerp(float3(x_290.x, x_290.y, x_290.z), float3(x_292.x, x_292.y, x_292.z), float3(x_295, x_295, x_295));
+        const float3 x_298 = mixed;
+        color = float4(x_298.x, x_298.y, x_298.z, alpha);
+      }
     }
   }
   const uint scalar_offset_10 = (112u) / 4;
diff --git a/test/bug/tint/949.wgsl.expected.hlsl b/test/bug/tint/949.wgsl.expected.hlsl
index 926e020..2bb720a 100644
--- a/test/bug/tint/949.wgsl.expected.hlsl
+++ b/test/bug/tint/949.wgsl.expected.hlsl
@@ -249,20 +249,22 @@
   lastSampledHeight = 1.0f;
   currSampledHeight = 1.0f;
   i = 0;
-  for(; (i < 15); i = (i + 1)) {
-    const float4 x_397 = TextureSamplerTexture.Sample(TextureSamplerSampler, (v_uv + vCurrOffset));
-    currSampledHeight = x_397.w;
-    if ((currSampledHeight > currRayHeight)) {
-      delta1 = (currSampledHeight - currRayHeight);
-      delta2 = ((currRayHeight + stepSize) - lastSampledHeight);
-      ratio = (delta1 / (delta1 + delta2));
-      vCurrOffset = ((vLastOffset * ratio) + (vCurrOffset * (1.0f - ratio)));
-      break;
-    } else {
-      currRayHeight = (currRayHeight - stepSize);
-      vLastOffset = vCurrOffset;
-      vCurrOffset = (vCurrOffset + (vMaxOffset * stepSize));
-      lastSampledHeight = currSampledHeight;
+  {
+    for(; (i < 15); i = (i + 1)) {
+      const float4 x_397 = TextureSamplerTexture.Sample(TextureSamplerSampler, (v_uv + vCurrOffset));
+      currSampledHeight = x_397.w;
+      if ((currSampledHeight > currRayHeight)) {
+        delta1 = (currSampledHeight - currRayHeight);
+        delta2 = ((currRayHeight + stepSize) - lastSampledHeight);
+        ratio = (delta1 / (delta1 + delta2));
+        vCurrOffset = ((vLastOffset * ratio) + (vCurrOffset * (1.0f - ratio)));
+        break;
+      } else {
+        currRayHeight = (currRayHeight - stepSize);
+        vLastOffset = vCurrOffset;
+        vCurrOffset = (vCurrOffset + (vMaxOffset * stepSize));
+        lastSampledHeight = currSampledHeight;
+      }
     }
   }
   parallaxOcclusion_0 = vCurrOffset;
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl
index 6496cf5..a222773 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.hlsl
@@ -32,8 +32,7 @@
 
 void foo() {
   {
-    int i = 0;
-    for(; !(!((i < 2))); i = (i + 1)) {
+    for(int i = 0; (i < 2); i = (i + 1)) {
       Set_float2(v2f, i, 1.0f);
       Set_int3(v3i, i, 1);
       Set_uint4(v4u, i, 1u);
@@ -45,8 +44,7 @@
 [numthreads(1, 1, 1)]
 void main() {
   {
-    int i = 0;
-    for(; !(!((i < 2))); i = (i + 1)) {
+    for(int i = 0; (i < 2); i = (i + 1)) {
       foo();
     }
   }
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
index ab45d31..2ab266f 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl
@@ -2,20 +2,11 @@
 
 using namespace metal;
 void foo(thread float2* const tint_symbol_1, thread int3* const tint_symbol_2, thread uint4* const tint_symbol_3, thread bool2* const tint_symbol_4) {
-  {
-    int i = 0;
-    while (true) {
-      if (!((i < 2))) {
-        break;
-      }
-      (*(tint_symbol_1))[i] = 1.0f;
-      (*(tint_symbol_2))[i] = 1;
-      (*(tint_symbol_3))[i] = 1u;
-      (*(tint_symbol_4))[i] = true;
-      {
-        i = (i + 1);
-      }
-    }
+  for(int i = 0; (i < 2); i = (i + 1)) {
+    (*(tint_symbol_1))[i] = 1.0f;
+    (*(tint_symbol_2))[i] = 1;
+    (*(tint_symbol_3))[i] = 1u;
+    (*(tint_symbol_4))[i] = true;
   }
 }
 
@@ -24,17 +15,8 @@
   thread int3 tint_symbol_6 = 0;
   thread uint4 tint_symbol_7 = 0u;
   thread bool2 tint_symbol_8 = false;
-  {
-    int i = 0;
-    while (true) {
-      if (!((i < 2))) {
-        break;
-      }
-      foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
-      {
-        i = (i + 1);
-      }
-    }
+  for(int i = 0; (i < 2); i = (i + 1)) {
+    foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
   }
   return;
 }
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl
index bc302f1..def10cb 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.wgsl
@@ -7,37 +7,17 @@
 var<private> v2b : vec2<bool>;
 
 fn foo() {
-  {
-    var i : i32 = 0;
-    loop {
-      if (!((i < 2))) {
-        break;
-      }
-      v2f[i] = 1.0;
-      v3i[i] = 1;
-      v4u[i] = 1u;
-      v2b[i] = true;
-
-      continuing {
-        i = (i + 1);
-      }
-    }
+  for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+    v2f[i] = 1.0;
+    v3i[i] = 1;
+    v4u[i] = 1u;
+    v2b[i] = true;
   }
 }
 
 [[stage(compute), workgroup_size(1, 1, 1)]]
 fn main() {
-  {
-    var i : i32 = 0;
-    loop {
-      if (!((i < 2))) {
-        break;
-      }
-      foo();
-
-      continuing {
-        i = (i + 1);
-      }
-    }
+  for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+    foo();
   }
 }
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl
index 33a99b9..ea05754 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.hlsl
@@ -14,8 +14,7 @@
 [numthreads(1, 1, 1)]
 void main() {
   {
-    int i = 0;
-    for(; !(!((i < 2))); i = (i + 1)) {
+    for(int i = 0; (i < 2); i = (i + 1)) {
       foo();
     }
   }
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
index 2685e75..44a4fa0 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl
@@ -14,17 +14,8 @@
   thread int3 tint_symbol_6 = 0;
   thread uint4 tint_symbol_7 = 0u;
   thread bool2 tint_symbol_8 = false;
-  {
-    int i = 0;
-    while (true) {
-      if (!((i < 2))) {
-        break;
-      }
-      foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
-      {
-        i = (i + 1);
-      }
-    }
+  for(int i = 0; (i < 2); i = (i + 1)) {
+    foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
   }
   return;
 }
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl
index ae1667c4..21db9dd 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.wgsl
@@ -16,17 +16,7 @@
 
 [[stage(compute), workgroup_size(1, 1, 1)]]
 fn main() {
-  {
-    var i : i32 = 0;
-    loop {
-      if (!((i < 2))) {
-        break;
-      }
-      foo();
-
-      continuing {
-        i = (i + 1);
-      }
-    }
+  for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+    foo();
   }
 }
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl
index c8fa518..85ffc25 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.hlsl
@@ -97,8 +97,7 @@
   bool3 v3b = bool3(false, false, false);
   bool4 v4b = bool4(false, false, false, false);
   {
-    int i = 0;
-    for(; !(!((i < 2))); i = (i + 1)) {
+    for(int i = 0; (i < 2); i = (i + 1)) {
       Set_float2(v2f, i, 1.0f);
       Set_float3(v3f, i, 1.0f);
       Set_float4(v4f, i, 1.0f);
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
index a460250..b622316 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl
@@ -14,28 +14,19 @@
   bool2 v2b = false;
   bool3 v3b = false;
   bool4 v4b = false;
-  {
-    int i = 0;
-    while (true) {
-      if (!((i < 2))) {
-        break;
-      }
-      v2f[i] = 1.0f;
-      v3f[i] = 1.0f;
-      v4f[i] = 1.0f;
-      v2i[i] = 1;
-      v3i[i] = 1;
-      v4i[i] = 1;
-      v2u[i] = 1u;
-      v3u[i] = 1u;
-      v4u[i] = 1u;
-      v2b[i] = true;
-      v3b[i] = true;
-      v4b[i] = true;
-      {
-        i = (i + 1);
-      }
-    }
+  for(int i = 0; (i < 2); i = (i + 1)) {
+    v2f[i] = 1.0f;
+    v3f[i] = 1.0f;
+    v4f[i] = 1.0f;
+    v2i[i] = 1;
+    v3i[i] = 1;
+    v4i[i] = 1;
+    v2u[i] = 1u;
+    v3u[i] = 1u;
+    v4u[i] = 1u;
+    v2b[i] = true;
+    v3b[i] = true;
+    v4b[i] = true;
   }
   return;
 }
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl
index 7d73628..9f8f2e0 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_all.wgsl.expected.wgsl
@@ -12,28 +12,18 @@
   var v2b : vec2<bool>;
   var v3b : vec3<bool>;
   var v4b : vec4<bool>;
-  {
-    var i : i32 = 0;
-    loop {
-      if (!((i < 2))) {
-        break;
-      }
-      v2f[i] = 1.0;
-      v3f[i] = 1.0;
-      v4f[i] = 1.0;
-      v2i[i] = 1;
-      v3i[i] = 1;
-      v4i[i] = 1;
-      v2u[i] = 1u;
-      v3u[i] = 1u;
-      v4u[i] = 1u;
-      v2b[i] = true;
-      v3b[i] = true;
-      v4b[i] = true;
-
-      continuing {
-        i = (i + 1);
-      }
-    }
+  for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+    v2f[i] = 1.0;
+    v3f[i] = 1.0;
+    v4f[i] = 1.0;
+    v2i[i] = 1;
+    v3i[i] = 1;
+    v4i[i] = 1;
+    v2u[i] = 1u;
+    v3u[i] = 1u;
+    v4u[i] = 1u;
+    v2b[i] = true;
+    v3b[i] = true;
+    v4b[i] = true;
   }
 }
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl
index cbc4e6a..9c60097 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.hlsl
@@ -36,8 +36,7 @@
   bool2 v2b = bool2(false, false);
   bool2 v2b_2 = bool2(false, false);
   {
-    int i = 0;
-    for(; !(!((i < 2))); i = (i + 1)) {
+    for(int i = 0; (i < 2); i = (i + 1)) {
       Set_float2(v2f, i, 1.0f);
       Set_int3(v3i, i, 1);
       Set_uint4(v4u, i, 1u);
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
index 9f28e5c..4c124ca 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl
@@ -10,24 +10,15 @@
   uint4 v4u_2 = 0u;
   bool2 v2b = false;
   bool2 v2b_2 = false;
-  {
-    int i = 0;
-    while (true) {
-      if (!((i < 2))) {
-        break;
-      }
-      v2f[i] = 1.0f;
-      v3i[i] = 1;
-      v4u[i] = 1u;
-      v2b[i] = true;
-      v2f_2[i] = 1.0f;
-      v3i_2[i] = 1;
-      v4u_2[i] = 1u;
-      v2b_2[i] = true;
-      {
-        i = (i + 1);
-      }
-    }
+  for(int i = 0; (i < 2); i = (i + 1)) {
+    v2f[i] = 1.0f;
+    v3i[i] = 1;
+    v4u[i] = 1u;
+    v2b[i] = true;
+    v2f_2[i] = 1.0f;
+    v3i_2[i] = 1;
+    v4u_2[i] = 1u;
+    v2b_2[i] = true;
   }
   return;
 }
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl
index 0229323..74950ca 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.wgsl
@@ -8,24 +8,14 @@
   var v4u_2 : vec4<u32>;
   var v2b : vec2<bool>;
   var v2b_2 : vec2<bool>;
-  {
-    var i : i32 = 0;
-    loop {
-      if (!((i < 2))) {
-        break;
-      }
-      v2f[i] = 1.0;
-      v3i[i] = 1;
-      v4u[i] = 1u;
-      v2b[i] = true;
-      v2f_2[i] = 1.0;
-      v3i_2[i] = 1;
-      v4u_2[i] = 1u;
-      v2b_2[i] = true;
-
-      continuing {
-        i = (i + 1);
-      }
-    }
+  for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+    v2f[i] = 1.0;
+    v3i[i] = 1;
+    v4u[i] = 1u;
+    v2b[i] = true;
+    v2f_2[i] = 1.0;
+    v3i_2[i] = 1;
+    v4u_2[i] = 1u;
+    v2b_2[i] = true;
   }
 }
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl
index aed923f..2e6de37 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.hlsl
@@ -37,8 +37,7 @@
   bool3 v3b = bool3(false, false, false);
   bool4 v4b = bool4(false, false, false, false);
   {
-    int i = 0;
-    for(; !(!((i < 2))); i = (i + 1)) {
+    for(int i = 0; (i < 2); i = (i + 1)) {
       Set_float2(v2f, i, 1.0f);
       Set_int2(v2i, i, 1);
       Set_uint2(v2u, i, 1u);
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
index fb53199..2113ef6 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl
@@ -14,20 +14,11 @@
   bool2 v2b = false;
   bool3 v3b = false;
   bool4 v4b = false;
-  {
-    int i = 0;
-    while (true) {
-      if (!((i < 2))) {
-        break;
-      }
-      v2f[i] = 1.0f;
-      v2i[i] = 1;
-      v2u[i] = 1u;
-      v2b[i] = true;
-      {
-        i = (i + 1);
-      }
-    }
+  for(int i = 0; (i < 2); i = (i + 1)) {
+    v2f[i] = 1.0f;
+    v2i[i] = 1;
+    v2u[i] = 1u;
+    v2b[i] = true;
   }
   int i = 0;
   v3f[i] = 1.0f;
diff --git a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl
index 8c54146..1fada75 100644
--- a/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl
+++ b/test/fxc_bugs/vector_assignment_in_loop/loop_types_some.wgsl.expected.wgsl
@@ -12,21 +12,11 @@
   var v2b : vec2<bool>;
   var v3b : vec3<bool>;
   var v4b : vec4<bool>;
-  {
-    var i : i32 = 0;
-    loop {
-      if (!((i < 2))) {
-        break;
-      }
-      v2f[i] = 1.0;
-      v2i[i] = 1;
-      v2u[i] = 1u;
-      v2b[i] = true;
-
-      continuing {
-        i = (i + 1);
-      }
-    }
+  for(var i : i32 = 0; (i < 2); i = (i + 1)) {
+    v2f[i] = 1.0;
+    v2i[i] = 1;
+    v2u[i] = 1u;
+    v2b[i] = true;
   }
   var i : i32 = 0;
   v3f[i] = 1.0;
diff --git a/test/samples/compute_boids.wgsl.expected.hlsl b/test/samples/compute_boids.wgsl.expected.hlsl
index d98a316..4643faa 100644
--- a/test/samples/compute_boids.wgsl.expected.hlsl
+++ b/test/samples/compute_boids.wgsl.expected.hlsl
@@ -53,8 +53,7 @@
   float2 pos = float2(0.0f, 0.0f);
   float2 vel = float2(0.0f, 0.0f);
   {
-    uint i = 0u;
-    for(; !(!((i < 5u))); i = (i + 1u)) {
+    for(uint i = 0u; (i < 5u); i = (i + 1u)) {
       if ((i == index)) {
         continue;
       }
diff --git a/test/samples/compute_boids.wgsl.expected.msl b/test/samples/compute_boids.wgsl.expected.msl
index dae2f1f..cc9a856 100644
--- a/test/samples/compute_boids.wgsl.expected.msl
+++ b/test/samples/compute_boids.wgsl.expected.msl
@@ -61,34 +61,22 @@
   int cVelCount = 0;
   float2 pos = 0.0f;
   float2 vel = 0.0f;
-  {
-    uint i = 0u;
-    while (true) {
-      if (!((i < 5u))) {
-        break;
-      }
-      if ((i == index)) {
-        {
-          i = (i + 1u);
-        }
-        continue;
-      }
-      pos = particlesA.particles.arr[i].pos.xy;
-      vel = particlesA.particles.arr[i].vel.xy;
-      if ((distance(pos, vPos) < params.rule1Distance)) {
-        cMass = (cMass + pos);
-        cMassCount = (cMassCount + 1);
-      }
-      if ((distance(pos, vPos) < params.rule2Distance)) {
-        colVel = (colVel - (pos - vPos));
-      }
-      if ((distance(pos, vPos) < params.rule3Distance)) {
-        cVel = (cVel + vel);
-        cVelCount = (cVelCount + 1);
-      }
-      {
-        i = (i + 1u);
-      }
+  for(uint i = 0u; (i < 5u); i = (i + 1u)) {
+    if ((i == index)) {
+      continue;
+    }
+    pos = particlesA.particles.arr[i].pos.xy;
+    vel = particlesA.particles.arr[i].vel.xy;
+    if ((distance(pos, vPos) < params.rule1Distance)) {
+      cMass = (cMass + pos);
+      cMassCount = (cMassCount + 1);
+    }
+    if ((distance(pos, vPos) < params.rule2Distance)) {
+      colVel = (colVel - (pos - vPos));
+    }
+    if ((distance(pos, vPos) < params.rule3Distance)) {
+      cVel = (cVel + vel);
+      cVelCount = (cVelCount + 1);
     }
   }
   if ((cMassCount > 0)) {
diff --git a/test/samples/compute_boids.wgsl.expected.wgsl b/test/samples/compute_boids.wgsl.expected.wgsl
index 09cb30f..dd981f4 100644
--- a/test/samples/compute_boids.wgsl.expected.wgsl
+++ b/test/samples/compute_boids.wgsl.expected.wgsl
@@ -52,32 +52,22 @@
   var cVelCount : i32 = 0;
   var pos : vec2<f32>;
   var vel : vec2<f32>;
-  {
-    var i : u32 = 0u;
-    loop {
-      if (!((i < 5u))) {
-        break;
-      }
-      if ((i == index)) {
-        continue;
-      }
-      pos = particlesA.particles[i].pos.xy;
-      vel = particlesA.particles[i].vel.xy;
-      if ((distance(pos, vPos) < params.rule1Distance)) {
-        cMass = (cMass + pos);
-        cMassCount = (cMassCount + 1);
-      }
-      if ((distance(pos, vPos) < params.rule2Distance)) {
-        colVel = (colVel - (pos - vPos));
-      }
-      if ((distance(pos, vPos) < params.rule3Distance)) {
-        cVel = (cVel + vel);
-        cVelCount = (cVelCount + 1);
-      }
-
-      continuing {
-        i = (i + 1u);
-      }
+  for(var i : u32 = 0u; (i < 5u); i = (i + 1u)) {
+    if ((i == index)) {
+      continue;
+    }
+    pos = particlesA.particles[i].pos.xy;
+    vel = particlesA.particles[i].vel.xy;
+    if ((distance(pos, vPos) < params.rule1Distance)) {
+      cMass = (cMass + pos);
+      cMassCount = (cMassCount + 1);
+    }
+    if ((distance(pos, vPos) < params.rule2Distance)) {
+      colVel = (colVel - (pos - vPos));
+    }
+    if ((distance(pos, vPos) < params.rule3Distance)) {
+      cVel = (cVel + vel);
+      cVelCount = (cVelCount + 1);
     }
   }
   if ((cMassCount > 0)) {
diff --git a/test/statements/for/basic.wgsl.expected.hlsl b/test/statements/for/basic.wgsl.expected.hlsl
index c1e6281..f215f4b 100644
--- a/test/statements/for/basic.wgsl.expected.hlsl
+++ b/test/statements/for/basic.wgsl.expected.hlsl
@@ -8,8 +8,7 @@
 
 void f() {
   {
-    int i = 0;
-    for(; !(!((i < 5))); i = (i + 1)) {
+    for(int i = 0; (i < 5); i = (i + 1)) {
       some_loop_body();
     }
   }
diff --git a/test/statements/for/basic.wgsl.expected.msl b/test/statements/for/basic.wgsl.expected.msl
index d2e498f..e36af9f 100644
--- a/test/statements/for/basic.wgsl.expected.msl
+++ b/test/statements/for/basic.wgsl.expected.msl
@@ -5,17 +5,8 @@
 }
 
 void f() {
-  {
-    int i = 0;
-    while (true) {
-      if (!((i < 5))) {
-        break;
-      }
-      some_loop_body();
-      {
-        i = (i + 1);
-      }
-    }
+  for(int i = 0; (i < 5); i = (i + 1)) {
+    some_loop_body();
   }
 }
 
diff --git a/test/statements/for/basic.wgsl.expected.wgsl b/test/statements/for/basic.wgsl.expected.wgsl
index 6a66161..3f9a88a 100644
--- a/test/statements/for/basic.wgsl.expected.wgsl
+++ b/test/statements/for/basic.wgsl.expected.wgsl
@@ -2,17 +2,7 @@
 }
 
 fn f() {
-  {
-    var i : i32 = 0;
-    loop {
-      if (!((i < 5))) {
-        break;
-      }
-      some_loop_body();
-
-      continuing {
-        i = (i + 1);
-      }
-    }
+  for(var i : i32 = 0; (i < 5); i = (i + 1)) {
+    some_loop_body();
   }
 }
diff --git a/test/statements/for/complex.wgsl.expected.hlsl b/test/statements/for/complex.wgsl.expected.hlsl
index a3197c1..b48727f 100644
--- a/test/statements/for/complex.wgsl.expected.hlsl
+++ b/test/statements/for/complex.wgsl.expected.hlsl
@@ -15,7 +15,7 @@
       if (tint_tmp) {
         tint_tmp = (j < 10);
       }
-      if (!(!(!((tint_tmp))))) { break; }
+      if (!((tint_tmp))) { break; }
       some_loop_body();
       j = (i * 30);
       i = (i + 1);
diff --git a/test/statements/for/complex.wgsl.expected.msl b/test/statements/for/complex.wgsl.expected.msl
index 2030618..9ed6331 100644
--- a/test/statements/for/complex.wgsl.expected.msl
+++ b/test/statements/for/complex.wgsl.expected.msl
@@ -6,18 +6,9 @@
 
 void f() {
   int j = 0;
-  {
-    int i = 0;
-    while (true) {
-      if (!(((i < 5) && (j < 10)))) {
-        break;
-      }
-      some_loop_body();
-      j = (i * 30);
-      {
-        i = (i + 1);
-      }
-    }
+  for(int i = 0; ((i < 5) && (j < 10)); i = (i + 1)) {
+    some_loop_body();
+    j = (i * 30);
   }
 }
 
diff --git a/test/statements/for/complex.wgsl.expected.wgsl b/test/statements/for/complex.wgsl.expected.wgsl
index 1f3ce92..32bc4e0 100644
--- a/test/statements/for/complex.wgsl.expected.wgsl
+++ b/test/statements/for/complex.wgsl.expected.wgsl
@@ -3,18 +3,8 @@
 
 fn f() {
   var j : i32;
-  {
-    var i : i32 = 0;
-    loop {
-      if (!(((i < 5) && (j < 10)))) {
-        break;
-      }
-      some_loop_body();
-      j = (i * 30);
-
-      continuing {
-        i = (i + 1);
-      }
-    }
+  for(var i : i32 = 0; ((i < 5) && (j < 10)); i = (i + 1)) {
+    some_loop_body();
+    j = (i * 30);
   }
 }
diff --git a/test/statements/for/condition.wgsl.expected.hlsl b/test/statements/for/condition.wgsl.expected.hlsl
index 5f51493..e4456e0 100644
--- a/test/statements/for/condition.wgsl.expected.hlsl
+++ b/test/statements/for/condition.wgsl.expected.hlsl
@@ -5,6 +5,8 @@
 
 void f() {
   int i = 0;
-  for(; !(!((i < 4))); ) {
+  {
+    for(; (i < 4); ) {
+    }
   }
 }
diff --git a/test/statements/for/condition.wgsl.expected.msl b/test/statements/for/condition.wgsl.expected.msl
index c99148a..839541b 100644
--- a/test/statements/for/condition.wgsl.expected.msl
+++ b/test/statements/for/condition.wgsl.expected.msl
@@ -3,10 +3,7 @@
 using namespace metal;
 void f() {
   int i = 0;
-  while (true) {
-    if (!((i < 4))) {
-      break;
-    }
+  for(; (i < 4); ) {
   }
 }
 
diff --git a/test/statements/for/condition.wgsl.expected.wgsl b/test/statements/for/condition.wgsl.expected.wgsl
index 14f48d5..7fc3c6f 100644
--- a/test/statements/for/condition.wgsl.expected.wgsl
+++ b/test/statements/for/condition.wgsl.expected.wgsl
@@ -1,8 +1,5 @@
 fn f() {
   var i : i32;
-  loop {
-    if (!((i < 4))) {
-      break;
-    }
+  for(; (i < 4); ;) {
   }
 }
diff --git a/test/statements/for/continuing.wgsl.expected.hlsl b/test/statements/for/continuing.wgsl.expected.hlsl
index 28b8b3b..b3ec736 100644
--- a/test/statements/for/continuing.wgsl.expected.hlsl
+++ b/test/statements/for/continuing.wgsl.expected.hlsl
@@ -5,9 +5,8 @@
 
 void f() {
   int i = 0;
-  while (true) {
-    {
-      i = (i + 1);
+  {
+    for(; ; i = (i + 1)) {
     }
   }
 }
diff --git a/test/statements/for/continuing.wgsl.expected.msl b/test/statements/for/continuing.wgsl.expected.msl
index 67a069f..a288b5d 100644
--- a/test/statements/for/continuing.wgsl.expected.msl
+++ b/test/statements/for/continuing.wgsl.expected.msl
@@ -3,10 +3,7 @@
 using namespace metal;
 void f() {
   int i = 0;
-  while (true) {
-    {
-      i = (i + 1);
-    }
+  for(; ; i = (i + 1)) {
   }
 }
 
diff --git a/test/statements/for/continuing.wgsl.expected.wgsl b/test/statements/for/continuing.wgsl.expected.wgsl
index f3faf03..c76b514 100644
--- a/test/statements/for/continuing.wgsl.expected.wgsl
+++ b/test/statements/for/continuing.wgsl.expected.wgsl
@@ -1,9 +1,5 @@
 fn f() {
   var i : i32;
-  loop {
-
-    continuing {
-      i = (i + 1);
-    }
+  for(; ; i = (i + 1)) {
   }
 }
diff --git a/test/statements/for/empty.wgsl.expected.hlsl b/test/statements/for/empty.wgsl.expected.hlsl
index b642595..cd85e1a 100644
--- a/test/statements/for/empty.wgsl.expected.hlsl
+++ b/test/statements/for/empty.wgsl.expected.hlsl
@@ -4,6 +4,8 @@
 }
 
 void f() {
-  while (true) {
+  {
+    for(; ; ) {
+    }
   }
 }
diff --git a/test/statements/for/empty.wgsl.expected.msl b/test/statements/for/empty.wgsl.expected.msl
index 622427b..e716521 100644
--- a/test/statements/for/empty.wgsl.expected.msl
+++ b/test/statements/for/empty.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 void f() {
-  while (true) {
+  for(; ; ) {
   }
 }
 
diff --git a/test/statements/for/empty.wgsl.expected.wgsl b/test/statements/for/empty.wgsl.expected.wgsl
index 3674ab1..bb4e48b 100644
--- a/test/statements/for/empty.wgsl.expected.wgsl
+++ b/test/statements/for/empty.wgsl.expected.wgsl
@@ -1,4 +1,4 @@
 fn f() {
-  loop {
+  for(; ; ;) {
   }
 }
diff --git a/test/statements/for/initializer.wgsl.expected.hlsl b/test/statements/for/initializer.wgsl.expected.hlsl
index 8311c7b..e614afd 100644
--- a/test/statements/for/initializer.wgsl.expected.hlsl
+++ b/test/statements/for/initializer.wgsl.expected.hlsl
@@ -5,8 +5,7 @@
 
 void f() {
   {
-    int i = 0;
-    while (true) {
+    for(int i = 0; ; ) {
     }
   }
 }
diff --git a/test/statements/for/initializer.wgsl.expected.msl b/test/statements/for/initializer.wgsl.expected.msl
index 70c5111..90082d6 100644
--- a/test/statements/for/initializer.wgsl.expected.msl
+++ b/test/statements/for/initializer.wgsl.expected.msl
@@ -2,10 +2,7 @@
 
 using namespace metal;
 void f() {
-  {
-    int i = 0;
-    while (true) {
-    }
+  for(int i = 0; ; ) {
   }
 }
 
diff --git a/test/statements/for/initializer.wgsl.expected.wgsl b/test/statements/for/initializer.wgsl.expected.wgsl
index e021556..ef9b991 100644
--- a/test/statements/for/initializer.wgsl.expected.wgsl
+++ b/test/statements/for/initializer.wgsl.expected.wgsl
@@ -1,7 +1,4 @@
 fn f() {
-  {
-    var i : i32 = 0;
-    loop {
-    }
+  for(var i : i32 = 0; ; ;) {
   }
 }
diff --git a/test/statements/for/scoping.wgsl.expected.hlsl b/test/statements/for/scoping.wgsl.expected.hlsl
index 459eb43..143b1e3 100644
--- a/test/statements/for/scoping.wgsl.expected.hlsl
+++ b/test/statements/for/scoping.wgsl.expected.hlsl
@@ -5,8 +5,7 @@
 
 void f() {
   {
-    int must_not_collide = 0;
-    while (true) {
+    for(int must_not_collide = 0; ; ) {
     }
   }
   int must_not_collide = 0;
diff --git a/test/statements/for/scoping.wgsl.expected.msl b/test/statements/for/scoping.wgsl.expected.msl
index d5c3f3a..0572921 100644
--- a/test/statements/for/scoping.wgsl.expected.msl
+++ b/test/statements/for/scoping.wgsl.expected.msl
@@ -2,10 +2,7 @@
 
 using namespace metal;
 void f() {
-  {
-    int must_not_collide = 0;
-    while (true) {
-    }
+  for(int must_not_collide = 0; ; ) {
   }
   int must_not_collide = 0;
 }
diff --git a/test/statements/for/scoping.wgsl.expected.wgsl b/test/statements/for/scoping.wgsl.expected.wgsl
index ec202ae..0cac8520 100644
--- a/test/statements/for/scoping.wgsl.expected.wgsl
+++ b/test/statements/for/scoping.wgsl.expected.wgsl
@@ -1,8 +1,5 @@
 fn f() {
-  {
-    var must_not_collide : i32 = 0;
-    loop {
-    }
+  for(var must_not_collide : i32 = 0; ; ;) {
   }
   var must_not_collide : i32;
 }