diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
index 896696e..943a64c 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
@@ -1,31 +1,20 @@
 #include <metal_stdlib>
 
 using namespace metal;
-struct tint_array_wrapper {
-  uint arr[1];
+
+template<typename T, size_t N>
+struct tint_array {
+    const constant T& operator[](size_t i) const constant { return elements[i]; }
+    device T& operator[](size_t i) device { return elements[i]; }
+    const device T& operator[](size_t i) const device { return elements[i]; }
+    thread T& operator[](size_t i) thread { return elements[i]; }
+    const thread T& operator[](size_t i) const thread { return elements[i]; }
+    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+    T elements[N];
 };
 
-struct tint_array_wrapper_1 {
-  tint_array_wrapper arr[2];
-};
-
-struct tint_array_wrapper_2 {
-  tint_array_wrapper_1 arr[3];
-};
-
-struct tint_array_wrapper_5 {
-  atomic_uint arr[1];
-};
-
-struct tint_array_wrapper_4 {
-  tint_array_wrapper_5 arr[2];
-};
-
-struct tint_array_wrapper_3 {
-  tint_array_wrapper_4 arr[3];
-};
-
-void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrapper_3* const tint_symbol) {
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
   uint idx = 0u;
   idx = local_invocation_index;
   while (true) {
@@ -36,29 +25,29 @@
     uint const x_31 = idx;
     uint const x_33 = idx;
     uint const x_35 = idx;
-    atomic_store_explicit(&((*(tint_symbol)).arr[(x_31 / 2u)].arr[(x_33 % 2u)].arr[(x_35 % 1u)]), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol))[(x_31 / 2u)][(x_33 % 2u)][(x_35 % 1u)]), 0u, memory_order_relaxed);
     {
       uint const x_42 = idx;
       idx = (x_42 + 1u);
     }
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  atomic_store_explicit(&((*(tint_symbol)).arr[2].arr[1].arr[0]), 1u, memory_order_relaxed);
+  atomic_store_explicit(&((*(tint_symbol))[2][1][0]), 1u, memory_order_relaxed);
   return;
 }
 
-void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array_wrapper_3* const tint_symbol_2) {
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_2) {
   uint const x_57 = *(tint_symbol_1);
   compute_main_inner(x_57, tint_symbol_2);
   return;
 }
 
-void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array_wrapper_3* const tint_symbol_3, thread uint* const tint_symbol_4) {
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_3, thread uint* const tint_symbol_4) {
   for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
     uint const i = (idx_1 / 2u);
     uint const i_1 = (idx_1 % 2u);
     uint const i_2 = (idx_1 % 1u);
-    atomic_store_explicit(&((*(tint_symbol_3)).arr[i].arr[i_1].arr[i_2]), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol_3))[i][i_1][i_2]), 0u, memory_order_relaxed);
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
   *(tint_symbol_4) = local_invocation_index_1_param;
@@ -66,7 +55,7 @@
 }
 
 kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
-  threadgroup tint_array_wrapper_3 tint_symbol_5;
+  threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_5;
   thread uint tint_symbol_6 = 0u;
   compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
   return;
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.glsl b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.glsl
new file mode 100644
index 0000000..598fde9
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.glsl
@@ -0,0 +1,21 @@
+#version 310 es
+
+shared uint wg[3][2][1];
+void compute_main(uint local_invocation_index) {
+  {
+    for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
+      uint i = (idx / 2u);
+      uint i_1 = (idx % 2u);
+      uint i_2 = (idx % 1u);
+      atomicExchange(wg[i][i_1][i_2], 0u);
+    }
+  }
+  barrier();
+  atomicExchange(wg[2][1][0], 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.hlsl
new file mode 100644
index 0000000..6f993d5
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.hlsl
@@ -0,0 +1,26 @@
+groupshared uint wg[3][2][1];
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+  {
+    [loop] for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
+      const uint i = (idx / 2u);
+      const uint i_1 = (idx % 2u);
+      const uint i_2 = (idx % 1u);
+      uint atomic_result = 0u;
+      InterlockedExchange(wg[i][i_1][i_2], 0u, atomic_result);
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg[2][1][0], 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner(tint_symbol.local_invocation_index);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl
new file mode 100644
index 0000000..51e5245
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#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];
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
+  for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
+    uint const i = (idx / 2u);
+    uint const i_1 = (idx % 2u);
+    uint const i_2 = (idx % 1u);
+    atomic_store_explicit(&((*(tint_symbol))[i][i_1][i_2]), 0u, memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol))[2][1][0]), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_1;
+  compute_main_inner(local_invocation_index, &(tint_symbol_1));
+  return;
+}
+
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.spvasm
new file mode 100644
index 0000000..0dbd964
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.spvasm
@@ -0,0 +1,88 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 58
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %idx "idx"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpDecorate %_arr_uint_uint_1 ArrayStride 4
+               OpDecorate %_arr__arr_uint_uint_1_uint_2 ArrayStride 4
+               OpDecorate %_arr__arr__arr_uint_uint_1_uint_2_uint_3 ArrayStride 8
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+     %uint_1 = OpConstant %uint 1
+%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
+     %uint_2 = OpConstant %uint 2
+%_arr__arr_uint_uint_1_uint_2 = OpTypeArray %_arr_uint_uint_1 %uint_2
+     %uint_3 = OpConstant %uint 3
+%_arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypeArray %_arr__arr_uint_uint_1_uint_2 %uint_3
+%_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypePointer Workgroup %_arr__arr__arr_uint_uint_1_uint_2_uint_3
+         %wg = OpVariable %_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 Workgroup
+       %void = OpTypeVoid
+         %12 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %19 = OpConstantNull %uint
+     %uint_6 = OpConstant %uint 6
+       %bool = OpTypeBool
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+        %int = OpTypeInt 32 1
+      %int_2 = OpConstant %int 2
+      %int_1 = OpConstant %int 1
+         %51 = OpConstantNull %int
+         %53 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %12
+%local_invocation_index = OpFunctionParameter %uint
+         %16 = OpLabel
+        %idx = OpVariable %_ptr_Function_uint Function %19
+               OpStore %idx %local_invocation_index
+               OpBranch %20
+         %20 = OpLabel
+               OpLoopMerge %21 %22 None
+               OpBranch %23
+         %23 = OpLabel
+         %25 = OpLoad %uint %idx
+         %27 = OpULessThan %bool %25 %uint_6
+         %24 = OpLogicalNot %bool %27
+               OpSelectionMerge %29 None
+               OpBranchConditional %24 %30 %29
+         %30 = OpLabel
+               OpBranch %21
+         %29 = OpLabel
+         %31 = OpLoad %uint %idx
+         %32 = OpUDiv %uint %31 %uint_2
+         %33 = OpLoad %uint %idx
+         %34 = OpUMod %uint %33 %uint_2
+         %35 = OpLoad %uint %idx
+         %36 = OpUMod %uint %35 %uint_1
+         %41 = OpAccessChain %_ptr_Workgroup_uint %wg %32 %34 %36
+               OpAtomicStore %41 %uint_2 %uint_0 %19
+               OpBranch %22
+         %22 = OpLabel
+         %42 = OpLoad %uint %idx
+         %43 = OpIAdd %uint %42 %uint_1
+               OpStore %idx %43
+               OpBranch %20
+         %21 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %52 = OpAccessChain %_ptr_Workgroup_uint %wg %int_2 %int_1 %51
+               OpAtomicStore %52 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %53
+         %55 = OpLabel
+         %57 = OpLoad %uint %local_invocation_index_1
+         %56 = OpFunctionCall %void %compute_main_inner %57
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.wgsl
new file mode 100644
index 0000000..0a5ac8a
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.wgsl
@@ -0,0 +1,14 @@
+type A0 = atomic<u32>;
+
+type A1 = array<A0, 1>;
+
+type A2 = array<A1, 2>;
+
+type A3 = array<A2, 3>;
+
+var<workgroup> wg : A3;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&(wg[2][1][0]), 1u);
+}
diff --git a/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl b/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
index d2ab1ec..8ef9d39 100644
--- a/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
@@ -1,15 +1,20 @@
 #include <metal_stdlib>
 
 using namespace metal;
