[msl][glsl] Fix incorrect code gen for continuing inside switch.
In the GLSL and MSL backends when generating a continue inside a switch
block, the continuing statement would sink up into the switch label.
This means that a `break` for the continuing would break the switch
instead of the loop.
This CL moves the RemoveContinuingInSwitch transform from the HLSL
backend up to being a generic AST transform and calls it from the MSL
and GLSL backends to write the AST to no longer continue from switches.
Bug: tint:2039
Change-Id: Ica4a8d2de6bf97a179dac9e9216d8ddce99b5976
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/175820
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
index c83273f..c40bcbb 100644
--- a/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/glsl/writer/ast_printer/ast_printer.cc
@@ -72,6 +72,7 @@
 #include "src/tint/lang/wgsl/ast/transform/preserve_padding.h"
 #include "src/tint/lang/wgsl/ast/transform/promote_initializers_to_let.h"
 #include "src/tint/lang/wgsl/ast/transform/promote_side_effects_to_decl.h"
+#include "src/tint/lang/wgsl/ast/transform/remove_continue_in_switch.h"
 #include "src/tint/lang/wgsl/ast/transform/remove_phonies.h"
 #include "src/tint/lang/wgsl/ast/transform/renamer.h"
 #include "src/tint/lang/wgsl/ast/transform/robustness.h"
@@ -240,6 +241,7 @@
     manager.Add<ast::transform::BindingRemapper>();
 
     manager.Add<ast::transform::PromoteInitializersToLet>();
+    manager.Add<ast::transform::RemoveContinueInSwitch>();
     manager.Add<ast::transform::AddEmptyEntryPoint>();
 
     // Std140 must come after PromoteSideEffectsToDecl and before SimplifyPointers.
diff --git a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
index 6ab05d3..ab2551a 100644
--- a/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/hlsl/writer/ast_printer/ast_printer.cc
@@ -48,7 +48,6 @@
 #include "src/tint/lang/hlsl/writer/ast_raise/localize_struct_array_assignment.h"
 #include "src/tint/lang/hlsl/writer/ast_raise/num_workgroups_from_uniform.h"
 #include "src/tint/lang/hlsl/writer/ast_raise/pixel_local.h"
-#include "src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch.h"
 #include "src/tint/lang/hlsl/writer/ast_raise/truncate_interstage_variables.h"
 #include "src/tint/lang/hlsl/writer/common/option_helpers.h"
 #include "src/tint/lang/wgsl/ast/call_statement.h"
@@ -67,6 +66,7 @@
 #include "src/tint/lang/wgsl/ast/transform/multiplanar_external_texture.h"
 #include "src/tint/lang/wgsl/ast/transform/promote_initializers_to_let.h"
 #include "src/tint/lang/wgsl/ast/transform/promote_side_effects_to_decl.h"
+#include "src/tint/lang/wgsl/ast/transform/remove_continue_in_switch.h"
 #include "src/tint/lang/wgsl/ast/transform/remove_phonies.h"
 #include "src/tint/lang/wgsl/ast/transform/robustness.h"
 #include "src/tint/lang/wgsl/ast/transform/simplify_pointers.h"
@@ -362,7 +362,7 @@
     manager.Add<CalculateArrayLength>();
     manager.Add<ast::transform::PromoteInitializersToLet>();
 
-    manager.Add<RemoveContinueInSwitch>();
+    manager.Add<ast::transform::RemoveContinueInSwitch>();
 
     manager.Add<ast::transform::AddEmptyEntryPoint>();
 
diff --git a/src/tint/lang/hlsl/writer/ast_raise/BUILD.bazel b/src/tint/lang/hlsl/writer/ast_raise/BUILD.bazel
index dea6056..678aec4 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/BUILD.bazel
+++ b/src/tint/lang/hlsl/writer/ast_raise/BUILD.bazel
@@ -44,7 +44,6 @@
     "localize_struct_array_assignment.cc",
     "num_workgroups_from_uniform.cc",
     "pixel_local.cc",
