[msl] Emit continuing inside a new block

The scoping rules used by RenameConflicts assume that the continuing
block is emitted as its own block, not inline inside the body.

Bug: 42251016
Change-Id: Ib9a587cadd5384eed4543847878c417a4e7901a4
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/191421
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index f3c0d79..996010e 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -591,7 +591,14 @@
         //   }
         // }
 
-        auto emit_continuing = [&] { EmitBlock(l->Continuing()); };
+        auto emit_continuing = [&] {
+            Line() << "{";
+            {
+                const ScopedIndent si(current_buffer_);
+                EmitBlock(l->Continuing());
+            }
+            Line() << "}";
+        };
         TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
 
         Line() << "{";
diff --git a/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl
index d2c6279..a9a21ac 100644
--- a/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_function_var.wgsl.expected.ir.msl
@@ -59,7 +59,9 @@
         break;
       }
       (*tint_module_vars.src_workgroup)[v_1] = int4(0);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl
index 65a63e9..2d69782 100644
--- a/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_private_var.wgsl.expected.ir.msl
@@ -59,7 +59,9 @@
         break;
       }
       (*tint_module_vars.src_workgroup)[v_1] = int4(0);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl
index 8314e8a..ae88516 100644
--- a/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_storage_var.wgsl.expected.ir.msl
@@ -62,7 +62,9 @@
         break;
       }
       (*tint_module_vars.src_workgroup)[v_1] = int4(0);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl b/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl
index 76ebf03..4072fb7d 100644
--- a/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_workgroup_var.wgsl.expected.ir.msl
@@ -62,7 +62,9 @@
       }
       (*tint_module_vars.src_workgroup)[v_1] = int4(0);
       (*tint_module_vars.dst)[v_1] = int4(0);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
@@ -75,7 +77,9 @@
         break;
       }
       (*tint_module_vars.dst_nested)[(v_3 / 6u)][((v_3 / 2u) % 3u)][(v_3 % 2u)] = 0;
-      v_2 = (v_3 + 1u);
+      {
+        v_2 = (v_3 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/array/strides.spvasm.expected.ir.msl b/test/tint/array/strides.spvasm.expected.ir.msl
index efe481d..3ffaec4 100644
--- a/test/tint/array/strides.spvasm.expected.ir.msl
+++ b/test/tint/array/strides.spvasm.expected.ir.msl
@@ -38,7 +38,9 @@
         break;
       }
       tint_store_and_preserve_padding_4((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
@@ -53,7 +55,9 @@
         break;
       }
       tint_store_and_preserve_padding_3((&(*target)[v_3]), value_param[v_3]);
-      v_2 = (v_3 + 1u);
+      {
+        v_2 = (v_3 + 1u);
+      }
       continue;
     }
   }
