diff --git a/src/tint/lang/msl/writer/printer/var_test.cc b/src/tint/lang/msl/writer/printer/var_test.cc
index b91c2bd..4b3c2f9 100644
--- a/src/tint/lang/msl/writer/printer/var_test.cc
+++ b/src/tint/lang/msl/writer/printer/var_test.cc
@@ -282,8 +282,7 @@
 )");
 }
 
-// TODO(jrprice): Requires ModuleScopeVarToEntryPointParam transform
-TEST_F(MslPrinterTest, DISABLED_VarGlobalWorkgroup) {
+TEST_F(MslPrinterTest, VarGlobalWorkgroup) {
     core::ir::Var* v = nullptr;
     b.Append(mod.root_block,
              [&] { v = b.Var("v", ty.ptr<core::AddressSpace::kWorkgroup, f32>()); });
@@ -297,10 +296,12 @@
     });
 
     ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
-threadgroup float v;
-void foo() {
-  float a = v;
+    EXPECT_EQ(output_, MetalHeader() + R"(struct tint_module_vars_struct {
+  threadgroup float* v;
+};
+
+void foo(tint_module_vars_struct tint_module_vars) {
+  float a = (*tint_module_vars.v);
 }
 )");
 }
diff --git a/src/tint/lang/msl/writer/raise/module_scope_vars.cc b/src/tint/lang/msl/writer/raise/module_scope_vars.cc
index 00a569f..6a4ecb2 100644
--- a/src/tint/lang/msl/writer/raise/module_scope_vars.cc
+++ b/src/tint/lang/msl/writer/raise/module_scope_vars.cc
@@ -174,6 +174,9 @@
         core::ir::Function* func,
         const core::ir::ReferencedModuleVars::VarSet& referenced_vars) {
         core::ir::Value* module_var_struct = nullptr;
+        core::ir::FunctionParam* workgroup_allocation_param = nullptr;
+        Vector<core::type::Manager::StructMemberDesc, 4> workgroup_struct_members;
+
         // Add parameters and insert instruction at the top of the entry point to set up the
         // module-scope variables structure.
         b.InsertBefore(func->Block()->Front(), [&] {  //
@@ -206,6 +209,24 @@
                         decl = param;
                         break;
                     }
+                    case core::AddressSpace::kWorkgroup: {
+                        // Workgroup variables are received as a function parameter (to workaround
+                        // an MSL compiler bug with threadgroup matrices), and we aggregate all
+                        // workgroup variables into a structure to avoid hitting MSL's limit for
+                        // threadgroup memory arguments.
+                        if (!workgroup_allocation_param) {
+                            workgroup_allocation_param = b.FunctionParam(nullptr);
+                            func->AppendParam(workgroup_allocation_param);
+                        }
+                        decl = b.Access(ptr, workgroup_allocation_param,
+                                        u32(workgroup_struct_members.Length()))
+                                   ->Result(0);
+                        workgroup_struct_members.Push(core::type::Manager::StructMemberDesc{
+                            ir.symbols.New(),
+                            ptr->StoreType(),
+                        });
+                        break;
+                    }
                     case core::AddressSpace::kHandle: {
                         // Handle types become function parameters and drop the pointer.
                         auto* param = b.FunctionParam(ptr->UnwrapPtr());
@@ -230,6 +251,14 @@
             auto* construct = b.Construct(struct_type, std::move(construct_args));
             module_var_struct = b.Let(kModuleVarsName, construct)->Result(0);
         });
+
+        // Create the workgroup variable structure if needed.
+        if (!workgroup_struct_members.IsEmpty()) {
+            auto* workgroup_struct =
+                ty.Struct(ir.symbols.New(), std::move(workgroup_struct_members));
+            workgroup_allocation_param->SetType(ty.ptr<workgroup>(workgroup_struct));
+        }
+
         return module_var_struct;
     }
 
diff --git a/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc b/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
index 3bef228..cf3fe76 100644
--- a/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
+++ b/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
@@ -364,6 +364,73 @@
     EXPECT_EQ(expect, str());
 }
 
