writer/hlsl: Simplify UBO accesses for static indexing

Use the new semantic constant value information to significantly reduce the complex indexing logic emitted for UBO accesses.
This will dramatically reduce the number of `for` loops that are decayed to `while` loops.

Change-Id: I1b0adb5edde2b4ed39c6beafc2e28106b86e0edd
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57701
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>

test/buffer/storage/dynamic_index/read.wgsl

@@ -0,0 +1,31 @@
+struct Inner {
+    a : vec3<i32>;
+    b : i32;
+    c : vec3<u32>;
+    d : u32;
+    e : vec3<f32>;
+    f : f32;
+    g : mat2x3<f32>;
+    h : mat3x2<f32>;
+    i : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+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/buffer/storage/dynamic_index/read.wgsl.expected.hlsl

@@ -0,0 +1,34 @@
+float2x3 tint_symbol_8(ByteAddressBuffer buffer, uint offset) {
+  return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u))));
+}
+
+float3x2 tint_symbol_10(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_12_ret[4];
+tint_symbol_12_ret tint_symbol_12(ByteAddressBuffer buffer, uint offset) {
+  const int4 tint_symbol_13[4] = {asint(buffer.Load4((offset + 0u))), asint(buffer.Load4((offset + 16u))), asint(buffer.Load4((offset + 32u))), asint(buffer.Load4((offset + 48u)))};
+  return tint_symbol_13;
+}
+
+ByteAddressBuffer s : register(t0, space0);
+
+struct tint_symbol_1 {
+  uint idx : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  const uint idx = tint_symbol.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_10(s, ((176u * idx) + 80u));
+  const int4 i[4] = tint_symbol_12(s, ((176u * idx) + 112u));
+  return;
+}

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

@@ -0,0 +1,35 @@
+#include <metal_stdlib>
+
+using namespace metal;
+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];
+};
+
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], const device S& s [[buffer(0)]]) {
+  int3 const a = s.arr[idx].a;
+  int const b = s.arr[idx].b;
+  uint3 const c = s.arr[idx].c;
+  uint const d = s.arr[idx].d;
+  float3 const e = s.arr[idx].e;
+  float const f = s.arr[idx].f;
+  float2x3 const g = s.arr[idx].g;
+  float3x2 const h = s.arr[idx].h;
+  tint_array_wrapper const i = s.arr[idx].i;
+  return;
+}
+

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

@@ -0,0 +1,114 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 68
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %tint_symbol
+               OpExecutionMode %main LocalSize 1 1 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 %tint_symbol "tint_symbol"
+               OpName %main "main"
+               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
+               OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+       %uint = OpTypeInt 32 0
+     %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
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+       %void = OpTypeVoid
+         %20 = OpTypeFunction %void
+     %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
+       %main = OpFunction %void None %20
+         %23 = OpLabel
+         %25 = OpLoad %uint %tint_symbol
+         %27 = OpAccessChain %_ptr_StorageBuffer_v3int %s %uint_0 %25 %uint_0
+         %28 = OpLoad %v3int %27
+         %29 = OpLoad %uint %tint_symbol
+         %32 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0 %29 %uint_1
+         %33 = OpLoad %int %32
+         %34 = OpLoad %uint %tint_symbol
+         %37 = OpAccessChain %_ptr_StorageBuffer_v3uint %s %uint_0 %34 %uint_2
+         %38 = OpLoad %v3uint %37
+         %39 = OpLoad %uint %tint_symbol
+         %42 = OpAccessChain %_ptr_StorageBuffer_uint %s %uint_0 %39 %uint_3
+         %43 = OpLoad %uint %42
+         %44 = OpLoad %uint %tint_symbol
+         %46 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_0 %44 %uint_4
+         %47 = OpLoad %v3float %46
+         %48 = OpLoad %uint %tint_symbol
+         %51 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %48 %uint_5
+         %52 = OpLoad %float %51
+         %53 = OpLoad %uint %tint_symbol
+         %56 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %s %uint_0 %53 %uint_6
+         %57 = OpLoad %mat2v3float %56
+         %58 = OpLoad %uint %tint_symbol
+         %61 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %s %uint_0 %58 %uint_7
+         %62 = OpLoad %mat3v2float %61
+         %63 = OpLoad %uint %tint_symbol
+         %66 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %s %uint_0 %63 %uint_8
+         %67 = OpLoad %_arr_v4int_uint_4 %66
+               OpReturn
+               OpFunctionEnd

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

@@ -0,0 +1,31 @@
+struct Inner {
+  a : vec3<i32>;
+  b : i32;
+  c : vec3<u32>;
+  d : u32;
+  e : vec3<f32>;
+  f : f32;
+  g : mat2x3<f32>;
+  h : mat3x2<f32>;
+  i : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+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/buffer/storage/dynamic_index/write.wgsl

@@ -0,0 +1,31 @@
+struct Inner {
+    a : vec3<i32>;
+    b : i32;
+    c : vec3<u32>;
+    d : u32;
+    e : vec3<f32>;
+    f : f32;
+    g : mat2x3<f32>;
+    h : mat3x2<f32>;
+    i : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+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 = [[stride(16)]] array<vec4<i32>, 4>();
+}

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

@@ -0,0 +1,39 @@
+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_10(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_12(RWByteAddressBuffer buffer, uint offset, int4 value[4]) {
+  buffer.Store4((offset + 0u), asuint(value[0u]));
+  buffer.Store4((offset + 16u), asuint(value[1u]));
+  buffer.Store4((offset + 32u), asuint(value[2u]));
+  buffer.Store4((offset + 48u), asuint(value[3u]));
+}
+
+RWByteAddressBuffer s : register(u0, space0);
+
+struct tint_symbol_1 {
+  uint idx : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  const uint idx = tint_symbol.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_10(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_12(s, ((176u * idx) + 112u), tint_symbol_13);
+  return;
+}

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

@@ -0,0 +1,36 @@
+#include <metal_stdlib>
+
+using namespace metal;
+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];
+};
+
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], device S& s [[buffer(0)]]) {
+  s.arr[idx].a = int3();
+  s.arr[idx].b = int();
+  s.arr[idx].c = uint3();
+  s.arr[idx].d = uint();
+  s.arr[idx].e = float3();
+  s.arr[idx].f = float();
+  s.arr[idx].g = float2x3();
+  s.arr[idx].h = float3x2();
+  tint_array_wrapper const tint_symbol_2 = {.arr={}};
+  s.arr[idx].i = tint_symbol_2;
+  return;
+}
+

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