@@ -71,7 +75,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_5]), value_param[v_5]);
-      v_4 = (v_5 + 1u);
+      {
+        v_4 = (v_5 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.ir.msl b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.ir.msl
index bd0f703..86cd047 100644
--- a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.ir.msl
@@ -53,7 +53,9 @@
         break;
       }
       (*target)[v_1] = value_param[v_1];
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.ir.msl b/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.ir.msl
index 25ab70b..89d5b89 100644
--- a/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/dynamic_index/write_f16.wgsl.expected.ir.msl
@@ -67,7 +67,9 @@
         break;
       }
       (*target)[v_1] = value_param[v_1];
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/storage/static_index/write.wgsl.expected.ir.msl b/test/tint/buffer/storage/static_index/write.wgsl.expected.ir.msl
index fb7d39b..89b9fc0 100644
--- a/test/tint/buffer/storage/static_index/write.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/static_index/write.wgsl.expected.ir.msl
@@ -56,7 +56,9 @@
         break;
       }
       (*target)[v_1] = value_param[v_1];
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.ir.msl b/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.ir.msl
index 0021c52..8a4d990 100644
--- a/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/static_index/write_f16.wgsl.expected.ir.msl
@@ -76,7 +76,9 @@
         break;
       }
       tint_store_and_preserve_padding_7((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
@@ -91,7 +93,9 @@
         break;
       }
       (*target)[v_3] = value_param[v_3];
-      v_2 = (v_3 + 1u);
+      {
+        v_2 = (v_3 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
index 996dcf4..e67e087 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -30,7 +30,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = float2x2(float2(0.0f), float2(0.0f));
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.ir.msl
index 1b2bbd6..57e4d16 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_storage.wgsl.expected.ir.msl
@@ -31,7 +31,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
index 0d6478f..da3487f 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -31,7 +31,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = half2x3(half3(0.0h), half3(0.0h));
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.ir.msl
index dbc46bc..3493263 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_storage.wgsl.expected.ir.msl
@@ -31,7 +31,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
index 34e56b0..f471557 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -30,7 +30,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = float2x3(float3(0.0f), float3(0.0f));
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
index ab908e1..e4a7364 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -30,7 +30,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = half2x4(half4(0.0h), half4(0.0h));
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
index 54c7ef8..57df245 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -30,7 +30,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = float2x4(float4(0.0f), float4(0.0f));
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.ir.msl
index e9184d7..5fdff5a 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_storage.wgsl.expected.ir.msl
@@ -32,7 +32,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
index d7abe28..b44e8d3 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -30,7 +30,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = float3x3(float3(0.0f), float3(0.0f), float3(0.0f));
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
index fb760c8..98c262b 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -30,7 +30,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = float3x4(float4(0.0f), float4(0.0f), float4(0.0f));
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
index 8b2ab62..2550d76 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -30,7 +30,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = half4x2(half2(0.0h), half2(0.0h), half2(0.0h), half2(0.0h));
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
index 312e1d9..4fd6b7f 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -30,7 +30,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = float4x2(float2(0.0f), float2(0.0f), float2(0.0f), float2(0.0f));
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.ir.msl
index bb3cd82..2bb28d0 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_storage.wgsl.expected.ir.msl
@@ -33,7 +33,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
index e49a258..3c01f91 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -30,7 +30,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = half4x3(half3(0.0h), half3(0.0h), half3(0.0h), half3(0.0h));
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.ir.msl
index f7d7131..a72dc5e 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_storage.wgsl.expected.ir.msl
@@ -33,7 +33,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
index 34af23e..b72113d 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -30,7 +30,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = float4x3(float3(0.0f), float3(0.0f), float3(0.0f), float3(0.0f));
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
index 7d31396..7735419 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -30,7 +30,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = half4x4(half4(0.0h), half4(0.0h), half4(0.0h), half4(0.0h));
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
index 8138ba1..421d2dd 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -30,7 +30,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = float4x4(float4(0.0f), float4(0.0f), float4(0.0f), float4(0.0f));
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.ir.msl
index 5b1e326..d0a05b0 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_storage.wgsl.expected.ir.msl
@@ -37,7 +37,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
index 325a422..aab2bcd 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.ir.msl
index f5d7e61..80dffc3 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_storage.wgsl.expected.ir.msl
@@ -37,7 +37,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
index 5361ac5..fb73b3c 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.ir.msl
index 416201a..85dd7c9 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_storage.wgsl.expected.ir.msl
@@ -41,7 +41,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
index bdea8cd..db2aebb 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.ir.msl
index cb8ded9..88c6b50 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_storage.wgsl.expected.ir.msl
@@ -41,7 +41,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
index dbc2ab2..6d6e56b 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.ir.msl
index af1c2d0..c0e5941 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_storage.wgsl.expected.ir.msl
@@ -37,7 +37,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
index b3f15ac..138b85c 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.ir.msl
index fc147b4..37626eb 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_storage.wgsl.expected.ir.msl
@@ -37,7 +37,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
index c3c35f0..1d5a52e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.ir.msl
index facce53..d94beff 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_storage.wgsl.expected.ir.msl
@@ -37,7 +37,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
index 0bb6c7c..24934e8 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.ir.msl
index e0009c6..05d74bf 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_storage.wgsl.expected.ir.msl
@@ -37,7 +37,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
index a08d9ef..04c7dee 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.ir.msl
index c7bc56c..cd3c4a8 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_storage.wgsl.expected.ir.msl
@@ -42,7 +42,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
index 50cf5f5..bd8162f 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.ir.msl
index 32e0afe..dd15e70 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_storage.wgsl.expected.ir.msl
@@ -42,7 +42,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
index 27f3aea..ae0657c 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.ir.msl
index 48e1dab..c32eccd 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_storage.wgsl.expected.ir.msl
@@ -37,7 +37,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
index 9d6eb26..589abc0 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.ir.msl
index 87f0b50..27015d4 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_storage.wgsl.expected.ir.msl
@@ -37,7 +37,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
index 243be21..7e549c6 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.ir.msl
index 72a064f..19e8dbc 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_storage.wgsl.expected.ir.msl
@@ -37,7 +37,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
index 47d79b2..da15cfd 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.ir.msl
index a878866..2a9b058 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_storage.wgsl.expected.ir.msl
@@ -37,7 +37,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
index cd33834..d8ab59e 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.ir.msl
index fed5e52..f70bd8f 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_storage.wgsl.expected.ir.msl
@@ -43,7 +43,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
index 40dca05..2ce1445 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.ir.msl
index fea6aec..421dcd5 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_storage.wgsl.expected.ir.msl
@@ -43,7 +43,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
index 528991d..262ac37 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.ir.msl
index 7110f20..ab0de45 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_storage.wgsl.expected.ir.msl
@@ -37,7 +37,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
index 1404357..d8cd693 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.ir.msl
index 4993720..dcac4be 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_storage.wgsl.expected.ir.msl
@@ -37,7 +37,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
index 645e2db..a7d58de 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.w)[v_1] = S{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl b/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl
index 1321fac..0362f07 100644
--- a/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/1403752.wgsl.expected.ir.msl
@@ -9,6 +9,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl b/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl
index 1a7f09c..d0ebe5e 100644
--- a/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/1449538.wgsl.expected.ir.msl
@@ -9,6 +9,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
@@ -19,6 +21,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
@@ -29,15 +33,7 @@
       } else {
         break;
       }
-      continue;
-    }
-  }
-  {
-    int i62 = 0;
-    while(true) {
-      if (false) {
-      } else {
-        break;
+      {
       }
       continue;
     }
@@ -49,6 +45,20 @@
       } else {
         break;
       }
+      {
+      }
+      continue;
+    }
+  }
+  {
+    int i62 = 0;
+    while(true) {
+      if (false) {
+      } else {
+        break;
+      }
+      {
+      }
       continue;
     }
   }
@@ -59,6 +69,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
@@ -69,6 +81,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
@@ -79,6 +93,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.ir.msl b/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.ir.msl
index b203d8a..7a56415 100644
--- a/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/read/workgroup.wgsl.expected.ir.msl
@@ -40,7 +40,9 @@
         break;
       }
       (*tint_module_vars.s).data[v_1] = 0;
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.ir.msl b/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.ir.msl
index dd9da5a..edcb041 100644
--- a/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/dyn_array_idx/write/workgroup.wgsl.expected.ir.msl
@@ -40,7 +40,9 @@
         break;
       }
       (*tint_module_vars.s).data[v_1] = 0;
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl
index cced0a6..d941909 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl
@@ -19,7 +19,9 @@
       (*tint_module_vars.v3i)[i] = 1;
       (*tint_module_vars.v4u)[i] = 1u;
       (*tint_module_vars.v2b)[i] = true;
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
@@ -38,7 +40,9 @@
         break;
       }
       foo(tint_module_vars);
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl
index 410a3d3..6ef7133 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl
@@ -28,7 +28,9 @@
         break;
       }
       foo(tint_module_vars);
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl
index 9bd7e16..a53d1ad 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl
@@ -33,7 +33,9 @@
       v2b[i] = true;
       v3b[i] = true;
       v4b[i] = true;
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl
index 5457395..5c6fc4c3 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl
@@ -25,7 +25,9 @@
       v3i_2[i] = 1;
       v4u_2[i] = 1u;
       v2b_2[i] = true;
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl
index 640fd1f..259100d 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl
@@ -25,7 +25,9 @@
       v2i[i] = 1;
       v2u[i] = 1u;
       v2b[i] = true;
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/tint/1081.wgsl.expected.ir.msl b/test/tint/bug/tint/1081.wgsl.expected.ir.msl
index 7c93222..16a13ca 100644
--- a/test/tint/bug/tint/1081.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1081.wgsl.expected.ir.msl
@@ -24,6 +24,8 @@
       if ((r == 0)) {
         break;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl b/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl
index 796d667..84d42b6 100644
--- a/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1474-a.wgsl.expected.ir.msl
@@ -13,6 +13,8 @@
       } else {
         return;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/tint/1538.wgsl.expected.ir.msl b/test/tint/bug/tint/1538.wgsl.expected.ir.msl
index 21deb74..f5690bc 100644
--- a/test/tint/bug/tint/1538.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1538.wgsl.expected.ir.msl
@@ -38,6 +38,8 @@
       }
       int s = f();
       (*tint_module_vars.buf)[0] = 0u;
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/tint/1557.wgsl.expected.ir.msl b/test/tint/bug/tint/1557.wgsl.expected.ir.msl
index c775000..c98dc8b 100644
--- a/test/tint/bug/tint/1557.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1557.wgsl.expected.ir.msl
@@ -16,6 +16,8 @@
       }
       j = (j + 1);
       int k = f();
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/tint/1605.wgsl.expected.ir.msl b/test/tint/bug/tint/1605.wgsl.expected.ir.msl
index 4379028..1f37c49 100644
--- a/test/tint/bug/tint/1605.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1605.wgsl.expected.ir.msl
@@ -22,7 +22,9 @@
           return false;
         }
       }
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/tint/1764.wgsl.expected.ir.msl b/test/tint/bug/tint/1764.wgsl.expected.ir.msl
index ccecc39..fded27c 100644
--- a/test/tint/bug/tint/1764.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1764.wgsl.expected.ir.msl
@@ -29,7 +29,9 @@
         break;
       }
       (*tint_module_vars.W)[v_1] = 0;
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/tint/2059.wgsl.expected.ir.msl b/test/tint/bug/tint/2059.wgsl.expected.ir.msl
index 1da6193..e23abed 100644
--- a/test/tint/bug/tint/2059.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/2059.wgsl.expected.ir.msl
@@ -50,7 +50,9 @@
         break;
       }
       tint_store_and_preserve_padding((&(*target)[v_1]), value_param[v_1]);
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
@@ -68,7 +70,9 @@
         break;
       }
       tint_store_and_preserve_padding_2((&(*target)[v_3]), value_param[v_3]);
