[spirv-reader][ir] Handle propagation through multiple exits.
It's possible to setup a block where a value which needs to be
propagated dominates multiple blocks which jump to a continuing block.
In this case, we need to pass the propagated value through all of the
exits instead of just a single exit.
Change-Id: I8e4c37f37340899159a3d838bb2f38287efa924f
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/245334
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: James Price <jrprice@google.com>
Auto-Submit: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/spirv/reader/parser/branch_test.cc b/src/tint/lang/spirv/reader/parser/branch_test.cc
index a10d2df..566c562 100644
--- a/src/tint/lang/spirv/reader/parser/branch_test.cc
+++ b/src/tint/lang/spirv/reader/parser/branch_test.cc
@@ -503,6 +503,78 @@
)");
}
+TEST_F(SpirvParserTest, BranchConditional_HoistingMultiExit) {
+ 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 %merge %cont None
+ OpBranch %20
+ %20 = OpLabel
+ %foo = OpIAdd %i32 %one %two
+ OpSelectionMerge %50 None
+ OpBranchConditional %true %30 %40
+ %30 = OpLabel
+ OpBranch %merge
+ %40 = OpLabel
+ OpBranch %merge
+ %50 = OpLabel
+ OpUnreachable
+ %cont = OpLabel
+ OpBranch %10
+ %merge = 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:i32 = spirv.add<i32> 1i, 2i
+ if true [t: $B4, f: $B5] { # if_1
+ $B4: { # true
+ exit_loop %3 # loop_1
+ }
+ $B5: { # false
+ exit_loop %3 # loop_1
+ }
+ }
+ if true [t: $B6, f: $B7] { # if_2
+ $B6: { # true
+ %4:i32 = let %3
+ ret
+ }
+ $B7: { # 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 893370d..7801b9c 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -750,19 +750,20 @@
cont->PushOperand(src);
- // Set the src to the param so we return param as the new
- // value.
+ // 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) {
break;
}
- // Add ourselves as part of the terminator return value
- blk->Terminator()->PushOperand(src);
+ for (auto& exit : ctrl->Exits()) {
+ exit->PushOperand(src);
+ }
+
// 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