-struct tint_array_wrapper {
-  uint arr[4];
+
+template<typename T, size_t N>
+struct tint_array {
+    const constant T& operator[](size_t i) const constant { return elements[i]; }
+    device T& operator[](size_t i) device { return elements[i]; }
+    const device T& operator[](size_t i) const device { return elements[i]; }
+    thread T& operator[](size_t i) thread { return elements[i]; }
+    const thread T& operator[](size_t i) const thread { return elements[i]; }
+    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+    T elements[N];
 };
 
-struct tint_array_wrapper_1 {
-  atomic_uint arr[4];
-};
-
-void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrapper_1* const tint_symbol) {
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<atomic_uint, 4>* const tint_symbol) {
   uint idx = 0u;
   idx = local_invocation_index;
   while (true) {
@@ -18,27 +23,27 @@
       break;
     }
     uint const x_26 = idx;
-    atomic_store_explicit(&((*(tint_symbol)).arr[x_26]), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol))[x_26]), 0u, memory_order_relaxed);
     {
       uint const x_33 = idx;
       idx = (x_33 + 1u);
     }
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  atomic_store_explicit(&((*(tint_symbol)).arr[1]), 1u, memory_order_relaxed);
+  atomic_store_explicit(&((*(tint_symbol))[1]), 1u, memory_order_relaxed);
   return;
 }
 
-void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array_wrapper_1* const tint_symbol_2) {
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array<atomic_uint, 4>* const tint_symbol_2) {
   uint const x_47 = *(tint_symbol_1);
   compute_main_inner(x_47, tint_symbol_2);
   return;
 }
 
-void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array_wrapper_1* const tint_symbol_3, thread uint* const tint_symbol_4) {
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<atomic_uint, 4>* const tint_symbol_3, thread uint* const tint_symbol_4) {
   for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 4u); idx_1 = (idx_1 + 1u)) {
     uint const i = idx_1;
-    atomic_store_explicit(&((*(tint_symbol_3)).arr[i]), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol_3))[i]), 0u, memory_order_relaxed);
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
   *(tint_symbol_4) = local_invocation_index_1_param;
@@ -46,7 +51,7 @@
 }
 
 kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
-  threadgroup tint_array_wrapper_1 tint_symbol_5;
+  threadgroup tint_array<atomic_uint, 4> tint_symbol_5;
   thread uint tint_symbol_6 = 0u;
   compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
   return;
diff --git a/test/tint/builtins/atomicStore/array/array.wgsl.expected.glsl b/test/tint/builtins/atomicStore/array/array.wgsl.expected.glsl
new file mode 100644
index 0000000..653bae6
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.glsl
@@ -0,0 +1,19 @@
+#version 310 es
+
+shared uint wg[4];
+void compute_main(uint local_invocation_index) {
+  {
+    for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+      uint i = idx;
+      atomicExchange(wg[i], 0u);
+    }
+  }
+  barrier();
+  atomicExchange(wg[1], 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/array/array.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/array/array.wgsl.expected.hlsl
new file mode 100644
index 0000000..d2ff575
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.hlsl
@@ -0,0 +1,24 @@
+groupshared uint wg[4];
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+  {
+    [loop] for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+      const uint i = idx;
+      uint atomic_result = 0u;
+      InterlockedExchange(wg[i], 0u, atomic_result);
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg[1], 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner(tint_symbol.local_invocation_index);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl b/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl
new file mode 100644
index 0000000..5f20d27
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl
@@ -0,0 +1,31 @@
+#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];
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<atomic_uint, 4>* const tint_symbol) {
+  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+    uint const i = idx;
+    atomic_store_explicit(&((*(tint_symbol))[i]), 0u, memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol))[1]), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup tint_array<atomic_uint, 4> tint_symbol_1;
+  compute_main_inner(local_invocation_index, &(tint_symbol_1));
+  return;
+}
+
diff --git a/test/tint/builtins/atomicStore/array/array.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/array/array.wgsl.expected.spvasm
new file mode 100644
index 0000000..f41ddf7
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.spvasm
@@ -0,0 +1,76 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %idx "idx"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpDecorate %_arr_uint_uint_4 ArrayStride 4
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+     %uint_4 = OpConstant %uint 4
+%_arr_uint_uint_4 = OpTypeArray %uint %uint_4
+%_ptr_Workgroup__arr_uint_uint_4 = OpTypePointer Workgroup %_arr_uint_uint_4
+         %wg = OpVariable %_ptr_Workgroup__arr_uint_uint_4 Workgroup
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %15 = OpConstantNull %uint
+       %bool = OpTypeBool
+     %uint_2 = OpConstant %uint 2
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+     %uint_1 = OpConstant %uint 1
+   %uint_264 = OpConstant %uint 264
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+         %43 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %8
+%local_invocation_index = OpFunctionParameter %uint
+         %12 = OpLabel
+        %idx = OpVariable %_ptr_Function_uint Function %15
+               OpStore %idx %local_invocation_index
+               OpBranch %16
+         %16 = OpLabel
+               OpLoopMerge %17 %18 None
+               OpBranch %19
+         %19 = OpLabel
+         %21 = OpLoad %uint %idx
+         %22 = OpULessThan %bool %21 %uint_4
+         %20 = OpLogicalNot %bool %22
+               OpSelectionMerge %24 None
+               OpBranchConditional %20 %25 %24
+         %25 = OpLabel
+               OpBranch %17
+         %24 = OpLabel
+         %26 = OpLoad %uint %idx
+         %32 = OpAccessChain %_ptr_Workgroup_uint %wg %26
+               OpAtomicStore %32 %uint_2 %uint_0 %15
+               OpBranch %18
+         %18 = OpLabel
+         %33 = OpLoad %uint %idx
+         %35 = OpIAdd %uint %33 %uint_1
+               OpStore %idx %35
+               OpBranch %16
+         %17 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %42 = OpAccessChain %_ptr_Workgroup_uint %wg %int_1
+               OpAtomicStore %42 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %43
+         %45 = OpLabel
+         %47 = OpLoad %uint %local_invocation_index_1
+         %46 = OpFunctionCall %void %compute_main_inner %47
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/array/array.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/array/array.wgsl.expected.wgsl
new file mode 100644
index 0000000..f6812d4
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+var<workgroup> wg : array<atomic<u32>, 4>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&(wg[1]), 1u);
+}
diff --git a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
index 896696e..943a64c 100644
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
@@ -1,31 +1,20 @@
 #include <metal_stdlib>
 
 using namespace metal;
