tint->dawn: Shuffle source tree in preperation of merging repos

docs/    -> docs/tint/
fuzzers/ -> src/tint/fuzzers/
samples/ -> src/tint/cmd/
src/     -> src/tint/
test/    -> test/tint/

BUG=tint:1418,tint:1433

Change-Id: Id2aa79f989aef3245b80ef4aa37a27ff16cd700b
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/80482
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ryan Harrison <rharrison@chromium.org>

test/tint/buffer/storage/dynamic_index/read.wgsl

@@ -0,0 +1,30 @@
+struct Inner {
+    a : vec3<i32>;
+    b : i32;
+    c : vec3<u32>;
+    d : u32;
+    e : vec3<f32>;
+    f : f32;
+    g : mat2x3<f32>;
+    h : mat3x2<f32>;
+    i : array<vec4<i32>, 4>;
+};
+
+struct S {
+    arr : array<Inner>;
+};
+
+@binding(0) @group(0) var<storage, read> s : S;
+
+@stage(compute) @workgroup_size(1)
+fn main(@builtin(local_invocation_index) idx : u32) {
+    let a = s.arr[idx].a;
+    let b = s.arr[idx].b;
+    let c = s.arr[idx].c;
+    let d = s.arr[idx].d;
+    let e = s.arr[idx].e;
+    let f = s.arr[idx].f;
+    let g = s.arr[idx].g;
+    let h = s.arr[idx].h;
+    let i = s.arr[idx].i;
+}

test/tint/buffer/storage/dynamic_index/read.wgsl.expected.glsl

@@ -0,0 +1,34 @@
+#version 310 es
+
+struct Inner {
+  ivec3 a;
+  int b;
+  uvec3 c;
+  uint d;
+  vec3 e;
+  float f;
+  mat2x3 g;
+  mat3x2 h;
+  ivec4 i[4];
+};
+
+layout(binding = 0, std430) buffer S_1 {
+  Inner arr[];
+} s;
+void tint_symbol(uint idx) {
+  ivec3 a = s.arr[idx].a;
+  int b = s.arr[idx].b;
+  uvec3 c = s.arr[idx].c;
+  uint d = s.arr[idx].d;
+  vec3 e = s.arr[idx].e;
+  float f = s.arr[idx].f;
+  mat2x3 g = s.arr[idx].g;
+  mat3x2 h = s.arr[idx].h;
+  ivec4 i[4] = s.arr[idx].i;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol(gl_LocalInvocationIndex);
+  return;
+}

test/tint/buffer/storage/dynamic_index/read.wgsl.expected.hlsl

@@ -0,0 +1,42 @@
+ByteAddressBuffer s : register(t0, space0);
+
+struct tint_symbol_1 {
+  uint idx : SV_GroupIndex;
+};
+
+float2x3 tint_symbol_8(ByteAddressBuffer buffer, uint offset) {
+  return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u))));
+}
+
+float3x2 tint_symbol_9(ByteAddressBuffer buffer, uint offset) {
+  return float3x2(asfloat(buffer.Load2((offset + 0u))), asfloat(buffer.Load2((offset + 8u))), asfloat(buffer.Load2((offset + 16u))));
+}
+
+typedef int4 tint_symbol_11_ret[4];
+tint_symbol_11_ret tint_symbol_11(ByteAddressBuffer buffer, uint offset) {
+  int4 arr_1[4] = (int4[4])0;
+  {
+    [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+      arr_1[i_1] = asint(buffer.Load4((offset + (i_1 * 16u))));
+    }
+  }
+  return arr_1;
+}
+
+void main_inner(uint idx) {
+  const int3 a = asint(s.Load3((176u * idx)));
+  const int b = asint(s.Load(((176u * idx) + 12u)));
+  const uint3 c = s.Load3(((176u * idx) + 16u));
+  const uint d = s.Load(((176u * idx) + 28u));
+  const float3 e = asfloat(s.Load3(((176u * idx) + 32u)));
+  const float f = asfloat(s.Load(((176u * idx) + 44u)));
+  const float2x3 g = tint_symbol_8(s, ((176u * idx) + 48u));
+  const float3x2 h = tint_symbol_9(s, ((176u * idx) + 80u));
+  const int4 i[4] = tint_symbol_11(s, ((176u * idx) + 112u));
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.idx);
+  return;
+}

