Allow non-struct buffer store types

For SPIR-V, wrap non-struct types in structs in the
AddSpirvBlockDecoration transform.

For MSL, wrap runtime-sized arrays in structs in the
ModuleScopeVarToEntryPointParam transform.

Bug: tint:1372
Change-Id: Icced5d77b4538e816aa9fab57a634a9f4c52fdab
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/76162
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

test/buffer/storage/types/array.wgsl

@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : array<f32, 4>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : array<f32, 4>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/array.wgsl.expected.hlsl

@@ -0,0 +1,28 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, float value[4]) {
+  float array[4] = value;
+  {
+    [loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+      buffer.Store((offset + (i * 4u)), asuint(array[i]));
+    }
+  }
+}
+
+typedef float tint_symbol_4_ret[4];
+tint_symbol_4_ret tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+  float arr[4] = (float[4])0;
+  {
+    [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+      arr[i_1] = asfloat(buffer.Load((offset + (i_1 * 4u))));
+    }
+  }
+  return arr;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_2(tint_symbol_1, 0u, tint_symbol_4(tint_symbol, 0u));
+  return;
+}

test/buffer/storage/types/array.wgsl.expected.msl

@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  /* 0x0000 */ float arr[4];
+};
+
+kernel void tint_symbol(device tint_array_wrapper* tint_symbol_1 [[buffer(1)]], const device tint_array_wrapper* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+

test/buffer/storage/types/array.wgsl.expected.spvasm

@@ -0,0 +1,42 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpDecorate %_arr_float_uint_4 ArrayStride 4
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+       %uint = OpTypeInt 32 0
+     %uint_4 = OpConstant %uint 4
+%_arr_float_uint_4 = OpTypeArray %float %uint_4
+   %in_block = OpTypeStruct %_arr_float_uint_4
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer__arr_float_uint_4 = OpTypePointer StorageBuffer %_arr_float_uint_4
+       %main = OpFunction %void None %9
+         %12 = OpLabel
+         %15 = OpAccessChain %_ptr_StorageBuffer__arr_float_uint_4 %out %uint_0
+         %16 = OpAccessChain %_ptr_StorageBuffer__arr_float_uint_4 %in %uint_0
+         %17 = OpLoad %_arr_float_uint_4 %16
+               OpStore %15 %17
+               OpReturn
+               OpFunctionEnd

test/buffer/storage/types/array.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : array<f32, 4>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : array<f32, 4>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/f32.wgsl

@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : f32;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/f32.wgsl.expected.hlsl

@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_1.Store(0u, asuint(asfloat(tint_symbol.Load(0u))));
+  return;
+}

test/buffer/storage/types/f32.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float* tint_symbol_1 [[buffer(1)]], const device float* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+

test/buffer/storage/types/f32.wgsl.expected.spvasm

@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+   %in_block = OpTypeStruct %float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %13 = OpAccessChain %_ptr_StorageBuffer_float %out %uint_0
+         %14 = OpAccessChain %_ptr_StorageBuffer_float %in %uint_0
+         %15 = OpLoad %float %14
+               OpStore %13 %15
+               OpReturn
+               OpFunctionEnd

test/buffer/storage/types/f32.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : f32;
+
+[[group(0), binding(1)]] var<storage, read_write> out : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/i32.wgsl

@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : i32;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/i32.wgsl.expected.hlsl

@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_1.Store(0u, asuint(asint(tint_symbol.Load(0u))));
+  return;
+}

test/buffer/storage/types/i32.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device int* tint_symbol_1 [[buffer(1)]], const device int* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+

test/buffer/storage/types/i32.wgsl.expected.spvasm

@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+        %int = OpTypeInt 32 1
+   %in_block = OpTypeStruct %int
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %13 = OpAccessChain %_ptr_StorageBuffer_int %out %uint_0
+         %14 = OpAccessChain %_ptr_StorageBuffer_int %in %uint_0
+         %15 = OpLoad %int %14
+               OpStore %13 %15
+               OpReturn
+               OpFunctionEnd

test/buffer/storage/types/i32.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : i32;
+
+[[group(0), binding(1)]] var<storage, read_write> out : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/mat2x2.wgsl

@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : mat2x2<f32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : mat2x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/mat2x2.wgsl.expected.hlsl