-    "remove_continue_in_switch.cc",
     "truncate_interstage_variables.cc",
   ],
   hdrs = [
@@ -53,7 +52,6 @@
     "localize_struct_array_assignment.h",
     "num_workgroups_from_uniform.h",
     "pixel_local.h",
-    "remove_continue_in_switch.h",
     "truncate_interstage_variables.h",
   ],
   deps = [
@@ -95,7 +93,6 @@
     "localize_struct_array_assignment_test.cc",
     "num_workgroups_from_uniform_test.cc",
     "pixel_local_test.cc",
-    "remove_continue_in_switch_test.cc",
     "truncate_interstage_variables_test.cc",
   ],
   deps = [
diff --git a/src/tint/lang/hlsl/writer/ast_raise/BUILD.cmake b/src/tint/lang/hlsl/writer/ast_raise/BUILD.cmake
index 71a814c..49f9147 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/BUILD.cmake
+++ b/src/tint/lang/hlsl/writer/ast_raise/BUILD.cmake
@@ -51,8 +51,6 @@
   lang/hlsl/writer/ast_raise/num_workgroups_from_uniform.h
   lang/hlsl/writer/ast_raise/pixel_local.cc
   lang/hlsl/writer/ast_raise/pixel_local.h
-  lang/hlsl/writer/ast_raise/remove_continue_in_switch.cc
-  lang/hlsl/writer/ast_raise/remove_continue_in_switch.h
   lang/hlsl/writer/ast_raise/truncate_interstage_variables.cc
   lang/hlsl/writer/ast_raise/truncate_interstage_variables.h
 )
@@ -98,7 +96,6 @@
   lang/hlsl/writer/ast_raise/localize_struct_array_assignment_test.cc
   lang/hlsl/writer/ast_raise/num_workgroups_from_uniform_test.cc
   lang/hlsl/writer/ast_raise/pixel_local_test.cc
-  lang/hlsl/writer/ast_raise/remove_continue_in_switch_test.cc
   lang/hlsl/writer/ast_raise/truncate_interstage_variables_test.cc
 )
 
diff --git a/src/tint/lang/hlsl/writer/ast_raise/BUILD.gn b/src/tint/lang/hlsl/writer/ast_raise/BUILD.gn
index 8ac7972..1fdadd9 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/BUILD.gn
+++ b/src/tint/lang/hlsl/writer/ast_raise/BUILD.gn
@@ -54,8 +54,6 @@
       "num_workgroups_from_uniform.h",
       "pixel_local.cc",
       "pixel_local.h",
-      "remove_continue_in_switch.cc",
-      "remove_continue_in_switch.h",
       "truncate_interstage_variables.cc",
       "truncate_interstage_variables.h",
     ]
@@ -98,7 +96,6 @@
         "localize_struct_array_assignment_test.cc",
         "num_workgroups_from_uniform_test.cc",
         "pixel_local_test.cc",