test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl

@@ -0,0 +1,52 @@
+#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 tint_array_wrapper {
+  /* 0x0000 */ int4 arr[4];
+};
+
+struct Inner {
+  /* 0x0000 */ packed_int3 a;
+  /* 0x000c */ int b;
+  /* 0x0010 */ packed_uint3 c;
+  /* 0x001c */ uint d;
+  /* 0x0020 */ packed_float3 e;
+  /* 0x002c */ float f;
+  /* 0x0030 */ float2x3 g;
+  /* 0x0050 */ float3x2 h;
+  /* 0x0068 */ int8_t tint_pad[8];
+  /* 0x0070 */ tint_array_wrapper i;
+};
+
+struct S {
+  /* 0x0000 */ Inner arr[1];
+};
+
+void tint_symbol_inner(uint idx, const device S* const tint_symbol_1) {
+  int3 const a = (*(tint_symbol_1)).arr[idx].a;
+  int const b = (*(tint_symbol_1)).arr[idx].b;
+  uint3 const c = (*(tint_symbol_1)).arr[idx].c;
+  uint const d = (*(tint_symbol_1)).arr[idx].d;
+  float3 const e = (*(tint_symbol_1)).arr[idx].e;
+  float const f = (*(tint_symbol_1)).arr[idx].f;
+  float2x3 const g = (*(tint_symbol_1)).arr[idx].g;
+  float3x2 const h = (*(tint_symbol_1)).arr[idx].h;
+  tint_array_wrapper const i = (*(tint_symbol_1)).arr[idx].i;
+}
+
+kernel void tint_symbol(const device S* tint_symbol_2 [[buffer(0)]], uint idx [[thread_index_in_threadgroup]]) {
+  tint_symbol_inner(idx, tint_symbol_2);
+  return;
+}
+

test/tint/buffer/storage/dynamic_index/read.wgsl.expected.spvasm

