diff --git a/src/dawn/native/Toggles.cpp b/src/dawn/native/Toggles.cpp
index dd87d12..7ab0247 100644
--- a/src/dawn/native/Toggles.cpp
+++ b/src/dawn/native/Toggles.cpp
@@ -573,8 +573,9 @@
     {Toggle::ScalarizeMaxMinClamp,
      {"scalarize_max_min_clamp", "Scalarize max, min, and clamp builtins.",
       "https://crbug.com/422144514", ToggleStage::Device}},
-    {Toggle::MetalEnableModuleConstant,
-     {"metal_enable_module_constant_transform", "Enable the module constant transform.",
+    {Toggle::MetalDisableModuleConstantF16,
+     {"metal_disable_module_constant_f16",
+      "Disable module constant hoisting for values that contain f16 types.",
       "https://crbug.com/419804339", ToggleStage::Device}},
     {Toggle::EnableImmediateErrorHandling,
      {"enable_immediate_error_handling",
diff --git a/src/dawn/native/Toggles.h b/src/dawn/native/Toggles.h
index e939459..4650c45 100644
--- a/src/dawn/native/Toggles.h
+++ b/src/dawn/native/Toggles.h
@@ -140,7 +140,7 @@
     ExposeWGSLExperimentalFeatures,
     DisablePolyfillsOnIntegerDivisonAndModulo,
     ScalarizeMaxMinClamp,
-    MetalEnableModuleConstant,
+    MetalDisableModuleConstantF16,
     EnableImmediateErrorHandling,
     VulkanUseStorageInputOutput16,
     D3D12DontUseShaderModel66OrHigher,
diff --git a/src/dawn/native/metal/PhysicalDeviceMTL.mm b/src/dawn/native/metal/PhysicalDeviceMTL.mm
index 1fa71ae..73c9fef 100644
--- a/src/dawn/native/metal/PhysicalDeviceMTL.mm
+++ b/src/dawn/native/metal/PhysicalDeviceMTL.mm
@@ -442,13 +442,10 @@
         deviceToggles->Default(Toggle::MetalRenderR8RG8UnormSmallMipToTempTexture, true);
     }
 
-    // chromium:419804339: Module constant hoisting is broadly available as a msl transform but
-    // there are execution correction issues with f16 for non apple silicon (Intel/AMD). Therefore
-    // we only enable for apple silicon for now.
-    // chromium:417519810: Mutiple cts tests will fail on AMD if module scope hoisting is not
-    // enabled on AMD. These failures will be internal compiler errors.
-    if (gpu_info::IsApple(vendorId) || gpu_info::IsAMD(vendorId)) {
-        deviceToggles->Default(Toggle::MetalEnableModuleConstant, true);
+    // chromium:419804339: Module constant hoisting is not supported for values containing f16
+    // types on Intel.
+    if (gpu_info::IsIntel(vendorId)) {
+        deviceToggles->Default(Toggle::MetalDisableModuleConstantF16, true);
     }
 
     // On some Intel GPUs vertex only render pipeline get wrong depth result if no fragment
diff --git a/src/dawn/native/metal/ShaderModuleMTL.mm b/src/dawn/native/metal/ShaderModuleMTL.mm
index ec35b39..d5af501 100644
--- a/src/dawn/native/metal/ShaderModuleMTL.mm
+++ b/src/dawn/native/metal/ShaderModuleMTL.mm
@@ -296,8 +296,8 @@
     req.tintOptions.disable_polyfill_integer_div_mod =
         device->IsToggleEnabled(Toggle::DisablePolyfillsOnIntegerDivisonAndModulo);
     req.tintOptions.scalarize_max_min_clamp = device->IsToggleEnabled(Toggle::ScalarizeMaxMinClamp);
-    req.tintOptions.enable_module_constant =
-        device->IsToggleEnabled(Toggle::MetalEnableModuleConstant);
+    req.tintOptions.disable_module_constant_f16 =
+        device->IsToggleEnabled(Toggle::MetalDisableModuleConstantF16);
     req.tintOptions.vertex_pulling_config = std::move(vertexPullingTransformConfig);
     req.tintOptions.enable_integer_range_analysis =
         device->IsToggleEnabled(Toggle::EnableIntegerRangeAnalysisInRobustness);
diff --git a/src/dawn/node/README.md b/src/dawn/node/README.md
index ed325b7..cea704d 100644
--- a/src/dawn/node/README.md
+++ b/src/dawn/node/README.md
@@ -272,7 +272,7 @@
 loop:nested_loops:preventValueOptimizations=false'
 <SNIP>
 Running:
-  Cmd: /home/user/src/dawn/third_party/node/node-linux-x64/bin/node -e "require('./out-node/common/runtime/cmdline.js');" -- placeholder-arg --gpu-provider /home/user/src/dawn/build-clang/cts.js --verbose --quiet --gpu-provider-flag verbose=1 --colors --unroll-const-eval-loops --gpu-provider-flag enable-dawn-features=allow_unsafe_apis "webgpu:shader,execution,flow_control,loop:nested_loops:preventValueOptimizations=false"
+  Cmd: /home/user/src/dawn/third_party/node/node-linux-x64/bin/node -e "require('./out-node/common/runtime/cmdline.js');" -- placeholder-arg --gpu-provider /home/user/src/dawn/build-clang/cts.js --verbose --quiet --gpu-provider-flag verbose=1 --colors --gpu-provider-flag enable-dawn-features=allow_unsafe_apis "webgpu:shader,execution,flow_control,loop:nested_loops:preventValueOptimizations=false"
   Dir: /home/user/src/dawn/third_party/webgpu-cts
 
   For VS Code launch.json:
@@ -289,7 +289,6 @@
         "--gpu-provider-flag",
         "verbose=1",
         "--colors",
-        "--unroll-const-eval-loops",
         "--gpu-provider-flag",
         "enable-dawn-features=allow_unsafe_apis",
         "webgpu:shader,execution,flow_control,loop:nested_loops:preventValueOptimizations=false"
diff --git a/src/tint/lang/msl/writer/common/options.h b/src/tint/lang/msl/writer/common/options.h
index 925dcfb..551ee4c 100644
--- a/src/tint/lang/msl/writer/common/options.h
+++ b/src/tint/lang/msl/writer/common/options.h
@@ -166,8 +166,8 @@
     /// Set to `true` to scalarize max min and clamp builtins.
     bool scalarize_max_min_clamp = false;
 
-    /// Set to `true` to enable the module constant transform
-    bool enable_module_constant = false;
+    /// Set to `true` to disable the module constant transform for f16
+    bool disable_module_constant_f16 = false;
 
     /// Emit argument buffers
     bool use_argument_buffers = false;
@@ -204,7 +204,7 @@
                  emit_vertex_point_size,
                  disable_polyfill_integer_div_mod,
                  scalarize_max_min_clamp,
-                 enable_module_constant,
+                 disable_module_constant_f16,
                  use_argument_buffers,
                  buffer_size_ubo_index,
                  fixed_sample_mask,
diff --git a/src/tint/lang/msl/writer/raise/module_constant.cc b/src/tint/lang/msl/writer/raise/module_constant.cc
index 6c9c8b6..3d31ae4 100644
--- a/src/tint/lang/msl/writer/raise/module_constant.cc
+++ b/src/tint/lang/msl/writer/raise/module_constant.cc
@@ -42,9 +42,32 @@
 /// PIMPL state for the transform.
 struct State {
     core::ir::Module& ir;
+    const ModuleConstantConfig& config;
     core::ir::Builder b{ir};
     core::type::Manager& ty{ir.Types()};
 
+    /// @returns true if @p type is or contains a f16 type
+    bool ContainsF16Type(const core::type::Type* type) {
+        if (type->Is<core::type::F16>()) {
+            return true;
+        }
+
+        if (type->IsScalar()) {
+            return false;
+        }
+
+        if (const auto* str = type->As<core::type::Struct>()) {
+            for (auto* member : str->Members()) {
+                if (ContainsF16Type(member->Type())) {
+                    return true;
+                }
+            }
+            return false;
+        }
+
+        return ContainsF16Type(type->DeepestElement());
+    }
+
     void Process() {
         Hashmap<core::ir::Value*, core::ir::Value*, 16> object_to_var;
         for (auto* inst : ir.Instructions()) {
@@ -64,6 +87,11 @@
             if (!curr_const->Type()->IsAnyOf<core::type::Array, core::type::Struct>()) {
                 continue;
             }
+
+            if (config.disable_module_constant_f16 && ContainsF16Type(curr_const->Type())) {
+                continue;
+            }
+
             // Declare a variable and copy the source object to it.
             auto* var = object_to_var.GetOrAdd(source_object, [&] {
                 // If the source object is a constant we use a module-scope variable
@@ -78,13 +106,13 @@
 
 }  // namespace
 
-Result<SuccessType> ModuleConstant(core::ir::Module& ir) {
+Result<SuccessType> ModuleConstant(core::ir::Module& ir, const ModuleConstantConfig& config) {
     auto result = ValidateAndDumpIfNeeded(ir, "msl.ModuleConstant", kModuleConstantCapabilities);
     if (result != Success) {
         return result;
     }
 
-    State{ir}.Process();
+    State{ir, config}.Process();
 
     return Success;
 }
diff --git a/src/tint/lang/msl/writer/raise/module_constant.h b/src/tint/lang/msl/writer/raise/module_constant.h
index cf8b5ab..33dfd6e 100644
--- a/src/tint/lang/msl/writer/raise/module_constant.h
+++ b/src/tint/lang/msl/writer/raise/module_constant.h
@@ -53,12 +53,18 @@
     core::ir::Capability::kAllowModuleScopeLets,
 };
 
+/// The set of polyfills that should be applied.
+struct ModuleConstantConfig {
+    // Set to true to disable module constant transform on constant data that has any f16.
+    bool disable_module_constant_f16 = false;
+};
+
 /// ModuleConstant is a transform that moves all const data associated with access to a module scope
 /// let. This transform is used to support 'program scope constants' in msl and thereby avoid the
 /// potential for copying of large const in nested loops.
 /// @param module the module to transform
 /// @returns success or failure
-Result<SuccessType> ModuleConstant(core::ir::Module& module);
+Result<SuccessType> ModuleConstant(core::ir::Module& module, const ModuleConstantConfig& config);
 
 }  // namespace tint::msl::writer::raise
 
diff --git a/src/tint/lang/msl/writer/raise/module_constant_test.cc b/src/tint/lang/msl/writer/raise/module_constant_test.cc
index 5e56a28..2a1b783 100644
--- a/src/tint/lang/msl/writer/raise/module_constant_test.cc
+++ b/src/tint/lang/msl/writer/raise/module_constant_test.cc
@@ -82,7 +82,8 @@
 }
 )";
 
-    Run(ModuleConstant);
+    ModuleConstantConfig cfg;
+    Run(ModuleConstant, cfg);
 
     EXPECT_EQ(expect, str());
 }
@@ -119,7 +120,8 @@
   }
 }
 )";
-    Run(ModuleConstant);
+    ModuleConstantConfig cfg;
+    Run(ModuleConstant, cfg);
 
     EXPECT_EQ(expect, str());
 }
