[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); + } + } +}