@@ -0,0 +1,122 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 68
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %tint_symbol
+               OpExecutionMode %main LocalSize 1 1 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 %tint_symbol "tint_symbol"
+               OpName %main "main"
+               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
+               OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+       %uint = OpTypeInt 32 0
+     %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
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+       %void = OpTypeVoid
+         %20 = OpTypeFunction %void
+     %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
+         %33 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+         %38 = OpConstantNull %v3uint
+     %uint_3 = OpConstant %uint 3
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+         %43 = OpConstantNull %uint
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+         %47 = OpConstantNull %v3float
+     %uint_5 = OpConstant %uint 5
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+         %52 = OpConstantNull %float
+     %uint_6 = OpConstant %uint 6
+%_ptr_StorageBuffer_mat2v3float = OpTypePointer StorageBuffer %mat2v3float
+         %57 = OpConstantNull %mat2v3float
+     %uint_7 = OpConstant %uint 7
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+         %62 = OpConstantNull %mat3v2float
+     %uint_8 = OpConstant %uint 8
+%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
+         %67 = OpConstantNull %_arr_v4int_uint_4
+       %main = OpFunction %void None %20
+         %23 = OpLabel
+         %25 = OpLoad %uint %tint_symbol
+         %27 = OpAccessChain %_ptr_StorageBuffer_v3int %s %uint_0 %25 %uint_0
+               OpStore %27 %28
+         %29 = OpLoad %uint %tint_symbol
+         %32 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0 %29 %uint_1
+               OpStore %32 %33
+         %34 = OpLoad %uint %tint_symbol
+         %37 = OpAccessChain %_ptr_StorageBuffer_v3uint %s %uint_0 %34 %uint_2
+               OpStore %37 %38
+         %39 = OpLoad %uint %tint_symbol
+         %42 = OpAccessChain %_ptr_StorageBuffer_uint %s %uint_0 %39 %uint_3
+               OpStore %42 %43
+         %44 = OpLoad %uint %tint_symbol
+         %46 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_0 %44 %uint_4
+               OpStore %46 %47
+         %48 = OpLoad %uint %tint_symbol
+         %51 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %48 %uint_5
+               OpStore %51 %52
+         %53 = OpLoad %uint %tint_symbol
+         %56 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %s %uint_0 %53 %uint_6
+               OpStore %56 %57
+         %58 = OpLoad %uint %tint_symbol
+         %61 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %s %uint_0 %58 %uint_7
+               OpStore %61 %62
+         %63 = OpLoad %uint %tint_symbol
+         %66 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %s %uint_0 %63 %uint_8
+               OpStore %66 %67
+               OpReturn
+               OpFunctionEnd

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

@@ -0,0 +1,31 @@
+struct Inner {
+  a : vec3<i32>;
+  b : i32;
+  c : vec3<u32>;
+  d : u32;
+  e : vec3<f32>;
+  f : f32;
+  g : mat2x3<f32>;
+  h : mat3x2<f32>;
+  i : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+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 = [[stride(16)]] array<vec4<i32>, 4>();
+}

test/buffer/storage/static_index/read.wgsl

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

test/buffer/storage/static_index/read.wgsl.expected.hlsl

@@ -0,0 +1,42 @@
+struct Inner {
+  int x;
+};
+struct tint_padded_array_element {
+  Inner el;
+};
+
+float2x3 tint_symbol_6(ByteAddressBuffer buffer, uint offset) {
+  return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u))));
+}
+
+float3x2 tint_symbol_8(ByteAddressBuffer buffer, uint offset) {
+  return float3x2(asfloat(buffer.Load2((offset + 0u))), asfloat(buffer.Load2((offset + 8u))), asfloat(buffer.Load2((offset + 16u))));
+}
+
+Inner tint_symbol_9(ByteAddressBuffer buffer, uint offset) {
+  const Inner tint_symbol_11 = {asint(buffer.Load((offset + 0u)))};
+  return tint_symbol_11;
+}
+
+typedef tint_padded_array_element tint_symbol_10_ret[4];
+tint_symbol_10_ret tint_symbol_10(ByteAddressBuffer buffer, uint offset) {
+  const tint_padded_array_element tint_symbol_12[4] = {{tint_symbol_9(buffer, (offset + 0u))}, {tint_symbol_9(buffer, (offset + 16u))}, {tint_symbol_9(buffer, (offset + 32u))}, {tint_symbol_9(buffer, (offset + 48u))}};
+  return tint_symbol_12;
+}
+
+ByteAddressBuffer s : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  const int3 a = asint(s.Load3(0u));
+  const int b = asint(s.Load(12u));
+  const uint3 c = s.Load3(16u);
+  const uint d = s.Load(28u);
+  const float3 e = asfloat(s.Load3(32u));
+  const float f = asfloat(s.Load(44u));
+  const float2x3 g = tint_symbol_6(s, 48u);
+  const float3x2 h = tint_symbol_8(s, 80u);
+  const Inner i = tint_symbol_9(s, 104u);
+  const tint_padded_array_element j[4] = tint_symbol_10(s, 108u);
+  return;
+}

