[msl][glsl] Fix incorrect code gen for continuing inside switch.
In the GLSL and MSL backends when generating a continue inside a switch
block, the continuing statement would sink up into the switch label.
This means that a `break` for the continuing would break the switch
instead of the loop.
This CL moves the RemoveContinuingInSwitch transform from the HLSL
backend up to being a generic AST transform and calls it from the MSL
and GLSL backends to write the AST to no longer continue from switches.
Bug: tint:2039
Change-Id: Ica4a8d2de6bf97a179dac9e9216d8ddce99b5976
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/175820
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/test/tint/bug/tint/2039.wgsl b/test/tint/bug/tint/2039.wgsl
new file mode 100644
index 0000000..419edb0
--- /dev/null
+++ b/test/tint/bug/tint/2039.wgsl
@@ -0,0 +1,17 @@
+
+@compute @workgroup_size(1,1,1)
+fn main() {
+ var out = 0u;
+ loop {
+ switch 2 {
+ case 1 {
+ continue;
+ }
+ default {}
+ }
+ out += 1u;
+ continuing {
+ break if true;
+ }
+ }
+}
diff --git a/test/tint/bug/tint/2039.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/2039.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..2707ef2
--- /dev/null
+++ b/test/tint/bug/tint/2039.wgsl.expected.dxc.hlsl
@@ -0,0 +1,28 @@
+[numthreads(1, 1, 1)]
+void main() {
+ uint tint_symbol = 0u;
+ bool tint_continue = false;
+ while (true) {
+ tint_continue = false;
+ switch(2) {
+ case 1: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ {
+ if (true) { break; }
+ }
+ continue;
+ }
+ tint_symbol = (tint_symbol + 1u);
+ {
+ if (true) { break; }
+ }
+ }
+ return;
+}
diff --git a/test/tint/bug/tint/2039.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/2039.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..2707ef2
--- /dev/null
+++ b/test/tint/bug/tint/2039.wgsl.expected.fxc.hlsl
@@ -0,0 +1,28 @@
+[numthreads(1, 1, 1)]
+void main() {
+ uint tint_symbol = 0u;
+ bool tint_continue = false;
+ while (true) {
+ tint_continue = false;
+ switch(2) {
+ case 1: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ {
+ if (true) { break; }
+ }
+ continue;
+ }
+ tint_symbol = (tint_symbol + 1u);
+ {
+ if (true) { break; }
+ }
+ }
+ return;
+}
diff --git a/test/tint/bug/tint/2039.wgsl.expected.glsl b/test/tint/bug/tint/2039.wgsl.expected.glsl
new file mode 100644
index 0000000..41aa478
--- /dev/null
+++ b/test/tint/bug/tint/2039.wgsl.expected.glsl
@@ -0,0 +1,34 @@
+#version 310 es
+
+void tint_symbol() {
+ uint tint_symbol_1 = 0u;
+ bool tint_continue = false;
+ while (true) {
+ tint_continue = false;
+ switch(2) {
+ case 1: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ {
+ if (true) { break; }
+ }
+ continue;
+ }
+ tint_symbol_1 = (tint_symbol_1 + 1u);
+ {
+ if (true) { 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/bug/tint/2039.wgsl.expected.msl b/test/tint/bug/tint/2039.wgsl.expected.msl
new file mode 100644
index 0000000..c756acb
--- /dev/null
+++ b/test/tint/bug/tint/2039.wgsl.expected.msl
@@ -0,0 +1,36 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
+kernel void tint_symbol() {
+ uint out = 0u;
+ bool tint_continue = false;
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ tint_continue = false;
+ switch(2) {
+ case 1: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ {
+ if (true) { break; }
+ }
+ continue;
+ }
+ out = (out + 1u);
+ {
+ if (true) { break; }
+ }
+ }
+ return;
+}
+
diff --git a/test/tint/bug/tint/2039.wgsl.expected.spvasm b/test/tint/bug/tint/2039.wgsl.expected.spvasm
new file mode 100644
index 0000000..ce5ae5e
--- /dev/null
+++ b/test/tint/bug/tint/2039.wgsl.expected.spvasm
@@ -0,0 +1,46 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 23
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ OpName %out "out"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %6 = OpConstantNull %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %int = OpTypeInt 32 1
+ %int_2 = OpConstant %int 2
+ %uint_1 = OpConstant %uint 1
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %main = OpFunction %void None %1
+ %4 = OpLabel
+ %out = OpVariable %_ptr_Function_uint Function %6
+ OpStore %out %6
+ OpBranch %9
+ %9 = OpLabel
+ OpLoopMerge %10 %11 None
+ OpBranch %12
+ %12 = OpLabel
+ OpSelectionMerge %13 None
+ OpSwitch %int_2 %16 1 %17
+ %17 = OpLabel
+ OpBranch %11
+ %16 = OpLabel
+ OpBranch %13
+ %13 = OpLabel
+ %18 = OpLoad %uint %out
+ %20 = OpIAdd %uint %18 %uint_1
+ OpStore %out %20
+ OpBranch %11
+ %11 = OpLabel
+ OpBranchConditional %true %10 %9
+ %10 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/2039.wgsl.expected.wgsl b/test/tint/bug/tint/2039.wgsl.expected.wgsl
new file mode 100644
index 0000000..0bf5ccd
--- /dev/null
+++ b/test/tint/bug/tint/2039.wgsl.expected.wgsl
@@ -0,0 +1,18 @@
+@compute @workgroup_size(1, 1, 1)
+fn main() {
+ var out = 0u;
+ loop {
+ switch(2) {
+ case 1: {
+ continue;
+ }
+ default: {
+ }
+ }
+ out += 1u;
+
+ continuing {
+ break if true;
+ }
+ }
+}
diff --git a/test/tint/loops/continue_in_switch.wgsl.expected.glsl b/test/tint/loops/continue_in_switch.wgsl.expected.glsl
index 9678a12..291f4cd 100644
--- a/test/tint/loops/continue_in_switch.wgsl.expected.glsl
+++ b/test/tint/loops/continue_in_switch.wgsl.expected.glsl
@@ -1,17 +1,22 @@
#version 310 es
void f() {
+ bool tint_continue = false;
{
for(int i = 0; (i < 4); i = (i + 1)) {
+ tint_continue = false;
switch(i) {
case 0: {
- continue;
+ tint_continue = true;
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ continue;
+ }
}
}
}
diff --git a/test/tint/loops/continue_in_switch.wgsl.expected.msl b/test/tint/loops/continue_in_switch.wgsl.expected.msl
index 890daca..e7349b0c 100644
--- a/test/tint/loops/continue_in_switch.wgsl.expected.msl
+++ b/test/tint/loops/continue_in_switch.wgsl.expected.msl
@@ -7,16 +7,21 @@
if (VOLATILE_NAME)
kernel void f() {
+ bool tint_continue = false;
TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ tint_continue = false;
switch(i) {
case 0: {
- continue;
+ 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
index 36c86d8..bbdf58c 100644
--- a/test/tint/loops/multiple_continues.wgsl.expected.glsl
+++ b/test/tint/loops/multiple_continues.wgsl.expected.glsl
@@ -1,25 +1,30 @@
#version 310 es
void tint_symbol() {
+ bool tint_continue = false;
{
for(int i = 0; (i < 2); i = (i + 1)) {
+ tint_continue = false;
switch(i) {
case 0: {
- continue;
+ tint_continue = true;
break;
}
case 1: {
- continue;
+ tint_continue = true;
break;
}
case 2: {
- continue;
+ tint_continue = true;
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ continue;
+ }
}
}
}
diff --git a/test/tint/loops/multiple_continues.wgsl.expected.msl b/test/tint/loops/multiple_continues.wgsl.expected.msl
index f8aaba6..76e84d7 100644
--- a/test/tint/loops/multiple_continues.wgsl.expected.msl
+++ b/test/tint/loops/multiple_continues.wgsl.expected.msl
@@ -7,24 +7,29 @@
if (VOLATILE_NAME)
kernel void tint_symbol() {
+ bool tint_continue = false;
TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ tint_continue = false;
switch(i) {
case 0: {
- continue;
+ tint_continue = true;
break;
}
case 1: {
- continue;
+ tint_continue = true;
break;
}
case 2: {
- continue;
+ 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
index 290f1e3..82abd76 100644
--- a/test/tint/loops/multiple_switch.wgsl.expected.glsl
+++ b/test/tint/loops/multiple_switch.wgsl.expected.glsl
@@ -2,26 +2,34 @@
void tint_symbol() {
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: {
- continue;
+ tint_continue = true;
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ continue;
+ }
switch(i_1) {
case 0: {
- continue;
+ tint_continue = true;
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ continue;
+ }
}
}
}
diff --git a/test/tint/loops/multiple_switch.wgsl.expected.msl b/test/tint/loops/multiple_switch.wgsl.expected.msl
index 0bad249..7728952 100644
--- a/test/tint/loops/multiple_switch.wgsl.expected.msl
+++ b/test/tint/loops/multiple_switch.wgsl.expected.msl
@@ -8,25 +8,33 @@
kernel void tint_symbol() {
int i = 0;
+ bool tint_continue = false;
TINT_ISOLATE_UB(tint_volatile_true) for(int i_1 = 0; (i_1 < 2); i_1 = as_type<int>((as_type<uint>(i_1) + as_type<uint>(1)))) {
+ tint_continue = false;
switch(i_1) {
case 0: {
- continue;
+ tint_continue = true;
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ continue;
+ }
switch(i_1) {
case 0: {
- continue;
+ 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
index a3ad0a3..63e7289 100644
--- a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.glsl
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.glsl
@@ -3,17 +3,22 @@
void tint_symbol() {
{
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: {
- continue;
+ tint_continue = true;
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ continue;
+ }
}
}
}
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
index 8940027..eb32ea2 100644
--- a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
@@ -8,16 +8,21 @@
kernel void tint_symbol() {
TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+ bool tint_continue = false;
TINT_ISOLATE_UB(tint_volatile_true_1) for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+ tint_continue = false;
switch(i) {
case 0: {
- continue;
+ tint_continue = true;
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ 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
index 1e45661..06552d5 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.glsl
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.glsl
@@ -1,30 +1,40 @@
#version 310 es
void tint_symbol() {
+ 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: {
- continue;
+ tint_continue = true;
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ continue;
+ }
}
}
- continue;
+ tint_continue_1 = true;
break;
}
default: {
break;
}
}
+ if (tint_continue_1) {
+ continue;
+ }
}
}
}
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
index 2662928..9ac0463 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
@@ -7,27 +7,37 @@
if (VOLATILE_NAME)
kernel void tint_symbol() {
+ bool tint_continue_1 = false;
TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+ tint_continue_1 = false;
switch(i) {
case 0: {
+ bool tint_continue = false;
TINT_ISOLATE_UB(tint_volatile_true_1) for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+ tint_continue = false;
switch(j) {
case 0: {
- continue;
+ tint_continue = true;
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ 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
index 3c59535..2b1cc90 100644
--- 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
@@ -2,42 +2,55 @@
void tint_symbol() {
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: {
- continue;
+ tint_continue = true;
break;
}
case 1: {
switch(k) {
case 0: {
- continue;
+ tint_continue = true;
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ break;
+ }
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ continue;
+ }
}
}
- continue;
+ tint_continue_1 = true;
break;
}
default: {
break;
}
}
+ if (tint_continue_1) {
+ continue;
+ }
}
}
}
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
index 47f158b..c3b1e93 100644
--- 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
@@ -8,39 +8,52 @@
kernel void tint_symbol() {
int k = 0;
+ bool tint_continue_1 = false;
TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+ tint_continue_1 = false;
switch(i) {
case 0: {
+ bool tint_continue = false;
TINT_ISOLATE_UB(tint_volatile_true_1) for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+ tint_continue = false;
switch(j) {
case 0: {
- continue;
+ tint_continue = true;
break;
}
case 1: {
switch(k) {
case 0: {
- continue;
+ tint_continue = true;
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ break;
+ }
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ continue;
+ }
}
- continue;
+ tint_continue_1 = true;
break;
}
default: {
break;
}
}
+ if (tint_continue_1) {
+ 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
index ec32b25..f7b3dd6 100644
--- a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.glsl
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.glsl
@@ -2,25 +2,33 @@
void tint_symbol() {
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: {
- continue;
+ tint_continue = true;
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ break;
+ }
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ continue;
+ }
}
}
}
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
index 0f78ee3..e122be4 100644
--- a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
@@ -8,24 +8,32 @@
kernel void tint_symbol() {
int j = 0;
+ bool tint_continue = false;
TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+ tint_continue = false;
switch(i) {
case 0: {
switch(j) {
case 0: {
- continue;
+ tint_continue = true;
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ break;
+ }
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
index 79afd54..58309a6 100644
--- a/test/tint/loops/single_continue.wgsl.expected.glsl
+++ b/test/tint/loops/single_continue.wgsl.expected.glsl
@@ -1,17 +1,22 @@
#version 310 es
void tint_symbol() {
+ bool tint_continue = false;
{
for(int i = 0; (i < 2); i = (i + 1)) {
+ tint_continue = false;
switch(i) {
case 0: {
- continue;
+ tint_continue = true;
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ continue;
+ }
}
}
}
diff --git a/test/tint/loops/single_continue.wgsl.expected.msl b/test/tint/loops/single_continue.wgsl.expected.msl
index e223c34..113c92c 100644
--- a/test/tint/loops/single_continue.wgsl.expected.msl
+++ b/test/tint/loops/single_continue.wgsl.expected.msl
@@ -7,16 +7,21 @@
if (VOLATILE_NAME)
kernel void tint_symbol() {
+ bool tint_continue = false;
TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+ tint_continue = false;
switch(i) {
case 0: {
- continue;
+ tint_continue = true;
break;
}
default: {
break;
}
}
+ if (tint_continue) {
+ continue;
+ }
}
return;
}