+TEST_F(MslWriter_ModuleScopeVarsTest, Workgroup) {
+    auto* var_a = b.Var("a", ty.ptr<workgroup, i32>());
+    auto* var_b = b.Var("b", ty.ptr<workgroup, i32>());
+    mod.root_block->Append(var_a);
+    mod.root_block->Append(var_b);
+
+    auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kCompute,
+                            std::array<uint32_t, 3>{1u, 1u, 1u});
+    b.Append(func->Block(), [&] {
+        auto* load_a = b.Load(var_a);
+        auto* load_b = b.Load(var_b);
+        b.Store(var_a, b.Add<i32>(load_a, load_b));
+        b.Return(func);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %a:ptr<workgroup, i32, read_write> = var
+  %b:ptr<workgroup, i32, read_write> = var
+}
+
+%foo = @compute @workgroup_size(1, 1, 1) func():void {
+  $B2: {
+    %4:i32 = load %a
+    %5:i32 = load %b
+    %6:i32 = add %4, %5
+    store %a, %6
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+tint_module_vars_struct = struct @align(1) {
+  a:ptr<workgroup, i32, read_write> @offset(0)
+  b:ptr<workgroup, i32, read_write> @offset(0)
+}
+
+tint_symbol_2 = struct @align(4) {
+  tint_symbol:i32 @offset(0)
+  tint_symbol_1:i32 @offset(4)
+}
+
+%foo = @compute @workgroup_size(1, 1, 1) func(%2:ptr<workgroup, tint_symbol_2, read_write>):void {
+  $B1: {
+    %a:ptr<workgroup, i32, read_write> = access %2, 0u
+    %b:ptr<workgroup, i32, read_write> = access %2, 1u
+    %5:tint_module_vars_struct = construct %a, %b
+    %tint_module_vars:tint_module_vars_struct = let %5
+    %7:ptr<workgroup, i32, read_write> = access %tint_module_vars, 0u
+    %8:i32 = load %7
+    %9:ptr<workgroup, i32, read_write> = access %tint_module_vars, 1u
+    %10:i32 = load %9
+    %11:i32 = add %8, %10
+    %12:ptr<workgroup, i32, read_write> = access %tint_module_vars, 0u
+    store %12, %11
+    ret
+  }
+}
+)";
+
+    Run(ModuleScopeVars);
+
+    EXPECT_EQ(expect, str());
+}
+
 TEST_F(MslWriter_ModuleScopeVarsTest, MultipleAddressSpaces) {
     auto* var_a = b.Var("a", ty.ptr<uniform, i32, core::Access::kRead>());
     auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
index 37f0e9b..db158da 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  half2x2 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant half2x2* u;
+  threadgroup half2x2* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat2x2<f16>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat2x2<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat2x2<f16>(vec2<f16>(0.0h))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat2x2<f16> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec2<f16>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec2<f16>, read> = access %u, 0i
-    %10:vec2<f16> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec2<f16>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec2<f16>, read> = access %u, 0i
-    %13:vec2<f16> = load %12
-    %14:vec2<f16> = swizzle %13, yx
-    store %11, %14
-    %15:ptr<workgroup, vec2<f16>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec2<f16>, read> = access %u, 1i
-    %17:f16 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half2x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = half2x2(half2(0.0h), half2(0.0h));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].yx;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
index 363c971..938d418 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  float2x2 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant float2x2* u;
+  threadgroup float2x2* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat2x2<f32>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat2x2<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat2x2<f32>(vec2<f32>(0.0f))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat2x2<f32> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec2<f32>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec2<f32>, read> = access %u, 0i
-    %10:vec2<f32> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec2<f32>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec2<f32>, read> = access %u, 0i
-    %13:vec2<f32> = load %12
-    %14:vec2<f32> = swizzle %13, yx
-    store %11, %14
-    %15:ptr<workgroup, vec2<f32>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec2<f32>, read> = access %u, 1i
-    %17:f32 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float2x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = float2x2(float2(0.0f), float2(0.0f));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].yx;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
index 2a286d3..f0e615f 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  half2x3 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant half2x3* u;
+  threadgroup half2x3* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat2x3<f16>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat2x3<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat2x3<f16>(vec3<f16>(0.0h))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat2x3<f16> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec3<f16>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec3<f16>, read> = access %u, 0i
-    %10:vec3<f16> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec3<f16>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec3<f16>, read> = access %u, 0i
-    %13:vec3<f16> = load %12
-    %14:vec3<f16> = swizzle %13, zxy
-    store %11, %14
-    %15:ptr<workgroup, vec3<f16>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec3<f16>, read> = access %u, 1i
-    %17:f16 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half2x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = half2x3(half3(0.0h), half3(0.0h));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].zxy;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
index 93c51a4..bb2ee0c 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  float2x3 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant float2x3* u;
+  threadgroup float2x3* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat2x3<f32>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat2x3<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat2x3<f32>(vec3<f32>(0.0f))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat2x3<f32> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec3<f32>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec3<f32>, read> = access %u, 0i
-    %10:vec3<f32> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec3<f32>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec3<f32>, read> = access %u, 0i
-    %13:vec3<f32> = load %12
-    %14:vec3<f32> = swizzle %13, zxy
-    store %11, %14
-    %15:ptr<workgroup, vec3<f32>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec3<f32>, read> = access %u, 1i
-    %17:f32 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float2x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = float2x3(float3(0.0f), float3(0.0f));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].zxy;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
index 9a8c6ea..85389c0 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  half2x4 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant half2x4* u;
+  threadgroup half2x4* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat2x4<f16>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat2x4<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat2x4<f16>(vec4<f16>(0.0h))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat2x4<f16> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec4<f16>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec4<f16>, read> = access %u, 0i
-    %10:vec4<f16> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec4<f16>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec4<f16>, read> = access %u, 0i
-    %13:vec4<f16> = load %12
-    %14:vec4<f16> = swizzle %13, ywxz
-    store %11, %14
-    %15:ptr<workgroup, vec4<f16>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec4<f16>, read> = access %u, 1i
-    %17:f16 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half2x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = half2x4(half4(0.0h), half4(0.0h));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].ywxz;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
index 0a303a5..52da3cb 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  float2x4 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant float2x4* u;
+  threadgroup float2x4* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat2x4<f32>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat2x4<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat2x4<f32>(vec4<f32>(0.0f))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat2x4<f32> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec4<f32>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec4<f32>, read> = access %u, 0i
-    %10:vec4<f32> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec4<f32>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec4<f32>, read> = access %u, 0i
-    %13:vec4<f32> = load %12
-    %14:vec4<f32> = swizzle %13, ywxz
-    store %11, %14
-    %15:ptr<workgroup, vec4<f32>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec4<f32>, read> = access %u, 1i
-    %17:f32 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float2x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = float2x4(float4(0.0f), float4(0.0f));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].ywxz;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
index d5fae9d..5893867 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  half3x2 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant half3x2* u;
+  threadgroup half3x2* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat3x2<f16>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat3x2<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat3x2<f16>(vec2<f16>(0.0h))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat3x2<f16> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec2<f16>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec2<f16>, read> = access %u, 0i
-    %10:vec2<f16> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec2<f16>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec2<f16>, read> = access %u, 0i
-    %13:vec2<f16> = load %12
-    %14:vec2<f16> = swizzle %13, yx
-    store %11, %14
-    %15:ptr<workgroup, vec2<f16>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec2<f16>, read> = access %u, 1i
-    %17:f16 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half3x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = half3x2(half2(0.0h), half2(0.0h), half2(0.0h));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].yx;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
index 45c60a8..da3e2fd 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  float3x2 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant float3x2* u;
+  threadgroup float3x2* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat3x2<f32>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat3x2<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat3x2<f32>(vec2<f32>(0.0f))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat3x2<f32> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec2<f32>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec2<f32>, read> = access %u, 0i
-    %10:vec2<f32> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec2<f32>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec2<f32>, read> = access %u, 0i
-    %13:vec2<f32> = load %12
-    %14:vec2<f32> = swizzle %13, yx
-    store %11, %14
-    %15:ptr<workgroup, vec2<f32>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec2<f32>, read> = access %u, 1i
-    %17:f32 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float3x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = float3x2(float2(0.0f), float2(0.0f), float2(0.0f));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].yx;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
index aef9a32..3276f46 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  half3x3 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant half3x3* u;
+  threadgroup half3x3* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat3x3<f16>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat3x3<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat3x3<f16>(vec3<f16>(0.0h))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat3x3<f16> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec3<f16>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec3<f16>, read> = access %u, 0i
-    %10:vec3<f16> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec3<f16>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec3<f16>, read> = access %u, 0i
-    %13:vec3<f16> = load %12
-    %14:vec3<f16> = swizzle %13, zxy
-    store %11, %14
-    %15:ptr<workgroup, vec3<f16>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec3<f16>, read> = access %u, 1i
-    %17:f16 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half3x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = half3x3(half3(0.0h), half3(0.0h), half3(0.0h));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].zxy;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
index db86b8c..105e86d 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  float3x3 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant float3x3* u;
+  threadgroup float3x3* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat3x3<f32>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat3x3<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat3x3<f32>(vec3<f32>(0.0f))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat3x3<f32> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec3<f32>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec3<f32>, read> = access %u, 0i
-    %10:vec3<f32> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec3<f32>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec3<f32>, read> = access %u, 0i
-    %13:vec3<f32> = load %12
-    %14:vec3<f32> = swizzle %13, zxy
-    store %11, %14
-    %15:ptr<workgroup, vec3<f32>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec3<f32>, read> = access %u, 1i
-    %17:f32 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float3x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = float3x3(float3(0.0f), float3(0.0f), float3(0.0f));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].zxy;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
index 6f09501..fce2da8 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  half3x4 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant half3x4* u;
+  threadgroup half3x4* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat3x4<f16>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat3x4<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat3x4<f16>(vec4<f16>(0.0h))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat3x4<f16> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec4<f16>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec4<f16>, read> = access %u, 0i
-    %10:vec4<f16> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec4<f16>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec4<f16>, read> = access %u, 0i
-    %13:vec4<f16> = load %12
-    %14:vec4<f16> = swizzle %13, ywxz
-    store %11, %14
-    %15:ptr<workgroup, vec4<f16>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec4<f16>, read> = access %u, 1i
-    %17:f16 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half3x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = half3x4(half4(0.0h), half4(0.0h), half4(0.0h));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].ywxz;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
index 675d4d3..bc3018c 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  float3x4 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant float3x4* u;
+  threadgroup float3x4* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat3x4<f32>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat3x4<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat3x4<f32>(vec4<f32>(0.0f))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat3x4<f32> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec4<f32>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec4<f32>, read> = access %u, 0i
-    %10:vec4<f32> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec4<f32>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec4<f32>, read> = access %u, 0i
-    %13:vec4<f32> = load %12
-    %14:vec4<f32> = swizzle %13, ywxz
-    store %11, %14
-    %15:ptr<workgroup, vec4<f32>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec4<f32>, read> = access %u, 1i
-    %17:f32 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float3x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = float3x4(float4(0.0f), float4(0.0f), float4(0.0f));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].ywxz;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
index ced9ab1..2a21f89 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  half4x2 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant half4x2* u;
+  threadgroup half4x2* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat4x2<f16>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat4x2<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat4x2<f16>(vec2<f16>(0.0h))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat4x2<f16> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec2<f16>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec2<f16>, read> = access %u, 0i
-    %10:vec2<f16> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec2<f16>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec2<f16>, read> = access %u, 0i
-    %13:vec2<f16> = load %12
-    %14:vec2<f16> = swizzle %13, yx
-    store %11, %14
-    %15:ptr<workgroup, vec2<f16>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec2<f16>, read> = access %u, 1i
-    %17:f16 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half4x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = half4x2(half2(0.0h), half2(0.0h), half2(0.0h), half2(0.0h));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].yx;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
index ff6fa47..5ffe4bd 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  float4x2 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant float4x2* u;
+  threadgroup float4x2* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat4x2<f32>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat4x2<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat4x2<f32>(vec2<f32>(0.0f))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat4x2<f32> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec2<f32>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec2<f32>, read> = access %u, 0i
-    %10:vec2<f32> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec2<f32>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec2<f32>, read> = access %u, 0i
-    %13:vec2<f32> = load %12
-    %14:vec2<f32> = swizzle %13, yx
-    store %11, %14
-    %15:ptr<workgroup, vec2<f32>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec2<f32>, read> = access %u, 1i
-    %17:f32 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float4x2* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = float4x2(float2(0.0f), float2(0.0f), float2(0.0f), float2(0.0f));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].yx;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
index 4ec8615..2a836c1 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  half4x3 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant half4x3* u;
+  threadgroup half4x3* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat4x3<f16>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat4x3<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat4x3<f16>(vec3<f16>(0.0h))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat4x3<f16> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec3<f16>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec3<f16>, read> = access %u, 0i
-    %10:vec3<f16> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec3<f16>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec3<f16>, read> = access %u, 0i
-    %13:vec3<f16> = load %12
-    %14:vec3<f16> = swizzle %13, zxy
-    store %11, %14
-    %15:ptr<workgroup, vec3<f16>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec3<f16>, read> = access %u, 1i
-    %17:f16 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half4x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = half4x3(half3(0.0h), half3(0.0h), half3(0.0h), half3(0.0h));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].zxy;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
index ded9db0..1d930a4 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  float4x3 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant float4x3* u;
+  threadgroup float4x3* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat4x3<f32>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat4x3<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat4x3<f32>(vec3<f32>(0.0f))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat4x3<f32> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec3<f32>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec3<f32>, read> = access %u, 0i
-    %10:vec3<f32> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec3<f32>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec3<f32>, read> = access %u, 0i
-    %13:vec3<f32> = load %12
-    %14:vec3<f32> = swizzle %13, zxy
-    store %11, %14
-    %15:ptr<workgroup, vec3<f32>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec3<f32>, read> = access %u, 1i
-    %17:f32 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float4x3* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = float4x3(float3(0.0f), float3(0.0f), float3(0.0f), float3(0.0f));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].zxy;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
index 68cc6fd..67b0ff7 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  half4x4 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant half4x4* u;
+  threadgroup half4x4* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat4x4<f16>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat4x4<f16>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat4x4<f16>(vec4<f16>(0.0h))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat4x4<f16> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec4<f16>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec4<f16>, read> = access %u, 0i
-    %10:vec4<f16> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec4<f16>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec4<f16>, read> = access %u, 0i
-    %13:vec4<f16> = load %12
-    %14:vec4<f16> = swizzle %13, ywxz
-    store %11, %14
-    %15:ptr<workgroup, vec4<f16>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec4<f16>, read> = access %u, 1i
-    %17:f16 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant half4x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = half4x4(half4(0.0h), half4(0.0h), half4(0.0h), half4(0.0h));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].ywxz;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
index 9818ee4..1143198 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/to_workgroup.wgsl.expected.ir.msl
@@ -1,43 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_1 {
+  float4x4 tint_symbol;
+};
+struct tint_module_vars_struct {
+  const constant float4x4* u;
+  threadgroup float4x4* w;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %u:ptr<uniform, mat4x4<f32>, read> = var @binding_point(0, 0)
-  %w:ptr<workgroup, mat4x4<f32>, read_write> = var
-}
-
-%f = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %5:bool = eq %tint_local_index, 0u
-    if %5 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %w, mat4x4<f32>(vec4<f32>(0.0f))
-        exit_if  # if_1
-      }
-    }
-    %6:void = msl.threadgroup_barrier 4u
-    %7:mat4x4<f32> = load %u
-    store %w, %7
-    %8:ptr<workgroup, vec4<f32>, read_write> = access %w, 1i
-    %9:ptr<uniform, vec4<f32>, read> = access %u, 0i
-    %10:vec4<f32> = load %9
-    store %8, %10
-    %11:ptr<workgroup, vec4<f32>, read_write> = access %w, 1i
-    %12:ptr<uniform, vec4<f32>, read> = access %u, 0i
-    %13:vec4<f32> = load %12
-    %14:vec4<f32> = swizzle %13, ywxz
-    store %11, %14
-    %15:ptr<workgroup, vec4<f32>, read_write> = access %w, 0i
-    %16:ptr<uniform, vec4<f32>, read> = access %u, 1i
-    %17:f32 = load_vector_element %16, 0i
-    store_vector_element %15, 1i, %17
-    ret
+kernel void f(uint tint_local_index [[thread_index_in_threadgroup]], const constant float4x4* u [[buffer(0)]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.u=u, .w=(&(*v).tint_symbol)};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = float4x4(float4(0.0f), float4(0.0f), float4(0.0f), float4(0.0f));
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (*tint_module_vars.w) = (*tint_module_vars.u);
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0];
+  (*tint_module_vars.w)[1] = (*tint_module_vars.u)[0].ywxz;
+  (*tint_module_vars.w)[0][1] = (*tint_module_vars.u)[1][0];
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/bug/chromium/40943165.wgsl.expected.ir.msl b/test/tint/bug/chromium/40943165.wgsl.expected.ir.msl
index fbd61a0..a7b5929 100644
--- a/test/tint/bug/chromium/40943165.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/40943165.wgsl.expected.ir.msl
@@ -1,27 +1,17 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_1 {
+  float2x2 tint_symbol;
+};
+struct tint_module_vars_struct {
+  threadgroup float2x2* W;
+};
 
-threadgroup float2x2 W;
-kernel void F(uint mat2x2 [[thread_index_in_threadgroup]]) {
+kernel void F(uint mat2x2 [[thread_index_in_threadgroup]], threadgroup tint_symbol_1* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.W=(&(*v).tint_symbol)};
   if ((mat2x2 == 0u)) {
-    W = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.W) = float2x2(float2(0.0f), float2(0.0f));
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  W[0] = (W[0] + 0.0f);
+  (*tint_module_vars.W)[0] = ((*tint_module_vars.W)[0] + 0.0f);
 }