-        "remove_continue_in_switch_test.cc",
         "truncate_interstage_variables_test.cc",
       ]
       deps = [
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
index 03a3ec6..5a3a574 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -76,6 +76,7 @@
 #include "src/tint/lang/wgsl/ast/transform/preserve_padding.h"
 #include "src/tint/lang/wgsl/ast/transform/promote_initializers_to_let.h"
 #include "src/tint/lang/wgsl/ast/transform/promote_side_effects_to_decl.h"
+#include "src/tint/lang/wgsl/ast/transform/remove_continue_in_switch.h"
 #include "src/tint/lang/wgsl/ast/transform/remove_phonies.h"
 #include "src/tint/lang/wgsl/ast/transform/robustness.h"
 #include "src/tint/lang/wgsl/ast/transform/simplify_pointers.h"
@@ -223,6 +224,7 @@
     data.Add<ast::transform::CanonicalizeEntryPointIO::Config>(std::move(entry_point_io_cfg));
 
     manager.Add<ast::transform::PromoteInitializersToLet>();
+    manager.Add<ast::transform::RemoveContinueInSwitch>();
 
     // DemoteToHelper must come after PromoteSideEffectsToDecl and ExpandCompoundAssignment.
     // TODO(crbug.com/tint/1752): This is only necessary for Metal versions older than 2.3.
diff --git a/src/tint/lang/wgsl/ast/transform/BUILD.bazel b/src/tint/lang/wgsl/ast/transform/BUILD.bazel
index 576a31a..8f4e427 100644
--- a/src/tint/lang/wgsl/ast/transform/BUILD.bazel
+++ b/src/tint/lang/wgsl/ast/transform/BUILD.bazel
@@ -61,6 +61,7 @@
     "promote_initializers_to_let.cc",
     "promote_side_effects_to_decl.cc",
     "push_constant_helper.cc",
+    "remove_continue_in_switch.cc",
     "remove_phonies.cc",
     "remove_unreachable_statements.cc",
     "renamer.cc",
@@ -98,6 +99,7 @@
     "promote_initializers_to_let.h",
     "promote_side_effects_to_decl.h",
     "push_constant_helper.h",
+    "remove_continue_in_switch.h",
     "remove_phonies.h",
     "remove_unreachable_statements.h",
     "renamer.h",
@@ -168,6 +170,7 @@
     "promote_initializers_to_let_test.cc",
     "promote_side_effects_to_decl_test.cc",
     "push_constant_helper_test.cc",
+    "remove_continue_in_switch_test.cc",
     "remove_phonies_test.cc",
     "remove_unreachable_statements_test.cc",
     "renamer_test.cc",
diff --git a/src/tint/lang/wgsl/ast/transform/BUILD.cmake b/src/tint/lang/wgsl/ast/transform/BUILD.cmake
index 84dc975..dea32eb 100644
--- a/src/tint/lang/wgsl/ast/transform/BUILD.cmake
+++ b/src/tint/lang/wgsl/ast/transform/BUILD.cmake
@@ -83,6 +83,8 @@
   lang/wgsl/ast/transform/promote_side_effects_to_decl.h
   lang/wgsl/ast/transform/push_constant_helper.cc
   lang/wgsl/ast/transform/push_constant_helper.h
+  lang/wgsl/ast/transform/remove_continue_in_switch.cc
+  lang/wgsl/ast/transform/remove_continue_in_switch.h
   lang/wgsl/ast/transform/remove_phonies.cc
   lang/wgsl/ast/transform/remove_phonies.h
   lang/wgsl/ast/transform/remove_unreachable_statements.cc
@@ -168,6 +170,7 @@
   lang/wgsl/ast/transform/promote_initializers_to_let_test.cc
   lang/wgsl/ast/transform/promote_side_effects_to_decl_test.cc
   lang/wgsl/ast/transform/push_constant_helper_test.cc
+  lang/wgsl/ast/transform/remove_continue_in_switch_test.cc
   lang/wgsl/ast/transform/remove_phonies_test.cc
   lang/wgsl/ast/transform/remove_unreachable_statements_test.cc
   lang/wgsl/ast/transform/renamer_test.cc
diff --git a/src/tint/lang/wgsl/ast/transform/BUILD.gn b/src/tint/lang/wgsl/ast/transform/BUILD.gn
index 6624981..70f6014 100644
--- a/src/tint/lang/wgsl/ast/transform/BUILD.gn
+++ b/src/tint/lang/wgsl/ast/transform/BUILD.gn
@@ -88,6 +88,8 @@
     "promote_side_effects_to_decl.h",
     "push_constant_helper.cc",
     "push_constant_helper.h",
+    "remove_continue_in_switch.cc",
+    "remove_continue_in_switch.h",
     "remove_phonies.cc",
     "remove_phonies.h",
     "remove_unreachable_statements.cc",
@@ -169,6 +171,7 @@
         "promote_initializers_to_let_test.cc",
         "promote_side_effects_to_decl_test.cc",
         "push_constant_helper_test.cc",
+        "remove_continue_in_switch_test.cc",
         "remove_phonies_test.cc",
         "remove_unreachable_statements_test.cc",
         "renamer_test.cc",
diff --git a/src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch.cc b/src/tint/lang/wgsl/ast/transform/remove_continue_in_switch.cc
similarity index 97%
rename from src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch.cc
rename to src/tint/lang/wgsl/ast/transform/remove_continue_in_switch.cc
index 899ea10..1ee9333 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch.cc
+++ b/src/tint/lang/wgsl/ast/transform/remove_continue_in_switch.cc
@@ -25,7 +25,7 @@
 // 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/hlsl/writer/ast_raise/remove_continue_in_switch.h"
+#include "src/tint/lang/wgsl/ast/transform/remove_continue_in_switch.h"
 
 #include <string>
 #include <utility>
@@ -42,9 +42,9 @@
 #include "src/tint/utils/containers/hashmap.h"
 #include "src/tint/utils/containers/hashset.h"
 
-TINT_INSTANTIATE_TYPEINFO(tint::hlsl::writer::RemoveContinueInSwitch);
+TINT_INSTANTIATE_TYPEINFO(tint::ast::transform::RemoveContinueInSwitch);
 
-namespace tint::hlsl::writer {
+namespace tint::ast::transform {
 
 /// PIMPL state for the transform
 struct RemoveContinueInSwitch::State {
@@ -197,4 +197,4 @@
     return state.Run();
 }
 
-}  // namespace tint::hlsl::writer
+}  // namespace tint::ast::transform
diff --git a/src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch.h b/src/tint/lang/wgsl/ast/transform/remove_continue_in_switch.h
similarity index 81%
rename from src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch.h
rename to src/tint/lang/wgsl/ast/transform/remove_continue_in_switch.h
index e29b0ee..6c62826 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch.h
+++ b/src/tint/lang/wgsl/ast/transform/remove_continue_in_switch.h
@@ -25,17 +25,18 @@
 // 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_HLSL_WRITER_AST_RAISE_REMOVE_CONTINUE_IN_SWITCH_H_
-#define SRC_TINT_LANG_HLSL_WRITER_AST_RAISE_REMOVE_CONTINUE_IN_SWITCH_H_
+#ifndef SRC_TINT_LANG_WGSL_AST_TRANSFORM_REMOVE_CONTINUE_IN_SWITCH_H_
+#define SRC_TINT_LANG_WGSL_AST_TRANSFORM_REMOVE_CONTINUE_IN_SWITCH_H_
 
 #include "src/tint/lang/wgsl/ast/transform/transform.h"
 
-namespace tint::hlsl::writer {
+namespace tint::ast::transform {
 
 /// This transform replaces continue statements in switch cases with setting a
 /// bool variable, and checking if the variable is set after the switch to
-/// continue. It is necessary to work around FXC "error X3708: continue cannot
-/// be used in a switch". See crbug.com/tint/1080.
+/// continue. It is necessary to work around various issues including:
+///  * FXC "error X3708: continue cannot be used in a switch". See crbug.com/tint/1080.
+///  * MSL and GLSL invalid code-gen. See crbug.com/tint/2039.
 class RemoveContinueInSwitch final
     : public Castable<RemoveContinueInSwitch, ast::transform::Transform> {
   public:
@@ -54,6 +55,6 @@
     struct State;
 };
 
-}  // namespace tint::hlsl::writer
+}  // namespace tint::ast::transform
 
-#endif  // SRC_TINT_LANG_HLSL_WRITER_AST_RAISE_REMOVE_CONTINUE_IN_SWITCH_H_
+#endif  // SRC_TINT_LANG_WGSL_AST_TRANSFORM_REMOVE_CONTINUE_IN_SWITCH_H_
diff --git a/src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch_test.cc b/src/tint/lang/wgsl/ast/transform/remove_continue_in_switch_test.cc
similarity index 98%
rename from src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch_test.cc
rename to src/tint/lang/wgsl/ast/transform/remove_continue_in_switch_test.cc
index e20839f..6f36aab 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/remove_continue_in_switch_test.cc
+++ b/src/tint/lang/wgsl/ast/transform/remove_continue_in_switch_test.cc
@@ -25,10 +25,10 @@
 // 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/hlsl/writer/ast_raise/remove_continue_in_switch.h"
+#include "src/tint/lang/wgsl/ast/transform/remove_continue_in_switch.h"
 #include "src/tint/lang/wgsl/ast/transform/helper_test.h"
 
-namespace tint::hlsl::writer {
+namespace tint::ast::transform {
 namespace {
 
 using RemoveContinueInSwitchTest = ast::transform::TransformTest;
@@ -767,4 +767,4 @@
 }
 
 }  // namespace
-}  // namespace tint::hlsl::writer
+}  // namespace tint::ast::transform
diff --git a/test/tint/bug/tint/2039.wgsl b/test/tint/bug/tint/2039.wgsl
new file mode 100644
index 0000000..419edb0
--- /dev/null
+++ b/test/tint/bug/tint/2039.wgsl
@@ -0,0 +1,17 @@
+
+@compute @workgroup_size(1,1,1)
+fn main() {
+  var out = 0u;
+  loop {
+    switch 2 {
+      case 1 {
+        continue;
+      }
+      default {}
+    }
+    out += 1u;
+    continuing {
+      break if true;
+    }
+  }
+}
diff --git a/test/tint/bug/tint/2039.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/2039.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..2707ef2
--- /dev/null
+++ b/test/tint/bug/tint/2039.wgsl.expected.dxc.hlsl
@@ -0,0 +1,28 @@
+[numthreads(1, 1, 1)]
+void main() {
+  uint tint_symbol = 0u;
+  bool tint_continue = false;
+  while (true) {
+    tint_continue = false;
+    switch(2) {
+      case 1: {
+        tint_continue = true;
+        break;
+      }
+      default: {
+        break;
+      }
+    }
+    if (tint_continue) {
+      {
+        if (true) { break; }
+      }
+      continue;
+    }
+    tint_symbol = (tint_symbol + 1u);
+    {
+      if (true) { break; }
+    }
+  }
+  return;
+}
diff --git a/test/tint/bug/tint/2039.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/2039.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..2707ef2
--- /dev/null
+++ b/test/tint/bug/tint/2039.wgsl.expected.fxc.hlsl
@@ -0,0 +1,28 @@
+[numthreads(1, 1, 1)]
+void main() {
+  uint tint_symbol = 0u;
+  bool tint_continue = false;
+  while (true) {
+    tint_continue = false;
+    switch(2) {
+      case 1: {
+        tint_continue = true;
+        break;
+      }
+      default: {
+        break;
+      }
+    }
+    if (tint_continue) {
+      {
+        if (true) { break; }
+      }
+      continue;
+    }
+    tint_symbol = (tint_symbol + 1u);
+    {
+      if (true) { break; }
+    }
+  }
+  return;
+}
diff --git a/test/tint/bug/tint/2039.wgsl.expected.glsl b/test/tint/bug/tint/2039.wgsl.expected.glsl
new file mode 100644
index 0000000..41aa478
--- /dev/null
+++ b/test/tint/bug/tint/2039.wgsl.expected.glsl
@@ -0,0 +1,34 @@
+#version 310 es
+
+void tint_symbol() {
+  uint tint_symbol_1 = 0u;
+  bool tint_continue = false;
+  while (true) {
+    tint_continue = false;
+    switch(2) {
+      case 1: {
+        tint_continue = true;
+        break;
+      }
+      default: {
+        break;
+      }
+    }
+    if (tint_continue) {
+      {
+        if (true) { break; }
+      }
+      continue;
+    }
+    tint_symbol_1 = (tint_symbol_1 + 1u);
+    {
+      if (true) { break; }
+    }
+  }
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}
diff --git a/test/tint/bug/tint/2039.wgsl.expected.msl b/test/tint/bug/tint/2039.wgsl.expected.msl
new file mode 100644
index 0000000..c756acb
--- /dev/null
+++ b/test/tint/bug/tint/2039.wgsl.expected.msl
@@ -0,0 +1,36 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+  volatile bool VOLATILE_NAME = true; \
+  if (VOLATILE_NAME)
+
+kernel void tint_symbol() {
+  uint out = 0u;
+  bool tint_continue = false;
+  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    tint_continue = false;
+    switch(2) {
+      case 1: {
+        tint_continue = true;
+        break;
+      }
+      default: {
+        break;
+      }
+    }
+    if (tint_continue) {
+      {
+        if (true) { break; }
+      }
+      continue;
+    }
+    out = (out + 1u);
+    {
+      if (true) { break; }
+    }
+  }
+  return;
+}
+
diff --git a/test/tint/bug/tint/2039.wgsl.expected.spvasm b/test/tint/bug/tint/2039.wgsl.expected.spvasm
new file mode 100644
index 0000000..ce5ae5e
--- /dev/null
+++ b/test/tint/bug/tint/2039.wgsl.expected.spvasm
@@ -0,0 +1,46 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 23
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %main "main"
+               OpName %out "out"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+          %6 = OpConstantNull %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+        %int = OpTypeInt 32 1
+      %int_2 = OpConstant %int 2
+     %uint_1 = OpConstant %uint 1
+       %bool = OpTypeBool
+       %true = OpConstantTrue %bool
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+        %out = OpVariable %_ptr_Function_uint Function %6
+               OpStore %out %6
+               OpBranch %9
+          %9 = OpLabel
+               OpLoopMerge %10 %11 None
+               OpBranch %12
+         %12 = OpLabel
+               OpSelectionMerge %13 None
+               OpSwitch %int_2 %16 1 %17
+         %17 = OpLabel
+               OpBranch %11
+         %16 = OpLabel
+               OpBranch %13
+         %13 = OpLabel
+         %18 = OpLoad %uint %out
+         %20 = OpIAdd %uint %18 %uint_1
+               OpStore %out %20
+               OpBranch %11
+         %11 = OpLabel
+               OpBranchConditional %true %10 %9
+         %10 = OpLabel
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/bug/tint/2039.wgsl.expected.wgsl b/test/tint/bug/tint/2039.wgsl.expected.wgsl
new file mode 100644
index 0000000..0bf5ccd
--- /dev/null
+++ b/test/tint/bug/tint/2039.wgsl.expected.wgsl
@@ -0,0 +1,18 @@
+@compute @workgroup_size(1, 1, 1)
+fn main() {
+  var out = 0u;
+  loop {
+    switch(2) {
+      case 1: {
+        continue;
+      }
+      default: {
+      }
+    }
+    out += 1u;
+
+    continuing {
+      break if true;
+    }
+  }
+}
diff --git a/test/tint/loops/continue_in_switch.wgsl.expected.glsl b/test/tint/loops/continue_in_switch.wgsl.expected.glsl
index 9678a12..291f4cd 100644
--- a/test/tint/loops/continue_in_switch.wgsl.expected.glsl
+++ b/test/tint/loops/continue_in_switch.wgsl.expected.glsl
@@ -1,17 +1,22 @@
 #version 310 es
 
 void f() {
+  bool tint_continue = false;
   {
     for(int i = 0; (i < 4); i = (i + 1)) {
+      tint_continue = false;
       switch(i) {
         case 0: {
-          continue;
+          tint_continue = true;
           break;
         }
         default: {
           break;
         }
       }
+      if (tint_continue) {
+        continue;
+      }
     }
   }
 }