@@ -0,0 +1,17 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, float2x2 value) {
+  buffer.Store2((offset + 0u), asuint(value[0u]));
+  buffer.Store2((offset + 8u), asuint(value[1u]));
+}
+
+float2x2 tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+  return float2x2(asfloat(buffer.Load2((offset + 0u))), asfloat(buffer.Load2((offset + 8u))));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_2(tint_symbol_1, 0u, tint_symbol_4(tint_symbol, 0u));
+  return;
+}

test/buffer/storage/types/mat2x2.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float2x2* tint_symbol_1 [[buffer(1)]], const device float2x2* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+

test/buffer/storage/types/mat2x2.wgsl.expected.spvasm

@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpMemberDecorate %in_block 0 ColMajor
+               OpMemberDecorate %in_block 0 MatrixStride 8
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%mat2v2float = OpTypeMatrix %v2float 2
+   %in_block = OpTypeStruct %mat2v2float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat2v2float = OpTypePointer StorageBuffer %mat2v2float
+       %main = OpFunction %void None %8
+         %11 = OpLabel
+         %15 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %out %uint_0
+         %16 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %in %uint_0
+         %17 = OpLoad %mat2v2float %16
+               OpStore %15 %17
+               OpReturn
+               OpFunctionEnd

test/buffer/storage/types/mat2x2.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : mat2x2<f32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : mat2x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/mat2x3.wgsl

@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : mat2x3<f32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/mat2x3.wgsl.expected.hlsl

@@ -0,0 +1,17 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, float2x3 value) {
+  buffer.Store3((offset + 0u), asuint(value[0u]));
+  buffer.Store3((offset + 16u), asuint(value[1u]));
+}
+
+float2x3 tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+  return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u))));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_2(tint_symbol_1, 0u, tint_symbol_4(tint_symbol, 0u));
+  return;
+}

test/buffer/storage/types/mat2x3.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float2x3* tint_symbol_1 [[buffer(1)]], const device float2x3* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+

test/buffer/storage/types/mat2x3.wgsl.expected.spvasm

@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpMemberDecorate %in_block 0 ColMajor
+               OpMemberDecorate %in_block 0 MatrixStride 16
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+   %in_block = OpTypeStruct %mat2v3float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat2v3float = OpTypePointer StorageBuffer %mat2v3float
+       %main = OpFunction %void None %8
+         %11 = OpLabel
+         %15 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %out %uint_0
+         %16 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %in %uint_0
+         %17 = OpLoad %mat2v3float %16
+               OpStore %15 %17
+               OpReturn
+               OpFunctionEnd

test/buffer/storage/types/mat2x3.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : mat2x3<f32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/mat3x2.wgsl

@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : mat3x2<f32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : mat3x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/mat3x2.wgsl.expected.hlsl

@@ -0,0 +1,18 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(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]));
+}
+
+float3x2 tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+  return float3x2(asfloat(buffer.Load2((offset + 0u))), asfloat(buffer.Load2((offset + 8u))), asfloat(buffer.Load2((offset + 16u))));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_2(tint_symbol_1, 0u, tint_symbol_4(tint_symbol, 0u));
+  return;
+}

test/buffer/storage/types/mat3x2.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float3x2* tint_symbol_1 [[buffer(1)]], const device float3x2* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+

test/buffer/storage/types/mat3x2.wgsl.expected.spvasm

@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpMemberDecorate %in_block 0 ColMajor
+               OpMemberDecorate %in_block 0 MatrixStride 8
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+   %in_block = OpTypeStruct %mat3v2float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+       %main = OpFunction %void None %8
+         %11 = OpLabel
+         %15 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %out %uint_0
+         %16 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %in %uint_0
+         %17 = OpLoad %mat3v2float %16
+               OpStore %15 %17
+               OpReturn
+               OpFunctionEnd

test/buffer/storage/types/mat3x2.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : mat3x2<f32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : mat3x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/mat4x4.wgsl

@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : mat4x4<f32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : mat4x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/mat4x4.wgsl.expected.hlsl

@@ -0,0 +1,19 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, float4x4 value) {
+  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]));
+}
+
+float4x4 tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+  return float4x4(asfloat(buffer.Load4((offset + 0u))), asfloat(buffer.Load4((offset + 16u))), asfloat(buffer.Load4((offset + 32u))), asfloat(buffer.Load4((offset + 48u))));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_2(tint_symbol_1, 0u, tint_symbol_4(tint_symbol, 0u));
+  return;
+}

test/buffer/storage/types/mat4x4.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float4x4* tint_symbol_1 [[buffer(1)]], const device float4x4* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+

test/buffer/storage/types/mat4x4.wgsl.expected.spvasm