-program_source:4:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 W;
-                     ^
-program_source:6:15: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((mat2x2 == 0u)) {
-       ~~~~~~~^~~~~
-program_source:6:15: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((mat2x2 == 0u)) {
-      ~       ^    ~
-program_source:6:15: note: use '=' to turn this equality comparison into an assignment
-  if ((mat2x2 == 0u)) {
-              ^~
-              =
-
diff --git a/test/tint/bug/tint/1926.wgsl.expected.ir.msl b/test/tint/bug/tint/1926.wgsl.expected.ir.msl
index c450d39..812684c 100644
--- a/test/tint/bug/tint/1926.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1926.wgsl.expected.ir.msl
@@ -1,41 +1,24 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_symbol_2 {
+  uint tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup uint* sh_atomic_failed;
+  device uint* output;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %sh_atomic_failed:ptr<workgroup, u32, read_write> = var
-  %output:ptr<storage, u32, read_write> = var @binding_point(0, 4)
-}
-
-%tint_symbol = @compute @workgroup_size(256, 1, 1) func(%global_id:vec3<u32> [@global_invocation_id], %local_id:vec3<u32> [@local_invocation_id], %tint_local_index:u32 [@local_invocation_index]):void {
-  $B2: {
-    %7:bool = eq %tint_local_index, 0u
-    if %7 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %sh_atomic_failed, 0u
-        exit_if  # if_1
-      }
-    }
-    %8:void = msl.threadgroup_barrier 4u
-    %9:void = msl.threadgroup_barrier 4u
-    %10:u32 = load %sh_atomic_failed
-    %11:u32 = let %10
-    %12:void = msl.threadgroup_barrier 4u
-    %failed:u32 = let %11
-    %14:u32 = access %local_id, 0u
-    %15:bool = eq %14, 0u
-    if %15 [t: $B4] {  # if_2
-      $B4: {  # true
-        store %output, %failed
-        exit_if  # if_2
-      }
-    }
-    ret
+kernel void tint_symbol(uint3 global_id [[thread_position_in_grid]], uint3 local_id [[thread_position_in_threadgroup]], uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v, device uint* output [[buffer(4)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.sh_atomic_failed=(&(*v).tint_symbol_1), .output=output};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.sh_atomic_failed) = 0u;
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  uint const v_1 = (*tint_module_vars.sh_atomic_failed);
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  uint const failed = v_1;
+  if ((local_id[0u] == 0u)) {
+    (*tint_module_vars.output) = failed;
   }
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.ir.msl
index 5a6429d..5f96aa4 100644
--- a/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.ir.msl
@@ -1,32 +1,19 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_2 {
+  int tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup int* i;
+};
 
-threadgroup int i;
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.i=(&(*v).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    i = 0;
+    (*tint_module_vars.i) = 0;
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  i = 123;
-  threadgroup int* const p = i;
-  int const u = (p + 1);
+  (*tint_module_vars.i) = 123;
+  threadgroup int* const p = tint_module_vars.i;
+  int const u = ((*p) + 1);
 }
-program_source:4:17: error: program scope variable must reside in constant address space
-threadgroup int i;
-                ^
-program_source:6:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:6:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:6:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-program_source:12:13: error: cannot initialize a variable of type 'const int' with an rvalue of type 'threadgroup int *'
-  int const u = (p + 1);
-            ^   ~~~~~~~
-
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.ir.msl
index 25d79e7..fbd0e7d 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32.wgsl.expected.ir.msl
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_2 {
+  int tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup int* S;
+};
 
-threadgroup int S;
 int func(threadgroup int* const pointer) {
-  return pointer;
+  return (*pointer);
 }
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    S = 0;
+    (*tint_module_vars.S) = 0;
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  int const r = func(S);
+  int const r = func(tint_module_vars.S);
 }
-program_source:4:17: error: program scope variable must reside in constant address space
-threadgroup int S;
-                ^
-program_source:6:10: error: cannot initialize return object of type 'int' with an lvalue of type 'threadgroup int *const'
-  return pointer;
-         ^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-program_source:13:13: warning: unused variable 'r' [-Wunused-variable]
-  int const r = func(S);
-            ^
-
diff --git a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.ir.msl
index 4be387c..0e25c93 100644
--- a/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/i32_in_struct.wgsl.expected.ir.msl
@@ -1,39 +1,23 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 struct str {
   int i;
 };
+struct tint_symbol_2 {
+  str tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup str* S;
+};
 
-threadgroup str S;
 int func(threadgroup int* const pointer) {
-  return pointer;
+  return (*pointer);
 }
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    S = str{};
+    (*tint_module_vars.S) = str{};
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  int const r = func(S.i);
+  int const r = func((&(*tint_module_vars.S).i));
 }
-program_source:7:17: error: program scope variable must reside in constant address space
-threadgroup str S;
-                ^
-program_source:9:10: error: cannot initialize return object of type 'int' with an lvalue of type 'threadgroup int *const'
-  return pointer;
-         ^~~~~~~
-program_source:12:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:12:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:12:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-program_source:16:13: warning: unused variable 'r' [-Wunused-variable]
-  int const r = func(S.i);
-            ^
-
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl
index 6866579..9e76b8b 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_2 {
+  float2x2 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup float2x2* S;
+};
 
-threadgroup float2x2 S;
 float2 func(threadgroup float2* const pointer) {
-  return pointer;
+  return (*pointer);
 }
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    S = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.S) = float2x2(float2(0.0f), float2(0.0f));
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  float2 const r = func(S[1]);
+  float2 const r = func((&(*tint_module_vars.S)[1]));
 }
-program_source:4:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 S;
-                     ^
-program_source:6:10: error: cannot initialize return object of type 'float2' (vector of 2 'float' values) with an lvalue of type 'threadgroup float2 *const'
-  return pointer;
-         ^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-program_source:13:16: warning: unused variable 'r' [-Wunused-variable]
-  float2 const r = func(S[1]);
-               ^
-
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.ir.msl
index 1e11dc4..3764937 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32.wgsl.expected.ir.msl
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_2 {
+  float4 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup float4* S;
+};
 
-threadgroup float4 S;
 float4 func(threadgroup float4* const pointer) {
-  return pointer;
+  return (*pointer);
 }
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    S = float4(0.0f);
+    (*tint_module_vars.S) = float4(0.0f);
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  float4 const r = func(S);
+  float4 const r = func(tint_module_vars.S);
 }
-program_source:4:20: error: program scope variable must reside in constant address space
-threadgroup float4 S;
-                   ^
-program_source:6:10: error: cannot initialize return object of type 'float4' (vector of 4 'float' values) with an lvalue of type 'threadgroup float4 *const'
-  return pointer;
-         ^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-program_source:13:16: warning: unused variable 'r' [-Wunused-variable]
-  float4 const r = func(S);
-               ^
-
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl
index a6f5b85..fb1ab2d 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_2 {
+  float2x4 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup float2x4* S;
+};
 
-threadgroup float2x4 S;
 float4 func(threadgroup float4* const pointer) {
-  return pointer;
+  return (*pointer);
 }
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    S = float2x4(float4(0.0f), float4(0.0f));
+    (*tint_module_vars.S) = float2x4(float4(0.0f), float4(0.0f));
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  float4 const r = func(S[1]);
+  float4 const r = func((&(*tint_module_vars.S)[1]));
 }
-program_source:4:22: error: program scope variable must reside in constant address space
-threadgroup float2x4 S;
-                     ^
-program_source:6:10: error: cannot initialize return object of type 'float4' (vector of 4 'float' values) with an lvalue of type 'threadgroup float4 *const'
-  return pointer;
-         ^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-program_source:13:16: warning: unused variable 'r' [-Wunused-variable]
-  float4 const r = func(S[1]);
-               ^
-
diff --git a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl
index 62142f9..96d683f 100644
--- a/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl
@@ -1,39 +1,23 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 struct str {
   float4 i;
 };
+struct tint_symbol_2 {
+  str tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup str* S;
+};
 
-threadgroup str S;
 float4 func(threadgroup float4* const pointer) {
-  return pointer;
+  return (*pointer);
 }
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    S = str{};
+    (*tint_module_vars.S) = str{};
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  float4 const r = func(S.i);
+  float4 const r = func((&(*tint_module_vars.S).i));
 }
-program_source:7:17: error: program scope variable must reside in constant address space
-threadgroup str S;
-                ^
-program_source:9:10: error: cannot initialize return object of type 'float4' (vector of 4 'float' values) with an lvalue of type 'threadgroup float4 *const'
-  return pointer;
-         ^~~~~~~
-program_source:12:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:12:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:12:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-program_source:16:16: warning: unused variable 'r' [-Wunused-variable]
-  float4 const r = func(S.i);
-               ^
-
diff --git a/test/tint/ptr_ref/store/param/workgroup/i32.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/i32.wgsl.expected.ir.msl
index b5cca68..57dbb78 100644
--- a/test/tint/ptr_ref/store/param/workgroup/i32.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/i32.wgsl.expected.ir.msl
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_2 {
+  int tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup int* S;
+};
 
-threadgroup int S;
 void func(threadgroup int* const pointer) {
-  pointer = 42;
+  (*pointer) = 42;
 }
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    S = 0;
+    (*tint_module_vars.S) = 0;
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  func(S);
+  func(tint_module_vars.S);
 }
-program_source:4:17: error: program scope variable must reside in constant address space
-threadgroup int S;
-                ^
-program_source:6:11: error: cannot assign to variable 'pointer' with const-qualified type 'threadgroup int *const'
-  pointer = 42;
-  ~~~~~~~ ^
-program_source:5:34: note: variable 'pointer' declared const here
-void func(threadgroup int* const pointer) {
-          ~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-
diff --git a/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl.expected.ir.msl
index cc3cff5..1defcd2 100644
--- a/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/i32_in_struct.wgsl.expected.ir.msl
@@ -1,39 +1,23 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 struct str {
   int i;
 };
+struct tint_symbol_2 {
+  str tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup str* S;
+};
 
-threadgroup str S;
 void func(threadgroup int* const pointer) {
-  pointer = 42;
+  (*pointer) = 42;
 }
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    S = str{};
+    (*tint_module_vars.S) = str{};
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  func(S.i);
+  func((&(*tint_module_vars.S).i));
 }
-program_source:7:17: error: program scope variable must reside in constant address space
-threadgroup str S;
-                ^
-program_source:9:11: error: cannot assign to variable 'pointer' with const-qualified type 'threadgroup int *const'
-  pointer = 42;
-  ~~~~~~~ ^
-program_source:8:34: note: variable 'pointer' declared const here
-void func(threadgroup int* const pointer) {
-          ~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~
-program_source:12:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:12:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:12:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl
index ef57e5a..f089435 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec2_f32_in_mat2x2.wgsl.expected.ir.msl
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_2 {
+  float2x2 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup float2x2* S;
+};
 
-threadgroup float2x2 S;
 void func(threadgroup float2* const pointer) {
-  pointer = float2(0.0f);
+  (*pointer) = float2(0.0f);
 }
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    S = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.S) = float2x2(float2(0.0f), float2(0.0f));
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  func(S[1]);
+  func((&(*tint_module_vars.S)[1]));
 }
-program_source:4:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 S;
-                     ^
-program_source:6:11: error: cannot assign to variable 'pointer' with const-qualified type 'threadgroup float2 *const'
-  pointer = float2(0.0f);
-  ~~~~~~~ ^
-program_source:5:37: note: variable 'pointer' declared const here
-void func(threadgroup float2* const pointer) {
-          ~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl.expected.ir.msl
index 2de34c8..416f51b 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec4_f32.wgsl.expected.ir.msl
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_2 {
+  float4 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup float4* S;
+};
 
-threadgroup float4 S;
 void func(threadgroup float4* const pointer) {
-  pointer = float4(0.0f);
+  (*pointer) = float4(0.0f);
 }
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    S = float4(0.0f);
+    (*tint_module_vars.S) = float4(0.0f);
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  func(S);
+  func(tint_module_vars.S);
 }