diff --git a/test/tint/loops/continue_in_switch.wgsl.expected.msl b/test/tint/loops/continue_in_switch.wgsl.expected.msl
index 890daca..e7349b0c 100644
--- a/test/tint/loops/continue_in_switch.wgsl.expected.msl
+++ b/test/tint/loops/continue_in_switch.wgsl.expected.msl
@@ -7,16 +7,21 @@
   if (VOLATILE_NAME)
 
 kernel void f() {
+  bool tint_continue = false;
   TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    tint_continue = false;
     switch(i) {
       case 0: {
-        continue;
+        tint_continue = true;
         break;
       }
       default: {
         break;
       }
     }
+    if (tint_continue) {
+      continue;
+    }
   }
   return;
 }
diff --git a/test/tint/loops/multiple_continues.wgsl.expected.glsl b/test/tint/loops/multiple_continues.wgsl.expected.glsl
index 36c86d8..bbdf58c 100644
--- a/test/tint/loops/multiple_continues.wgsl.expected.glsl
+++ b/test/tint/loops/multiple_continues.wgsl.expected.glsl
@@ -1,25 +1,30 @@
 #version 310 es
 
 void tint_symbol() {
+  bool tint_continue = false;
   {
     for(int i = 0; (i < 2); i = (i + 1)) {
+      tint_continue = false;
       switch(i) {
         case 0: {
-          continue;
+          tint_continue = true;
           break;
         }
         case 1: {
-          continue;
+          tint_continue = true;
           break;
         }
         case 2: {
-          continue;
+          tint_continue = true;
           break;
         }
         default: {
           break;
         }
       }
+      if (tint_continue) {
+        continue;
+      }
     }
   }
 }