test/buffer/storage/static_index/read.wgsl.expected.msl

@@ -0,0 +1,41 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Inner {
+  /* 0x0000 */ int x;
+};
+struct tint_padded_array_element {
+  /* 0x0000 */ Inner el;
+  /* 0x0004 */ int8_t tint_pad[12];
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ tint_padded_array_element arr[4];
+};
+struct S {
+  /* 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 */ Inner i;
+  /* 0x006c */ tint_array_wrapper j;
+  /* 0x00ac */ int8_t tint_pad_1[4];
+};
+
+kernel void tint_symbol(const device S& s [[buffer(0)]]) {
+  int3 const a = s.a;
+  int const b = s.b;
+  uint3 const c = s.c;
+  uint const d = s.d;
+  float3 const e = s.e;
+  float const f = s.f;
+  float2x3 const g = s.g;
+  float3x2 const h = s.h;
+  Inner const i = s.i;
+  tint_array_wrapper const j = s.j;
+  return;
+}
+

test/buffer/storage/static_index/read.wgsl.expected.spvasm

@@ -0,0 +1,104 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 59
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "a"
+               OpMemberName %S 1 "b"
+               OpMemberName %S 2 "c"
+               OpMemberName %S 3 "d"
+               OpMemberName %S 4 "e"
+               OpMemberName %S 5 "f"
+               OpMemberName %S 6 "g"
+               OpMemberName %S 7 "h"
+               OpMemberName %S 8 "i"
+               OpName %Inner "Inner"
+               OpMemberName %Inner 0 "x"
+               OpMemberName %S 9 "j"
+               OpName %s "s"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 12
+               OpMemberDecorate %S 2 Offset 16
+               OpMemberDecorate %S 3 Offset 28
+               OpMemberDecorate %S 4 Offset 32
+               OpMemberDecorate %S 5 Offset 44
+               OpMemberDecorate %S 6 Offset 48
+               OpMemberDecorate %S 6 ColMajor
+               OpMemberDecorate %S 6 MatrixStride 16
+               OpMemberDecorate %S 7 Offset 80
+               OpMemberDecorate %S 7 ColMajor
+               OpMemberDecorate %S 7 MatrixStride 8
+               OpMemberDecorate %S 8 Offset 104
+               OpMemberDecorate %Inner 0 Offset 0
+               OpMemberDecorate %S 9 Offset 108
+               OpDecorate %_arr_Inner_uint_4 ArrayStride 16
+               OpDecorate %s NonWritable
+               OpDecorate %s Binding 0
+               OpDecorate %s DescriptorSet 0
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+      %Inner = OpTypeStruct %int
+     %uint_4 = OpConstant %uint 4
+%_arr_Inner_uint_4 = OpTypeArray %Inner %uint_4
+          %S = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %mat2v3float %mat3v2float %Inner %_arr_Inner_uint_4
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+          %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+     %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_Inner = OpTypePointer StorageBuffer %Inner
+     %uint_9 = OpConstant %uint 9
+%_ptr_StorageBuffer__arr_Inner_uint_4 = OpTypePointer StorageBuffer %_arr_Inner_uint_4
+       %main = OpFunction %void None %16
+         %19 = OpLabel
+         %22 = OpAccessChain %_ptr_StorageBuffer_v3int %s %uint_0
+         %23 = OpLoad %v3int %22
+         %26 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_1
+         %27 = OpLoad %int %26
+         %30 = OpAccessChain %_ptr_StorageBuffer_v3uint %s %uint_2
+         %31 = OpLoad %v3uint %30
+         %34 = OpAccessChain %_ptr_StorageBuffer_uint %s %uint_3
+         %35 = OpLoad %uint %34
+         %37 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_4
+         %38 = OpLoad %v3float %37
+         %41 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_5
+         %42 = OpLoad %float %41
+         %45 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %s %uint_6
+         %46 = OpLoad %mat2v3float %45
+         %49 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %s %uint_7
+         %50 = OpLoad %mat3v2float %49
+         %53 = OpAccessChain %_ptr_StorageBuffer_Inner %s %uint_8
+         %54 = OpLoad %Inner %53
+         %57 = OpAccessChain %_ptr_StorageBuffer__arr_Inner_uint_4 %s %uint_9
+         %58 = OpLoad %_arr_Inner_uint_4 %57
+               OpReturn
+               OpFunctionEnd

test/buffer/storage/static_index/read.wgsl.expected.wgsl

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

test/buffer/storage/static_index/write.wgsl

@@ -0,0 +1,33 @@
+struct Inner {
+    x : i32;
+};
+
+[[block]]
+struct S {
+    a : vec3<i32>;
+    b : i32;
+    c : vec3<u32>;
+    d : u32;
+    e : vec3<f32>;
+    f : f32;
+    g : mat2x3<f32>;
+    h : mat3x2<f32>;
+    i : Inner;
+    j : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<storage, write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    s.a = vec3<i32>();
+    s.b = i32();
+    s.c = vec3<u32>();
+    s.d = u32();
+    s.e = vec3<f32>();
+    s.f = f32();
+    s.g = mat2x3<f32>();
+    s.h = mat3x2<f32>();
+    s.i = Inner();
+    s.j = [[stride(16)]] array<Inner, 4>();
+}