-program_source:4:20: error: program scope variable must reside in constant address space
-threadgroup float4 S;
-                   ^
-program_source:6:11: error: cannot assign to variable 'pointer' with const-qualified type 'threadgroup float4 *const'
-  pointer = float4(0.0f);
-  ~~~~~~~ ^
-program_source:5:37: note: variable 'pointer' declared const here
-void func(threadgroup float4* const pointer) {
-          ~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl
index 0f5c1cc..c4d6342 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_mat2x4.wgsl.expected.ir.msl
@@ -1,36 +1,20 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_2 {
+  float2x4 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup float2x4* S;
+};
 
-threadgroup float2x4 S;
 void func(threadgroup float4* const pointer) {
-  pointer = float4(0.0f);
+  (*pointer) = float4(0.0f);
 }
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    S = float2x4(float4(0.0f), float4(0.0f));
+    (*tint_module_vars.S) = float2x4(float4(0.0f), float4(0.0f));
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  func(S[1]);
+  func((&(*tint_module_vars.S)[1]));
 }
-program_source:4:22: error: program scope variable must reside in constant address space
-threadgroup float2x4 S;
-                     ^
-program_source:6:11: error: cannot assign to variable 'pointer' with const-qualified type 'threadgroup float4 *const'
-  pointer = float4(0.0f);
-  ~~~~~~~ ^
-program_source:5:37: note: variable 'pointer' declared const here
-void func(threadgroup float4* const pointer) {
-          ~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~
-program_source:9:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:9:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:9:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-
diff --git a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl
index 9017799..e943602 100644
--- a/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/store/param/workgroup/vec4_f32_in_struct.wgsl.expected.ir.msl
@@ -1,39 +1,23 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 struct str {
   float4 i;
 };
+struct tint_symbol_2 {
+  str tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup str* S;
+};
 
-threadgroup str S;
 void func(threadgroup float4* const pointer) {
-  pointer = float4(0.0f);
+  (*pointer) = float4(0.0f);
 }
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.S=(&(*v).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    S = str{};
+    (*tint_module_vars.S) = str{};
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  func(S.i);
+  func((&(*tint_module_vars.S).i));
 }
-program_source:7:17: error: program scope variable must reside in constant address space
-threadgroup str S;
-                ^
-program_source:9:11: error: cannot assign to variable 'pointer' with const-qualified type 'threadgroup float4 *const'
-  pointer = float4(0.0f);
-  ~~~~~~~ ^
-program_source:8:37: note: variable 'pointer' declared const here
-void func(threadgroup float4* const pointer) {
-          ~~~~~~~~~~~~~~~~~~~~~~~~~~^~~~~~~
-program_source:12:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:12:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:12:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-
diff --git a/test/tint/types/module_scope_used_in_functions.wgsl.expected.ir.msl b/test/tint/types/module_scope_used_in_functions.wgsl.expected.ir.msl
index 85e0261..2b6c649 100644
--- a/test/tint/types/module_scope_used_in_functions.wgsl.expected.ir.msl
+++ b/test/tint/types/module_scope_used_in_functions.wgsl.expected.ir.msl
@@ -1,63 +1,49 @@
-SKIP: FAILED
+#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];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %p:ptr<private, f32, read_write> = var
-  %w:ptr<workgroup, f32, read_write> = var
-  %uniforms:ptr<storage, vec2<f32>, read> = var @binding_point(0, 1)
-  %storages:ptr<storage, array<f32>, read_write> = var @binding_point(0, 0)
-}
+struct tint_module_vars_struct {
+  thread float* p;
+  threadgroup float* w;
+  const device float2* uniforms;
+  device tint_array<float, 1>* storages;
+};
+struct tint_symbol_2 {
+  float tint_symbol_1;
+};
 
-%no_uses = func():void {
-  $B2: {
-    ret
-  }
+void no_uses() {
 }
-%zoo = func():void {
-  $B3: {
-    %7:f32 = load %p
-    %8:f32 = mul %7, 2.0f
-    store %p, %8
-    ret
-  }
+void zoo(tint_module_vars_struct tint_module_vars) {
+  (*tint_module_vars.p) = ((*tint_module_vars.p) * 2.0f);
 }
-%bar = func(%a:f32, %b:f32):void {
-  $B4: {
-    store %p, %a
-    store %w, %b
-    %12:ptr<storage, f32, read_write> = access %storages, 0i
-    %13:f32 = load_vector_element %uniforms, 0u
-    store %12, %13
-    %14:void = call %zoo
-    ret
-  }
+void bar(float a, float b, tint_module_vars_struct tint_module_vars) {
+  (*tint_module_vars.p) = a;
+  (*tint_module_vars.w) = b;
+  (*tint_module_vars.storages)[0] = (*tint_module_vars.uniforms)[0u];
+  zoo(tint_module_vars);
 }
-%foo = func(%a_1:f32):void {  # %a_1: 'a'
-  $B5: {
-    %b_1:f32 = let 2.0f  # %b_1: 'b'
-    %18:void = call %bar, %a_1, %b_1
-    %19:void = call %no_uses
-    ret
-  }
+void foo(float a, tint_module_vars_struct tint_module_vars) {
+  float const b = 2.0f;
+  bar(a, b, tint_module_vars);
+  no_uses();
 }
-%tint_symbol = @compute @workgroup_size(1, 1, 1) func(%tint_local_index:u32 [@local_invocation_index]):void {
-  $B6: {
-    %22:bool = eq %tint_local_index, 0u
-    if %22 [t: $B7] {  # if_1
-      $B7: {  # true
-        store %w, 0.0f
-        exit_if  # if_1
-      }
-    }
-    %23:void = msl.threadgroup_barrier 4u
-    %24:void = call %foo, 1.0f
-    ret
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v, const device float2* uniforms [[buffer(1)]], device tint_array<float, 1>* storages [[buffer(0)]]) {
+  thread float p = 0.0f;
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.p=(&p), .w=(&(*v).tint_symbol_1), .uniforms=uniforms, .storages=storages};
+  if ((tint_local_index == 0u)) {
+    (*tint_module_vars.w) = 0.0f;
   }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  foo(1.0f, tint_module_vars);
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/types/module_scope_var.wgsl.expected.ir.msl b/test/tint/types/module_scope_var.wgsl.expected.ir.msl
index e3cf1cd..b44443a 100644
--- a/test/tint/types/module_scope_var.wgsl.expected.ir.msl
+++ b/test/tint/types/module_scope_var.wgsl.expected.ir.msl
@@ -1,7 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_2 {
+  float tint_symbol_1;
+};
 template<typename T, size_t N>
 struct tint_array {
   const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -17,76 +18,45 @@
 struct S {
   float a;
 };
+struct tint_module_vars_struct {
+  thread bool* bool_var;
+  thread int* i32_var;
+  thread uint* u32_var;
+  thread float* f32_var;
+  thread int2* v2i32_var;
+  thread uint3* v3u32_var;
+  thread float4* v4f32_var;
+  thread float2x3* m2x3_var;
+  thread tint_array<float, 4>* arr_var;
+  thread S* struct_var;
+  threadgroup float* wg_var;
+};
 
-thread bool bool_var = false;
-thread int i32_var = 0;
-thread uint u32_var = 0u;
-thread float f32_var = 0.0f;
-thread int2 v2i32_var = 0;
-thread uint3 v3u32_var = 0u;
-thread float4 v4f32_var = 0.0f;
-thread float2x3 m2x3_var = float2x3(0.0f);
-thread tint_array<float, 4> arr_var = {};
-thread S struct_var = {};
-threadgroup float wg_var;
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+  thread bool bool_var = false;
+  thread int i32_var = 0;
+  thread uint u32_var = 0u;
+  thread float f32_var = 0.0f;
+  thread int2 v2i32_var = 0;
+  thread uint3 v3u32_var = 0u;
+  thread float4 v4f32_var = 0.0f;
+  thread float2x3 m2x3_var = float2x3(0.0f);
+  thread tint_array<float, 4> arr_var = {};
+  thread S struct_var = {};
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.bool_var=(&bool_var), .i32_var=(&i32_var), .u32_var=(&u32_var), .f32_var=(&f32_var), .v2i32_var=(&v2i32_var), .v3u32_var=(&v3u32_var), .v4f32_var=(&v4f32_var), .m2x3_var=(&m2x3_var), .arr_var=(&arr_var), .struct_var=(&struct_var), .wg_var=(&(*v).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    wg_var = 0.0f;
+    (*tint_module_vars.wg_var) = 0.0f;
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  bool_var = false;
-  i32_var = 0;
-  u32_var = 0u;
-  f32_var = 0.0f;
-  v2i32_var = int2(0);
-  v3u32_var = uint3(0u);
-  v4f32_var = float4(0.0f);
-  m2x3_var = float2x3(float3(0.0f), float3(0.0f));
-  arr_var = tint_array<float, 4>{};
-  struct_var = S{};
-  wg_var = 42.0f;
+  (*tint_module_vars.bool_var) = false;
+  (*tint_module_vars.i32_var) = 0;
+  (*tint_module_vars.u32_var) = 0u;
+  (*tint_module_vars.f32_var) = 0.0f;
+  (*tint_module_vars.v2i32_var) = int2(0);
+  (*tint_module_vars.v3u32_var) = uint3(0u);
+  (*tint_module_vars.v4f32_var) = float4(0.0f);
+  (*tint_module_vars.m2x3_var) = float2x3(float3(0.0f), float3(0.0f));
+  (*tint_module_vars.arr_var) = tint_array<float, 4>{};
+  (*tint_module_vars.struct_var) = S{};
+  (*tint_module_vars.wg_var) = 42.0f;
 }
-program_source:19:13: error: program scope variable must reside in constant address space
-thread bool bool_var = false;
-            ^
-program_source:20:12: error: program scope variable must reside in constant address space
-thread int i32_var = 0;
-           ^
-program_source:21:13: error: program scope variable must reside in constant address space
-thread uint u32_var = 0u;
-            ^
-program_source:22:14: error: program scope variable must reside in constant address space
-thread float f32_var = 0.0f;
-             ^
-program_source:23:13: error: program scope variable must reside in constant address space
-thread int2 v2i32_var = 0;
-            ^
-program_source:24:14: error: program scope variable must reside in constant address space
-thread uint3 v3u32_var = 0u;
-             ^
-program_source:25:15: error: program scope variable must reside in constant address space
-thread float4 v4f32_var = 0.0f;
-              ^
-program_source:26:17: error: program scope variable must reside in constant address space
-thread float2x3 m2x3_var = float2x3(0.0f);
-                ^
-program_source:27:29: error: program scope variable must reside in constant address space
-thread tint_array<float, 4> arr_var = {};
-                            ^
-program_source:28:10: error: program scope variable must reside in constant address space
-thread S struct_var = {};
-         ^
-program_source:29:19: error: program scope variable must reside in constant address space
-threadgroup float wg_var;
-                  ^
-program_source:31:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:31:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:31:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-
diff --git a/test/tint/types/module_scope_vars_pointers.wgsl.expected.ir.msl b/test/tint/types/module_scope_vars_pointers.wgsl.expected.ir.msl
index 68d0820..bc857e6 100644
--- a/test/tint/types/module_scope_vars_pointers.wgsl.expected.ir.msl
+++ b/test/tint/types/module_scope_vars_pointers.wgsl.expected.ir.msl
@@ -1,43 +1,22 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_2 {
+  float tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  thread float* p;
+  threadgroup float* w;
+};
 
-thread float p = 0.0f;
-threadgroup float w;
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v) {
+  thread float p = 0.0f;
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.p=(&p), .w=(&(*v).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    w = 0.0f;
+    (*tint_module_vars.w) = 0.0f;
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  thread float* const p_ptr = p;
-  threadgroup float* const w_ptr = w;
-  float const x = (p_ptr + w_ptr);
-  p_ptr = x;
+  thread float* const p_ptr = tint_module_vars.p;
+  threadgroup float* const w_ptr = tint_module_vars.w;
+  float const x = ((*p_ptr) + (*w_ptr));
+  (*p_ptr) = x;
 }
-program_source:4:14: error: program scope variable must reside in constant address space
-thread float p = 0.0f;
-             ^
-program_source:5:19: error: program scope variable must reside in constant address space
-threadgroup float w;
-                  ^
-program_source:7:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:7:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:7:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-program_source:13:26: error: invalid operands to binary expression ('float *const' and 'threadgroup float *const')
-  float const x = (p_ptr + w_ptr);
-                   ~~~~~ ^ ~~~~~
-program_source:14:9: error: cannot assign to variable 'p_ptr' with const-qualified type 'float *const'
-  p_ptr = x;
-  ~~~~~ ^
-program_source:11:23: note: variable 'p_ptr' declared const here
-  thread float* const p_ptr = p;
-  ~~~~~~~~~~~~~~~~~~~~^~~~~~~~~
-
diff --git a/test/tint/var/initialization/workgroup/matrix.wgsl.expected.ir.msl b/test/tint/var/initialization/workgroup/matrix.wgsl.expected.ir.msl
index a1c80e0..44e6f0a 100644
--- a/test/tint/var/initialization/workgroup/matrix.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/matrix.wgsl.expected.ir.msl
@@ -1,26 +1,16 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_2 {
+  float2x3 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup float2x3* v;
+};
 
-threadgroup float2x3 v;
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v=(&(*v_1).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    v = float2x3(float3(0.0f), float3(0.0f));
+    (*tint_module_vars.v) = float2x3(float3(0.0f), float3(0.0f));
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
 }
-program_source:4:22: error: program scope variable must reside in constant address space
-threadgroup float2x3 v;
-                     ^
-program_source:6:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:6:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:6:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-
diff --git a/test/tint/var/initialization/workgroup/scalar.wgsl.expected.ir.msl b/test/tint/var/initialization/workgroup/scalar.wgsl.expected.ir.msl
index e1f46e9..97b64c8 100644
--- a/test/tint/var/initialization/workgroup/scalar.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/scalar.wgsl.expected.ir.msl
@@ -1,30 +1,17 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_2 {
+  int tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup int* v;
+};
 
-threadgroup int v;
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v=(&(*v_1).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    v = 0;
+    (*tint_module_vars.v) = 0;
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  int const i = v;
+  int const i = (*tint_module_vars.v);
 }
-program_source:4:17: error: program scope variable must reside in constant address space
-threadgroup int v;
-                ^
-program_source:6:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:6:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:6:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-program_source:10:13: warning: unused variable 'i' [-Wunused-variable]
-  int const i = v;
-            ^
-
diff --git a/test/tint/var/initialization/workgroup/struct.wgsl.expected.ir.msl b/test/tint/var/initialization/workgroup/struct.wgsl.expected.ir.msl
index 6655af9..e974d43 100644
--- a/test/tint/var/initialization/workgroup/struct.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/struct.wgsl.expected.ir.msl
@@ -1,30 +1,20 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 struct S {
   int a;
   float b;
 };
+struct tint_symbol_2 {
+  S tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup S* v;
+};
 
-threadgroup S v;
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v=(&(*v_1).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    v = S{};
+    (*tint_module_vars.v) = S{};
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
 }
-program_source:8:15: error: program scope variable must reside in constant address space
-threadgroup S v;
-              ^
-program_source:10:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:10:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:10:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-
diff --git a/test/tint/var/initialization/workgroup/vector.wgsl.expected.ir.msl b/test/tint/var/initialization/workgroup/vector.wgsl.expected.ir.msl
index a8308b3..d24feeb 100644
--- a/test/tint/var/initialization/workgroup/vector.wgsl.expected.ir.msl
+++ b/test/tint/var/initialization/workgroup/vector.wgsl.expected.ir.msl
@@ -1,26 +1,16 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_2 {
+  int3 tint_symbol_1;
+};
+struct tint_module_vars_struct {
+  threadgroup int3* v;
+};
 
-threadgroup int3 v;
-kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v_1) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v=(&(*v_1).tint_symbol_1)};
   if ((tint_local_index == 0u)) {
-    v = int3(0);
+    (*tint_module_vars.v) = int3(0);
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
 }
-program_source:4:18: error: program scope variable must reside in constant address space
-threadgroup int3 v;
-                 ^
-program_source:6:25: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((tint_local_index == 0u)) {
-       ~~~~~~~~~~~~~~~~~^~~~~
-program_source:6:25: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((tint_local_index == 0u)) {
-      ~                 ^    ~
-program_source:6:25: note: use '=' to turn this equality comparison into an assignment
-  if ((tint_local_index == 0u)) {
-                        ^~
-                        =
-
diff --git a/test/tint/var/uses/many_workgroup_vars.wgsl.expected.ir.msl b/test/tint/var/uses/many_workgroup_vars.wgsl.expected.ir.msl
index 81f00cc..0a83ab7 100644
--- a/test/tint/var/uses/many_workgroup_vars.wgsl.expected.ir.msl
+++ b/test/tint/var/uses/many_workgroup_vars.wgsl.expected.ir.msl
@@ -1,621 +1,413 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
+struct tint_symbol_101 {
+  float2x2 tint_symbol_1;
+  float2x2 tint_symbol_2;
+  float2x2 tint_symbol_3;
+  float2x2 tint_symbol_4;
+  float2x2 tint_symbol_5;
+  float2x2 tint_symbol_6;
+  float2x2 tint_symbol_7;
+  float2x2 tint_symbol_8;
+  float2x2 tint_symbol_9;
+  float2x2 tint_symbol_10;
+  float2x2 tint_symbol_11;
+  float2x2 tint_symbol_12;
+  float2x2 tint_symbol_13;
+  float2x2 tint_symbol_14;
+  float2x2 tint_symbol_15;
+  float2x2 tint_symbol_16;
+  float2x2 tint_symbol_17;
+  float2x2 tint_symbol_18;
+  float2x2 tint_symbol_19;
+  float2x2 tint_symbol_20;
+  float2x2 tint_symbol_21;
+  float2x2 tint_symbol_22;
+  float2x2 tint_symbol_23;
+  float2x2 tint_symbol_24;
+  float2x2 tint_symbol_25;
+  float2x2 tint_symbol_26;
+  float2x2 tint_symbol_27;
+  float2x2 tint_symbol_28;
+  float2x2 tint_symbol_29;
+  float2x2 tint_symbol_30;
+  float2x2 tint_symbol_31;
+  float2x2 tint_symbol_32;
+  float2x2 tint_symbol_33;
+  float2x2 tint_symbol_34;
+  float2x2 tint_symbol_35;
+  float2x2 tint_symbol_36;
+  float2x2 tint_symbol_37;
+  float2x2 tint_symbol_38;
+  float2x2 tint_symbol_39;
+  float2x2 tint_symbol_40;
+  float2x2 tint_symbol_41;
+  float2x2 tint_symbol_42;
+  float2x2 tint_symbol_43;
+  float2x2 tint_symbol_44;
+  float2x2 tint_symbol_45;
+  float2x2 tint_symbol_46;
+  float2x2 tint_symbol_47;
+  float2x2 tint_symbol_48;
+  float2x2 tint_symbol_49;
+  float2x2 tint_symbol_50;
+  float2x2 tint_symbol_51;
+  float2x2 tint_symbol_52;
+  float2x2 tint_symbol_53;
+  float2x2 tint_symbol_54;
+  float2x2 tint_symbol_55;
+  float2x2 tint_symbol_56;
+  float2x2 tint_symbol_57;
+  float2x2 tint_symbol_58;
+  float2x2 tint_symbol_59;
+  float2x2 tint_symbol_60;
+  float2x2 tint_symbol_61;
+  float2x2 tint_symbol_62;
+  float2x2 tint_symbol_63;
+  float2x2 tint_symbol_64;
+  float2x2 tint_symbol_65;
+  float2x2 tint_symbol_66;
+  float2x2 tint_symbol_67;
+  float2x2 tint_symbol_68;
+  float2x2 tint_symbol_69;
+  float2x2 tint_symbol_70;
+  float2x2 tint_symbol_71;
+  float2x2 tint_symbol_72;
+  float2x2 tint_symbol_73;
+  float2x2 tint_symbol_74;
+  float2x2 tint_symbol_75;
+  float2x2 tint_symbol_76;
+  float2x2 tint_symbol_77;
+  float2x2 tint_symbol_78;
+  float2x2 tint_symbol_79;
+  float2x2 tint_symbol_80;
+  float2x2 tint_symbol_81;
+  float2x2 tint_symbol_82;
+  float2x2 tint_symbol_83;
+  float2x2 tint_symbol_84;
+  float2x2 tint_symbol_85;
+  float2x2 tint_symbol_86;
+  float2x2 tint_symbol_87;
+  float2x2 tint_symbol_88;
+  float2x2 tint_symbol_89;
+  float2x2 tint_symbol_90;
+  float2x2 tint_symbol_91;
+  float2x2 tint_symbol_92;
+  float2x2 tint_symbol_93;
+  float2x2 tint_symbol_94;
+  float2x2 tint_symbol_95;
+  float2x2 tint_symbol_96;
+  float2x2 tint_symbol_97;
+  float2x2 tint_symbol_98;
+  float2x2 tint_symbol_99;
+  float2x2 tint_symbol_100;
+};
+struct tint_module_vars_struct {
+  threadgroup float2x2* m00;
+  threadgroup float2x2* m01;
+  threadgroup float2x2* m02;
+  threadgroup float2x2* m03;
+  threadgroup float2x2* m04;
+  threadgroup float2x2* m05;
+  threadgroup float2x2* m06;
+  threadgroup float2x2* m07;
+  threadgroup float2x2* m08;
+  threadgroup float2x2* m09;
+  threadgroup float2x2* m10;
+  threadgroup float2x2* m11;
+  threadgroup float2x2* m12;
+  threadgroup float2x2* m13;
+  threadgroup float2x2* m14;
+  threadgroup float2x2* m15;
+  threadgroup float2x2* m16;
+  threadgroup float2x2* m17;
+  threadgroup float2x2* m18;
+  threadgroup float2x2* m19;
+  threadgroup float2x2* m20;
+  threadgroup float2x2* m21;
+  threadgroup float2x2* m22;
+  threadgroup float2x2* m23;
+  threadgroup float2x2* m24;
+  threadgroup float2x2* m25;
+  threadgroup float2x2* m26;
+  threadgroup float2x2* m27;
+  threadgroup float2x2* m28;
+  threadgroup float2x2* m29;
+  threadgroup float2x2* m30;
+  threadgroup float2x2* m31;
+  threadgroup float2x2* m32;
+  threadgroup float2x2* m33;
+  threadgroup float2x2* m34;
+  threadgroup float2x2* m35;
+  threadgroup float2x2* m36;
+  threadgroup float2x2* m37;
+  threadgroup float2x2* m38;
+  threadgroup float2x2* m39;
+  threadgroup float2x2* m40;
+  threadgroup float2x2* m41;
+  threadgroup float2x2* m42;
+  threadgroup float2x2* m43;
+  threadgroup float2x2* m44;
+  threadgroup float2x2* m45;
+  threadgroup float2x2* m46;
+  threadgroup float2x2* m47;
+  threadgroup float2x2* m48;
+  threadgroup float2x2* m49;
+  threadgroup float2x2* m50;
+  threadgroup float2x2* m51;
+  threadgroup float2x2* m52;
+  threadgroup float2x2* m53;
+  threadgroup float2x2* m54;
+  threadgroup float2x2* m55;
+  threadgroup float2x2* m56;
+  threadgroup float2x2* m57;
+  threadgroup float2x2* m58;
+  threadgroup float2x2* m59;
+  threadgroup float2x2* m60;
+  threadgroup float2x2* m61;
+  threadgroup float2x2* m62;
+  threadgroup float2x2* m63;
+  threadgroup float2x2* m64;
+  threadgroup float2x2* m65;
+  threadgroup float2x2* m66;
+  threadgroup float2x2* m67;
+  threadgroup float2x2* m68;
+  threadgroup float2x2* m69;
+  threadgroup float2x2* m70;
+  threadgroup float2x2* m71;
+  threadgroup float2x2* m72;
+  threadgroup float2x2* m73;
+  threadgroup float2x2* m74;
+  threadgroup float2x2* m75;
+  threadgroup float2x2* m76;
+  threadgroup float2x2* m77;
+  threadgroup float2x2* m78;
+  threadgroup float2x2* m79;
+  threadgroup float2x2* m80;
+  threadgroup float2x2* m81;
+  threadgroup float2x2* m82;
+  threadgroup float2x2* m83;
+  threadgroup float2x2* m84;
+  threadgroup float2x2* m85;
+  threadgroup float2x2* m86;
+  threadgroup float2x2* m87;
+  threadgroup float2x2* m88;
+  threadgroup float2x2* m89;
+  threadgroup float2x2* m90;
+  threadgroup float2x2* m91;
+  threadgroup float2x2* m92;
+  threadgroup float2x2* m93;
+  threadgroup float2x2* m94;
+  threadgroup float2x2* m95;
+  threadgroup float2x2* m96;
+  threadgroup float2x2* m97;
+  threadgroup float2x2* m98;
+  threadgroup float2x2* m99;
+};
 
-threadgroup float2x2 m00;
-threadgroup float2x2 m01;
-threadgroup float2x2 m02;
-threadgroup float2x2 m03;
-threadgroup float2x2 m04;
-threadgroup float2x2 m05;
-threadgroup float2x2 m06;
-threadgroup float2x2 m07;
-threadgroup float2x2 m08;
-threadgroup float2x2 m09;
-threadgroup float2x2 m10;
-threadgroup float2x2 m11;
-threadgroup float2x2 m12;
-threadgroup float2x2 m13;
-threadgroup float2x2 m14;
-threadgroup float2x2 m15;
-threadgroup float2x2 m16;
-threadgroup float2x2 m17;
-threadgroup float2x2 m18;
-threadgroup float2x2 m19;
-threadgroup float2x2 m20;
-threadgroup float2x2 m21;
-threadgroup float2x2 m22;
-threadgroup float2x2 m23;
-threadgroup float2x2 m24;
-threadgroup float2x2 m25;
-threadgroup float2x2 m26;
-threadgroup float2x2 m27;
-threadgroup float2x2 m28;
-threadgroup float2x2 m29;
-threadgroup float2x2 m30;
-threadgroup float2x2 m31;
-threadgroup float2x2 m32;
-threadgroup float2x2 m33;
-threadgroup float2x2 m34;
-threadgroup float2x2 m35;
-threadgroup float2x2 m36;
-threadgroup float2x2 m37;
-threadgroup float2x2 m38;
-threadgroup float2x2 m39;
-threadgroup float2x2 m40;
-threadgroup float2x2 m41;
-threadgroup float2x2 m42;
-threadgroup float2x2 m43;
-threadgroup float2x2 m44;
-threadgroup float2x2 m45;
-threadgroup float2x2 m46;
-threadgroup float2x2 m47;
-threadgroup float2x2 m48;
-threadgroup float2x2 m49;
-threadgroup float2x2 m50;
-threadgroup float2x2 m51;
-threadgroup float2x2 m52;
-threadgroup float2x2 m53;
-threadgroup float2x2 m54;
-threadgroup float2x2 m55;
-threadgroup float2x2 m56;
-threadgroup float2x2 m57;
-threadgroup float2x2 m58;
-threadgroup float2x2 m59;
-threadgroup float2x2 m60;
-threadgroup float2x2 m61;
-threadgroup float2x2 m62;
-threadgroup float2x2 m63;
-threadgroup float2x2 m64;
-threadgroup float2x2 m65;
-threadgroup float2x2 m66;
-threadgroup float2x2 m67;
-threadgroup float2x2 m68;
-threadgroup float2x2 m69;
-threadgroup float2x2 m70;
-threadgroup float2x2 m71;
-threadgroup float2x2 m72;
-threadgroup float2x2 m73;
-threadgroup float2x2 m74;
-threadgroup float2x2 m75;
-threadgroup float2x2 m76;
-threadgroup float2x2 m77;
-threadgroup float2x2 m78;
-threadgroup float2x2 m79;
-threadgroup float2x2 m80;
-threadgroup float2x2 m81;
-threadgroup float2x2 m82;
-threadgroup float2x2 m83;
-threadgroup float2x2 m84;
-threadgroup float2x2 m85;
-threadgroup float2x2 m86;
-threadgroup float2x2 m87;
-threadgroup float2x2 m88;
-threadgroup float2x2 m89;
-threadgroup float2x2 m90;
-threadgroup float2x2 m91;
-threadgroup float2x2 m92;
-threadgroup float2x2 m93;
-threadgroup float2x2 m94;
-threadgroup float2x2 m95;
-threadgroup float2x2 m96;
-threadgroup float2x2 m97;
-threadgroup float2x2 m98;
-threadgroup float2x2 m99;
-kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]]) {
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], threadgroup tint_symbol_101* v) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.m00=(&(*v).tint_symbol_1), .m01=(&(*v).tint_symbol_2), .m02=(&(*v).tint_symbol_3), .m03=(&(*v).tint_symbol_4), .m04=(&(*v).tint_symbol_5), .m05=(&(*v).tint_symbol_6), .m06=(&(*v).tint_symbol_7), .m07=(&(*v).tint_symbol_8), .m08=(&(*v).tint_symbol_9), .m09=(&(*v).tint_symbol_10), .m10=(&(*v).tint_symbol_11), .m11=(&(*v).tint_symbol_12), .m12=(&(*v).tint_symbol_13), .m13=(&(*v).tint_symbol_14), .m14=(&(*v).tint_symbol_15), .m15=(&(*v).tint_symbol_16), .m16=(&(*v).tint_symbol_17), .m17=(&(*v).tint_symbol_18), .m18=(&(*v).tint_symbol_19), .m19=(&(*v).tint_symbol_20), .m20=(&(*v).tint_symbol_21), .m21=(&(*v).tint_symbol_22), .m22=(&(*v).tint_symbol_23), .m23=(&(*v).tint_symbol_24), .m24=(&(*v).tint_symbol_25), .m25=(&(*v).tint_symbol_26), .m26=(&(*v).tint_symbol_27), .m27=(&(*v).tint_symbol_28), .m28=(&(*v).tint_symbol_29), .m29=(&(*v).tint_symbol_30), .m30=(&(*v).tint_symbol_31), .m31=(&(*v).tint_symbol_32), .m32=(&(*v).tint_symbol_33), .m33=(&(*v).tint_symbol_34), .m34=(&(*v).tint_symbol_35), .m35=(&(*v).tint_symbol_36), .m36=(&(*v).tint_symbol_37), .m37=(&(*v).tint_symbol_38), .m38=(&(*v).tint_symbol_39), .m39=(&(*v).tint_symbol_40), .m40=(&(*v).tint_symbol_41), .m41=(&(*v).tint_symbol_42), .m42=(&(*v).tint_symbol_43), .m43=(&(*v).tint_symbol_44), .m44=(&(*v).tint_symbol_45), .m45=(&(*v).tint_symbol_46), .m46=(&(*v).tint_symbol_47), .m47=(&(*v).tint_symbol_48), .m48=(&(*v).tint_symbol_49), .m49=(&(*v).tint_symbol_50), .m50=(&(*v).tint_symbol_51), .m51=(&(*v).tint_symbol_52), .m52=(&(*v).tint_symbol_53), .m53=(&(*v).tint_symbol_54), .m54=(&(*v).tint_symbol_55), .m55=(&(*v).tint_symbol_56), .m56=(&(*v).tint_symbol_57), .m57=(&(*v).tint_symbol_58), .m58=(&(*v).tint_symbol_59), .m59=(&(*v).tint_symbol_60), .m60=(&(*v).tint_symbol_61), .m61=(&(*v).tint_symbol_62), .m62=(&(*v).tint_symbol_63), .m63=(&(*v).tint_symbol_64), .m64=(&(*v).tint_symbol_65), .m65=(&(*v).tint_symbol_66), .m66=(&(*v).tint_symbol_67), .m67=(&(*v).tint_symbol_68), .m68=(&(*v).tint_symbol_69), .m69=(&(*v).tint_symbol_70), .m70=(&(*v).tint_symbol_71), .m71=(&(*v).tint_symbol_72), .m72=(&(*v).tint_symbol_73), .m73=(&(*v).tint_symbol_74), .m74=(&(*v).tint_symbol_75), .m75=(&(*v).tint_symbol_76), .m76=(&(*v).tint_symbol_77), .m77=(&(*v).tint_symbol_78), .m78=(&(*v).tint_symbol_79), .m79=(&(*v).tint_symbol_80), .m80=(&(*v).tint_symbol_81), .m81=(&(*v).tint_symbol_82), .m82=(&(*v).tint_symbol_83), .m83=(&(*v).tint_symbol_84), .m84=(&(*v).tint_symbol_85), .m85=(&(*v).tint_symbol_86), .m86=(&(*v).tint_symbol_87), .m87=(&(*v).tint_symbol_88), .m88=(&(*v).tint_symbol_89), .m89=(&(*v).tint_symbol_90), .m90=(&(*v).tint_symbol_91), .m91=(&(*v).tint_symbol_92), .m92=(&(*v).tint_symbol_93), .m93=(&(*v).tint_symbol_94), .m94=(&(*v).tint_symbol_95), .m95=(&(*v).tint_symbol_96), .m96=(&(*v).tint_symbol_97), .m97=(&(*v).tint_symbol_98), .m98=(&(*v).tint_symbol_99), .m99=(&(*v).tint_symbol_100)};
   if ((idx == 0u)) {
-    m00 = float2x2(float2(0.0f), float2(0.0f));
-    m01 = float2x2(float2(0.0f), float2(0.0f));
-    m02 = float2x2(float2(0.0f), float2(0.0f));
-    m03 = float2x2(float2(0.0f), float2(0.0f));
-    m04 = float2x2(float2(0.0f), float2(0.0f));
-    m05 = float2x2(float2(0.0f), float2(0.0f));
-    m06 = float2x2(float2(0.0f), float2(0.0f));
-    m07 = float2x2(float2(0.0f), float2(0.0f));
-    m08 = float2x2(float2(0.0f), float2(0.0f));
-    m09 = float2x2(float2(0.0f), float2(0.0f));
-    m10 = float2x2(float2(0.0f), float2(0.0f));
-    m11 = float2x2(float2(0.0f), float2(0.0f));
-    m12 = float2x2(float2(0.0f), float2(0.0f));
-    m13 = float2x2(float2(0.0f), float2(0.0f));
-    m14 = float2x2(float2(0.0f), float2(0.0f));
-    m15 = float2x2(float2(0.0f), float2(0.0f));
-    m16 = float2x2(float2(0.0f), float2(0.0f));
-    m17 = float2x2(float2(0.0f), float2(0.0f));
-    m18 = float2x2(float2(0.0f), float2(0.0f));
-    m19 = float2x2(float2(0.0f), float2(0.0f));
-    m20 = float2x2(float2(0.0f), float2(0.0f));
-    m21 = float2x2(float2(0.0f), float2(0.0f));
-    m22 = float2x2(float2(0.0f), float2(0.0f));
-    m23 = float2x2(float2(0.0f), float2(0.0f));
-    m24 = float2x2(float2(0.0f), float2(0.0f));
-    m25 = float2x2(float2(0.0f), float2(0.0f));
-    m26 = float2x2(float2(0.0f), float2(0.0f));
-    m27 = float2x2(float2(0.0f), float2(0.0f));
-    m28 = float2x2(float2(0.0f), float2(0.0f));
-    m29 = float2x2(float2(0.0f), float2(0.0f));
-    m30 = float2x2(float2(0.0f), float2(0.0f));
-    m31 = float2x2(float2(0.0f), float2(0.0f));
-    m32 = float2x2(float2(0.0f), float2(0.0f));
-    m33 = float2x2(float2(0.0f), float2(0.0f));
-    m34 = float2x2(float2(0.0f), float2(0.0f));
-    m35 = float2x2(float2(0.0f), float2(0.0f));
-    m36 = float2x2(float2(0.0f), float2(0.0f));
-    m37 = float2x2(float2(0.0f), float2(0.0f));
-    m38 = float2x2(float2(0.0f), float2(0.0f));
-    m39 = float2x2(float2(0.0f), float2(0.0f));
-    m40 = float2x2(float2(0.0f), float2(0.0f));
-    m41 = float2x2(float2(0.0f), float2(0.0f));
-    m42 = float2x2(float2(0.0f), float2(0.0f));
-    m43 = float2x2(float2(0.0f), float2(0.0f));
-    m44 = float2x2(float2(0.0f), float2(0.0f));
-    m45 = float2x2(float2(0.0f), float2(0.0f));
-    m46 = float2x2(float2(0.0f), float2(0.0f));
-    m47 = float2x2(float2(0.0f), float2(0.0f));
-    m48 = float2x2(float2(0.0f), float2(0.0f));
-    m49 = float2x2(float2(0.0f), float2(0.0f));
-    m50 = float2x2(float2(0.0f), float2(0.0f));
-    m51 = float2x2(float2(0.0f), float2(0.0f));
-    m52 = float2x2(float2(0.0f), float2(0.0f));
-    m53 = float2x2(float2(0.0f), float2(0.0f));
-    m54 = float2x2(float2(0.0f), float2(0.0f));
-    m55 = float2x2(float2(0.0f), float2(0.0f));
-    m56 = float2x2(float2(0.0f), float2(0.0f));
-    m57 = float2x2(float2(0.0f), float2(0.0f));
-    m58 = float2x2(float2(0.0f), float2(0.0f));
-    m59 = float2x2(float2(0.0f), float2(0.0f));
-    m60 = float2x2(float2(0.0f), float2(0.0f));
-    m61 = float2x2(float2(0.0f), float2(0.0f));
-    m62 = float2x2(float2(0.0f), float2(0.0f));
-    m63 = float2x2(float2(0.0f), float2(0.0f));
-    m64 = float2x2(float2(0.0f), float2(0.0f));
-    m65 = float2x2(float2(0.0f), float2(0.0f));
-    m66 = float2x2(float2(0.0f), float2(0.0f));
-    m67 = float2x2(float2(0.0f), float2(0.0f));
-    m68 = float2x2(float2(0.0f), float2(0.0f));
-    m69 = float2x2(float2(0.0f), float2(0.0f));
-    m70 = float2x2(float2(0.0f), float2(0.0f));
-    m71 = float2x2(float2(0.0f), float2(0.0f));
-    m72 = float2x2(float2(0.0f), float2(0.0f));
-    m73 = float2x2(float2(0.0f), float2(0.0f));
-    m74 = float2x2(float2(0.0f), float2(0.0f));
-    m75 = float2x2(float2(0.0f), float2(0.0f));
-    m76 = float2x2(float2(0.0f), float2(0.0f));
-    m77 = float2x2(float2(0.0f), float2(0.0f));
-    m78 = float2x2(float2(0.0f), float2(0.0f));
-    m79 = float2x2(float2(0.0f), float2(0.0f));
-    m80 = float2x2(float2(0.0f), float2(0.0f));
-    m81 = float2x2(float2(0.0f), float2(0.0f));
-    m82 = float2x2(float2(0.0f), float2(0.0f));
-    m83 = float2x2(float2(0.0f), float2(0.0f));
-    m84 = float2x2(float2(0.0f), float2(0.0f));
-    m85 = float2x2(float2(0.0f), float2(0.0f));
-    m86 = float2x2(float2(0.0f), float2(0.0f));
-    m87 = float2x2(float2(0.0f), float2(0.0f));
-    m88 = float2x2(float2(0.0f), float2(0.0f));
-    m89 = float2x2(float2(0.0f), float2(0.0f));
-    m90 = float2x2(float2(0.0f), float2(0.0f));
-    m91 = float2x2(float2(0.0f), float2(0.0f));
-    m92 = float2x2(float2(0.0f), float2(0.0f));
-    m93 = float2x2(float2(0.0f), float2(0.0f));
-    m94 = float2x2(float2(0.0f), float2(0.0f));
-    m95 = float2x2(float2(0.0f), float2(0.0f));
-    m96 = float2x2(float2(0.0f), float2(0.0f));
-    m97 = float2x2(float2(0.0f), float2(0.0f));
-    m98 = float2x2(float2(0.0f), float2(0.0f));
-    m99 = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m00) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m01) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m02) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m03) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m04) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m05) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m06) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m07) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m08) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m09) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m10) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m11) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m12) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m13) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m14) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m15) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m16) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m17) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m18) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m19) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m20) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m21) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m22) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m23) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m24) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m25) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m26) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m27) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m28) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m29) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m30) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m31) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m32) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m33) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m34) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m35) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m36) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m37) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m38) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m39) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m40) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m41) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m42) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m43) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m44) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m45) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m46) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m47) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m48) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m49) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m50) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m51) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m52) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m53) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m54) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m55) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m56) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m57) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m58) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m59) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m60) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m61) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m62) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m63) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m64) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m65) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m66) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m67) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m68) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m69) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m70) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m71) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m72) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m73) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m74) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m75) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m76) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m77) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m78) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m79) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m80) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m81) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m82) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m83) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m84) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m85) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m86) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m87) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m88) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m89) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m90) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m91) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m92) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m93) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m94) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m95) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m96) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m97) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m98) = float2x2(float2(0.0f), float2(0.0f));
+    (*tint_module_vars.m99) = float2x2(float2(0.0f), float2(0.0f));
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  m00[0][0] = 1.0f;
-  m01[0][0] = 1.0f;
-  m02[0][0] = 1.0f;
-  m03[0][0] = 1.0f;
-  m04[0][0] = 1.0f;
-  m05[0][0] = 1.0f;
-  m06[0][0] = 1.0f;
-  m07[0][0] = 1.0f;
-  m08[0][0] = 1.0f;
-  m09[0][0] = 1.0f;
-  m10[0][0] = 1.0f;
-  m11[0][0] = 1.0f;
-  m12[0][0] = 1.0f;
-  m13[0][0] = 1.0f;
-  m14[0][0] = 1.0f;
-  m15[0][0] = 1.0f;
-  m16[0][0] = 1.0f;
-  m17[0][0] = 1.0f;
-  m18[0][0] = 1.0f;
-  m19[0][0] = 1.0f;
-  m20[0][0] = 1.0f;
-  m21[0][0] = 1.0f;
-  m22[0][0] = 1.0f;
-  m23[0][0] = 1.0f;
-  m24[0][0] = 1.0f;
-  m25[0][0] = 1.0f;
-  m26[0][0] = 1.0f;
-  m27[0][0] = 1.0f;
-  m28[0][0] = 1.0f;
-  m29[0][0] = 1.0f;
-  m30[0][0] = 1.0f;
-  m31[0][0] = 1.0f;
-  m32[0][0] = 1.0f;
-  m33[0][0] = 1.0f;
-  m34[0][0] = 1.0f;
-  m35[0][0] = 1.0f;
-  m36[0][0] = 1.0f;
-  m37[0][0] = 1.0f;
-  m38[0][0] = 1.0f;
-  m39[0][0] = 1.0f;
-  m40[0][0] = 1.0f;
-  m41[0][0] = 1.0f;
-  m42[0][0] = 1.0f;
-  m43[0][0] = 1.0f;
-  m44[0][0] = 1.0f;
-  m45[0][0] = 1.0f;
-  m46[0][0] = 1.0f;
-  m47[0][0] = 1.0f;
-  m48[0][0] = 1.0f;
-  m49[0][0] = 1.0f;
-  m50[0][0] = 1.0f;
-  m51[0][0] = 1.0f;
-  m52[0][0] = 1.0f;
-  m53[0][0] = 1.0f;
-  m54[0][0] = 1.0f;
-  m55[0][0] = 1.0f;
-  m56[0][0] = 1.0f;
-  m57[0][0] = 1.0f;
-  m58[0][0] = 1.0f;
-  m59[0][0] = 1.0f;
-  m60[0][0] = 1.0f;
-  m61[0][0] = 1.0f;
-  m62[0][0] = 1.0f;
-  m63[0][0] = 1.0f;
-  m64[0][0] = 1.0f;
-  m65[0][0] = 1.0f;
-  m66[0][0] = 1.0f;
-  m67[0][0] = 1.0f;
-  m68[0][0] = 1.0f;
-  m69[0][0] = 1.0f;
-  m70[0][0] = 1.0f;
-  m71[0][0] = 1.0f;
-  m72[0][0] = 1.0f;
-  m73[0][0] = 1.0f;
-  m74[0][0] = 1.0f;
-  m75[0][0] = 1.0f;
-  m76[0][0] = 1.0f;
-  m77[0][0] = 1.0f;
-  m78[0][0] = 1.0f;
-  m79[0][0] = 1.0f;
-  m80[0][0] = 1.0f;
-  m81[0][0] = 1.0f;
-  m82[0][0] = 1.0f;
-  m83[0][0] = 1.0f;
-  m84[0][0] = 1.0f;
-  m85[0][0] = 1.0f;
-  m86[0][0] = 1.0f;
-  m87[0][0] = 1.0f;
-  m88[0][0] = 1.0f;
-  m89[0][0] = 1.0f;
-  m90[0][0] = 1.0f;
-  m91[0][0] = 1.0f;
-  m92[0][0] = 1.0f;
-  m93[0][0] = 1.0f;
-  m94[0][0] = 1.0f;
-  m95[0][0] = 1.0f;
-  m96[0][0] = 1.0f;
-  m97[0][0] = 1.0f;
-  m98[0][0] = 1.0f;
-  m99[0][0] = 1.0f;
+  (*tint_module_vars.m00)[0][0] = 1.0f;
+  (*tint_module_vars.m01)[0][0] = 1.0f;
+  (*tint_module_vars.m02)[0][0] = 1.0f;
+  (*tint_module_vars.m03)[0][0] = 1.0f;
+  (*tint_module_vars.m04)[0][0] = 1.0f;
+  (*tint_module_vars.m05)[0][0] = 1.0f;
+  (*tint_module_vars.m06)[0][0] = 1.0f;
+  (*tint_module_vars.m07)[0][0] = 1.0f;
+  (*tint_module_vars.m08)[0][0] = 1.0f;
+  (*tint_module_vars.m09)[0][0] = 1.0f;
+  (*tint_module_vars.m10)[0][0] = 1.0f;
+  (*tint_module_vars.m11)[0][0] = 1.0f;
+  (*tint_module_vars.m12)[0][0] = 1.0f;
+  (*tint_module_vars.m13)[0][0] = 1.0f;
+  (*tint_module_vars.m14)[0][0] = 1.0f;
+  (*tint_module_vars.m15)[0][0] = 1.0f;
+  (*tint_module_vars.m16)[0][0] = 1.0f;
+  (*tint_module_vars.m17)[0][0] = 1.0f;
+  (*tint_module_vars.m18)[0][0] = 1.0f;
+  (*tint_module_vars.m19)[0][0] = 1.0f;
+  (*tint_module_vars.m20)[0][0] = 1.0f;
+  (*tint_module_vars.m21)[0][0] = 1.0f;
+  (*tint_module_vars.m22)[0][0] = 1.0f;
+  (*tint_module_vars.m23)[0][0] = 1.0f;
+  (*tint_module_vars.m24)[0][0] = 1.0f;
+  (*tint_module_vars.m25)[0][0] = 1.0f;
+  (*tint_module_vars.m26)[0][0] = 1.0f;
+  (*tint_module_vars.m27)[0][0] = 1.0f;
+  (*tint_module_vars.m28)[0][0] = 1.0f;
+  (*tint_module_vars.m29)[0][0] = 1.0f;
+  (*tint_module_vars.m30)[0][0] = 1.0f;
+  (*tint_module_vars.m31)[0][0] = 1.0f;
+  (*tint_module_vars.m32)[0][0] = 1.0f;
+  (*tint_module_vars.m33)[0][0] = 1.0f;
+  (*tint_module_vars.m34)[0][0] = 1.0f;
+  (*tint_module_vars.m35)[0][0] = 1.0f;
+  (*tint_module_vars.m36)[0][0] = 1.0f;
+  (*tint_module_vars.m37)[0][0] = 1.0f;
+  (*tint_module_vars.m38)[0][0] = 1.0f;
+  (*tint_module_vars.m39)[0][0] = 1.0f;
+  (*tint_module_vars.m40)[0][0] = 1.0f;
+  (*tint_module_vars.m41)[0][0] = 1.0f;
+  (*tint_module_vars.m42)[0][0] = 1.0f;
+  (*tint_module_vars.m43)[0][0] = 1.0f;
+  (*tint_module_vars.m44)[0][0] = 1.0f;
+  (*tint_module_vars.m45)[0][0] = 1.0f;
+  (*tint_module_vars.m46)[0][0] = 1.0f;
+  (*tint_module_vars.m47)[0][0] = 1.0f;
+  (*tint_module_vars.m48)[0][0] = 1.0f;
+  (*tint_module_vars.m49)[0][0] = 1.0f;
+  (*tint_module_vars.m50)[0][0] = 1.0f;
+  (*tint_module_vars.m51)[0][0] = 1.0f;
+  (*tint_module_vars.m52)[0][0] = 1.0f;
+  (*tint_module_vars.m53)[0][0] = 1.0f;
+  (*tint_module_vars.m54)[0][0] = 1.0f;
+  (*tint_module_vars.m55)[0][0] = 1.0f;
+  (*tint_module_vars.m56)[0][0] = 1.0f;
+  (*tint_module_vars.m57)[0][0] = 1.0f;
+  (*tint_module_vars.m58)[0][0] = 1.0f;
+  (*tint_module_vars.m59)[0][0] = 1.0f;
+  (*tint_module_vars.m60)[0][0] = 1.0f;
+  (*tint_module_vars.m61)[0][0] = 1.0f;
+  (*tint_module_vars.m62)[0][0] = 1.0f;
+  (*tint_module_vars.m63)[0][0] = 1.0f;
+  (*tint_module_vars.m64)[0][0] = 1.0f;
+  (*tint_module_vars.m65)[0][0] = 1.0f;
+  (*tint_module_vars.m66)[0][0] = 1.0f;
+  (*tint_module_vars.m67)[0][0] = 1.0f;
+  (*tint_module_vars.m68)[0][0] = 1.0f;
+  (*tint_module_vars.m69)[0][0] = 1.0f;
+  (*tint_module_vars.m70)[0][0] = 1.0f;
+  (*tint_module_vars.m71)[0][0] = 1.0f;
+  (*tint_module_vars.m72)[0][0] = 1.0f;
+  (*tint_module_vars.m73)[0][0] = 1.0f;
+  (*tint_module_vars.m74)[0][0] = 1.0f;
+  (*tint_module_vars.m75)[0][0] = 1.0f;
+  (*tint_module_vars.m76)[0][0] = 1.0f;
+  (*tint_module_vars.m77)[0][0] = 1.0f;
+  (*tint_module_vars.m78)[0][0] = 1.0f;
+  (*tint_module_vars.m79)[0][0] = 1.0f;
+  (*tint_module_vars.m80)[0][0] = 1.0f;
+  (*tint_module_vars.m81)[0][0] = 1.0f;
+  (*tint_module_vars.m82)[0][0] = 1.0f;
+  (*tint_module_vars.m83)[0][0] = 1.0f;
+  (*tint_module_vars.m84)[0][0] = 1.0f;
+  (*tint_module_vars.m85)[0][0] = 1.0f;
+  (*tint_module_vars.m86)[0][0] = 1.0f;
+  (*tint_module_vars.m87)[0][0] = 1.0f;
+  (*tint_module_vars.m88)[0][0] = 1.0f;
+  (*tint_module_vars.m89)[0][0] = 1.0f;
+  (*tint_module_vars.m90)[0][0] = 1.0f;
+  (*tint_module_vars.m91)[0][0] = 1.0f;
+  (*tint_module_vars.m92)[0][0] = 1.0f;
+  (*tint_module_vars.m93)[0][0] = 1.0f;
+  (*tint_module_vars.m94)[0][0] = 1.0f;
+  (*tint_module_vars.m95)[0][0] = 1.0f;
+  (*tint_module_vars.m96)[0][0] = 1.0f;
+  (*tint_module_vars.m97)[0][0] = 1.0f;
+  (*tint_module_vars.m98)[0][0] = 1.0f;
+  (*tint_module_vars.m99)[0][0] = 1.0f;
 }
