transform: Fix PromoteInitializersToConstVar handling of for-loops

PromoteInitializersToConstVar was erroring on for loops that contained array or structure constructor expressions.

Added lots more tests.

Fixed: tint:1364
Change-Id: I033eaad94756ea496fc8bc5f03f39c6dba4e3a88
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/75580
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>

docs/origin-trial-changes.md

@@ -15,6 +15,10 @@ The following features have been deprecated and will be removed in M102:
 * Added builtins `degrees()` and `radians()` for converting between degrees and radians. [tint:1329](https://crbug.com/tint/1329)
 * `let` arrays and matrices can now be dynamically indexed. [tint:1352](https://crbug.com/tint/1352)
 
+### Fixes
+
+* Fixed an issue where for-loops that contain array or structure constructors in the loop initializer statements, condition expressions or continuing statements could fail to compile. [tint:1364](https://crbug.com/tint/1364)
+
 ## Changes for M98
 
 ### Breaking Changes

src/resolver/resolver_behavior_test.cc

@@ -17,6 +17,7 @@
 #include "gtest/gtest.h"
 #include "src/resolver/resolver_test_helper.h"
 #include "src/sem/expression.h"
+#include "src/sem/for_loop_statement.h"
 
 namespace tint {
 namespace resolver {

src/sem/type_mappings.h

@@ -23,6 +23,7 @@ namespace tint {
 namespace ast {
 class CallExpression;
 class Expression;
+class ForLoopStatement;
 class Function;
 class MemberAccessorExpression;
 class Node;
@@ -39,6 +40,7 @@ namespace sem {
 class Array;
 class Call;
 class Expression;
+class ForLoopStatement;
 class Function;
 class MemberAccessorExpression;
 class Node;
@@ -56,6 +58,7 @@ struct TypeMappings {
   //! @cond Doxygen_Suppress
   Call* operator()(ast::CallExpression*);
   Expression* operator()(ast::Expression*);
+  ForLoopStatement* operator()(ast::ForLoopStatement*);
   Function* operator()(ast::Function*);
   MemberAccessorExpression* operator()(ast::MemberAccessorExpression*);
   Node* operator()(ast::Node*);

src/transform/promote_initializers_to_const_var.cc

@@ -14,12 +14,14 @@
 
 #include "src/transform/promote_initializers_to_const_var.h"
 
+#include <unordered_map>
 #include <utility>
 
 #include "src/program_builder.h"
 #include "src/sem/block_statement.h"
 #include "src/sem/call.h"
 #include "src/sem/expression.h"
+#include "src/sem/for_loop_statement.h"
 #include "src/sem/statement.h"
 #include "src/sem/type_constructor.h"
 
@@ -28,6 +30,18 @@ TINT_INSTANTIATE_TYPEINFO(tint::transform::PromoteInitializersToConstVar);
 namespace tint {
 namespace transform {
 
+namespace {
+
+/// Holds information about a for-loop that needs to be decomposed into a loop,
+/// so that initializer declaration statements can be inserted before the
+/// condition expression or continuing statement.
+struct LoopInfo {
+  ast::StatementList cond_decls;
+  ast::StatementList cont_decls;
+};
+
+}  // namespace
+
 PromoteInitializersToConstVar::PromoteInitializersToConstVar() = default;
 
 PromoteInitializersToConstVar::~PromoteInitializersToConstVar() = default;
@@ -35,6 +49,7 @@ PromoteInitializersToConstVar::~PromoteInitializersToConstVar() = default;
 void PromoteInitializersToConstVar::Run(CloneContext& ctx,
                                         const DataMap&,
                                         DataMap&) {
+  auto& sem = ctx.src->Sem();
   // Scan the AST nodes for array and structure initializers which
   // need to be promoted to their own constant declaration.
 
@@ -51,24 +66,27 @@ void PromoteInitializersToConstVar::Run(CloneContext& ctx,
   // immutable and require their children to be constructed first so their
   // pointer can be passed to the parent's constructor.
 
-  for (auto* src_node : ctx.src->ASTNodes().Objects()) {
+  // For-loops that need to be decomposed to loops.
-    if (auto* src_init = src_node->As<ast::CallExpression>()) {
+  std::unordered_map<const sem::ForLoopStatement*, LoopInfo> loops;
-      auto* call = ctx.src->Sem().Get(src_init);
+
-      if (!call->Target()->Is<sem::TypeConstructor>()) {
+  for (auto* node : ctx.src->ASTNodes().Objects()) {
+    if (auto* expr = node->As<ast::CallExpression>()) {
+      auto* ctor = ctx.src->Sem().Get(expr);
+      if (!ctor->Target()->Is<sem::TypeConstructor>()) {
         continue;
       }
-      auto* src_sem_stmt = call->Stmt();
+      auto* sem_stmt = ctor->Stmt();
-      if (!src_sem_stmt) {
+      if (!sem_stmt) {
         // Expression is outside of a statement. This usually means the
         // expression is part of a global (module-scope) constant declaration.
         // These must be constexpr, and so cannot contain the type of
         // expressions that must be sanitized.
         continue;
       }
-      auto* src_stmt = src_sem_stmt->Declaration();
+      auto* stmt = sem_stmt->Declaration();
 
-      if (auto* src_var_decl = src_stmt->As<ast::VariableDeclStatement>()) {
+      if (auto* src_var_decl = stmt->As<ast::VariableDeclStatement>()) {
-        if (src_var_decl->variable->constructor == src_init) {
+        if (src_var_decl->variable->constructor == expr) {
           // This statement is just a variable declaration with the initializer
           // as the constructor value. This is what we're attempting to
           // transform to, and so ignore.
@@ -76,30 +94,108 @@ void PromoteInitializersToConstVar::Run(CloneContext& ctx,
         }
       }
 
-      auto* src_ty = call->Type();
+      auto* src_ty = ctor->Type();
       if (src_ty->IsAnyOf<sem::Array, sem::Struct>()) {
-        // Create a new symbol for the constant
+        // Create a new symbol for the let
-        auto dst_symbol = ctx.dst->Sym();
+        auto name = ctx.dst->Sym();
-        // Clone the type
+        // Construct the let that holds the hoisted initializer
-        auto* dst_ty = CreateASTTypeFor(ctx, call->Type());
+        auto* let = ctx.dst->Const(name, nullptr, ctx.Clone(expr));
-        // Clone the initializer
+        // Construct the let declaration statement
-        auto* dst_init = ctx.Clone(src_init);
+        auto* let_decl = ctx.dst->Decl(let);
-        // Construct the constant that holds the hoisted initializer
+        // Replace the initializer expression with a reference to the let
-        auto* dst_var = ctx.dst->Const(dst_symbol, dst_ty, dst_init);
+        ctx.Replace(expr, ctx.dst->Expr(name));
-        // Construct the variable declaration statement
-        auto* dst_var_decl = ctx.dst->Decl(dst_var);
-        // Construct the identifier for referencing the constant
-        auto* dst_ident = ctx.dst->Expr(dst_symbol);
 
-        // Insert the constant before the usage
+        if (auto* fl = sem_stmt->As<sem::ForLoopStatement>()) {
-        ctx.InsertBefore(src_sem_stmt->Block()->Declaration()->statements,
+          // Expression used in for-loop condition.
-                         src_stmt, dst_var_decl);
+          // For-loop needs to be decomposed to a loop.
-        // Replace the inlined initializer with a reference to the constant
+          loops[fl].cond_decls.emplace_back(let_decl);
-        ctx.Replace(src_init, dst_ident);
+          continue;
+        }
+
+        auto* parent = sem_stmt->Parent();  // The statement's parent
+        if (auto* block = parent->As<sem::BlockStatement>()) {
+          // Expression's statement sits in a block. Simple case.
+          // Insert the let before the parent statement
+          ctx.InsertBefore(block->Declaration()->statements, stmt, let_decl);
+          continue;
+        }
+        if (auto* fl = parent->As<sem::ForLoopStatement>()) {
+          // Expression is used in a for-loop. These require special care.
+          if (fl->Declaration()->initializer == stmt) {
+            // Expression used in for-loop initializer.
+            // Insert the let above the for-loop.
+            ctx.InsertBefore(fl->Block()->Declaration()->statements,
+                             fl->Declaration(), let_decl);
+            continue;
+          }
+          if (fl->Declaration()->continuing == stmt) {
+            // Expression used in for-loop continuing.
+            // For-loop needs to be decomposed to a loop.
+            loops[fl].cont_decls.emplace_back(let_decl);
+            continue;
+          }
+          TINT_ICE(Transform, ctx.dst->Diagnostics())
+              << "unhandled use of expression in for-loop";
+        }
+
+        TINT_ICE(Transform, ctx.dst->Diagnostics())
+            << "unhandled expression parent statement type: "
+            << parent->TypeInfo().name;
       }
     }
   }
 
+  if (!loops.empty()) {
+    // At least one for-loop needs to be transformed into a loop.
+    ctx.ReplaceAll(
+        [&](const ast::ForLoopStatement* stmt) -> const ast::Statement* {
+          if (auto* fl = sem.Get(stmt)) {
+            if (auto it = loops.find(fl); it != loops.end()) {
+              auto& info = it->second;
+              auto* for_loop = fl->Declaration();
+              // For-loop needs to be decomposed to a loop.
+              // Build the loop body's statements.
+              // Start with any let declarations for the conditional expression.
+              auto body_stmts = info.cond_decls;
+              // If the for-loop has a condition, emit this next as:
+              //   if (!cond) { break; }
+              if (auto* cond = for_loop->condition) {
+                // !condition
+                auto* not_cond = ctx.dst->create<ast::UnaryOpExpression>(
+                    ast::UnaryOp::kNot, ctx.Clone(cond));
+                // { break; }
+                auto* break_body =
+                    ctx.dst->Block(ctx.dst->create<ast::BreakStatement>());
+                // if (!condition) { break; }
+                body_stmts.emplace_back(ctx.dst->If(not_cond, break_body));
+              }
+              // Next emit the for-loop body
+              for (auto* body_stmt : for_loop->body->statements) {
+                body_stmts.emplace_back(ctx.Clone(body_stmt));
+              }
+
+              // Finally create the continuing block if there was one.
+              const ast::BlockStatement* continuing = nullptr;
+              if (auto* cont = for_loop->continuing) {
+                // Continuing block starts with any let declarations used by the
+                // continuing.
+                auto cont_stmts = info.cont_decls;
+                cont_stmts.emplace_back(ctx.Clone(cont));
+                continuing = ctx.dst->Block(cont_stmts);
+              }
+
+              auto* body = ctx.dst->Block(body_stmts);
+              auto* loop = ctx.dst->Loop(body, continuing);
+              if (auto* init = for_loop->initializer) {
+                return ctx.dst->Block(ctx.Clone(init), loop);
+              }
+              return loop;
+            }
+          }
+          return nullptr;
+        });
+  }
+
   ctx.Clone();
 }
 

src/transform/promote_initializers_to_const_var.h

@@ -21,8 +21,10 @@ namespace tint {
 namespace transform {
 
 /// A transform that hoists the array and structure initializers to a constant
-/// variable, declared just before the statement of usage. See
+/// variable, declared just before the statement of usage. This transform may
-/// crbug.com/tint/406 for more details.
+/// also decompose for-loops into loops so that let declarations can be emitted
+/// before loop condition expressions and/or continuing statements.
+/// @see crbug.com/tint/406
 class PromoteInitializersToConstVar
     : public Castable<PromoteInitializersToConstVar, Transform> {
  public:

src/transform/promote_initializers_to_const_var_test.cc

@@ -24,25 +24,23 @@ using PromoteInitializersToConstVarTest = TransformTest;
 
 TEST_F(PromoteInitializersToConstVarTest, BasicArray) {
   auto* src = R"(
-[[stage(compute), workgroup_size(1)]]
+fn f() {
-fn main() {
+  var f0 = 1.0;
-  var f0 : f32 = 1.0;
+  var f1 = 2.0;
-  var f1 : f32 = 2.0;
+  var f2 = 3.0;
-  var f2 : f32 = 3.0;
+  var f3 = 4.0;
-  var f3 : f32 = 4.0;
+  var i = array<f32, 4u>(f0, f1, f2, f3)[2];
-  var i : f32 = array<f32, 4u>(f0, f1, f2, f3)[2];
 }
 )";
 
   auto* expect = R"(
-[[stage(compute), workgroup_size(1)]]
+fn f() {
-fn main() {
+  var f0 = 1.0;
-  var f0 : f32 = 1.0;
+  var f1 = 2.0;
-  var f1 : f32 = 2.0;
+  var f2 = 3.0;
-  var f2 : f32 = 3.0;
+  var f3 = 4.0;
-  var f3 : f32 = 4.0;
+  let tint_symbol = array<f32, 4u>(f0, f1, f2, f3);
-  let tint_symbol : array<f32, 4u> = array<f32, 4u>(f0, f1, f2, f3);
+  var i = tint_symbol[2];
-  var i : f32 = tint_symbol[2];
 }
 )";
 
@@ -59,9 +57,8 @@ struct S {
   c : vec3<f32>;
 };
 
-[[stage(compute), workgroup_size(1)]]
+fn f() {
-fn main() {
+  var x = S(1, 2.0, vec3<f32>()).b;
-  var x : f32 = S(1, 2.0, vec3<f32>()).b;
 }
 )";
 
@@ -72,10 +69,168 @@ struct S {
   c : vec3<f32>;
 };
 
-[[stage(compute), workgroup_size(1)]]
+fn f() {
-fn main() {
+  let tint_symbol = S(1, 2.0, vec3<f32>());
-  let tint_symbol : S = S(1, 2.0, vec3<f32>());
+  var x = tint_symbol.b;
-  var x : f32 = tint_symbol.b;
+}
+)";
+
+  auto got = Run<PromoteInitializersToConstVar>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PromoteInitializersToConstVarTest, ArrayInForLoopInit) {
+  auto* src = R"(
+fn f() {
+  var insert_after = 1;
+  for(var i = array<f32, 4u>(0.0, 1.0, 2.0, 3.0)[2]; ; ) {
+  }
+}
+)";
+
+  auto* expect = R"(
+fn f() {
+  var insert_after = 1;
+  let tint_symbol = array<f32, 4u>(0.0, 1.0, 2.0, 3.0);
+  for(var i = tint_symbol[2]; ; ) {
+  }
+}
+)";
+
+  auto got = Run<PromoteInitializersToConstVar>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PromoteInitializersToConstVarTest, StructInForLoopInit) {
+  auto* src = R"(
+struct S {
+  a : i32;
+  b : f32;
+  c : vec3<f32>;
+};
+
+fn f() {
+  var insert_after = 1;
+  for(var x = S(1, 2.0, vec3<f32>()).b; ; ) {
+  }
+}
+)";
+
+  auto* expect = R"(
+struct S {
+  a : i32;
+  b : f32;
+  c : vec3<f32>;
+};
+
+fn f() {
+  var insert_after = 1;
+  let tint_symbol = S(1, 2.0, vec3<f32>());
+  for(var x = tint_symbol.b; ; ) {
+  }
+}
+)";
+
+  auto got = Run<PromoteInitializersToConstVar>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PromoteInitializersToConstVarTest, ArrayInForLoopCond) {
+  auto* src = R"(
+fn f() {
+  var f = 1.0;
+  for(; f == array<f32, 1u>(f)[0]; f = f + 1.0) {
+    var marker = 1;
+  }
+}
+)";
+
+  auto* expect = R"(
+fn f() {
+  var f = 1.0;
+  loop {
+    let tint_symbol = array<f32, 1u>(f);
+    if (!((f == tint_symbol[0]))) {
+      break;
+    }
+    var marker = 1;
+
+    continuing {
+      f = (f + 1.0);
+    }
+  }
+}
+)";
+
+  auto got = Run<PromoteInitializersToConstVar>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PromoteInitializersToConstVarTest, ArrayInForLoopCont) {
+  auto* src = R"(
+fn f() {
+  var f = 0.0;
+  for(; f < 10.0; f = f + array<f32, 1u>(1.0)[0]) {
+    var marker = 1;
+  }
+}
+)";
+
+  auto* expect = R"(
+fn f() {
+  var f = 0.0;
+  loop {
+    if (!((f < 10.0))) {
+      break;
+    }
+    var marker = 1;
+
+    continuing {
+      let tint_symbol = array<f32, 1u>(1.0);
+      f = (f + tint_symbol[0]);
+    }
+  }
+}
+)";
+
+  auto got = Run<PromoteInitializersToConstVar>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PromoteInitializersToConstVarTest, ArrayInForLoopInitCondCont) {
+  auto* src = R"(
+fn f() {
+  for(var f = array<f32, 1u>(0.0)[0];
+      f < array<f32, 1u>(1.0)[0];
+      f = f + array<f32, 1u>(2.0)[0]) {
+    var marker = 1;
+  }
+}
+)";
+
+  auto* expect = R"(
+fn f() {
+  let tint_symbol = array<f32, 1u>(0.0);
+  {
+    var f = tint_symbol[0];
+    loop {
+      let tint_symbol_1 = array<f32, 1u>(1.0);
+      if (!((f < tint_symbol_1[0]))) {
+        break;
+      }
+      var marker = 1;
+
+      continuing {
+        let tint_symbol_2 = array<f32, 1u>(2.0);
+        f = (f + tint_symbol_2[0]);
+      }
+    }
+  }
 }
 )";
 
@@ -86,19 +241,17 @@ fn main() {
 
 TEST_F(PromoteInitializersToConstVarTest, ArrayInArrayArray) {
   auto* src = R"(
-[[stage(compute), workgroup_size(1)]]
+fn f() {
-fn main() {
+  var i = array<array<f32, 2u>, 2u>(array<f32, 2u>(1.0, 2.0), array<f32, 2u>(3.0, 4.0))[0][1];
-  var i : f32 = array<array<f32, 2u>, 2u>(array<f32, 2u>(1.0, 2.0), array<f32, 2u>(3.0, 4.0))[0][1];
 }
 )";
 
   auto* expect = R"(
-[[stage(compute), workgroup_size(1)]]
+fn f() {
-fn main() {
+  let tint_symbol = array<f32, 2u>(1.0, 2.0);
-  let tint_symbol : array<f32, 2u> = array<f32, 2u>(1.0, 2.0);
+  let tint_symbol_1 = array<f32, 2u>(3.0, 4.0);
-  let tint_symbol_1 : array<f32, 2u> = array<f32, 2u>(3.0, 4.0);
+  let tint_symbol_2 = array<array<f32, 2u>, 2u>(tint_symbol, tint_symbol_1);
-  let tint_symbol_2 : array<array<f32, 2u>, 2u> = array<array<f32, 2u>, 2u>(tint_symbol, tint_symbol_1);
+  var i = tint_symbol_2[0][1];
-  var i : f32 = tint_symbol_2[0][1];
 }
 )";
 
@@ -123,9 +276,8 @@ struct S3 {
   a : S2;
 };
 
-[[stage(compute), workgroup_size(1)]]
+fn f() {
-fn main() {
+  var x = S3(S2(1, S1(2), 3)).a.b.a;
-  var x : i32 = S3(S2(1, S1(2), 3)).a.b.a;
 }
 )";
 
@@ -144,12 +296,11 @@ struct S3 {
   a : S2;
 };
 
-[[stage(compute), workgroup_size(1)]]
+fn f() {
-fn main() {
+  let tint_symbol = S1(2);
-  let tint_symbol : S1 = S1(2);
+  let tint_symbol_1 = S2(1, tint_symbol, 3);
-  let tint_symbol_1 : S2 = S2(1, tint_symbol, 3);
+  let tint_symbol_2 = S3(tint_symbol_1);
-  let tint_symbol_2 : S3 = S3(tint_symbol_1);
+  var x = tint_symbol_2.a.b.a;
-  var x : i32 = tint_symbol_2.a.b.a;
 }
 )";
 
@@ -168,9 +319,8 @@ struct S2 {
   a : array<S1, 3u>;
 };
 
-[[stage(compute), workgroup_size(1)]]
+fn f() {
-fn main() {
+  var x = S2(array<S1, 3u>(S1(1), S1(2), S1(3))).a[1].a;
-  var x : i32 = S2(array<S1, 3u>(S1(1), S1(2), S1(3))).a[1].a;
 }
 )";
 
@@ -183,14 +333,13 @@ struct S2 {
   a : array<S1, 3u>;
 };
 
-[[stage(compute), workgroup_size(1)]]
+fn f() {
-fn main() {
+  let tint_symbol = S1(1);
-  let tint_symbol : S1 = S1(1);
+  let tint_symbol_1 = S1(2);
-  let tint_symbol_1 : S1 = S1(2);
+  let tint_symbol_2 = S1(3);
-  let tint_symbol_2 : S1 = S1(3);
+  let tint_symbol_3 = array<S1, 3u>(tint_symbol, tint_symbol_1, tint_symbol_2);
-  let tint_symbol_3 : array<S1, 3u> = array<S1, 3u>(tint_symbol, tint_symbol_1, tint_symbol_2);
+  let tint_symbol_4 = S2(tint_symbol_3);
-  let tint_symbol_4 : S2 = S2(tint_symbol_3);
+  var x = tint_symbol_4.a[1].a;
-  var x : i32 = tint_symbol_4.a[1].a;
 }
 )";
 
@@ -207,10 +356,9 @@ struct S {
   c : i32;
 };
 
-[[stage(compute), workgroup_size(1)]]
+fn f() {
-fn main() {
+  var local_arr = array<f32, 4u>(0.0, 1.0, 2.0, 3.0);
-  var local_arr : array<f32, 4u> = array<f32, 4u>(0.0, 1.0, 2.0, 3.0);
+  var local_str = S(1, 2.0, 3);
-  var local_str : S = S(1, 2.0, 3);
 }
 
 let module_arr : array<f32, 4u> = array<f32, 4u>(0.0, 1.0, 2.0, 3.0);

test/statements/for/condition.wgsl

@@ -1,5 +0,0 @@
-fn f() {
-    var i : i32;
-    for (;i < 4;) {
-    }
-}

test/statements/for/condition/array_ctor.wgsl

@@ -0,0 +1,4 @@
+fn f() {
+  var i : i32;
+  for (;i < array<i32, 1>(1)[0];) {}
+}

test/statements/for/condition/array_ctor.wgsl.expected.hlsl

@@ -0,0 +1,14 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+void f() {
+  int i = 0;
+  [loop] while (true) {
+    const int tint_symbol[1] = {1};
+    if (!((i < tint_symbol[0]))) {
+      break;
+    }
+  }
+}

test/statements/for/condition/array_ctor.wgsl.expected.msl

@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  int arr[1];
+};
+
+void f() {
+  int i = 0;
+  while (true) {
+    tint_array_wrapper const tint_symbol = {.arr={1}};
+    if (!((i < tint_symbol.arr[0]))) {
+      break;
+    }
+  }
+}
+

test/statements/for/condition/array_ctor.wgsl.expected.spvasm

@@ -0,0 +1,52 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 28
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %f "f"
+               OpName %i "i"
+               OpDecorate %_arr_int_uint_1 ArrayStride 4
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+%_ptr_Function_int = OpTypePointer Function %int
+         %10 = OpConstantNull %int
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+%_arr_int_uint_1 = OpTypeArray %int %uint_1
+      %int_1 = OpConstant %int 1
+         %21 = OpConstantComposite %_arr_int_uint_1 %int_1
+      %int_0 = OpConstant %int 0
+       %bool = OpTypeBool
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %1
+          %6 = OpLabel
+          %i = OpVariable %_ptr_Function_int Function %10
+               OpBranch %11
+         %11 = OpLabel
+               OpLoopMerge %12 %13 None
+               OpBranch %14
+         %14 = OpLabel
+         %16 = OpLoad %int %i
+         %23 = OpCompositeExtract %int %21 0
+         %24 = OpSLessThan %bool %16 %23
+         %15 = OpLogicalNot %bool %24
+               OpSelectionMerge %26 None
+               OpBranchConditional %15 %27 %26
+         %27 = OpLabel
+               OpBranch %12
+         %26 = OpLabel
+               OpBranch %13
+         %13 = OpLabel
+               OpBranch %11
+         %12 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/statements/for/condition/array_ctor.wgsl.expected.wgsl

@@ -0,0 +1,5 @@
+fn f() {
+  var i : i32;
+  for(; (i < array<i32, 1>(1)[0]); ) {
+  }
+}

test/statements/for/condition/basic.wgsl

@@ -0,0 +1,5 @@
+fn f() {
+  var i : i32;
+  for (;i < 4;) {
+  }
+}

test/statements/for/condition/struct_ctor.wgsl

@@ -0,0 +1,8 @@
+struct S {
+  i : i32;
+};
+
+fn f() {
+  var i : i32;
+  for (; i < S(1).i;) {}
+}

test/statements/for/condition/struct_ctor.wgsl.expected.hlsl

@@ -0,0 +1,18 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+struct S {
+  int i;
+};
+
+void f() {
+  int i = 0;
+  [loop] while (true) {
+    const S tint_symbol = {1};
+    if (!((i < tint_symbol.i))) {
+      break;
+    }
+  }
+}

test/statements/for/condition/struct_ctor.wgsl.expected.msl

@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  int i;
+};
+
+void f() {
+  int i = 0;
+  while (true) {
+    S const tint_symbol = {.i=1};
+    if (!((i < tint_symbol.i))) {
+      break;
+    }
+  }
+}
+

test/statements/for/condition/struct_ctor.wgsl.expected.spvasm

@@ -0,0 +1,51 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 25
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %f "f"
+               OpName %i "i"
+               OpName %S "S"
+               OpMemberName %S 0 "i"
+               OpMemberDecorate %S 0 Offset 0
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+%_ptr_Function_int = OpTypePointer Function %int
+         %10 = OpConstantNull %int
+          %S = OpTypeStruct %int
+      %int_1 = OpConstant %int 1
+         %19 = OpConstantComposite %S %int_1
+       %bool = OpTypeBool
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %1
+          %6 = OpLabel
+          %i = OpVariable %_ptr_Function_int Function %10
+               OpBranch %11
+         %11 = OpLabel
+               OpLoopMerge %12 %13 None
+               OpBranch %14
+         %14 = OpLabel
+         %16 = OpLoad %int %i
+         %20 = OpCompositeExtract %int %19 0
+         %21 = OpSLessThan %bool %16 %20
+         %15 = OpLogicalNot %bool %21
+               OpSelectionMerge %23 None
+               OpBranchConditional %15 %24 %23
+         %24 = OpLabel
+               OpBranch %12
+         %23 = OpLabel
+               OpBranch %13
+         %13 = OpLabel
+               OpBranch %11
+         %12 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/statements/for/condition/struct_ctor.wgsl.expected.wgsl

@@ -0,0 +1,9 @@
+struct S {
+  i : i32;
+};
+
+fn f() {
+  var i : i32;
+  for(; (i < S(1).i); ) {
+  }
+}

test/statements/for/continuing.wgsl

@@ -1,4 +0,0 @@
-fn f() {
-    var i : i32;
-    for (;;i = i + 1) {}
-}

test/statements/for/continuing/array_ctor.wgsl

@@ -0,0 +1,4 @@
+fn f() {
+  var i : i32;
+  for (;;i = i + array<i32, 1>(1)[0]) {}
+}

test/statements/for/continuing/array_ctor.wgsl.expected.hlsl

@@ -0,0 +1,14 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+void f() {
+  int i = 0;
+  [loop] while (true) {
+    {
+      const int tint_symbol[1] = {1};
+      i = (i + tint_symbol[0]);
+    }
+  }
+}

test/statements/for/continuing/array_ctor.wgsl.expected.msl

@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  int arr[1];
+};
+
+void f() {
+  int i = 0;
+  while (true) {
+    {
+      tint_array_wrapper const tint_symbol = {.arr={1}};
+      i = as_type<int>((as_type<uint>(i) + as_type<uint>(tint_symbol.arr[0])));
+    }
+  }
+}
+

test/statements/for/continuing/array_ctor.wgsl.expected.spvasm

@@ -0,0 +1,46 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %f "f"
+               OpName %i "i"
+               OpDecorate %_arr_int_uint_1 ArrayStride 4
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+%_ptr_Function_int = OpTypePointer Function %int
+         %10 = OpConstantNull %int
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+%_arr_int_uint_1 = OpTypeArray %int %uint_1
+      %int_1 = OpConstant %int 1
+         %20 = OpConstantComposite %_arr_int_uint_1 %int_1
+      %int_0 = OpConstant %int 0
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %1
+          %6 = OpLabel
+          %i = OpVariable %_ptr_Function_int Function %10
+               OpBranch %11
+         %11 = OpLabel
+               OpLoopMerge %12 %13 None
+               OpBranch %14
+         %14 = OpLabel
+               OpBranch %13
+         %13 = OpLabel
+         %15 = OpLoad %int %i
+         %22 = OpCompositeExtract %int %20 0
+         %23 = OpIAdd %int %15 %22
+               OpStore %i %23
+               OpBranch %11
+         %12 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/statements/for/continuing/array_ctor.wgsl.expected.wgsl

@@ -0,0 +1,5 @@
+fn f() {
+  var i : i32;
+  for(; ; i = (i + array<i32, 1>(1)[0])) {
+  }
+}

test/statements/for/continuing/basic.wgsl

@@ -0,0 +1,4 @@
+fn f() {
+  var i : i32;
+  for (;;i = i + 1) {}
+}

test/statements/for/continuing/struct_ctor.wgsl

@@ -0,0 +1,7 @@
+struct S {
+  i : i32;
+};
+
+fn f() {
+  for (var i = 0;; i = i + S(1).i) {}
+}

test/statements/for/continuing/struct_ctor.wgsl.expected.hlsl

@@ -0,0 +1,20 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+struct S {
+  int i;
+};
+
+void f() {
+  {
+    int i = 0;
+    [loop] while (true) {
+      {
+        const S tint_symbol = {1};
+        i = (i + tint_symbol.i);
+      }
+    }
+  }
+}

test/statements/for/continuing/struct_ctor.wgsl.expected.msl

@@ -0,0 +1,19 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  int i;
+};
+
+void f() {
+  {
+    int i = 0;
+    while (true) {
+      {
+        S const tint_symbol = {.i=1};
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(tint_symbol.i)));
+      }
+    }
+  }
+}
+

test/statements/for/continuing/struct_ctor.wgsl.expected.spvasm

@@ -0,0 +1,47 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 22
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %f "f"
+               OpName %i "i"
+               OpName %S "S"
+               OpMemberName %S 0 "i"
+               OpMemberDecorate %S 0 Offset 0
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+         %11 = OpConstantNull %int
+          %S = OpTypeStruct %int
+      %int_1 = OpConstant %int 1
+         %19 = OpConstantComposite %S %int_1
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %1
+          %6 = OpLabel
+          %i = OpVariable %_ptr_Function_int Function %11
+               OpStore %i %int_0
+               OpBranch %12
+         %12 = OpLabel
+               OpLoopMerge %13 %14 None
+               OpBranch %15
+         %15 = OpLabel
+               OpBranch %14
+         %14 = OpLabel
+         %16 = OpLoad %int %i
+         %20 = OpCompositeExtract %int %19 0
+         %21 = OpIAdd %int %16 %20
+               OpStore %i %21
+               OpBranch %12
+         %13 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/statements/for/continuing/struct_ctor.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+struct S {
+  i : i32;
+};
+
+fn f() {
+  for(var i = 0; ; i = (i + S(1).i)) {
+  }
+}

test/statements/for/initializer.wgsl

@@ -1,3 +0,0 @@
-fn f() {
-    for (var i : i32 = 0;;) {}
-}

test/statements/for/initializer/array_ctor.wgsl

@@ -0,0 +1,3 @@
+fn f() {
+    for (var i : i32 = array<i32, 1>(1)[0];;) {}
+}

test/statements/for/initializer/array_ctor.wgsl.expected.hlsl

@@ -0,0 +1,12 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+void f() {
+  const int tint_symbol[1] = {1};
+  {
+    [loop] for(int i = tint_symbol[0]; ; ) {
+    }
+  }
+}

test/statements/for/initializer/array_ctor.wgsl.expected.msl

@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  int arr[1];
+};
+
+void f() {
+  tint_array_wrapper const tint_symbol = {.arr={1}};
+  for(int i = tint_symbol.arr[0]; ; ) {
+  }
+}
+

test/statements/for/initializer/array_ctor.wgsl.expected.spvasm

@@ -0,0 +1,44 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 22
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %f "f"
+               OpName %i "i"
+               OpDecorate %_arr_int_uint_1 ArrayStride 4
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+%_arr_int_uint_1 = OpTypeArray %int %uint_1
+      %int_1 = OpConstant %int 1
+         %12 = OpConstantComposite %_arr_int_uint_1 %int_1
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+         %17 = OpConstantNull %int
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %1
+          %6 = OpLabel
+          %i = OpVariable %_ptr_Function_int Function %17
+         %14 = OpCompositeExtract %int %12 0
+               OpStore %i %14
+               OpBranch %18
+         %18 = OpLabel
+               OpLoopMerge %19 %20 None
+               OpBranch %21
+         %21 = OpLabel
+               OpBranch %20
+         %20 = OpLabel
+               OpBranch %18
+         %19 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/statements/for/initializer/array_ctor.wgsl.expected.wgsl

@@ -0,0 +1,4 @@
+fn f() {
+  for(var i : i32 = array<i32, 1>(1)[0]; ; ) {
+  }
+}

test/statements/for/initializer/basic.wgsl

@@ -0,0 +1,3 @@
+fn f() {
+  for (var i : i32 = 0;;) {}
+}

test/statements/for/initializer/struct_ctor.wgsl

@@ -0,0 +1,7 @@
+struct S {
+  i : i32;
+};
+
+fn f() {
+  for (var i : i32 = S(1).i;;) {}
+}

test/statements/for/initializer/struct_ctor.wgsl.expected.hlsl

@@ -0,0 +1,16 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+struct S {
+  int i;
+};
+
+void f() {
+  const S tint_symbol = {1};
+  {
+    [loop] for(int i = tint_symbol.i; ; ) {
+    }
+  }
+}

test/statements/for/initializer/struct_ctor.wgsl.expected.msl

@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  int i;
+};
+
+void f() {
+  S const tint_symbol = {.i=1};
+  for(int i = tint_symbol.i; ; ) {
+  }
+}
+

test/statements/for/initializer/struct_ctor.wgsl.expected.spvasm

@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 19
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %f "f"
+               OpName %S "S"
+               OpMemberName %S 0 "i"
+               OpName %i "i"
+               OpMemberDecorate %S 0 Offset 0
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %S = OpTypeStruct %int
+      %int_1 = OpConstant %int 1
+         %10 = OpConstantComposite %S %int_1
+%_ptr_Function_int = OpTypePointer Function %int
+         %14 = OpConstantNull %int
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %void None %1
+          %6 = OpLabel
+          %i = OpVariable %_ptr_Function_int Function %14
+         %11 = OpCompositeExtract %int %10 0
+               OpStore %i %11
+               OpBranch %15
+         %15 = OpLabel
+               OpLoopMerge %16 %17 None
+               OpBranch %18
+         %18 = OpLabel
+               OpBranch %17
+         %17 = OpLabel
+               OpBranch %15
+         %16 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/statements/for/initializer/struct_ctor.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+struct S {
+  i : i32;
+};
+
+fn f() {
+  for(var i : i32 = S(1).i; ; ) {
+  }
+}