e2e tests: generate missing expected files

Change-Id: I6b4aee4bb08b0f4c02c015f469edc24be1623fc3
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/73340
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>

test/bug/chromium/1273451.wgsl.expected.hlsl

@@ -0,0 +1,16 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+struct A {
+  int a;
+};
+struct B {
+  int b;
+};
+
+B f(A a) {
+  const B tint_symbol = (B)0;
+  return tint_symbol;
+}

test/bug/chromium/1273451.wgsl.expected.spvasm

@@ -0,0 +1,34 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 13
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %B "B"
+               OpMemberName %B 0 "b"
+               OpName %A "A"
+               OpMemberName %A 0 "a"
+               OpName %f "f"
+               OpName %a "a"
+               OpMemberDecorate %B 0 Offset 0
+               OpMemberDecorate %A 0 Offset 0
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %B = OpTypeStruct %int
+          %A = OpTypeStruct %int
+          %5 = OpTypeFunction %B %A
+         %12 = OpConstantNull %B
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %f = OpFunction %B None %5
+          %a = OpFunctionParameter %A
+         %11 = OpLabel
+               OpReturnValue %12
+               OpFunctionEnd

test/bug/chromium/1273451.wgsl.expected.wgsl

@@ -0,0 +1,11 @@
+struct A {
+  a : i32;
+};
+
+struct B {
+  b : i32;
+};
+
+fn f(a : A) -> B {
+  return B();
+}

test/bug/fxc/dyn_array_idx/write/function.wgsl.expected.hlsl

@@ -0,0 +1,21 @@
+cbuffer cbuffer_ubo : register(b0, space0) {
+  uint4 ubo[1];
+};
+
+struct S {
+  int data[64];
+};
+
+RWByteAddressBuffer result : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void f() {
+  S s = (S)0;
+  {
+    int tint_symbol_2[64] = s.data;
+    tint_symbol_2[asint(ubo[0].x)] = 1;
+    s.data = tint_symbol_2;
+  }
+  result.Store(0u, asuint(s.data[3]));
+  return;
+}

test/bug/fxc/dyn_array_idx/write/function_via_param.wgsl.expected.hlsl

@@ -0,0 +1,25 @@
+cbuffer cbuffer_ubo : register(b0, space0) {
+  uint4 ubo[1];
+};
+
+struct S {
+  int data[64];
+};
+
+RWByteAddressBuffer result : register(u1, space0);
+
+void x(inout S p) {
+  {
+    int tint_symbol_2[64] = p.data;
+    tint_symbol_2[asint(ubo[0].x)] = 1;
+    p.data = tint_symbol_2;
+  }
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  S s = (S)0;
+  x(s);
+  result.Store(0u, asuint(s.data[3]));
+  return;
+}

test/bug/fxc/dyn_array_idx/write/private.wgsl.expected.hlsl

@@ -0,0 +1,21 @@
+cbuffer cbuffer_ubo : register(b0, space0) {
+  uint4 ubo[1];
+};
+
+struct S {
+  int data[64];
+};
+
+RWByteAddressBuffer result : register(u1, space0);
+static S s = (S)0;
+
+[numthreads(1, 1, 1)]
+void f() {
+  {
+    int tint_symbol_2[64] = s.data;
+    tint_symbol_2[asint(ubo[0].x)] = 1;
+    s.data = tint_symbol_2;
+  }
+  result.Store(0u, asuint(s.data[3]));
+  return;
+}

test/bug/fxc/dyn_array_idx/write/private_via_param.wgsl.expected.hlsl

@@ -0,0 +1,25 @@
+cbuffer cbuffer_ubo : register(b0, space0) {
+  uint4 ubo[1];
+};
+
+struct S {
+  int data[64];
+};
+
+RWByteAddressBuffer result : register(u1, space0);
+static S s = (S)0;
+
+void x(inout S p) {
+  {
+    int tint_symbol_2[64] = p.data;
+    tint_symbol_2[asint(ubo[0].x)] = 1;
+    p.data = tint_symbol_2;
+  }
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+  x(s);
+  result.Store(0u, asuint(s.data[3]));
+  return;
+}

test/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl

@@ -0,0 +1,38 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, int N, int M>
+inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
+  return lhs * vec<T, N>(rhs);
+}
+
+template<typename T, int N, int M>
+inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
+  return vec<T, M>(lhs) * rhs;
+}
+
+struct Simulation {
+  /* 0x0000 */ uint i;
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ float3 arr[8];
+};
+struct Particle {
+  /* 0x0000 */ tint_array_wrapper position;
+  /* 0x0080 */ float lifetime;
+  /* 0x0084 */ int8_t tint_pad[12];
+  /* 0x0090 */ float4 color;
+  /* 0x00a0 */ packed_float3 velocity;
+  /* 0x00ac */ int8_t tint_pad_1[4];
+};
+struct Particles {
+  /* 0x0000 */ Particle p[1];
+};
+
+kernel void tint_symbol(const device Particles* tint_symbol_1 [[buffer(1)]], const constant Simulation* tint_symbol_2 [[buffer(0)]]) {
+  Particle particle = (*(tint_symbol_1)).p[0];
+  particle.position.arr[(*(tint_symbol_2)).i] = particle.position.arr[(*(tint_symbol_2)).i];
+  return;
+}
+

test/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.spvasm

