[spirv-reader][ir] Fix issue where true branches to false
In the case where one branch from `OpBranchConditional` branches to the
other branch we were incorrectly determining that there was no
pre-merge. In this case, we need to treat the entire branch as the
premerge.
Bug: 391487552
Change-Id: Id6b86ae8f5312fb974c929c49bcbacc1c559b635
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/230454
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: 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 0760a87..2c004cb 100644
--- a/src/tint/lang/spirv/reader/parser/branch_test.cc
+++ b/src/tint/lang/spirv/reader/parser/branch_test.cc
@@ -400,7 +400,7 @@
)");
}
-TEST_F(SpirvParserTest, DISABLED_BranchConditional_TrueBranchesToFalse) {
+TEST_F(SpirvParserTest, BranchConditional_TrueBranchesToFalse) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -412,12 +412,12 @@
%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
%20 = OpLabel
OpSelectionMerge %99 None
- OpBranchConditional %true %40 %30
+ OpBranchConditional %false %40 %30
%30 = OpLabel
%50 = OpIAdd %i32 %one %one
OpReturn
@@ -431,7 +431,7 @@
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- if true [t: $B2, f: $B3] { # if_1
+ if false [t: $B2, f: $B3] { # if_1
$B2: { # true
%2:i32 = spirv.add<i32> 2i, 2i
exit_if # if_1
@@ -440,7 +440,15 @@
exit_if # if_1
}
}
- %3:i32 = spirv.add<i32> 1i, 1i
+ if true [t: $B4, f: $B5] { # if_2
+ $B4: { # true
+ %3:i32 = spirv.add<i32> 1i, 1i
+ ret
+ }
+ $B5: { # false
+ unreachable
+ }
+ }
ret
}
}
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 4c3fcc5..a364a3f 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -1190,15 +1190,30 @@
current_spirv_function_, &*(current_spirv_function_->FindBlock(false_id)),
&*(current_spirv_function_->FindBlock(merge_id.value())), &false_blocks);
- // It's possible one of the blocks returns, so bail early if they don't both end at the
- // merge.
- if (true_blocks.back()->id() != merge_id || false_blocks.back()->id() != merge_id) {
+ auto& true_end = true_blocks.back();
+ auto& false_end = false_blocks.back();
+
+ // We only consider the block as returning if it didn't return through
+ // the merge block. (I.e. it's a direct exit from inside the branch
+ // itself.
+ bool true_returns = true_end->id() != merge_id && true_end->IsReturn();
+ bool false_returns = false_end->id() != merge_id && false_end->IsReturn();
+ // If one of the blocks returns but the other doesn't, then we can't
+ // have a premerge block.
+ if (true_returns != false_returns) {
return std::nullopt;
}
- // Remove the merge blocks.
- true_blocks.pop_back();
- false_blocks.pop_back();
+ // If they don't return, both blocks must merge to the same place.
+ if (!true_returns && (true_end->id() != false_end->id())) {
+ return std::nullopt;
+ }
+
+ // If these aren't returns, then remove the merge blocks.
+ if (!true_returns) {
+ true_blocks.pop_back();
+ false_blocks.pop_back();
+ }
std::optional<uint32_t> id = std::nullopt;
while (!true_blocks.empty() && !false_blocks.empty()) {