diff --git a/test/tint/loops/multiple_continues.wgsl.expected.msl b/test/tint/loops/multiple_continues.wgsl.expected.msl
index f8aaba6..76e84d7 100644
--- a/test/tint/loops/multiple_continues.wgsl.expected.msl
+++ b/test/tint/loops/multiple_continues.wgsl.expected.msl
@@ -7,24 +7,29 @@
   if (VOLATILE_NAME)
 
 kernel void tint_symbol() {
+  bool tint_continue = false;
   TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    tint_continue = false;
     switch(i) {
       case 0: {
-        continue;
+        tint_continue = true;
         break;
       }
       case 1: {
-        continue;
+        tint_continue = true;
         break;
       }
       case 2: {
-        continue;
+        tint_continue = true;
         break;
       }
       default: {
         break;
       }
     }
+    if (tint_continue) {
+      continue;
+    }
   }
   return;
 }
diff --git a/test/tint/loops/multiple_switch.wgsl.expected.glsl b/test/tint/loops/multiple_switch.wgsl.expected.glsl
index 290f1e3..82abd76 100644
--- a/test/tint/loops/multiple_switch.wgsl.expected.glsl
+++ b/test/tint/loops/multiple_switch.wgsl.expected.glsl
@@ -2,26 +2,34 @@
 
 void tint_symbol() {
   int i = 0;
+  bool tint_continue = false;
   {
     for(int i_1 = 0; (i_1 < 2); i_1 = (i_1 + 1)) {
+      tint_continue = false;
       switch(i_1) {
         case 0: {
-          continue;
+          tint_continue = true;
           break;
         }
         default: {
           break;
         }
       }
+      if (tint_continue) {
+        continue;
+      }
       switch(i_1) {
         case 0: {
-          continue;
+          tint_continue = true;
           break;
         }
         default: {
           break;
         }
       }
+      if (tint_continue) {
+        continue;
+      }
     }
   }
 }
