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/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: {
+ }
+ }
+ }
+}