@@ -0,0 +1,115 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 65
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %idx_1
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %idx_1 "idx_1"
+               OpName %S "S"
+               OpMemberName %S 0 "arr"
+               OpName %Inner "Inner"
+               OpMemberName %Inner 0 "a"
+               OpMemberName %Inner 1 "b"
+               OpMemberName %Inner 2 "c"
+               OpMemberName %Inner 3 "d"
+               OpMemberName %Inner 4 "e"
+               OpMemberName %Inner 5 "f"
+               OpMemberName %Inner 6 "g"
+               OpMemberName %Inner 7 "h"
+               OpMemberName %Inner 8 "i"
+               OpName %s "s"
+               OpName %main_inner "main_inner"
+               OpName %idx "idx"
+               OpName %main "main"
+               OpDecorate %idx_1 BuiltIn LocalInvocationIndex
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %Inner 0 Offset 0
+               OpMemberDecorate %Inner 1 Offset 12
+               OpMemberDecorate %Inner 2 Offset 16
+               OpMemberDecorate %Inner 3 Offset 28
+               OpMemberDecorate %Inner 4 Offset 32
+               OpMemberDecorate %Inner 5 Offset 44
+               OpMemberDecorate %Inner 6 Offset 48
+               OpMemberDecorate %Inner 6 ColMajor
+               OpMemberDecorate %Inner 6 MatrixStride 16
+               OpMemberDecorate %Inner 7 Offset 80
+               OpMemberDecorate %Inner 7 ColMajor
+               OpMemberDecorate %Inner 7 MatrixStride 8
+               OpMemberDecorate %Inner 8 Offset 112
+               OpDecorate %_arr_v4int_uint_4 ArrayStride 16
+               OpDecorate %_runtimearr_Inner ArrayStride 176
+               OpDecorate %s NonWritable
+               OpDecorate %s Binding 0
+               OpDecorate %s DescriptorSet 0
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+      %idx_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+     %v3uint = OpTypeVector %uint 3
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+      %v4int = OpTypeVector %int 4
+     %uint_4 = OpConstant %uint 4
+%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
+      %Inner = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %mat2v3float %mat3v2float %_arr_v4int_uint_4
+%_runtimearr_Inner = OpTypeRuntimeArray %Inner
+          %S = OpTypeStruct %_runtimearr_Inner
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+          %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+       %void = OpTypeVoid
+         %20 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int
+     %uint_1 = OpConstant %uint 1
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+     %uint_2 = OpConstant %uint 2
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+     %uint_3 = OpConstant %uint 3
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+     %uint_5 = OpConstant %uint 5
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+     %uint_6 = OpConstant %uint 6
+%_ptr_StorageBuffer_mat2v3float = OpTypePointer StorageBuffer %mat2v3float
+     %uint_7 = OpConstant %uint 7
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+     %uint_8 = OpConstant %uint 8
+%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
+         %60 = OpTypeFunction %void
+ %main_inner = OpFunction %void None %20
+        %idx = OpFunctionParameter %uint
+         %24 = OpLabel
+         %27 = OpAccessChain %_ptr_StorageBuffer_v3int %s %uint_0 %idx %uint_0
+         %28 = OpLoad %v3int %27
+         %31 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0 %idx %uint_1
+         %32 = OpLoad %int %31
+         %35 = OpAccessChain %_ptr_StorageBuffer_v3uint %s %uint_0 %idx %uint_2
+         %36 = OpLoad %v3uint %35
+         %39 = OpAccessChain %_ptr_StorageBuffer_uint %s %uint_0 %idx %uint_3
+         %40 = OpLoad %uint %39
+         %42 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_0 %idx %uint_4
+         %43 = OpLoad %v3float %42
+         %46 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %idx %uint_5
+         %47 = OpLoad %float %46
+         %50 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %s %uint_0 %idx %uint_6
+         %51 = OpLoad %mat2v3float %50
+         %54 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %s %uint_0 %idx %uint_7
+         %55 = OpLoad %mat3v2float %54
+         %58 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %s %uint_0 %idx %uint_8
+         %59 = OpLoad %_arr_v4int_uint_4 %58
+               OpReturn
+               OpFunctionEnd
+       %main = OpFunction %void None %60
+         %62 = OpLabel
+         %64 = OpLoad %uint %idx_1
+         %63 = OpFunctionCall %void %main_inner %64
+               OpReturn
+               OpFunctionEnd

test/tint/buffer/storage/dynamic_index/read.wgsl.expected.wgsl

@@ -0,0 +1,30 @@
+struct Inner {
+  a : vec3<i32>;
+  b : i32;
+  c : vec3<u32>;
+  d : u32;
+  e : vec3<f32>;
+  f : f32;
+  g : mat2x3<f32>;
+  h : mat3x2<f32>;
+  i : array<vec4<i32>, 4>;
+}
+
+struct S {
+  arr : array<Inner>;
+}
+
+@binding(0) @group(0) var<storage, read> s : S;
+
+@stage(compute) @workgroup_size(1)
+fn main(@builtin(local_invocation_index) idx : u32) {
+  let a = s.arr[idx].a;
+  let b = s.arr[idx].b;
+  let c = s.arr[idx].c;
+  let d = s.arr[idx].d;
+  let e = s.arr[idx].e;
+  let f = s.arr[idx].f;
+  let g = s.arr[idx].g;
+  let h = s.arr[idx].h;
+  let i = s.arr[idx].i;
+}

test/tint/buffer/storage/dynamic_index/write.wgsl