diff --git a/test/tint/loops/multiple_switch.wgsl.expected.msl b/test/tint/loops/multiple_switch.wgsl.expected.msl
index 0bad249..7728952 100644
--- a/test/tint/loops/multiple_switch.wgsl.expected.msl
+++ b/test/tint/loops/multiple_switch.wgsl.expected.msl
@@ -8,25 +8,33 @@
 
 kernel void tint_symbol() {
   int i = 0;
+  bool tint_continue = false;
   TINT_ISOLATE_UB(tint_volatile_true) for(int i_1 = 0; (i_1 < 2); i_1 = as_type<int>((as_type<uint>(i_1) + as_type<uint>(1)))) {
+    tint_continue = false;
     switch(i_1) {
       case 0: {
-        continue;
+        tint_continue = true;
         break;
       }
       default: {
         break;
       }
     }
+    if (tint_continue) {
+      continue;
+    }
     switch(i_1) {
       case 0: {
-        continue;
+        tint_continue = true;
         break;
       }
       default: {
         break;
       }
     }
+    if (tint_continue) {
+      continue;
+    }
   }
   return;
 }
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.glsl b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.glsl
index a3ad0a3..63e7289 100644
--- a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.glsl
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.glsl
@@ -3,17 +3,22 @@
 void tint_symbol() {
   {
     for(int i = 0; (i < 2); i = (i + 2)) {
+      bool tint_continue = false;
       {
         for(int j = 0; (j < 2); j = (j + 2)) {
+          tint_continue = false;
           switch(i) {
             case 0: {
-              continue;
+              tint_continue = true;
               break;
             }
             default: {
               break;
             }
           }
+          if (tint_continue) {
+            continue;
+          }
         }
       }
     }
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
index 8940027..eb32ea2 100644
--- a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.msl
@@ -8,16 +8,21 @@
 
 kernel void tint_symbol() {
   TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+    bool tint_continue = false;
     TINT_ISOLATE_UB(tint_volatile_true_1) for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+      tint_continue = false;
       switch(i) {
         case 0: {
-          continue;
+          tint_continue = true;
           break;
         }
         default: {
           break;
         }
       }
+      if (tint_continue) {
+        continue;
+      }
     }
   }
   return;
diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.glsl b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.glsl
index 1e45661..06552d5 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.glsl
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.glsl
@@ -1,30 +1,40 @@
 #version 310 es
 
 void tint_symbol() {
+  bool tint_continue_1 = false;
   {
     for(int i = 0; (i < 2); i = (i + 2)) {
+      tint_continue_1 = false;
       switch(i) {
         case 0: {
+          bool tint_continue = false;
           {
             for(int j = 0; (j < 2); j = (j + 2)) {
+              tint_continue = false;
               switch(j) {
                 case 0: {
-                  continue;
+                  tint_continue = true;
                   break;
                 }
                 default: {
                   break;
                 }
               }
+              if (tint_continue) {
+                continue;
+              }
             }
           }
-          continue;
+          tint_continue_1 = true;
           break;
         }
         default: {
           break;
         }
       }
+      if (tint_continue_1) {
+        continue;
+      }
     }
   }
 }
diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
index 2662928..9ac0463 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.msl
@@ -7,27 +7,37 @@
   if (VOLATILE_NAME)
 
 kernel void tint_symbol() {
+  bool tint_continue_1 = false;
   TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+    tint_continue_1 = false;
     switch(i) {
       case 0: {
+        bool tint_continue = false;
         TINT_ISOLATE_UB(tint_volatile_true_1) for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+          tint_continue = false;
           switch(j) {
             case 0: {
-              continue;
+              tint_continue = true;
               break;
             }
             default: {
               break;
             }
           }
+          if (tint_continue) {
+            continue;
+          }
         }
-        continue;
+        tint_continue_1 = true;
         break;
       }
       default: {
         break;
       }
     }
+    if (tint_continue_1) {
+      continue;
+    }
   }
   return;
 }
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.glsl b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.glsl
index 3c59535..2b1cc90 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.glsl
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.glsl
@@ -2,42 +2,55 @@
 
 void tint_symbol() {
   int k = 0;
+  bool tint_continue_1 = false;
   {
     for(int i = 0; (i < 2); i = (i + 2)) {
+      tint_continue_1 = false;
       switch(i) {
         case 0: {
+          bool tint_continue = false;
           {
             for(int j = 0; (j < 2); j = (j + 2)) {
+              tint_continue = false;
               switch(j) {
                 case 0: {
-                  continue;
+                  tint_continue = true;
                   break;
                 }
                 case 1: {
                   switch(k) {
                     case 0: {
-                      continue;
+                      tint_continue = true;
                       break;
                     }
                     default: {
                       break;
                     }
                   }
+                  if (tint_continue) {
+                    break;
+                  }
                   break;
                 }
                 default: {
                   break;
                 }
               }
+              if (tint_continue) {
+                continue;
+              }
             }
           }
-          continue;
+          tint_continue_1 = true;
           break;
         }
         default: {
           break;
         }
       }
+      if (tint_continue_1) {
+        continue;
+      }
     }
   }
 }
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl
index 47f158b..c3b1e93 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.msl
@@ -8,39 +8,52 @@
 
 kernel void tint_symbol() {
   int k = 0;
+  bool tint_continue_1 = false;
   TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+    tint_continue_1 = false;
     switch(i) {
       case 0: {
+        bool tint_continue = false;
         TINT_ISOLATE_UB(tint_volatile_true_1) for(int j = 0; (j < 2); j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)))) {
+          tint_continue = false;
           switch(j) {
             case 0: {
-              continue;
+              tint_continue = true;
               break;
             }
             case 1: {
               switch(k) {
                 case 0: {
-                  continue;
+                  tint_continue = true;
                   break;
                 }
                 default: {
                   break;
                 }
               }
+              if (tint_continue) {
+                break;
+              }
               break;
             }
             default: {
               break;
             }
           }
+          if (tint_continue) {
+            continue;
+          }
         }