test/buffer/storage/static_index/write.wgsl.expected.hlsl

@@ -0,0 +1,47 @@
+struct Inner {
+  int x;
+};
+struct tint_padded_array_element {
+  Inner el;
+};
+
+void tint_symbol_6(RWByteAddressBuffer buffer, uint offset, float2x3 value) {
+  buffer.Store3((offset + 0u), asuint(value[0u]));
+  buffer.Store3((offset + 16u), asuint(value[1u]));
+}
+
+void tint_symbol_8(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_9(RWByteAddressBuffer buffer, uint offset, Inner value) {
+  buffer.Store((offset + 0u), asuint(value.x));
+}
+
+void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) {
+  tint_symbol_9(buffer, (offset + 0u), value[0u].el);
+  tint_symbol_9(buffer, (offset + 16u), value[1u].el);
+  tint_symbol_9(buffer, (offset + 32u), value[2u].el);
+  tint_symbol_9(buffer, (offset + 48u), value[3u].el);
+}
+
+RWByteAddressBuffer s : register(u0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  s.Store3(0u, asuint(int3(0, 0, 0)));
+  s.Store(12u, asuint(0));
+  s.Store3(16u, asuint(uint3(0u, 0u, 0u)));
+  s.Store(28u, asuint(0u));
+  s.Store3(32u, asuint(float3(0.0f, 0.0f, 0.0f)));
+  s.Store(44u, asuint(0.0f));
+  tint_symbol_6(s, 48u, float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
+  tint_symbol_8(s, 80u, float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
+  const Inner tint_symbol_11 = (Inner)0;
+  tint_symbol_9(s, 104u, tint_symbol_11);
+  const tint_padded_array_element tint_symbol_12[4] = (tint_padded_array_element[4])0;
+  tint_symbol_10(s, 108u, tint_symbol_12);
+  return;
+}

test/buffer/storage/static_index/write.wgsl.expected.msl

@@ -0,0 +1,43 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Inner {
+  /* 0x0000 */ int x;
+};
+struct tint_padded_array_element {
+  /* 0x0000 */ Inner el;
+  /* 0x0004 */ int8_t tint_pad[12];
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ tint_padded_array_element arr[4];
+};
+struct S {
+  /* 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 */ Inner i;
+  /* 0x006c */ tint_array_wrapper j;
+  /* 0x00ac */ int8_t tint_pad_1[4];
+};
+
+kernel void tint_symbol(device S& s [[buffer(0)]]) {
+  s.a = int3();
+  s.b = int();
+  s.c = uint3();
+  s.d = uint();
+  s.e = float3();
+  s.f = float();
+  s.g = float2x3();
+  s.h = float3x2();
+  Inner const tint_symbol_1 = {};
+  s.i = tint_symbol_1;
+  tint_array_wrapper const tint_symbol_2 = {.arr={}};
+  s.j = tint_symbol_2;
+  return;
+}
+

test/buffer/storage/static_index/write.wgsl.expected.spvasm

@@ -0,0 +1,114 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 59
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "a"
+               OpMemberName %S 1 "b"
+               OpMemberName %S 2 "c"
+               OpMemberName %S 3 "d"
+               OpMemberName %S 4 "e"
+               OpMemberName %S 5 "f"
+               OpMemberName %S 6 "g"
+               OpMemberName %S 7 "h"
+               OpMemberName %S 8 "i"
+               OpName %Inner "Inner"
+               OpMemberName %Inner 0 "x"
+               OpMemberName %S 9 "j"
+               OpName %s "s"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 12
+               OpMemberDecorate %S 2 Offset 16
+               OpMemberDecorate %S 3 Offset 28
+               OpMemberDecorate %S 4 Offset 32
+               OpMemberDecorate %S 5 Offset 44
+               OpMemberDecorate %S 6 Offset 48
+               OpMemberDecorate %S 6 ColMajor
+               OpMemberDecorate %S 6 MatrixStride 16
+               OpMemberDecorate %S 7 Offset 80
+               OpMemberDecorate %S 7 ColMajor
+               OpMemberDecorate %S 7 MatrixStride 8
+               OpMemberDecorate %S 8 Offset 104
+               OpMemberDecorate %Inner 0 Offset 0
+               OpMemberDecorate %S 9 Offset 108
+               OpDecorate %_arr_Inner_uint_4 ArrayStride 16
+               OpDecorate %s NonReadable
+               OpDecorate %s Binding 0
+               OpDecorate %s DescriptorSet 0
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+      %Inner = OpTypeStruct %int
+     %uint_4 = OpConstant %uint 4
+%_arr_Inner_uint_4 = OpTypeArray %Inner %uint_4
+          %S = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %mat2v3float %mat3v2float %Inner %_arr_Inner_uint_4
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+          %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int
+         %23 = OpConstantNull %v3int
+     %uint_1 = OpConstant %uint 1
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+         %27 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+         %31 = OpConstantNull %v3uint
+     %uint_3 = OpConstant %uint 3
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+         %35 = OpConstantNull %uint
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+         %38 = OpConstantNull %v3float
+     %uint_5 = OpConstant %uint 5
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+         %42 = OpConstantNull %float
+     %uint_6 = OpConstant %uint 6
+%_ptr_StorageBuffer_mat2v3float = OpTypePointer StorageBuffer %mat2v3float
+         %46 = OpConstantNull %mat2v3float
+     %uint_7 = OpConstant %uint 7
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+         %50 = OpConstantNull %mat3v2float
+     %uint_8 = OpConstant %uint 8
+%_ptr_StorageBuffer_Inner = OpTypePointer StorageBuffer %Inner
+         %54 = OpConstantNull %Inner
+     %uint_9 = OpConstant %uint 9
+%_ptr_StorageBuffer__arr_Inner_uint_4 = OpTypePointer StorageBuffer %_arr_Inner_uint_4
+         %58 = OpConstantNull %_arr_Inner_uint_4
+       %main = OpFunction %void None %16
+         %19 = OpLabel
+         %22 = OpAccessChain %_ptr_StorageBuffer_v3int %s %uint_0
+               OpStore %22 %23
+         %26 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_1
+               OpStore %26 %27
+         %30 = OpAccessChain %_ptr_StorageBuffer_v3uint %s %uint_2
+               OpStore %30 %31
+         %34 = OpAccessChain %_ptr_StorageBuffer_uint %s %uint_3
+               OpStore %34 %35
+         %37 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_4
+               OpStore %37 %38
+         %41 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_5
+               OpStore %41 %42
+         %45 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %s %uint_6
+               OpStore %45 %46
+         %49 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %s %uint_7
+               OpStore %49 %50
+         %53 = OpAccessChain %_ptr_StorageBuffer_Inner %s %uint_8
+               OpStore %53 %54
+         %57 = OpAccessChain %_ptr_StorageBuffer__arr_Inner_uint_4 %s %uint_9
+               OpStore %57 %58
+               OpReturn
+               OpFunctionEnd

test/buffer/storage/static_index/write.wgsl.expected.wgsl

@@ -0,0 +1,33 @@
+struct Inner {
+  x : i32;
+};
+
+[[block]]
+struct S {
+  a : vec3<i32>;
+  b : i32;
+  c : vec3<u32>;
+  d : u32;
+  e : vec3<f32>;
+  f : f32;
+  g : mat2x3<f32>;
+  h : mat3x2<f32>;
+  i : Inner;
+  j : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<storage, write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  s.a = vec3<i32>();
+  s.b = i32();
+  s.c = vec3<u32>();
+  s.d = u32();
+  s.e = vec3<f32>();
+  s.f = f32();
+  s.g = mat2x3<f32>();
+  s.h = mat3x2<f32>();
+  s.i = Inner();
+  s.j = [[stride(16)]] array<Inner, 4>();
+}

test/buffer/uniform/dynamic_index/read.wgsl

@@ -0,0 +1,33 @@
+struct Inner {
+    a : vec3<i32>;
+    b : i32;
+    c : vec3<u32>;
+    d : u32;
+    e : vec3<f32>;
+    f : f32;
+    g : vec2<i32>;
+    h : vec2<i32>;
+    i : mat2x3<f32>;
+    [[align(16)]] j : mat3x2<f32>;
+    [[align(16)]] k : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+struct S {
+    arr : array<Inner, 8>;
+};
+
+[[binding(0), group(0)]] var<uniform> 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/buffer/uniform/dynamic_index/read.wgsl.expected.hlsl

@@ -0,0 +1,38 @@
+float2x3 tint_symbol_9(uint4 buffer[96], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  return float2x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz));
+}
+
+cbuffer cbuffer_s : register(b0, space0) {
+  uint4 s[96];
+};
+
+struct tint_symbol_1 {
+  uint idx : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  const uint idx = tint_symbol.idx;
+  const uint scalar_offset_2 = ((192u * idx)) / 4;
+  const int3 a = asint(s[scalar_offset_2 / 4].xyz);
+  const uint scalar_offset_3 = (((192u * idx) + 12u)) / 4;
+  const int b = asint(s[scalar_offset_3 / 4][scalar_offset_3 % 4]);
+  const uint scalar_offset_4 = (((192u * idx) + 16u)) / 4;
+  const uint3 c = s[scalar_offset_4 / 4].xyz;
+  const uint scalar_offset_5 = (((192u * idx) + 28u)) / 4;
+  const uint d = s[scalar_offset_5 / 4][scalar_offset_5 % 4];
+  const uint scalar_offset_6 = (((192u * idx) + 32u)) / 4;
+  const float3 e = asfloat(s[scalar_offset_6 / 4].xyz);
+  const uint scalar_offset_7 = (((192u * idx) + 44u)) / 4;
+  const float f = asfloat(s[scalar_offset_7 / 4][scalar_offset_7 % 4]);
+  const uint scalar_offset_8 = (((192u * idx) + 48u)) / 4;
+  uint4 ubo_load = s[scalar_offset_8 / 4];
+  const int2 g = asint(((scalar_offset_8 & 2) ? ubo_load.zw : ubo_load.xy));
+  const uint scalar_offset_9 = (((192u * idx) + 56u)) / 4;
+  uint4 ubo_load_1 = s[scalar_offset_9 / 4];
+  const int2 h = asint(((scalar_offset_9 & 2) ? ubo_load_1.zw : ubo_load_1.xy));
+  const float2x3 i = tint_symbol_9(s, ((192u * idx) + 64u));
+  return;
+}

test/buffer/uniform/dynamic_index/read.wgsl.expected.msl

@@ -0,0 +1,40 @@
+#include <metal_stdlib>
+
+using namespace metal;
+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 */ packed_int2 g;
+  /* 0x0038 */ packed_int2 h;
+  /* 0x0040 */ float2x3 i;
+  /* 0x0060 */ float3x2 j;
+  /* 0x0078 */ int8_t tint_pad[8];
+  /* 0x0080 */ tint_array_wrapper k;
+};
+struct tint_array_wrapper_1 {
+  /* 0x0000 */ Inner arr[8];
+};
+struct S {
+  /* 0x0000 */ tint_array_wrapper_1 arr;
+};
+
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], constant S& s [[buffer(0)]]) {
+  int3 const a = s.arr.arr[idx].a;
+  int const b = s.arr.arr[idx].b;
+  uint3 const c = s.arr.arr[idx].c;
+  uint const d = s.arr.arr[idx].d;
+  float3 const e = s.arr.arr[idx].e;
+  float const f = s.arr.arr[idx].f;
+  int2 const g = s.arr.arr[idx].g;
+  int2 const h = s.arr.arr[idx].h;
+  float2x3 const i = s.arr.arr[idx].i;
+  return;
+}
+