-program_source:4:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m00;
-                     ^
-program_source:5:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m01;
-                     ^
-program_source:6:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m02;
-                     ^
-program_source:7:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m03;
-                     ^
-program_source:8:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m04;
-                     ^
-program_source:9:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m05;
-                     ^
-program_source:10:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m06;
-                     ^
-program_source:11:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m07;
-                     ^
-program_source:12:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m08;
-                     ^
-program_source:13:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m09;
-                     ^
-program_source:14:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m10;
-                     ^
-program_source:15:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m11;
-                     ^
-program_source:16:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m12;
-                     ^
-program_source:17:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m13;
-                     ^
-program_source:18:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m14;
-                     ^
-program_source:19:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m15;
-                     ^
-program_source:20:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m16;
-                     ^
-program_source:21:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m17;
-                     ^
-program_source:22:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m18;
-                     ^
-program_source:23:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m19;
-                     ^
-program_source:24:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m20;
-                     ^
-program_source:25:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m21;
-                     ^
-program_source:26:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m22;
-                     ^
-program_source:27:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m23;
-                     ^
-program_source:28:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m24;
-                     ^
-program_source:29:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m25;
-                     ^
-program_source:30:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m26;
-                     ^
-program_source:31:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m27;
-                     ^
-program_source:32:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m28;
-                     ^
-program_source:33:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m29;
-                     ^
-program_source:34:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m30;
-                     ^
-program_source:35:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m31;
-                     ^
-program_source:36:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m32;
-                     ^
-program_source:37:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m33;
-                     ^
-program_source:38:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m34;
-                     ^
-program_source:39:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m35;
-                     ^
-program_source:40:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m36;
-                     ^
-program_source:41:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m37;
-                     ^
-program_source:42:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m38;
-                     ^
-program_source:43:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m39;
-                     ^
-program_source:44:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m40;
-                     ^
-program_source:45:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m41;
-                     ^
-program_source:46:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m42;
-                     ^
-program_source:47:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m43;
-                     ^
-program_source:48:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m44;
-                     ^
-program_source:49:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m45;
-                     ^
-program_source:50:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m46;
-                     ^
-program_source:51:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m47;
-                     ^
-program_source:52:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m48;
-                     ^
-program_source:53:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m49;
-                     ^
-program_source:54:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m50;
-                     ^
-program_source:55:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m51;
-                     ^
-program_source:56:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m52;
-                     ^
-program_source:57:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m53;
-                     ^
-program_source:58:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m54;
-                     ^
-program_source:59:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m55;
-                     ^
-program_source:60:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m56;
-                     ^
-program_source:61:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m57;
-                     ^
-program_source:62:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m58;
-                     ^
-program_source:63:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m59;
-                     ^
-program_source:64:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m60;
-                     ^
-program_source:65:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m61;
-                     ^
-program_source:66:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m62;
-                     ^
-program_source:67:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m63;
-                     ^
-program_source:68:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m64;
-                     ^
-program_source:69:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m65;
-                     ^
-program_source:70:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m66;
-                     ^
-program_source:71:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m67;
-                     ^
-program_source:72:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m68;
-                     ^
-program_source:73:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m69;
-                     ^
-program_source:74:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m70;
-                     ^
-program_source:75:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m71;
-                     ^
-program_source:76:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m72;
-                     ^
-program_source:77:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m73;
-                     ^
-program_source:78:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m74;
-                     ^
-program_source:79:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m75;
-                     ^
-program_source:80:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m76;
-                     ^
-program_source:81:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m77;
-                     ^
-program_source:82:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m78;
-                     ^
-program_source:83:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m79;
-                     ^
-program_source:84:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m80;
-                     ^
-program_source:85:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m81;
-                     ^
-program_source:86:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m82;
-                     ^
-program_source:87:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m83;
-                     ^
-program_source:88:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m84;
-                     ^
-program_source:89:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m85;
-                     ^
-program_source:90:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m86;
-                     ^
-program_source:91:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m87;
-                     ^
-program_source:92:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m88;
-                     ^
-program_source:93:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m89;
-                     ^
-program_source:94:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m90;
-                     ^
-program_source:95:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m91;
-                     ^
-program_source:96:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m92;
-                     ^
-program_source:97:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m93;
-                     ^
-program_source:98:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m94;
-                     ^
-program_source:99:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m95;
-                     ^
-program_source:100:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m96;
-                     ^
-program_source:101:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m97;
-                     ^
-program_source:102:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m98;
-                     ^
-program_source:103:22: error: program scope variable must reside in constant address space
-threadgroup float2x2 m99;
-                     ^
-program_source:105:12: warning: equality comparison with extraneous parentheses [-Wparentheses-equality]
-  if ((idx == 0u)) {
-       ~~~~^~~~~
-program_source:105:12: note: remove extraneous parentheses around the comparison to silence this warning
-  if ((idx == 0u)) {
-      ~    ^    ~
-program_source:105:12: note: use '=' to turn this equality comparison into an assignment
-  if ((idx == 0u)) {
-           ^~
-           =
-