-      v_2 = (v_3 + 1u);
+      {
+        v_2 = (v_3 + 1u);
+      }
       continue;
     }
   }
@@ -86,7 +90,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_5]), value_param[v_5]);
-      v_4 = (v_5 + 1u);
+      {
+        v_4 = (v_5 + 1u);
+      }
       continue;
     }
   }
@@ -111,7 +117,9 @@
       float const v_7 = float(((c * 3u) + 1u));
       float const v_8 = float(((c * 3u) + 2u));
       (*v_6) = float3(v_7, v_8, float(((c * 3u) + 3u)));
-      c = (c + 1u);
+      {
+        c = (c + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/tint/221.wgsl.expected.ir.msl b/test/tint/bug/tint/221.wgsl.expected.ir.msl
index 24bb006..9d14fce 100644
--- a/test/tint/bug/tint/221.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/221.wgsl.expected.ir.msl
@@ -34,13 +34,17 @@
       }
       device uint* const p = (&(*tint_module_vars.b).data[i]);
       if ((tint_mod_u32(i, 2u) == 0u)) {
-        (*p) = ((*p) * 2u);
-        i = (i + 1u);
+        {
+          (*p) = ((*p) * 2u);
+          i = (i + 1u);
+        }
         continue;
       }
       (*p) = 0u;
-      (*p) = ((*p) * 2u);
-      i = (i + 1u);
+      {
+        (*p) = ((*p) * 2u);
+        i = (i + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/tint/744.wgsl.expected.ir.msl b/test/tint/bug/tint/744.wgsl.expected.ir.msl
index ee964c8..34fbae0 100644
--- a/test/tint/bug/tint/744.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/744.wgsl.expected.ir.msl
@@ -42,7 +42,9 @@
       uint const a = (i + (resultCell[0u] * dimInner));
       uint const b = (resultCell[1u] + (i * dimOutter));
       result = (result + ((*tint_module_vars.firstMatrix).numbers[a] * (*tint_module_vars.secondMatrix).numbers[b]));
-      i = (i + 1u);
+      {
+        i = (i + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/tint/914.wgsl.expected.ir.msl b/test/tint/bug/tint/914.wgsl.expected.ir.msl
index 73bf394..9d9ab28 100644
--- a/test/tint/bug/tint/914.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/914.wgsl.expected.ir.msl
@@ -85,7 +85,9 @@
       }
       (*tint_module_vars.mm_Asub)[(v_4 / 64u)][(v_4 % 64u)] = 0.0f;
       (*tint_module_vars.mm_Bsub)[(v_4 / 64u)][(v_4 % 64u)] = 0.0f;
-      v_3 = (v_4 + 256u);
+      {
+        v_3 = (v_4 + 256u);
+      }
       continue;
     }
   }
@@ -106,7 +108,9 @@
         break;
       }
       acc[index] = 0.0f;
-      index = (index + 1u);
+      {
+        index = (index + 1u);
+      }
       continue;
     }
   }
@@ -138,11 +142,15 @@
               uint const inputRow = (tileRow + innerRow);
               uint const inputCol = (tileColA + innerCol);
               (*tint_module_vars.mm_Asub)[inputRow][inputCol] = mm_readA((globalRow + innerRow), ((t * 64u) + inputCol), tint_module_vars);
-              innerCol = (innerCol + 1u);
+              {
+                innerCol = (innerCol + 1u);
+              }
               continue;
             }
           }
-          innerRow = (innerRow + 1u);
+          {
+            innerRow = (innerRow + 1u);
+          }
           continue;
         }
       }
@@ -164,11 +172,15 @@
               uint const inputCol = (tileCol + innerCol);
               threadgroup float* const v_5 = (&(*tint_module_vars.mm_Bsub)[innerCol][inputCol]);
               (*v_5) = mm_readB(((t * 64u) + inputRow), (globalCol + innerCol), tint_module_vars);
-              innerCol = (innerCol + 1u);
+              {
+                innerCol = (innerCol + 1u);
+              }
               continue;
             }
           }
-          innerRow = (innerRow + 1u);
+          {
+            innerRow = (innerRow + 1u);
+          }
           continue;
         }
       }
@@ -188,7 +200,9 @@
                 break;
               }
               BCached[inner] = (*tint_module_vars.mm_Bsub)[k][(tileCol + inner)];
-              inner = (inner + 1u);
+              {
+                inner = (inner + 1u);
+              }
               continue;
             }
           }
@@ -209,20 +223,28 @@
                   }
                   uint const index = ((innerRow * 4u) + innerCol);
                   acc[index] = (acc[index] + (ACached * BCached[innerCol]));