test/buffer/uniform/dynamic_index/read.wgsl.expected.spvasm

@@ -0,0 +1,118 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 68
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %tint_symbol
+               OpExecutionMode %main LocalSize 1 1 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"
+               OpMemberName %Inner 9 "j"
+               OpMemberName %Inner 10 "k"
+               OpName %s "s"
+               OpName %tint_symbol "tint_symbol"
+               OpName %main "main"
+               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 7 Offset 56
+               OpMemberDecorate %Inner 8 Offset 64
+               OpMemberDecorate %Inner 8 ColMajor
+               OpMemberDecorate %Inner 8 MatrixStride 16
+               OpMemberDecorate %Inner 9 Offset 96
+               OpMemberDecorate %Inner 9 ColMajor
+               OpMemberDecorate %Inner 9 MatrixStride 8
+               OpMemberDecorate %Inner 10 Offset 128
+               OpDecorate %_arr_v4int_uint_4 ArrayStride 16
+               OpDecorate %_arr_Inner_uint_8 ArrayStride 192
+               OpDecorate %s NonWritable
+               OpDecorate %s Binding 0
+               OpDecorate %s DescriptorSet 0
+               OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+      %v2int = OpTypeVector %int 2
+%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 %v2int %v2int %mat2v3float %mat3v2float %_arr_v4int_uint_4
+     %uint_8 = OpConstant %uint 8
+%_arr_Inner_uint_8 = OpTypeArray %Inner %uint_8
+          %S = OpTypeStruct %_arr_Inner_uint_8
+%_ptr_Uniform_S = OpTypePointer Uniform %S
+          %s = OpVariable %_ptr_Uniform_S Uniform
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+       %void = OpTypeVoid
+         %22 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_v3int = OpTypePointer Uniform %v3int
+     %uint_1 = OpConstant %uint 1
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+     %uint_2 = OpConstant %uint 2
+%_ptr_Uniform_v3uint = OpTypePointer Uniform %v3uint
+     %uint_3 = OpConstant %uint 3
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
+     %uint_5 = OpConstant %uint 5
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+     %uint_6 = OpConstant %uint 6
+%_ptr_Uniform_v2int = OpTypePointer Uniform %v2int
+     %uint_7 = OpConstant %uint 7
+%_ptr_Uniform_mat2v3float = OpTypePointer Uniform %mat2v3float
+       %main = OpFunction %void None %22
+         %25 = OpLabel
+         %27 = OpLoad %uint %tint_symbol
+         %29 = OpAccessChain %_ptr_Uniform_v3int %s %uint_0 %27 %uint_0
+         %30 = OpLoad %v3int %29
+         %31 = OpLoad %uint %tint_symbol
+         %34 = OpAccessChain %_ptr_Uniform_int %s %uint_0 %31 %uint_1
+         %35 = OpLoad %int %34
+         %36 = OpLoad %uint %tint_symbol
+         %39 = OpAccessChain %_ptr_Uniform_v3uint %s %uint_0 %36 %uint_2
+         %40 = OpLoad %v3uint %39
+         %41 = OpLoad %uint %tint_symbol
+         %44 = OpAccessChain %_ptr_Uniform_uint %s %uint_0 %41 %uint_3
+         %45 = OpLoad %uint %44
+         %46 = OpLoad %uint %tint_symbol
+         %48 = OpAccessChain %_ptr_Uniform_v3float %s %uint_0 %46 %uint_4
+         %49 = OpLoad %v3float %48
+         %50 = OpLoad %uint %tint_symbol
+         %53 = OpAccessChain %_ptr_Uniform_float %s %uint_0 %50 %uint_5
+         %54 = OpLoad %float %53
+         %55 = OpLoad %uint %tint_symbol
+         %58 = OpAccessChain %_ptr_Uniform_v2int %s %uint_0 %55 %uint_6
+         %59 = OpLoad %v2int %58
+         %60 = OpLoad %uint %tint_symbol
+         %62 = OpAccessChain %_ptr_Uniform_v2int %s %uint_0 %60 %uint_7
+         %63 = OpLoad %v2int %62
+         %64 = OpLoad %uint %tint_symbol
+         %66 = OpAccessChain %_ptr_Uniform_mat2v3float %s %uint_0 %64 %uint_8
+         %67 = OpLoad %mat2v3float %66
+               OpReturn
+               OpFunctionEnd

