[tint][ir] Skip unreachable continuing blocks
When converting AST to IR, skip over unreachable continuing blocks as
they may contain references to unreachable declarations that were also
skipped.
Note that this issue is also present in the legacy SPIR-V writer, and
is not fixed in this CL, hence the SKIP expectations for that path.
Fixed: tint:2201, tint:2202
Change-Id: I2689d412e340508b005a33e251b669948d986cd5
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/179240
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
index c79c6c8..b838c02 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
@@ -592,11 +592,14 @@
}
}
- {
+ // Emit a continuing block if it is reachable.
+ if (body_behaviors.Contains(sem::Behavior::kNext) ||
+ body_behaviors.Contains(sem::Behavior::kContinue)) {
TINT_SCOPED_ASSIGNMENT(current_block_, loop_inst->Continuing());
if (stmt->continuing) {
EmitBlock(stmt->continuing);
}
+
// Branch back to the start block if the continue target didn't terminate already
if (NeedTerminator()) {
SetTerminator(builder_.NextIteration(loop_inst));
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
index e790c55..8ae7aa1 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
@@ -267,13 +267,10 @@
%b1 = block {
if true [t: %b2] { # if_1
%b2 = block { # true
- loop [b: %b3, c: %b4] { # loop_1
+ loop [b: %b3] { # loop_1
%b3 = block { # body
exit_loop # loop_1
}
- %b4 = block { # continuing
- next_iteration %b3
- }
}
exit_if # if_1
}
@@ -296,19 +293,16 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length());
EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2, c: %b3] { # loop_1
+ loop [b: %b2] { # loop_1
%b2 = block { # body
exit_loop # loop_1
}
- %b3 = block { # continuing
- next_iteration %b2
- }
}
ret
}
@@ -465,19 +459,16 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length());
EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2, c: %b3] { # loop_1
+ loop [b: %b2] { # loop_1
%b2 = block { # body
ret
}
- %b3 = block { # continuing
- next_iteration %b2
- }
}
unreachable
}
@@ -506,22 +497,19 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length());
EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2, c: %b3] { # loop_1
+ loop [b: %b2] { # loop_1
%b2 = block { # body
ret
}
- %b3 = block { # continuing
- break_if true %b2
- }
}
- if true [t: %b4] { # if_1
- %b4 = block { # true
+ if true [t: %b3] { # if_1
+ %b3 = block { # true
ret
}
}
@@ -544,27 +532,24 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length());
EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2, c: %b3] { # loop_1
+ loop [b: %b2] { # loop_1
%b2 = block { # body
- if true [t: %b4, f: %b5] { # if_1
- %b4 = block { # true
+ if true [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
exit_loop # loop_1
}
- %b5 = block { # false
+ %b4 = block { # false
exit_loop # loop_1
}
}
unreachable
}
- %b3 = block { # continuing
- next_iteration %b2
- }
}
ret
}
@@ -609,27 +594,24 @@
continue %b5
}
%b5 = block { # continuing
- loop [b: %b8, c: %b9] { # loop_3
+ loop [b: %b8] { # loop_3
%b8 = block { # body
exit_loop # loop_3
}
- %b9 = block { # continuing
- next_iteration %b8
- }
}
- loop [b: %b10, c: %b11] { # loop_4
- %b10 = block { # body
- continue %b11
+ loop [b: %b9, c: %b10] { # loop_4
+ %b9 = block { # body
+ continue %b10
}
- %b11 = block { # continuing
- break_if true %b10
+ %b10 = block { # continuing
+ break_if true %b9
}
}
next_iteration %b4
}
}
- if true [t: %b12] { # if_3
- %b12 = block { # true
+ if true [t: %b11] { # if_3
+ %b11 = block { # true
exit_loop # loop_1
}
}
diff --git a/test/tint/bug/tint/2201.wgsl b/test/tint/bug/tint/2201.wgsl
new file mode 100644
index 0000000..701a748
--- /dev/null
+++ b/test/tint/bug/tint/2201.wgsl
@@ -0,0 +1,14 @@
+@compute @workgroup_size(1)
+fn main() {
+ loop {
+ if true {
+ break;
+ } else {
+ break;
+ }
+ let _e16_ = vec2(false, false);
+ continuing {
+ break if _e16_.x;
+ }
+ }
+}
diff --git a/test/tint/bug/tint/2201.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/2201.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..232a593
--- /dev/null
+++ b/test/tint/bug/tint/2201.wgsl.expected.dxc.hlsl
@@ -0,0 +1,19 @@
+<dawn>/test/tint/bug/tint/2201.wgsl:9:9 warning: code is unreachable
+ let _e16_ = vec2(false, false);
+ ^^^^^^^^^
+
+[numthreads(1, 1, 1)]
+void main() {
+ while (true) {
+ if (true) {
+ break;
+ } else {
+ break;
+ }
+ bool2 _e16_ = (false).xx;
+ {
+ if (_e16_.x) { break; }
+ }
+ }
+ return;
+}
diff --git a/test/tint/bug/tint/2201.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/2201.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..232a593
--- /dev/null
+++ b/test/tint/bug/tint/2201.wgsl.expected.fxc.hlsl
@@ -0,0 +1,19 @@
+<dawn>/test/tint/bug/tint/2201.wgsl:9:9 warning: code is unreachable
+ let _e16_ = vec2(false, false);
+ ^^^^^^^^^
+
+[numthreads(1, 1, 1)]
+void main() {
+ while (true) {
+ if (true) {
+ break;
+ } else {
+ break;
+ }
+ bool2 _e16_ = (false).xx;
+ {
+ if (_e16_.x) { break; }
+ }
+ }
+ return;
+}
diff --git a/test/tint/bug/tint/2201.wgsl.expected.glsl b/test/tint/bug/tint/2201.wgsl.expected.glsl
new file mode 100644
index 0000000..86604a3
--- /dev/null
+++ b/test/tint/bug/tint/2201.wgsl.expected.glsl
@@ -0,0 +1,25 @@
+<dawn>/test/tint/bug/tint/2201.wgsl:9:9 warning: code is unreachable
+ let _e16_ = vec2(false, false);
+ ^^^^^^^^^
+
+#version 310 es
+
+void tint_symbol() {
+ while (true) {
+ if (true) {
+ break;
+ } else {
+ break;
+ }
+ bvec2 _e16_ = bvec2(false);
+ {
+ if (_e16_.x) { 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/2201.wgsl.expected.ir.spvasm b/test/tint/bug/tint/2201.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..49413ac
--- /dev/null
+++ b/test/tint/bug/tint/2201.wgsl.expected.ir.spvasm
@@ -0,0 +1,38 @@
+<dawn>/test/tint/bug/tint/2201.wgsl:9:9 warning: code is unreachable
+ let _e16_ = vec2(false, false);
+ ^^^^^^^^^
+
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 14
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %main = OpFunction %void None %3
+ %4 = OpLabel
+ OpBranch %7
+ %7 = OpLabel
+ OpLoopMerge %8 %6 None
+ OpBranch %5
+ %5 = OpLabel
+ OpSelectionMerge %9 None
+ OpBranchConditional %true %10 %11
+ %10 = OpLabel
+ OpBranch %8
+ %11 = OpLabel
+ OpBranch %8
+ %9 = OpLabel
+ OpUnreachable
+ %6 = OpLabel
+ OpBranch %7
+ %8 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/2201.wgsl.expected.msl b/test/tint/bug/tint/2201.wgsl.expected.msl
new file mode 100644
index 0000000..5852fb6
--- /dev/null
+++ b/test/tint/bug/tint/2201.wgsl.expected.msl
@@ -0,0 +1,27 @@
+<dawn>/test/tint/bug/tint/2201.wgsl:9:9 warning: code is unreachable
+ let _e16_ = vec2(false, false);
+ ^^^^^^^^^
+
+#include <metal_stdlib>
+
+using namespace metal;
+
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
+kernel void tint_symbol() {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ if (true) {
+ break;
+ } else {
+ break;
+ }
+ bool2 const _e16_ = bool2(false);
+ {
+ if (_e16_[0]) { break; }
+ }
+ }
+ return;
+}
+
diff --git a/test/tint/bug/tint/2201.wgsl.expected.spvasm b/test/tint/bug/tint/2201.wgsl.expected.spvasm
new file mode 100644
index 0000000..2fed81e
--- /dev/null
+++ b/test/tint/bug/tint/2201.wgsl.expected.spvasm
@@ -0,0 +1,27 @@
+SKIP: FAILED
+
+<dawn>/test/tint/bug/tint/2201.wgsl:9:9 warning: code is unreachable
+ let _e16_ = vec2(false, false);
+ ^^^^^^^^^
+
+
+@compute @workgroup_size(1)
+fn main() {
+ loop {
+ if (true) {
+ break;
+ } else {
+ break;
+ }
+ let _e16_ = vec2(false, false);
+
+ continuing {
+ break if _e16_.x;
+ }
+ }
+}
+
+Failed to generate: <dawn>/test/tint/bug/tint/2201.wgsl:11:22 error: unresolved value '_e16_'
+ break if _e16_.x;
+ ^^^^^
+
diff --git a/test/tint/bug/tint/2201.wgsl.expected.wgsl b/test/tint/bug/tint/2201.wgsl.expected.wgsl
new file mode 100644
index 0000000..81af241
--- /dev/null
+++ b/test/tint/bug/tint/2201.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+<dawn>/test/tint/bug/tint/2201.wgsl:9:9 warning: code is unreachable
+ let _e16_ = vec2(false, false);
+ ^^^^^^^^^
+
+@compute @workgroup_size(1)
+fn main() {
+ loop {
+ if (true) {
+ break;
+ } else {
+ break;
+ }
+ let _e16_ = vec2(false, false);
+
+ continuing {
+ break if _e16_.x;
+ }
+ }
+}
diff --git a/test/tint/bug/tint/2202.wgsl b/test/tint/bug/tint/2202.wgsl
new file mode 100644
index 0000000..99b43da
--- /dev/null
+++ b/test/tint/bug/tint/2202.wgsl
@@ -0,0 +1,12 @@
+@compute @workgroup_size(1)
+fn main() {
+ loop {
+ loop {
+ return;
+ }
+ let _e9 = (vec3<i32>().y >= vec3<i32>().y);
+ continuing {
+ break if (_e9 || _e9);
+ }
+ }
+}
diff --git a/test/tint/bug/tint/2202.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/2202.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..2e1b43a
--- /dev/null
+++ b/test/tint/bug/tint/2202.wgsl.expected.dxc.hlsl
@@ -0,0 +1,21 @@
+<dawn>/test/tint/bug/tint/2202.wgsl:7:9 warning: code is unreachable
+ let _e9 = (vec3<i32>().y >= vec3<i32>().y);
+ ^^^^^^^
+
+[numthreads(1, 1, 1)]
+void main() {
+ while (true) {
+ while (true) {
+ return;
+ }
+ bool _e9 = true;
+ {
+ bool tint_tmp = _e9;
+ if (!tint_tmp) {
+ tint_tmp = _e9;
+ }
+ if ((tint_tmp)) { break; }
+ }
+ }
+ return;
+}
diff --git a/test/tint/bug/tint/2202.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/2202.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..2e1b43a
--- /dev/null
+++ b/test/tint/bug/tint/2202.wgsl.expected.fxc.hlsl
@@ -0,0 +1,21 @@
+<dawn>/test/tint/bug/tint/2202.wgsl:7:9 warning: code is unreachable
+ let _e9 = (vec3<i32>().y >= vec3<i32>().y);
+ ^^^^^^^
+
+[numthreads(1, 1, 1)]
+void main() {
+ while (true) {
+ while (true) {
+ return;
+ }
+ bool _e9 = true;
+ {
+ bool tint_tmp = _e9;
+ if (!tint_tmp) {
+ tint_tmp = _e9;
+ }
+ if ((tint_tmp)) { break; }
+ }
+ }
+ return;
+}
diff --git a/test/tint/bug/tint/2202.wgsl.expected.glsl b/test/tint/bug/tint/2202.wgsl.expected.glsl
new file mode 100644
index 0000000..feda98a
--- /dev/null
+++ b/test/tint/bug/tint/2202.wgsl.expected.glsl
@@ -0,0 +1,27 @@
+<dawn>/test/tint/bug/tint/2202.wgsl:7:9 warning: code is unreachable
+ let _e9 = (vec3<i32>().y >= vec3<i32>().y);
+ ^^^^^^^
+
+#version 310 es
+
+void tint_symbol() {
+ while (true) {
+ while (true) {
+ return;
+ }
+ bool _e9 = true;
+ {
+ bool tint_tmp = _e9;
+ if (!tint_tmp) {
+ tint_tmp = _e9;
+ }
+ if ((tint_tmp)) { 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/2202.wgsl.expected.ir.spvasm b/test/tint/bug/tint/2202.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..7c83393
--- /dev/null
+++ b/test/tint/bug/tint/2202.wgsl.expected.ir.spvasm
@@ -0,0 +1,38 @@
+<dawn>/test/tint/bug/tint/2202.wgsl:7:9 warning: code is unreachable
+ let _e9 = (vec3<i32>().y >= vec3<i32>().y);
+ ^^^^^^^
+
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 13
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %main = OpFunction %void None %3
+ %4 = OpLabel
+ OpBranch %7
+ %7 = OpLabel
+ OpLoopMerge %8 %6 None
+ OpBranch %5
+ %5 = OpLabel
+ OpBranch %11
+ %11 = OpLabel
+ OpLoopMerge %12 %10 None
+ OpBranch %9
+ %9 = OpLabel
+ OpBranch %12
+ %10 = OpLabel
+ OpBranch %11
+ %12 = OpLabel
+ OpBranch %8
+ %6 = OpLabel
+ OpBranch %7
+ %8 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/2202.wgsl.expected.msl b/test/tint/bug/tint/2202.wgsl.expected.msl
new file mode 100644
index 0000000..2d601f3
--- /dev/null
+++ b/test/tint/bug/tint/2202.wgsl.expected.msl
@@ -0,0 +1,25 @@
+<dawn>/test/tint/bug/tint/2202.wgsl:7:9 warning: code is unreachable
+ let _e9 = (vec3<i32>().y >= vec3<i32>().y);
+ ^^^^^^^
+
+#include <metal_stdlib>
+
+using namespace metal;
+
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
+kernel void tint_symbol() {
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ return;
+ }
+ bool const _e9 = true;
+ {
+ if ((_e9 || _e9)) { break; }
+ }
+ }
+ return;
+}
+
diff --git a/test/tint/bug/tint/2202.wgsl.expected.spvasm b/test/tint/bug/tint/2202.wgsl.expected.spvasm
new file mode 100644
index 0000000..cb2724b
--- /dev/null
+++ b/test/tint/bug/tint/2202.wgsl.expected.spvasm
@@ -0,0 +1,29 @@
+SKIP: FAILED
+
+<dawn>/test/tint/bug/tint/2202.wgsl:7:9 warning: code is unreachable
+ let _e9 = (vec3<i32>().y >= vec3<i32>().y);
+ ^^^^^^^
+
+
+@compute @workgroup_size(1)
+fn main() {
+ loop {
+ loop {
+ return;
+ }
+ let _e9 = (vec3<i32>().y >= vec3<i32>().y);
+
+ continuing {
+ break if (_e9 || _e9);
+ }
+ }
+}
+
+Failed to generate: <dawn>/test/tint/bug/tint/2202.wgsl:9:23 error: unresolved value '_e9'
+ break if (_e9 || _e9);
+ ^^^
+
+<dawn>/test/tint/bug/tint/2202.wgsl:9:30 error: unresolved value '_e9'
+ break if (_e9 || _e9);
+ ^^^
+
diff --git a/test/tint/bug/tint/2202.wgsl.expected.wgsl b/test/tint/bug/tint/2202.wgsl.expected.wgsl
new file mode 100644
index 0000000..5fa536e
--- /dev/null
+++ b/test/tint/bug/tint/2202.wgsl.expected.wgsl
@@ -0,0 +1,17 @@
+<dawn>/test/tint/bug/tint/2202.wgsl:7:9 warning: code is unreachable
+ let _e9 = (vec3<i32>().y >= vec3<i32>().y);
+ ^^^^^^^
+
+@compute @workgroup_size(1)
+fn main() {
+ loop {
+ loop {
+ return;
+ }
+ let _e9 = (vec3<i32>().y >= vec3<i32>().y);
+
+ continuing {
+ break if (_e9 || _e9);
+ }
+ }
+}