@@ -0,0 +1,30 @@
+struct Inner {
+    a : vec3<i32>;
+    b : i32;
+    c : vec3<u32>;
+    d : u32;
+    e : vec3<f32>;
+    f : f32;
+    g : mat2x3<f32>;
+    h : mat3x2<f32>;
+    i : array<vec4<i32>, 4>;
+};
+
+struct S {
+    arr : array<Inner>;
+};
+
+@binding(0) @group(0) var<storage, read_write> s : S;
+
+@stage(compute) @workgroup_size(1)
+fn main(@builtin(local_invocation_index) idx : u32) {
+    s.arr[idx].a = vec3<i32>();
+    s.arr[idx].b = i32();
+    s.arr[idx].c = vec3<u32>();
+    s.arr[idx].d = u32();
+    s.arr[idx].e = vec3<f32>();
+    s.arr[idx].f = f32();
+    s.arr[idx].g = mat2x3<f32>();
+    s.arr[idx].h = mat3x2<f32>();
+    s.arr[idx].i = array<vec4<i32>, 4>();
+}

test/tint/buffer/storage/dynamic_index/write.wgsl.expected.glsl

@@ -0,0 +1,35 @@
+#version 310 es
+
+struct Inner {
+  ivec3 a;
+  int b;
+  uvec3 c;
+  uint d;
+  vec3 e;
+  float f;
+  mat2x3 g;
+  mat3x2 h;
+  ivec4 i[4];
+};
+
+layout(binding = 0, std430) buffer S_1 {
+  Inner arr[];
+} s;
+void tint_symbol(uint idx) {
+  s.arr[idx].a = ivec3(0, 0, 0);
+  s.arr[idx].b = 0;
+  s.arr[idx].c = uvec3(0u, 0u, 0u);
+  s.arr[idx].d = 0u;
+  s.arr[idx].e = vec3(0.0f, 0.0f, 0.0f);
+  s.arr[idx].f = 0.0f;
+  s.arr[idx].g = mat2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+  s.arr[idx].h = mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+  ivec4 tint_symbol_1[4] = ivec4[4](ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0), ivec4(0, 0, 0, 0));
+  s.arr[idx].i = tint_symbol_1;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol(gl_LocalInvocationIndex);
+  return;
+}

test/tint/buffer/storage/dynamic_index/write.wgsl.expected.hlsl

@@ -0,0 +1,44 @@
+RWByteAddressBuffer s : register(u0, space0);
+
+struct tint_symbol_1 {
+  uint idx : SV_GroupIndex;
+};
+
+void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, float2x3 value) {
+  buffer.Store3((offset + 0u), asuint(value[0u]));
+  buffer.Store3((offset + 16u), asuint(value[1u]));
+}
+
+void tint_symbol_9(RWByteAddressBuffer buffer, uint offset, float3x2 value) {
+  buffer.Store2((offset + 0u), asuint(value[0u]));
+  buffer.Store2((offset + 8u), asuint(value[1u]));
+  buffer.Store2((offset + 16u), asuint(value[2u]));
+}
+
+void tint_symbol_11(RWByteAddressBuffer buffer, uint offset, int4 value[4]) {
+  int4 array[4] = value;
+  {
+    [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+      buffer.Store4((offset + (i_1 * 16u)), asuint(array[i_1]));
+    }
+  }
+}
+
+void main_inner(uint idx) {
+  s.Store3((176u * idx), asuint(int3(0, 0, 0)));
+  s.Store(((176u * idx) + 12u), asuint(0));
+  s.Store3(((176u * idx) + 16u), asuint(uint3(0u, 0u, 0u)));
+  s.Store(((176u * idx) + 28u), asuint(0u));
+  s.Store3(((176u * idx) + 32u), asuint(float3(0.0f, 0.0f, 0.0f)));
+  s.Store(((176u * idx) + 44u), asuint(0.0f));
+  tint_symbol_8(s, ((176u * idx) + 48u), float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
+  tint_symbol_9(s, ((176u * idx) + 80u), float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
+  const int4 tint_symbol_13[4] = (int4[4])0;
+  tint_symbol_11(s, ((176u * idx) + 112u), tint_symbol_13);
+}
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.idx);
+  return;
+}