test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl

@@ -0,0 +1,35 @@
+struct Inner {
+  a : vec3<i32>;
+  b : i32;
+  c : vec3<u32>;
+  d : u32;
+  e : vec3<f32>;
+  f : f32;
+  g : vec2<i32>;
+  h : vec2<i32>;
+  i : mat2x3<f32>;
+  [[align(16)]]
+  j : mat3x2<f32>;
+  [[align(16)]]
+  k : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+struct S {
+  arr : array<Inner, 8>;
+};
+
+[[binding(0), group(0)]] var<uniform> 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/buffer/uniform/static_index/read.wgsl

@@ -0,0 +1,37 @@
+struct Inner {
+    x : i32;
+};
+
+[[block]]
+struct S {
+    a : vec3<i32>;
+    b : i32;
+    c : vec3<u32>;
+    d : u32;
+    e : vec3<f32>;
+    f : f32;
+    g : vec2<i32>;
+    h : vec2<i32>;
+    i : mat2x3<f32>;
+    j : mat3x2<f32>;
+    [[align(16)]] k : Inner;
+    [[align(16)]] l : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<uniform> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    let a = s.a;
+    let b = s.b;
+    let c = s.c;
+    let d = s.d;
+    let e = s.e;
+    let f = s.f;
+    let g = s.g;
+    let h = s.h;
+    let i = s.i;
+    let j = s.j;
+    let k = s.k;
+    let l = s.l;
+}

test/buffer/uniform/static_index/read.wgsl.expected.hlsl

@@ -0,0 +1,55 @@
+struct Inner {
+  int x;
+};
+struct tint_padded_array_element {
+  Inner el;
+};
+
+float2x3 tint_symbol_7(uint4 buffer[13], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  return float2x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz));
+}
+
+float3x2 tint_symbol_9(uint4 buffer[13], uint offset) {
+  const uint scalar_offset_2 = ((offset + 0u)) / 4;
+  uint4 ubo_load = buffer[scalar_offset_2 / 4];
+  const uint scalar_offset_3 = ((offset + 8u)) / 4;
+  uint4 ubo_load_1 = buffer[scalar_offset_3 / 4];
+  const uint scalar_offset_4 = ((offset + 16u)) / 4;
+  uint4 ubo_load_2 = buffer[scalar_offset_4 / 4];
+  return float3x2(asfloat(((scalar_offset_2 & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_4 & 2) ? ubo_load_2.zw : ubo_load_2.xy)));
+}
+
+Inner tint_symbol_10(uint4 buffer[13], uint offset) {
+  const uint scalar_offset_5 = ((offset + 0u)) / 4;
+  const Inner tint_symbol_12 = {asint(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4])};
+  return tint_symbol_12;
+}
+
+typedef tint_padded_array_element tint_symbol_11_ret[4];
+tint_symbol_11_ret tint_symbol_11(uint4 buffer[13], uint offset) {
+  const tint_padded_array_element tint_symbol_13[4] = {{tint_symbol_10(buffer, (offset + 0u))}, {tint_symbol_10(buffer, (offset + 16u))}, {tint_symbol_10(buffer, (offset + 32u))}, {tint_symbol_10(buffer, (offset + 48u))}};
+  return tint_symbol_13;
+}
+
+cbuffer cbuffer_s : register(b0, space0) {
+  uint4 s[13];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  const int3 a = asint(s[0].xyz);
+  const int b = asint(s[0].w);
+  const uint3 c = s[1].xyz;
+  const uint d = s[1].w;
+  const float3 e = asfloat(s[2].xyz);
+  const float f = asfloat(s[2].w);
+  const int2 g = asint(s[3].xy);
+  const int2 h = asint(s[3].zw);
+  const float2x3 i = tint_symbol_7(s, 64u);
+  const float3x2 j = tint_symbol_9(s, 96u);
+  const Inner k = tint_symbol_10(s, 128u);
+  const tint_padded_array_element l[4] = tint_symbol_11(s, 144u);
+  return;
+}

