reader/spirv: Decompose arrays with strides
Transform any SPIR-V that has an array with a custom stride:
@stride(S) array<T, N>
into:
struct strided_arr {
@size(S) er : T;
};
array<strided_arr, N>
Also remove any @stride decorations that match the default array stride.
Bug: tint:1394
Bug: tint:1381
Change-Id: I8be8f3a76c5335fdb2bc5183388366091dbc7642
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/78781
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/test/BUILD.gn b/test/BUILD.gn
index 538c4a9..2eb530a 100644
--- a/test/BUILD.gn
+++ b/test/BUILD.gn
@@ -310,6 +310,7 @@
"../src/transform/canonicalize_entry_point_io_test.cc",
"../src/transform/combine_samplers_test.cc",
"../src/transform/decompose_memory_access_test.cc",
+ "../src/transform/decompose_strided_array_test.cc",
"../src/transform/decompose_strided_matrix_test.cc",
"../src/transform/external_texture_transform_test.cc",
"../src/transform/first_index_offset_test.cc",
diff --git a/test/array/strides.spvasm b/test/array/strides.spvasm
new file mode 100644
index 0000000..68fc329
--- /dev/null
+++ b/test/array/strides.spvasm
@@ -0,0 +1,71 @@
+; type ARR_A = @stride(8) array<f32, 2>;
+; type ARR_B = @stride(128) array<@stride(16) array<ARR_A, 4>, 3>;
+; struct S {
+; a : ARR_B;
+; };
+; @group(0) @binding(0) var<storage, read_write> s : S;
+;
+; @stage(compute) @workgroup_size(1)
+; fn f() {
+; let a : ARR_B = s.a;
+; let b : array<@stride(8) array<f32, 2>, 3> = s.a[3];
+; let c = s.a[3][2];
+; let d = s.a[3][2][1];
+; s.a = ARR_B();
+; s.a[3][2][1] = 5.0;
+; }
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %f "f"
+ OpExecutionMode %f LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "a"
+ OpName %s "s"
+ OpName %f "f"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpDecorate %_arr_float_uint_2 ArrayStride 8
+ OpDecorate %_arr__arr_float_uint_2_uint_3 ArrayStride 16
+ OpDecorate %_arr__arr__arr_float_uint_2_uint_3_uint_4 ArrayStride 128
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 0
+ %float = OpTypeFloat 32
+ %uint = OpTypeInt 32 0
+ %uint_2 = OpConstant %uint 2
+ %_arr_float_uint_2 = OpTypeArray %float %uint_2
+ %uint_3 = OpConstant %uint 3
+ %_arr__arr_float_uint_2_uint_3 = OpTypeArray %_arr_float_uint_2 %uint_3
+ %uint_4 = OpConstant %uint 4
+ %_arr__arr__arr_float_uint_2_uint_3_uint_4 = OpTypeArray %_arr__arr_float_uint_2_uint_3 %uint_4
+ %S = OpTypeStruct %_arr__arr__arr_float_uint_2_uint_3_uint_4
+ %_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+ %void = OpTypeVoid
+ %12 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer__arr__arr__arr_float_uint_2_uint_3_uint_4 = OpTypePointer StorageBuffer %_arr__arr__arr_float_uint_2_uint_3_uint_4
+ %int = OpTypeInt 32 1
+ %int_3 = OpConstant %int 3
+ %_ptr_StorageBuffer__arr__arr_float_uint_2_uint_3 = OpTypePointer StorageBuffer %_arr__arr_float_uint_2_uint_3
+ %int_2 = OpConstant %int 2
+ %_ptr_StorageBuffer__arr_float_uint_2 = OpTypePointer StorageBuffer %_arr_float_uint_2
+ %int_1 = OpConstant %int 1
+ %_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %34 = OpConstantNull %_arr__arr__arr_float_uint_2_uint_3_uint_4
+ %float_5 = OpConstant %float 5
+ %f = OpFunction %void None %12
+ %15 = OpLabel
+ %18 = OpAccessChain %_ptr_StorageBuffer__arr__arr__arr_float_uint_2_uint_3_uint_4 %s %uint_0
+ %19 = OpLoad %_arr__arr__arr_float_uint_2_uint_3_uint_4 %18
+ %23 = OpAccessChain %_ptr_StorageBuffer__arr__arr_float_uint_2_uint_3 %s %uint_0 %int_3
+ %24 = OpLoad %_arr__arr_float_uint_2_uint_3 %23
+ %27 = OpAccessChain %_ptr_StorageBuffer__arr_float_uint_2 %s %uint_0 %int_3 %int_2
+ %28 = OpLoad %_arr_float_uint_2 %27
+ %31 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %int_3 %int_2 %int_1
+ %32 = OpLoad %float %31
+ %33 = OpAccessChain %_ptr_StorageBuffer__arr__arr__arr_float_uint_2_uint_3_uint_4 %s %uint_0
+ OpStore %33 %34
+ %35 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %int_3 %int_2 %int_1
+ OpStore %35 %float_5
+ OpReturn
+ OpFunctionEnd
diff --git a/test/array/strides.spvasm.expected.glsl b/test/array/strides.spvasm.expected.glsl
new file mode 100644
index 0000000..21fc3e8
--- /dev/null
+++ b/test/array/strides.spvasm.expected.glsl
@@ -0,0 +1,38 @@
+#version 310 es
+precision mediump float;
+
+struct strided_arr {
+ float el;
+};
+
+struct strided_arr_1 {
+ strided_arr el[3][2];
+};
+
+struct S {
+ strided_arr_1 a[4];
+};
+
+layout(binding = 0) buffer S_1 {
+ strided_arr_1 a[4];
+} s;
+void f_1() {
+ strided_arr_1 x_19[4] = s.a;
+ strided_arr x_24[3][2] = s.a[3].el;
+ strided_arr x_28[2] = s.a[3].el[2];
+ float x_32 = s.a[3].el[2][1].el;
+ strided_arr_1 tint_symbol[4] = strided_arr_1[4](strided_arr_1(strided_arr[3][2](strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)))), strided_arr_1(strided_arr[3][2](strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)))), strided_arr_1(strided_arr[3][2](strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)))), strided_arr_1(strided_arr[3][2](strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)))));
+ s.a = tint_symbol;
+ s.a[3].el[2][1].el = 5.0f;
+ return;
+}
+
+void f() {
+ f_1();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ f();
+ return;
+}
diff --git a/test/array/strides.spvasm.expected.hlsl b/test/array/strides.spvasm.expected.hlsl
new file mode 100644
index 0000000..910a49a
--- /dev/null
+++ b/test/array/strides.spvasm.expected.hlsl
@@ -0,0 +1,103 @@
+struct strided_arr {
+ float el;
+};
+struct strided_arr_1 {
+ strided_arr el[3][2];
+};
+
+RWByteAddressBuffer s : register(u0, space0);
+
+strided_arr tint_symbol_4(RWByteAddressBuffer buffer, uint offset) {
+ const strided_arr tint_symbol_12 = {asfloat(buffer.Load((offset + 0u)))};
+ return tint_symbol_12;
+}
+
+typedef strided_arr tint_symbol_3_ret[2];
+tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
+ strided_arr arr[2] = (strided_arr[2])0;
+ {
+ [loop] for(uint i = 0u; (i < 2u); i = (i + 1u)) {
+ arr[i] = tint_symbol_4(buffer, (offset + (i * 8u)));
+ }
+ }
+ return arr;
+}
+
+typedef strided_arr tint_symbol_2_ret[3][2];
+tint_symbol_2_ret tint_symbol_2(RWByteAddressBuffer buffer, uint offset) {
+ strided_arr arr_1[3][2] = (strided_arr[3][2])0;
+ {
+ [loop] for(uint i_1 = 0u; (i_1 < 3u); i_1 = (i_1 + 1u)) {
+ arr_1[i_1] = tint_symbol_3(buffer, (offset + (i_1 * 16u)));
+ }
+ }
+ return arr_1;
+}
+
+strided_arr_1 tint_symbol_1(RWByteAddressBuffer buffer, uint offset) {
+ const strided_arr_1 tint_symbol_13 = {tint_symbol_2(buffer, (offset + 0u))};
+ return tint_symbol_13;
+}
+
+typedef strided_arr_1 tint_symbol_ret[4];
+tint_symbol_ret tint_symbol(RWByteAddressBuffer buffer, uint offset) {
+ strided_arr_1 arr_2[4] = (strided_arr_1[4])0;
+ {
+ [loop] for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
+ arr_2[i_2] = tint_symbol_1(buffer, (offset + (i_2 * 128u)));
+ }
+ }
+ return arr_2;
+}
+
+void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, strided_arr value) {
+ buffer.Store((offset + 0u), asuint(value.el));
+}
+
+void tint_symbol_9(RWByteAddressBuffer buffer, uint offset, strided_arr value[2]) {
+ strided_arr array_2[2] = value;
+ {
+ [loop] for(uint i_3 = 0u; (i_3 < 2u); i_3 = (i_3 + 1u)) {
+ tint_symbol_10(buffer, (offset + (i_3 * 8u)), array_2[i_3]);
+ }
+ }
+}
+
+void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, strided_arr value[3][2]) {
+ strided_arr array_1[3][2] = value;
+ {
+ [loop] for(uint i_4 = 0u; (i_4 < 3u); i_4 = (i_4 + 1u)) {
+ tint_symbol_9(buffer, (offset + (i_4 * 16u)), array_1[i_4]);
+ }
+ }
+}
+
+void tint_symbol_7(RWByteAddressBuffer buffer, uint offset, strided_arr_1 value) {
+ tint_symbol_8(buffer, (offset + 0u), value.el);
+}
+
+void tint_symbol_6(RWByteAddressBuffer buffer, uint offset, strided_arr_1 value[4]) {
+ strided_arr_1 array[4] = value;
+ {
+ [loop] for(uint i_5 = 0u; (i_5 < 4u); i_5 = (i_5 + 1u)) {
+ tint_symbol_7(buffer, (offset + (i_5 * 128u)), array[i_5]);
+ }
+ }
+}
+
+void f_1() {
+ const strided_arr_1 x_19[4] = tint_symbol(s, 0u);
+ const strided_arr x_24[3][2] = tint_symbol_2(s, 384u);
+ const strided_arr x_28[2] = tint_symbol_3(s, 416u);
+ const float x_32 = asfloat(s.Load(424u));
+ const strided_arr_1 tint_symbol_14[4] = (strided_arr_1[4])0;
+ tint_symbol_6(s, 0u, tint_symbol_14);
+ s.Store(424u, asuint(5.0f));
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void f() {
+ f_1();
+ return;
+}
diff --git a/test/array/strides.spvasm.expected.msl b/test/array/strides.spvasm.expected.msl
new file mode 100644
index 0000000..62f817c
--- /dev/null
+++ b/test/array/strides.spvasm.expected.msl
@@ -0,0 +1,40 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct strided_arr {
+ /* 0x0000 */ float el;
+ /* 0x0004 */ int8_t tint_pad[4];
+};
+struct tint_array_wrapper {
+ /* 0x0000 */ strided_arr arr[2];
+};
+struct tint_array_wrapper_1 {
+ /* 0x0000 */ tint_array_wrapper arr[3];
+};
+struct strided_arr_1 {
+ /* 0x0000 */ tint_array_wrapper_1 el;
+ /* 0x0030 */ int8_t tint_pad_1[80];
+};
+struct tint_array_wrapper_2 {
+ /* 0x0000 */ strided_arr_1 arr[4];
+};
+struct S {
+ /* 0x0000 */ tint_array_wrapper_2 a;
+};
+
+void f_1(device S* const tint_symbol_1) {
+ tint_array_wrapper_2 const x_19 = (*(tint_symbol_1)).a;
+ tint_array_wrapper_1 const x_24 = (*(tint_symbol_1)).a.arr[3].el;
+ tint_array_wrapper const x_28 = (*(tint_symbol_1)).a.arr[3].el.arr[2];
+ float const x_32 = (*(tint_symbol_1)).a.arr[3].el.arr[2].arr[1].el;
+ tint_array_wrapper_2 const tint_symbol = {.arr={}};
+ (*(tint_symbol_1)).a = tint_symbol;
+ (*(tint_symbol_1)).a.arr[3].el.arr[2].arr[1].el = 5.0f;
+ return;
+}
+
+kernel void f(device S* tint_symbol_2 [[buffer(0)]]) {
+ f_1(tint_symbol_2);
+ return;
+}
+
diff --git a/test/array/strides.spvasm.expected.spvasm b/test/array/strides.spvasm.expected.spvasm
new file mode 100644
index 0000000..3108b19
--- /dev/null
+++ b/test/array/strides.spvasm.expected.spvasm
@@ -0,0 +1,74 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 42
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %f "f"
+ OpExecutionMode %f LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "a"
+ OpName %strided_arr_1 "strided_arr_1"
+ OpMemberName %strided_arr_1 0 "el"
+ OpName %strided_arr "strided_arr"
+ OpMemberName %strided_arr 0 "el"
+ OpName %s "s"
+ OpName %f_1 "f_1"
+ OpName %f "f"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %strided_arr_1 0 Offset 0
+ OpMemberDecorate %strided_arr 0 Offset 0
+ OpDecorate %_arr_strided_arr_uint_2 ArrayStride 8
+ OpDecorate %_arr__arr_strided_arr_uint_2_uint_3 ArrayStride 16
+ OpDecorate %_arr_strided_arr_1_uint_4 ArrayStride 128
+ OpDecorate %s DescriptorSet 0
+ OpDecorate %s Binding 0
+ %float = OpTypeFloat 32
+%strided_arr = OpTypeStruct %float
+ %uint = OpTypeInt 32 0
+ %uint_2 = OpConstant %uint 2
+%_arr_strided_arr_uint_2 = OpTypeArray %strided_arr %uint_2
+ %uint_3 = OpConstant %uint 3
+%_arr__arr_strided_arr_uint_2_uint_3 = OpTypeArray %_arr_strided_arr_uint_2 %uint_3
+%strided_arr_1 = OpTypeStruct %_arr__arr_strided_arr_uint_2_uint_3
+ %uint_4 = OpConstant %uint 4
+%_arr_strided_arr_1_uint_4 = OpTypeArray %strided_arr_1 %uint_4
+ %S = OpTypeStruct %_arr_strided_arr_1_uint_4
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+ %void = OpTypeVoid
+ %14 = OpTypeFunction %void
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer__arr_strided_arr_1_uint_4 = OpTypePointer StorageBuffer %_arr_strided_arr_1_uint_4
+ %int = OpTypeInt 32 1
+ %int_3 = OpConstant %int 3
+%_ptr_StorageBuffer__arr__arr_strided_arr_uint_2_uint_3 = OpTypePointer StorageBuffer %_arr__arr_strided_arr_uint_2_uint_3
+ %int_2 = OpConstant %int 2
+%_ptr_StorageBuffer__arr_strided_arr_uint_2 = OpTypePointer StorageBuffer %_arr_strided_arr_uint_2
+ %int_1 = OpConstant %int 1
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+ %36 = OpConstantNull %_arr_strided_arr_1_uint_4
+ %float_5 = OpConstant %float 5
+ %f_1 = OpFunction %void None %14
+ %17 = OpLabel
+ %20 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_1_uint_4 %s %uint_0
+ %21 = OpLoad %_arr_strided_arr_1_uint_4 %20
+ %25 = OpAccessChain %_ptr_StorageBuffer__arr__arr_strided_arr_uint_2_uint_3 %s %uint_0 %int_3 %uint_0
+ %26 = OpLoad %_arr__arr_strided_arr_uint_2_uint_3 %25
+ %29 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_uint_2 %s %uint_0 %int_3 %uint_0 %int_2
+ %30 = OpLoad %_arr_strided_arr_uint_2 %29
+ %33 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %int_3 %uint_0 %int_2 %int_1 %uint_0
+ %34 = OpLoad %float %33
+ %35 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_1_uint_4 %s %uint_0
+ OpStore %35 %36
+ %37 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %int_3 %uint_0 %int_2 %int_1 %uint_0
+ OpStore %37 %float_5
+ OpReturn
+ OpFunctionEnd
+ %f = OpFunction %void None %14
+ %40 = OpLabel
+ %41 = OpFunctionCall %void %f_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/array/strides.spvasm.expected.wgsl b/test/array/strides.spvasm.expected.wgsl
new file mode 100644
index 0000000..799ce0d
--- /dev/null
+++ b/test/array/strides.spvasm.expected.wgsl
@@ -0,0 +1,36 @@
+struct strided_arr {
+ @size(8)
+ el : f32;
+}
+
+type Arr = array<strided_arr, 2u>;
+
+type Arr_1 = array<Arr, 3u>;
+
+struct strided_arr_1 {
+ @size(128)
+ el : Arr_1;
+}
+
+type Arr_2 = array<strided_arr_1, 4u>;
+
+struct S {
+ a : Arr_2;
+}
+
+@group(0) @binding(0) var<storage, read_write> s : S;
+
+fn f_1() {
+ let x_19 : Arr_2 = s.a;
+ let x_24 : Arr_1 = s.a[3].el;
+ let x_28 : Arr = s.a[3].el[2];
+ let x_32 : f32 = s.a[3].el[2][1].el;
+ s.a = array<strided_arr_1, 4u>();
+ s.a[3].el[2][1].el = 5.0;
+ return;
+}
+
+@stage(compute) @workgroup_size(1, 1, 1)
+fn f() {
+ f_1();
+}
diff --git a/test/bug/tint/1088.spvasm.expected.glsl b/test/bug/tint/1088.spvasm.expected.glsl
index a94d391..dea6c16 100644
--- a/test/bug/tint/1088.spvasm.expected.glsl
+++ b/test/bug/tint/1088.spvasm.expected.glsl
@@ -5,7 +5,7 @@
layout(location = 2) in vec2 uv_param_1;
layout(location = 1) in vec3 normal_param_1;
layout(location = 0) out vec2 vUV_1_1;
-struct tint_padded_array_element {
+struct strided_arr {
float el;
};
@@ -13,7 +13,7 @@
mat4 worldViewProjection;
float time;
mat4 test2[2];
- tint_padded_array_element test[4];
+ strided_arr test[4];
};
vec3 position = vec3(0.0f, 0.0f, 0.0f);
@@ -21,7 +21,7 @@
mat4 worldViewProjection;
float time;
mat4 test2[2];
- tint_padded_array_element test[4];
+ strided_arr test[4];
} x_14;
vec2 vUV = vec2(0.0f, 0.0f);
diff --git a/test/bug/tint/1088.spvasm.expected.msl b/test/bug/tint/1088.spvasm.expected.msl
index c0dd650..1f074d5 100644
--- a/test/bug/tint/1088.spvasm.expected.msl
+++ b/test/bug/tint/1088.spvasm.expected.msl
@@ -4,12 +4,12 @@
struct tint_array_wrapper {
/* 0x0000 */ float4x4 arr[2];
};
-struct tint_padded_array_element {
+struct strided_arr {
/* 0x0000 */ float el;
/* 0x0004 */ int8_t tint_pad[12];
};
struct tint_array_wrapper_1 {
- /* 0x0000 */ tint_padded_array_element arr[4];
+ /* 0x0000 */ strided_arr arr[4];
};
struct LeftOver {
/* 0x0000 */ float4x4 worldViewProjection;
diff --git a/test/bug/tint/1088.spvasm.expected.spvasm b/test/bug/tint/1088.spvasm.expected.spvasm
index 721a011..f157cd1 100644
--- a/test/bug/tint/1088.spvasm.expected.spvasm
+++ b/test/bug/tint/1088.spvasm.expected.spvasm
@@ -1,10 +1,10 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 121
+; Bound: 122
; Schema: 0
OpCapability Shader
- %74 = OpExtInstImport "GLSL.std.450"
+ %75 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %main "main" %position_param_1 %uv_param_1 %normal_param_1 %gl_Position_1 %vUV_1_1 %vertex_point_size
OpName %position_param_1 "position_param_1"
@@ -19,6 +19,8 @@
OpMemberName %LeftOver 1 "time"
OpMemberName %LeftOver 2 "test2"
OpMemberName %LeftOver 3 "test"
+ OpName %strided_arr "strided_arr"
+ OpMemberName %strided_arr 0 "el"
OpName %x_14 "x_14"
OpName %vUV "vUV"
OpName %uv "uv"
@@ -51,7 +53,8 @@
OpMemberDecorate %LeftOver 2 MatrixStride 16
OpDecorate %_arr_mat4v4float_uint_2 ArrayStride 64
OpMemberDecorate %LeftOver 3 Offset 208
- OpDecorate %_arr_float_uint_4 ArrayStride 16
+ OpMemberDecorate %strided_arr 0 Offset 0
+ OpDecorate %_arr_strided_arr_uint_4 ArrayStride 16
OpDecorate %x_14 NonWritable
OpDecorate %x_14 DescriptorSet 2
OpDecorate %x_14 Binding 2
@@ -82,9 +85,10 @@
%uint = OpTypeInt 32 0
%uint_2 = OpConstant %uint 2
%_arr_mat4v4float_uint_2 = OpTypeArray %mat4v4float %uint_2
+%strided_arr = OpTypeStruct %float
%uint_4 = OpConstant %uint 4
-%_arr_float_uint_4 = OpTypeArray %float %uint_4
- %LeftOver = OpTypeStruct %mat4v4float %float %_arr_mat4v4float_uint_2 %_arr_float_uint_4
+%_arr_strided_arr_uint_4 = OpTypeArray %strided_arr %uint_4
+ %LeftOver = OpTypeStruct %mat4v4float %float %_arr_mat4v4float_uint_2 %_arr_strided_arr_uint_4
%_ptr_Uniform_LeftOver = OpTypePointer Uniform %LeftOver
%x_14 = OpVariable %_ptr_Uniform_LeftOver Uniform
%_ptr_Private_v2float = OpTypePointer Private %v2float
@@ -94,7 +98,7 @@
%_ptr_Private_v4float = OpTypePointer Private %v4float
%gl_Position = OpVariable %_ptr_Private_v4float Private %12
%void = OpTypeVoid
- %37 = OpTypeFunction %void
+ %38 = OpTypeFunction %void
%_ptr_Function_v4float = OpTypePointer Function %v4float
%_ptr_Function_v3float = OpTypePointer Function %v3float
%float_1 = OpConstant %float 1
@@ -110,88 +114,88 @@
%_ptr_Uniform_mat4v4float = OpTypePointer Uniform %mat4v4float
%float_n1 = OpConstant %float -1
%main_out = OpTypeStruct %v4float %v2float
- %102 = OpTypeFunction %main_out %v3float %v2float %v3float
- %main_1 = OpFunction %void None %37
- %40 = OpLabel
+ %103 = OpTypeFunction %main_out %v3float %v2float %v3float
+ %main_1 = OpFunction %void None %38
+ %41 = OpLabel
%q = OpVariable %_ptr_Function_v4float Function %12
%p = OpVariable %_ptr_Function_v3float Function %21
- %45 = OpLoad %v3float %position
- %46 = OpCompositeExtract %float %45 0
- %47 = OpCompositeExtract %float %45 1
- %48 = OpCompositeExtract %float %45 2
- %50 = OpCompositeConstruct %v4float %46 %47 %48 %float_1
- OpStore %q %50
- %51 = OpLoad %v4float %q
- %52 = OpCompositeExtract %float %51 0
- %53 = OpCompositeExtract %float %51 1
- %54 = OpCompositeExtract %float %51 2
- %55 = OpCompositeConstruct %v3float %52 %53 %54
- OpStore %p %55
- %58 = OpAccessChain %_ptr_Function_float %p %uint_0
- %59 = OpLoad %float %58
- %64 = OpAccessChain %_ptr_Uniform_float %x_14 %uint_3 %int_0
- %65 = OpLoad %float %64
- %68 = OpAccessChain %_ptr_Private_float %position %uint_1
- %69 = OpLoad %float %68
- %70 = OpAccessChain %_ptr_Uniform_float %x_14 %uint_1
- %71 = OpLoad %float %70
- %72 = OpAccessChain %_ptr_Function_float %p %uint_0
- %75 = OpFMul %float %65 %69
- %76 = OpFAdd %float %75 %71
- %73 = OpExtInst %float %74 Sin %76
- %77 = OpFAdd %float %59 %73
- OpStore %72 %77
- %78 = OpAccessChain %_ptr_Function_float %p %uint_1
- %79 = OpLoad %float %78
- %80 = OpAccessChain %_ptr_Uniform_float %x_14 %uint_1
- %81 = OpLoad %float %80
- %82 = OpAccessChain %_ptr_Function_float %p %uint_1
- %85 = OpFAdd %float %81 %float_4
- %83 = OpExtInst %float %74 Sin %85
- %86 = OpFAdd %float %79 %83
- OpStore %82 %86
- %88 = OpAccessChain %_ptr_Uniform_mat4v4float %x_14 %uint_0
- %89 = OpLoad %mat4v4float %88
- %90 = OpLoad %v3float %p
- %91 = OpCompositeExtract %float %90 0
- %92 = OpCompositeExtract %float %90 1
- %93 = OpCompositeExtract %float %90 2
- %94 = OpCompositeConstruct %v4float %91 %92 %93 %float_1
- %95 = OpMatrixTimesVector %v4float %89 %94
- OpStore %gl_Position %95
- %96 = OpLoad %v2float %uv
- OpStore %vUV %96
- %97 = OpAccessChain %_ptr_Private_float %gl_Position %uint_1
- %98 = OpLoad %float %97
- %99 = OpAccessChain %_ptr_Private_float %gl_Position %uint_1
- %101 = OpFMul %float %98 %float_n1
- OpStore %99 %101
+ %46 = OpLoad %v3float %position
+ %47 = OpCompositeExtract %float %46 0
+ %48 = OpCompositeExtract %float %46 1
+ %49 = OpCompositeExtract %float %46 2
+ %51 = OpCompositeConstruct %v4float %47 %48 %49 %float_1
+ OpStore %q %51
+ %52 = OpLoad %v4float %q
+ %53 = OpCompositeExtract %float %52 0
+ %54 = OpCompositeExtract %float %52 1
+ %55 = OpCompositeExtract %float %52 2
+ %56 = OpCompositeConstruct %v3float %53 %54 %55
+ OpStore %p %56
+ %59 = OpAccessChain %_ptr_Function_float %p %uint_0
+ %60 = OpLoad %float %59
+ %65 = OpAccessChain %_ptr_Uniform_float %x_14 %uint_3 %int_0 %uint_0
+ %66 = OpLoad %float %65
+ %69 = OpAccessChain %_ptr_Private_float %position %uint_1
+ %70 = OpLoad %float %69
+ %71 = OpAccessChain %_ptr_Uniform_float %x_14 %uint_1
+ %72 = OpLoad %float %71
+ %73 = OpAccessChain %_ptr_Function_float %p %uint_0
+ %76 = OpFMul %float %66 %70
+ %77 = OpFAdd %float %76 %72
+ %74 = OpExtInst %float %75 Sin %77
+ %78 = OpFAdd %float %60 %74
+ OpStore %73 %78
+ %79 = OpAccessChain %_ptr_Function_float %p %uint_1
+ %80 = OpLoad %float %79
+ %81 = OpAccessChain %_ptr_Uniform_float %x_14 %uint_1
+ %82 = OpLoad %float %81
+ %83 = OpAccessChain %_ptr_Function_float %p %uint_1
+ %86 = OpFAdd %float %82 %float_4
+ %84 = OpExtInst %float %75 Sin %86
+ %87 = OpFAdd %float %80 %84
+ OpStore %83 %87
+ %89 = OpAccessChain %_ptr_Uniform_mat4v4float %x_14 %uint_0
+ %90 = OpLoad %mat4v4float %89
+ %91 = OpLoad %v3float %p
+ %92 = OpCompositeExtract %float %91 0
+ %93 = OpCompositeExtract %float %91 1
+ %94 = OpCompositeExtract %float %91 2
+ %95 = OpCompositeConstruct %v4float %92 %93 %94 %float_1
+ %96 = OpMatrixTimesVector %v4float %90 %95
+ OpStore %gl_Position %96
+ %97 = OpLoad %v2float %uv
+ OpStore %vUV %97
+ %98 = OpAccessChain %_ptr_Private_float %gl_Position %uint_1
+ %99 = OpLoad %float %98
+ %100 = OpAccessChain %_ptr_Private_float %gl_Position %uint_1
+ %102 = OpFMul %float %99 %float_n1
+ OpStore %100 %102
OpReturn
OpFunctionEnd
- %main_inner = OpFunction %main_out None %102
+ %main_inner = OpFunction %main_out None %103
%position_param = OpFunctionParameter %v3float
%uv_param = OpFunctionParameter %v2float
%normal_param = OpFunctionParameter %v3float
- %108 = OpLabel
+ %109 = OpLabel
OpStore %position %position_param
OpStore %uv %uv_param
OpStore %normal %normal_param
- %109 = OpFunctionCall %void %main_1
- %110 = OpLoad %v4float %gl_Position
- %111 = OpLoad %v2float %vUV
- %112 = OpCompositeConstruct %main_out %110 %111
- OpReturnValue %112
+ %110 = OpFunctionCall %void %main_1
+ %111 = OpLoad %v4float %gl_Position
+ %112 = OpLoad %v2float %vUV
+ %113 = OpCompositeConstruct %main_out %111 %112
+ OpReturnValue %113
OpFunctionEnd
- %main = OpFunction %void None %37
- %114 = OpLabel
- %116 = OpLoad %v3float %position_param_1
- %117 = OpLoad %v2float %uv_param_1
- %118 = OpLoad %v3float %normal_param_1
- %115 = OpFunctionCall %main_out %main_inner %116 %117 %118
- %119 = OpCompositeExtract %v4float %115 0
- OpStore %gl_Position_1 %119
- %120 = OpCompositeExtract %v2float %115 1
- OpStore %vUV_1_1 %120
+ %main = OpFunction %void None %38
+ %115 = OpLabel
+ %117 = OpLoad %v3float %position_param_1
+ %118 = OpLoad %v2float %uv_param_1
+ %119 = OpLoad %v3float %normal_param_1
+ %116 = OpFunctionCall %main_out %main_inner %117 %118 %119
+ %120 = OpCompositeExtract %v4float %116 0
+ OpStore %gl_Position_1 %120
+ %121 = OpCompositeExtract %v2float %116 1
+ OpStore %vUV_1_1 %121
OpStore %vertex_point_size %float_1
OpReturn
OpFunctionEnd
diff --git a/test/bug/tint/1088.spvasm.expected.wgsl b/test/bug/tint/1088.spvasm.expected.wgsl
index c67da67..77f1063 100644
--- a/test/bug/tint/1088.spvasm.expected.wgsl
+++ b/test/bug/tint/1088.spvasm.expected.wgsl
@@ -1,6 +1,11 @@
-type Arr = @stride(64) array<mat4x4<f32>, 2u>;
+type Arr = array<mat4x4<f32>, 2u>;
-type Arr_1 = @stride(16) array<f32, 4u>;
+struct strided_arr {
+ @size(16)
+ el : f32;
+}
+
+type Arr_1 = array<strided_arr, 4u>;
struct LeftOver {
worldViewProjection : mat4x4<f32>;
@@ -31,7 +36,7 @@
let x_21 : vec4<f32> = q;
p = vec3<f32>(x_21.x, x_21.y, x_21.z);
let x_27 : f32 = p.x;
- let x_41 : f32 = x_14.test[0];
+ let x_41 : f32 = x_14.test[0].el;
let x_45 : f32 = position.y;
let x_49 : f32 = x_14.time;
p.x = (x_27 + sin(((x_41 * x_45) + x_49)));
diff --git a/test/bug/tint/870.spvasm.expected.wgsl b/test/bug/tint/870.spvasm.expected.wgsl
index 4af74ac..a04e265 100644
--- a/test/bug/tint/870.spvasm.expected.wgsl
+++ b/test/bug/tint/870.spvasm.expected.wgsl
@@ -1,4 +1,4 @@
-type Arr = @stride(4) array<i32, 6u>;
+type Arr = array<i32, 6u>;
struct sspp962805860buildInformationS {
footprint : vec4<f32>;
diff --git a/test/bug/tint/943.spvasm.expected.wgsl b/test/bug/tint/943.spvasm.expected.wgsl
index ea2d6ad..e587782 100644
--- a/test/bug/tint/943.spvasm.expected.wgsl
+++ b/test/bug/tint/943.spvasm.expected.wgsl
@@ -14,15 +14,15 @@
outShapeStrides : vec2<i32>;
}
-type RTArr = @stride(4) array<f32>;
+type RTArr = array<f32>;
-type RTArr_1 = @stride(4) array<f32>;
+type RTArr_1 = array<f32>;
struct ssbOut {
result : RTArr_1;
}
-type RTArr_2 = @stride(4) array<f32>;
+type RTArr_2 = array<f32>;
struct ssbA {
A : RTArr_1;
diff --git a/test/bug/tint/951.spvasm.expected.wgsl b/test/bug/tint/951.spvasm.expected.wgsl
index 63fd79e..865bac8 100644
--- a/test/bug/tint/951.spvasm.expected.wgsl
+++ b/test/bug/tint/951.spvasm.expected.wgsl
@@ -1,6 +1,6 @@
-type RTArr = @stride(4) array<f32>;
+type RTArr = array<f32>;
-type RTArr_1 = @stride(4) array<f32>;
+type RTArr_1 = array<f32>;
struct ssbOut {
result : RTArr_1;
diff --git a/test/bug/tint/977.spvasm.expected.wgsl b/test/bug/tint/977.spvasm.expected.wgsl
index 8b5ad53..ac33a71 100644
--- a/test/bug/tint/977.spvasm.expected.wgsl
+++ b/test/bug/tint/977.spvasm.expected.wgsl
@@ -1,12 +1,12 @@
-type RTArr = @stride(4) array<f32>;
+type RTArr = array<f32>;
-type RTArr_1 = @stride(4) array<f32>;
+type RTArr_1 = array<f32>;
struct ResultMatrix {
numbers : RTArr_1;
}
-type RTArr_2 = @stride(4) array<f32>;
+type RTArr_2 = array<f32>;
struct FirstMatrix {
numbers : RTArr_1;
diff --git a/test/layout/storage/mat2x2/stride/16.spvasm.expected.glsl b/test/layout/storage/mat2x2/stride/16.spvasm.expected.glsl
index aca2256..e82307f 100644
--- a/test/layout/storage/mat2x2/stride/16.spvasm.expected.glsl
+++ b/test/layout/storage/mat2x2/stride/16.spvasm.expected.glsl
@@ -1,24 +1,26 @@
#version 310 es
precision mediump float;
-struct tint_padded_array_element {
+struct strided_arr {
vec2 el;
};
struct SSBO {
- tint_padded_array_element m[2];
+ strided_arr m[2];
};
layout(binding = 0) buffer SSBO_1 {
- tint_padded_array_element m[2];
+ strided_arr m[2];
} ssbo;
-mat2 arr_to_mat2x2_stride_16(tint_padded_array_element arr[2]) {
+mat2 arr_to_mat2x2_stride_16(strided_arr arr[2]) {
return mat2(arr[0u].el, arr[1u].el);
}
-tint_padded_array_element[2] mat2x2_stride_16_to_arr(mat2 mat) {
- tint_padded_array_element tint_symbol[2] = tint_padded_array_element[2](tint_padded_array_element(mat[0u]), tint_padded_array_element(mat[1u]));
- return tint_symbol;
+strided_arr[2] mat2x2_stride_16_to_arr(mat2 mat) {
+ strided_arr tint_symbol = strided_arr(mat[0u]);
+ strided_arr tint_symbol_1 = strided_arr(mat[1u]);
+ strided_arr tint_symbol_2[2] = strided_arr[2](tint_symbol, tint_symbol_1);
+ return tint_symbol_2;
}
void f_1() {
diff --git a/test/layout/storage/mat2x2/stride/16.spvasm.expected.hlsl b/test/layout/storage/mat2x2/stride/16.spvasm.expected.hlsl
index 18bd54d..7e88aa7 100644
--- a/test/layout/storage/mat2x2/stride/16.spvasm.expected.hlsl
+++ b/test/layout/storage/mat2x2/stride/16.spvasm.expected.hlsl
@@ -1,42 +1,53 @@
-struct tint_padded_array_element {
+struct strided_arr {
float2 el;
};
RWByteAddressBuffer ssbo : register(u0, space0);
-float2x2 arr_to_mat2x2_stride_16(tint_padded_array_element arr[2]) {
+float2x2 arr_to_mat2x2_stride_16(strided_arr arr[2]) {
return float2x2(arr[0u].el, arr[1u].el);
}
-typedef tint_padded_array_element mat2x2_stride_16_to_arr_ret[2];
+typedef strided_arr mat2x2_stride_16_to_arr_ret[2];
mat2x2_stride_16_to_arr_ret mat2x2_stride_16_to_arr(float2x2 mat) {
- const tint_padded_array_element tint_symbol_4[2] = {{mat[0u]}, {mat[1u]}};
- return tint_symbol_4;
+ const strided_arr tint_symbol_6 = {mat[0u]};
+ const strided_arr tint_symbol_7 = {mat[1u]};
+ const strided_arr tint_symbol_8[2] = {tint_symbol_6, tint_symbol_7};
+ return tint_symbol_8;
}
-typedef tint_padded_array_element tint_symbol_ret[2];
+strided_arr tint_symbol_1(RWByteAddressBuffer buffer, uint offset) {
+ const strided_arr tint_symbol_9 = {asfloat(buffer.Load2((offset + 0u)))};
+ return tint_symbol_9;
+}
+
+typedef strided_arr tint_symbol_ret[2];
tint_symbol_ret tint_symbol(RWByteAddressBuffer buffer, uint offset) {
- tint_padded_array_element arr_1[2] = (tint_padded_array_element[2])0;
+ strided_arr arr_1[2] = (strided_arr[2])0;
{
[loop] for(uint i = 0u; (i < 2u); i = (i + 1u)) {
- arr_1[i].el = asfloat(buffer.Load2((offset + (i * 16u))));
+ arr_1[i] = tint_symbol_1(buffer, (offset + (i * 16u)));
}
}
return arr_1;
}
-void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[2]) {
- tint_padded_array_element array[2] = value;
+void tint_symbol_4(RWByteAddressBuffer buffer, uint offset, strided_arr value) {
+ buffer.Store2((offset + 0u), asuint(value.el));
+}
+
+void tint_symbol_3(RWByteAddressBuffer buffer, uint offset, strided_arr value[2]) {
+ strided_arr array[2] = value;
{
[loop] for(uint i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
- buffer.Store2((offset + (i_1 * 16u)), asuint(array[i_1].el));
+ tint_symbol_4(buffer, (offset + (i_1 * 16u)), array[i_1]);
}
}
}
void f_1() {
const float2x2 x_15 = arr_to_mat2x2_stride_16(tint_symbol(ssbo, 0u));
- tint_symbol_2(ssbo, 0u, mat2x2_stride_16_to_arr(x_15));
+ tint_symbol_3(ssbo, 0u, mat2x2_stride_16_to_arr(x_15));
return;
}
diff --git a/test/layout/storage/mat2x2/stride/16.spvasm.expected.msl b/test/layout/storage/mat2x2/stride/16.spvasm.expected.msl
index ec498f4..6a0ca3e 100644
--- a/test/layout/storage/mat2x2/stride/16.spvasm.expected.msl
+++ b/test/layout/storage/mat2x2/stride/16.spvasm.expected.msl
@@ -1,12 +1,12 @@
#include <metal_stdlib>
using namespace metal;
-struct tint_padded_array_element {
+struct strided_arr {
/* 0x0000 */ float2 el;
/* 0x0008 */ int8_t tint_pad[8];
};
struct tint_array_wrapper {
- /* 0x0000 */ tint_padded_array_element arr[2];
+ /* 0x0000 */ strided_arr arr[2];
};
struct SSBO {
/* 0x0000 */ tint_array_wrapper m;
@@ -17,18 +17,20 @@
}
tint_array_wrapper mat2x2_stride_16_to_arr(float2x2 mat) {
- tint_array_wrapper const tint_symbol = {.arr={{.el=mat[0u]}, {.el=mat[1u]}}};
- return tint_symbol;
+ strided_arr const tint_symbol = {.el=mat[0u]};
+ strided_arr const tint_symbol_1 = {.el=mat[1u]};
+ tint_array_wrapper const tint_symbol_2 = {.arr={tint_symbol, tint_symbol_1}};
+ return tint_symbol_2;
}
-void f_1(device SSBO* const tint_symbol_1) {
- float2x2 const x_15 = arr_to_mat2x2_stride_16((*(tint_symbol_1)).m);
- (*(tint_symbol_1)).m = mat2x2_stride_16_to_arr(x_15);
+void f_1(device SSBO* const tint_symbol_3) {
+ float2x2 const x_15 = arr_to_mat2x2_stride_16((*(tint_symbol_3)).m);
+ (*(tint_symbol_3)).m = mat2x2_stride_16_to_arr(x_15);
return;
}
-kernel void f(device SSBO* tint_symbol_2 [[buffer(0)]]) {
- f_1(tint_symbol_2);
+kernel void f(device SSBO* tint_symbol_4 [[buffer(0)]]) {
+ f_1(tint_symbol_4);
return;
}
diff --git a/test/layout/storage/mat2x2/stride/16.spvasm.expected.spvasm b/test/layout/storage/mat2x2/stride/16.spvasm.expected.spvasm
index 94c280b..f5ee801 100644
--- a/test/layout/storage/mat2x2/stride/16.spvasm.expected.spvasm
+++ b/test/layout/storage/mat2x2/stride/16.spvasm.expected.spvasm
@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
-; Bound: 39
+; Bound: 44
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -9,6 +9,8 @@
OpExecutionMode %f LocalSize 1 1 1
OpName %SSBO "SSBO"
OpMemberName %SSBO 0 "m"
+ OpName %strided_arr "strided_arr"
+ OpMemberName %strided_arr 0 "el"
OpName %ssbo "ssbo"
OpName %arr_to_mat2x2_stride_16 "arr_to_mat2x2_stride_16"
OpName %arr "arr"
@@ -18,53 +20,59 @@
OpName %f "f"
OpDecorate %SSBO Block
OpMemberDecorate %SSBO 0 Offset 0
- OpDecorate %_arr_v2float_uint_2 ArrayStride 16
+ OpMemberDecorate %strided_arr 0 Offset 0
+ OpDecorate %_arr_strided_arr_uint_2 ArrayStride 16
OpDecorate %ssbo DescriptorSet 0
OpDecorate %ssbo Binding 0
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
+%strided_arr = OpTypeStruct %v2float
%uint = OpTypeInt 32 0
%uint_2 = OpConstant %uint 2
-%_arr_v2float_uint_2 = OpTypeArray %v2float %uint_2
- %SSBO = OpTypeStruct %_arr_v2float_uint_2
+%_arr_strided_arr_uint_2 = OpTypeArray %strided_arr %uint_2
+ %SSBO = OpTypeStruct %_arr_strided_arr_uint_2
%_ptr_StorageBuffer_SSBO = OpTypePointer StorageBuffer %SSBO
%ssbo = OpVariable %_ptr_StorageBuffer_SSBO StorageBuffer
%mat2v2float = OpTypeMatrix %v2float 2
- %9 = OpTypeFunction %mat2v2float %_arr_v2float_uint_2
+ %10 = OpTypeFunction %mat2v2float %_arr_strided_arr_uint_2
%uint_0 = OpConstant %uint 0
%uint_1 = OpConstant %uint 1
- %19 = OpTypeFunction %_arr_v2float_uint_2 %mat2v2float
+ %22 = OpTypeFunction %_arr_strided_arr_uint_2 %mat2v2float
%void = OpTypeVoid
- %26 = OpTypeFunction %void
-%_ptr_StorageBuffer__arr_v2float_uint_2 = OpTypePointer StorageBuffer %_arr_v2float_uint_2
-%arr_to_mat2x2_stride_16 = OpFunction %mat2v2float None %9
- %arr = OpFunctionParameter %_arr_v2float_uint_2
- %13 = OpLabel
- %15 = OpCompositeExtract %v2float %arr 0
- %17 = OpCompositeExtract %v2float %arr 1
- %18 = OpCompositeConstruct %mat2v2float %15 %17
- OpReturnValue %18
+ %31 = OpTypeFunction %void
+%_ptr_StorageBuffer__arr_strided_arr_uint_2 = OpTypePointer StorageBuffer %_arr_strided_arr_uint_2
+%arr_to_mat2x2_stride_16 = OpFunction %mat2v2float None %10
+ %arr = OpFunctionParameter %_arr_strided_arr_uint_2
+ %14 = OpLabel
+ %16 = OpCompositeExtract %strided_arr %arr 0
+ %17 = OpCompositeExtract %v2float %16 0
+ %19 = OpCompositeExtract %strided_arr %arr 1
+ %20 = OpCompositeExtract %v2float %19 0
+ %21 = OpCompositeConstruct %mat2v2float %17 %20
+ OpReturnValue %21
OpFunctionEnd
-%mat2x2_stride_16_to_arr = OpFunction %_arr_v2float_uint_2 None %19
+%mat2x2_stride_16_to_arr = OpFunction %_arr_strided_arr_uint_2 None %22
%mat = OpFunctionParameter %mat2v2float
- %22 = OpLabel
- %23 = OpCompositeExtract %v2float %mat 0
- %24 = OpCompositeExtract %v2float %mat 1
- %25 = OpCompositeConstruct %_arr_v2float_uint_2 %23 %24
- OpReturnValue %25
+ %25 = OpLabel
+ %26 = OpCompositeExtract %v2float %mat 0
+ %27 = OpCompositeConstruct %strided_arr %26
+ %28 = OpCompositeExtract %v2float %mat 1
+ %29 = OpCompositeConstruct %strided_arr %28
+ %30 = OpCompositeConstruct %_arr_strided_arr_uint_2 %27 %29
+ OpReturnValue %30
OpFunctionEnd
- %f_1 = OpFunction %void None %26
- %29 = OpLabel
- %32 = OpAccessChain %_ptr_StorageBuffer__arr_v2float_uint_2 %ssbo %uint_0
- %33 = OpLoad %_arr_v2float_uint_2 %32
- %30 = OpFunctionCall %mat2v2float %arr_to_mat2x2_stride_16 %33
- %34 = OpAccessChain %_ptr_StorageBuffer__arr_v2float_uint_2 %ssbo %uint_0
- %35 = OpFunctionCall %_arr_v2float_uint_2 %mat2x2_stride_16_to_arr %30
- OpStore %34 %35
+ %f_1 = OpFunction %void None %31
+ %34 = OpLabel
+ %37 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_uint_2 %ssbo %uint_0
+ %38 = OpLoad %_arr_strided_arr_uint_2 %37
+ %35 = OpFunctionCall %mat2v2float %arr_to_mat2x2_stride_16 %38
+ %39 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_uint_2 %ssbo %uint_0
+ %40 = OpFunctionCall %_arr_strided_arr_uint_2 %mat2x2_stride_16_to_arr %35
+ OpStore %39 %40
OpReturn
OpFunctionEnd
- %f = OpFunction %void None %26
- %37 = OpLabel
- %38 = OpFunctionCall %void %f_1
+ %f = OpFunction %void None %31
+ %42 = OpLabel
+ %43 = OpFunctionCall %void %f_1
OpReturn
OpFunctionEnd
diff --git a/test/layout/storage/mat2x2/stride/16.spvasm.expected.wgsl b/test/layout/storage/mat2x2/stride/16.spvasm.expected.wgsl
index a368ee1..669f8cc 100644
--- a/test/layout/storage/mat2x2/stride/16.spvasm.expected.wgsl
+++ b/test/layout/storage/mat2x2/stride/16.spvasm.expected.wgsl
@@ -1,15 +1,20 @@
+struct strided_arr {
+ @size(16)
+ el : vec2<f32>;
+}
+
struct SSBO {
- m : @stride(16) array<vec2<f32>, 2u>;
+ m : array<strided_arr, 2u>;
}
@group(0) @binding(0) var<storage, read_write> ssbo : SSBO;
-fn arr_to_mat2x2_stride_16(arr : @stride(16) array<vec2<f32>, 2u>) -> mat2x2<f32> {
- return mat2x2<f32>(arr[0u], arr[1u]);
+fn arr_to_mat2x2_stride_16(arr : array<strided_arr, 2u>) -> mat2x2<f32> {
+ return mat2x2<f32>(arr[0u].el, arr[1u].el);
}
-fn mat2x2_stride_16_to_arr(mat : mat2x2<f32>) -> @stride(16) array<vec2<f32>, 2u> {
- return @stride(16) array<vec2<f32>, 2u>(mat[0u], mat[1u]);
+fn mat2x2_stride_16_to_arr(mat : mat2x2<f32>) -> array<strided_arr, 2u> {
+ return array<strided_arr, 2u>(strided_arr(mat[0u]), strided_arr(mat[1u]));
}
fn f_1() {