@@ -0,0 +1,78 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 37
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Particles "Particles"
+               OpMemberName %Particles 0 "p"
+               OpName %Particle "Particle"
+               OpMemberName %Particle 0 "position"
+               OpMemberName %Particle 1 "lifetime"
+               OpMemberName %Particle 2 "color"
+               OpMemberName %Particle 3 "velocity"
+               OpName %particles "particles"
+               OpName %Simulation "Simulation"
+               OpMemberName %Simulation 0 "i"
+               OpName %sim "sim"
+               OpName %main "main"
+               OpName %particle "particle"
+               OpDecorate %Particles Block
+               OpMemberDecorate %Particles 0 Offset 0
+               OpMemberDecorate %Particle 0 Offset 0
+               OpDecorate %_arr_v3float_uint_8 ArrayStride 16
+               OpMemberDecorate %Particle 1 Offset 128
+               OpMemberDecorate %Particle 2 Offset 144
+               OpMemberDecorate %Particle 3 Offset 160
+               OpDecorate %_runtimearr_Particle ArrayStride 176
+               OpDecorate %particles NonWritable
+               OpDecorate %particles DescriptorSet 1
+               OpDecorate %particles Binding 3
+               OpDecorate %Simulation Block
+               OpMemberDecorate %Simulation 0 Offset 0
+               OpDecorate %sim NonWritable
+               OpDecorate %sim DescriptorSet 1
+               OpDecorate %sim Binding 4
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+       %uint = OpTypeInt 32 0
+     %uint_8 = OpConstant %uint 8
+%_arr_v3float_uint_8 = OpTypeArray %v3float %uint_8
+    %v4float = OpTypeVector %float 4
+   %Particle = OpTypeStruct %_arr_v3float_uint_8 %float %v4float %v3float
+%_runtimearr_Particle = OpTypeRuntimeArray %Particle
+  %Particles = OpTypeStruct %_runtimearr_Particle
+%_ptr_StorageBuffer_Particles = OpTypePointer StorageBuffer %Particles
+  %particles = OpVariable %_ptr_StorageBuffer_Particles StorageBuffer
+ %Simulation = OpTypeStruct %uint
+%_ptr_Uniform_Simulation = OpTypePointer Uniform %Simulation
+        %sim = OpVariable %_ptr_Uniform_Simulation Uniform
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_StorageBuffer_Particle = OpTypePointer StorageBuffer %Particle
+%_ptr_Function_Particle = OpTypePointer Function %Particle
+         %27 = OpConstantNull %Particle
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+       %main = OpFunction %void None %15
+         %18 = OpLabel
+   %particle = OpVariable %_ptr_Function_Particle Function %27
+         %23 = OpAccessChain %_ptr_StorageBuffer_Particle %particles %uint_0 %int_0
+         %24 = OpLoad %Particle %23
+               OpStore %particle %24
+         %29 = OpAccessChain %_ptr_Uniform_uint %sim %uint_0
+         %30 = OpLoad %uint %29
+         %32 = OpAccessChain %_ptr_Function_v3float %particle %uint_0 %30
+         %33 = OpAccessChain %_ptr_Uniform_uint %sim %uint_0
+         %34 = OpLoad %uint %33
+         %35 = OpAccessChain %_ptr_Function_v3float %particle %uint_0 %34
+         %36 = OpLoad %v3float %35
+               OpStore %32 %36
+               OpReturn
+               OpFunctionEnd

test/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.wgsl

@@ -0,0 +1,24 @@
+struct Simulation {
+  i : u32;
+};
+
+struct Particle {
+  position : array<vec3<f32>, 8>;
+  lifetime : f32;
+  color : vec4<f32>;
+  velocity : vec3<f32>;
+};
+
+struct Particles {
+  p : array<Particle>;
+};
+
+[[group(1), binding(3)]] var<storage, read> particles : Particles;
+
+[[group(1), binding(4)]] var<uniform> sim : Simulation;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var particle = particles.p[0];
+  particle.position[sim.i] = particle.position[sim.i];
+}

test/bug/tint/998.wgsl.expected.hlsl

@@ -0,0 +1,21 @@
+cbuffer cbuffer_constants : register(b0, space1) {
+  uint4 constants[1];
+};
+
+RWByteAddressBuffer result : register(u1, space1);
+
+struct S {
+  uint data[3];
+};
+
+static S s = (S)0;
+
+[numthreads(1, 1, 1)]
+void main() {
+  {
+    uint tint_symbol_1[3] = s.data;
+    tint_symbol_1[constants[0].x] = 0u;
+    s.data = tint_symbol_1;
+  }
+  return;
+}

test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.msl