@@ -172,7 +174,8 @@
   }
 }
 )";
-    Run(ModuleConstant);
+    ModuleConstantConfig cfg;
+    Run(ModuleConstant, cfg);
 
     EXPECT_EQ(expect, str());
 }
@@ -215,7 +218,8 @@
 }
 )";
 
-    Run(ModuleConstant);
+    ModuleConstantConfig cfg;
+    Run(ModuleConstant, cfg);
 
     EXPECT_EQ(expect, str());
 }
@@ -273,7 +277,109 @@
 }
 )";
 
-    Run(ModuleConstant);
+    ModuleConstantConfig cfg;
+    Run(ModuleConstant, cfg);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_ModuleConstantTest, DisableF16_ConstArrayStruct) {
+    auto* func = b.Function("foo", ty.u32());
+    auto* s = ty.Struct(mod.symbols.New("S"), {
+                                                  {mod.symbols.Register("a"), ty.u32()},
+                                                  {mod.symbols.Register("b"), ty.f16()},
+                                              });
+
+    b.Append(func->Block(), [&] {
+        auto array_struct_type = ty.array(s, 2);
+        auto* c = b.Composite(array_struct_type, b.Splat(s, 1_u), b.Splat(s, 2_u));
+        auto* index = b.Let(1_u);
+        auto* access = b.Access(ty.u32(), c, index, 0_u);
+        auto* r = b.Let("q", access);
+        b.Return(func, r);
+    });
+
+    auto* src = R"(
+S = struct @align(4) {
+  a:u32 @offset(0)
+  b:f16 @offset(4)
+}
+
+%foo = func():u32 {
+  $B1: {
+    %2:u32 = let 1u
+    %3:u32 = access array<S, 2>(S(1u), S(2u)), %2, 0u
+    %q:u32 = let %3
+    ret %q
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = src;
+
+    ModuleConstantConfig cfg{.disable_module_constant_f16 = true};
+    Run(ModuleConstant, cfg);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_ModuleConstantTest, EnableF16_ConstArrayStruct) {
+    auto* func = b.Function("foo", ty.u32());
+    auto* s = ty.Struct(mod.symbols.New("S"), {
+                                                  {mod.symbols.Register("a"), ty.u32()},
+                                                  {mod.symbols.Register("b"), ty.f16()},
+                                              });
+
+    b.Append(func->Block(), [&] {
+        auto array_struct_type = ty.array(s, 2);
+        auto* c = b.Composite(array_struct_type, b.Splat(s, 1_u), b.Splat(s, 2_u));
+        auto* index = b.Let(1_u);
+        auto* access = b.Access(ty.u32(), c, index, 0_u);
+        auto* r = b.Let("q", access);
+        b.Return(func, r);
+    });
+
+    auto* src = R"(
+S = struct @align(4) {
+  a:u32 @offset(0)
+  b:f16 @offset(4)
+}
+
+%foo = func():u32 {
+  $B1: {
+    %2:u32 = let 1u
+    %3:u32 = access array<S, 2>(S(1u), S(2u)), %2, 0u
+    %q:u32 = let %3
+    ret %q
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+S = struct @align(4) {
+  a:u32 @offset(0)
+  b:f16 @offset(4)
+}
+
+$B1: {  # root
+  %1:array<S, 2> = let array<S, 2>(S(1u), S(2u))
+}
+
+%foo = func():u32 {
+  $B2: {
+    %3:u32 = let 1u
+    %4:u32 = access %1, %3, 0u
+    %q:u32 = let %4
+    ret %q
+  }
+}
+)";
+
+    // Note the disable f16 is false by default.
+    ModuleConstantConfig cfg;
+    Run(ModuleConstant, cfg);
 
     EXPECT_EQ(expect, str());
 }
@@ -331,7 +437,49 @@
 }
 )";
 
-    Run(ModuleConstant);
+    ModuleConstantConfig cfg;
+    Run(ModuleConstant, cfg);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_ModuleConstantTest, DisableF16_ConstStructArray) {
+    auto* func = b.Function("foo", ty.u32());
+    auto* s =
+        ty.Struct(mod.symbols.New("S"), {
+                                            {mod.symbols.Register("a"), ty.array(ty.u32(), 2)},
+                                            {mod.symbols.Register("b"), ty.array(ty.f16(), 2)},
+                                        });
+
+    b.Append(func->Block(), [&] {
+        auto* c = b.Splat(s, 1_u);
+        auto* index = b.Let(1_u);
+        auto* access = b.Access(ty.u32(), c, 0_u, index);
+        auto* r = b.Let("q", access);
+        b.Return(func, r);
+    });
+
+    auto* src = R"(
+S = struct @align(4) {
+  a:array<u32, 2> @offset(0)
+  b:array<f16, 2> @offset(8)
+}
+
+%foo = func():u32 {
+  $B1: {
+    %2:u32 = let 1u
+    %3:u32 = access S(1u), 0u, %2
+    %q:u32 = let %3
+    ret %q
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = src;
+
+    ModuleConstantConfig cfg{.disable_module_constant_f16 = true};
+    Run(ModuleConstant, cfg);
 
     EXPECT_EQ(expect, str());
 }
@@ -377,7 +525,211 @@
 }
 )";
 