-        continue;
+        tint_continue_1 = true;
         break;
       }
       default: {
         break;
       }
     }
+    if (tint_continue_1) {
+      continue;
+    }
   }
   return;
 }
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.glsl b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.glsl
index ec32b25..f7b3dd6 100644
--- a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.glsl
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.glsl
@@ -2,25 +2,33 @@
 
 void tint_symbol() {
   int j = 0;
+  bool tint_continue = false;
   {
     for(int i = 0; (i < 2); i = (i + 2)) {
+      tint_continue = false;
       switch(i) {
         case 0: {
           switch(j) {
             case 0: {
-              continue;
+              tint_continue = true;
               break;
             }
             default: {
               break;
             }
           }
+          if (tint_continue) {
+            break;
+          }
           break;
         }
         default: {
           break;
         }
       }
+      if (tint_continue) {
+        continue;
+      }
     }
   }
 }
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
index 0f78ee3..e122be4 100644
--- a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.msl
@@ -8,24 +8,32 @@
 
 kernel void tint_symbol() {
   int j = 0;
+  bool tint_continue = false;
   TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)))) {
+    tint_continue = false;
     switch(i) {
       case 0: {
         switch(j) {
           case 0: {
-            continue;
+            tint_continue = true;
             break;
           }
           default: {
             break;
           }
         }
+        if (tint_continue) {
+          break;
+        }
         break;
       }
       default: {
         break;
       }
     }
+    if (tint_continue) {
+      continue;
+    }
   }
   return;
 }
