tint/hlsl: fix RemoveContinueInSwitch transform to handle nested switch statements
Rework the transform so that it handles a continue within a switch
nested within switches within a loop. Also make the transform only
generate a single bool variable per loop, rather than per continue
statement.
Bug: tint:1080
Change-Id: I626cf4aa35e900ccc8ad08712ff1393c91587f10
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/151041
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Ben Clayton <bclayton@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch.cc b/src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch.cc
index b091b91..4f1808b 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch.cc
@@ -15,7 +15,6 @@
#include "src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch.h"
#include <string>
-#include <unordered_map>
#include <utility>
#include "src/tint/lang/wgsl/ast/continue_statement.h"
@@ -25,11 +24,10 @@
#include "src/tint/lang/wgsl/program/program_builder.h"
#include "src/tint/lang/wgsl/resolver/resolve.h"
#include "src/tint/lang/wgsl/sem/block_statement.h"
-#include "src/tint/lang/wgsl/sem/for_loop_statement.h"
#include "src/tint/lang/wgsl/sem/loop_statement.h"
#include "src/tint/lang/wgsl/sem/switch_statement.h"
-#include "src/tint/lang/wgsl/sem/while_statement.h"
-#include "src/tint/utils/containers/map.h"
+#include "src/tint/utils/containers/hashmap.h"
+#include "src/tint/utils/containers/hashset.h"
TINT_INSTANTIATE_TYPEINFO(tint::hlsl::writer::RemoveContinueInSwitch);
@@ -44,8 +42,7 @@
/// Runs the transform
/// @returns the new program or SkipTransform if the transform is not required
ApplyResult Run() {
- bool made_changes = false;
-
+ // First collect all switch statements within loops that contain a continue statement.
for (auto* node : src->ASTNodes().Objects()) {
auto* cont = node->As<ast::ContinueStatement>();
if (!cont) {
@@ -58,37 +55,71 @@
continue;
}
- made_changes = true;
-
- auto cont_var_name = tint::GetOrCreate(switch_to_cont_var_name, switch_stmt, [&] {
- // Create and insert 'var tint_continue : bool = false;' before the
- // switch.
- auto var_name = b.Symbols().New("tint_continue");
- auto* decl = b.Decl(b.Var(var_name, b.ty.bool_(), b.Expr(false)));
- auto ip = ast::transform::utils::GetInsertionPoint(ctx, switch_stmt);
- ctx.InsertBefore(ip.first->Declaration()->statements, ip.second, decl);
-
- // Create and insert 'if (tint_continue) { continue; }' after
- // switch.
- auto* if_stmt = b.If(b.Expr(var_name), b.Block(b.Continue()));
- ctx.InsertAfter(ip.first->Declaration()->statements, ip.second, if_stmt);
-
- // Return the new var name
- return var_name;
+ auto& info = switch_infos.GetOrCreate(switch_stmt, [&] {
+ switch_stmts.Push(switch_stmt);
+ auto* block = sem.Get(switch_stmt)->FindFirstParent<sem::LoopBlockStatement>();
+ return SwitchInfo{/* loop_block */ block, /* continues */ Empty};
});
-
- // Replace 'continue;' with '{ tint_continue = true; break; }'
- auto* new_stmt = b.Block( //
- b.Assign(b.Expr(cont_var_name), true), //
- b.Break());
-
- ctx.Replace(cont, new_stmt);
+ info.continues.Push(cont);
}
- if (!made_changes) {
+ if (switch_stmts.IsEmpty()) {
return SkipTransform;
}
+ // For each switch statement:
+ // 1. Declare a 'tint_continue' var just before the parent loop, and reset it to false at
+ // the top of the loop body.
+ // 2. Replace 'continue' with 'tint_continue = true; break;'
+ // 3. Insert 'if (tint_continue) { break; }' after switch, and all parent switches, except
+ // for the parent-most, for which we insert 'if (tint_continue) { continue; }'
+ for (auto* switch_stmt : switch_stmts) {
+ const auto& info = switch_infos.Get(switch_stmt);
+
+ auto var_name = loop_to_var.GetOrCreate(info->loop_block, [&] {
+ // Create and insert 'var tint_continue : bool;' before loop
+ auto var = b.Symbols().New("tint_continue");
+ auto* decl = b.Decl(b.Var(var, b.ty.bool_()));
+ auto ip = ast::transform::utils::GetInsertionPoint(
+ ctx, info->loop_block->Parent()->Declaration());
+ ctx.InsertBefore(ip.first->Declaration()->statements, ip.second, decl);
+
+ // Insert 'tint_continue = false' at top of loop body
+ auto assign_false = b.Assign(var, false);
+ ctx.InsertFront(info->loop_block->Declaration()->statements, assign_false);
+
+ return var;
+ });
+
+ for (auto& c : info->continues) {
+ // Replace 'continue;' with 'tint_continue = true; break;'
+ ctx.Replace(c, b.Assign(b.Expr(var_name), true));
+ ctx.InsertAfter(sem.Get(c)->Block()->Declaration()->statements, c, b.Break());
+ }
+
+ // Insert 'if (tint_continue) { break; }' after switch, and all parent switches,
+ // except for the parent-most, for which we insert 'if (tint_continue) { continue; }'
+ auto* curr_switch = switch_stmt;
+ while (curr_switch) {
+ auto* curr_switch_sem = sem.Get(curr_switch);
+ auto* parent = curr_switch_sem->Parent()->Declaration();
+ auto* next_switch = GetParentSwitchInLoop(sem, parent);
+
+ if (switch_handles_continue.Add(curr_switch)) {
+ const ast::IfStatement* if_stmt = nullptr;
+ if (next_switch) {
+ if_stmt = b.If(b.Expr(var_name), b.Block(b.Break()));
+ } else {
+ if_stmt = b.If(b.Expr(var_name), b.Block(b.Continue()));
+ }
+ ctx.InsertAfter(curr_switch_sem->Block()->Declaration()->statements,
+ curr_switch, if_stmt);
+ }
+
+ curr_switch = next_switch;
+ }
+ }
+
ctx.Clone();
return resolver::Resolve(b);
}
@@ -103,17 +134,38 @@
/// Alias to src->sem
const sem::Info& sem = src->Sem();
- // Map of switch statement to 'tint_continue' variable.
- std::unordered_map<const ast::SwitchStatement*, Symbol> switch_to_cont_var_name;
+ // Vector of switch statements within a loop that contains at least one continue statement.
+ Vector<const ast::SwitchStatement*, 4> switch_stmts;
- // If `cont` is within a switch statement within a loop, returns a pointer to
+ // Info for each switch statement within a loop that contains at least one continue statement.
+ struct SwitchInfo {
+ // Loop block containing this switch
+ const sem::LoopBlockStatement* loop_block;
+ // Continue statements within this switch
+ Vector<const ast::ContinueStatement*, 4> continues;
+ };
+
+ // Map of switch statements to per-switch info for switch statements within a loop that contains
+ // at least one continue statement.
+ Hashmap<const ast::SwitchStatement*, SwitchInfo, 4> switch_infos;
+
+ // Map of loop block statement to the single 'tint_continue' variable used to replace 'continue'
+ // control flow.
+ Hashmap<const sem::LoopBlockStatement*, Symbol, 4> loop_to_var;
+
+ // Set used to avoid duplicating 'if (tint_continue) { break/continue; }' after each switch
+ // within a loop.
+ Hashset<const ast::SwitchStatement*, 4> switch_handles_continue;
+
+ // If `stmt` is within a switch statement within a loop, returns a pointer to
// that switch statement.
static const ast::SwitchStatement* GetParentSwitchInLoop(const sem::Info& sem,
- const ast::ContinueStatement* cont) {
+ const ast::Statement* stmt) {
// Find whether first parent is a switch or a loop
- auto* sem_stmt = sem.Get(cont);
- auto* sem_parent = sem_stmt->FindFirstParent<sem::SwitchStatement, sem::LoopBlockStatement,
- sem::ForLoopStatement, sem::WhileStatement>();
+ auto* sem_stmt = sem.Get(stmt);
+ auto* sem_parent =
+ sem_stmt->FindFirstParent<sem::SwitchStatement, sem::LoopBlockStatement>();
+
if (!sem_parent) {
return nullptr;
}
diff --git a/src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch_test.cc b/src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch_test.cc
index 60153b1..e03e734 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch_test.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch_test.cc
@@ -28,7 +28,6 @@
switch(i) {
case 0: {
continue;
- break;
}
default: {
break;
@@ -116,7 +115,6 @@
switch(i) {
case 0: {
continue;
- break;
}
default: {
break;
@@ -135,15 +133,13 @@
auto* expect = R"(
fn f() {
var i = 0;
+ var tint_continue : bool;
loop {
+ tint_continue = false;
let marker1 = 0;
- var tint_continue : bool = false;
switch(i) {
case 0: {
- {
- tint_continue = true;
- break;
- }
+ tint_continue = true;
break;
}
default: {
@@ -178,15 +174,12 @@
switch(i) {
case 0: {
continue;
- break;
}
case 1: {
continue;
- break;
}
case 2: {
continue;
- break;
}
default: {
break;
@@ -205,29 +198,21 @@
auto* expect = R"(
fn f() {
var i = 0;
+ var tint_continue : bool;
loop {
+ tint_continue = false;
let marker1 = 0;
- var tint_continue : bool = false;
switch(i) {
case 0: {
- {
- tint_continue = true;
- break;
- }
+ tint_continue = true;
break;
}
case 1: {
- {
- tint_continue = true;
- break;
- }
+ tint_continue = true;
break;
}
case 2: {
- {
- tint_continue = true;
- break;
- }
+ tint_continue = true;
break;
}
default: {
@@ -262,7 +247,6 @@
switch(i) {
case 0: {
continue;
- break;
}
default: {
break;
@@ -274,7 +258,6 @@
switch(i) {
case 0: {
continue;
- break;
}
default: {
break;
@@ -290,15 +273,13 @@
auto* expect = R"(
fn f() {
var i = 0;
+ var tint_continue : bool;
loop {
+ tint_continue = false;
let marker1 = 0;
- var tint_continue : bool = false;
switch(i) {
case 0: {
- {
- tint_continue = true;
- break;
- }
+ tint_continue = true;
break;
}
default: {
@@ -310,20 +291,16 @@
}
let marker2 = 0;
let marker3 = 0;
- var tint_continue_1 : bool = false;
switch(i) {
case 0: {
- {
- tint_continue_1 = true;
- break;
- }
+ tint_continue = true;
break;
}
default: {
break;
}
}
- if (tint_continue_1) {
+ if (tint_continue) {
continue;
}
let marker4 = 0;
@@ -338,87 +315,253 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(RemoveContinueInSwitchTest, NestedLoopSwitch) {
+TEST_F(RemoveContinueInSwitchTest, NestedLoopSwitchSwitch) {
auto* src = R"(
fn f() {
- var i = 0;
- loop {
- let marker1 = 0;
+ var j = 0;
+ for (var i = 0; i < 2; i += 2) {
switch(i) {
case 0: {
- var j = 0;
- loop {
- let marker3 = 0;
- switch(j) {
- case 0: {
- continue;
- break;
- }
- default: {
- break;
- }
+ switch(j) {
+ case 0: {
+ continue;
}
- let marker4 = 0;
- break;
+ default: {
+ }
}
- continue;
- break;
}
default: {
- break;
}
}
- let marker2 = 0;
- break;
}
}
)";
auto* expect = R"(
fn f() {
- var i = 0;
- loop {
- let marker1 = 0;
- var tint_continue_1 : bool = false;
+ var j = 0;
+ var tint_continue : bool;
+ for(var i = 0; (i < 2); i += 2) {
+ tint_continue = false;
switch(i) {
case 0: {
- var j = 0;
- loop {
- let marker3 = 0;
- var tint_continue : bool = false;
+ switch(j) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ }
+ }
+ if (tint_continue) {
+ break;
+ }
+ }
+ default: {
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<RemoveContinueInSwitch>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(RemoveContinueInSwitchTest, NestedLoopLoopSwitch) {
+ auto* src = R"(
+fn f() {
+ for (var i = 0; i < 2; i += 2) {
+ for (var j = 0; j < 2; j += 2) {
+ switch(i) {
+ case 0: {
+ continue;
+ }
+ default: {
+ }
+ }
+ }
+ }
+}
+)";
+
+ auto* expect = R"(
+fn f() {
+ for(var i = 0; (i < 2); i += 2) {
+ var tint_continue : bool;
+ for(var j = 0; (j < 2); j += 2) {
+ tint_continue = false;
+ switch(i) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ }
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<RemoveContinueInSwitch>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(RemoveContinueInSwitchTest, NestedLoopSwitchLoopSwitch) {
+ auto* src = R"(
+fn f() {
+ for (var i = 0; i < 2; i += 2) {
+ switch(i) {
+ case 0: {
+ for (var j = 0; j < 2; j += 2) {
switch(j) {
case 0: {
- {
- tint_continue = true;
- break;
- }
+ continue; // j loop
+ }
+ default: {
+ }
+ }
+ }
+ continue; // i loop
+ }
+ default: {
+ }
+ }
+ }
+}
+)";
+
+ auto* expect = R"(
+fn f() {
+ var tint_continue_1 : bool;
+ for(var i = 0; (i < 2); i += 2) {
+ tint_continue_1 = false;
+ switch(i) {
+ case 0: {
+ var tint_continue : bool;
+ for(var j = 0; (j < 2); j += 2) {
+ tint_continue = false;
+ switch(j) {
+ case 0: {
+ tint_continue = true;
break;
}
default: {
- break;
}
}
if (tint_continue) {
continue;
}
- let marker4 = 0;
- break;
}
- {
- tint_continue_1 = true;
- break;
- }
+ tint_continue_1 = true;
break;
}
default: {
- break;
}
}
if (tint_continue_1) {
continue;
}
- let marker2 = 0;
- break;
+ }
+}
+)";
+
+ ast::transform::DataMap data;
+ auto got = Run<RemoveContinueInSwitch>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(RemoveContinueInSwitchTest, NestedLoopSwitchLoopSwitchSwitch) {
+ auto* src = R"(
+fn f() {
+ var k = 0;
+ for (var i = 0; i < 2; i += 2) {
+ switch(i) {
+ case 0: {
+ for (var j = 0; j < 2; j += 2) {
+ switch(j) {
+ case 0: {
+ continue; // j loop
+ }
+ case 1: {
+ switch (k) {
+ case 0: {
+ continue; // j loop
+ }
+ default: {
+ }
+ }
+ }
+ default: {
+ }
+ }
+ }
+ continue; // i loop
+ }
+ default: {
+ }
+ }
+ }
+}
+)";
+
+ auto* expect = R"(
+fn f() {
+ var k = 0;
+ var tint_continue_1 : bool;
+ for(var i = 0; (i < 2); i += 2) {
+ tint_continue_1 = false;
+ switch(i) {
+ case 0: {
+ var tint_continue : bool;
+ for(var j = 0; (j < 2); j += 2) {
+ tint_continue = false;
+ switch(j) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ case 1: {
+ switch(k) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ }
+ }
+ if (tint_continue) {
+ break;
+ }
+ }
+ default: {
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ tint_continue_1 = true;
+ break;
+ }
+ default: {
+ }
+ }
+ if (tint_continue_1) {
+ continue;
+ }
}
}
)";
@@ -469,19 +612,18 @@
var b = true;
var c = true;
var d = true;
+ var tint_continue : bool;
loop {
+ tint_continue = false;
if (a) {
if (b) {
let marker1 = 0;
- var tint_continue : bool = false;
switch(i) {
case 0: {
if (c) {
if (d) {
- {
- tint_continue = true;
- break;
- }
+ tint_continue = true;
+ break;
}
}
break;
@@ -529,15 +671,14 @@
auto* expect = R"(
fn f() {
+ var tint_continue : bool;
for(var i = 0; (i < 4); i = (i + 1)) {
+ tint_continue = false;
let marker1 = 0;
- var tint_continue : bool = false;
switch(i) {
case 0: {
- {
- tint_continue = true;
- break;
- }
+ tint_continue = true;
+ break;
break;
}
default: {
@@ -583,15 +724,14 @@
auto* expect = R"(
fn f() {
var i = 0;
+ var tint_continue : bool;
while((i < 4)) {
+ tint_continue = false;
let marker1 = 0;
- var tint_continue : bool = false;
switch(i) {
case 0: {
- {
- tint_continue = true;
- break;
- }
+ tint_continue = true;
+ break;
break;
}
default: {
diff --git a/test/tint/loops/continue_in_switch.wgsl.expected.dxc.hlsl b/test/tint/loops/continue_in_switch.wgsl.expected.dxc.hlsl
index ca926df..3ed9dc4 100644
--- a/test/tint/loops/continue_in_switch.wgsl.expected.dxc.hlsl
+++ b/test/tint/loops/continue_in_switch.wgsl.expected.dxc.hlsl
@@ -1,14 +1,12 @@
[numthreads(1, 1, 1)]
void f() {
+ bool tint_continue = false;
{
for(int i = 0; (i < 4); i = (i + 1)) {
- bool tint_continue = false;
+ tint_continue = false;
switch(i) {
case 0: {
- {
- tint_continue = true;
- break;
- }
+ tint_continue = true;
break;
}
default: {
diff --git a/test/tint/loops/continue_in_switch.wgsl.expected.fxc.hlsl b/test/tint/loops/continue_in_switch.wgsl.expected.fxc.hlsl
index ca926df..3ed9dc4 100644
--- a/test/tint/loops/continue_in_switch.wgsl.expected.fxc.hlsl
+++ b/test/tint/loops/continue_in_switch.wgsl.expected.fxc.hlsl
@@ -1,14 +1,12 @@
[numthreads(1, 1, 1)]
void f() {
+ bool tint_continue = false;
{
for(int i = 0; (i < 4); i = (i + 1)) {
- bool tint_continue = false;
+ tint_continue = false;
switch(i) {
case 0: {
- {
- tint_continue = true;
- break;
- }
+ tint_continue = true;
break;
}
default: {
diff --git a/test/tint/loops/multiple_continues.wgsl b/test/tint/loops/multiple_continues.wgsl
new file mode 100644
index 0000000..a57f773
--- /dev/null
+++ b/test/tint/loops/multiple_continues.wgsl
@@ -0,0 +1,18 @@
+@compute @workgroup_size(1)
+fn main() {
+ for (var i = 0; i < 2; i += 1) {
+ switch(i) {
+ case 0: {
+ continue;
+ }
+ case 1: {
+ continue;
+ }
+ case 2: {
+ continue;
+ }
+ default: {
+ }
+ }
+ }
+}
diff --git a/test/tint/loops/multiple_continues.wgsl.expected.dxc.hlsl b/test/tint/loops/multiple_continues.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..21f9f90
--- /dev/null
+++ b/test/tint/loops/multiple_continues.wgsl.expected.dxc.hlsl
@@ -0,0 +1,30 @@
+[numthreads(1, 1, 1)]
+void main() {
+ bool tint_continue = false;
+ {
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ tint_continue = false;
+ switch(i) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ case 1: {
+ tint_continue = true;
+ break;
+ }
+ case 2: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/multiple_continues.wgsl.expected.fxc.hlsl b/test/tint/loops/multiple_continues.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..21f9f90
--- /dev/null
+++ b/test/tint/loops/multiple_continues.wgsl.expected.fxc.hlsl
@@ -0,0 +1,30 @@
+[numthreads(1, 1, 1)]
+void main() {
+ bool tint_continue = false;
+ {
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ tint_continue = false;
+ switch(i) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ case 1: {
+ tint_continue = true;
+ break;
+ }
+ case 2: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/multiple_continues.wgsl.expected.glsl b/test/tint/loops/multiple_continues.wgsl.expected.glsl
new file mode 100644
index 0000000..36c86d8
--- /dev/null
+++ b/test/tint/loops/multiple_continues.wgsl.expected.glsl
@@ -0,0 +1,31 @@
+#version 310 es
+
+void tint_symbol() {
+ {
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ switch(i) {
+ case 0: {
+ continue;
+ break;
+ }
+ case 1: {
+ continue;
+ break;
+ }
+ case 2: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ }
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/loops/multiple_continues.wgsl.expected.msl b/test/tint/loops/multiple_continues.wgsl.expected.msl
new file mode 100644
index 0000000..ed8fcc1
--- /dev/null
+++ b/test/tint/loops/multiple_continues.wgsl.expected.msl
@@ -0,0 +1,26 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+ for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ switch(i) {
+ case 0: {
+ continue;
+ break;
+ }
+ case 1: {
+ continue;
+ break;
+ }
+ case 2: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ return;
+}
+
diff --git a/test/tint/loops/multiple_continues.wgsl.expected.spvasm b/test/tint/loops/multiple_continues.wgsl.expected.spvasm
new file mode 100644
index 0000000..4d6a668
--- /dev/null
+++ b/test/tint/loops/multiple_continues.wgsl.expected.spvasm
@@ -0,0 +1,57 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 29
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ OpName %i "i"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %6 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_2 = OpConstant %int 2
+ %bool = OpTypeBool
+ %int_1 = OpConstant %int 1
+ %main = OpFunction %void None %1
+ %4 = OpLabel
+ %i = OpVariable %_ptr_Function_int Function %6
+ OpStore %i %6
+ OpBranch %9
+ %9 = OpLabel
+ OpLoopMerge %10 %11 None
+ OpBranch %12
+ %12 = OpLabel
+ %14 = OpLoad %int %i
+ %16 = OpSLessThan %bool %14 %int_2
+ %13 = OpLogicalNot %bool %16
+ OpSelectionMerge %18 None
+ OpBranchConditional %13 %19 %18
+ %19 = OpLabel
+ OpBranch %10
+ %18 = OpLabel
+ %21 = OpLoad %int %i
+ OpSelectionMerge %20 None
+ OpSwitch %21 %22 0 %23 1 %24 2 %25
+ %23 = OpLabel
+ OpBranch %11
+ %24 = OpLabel
+ OpBranch %11
+ %25 = OpLabel
+ OpBranch %11
+ %22 = OpLabel
+ OpBranch %20
+ %20 = OpLabel
+ OpBranch %11
+ %11 = OpLabel
+ %26 = OpLoad %int %i
+ %28 = OpIAdd %int %26 %int_1
+ OpStore %i %28
+ OpBranch %9
+ %10 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/loops/multiple_continues.wgsl.expected.wgsl b/test/tint/loops/multiple_continues.wgsl.expected.wgsl
new file mode 100644
index 0000000..a80bcaf
--- /dev/null
+++ b/test/tint/loops/multiple_continues.wgsl.expected.wgsl
@@ -0,0 +1,18 @@
+@compute @workgroup_size(1)
+fn main() {
+ for(var i = 0; (i < 2); i += 1) {
+ switch(i) {
+ case 0: {
+ continue;
+ }
+ case 1: {
+ continue;
+ }
+ case 2: {
+ continue;
+ }
+ default: {
+ }
+ }
+ }
+}
diff --git a/test/tint/loops/multiple_switch.wgsl b/test/tint/loops/multiple_switch.wgsl
new file mode 100644
index 0000000..b3acc1f
--- /dev/null
+++ b/test/tint/loops/multiple_switch.wgsl
@@ -0,0 +1,21 @@
+@compute @workgroup_size(1)
+fn main() {
+ var i = 0;
+ for (var i = 0; i < 2; i += 1) {
+ switch(i) {
+ case 0: {
+ continue;
+ }
+ default: {
+ }
+ }
+
+ switch(i) {
+ case 0: {
+ continue;
+ }
+ default: {
+ }
+ }
+ }
+}
diff --git a/test/tint/loops/multiple_switch.wgsl.expected.dxc.hlsl b/test/tint/loops/multiple_switch.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..b61b6b5
--- /dev/null
+++ b/test/tint/loops/multiple_switch.wgsl.expected.dxc.hlsl
@@ -0,0 +1,35 @@
+[numthreads(1, 1, 1)]
+void main() {
+ int i = 0;
+ bool tint_continue = false;
+ {
+ for(int i_1 = 0; (i_1 < 2); i_1 = (i_1 + 1)) {
+ tint_continue = false;
+ switch(i_1) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ switch(i_1) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/multiple_switch.wgsl.expected.fxc.hlsl b/test/tint/loops/multiple_switch.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..b61b6b5
--- /dev/null
+++ b/test/tint/loops/multiple_switch.wgsl.expected.fxc.hlsl
@@ -0,0 +1,35 @@
+[numthreads(1, 1, 1)]
+void main() {
+ int i = 0;
+ bool tint_continue = false;
+ {
+ for(int i_1 = 0; (i_1 < 2); i_1 = (i_1 + 1)) {
+ tint_continue = false;
+ switch(i_1) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ switch(i_1) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/multiple_switch.wgsl.expected.glsl b/test/tint/loops/multiple_switch.wgsl.expected.glsl
new file mode 100644
index 0000000..290f1e3
--- /dev/null
+++ b/test/tint/loops/multiple_switch.wgsl.expected.glsl
@@ -0,0 +1,33 @@
+#version 310 es
+
+void tint_symbol() {
+ int i = 0;
+ {
+ for(int i_1 = 0; (i_1 < 2); i_1 = (i_1 + 1)) {
+ switch(i_1) {
+ case 0: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ switch(i_1) {
+ case 0: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ }
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/loops/multiple_switch.wgsl.expected.msl b/test/tint/loops/multiple_switch.wgsl.expected.msl
new file mode 100644
index 0000000..342db50
--- /dev/null
+++ b/test/tint/loops/multiple_switch.wgsl.expected.msl
@@ -0,0 +1,28 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+ int i = 0;
+ for(int i_1 = 0; (i_1 < 2); i_1 = as_type<int>((as_type<uint>(i_1) + as_type<uint>(1)))) {
+ switch(i_1) {
+ case 0: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ switch(i_1) {
+ case 0: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ return;
+}
+
diff --git a/test/tint/loops/multiple_switch.wgsl.expected.spvasm b/test/tint/loops/multiple_switch.wgsl.expected.spvasm
new file mode 100644
index 0000000..edf7850
--- /dev/null
+++ b/test/tint/loops/multiple_switch.wgsl.expected.spvasm
@@ -0,0 +1,64 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 32
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ OpName %i "i"
+ OpName %i_1 "i_1"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %6 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_2 = OpConstant %int 2
+ %bool = OpTypeBool
+ %int_1 = OpConstant %int 1
+ %main = OpFunction %void None %1
+ %4 = OpLabel
+ %i = OpVariable %_ptr_Function_int Function %6
+ %i_1 = OpVariable %_ptr_Function_int Function %6
+ OpStore %i %6
+ OpStore %i_1 %6
+ OpBranch %10
+ %10 = OpLabel
+ OpLoopMerge %11 %12 None
+ OpBranch %13
+ %13 = OpLabel
+ %15 = OpLoad %int %i_1
+ %17 = OpSLessThan %bool %15 %int_2
+ %14 = OpLogicalNot %bool %17
+ OpSelectionMerge %19 None
+ OpBranchConditional %14 %20 %19
+ %20 = OpLabel
+ OpBranch %11
+ %19 = OpLabel
+ %22 = OpLoad %int %i_1
+ OpSelectionMerge %21 None
+ OpSwitch %22 %23 0 %24
+ %24 = OpLabel
+ OpBranch %12
+ %23 = OpLabel
+ OpBranch %21
+ %21 = OpLabel
+ %26 = OpLoad %int %i_1
+ OpSelectionMerge %25 None
+ OpSwitch %26 %27 0 %28
+ %28 = OpLabel
+ OpBranch %12
+ %27 = OpLabel
+ OpBranch %25
+ %25 = OpLabel
+ OpBranch %12
+ %12 = OpLabel
+ %29 = OpLoad %int %i_1
+ %31 = OpIAdd %int %29 %int_1
+ OpStore %i_1 %31
+ OpBranch %10
+ %11 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/loops/multiple_switch.wgsl.expected.wgsl b/test/tint/loops/multiple_switch.wgsl.expected.wgsl
new file mode 100644
index 0000000..1ef17ce
--- /dev/null
+++ b/test/tint/loops/multiple_switch.wgsl.expected.wgsl
@@ -0,0 +1,20 @@
+@compute @workgroup_size(1)
+fn main() {
+ var i = 0;
+ for(var i = 0; (i < 2); i += 1) {
+ switch(i) {
+ case 0: {
+ continue;
+ }
+ default: {
+ }
+ }
+ switch(i) {
+ case 0: {
+ continue;
+ }
+ default: {
+ }
+ }
+ }
+}
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl b/test/tint/loops/nested_loop_loop_switch.wgsl
new file mode 100644
index 0000000..f925971
--- /dev/null
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl
@@ -0,0 +1,14 @@
+@compute @workgroup_size(1)
+fn main() {
+ for (var i = 0; i < 2; i += 2) {
+ for (var j = 0; j < 2; j += 2) {
+ switch(i) {
+ case 0: {
+ continue;
+ }
+ default: {
+ }
+ }
+ }
+ }
+}
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.dxc.hlsl b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..01759c8
--- /dev/null
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.dxc.hlsl
@@ -0,0 +1,26 @@
+[numthreads(1, 1, 1)]
+void main() {
+ {
+ for(int i = 0; (i < 2); i = (i + 2)) {
+ bool tint_continue = false;
+ {
+ for(int j = 0; (j < 2); j = (j + 2)) {
+ tint_continue = false;
+ switch(i) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.fxc.hlsl b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..01759c8
--- /dev/null
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.fxc.hlsl
@@ -0,0 +1,26 @@
+[numthreads(1, 1, 1)]
+void main() {
+ {
+ for(int i = 0; (i < 2); i = (i + 2)) {
+ bool tint_continue = false;
+ {
+ for(int j = 0; (j < 2); j = (j + 2)) {
+ tint_continue = false;
+ switch(i) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.glsl b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.glsl
new file mode 100644
index 0000000..a3ad0a3
--- /dev/null
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.glsl
@@ -0,0 +1,27 @@
+#version 310 es
+
+void tint_symbol() {
+ {
+ for(int i = 0; (i < 2); i = (i + 2)) {
+ {
+ for(int j = 0; (j < 2); j = (j + 2)) {
+ switch(i) {
+ case 0: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ }
+ }
+ }
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
new file mode 100644
index 0000000..a1f0843
--- /dev/null
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
@@ -0,0 +1,20 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+ for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+ for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+ switch(i) {
+ case 0: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ }
+ return;
+}
+
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.spvasm b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.spvasm
new file mode 100644
index 0000000..16f154d
--- /dev/null
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.spvasm
@@ -0,0 +1,75 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 38
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ OpName %i "i"
+ OpName %j "j"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %6 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_2 = OpConstant %int 2
+ %bool = OpTypeBool
+ %main = OpFunction %void None %1
+ %4 = OpLabel
+ %i = OpVariable %_ptr_Function_int Function %6
+ %j = OpVariable %_ptr_Function_int Function %6
+ OpStore %i %6
+ OpBranch %9
+ %9 = OpLabel
+ OpLoopMerge %10 %11 None
+ OpBranch %12
+ %12 = OpLabel
+ %14 = OpLoad %int %i
+ %16 = OpSLessThan %bool %14 %int_2
+ %13 = OpLogicalNot %bool %16
+ OpSelectionMerge %18 None
+ OpBranchConditional %13 %19 %18
+ %19 = OpLabel
+ OpBranch %10
+ %18 = OpLabel
+ OpStore %j %6
+ OpBranch %21
+ %21 = OpLabel
+ OpLoopMerge %22 %23 None
+ OpBranch %24
+ %24 = OpLabel
+ %26 = OpLoad %int %j
+ %27 = OpSLessThan %bool %26 %int_2
+ %25 = OpLogicalNot %bool %27
+ OpSelectionMerge %28 None
+ OpBranchConditional %25 %29 %28
+ %29 = OpLabel
+ OpBranch %22
+ %28 = OpLabel
+ %31 = OpLoad %int %i
+ OpSelectionMerge %30 None
+ OpSwitch %31 %32 0 %33
+ %33 = OpLabel
+ OpBranch %23
+ %32 = OpLabel
+ OpBranch %30
+ %30 = OpLabel
+ OpBranch %23
+ %23 = OpLabel
+ %34 = OpLoad %int %j
+ %35 = OpIAdd %int %34 %int_2
+ OpStore %j %35
+ OpBranch %21
+ %22 = OpLabel
+ OpBranch %11
+ %11 = OpLabel
+ %36 = OpLoad %int %i
+ %37 = OpIAdd %int %36 %int_2
+ OpStore %i %37
+ OpBranch %9
+ %10 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.wgsl b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.wgsl
new file mode 100644
index 0000000..8026320
--- /dev/null
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.wgsl
@@ -0,0 +1,14 @@
+@compute @workgroup_size(1)
+fn main() {
+ for(var i = 0; (i < 2); i += 2) {
+ for(var j = 0; (j < 2); j += 2) {
+ switch(i) {
+ case 0: {
+ continue;
+ }
+ default: {
+ }
+ }
+ }
+ }
+}
diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl b/test/tint/loops/nested_loop_switch_loop_switch.wgsl
new file mode 100644
index 0000000..020d6df
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl
@@ -0,0 +1,21 @@
+@compute @workgroup_size(1)
+fn main() {
+ for (var i = 0; i < 2; i += 2) {
+ switch(i) {
+ case 0: {
+ for (var j = 0; j < 2; j += 2) {
+ switch(j) {
+ case 0: {
+ continue; // j loop
+ }
+ default: {
+ }
+ }
+ }
+ continue; // i loop
+ }
+ default: {
+ }
+ }
+ }
+}
diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.dxc.hlsl b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..305bfbd
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.dxc.hlsl
@@ -0,0 +1,40 @@
+[numthreads(1, 1, 1)]
+void main() {
+ bool tint_continue_1 = false;
+ {
+ for(int i = 0; (i < 2); i = (i + 2)) {
+ tint_continue_1 = false;
+ switch(i) {
+ case 0: {
+ bool tint_continue = false;
+ {
+ for(int j = 0; (j < 2); j = (j + 2)) {
+ tint_continue = false;
+ switch(j) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ }
+ tint_continue_1 = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue_1) {
+ continue;
+ }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.fxc.hlsl b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..305bfbd
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.fxc.hlsl
@@ -0,0 +1,40 @@
+[numthreads(1, 1, 1)]
+void main() {
+ bool tint_continue_1 = false;
+ {
+ for(int i = 0; (i < 2); i = (i + 2)) {
+ tint_continue_1 = false;
+ switch(i) {
+ case 0: {
+ bool tint_continue = false;
+ {
+ for(int j = 0; (j < 2); j = (j + 2)) {
+ tint_continue = false;
+ switch(j) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ }
+ tint_continue_1 = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue_1) {
+ continue;
+ }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.glsl b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.glsl
new file mode 100644
index 0000000..1e45661
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.glsl
@@ -0,0 +1,36 @@
+#version 310 es
+
+void tint_symbol() {
+ {
+ for(int i = 0; (i < 2); i = (i + 2)) {
+ switch(i) {
+ case 0: {
+ {
+ for(int j = 0; (j < 2); j = (j + 2)) {
+ switch(j) {
+ case 0: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ }
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ }
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
new file mode 100644
index 0000000..a5e768d
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
@@ -0,0 +1,29 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+ for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+ switch(i) {
+ case 0: {
+ for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+ switch(j) {
+ case 0: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ return;
+}
+
diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.spvasm b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.spvasm
new file mode 100644
index 0000000..9509f0d
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.spvasm
@@ -0,0 +1,83 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 42
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ OpName %i "i"
+ OpName %j "j"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %6 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_2 = OpConstant %int 2
+ %bool = OpTypeBool
+ %main = OpFunction %void None %1
+ %4 = OpLabel
+ %i = OpVariable %_ptr_Function_int Function %6
+ %j = OpVariable %_ptr_Function_int Function %6
+ OpStore %i %6
+ OpBranch %9
+ %9 = OpLabel
+ OpLoopMerge %10 %11 None
+ OpBranch %12
+ %12 = OpLabel
+ %14 = OpLoad %int %i
+ %16 = OpSLessThan %bool %14 %int_2
+ %13 = OpLogicalNot %bool %16
+ OpSelectionMerge %18 None
+ OpBranchConditional %13 %19 %18
+ %19 = OpLabel
+ OpBranch %10
+ %18 = OpLabel
+ %21 = OpLoad %int %i
+ OpSelectionMerge %20 None
+ OpSwitch %21 %22 0 %23
+ %23 = OpLabel
+ OpStore %j %6
+ OpBranch %25
+ %25 = OpLabel
+ OpLoopMerge %26 %27 None
+ OpBranch %28
+ %28 = OpLabel
+ %30 = OpLoad %int %j
+ %31 = OpSLessThan %bool %30 %int_2
+ %29 = OpLogicalNot %bool %31
+ OpSelectionMerge %32 None
+ OpBranchConditional %29 %33 %32
+ %33 = OpLabel
+ OpBranch %26
+ %32 = OpLabel
+ %35 = OpLoad %int %j
+ OpSelectionMerge %34 None
+ OpSwitch %35 %36 0 %37
+ %37 = OpLabel
+ OpBranch %27
+ %36 = OpLabel
+ OpBranch %34
+ %34 = OpLabel
+ OpBranch %27
+ %27 = OpLabel
+ %38 = OpLoad %int %j
+ %39 = OpIAdd %int %38 %int_2
+ OpStore %j %39
+ OpBranch %25
+ %26 = OpLabel
+ OpBranch %11
+ %22 = OpLabel
+ OpBranch %20
+ %20 = OpLabel
+ OpBranch %11
+ %11 = OpLabel
+ %40 = OpLoad %int %i
+ %41 = OpIAdd %int %40 %int_2
+ OpStore %i %41
+ OpBranch %9
+ %10 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.wgsl b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.wgsl
new file mode 100644
index 0000000..0398609
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.wgsl
@@ -0,0 +1,21 @@
+@compute @workgroup_size(1)
+fn main() {
+ for(var i = 0; (i < 2); i += 2) {
+ switch(i) {
+ case 0: {
+ for(var j = 0; (j < 2); j += 2) {
+ switch(j) {
+ case 0: {
+ continue;
+ }
+ default: {
+ }
+ }
+ }
+ continue;
+ }
+ default: {
+ }
+ }
+ }
+}
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl
new file mode 100644
index 0000000..c152681
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl
@@ -0,0 +1,31 @@
+@compute @workgroup_size(1)
+fn main() {
+ var k = 0;
+ for (var i = 0; i < 2; i += 2) {
+ switch(i) {
+ case 0: {
+ for (var j = 0; j < 2; j += 2) {
+ switch(j) {
+ case 0: {
+ continue; // j loop
+ }
+ case 1: {
+ switch (k) {
+ case 0: {
+ continue; // j loop
+ }
+ default: {
+ }
+ }
+ }
+ default: {
+ }
+ }
+ }
+ continue; // i loop
+ }
+ default: {
+ }
+ }
+ }
+}
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.dxc.hlsl b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..7d646a9
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.dxc.hlsl
@@ -0,0 +1,56 @@
+[numthreads(1, 1, 1)]
+void main() {
+ int k = 0;
+ bool tint_continue_1 = false;
+ {
+ for(int i = 0; (i < 2); i = (i + 2)) {
+ tint_continue_1 = false;
+ switch(i) {
+ case 0: {
+ bool tint_continue = false;
+ {
+ for(int j = 0; (j < 2); j = (j + 2)) {
+ tint_continue = false;
+ switch(j) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ case 1: {
+ switch(k) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ break;
+ }
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ }
+ tint_continue_1 = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue_1) {
+ continue;
+ }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.fxc.hlsl b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..7d646a9
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.fxc.hlsl
@@ -0,0 +1,56 @@
+[numthreads(1, 1, 1)]
+void main() {
+ int k = 0;
+ bool tint_continue_1 = false;
+ {
+ for(int i = 0; (i < 2); i = (i + 2)) {
+ tint_continue_1 = false;
+ switch(i) {
+ case 0: {
+ bool tint_continue = false;
+ {
+ for(int j = 0; (j < 2); j = (j + 2)) {
+ tint_continue = false;
+ switch(j) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ case 1: {
+ switch(k) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ break;
+ }
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ }
+ tint_continue_1 = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue_1) {
+ continue;
+ }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.glsl b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.glsl
new file mode 100644
index 0000000..3c59535
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.glsl
@@ -0,0 +1,49 @@
+#version 310 es
+
+void tint_symbol() {
+ int k = 0;
+ {
+ for(int i = 0; (i < 2); i = (i + 2)) {
+ switch(i) {
+ case 0: {
+ {
+ for(int j = 0; (j < 2); j = (j + 2)) {
+ switch(j) {
+ case 0: {
+ continue;
+ break;
+ }
+ case 1: {
+ switch(k) {
+ case 0: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ }
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ }
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl
new file mode 100644
index 0000000..b3a1fd6
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl
@@ -0,0 +1,42 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+ int k = 0;
+ for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+ switch(i) {
+ case 0: {
+ for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+ switch(j) {
+ case 0: {
+ continue;
+ break;
+ }
+ case 1: {
+ switch(k) {
+ case 0: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ return;
+}
+
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.spvasm b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.spvasm
new file mode 100644
index 0000000..b32b253
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.spvasm
@@ -0,0 +1,96 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ OpName %k "k"
+ OpName %i "i"
+ OpName %j "j"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %6 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_2 = OpConstant %int 2
+ %bool = OpTypeBool
+ %main = OpFunction %void None %1
+ %4 = OpLabel
+ %k = OpVariable %_ptr_Function_int Function %6
+ %i = OpVariable %_ptr_Function_int Function %6
+ %j = OpVariable %_ptr_Function_int Function %6
+ OpStore %k %6
+ OpStore %i %6
+ OpBranch %10
+ %10 = OpLabel
+ OpLoopMerge %11 %12 None
+ OpBranch %13
+ %13 = OpLabel
+ %15 = OpLoad %int %i
+ %17 = OpSLessThan %bool %15 %int_2
+ %14 = OpLogicalNot %bool %17
+ OpSelectionMerge %19 None
+ OpBranchConditional %14 %20 %19
+ %20 = OpLabel
+ OpBranch %11
+ %19 = OpLabel
+ %22 = OpLoad %int %i
+ OpSelectionMerge %21 None
+ OpSwitch %22 %23 0 %24
+ %24 = OpLabel
+ OpStore %j %6
+ OpBranch %26
+ %26 = OpLabel
+ OpLoopMerge %27 %28 None
+ OpBranch %29
+ %29 = OpLabel
+ %31 = OpLoad %int %j
+ %32 = OpSLessThan %bool %31 %int_2
+ %30 = OpLogicalNot %bool %32
+ OpSelectionMerge %33 None
+ OpBranchConditional %30 %34 %33
+ %34 = OpLabel
+ OpBranch %27
+ %33 = OpLabel
+ %36 = OpLoad %int %j
+ OpSelectionMerge %35 None
+ OpSwitch %36 %37 0 %38 1 %39
+ %38 = OpLabel
+ OpBranch %28
+ %39 = OpLabel
+ %41 = OpLoad %int %k
+ OpSelectionMerge %40 None
+ OpSwitch %41 %42 0 %43
+ %43 = OpLabel
+ OpBranch %28
+ %42 = OpLabel
+ OpBranch %40
+ %40 = OpLabel
+ OpBranch %35
+ %37 = OpLabel
+ OpBranch %35
+ %35 = OpLabel
+ OpBranch %28
+ %28 = OpLabel
+ %44 = OpLoad %int %j
+ %45 = OpIAdd %int %44 %int_2
+ OpStore %j %45
+ OpBranch %26
+ %27 = OpLabel
+ OpBranch %12
+ %23 = OpLabel
+ OpBranch %21
+ %21 = OpLabel
+ OpBranch %12
+ %12 = OpLabel
+ %46 = OpLoad %int %i
+ %47 = OpIAdd %int %46 %int_2
+ OpStore %i %47
+ OpBranch %10
+ %11 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.wgsl b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.wgsl
new file mode 100644
index 0000000..7c7cdb9
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.wgsl
@@ -0,0 +1,31 @@
+@compute @workgroup_size(1)
+fn main() {
+ var k = 0;
+ for(var i = 0; (i < 2); i += 2) {
+ switch(i) {
+ case 0: {
+ for(var j = 0; (j < 2); j += 2) {
+ switch(j) {
+ case 0: {
+ continue;
+ }
+ case 1: {
+ switch(k) {
+ case 0: {
+ continue;
+ }
+ default: {
+ }
+ }
+ }
+ default: {
+ }
+ }
+ }
+ continue;
+ }
+ default: {
+ }
+ }
+ }
+}
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl b/test/tint/loops/nested_loop_switch_switch.wgsl
new file mode 100644
index 0000000..9949f96
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl
@@ -0,0 +1,19 @@
+@compute @workgroup_size(1)
+fn main() {
+ var j = 0;
+ for (var i = 0; i < 2; i += 2) {
+ switch(i) {
+ case 0: {
+ switch(j) {
+ case 0: {
+ continue;
+ }
+ default: {
+ }
+ }
+ }
+ default: {
+ }
+ }
+ }
+}
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.dxc.hlsl b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..160d4d6
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.dxc.hlsl
@@ -0,0 +1,34 @@
+[numthreads(1, 1, 1)]
+void main() {
+ int j = 0;
+ bool tint_continue = false;
+ {
+ for(int i = 0; (i < 2); i = (i + 2)) {
+ tint_continue = false;
+ switch(i) {
+ case 0: {
+ switch(j) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ break;
+ }
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.fxc.hlsl b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..160d4d6
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.fxc.hlsl
@@ -0,0 +1,34 @@
+[numthreads(1, 1, 1)]
+void main() {
+ int j = 0;
+ bool tint_continue = false;
+ {
+ for(int i = 0; (i < 2); i = (i + 2)) {
+ tint_continue = false;
+ switch(i) {
+ case 0: {
+ switch(j) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ break;
+ }
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.glsl b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.glsl
new file mode 100644
index 0000000..ec32b25
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.glsl
@@ -0,0 +1,32 @@
+#version 310 es
+
+void tint_symbol() {
+ int j = 0;
+ {
+ for(int i = 0; (i < 2); i = (i + 2)) {
+ switch(i) {
+ case 0: {
+ switch(j) {
+ case 0: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ }
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
new file mode 100644
index 0000000..55f4436
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+ int j = 0;
+ for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+ switch(i) {
+ case 0: {
+ switch(j) {
+ case 0: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ return;
+}
+
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.spvasm b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.spvasm
new file mode 100644
index 0000000..f6add1f
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.spvasm
@@ -0,0 +1,63 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ OpName %j "j"
+ OpName %i "i"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %6 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_2 = OpConstant %int 2
+ %bool = OpTypeBool
+ %main = OpFunction %void None %1
+ %4 = OpLabel
+ %j = OpVariable %_ptr_Function_int Function %6
+ %i = OpVariable %_ptr_Function_int Function %6
+ OpStore %j %6
+ OpStore %i %6
+ OpBranch %10
+ %10 = OpLabel
+ OpLoopMerge %11 %12 None
+ OpBranch %13
+ %13 = OpLabel
+ %15 = OpLoad %int %i
+ %17 = OpSLessThan %bool %15 %int_2
+ %14 = OpLogicalNot %bool %17
+ OpSelectionMerge %19 None
+ OpBranchConditional %14 %20 %19
+ %20 = OpLabel
+ OpBranch %11
+ %19 = OpLabel
+ %22 = OpLoad %int %i
+ OpSelectionMerge %21 None
+ OpSwitch %22 %23 0 %24
+ %24 = OpLabel
+ %26 = OpLoad %int %j
+ OpSelectionMerge %25 None
+ OpSwitch %26 %27 0 %28
+ %28 = OpLabel
+ OpBranch %12
+ %27 = OpLabel
+ OpBranch %25
+ %25 = OpLabel
+ OpBranch %21
+ %23 = OpLabel
+ OpBranch %21
+ %21 = OpLabel
+ OpBranch %12
+ %12 = OpLabel
+ %29 = OpLoad %int %i
+ %30 = OpIAdd %int %29 %int_2
+ OpStore %i %30
+ OpBranch %10
+ %11 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.wgsl b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.wgsl
new file mode 100644
index 0000000..4d8f604
--- /dev/null
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+@compute @workgroup_size(1)
+fn main() {
+ var j = 0;
+ for(var i = 0; (i < 2); i += 2) {
+ switch(i) {
+ case 0: {
+ switch(j) {
+ case 0: {
+ continue;
+ }
+ default: {
+ }
+ }
+ }
+ default: {
+ }
+ }
+ }
+}
diff --git a/test/tint/loops/single_continue.wgsl b/test/tint/loops/single_continue.wgsl
new file mode 100644
index 0000000..d59d2fa
--- /dev/null
+++ b/test/tint/loops/single_continue.wgsl
@@ -0,0 +1,12 @@
+@compute @workgroup_size(1)
+fn main() {
+ for (var i = 0; i < 2; i += 1) {
+ switch(i) {
+ case 0: {
+ continue;
+ }
+ default: {
+ }
+ }
+ }
+}
diff --git a/test/tint/loops/single_continue.wgsl.expected.dxc.hlsl b/test/tint/loops/single_continue.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..5955f1e
--- /dev/null
+++ b/test/tint/loops/single_continue.wgsl.expected.dxc.hlsl
@@ -0,0 +1,22 @@
+[numthreads(1, 1, 1)]
+void main() {
+ bool tint_continue = false;
+ {
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ tint_continue = false;
+ switch(i) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/single_continue.wgsl.expected.fxc.hlsl b/test/tint/loops/single_continue.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..5955f1e
--- /dev/null
+++ b/test/tint/loops/single_continue.wgsl.expected.fxc.hlsl
@@ -0,0 +1,22 @@
+[numthreads(1, 1, 1)]
+void main() {
+ bool tint_continue = false;
+ {
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ tint_continue = false;
+ switch(i) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ continue;
+ }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/single_continue.wgsl.expected.glsl b/test/tint/loops/single_continue.wgsl.expected.glsl
new file mode 100644
index 0000000..79afd54
--- /dev/null
+++ b/test/tint/loops/single_continue.wgsl.expected.glsl
@@ -0,0 +1,23 @@
+#version 310 es
+
+void tint_symbol() {
+ {
+ for(int i = 0; (i < 2); i = (i + 1)) {
+ switch(i) {
+ case 0: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ }
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol();
+ return;
+}
diff --git a/test/tint/loops/single_continue.wgsl.expected.msl b/test/tint/loops/single_continue.wgsl.expected.msl
new file mode 100644
index 0000000..3ce3ba5
--- /dev/null
+++ b/test/tint/loops/single_continue.wgsl.expected.msl
@@ -0,0 +1,18 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+ for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ switch(i) {
+ case 0: {
+ continue;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ }
+ return;
+}
+
diff --git a/test/tint/loops/single_continue.wgsl.expected.spvasm b/test/tint/loops/single_continue.wgsl.expected.spvasm
new file mode 100644
index 0000000..7b4bed2
--- /dev/null
+++ b/test/tint/loops/single_continue.wgsl.expected.spvasm
@@ -0,0 +1,53 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 27
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ OpName %i "i"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %6 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_2 = OpConstant %int 2
+ %bool = OpTypeBool
+ %int_1 = OpConstant %int 1
+ %main = OpFunction %void None %1
+ %4 = OpLabel
+ %i = OpVariable %_ptr_Function_int Function %6
+ OpStore %i %6
+ OpBranch %9
+ %9 = OpLabel
+ OpLoopMerge %10 %11 None
+ OpBranch %12
+ %12 = OpLabel
+ %14 = OpLoad %int %i
+ %16 = OpSLessThan %bool %14 %int_2
+ %13 = OpLogicalNot %bool %16
+ OpSelectionMerge %18 None
+ OpBranchConditional %13 %19 %18
+ %19 = OpLabel
+ OpBranch %10
+ %18 = OpLabel
+ %21 = OpLoad %int %i
+ OpSelectionMerge %20 None
+ OpSwitch %21 %22 0 %23
+ %23 = OpLabel
+ OpBranch %11
+ %22 = OpLabel
+ OpBranch %20
+ %20 = OpLabel
+ OpBranch %11
+ %11 = OpLabel
+ %24 = OpLoad %int %i
+ %26 = OpIAdd %int %24 %int_1
+ OpStore %i %26
+ OpBranch %9
+ %10 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/loops/single_continue.wgsl.expected.wgsl b/test/tint/loops/single_continue.wgsl.expected.wgsl
new file mode 100644
index 0000000..6554849
--- /dev/null
+++ b/test/tint/loops/single_continue.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+@compute @workgroup_size(1)
+fn main() {
+ for(var i = 0; (i < 2); i += 1) {
+ switch(i) {
+ case 0: {
+ continue;
+ }
+ default: {
+ }
+ }
+ }
+}