-    Run(ModuleConstant);
+    ModuleConstantConfig cfg;
+    Run(ModuleConstant, cfg);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_ModuleConstantTest, DisableF16_ConstVecArrayStruct) {
+    auto* func = b.Function("foo", ty.f16());
+    auto* s = ty.Struct(mod.symbols.New("S"), {
+                                                  {mod.symbols.Register("a"), ty.vec2(ty.f16())},
+                                                  {mod.symbols.Register("b"), ty.u32()},
+                                              });
+
+    b.Append(func->Block(), [&] {
+        auto array_struct_type = ty.array(s, 2);
+        auto* c = b.Composite(array_struct_type, b.Splat(s, 1_u), b.Splat(s, 2_u));
+        auto* index = b.Let(1_u);
+        auto* access = b.Access(ty.f16(), c, index, 0_u, 0_u);
+        auto* r = b.Let("q", access);
+        b.Return(func, r);
+    });
+
+    auto* src = R"(
+S = struct @align(4) {
+  a:vec2<f16> @offset(0)
+  b:u32 @offset(4)
+}
+
+%foo = func():f16 {
+  $B1: {
+    %2:u32 = let 1u
+    %3:f16 = access array<S, 2>(S(1u), S(2u)), %2, 0u, 0u
+    %q:f16 = let %3
+    ret %q
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = src;
+
+    ModuleConstantConfig cfg{.disable_module_constant_f16 = true};
+    Run(ModuleConstant, cfg);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_ModuleConstantTest, EnableF16_ConstVecArrayStruct) {
+    auto* func = b.Function("foo", ty.f16());
+    auto* s = ty.Struct(mod.symbols.New("S"), {
+                                                  {mod.symbols.Register("a"), ty.vec2(ty.f16())},
+                                                  {mod.symbols.Register("b"), ty.f16()},
+                                              });
+
+    b.Append(func->Block(), [&] {
+        auto array_struct_type = ty.array(s, 2);
+        auto* c = b.Composite(array_struct_type, b.Splat(s, 1_u), b.Splat(s, 2_u));
+        auto* index = b.Let(1_u);
+        auto* access = b.Access(ty.f16(), c, index, 0_u, 0_u);
+        auto* r = b.Let("q", access);
+        b.Return(func, r);
+    });
+
+    auto* src = R"(
+S = struct @align(4) {
+  a:vec2<f16> @offset(0)
+  b:f16 @offset(4)
+}
+
+%foo = func():f16 {
+  $B1: {
+    %2:u32 = let 1u
+    %3:f16 = access array<S, 2>(S(1u), S(2u)), %2, 0u, 0u
+    %q:f16 = let %3
+    ret %q
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+S = struct @align(4) {
+  a:vec2<f16> @offset(0)
+  b:f16 @offset(4)
+}
+
+$B1: {  # root
+  %1:array<S, 2> = let array<S, 2>(S(1u), S(2u))
+}
+
+%foo = func():f16 {
+  $B2: {
+    %3:u32 = let 1u
+    %4:f16 = access %1, %3, 0u, 0u
+    %q:f16 = let %4
+    ret %q
+  }
+}
+)";
+
+    // Note the disable f16 is false by default.
+    ModuleConstantConfig cfg;
+    Run(ModuleConstant, cfg);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_ModuleConstantTest, DisableF16_ConstMatStructArray) {
+    auto* func = b.Function("foo", ty.f16());
+    auto* s = ty.Struct(mod.symbols.New("S"),
+                        {
+                            {mod.symbols.Register("a"), ty.array(ty.u32(), 2)},
+                            {mod.symbols.Register("b"), ty.array(ty.mat3x3(ty.f16()), 3)},
+                        });
+
+    b.Append(func->Block(), [&] {
+        auto* c = b.Splat(s, 1_u);
+        auto* index = b.Let(1_u);
+        auto* access = b.Access(ty.f16(), c, 1_u, index, 0_u, 0_u);
+        auto* r = b.Let("q", access);
+        b.Return(func, r);
+    });
+
+    auto* src = R"(
+S = struct @align(8) {
+  a:array<u32, 2> @offset(0)
+  b:array<mat3x3<f16>, 3> @offset(8)
+}
+
+%foo = func():f16 {
+  $B1: {
+    %2:u32 = let 1u
+    %3:f16 = access S(1u), 1u, %2, 0u, 0u
+    %q:f16 = let %3
+    ret %q
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = src;
+
+    ModuleConstantConfig cfg{.disable_module_constant_f16 = true};
+    Run(ModuleConstant, cfg);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_ModuleConstantTest, EnabledF16_ConstMatStructArray) {
+    auto* func = b.Function("foo", ty.f16());
+    auto* s = ty.Struct(mod.symbols.New("S"),
+                        {
+                            {mod.symbols.Register("a"), ty.array(ty.u32(), 2)},
+                            {mod.symbols.Register("b"), ty.array(ty.mat3x3(ty.f16()), 3)},
+                        });
+
+    b.Append(func->Block(), [&] {
+        auto* c = b.Splat(s, 1_u);
+        auto* index = b.Let(1_u);
+        auto* access = b.Access(ty.f16(), c, 1_u, index, 0_u, 0_u);
+        auto* r = b.Let("q", access);
+        b.Return(func, r);
+    });
+
+    auto* src = R"(
+S = struct @align(8) {
+  a:array<u32, 2> @offset(0)
+  b:array<mat3x3<f16>, 3> @offset(8)
+}
+
+%foo = func():f16 {
+  $B1: {
+    %2:u32 = let 1u
+    %3:f16 = access S(1u), 1u, %2, 0u, 0u
+    %q:f16 = let %3
+    ret %q
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+S = struct @align(8) {
+  a:array<u32, 2> @offset(0)
+  b:array<mat3x3<f16>, 3> @offset(8)
+}
+
+$B1: {  # root
+  %1:S = let S(1u)
+}
+
+%foo = func():f16 {
+  $B2: {
+    %3:u32 = let 1u
+    %4:f16 = access %1, 1u, %3, 0u, 0u
+    %q:f16 = let %4
+    ret %q
+  }
+}
+)";
+
+    EXPECT_EQ(src, str());
+    // Note the disable f16 is false by default.
+    ModuleConstantConfig cfg;
+    Run(ModuleConstant, cfg);
 
     EXPECT_EQ(expect, str());
 }
diff --git a/src/tint/lang/msl/writer/raise/raise.cc b/src/tint/lang/msl/writer/raise/raise.cc
index 2a59e32..aead774 100644
--- a/src/tint/lang/msl/writer/raise/raise.cc
+++ b/src/tint/lang/msl/writer/raise/raise.cc
@@ -168,9 +168,8 @@
     };
     RUN_TRANSFORM(core::ir::transform::BuiltinScalarize, module, scalarize_config);
 
-    if (options.enable_module_constant) {
-        RUN_TRANSFORM(raise::ModuleConstant, module);
-    }
+    raise::ModuleConstantConfig module_const_config{options.disable_module_constant_f16};
+    RUN_TRANSFORM(raise::ModuleConstant, module, module_const_config);
 
     // These transforms need to be run last as various transforms introduce terminator arguments,
     // naming conflicts, and expressions that need to be explicitly not inlined.
diff --git a/test/tint/bug/tint/1641.wgsl.expected.msl b/test/tint/bug/tint/1641.wgsl.expected.msl
index 1525b08..d0ac73b 100644
--- a/test/tint/bug/tint/1641.wgsl.expected.msl
+++ b/test/tint/bug/tint/1641.wgsl.expected.msl
@@ -20,13 +20,14 @@
 struct main_outputs {
   float4 tint_symbol [[position]];
 };
+const constant tint_array<Normals, 1> v = tint_array<Normals, 1>{Normals{.f=float3(0.0f, 0.0f, 1.0f)}};
 
 float4 main_inner() {
   int const zero = 0;
-  return float4(tint_array<Normals, 1>{Normals{.f=float3(0.0f, 0.0f, 1.0f)}}[min(uint(zero), 0u)].f, 1.0f);
+  return float4(v[min(uint(zero), 0u)].f, 1.0f);
 }
 
-vertex main_outputs v() {
+vertex main_outputs v_1() {
   main_outputs tint_wrapper_result = {};
   tint_wrapper_result.tint_symbol = main_inner();
   return tint_wrapper_result;
diff --git a/test/tint/bug/tint/1653.wgsl.expected.msl b/test/tint/bug/tint/1653.wgsl.expected.msl
index 3c39813..317c598 100644
--- a/test/tint/bug/tint/1653.wgsl.expected.msl
+++ b/test/tint/bug/tint/1653.wgsl.expected.msl
@@ -16,9 +16,10 @@
 struct vs_main_outputs {
   float4 tint_symbol [[position]];
 };
+const constant tint_array<float4, 3> v = tint_array<float4, 3>{float4(0.0f, 0.0f, 0.0f, 1.0f), float4(0.0f, 1.0f, 0.0f, 1.0f), float4(1.0f, 1.0f, 0.0f, 1.0f)};
 
 float4 vs_main_inner(uint in_vertex_index) {
-  return tint_array<float4, 3>{float4(0.0f, 0.0f, 0.0f, 1.0f), float4(0.0f, 1.0f, 0.0f, 1.0f), float4(1.0f, 1.0f, 0.0f, 1.0f)}[min(in_vertex_index, 2u)];
+  return v[min(in_vertex_index, 2u)];
 }
 
 vertex vs_main_outputs vs_main(uint in_vertex_index [[vertex_id]]) {
diff --git a/test/tint/bug/tint/2237.wgsl.expected.msl b/test/tint/bug/tint/2237.wgsl.expected.msl
index f9846ac..3301641 100644
--- a/test/tint/bug/tint/2237.wgsl.expected.msl
+++ b/test/tint/bug/tint/2237.wgsl.expected.msl
@@ -1,10 +1,6 @@
 #include <metal_stdlib>
 using namespace metal;
 
-struct tint_module_vars_struct {
-  device uint* tint_member;
-};
-
 template<typename T, size_t N>
 struct tint_array {
   const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -17,12 +13,17 @@
   T elements[N];
 };
 
+struct tint_module_vars_struct {
+  device uint* tint_member;
+};
+const constant tint_array<uint, 4> v_1 = tint_array<uint, 4>{0u, 1u, 2u, 4u};
+
 uint foo(tint_module_vars_struct tint_module_vars) {
-  return tint_array<uint, 4>{0u, 1u, 2u, 4u}[min((*tint_module_vars.tint_member), 3u)];
+  return v_1[min((*tint_module_vars.tint_member), 3u)];
 }
 
-kernel void v_1(device uint* v_2 [[buffer(0)]]) {
-  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.tint_member=v_2};
-  uint const v = tint_array<uint, 4>{0u, 1u, 2u, 4u}[min((*tint_module_vars.tint_member), 3u)];
+kernel void v_2(device uint* v_3 [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.tint_member=v_3};
+  uint const v = v_1[min((*tint_module_vars.tint_member), 3u)];
   (*tint_module_vars.tint_member) = (v + foo(tint_module_vars));
 }
diff --git a/test/tint/const/array/array_matrix_f16.wgsl.expected.msl b/test/tint/const/array/array_matrix_f16.wgsl.expected.msl
index a798803..223f418 100644
--- a/test/tint/const/array/array_matrix_f16.wgsl.expected.msl
+++ b/test/tint/const/array/array_matrix_f16.wgsl.expected.msl
@@ -21,15 +21,16 @@
 struct tint_array_lengths_struct {
   uint tint_array_length_0_0;
 };
+const constant tint_array<half3x2, 2> v = tint_array<half3x2, 2>{half3x2(half2(0.0h, 1.0h), half2(2.0h, 3.0h), half2(2.0h, 3.0h)), half3x2(half2(0.0h, 1.0h), half2(2.0h, 3.0h), half2(2.0h, 3.0h))};
 
 uint tint_f16_to_u32(half value) {
   return uint(clamp(value, 0.0h, 65504.0h));
 }
 
-kernel void v(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
   uint q = 0u;
-  uint const v_1 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
-  device uint* const v_2 = (&(*tint_module_vars.s)[min(uint(0), v_1)]);
-  (*v_2) = tint_f16_to_u32(tint_array<half3x2, 2>{half3x2(half2(0.0h, 1.0h), half2(2.0h, 3.0h), half2(2.0h, 3.0h)), half3x2(half2(0.0h, 1.0h), half2(2.0h, 3.0h), half2(2.0h, 3.0h))}[min(q, 1u)][0u].x);
+  uint const v_2 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
+  device uint* const v_3 = (&(*tint_module_vars.s)[min(uint(0), v_2)]);
+  (*v_3) = tint_f16_to_u32(v[min(q, 1u)][0u].x);
 }
diff --git a/test/tint/const/array/array_of_array_of_array.wgsl.expected.msl b/test/tint/const/array/array_of_array_of_array.wgsl.expected.msl
index cf141d3..8c998ac 100644
--- a/test/tint/const/array/array_of_array_of_array.wgsl.expected.msl
+++ b/test/tint/const/array/array_of_array_of_array.wgsl.expected.msl
@@ -21,11 +21,12 @@
 struct tint_array_lengths_struct {
   uint tint_array_length_0_0;
 };
+const constant tint_array<tint_array<tint_array<uint, 2>, 2>, 2> v = tint_array<tint_array<tint_array<uint, 2>, 2>, 2>{tint_array<tint_array<uint, 2>, 2>{tint_array<uint, 2>{0u, 1u}, tint_array<uint, 2>{2u, 3u}}, tint_array<tint_array<uint, 2>, 2>{tint_array<uint, 2>{4u, 5u}, tint_array<uint, 2>{6u, 7u}}};
 
-kernel void v(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
   uint q = 0u;
-  uint const v_1 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
-  device uint* const v_2 = (&(*tint_module_vars.s)[min(uint(0), v_1)]);
-  (*v_2) = tint_array<tint_array<tint_array<uint, 2>, 2>, 2>{tint_array<tint_array<uint, 2>, 2>{tint_array<uint, 2>{0u, 1u}, tint_array<uint, 2>{2u, 3u}}, tint_array<tint_array<uint, 2>, 2>{tint_array<uint, 2>{4u, 5u}, tint_array<uint, 2>{6u, 7u}}}[min(q, 1u)][min(q, 1u)][min(q, 1u)];
+  uint const v_2 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
+  device uint* const v_3 = (&(*tint_module_vars.s)[min(uint(0), v_2)]);
+  (*v_3) = v[min(q, 1u)][min(q, 1u)][min(q, 1u)];
 }
diff --git a/test/tint/const/array/array_of_struct.wgsl.expected.msl b/test/tint/const/array/array_of_struct.wgsl.expected.msl
index d81bcdd..e8677fb 100644
--- a/test/tint/const/array/array_of_struct.wgsl.expected.msl
+++ b/test/tint/const/array/array_of_struct.wgsl.expected.msl
@@ -13,6 +13,11 @@
   T elements[N];
 };
 
+struct A {
+  uint2 b;
+  uint c;
+};
+
 struct tint_module_vars_struct {
   device tint_array<uint, 1>* s;
   const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
@@ -21,16 +26,12 @@
 struct tint_array_lengths_struct {
   uint tint_array_length_0_0;
 };
+const constant tint_array<A, 2> v = tint_array<A, 2>{A{.b=uint2(1u, 2u), .c=3u}, A{.b=uint2(4u, 5u), .c=6u}};
 
-struct A {
-  uint2 b;
-  uint c;
-};
-
-kernel void v(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
   uint q = 0u;
-  uint const v_1 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
-  device uint* const v_2 = (&(*tint_module_vars.s)[min(uint(0), v_1)]);
-  (*v_2) = tint_array<A, 2>{A{.b=uint2(1u, 2u), .c=3u}, A{.b=uint2(4u, 5u), .c=6u}}[min(q, 1u)].b.x;
+  uint const v_2 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
+  device uint* const v_3 = (&(*tint_module_vars.s)[min(uint(0), v_2)]);
+  (*v_3) = v[min(q, 1u)].b.x;
 }
diff --git a/test/tint/const/array/array_of_struct_exp_val.wgsl.expected.msl b/test/tint/const/array/array_of_struct_exp_val.wgsl.expected.msl
index d81bcdd..e8677fb 100644
--- a/test/tint/const/array/array_of_struct_exp_val.wgsl.expected.msl
+++ b/test/tint/const/array/array_of_struct_exp_val.wgsl.expected.msl
@@ -13,6 +13,11 @@
   T elements[N];
 };
 
+struct A {
+  uint2 b;
+  uint c;
+};
+
 struct tint_module_vars_struct {
   device tint_array<uint, 1>* s;
   const constant tint_array<uint4, 1>* tint_storage_buffer_sizes;
@@ -21,16 +26,12 @@
 struct tint_array_lengths_struct {
   uint tint_array_length_0_0;
 };
+const constant tint_array<A, 2> v = tint_array<A, 2>{A{.b=uint2(1u, 2u), .c=3u}, A{.b=uint2(4u, 5u), .c=6u}};
 
-struct A {
-  uint2 b;
-  uint c;
-};
-
-kernel void v(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
   uint q = 0u;
-  uint const v_1 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
-  device uint* const v_2 = (&(*tint_module_vars.s)[min(uint(0), v_1)]);
-  (*v_2) = tint_array<A, 2>{A{.b=uint2(1u, 2u), .c=3u}, A{.b=uint2(4u, 5u), .c=6u}}[min(q, 1u)].b.x;
+  uint const v_2 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
+  device uint* const v_3 = (&(*tint_module_vars.s)[min(uint(0), v_2)]);
+  (*v_3) = v[min(q, 1u)].b.x;
 }
diff --git a/test/tint/const/array/large_array.wgsl.expected.msl b/test/tint/const/array/large_array.wgsl.expected.msl
index ba6f9f3..c61cb31 100644
--- a/test/tint/const/array/large_array.wgsl.expected.msl
+++ b/test/tint/const/array/large_array.wgsl.expected.msl
@@ -21,11 +21,12 @@
 struct tint_array_lengths_struct {
   uint tint_array_length_0_0;
 };
+const constant tint_array<uint, 2310> v = tint_array<uint, 2310>{0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u};
 
-kernel void v(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
   uint q = 0u;
-  uint const v_1 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
-  device uint* const v_2 = (&(*tint_module_vars.s)[min(uint(0), v_1)]);
-  (*v_2) = tint_array<uint, 2310>{0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u, 0u, 1u, 2u, 3u, 4u, 5u, 6u, 7u, 8u, 9u, 10u}[min(q, 2309u)];
+  uint const v_2 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
+  device uint* const v_3 = (&(*tint_module_vars.s)[min(uint(0), v_2)]);
+  (*v_3) = v[min(q, 2309u)];
 }
diff --git a/test/tint/const/struct/struct_array_exp_value.wgsl.expected.msl b/test/tint/const/struct/struct_array_exp_value.wgsl.expected.msl
index 6ba612c..407d525 100644
--- a/test/tint/const/struct/struct_array_exp_value.wgsl.expected.msl
+++ b/test/tint/const/struct/struct_array_exp_value.wgsl.expected.msl
@@ -21,11 +21,12 @@
 struct tint_array_lengths_struct {
   uint tint_array_length_0_0;
 };
+const constant tint_array<uint2, 2> v = tint_array<uint2, 2>{uint2(1u, 2u), uint2(3u, 4u)};
 
-kernel void v(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
   uint q = 0u;
-  uint const v_1 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
-  device uint* const v_2 = (&(*tint_module_vars.s)[min(uint(0), v_1)]);
-  (*v_2) = tint_array<uint2, 2>{uint2(1u, 2u), uint2(3u, 4u)}[min(q, 1u)].x;
+  uint const v_2 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
+  device uint* const v_3 = (&(*tint_module_vars.s)[min(uint(0), v_2)]);
+  (*v_3) = v[min(q, 1u)].x;
 }
diff --git a/test/tint/const/struct/struct_of_array.wgsl.expected.msl b/test/tint/const/struct/struct_of_array.wgsl.expected.msl
index 6ba612c..407d525 100644
--- a/test/tint/const/struct/struct_of_array.wgsl.expected.msl
+++ b/test/tint/const/struct/struct_of_array.wgsl.expected.msl
@@ -21,11 +21,12 @@
 struct tint_array_lengths_struct {
   uint tint_array_length_0_0;
 };
+const constant tint_array<uint2, 2> v = tint_array<uint2, 2>{uint2(1u, 2u), uint2(3u, 4u)};
 
-kernel void v(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
+kernel void v_1(device tint_array<uint, 1>* s [[buffer(0)]], const constant tint_array<uint4, 1>* tint_storage_buffer_sizes [[buffer(30)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .tint_storage_buffer_sizes=tint_storage_buffer_sizes};
   uint q = 0u;
-  uint const v_1 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
-  device uint* const v_2 = (&(*tint_module_vars.s)[min(uint(0), v_1)]);
-  (*v_2) = tint_array<uint2, 2>{uint2(1u, 2u), uint2(3u, 4u)}[min(q, 1u)].x;
+  uint const v_2 = (tint_array_lengths_struct{.tint_array_length_0_0=((*tint_module_vars.tint_storage_buffer_sizes)[0u].x / 4u)}.tint_array_length_0_0 - 1u);
+  device uint* const v_3 = (&(*tint_module_vars.s)[min(uint(0), v_2)]);
+  (*v_3) = v[min(q, 1u)].x;
 }
diff --git a/test/tint/samples/triangle.wgsl.expected.msl b/test/tint/samples/triangle.wgsl.expected.msl
index b8152ed..469449e 100644
--- a/test/tint/samples/triangle.wgsl.expected.msl
+++ b/test/tint/samples/triangle.wgsl.expected.msl
@@ -19,9 +19,10 @@
 struct vtx_main_outputs {
   float4 tint_symbol [[position]];
 };
+const constant tint_array<float2, 3> v = tint_array<float2, 3>{float2(0.0f, 0.5f), float2(-0.5f), float2(0.5f, -0.5f)};
 
 float4 vtx_main_inner(uint VertexIndex) {
-  return float4(tint_array<float2, 3>{float2(0.0f, 0.5f), float2(-0.5f), float2(0.5f, -0.5f)}[min(VertexIndex, 2u)], 0.0f, 1.0f);
+  return float4(v[min(VertexIndex, 2u)], 0.0f, 1.0f);
 }
 
 vertex vtx_main_outputs vtx_main(uint VertexIndex [[vertex_id]]) {
diff --git a/tools/src/cmd/run-cts/node/cmd.go b/tools/src/cmd/run-cts/node/cmd.go
index f7fb281..960c6da 100644
--- a/tools/src/cmd/run-cts/node/cmd.go
+++ b/tools/src/cmd/run-cts/node/cmd.go
@@ -87,7 +87,7 @@
 }
 
 func (c *cmd) RegisterFlags(ctx context.Context, cfg common.Config) ([]string, error) {
-	unrollConstEvalLoopsDefault := runtime.GOOS != "windows"
+	unrollConstEvalLoopsDefault := false
 
 	backendDefault := "default"
 	if vkIcdFilenames := cfg.OsWrapper.Getenv("VK_ICD_FILENAMES"); vkIcdFilenames != "" {
diff --git a/webgpu-cts/compat-expectations.txt b/webgpu-cts/compat-expectations.txt
index 463b443..9596069 100644
--- a/webgpu-cts/compat-expectations.txt
+++ b/webgpu-cts/compat-expectations.txt
@@ -610,6 +610,14 @@
 crbug.com/373670504 [ android arm ] webgpu:shader,execution,expression,constructor,non_zero:concrete_array_elements:inputSource="uniform";type="bool";length=5;infer_type=false [ Failure ]
 crbug.com/373670504 [ android arm ] webgpu:shader,execution,expression,constructor,non_zero:concrete_array_elements:inputSource="uniform";type="bool";length=5;infer_type=true [ Failure ]
 
+# Const bool arrays on Pixel 6
+crbug.com/422510023 [ android arm ] webgpu:shader,execution,expression,constructor,non_zero:concrete_array_elements:inputSource="const";type="bool";length=10;infer_type=false [ Failure ]
+crbug.com/422510023 [ android arm ] webgpu:shader,execution,expression,constructor,non_zero:concrete_array_elements:inputSource="const";type="bool";length=10;infer_type=true [ Failure ]
+crbug.com/422510023 [ android arm ] webgpu:shader,execution,expression,constructor,non_zero:concrete_array_elements:inputSource="const";type="bool";length=5;infer_type=false [ Failure ]
+crbug.com/422510023 [ android arm ] webgpu:shader,execution,expression,constructor,non_zero:concrete_array_elements:inputSource="const";type="bool";length=5;infer_type=true [ Failure ]
+crbug.com/422510023 [ android arm ] webgpu:shader,execution,expression,constructor,zero_value:array:type="bool";length=10 [ Failure ]
+crbug.com/422510023 [ android arm ] webgpu:shader,execution,expression,constructor,zero_value:array:type="bool";length=5 [ Failure ]
+
 # set_index_buffer_without_changing_buffer failing on Pixel 6
 crbug.com/373664074 [ android-t arm ] webgpu:api,operation,command_buffer,render,state_tracking:set_index_buffer_without_changing_buffer: [ Failure ]
 
diff --git a/webgpu-cts/test_runner.js b/webgpu-cts/test_runner.js
index 5992de8..aa977ce 100644
--- a/webgpu-cts/test_runner.js
+++ b/webgpu-cts/test_runner.js
@@ -167,14 +167,9 @@
 // simultaneously on a 32-bit system easily runs out of memory).
 globalTestConfig.maxSubcasesInFlight = 100;
 
-// FXC is very slow to compile unrolled const-eval loops, where the metal shader
-// compiler (Intel GPU) is very slow to compile rolled loops. Intel drivers for
-// linux may also suffer the same performance issues, so unroll const-eval loops
-// if we're not running on Windows.
-const isWindows = navigator.userAgent.includes("Windows");
-if (!isWindows) {
-  globalTestConfig.unrollConstEvalLoops = true;
-}
+
+
+
 
 let lastOptionsKey, testWorker;