@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpMemberDecorate %in_block 0 ColMajor
+               OpMemberDecorate %in_block 0 MatrixStride 16
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%mat4v4float = OpTypeMatrix %v4float 4
+   %in_block = OpTypeStruct %mat4v4float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat4v4float = OpTypePointer StorageBuffer %mat4v4float
+       %main = OpFunction %void None %8
+         %11 = OpLabel
+         %15 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %out %uint_0
+         %16 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %in %uint_0
+         %17 = OpLoad %mat4v4float %16
+               OpStore %15 %17
+               OpReturn
+               OpFunctionEnd

test/buffer/storage/types/mat4x4.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : mat4x4<f32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : mat4x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/runtime_array.wgsl

@@ -0,0 +1,14 @@
+struct S {
+  f : f32;
+};
+
+[[group(0), binding(0)]]
+var<storage, read> in : array<S>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : array<S>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out[0] = in[0];
+}

test/buffer/storage/types/runtime_array.wgsl.expected.hlsl

@@ -0,0 +1,21 @@
+struct S {
+  float f;
+};
+
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, S value) {
+  buffer.Store((offset + 0u), asuint(value.f));
+}
+
+S tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+  const S tint_symbol_6 = {asfloat(buffer.Load((offset + 0u)))};
+  return tint_symbol_6;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_2(tint_symbol_1, (4u * uint(0)), tint_symbol_4(tint_symbol, (4u * uint(0))));
+  return;
+}

test/buffer/storage/types/runtime_array.wgsl.expected.msl

@@ -0,0 +1,18 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  /* 0x0000 */ float f;
+};
+struct tint_symbol_2 {
+  /* 0x0000 */ S arr[1];
+};
+struct tint_symbol_4 {
+  /* 0x0000 */ S arr[1];
+};
+
+kernel void tint_symbol(device tint_symbol_2* tint_symbol_1 [[buffer(1)]], const device tint_symbol_4* tint_symbol_3 [[buffer(0)]]) {
+  (*(tint_symbol_1)).arr[0] = (*(tint_symbol_3)).arr[0];
+  return;
+}
+

test/buffer/storage/types/runtime_array.wgsl.expected.spvasm

@@ -0,0 +1,47 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 20
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %S "S"
+               OpMemberName %S 0 "f"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_runtimearr_S ArrayStride 4
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+          %S = OpTypeStruct %float
+%_runtimearr_S = OpTypeRuntimeArray %S
+   %in_block = OpTypeStruct %_runtimearr_S
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+       %main = OpFunction %void None %8
+         %11 = OpLabel
+         %17 = OpAccessChain %_ptr_StorageBuffer_S %out %uint_0 %int_0
+         %18 = OpAccessChain %_ptr_StorageBuffer_S %in %uint_0 %int_0
+         %19 = OpLoad %S %18
+               OpStore %17 %19
+               OpReturn
+               OpFunctionEnd

test/buffer/storage/types/runtime_array.wgsl.expected.wgsl

@@ -0,0 +1,12 @@
+struct S {
+  f : f32;
+};
+
+[[group(0), binding(0)]] var<storage, read> in : array<S>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : array<S>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out[0] = in[0];
+}

test/buffer/storage/types/struct.wgsl

@@ -0,0 +1,17 @@
+struct Inner {
+  f : f32;
+};
+struct S {
+  inner : Inner;
+};
+
+[[group(0), binding(0)]]
+var<storage, read> in : S;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/struct.wgsl.expected.hlsl

@@ -0,0 +1,33 @@
+struct Inner {
+  float f;
+};
+struct S {
+  Inner inner;
+};
+
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_3(RWByteAddressBuffer buffer, uint offset, Inner value) {
+  buffer.Store((offset + 0u), asuint(value.f));
+}
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, S value) {
+  tint_symbol_3(buffer, (offset + 0u), value.inner);
+}
+
+Inner tint_symbol_6(ByteAddressBuffer buffer, uint offset) {
+  const Inner tint_symbol_8 = {asfloat(buffer.Load((offset + 0u)))};
+  return tint_symbol_8;
+}
+
+S tint_symbol_5(ByteAddressBuffer buffer, uint offset) {
+  const S tint_symbol_9 = {tint_symbol_6(buffer, (offset + 0u))};
+  return tint_symbol_9;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_2(tint_symbol_1, 0u, tint_symbol_5(tint_symbol, 0u));
+  return;
+}

test/buffer/storage/types/struct.wgsl.expected.msl