diff --git a/test/tint/loops/single_continue.wgsl.expected.glsl b/test/tint/loops/single_continue.wgsl.expected.glsl
index 79afd54..58309a6 100644
--- a/test/tint/loops/single_continue.wgsl.expected.glsl
+++ b/test/tint/loops/single_continue.wgsl.expected.glsl
@@ -1,17 +1,22 @@
 #version 310 es
 
 void tint_symbol() {
+  bool tint_continue = false;
   {
     for(int i = 0; (i < 2); i = (i + 1)) {
+      tint_continue = false;
       switch(i) {
         case 0: {
-          continue;
+          tint_continue = true;
           break;
         }
         default: {
           break;
         }
       }
+      if (tint_continue) {
+        continue;
+      }
     }
   }
 }
diff --git a/test/tint/loops/single_continue.wgsl.expected.msl b/test/tint/loops/single_continue.wgsl.expected.msl
index e223c34..113c92c 100644
--- a/test/tint/loops/single_continue.wgsl.expected.msl
+++ b/test/tint/loops/single_continue.wgsl.expected.msl
@@ -7,16 +7,21 @@
   if (VOLATILE_NAME)
 
 kernel void tint_symbol() {
+  bool tint_continue = false;
   TINT_ISOLATE_UB(tint_volatile_true) for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    tint_continue = false;
     switch(i) {
       case 0: {
-        continue;
+        tint_continue = true;
         break;
       }
       default: {
         break;
       }
     }
+    if (tint_continue) {
+      continue;
+    }
   }
   return;
 }