-struct tint_array_wrapper {
-  uint arr[1];
+
+template<typename T, size_t N>
+struct tint_array {
+    const constant T& operator[](size_t i) const constant { return elements[i]; }
+    device T& operator[](size_t i) device { return elements[i]; }
+    const device T& operator[](size_t i) const device { return elements[i]; }
+    thread T& operator[](size_t i) thread { return elements[i]; }
+    const thread T& operator[](size_t i) const thread { return elements[i]; }
+    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+    T elements[N];
 };
 
-struct tint_array_wrapper_1 {
-  tint_array_wrapper arr[2];
-};
-
-struct tint_array_wrapper_2 {
-  tint_array_wrapper_1 arr[3];
-};
-
-struct tint_array_wrapper_5 {
-  atomic_uint arr[1];
-};
-
-struct tint_array_wrapper_4 {
-  tint_array_wrapper_5 arr[2];
-};
-
-struct tint_array_wrapper_3 {
-  tint_array_wrapper_4 arr[3];
-};
-
-void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrapper_3* const tint_symbol) {
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
   uint idx = 0u;
   idx = local_invocation_index;
   while (true) {
@@ -36,29 +25,29 @@
     uint const x_31 = idx;
     uint const x_33 = idx;
     uint const x_35 = idx;
-    atomic_store_explicit(&((*(tint_symbol)).arr[(x_31 / 2u)].arr[(x_33 % 2u)].arr[(x_35 % 1u)]), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol))[(x_31 / 2u)][(x_33 % 2u)][(x_35 % 1u)]), 0u, memory_order_relaxed);
     {
       uint const x_42 = idx;
       idx = (x_42 + 1u);
     }
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  atomic_store_explicit(&((*(tint_symbol)).arr[2].arr[1].arr[0]), 1u, memory_order_relaxed);
+  atomic_store_explicit(&((*(tint_symbol))[2][1][0]), 1u, memory_order_relaxed);
   return;
 }
 
-void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array_wrapper_3* const tint_symbol_2) {
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_2) {
   uint const x_57 = *(tint_symbol_1);
   compute_main_inner(x_57, tint_symbol_2);
   return;
 }
 
-void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array_wrapper_3* const tint_symbol_3, thread uint* const tint_symbol_4) {
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_3, thread uint* const tint_symbol_4) {
   for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
     uint const i = (idx_1 / 2u);
     uint const i_1 = (idx_1 % 2u);
     uint const i_2 = (idx_1 % 1u);
-    atomic_store_explicit(&((*(tint_symbol_3)).arr[i].arr[i_1].arr[i_2]), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol_3))[i][i_1][i_2]), 0u, memory_order_relaxed);
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
   *(tint_symbol_4) = local_invocation_index_1_param;
@@ -66,7 +55,7 @@
 }
 
 kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
-  threadgroup tint_array_wrapper_3 tint_symbol_5;
+  threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_5;
   thread uint tint_symbol_6 = 0u;
   compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
   return;