-                  innerCol = (innerCol + 1u);
+                  {
+                    innerCol = (innerCol + 1u);
+                  }
                   continue;
                 }
               }
-              innerRow = (innerRow + 1u);
+              {
+                innerRow = (innerRow + 1u);
+              }
               continue;
             }
           }
-          k = (k + 1u);
+          {
+            k = (k + 1u);
+          }
           continue;
         }
       }
       threadgroup_barrier(mem_flags::mem_threadgroup);
-      t = (t + 1u);
+      {
+        t = (t + 1u);
+      }
       continue;
     }
   }
@@ -242,11 +264,15 @@
           }
           uint const index = ((innerRow * 4u) + innerCol);
           mm_write((globalRow + innerRow), (globalCol + innerCol), acc[index], tint_module_vars);
-          innerCol = (innerCol + 1u);
+          {
+            innerCol = (innerCol + 1u);
+          }
           continue;
         }
       }
-      innerRow = (innerRow + 1u);
+      {
+        innerRow = (innerRow + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/bug/tint/990.wgsl.expected.ir.msl b/test/tint/bug/tint/990.wgsl.expected.ir.msl
index aa031c2..c72f776 100644
--- a/test/tint/bug/tint/990.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/990.wgsl.expected.ir.msl
@@ -10,6 +10,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl
index d39052c..8f1eafe 100644
--- a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl
@@ -17,10 +17,12 @@
       } else {
         break;
       }
-      threadgroup_barrier(mem_flags::mem_threadgroup);
-      int const v_2 = (*tint_module_vars.b);
-      threadgroup_barrier(mem_flags::mem_threadgroup);
-      i = (i + v_2);
+      {
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+        int const v_2 = (*tint_module_vars.b);
+        threadgroup_barrier(mem_flags::mem_threadgroup);
+        i = (i + v_2);
+      }
       continue;
     }
   }
diff --git a/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.ir.msl
index 66c7b64..a6d9b15 100644
--- a/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/for_loop_attribute.wgsl.expected.ir.msl
@@ -30,6 +30,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl
index be50a02..d5d5332 100644
--- a/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/loop_attribute.wgsl.expected.ir.msl
@@ -20,7 +20,9 @@
   {
     while(true) {
       dfdx(1.0f);
-      if (x > 0.0f) { break; }
+      {
+        if (x > 0.0f) { break; }
+      }
       continue;
     }
   }
diff --git a/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.ir.msl
index 8929f16..8989ea2 100644
--- a/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/loop_body_attribute.wgsl.expected.ir.msl
@@ -20,7 +20,9 @@
   {
     while(true) {
       dfdx(1.0f);
-      if (x > 0.0f) { break; }
+      {
+        if (x > 0.0f) { break; }
+      }
       continue;
     }
   }
diff --git a/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.ir.msl
index ec61c6b..974c13a 100644
--- a/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/loop_continuing_attribute.wgsl.expected.ir.msl
@@ -19,8 +19,10 @@
 void tint_symbol_inner(float x) {
   {
     while(true) {
-      dfdx(1.0f);
-      if (x > 0.0f) { break; }
+      {
+        dfdx(1.0f);
+        if (x > 0.0f) { break; }
+      }
       continue;
     }
   }
diff --git a/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.ir.msl
index c92eacd..bc1de81 100644
--- a/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/while_loop_attribute.wgsl.expected.ir.msl
@@ -30,6 +30,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.ir.msl b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.ir.msl
index 633fdb3..4add4ed 100644
--- a/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.ir.msl
+++ b/test/tint/layout/storage/mat2x2/stride/16.spvasm.expected.ir.msl
@@ -42,7 +42,9 @@
         break;
       }
       tint_store_and_preserve_padding_1((&(*target)[v_2]), value_param[v_2]);
-      v_1 = (v_2 + 1u);
+      {
+        v_1 = (v_2 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl b/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl
index 88c3be7..9114b25 100644
--- a/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl
@@ -12,7 +12,9 @@
       switch(i) {
         case 0:
         {
-          i = (i + 1);
+          {
+            i = (i + 1);
+          }
           continue;
         }
         default:
@@ -20,7 +22,9 @@
           break;
         }
       }
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/loops/loop.wgsl.expected.ir.msl b/test/tint/loops/loop.wgsl.expected.ir.msl
index e696a11..777ab74 100644
--- a/test/tint/loops/loop.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop.wgsl.expected.ir.msl
@@ -9,6 +9,8 @@
       if ((i > 4)) {
         return i;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl b/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl
index 118d545..a0cf3db 100644
--- a/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl
@@ -8,8 +8,10 @@
       if ((i > 4)) {
         return i;
       }
-      i = (i + 1);
-      if (i == 4) { break; }
+      {
+        i = (i + 1);
+        if (i == 4) { break; }
+      }
       continue;
     }
   }
diff --git a/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl b/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl
index 327c76a..66d9c1c 100644
--- a/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl
@@ -8,7 +8,9 @@
       if ((i > 4)) {
         return i;
       }
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/loops/multiple_continues.wgsl.expected.ir.msl b/test/tint/loops/multiple_continues.wgsl.expected.ir.msl
index de48573..c5e78ae 100644
--- a/test/tint/loops/multiple_continues.wgsl.expected.ir.msl
+++ b/test/tint/loops/multiple_continues.wgsl.expected.ir.msl
@@ -12,17 +12,23 @@
       switch(i) {
         case 0:
         {
-          i = (i + 1);
+          {
+            i = (i + 1);
+          }
           continue;
         }
         case 1:
         {
-          i = (i + 1);
+          {
+            i = (i + 1);
+          }
           continue;
         }
         case 2:
         {
-          i = (i + 1);
+          {
+            i = (i + 1);
+          }
           continue;
         }
         default:
@@ -30,7 +36,9 @@
           break;
         }
       }
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/loops/multiple_switch.wgsl.expected.ir.msl b/test/tint/loops/multiple_switch.wgsl.expected.ir.msl
index 6a761f7..d8a8569 100644
--- a/test/tint/loops/multiple_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/multiple_switch.wgsl.expected.ir.msl
@@ -13,7 +13,9 @@
       switch(i) {
         case 0:
         {
-          i = (i + 1);
+          {
+            i = (i + 1);
+          }
           continue;
         }
         default:
@@ -24,7 +26,9 @@
       switch(i) {
         case 0:
         {
-          i = (i + 1);
+          {
+            i = (i + 1);
+          }
           continue;
         }
         default:
@@ -32,7 +36,9 @@
           break;
         }
       }
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl
index 8921168..548ba27 100644
--- a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl
@@ -19,7 +19,9 @@
           switch(i) {
             case 0:
             {
-              j = (j + 2);
+              {
+                j = (j + 2);
+              }
               continue;
             }
             default:
@@ -27,11 +29,15 @@
               break;
             }
           }