@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Inner {
+  /* 0x0000 */ float f;
+};
+struct S {
+  /* 0x0000 */ Inner inner;
+};
+
+kernel void tint_symbol(device S* tint_symbol_1 [[buffer(0)]], const device S* tint_symbol_2 [[buffer(1)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+

test/buffer/storage/types/struct.wgsl.expected.spvasm

@@ -0,0 +1,38 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "inner"
+               OpName %Inner "Inner"
+               OpMemberName %Inner 0 "f"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %Inner 0 Offset 0
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+      %Inner = OpTypeStruct %float
+          %S = OpTypeStruct %Inner
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+         %in = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %11 = OpLoad %S %in
+               OpStore %out %11
+               OpReturn
+               OpFunctionEnd

test/buffer/storage/types/struct.wgsl.expected.wgsl

@@ -0,0 +1,16 @@
+struct Inner {
+  f : f32;
+};
+
+struct S {
+  inner : Inner;
+};
+
+[[group(0), binding(0)]] var<storage, read> in : S;
+
+[[group(0), binding(1)]] var<storage, read_write> out : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/u32.wgsl

@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : u32;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : u32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/u32.wgsl.expected.hlsl

@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_1.Store(0u, asuint(tint_symbol.Load(0u)));
+  return;
+}

test/buffer/storage/types/u32.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device uint* tint_symbol_1 [[buffer(1)]], const device uint* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+

test/buffer/storage/types/u32.wgsl.expected.spvasm

@@ -0,0 +1,38 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 15
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+       %uint = OpTypeInt 32 0
+   %in_block = OpTypeStruct %uint
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %12 = OpAccessChain %_ptr_StorageBuffer_uint %out %uint_0
+         %13 = OpAccessChain %_ptr_StorageBuffer_uint %in %uint_0
+         %14 = OpLoad %uint %13
+               OpStore %12 %14
+               OpReturn
+               OpFunctionEnd

test/buffer/storage/types/u32.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : u32;
+
+[[group(0), binding(1)]] var<storage, read_write> out : u32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/vec2.wgsl

@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : vec2<i32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : vec2<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/vec2.wgsl.expected.hlsl

@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_1.Store2(0u, asuint(asint(tint_symbol.Load2(0u))));
+  return;
+}

test/buffer/storage/types/vec2.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device int2* tint_symbol_1 [[buffer(1)]], const device int2* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+

test/buffer/storage/types/vec2.wgsl.expected.spvasm

@@ -0,0 +1,40 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+   %in_block = OpTypeStruct %v2int
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v2int = OpTypePointer StorageBuffer %v2int
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_StorageBuffer_v2int %out %uint_0
+         %15 = OpAccessChain %_ptr_StorageBuffer_v2int %in %uint_0
+         %16 = OpLoad %v2int %15
+               OpStore %14 %16
+               OpReturn
+               OpFunctionEnd

test/buffer/storage/types/vec2.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : vec2<i32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : vec2<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/vec3.wgsl

@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : vec3<u32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : vec3<u32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/vec3.wgsl.expected.hlsl

@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_1.Store3(0u, asuint(tint_symbol.Load3(0u)));
+  return;
+}

test/buffer/storage/types/vec3.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device uint3* tint_symbol_1 [[buffer(1)]], const device uint3* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+

test/buffer/storage/types/vec3.wgsl.expected.spvasm

@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+   %in_block = OpTypeStruct %v3uint
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %13 = OpAccessChain %_ptr_StorageBuffer_v3uint %out %uint_0
+         %14 = OpAccessChain %_ptr_StorageBuffer_v3uint %in %uint_0
+         %15 = OpLoad %v3uint %14
+               OpStore %13 %15
+               OpReturn
+               OpFunctionEnd

test/buffer/storage/types/vec3.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : vec3<u32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : vec3<u32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/vec4.wgsl

@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : vec4<f32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : vec4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/storage/types/vec4.wgsl.expected.hlsl

@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_1.Store4(0u, asuint(asfloat(tint_symbol.Load4(0u))));
+  return;
+}

test/buffer/storage/types/vec4.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float4* tint_symbol_1 [[buffer(1)]], const device float4* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+

test/buffer/storage/types/vec4.wgsl.expected.spvasm

@@ -0,0 +1,40 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+   %in_block = OpTypeStruct %v4float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v4float = OpTypePointer StorageBuffer %v4float
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_StorageBuffer_v4float %out %uint_0
+         %15 = OpAccessChain %_ptr_StorageBuffer_v4float %in %uint_0
+         %16 = OpLoad %v4float %15
+               OpStore %14 %16
+               OpReturn
+               OpFunctionEnd