test/buffer/uniform/static_index/read.wgsl.expected.msl

@@ -0,0 +1,46 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Inner {
+  /* 0x0000 */ int x;
+};
+struct tint_padded_array_element {
+  /* 0x0000 */ Inner el;
+  /* 0x0004 */ int8_t tint_pad[12];
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ tint_padded_array_element arr[4];
+};
+struct S {
+  /* 0x0000 */ packed_int3 a;
+  /* 0x000c */ int b;
+  /* 0x0010 */ packed_uint3 c;
+  /* 0x001c */ uint d;
+  /* 0x0020 */ packed_float3 e;
+  /* 0x002c */ float f;
+  /* 0x0030 */ packed_int2 g;
+  /* 0x0038 */ packed_int2 h;
+  /* 0x0040 */ float2x3 i;
+  /* 0x0060 */ float3x2 j;
+  /* 0x0078 */ int8_t tint_pad_1[8];
+  /* 0x0080 */ Inner k;
+  /* 0x0084 */ int8_t tint_pad_2[12];
+  /* 0x0090 */ tint_array_wrapper l;
+};
+
+kernel void tint_symbol(constant S& s [[buffer(0)]]) {
+  int3 const a = s.a;
+  int const b = s.b;
+  uint3 const c = s.c;
+  uint const d = s.d;
+  float3 const e = s.e;
+  float const f = s.f;
+  int2 const g = s.g;
+  int2 const h = s.h;
+  float2x3 const i = s.i;
+  float3x2 const j = s.j;
+  Inner const k = s.k;
+  tint_array_wrapper const l = s.l;
+  return;
+}
+

test/buffer/uniform/static_index/read.wgsl.expected.spvasm