-          j = (j + 2);
+          {
+            j = (j + 2);
+          }
           continue;
         }
       }
-      i = (i + 2);
+      {
+        i = (i + 2);
+      }
       continue;
     }
   }
diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl
index 6dbcb48..ae0b853 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl
@@ -22,7 +22,9 @@
               switch(j) {
                 case 0:
                 {
-                  j = (j + 2);
+                  {
+                    j = (j + 2);
+                  }
                   continue;
                 }
                 default:
@@ -30,11 +32,15 @@
                   break;
                 }
               }
-              j = (j + 2);
+              {
+                j = (j + 2);
+              }
               continue;
             }
           }
-          i = (i + 2);
+          {
+            i = (i + 2);
+          }
           continue;
         }
         default:
@@ -42,7 +48,9 @@
           break;
         }
       }
-      i = (i + 2);
+      {
+        i = (i + 2);
+      }
       continue;
     }
   }
diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl
index 2fea1a7..6c06c1e 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl
@@ -23,7 +23,9 @@
               switch(j) {
                 case 0:
                 {
-                  j = (j + 2);
+                  {
+                    j = (j + 2);
+                  }
                   continue;
                 }
                 case 1:
@@ -31,7 +33,9 @@
                   switch(k) {
                     case 0:
                     {
-                      j = (j + 2);
+                      {
+                        j = (j + 2);
+                      }
                       continue;
                     }
                     default:
@@ -46,11 +50,15 @@
                   break;
                 }
               }
-              j = (j + 2);
+              {
+                j = (j + 2);
+              }
               continue;
             }
           }
-          i = (i + 2);
+          {
+            i = (i + 2);
+          }
           continue;
         }
         default:
@@ -58,7 +66,9 @@
           break;
         }
       }
-      i = (i + 2);
+      {
+        i = (i + 2);
+      }
       continue;
     }
   }
diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl
index a4f719e..703e006 100644
--- a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl
@@ -16,7 +16,9 @@
           switch(j) {
             case 0:
             {
-              i = (i + 2);
+              {
+                i = (i + 2);
+              }
               continue;
             }
             default:
@@ -31,7 +33,9 @@
           break;
         }
       }
-      i = (i + 2);
+      {
+        i = (i + 2);
+      }
       continue;
     }
   }
diff --git a/test/tint/loops/nested_loops.wgsl.expected.ir.msl b/test/tint/loops/nested_loops.wgsl.expected.ir.msl
index 70b1c91..939ab1f 100644
--- a/test/tint/loops/nested_loops.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loops.wgsl.expected.ir.msl
@@ -16,6 +16,8 @@
           if ((j > 4)) {
             return 2;
           }
+          {
+          }
           continue;
         }
       }
diff --git a/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl
index 8f07396..53fe43a 100644
--- a/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl
@@ -14,7 +14,9 @@
           if ((j > 4)) {
             return 2;
           }
-          j = (j + 1);
+          {
+            j = (j + 1);
+          }
           continue;
         }
       }
diff --git a/test/tint/loops/single_continue.wgsl.expected.ir.msl b/test/tint/loops/single_continue.wgsl.expected.ir.msl
index 573c2b5..7fecaf4 100644
--- a/test/tint/loops/single_continue.wgsl.expected.ir.msl
+++ b/test/tint/loops/single_continue.wgsl.expected.ir.msl
@@ -12,7 +12,9 @@
       switch(i) {
         case 0:
         {
-          i = (i + 1);
+          {
+            i = (i + 1);
+          }
           continue;
         }
         default:
@@ -20,7 +22,9 @@
           break;
         }
       }
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/loops/while.wgsl.expected.ir.msl b/test/tint/loops/while.wgsl.expected.ir.msl
index 9686db1..fdf1dd6 100644
--- a/test/tint/loops/while.wgsl.expected.ir.msl
+++ b/test/tint/loops/while.wgsl.expected.ir.msl
@@ -10,6 +10,8 @@
         break;
       }
       i = (i + 1);
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/loops/while_with_continue.wgsl.expected.ir.msl b/test/tint/loops/while_with_continue.wgsl.expected.ir.msl
index 9686db1..fdf1dd6 100644
--- a/test/tint/loops/while_with_continue.wgsl.expected.ir.msl
+++ b/test/tint/loops/while_with_continue.wgsl.expected.ir.msl
@@ -10,6 +10,8 @@
         break;
       }
       i = (i + 1);
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.ir.msl
index cf7ef8e..06d5e1e 100644
--- a/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/array_in_struct.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.S).arr[v_1] = 0;
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.ir.msl
index 7bdbc34..7b1e624 100644
--- a/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/struct_in_array.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.S)[v_1] = str{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.ir.msl
index 787c1b1..32bd669 100644
--- a/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/array_in_struct.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.S).arr[v_1] = 0;
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.ir.msl
index e0e8d90..945733b 100644
--- a/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/struct_in_array.wgsl.expected.ir.msl
@@ -35,7 +35,9 @@
         break;
       }
       (*tint_module_vars.S)[v_1] = str{};