test/buffer/storage/types/vec4.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : vec4<f32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : vec4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}

test/buffer/uniform/types/array.wgsl

@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : array<vec4<f32>, 4>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}

test/buffer/uniform/types/array.wgsl.expected.hlsl

@@ -0,0 +1,21 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[4];
+};
+
+typedef float4 tint_symbol_ret[4];
+tint_symbol_ret tint_symbol(uint4 buffer[4], uint offset) {
+  float4 arr[4] = (float4[4])0;
+  {
+    [loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+      const uint scalar_offset = ((offset + (i * 16u))) / 4;
+      arr[i] = asfloat(buffer[scalar_offset / 4]);
+    }
+  }
+  return arr;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  const float4 x[4] = tint_symbol(u, 0u);
+  return;
+}

test/buffer/uniform/types/array.wgsl.expected.msl

@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  /* 0x0000 */ float4 arr[4];
+};
+
+kernel void tint_symbol(const constant tint_array_wrapper* tint_symbol_1 [[buffer(0)]]) {
+  tint_array_wrapper const x = *(tint_symbol_1);
+  return;
+}
+

test/buffer/uniform/types/array.wgsl.expected.spvasm

@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpDecorate %_arr_v4float_uint_4 ArrayStride 16
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+       %uint = OpTypeInt 32 0
+     %uint_4 = OpConstant %uint 4
+%_arr_v4float_uint_4 = OpTypeArray %v4float %uint_4
+    %u_block = OpTypeStruct %_arr_v4float_uint_4
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform__arr_v4float_uint_4 = OpTypePointer Uniform %_arr_v4float_uint_4
+       %main = OpFunction %void None %9
+         %12 = OpLabel
+         %15 = OpAccessChain %_ptr_Uniform__arr_v4float_uint_4 %u %uint_0
+         %16 = OpLoad %_arr_v4float_uint_4 %15
+               OpReturn
+               OpFunctionEnd

test/buffer/uniform/types/array.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : array<vec4<f32>, 4>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}

test/buffer/uniform/types/f32.wgsl

@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}

test/buffer/uniform/types/f32.wgsl.expected.hlsl

@@ -0,0 +1,9 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  const float x = asfloat(u[0].x);
+  return;
+}

test/buffer/uniform/types/f32.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float* tint_symbol_1 [[buffer(0)]]) {
+  float const x = *(tint_symbol_1);
+  return;
+}
+

test/buffer/uniform/types/f32.wgsl.expected.spvasm

@@ -0,0 +1,33 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+      %float = OpTypeFloat 32
+    %u_block = OpTypeStruct %float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+         %12 = OpAccessChain %_ptr_Uniform_float %u %uint_0
+         %13 = OpLoad %float %12
+               OpReturn
+               OpFunctionEnd

test/buffer/uniform/types/f32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}

test/buffer/uniform/types/i32.wgsl

@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}

test/buffer/uniform/types/i32.wgsl.expected.hlsl

@@ -0,0 +1,9 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  const int x = asint(u[0].x);
+  return;
+}

test/buffer/uniform/types/i32.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant int* tint_symbol_1 [[buffer(0)]]) {
+  int const x = *(tint_symbol_1);
+  return;
+}
+

test/buffer/uniform/types/i32.wgsl.expected.spvasm

@@ -0,0 +1,33 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+        %int = OpTypeInt 32 1
+    %u_block = OpTypeStruct %int
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+         %12 = OpAccessChain %_ptr_Uniform_int %u %uint_0
+         %13 = OpLoad %int %12
+               OpReturn
+               OpFunctionEnd

test/buffer/uniform/types/i32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}

test/buffer/uniform/types/mat2x2.wgsl

@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : mat2x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}

test/buffer/uniform/types/mat2x2.wgsl.expected.hlsl

@@ -0,0 +1,17 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[1];
+};
+
+float2x2 tint_symbol(uint4 buffer[1], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  uint4 ubo_load = buffer[scalar_offset / 4];
+  const uint scalar_offset_1 = ((offset + 8u)) / 4;
+  uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
+  return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  const float2x2 x = tint_symbol(u, 0u);
+  return;
+}

test/buffer/uniform/types/mat2x2.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float2x2* tint_symbol_1 [[buffer(0)]]) {
+  float2x2 const x = *(tint_symbol_1);
+  return;
+}
+