diff --git a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.glsl b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.glsl
new file mode 100644
index 0000000..598fde9
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.glsl
@@ -0,0 +1,21 @@
+#version 310 es
+
+shared uint wg[3][2][1];
+void compute_main(uint local_invocation_index) {
+  {
+    for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
+      uint i = (idx / 2u);
+      uint i_1 = (idx % 2u);
+      uint i_2 = (idx % 1u);
+      atomicExchange(wg[i][i_1][i_2], 0u);
+    }
+  }
+  barrier();
+  atomicExchange(wg[2][1][0], 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.hlsl
new file mode 100644
index 0000000..6f993d5
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.hlsl
@@ -0,0 +1,26 @@
+groupshared uint wg[3][2][1];
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+  {
+    [loop] for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
+      const uint i = (idx / 2u);
+      const uint i_1 = (idx % 2u);
+      const uint i_2 = (idx % 1u);
+      uint atomic_result = 0u;
+      InterlockedExchange(wg[i][i_1][i_2], 0u, atomic_result);
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg[2][1][0], 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner(tint_symbol.local_invocation_index);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl
new file mode 100644
index 0000000..51e5245
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#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];
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
+  for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
+    uint const i = (idx / 2u);
+    uint const i_1 = (idx % 2u);
+    uint const i_2 = (idx % 1u);
+    atomic_store_explicit(&((*(tint_symbol))[i][i_1][i_2]), 0u, memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol))[2][1][0]), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_1;
+  compute_main_inner(local_invocation_index, &(tint_symbol_1));
+  return;
+}
+
diff --git a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.spvasm
new file mode 100644
index 0000000..0dbd964
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.spvasm
@@ -0,0 +1,88 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 58
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %idx "idx"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpDecorate %_arr_uint_uint_1 ArrayStride 4
+               OpDecorate %_arr__arr_uint_uint_1_uint_2 ArrayStride 4
+               OpDecorate %_arr__arr__arr_uint_uint_1_uint_2_uint_3 ArrayStride 8
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+     %uint_1 = OpConstant %uint 1
+%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
+     %uint_2 = OpConstant %uint 2
+%_arr__arr_uint_uint_1_uint_2 = OpTypeArray %_arr_uint_uint_1 %uint_2
+     %uint_3 = OpConstant %uint 3
+%_arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypeArray %_arr__arr_uint_uint_1_uint_2 %uint_3
+%_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypePointer Workgroup %_arr__arr__arr_uint_uint_1_uint_2_uint_3
+         %wg = OpVariable %_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 Workgroup
+       %void = OpTypeVoid
+         %12 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %19 = OpConstantNull %uint
+     %uint_6 = OpConstant %uint 6
+       %bool = OpTypeBool
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+        %int = OpTypeInt 32 1
+      %int_2 = OpConstant %int 2
+      %int_1 = OpConstant %int 1
+         %51 = OpConstantNull %int
+         %53 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %12
+%local_invocation_index = OpFunctionParameter %uint
+         %16 = OpLabel
+        %idx = OpVariable %_ptr_Function_uint Function %19
+               OpStore %idx %local_invocation_index
+               OpBranch %20
+         %20 = OpLabel
+               OpLoopMerge %21 %22 None
+               OpBranch %23
+         %23 = OpLabel
+         %25 = OpLoad %uint %idx
+         %27 = OpULessThan %bool %25 %uint_6
+         %24 = OpLogicalNot %bool %27
+               OpSelectionMerge %29 None
+               OpBranchConditional %24 %30 %29
+         %30 = OpLabel
+               OpBranch %21
+         %29 = OpLabel
+         %31 = OpLoad %uint %idx
+         %32 = OpUDiv %uint %31 %uint_2
+         %33 = OpLoad %uint %idx
+         %34 = OpUMod %uint %33 %uint_2
+         %35 = OpLoad %uint %idx
+         %36 = OpUMod %uint %35 %uint_1
+         %41 = OpAccessChain %_ptr_Workgroup_uint %wg %32 %34 %36
+               OpAtomicStore %41 %uint_2 %uint_0 %19
+               OpBranch %22
+         %22 = OpLabel
+         %42 = OpLoad %uint %idx
+         %43 = OpIAdd %uint %42 %uint_1
+               OpStore %idx %43
+               OpBranch %20
+         %21 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %52 = OpAccessChain %_ptr_Workgroup_uint %wg %int_2 %int_1 %51
+               OpAtomicStore %52 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %53
+         %55 = OpLabel
+         %57 = OpLoad %uint %local_invocation_index_1
+         %56 = OpFunctionCall %void %compute_main_inner %57
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.wgsl
new file mode 100644
index 0000000..3ca55977
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+var<workgroup> wg : array<array<array<atomic<u32>, 1>, 2>, 3>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&(wg[2][1][0]), 1u);
+}
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
index 5360fda..f4fd526 100644
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
@@ -1,6 +1,19 @@
 #include <metal_stdlib>
 
 using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+    const constant T& operator[](size_t i) const constant { return elements[i]; }
+    device T& operator[](size_t i) device { return elements[i]; }
+    const device T& operator[](size_t i) const device { return elements[i]; }
+    thread T& operator[](size_t i) thread { return elements[i]; }
+    const thread T& operator[](size_t i) const thread { return elements[i]; }
+    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+    T elements[N];
+};
+
 struct S_atomic {
   int x;
   atomic_uint a;
@@ -13,15 +26,7 @@
   uint y;
 };
 
-struct tint_array_wrapper {
-  S arr[10];
-};
-
-struct tint_array_wrapper_1 {
-  S_atomic arr[10];
-};
-
-void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrapper_1* const tint_symbol) {
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<S_atomic, 10>* const tint_symbol) {
   uint idx = 0u;
   idx = local_invocation_index;
   while (true) {
@@ -30,31 +35,31 @@
       break;
     }
     uint const x_28 = idx;
-    (*(tint_symbol)).arr[x_28].x = 0;
-    atomic_store_explicit(&((*(tint_symbol)).arr[x_28].a), 0u, memory_order_relaxed);
-    (*(tint_symbol)).arr[x_28].y = 0u;
+    (*(tint_symbol))[x_28].x = 0;
+    atomic_store_explicit(&((*(tint_symbol))[x_28].a), 0u, memory_order_relaxed);
+    (*(tint_symbol))[x_28].y = 0u;
     {
       uint const x_41 = idx;
       idx = (x_41 + 1u);
     }
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  atomic_store_explicit(&((*(tint_symbol)).arr[4].a), 1u, memory_order_relaxed);
+  atomic_store_explicit(&((*(tint_symbol))[4].a), 1u, memory_order_relaxed);
   return;
 }
 
-void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array_wrapper_1* const tint_symbol_2) {
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array<S_atomic, 10>* const tint_symbol_2) {
   uint const x_53 = *(tint_symbol_1);
   compute_main_inner(x_53, tint_symbol_2);
   return;
 }
 
-void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array_wrapper_1* const tint_symbol_3, thread uint* const tint_symbol_4) {
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<S_atomic, 10>* const tint_symbol_3, thread uint* const tint_symbol_4) {
   for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
     uint const i = idx_1;
-    (*(tint_symbol_3)).arr[i].x = 0;
-    atomic_store_explicit(&((*(tint_symbol_3)).arr[i].a), 0u, memory_order_relaxed);
-    (*(tint_symbol_3)).arr[i].y = 0u;
+    (*(tint_symbol_3))[i].x = 0;
+    atomic_store_explicit(&((*(tint_symbol_3))[i].a), 0u, memory_order_relaxed);
+    (*(tint_symbol_3))[i].y = 0u;
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
   *(tint_symbol_4) = local_invocation_index_1_param;
@@ -62,7 +67,7 @@
 }
 
 kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