@@ -0,0 +1,116 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 67
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "a"
+               OpMemberName %S 1 "b"
+               OpMemberName %S 2 "c"
+               OpMemberName %S 3 "d"
+               OpMemberName %S 4 "e"
+               OpMemberName %S 5 "f"
+               OpMemberName %S 6 "g"
+               OpMemberName %S 7 "h"
+               OpMemberName %S 8 "i"
+               OpMemberName %S 9 "j"
+               OpMemberName %S 10 "k"
+               OpName %Inner "Inner"
+               OpMemberName %Inner 0 "x"
+               OpMemberName %S 11 "l"
+               OpName %s "s"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 12
+               OpMemberDecorate %S 2 Offset 16
+               OpMemberDecorate %S 3 Offset 28
+               OpMemberDecorate %S 4 Offset 32
+               OpMemberDecorate %S 5 Offset 44
+               OpMemberDecorate %S 6 Offset 48
+               OpMemberDecorate %S 7 Offset 56
+               OpMemberDecorate %S 8 Offset 64
+               OpMemberDecorate %S 8 ColMajor
+               OpMemberDecorate %S 8 MatrixStride 16
+               OpMemberDecorate %S 9 Offset 96
+               OpMemberDecorate %S 9 ColMajor
+               OpMemberDecorate %S 9 MatrixStride 8
+               OpMemberDecorate %S 10 Offset 128
+               OpMemberDecorate %Inner 0 Offset 0
+               OpMemberDecorate %S 11 Offset 144
+               OpDecorate %_arr_Inner_uint_4 ArrayStride 16
+               OpDecorate %s NonWritable
+               OpDecorate %s Binding 0
+               OpDecorate %s DescriptorSet 0
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+      %v2int = OpTypeVector %int 2
+%mat2v3float = OpTypeMatrix %v3float 2
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+      %Inner = OpTypeStruct %int
+     %uint_4 = OpConstant %uint 4
+%_arr_Inner_uint_4 = OpTypeArray %Inner %uint_4
+          %S = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %v2int %v2int %mat2v3float %mat3v2float %Inner %_arr_Inner_uint_4
+%_ptr_Uniform_S = OpTypePointer Uniform %S
+          %s = OpVariable %_ptr_Uniform_S Uniform
+       %void = OpTypeVoid
+         %17 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_v3int = OpTypePointer Uniform %v3int
+     %uint_1 = OpConstant %uint 1
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+     %uint_2 = OpConstant %uint 2
+%_ptr_Uniform_v3uint = OpTypePointer Uniform %v3uint
+     %uint_3 = OpConstant %uint 3
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
+     %uint_5 = OpConstant %uint 5
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+     %uint_6 = OpConstant %uint 6
+%_ptr_Uniform_v2int = OpTypePointer Uniform %v2int
+     %uint_7 = OpConstant %uint 7
+     %uint_8 = OpConstant %uint 8
+%_ptr_Uniform_mat2v3float = OpTypePointer Uniform %mat2v3float
+     %uint_9 = OpConstant %uint 9
+%_ptr_Uniform_mat3v2float = OpTypePointer Uniform %mat3v2float
+    %uint_10 = OpConstant %uint 10
+%_ptr_Uniform_Inner = OpTypePointer Uniform %Inner
+    %uint_11 = OpConstant %uint 11
+%_ptr_Uniform__arr_Inner_uint_4 = OpTypePointer Uniform %_arr_Inner_uint_4
+       %main = OpFunction %void None %17
+         %20 = OpLabel
+         %23 = OpAccessChain %_ptr_Uniform_v3int %s %uint_0
+         %24 = OpLoad %v3int %23
+         %27 = OpAccessChain %_ptr_Uniform_int %s %uint_1
+         %28 = OpLoad %int %27
+         %31 = OpAccessChain %_ptr_Uniform_v3uint %s %uint_2
+         %32 = OpLoad %v3uint %31
+         %35 = OpAccessChain %_ptr_Uniform_uint %s %uint_3
+         %36 = OpLoad %uint %35
+         %38 = OpAccessChain %_ptr_Uniform_v3float %s %uint_4
+         %39 = OpLoad %v3float %38
+         %42 = OpAccessChain %_ptr_Uniform_float %s %uint_5
+         %43 = OpLoad %float %42
+         %46 = OpAccessChain %_ptr_Uniform_v2int %s %uint_6
+         %47 = OpLoad %v2int %46
+         %49 = OpAccessChain %_ptr_Uniform_v2int %s %uint_7
+         %50 = OpLoad %v2int %49
+         %53 = OpAccessChain %_ptr_Uniform_mat2v3float %s %uint_8
+         %54 = OpLoad %mat2v3float %53
+         %57 = OpAccessChain %_ptr_Uniform_mat3v2float %s %uint_9
+         %58 = OpLoad %mat3v2float %57
+         %61 = OpAccessChain %_ptr_Uniform_Inner %s %uint_10
+         %62 = OpLoad %Inner %61
+         %65 = OpAccessChain %_ptr_Uniform__arr_Inner_uint_4 %s %uint_11
+         %66 = OpLoad %_arr_Inner_uint_4 %65
+               OpReturn
+               OpFunctionEnd

test/buffer/uniform/static_index/read.wgsl.expected.wgsl

@@ -0,0 +1,39 @@
+struct Inner {
+  x : i32;
+};
+
+[[block]]
+struct S {
+  a : vec3<i32>;
+  b : i32;
+  c : vec3<u32>;
+  d : u32;
+  e : vec3<f32>;
+  f : f32;
+  g : vec2<i32>;
+  h : vec2<i32>;
+  i : mat2x3<f32>;
+  j : mat3x2<f32>;
+  [[align(16)]]
+  k : Inner;
+  [[align(16)]]
+  l : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<uniform> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let a = s.a;
+  let b = s.b;
+  let c = s.c;
+  let d = s.d;
+  let e = s.e;
+  let f = s.f;
+  let g = s.g;
+  let h = s.h;
+  let i = s.i;
+  let j = s.j;
+  let k = s.k;
+  let l = s.l;
+}