test/buffer/uniform/types/mat2x2.wgsl.expected.spvasm

@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpMemberDecorate %u_block 0 ColMajor
+               OpMemberDecorate %u_block 0 MatrixStride 8
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%mat2v2float = OpTypeMatrix %v2float 2
+    %u_block = OpTypeStruct %mat2v2float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat2v2float = OpTypePointer Uniform %mat2v2float
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_Uniform_mat2v2float %u %uint_0
+         %15 = OpLoad %mat2v2float %14
+               OpReturn
+               OpFunctionEnd

test/buffer/uniform/types/mat2x2.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : mat2x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}

test/buffer/uniform/types/mat2x3.wgsl

@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}

test/buffer/uniform/types/mat2x3.wgsl.expected.hlsl

@@ -0,0 +1,15 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[2];
+};
+
+float2x3 tint_symbol(uint4 buffer[2], 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));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  const float2x3 x = tint_symbol(u, 0u);
+  return;
+}

test/buffer/uniform/types/mat2x3.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float2x3* tint_symbol_1 [[buffer(0)]]) {
+  float2x3 const x = *(tint_symbol_1);
+  return;
+}
+

test/buffer/uniform/types/mat2x3.wgsl.expected.spvasm

@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpMemberDecorate %u_block 0 ColMajor
+               OpMemberDecorate %u_block 0 MatrixStride 16
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+    %u_block = OpTypeStruct %mat2v3float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat2v3float = OpTypePointer Uniform %mat2v3float
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_Uniform_mat2v3float %u %uint_0
+         %15 = OpLoad %mat2v3float %14
+               OpReturn
+               OpFunctionEnd

test/buffer/uniform/types/mat2x3.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}

test/buffer/uniform/types/mat3x2.wgsl

@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : mat3x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}

test/buffer/uniform/types/mat3x2.wgsl.expected.hlsl

@@ -0,0 +1,19 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[2];
+};
+
+float3x2 tint_symbol(uint4 buffer[2], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  uint4 ubo_load = buffer[scalar_offset / 4];
+  const uint scalar_offset_1 = ((offset + 8u)) / 4;
+  uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
+  const uint scalar_offset_2 = ((offset + 16u)) / 4;
+  uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
+  return float3x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  const float3x2 x = tint_symbol(u, 0u);
+  return;
+}

test/buffer/uniform/types/mat3x2.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float3x2* tint_symbol_1 [[buffer(0)]]) {
+  float3x2 const x = *(tint_symbol_1);
+  return;
+}
+

test/buffer/uniform/types/mat3x2.wgsl.expected.spvasm

@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpMemberDecorate %u_block 0 ColMajor
+               OpMemberDecorate %u_block 0 MatrixStride 8
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+    %u_block = OpTypeStruct %mat3v2float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat3v2float = OpTypePointer Uniform %mat3v2float
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_Uniform_mat3v2float %u %uint_0
+         %15 = OpLoad %mat3v2float %14
+               OpReturn
+               OpFunctionEnd

test/buffer/uniform/types/mat3x2.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : mat3x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}

test/buffer/uniform/types/mat4x4.wgsl

@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : mat4x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}

test/buffer/uniform/types/mat4x4.wgsl.expected.hlsl

@@ -0,0 +1,17 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[4];
+};
+
+float4x4 tint_symbol(uint4 buffer[4], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  const uint scalar_offset_2 = ((offset + 32u)) / 4;
+  const uint scalar_offset_3 = ((offset + 48u)) / 4;
+  return float4x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]), asfloat(buffer[scalar_offset_3 / 4]));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  const float4x4 x = tint_symbol(u, 0u);
+  return;
+}

test/buffer/uniform/types/mat4x4.wgsl.expected.msl

@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float4x4* tint_symbol_1 [[buffer(0)]]) {
+  float4x4 const x = *(tint_symbol_1);
+  return;
+}
+

test/buffer/uniform/types/mat4x4.wgsl.expected.spvasm

@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpMemberDecorate %u_block 0 ColMajor
+               OpMemberDecorate %u_block 0 MatrixStride 16
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%mat4v4float = OpTypeMatrix %v4float 4
+    %u_block = OpTypeStruct %mat4v4float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat4v4float = OpTypePointer Uniform %mat4v4float
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_Uniform_mat4v4float %u %uint_0
+         %15 = OpLoad %mat4v4float %14
+               OpReturn
+               OpFunctionEnd

test/buffer/uniform/types/mat4x4.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : mat4x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}