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>
diff --git a/docs/origin-trial-changes.md b/docs/origin-trial-changes.md
index 64c5a3e..e83ef1b 100644
--- a/docs/origin-trial-changes.md
+++ b/docs/origin-trial-changes.md
@@ -15,6 +15,10 @@
* 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
diff --git a/src/resolver/resolver_behavior_test.cc b/src/resolver/resolver_behavior_test.cc
index 9099dd7..9c5da9a 100644
--- a/src/resolver/resolver_behavior_test.cc
+++ b/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 {
diff --git a/src/sem/type_mappings.h b/src/sem/type_mappings.h
index ecc6ea3..3fb8fe8 100644
--- a/src/sem/type_mappings.h
+++ b/src/sem/type_mappings.h
@@ -23,6 +23,7 @@
namespace ast {
class CallExpression;
class Expression;
+class ForLoopStatement;
class Function;
class MemberAccessorExpression;
class Node;
@@ -39,6 +40,7 @@
class Array;
class Call;
class Expression;
+class ForLoopStatement;
class Function;
class MemberAccessorExpression;
class Node;
@@ -56,6 +58,7 @@
//! @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*);
diff --git a/src/transform/promote_initializers_to_const_var.cc b/src/transform/promote_initializers_to_const_var.cc
index 9a9001a..e672947 100644
--- a/src/transform/promote_initializers_to_const_var.cc
+++ b/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 @@
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 @@
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 @@
// 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()) {
- if (auto* src_init = src_node->As<ast::CallExpression>()) {
- auto* call = ctx.src->Sem().Get(src_init);
- if (!call->Target()->Is<sem::TypeConstructor>()) {
+ // For-loops that need to be decomposed to loops.
+ std::unordered_map<const sem::ForLoopStatement*, LoopInfo> loops;
+
+ 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();
- if (!src_sem_stmt) {
+ auto* sem_stmt = ctor->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 (src_var_decl->variable->constructor == src_init) {
+ if (auto* src_var_decl = stmt->As<ast::VariableDeclStatement>()) {
+ 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 @@
}
}
- 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
- auto dst_symbol = ctx.dst->Sym();
- // Clone the type
- auto* dst_ty = CreateASTTypeFor(ctx, call->Type());
- // Clone the initializer
- auto* dst_init = ctx.Clone(src_init);
- // Construct the constant that holds the hoisted initializer
- auto* dst_var = ctx.dst->Const(dst_symbol, dst_ty, dst_init);
- // 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);
+ // Create a new symbol for the let
+ auto name = ctx.dst->Sym();
+ // Construct the let that holds the hoisted initializer
+ auto* let = ctx.dst->Const(name, nullptr, ctx.Clone(expr));
+ // Construct the let declaration statement
+ auto* let_decl = ctx.dst->Decl(let);
+ // Replace the initializer expression with a reference to the let
+ ctx.Replace(expr, ctx.dst->Expr(name));
- // Insert the constant before the usage
- ctx.InsertBefore(src_sem_stmt->Block()->Declaration()->statements,
- src_stmt, dst_var_decl);
- // Replace the inlined initializer with a reference to the constant
- ctx.Replace(src_init, dst_ident);
+ if (auto* fl = sem_stmt->As<sem::ForLoopStatement>()) {
+ // Expression used in for-loop condition.
+ // For-loop needs to be decomposed to a loop.
+ loops[fl].cond_decls.emplace_back(let_decl);
+ 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();
}
diff --git a/src/transform/promote_initializers_to_const_var.h b/src/transform/promote_initializers_to_const_var.h
index 59adfb7..b6150d7 100644
--- a/src/transform/promote_initializers_to_const_var.h
+++ b/src/transform/promote_initializers_to_const_var.h
@@ -21,8 +21,10 @@
namespace transform {
/// A transform that hoists the array and structure initializers to a constant
-/// variable, declared just before the statement of usage. See
-/// crbug.com/tint/406 for more details.
+/// variable, declared just before the statement of usage. This transform may
+/// 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:
diff --git a/src/transform/promote_initializers_to_const_var_test.cc b/src/transform/promote_initializers_to_const_var_test.cc
index 2af5404..c303fce 100644
--- a/src/transform/promote_initializers_to_const_var_test.cc
+++ b/src/transform/promote_initializers_to_const_var_test.cc
@@ -24,25 +24,23 @@
TEST_F(PromoteInitializersToConstVarTest, BasicArray) {
auto* src = R"(
-[[stage(compute), workgroup_size(1)]]
-fn main() {
- var f0 : f32 = 1.0;
- var f1 : f32 = 2.0;
- var f2 : f32 = 3.0;
- var f3 : f32 = 4.0;
- var i : f32 = array<f32, 4u>(f0, f1, f2, f3)[2];
+fn f() {
+ var f0 = 1.0;
+ var f1 = 2.0;
+ var f2 = 3.0;
+ var f3 = 4.0;
+ var i = array<f32, 4u>(f0, f1, f2, f3)[2];
}
)";
auto* expect = R"(
-[[stage(compute), workgroup_size(1)]]
-fn main() {
- var f0 : f32 = 1.0;
- var f1 : f32 = 2.0;
- var f2 : f32 = 3.0;
- var f3 : f32 = 4.0;
- let tint_symbol : array<f32, 4u> = array<f32, 4u>(f0, f1, f2, f3);
- var i : f32 = tint_symbol[2];
+fn f() {
+ var f0 = 1.0;
+ var f1 = 2.0;
+ var f2 = 3.0;
+ var f3 = 4.0;
+ let tint_symbol = array<f32, 4u>(f0, f1, f2, f3);
+ var i = tint_symbol[2];
}
)";
@@ -59,9 +57,8 @@
c : vec3<f32>;
};
-[[stage(compute), workgroup_size(1)]]
-fn main() {
- var x : f32 = S(1, 2.0, vec3<f32>()).b;
+fn f() {
+ var x = S(1, 2.0, vec3<f32>()).b;
}
)";
@@ -72,10 +69,168 @@
c : vec3<f32>;
};
-[[stage(compute), workgroup_size(1)]]
-fn main() {
- let tint_symbol : S = S(1, 2.0, vec3<f32>());
- var x : f32 = tint_symbol.b;
+fn f() {
+ let tint_symbol = S(1, 2.0, vec3<f32>());
+ var x = 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 @@
TEST_F(PromoteInitializersToConstVarTest, ArrayInArrayArray) {
auto* src = R"(
-[[stage(compute), workgroup_size(1)]]
-fn main() {
- var i : f32 = array<array<f32, 2u>, 2u>(array<f32, 2u>(1.0, 2.0), array<f32, 2u>(3.0, 4.0))[0][1];
+fn f() {
+ var i = 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 main() {
- let tint_symbol : array<f32, 2u> = array<f32, 2u>(1.0, 2.0);
- let tint_symbol_1 : array<f32, 2u> = array<f32, 2u>(3.0, 4.0);
- let tint_symbol_2 : array<array<f32, 2u>, 2u> = array<array<f32, 2u>, 2u>(tint_symbol, tint_symbol_1);
- var i : f32 = tint_symbol_2[0][1];
+fn f() {
+ let tint_symbol = array<f32, 2u>(1.0, 2.0);
+ let tint_symbol_1 = array<f32, 2u>(3.0, 4.0);
+ let tint_symbol_2 = array<array<f32, 2u>, 2u>(tint_symbol, tint_symbol_1);
+ var i = tint_symbol_2[0][1];
}
)";
@@ -123,9 +276,8 @@
a : S2;
};
-[[stage(compute), workgroup_size(1)]]
-fn main() {
- var x : i32 = S3(S2(1, S1(2), 3)).a.b.a;
+fn f() {
+ var x = S3(S2(1, S1(2), 3)).a.b.a;
}
)";
@@ -144,12 +296,11 @@
a : S2;
};
-[[stage(compute), workgroup_size(1)]]
-fn main() {
- let tint_symbol : S1 = S1(2);
- let tint_symbol_1 : S2 = S2(1, tint_symbol, 3);
- let tint_symbol_2 : S3 = S3(tint_symbol_1);
- var x : i32 = tint_symbol_2.a.b.a;
+fn f() {
+ let tint_symbol = S1(2);
+ let tint_symbol_1 = S2(1, tint_symbol, 3);
+ let tint_symbol_2 = S3(tint_symbol_1);
+ var x = tint_symbol_2.a.b.a;
}
)";
@@ -168,9 +319,8 @@
a : array<S1, 3u>;
};
-[[stage(compute), workgroup_size(1)]]
-fn main() {
- var x : i32 = S2(array<S1, 3u>(S1(1), S1(2), S1(3))).a[1].a;
+fn f() {
+ var x = S2(array<S1, 3u>(S1(1), S1(2), S1(3))).a[1].a;
}
)";
@@ -183,14 +333,13 @@
a : array<S1, 3u>;
};
-[[stage(compute), workgroup_size(1)]]
-fn main() {
- let tint_symbol : S1 = S1(1);
- let tint_symbol_1 : S1 = S1(2);
- let tint_symbol_2 : S1 = S1(3);
- let tint_symbol_3 : array<S1, 3u> = array<S1, 3u>(tint_symbol, tint_symbol_1, tint_symbol_2);
- let tint_symbol_4 : S2 = S2(tint_symbol_3);
- var x : i32 = tint_symbol_4.a[1].a;
+fn f() {
+ let tint_symbol = S1(1);
+ let tint_symbol_1 = S1(2);
+ let tint_symbol_2 = S1(3);
+ let tint_symbol_3 = array<S1, 3u>(tint_symbol, tint_symbol_1, tint_symbol_2);
+ let tint_symbol_4 = S2(tint_symbol_3);
+ var x = tint_symbol_4.a[1].a;
}
)";
@@ -207,10 +356,9 @@
c : i32;
};
-[[stage(compute), workgroup_size(1)]]
-fn main() {
- var local_arr : array<f32, 4u> = array<f32, 4u>(0.0, 1.0, 2.0, 3.0);
- var local_str : S = S(1, 2.0, 3);
+fn f() {
+ var local_arr = array<f32, 4u>(0.0, 1.0, 2.0, 3.0);
+ var local_str = S(1, 2.0, 3);
}
let module_arr : array<f32, 4u> = array<f32, 4u>(0.0, 1.0, 2.0, 3.0);
diff --git a/test/statements/for/condition.wgsl b/test/statements/for/condition.wgsl
deleted file mode 100644
index 7b577dc..0000000
--- a/test/statements/for/condition.wgsl
+++ /dev/null
@@ -1,5 +0,0 @@
-fn f() {
- var i : i32;
- for (;i < 4;) {
- }
-}
diff --git a/test/statements/for/condition/array_ctor.wgsl b/test/statements/for/condition/array_ctor.wgsl
new file mode 100644
index 0000000..b6ff835
--- /dev/null
+++ b/test/statements/for/condition/array_ctor.wgsl
@@ -0,0 +1,4 @@
+fn f() {
+ var i : i32;
+ for (;i < array<i32, 1>(1)[0];) {}
+}
diff --git a/test/statements/for/initializer.wgsl.expected.glsl b/test/statements/for/condition/array_ctor.wgsl.expected.glsl
similarity index 100%
rename from test/statements/for/initializer.wgsl.expected.glsl
rename to test/statements/for/condition/array_ctor.wgsl.expected.glsl
diff --git a/test/statements/for/condition/array_ctor.wgsl.expected.hlsl b/test/statements/for/condition/array_ctor.wgsl.expected.hlsl
new file mode 100644
index 0000000..409e7c4
--- /dev/null
+++ b/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;
+ }
+ }
+}
diff --git a/test/statements/for/condition/array_ctor.wgsl.expected.msl b/test/statements/for/condition/array_ctor.wgsl.expected.msl
new file mode 100644
index 0000000..2ec1a68
--- /dev/null
+++ b/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;
+ }
+ }
+}
+
diff --git a/test/statements/for/condition/array_ctor.wgsl.expected.spvasm b/test/statements/for/condition/array_ctor.wgsl.expected.spvasm
new file mode 100644
index 0000000..b86a19f
--- /dev/null
+++ b/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
diff --git a/test/statements/for/condition/array_ctor.wgsl.expected.wgsl b/test/statements/for/condition/array_ctor.wgsl.expected.wgsl
new file mode 100644
index 0000000..a21e739
--- /dev/null
+++ b/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]); ) {
+ }
+}
diff --git a/test/statements/for/condition/basic.wgsl b/test/statements/for/condition/basic.wgsl
new file mode 100644
index 0000000..a897474
--- /dev/null
+++ b/test/statements/for/condition/basic.wgsl
@@ -0,0 +1,5 @@
+fn f() {
+ var i : i32;
+ for (;i < 4;) {
+ }
+}
diff --git a/test/statements/for/condition.wgsl.expected.glsl b/test/statements/for/condition/basic.wgsl.expected.glsl
similarity index 100%
rename from test/statements/for/condition.wgsl.expected.glsl
rename to test/statements/for/condition/basic.wgsl.expected.glsl
diff --git a/test/statements/for/condition.wgsl.expected.hlsl b/test/statements/for/condition/basic.wgsl.expected.hlsl
similarity index 100%
rename from test/statements/for/condition.wgsl.expected.hlsl
rename to test/statements/for/condition/basic.wgsl.expected.hlsl
diff --git a/test/statements/for/condition.wgsl.expected.msl b/test/statements/for/condition/basic.wgsl.expected.msl
similarity index 100%
rename from test/statements/for/condition.wgsl.expected.msl
rename to test/statements/for/condition/basic.wgsl.expected.msl
diff --git a/test/statements/for/condition.wgsl.expected.spvasm b/test/statements/for/condition/basic.wgsl.expected.spvasm
similarity index 100%
rename from test/statements/for/condition.wgsl.expected.spvasm
rename to test/statements/for/condition/basic.wgsl.expected.spvasm
diff --git a/test/statements/for/condition.wgsl.expected.wgsl b/test/statements/for/condition/basic.wgsl.expected.wgsl
similarity index 100%
rename from test/statements/for/condition.wgsl.expected.wgsl
rename to test/statements/for/condition/basic.wgsl.expected.wgsl
diff --git a/test/statements/for/condition/struct_ctor.wgsl b/test/statements/for/condition/struct_ctor.wgsl
new file mode 100644
index 0000000..1d2396b
--- /dev/null
+++ b/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;) {}
+}
diff --git a/test/statements/for/initializer.wgsl.expected.glsl b/test/statements/for/condition/struct_ctor.wgsl.expected.glsl
similarity index 100%
copy from test/statements/for/initializer.wgsl.expected.glsl
copy to test/statements/for/condition/struct_ctor.wgsl.expected.glsl
diff --git a/test/statements/for/condition/struct_ctor.wgsl.expected.hlsl b/test/statements/for/condition/struct_ctor.wgsl.expected.hlsl
new file mode 100644
index 0000000..c9d9ff0
--- /dev/null
+++ b/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;
+ }
+ }
+}
diff --git a/test/statements/for/condition/struct_ctor.wgsl.expected.msl b/test/statements/for/condition/struct_ctor.wgsl.expected.msl
new file mode 100644
index 0000000..a370f25
--- /dev/null
+++ b/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;
+ }
+ }
+}
+
diff --git a/test/statements/for/condition/struct_ctor.wgsl.expected.spvasm b/test/statements/for/condition/struct_ctor.wgsl.expected.spvasm
new file mode 100644
index 0000000..0ad5b63
--- /dev/null
+++ b/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
diff --git a/test/statements/for/condition/struct_ctor.wgsl.expected.wgsl b/test/statements/for/condition/struct_ctor.wgsl.expected.wgsl
new file mode 100644
index 0000000..dd9b9c5
--- /dev/null
+++ b/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); ) {
+ }
+}
diff --git a/test/statements/for/continuing.wgsl b/test/statements/for/continuing.wgsl
deleted file mode 100644
index e174e9a..0000000
--- a/test/statements/for/continuing.wgsl
+++ /dev/null
@@ -1,4 +0,0 @@
-fn f() {
- var i : i32;
- for (;;i = i + 1) {}
-}
diff --git a/test/statements/for/continuing.wgsl.expected.glsl b/test/statements/for/continuing.wgsl.expected.glsl
deleted file mode 100644
index e69de29..0000000
--- a/test/statements/for/continuing.wgsl.expected.glsl
+++ /dev/null
diff --git a/test/statements/for/continuing/array_ctor.wgsl b/test/statements/for/continuing/array_ctor.wgsl
new file mode 100644
index 0000000..821777d
--- /dev/null
+++ b/test/statements/for/continuing/array_ctor.wgsl
@@ -0,0 +1,4 @@
+fn f() {
+ var i : i32;
+ for (;;i = i + array<i32, 1>(1)[0]) {}
+}
diff --git a/test/statements/for/condition.wgsl.expected.glsl b/test/statements/for/continuing/array_ctor.wgsl.expected.glsl
similarity index 100%
copy from test/statements/for/condition.wgsl.expected.glsl
copy to test/statements/for/continuing/array_ctor.wgsl.expected.glsl
diff --git a/test/statements/for/continuing/array_ctor.wgsl.expected.hlsl b/test/statements/for/continuing/array_ctor.wgsl.expected.hlsl
new file mode 100644
index 0000000..df66919
--- /dev/null
+++ b/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]);
+ }
+ }
+}
diff --git a/test/statements/for/continuing/array_ctor.wgsl.expected.msl b/test/statements/for/continuing/array_ctor.wgsl.expected.msl
new file mode 100644
index 0000000..cb9f3f6
--- /dev/null
+++ b/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])));
+ }
+ }
+}
+
diff --git a/test/statements/for/continuing/array_ctor.wgsl.expected.spvasm b/test/statements/for/continuing/array_ctor.wgsl.expected.spvasm
new file mode 100644
index 0000000..8c32307
--- /dev/null
+++ b/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
diff --git a/test/statements/for/continuing/array_ctor.wgsl.expected.wgsl b/test/statements/for/continuing/array_ctor.wgsl.expected.wgsl
new file mode 100644
index 0000000..21d2c5a
--- /dev/null
+++ b/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])) {
+ }
+}
diff --git a/test/statements/for/continuing/basic.wgsl b/test/statements/for/continuing/basic.wgsl
new file mode 100644
index 0000000..745661c
--- /dev/null
+++ b/test/statements/for/continuing/basic.wgsl
@@ -0,0 +1,4 @@
+fn f() {
+ var i : i32;
+ for (;;i = i + 1) {}
+}
diff --git a/test/statements/for/condition.wgsl.expected.glsl b/test/statements/for/continuing/basic.wgsl.expected.glsl
similarity index 100%
copy from test/statements/for/condition.wgsl.expected.glsl
copy to test/statements/for/continuing/basic.wgsl.expected.glsl
diff --git a/test/statements/for/continuing.wgsl.expected.hlsl b/test/statements/for/continuing/basic.wgsl.expected.hlsl
similarity index 100%
rename from test/statements/for/continuing.wgsl.expected.hlsl
rename to test/statements/for/continuing/basic.wgsl.expected.hlsl
diff --git a/test/statements/for/continuing.wgsl.expected.msl b/test/statements/for/continuing/basic.wgsl.expected.msl
similarity index 100%
rename from test/statements/for/continuing.wgsl.expected.msl
rename to test/statements/for/continuing/basic.wgsl.expected.msl
diff --git a/test/statements/for/continuing.wgsl.expected.spvasm b/test/statements/for/continuing/basic.wgsl.expected.spvasm
similarity index 100%
rename from test/statements/for/continuing.wgsl.expected.spvasm
rename to test/statements/for/continuing/basic.wgsl.expected.spvasm
diff --git a/test/statements/for/continuing.wgsl.expected.wgsl b/test/statements/for/continuing/basic.wgsl.expected.wgsl
similarity index 100%
rename from test/statements/for/continuing.wgsl.expected.wgsl
rename to test/statements/for/continuing/basic.wgsl.expected.wgsl
diff --git a/test/statements/for/continuing/struct_ctor.wgsl b/test/statements/for/continuing/struct_ctor.wgsl
new file mode 100644
index 0000000..cf02a48
--- /dev/null
+++ b/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) {}
+}
diff --git a/test/statements/for/initializer.wgsl.expected.glsl b/test/statements/for/continuing/struct_ctor.wgsl.expected.glsl
similarity index 100%
copy from test/statements/for/initializer.wgsl.expected.glsl
copy to test/statements/for/continuing/struct_ctor.wgsl.expected.glsl
diff --git a/test/statements/for/continuing/struct_ctor.wgsl.expected.hlsl b/test/statements/for/continuing/struct_ctor.wgsl.expected.hlsl
new file mode 100644
index 0000000..1612668
--- /dev/null
+++ b/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);
+ }
+ }
+ }
+}
diff --git a/test/statements/for/continuing/struct_ctor.wgsl.expected.msl b/test/statements/for/continuing/struct_ctor.wgsl.expected.msl
new file mode 100644
index 0000000..d19c0d2
--- /dev/null
+++ b/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)));
+ }
+ }
+ }
+}
+
diff --git a/test/statements/for/continuing/struct_ctor.wgsl.expected.spvasm b/test/statements/for/continuing/struct_ctor.wgsl.expected.spvasm
new file mode 100644
index 0000000..a7942c3
--- /dev/null
+++ b/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
diff --git a/test/statements/for/continuing/struct_ctor.wgsl.expected.wgsl b/test/statements/for/continuing/struct_ctor.wgsl.expected.wgsl
new file mode 100644
index 0000000..c51ce52
--- /dev/null
+++ b/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)) {
+ }
+}
diff --git a/test/statements/for/initializer.wgsl b/test/statements/for/initializer.wgsl
deleted file mode 100644
index 5f6ef9c..0000000
--- a/test/statements/for/initializer.wgsl
+++ /dev/null
@@ -1,3 +0,0 @@
-fn f() {
- for (var i : i32 = 0;;) {}
-}
diff --git a/test/statements/for/initializer/array_ctor.wgsl b/test/statements/for/initializer/array_ctor.wgsl
new file mode 100644
index 0000000..06f1d9b
--- /dev/null
+++ b/test/statements/for/initializer/array_ctor.wgsl
@@ -0,0 +1,3 @@
+fn f() {
+ for (var i : i32 = array<i32, 1>(1)[0];;) {}
+}
diff --git a/test/statements/for/initializer.wgsl.expected.glsl b/test/statements/for/initializer/array_ctor.wgsl.expected.glsl
similarity index 100%
copy from test/statements/for/initializer.wgsl.expected.glsl
copy to test/statements/for/initializer/array_ctor.wgsl.expected.glsl
diff --git a/test/statements/for/initializer/array_ctor.wgsl.expected.hlsl b/test/statements/for/initializer/array_ctor.wgsl.expected.hlsl
new file mode 100644
index 0000000..4f60bdb
--- /dev/null
+++ b/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]; ; ) {
+ }
+ }
+}
diff --git a/test/statements/for/initializer/array_ctor.wgsl.expected.msl b/test/statements/for/initializer/array_ctor.wgsl.expected.msl
new file mode 100644
index 0000000..d4427b2
--- /dev/null
+++ b/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]; ; ) {
+ }
+}
+
diff --git a/test/statements/for/initializer/array_ctor.wgsl.expected.spvasm b/test/statements/for/initializer/array_ctor.wgsl.expected.spvasm
new file mode 100644
index 0000000..5587397
--- /dev/null
+++ b/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
diff --git a/test/statements/for/initializer/array_ctor.wgsl.expected.wgsl b/test/statements/for/initializer/array_ctor.wgsl.expected.wgsl
new file mode 100644
index 0000000..91b2f5b
--- /dev/null
+++ b/test/statements/for/initializer/array_ctor.wgsl.expected.wgsl
@@ -0,0 +1,4 @@
+fn f() {
+ for(var i : i32 = array<i32, 1>(1)[0]; ; ) {
+ }
+}
diff --git a/test/statements/for/initializer/basic.wgsl b/test/statements/for/initializer/basic.wgsl
new file mode 100644
index 0000000..cf5a1fd
--- /dev/null
+++ b/test/statements/for/initializer/basic.wgsl
@@ -0,0 +1,3 @@
+fn f() {
+ for (var i : i32 = 0;;) {}
+}
diff --git a/test/statements/for/condition.wgsl.expected.glsl b/test/statements/for/initializer/basic.wgsl.expected.glsl
similarity index 100%
copy from test/statements/for/condition.wgsl.expected.glsl
copy to test/statements/for/initializer/basic.wgsl.expected.glsl
diff --git a/test/statements/for/initializer.wgsl.expected.hlsl b/test/statements/for/initializer/basic.wgsl.expected.hlsl
similarity index 100%
rename from test/statements/for/initializer.wgsl.expected.hlsl
rename to test/statements/for/initializer/basic.wgsl.expected.hlsl
diff --git a/test/statements/for/initializer.wgsl.expected.msl b/test/statements/for/initializer/basic.wgsl.expected.msl
similarity index 100%
rename from test/statements/for/initializer.wgsl.expected.msl
rename to test/statements/for/initializer/basic.wgsl.expected.msl
diff --git a/test/statements/for/initializer.wgsl.expected.spvasm b/test/statements/for/initializer/basic.wgsl.expected.spvasm
similarity index 100%
rename from test/statements/for/initializer.wgsl.expected.spvasm
rename to test/statements/for/initializer/basic.wgsl.expected.spvasm
diff --git a/test/statements/for/initializer.wgsl.expected.wgsl b/test/statements/for/initializer/basic.wgsl.expected.wgsl
similarity index 100%
rename from test/statements/for/initializer.wgsl.expected.wgsl
rename to test/statements/for/initializer/basic.wgsl.expected.wgsl
diff --git a/test/statements/for/initializer/struct_ctor.wgsl b/test/statements/for/initializer/struct_ctor.wgsl
new file mode 100644
index 0000000..5b0b5ba
--- /dev/null
+++ b/test/statements/for/initializer/struct_ctor.wgsl
@@ -0,0 +1,7 @@
+struct S {
+ i : i32;
+};
+
+fn f() {
+ for (var i : i32 = S(1).i;;) {}
+}
diff --git a/test/statements/for/initializer.wgsl.expected.glsl b/test/statements/for/initializer/struct_ctor.wgsl.expected.glsl
similarity index 100%
copy from test/statements/for/initializer.wgsl.expected.glsl
copy to test/statements/for/initializer/struct_ctor.wgsl.expected.glsl
diff --git a/test/statements/for/initializer/struct_ctor.wgsl.expected.hlsl b/test/statements/for/initializer/struct_ctor.wgsl.expected.hlsl
new file mode 100644
index 0000000..9cbe791
--- /dev/null
+++ b/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; ; ) {
+ }
+ }
+}
diff --git a/test/statements/for/initializer/struct_ctor.wgsl.expected.msl b/test/statements/for/initializer/struct_ctor.wgsl.expected.msl
new file mode 100644
index 0000000..f0ebc83
--- /dev/null
+++ b/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; ; ) {
+ }
+}
+
diff --git a/test/statements/for/initializer/struct_ctor.wgsl.expected.spvasm b/test/statements/for/initializer/struct_ctor.wgsl.expected.spvasm
new file mode 100644
index 0000000..1cabffb
--- /dev/null
+++ b/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
diff --git a/test/statements/for/initializer/struct_ctor.wgsl.expected.wgsl b/test/statements/for/initializer/struct_ctor.wgsl.expected.wgsl
new file mode 100644
index 0000000..9e32474
--- /dev/null
+++ b/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; ; ) {
+ }
+}