-  threadgroup tint_array_wrapper_1 tint_symbol_5;
+  threadgroup tint_array<S_atomic, 10> tint_symbol_5;
   thread uint tint_symbol_6 = 0u;
   compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
   return;
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.glsl b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.glsl
new file mode 100644
index 0000000..bd0a3ba
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.glsl
@@ -0,0 +1,27 @@
+#version 310 es
+
+struct S {
+  int x;
+  uint a;
+  uint y;
+};
+
+shared S wg[10];
+void compute_main(uint local_invocation_index) {
+  {
+    for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
+      uint i = idx;
+      wg[i].x = 0;
+      atomicExchange(wg[i].a, 0u);
+      wg[i].y = 0u;
+    }
+  }
+  barrier();
+  atomicExchange(wg[4].a, 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.hlsl
new file mode 100644
index 0000000..4642386
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.hlsl
@@ -0,0 +1,32 @@
+struct S {
+  int x;
+  uint a;
+  uint y;
+};
+
+groupshared S wg[10];
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+  {
+    [loop] for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
+      const uint i = idx;
+      wg[i].x = 0;
+      uint atomic_result = 0u;
+      InterlockedExchange(wg[i].a, 0u, atomic_result);
+      wg[i].y = 0u;
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg[4].a, 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner(tint_symbol.local_invocation_index);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl
new file mode 100644
index 0000000..978edb1
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl
@@ -0,0 +1,39 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+    const constant T& operator[](size_t i) const constant { return elements[i]; }
+    device T& operator[](size_t i) device { return elements[i]; }
+    const device T& operator[](size_t i) const device { return elements[i]; }
+    thread T& operator[](size_t i) thread { return elements[i]; }
+    const thread T& operator[](size_t i) const thread { return elements[i]; }
+    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+    T elements[N];
+};
+
+struct S {
+  int x;
+  atomic_uint a;
+  uint y;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<S, 10>* const tint_symbol) {
+  for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
+    uint const i = idx;
+    (*(tint_symbol))[i].x = 0;
+    atomic_store_explicit(&((*(tint_symbol))[i].a), 0u, memory_order_relaxed);
+    (*(tint_symbol))[i].y = 0u;
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol))[4].a), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup tint_array<S, 10> tint_symbol_1;
+  compute_main_inner(local_invocation_index, &(tint_symbol_1));
+  return;
+}
+
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..b9e0fed
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.spvasm
@@ -0,0 +1,91 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 54
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S "S"
+               OpMemberName %S 0 "x"
+               OpMemberName %S 1 "a"
+               OpMemberName %S 2 "y"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %idx "idx"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 4
+               OpMemberDecorate %S 2 Offset 8
+               OpDecorate %_arr_S_uint_10 ArrayStride 12
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+          %S = OpTypeStruct %int %uint %uint
+    %uint_10 = OpConstant %uint 10
+%_arr_S_uint_10 = OpTypeArray %S %uint_10
+%_ptr_Workgroup__arr_S_uint_10 = OpTypePointer Workgroup %_arr_S_uint_10
+         %wg = OpVariable %_ptr_Workgroup__arr_S_uint_10 Workgroup
+       %void = OpTypeVoid
+         %10 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %17 = OpConstantNull %uint
+       %bool = OpTypeBool
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %32 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+      %int_4 = OpConstant %int 4
+         %49 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %10
+%local_invocation_index = OpFunctionParameter %uint
+         %14 = OpLabel
+        %idx = OpVariable %_ptr_Function_uint Function %17
+               OpStore %idx %local_invocation_index
+               OpBranch %18
+         %18 = OpLabel
+               OpLoopMerge %19 %20 None
+               OpBranch %21
+         %21 = OpLabel
+         %23 = OpLoad %uint %idx
+         %24 = OpULessThan %bool %23 %uint_10
+         %22 = OpLogicalNot %bool %24
+               OpSelectionMerge %26 None
+               OpBranchConditional %22 %27 %26
+         %27 = OpLabel
+               OpBranch %19
+         %26 = OpLabel
+         %28 = OpLoad %uint %idx
+         %31 = OpAccessChain %_ptr_Workgroup_int %wg %28 %uint_0
+               OpStore %31 %32
+         %38 = OpAccessChain %_ptr_Workgroup_uint %wg %28 %uint_1
+               OpAtomicStore %38 %uint_2 %uint_0 %17
+         %40 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %28 %uint_2
+               OpStore %40 %17
+               OpBranch %20
+         %20 = OpLabel
+         %41 = OpLoad %uint %idx
+         %42 = OpIAdd %uint %41 %uint_1
+               OpStore %idx %42
+               OpBranch %18
+         %19 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %48 = OpAccessChain %_ptr_Workgroup_uint %wg %int_4 %uint_1
+               OpAtomicStore %48 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %49
+         %51 = OpLabel
+         %53 = OpLoad %uint %local_invocation_index_1
+         %52 = OpFunctionCall %void %compute_main_inner %53
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..1def9bd
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+struct S {
+  x : i32,
+  a : atomic<u32>,
+  y : u32,
+}
+
+var<workgroup> wg : array<S, 10>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&(wg[4].a), 1u);
+}
diff --git a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.glsl b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.glsl
new file mode 100644
index 0000000..d3a4b83
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.glsl
@@ -0,0 +1,25 @@
+#version 310 es
+
+struct S {
+  int x;
+  uint a;
+  uint b;
+};
+
+shared S wg;
+void compute_main(uint local_invocation_index) {
+  {
+    wg.x = 0;
+    atomicExchange(wg.a, 0u);
+    atomicExchange(wg.b, 0u);
+  }
+  barrier();
+  atomicExchange(wg.a, 1u);
+  atomicExchange(wg.b, 2u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.hlsl
new file mode 100644
index 0000000..35b428f
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.hlsl
@@ -0,0 +1,32 @@
+struct S {
+  int x;
+  uint a;
+  uint b;
+};
+
+groupshared S wg;
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+  {
+    wg.x = 0;
+    uint atomic_result = 0u;
+    InterlockedExchange(wg.a, 0u, atomic_result);
+    uint atomic_result_1 = 0u;
+    InterlockedExchange(wg.b, 0u, atomic_result_1);
+  }
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_2 = 0u;
+  InterlockedExchange(wg.a, 1u, atomic_result_2);
+  uint atomic_result_3 = 0u;
+  InterlockedExchange(wg.b, 2u, atomic_result_3);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner(tint_symbol.local_invocation_index);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.msl
new file mode 100644
index 0000000..3ea40fa
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.msl
@@ -0,0 +1,26 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  int x;
+  atomic_uint a;
+  atomic_uint b;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup S* const tint_symbol) {
+  {
+    (*(tint_symbol)).x = 0;
+    atomic_store_explicit(&((*(tint_symbol)).a), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol)).b), 0u, memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol)).a), 1u, memory_order_relaxed);
+  atomic_store_explicit(&((*(tint_symbol)).b), 2u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup S tint_symbol_1;
+  compute_main_inner(local_invocation_index, &(tint_symbol_1));
+  return;
+}
+
diff --git a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.spvasm
new file mode 100644
index 0000000..f61269c
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.spvasm
@@ -0,0 +1,62 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 40
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S "S"
+               OpMemberName %S 0 "x"
+               OpMemberName %S 1 "a"
+               OpMemberName %S 2 "b"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 4
+               OpMemberDecorate %S 2 Offset 8
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+          %S = OpTypeStruct %int %uint %uint
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+         %wg = OpVariable %_ptr_Workgroup_S Workgroup
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %16 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+         %23 = OpConstantNull %uint
+   %uint_264 = OpConstant %uint 264
+         %35 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %8
+%local_invocation_index = OpFunctionParameter %uint
+         %12 = OpLabel
+         %15 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %15 %16
+         %22 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %22 %uint_2 %uint_0 %23
+         %26 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_2
+               OpAtomicStore %26 %uint_2 %uint_0 %23
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %31 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %31 %uint_2 %uint_0 %uint_1
+         %34 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_2
+               OpAtomicStore %34 %uint_2 %uint_0 %uint_2
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %35
+         %37 = OpLabel
+         %39 = OpLoad %uint %local_invocation_index_1
+         %38 = OpFunctionCall %void %compute_main_inner %39
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.wgsl
new file mode 100644
index 0000000..1d25d8a
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.wgsl
@@ -0,0 +1,13 @@
+struct S {
+  x : i32,
+  a : atomic<u32>,
+  b : atomic<u32>,
+}
+
+var<workgroup> wg : S;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&(wg.a), 1u);
+  atomicStore(&(wg.b), 2u);
+}
diff --git a/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.glsl b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.glsl
new file mode 100644
index 0000000..39f13a5
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.glsl
@@ -0,0 +1,24 @@
+#version 310 es
+
+struct S {
+  int x;
+  uint a;
+  uint y;
+};
+
+shared S wg;
+void compute_main(uint local_invocation_index) {
+  {
+    wg.x = 0;
+    atomicExchange(wg.a, 0u);
+    wg.y = 0u;
+  }
+  barrier();
+  atomicExchange(wg.a, 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.hlsl
new file mode 100644
index 0000000..fb93daa
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.hlsl
@@ -0,0 +1,29 @@
+struct S {
+  int x;
+  uint a;
+  uint y;
+};
+
+groupshared S wg;
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+  {
+    wg.x = 0;
+    uint atomic_result = 0u;
+    InterlockedExchange(wg.a, 0u, atomic_result);
+    wg.y = 0u;
+  }
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg.a, 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner(tint_symbol.local_invocation_index);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.msl
new file mode 100644
index 0000000..ccec811b
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.msl
@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  int x;
+  atomic_uint a;
+  uint y;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup S* const tint_symbol) {
+  {
+    (*(tint_symbol)).x = 0;
+    atomic_store_explicit(&((*(tint_symbol)).a), 0u, memory_order_relaxed);
+    (*(tint_symbol)).y = 0u;
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol)).a), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup S tint_symbol_1;
+  compute_main_inner(local_invocation_index, &(tint_symbol_1));
+  return;
+}
+
diff --git a/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.spvasm
new file mode 100644
index 0000000..acad988
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.spvasm
@@ -0,0 +1,61 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 36
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S "S"
+               OpMemberName %S 0 "x"
+               OpMemberName %S 1 "a"
+               OpMemberName %S 2 "y"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 4
+               OpMemberDecorate %S 2 Offset 8
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+          %S = OpTypeStruct %int %uint %uint
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+         %wg = OpVariable %_ptr_Workgroup_S Workgroup
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %16 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+         %23 = OpConstantNull %uint
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+         %31 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %8
+%local_invocation_index = OpFunctionParameter %uint
+         %12 = OpLabel
+         %15 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %15 %16
+         %22 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %22 %uint_2 %uint_0 %23
+         %25 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_2
+               OpStore %25 %23
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %30 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %30 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %31
+         %33 = OpLabel
+         %35 = OpLoad %uint %local_invocation_index_1
+         %34 = OpFunctionCall %void %compute_main_inner %35
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.wgsl
new file mode 100644
index 0000000..164ad73
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+struct S {
+  x : i32,
+  a : atomic<u32>,
+  y : u32,
+}
+
+var<workgroup> wg : S;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&(wg.a), 1u);
+}
diff --git a/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.glsl b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.glsl
new file mode 100644
index 0000000..10f2103
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.glsl
@@ -0,0 +1,46 @@
+#version 310 es
+
+struct S0 {
+  int x;
+  uint a;
+  int y;
+  int z;
+};
+
+struct S1 {
+  int x;
+  S0 a;
+  int y;
+  int z;
+};
+
+struct S2 {
+  int x;
+  int y;
+  int z;
+  S1 a;
+};
+
+shared S2 wg;
+void compute_main(uint local_invocation_index) {
+  {
+    wg.x = 0;
+    wg.y = 0;
+    wg.z = 0;
+    wg.a.x = 0;
+    wg.a.a.x = 0;
+    atomicExchange(wg.a.a.a, 0u);
+    wg.a.a.y = 0;
+    wg.a.a.z = 0;
+    wg.a.y = 0;
+    wg.a.z = 0;
+  }
+  barrier();
+  atomicExchange(wg.a.a.a, 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.hlsl
new file mode 100644
index 0000000..4ceb2ff
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.hlsl
@@ -0,0 +1,49 @@
+struct S0 {
+  int x;
+  uint a;
+  int y;
+  int z;
+};
+struct S1 {
+  int x;
+  S0 a;
+  int y;
+  int z;
+};
+struct S2 {
+  int x;
+  int y;
+  int z;
+  S1 a;
+};
+
+groupshared S2 wg;
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+  {
+    wg.x = 0;
+    wg.y = 0;
+    wg.z = 0;
+    wg.a.x = 0;
+    wg.a.a.x = 0;
+    uint atomic_result = 0u;
+    InterlockedExchange(wg.a.a.a, 0u, atomic_result);
+    wg.a.a.y = 0;
+    wg.a.a.z = 0;
+    wg.a.y = 0;
+    wg.a.z = 0;
+  }
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg.a.a.a, 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner(tint_symbol.local_invocation_index);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.msl
new file mode 100644
index 0000000..dbe7cc0
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.msl
@@ -0,0 +1,47 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S0 {
+  int x;
+  atomic_uint a;
+  int y;
+  int z;
+};
+
+struct S1 {
+  int x;
+  S0 a;
+  int y;
+  int z;
+};
+
+struct S2 {
+  int x;
+  int y;
+  int z;
+  S1 a;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup S2* const tint_symbol) {
+  {
+    (*(tint_symbol)).x = 0;
+    (*(tint_symbol)).y = 0;
+    (*(tint_symbol)).z = 0;
+    (*(tint_symbol)).a.x = 0;
+    (*(tint_symbol)).a.a.x = 0;
+    atomic_store_explicit(&((*(tint_symbol)).a.a.a), 0u, memory_order_relaxed);
+    (*(tint_symbol)).a.a.y = 0;
+    (*(tint_symbol)).a.a.z = 0;
+    (*(tint_symbol)).a.y = 0;
+    (*(tint_symbol)).a.z = 0;
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol)).a.a.a), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup S2 tint_symbol_1;
+  compute_main_inner(local_invocation_index, &(tint_symbol_1));
+  return;
+}
+
diff --git a/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.spvasm
new file mode 100644
index 0000000..a8aaff3
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.spvasm
@@ -0,0 +1,97 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 45
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S2 "S2"
+               OpMemberName %S2 0 "x"
+               OpMemberName %S2 1 "y"
+               OpMemberName %S2 2 "z"
+               OpMemberName %S2 3 "a"
+               OpName %S1 "S1"
+               OpMemberName %S1 0 "x"
+               OpMemberName %S1 1 "a"
+               OpName %S0 "S0"
+               OpMemberName %S0 0 "x"
+               OpMemberName %S0 1 "a"
+               OpMemberName %S0 2 "y"
+               OpMemberName %S0 3 "z"
+               OpMemberName %S1 2 "y"
+               OpMemberName %S1 3 "z"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S2 0 Offset 0
+               OpMemberDecorate %S2 1 Offset 4
+               OpMemberDecorate %S2 2 Offset 8
+               OpMemberDecorate %S2 3 Offset 12
+               OpMemberDecorate %S1 0 Offset 0
+               OpMemberDecorate %S1 1 Offset 4
+               OpMemberDecorate %S0 0 Offset 0
+               OpMemberDecorate %S0 1 Offset 4
+               OpMemberDecorate %S0 2 Offset 8
+               OpMemberDecorate %S0 3 Offset 12
+               OpMemberDecorate %S1 2 Offset 20
+               OpMemberDecorate %S1 3 Offset 24
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+         %S0 = OpTypeStruct %int %uint %int %int
+         %S1 = OpTypeStruct %int %S0 %int %int
+         %S2 = OpTypeStruct %int %int %int %S1
+%_ptr_Workgroup_S2 = OpTypePointer Workgroup %S2
+         %wg = OpVariable %_ptr_Workgroup_S2 Workgroup
+       %void = OpTypeVoid
+         %10 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %18 = OpConstantNull %int
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+         %30 = OpConstantNull %uint
+   %uint_264 = OpConstant %uint 264
+         %40 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %10
+%local_invocation_index = OpFunctionParameter %uint
+         %14 = OpLabel
+         %17 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %17 %18
+         %20 = OpAccessChain %_ptr_Workgroup_int %wg %uint_1
+               OpStore %20 %18
+         %22 = OpAccessChain %_ptr_Workgroup_int %wg %uint_2
+               OpStore %22 %18
+         %24 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_0
+               OpStore %24 %18
+         %25 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_1 %uint_0
+               OpStore %25 %18
+         %29 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_3 %uint_1 %uint_1
+               OpAtomicStore %29 %uint_2 %uint_0 %30
+         %31 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_1 %uint_2
+               OpStore %31 %18
+         %32 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_1 %uint_3
+               OpStore %32 %18
+         %33 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_2
+               OpStore %33 %18
+         %34 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_3
+               OpStore %34 %18
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %39 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_3 %uint_1 %uint_1
+               OpAtomicStore %39 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %40
+         %42 = OpLabel
+         %44 = OpLoad %uint %local_invocation_index_1
+         %43 = OpFunctionCall %void %compute_main_inner %44
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.wgsl
new file mode 100644
index 0000000..6a9bdac
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.wgsl
@@ -0,0 +1,27 @@
+struct S0 {
+  x : i32,
+  a : atomic<u32>,
+  y : i32,
+  z : i32,
+}
+
+struct S1 {
+  x : i32,
+  a : S0,
+  y : i32,
+  z : i32,
+}
+
+struct S2 {
+  x : i32,
+  y : i32,
+  z : i32,
+  a : S1,
+}
+
+var<workgroup> wg : S2;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&(wg.a.a.a), 1u);
+}
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
index 089c011..69d565e 100644
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
@@ -1,23 +1,28 @@
 #include <metal_stdlib>
 
 using namespace metal;
