[spirv] Fix unreachable in loop continuing blocks
SPIR-V requires that continue blocks are structurally post-dominated
by back-edge blocks, and the presence of OpUnreachable (a function
terminator) can trip up this validation.
Use a transform to replace unreachable instructions nested inside loop
continuing blocks with regular branches.
Fixed: 354627692
Change-Id: Idd6f58e755f761ba8b8c966ac01619df53974177
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/199896
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: David Neto <dneto@google.com>
diff --git a/src/tint/lang/spirv/writer/loop_test.cc b/src/tint/lang/spirv/writer/loop_test.cc
index a02873e..524d7de 100644
--- a/src/tint/lang/spirv/writer/loop_test.cc
+++ b/src/tint/lang/spirv/writer/loop_test.cc
@@ -336,6 +336,146 @@
)");
}
+// Test that we generate valid SPIR-V when there is an unreachable instruction in the body of a
+// loop nested inside another loop's continuing block. SPIR-V requires that continue blocks are
+// structurally post-dominated by back-edge blocks, and the presence of OpUnreachable (a function
+// terminator) can trip up this validation. See crbug.com/354627692.
+TEST_F(SpirvWriterTest, Loop_NestedLoopInContinuing_UnreachableInNestedBody) {
+ auto* func = b.Function("foo", ty.void_());
+ b.Append(func->Block(), [&] {
+ auto* outer_loop = b.Loop();
+ b.Append(outer_loop->Body(), [&] {
+ b.Continue(outer_loop);
+
+ b.Append(outer_loop->Continuing(), [&] {
+ auto* inner_loop = b.Loop();
+ b.Append(inner_loop->Body(), [&] {
+ auto* ifelse = b.If(true);
+ b.Append(ifelse->True(), [&] { //
+ b.ExitLoop(inner_loop);
+ });
+ b.Append(ifelse->False(), [&] { //
+ b.ExitLoop(inner_loop);
+ });
+ b.Unreachable();
+
+ b.Append(inner_loop->Continuing(), [&] { //
+ b.BreakIf(inner_loop, true);
+ });
+ });
+ b.BreakIf(outer_loop, true);
+ });
+ });
+ b.Return(func);
+ });
+
+ ASSERT_TRUE(Generate()) << Error() << output_;
+ EXPECT_INST(R"(
+ %4 = OpLabel
+ OpBranch %7
+ %7 = OpLabel
+ OpLoopMerge %8 %6 None
+ OpBranch %5
+ %5 = OpLabel
+ OpBranch %6
+ %6 = OpLabel
+ OpBranch %11
+ %11 = OpLabel
+ OpLoopMerge %12 %10 None
+ OpBranch %9
+ %9 = OpLabel
+ OpSelectionMerge %13 None
+ OpBranchConditional %true %14 %15
+ %14 = OpLabel
+ OpBranch %12
+ %15 = OpLabel
+ OpBranch %12
+ %13 = OpLabel
+ OpBranch %12
+ %10 = OpLabel
+ OpBranchConditional %true %12 %11
+ %12 = OpLabel
+ OpBranchConditional %true %8 %7
+ %8 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)");
+}
+
+TEST_F(SpirvWriterTest, Loop_NestedLoopInContinuing_UnreachableInNestedBody_WithResults) {
+ auto* func = b.Function("foo", ty.i32());
+ b.Append(func->Block(), [&] {
+ auto* outer_result = b.InstructionResult(ty.i32());
+ auto* outer_loop = b.Loop();
+ outer_loop->SetResults(Vector{outer_result});
+ b.Append(outer_loop->Body(), [&] {
+ b.Continue(outer_loop);
+
+ b.Append(outer_loop->Continuing(), [&] {
+ auto* inner_result = b.InstructionResult(ty.i32());
+ auto* inner_loop = b.Loop();
+ inner_loop->SetResults(Vector{inner_result});
+ b.Append(inner_loop->Body(), [&] {
+ auto* ifelse = b.If(true);
+ b.Append(ifelse->True(), [&] { //
+ b.ExitLoop(inner_loop, 1_i);
+ });
+ b.Append(ifelse->False(), [&] { //
+ b.ExitLoop(inner_loop, 2_i);
+ });
+ b.Unreachable();
+
+ b.Append(inner_loop->Continuing(), [&] { //
+ b.BreakIf(inner_loop, true, Empty, 3_i);
+ });
+ });
+ b.BreakIf(outer_loop, true, Empty, inner_result);
+ });
+ });
+ b.Return(func, outer_result);
+ });
+
+ ASSERT_TRUE(Generate()) << Error() << output_;
+ EXPECT_INST(R"(
+ %4 = OpLabel
+ OpBranch %7
+ %7 = OpLabel
+ OpLoopMerge %8 %6 None
+ OpBranch %5
+ %5 = OpLabel
+ OpBranch %6
+ %6 = OpLabel
+ OpBranch %11
+ %11 = OpLabel
+ OpLoopMerge %12 %10 None
+ OpBranch %9
+ %9 = OpLabel
+ OpSelectionMerge %13 None
+ OpBranchConditional %true %14 %15
+ %14 = OpLabel
+ OpBranch %12
+ %15 = OpLabel
+ OpBranch %12
+ %13 = OpLabel
+ OpBranch %12
+ %10 = OpLabel
+ OpBranchConditional %true %12 %11
+ %12 = OpLabel
+ %18 = OpPhi %int %int_3 %10 %20 %13 %int_1 %14 %int_2 %15
+ OpBranchConditional %true %8 %7
+ %8 = OpLabel
+ %23 = OpPhi %int %18 %12
+ OpReturnValue %23
+ OpFunctionEnd
+
+ ; Function unused_entry_point
+%unused_entry_point = OpFunction %void None %26
+ %27 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)");
+}
+
TEST_F(SpirvWriterTest, Loop_Phi_SingleValue) {
auto* func = b.Function("foo", ty.void_());
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index e403d0f..647ee0a 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -896,6 +896,7 @@
/// Emit all instructions of @p block.
/// @param block the block's instructions to emit
void EmitBlockInstructions(core::ir::Block* block) {
+ TINT_ASSERT(!block->IsEmpty());
for (auto* inst : *block) {
Switch(
inst, //
@@ -928,11 +929,6 @@
}
}
}
-
- if (block->IsEmpty()) {
- // If the last emitted instruction is not a branch, then this should be unreachable.
- current_function_.push_inst(spv::Op::OpUnreachable, {});
- }
}
/// Emit a terminator instruction.
diff --git a/src/tint/lang/spirv/writer/raise/BUILD.bazel b/src/tint/lang/spirv/writer/raise/BUILD.bazel
index 970bbab..3895fd0 100644
--- a/src/tint/lang/spirv/writer/raise/BUILD.bazel
+++ b/src/tint/lang/spirv/writer/raise/BUILD.bazel
@@ -45,6 +45,7 @@
"merge_return.cc",
"pass_matrix_by_pointer.cc",
"raise.cc",
+ "remove_unreachable_in_loop_continuing.cc",
"shader_io.cc",
"var_for_dynamic_index.cc",
],
@@ -55,6 +56,7 @@
"merge_return.h",
"pass_matrix_by_pointer.h",
"raise.h",
+ "remove_unreachable_in_loop_continuing.h",
"shader_io.h",
"var_for_dynamic_index.h",
],
@@ -107,6 +109,7 @@
"handle_matrix_arithmetic_test.cc",
"merge_return_test.cc",
"pass_matrix_by_pointer_test.cc",
+ "remove_unreachable_in_loop_continuing_test.cc",
"shader_io_test.cc",
"var_for_dynamic_index_test.cc",
],
diff --git a/src/tint/lang/spirv/writer/raise/BUILD.cmake b/src/tint/lang/spirv/writer/raise/BUILD.cmake
index 33ff0e9..71c2ef4 100644
--- a/src/tint/lang/spirv/writer/raise/BUILD.cmake
+++ b/src/tint/lang/spirv/writer/raise/BUILD.cmake
@@ -53,6 +53,8 @@
lang/spirv/writer/raise/pass_matrix_by_pointer.h
lang/spirv/writer/raise/raise.cc
lang/spirv/writer/raise/raise.h
+ lang/spirv/writer/raise/remove_unreachable_in_loop_continuing.cc
+ lang/spirv/writer/raise/remove_unreachable_in_loop_continuing.h
lang/spirv/writer/raise/shader_io.cc
lang/spirv/writer/raise/shader_io.h
lang/spirv/writer/raise/var_for_dynamic_index.cc
@@ -112,6 +114,7 @@
lang/spirv/writer/raise/handle_matrix_arithmetic_test.cc
lang/spirv/writer/raise/merge_return_test.cc
lang/spirv/writer/raise/pass_matrix_by_pointer_test.cc
+ lang/spirv/writer/raise/remove_unreachable_in_loop_continuing_test.cc
lang/spirv/writer/raise/shader_io_test.cc
lang/spirv/writer/raise/var_for_dynamic_index_test.cc
)
diff --git a/src/tint/lang/spirv/writer/raise/BUILD.gn b/src/tint/lang/spirv/writer/raise/BUILD.gn
index a6a8303..2927a46 100644
--- a/src/tint/lang/spirv/writer/raise/BUILD.gn
+++ b/src/tint/lang/spirv/writer/raise/BUILD.gn
@@ -56,6 +56,8 @@
"pass_matrix_by_pointer.h",
"raise.cc",
"raise.h",
+ "remove_unreachable_in_loop_continuing.cc",
+ "remove_unreachable_in_loop_continuing.h",
"shader_io.cc",
"shader_io.h",
"var_for_dynamic_index.cc",
@@ -107,6 +109,7 @@
"handle_matrix_arithmetic_test.cc",
"merge_return_test.cc",
"pass_matrix_by_pointer_test.cc",
+ "remove_unreachable_in_loop_continuing_test.cc",
"shader_io_test.cc",
"var_for_dynamic_index_test.cc",
]
diff --git a/src/tint/lang/spirv/writer/raise/raise.cc b/src/tint/lang/spirv/writer/raise/raise.cc
index 38b8e63..6314869 100644
--- a/src/tint/lang/spirv/writer/raise/raise.cc
+++ b/src/tint/lang/spirv/writer/raise/raise.cc
@@ -51,6 +51,7 @@
#include "src/tint/lang/spirv/writer/raise/handle_matrix_arithmetic.h"
#include "src/tint/lang/spirv/writer/raise/merge_return.h"
#include "src/tint/lang/spirv/writer/raise/pass_matrix_by_pointer.h"
+#include "src/tint/lang/spirv/writer/raise/remove_unreachable_in_loop_continuing.h"
#include "src/tint/lang/spirv/writer/raise/shader_io.h"
#include "src/tint/lang/spirv/writer/raise/var_for_dynamic_index.h"
@@ -142,6 +143,7 @@
RUN_TRANSFORM(raise::ExpandImplicitSplats, module);
RUN_TRANSFORM(raise::HandleMatrixArithmetic, module);
RUN_TRANSFORM(raise::MergeReturn, module);
+ RUN_TRANSFORM(raise::RemoveUnreachableInLoopContinuing, module);
RUN_TRANSFORM(raise::ShaderIO, module,
raise::ShaderIOConfig{options.clamp_frag_depth, options.emit_vertex_point_size,
!options.use_storage_input_output_16});
diff --git a/src/tint/lang/spirv/writer/raise/remove_unreachable_in_loop_continuing.cc b/src/tint/lang/spirv/writer/raise/remove_unreachable_in_loop_continuing.cc
new file mode 100644
index 0000000..4f1fcdb
--- /dev/null
+++ b/src/tint/lang/spirv/writer/raise/remove_unreachable_in_loop_continuing.cc
@@ -0,0 +1,107 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/spirv/writer/raise/remove_unreachable_in_loop_continuing.h"
+
+#include <utility>
+
+#include "src/tint/lang/core/ir/builder.h"
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+using namespace tint::core::number_suffixes; // NOLINT
+using namespace tint::core::fluent_types; // NOLINT
+
+namespace tint::spirv::writer::raise {
+
+namespace {
+
+/// PIMPL state for the transform.
+struct State {
+ /// The IR module.
+ core::ir::Module& ir;
+
+ /// The IR builder.
+ core::ir::Builder b{ir};
+
+ /// Process the module.
+ void Process() {
+ // Find all unreachable instructions.
+ for (auto* inst : ir.Instructions()) {
+ if (auto* unreachable = inst->As<core::ir::Unreachable>()) {
+ Process(unreachable);
+ }
+ }
+ }
+
+ /// Check and replace an unreachable instruction if necessary.
+ /// @param unreachable the instruction to check and maybe replace
+ void Process(core::ir::Unreachable* unreachable) {
+ // Walk up the control stack to see if we are inside a loop continuing block.
+ auto* block = unreachable->Block();
+ while (block->Parent()) {
+ auto* control = block->Parent();
+ if (auto* loop = control->As<core::ir::Loop>()) {
+ if (loop->Continuing() == block) {
+ Replace(unreachable);
+ return;
+ }
+ }
+ block = control->Block();
+ }
+ }
+
+ /// Replace an unreachable instruction.
+ /// @param unreachable the instruction to replace
+ void Replace(core::ir::Unreachable* unreachable) {
+ auto* control = unreachable->Block()->Parent();
+
+ // Fill the exit argument with `undef`.
+ Vector<core::ir::Value*, 4> exit_args;
+ exit_args.Resize(control->Results().Length());
+
+ // Replace the `unreachable` with an instruction that exits from the control construct.
+ auto* exit = b.Exit(unreachable->Block()->Parent(), std::move(exit_args));
+ unreachable->ReplaceWith(exit);
+ unreachable->Destroy();
+ }
+};
+
+} // namespace
+
+Result<SuccessType> RemoveUnreachableInLoopContinuing(core::ir::Module& ir) {
+ auto result = ValidateAndDumpIfNeeded(ir, "RemoveUnreachableInLoopContinuing transform");
+ if (result != Success) {
+ return result;
+ }
+
+ State{ir}.Process();
+
+ return Success;
+}
+
+} // namespace tint::spirv::writer::raise
diff --git a/src/tint/lang/spirv/writer/raise/remove_unreachable_in_loop_continuing.h b/src/tint/lang/spirv/writer/raise/remove_unreachable_in_loop_continuing.h
new file mode 100644
index 0000000..17602fc
--- /dev/null
+++ b/src/tint/lang/spirv/writer/raise/remove_unreachable_in_loop_continuing.h
@@ -0,0 +1,48 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#ifndef SRC_TINT_LANG_SPIRV_WRITER_RAISE_REMOVE_UNREACHABLE_IN_LOOP_CONTINUING_H_
+#define SRC_TINT_LANG_SPIRV_WRITER_RAISE_REMOVE_UNREACHABLE_IN_LOOP_CONTINUING_H_
+
+#include "src/tint/utils/result/result.h"
+
+// Forward declarations.
+namespace tint::core::ir {
+class Module;
+}
+
+namespace tint::spirv::writer::raise {
+
+/// RemoveUnreachableInLoopContinuing is a transform that replaces unreachable statements that are
+/// nested inside a loop continuing block, as SPIR-V's structured control flow rules prohibit this.
+/// @param module the module to transform
+/// @returns success or failure
+Result<SuccessType> RemoveUnreachableInLoopContinuing(core::ir::Module& module);
+
+} // namespace tint::spirv::writer::raise
+
+#endif // SRC_TINT_LANG_SPIRV_WRITER_RAISE_REMOVE_UNREACHABLE_IN_LOOP_CONTINUING_H_
diff --git a/src/tint/lang/spirv/writer/raise/remove_unreachable_in_loop_continuing_test.cc b/src/tint/lang/spirv/writer/raise/remove_unreachable_in_loop_continuing_test.cc
new file mode 100644
index 0000000..196a744
--- /dev/null
+++ b/src/tint/lang/spirv/writer/raise/remove_unreachable_in_loop_continuing_test.cc
@@ -0,0 +1,624 @@
+// Copyright 2023 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/spirv/writer/raise/remove_unreachable_in_loop_continuing.h"
+
+#include <utility>
+
+#include "src/tint/lang/core/ir/transform/helper_test.h"
+
+namespace tint::spirv::writer::raise {
+namespace {
+
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
+
+using SpirvWriter_RemoveUnreachableInLoopContinuingTest = core::ir::transform::TransformTest;
+
+TEST_F(SpirvWriter_RemoveUnreachableInLoopContinuingTest, NoModify_TopLevel) {
+ auto* func = b.Function("foo", ty.void_());
+ b.Append(func->Block(), [&] {
+ auto* ifelse = b.If(true);
+ b.Append(ifelse->True(), [&] { //
+ b.Return(func);
+ });
+ b.Append(ifelse->False(), [&] { //
+ b.Return(func);
+ });
+ b.Unreachable();
+ });
+
+ auto* src = R"(
+%foo = func():void {
+ $B1: {
+ if true [t: $B2, f: $B3] { # if_1
+ $B2: { # true
+ ret
+ }
+ $B3: { # false
+ ret
+ }
+ }
+ unreachable
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ Run(RemoveUnreachableInLoopContinuing);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_RemoveUnreachableInLoopContinuingTest, NoModify_InLoopBody) {
+ auto* func = b.Function("foo", ty.void_());
+ b.Append(func->Block(), [&] {
+ auto* outer_loop = b.Loop();
+ b.Append(outer_loop->Body(), [&] {
+ auto* inner_loop = b.Loop();
+ b.Append(inner_loop->Body(), [&] {
+ auto* ifelse = b.If(true);
+ b.Append(ifelse->True(), [&] { //
+ b.ExitLoop(inner_loop);
+ });
+ b.Append(ifelse->False(), [&] { //
+ b.ExitLoop(inner_loop);
+ });
+ b.Unreachable();
+
+ b.Append(inner_loop->Continuing(), [&] { //
+ b.BreakIf(inner_loop, true);
+ });
+ });
+ b.ExitLoop(outer_loop);
+ });
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%foo = func():void {
+ $B1: {
+ loop [b: $B2] { # loop_1
+ $B2: { # body
+ loop [b: $B3, c: $B4] { # loop_2
+ $B3: { # body
+ if true [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_loop # loop_2
+ }
+ $B6: { # false
+ exit_loop # loop_2
+ }
+ }
+ unreachable
+ }
+ $B4: { # continuing
+ break_if true # -> [t: exit_loop loop_2, f: $B3]
+ }
+ }
+ exit_loop # loop_1
+ }
+ }
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ Run(RemoveUnreachableInLoopContinuing);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_RemoveUnreachableInLoopContinuingTest, InContinuing_NestedInIfBlock) {
+ auto* func = b.Function("foo", ty.void_());
+ b.Append(func->Block(), [&] {
+ auto* outer_loop = b.Loop();
+ b.Append(outer_loop->Body(), [&] {
+ b.Continue(outer_loop);
+
+ b.Append(outer_loop->Continuing(), [&] {
+ auto* ifelse = b.If(true);
+ b.Append(ifelse->True(), [&] { //
+ b.ExitIf(ifelse);
+ });
+ b.Append(ifelse->False(), [&] { //
+ b.Unreachable();
+ });
+ b.BreakIf(outer_loop, true);
+ });
+ });
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%foo = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ if true [t: $B4, f: $B5] { # if_1
+ $B4: { # true
+ exit_if # if_1
+ }
+ $B5: { # false
+ unreachable
+ }
+ }
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ if true [t: $B4, f: $B5] { # if_1
+ $B4: { # true
+ exit_if # if_1
+ }
+ $B5: { # false
+ exit_if # if_1
+ }
+ }
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+
+ Run(RemoveUnreachableInLoopContinuing);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_RemoveUnreachableInLoopContinuingTest, InContinuing_NestedInSwitchCase) {
+ auto* func = b.Function("foo", ty.void_());
+ b.Append(func->Block(), [&] {
+ auto* outer_loop = b.Loop();
+ b.Append(outer_loop->Body(), [&] {
+ b.Continue(outer_loop);
+
+ b.Append(outer_loop->Continuing(), [&] {
+ auto* swtch = b.Switch(1_i);
+ auto* case_1 = b.Case(swtch, {b.Constant(1_i)});
+ auto* def = b.DefaultCase(swtch);
+ b.Append(case_1, [&] { //
+ b.ExitSwitch(swtch);
+ });
+ b.Append(def, [&] { //
+ b.Unreachable();
+ });
+ b.BreakIf(outer_loop, true);
+ });
+ });
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%foo = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ switch 1i [c: (1i, $B4), c: (default, $B5)] { # switch_1
+ $B4: { # case
+ exit_switch # switch_1
+ }
+ $B5: { # case
+ unreachable
+ }
+ }
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ switch 1i [c: (1i, $B4), c: (default, $B5)] { # switch_1
+ $B4: { # case
+ exit_switch # switch_1
+ }
+ $B5: { # case
+ exit_switch # switch_1
+ }
+ }
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+
+ Run(RemoveUnreachableInLoopContinuing);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_RemoveUnreachableInLoopContinuingTest, InContinuing_NestedInLoopBody) {
+ auto* func = b.Function("foo", ty.void_());
+ b.Append(func->Block(), [&] {
+ auto* outer_loop = b.Loop();
+ b.Append(outer_loop->Body(), [&] {
+ b.Continue(outer_loop);
+
+ b.Append(outer_loop->Continuing(), [&] {
+ auto* inner_loop = b.Loop();
+ b.Append(inner_loop->Body(), [&] {
+ auto* ifelse = b.If(true);
+ b.Append(ifelse->True(), [&] { //
+ b.ExitLoop(inner_loop);
+ });
+ b.Append(ifelse->False(), [&] { //
+ b.ExitLoop(inner_loop);
+ });
+ b.Unreachable();
+
+ b.Append(inner_loop->Continuing(), [&] { //
+ b.BreakIf(inner_loop, true);
+ });
+ });
+ b.BreakIf(outer_loop, true);
+ });
+ });
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%foo = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ loop [b: $B4, c: $B5] { # loop_2
+ $B4: { # body
+ if true [t: $B6, f: $B7] { # if_1
+ $B6: { # true
+ exit_loop # loop_2
+ }
+ $B7: { # false
+ exit_loop # loop_2
+ }
+ }
+ unreachable
+ }
+ $B5: { # continuing
+ break_if true # -> [t: exit_loop loop_2, f: $B4]
+ }
+ }
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ loop [b: $B4, c: $B5] { # loop_2
+ $B4: { # body
+ if true [t: $B6, f: $B7] { # if_1
+ $B6: { # true
+ exit_loop # loop_2
+ }
+ $B7: { # false
+ exit_loop # loop_2
+ }
+ }
+ exit_loop # loop_2
+ }
+ $B5: { # continuing
+ break_if true # -> [t: exit_loop loop_2, f: $B4]
+ }
+ }
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+
+ Run(RemoveUnreachableInLoopContinuing);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_RemoveUnreachableInLoopContinuingTest,
+ InContinuing_NestedInLoopBody_WithResults) {
+ auto* func = b.Function("foo", ty.i32());
+ b.Append(func->Block(), [&] {
+ auto* outer_result = b.InstructionResult(ty.i32());
+ auto* outer_loop = b.Loop();
+ outer_loop->SetResults(Vector{outer_result});
+ b.Append(outer_loop->Body(), [&] {
+ b.Continue(outer_loop);
+
+ b.Append(outer_loop->Continuing(), [&] {
+ auto* inner_result = b.InstructionResult(ty.i32());
+ auto* inner_loop = b.Loop();
+ inner_loop->SetResults(Vector{inner_result});
+ b.Append(inner_loop->Body(), [&] {
+ auto* ifelse = b.If(true);
+ b.Append(ifelse->True(), [&] { //
+ b.ExitLoop(inner_loop, 1_i);
+ });
+ b.Append(ifelse->False(), [&] { //
+ b.ExitLoop(inner_loop, 2_i);
+ });
+ b.Unreachable();
+
+ b.Append(inner_loop->Continuing(), [&] { //
+ b.BreakIf(inner_loop, true, Empty, 3_i);
+ });
+ });
+ b.BreakIf(outer_loop, true, Empty, inner_result);
+ });
+ });
+ b.Return(func, outer_result);
+ });
+
+ auto* src = R"(
+%foo = func():i32 {
+ $B1: {
+ %2:i32 = loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ %3:i32 = loop [b: $B4, c: $B5] { # loop_2
+ $B4: { # body
+ if true [t: $B6, f: $B7] { # if_1
+ $B6: { # true
+ exit_loop 1i # loop_2
+ }
+ $B7: { # false
+ exit_loop 2i # loop_2
+ }
+ }
+ unreachable
+ }
+ $B5: { # continuing
+ break_if true exit_loop: [ 3i ] # -> [t: exit_loop loop_2, f: $B4]
+ }
+ }
+ break_if true exit_loop: [ %3 ] # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret %2
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func():i32 {
+ $B1: {
+ %2:i32 = loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ %3:i32 = loop [b: $B4, c: $B5] { # loop_2
+ $B4: { # body
+ if true [t: $B6, f: $B7] { # if_1
+ $B6: { # true
+ exit_loop 1i # loop_2
+ }
+ $B7: { # false
+ exit_loop 2i # loop_2
+ }
+ }
+ exit_loop undef # loop_2
+ }
+ $B5: { # continuing
+ break_if true exit_loop: [ 3i ] # -> [t: exit_loop loop_2, f: $B4]
+ }
+ }
+ break_if true exit_loop: [ %3 ] # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret %2
+ }
+}
+)";
+
+ Run(RemoveUnreachableInLoopContinuing);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvWriter_RemoveUnreachableInLoopContinuingTest, InContinuing_DeeplyNestedInLoopBody) {
+ auto* func = b.Function("foo", ty.void_());
+ b.Append(func->Block(), [&] {
+ auto* outer_loop = b.Loop();
+ b.Append(outer_loop->Body(), [&] {
+ b.Continue(outer_loop);
+
+ b.Append(outer_loop->Continuing(), [&] {
+ auto* outer_if = b.If(true);
+ b.Append(outer_if->True(), [&] {
+ auto* inner_loop = b.Loop();
+ b.Append(inner_loop->Body(), [&] {
+ auto* inner_if = b.If(true);
+ b.Append(inner_if->True(), [&] {
+ auto* ifelse = b.If(true);
+ b.Append(ifelse->True(), [&] { //
+ b.ExitLoop(inner_loop);
+ });
+ b.Append(ifelse->False(), [&] { //
+ b.ExitLoop(inner_loop);
+ });
+ b.Unreachable();
+ });
+ b.Unreachable();
+
+ b.Append(inner_loop->Continuing(), [&] { //
+ b.BreakIf(inner_loop, true);
+ });
+ });
+ b.ExitIf(outer_if);
+ });
+ b.BreakIf(outer_loop, true);
+ });
+ });
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%foo = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ if true [t: $B4] { # if_1
+ $B4: { # true
+ loop [b: $B5, c: $B6] { # loop_2
+ $B5: { # body
+ if true [t: $B7] { # if_2
+ $B7: { # true
+ if true [t: $B8, f: $B9] { # if_3
+ $B8: { # true
+ exit_loop # loop_2
+ }
+ $B9: { # false
+ exit_loop # loop_2
+ }
+ }
+ unreachable
+ }
+ }
+ unreachable
+ }
+ $B6: { # continuing
+ break_if true # -> [t: exit_loop loop_2, f: $B5]
+ }
+ }
+ exit_if # if_1
+ }
+ }
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ if true [t: $B4] { # if_1
+ $B4: { # true
+ loop [b: $B5, c: $B6] { # loop_2
+ $B5: { # body
+ if true [t: $B7] { # if_2
+ $B7: { # true
+ if true [t: $B8, f: $B9] { # if_3
+ $B8: { # true
+ exit_loop # loop_2
+ }
+ $B9: { # false
+ exit_loop # loop_2
+ }
+ }
+ exit_if # if_2
+ }
+ }
+ exit_loop # loop_2
+ }
+ $B6: { # continuing
+ break_if true # -> [t: exit_loop loop_2, f: $B5]
+ }
+ }
+ exit_if # if_1
+ }
+ }
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+
+ Run(RemoveUnreachableInLoopContinuing);
+
+ EXPECT_EQ(expect, str());
+}
+
+} // namespace
+} // namespace tint::spirv::writer::raise
diff --git a/test/tint/bug/tint/354627692.wgsl b/test/tint/bug/tint/354627692.wgsl
new file mode 100644
index 0000000..b3f18aa
--- /dev/null
+++ b/test/tint/bug/tint/354627692.wgsl
@@ -0,0 +1,22 @@
+@group(0) @binding(0)
+var<storage, read_write> buffer : i32;
+
+@compute @workgroup_size(1)
+fn main() {
+ var i : i32 = buffer;
+ loop {
+ continuing {
+ loop {
+ if (i > 5) {
+ i = i * 2;
+ break;
+ } else {
+ i = i * 2;
+ break;
+ }
+ }
+ break if i > 10;
+ }
+ }
+ buffer = i;
+}
diff --git a/test/tint/bug/tint/354627692.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/354627692.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..aacc43f
--- /dev/null
+++ b/test/tint/bug/tint/354627692.wgsl.expected.dxc.hlsl
@@ -0,0 +1,22 @@
+RWByteAddressBuffer buffer : register(u0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ int i = asint(buffer.Load(0u));
+ while (true) {
+ {
+ while (true) {
+ if ((i > 5)) {
+ i = (i * 2);
+ break;
+ } else {
+ i = (i * 2);
+ break;
+ }
+ }
+ if ((i > 10)) { break; }
+ }
+ }
+ buffer.Store(0u, asuint(i));
+ return;
+}
diff --git a/test/tint/bug/tint/354627692.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/354627692.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..aacc43f
--- /dev/null
+++ b/test/tint/bug/tint/354627692.wgsl.expected.fxc.hlsl
@@ -0,0 +1,22 @@
+RWByteAddressBuffer buffer : register(u0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ int i = asint(buffer.Load(0u));
+ while (true) {
+ {
+ while (true) {
+ if ((i > 5)) {
+ i = (i * 2);
+ break;
+ } else {
+ i = (i * 2);
+ break;
+ }
+ }
+ if ((i > 10)) { break; }
+ }
+ }
+ buffer.Store(0u, asuint(i));
+ return;
+}
diff --git a/test/tint/bug/tint/354627692.wgsl.expected.glsl b/test/tint/bug/tint/354627692.wgsl.expected.glsl
new file mode 100644
index 0000000..b2c6d01
--- /dev/null
+++ b/test/tint/bug/tint/354627692.wgsl.expected.glsl
@@ -0,0 +1,30 @@
+#version 310 es
+
+layout(binding = 0, std430) buffer tint_symbol_block_ssbo {
+ int inner;
+} tint_symbol;
+
+void tint_symbol_1() {
+ int i = tint_symbol.inner;
+ while (true) {
+ {
+ while (true) {
+ if ((i > 5)) {
+ i = (i * 2);
+ break;
+ } else {
+ i = (i * 2);
+ break;
+ }
+ }
+ if ((i > 10)) { break; }
+ }
+ }
+ tint_symbol.inner = i;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol_1();
+ return;
+}
diff --git a/test/tint/bug/tint/354627692.wgsl.expected.ir.dxc.hlsl b/test/tint/bug/tint/354627692.wgsl.expected.ir.dxc.hlsl
new file mode 100644
index 0000000..7b45bb1
--- /dev/null
+++ b/test/tint/bug/tint/354627692.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,28 @@
+
+RWByteAddressBuffer buffer : register(u0);
+[numthreads(1, 1, 1)]
+void main() {
+ int i = asint(buffer.Load(0u));
+ {
+ while(true) {
+ {
+ {
+ while(true) {
+ if ((i > 5)) {
+ i = (i * 2);
+ break;
+ } else {
+ i = (i * 2);
+ break;
+ }
+ /* unreachable */
+ }
+ }
+ if ((i > 10)) { break; }
+ }
+ continue;
+ }
+ }
+ buffer.Store(0u, asuint(i));
+}
+
diff --git a/test/tint/bug/tint/354627692.wgsl.expected.ir.fxc.hlsl b/test/tint/bug/tint/354627692.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..7b45bb1
--- /dev/null
+++ b/test/tint/bug/tint/354627692.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,28 @@
+
+RWByteAddressBuffer buffer : register(u0);
+[numthreads(1, 1, 1)]
+void main() {
+ int i = asint(buffer.Load(0u));
+ {
+ while(true) {
+ {
+ {
+ while(true) {
+ if ((i > 5)) {
+ i = (i * 2);
+ break;
+ } else {
+ i = (i * 2);
+ break;
+ }
+ /* unreachable */
+ }
+ }
+ if ((i > 10)) { break; }
+ }
+ continue;
+ }
+ }
+ buffer.Store(0u, asuint(i));
+}
+
diff --git a/test/tint/bug/tint/354627692.wgsl.expected.ir.msl b/test/tint/bug/tint/354627692.wgsl.expected.ir.msl
new file mode 100644
index 0000000..806cf8e
--- /dev/null
+++ b/test/tint/bug/tint/354627692.wgsl.expected.ir.msl
@@ -0,0 +1,32 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_module_vars_struct {
+ device int* tint_symbol;
+};
+
+kernel void tint_symbol_1(device int* tint_symbol [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.tint_symbol=tint_symbol};
+ int i = (*tint_module_vars.tint_symbol);
+ {
+ while(true) {
+ {
+ {
+ while(true) {
+ if ((i > 5)) {
+ i = (i * 2);
+ break;
+ } else {
+ i = (i * 2);
+ break;
+ }
+ /* unreachable */
+ }
+ }
+ if ((i > 10)) { break; }
+ }
+ continue;
+ }
+ }
+ (*tint_module_vars.tint_symbol) = i;
+}
diff --git a/test/tint/bug/tint/354627692.wgsl.expected.ir.spvasm b/test/tint/bug/tint/354627692.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..47b9c26
--- /dev/null
+++ b/test/tint/bug/tint/354627692.wgsl.expected.ir.spvasm
@@ -0,0 +1,77 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 41
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpMemberName %tint_symbol_1 0 "tint_symbol"
+ OpName %tint_symbol_1 "tint_symbol_1"
+ OpName %main "main"
+ OpName %i "i"
+ OpMemberDecorate %tint_symbol_1 0 Offset 0
+ OpDecorate %tint_symbol_1 Block
+ OpDecorate %1 DescriptorSet 0
+ OpDecorate %1 Binding 0
+ %int = OpTypeInt 32 1
+%tint_symbol_1 = OpTypeStruct %int
+%_ptr_StorageBuffer_tint_symbol_1 = OpTypePointer StorageBuffer %tint_symbol_1
+ %1 = OpVariable %_ptr_StorageBuffer_tint_symbol_1 StorageBuffer
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_5 = OpConstant %int 5
+ %bool = OpTypeBool
+ %int_2 = OpConstant %int 2
+ %int_10 = OpConstant %int 10
+ %main = OpFunction %void None %7
+ %8 = OpLabel
+ %i = OpVariable %_ptr_Function_int Function
+ %9 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0
+ %13 = OpLoad %int %9
+ OpStore %i %13
+ OpBranch %18
+ %18 = OpLabel
+ OpLoopMerge %19 %17 None
+ OpBranch %16
+ %16 = OpLabel
+ OpBranch %17
+ %17 = OpLabel
+ OpBranch %22
+ %22 = OpLabel
+ OpLoopMerge %23 %21 None
+ OpBranch %20
+ %20 = OpLabel
+ %24 = OpLoad %int %i
+ %25 = OpSGreaterThan %bool %24 %int_5
+ OpSelectionMerge %28 None
+ OpBranchConditional %25 %29 %30
+ %29 = OpLabel
+ %31 = OpLoad %int %i
+ %32 = OpIMul %int %31 %int_2
+ OpStore %i %32
+ OpBranch %23
+ %30 = OpLabel
+ %34 = OpLoad %int %i
+ %35 = OpIMul %int %34 %int_2
+ OpStore %i %35
+ OpBranch %23
+ %28 = OpLabel
+ OpBranch %23
+ %21 = OpLabel
+ OpBranch %22
+ %23 = OpLabel
+ %36 = OpLoad %int %i
+ %37 = OpSGreaterThan %bool %36 %int_10
+ OpBranchConditional %37 %19 %18
+ %19 = OpLabel
+ %39 = OpLoad %int %i
+ %40 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0
+ OpStore %40 %39
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/354627692.wgsl.expected.msl b/test/tint/bug/tint/354627692.wgsl.expected.msl
new file mode 100644
index 0000000..977f57a
--- /dev/null
+++ b/test/tint/bug/tint/354627692.wgsl.expected.msl
@@ -0,0 +1,28 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
+kernel void tint_symbol_1(device int* tint_symbol_2 [[buffer(0)]]) {
+ int i = *(tint_symbol_2);
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ {
+ TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
+ if ((i > 5)) {
+ i = as_type<int>((as_type<uint>(i) * as_type<uint>(2)));
+ break;
+ } else {
+ i = as_type<int>((as_type<uint>(i) * as_type<uint>(2)));
+ break;
+ }
+ }
+ if ((i > 10)) { break; }
+ }
+ }
+ *(tint_symbol_2) = i;
+ return;
+}
+
diff --git a/test/tint/bug/tint/354627692.wgsl.expected.spvasm b/test/tint/bug/tint/354627692.wgsl.expected.spvasm
new file mode 100644
index 0000000..2fee22e
--- /dev/null
+++ b/test/tint/bug/tint/354627692.wgsl.expected.spvasm
@@ -0,0 +1,79 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 42
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %buffer_block "buffer_block"
+ OpMemberName %buffer_block 0 "inner"
+ OpName %buffer "buffer"
+ OpName %main "main"
+ OpName %i "i"
+ OpDecorate %buffer_block Block
+ OpMemberDecorate %buffer_block 0 Offset 0
+ OpDecorate %buffer DescriptorSet 0
+ OpDecorate %buffer Binding 0
+ %int = OpTypeInt 32 1
+%buffer_block = OpTypeStruct %int
+%_ptr_StorageBuffer_buffer_block = OpTypePointer StorageBuffer %buffer_block
+ %buffer = OpVariable %_ptr_StorageBuffer_buffer_block StorageBuffer
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+%_ptr_Function_int = OpTypePointer Function %int
+ %16 = OpConstantNull %int
+ %int_5 = OpConstant %int 5
+ %bool = OpTypeBool
+ %int_2 = OpConstant %int 2
+ %int_10 = OpConstant %int 10
+ %main = OpFunction %void None %5
+ %8 = OpLabel
+ %i = OpVariable %_ptr_Function_int Function %16
+ %12 = OpAccessChain %_ptr_StorageBuffer_int %buffer %uint_0
+ %13 = OpLoad %int %12
+ OpStore %i %13
+ OpBranch %17
+ %17 = OpLabel
+ OpLoopMerge %18 %19 None
+ OpBranch %20
+ %20 = OpLabel
+ OpBranch %19
+ %19 = OpLabel
+ OpBranch %21
+ %21 = OpLabel
+ OpLoopMerge %22 %23 None
+ OpBranch %24
+ %24 = OpLabel
+ %25 = OpLoad %int %i
+ %27 = OpSGreaterThan %bool %25 %int_5
+ OpSelectionMerge %29 None
+ OpBranchConditional %27 %30 %31
+ %30 = OpLabel
+ %32 = OpLoad %int %i
+ %34 = OpIMul %int %32 %int_2
+ OpStore %i %34
+ OpBranch %22
+ %31 = OpLabel
+ %35 = OpLoad %int %i
+ %36 = OpIMul %int %35 %int_2
+ OpStore %i %36
+ OpBranch %22
+ %29 = OpLabel
+ OpBranch %23
+ %23 = OpLabel
+ OpBranch %21
+ %22 = OpLabel
+ %37 = OpLoad %int %i
+ %39 = OpSGreaterThan %bool %37 %int_10
+ OpBranchConditional %39 %18 %17
+ %18 = OpLabel
+ %40 = OpAccessChain %_ptr_StorageBuffer_int %buffer %uint_0
+ %41 = OpLoad %int %i
+ OpStore %40 %41
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/354627692.wgsl.expected.wgsl b/test/tint/bug/tint/354627692.wgsl.expected.wgsl
new file mode 100644
index 0000000..1543167
--- /dev/null
+++ b/test/tint/bug/tint/354627692.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+@group(0) @binding(0) var<storage, read_write> buffer : i32;
+
+@compute @workgroup_size(1)
+fn main() {
+ var i : i32 = buffer;
+ loop {
+
+ continuing {
+ loop {
+ if ((i > 5)) {
+ i = (i * 2);
+ break;
+ } else {
+ i = (i * 2);
+ break;
+ }
+ }
+ break if (i > 10);
+ }
+ }
+ buffer = i;
+}