@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+};
+struct InnerS {
+  int v;
+};
+struct tint_array_wrapper {
+  InnerS arr[8];
+};
+struct OuterS {
+  tint_array_wrapper a1;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  InnerS v = {};
+  OuterS s1 = {};
+  for(int i = 0; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+    s1.a1.arr[(*(tint_symbol_1)).i] = v;
+  }
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.spvasm

@@ -0,0 +1,84 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 44
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpName %uniforms "uniforms"
+               OpName %main "main"
+               OpName %InnerS "InnerS"
+               OpMemberName %InnerS 0 "v"
+               OpName %v "v"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "a1"
+               OpName %s1 "s1"
+               OpName %i "i"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpMemberDecorate %InnerS 0 Offset 0
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+     %InnerS = OpTypeStruct %int
+%_ptr_Function_InnerS = OpTypePointer Function %InnerS
+         %13 = OpConstantNull %InnerS
+     %uint_8 = OpConstant %uint 8
+%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
+     %OuterS = OpTypeStruct %_arr_InnerS_uint_8
+%_ptr_Function_OuterS = OpTypePointer Function %OuterS
+         %19 = OpConstantNull %OuterS
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+         %23 = OpConstantNull %int
+      %int_4 = OpConstant %int 4
+       %bool = OpTypeBool
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+      %int_1 = OpConstant %int 1
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+          %v = OpVariable %_ptr_Function_InnerS Function %13
+         %s1 = OpVariable %_ptr_Function_OuterS Function %19
+          %i = OpVariable %_ptr_Function_int Function %23
+               OpStore %i %int_0
+               OpBranch %24
+         %24 = OpLabel
+               OpLoopMerge %25 %26 None
+               OpBranch %27
+         %27 = OpLabel
+         %29 = OpLoad %int %i
+         %31 = OpSLessThan %bool %29 %int_4
+         %28 = OpLogicalNot %bool %31
+               OpSelectionMerge %33 None
+               OpBranchConditional %28 %34 %33
+         %34 = OpLabel
+               OpBranch %25
+         %33 = OpLabel
+         %37 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %38 = OpLoad %uint %37
+         %39 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %38
+         %40 = OpLoad %InnerS %v
+               OpStore %39 %40
+               OpBranch %26
+         %26 = OpLabel
+         %41 = OpLoad %int %i
+         %43 = OpIAdd %int %41 %int_1
+               OpStore %i %43
+               OpBranch %24
+         %25 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.wgsl

@@ -0,0 +1,22 @@
+struct Uniforms {
+  i : u32;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+struct OuterS {
+  a1 : array<InnerS, 8>;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : InnerS;
+  var s1 : OuterS;
+  for(var i : i32 = 0; (i < 4); i = (i + 1)) {
+    s1.a1[uniforms.i] = v;
+  }
+}

test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.msl

@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+};
+struct InnerS {
+  int v;
+};
+struct tint_array_wrapper {
+  InnerS arr[8];
+};
+struct OuterS {
+  tint_array_wrapper a1;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  InnerS v = {};
+  OuterS s1 = {};
+  for(int i = 0; (i < 4); s1.a1.arr[(*(tint_symbol_1)).i] = v) {
+    i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
+  }
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.spvasm

@@ -0,0 +1,84 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 44
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpName %uniforms "uniforms"
+               OpName %main "main"
+               OpName %InnerS "InnerS"
+               OpMemberName %InnerS 0 "v"
+               OpName %v "v"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "a1"
+               OpName %s1 "s1"
+               OpName %i "i"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpMemberDecorate %InnerS 0 Offset 0
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+     %InnerS = OpTypeStruct %int
+%_ptr_Function_InnerS = OpTypePointer Function %InnerS
+         %13 = OpConstantNull %InnerS
+     %uint_8 = OpConstant %uint 8
+%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
+     %OuterS = OpTypeStruct %_arr_InnerS_uint_8
+%_ptr_Function_OuterS = OpTypePointer Function %OuterS
+         %19 = OpConstantNull %OuterS
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+         %23 = OpConstantNull %int
+      %int_4 = OpConstant %int 4
+       %bool = OpTypeBool
+      %int_1 = OpConstant %int 1
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+          %v = OpVariable %_ptr_Function_InnerS Function %13
+         %s1 = OpVariable %_ptr_Function_OuterS Function %19
+          %i = OpVariable %_ptr_Function_int Function %23
+               OpStore %i %int_0
+               OpBranch %24
+         %24 = OpLabel
+               OpLoopMerge %25 %26 None
+               OpBranch %27
+         %27 = OpLabel
+         %29 = OpLoad %int %i
+         %31 = OpSLessThan %bool %29 %int_4
+         %28 = OpLogicalNot %bool %31
+               OpSelectionMerge %33 None
+               OpBranchConditional %28 %34 %33
+         %34 = OpLabel
+               OpBranch %25
+         %33 = OpLabel
+         %35 = OpLoad %int %i
+         %37 = OpIAdd %int %35 %int_1
+               OpStore %i %37
+               OpBranch %26
+         %26 = OpLabel
+         %40 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %41 = OpLoad %uint %40
+         %42 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %41
+         %43 = OpLoad %InnerS %v
+               OpStore %42 %43
+               OpBranch %24
+         %25 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.wgsl

@@ -0,0 +1,22 @@
+struct Uniforms {
+  i : u32;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+struct OuterS {
+  a1 : array<InnerS, 8>;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : InnerS;
+  var s1 : OuterS;
+  for(var i : i32 = 0; (i < 4); s1.a1[uniforms.i] = v) {
+    i = (i + 1);
+  }
+}

test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.msl

@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+};
+struct InnerS {
+  int v;
+};
+struct tint_array_wrapper {
+  InnerS arr[8];
+};
+struct OuterS {
+  tint_array_wrapper a1;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  InnerS v = {};
+  OuterS s1 = {};
+  int i = 0;
+  for(s1.a1.arr[(*(tint_symbol_1)).i] = v; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  }
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.spvasm

@@ -0,0 +1,84 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 44
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpName %uniforms "uniforms"
+               OpName %main "main"
+               OpName %InnerS "InnerS"
+               OpMemberName %InnerS 0 "v"
+               OpName %v "v"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "a1"
+               OpName %s1 "s1"
+               OpName %i "i"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpMemberDecorate %InnerS 0 Offset 0
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+     %InnerS = OpTypeStruct %int
+%_ptr_Function_InnerS = OpTypePointer Function %InnerS
+         %13 = OpConstantNull %InnerS
+     %uint_8 = OpConstant %uint 8
+%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
+     %OuterS = OpTypeStruct %_arr_InnerS_uint_8
+%_ptr_Function_OuterS = OpTypePointer Function %OuterS
+         %19 = OpConstantNull %OuterS
+      %int_0 = OpConstant %int 0
+%_ptr_Function_int = OpTypePointer Function %int
+         %23 = OpConstantNull %int
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+      %int_4 = OpConstant %int 4
+       %bool = OpTypeBool
+      %int_1 = OpConstant %int 1
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+          %v = OpVariable %_ptr_Function_InnerS Function %13
+         %s1 = OpVariable %_ptr_Function_OuterS Function %19
+          %i = OpVariable %_ptr_Function_int Function %23
+               OpStore %i %int_0
+         %26 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %27 = OpLoad %uint %26
+         %28 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %27
+         %29 = OpLoad %InnerS %v
+               OpStore %28 %29
+               OpBranch %30
+         %30 = OpLabel
+               OpLoopMerge %31 %32 None
+               OpBranch %33
+         %33 = OpLabel
+         %35 = OpLoad %int %i
+         %37 = OpSLessThan %bool %35 %int_4
+         %34 = OpLogicalNot %bool %37
+               OpSelectionMerge %39 None
+               OpBranchConditional %34 %40 %39
+         %40 = OpLabel
+               OpBranch %31
+         %39 = OpLabel
+               OpBranch %32
+         %32 = OpLabel
+         %41 = OpLoad %int %i
+         %43 = OpIAdd %int %41 %int_1
+               OpStore %i %43
+               OpBranch %30
+         %31 = OpLabel
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.wgsl

@@ -0,0 +1,22 @@
+struct Uniforms {
+  i : u32;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+struct OuterS {
+  a1 : array<InnerS, 8>;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : InnerS;
+  var s1 : OuterS;
+  var i : i32 = 0;
+  for(s1.a1[uniforms.i] = v; (i < 4); i = (i + 1)) {
+  }
+}

test/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.msl

@@ -0,0 +1,36 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+  /* 0x0004 */ uint j;
+};
+struct InnerS {
+  int v;
+};
+struct tint_array_wrapper {
+  InnerS arr[8];
+};
+struct S1 {
+  tint_array_wrapper a2;
+};
+struct tint_array_wrapper_1 {
+  S1 arr[8];
+};
+struct OuterS {
+  tint_array_wrapper_1 a1;
+};
+
+uint getNextIndex(thread uint* const tint_symbol_1) {
+  *(tint_symbol_1) = (*(tint_symbol_1) + 1u);
+  return *(tint_symbol_1);
+}
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_3 [[buffer(0)]]) {
+  thread uint tint_symbol_2 = 0u;
+  InnerS v = {};
+  OuterS s = {};
+  s.a1.arr[getNextIndex(&(tint_symbol_2))].a2.arr[(*(tint_symbol_3)).j] = v;
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.spvasm

@@ -0,0 +1,79 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 39
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %nextIndex "nextIndex"
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpMemberName %Uniforms 1 "j"
+               OpName %uniforms "uniforms"
+               OpName %getNextIndex "getNextIndex"
+               OpName %main "main"
+               OpName %InnerS "InnerS"
+               OpMemberName %InnerS 0 "v"
+               OpName %v "v"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "a1"
+               OpName %S1 "S1"
+               OpMemberName %S1 0 "a2"
+               OpName %s "s"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpMemberDecorate %Uniforms 1 Offset 4
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpMemberDecorate %InnerS 0 Offset 0
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpMemberDecorate %S1 0 Offset 0
+               OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
+               OpDecorate %_arr_S1_uint_8 ArrayStride 32
+       %uint = OpTypeInt 32 0
+%_ptr_Private_uint = OpTypePointer Private %uint
+          %4 = OpConstantNull %uint
+  %nextIndex = OpVariable %_ptr_Private_uint Private %4
+   %Uniforms = OpTypeStruct %uint %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+          %8 = OpTypeFunction %uint
+     %uint_1 = OpConstant %uint 1
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+     %InnerS = OpTypeStruct %int
+%_ptr_Function_InnerS = OpTypePointer Function %InnerS
+         %23 = OpConstantNull %InnerS
+     %uint_8 = OpConstant %uint 8
+%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
+         %S1 = OpTypeStruct %_arr_InnerS_uint_8
+%_arr_S1_uint_8 = OpTypeArray %S1 %uint_8
+     %OuterS = OpTypeStruct %_arr_S1_uint_8
+%_ptr_Function_OuterS = OpTypePointer Function %OuterS
+         %31 = OpConstantNull %OuterS
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%getNextIndex = OpFunction %uint None %8
+         %10 = OpLabel
+         %11 = OpLoad %uint %nextIndex
+         %13 = OpIAdd %uint %11 %uint_1
+               OpStore %nextIndex %13
+         %14 = OpLoad %uint %nextIndex
+               OpReturnValue %14
+               OpFunctionEnd
+       %main = OpFunction %void None %15
+         %18 = OpLabel
+          %v = OpVariable %_ptr_Function_InnerS Function %23
+          %s = OpVariable %_ptr_Function_OuterS Function %31
+         %33 = OpFunctionCall %uint %getNextIndex
+         %35 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+         %36 = OpLoad %uint %35
+         %37 = OpAccessChain %_ptr_Function_InnerS %s %uint_0 %33 %uint_0 %36
+         %38 = OpLoad %InnerS %v
+               OpStore %37 %38
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/indexing_with_side_effect_func.wgsl.expected.wgsl

@@ -0,0 +1,32 @@
+struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+struct S1 {
+  a2 : array<InnerS, 8>;
+};
+
+struct OuterS {
+  a1 : array<S1, 8>;
+};
+
+var<private> nextIndex : u32;
+
+fn getNextIndex() -> u32 {
+  nextIndex = (nextIndex + 1u);
+  return nextIndex;
+}
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : InnerS;
+  var s : OuterS;
+  s.a1[getNextIndex()].a2[uniforms.j] = v;
+}

test/statements/assign/indexed_assign_to_array_in_struct/struct_array.wgsl.expected.msl

@@ -0,0 +1,23 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+};
+struct InnerS {
+  int v;
+};
+struct tint_array_wrapper {
+  InnerS arr[8];
+};
+struct OuterS {
+  tint_array_wrapper a1;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  InnerS v = {};
+  OuterS s1 = {};
+  s1.a1.arr[(*(tint_symbol_1)).i] = v;
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/struct_array.wgsl.expected.spvasm

@@ -0,0 +1,55 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 26
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpName %uniforms "uniforms"
+               OpName %main "main"
+               OpName %InnerS "InnerS"
+               OpMemberName %InnerS 0 "v"
+               OpName %v "v"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "a1"
+               OpName %s1 "s1"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpMemberDecorate %InnerS 0 Offset 0
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+     %InnerS = OpTypeStruct %int
+%_ptr_Function_InnerS = OpTypePointer Function %InnerS
+         %13 = OpConstantNull %InnerS
+     %uint_8 = OpConstant %uint 8
+%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
+     %OuterS = OpTypeStruct %_arr_InnerS_uint_8
+%_ptr_Function_OuterS = OpTypePointer Function %OuterS
+         %19 = OpConstantNull %OuterS
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+          %v = OpVariable %_ptr_Function_InnerS Function %13
+         %s1 = OpVariable %_ptr_Function_OuterS Function %19
+         %22 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %23 = OpLoad %uint %22
+         %24 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %23
+         %25 = OpLoad %InnerS %v
+               OpStore %24 %25
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/struct_array.wgsl.expected.wgsl

@@ -0,0 +1,20 @@
+struct Uniforms {
+  i : u32;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+struct OuterS {
+  a1 : array<InnerS, 8>;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : InnerS;
+  var s1 : OuterS;
+  s1.a1[uniforms.i] = v;
+}

test/statements/assign/indexed_assign_to_array_in_struct/struct_array_array.wgsl.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+  /* 0x0004 */ uint j;
+};
+struct InnerS {
+  int v;
+};
+struct tint_array_wrapper_1 {
+  InnerS arr[8];
+};
+struct tint_array_wrapper {
+  tint_array_wrapper_1 arr[8];
+};
+struct OuterS {
+  tint_array_wrapper a1;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  InnerS v = {};
+  OuterS s1 = {};
+  s1.a1.arr[(*(tint_symbol_1)).i].arr[(*(tint_symbol_1)).j] = v;
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/struct_array_array.wgsl.expected.spvasm

@@ -0,0 +1,62 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 30
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpMemberName %Uniforms 1 "j"
+               OpName %uniforms "uniforms"
+               OpName %main "main"
+               OpName %InnerS "InnerS"
+               OpMemberName %InnerS 0 "v"
+               OpName %v "v"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "a1"
+               OpName %s1 "s1"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpMemberDecorate %Uniforms 1 Offset 4
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpMemberDecorate %InnerS 0 Offset 0
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
+               OpDecorate %_arr__arr_InnerS_uint_8_uint_8 ArrayStride 32
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+     %InnerS = OpTypeStruct %int
+%_ptr_Function_InnerS = OpTypePointer Function %InnerS
+         %13 = OpConstantNull %InnerS
+     %uint_8 = OpConstant %uint 8
+%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
+%_arr__arr_InnerS_uint_8_uint_8 = OpTypeArray %_arr_InnerS_uint_8 %uint_8
+     %OuterS = OpTypeStruct %_arr__arr_InnerS_uint_8_uint_8
+%_ptr_Function_OuterS = OpTypePointer Function %OuterS
+         %20 = OpConstantNull %OuterS
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+     %uint_1 = OpConstant %uint 1
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+          %v = OpVariable %_ptr_Function_InnerS Function %13
+         %s1 = OpVariable %_ptr_Function_OuterS Function %20
+         %23 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %24 = OpLoad %uint %23
+         %26 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+         %27 = OpLoad %uint %26
+         %28 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %24 %27
+         %29 = OpLoad %InnerS %v
+               OpStore %28 %29
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/struct_array_array.wgsl.expected.wgsl

@@ -0,0 +1,21 @@
+struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+struct OuterS {
+  a1 : array<array<InnerS, 8>, 8>;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : InnerS;
+  var s1 : OuterS;
+  s1.a1[uniforms.i][uniforms.j] = v;
+}

test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct.wgsl.expected.msl

@@ -0,0 +1,26 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+};
+struct InnerS {
+  int v;
+};
+struct S1 {
+  InnerS s2;
+};
+struct tint_array_wrapper {
+  S1 arr[8];
+};
+struct OuterS {
+  tint_array_wrapper a1;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  InnerS v = {};
+  OuterS s1 = {};
+  s1.a1.arr[(*(tint_symbol_1)).i].s2 = v;
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct.wgsl.expected.spvasm

@@ -0,0 +1,59 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 27
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpName %uniforms "uniforms"
+               OpName %main "main"
+               OpName %InnerS "InnerS"
+               OpMemberName %InnerS 0 "v"
+               OpName %v "v"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "a1"
+               OpName %S1 "S1"
+               OpMemberName %S1 0 "s2"
+               OpName %s1 "s1"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpMemberDecorate %InnerS 0 Offset 0
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpMemberDecorate %S1 0 Offset 0
+               OpDecorate %_arr_S1_uint_8 ArrayStride 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+     %InnerS = OpTypeStruct %int
+%_ptr_Function_InnerS = OpTypePointer Function %InnerS
+         %13 = OpConstantNull %InnerS
+         %S1 = OpTypeStruct %InnerS
+     %uint_8 = OpConstant %uint 8
+%_arr_S1_uint_8 = OpTypeArray %S1 %uint_8
+     %OuterS = OpTypeStruct %_arr_S1_uint_8
+%_ptr_Function_OuterS = OpTypePointer Function %OuterS
+         %20 = OpConstantNull %OuterS
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+          %v = OpVariable %_ptr_Function_InnerS Function %13
+         %s1 = OpVariable %_ptr_Function_OuterS Function %20
+         %23 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %24 = OpLoad %uint %23
+         %25 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %24 %uint_0
+         %26 = OpLoad %InnerS %v
+               OpStore %25 %26
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct.wgsl.expected.wgsl

@@ -0,0 +1,24 @@
+struct Uniforms {
+  i : u32;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+struct S1 {
+  s2 : InnerS;
+};
+
+struct OuterS {
+  a1 : array<S1, 8>;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : InnerS;
+  var s1 : OuterS;
+  s1.a1[uniforms.i].s2 = v;
+}

test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct_array.wgsl.expected.msl

@@ -0,0 +1,30 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+  /* 0x0004 */ uint j;
+};
+struct InnerS {
+  int v;
+};
+struct tint_array_wrapper {
+  InnerS arr[8];
+};
+struct S1 {
+  tint_array_wrapper a2;
+};
+struct tint_array_wrapper_1 {
+  S1 arr[8];
+};
+struct OuterS {
+  tint_array_wrapper_1 a1;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  InnerS v = {};
+  OuterS s = {};
+  s.a1.arr[(*(tint_symbol_1)).i].a2.arr[(*(tint_symbol_1)).j] = v;
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct_array.wgsl.expected.spvasm

@@ -0,0 +1,66 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpMemberName %Uniforms 1 "j"
+               OpName %uniforms "uniforms"
+               OpName %main "main"
+               OpName %InnerS "InnerS"
+               OpMemberName %InnerS 0 "v"
+               OpName %v "v"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "a1"
+               OpName %S1 "S1"
+               OpMemberName %S1 0 "a2"
+               OpName %s "s"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpMemberDecorate %Uniforms 1 Offset 4
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpMemberDecorate %InnerS 0 Offset 0
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpMemberDecorate %S1 0 Offset 0
+               OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
+               OpDecorate %_arr_S1_uint_8 ArrayStride 32
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+     %InnerS = OpTypeStruct %int
+%_ptr_Function_InnerS = OpTypePointer Function %InnerS
+         %13 = OpConstantNull %InnerS
+     %uint_8 = OpConstant %uint 8
+%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
+         %S1 = OpTypeStruct %_arr_InnerS_uint_8
+%_arr_S1_uint_8 = OpTypeArray %S1 %uint_8
+     %OuterS = OpTypeStruct %_arr_S1_uint_8
+%_ptr_Function_OuterS = OpTypePointer Function %OuterS
+         %21 = OpConstantNull %OuterS
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+     %uint_1 = OpConstant %uint 1
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+          %v = OpVariable %_ptr_Function_InnerS Function %13
+          %s = OpVariable %_ptr_Function_OuterS Function %21
+         %24 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %25 = OpLoad %uint %24
+         %27 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+         %28 = OpLoad %uint %27
+         %29 = OpAccessChain %_ptr_Function_InnerS %s %uint_0 %25 %uint_0 %28
+         %30 = OpLoad %InnerS %v
+               OpStore %29 %30
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/struct_array_struct_array.wgsl.expected.wgsl

@@ -0,0 +1,25 @@
+struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+struct S1 {
+  a2 : array<InnerS, 8>;
+};
+
+struct OuterS {
+  a1 : array<S1, 8>;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : InnerS;
+  var s : OuterS;
+  s.a1[uniforms.i].a2[uniforms.j] = v;
+}

test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.msl

@@ -0,0 +1,19 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+};
+struct InnerS {
+  /* 0x0000 */ int v;
+};
+struct OuterS {
+  /* 0x0000 */ InnerS a1[1];
+};
+
+kernel void tint_symbol(device OuterS* tint_symbol_1 [[buffer(1)]], const constant Uniforms* tint_symbol_2 [[buffer(0)]]) {
+  InnerS v = {};
+  (*(tint_symbol_1)).a1[(*(tint_symbol_2)).i] = v;
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.spvasm

@@ -0,0 +1,57 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 25
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpName %uniforms "uniforms"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "a1"
+               OpName %InnerS "InnerS"
+               OpMemberName %InnerS 0 "v"
+               OpName %s1 "s1"
+               OpName %main "main"
+               OpName %v "v"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpDecorate %OuterS Block
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpMemberDecorate %InnerS 0 Offset 0
+               OpDecorate %_runtimearr_InnerS ArrayStride 4
+               OpDecorate %s1 Binding 0
+               OpDecorate %s1 DescriptorSet 0
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+        %int = OpTypeInt 32 1
+     %InnerS = OpTypeStruct %int
+%_runtimearr_InnerS = OpTypeRuntimeArray %InnerS
+     %OuterS = OpTypeStruct %_runtimearr_InnerS
+%_ptr_StorageBuffer_OuterS = OpTypePointer StorageBuffer %OuterS
+         %s1 = OpVariable %_ptr_StorageBuffer_OuterS StorageBuffer
+       %void = OpTypeVoid
+         %11 = OpTypeFunction %void
+%_ptr_Function_InnerS = OpTypePointer Function %InnerS
+         %17 = OpConstantNull %InnerS
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_StorageBuffer_InnerS = OpTypePointer StorageBuffer %InnerS
+       %main = OpFunction %void None %11
+         %14 = OpLabel
+          %v = OpVariable %_ptr_Function_InnerS Function %17
+         %20 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %21 = OpLoad %uint %20
+         %23 = OpAccessChain %_ptr_StorageBuffer_InnerS %s1 %uint_0 %21
+         %24 = OpLoad %InnerS %v
+               OpStore %23 %24
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array.wgsl.expected.wgsl

@@ -0,0 +1,21 @@
+struct Uniforms {
+  i : u32;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+struct OuterS {
+  a1 : array<InnerS>;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[binding(0), group(0)]] var<storage, read_write> s1 : OuterS;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : InnerS;
+  s1.a1[uniforms.i] = v;
+}

test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.msl

@@ -0,0 +1,26 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+  /* 0x0004 */ uint j;
+};
+struct InnerS {
+  /* 0x0000 */ int v;
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ InnerS arr[8];
+};
+struct S1 {
+  /* 0x0000 */ tint_array_wrapper a2;
+};
+struct OuterS {
+  /* 0x0000 */ S1 a1[1];
+};
+
+kernel void tint_symbol(device OuterS* tint_symbol_1 [[buffer(1)]], const constant Uniforms* tint_symbol_2 [[buffer(0)]]) {
+  InnerS v = {};
+  (*(tint_symbol_1)).a1[(*(tint_symbol_2)).i].a2.arr[(*(tint_symbol_2)).j] = v;
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.spvasm

@@ -0,0 +1,69 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpMemberName %Uniforms 1 "j"
+               OpName %uniforms "uniforms"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "a1"
+               OpName %S1 "S1"
+               OpMemberName %S1 0 "a2"
+               OpName %InnerS "InnerS"
+               OpMemberName %InnerS 0 "v"
+               OpName %s "s"
+               OpName %main "main"
+               OpName %v "v"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpMemberDecorate %Uniforms 1 Offset 4
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpDecorate %OuterS Block
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpMemberDecorate %S1 0 Offset 0
+               OpMemberDecorate %InnerS 0 Offset 0
+               OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
+               OpDecorate %_runtimearr_S1 ArrayStride 32
+               OpDecorate %s Binding 0
+               OpDecorate %s DescriptorSet 0
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+        %int = OpTypeInt 32 1
+     %InnerS = OpTypeStruct %int
+     %uint_8 = OpConstant %uint 8
+%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
+         %S1 = OpTypeStruct %_arr_InnerS_uint_8
+%_runtimearr_S1 = OpTypeRuntimeArray %S1
+     %OuterS = OpTypeStruct %_runtimearr_S1
+%_ptr_StorageBuffer_OuterS = OpTypePointer StorageBuffer %OuterS
+          %s = OpVariable %_ptr_StorageBuffer_OuterS StorageBuffer
+       %void = OpTypeVoid
+         %14 = OpTypeFunction %void
+%_ptr_Function_InnerS = OpTypePointer Function %InnerS
+         %20 = OpConstantNull %InnerS
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+     %uint_1 = OpConstant %uint 1
+%_ptr_StorageBuffer_InnerS = OpTypePointer StorageBuffer %InnerS
+       %main = OpFunction %void None %14
+         %17 = OpLabel
+          %v = OpVariable %_ptr_Function_InnerS Function %20
+         %23 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %24 = OpLoad %uint %23
+         %26 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_1
+         %27 = OpLoad %uint %26
+         %29 = OpAccessChain %_ptr_StorageBuffer_InnerS %s %uint_0 %24 %uint_0 %27
+         %30 = OpLoad %InnerS %v
+               OpStore %29 %30
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/struct_dynamic_array_struct_array.wgsl.expected.wgsl

@@ -0,0 +1,26 @@
+struct Uniforms {
+  i : u32;
+  j : u32;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+struct S1 {
+  a2 : array<InnerS, 8>;
+};
+
+struct OuterS {
+  a1 : array<S1>;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[binding(0), group(0)]] var<storage, read_write> s : OuterS;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : InnerS;
+  s.a1[uniforms.i].a2[uniforms.j] = v;
+}

test/statements/assign/indexed_assign_to_array_in_struct/struct_matrix.wgsl.expected.msl

@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+};
+struct OuterS {
+  float2x4 m1;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  OuterS s1 = {};
+  s1.m1[(*(tint_symbol_1)).i] = float4(1.0f);
+  s1.m1[(*(tint_symbol_1)).i][(*(tint_symbol_1)).i] = 1.0f;
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/struct_matrix.wgsl.expected.spvasm

@@ -0,0 +1,57 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 30
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpName %uniforms "uniforms"
+               OpName %main "main"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "m1"
+               OpName %s1 "s1"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpMemberDecorate %OuterS 0 ColMajor
+               OpMemberDecorate %OuterS 0 MatrixStride 16
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%mat2v4float = OpTypeMatrix %v4float 2
+     %OuterS = OpTypeStruct %mat2v4float
+%_ptr_Function_OuterS = OpTypePointer Function %OuterS
+         %15 = OpConstantNull %OuterS
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_Function_v4float = OpTypePointer Function %v4float
+    %float_1 = OpConstant %float 1
+         %23 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
+%_ptr_Function_float = OpTypePointer Function %float
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+         %s1 = OpVariable %_ptr_Function_OuterS Function %15
+         %18 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %19 = OpLoad %uint %18
+         %21 = OpAccessChain %_ptr_Function_v4float %s1 %uint_0 %19
+               OpStore %21 %23
+         %24 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %25 = OpLoad %uint %24
+         %26 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %27 = OpLoad %uint %26
+         %29 = OpAccessChain %_ptr_Function_float %s1 %uint_0 %25 %27
+               OpStore %29 %float_1
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/struct_matrix.wgsl.expected.wgsl

@@ -0,0 +1,16 @@
+struct Uniforms {
+  i : u32;
+};
+
+struct OuterS {
+  m1 : mat2x4<f32>;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var s1 : OuterS;
+  s1.m1[uniforms.i] = vec4<f32>(1.0);
+  s1.m1[uniforms.i][uniforms.i] = 1.0;
+}

test/statements/assign/indexed_assign_to_array_in_struct/struct_multiple_arrays.wgsl.expected.msl

@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+};
+struct InnerS {
+  int v;
+};
+struct tint_array_wrapper {
+  InnerS arr[8];
+};
+struct OuterS {
+  tint_array_wrapper a1;
+  tint_array_wrapper a2;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  InnerS v = {};
+  OuterS s1 = {};
+  s1.a1.arr[(*(tint_symbol_1)).i] = v;
+  s1.a2.arr[(*(tint_symbol_1)).i] = v;
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/struct_multiple_arrays.wgsl.expected.spvasm

@@ -0,0 +1,63 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 31
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpName %uniforms "uniforms"
+               OpName %main "main"
+               OpName %InnerS "InnerS"
+               OpMemberName %InnerS 0 "v"
+               OpName %v "v"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "a1"
+               OpMemberName %OuterS 1 "a2"
+               OpName %s1 "s1"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpMemberDecorate %InnerS 0 Offset 0
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
+               OpMemberDecorate %OuterS 1 Offset 32
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+     %InnerS = OpTypeStruct %int
+%_ptr_Function_InnerS = OpTypePointer Function %InnerS
+         %13 = OpConstantNull %InnerS
+     %uint_8 = OpConstant %uint 8
+%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
+     %OuterS = OpTypeStruct %_arr_InnerS_uint_8 %_arr_InnerS_uint_8
+%_ptr_Function_OuterS = OpTypePointer Function %OuterS
+         %19 = OpConstantNull %OuterS
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+     %uint_1 = OpConstant %uint 1
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+          %v = OpVariable %_ptr_Function_InnerS Function %13
+         %s1 = OpVariable %_ptr_Function_OuterS Function %19
+         %22 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %23 = OpLoad %uint %22
+         %24 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %23
+         %25 = OpLoad %InnerS %v
+               OpStore %24 %25
+         %27 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %28 = OpLoad %uint %27
+         %29 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_1 %28
+         %30 = OpLoad %InnerS %v
+               OpStore %29 %30
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/struct_multiple_arrays.wgsl.expected.wgsl

@@ -0,0 +1,22 @@
+struct Uniforms {
+  i : u32;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+struct OuterS {
+  a1 : array<InnerS, 8>;
+  a2 : array<InnerS, 8>;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : InnerS;
+  var s1 : OuterS;
+  s1.a1[uniforms.i] = v;
+  s1.a2[uniforms.i] = v;
+}

test/statements/assign/indexed_assign_to_array_in_struct/struct_struct_array.wgsl.expected.msl

@@ -0,0 +1,26 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+};
+struct InnerS {
+  int v;
+};
+struct tint_array_wrapper {
+  InnerS arr[8];
+};
+struct S1 {
+  tint_array_wrapper a;
+};
+struct OuterS {
+  S1 s2;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  InnerS v = {};
+  OuterS s1 = {};
+  s1.s2.a.arr[(*(tint_symbol_1)).i] = v;
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/struct_struct_array.wgsl.expected.spvasm

@@ -0,0 +1,59 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 27
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpName %uniforms "uniforms"
+               OpName %main "main"
+               OpName %InnerS "InnerS"
+               OpMemberName %InnerS 0 "v"
+               OpName %v "v"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "s2"
+               OpName %S1 "S1"
+               OpMemberName %S1 0 "a"
+               OpName %s1 "s1"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpMemberDecorate %InnerS 0 Offset 0
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpMemberDecorate %S1 0 Offset 0
+               OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+     %InnerS = OpTypeStruct %int
+%_ptr_Function_InnerS = OpTypePointer Function %InnerS
+         %13 = OpConstantNull %InnerS
+     %uint_8 = OpConstant %uint 8
+%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
+         %S1 = OpTypeStruct %_arr_InnerS_uint_8
+     %OuterS = OpTypeStruct %S1
+%_ptr_Function_OuterS = OpTypePointer Function %OuterS
+         %20 = OpConstantNull %OuterS
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+          %v = OpVariable %_ptr_Function_InnerS Function %13
+         %s1 = OpVariable %_ptr_Function_OuterS Function %20
+         %23 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %24 = OpLoad %uint %23
+         %25 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %uint_0 %24
+         %26 = OpLoad %InnerS %v
+               OpStore %25 %26
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/struct_struct_array.wgsl.expected.wgsl

@@ -0,0 +1,24 @@
+struct Uniforms {
+  i : u32;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+struct S1 {
+  a : array<InnerS, 8>;
+};
+
+struct OuterS {
+  s2 : S1;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : InnerS;
+  var s1 : OuterS;
+  s1.s2.a[uniforms.i] = v;
+}

test/statements/assign/indexed_assign_to_array_in_struct/struct_vector.wgsl.expected.msl

@@ -0,0 +1,16 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+};
+struct OuterS {
+  float3 v1;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  OuterS s1 = {};
+  s1.v1[(*(tint_symbol_1)).i] = 1.0f;
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/struct_vector.wgsl.expected.spvasm

@@ -0,0 +1,46 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 22
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpName %uniforms "uniforms"
+               OpName %main "main"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "v1"
+               OpName %s1 "s1"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpMemberDecorate %OuterS 0 Offset 0
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+     %OuterS = OpTypeStruct %v3float
+%_ptr_Function_OuterS = OpTypePointer Function %OuterS
+         %14 = OpConstantNull %OuterS
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_Function_float = OpTypePointer Function %float
+    %float_1 = OpConstant %float 1
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+         %s1 = OpVariable %_ptr_Function_OuterS Function %14
+         %17 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %18 = OpLoad %uint %17
+         %20 = OpAccessChain %_ptr_Function_float %s1 %uint_0 %18
+               OpStore %20 %float_1
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/struct_vector.wgsl.expected.wgsl

@@ -0,0 +1,15 @@
+struct Uniforms {
+  i : u32;
+};
+
+struct OuterS {
+  v1 : vec3<f32>;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var s1 : OuterS;
+  s1.v1[uniforms.i] = 1.0;
+}

test/statements/assign/indexed_assign_to_array_in_struct/vector_assign.wgsl.expected.msl

@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+};
+struct tint_array_wrapper {
+  uint arr[8];
+};
+struct OuterS {
+  tint_array_wrapper a1;
+};
+
+uint f(uint i) {
+  return (i + 1u);
+}
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  OuterS s1 = {};
+  float3 v = 0.0f;
+  v[s1.a1.arr[(*(tint_symbol_1)).i]] = 1.0f;
+  v[f(s1.a1.arr[(*(tint_symbol_1)).i])] = 1.0f;
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/vector_assign.wgsl.expected.spvasm

@@ -0,0 +1,73 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 42
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpName %uniforms "uniforms"
+               OpName %f "f"
+               OpName %i "i"
+               OpName %main "main"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "a1"
+               OpName %s1 "s1"
+               OpName %v "v"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpDecorate %_arr_uint_uint_8 ArrayStride 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+          %5 = OpTypeFunction %uint %uint
+     %uint_1 = OpConstant %uint 1
+       %void = OpTypeVoid
+         %11 = OpTypeFunction %void
+     %uint_8 = OpConstant %uint 8
+%_arr_uint_uint_8 = OpTypeArray %uint %uint_8
+     %OuterS = OpTypeStruct %_arr_uint_uint_8
+%_ptr_Function_OuterS = OpTypePointer Function %OuterS
+         %20 = OpConstantNull %OuterS
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+         %25 = OpConstantNull %v3float
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+%_ptr_Function_float = OpTypePointer Function %float
+    %float_1 = OpConstant %float 1
+          %f = OpFunction %uint None %5
+          %i = OpFunctionParameter %uint
+          %8 = OpLabel
+         %10 = OpIAdd %uint %i %uint_1
+               OpReturnValue %10
+               OpFunctionEnd
+       %main = OpFunction %void None %11
+         %14 = OpLabel
+         %s1 = OpVariable %_ptr_Function_OuterS Function %20
+          %v = OpVariable %_ptr_Function_v3float Function %25
+         %28 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %29 = OpLoad %uint %28
+         %31 = OpAccessChain %_ptr_Function_uint %s1 %uint_0 %29
+         %32 = OpLoad %uint %31
+         %34 = OpAccessChain %_ptr_Function_float %v %32
+               OpStore %34 %float_1
+         %37 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %38 = OpLoad %uint %37
+         %39 = OpAccessChain %_ptr_Function_uint %s1 %uint_0 %38
+         %40 = OpLoad %uint %39
+         %36 = OpFunctionCall %uint %f %40
+         %41 = OpAccessChain %_ptr_Function_float %v %36
+               OpStore %41 %float_1
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/vector_assign.wgsl.expected.wgsl

@@ -0,0 +1,21 @@
+struct Uniforms {
+  i : u32;
+};
+
+struct OuterS {
+  a1 : array<u32, 8>;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+fn f(i : u32) -> u32 {
+  return (i + 1u);
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var s1 : OuterS;
+  var v : vec3<f32>;
+  v[s1.a1[uniforms.i]] = 1.0;
+  v[f(s1.a1[uniforms.i])] = 1.0;
+}

test/statements/assign/indexed_assign_to_array_in_struct/via_pointer.wgsl.expected.msl

@@ -0,0 +1,24 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+};
+struct InnerS {
+  int v;
+};
+struct tint_array_wrapper {
+  InnerS arr[8];
+};
+struct OuterS {
+  tint_array_wrapper a1;
+};
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_1 [[buffer(0)]]) {
+  InnerS v = {};
+  OuterS s1 = {};
+  uint const p_save = (*(tint_symbol_1)).i;
+  s1.a1.arr[p_save] = v;
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/via_pointer.wgsl.expected.spvasm

@@ -0,0 +1,55 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 26
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpName %uniforms "uniforms"
+               OpName %main "main"
+               OpName %InnerS "InnerS"
+               OpMemberName %InnerS 0 "v"
+               OpName %v "v"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "a1"
+               OpName %s1 "s1"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpMemberDecorate %InnerS 0 Offset 0
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+     %InnerS = OpTypeStruct %int
+%_ptr_Function_InnerS = OpTypePointer Function %InnerS
+         %13 = OpConstantNull %InnerS
+     %uint_8 = OpConstant %uint 8
+%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
+     %OuterS = OpTypeStruct %_arr_InnerS_uint_8
+%_ptr_Function_OuterS = OpTypePointer Function %OuterS
+         %19 = OpConstantNull %OuterS
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+          %v = OpVariable %_ptr_Function_InnerS Function %13
+         %s1 = OpVariable %_ptr_Function_OuterS Function %19
+         %22 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %23 = OpLoad %uint %22
+         %24 = OpAccessChain %_ptr_Function_InnerS %s1 %uint_0 %23
+         %25 = OpLoad %InnerS %v
+               OpStore %24 %25
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/via_pointer.wgsl.expected.wgsl

@@ -0,0 +1,21 @@
+struct Uniforms {
+  i : u32;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+struct OuterS {
+  a1 : array<InnerS, 8>;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : InnerS;
+  var s1 : OuterS;
+  let p = &(s1.a1[uniforms.i]);
+  *(p) = v;
+}

test/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Uniforms {
+  /* 0x0000 */ uint i;
+};
+struct InnerS {
+  int v;
+};
+struct tint_array_wrapper {
+  InnerS arr[8];
+};
+struct OuterS {
+  tint_array_wrapper a1;
+};
+
+void f(thread OuterS* const p, const constant Uniforms* const tint_symbol_1) {
+  InnerS v = {};
+  (*(p)).a1.arr[(*(tint_symbol_1)).i] = v;
+}
+
+kernel void tint_symbol(const constant Uniforms* tint_symbol_2 [[buffer(0)]]) {
+  OuterS s1 = {};
+  f(&(s1), tint_symbol_2);
+  return;
+}
+

test/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.spvasm

@@ -0,0 +1,64 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 33
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %Uniforms "Uniforms"
+               OpMemberName %Uniforms 0 "i"
+               OpName %uniforms "uniforms"
+               OpName %OuterS "OuterS"
+               OpMemberName %OuterS 0 "a1"
+               OpName %InnerS "InnerS"
+               OpMemberName %InnerS 0 "v"
+               OpName %f "f"
+               OpName %p "p"
+               OpName %v "v"
+               OpName %main "main"
+               OpName %s1 "s1"
+               OpDecorate %Uniforms Block
+               OpMemberDecorate %Uniforms 0 Offset 0
+               OpDecorate %uniforms NonWritable
+               OpDecorate %uniforms DescriptorSet 1
+               OpDecorate %uniforms Binding 4
+               OpMemberDecorate %OuterS 0 Offset 0
+               OpMemberDecorate %InnerS 0 Offset 0
+               OpDecorate %_arr_InnerS_uint_8 ArrayStride 4
+       %uint = OpTypeInt 32 0
+   %Uniforms = OpTypeStruct %uint
+%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms
+   %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform
+       %void = OpTypeVoid
+        %int = OpTypeInt 32 1
+     %InnerS = OpTypeStruct %int
+     %uint_8 = OpConstant %uint 8
+%_arr_InnerS_uint_8 = OpTypeArray %InnerS %uint_8
+     %OuterS = OpTypeStruct %_arr_InnerS_uint_8
+%_ptr_Function_OuterS = OpTypePointer Function %OuterS
+          %5 = OpTypeFunction %void %_ptr_Function_OuterS
+%_ptr_Function_InnerS = OpTypePointer Function %InnerS
+         %18 = OpConstantNull %InnerS
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+         %26 = OpTypeFunction %void
+         %30 = OpConstantNull %OuterS
+          %f = OpFunction %void None %5
+          %p = OpFunctionParameter %_ptr_Function_OuterS
+         %15 = OpLabel
+          %v = OpVariable %_ptr_Function_InnerS Function %18
+         %22 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0
+         %23 = OpLoad %uint %22
+         %24 = OpAccessChain %_ptr_Function_InnerS %p %uint_0 %23
+         %25 = OpLoad %InnerS %v
+               OpStore %24 %25
+               OpReturn
+               OpFunctionEnd
+       %main = OpFunction %void None %26
+         %28 = OpLabel
+         %s1 = OpVariable %_ptr_Function_OuterS Function %30
+         %31 = OpFunctionCall %void %f %s1
+               OpReturn
+               OpFunctionEnd

test/statements/assign/indexed_assign_to_array_in_struct/via_pointer_arg.wgsl.expected.wgsl

@@ -0,0 +1,24 @@
+struct Uniforms {
+  i : u32;
+};
+
+struct InnerS {
+  v : i32;
+};
+
+struct OuterS {
+  a1 : array<InnerS, 8>;
+};
+
+[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
+
+fn f(p : ptr<function, OuterS>) {
+  var v : InnerS;
+  (*(p)).a1[uniforms.i] = v;
+}
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var s1 : OuterS;
+  f(&(s1));
+}