[ir] Emit unreachable for loop body exits
When the behavior of the loop body does not contain 'Next', emit
`unreachable` instead of `continue`.
Change-Id: I736829b6e27f4b2a1ae5e5fbe4a1aa2e7bfd4c47
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/179301
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@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 f54aba6..c79c6c8 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
@@ -575,14 +575,20 @@
ControlStackScope scope(this, loop_inst);
+ auto& body_behaviors = program_.Sem().Get(stmt->body)->Behaviors();
{
TINT_SCOPED_ASSIGNMENT(current_block_, loop_inst->Body());
EmitStatements(stmt->body->statements);
- // The current block didn't `break`, `return` or `continue`, go to the continuing block.
+ // The current block didn't `break`, `return` or `continue`, go to the continuing block
+ // or mark the end of the block as unreachable.
if (NeedTerminator()) {
- SetTerminator(builder_.Continue(loop_inst));
+ if (body_behaviors.Contains(sem::Behavior::kNext)) {
+ SetTerminator(builder_.Continue(loop_inst));
+ } else {
+ SetTerminator(builder_.Unreachable());
+ }
}
}
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 9867cb4..e790c55 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
@@ -545,7 +545,7 @@
ASSERT_EQ(1u, m.functions.Length());
EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
- EXPECT_EQ(1u, loop->Continuing()->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 {
@@ -560,7 +560,7 @@
exit_loop # loop_1
}
}
- continue %b3
+ unreachable
}
%b3 = block { # continuing
next_iteration %b2
