[spirv-reader][ir] Handle value propagation into the continue block.
A value declared in the body of a loop can propagate into the continuing
block without an `OpPhi`, in the case there is only one branch into the
continuing block. This currently breaks the propagation logic as we have
to pass the value as a block parameter and add as an operand to the
terminator instead of pushing through the control instructions.
Bug: 42250952
Change-Id: I348e3c5984907d7d9368054d2d3dd17c685b279b
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/244154
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/core/ir/multi_in_block.cc b/src/tint/lang/core/ir/multi_in_block.cc
index a26ea2d..88d0188 100644
--- a/src/tint/lang/core/ir/multi_in_block.cc
+++ b/src/tint/lang/core/ir/multi_in_block.cc
@@ -75,6 +75,11 @@
}
}
+void MultiInBlock::AddParam(BlockParam* param) {
+ params_.Push(param);
+ param->SetBlock(this);
+}
+
void MultiInBlock::AddInboundSiblingBranch(ir::Terminator* node) {
TINT_ASSERT(node != nullptr);
inbound_sibling_branches_.Push(node);
diff --git a/src/tint/lang/core/ir/multi_in_block.h b/src/tint/lang/core/ir/multi_in_block.h
index 443e716..ecc0e5b 100644
--- a/src/tint/lang/core/ir/multi_in_block.h
+++ b/src/tint/lang/core/ir/multi_in_block.h
@@ -58,6 +58,10 @@
/// @param params the params for the block
void SetParams(std::initializer_list<BlockParam*> params);
+ /// Adds the param to the block
+ /// @param param the param for the block
+ void AddParam(BlockParam* param);
+
/// @returns the params to the block
VectorRef<BlockParam*> Params() { return params_; }
diff --git a/src/tint/lang/spirv/reader/parser/branch_test.cc b/src/tint/lang/spirv/reader/parser/branch_test.cc
index 8767950..d0fc55d 100644
--- a/src/tint/lang/spirv/reader/parser/branch_test.cc
+++ b/src/tint/lang/spirv/reader/parser/branch_test.cc
@@ -2261,6 +2261,64 @@
)");
}
+TEST_F(SpirvParserTest, Loop_ContinueUseBodyValue) {
+ 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
+ %three = OpConstant %i32 3
+ %true = OpConstantTrue %bool
+ %false = OpConstantFalse %bool
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %10 = OpLabel
+ OpBranch %20
+ %20 = OpLabel
+ OpLoopMerge %99 %50 None
+ OpBranchConditional %true %30 %99
+ %30 = OpLabel
+ %40 = OpIAdd %i32 %one %two
+ OpBranch %50
+ %50 = OpLabel
+ %41 = OpIAdd %i32 %40 %40
+ OpBranch %20
+ %99 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ if true [t: $B4, f: $B5] { # if_1
+ $B4: { # true
+ %2:i32 = spirv.add<i32> 1i, 2i
+ continue %2 # -> $B3
+ }
+ $B5: { # false
+ exit_loop # loop_1
+ }
+ }
+ unreachable
+ }
+ $B3 (%3:i32): { # continuing
+ %4:i32 = spirv.add<i32> %3, %3
+ next_iteration # -> $B2
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
TEST_F(SpirvParserTest, Loop_Body_Switch) {
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 21605a1..872c48b 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -725,10 +725,42 @@
break;
}
- auto* ctrl = blk->Parent();
-
TINT_ASSERT(blk->Terminator());
+ core::ir::ControlInstruction* ctrl = tint::Switch(
+ blk->Terminator(), //
+ [&](core::ir::ExitIf* ei) { return ei->If(); },
+ [&](core::ir::ExitSwitch* es) { return es->Switch(); },
+ [&](core::ir::ExitLoop* el) { return el->Loop(); },
+ [&](core::ir::Continue* cont) {
+ // The propagation is going through a `continue`. This means
+ // this is the only path to the continuing block, but it also
+ // means we're current in the continuing block. We can't do
+ // normal propagation here, we have to pass a block param
+ // instead.
+
+ auto* param = b_.BlockParam(src->Type());
+
+ // We're in the continuing block, so make the block param available in the
+ // scope.
+ id_stack_.back().insert(id);
+
+ auto* loop = cont->Loop();
+ loop->Continuing()->AddParam(param);
+
+ cont->PushOperand(src);
+
+ // Set the src to the param so we return param as the new
+ // value.
+ src = param;
+ return nullptr;
+ }, //
+ TINT_ICE_ON_NO_MATCH);
+
+ if (!ctrl) {
+ break;
+ }
+
// Add ourselves as part of the terminator return value
blk->Terminator()->PushOperand(src);
// Add a new result to the control instruction
@@ -1441,8 +1473,25 @@
auto* loop = StopWalkingAt(src.id())->As<core::ir::Loop>();
TINT_ASSERT(loop);
+ auto continue_id = loop_merge_inst->GetSingleWordInOperand(1);
+ if (continue_id != src.id()) {
+ const auto& bb_continue = current_spirv_function_->FindBlock(continue_id);
+
+ // Emit the continuing block.
+ EmitBlockParent(loop->Continuing(), *bb_continue);
+ }
+
+ // Add the body terminator after processing the continuing block in
+ // case we've propagated values through the continue.
if (!loop->Body()->Terminator()) {
- loop->Body()->Append(b_.Continue(loop));
+ if (loop->Continuing()->Params().IsEmpty()) {
+ loop->Body()->Append(b_.Continue(loop));
+ } else {
+ loop->Body()->Append(b_.Unreachable());
+ }
+ }
+ if (!loop->Continuing()->Terminator()) {
+ loop->Continuing()->Append(b_.NextIteration(loop));
}
current_blocks_.erase(loop->Body());
@@ -2037,15 +2086,6 @@
walk_stop_blocks_.insert({merge_id, loop});
if (continue_id != header_id) {
walk_stop_blocks_.insert({continue_id, loop});
-
- const auto& bb_continue = current_spirv_function_->FindBlock(continue_id);
-
- // Emit the continuing block.
- EmitBlockParent(loop->Continuing(), *bb_continue);
- }
-
- if (!loop->Continuing()->Terminator()) {
- loop->Continuing()->Append(b_.NextIteration(loop));
}
// The remainder of the loop body will process when we hit the