test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl

@@ -0,0 +1,53 @@
+#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 tint_array_wrapper {
+  /* 0x0000 */ int4 arr[4];
+};
+
+struct Inner {
+  /* 0x0000 */ packed_int3 a;
+  /* 0x000c */ int b;
+  /* 0x0010 */ packed_uint3 c;
+  /* 0x001c */ uint d;
+  /* 0x0020 */ packed_float3 e;
+  /* 0x002c */ float f;
+  /* 0x0030 */ float2x3 g;
+  /* 0x0050 */ float3x2 h;
+  /* 0x0068 */ int8_t tint_pad[8];
+  /* 0x0070 */ tint_array_wrapper i;
+};
+
+struct S {
+  /* 0x0000 */ Inner arr[1];
+};
+
+void tint_symbol_inner(uint idx, device S* const tint_symbol_2) {
+  (*(tint_symbol_2)).arr[idx].a = int3();
+  (*(tint_symbol_2)).arr[idx].b = int();
+  (*(tint_symbol_2)).arr[idx].c = uint3();
+  (*(tint_symbol_2)).arr[idx].d = uint();
+  (*(tint_symbol_2)).arr[idx].e = float3();
+  (*(tint_symbol_2)).arr[idx].f = float();
+  (*(tint_symbol_2)).arr[idx].g = float2x3();
+  (*(tint_symbol_2)).arr[idx].h = float3x2();
+  tint_array_wrapper const tint_symbol_1 = {.arr={}};
+  (*(tint_symbol_2)).arr[idx].i = tint_symbol_1;
+}
+
+kernel void tint_symbol(device S* tint_symbol_3 [[buffer(0)]], uint idx [[thread_index_in_threadgroup]]) {
+  tint_symbol_inner(idx, tint_symbol_3);
+  return;
+}
+

test/tint/buffer/storage/dynamic_index/write.wgsl.expected.spvasm

