[spirv-reader][ir] Fix propagation when jumping if controls.
It's possible to jump the exit of an `if` and go straight to the exit of
a `loop` (a `break`). In that case, if we're propagating something from
inside the `if` we do not pass the result up through the `if` but
instead pass it directly through the `exit_loop` and return from the
`loop`. This is done by checking if the control has any exits before
attaching the propagated result.
Bug: 422539666
Change-Id: I03484a38c2758c7d39c0fe2310c649042e6e40b6
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/245757
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Auto-Submit: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/reader/parser/branch_test.cc b/src/tint/lang/spirv/reader/parser/branch_test.cc
index a0d7b17..f1d14c9 100644
--- a/src/tint/lang/spirv/reader/parser/branch_test.cc
+++ b/src/tint/lang/spirv/reader/parser/branch_test.cc
@@ -566,6 +566,83 @@
)");
}
+TEST_F(SpirvParserTest, BranchConditional_HoistingMultiExit_FromMiddle) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %i32 = OpTypeInt 32 1
+ %bool = OpTypeBool
+ %one = OpConstant %i32 1
+ %two = OpConstant %i32 2
+ %true = OpConstantTrue %bool
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ OpBranch %10
+ %10 = OpLabel
+ OpLoopMerge %99 %80 None
+ OpBranch %20
+ %20 = OpLabel
+ OpSelectionMerge %50 None
+ OpBranchConditional %true %30 %30
+ %30 = OpLabel
+ %foo = OpIAdd %i32 %one %two
+ OpSelectionMerge %60 None
+ OpBranchConditional %true %65 %70
+ %65 = OpLabel
+ OpBranch %99
+ %70 = OpLabel
+ OpBranch %99
+ %60 = OpLabel
+ OpUnreachable
+ %50 = OpLabel
+ OpUnreachable
+ %80 = OpLabel
+ OpBranch %10
+ %99 = OpLabel
+ %foo2 = OpCopyObject %i32 %foo
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ %3:bool = or true, true
+ if %3 [t: $B4, f: $B5] { # if_1
+ $B4: { # true
+ %4:i32 = spirv.add<i32> 1i, 2i
+ if true [t: $B6, f: $B7] { # if_2
+ $B6: { # true
+ exit_loop %4 # loop_1
+ }
+ $B7: { # false
+ exit_loop %4 # loop_1
+ }
+ }
+ unreachable
+ }
+ $B5: { # false
+ unreachable
+ }
+ }
+ unreachable
+ }
+ $B3: { # continuing
+ next_iteration # -> $B2
+ }
+ }
+ %5:i32 = let %2
+ ret
+ }
+}
+)");
+}
+
TEST_F(SpirvParserTest, BranchConditional_HoistingIntoNested) {
EXPECT_IR(R"(
OpCapability Shader
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 84a3a22..54be1de 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -839,7 +839,8 @@
// Set `src` as the `param` so it's returned as the new value
src = param;
return nullptr;
- }, //
+ }, //
+ [&](core::ir::Unreachable*) { return blk->Parent(); }, //
TINT_ICE_ON_NO_MATCH);
if (!ctrl) {
@@ -851,6 +852,15 @@
}
}
+ // If this control has no exits, then we don't need to add the result through the
+ // control as we're jumping over the control to its parent control. (This is an
+ // `if` inside a `loop` where the `if` is doing a `break`).
+ if (ctrl->Exits().IsEmpty()) {
+ TINT_ASSERT(ctrl->Is<core::ir::If>());
+ blk = ctrl->Block();
+ continue;
+ }
+
// Add a new result to the control instruction
ctrl->AddResult(b_.InstructionResult(src->Type()));
// The source instruction is now the control result we just inserted