-struct tint_array_wrapper {
-  uint arr[10];
-};
 
-struct tint_array_wrapper_1 {
-  atomic_uint arr[10];
+template<typename T, size_t N>
+struct tint_array {
+    const constant T& operator[](size_t i) const constant { return elements[i]; }
+    device T& operator[](size_t i) device { return elements[i]; }
+    const device T& operator[](size_t i) const device { return elements[i]; }
+    thread T& operator[](size_t i) thread { return elements[i]; }
+    const thread T& operator[](size_t i) const thread { return elements[i]; }
+    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+    T elements[N];
 };
 
 struct S_atomic {
   int x;
-  tint_array_wrapper_1 a;
+  tint_array<atomic_uint, 10> a;
   uint y;
 };
 
 struct S {
   int x;
-  tint_array_wrapper a;
+  tint_array<uint, 10> a;
   uint y;
 };
 
@@ -32,14 +37,14 @@
       break;
     }
     uint const x_35 = idx;
-    atomic_store_explicit(&((*(tint_symbol)).a.arr[x_35]), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol)).a[x_35]), 0u, memory_order_relaxed);
     {
       uint const x_41 = idx;
       idx = (x_41 + 1u);
     }
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
-  atomic_store_explicit(&((*(tint_symbol)).a.arr[4]), 1u, memory_order_relaxed);
+  atomic_store_explicit(&((*(tint_symbol)).a[4]), 1u, memory_order_relaxed);
   return;
 }
 