@@ -0,0 +1,123 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 65
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %idx_1
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %idx_1 "idx_1"
+               OpName %S "S"
+               OpMemberName %S 0 "arr"
+               OpName %Inner "Inner"
+               OpMemberName %Inner 0 "a"
+               OpMemberName %Inner 1 "b"
+               OpMemberName %Inner 2 "c"
+               OpMemberName %Inner 3 "d"
+               OpMemberName %Inner 4 "e"
+               OpMemberName %Inner 5 "f"
+               OpMemberName %Inner 6 "g"
+               OpMemberName %Inner 7 "h"
+               OpMemberName %Inner 8 "i"
+               OpName %s "s"
+               OpName %main_inner "main_inner"
+               OpName %idx "idx"
+               OpName %main "main"
+               OpDecorate %idx_1 BuiltIn LocalInvocationIndex
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %Inner 0 Offset 0
+               OpMemberDecorate %Inner 1 Offset 12
+               OpMemberDecorate %Inner 2 Offset 16
+               OpMemberDecorate %Inner 3 Offset 28
+               OpMemberDecorate %Inner 4 Offset 32
+               OpMemberDecorate %Inner 5 Offset 44
+               OpMemberDecorate %Inner 6 Offset 48
+               OpMemberDecorate %Inner 6 ColMajor
+               OpMemberDecorate %Inner 6 MatrixStride 16
+               OpMemberDecorate %Inner 7 Offset 80
+               OpMemberDecorate %Inner 7 ColMajor
+               OpMemberDecorate %Inner 7 MatrixStride 8
+               OpMemberDecorate %Inner 8 Offset 112
+               OpDecorate %_arr_v4int_uint_4 ArrayStride 16
+               OpDecorate %_runtimearr_Inner ArrayStride 176
+               OpDecorate %s Binding 0
+               OpDecorate %s DescriptorSet 0
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+      %idx_1 = OpVariable %_ptr_Input_uint Input
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+     %v3uint = OpTypeVector %uint 3
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+      %v4int = OpTypeVector %int 4
+     %uint_4 = OpConstant %uint 4
+%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
+      %Inner = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %mat2v3float %mat3v2float %_arr_v4int_uint_4
+%_runtimearr_Inner = OpTypeRuntimeArray %Inner
+          %S = OpTypeStruct %_runtimearr_Inner
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+          %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+       %void = OpTypeVoid
+         %20 = OpTypeFunction %void %uint
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int
+         %28 = OpConstantNull %v3int
+     %uint_1 = OpConstant %uint 1
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+         %32 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+         %36 = OpConstantNull %v3uint
+     %uint_3 = OpConstant %uint 3
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+         %40 = OpConstantNull %uint
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+         %43 = OpConstantNull %v3float
+     %uint_5 = OpConstant %uint 5
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+         %47 = OpConstantNull %float
+     %uint_6 = OpConstant %uint 6
+%_ptr_StorageBuffer_mat2v3float = OpTypePointer StorageBuffer %mat2v3float
+         %51 = OpConstantNull %mat2v3float
+     %uint_7 = OpConstant %uint 7
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+         %55 = OpConstantNull %mat3v2float
+     %uint_8 = OpConstant %uint 8
+%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
+         %59 = OpConstantNull %_arr_v4int_uint_4
+         %60 = OpTypeFunction %void
+ %main_inner = OpFunction %void None %20
+        %idx = OpFunctionParameter %uint
+         %24 = OpLabel
+         %27 = OpAccessChain %_ptr_StorageBuffer_v3int %s %uint_0 %idx %uint_0
+               OpStore %27 %28
+         %31 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0 %idx %uint_1
+               OpStore %31 %32
+         %35 = OpAccessChain %_ptr_StorageBuffer_v3uint %s %uint_0 %idx %uint_2
+               OpStore %35 %36
+         %39 = OpAccessChain %_ptr_StorageBuffer_uint %s %uint_0 %idx %uint_3
+               OpStore %39 %40
+         %42 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_0 %idx %uint_4
+               OpStore %42 %43
+         %46 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %idx %uint_5
+               OpStore %46 %47
+         %50 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %s %uint_0 %idx %uint_6
+               OpStore %50 %51
+         %54 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %s %uint_0 %idx %uint_7
+               OpStore %54 %55
+         %58 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %s %uint_0 %idx %uint_8
+               OpStore %58 %59
+               OpReturn
+               OpFunctionEnd
+       %main = OpFunction %void None %60
+         %62 = OpLabel
+         %64 = OpLoad %uint %idx_1
+         %63 = OpFunctionCall %void %main_inner %64
+               OpReturn
+               OpFunctionEnd

test/tint/buffer/storage/dynamic_index/write.wgsl.expected.wgsl

@@ -0,0 +1,30 @@
+struct Inner {
+  a : vec3<i32>;
+  b : i32;
+  c : vec3<u32>;
+  d : u32;
+  e : vec3<f32>;
+  f : f32;
+  g : mat2x3<f32>;
+  h : mat3x2<f32>;
+  i : array<vec4<i32>, 4>;
+}
+
+struct S {
+  arr : array<Inner>;
+}
+
+@binding(0) @group(0) var<storage, read_write> s : S;
+
+@stage(compute) @workgroup_size(1)
+fn main(@builtin(local_invocation_index) idx : u32) {
+  s.arr[idx].a = vec3<i32>();
+  s.arr[idx].b = i32();
+  s.arr[idx].c = vec3<u32>();
+  s.arr[idx].d = u32();
+  s.arr[idx].e = vec3<f32>();
+  s.arr[idx].f = f32();
+  s.arr[idx].g = mat2x3<f32>();
+  s.arr[idx].h = mat3x2<f32>();
+  s.arr[idx].i = array<vec4<i32>, 4>();
+}