-      v = (v_1 + 1u);
+      {
+        v = (v_1 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/shadowing/loop.wgsl b/test/tint/shadowing/loop.wgsl
new file mode 100644
index 0000000..4a4f483
--- /dev/null
+++ b/test/tint/shadowing/loop.wgsl
@@ -0,0 +1,16 @@
+@group(0) @binding(0)
+var<storage, read_write> output : array<i32, 10>;
+
+@compute @workgroup_size(1)
+fn foo() {
+  var i = 0;
+  loop {
+    var x = output[i];
+    continuing {
+      var x = output[x];
+      i += x;
+      break if i > 10;
+    }
+  }
+  output[0] = i;
+}
diff --git a/test/tint/shadowing/loop.wgsl.expected.dxc.hlsl b/test/tint/shadowing/loop.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..8be1d6f
--- /dev/null
+++ b/test/tint/shadowing/loop.wgsl.expected.dxc.hlsl
@@ -0,0 +1,16 @@
+RWByteAddressBuffer output : register(u0);
+
+[numthreads(1, 1, 1)]
+void foo() {
+  int i = 0;
+  while (true) {
+    int x = asint(output.Load((4u * uint(i))));
+    {
+      int x_1 = asint(output.Load((4u * uint(x))));
+      i = (i + x_1);
+      if ((i > 10)) { break; }
+    }
+  }
+  output.Store(0u, asuint(i));
+  return;
+}
diff --git a/test/tint/shadowing/loop.wgsl.expected.fxc.hlsl b/test/tint/shadowing/loop.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..8be1d6f
--- /dev/null
+++ b/test/tint/shadowing/loop.wgsl.expected.fxc.hlsl
@@ -0,0 +1,16 @@
+RWByteAddressBuffer output : register(u0);
+
+[numthreads(1, 1, 1)]
+void foo() {
+  int i = 0;
+  while (true) {
+    int x = asint(output.Load((4u * uint(i))));
+    {
+      int x_1 = asint(output.Load((4u * uint(x))));
+      i = (i + x_1);
+      if ((i > 10)) { break; }
+    }
+  }
+  output.Store(0u, asuint(i));
+  return;
+}
diff --git a/test/tint/shadowing/loop.wgsl.expected.glsl b/test/tint/shadowing/loop.wgsl.expected.glsl
new file mode 100644
index 0000000..d106264
--- /dev/null
+++ b/test/tint/shadowing/loop.wgsl.expected.glsl
@@ -0,0 +1,24 @@
+#version 310 es
+
+layout(binding = 0, std430) buffer tint_symbol_block_ssbo {
+  int inner[10];
+} tint_symbol;
+
+void foo() {
+  int i = 0;
+  while (true) {
+    int x = tint_symbol.inner[i];
+    {
+      int x_1 = tint_symbol.inner[x];
+      i = (i + x_1);
+      if ((i > 10)) { break; }
+    }
+  }
+  tint_symbol.inner[0] = i;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  foo();
+  return;
+}
diff --git a/test/tint/shadowing/loop.wgsl.expected.ir.msl b/test/tint/shadowing/loop.wgsl.expected.ir.msl
new file mode 100644
index 0000000..81b5e39
--- /dev/null
+++ b/test/tint/shadowing/loop.wgsl.expected.ir.msl
@@ -0,0 +1,34 @@
+#include <metal_stdlib>
+using namespace metal;
+template<typename T, size_t N>
+struct tint_array {
+  const constant T& operator[](size_t i) const constant { return elements[i]; }
+  device T& operator[](size_t i) device { return elements[i]; }
+  const device T& operator[](size_t i) const device { return elements[i]; }
+  thread T& operator[](size_t i) thread { return elements[i]; }
+  const thread T& operator[](size_t i) const thread { return elements[i]; }
+  threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+  const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+  T elements[N];
+};
+
+struct tint_module_vars_struct {
+  device tint_array<int, 10>* output;
+};
+
+kernel void foo(device tint_array<int, 10>* output [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.output=output};
+  int i = 0;
+  {
+    while(true) {
+      int x = (*tint_module_vars.output)[i];
+      {
+        int x = (*tint_module_vars.output)[x];
+        i = (i + x);
+        if (i > 10) { break; }
+      }
+      continue;
+    }
+  }
+  (*tint_module_vars.output)[0] = i;
+}
diff --git a/test/tint/shadowing/loop.wgsl.expected.ir.spvasm b/test/tint/shadowing/loop.wgsl.expected.ir.spvasm
new file mode 100644
index 0000000..4b6635d
--- /dev/null
+++ b/test/tint/shadowing/loop.wgsl.expected.ir.spvasm
@@ -0,0 +1,69 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 38
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %foo "foo"
+               OpExecutionMode %foo LocalSize 1 1 1
+               OpMemberName %tint_symbol_1 0 "tint_symbol"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %foo "foo"
+               OpName %i "i"
+               OpName %x "x"
+               OpName %x_0 "x"
+               OpDecorate %_arr_int_uint_10 ArrayStride 4
+               OpMemberDecorate %tint_symbol_1 0 Offset 0
+               OpDecorate %tint_symbol_1 Block
+               OpDecorate %1 DescriptorSet 0
+               OpDecorate %1 Binding 0
+        %int = OpTypeInt 32 1
+       %uint = OpTypeInt 32 0
+    %uint_10 = OpConstant %uint 10
+%_arr_int_uint_10 = OpTypeArray %int %uint_10
+%tint_symbol_1 = OpTypeStruct %_arr_int_uint_10
+%_ptr_StorageBuffer_tint_symbol_1 = OpTypePointer StorageBuffer %tint_symbol_1
+          %1 = OpVariable %_ptr_StorageBuffer_tint_symbol_1 StorageBuffer
+       %void = OpTypeVoid
+         %10 = OpTypeFunction %void
+%_ptr_Function_int = OpTypePointer Function %int
+      %int_0 = OpConstant %int 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+     %uint_0 = OpConstant %uint 0
+     %int_10 = OpConstant %int 10
+       %bool = OpTypeBool
+        %foo = OpFunction %void None %10
+         %11 = OpLabel
+          %i = OpVariable %_ptr_Function_int Function
+          %x = OpVariable %_ptr_Function_int Function
+        %x_0 = OpVariable %_ptr_Function_int Function
+               OpStore %i %int_0
+               OpBranch %17
+         %17 = OpLabel
+               OpLoopMerge %18 %16 None
+               OpBranch %15
+         %15 = OpLabel
+         %19 = OpLoad %int %i
+         %20 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0 %19
+         %23 = OpLoad %int %20
+               OpStore %x %23
+               OpBranch %16
+         %16 = OpLabel
+         %25 = OpLoad %int %x
+         %26 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0 %25
+         %27 = OpLoad %int %26
+               OpStore %x_0 %27
+         %29 = OpLoad %int %x_0
+         %30 = OpLoad %int %i
+         %31 = OpIAdd %int %30 %29
+               OpStore %i %31
+         %32 = OpLoad %int %i
+         %33 = OpSGreaterThan %bool %32 %int_10
+               OpBranchConditional %33 %18 %17
+         %18 = OpLabel
+         %36 = OpAccessChain %_ptr_StorageBuffer_int %1 %uint_0 %int_0
+         %37 = OpLoad %int %i
+               OpStore %36 %37
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/shadowing/loop.wgsl.expected.msl b/test/tint/shadowing/loop.wgsl.expected.msl
new file mode 100644
index 0000000..5a225bd
--- /dev/null
+++ b/test/tint/shadowing/loop.wgsl.expected.msl
@@ -0,0 +1,34 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+    const constant T& operator[](size_t i) const constant { return elements[i]; }
+    device T& operator[](size_t i) device { return elements[i]; }
+    const device T& operator[](size_t i) const device { return elements[i]; }
+    thread T& operator[](size_t i) thread { return elements[i]; }
+    const thread T& operator[](size_t i) const thread { return elements[i]; }
+    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+    T elements[N];
+};
+
+#define TINT_ISOLATE_UB(VOLATILE_NAME) \
+  volatile bool VOLATILE_NAME = true; \
+  if (VOLATILE_NAME)
+
+kernel void foo(device tint_array<int, 10>* tint_symbol [[buffer(0)]]) {
+  int i = 0;
+  TINT_ISOLATE_UB(tint_volatile_true) while(true) {
+    int x = (*(tint_symbol))[i];
+    {
+      int x_1 = (*(tint_symbol))[x];
+      i = as_type<int>((as_type<uint>(i) + as_type<uint>(x_1)));
+      if ((i > 10)) { break; }
+    }
+  }
+  (*(tint_symbol))[0] = i;
+  return;
+}
+
diff --git a/test/tint/shadowing/loop.wgsl.expected.spvasm b/test/tint/shadowing/loop.wgsl.expected.spvasm
new file mode 100644
index 0000000..4b7a4b1
--- /dev/null
+++ b/test/tint/shadowing/loop.wgsl.expected.spvasm
@@ -0,0 +1,70 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 38
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %foo "foo"
+               OpExecutionMode %foo LocalSize 1 1 1
+               OpName %output_block "output_block"
+               OpMemberName %output_block 0 "inner"
+               OpName %output "output"
+               OpName %foo "foo"
+               OpName %i "i"
+               OpName %x "x"
+               OpName %x_1 "x_1"
+               OpDecorate %output_block Block
+               OpMemberDecorate %output_block 0 Offset 0
+               OpDecorate %_arr_int_uint_10 ArrayStride 4
+               OpDecorate %output DescriptorSet 0
+               OpDecorate %output Binding 0
+        %int = OpTypeInt 32 1
+       %uint = OpTypeInt 32 0
+    %uint_10 = OpConstant %uint 10
+%_arr_int_uint_10 = OpTypeArray %int %uint_10
+%output_block = OpTypeStruct %_arr_int_uint_10
+%_ptr_StorageBuffer_output_block = OpTypePointer StorageBuffer %output_block
+     %output = OpVariable %_ptr_StorageBuffer_output_block StorageBuffer
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+         %12 = OpConstantNull %int
+%_ptr_Function_int = OpTypePointer Function %int
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+     %int_10 = OpConstant %int 10
+       %bool = OpTypeBool
+        %foo = OpFunction %void None %8
+         %11 = OpLabel
+          %i = OpVariable %_ptr_Function_int Function %12
+          %x = OpVariable %_ptr_Function_int Function %12
+        %x_1 = OpVariable %_ptr_Function_int Function %12
+               OpStore %i %12
+               OpBranch %15
+         %15 = OpLabel
+               OpLoopMerge %16 %17 None
+               OpBranch %18
+         %18 = OpLabel
+         %20 = OpLoad %int %i
+         %22 = OpAccessChain %_ptr_StorageBuffer_int %output %uint_0 %20
+         %23 = OpLoad %int %22
+               OpStore %x %23
+               OpBranch %17
+         %17 = OpLabel
+         %25 = OpLoad %int %x
+         %26 = OpAccessChain %_ptr_StorageBuffer_int %output %uint_0 %25
+         %27 = OpLoad %int %26
+               OpStore %x_1 %27
+         %29 = OpLoad %int %i
+         %30 = OpLoad %int %x_1
+         %31 = OpIAdd %int %29 %30
+               OpStore %i %31
+         %32 = OpLoad %int %i
+         %34 = OpSGreaterThan %bool %32 %int_10
+               OpBranchConditional %34 %16 %15
+         %16 = OpLabel
+         %36 = OpAccessChain %_ptr_StorageBuffer_int %output %uint_0 %12
+         %37 = OpLoad %int %i
+               OpStore %36 %37
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/shadowing/loop.wgsl.expected.wgsl b/test/tint/shadowing/loop.wgsl.expected.wgsl
new file mode 100644
index 0000000..5447b18
--- /dev/null
+++ b/test/tint/shadowing/loop.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+@group(0) @binding(0) var<storage, read_write> output : array<i32, 10>;
+
+@compute @workgroup_size(1)
+fn foo() {
+  var i = 0;
+  loop {
+    var x = output[i];
+
+    continuing {
+      var x = output[x];
+      i += x;
+      break if (i > 10);
+    }
+  }
+  output[0] = i;
+}
diff --git a/test/tint/statements/compound_assign/for_loop.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/for_loop.wgsl.expected.ir.msl
index 05a8d5f..d371ac5 100644
--- a/test/tint/statements/compound_assign/for_loop.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/for_loop.wgsl.expected.ir.msl
@@ -44,8 +44,10 @@
       } else {
         break;
       }