@@ -56,7 +61,7 @@
   }
   for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
     uint const i = idx_1;
-    atomic_store_explicit(&((*(tint_symbol_3)).a.arr[i]), 0u, memory_order_relaxed);
+    atomic_store_explicit(&((*(tint_symbol_3)).a[i]), 0u, memory_order_relaxed);
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
   *(tint_symbol_4) = local_invocation_index_1_param;
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.glsl b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.glsl
new file mode 100644
index 0000000..287b69b
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.glsl
@@ -0,0 +1,29 @@
+#version 310 es
+
+struct S {
+  int x;
+  uint a[10];
+  uint y;
+};
+
+shared S wg;
+void compute_main(uint local_invocation_index) {
+  {
+    wg.x = 0;
+    wg.y = 0u;
+  }
+  {
+    for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
+      uint i = idx;
+      atomicExchange(wg.a[i], 0u);
+    }
+  }
+  barrier();
+  atomicExchange(wg.a[4], 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.hlsl
new file mode 100644
index 0000000..855c198
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.hlsl
@@ -0,0 +1,34 @@
+struct S {
+  int x;
+  uint a[10];
+  uint y;
+};
+
+groupshared S wg;
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+  {
+    wg.x = 0;
+    wg.y = 0u;
+  }
+  {
+    [loop] for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
+      const uint i = idx;
+      uint atomic_result = 0u;
+      InterlockedExchange(wg.a[i], 0u, atomic_result);
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg.a[4], 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner(tint_symbol.local_invocation_index);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl
new file mode 100644
index 0000000..00fa7d5
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl
@@ -0,0 +1,41 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+    const constant T& operator[](size_t i) const constant { return elements[i]; }
+    device T& operator[](size_t i) device { return elements[i]; }
+    const device T& operator[](size_t i) const device { return elements[i]; }
+    thread T& operator[](size_t i) thread { return elements[i]; }
+    const thread T& operator[](size_t i) const thread { return elements[i]; }
+    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+    T elements[N];
+};
+
+struct S {
+  int x;
+  tint_array<atomic_uint, 10> a;
+  uint y;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup S* const tint_symbol) {
+  {
+    (*(tint_symbol)).x = 0;
+    (*(tint_symbol)).y = 0u;
+  }
+  for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
+    uint const i = idx;
+    atomic_store_explicit(&((*(tint_symbol)).a[i]), 0u, memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol)).a[4]), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup S tint_symbol_1;
+  compute_main_inner(local_invocation_index, &(tint_symbol_1));
+  return;
+}
+
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.spvasm
new file mode 100644
index 0000000..9d06d13
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.spvasm
@@ -0,0 +1,91 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 54
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S "S"
+               OpMemberName %S 0 "x"
+               OpMemberName %S 1 "a"
+               OpMemberName %S 2 "y"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %idx "idx"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 4
+               OpDecorate %_arr_uint_uint_10 ArrayStride 4
+               OpMemberDecorate %S 2 Offset 44
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+    %uint_10 = OpConstant %uint 10
+%_arr_uint_uint_10 = OpTypeArray %uint %uint_10
+          %S = OpTypeStruct %int %_arr_uint_uint_10 %uint
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+         %wg = OpVariable %_ptr_Workgroup_S Workgroup
+       %void = OpTypeVoid
+         %10 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %18 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+         %22 = OpConstantNull %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+       %bool = OpTypeBool
+     %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+      %int_4 = OpConstant %int 4
+         %49 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %10
+%local_invocation_index = OpFunctionParameter %uint
+         %14 = OpLabel
+        %idx = OpVariable %_ptr_Function_uint Function %22
+         %17 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %17 %18
+         %21 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_2
+               OpStore %21 %22
+               OpStore %idx %local_invocation_index
+               OpBranch %25
+         %25 = OpLabel
+               OpLoopMerge %26 %27 None
+               OpBranch %28
+         %28 = OpLabel
+         %30 = OpLoad %uint %idx
+         %31 = OpULessThan %bool %30 %uint_10
+         %29 = OpLogicalNot %bool %31
+               OpSelectionMerge %33 None
+               OpBranchConditional %29 %34 %33
+         %34 = OpLabel
+               OpBranch %26
+         %33 = OpLabel
+         %35 = OpLoad %uint %idx
+         %40 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_1 %35
+               OpAtomicStore %40 %uint_2 %uint_0 %22
+               OpBranch %27
+         %27 = OpLabel
+         %41 = OpLoad %uint %idx
+         %42 = OpIAdd %uint %41 %uint_1
+               OpStore %idx %42
+               OpBranch %25
+         %26 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %48 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_1 %int_4
+               OpAtomicStore %48 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %49
+         %51 = OpLabel
+         %53 = OpLoad %uint %local_invocation_index_1
+         %52 = OpFunctionCall %void %compute_main_inner %53
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.wgsl
new file mode 100644
index 0000000..793d8c8
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+struct S {
+  x : i32,
+  a : array<atomic<u32>, 10>,
+  y : u32,
+}
+
+var<workgroup> wg : S;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  atomicStore(&(wg.a[4]), 1u);
+}
diff --git a/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.glsl b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.glsl
new file mode 100644
index 0000000..39f13a5
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.glsl
@@ -0,0 +1,24 @@
+#version 310 es
+
+struct S {
+  int x;
+  uint a;
+  uint y;
+};
+
+shared S wg;
+void compute_main(uint local_invocation_index) {
+  {
+    wg.x = 0;
+    atomicExchange(wg.a, 0u);
+    wg.y = 0u;
+  }
+  barrier();
+  atomicExchange(wg.a, 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.hlsl
new file mode 100644
index 0000000..fb93daa
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.hlsl
@@ -0,0 +1,29 @@
+struct S {
+  int x;
+  uint a;
+  uint y;
+};
+
+groupshared S wg;
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+  {
+    wg.x = 0;
+    uint atomic_result = 0u;
+    InterlockedExchange(wg.a, 0u, atomic_result);
+    wg.y = 0u;
+  }
+  GroupMemoryBarrierWithGroupSync();
+  uint atomic_result_1 = 0u;
+  InterlockedExchange(wg.a, 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+  compute_main_inner(tint_symbol.local_invocation_index);
+  return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.msl
new file mode 100644
index 0000000..ccec811b
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.msl
@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  int x;
+  atomic_uint a;
+  uint y;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup S* const tint_symbol) {
+  {
+    (*(tint_symbol)).x = 0;
+    atomic_store_explicit(&((*(tint_symbol)).a), 0u, memory_order_relaxed);
+    (*(tint_symbol)).y = 0u;
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomic_store_explicit(&((*(tint_symbol)).a), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup S tint_symbol_1;
+  compute_main_inner(local_invocation_index, &(tint_symbol_1));
+  return;
+}
+
diff --git a/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.spvasm
new file mode 100644
index 0000000..acad988
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.spvasm
@@ -0,0 +1,61 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 36
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+               OpExecutionMode %compute_main LocalSize 1 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %S "S"
+               OpMemberName %S 0 "x"
+               OpMemberName %S 1 "a"
+               OpMemberName %S 2 "y"
+               OpName %wg "wg"
+               OpName %compute_main_inner "compute_main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %compute_main "compute_main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 4
+               OpMemberDecorate %S 2 Offset 8
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+          %S = OpTypeStruct %int %uint %uint
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+         %wg = OpVariable %_ptr_Workgroup_S Workgroup
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+         %16 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+     %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+         %23 = OpConstantNull %uint
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+   %uint_264 = OpConstant %uint 264
+         %31 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %8
+%local_invocation_index = OpFunctionParameter %uint
+         %12 = OpLabel
+         %15 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+               OpStore %15 %16
+         %22 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %22 %uint_2 %uint_0 %23
+         %25 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_2
+               OpStore %25 %23
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %30 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+               OpAtomicStore %30 %uint_2 %uint_0 %uint_1
+               OpReturn
+               OpFunctionEnd
+%compute_main = OpFunction %void None %31
+         %33 = OpLabel
+         %35 = OpLoad %uint %local_invocation_index_1
+         %34 = OpFunctionCall %void %compute_main_inner %35
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.wgsl
new file mode 100644
index 0000000..5076b67
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.wgsl
@@ -0,0 +1,14 @@
+struct S {
+  x : i32,
+  a : atomic<u32>,
+  y : u32,
+}
+
+var<workgroup> wg : S;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  let p0 = &(wg);
+  let p1 = &((*(p0)).a);
+  atomicStore(p1, 1u);
+}