[msl] Add RemoveContinueInSwitch transform
The textual backends inline continuing blocks at all `continue`
callsites, which means that they may generate a `break` that breaks
from a switch instead of the intended loop.
This transforms sets a flag and performs the `continue` after switch
instruction instead.
Bug: 42251016
Change-Id: I3826dc21de9f1035afa7de2baac94a14bcbf3bf2
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/204377
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/core/ir/transform/BUILD.bazel b/src/tint/lang/core/ir/transform/BUILD.bazel
index e562c61..b90c54f 100644
--- a/src/tint/lang/core/ir/transform/BUILD.bazel
+++ b/src/tint/lang/core/ir/transform/BUILD.bazel
@@ -52,6 +52,7 @@
"direct_variable_access.cc",
"multiplanar_external_texture.cc",
"preserve_padding.cc",
+ "remove_continue_in_switch.cc",
"remove_terminator_args.cc",
"rename_conflicts.cc",
"robustness.cc",
@@ -75,6 +76,7 @@
"direct_variable_access.h",
"multiplanar_external_texture.h",
"preserve_padding.h",
+ "remove_continue_in_switch.h",
"remove_terminator_args.h",
"rename_conflicts.h",
"robustness.h",
@@ -129,6 +131,7 @@
"helper_test.h",
"multiplanar_external_texture_test.cc",
"preserve_padding_test.cc",
+ "remove_continue_in_switch_test.cc",
"remove_terminator_args_test.cc",
"rename_conflicts_test.cc",
"robustness_test.cc",
diff --git a/src/tint/lang/core/ir/transform/BUILD.cmake b/src/tint/lang/core/ir/transform/BUILD.cmake
index 046016d..1687918 100644
--- a/src/tint/lang/core/ir/transform/BUILD.cmake
+++ b/src/tint/lang/core/ir/transform/BUILD.cmake
@@ -67,6 +67,8 @@
lang/core/ir/transform/multiplanar_external_texture.h
lang/core/ir/transform/preserve_padding.cc
lang/core/ir/transform/preserve_padding.h
+ lang/core/ir/transform/remove_continue_in_switch.cc
+ lang/core/ir/transform/remove_continue_in_switch.h
lang/core/ir/transform/remove_terminator_args.cc
lang/core/ir/transform/remove_terminator_args.h
lang/core/ir/transform/rename_conflicts.cc
@@ -132,6 +134,7 @@
lang/core/ir/transform/helper_test.h
lang/core/ir/transform/multiplanar_external_texture_test.cc
lang/core/ir/transform/preserve_padding_test.cc
+ lang/core/ir/transform/remove_continue_in_switch_test.cc
lang/core/ir/transform/remove_terminator_args_test.cc
lang/core/ir/transform/rename_conflicts_test.cc
lang/core/ir/transform/robustness_test.cc
diff --git a/src/tint/lang/core/ir/transform/BUILD.gn b/src/tint/lang/core/ir/transform/BUILD.gn
index 61b989a..9387529 100644
--- a/src/tint/lang/core/ir/transform/BUILD.gn
+++ b/src/tint/lang/core/ir/transform/BUILD.gn
@@ -71,6 +71,8 @@
"multiplanar_external_texture.h",
"preserve_padding.cc",
"preserve_padding.h",
+ "remove_continue_in_switch.cc",
+ "remove_continue_in_switch.h",
"remove_terminator_args.cc",
"remove_terminator_args.h",
"rename_conflicts.cc",
@@ -130,6 +132,7 @@
"helper_test.h",
"multiplanar_external_texture_test.cc",
"preserve_padding_test.cc",
+ "remove_continue_in_switch_test.cc",
"remove_terminator_args_test.cc",
"rename_conflicts_test.cc",
"robustness_test.cc",
diff --git a/src/tint/lang/core/ir/transform/remove_continue_in_switch.cc b/src/tint/lang/core/ir/transform/remove_continue_in_switch.cc
new file mode 100644
index 0000000..d056862
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/remove_continue_in_switch.cc
@@ -0,0 +1,136 @@
+// 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/core/ir/transform/remove_continue_in_switch.h"
+
+#include "src/tint/lang/core/ir/builder.h"
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/utils/ice/ice.h"
+
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
+
+namespace tint::core::ir::transform {
+
+namespace {
+
+/// PIMPL state for the transform.
+struct State {
+ /// The IR module.
+ Module& ir;
+
+ /// The IR builder.
+ Builder b{ir};
+
+ /// The type manager.
+ core::type::Manager& ty{ir.Types()};
+
+ /// A map from `switch` instruction to the flag used to indicate whether a `continue` was hit.
+ Hashmap<Switch*, Var*, 4> continue_flag_for_switch{};
+
+ /// Process the module.
+ void Process() {
+ // Look for `continue` instructions.
+ for (auto* inst : ir.Instructions()) {
+ auto* cont = inst->As<Continue>();
+ if (!cont) {
+ continue;
+ }
+
+ // Skip `continue` instructions in loops whose continuing blocks do not end with
+ // `breakif`, as these continuing blocks can be safely inlined.
+ if (!tint::Is<BreakIf>(cont->Loop()->Continuing()->Terminator())) {
+ continue;
+ }
+
+ // Check if this `continue` is inside a `switch` that is inside the `loop`.
+ // Do this by walking up the stack of control flow instructions until we see a `loop`.
+ // If we hit a `switch` before we see the loop, we need to replace the `continue`.
+ auto* parent = cont->Block()->Parent();
+ while (!parent->Is<Loop>()) {
+ if (auto* swtch = parent->As<Switch>()) {
+ ReplaceContinue(cont, swtch);
+ break;
+ }
+ parent = parent->Block()->Parent();
+ }
+ }
+ }
+
+ /// Replace a `continue` instruction.
+ /// @param cont the `continue` to replace
+ /// @param swtch the `switch` instruction that it is nested inside
+ void ReplaceContinue(Continue* cont, Switch* swtch) {
+ auto* flag = GetContinueFlag(swtch, cont->Loop());
+ b.InsertBefore(cont, [&] {
+ b.Store(flag, true);
+ b.ExitSwitch(swtch);
+ });
+ cont->Destroy();
+ }
+
+ /// Get or create the flag used to indicate whether a `continue` was hit.
+ /// @param swtch the `switch` instruction to get the flag for
+ /// @param loop the `loop` that is the target of `continue` instruction in this switch
+ /// @returns the flag variable
+ Var* GetContinueFlag(Switch* swtch, Loop* loop) {
+ return continue_flag_for_switch.GetOrAdd(swtch, [&] {
+ // Declare the flag before the switch statement.
+ auto* flag = b.Var<function, bool>("tint_continue");
+ flag->InsertBefore(swtch);
+
+ // Check the flag after the `switch` instruction and `continue` if it was set.
+ b.InsertAfter(swtch, [&] {
+ auto* check = b.If(b.Load(flag));
+ b.Append(check->True(), [&] { //
+ b.Continue(loop);
+ });
+ });
+
+ return flag;
+ });
+ }
+};
+
+} // namespace
+
+Result<SuccessType> RemoveContinueInSwitch(Module& ir) {
+ auto result = ValidateAndDumpIfNeeded(ir, "RemoveContinueInSwitch transform",
+ core::ir::Capabilities{
+ core::ir::Capability::kAllowVectorElementPointer,
+ });
+ if (result != Success) {
+ return result;
+ }
+
+ State{ir}.Process();
+
+ return Success;
+}
+
+} // namespace tint::core::ir::transform
diff --git a/src/tint/lang/core/ir/transform/remove_continue_in_switch.h b/src/tint/lang/core/ir/transform/remove_continue_in_switch.h
new file mode 100644
index 0000000..ff12e61
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/remove_continue_in_switch.h
@@ -0,0 +1,52 @@
+// 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_CORE_IR_TRANSFORM_REMOVE_CONTINUE_IN_SWITCH_H_
+#define SRC_TINT_LANG_CORE_IR_TRANSFORM_REMOVE_CONTINUE_IN_SWITCH_H_
+
+#include "src/tint/utils/result/result.h"
+
+// Forward declarations.
+namespace tint::core::ir {
+class Module;
+}
+
+namespace tint::core::ir::transform {
+
+/// RemoveContinueInSwitch is a transform that removes `continue` instructions that are nested
+/// inside `switch` instructions, and instead sets a flag which is checked after the `switch`.
+///
+/// This is needed because textual backends inline loop continuing blocks at the location of each
+/// `continue`, which would lead to a `break` unintentionally breaking from the switch instead of
+/// the loop.
+/// @param module the module to transform
+/// @returns success or failure
+Result<SuccessType> RemoveContinueInSwitch(Module& module);
+
+} // namespace tint::core::ir::transform
+
+#endif // SRC_TINT_LANG_CORE_IR_TRANSFORM_REMOVE_CONTINUE_IN_SWITCH_H_
diff --git a/src/tint/lang/core/ir/transform/remove_continue_in_switch_test.cc b/src/tint/lang/core/ir/transform/remove_continue_in_switch_test.cc
new file mode 100644
index 0000000..02c47e2
--- /dev/null
+++ b/src/tint/lang/core/ir/transform/remove_continue_in_switch_test.cc
@@ -0,0 +1,614 @@
+// 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/core/ir/transform/remove_continue_in_switch.h"
+
+#include <utility>
+
+#include "src/tint/lang/core/ir/transform/helper_test.h"
+
+namespace tint::core::ir::transform {
+namespace {
+
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
+
+using IR_RemoveContinueInSwitchTest = TransformTest;
+
+TEST_F(IR_RemoveContinueInSwitchTest, NoModify_ContinueNotInSwitch) {
+ auto* func = b.Function("func", ty.void_());
+ b.Append(func->Block(), [&] { //
+ auto* loop = b.Loop();
+ b.Append(loop->Body(), [&] {
+ auto* swtch = b.Switch(42_i);
+ auto* def_case = b.DefaultCase(swtch);
+ b.Append(def_case, [&] { //
+ b.ExitSwitch(swtch);
+ });
+ b.Continue(loop);
+ });
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%func = func():void {
+ $B1: {
+ loop [b: $B2] { # loop_1
+ $B2: { # body
+ switch 42i [c: (default, $B3)] { # switch_1
+ $B3: { # case
+ exit_switch # switch_1
+ }
+ }
+ continue # -> $B4
+ }
+ }
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ Run(RemoveContinueInSwitch);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RemoveContinueInSwitchTest, NoModify_ContinueInSwitchCase_WithoutBreakIf) {
+ auto* func = b.Function("func", ty.void_());
+ b.Append(func->Block(), [&] { //
+ auto* loop = b.Loop();
+ b.Append(loop->Body(), [&] {
+ auto* swtch = b.Switch(42_i);
+ auto* def_case = b.DefaultCase(swtch);
+ b.Append(def_case, [&] { //
+ b.Continue(loop);
+ });
+ b.Continue(loop);
+ });
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%func = func():void {
+ $B1: {
+ loop [b: $B2] { # loop_1
+ $B2: { # body
+ switch 42i [c: (default, $B3)] { # switch_1
+ $B3: { # case
+ continue # -> $B4
+ }
+ }
+ continue # -> $B4
+ }
+ }
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ Run(RemoveContinueInSwitch);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RemoveContinueInSwitchTest, ContinueInSwitchCase_WithBreakIf) {
+ auto* func = b.Function("func", ty.void_());
+ b.Append(func->Block(), [&] { //
+ auto* loop = b.Loop();
+ b.Append(loop->Body(), [&] {
+ auto* swtch = b.Switch(42_i);
+ auto* def_case = b.DefaultCase(swtch);
+ b.Append(def_case, [&] { //
+ b.Continue(loop);
+ });
+ b.Continue(loop);
+ b.Append(loop->Continuing(), [&] { //
+ b.BreakIf(loop, true);
+ });
+ });
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%func = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ switch 42i [c: (default, $B4)] { # switch_1
+ $B4: { # case
+ continue # -> $B3
+ }
+ }
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%func = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ %tint_continue:ptr<function, bool, read_write> = var
+ switch 42i [c: (default, $B4)] { # switch_1
+ $B4: { # case
+ store %tint_continue, true
+ exit_switch # switch_1
+ }
+ }
+ %3:bool = load %tint_continue
+ if %3 [t: $B5] { # if_1
+ $B5: { # true
+ continue # -> $B3
+ }
+ }
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+
+ Run(RemoveContinueInSwitch);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RemoveContinueInSwitchTest, ContinueInMultipleCases) {
+ auto* func = b.Function("func", ty.void_());
+ b.Append(func->Block(), [&] { //
+ auto* loop = b.Loop();
+ b.Append(loop->Body(), [&] {
+ auto* swtch = b.Switch(42_i);
+ auto* case_a = b.Case(swtch, Vector{b.Constant(1_i)});
+ b.Append(case_a, [&] { //
+ b.Continue(loop);
+ });
+ auto* case_b = b.Case(swtch, Vector{b.Constant(2_i)});
+ b.Append(case_b, [&] { //
+ b.Continue(loop);
+ });
+ auto* def_case = b.DefaultCase(swtch);
+ b.Append(def_case, [&] { //
+ b.Continue(loop);
+ });
+ b.Continue(loop);
+ b.Append(loop->Continuing(), [&] { //
+ b.BreakIf(loop, true);
+ });
+ });
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%func = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ switch 42i [c: (1i, $B4), c: (2i, $B5), c: (default, $B6)] { # switch_1
+ $B4: { # case
+ continue # -> $B3
+ }
+ $B5: { # case
+ continue # -> $B3
+ }
+ $B6: { # case
+ continue # -> $B3
+ }
+ }
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%func = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ %tint_continue:ptr<function, bool, read_write> = var
+ switch 42i [c: (1i, $B4), c: (2i, $B5), c: (default, $B6)] { # switch_1
+ $B4: { # case
+ store %tint_continue, true
+ exit_switch # switch_1
+ }
+ $B5: { # case
+ store %tint_continue, true
+ exit_switch # switch_1
+ }
+ $B6: { # case
+ store %tint_continue, true
+ exit_switch # switch_1
+ }
+ }
+ %3:bool = load %tint_continue
+ if %3 [t: $B7] { # if_1
+ $B7: { # true
+ continue # -> $B3
+ }
+ }
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+
+ Run(RemoveContinueInSwitch);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RemoveContinueInSwitchTest, ContinueInMultipleSwitches) {
+ auto* func = b.Function("func", ty.void_());
+ b.Append(func->Block(), [&] { //
+ auto* loop = b.Loop();
+ b.Append(loop->Body(), [&] {
+ auto* swtch_a = b.Switch(42_i);
+ auto* def_case_a = b.DefaultCase(swtch_a);
+ b.Append(def_case_a, [&] { //
+ b.Continue(loop);
+ });
+ auto* swtch_b = b.Switch(43_i);
+ auto* def_case_b = b.DefaultCase(swtch_b);
+ b.Append(def_case_b, [&] { //
+ b.Continue(loop);
+ });
+ auto* swtch_c = b.Switch(44_i);
+ auto* def_case_c = b.DefaultCase(swtch_c);
+ b.Append(def_case_c, [&] { //
+ b.Continue(loop);
+ });
+ b.Continue(loop);
+ b.Append(loop->Continuing(), [&] { //
+ b.BreakIf(loop, true);
+ });
+ });
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%func = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ switch 42i [c: (default, $B4)] { # switch_1
+ $B4: { # case
+ continue # -> $B3
+ }
+ }
+ switch 43i [c: (default, $B5)] { # switch_2
+ $B5: { # case
+ continue # -> $B3
+ }
+ }
+ switch 44i [c: (default, $B6)] { # switch_3
+ $B6: { # case
+ continue # -> $B3
+ }
+ }
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%func = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ %tint_continue:ptr<function, bool, read_write> = var
+ switch 42i [c: (default, $B4)] { # switch_1
+ $B4: { # case
+ store %tint_continue, true
+ exit_switch # switch_1
+ }
+ }
+ %3:bool = load %tint_continue
+ if %3 [t: $B5] { # if_1
+ $B5: { # true
+ continue # -> $B3
+ }
+ }
+ %tint_continue_1:ptr<function, bool, read_write> = var # %tint_continue_1: 'tint_continue'
+ switch 43i [c: (default, $B6)] { # switch_2
+ $B6: { # case
+ store %tint_continue_1, true
+ exit_switch # switch_2
+ }
+ }
+ %5:bool = load %tint_continue_1
+ if %5 [t: $B7] { # if_2
+ $B7: { # true
+ continue # -> $B3
+ }
+ }
+ %tint_continue_2:ptr<function, bool, read_write> = var # %tint_continue_2: 'tint_continue'
+ switch 44i [c: (default, $B8)] { # switch_3
+ $B8: { # case
+ store %tint_continue_2, true
+ exit_switch # switch_3
+ }
+ }
+ %7:bool = load %tint_continue_2
+ if %7 [t: $B9] { # if_3
+ $B9: { # true
+ continue # -> $B3
+ }
+ }
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+
+ Run(RemoveContinueInSwitch);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RemoveContinueInSwitchTest, ContinueInSwitchCaseNestedInsideIf) {
+ auto* func = b.Function("func", ty.void_());
+ b.Append(func->Block(), [&] { //
+ auto* loop = b.Loop();
+ b.Append(loop->Body(), [&] {
+ auto* outer_if = b.If(true);
+ b.Append(outer_if->True(), [&] {
+ auto* swtch = b.Switch(42_i);
+ auto* def_case = b.DefaultCase(swtch);
+ b.Append(def_case, [&] { //
+ auto* inner_if = b.If(true);
+ b.Append(inner_if->True(), [&] { //
+ b.Continue(loop);
+ });
+ b.Unreachable();
+ });
+ b.Unreachable();
+ });
+ b.Continue(loop);
+ b.Append(loop->Continuing(), [&] { //
+ b.BreakIf(loop, true);
+ });
+ });
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%func = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ if true [t: $B4] { # if_1
+ $B4: { # true
+ switch 42i [c: (default, $B5)] { # switch_1
+ $B5: { # case
+ if true [t: $B6] { # if_2
+ $B6: { # true
+ continue # -> $B3
+ }
+ }
+ unreachable
+ }
+ }
+ unreachable
+ }
+ }
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%func = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ if true [t: $B4] { # if_1
+ $B4: { # true
+ %tint_continue:ptr<function, bool, read_write> = var
+ switch 42i [c: (default, $B5)] { # switch_1
+ $B5: { # case
+ if true [t: $B6] { # if_2
+ $B6: { # true
+ store %tint_continue, true
+ exit_switch # switch_1
+ }
+ }
+ unreachable
+ }
+ }
+ %3:bool = load %tint_continue
+ if %3 [t: $B7] { # if_3
+ $B7: { # true
+ continue # -> $B3
+ }
+ }
+ unreachable
+ }
+ }
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+
+ Run(RemoveContinueInSwitch);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_RemoveContinueInSwitchTest, ContinueInSwitchInsideAnotherSwitch) {
+ auto* func = b.Function("func", ty.void_());
+ b.Append(func->Block(), [&] { //
+ auto* loop = b.Loop();
+ b.Append(loop->Body(), [&] {
+ auto* outer_switch = b.Switch(42_i);
+ auto* outer_def_case = b.DefaultCase(outer_switch);
+ b.Append(outer_def_case, [&] { //
+ auto* inner_switch = b.Switch(42_i);
+ auto* inner_def_case = b.DefaultCase(inner_switch);
+ b.Append(inner_def_case, [&] { //
+ b.Continue(loop);
+ });
+ b.ExitSwitch(outer_switch);
+ });
+ b.Continue(loop);
+ b.Append(loop->Continuing(), [&] { //
+ b.BreakIf(loop, true);
+ });
+ });
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%func = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ switch 42i [c: (default, $B4)] { # switch_1
+ $B4: { # case
+ switch 42i [c: (default, $B5)] { # switch_2
+ $B5: { # case
+ continue # -> $B3
+ }
+ }
+ exit_switch # switch_1
+ }
+ }
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%func = func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ %tint_continue:ptr<function, bool, read_write> = var
+ switch 42i [c: (default, $B4)] { # switch_1
+ $B4: { # case
+ %tint_continue_1:ptr<function, bool, read_write> = var # %tint_continue_1: 'tint_continue'
+ switch 42i [c: (default, $B5)] { # switch_2
+ $B5: { # case
+ store %tint_continue_1, true
+ exit_switch # switch_2
+ }
+ }
+ %4:bool = load %tint_continue_1
+ if %4 [t: $B6] { # if_1
+ $B6: { # true
+ store %tint_continue, true
+ exit_switch # switch_1
+ }
+ }
+ exit_switch # switch_1
+ }
+ }
+ %5:bool = load %tint_continue
+ if %5 [t: $B7] { # if_2
+ $B7: { # true
+ continue # -> $B3
+ }
+ }
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)";
+
+ Run(RemoveContinueInSwitch);
+
+ EXPECT_EQ(expect, str());
+}
+
+} // namespace
+} // namespace tint::core::ir::transform
diff --git a/src/tint/lang/msl/writer/raise/raise.cc b/src/tint/lang/msl/writer/raise/raise.cc
index 8a2309f..1e51dad 100644
--- a/src/tint/lang/msl/writer/raise/raise.cc
+++ b/src/tint/lang/msl/writer/raise/raise.cc
@@ -38,6 +38,7 @@
#include "src/tint/lang/core/ir/transform/demote_to_helper.h"
#include "src/tint/lang/core/ir/transform/multiplanar_external_texture.h"
#include "src/tint/lang/core/ir/transform/preserve_padding.h"
+#include "src/tint/lang/core/ir/transform/remove_continue_in_switch.h"
#include "src/tint/lang/core/ir/transform/remove_terminator_args.h"
#include "src/tint/lang/core/ir/transform/rename_conflicts.h"
#include "src/tint/lang/core/ir/transform/robustness.h"
@@ -121,6 +122,7 @@
RUN_TRANSFORM(core::ir::transform::PreservePadding, module);
RUN_TRANSFORM(core::ir::transform::VectorizeScalarMatrixConstructors, module);
+ RUN_TRANSFORM(core::ir::transform::RemoveContinueInSwitch, module);
// DemoteToHelper must come before any transform that introduces non-core instructions.
RUN_TRANSFORM(core::ir::transform::DemoteToHelper, module);
diff --git a/test/tint/bug/tint/2039.wgsl.expected.ir.msl b/test/tint/bug/tint/2039.wgsl.expected.ir.msl
index 495d616..fd6830b 100644
--- a/test/tint/bug/tint/2039.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2039.wgsl.expected.ir.msl
@@ -5,19 +5,24 @@
uint out = 0u;
{
while(true) {
+ bool tint_continue = false;
switch(2) {
case 1:
{
- {
- if (true) { break; }
- }
- continue;
+ tint_continue = true;
+ break;
}
default:
{
break;
}
}
+ if (tint_continue) {
+ {
+ if (true) { break; }
+ }
+ continue;
+ }
out = (out + 1u);
{
if (true) { break; }
diff --git a/test/tint/loops/continue_in_switch_with_breakif.wgsl b/test/tint/loops/continue_in_switch_with_breakif.wgsl
new file mode 100644
index 0000000..0b65b9c
--- /dev/null
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl
@@ -0,0 +1,18 @@
+@compute @workgroup_size(1)
+fn f() {
+ var i : i32 = 0;
+ loop {
+ switch (i) {
+ case 0: {
+ continue;
+ }
+ default:{
+ break;
+ }
+ }
+ continuing {
+ i = i + 1;
+ break if i >= 4;
+ }
+ }
+}
diff --git a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.dxc.hlsl b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..3aebdc6
--- /dev/null
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.dxc.hlsl
@@ -0,0 +1,29 @@
+[numthreads(1, 1, 1)]
+void f() {
+ int i = 0;
+ bool tint_continue = false;
+ while (true) {
+ tint_continue = false;
+ switch(i) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ {
+ i = (i + 1);
+ if ((i >= 4)) { break; }
+ }
+ continue;
+ }
+ {
+ i = (i + 1);
+ if ((i >= 4)) { break; }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.fxc.hlsl b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..3aebdc6
--- /dev/null
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.fxc.hlsl
@@ -0,0 +1,29 @@
+[numthreads(1, 1, 1)]
+void f() {
+ int i = 0;
+ bool tint_continue = false;
+ while (true) {
+ tint_continue = false;
+ switch(i) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ {
+ i = (i + 1);
+ if ((i >= 4)) { break; }
+ }
+ continue;
+ }
+ {
+ i = (i + 1);
+ if ((i >= 4)) { break; }
+ }
+ }
+ return;
+}
diff --git a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.glsl b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.glsl
new file mode 100644
index 0000000..8de39f2
--- /dev/null
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.glsl
@@ -0,0 +1,35 @@
+#version 310 es
+
+void f() {
+ int i = 0;
+ bool tint_continue = false;
+ while (true) {
+ tint_continue = false;
+ switch(i) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ {
+ i = (i + 1);
+ if ((i >= 4)) { break; }
+ }
+ continue;
+ }
+ {
+ i = (i + 1);
+ if ((i >= 4)) { break; }
+ }
+ }
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.dxc.hlsl b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.dxc.hlsl
new file mode 100644
index 0000000..4edaaba
--- /dev/null
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,29 @@
+
+[numthreads(1, 1, 1)]
+void f() {
+ int i = 0;
+ {
+ while(true) {
+ switch(i) {
+ case 0:
+ {
+ {
+ i = (i + 1);
+ if ((i >= 4)) { break; }
+ }
+ continue;
+ }
+ default:
+ {
+ break;
+ }
+ }
+ {
+ i = (i + 1);
+ if ((i >= 4)) { break; }
+ }
+ continue;
+ }
+ }
+}
+
diff --git a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.fxc.hlsl b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..4edaaba
--- /dev/null
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,29 @@
+
+[numthreads(1, 1, 1)]
+void f() {
+ int i = 0;
+ {
+ while(true) {
+ switch(i) {
+ case 0:
+ {
+ {
+ i = (i + 1);
+ if ((i >= 4)) { break; }
+ }
+ continue;
+ }
+ default:
+ {
+ break;
+ }
+ }
+ {
+ i = (i + 1);
+ if ((i >= 4)) { break; }
+ }
+ continue;
+ }
+ }
+}
+
diff --git a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.glsl b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.glsl
new file mode 100644
index 0000000..c48695f
--- /dev/null
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.glsl
@@ -0,0 +1,11 @@
+SKIP: FAILED
+
+../../src/tint/lang/glsl/writer/printer/printer.cc:252 internal compiler error: Switch() matched no cases. Type: tint::core::ir::Loop
+********************************************************************
+* The tint shader compiler has encountered an unexpected error. *
+* *
+* Please help us fix this issue by submitting a bug report at *
+* crbug.com/tint with the source program that triggered the bug. *
+********************************************************************
+
+tint executable returned error: signal: trace/BPT trap
diff --git a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl
new file mode 100644
index 0000000..d73f2a3
--- /dev/null
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl
@@ -0,0 +1,34 @@
+#include <metal_stdlib>
+using namespace metal;
+
+kernel void f() {
+ int i = 0;
+ {
+ while(true) {
+ bool tint_continue = false;
+ switch(i) {
+ case 0:
+ {
+ tint_continue = true;
+ break;
+ }
+ default:
+ {
+ break;
+ }
+ }
+ if (tint_continue) {
+ {
+ i = (i + 1);
+ if ((i >= 4)) { break; }
+ }
+ continue;
+ }
+ {
+ i = (i + 1);
+ if ((i >= 4)) { break; }
+ }
+ continue;
+ }
+ }
+}
diff --git a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.msl b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.msl
new file mode 100644
index 0000000..e0ef3e1
--- /dev/null
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.msl
@@ -0,0 +1,37 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+ volatile bool VOLATILE_NAME = true; \
+ if (VOLATILE_NAME)
+
+kernel void f() {
+ int i = 0;
+ bool tint_continue = false;
+ TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+ tint_continue = false;
+ switch(i) {
+ case 0: {
+ tint_continue = true;
+ break;
+ }
+ default: {
+ break;
+ }
+ }
+ if (tint_continue) {
+ {
+ i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
+ if ((i >= 4)) { break; }
+ }
+ continue;
+ }
+ {
+ i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
+ if ((i >= 4)) { break; }
+ }
+ }
+ return;
+}
+
diff --git a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.spvasm b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.spvasm
new file mode 100644
index 0000000..055d49b
--- /dev/null
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.spvasm
@@ -0,0 +1,47 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 24
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %f "f"
+ OpExecutionMode %f LocalSize 1 1 1
+ OpName %f "f"
+ OpName %i "i"
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+%_ptr_Function_int = OpTypePointer Function %int
+ %int_0 = OpConstant %int 0
+ %int_1 = OpConstant %int 1
+ %int_4 = OpConstant %int 4
+ %bool = OpTypeBool
+ %f = OpFunction %void None %3
+ %4 = OpLabel
+ %i = OpVariable %_ptr_Function_int Function
+ OpStore %i %int_0
+ OpBranch %11
+ %11 = OpLabel
+ OpLoopMerge %12 %10 None
+ OpBranch %9
+ %9 = OpLabel
+ %13 = OpLoad %int %i None
+ OpSelectionMerge %16 None
+ OpSwitch %13 %14 0 %15
+ %15 = OpLabel
+ OpBranch %10
+ %14 = OpLabel
+ OpBranch %16
+ %16 = OpLabel
+ OpBranch %10
+ %10 = OpLabel
+ %17 = OpLoad %int %i None
+ %18 = OpIAdd %int %17 %int_1
+ OpStore %i %18 None
+ %20 = OpLoad %int %i None
+ %21 = OpSGreaterThanEqual %bool %20 %int_4
+ OpBranchConditional %21 %12 %11
+ %12 = OpLabel
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.wgsl b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.wgsl
new file mode 100644
index 0000000..1b15b09
--- /dev/null
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.wgsl
@@ -0,0 +1,19 @@
+@compute @workgroup_size(1)
+fn f() {
+ var i : i32 = 0;
+ loop {
+ switch(i) {
+ case 0: {
+ continue;
+ }
+ default: {
+ break;
+ }
+ }
+
+ continuing {
+ i = (i + 1);
+ break if (i >= 4);
+ }
+ }
+}