-      thread float* const v_2 = (&a[idx3(tint_module_vars)]);
-      (*v_2) = ((*v_2) + 1.0f);
+      {
+        thread float* const v_2 = (&a[idx3(tint_module_vars)]);
+        (*v_2) = ((*v_2) + 1.0f);
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/decrement/complex.wgsl.expected.ir.msl b/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
index 06f3397..ec9335f 100644
--- a/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
@@ -55,10 +55,12 @@
       } else {
         break;
       }
-      int const v_4 = idx4(tint_module_vars);
-      device int4* const v_5 = (&(*tint_module_vars.tint_symbol)[v_4].a[idx5(tint_module_vars)]);
-      int const v_6 = idx6(tint_module_vars);
-      (*v_5)[v_6] = ((*v_5)[v_6] - 1);
+      {
+        int const v_4 = idx4(tint_module_vars);
+        device int4* const v_5 = (&(*tint_module_vars.tint_symbol)[v_4].a[idx5(tint_module_vars)]);
+        int const v_6 = idx6(tint_module_vars);
+        (*v_5)[v_6] = ((*v_5)[v_6] - 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.ir.msl b/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.ir.msl
index d52fe70..8b97889 100644
--- a/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/for_loop_continuing.wgsl.expected.ir.msl
@@ -11,7 +11,9 @@
       } else {
         break;
       }
-      (*tint_module_vars.i) = ((*tint_module_vars.i) - 1u);
+      {
+        (*tint_module_vars.i) = ((*tint_module_vars.i) - 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.ir.msl b/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.ir.msl
index 5127afa..c8e815f 100644
--- a/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/for_loop_initializer.wgsl.expected.ir.msl
@@ -12,6 +12,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl b/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl
index 22fc778..477d758 100644
--- a/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl
+++ b/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl
@@ -31,8 +31,10 @@
           }
           return;
         }
-        i = (i + 1);
-        if (i == 5) { break; }
+        {
+          i = (i + 1);
+          if (i == 5) { break; }
+        }
         continue;
       }
     }
diff --git a/test/tint/statements/for/basic.wgsl.expected.ir.msl b/test/tint/statements/for/basic.wgsl.expected.ir.msl
index b98ea67..751223d 100644
--- a/test/tint/statements/for/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/basic.wgsl.expected.ir.msl
@@ -12,7 +12,9 @@
         break;
       }
       some_loop_body();
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/for/complex.wgsl.expected.ir.msl b/test/tint/statements/for/complex.wgsl.expected.ir.msl
index 8ce0549..100b9ed 100644
--- a/test/tint/statements/for/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/complex.wgsl.expected.ir.msl
@@ -20,7 +20,9 @@
       }
       some_loop_body();
       j = (i * 30);
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/for/condition/array_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/condition/array_ctor.wgsl.expected.ir.msl
index 6c73748..0acd4a1 100644
--- a/test/tint/statements/for/condition/array_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/condition/array_ctor.wgsl.expected.ir.msl
@@ -9,6 +9,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/for/condition/basic.wgsl.expected.ir.msl b/test/tint/statements/for/condition/basic.wgsl.expected.ir.msl
index 20ee970..c42baba 100644
--- a/test/tint/statements/for/condition/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/condition/basic.wgsl.expected.ir.msl
@@ -9,6 +9,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/for/condition/struct_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/condition/struct_ctor.wgsl.expected.ir.msl
index 6c73748..0acd4a1 100644
--- a/test/tint/statements/for/condition/struct_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/condition/struct_ctor.wgsl.expected.ir.msl
@@ -9,6 +9,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl
index 40a61e0..ec98763 100644
--- a/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl
@@ -9,7 +9,9 @@
       } else {
         break;
       }
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl b/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl
index 40a61e0..ec98763 100644
--- a/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl
@@ -9,7 +9,9 @@
       } else {
         break;
       }
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl
index 4285951..bce67aa 100644
--- a/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl
@@ -9,7 +9,9 @@
       } else {
         break;
       }
-      i = (i + 1);
+      {
+        i = (i + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/for/empty.wgsl.expected.ir.msl b/test/tint/statements/for/empty.wgsl.expected.ir.msl
index 7a17453..47641f0 100644
--- a/test/tint/statements/for/empty.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/empty.wgsl.expected.ir.msl
@@ -8,6 +8,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/for/initializer/array_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/initializer/array_ctor.wgsl.expected.ir.msl
index c0e53bb..db40de0 100644
--- a/test/tint/statements/for/initializer/array_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/initializer/array_ctor.wgsl.expected.ir.msl
@@ -9,6 +9,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/for/initializer/basic.wgsl.expected.ir.msl b/test/tint/statements/for/initializer/basic.wgsl.expected.ir.msl
index 60d3ea2..61e7a43 100644
--- a/test/tint/statements/for/initializer/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/initializer/basic.wgsl.expected.ir.msl
@@ -9,6 +9,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.ir.msl
index c0e53bb..db40de0 100644
--- a/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/initializer/struct_ctor.wgsl.expected.ir.msl
@@ -9,6 +9,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/increment/complex.wgsl.expected.ir.msl b/test/tint/statements/increment/complex.wgsl.expected.ir.msl
index 309a3b0..6b1f5c0 100644
--- a/test/tint/statements/increment/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/complex.wgsl.expected.ir.msl
@@ -55,10 +55,12 @@
       } else {
         break;
       }
-      int const v_4 = idx4(tint_module_vars);
-      device int4* const v_5 = (&(*tint_module_vars.tint_symbol)[v_4].a[idx5(tint_module_vars)]);
-      int const v_6 = idx6(tint_module_vars);
-      (*v_5)[v_6] = ((*v_5)[v_6] + 1);
+      {
+        int const v_4 = idx4(tint_module_vars);
+        device int4* const v_5 = (&(*tint_module_vars.tint_symbol)[v_4].a[idx5(tint_module_vars)]);
+        int const v_6 = idx6(tint_module_vars);
+        (*v_5)[v_6] = ((*v_5)[v_6] + 1);
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/increment/for_loop_continuing.wgsl.expected.ir.msl b/test/tint/statements/increment/for_loop_continuing.wgsl.expected.ir.msl
index d286201..d066265 100644
--- a/test/tint/statements/increment/for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/for_loop_continuing.wgsl.expected.ir.msl
@@ -11,7 +11,9 @@
       } else {
         break;
       }
-      (*tint_module_vars.i) = ((*tint_module_vars.i) + 1u);
+      {
+        (*tint_module_vars.i) = ((*tint_module_vars.i) + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/statements/increment/for_loop_initializer.wgsl.expected.ir.msl b/test/tint/statements/increment/for_loop_initializer.wgsl.expected.ir.msl
index 01c62af..7cd462d 100644
--- a/test/tint/statements/increment/for_loop_initializer.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/for_loop_initializer.wgsl.expected.ir.msl
@@ -12,6 +12,8 @@
       } else {
         break;
       }
+      {
+      }
       continue;
     }
   }
diff --git a/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.ir.msl b/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.ir.msl
index 269da23..cd3a51a 100644
--- a/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/array/array_i32.wgsl.expected.ir.msl
@@ -29,7 +29,9 @@
         break;
       }
       (*tint_module_vars.zero)[(v_2 / 3u)][(v_2 % 3u)] = 0;
-      v_1 = (v_2 + 1u);
+      {
+        v_1 = (v_2 + 1u);
+      }
       continue;
     }
   }
diff --git a/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.ir.msl b/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.ir.msl
index 531dd07..300fb5e 100644
--- a/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/array/i32.wgsl.expected.ir.msl
@@ -29,7 +29,9 @@
         break;
       }
       (*tint_module_vars.zero)[v_2] = 0;
-      v_1 = (v_2 + 1u);
+      {
+        v_1 = (v_2 + 1u);
+      }
       continue;
     }
   }