tint/test: Regenerate expectations
Fix collision of two CLs landing with different expectations.
Change-Id: I44eb904b552f635e37dd51dcc94329fbc34af031
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/94685
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Austin Eng <enga@chromium.org>
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
index 896696e..943a64c 100644
--- a/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.spvasm.expected.msl
@@ -1,31 +1,20 @@
#include <metal_stdlib>
using namespace metal;
-struct tint_array_wrapper {
- uint arr[1];
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
};
-struct tint_array_wrapper_1 {
- tint_array_wrapper arr[2];
-};
-
-struct tint_array_wrapper_2 {
- tint_array_wrapper_1 arr[3];
-};
-
-struct tint_array_wrapper_5 {
- atomic_uint arr[1];
-};
-
-struct tint_array_wrapper_4 {
- tint_array_wrapper_5 arr[2];
-};
-
-struct tint_array_wrapper_3 {
- tint_array_wrapper_4 arr[3];
-};
-
-void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrapper_3* const tint_symbol) {
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
uint idx = 0u;
idx = local_invocation_index;
while (true) {
@@ -36,29 +25,29 @@
uint const x_31 = idx;
uint const x_33 = idx;
uint const x_35 = idx;
- atomic_store_explicit(&((*(tint_symbol)).arr[(x_31 / 2u)].arr[(x_33 % 2u)].arr[(x_35 % 1u)]), 0u, memory_order_relaxed);
+ atomic_store_explicit(&((*(tint_symbol))[(x_31 / 2u)][(x_33 % 2u)][(x_35 % 1u)]), 0u, memory_order_relaxed);
{
uint const x_42 = idx;
idx = (x_42 + 1u);
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- atomic_store_explicit(&((*(tint_symbol)).arr[2].arr[1].arr[0]), 1u, memory_order_relaxed);
+ atomic_store_explicit(&((*(tint_symbol))[2][1][0]), 1u, memory_order_relaxed);
return;
}
-void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array_wrapper_3* const tint_symbol_2) {
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_2) {
uint const x_57 = *(tint_symbol_1);
compute_main_inner(x_57, tint_symbol_2);
return;
}
-void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array_wrapper_3* const tint_symbol_3, thread uint* const tint_symbol_4) {
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_3, thread uint* const tint_symbol_4) {
for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
uint const i = (idx_1 / 2u);
uint const i_1 = (idx_1 % 2u);
uint const i_2 = (idx_1 % 1u);
- atomic_store_explicit(&((*(tint_symbol_3)).arr[i].arr[i_1].arr[i_2]), 0u, memory_order_relaxed);
+ atomic_store_explicit(&((*(tint_symbol_3))[i][i_1][i_2]), 0u, memory_order_relaxed);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
*(tint_symbol_4) = local_invocation_index_1_param;
@@ -66,7 +55,7 @@
}
kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
- threadgroup tint_array_wrapper_3 tint_symbol_5;
+ threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_5;
thread uint tint_symbol_6 = 0u;
compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
return;
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.glsl b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.glsl
new file mode 100644
index 0000000..598fde9
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.glsl
@@ -0,0 +1,21 @@
+#version 310 es
+
+shared uint wg[3][2][1];
+void compute_main(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
+ uint i = (idx / 2u);
+ uint i_1 = (idx % 2u);
+ uint i_2 = (idx % 1u);
+ atomicExchange(wg[i][i_1][i_2], 0u);
+ }
+ }
+ barrier();
+ atomicExchange(wg[2][1][0], 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.hlsl
new file mode 100644
index 0000000..6f993d5
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.hlsl
@@ -0,0 +1,26 @@
+groupshared uint wg[3][2][1];
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ [loop] for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
+ const uint i = (idx / 2u);
+ const uint i_1 = (idx % 2u);
+ const uint i_2 = (idx % 1u);
+ uint atomic_result = 0u;
+ InterlockedExchange(wg[i][i_1][i_2], 0u, atomic_result);
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ uint atomic_result_1 = 0u;
+ InterlockedExchange(wg[2][1][0], 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl
new file mode 100644
index 0000000..51e5245
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
+ for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
+ uint const i = (idx / 2u);
+ uint const i_1 = (idx % 2u);
+ uint const i_2 = (idx % 1u);
+ atomic_store_explicit(&((*(tint_symbol))[i][i_1][i_2]), 0u, memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomic_store_explicit(&((*(tint_symbol))[2][1][0]), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_1;
+ compute_main_inner(local_invocation_index, &(tint_symbol_1));
+ return;
+}
+
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.spvasm
new file mode 100644
index 0000000..0dbd964
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.spvasm
@@ -0,0 +1,88 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 58
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %wg "wg"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %idx "idx"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ OpDecorate %_arr_uint_uint_1 ArrayStride 4
+ OpDecorate %_arr__arr_uint_uint_1_uint_2 ArrayStride 4
+ OpDecorate %_arr__arr__arr_uint_uint_1_uint_2_uint_3 ArrayStride 8
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %uint_1 = OpConstant %uint 1
+%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
+ %uint_2 = OpConstant %uint 2
+%_arr__arr_uint_uint_1_uint_2 = OpTypeArray %_arr_uint_uint_1 %uint_2
+ %uint_3 = OpConstant %uint 3
+%_arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypeArray %_arr__arr_uint_uint_1_uint_2 %uint_3
+%_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypePointer Workgroup %_arr__arr__arr_uint_uint_1_uint_2_uint_3
+ %wg = OpVariable %_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 Workgroup
+ %void = OpTypeVoid
+ %12 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %19 = OpConstantNull %uint
+ %uint_6 = OpConstant %uint 6
+ %bool = OpTypeBool
+ %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %uint_264 = OpConstant %uint 264
+ %int = OpTypeInt 32 1
+ %int_2 = OpConstant %int 2
+ %int_1 = OpConstant %int 1
+ %51 = OpConstantNull %int
+ %53 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %12
+%local_invocation_index = OpFunctionParameter %uint
+ %16 = OpLabel
+ %idx = OpVariable %_ptr_Function_uint Function %19
+ OpStore %idx %local_invocation_index
+ OpBranch %20
+ %20 = OpLabel
+ OpLoopMerge %21 %22 None
+ OpBranch %23
+ %23 = OpLabel
+ %25 = OpLoad %uint %idx
+ %27 = OpULessThan %bool %25 %uint_6
+ %24 = OpLogicalNot %bool %27
+ OpSelectionMerge %29 None
+ OpBranchConditional %24 %30 %29
+ %30 = OpLabel
+ OpBranch %21
+ %29 = OpLabel
+ %31 = OpLoad %uint %idx
+ %32 = OpUDiv %uint %31 %uint_2
+ %33 = OpLoad %uint %idx
+ %34 = OpUMod %uint %33 %uint_2
+ %35 = OpLoad %uint %idx
+ %36 = OpUMod %uint %35 %uint_1
+ %41 = OpAccessChain %_ptr_Workgroup_uint %wg %32 %34 %36
+ OpAtomicStore %41 %uint_2 %uint_0 %19
+ OpBranch %22
+ %22 = OpLabel
+ %42 = OpLoad %uint %idx
+ %43 = OpIAdd %uint %42 %uint_1
+ OpStore %idx %43
+ OpBranch %20
+ %21 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %52 = OpAccessChain %_ptr_Workgroup_uint %wg %int_2 %int_1 %51
+ OpAtomicStore %52 %uint_2 %uint_0 %uint_1
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %53
+ %55 = OpLabel
+ %57 = OpLoad %uint %local_invocation_index_1
+ %56 = OpFunctionCall %void %compute_main_inner %57
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.wgsl
new file mode 100644
index 0000000..0a5ac8a
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.wgsl
@@ -0,0 +1,14 @@
+type A0 = atomic<u32>;
+
+type A1 = array<A0, 1>;
+
+type A2 = array<A1, 2>;
+
+type A3 = array<A2, 3>;
+
+var<workgroup> wg : A3;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ atomicStore(&(wg[2][1][0]), 1u);
+}
diff --git a/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl b/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
index d2ab1ec..8ef9d39 100644
--- a/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/array.spvasm.expected.msl
@@ -1,15 +1,20 @@
#include <metal_stdlib>
using namespace metal;
-struct tint_array_wrapper {
- uint arr[4];
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
};
-struct tint_array_wrapper_1 {
- atomic_uint arr[4];
-};
-
-void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrapper_1* const tint_symbol) {
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<atomic_uint, 4>* const tint_symbol) {
uint idx = 0u;
idx = local_invocation_index;
while (true) {
@@ -18,27 +23,27 @@
break;
}
uint const x_26 = idx;
- atomic_store_explicit(&((*(tint_symbol)).arr[x_26]), 0u, memory_order_relaxed);
+ atomic_store_explicit(&((*(tint_symbol))[x_26]), 0u, memory_order_relaxed);
{
uint const x_33 = idx;
idx = (x_33 + 1u);
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- atomic_store_explicit(&((*(tint_symbol)).arr[1]), 1u, memory_order_relaxed);
+ atomic_store_explicit(&((*(tint_symbol))[1]), 1u, memory_order_relaxed);
return;
}
-void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array_wrapper_1* const tint_symbol_2) {
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array<atomic_uint, 4>* const tint_symbol_2) {
uint const x_47 = *(tint_symbol_1);
compute_main_inner(x_47, tint_symbol_2);
return;
}
-void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array_wrapper_1* const tint_symbol_3, thread uint* const tint_symbol_4) {
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<atomic_uint, 4>* const tint_symbol_3, thread uint* const tint_symbol_4) {
for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 4u); idx_1 = (idx_1 + 1u)) {
uint const i = idx_1;
- atomic_store_explicit(&((*(tint_symbol_3)).arr[i]), 0u, memory_order_relaxed);
+ atomic_store_explicit(&((*(tint_symbol_3))[i]), 0u, memory_order_relaxed);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
*(tint_symbol_4) = local_invocation_index_1_param;
@@ -46,7 +51,7 @@
}
kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
- threadgroup tint_array_wrapper_1 tint_symbol_5;
+ threadgroup tint_array<atomic_uint, 4> tint_symbol_5;
thread uint tint_symbol_6 = 0u;
compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
return;
diff --git a/test/tint/builtins/atomicStore/array/array.wgsl.expected.glsl b/test/tint/builtins/atomicStore/array/array.wgsl.expected.glsl
new file mode 100644
index 0000000..653bae6
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.glsl
@@ -0,0 +1,19 @@
+#version 310 es
+
+shared uint wg[4];
+void compute_main(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ uint i = idx;
+ atomicExchange(wg[i], 0u);
+ }
+ }
+ barrier();
+ atomicExchange(wg[1], 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/array/array.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/array/array.wgsl.expected.hlsl
new file mode 100644
index 0000000..d2ff575
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.hlsl
@@ -0,0 +1,24 @@
+groupshared uint wg[4];
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ [loop] for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ const uint i = idx;
+ uint atomic_result = 0u;
+ InterlockedExchange(wg[i], 0u, atomic_result);
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ uint atomic_result_1 = 0u;
+ InterlockedExchange(wg[1], 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl b/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl
new file mode 100644
index 0000000..5f20d27
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl
@@ -0,0 +1,31 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<atomic_uint, 4>* const tint_symbol) {
+ for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
+ uint const i = idx;
+ atomic_store_explicit(&((*(tint_symbol))[i]), 0u, memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomic_store_explicit(&((*(tint_symbol))[1]), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup tint_array<atomic_uint, 4> tint_symbol_1;
+ compute_main_inner(local_invocation_index, &(tint_symbol_1));
+ return;
+}
+
diff --git a/test/tint/builtins/atomicStore/array/array.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/array/array.wgsl.expected.spvasm
new file mode 100644
index 0000000..f41ddf7
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.spvasm
@@ -0,0 +1,76 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 48
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %wg "wg"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %idx "idx"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ OpDecorate %_arr_uint_uint_4 ArrayStride 4
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %uint_4 = OpConstant %uint 4
+%_arr_uint_uint_4 = OpTypeArray %uint %uint_4
+%_ptr_Workgroup__arr_uint_uint_4 = OpTypePointer Workgroup %_arr_uint_uint_4
+ %wg = OpVariable %_ptr_Workgroup__arr_uint_uint_4 Workgroup
+ %void = OpTypeVoid
+ %8 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %15 = OpConstantNull %uint
+ %bool = OpTypeBool
+ %uint_2 = OpConstant %uint 2
+ %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %uint_1 = OpConstant %uint 1
+ %uint_264 = OpConstant %uint 264
+ %int = OpTypeInt 32 1
+ %int_1 = OpConstant %int 1
+ %43 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %8
+%local_invocation_index = OpFunctionParameter %uint
+ %12 = OpLabel
+ %idx = OpVariable %_ptr_Function_uint Function %15
+ OpStore %idx %local_invocation_index
+ OpBranch %16
+ %16 = OpLabel
+ OpLoopMerge %17 %18 None
+ OpBranch %19
+ %19 = OpLabel
+ %21 = OpLoad %uint %idx
+ %22 = OpULessThan %bool %21 %uint_4
+ %20 = OpLogicalNot %bool %22
+ OpSelectionMerge %24 None
+ OpBranchConditional %20 %25 %24
+ %25 = OpLabel
+ OpBranch %17
+ %24 = OpLabel
+ %26 = OpLoad %uint %idx
+ %32 = OpAccessChain %_ptr_Workgroup_uint %wg %26
+ OpAtomicStore %32 %uint_2 %uint_0 %15
+ OpBranch %18
+ %18 = OpLabel
+ %33 = OpLoad %uint %idx
+ %35 = OpIAdd %uint %33 %uint_1
+ OpStore %idx %35
+ OpBranch %16
+ %17 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %42 = OpAccessChain %_ptr_Workgroup_uint %wg %int_1
+ OpAtomicStore %42 %uint_2 %uint_0 %uint_1
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %43
+ %45 = OpLabel
+ %47 = OpLoad %uint %local_invocation_index_1
+ %46 = OpFunctionCall %void %compute_main_inner %47
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/array/array.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/array/array.wgsl.expected.wgsl
new file mode 100644
index 0000000..f6812d4
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+var<workgroup> wg : array<atomic<u32>, 4>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ atomicStore(&(wg[1]), 1u);
+}
diff --git a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
index 896696e..943a64c 100644
--- a/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/array/arrays.spvasm.expected.msl
@@ -1,31 +1,20 @@
#include <metal_stdlib>
using namespace metal;
-struct tint_array_wrapper {
- uint arr[1];
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
};
-struct tint_array_wrapper_1 {
- tint_array_wrapper arr[2];
-};
-
-struct tint_array_wrapper_2 {
- tint_array_wrapper_1 arr[3];
-};
-
-struct tint_array_wrapper_5 {
- atomic_uint arr[1];
-};
-
-struct tint_array_wrapper_4 {
- tint_array_wrapper_5 arr[2];
-};
-
-struct tint_array_wrapper_3 {
- tint_array_wrapper_4 arr[3];
-};
-
-void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrapper_3* const tint_symbol) {
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
uint idx = 0u;
idx = local_invocation_index;
while (true) {
@@ -36,29 +25,29 @@
uint const x_31 = idx;
uint const x_33 = idx;
uint const x_35 = idx;
- atomic_store_explicit(&((*(tint_symbol)).arr[(x_31 / 2u)].arr[(x_33 % 2u)].arr[(x_35 % 1u)]), 0u, memory_order_relaxed);
+ atomic_store_explicit(&((*(tint_symbol))[(x_31 / 2u)][(x_33 % 2u)][(x_35 % 1u)]), 0u, memory_order_relaxed);
{
uint const x_42 = idx;
idx = (x_42 + 1u);
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- atomic_store_explicit(&((*(tint_symbol)).arr[2].arr[1].arr[0]), 1u, memory_order_relaxed);
+ atomic_store_explicit(&((*(tint_symbol))[2][1][0]), 1u, memory_order_relaxed);
return;
}
-void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array_wrapper_3* const tint_symbol_2) {
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_2) {
uint const x_57 = *(tint_symbol_1);
compute_main_inner(x_57, tint_symbol_2);
return;
}
-void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array_wrapper_3* const tint_symbol_3, thread uint* const tint_symbol_4) {
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol_3, thread uint* const tint_symbol_4) {
for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 6u); idx_1 = (idx_1 + 1u)) {
uint const i = (idx_1 / 2u);
uint const i_1 = (idx_1 % 2u);
uint const i_2 = (idx_1 % 1u);
- atomic_store_explicit(&((*(tint_symbol_3)).arr[i].arr[i_1].arr[i_2]), 0u, memory_order_relaxed);
+ atomic_store_explicit(&((*(tint_symbol_3))[i][i_1][i_2]), 0u, memory_order_relaxed);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
*(tint_symbol_4) = local_invocation_index_1_param;
@@ -66,7 +55,7 @@
}
kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
- threadgroup tint_array_wrapper_3 tint_symbol_5;
+ threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_5;
thread uint tint_symbol_6 = 0u;
compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
return;
diff --git a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.glsl b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.glsl
new file mode 100644
index 0000000..598fde9
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.glsl
@@ -0,0 +1,21 @@
+#version 310 es
+
+shared uint wg[3][2][1];
+void compute_main(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
+ uint i = (idx / 2u);
+ uint i_1 = (idx % 2u);
+ uint i_2 = (idx % 1u);
+ atomicExchange(wg[i][i_1][i_2], 0u);
+ }
+ }
+ barrier();
+ atomicExchange(wg[2][1][0], 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.hlsl
new file mode 100644
index 0000000..6f993d5
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.hlsl
@@ -0,0 +1,26 @@
+groupshared uint wg[3][2][1];
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ [loop] for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
+ const uint i = (idx / 2u);
+ const uint i_1 = (idx % 2u);
+ const uint i_2 = (idx % 1u);
+ uint atomic_result = 0u;
+ InterlockedExchange(wg[i][i_1][i_2], 0u, atomic_result);
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ uint atomic_result_1 = 0u;
+ InterlockedExchange(wg[2][1][0], 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl
new file mode 100644
index 0000000..51e5245
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl
@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3>* const tint_symbol) {
+ for(uint idx = local_invocation_index; (idx < 6u); idx = (idx + 1u)) {
+ uint const i = (idx / 2u);
+ uint const i_1 = (idx % 2u);
+ uint const i_2 = (idx % 1u);
+ atomic_store_explicit(&((*(tint_symbol))[i][i_1][i_2]), 0u, memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomic_store_explicit(&((*(tint_symbol))[2][1][0]), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup tint_array<tint_array<tint_array<atomic_uint, 1>, 2>, 3> tint_symbol_1;
+ compute_main_inner(local_invocation_index, &(tint_symbol_1));
+ return;
+}
+
diff --git a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.spvasm
new file mode 100644
index 0000000..0dbd964
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.spvasm
@@ -0,0 +1,88 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 58
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %wg "wg"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %idx "idx"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ OpDecorate %_arr_uint_uint_1 ArrayStride 4
+ OpDecorate %_arr__arr_uint_uint_1_uint_2 ArrayStride 4
+ OpDecorate %_arr__arr__arr_uint_uint_1_uint_2_uint_3 ArrayStride 8
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %uint_1 = OpConstant %uint 1
+%_arr_uint_uint_1 = OpTypeArray %uint %uint_1
+ %uint_2 = OpConstant %uint 2
+%_arr__arr_uint_uint_1_uint_2 = OpTypeArray %_arr_uint_uint_1 %uint_2
+ %uint_3 = OpConstant %uint 3
+%_arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypeArray %_arr__arr_uint_uint_1_uint_2 %uint_3
+%_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 = OpTypePointer Workgroup %_arr__arr__arr_uint_uint_1_uint_2_uint_3
+ %wg = OpVariable %_ptr_Workgroup__arr__arr__arr_uint_uint_1_uint_2_uint_3 Workgroup
+ %void = OpTypeVoid
+ %12 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %19 = OpConstantNull %uint
+ %uint_6 = OpConstant %uint 6
+ %bool = OpTypeBool
+ %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %uint_264 = OpConstant %uint 264
+ %int = OpTypeInt 32 1
+ %int_2 = OpConstant %int 2
+ %int_1 = OpConstant %int 1
+ %51 = OpConstantNull %int
+ %53 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %12
+%local_invocation_index = OpFunctionParameter %uint
+ %16 = OpLabel
+ %idx = OpVariable %_ptr_Function_uint Function %19
+ OpStore %idx %local_invocation_index
+ OpBranch %20
+ %20 = OpLabel
+ OpLoopMerge %21 %22 None
+ OpBranch %23
+ %23 = OpLabel
+ %25 = OpLoad %uint %idx
+ %27 = OpULessThan %bool %25 %uint_6
+ %24 = OpLogicalNot %bool %27
+ OpSelectionMerge %29 None
+ OpBranchConditional %24 %30 %29
+ %30 = OpLabel
+ OpBranch %21
+ %29 = OpLabel
+ %31 = OpLoad %uint %idx
+ %32 = OpUDiv %uint %31 %uint_2
+ %33 = OpLoad %uint %idx
+ %34 = OpUMod %uint %33 %uint_2
+ %35 = OpLoad %uint %idx
+ %36 = OpUMod %uint %35 %uint_1
+ %41 = OpAccessChain %_ptr_Workgroup_uint %wg %32 %34 %36
+ OpAtomicStore %41 %uint_2 %uint_0 %19
+ OpBranch %22
+ %22 = OpLabel
+ %42 = OpLoad %uint %idx
+ %43 = OpIAdd %uint %42 %uint_1
+ OpStore %idx %43
+ OpBranch %20
+ %21 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %52 = OpAccessChain %_ptr_Workgroup_uint %wg %int_2 %int_1 %51
+ OpAtomicStore %52 %uint_2 %uint_0 %uint_1
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %53
+ %55 = OpLabel
+ %57 = OpLoad %uint %local_invocation_index_1
+ %56 = OpFunctionCall %void %compute_main_inner %57
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.wgsl
new file mode 100644
index 0000000..3ca55977
--- /dev/null
+++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+var<workgroup> wg : array<array<array<atomic<u32>, 1>, 2>, 3>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ atomicStore(&(wg[2][1][0]), 1u);
+}
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
index 5360fda..f4fd526 100644
--- a/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.spvasm.expected.msl
@@ -1,6 +1,19 @@
#include <metal_stdlib>
using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
struct S_atomic {
int x;
atomic_uint a;
@@ -13,15 +26,7 @@
uint y;
};
-struct tint_array_wrapper {
- S arr[10];
-};
-
-struct tint_array_wrapper_1 {
- S_atomic arr[10];
-};
-
-void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrapper_1* const tint_symbol) {
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<S_atomic, 10>* const tint_symbol) {
uint idx = 0u;
idx = local_invocation_index;
while (true) {
@@ -30,31 +35,31 @@
break;
}
uint const x_28 = idx;
- (*(tint_symbol)).arr[x_28].x = 0;
- atomic_store_explicit(&((*(tint_symbol)).arr[x_28].a), 0u, memory_order_relaxed);
- (*(tint_symbol)).arr[x_28].y = 0u;
+ (*(tint_symbol))[x_28].x = 0;
+ atomic_store_explicit(&((*(tint_symbol))[x_28].a), 0u, memory_order_relaxed);
+ (*(tint_symbol))[x_28].y = 0u;
{
uint const x_41 = idx;
idx = (x_41 + 1u);
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- atomic_store_explicit(&((*(tint_symbol)).arr[4].a), 1u, memory_order_relaxed);
+ atomic_store_explicit(&((*(tint_symbol))[4].a), 1u, memory_order_relaxed);
return;
}
-void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array_wrapper_1* const tint_symbol_2) {
+void compute_main_1(thread uint* const tint_symbol_1, threadgroup tint_array<S_atomic, 10>* const tint_symbol_2) {
uint const x_53 = *(tint_symbol_1);
compute_main_inner(x_53, tint_symbol_2);
return;
}
-void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array_wrapper_1* const tint_symbol_3, thread uint* const tint_symbol_4) {
+void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_array<S_atomic, 10>* const tint_symbol_3, thread uint* const tint_symbol_4) {
for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
uint const i = idx_1;
- (*(tint_symbol_3)).arr[i].x = 0;
- atomic_store_explicit(&((*(tint_symbol_3)).arr[i].a), 0u, memory_order_relaxed);
- (*(tint_symbol_3)).arr[i].y = 0u;
+ (*(tint_symbol_3))[i].x = 0;
+ atomic_store_explicit(&((*(tint_symbol_3))[i].a), 0u, memory_order_relaxed);
+ (*(tint_symbol_3))[i].y = 0u;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
*(tint_symbol_4) = local_invocation_index_1_param;
@@ -62,7 +67,7 @@
}
kernel void compute_main(uint local_invocation_index_1_param [[thread_index_in_threadgroup]]) {
- threadgroup tint_array_wrapper_1 tint_symbol_5;
+ threadgroup tint_array<S_atomic, 10> tint_symbol_5;
thread uint tint_symbol_6 = 0u;
compute_main_inner_1(local_invocation_index_1_param, &(tint_symbol_5), &(tint_symbol_6));
return;
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.glsl b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.glsl
new file mode 100644
index 0000000..bd0a3ba
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.glsl
@@ -0,0 +1,27 @@
+#version 310 es
+
+struct S {
+ int x;
+ uint a;
+ uint y;
+};
+
+shared S wg[10];
+void compute_main(uint local_invocation_index) {
+ {
+ for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
+ uint i = idx;
+ wg[i].x = 0;
+ atomicExchange(wg[i].a, 0u);
+ wg[i].y = 0u;
+ }
+ }
+ barrier();
+ atomicExchange(wg[4].a, 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.hlsl
new file mode 100644
index 0000000..4642386
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.hlsl
@@ -0,0 +1,32 @@
+struct S {
+ int x;
+ uint a;
+ uint y;
+};
+
+groupshared S wg[10];
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ [loop] for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
+ const uint i = idx;
+ wg[i].x = 0;
+ uint atomic_result = 0u;
+ InterlockedExchange(wg[i].a, 0u, atomic_result);
+ wg[i].y = 0u;
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ uint atomic_result_1 = 0u;
+ InterlockedExchange(wg[4].a, 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl
new file mode 100644
index 0000000..978edb1
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl
@@ -0,0 +1,39 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct S {
+ int x;
+ atomic_uint a;
+ uint y;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup tint_array<S, 10>* const tint_symbol) {
+ for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
+ uint const i = idx;
+ (*(tint_symbol))[i].x = 0;
+ atomic_store_explicit(&((*(tint_symbol))[i].a), 0u, memory_order_relaxed);
+ (*(tint_symbol))[i].y = 0u;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomic_store_explicit(&((*(tint_symbol))[4].a), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup tint_array<S, 10> tint_symbol_1;
+ compute_main_inner(local_invocation_index, &(tint_symbol_1));
+ return;
+}
+
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..b9e0fed
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.spvasm
@@ -0,0 +1,91 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 54
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %S "S"
+ OpMemberName %S 0 "x"
+ OpMemberName %S 1 "a"
+ OpMemberName %S 2 "y"
+ OpName %wg "wg"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %idx "idx"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 1 Offset 4
+ OpMemberDecorate %S 2 Offset 8
+ OpDecorate %_arr_S_uint_10 ArrayStride 12
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %int = OpTypeInt 32 1
+ %S = OpTypeStruct %int %uint %uint
+ %uint_10 = OpConstant %uint 10
+%_arr_S_uint_10 = OpTypeArray %S %uint_10
+%_ptr_Workgroup__arr_S_uint_10 = OpTypePointer Workgroup %_arr_S_uint_10
+ %wg = OpVariable %_ptr_Workgroup__arr_S_uint_10 Workgroup
+ %void = OpTypeVoid
+ %10 = OpTypeFunction %void %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %17 = OpConstantNull %uint
+ %bool = OpTypeBool
+ %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+ %32 = OpConstantNull %int
+ %uint_2 = OpConstant %uint 2
+ %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+ %uint_264 = OpConstant %uint 264
+ %int_4 = OpConstant %int 4
+ %49 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %10
+%local_invocation_index = OpFunctionParameter %uint
+ %14 = OpLabel
+ %idx = OpVariable %_ptr_Function_uint Function %17
+ OpStore %idx %local_invocation_index
+ OpBranch %18
+ %18 = OpLabel
+ OpLoopMerge %19 %20 None
+ OpBranch %21
+ %21 = OpLabel
+ %23 = OpLoad %uint %idx
+ %24 = OpULessThan %bool %23 %uint_10
+ %22 = OpLogicalNot %bool %24
+ OpSelectionMerge %26 None
+ OpBranchConditional %22 %27 %26
+ %27 = OpLabel
+ OpBranch %19
+ %26 = OpLabel
+ %28 = OpLoad %uint %idx
+ %31 = OpAccessChain %_ptr_Workgroup_int %wg %28 %uint_0
+ OpStore %31 %32
+ %38 = OpAccessChain %_ptr_Workgroup_uint %wg %28 %uint_1
+ OpAtomicStore %38 %uint_2 %uint_0 %17
+ %40 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %28 %uint_2
+ OpStore %40 %17
+ OpBranch %20
+ %20 = OpLabel
+ %41 = OpLoad %uint %idx
+ %42 = OpIAdd %uint %41 %uint_1
+ OpStore %idx %42
+ OpBranch %18
+ %19 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %48 = OpAccessChain %_ptr_Workgroup_uint %wg %int_4 %uint_1
+ OpAtomicStore %48 %uint_2 %uint_0 %uint_1
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %49
+ %51 = OpLabel
+ %53 = OpLoad %uint %local_invocation_index_1
+ %52 = OpFunctionCall %void %compute_main_inner %53
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..1def9bd
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+struct S {
+ x : i32,
+ a : atomic<u32>,
+ y : u32,
+}
+
+var<workgroup> wg : array<S, 10>;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ atomicStore(&(wg[4].a), 1u);
+}
diff --git a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.glsl b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.glsl
new file mode 100644
index 0000000..d3a4b83
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.glsl
@@ -0,0 +1,25 @@
+#version 310 es
+
+struct S {
+ int x;
+ uint a;
+ uint b;
+};
+
+shared S wg;
+void compute_main(uint local_invocation_index) {
+ {
+ wg.x = 0;
+ atomicExchange(wg.a, 0u);
+ atomicExchange(wg.b, 0u);
+ }
+ barrier();
+ atomicExchange(wg.a, 1u);
+ atomicExchange(wg.b, 2u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.hlsl
new file mode 100644
index 0000000..35b428f
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.hlsl
@@ -0,0 +1,32 @@
+struct S {
+ int x;
+ uint a;
+ uint b;
+};
+
+groupshared S wg;
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ wg.x = 0;
+ uint atomic_result = 0u;
+ InterlockedExchange(wg.a, 0u, atomic_result);
+ uint atomic_result_1 = 0u;
+ InterlockedExchange(wg.b, 0u, atomic_result_1);
+ }
+ GroupMemoryBarrierWithGroupSync();
+ uint atomic_result_2 = 0u;
+ InterlockedExchange(wg.a, 1u, atomic_result_2);
+ uint atomic_result_3 = 0u;
+ InterlockedExchange(wg.b, 2u, atomic_result_3);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.msl
new file mode 100644
index 0000000..3ea40fa
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.msl
@@ -0,0 +1,26 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+ int x;
+ atomic_uint a;
+ atomic_uint b;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup S* const tint_symbol) {
+ {
+ (*(tint_symbol)).x = 0;
+ atomic_store_explicit(&((*(tint_symbol)).a), 0u, memory_order_relaxed);
+ atomic_store_explicit(&((*(tint_symbol)).b), 0u, memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomic_store_explicit(&((*(tint_symbol)).a), 1u, memory_order_relaxed);
+ atomic_store_explicit(&((*(tint_symbol)).b), 2u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup S tint_symbol_1;
+ compute_main_inner(local_invocation_index, &(tint_symbol_1));
+ return;
+}
+
diff --git a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.spvasm
new file mode 100644
index 0000000..f61269c
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.spvasm
@@ -0,0 +1,62 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 40
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %S "S"
+ OpMemberName %S 0 "x"
+ OpMemberName %S 1 "a"
+ OpMemberName %S 2 "b"
+ OpName %wg "wg"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 1 Offset 4
+ OpMemberDecorate %S 2 Offset 8
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %int = OpTypeInt 32 1
+ %S = OpTypeStruct %int %uint %uint
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+ %wg = OpVariable %_ptr_Workgroup_S Workgroup
+ %void = OpTypeVoid
+ %8 = OpTypeFunction %void %uint
+ %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+ %16 = OpConstantNull %int
+ %uint_2 = OpConstant %uint 2
+ %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %23 = OpConstantNull %uint
+ %uint_264 = OpConstant %uint 264
+ %35 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %8
+%local_invocation_index = OpFunctionParameter %uint
+ %12 = OpLabel
+ %15 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+ OpStore %15 %16
+ %22 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+ OpAtomicStore %22 %uint_2 %uint_0 %23
+ %26 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_2
+ OpAtomicStore %26 %uint_2 %uint_0 %23
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %31 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+ OpAtomicStore %31 %uint_2 %uint_0 %uint_1
+ %34 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_2
+ OpAtomicStore %34 %uint_2 %uint_0 %uint_2
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %35
+ %37 = OpLabel
+ %39 = OpLoad %uint %local_invocation_index_1
+ %38 = OpFunctionCall %void %compute_main_inner %39
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.wgsl
new file mode 100644
index 0000000..1d25d8a
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.wgsl
@@ -0,0 +1,13 @@
+struct S {
+ x : i32,
+ a : atomic<u32>,
+ b : atomic<u32>,
+}
+
+var<workgroup> wg : S;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ atomicStore(&(wg.a), 1u);
+ atomicStore(&(wg.b), 2u);
+}
diff --git a/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.glsl b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.glsl
new file mode 100644
index 0000000..39f13a5
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.glsl
@@ -0,0 +1,24 @@
+#version 310 es
+
+struct S {
+ int x;
+ uint a;
+ uint y;
+};
+
+shared S wg;
+void compute_main(uint local_invocation_index) {
+ {
+ wg.x = 0;
+ atomicExchange(wg.a, 0u);
+ wg.y = 0u;
+ }
+ barrier();
+ atomicExchange(wg.a, 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.hlsl
new file mode 100644
index 0000000..fb93daa
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.hlsl
@@ -0,0 +1,29 @@
+struct S {
+ int x;
+ uint a;
+ uint y;
+};
+
+groupshared S wg;
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ wg.x = 0;
+ uint atomic_result = 0u;
+ InterlockedExchange(wg.a, 0u, atomic_result);
+ wg.y = 0u;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ uint atomic_result_1 = 0u;
+ InterlockedExchange(wg.a, 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.msl
new file mode 100644
index 0000000..ccec811b
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.msl
@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+ int x;
+ atomic_uint a;
+ uint y;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup S* const tint_symbol) {
+ {
+ (*(tint_symbol)).x = 0;
+ atomic_store_explicit(&((*(tint_symbol)).a), 0u, memory_order_relaxed);
+ (*(tint_symbol)).y = 0u;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomic_store_explicit(&((*(tint_symbol)).a), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup S tint_symbol_1;
+ compute_main_inner(local_invocation_index, &(tint_symbol_1));
+ return;
+}
+
diff --git a/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.spvasm
new file mode 100644
index 0000000..acad988
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.spvasm
@@ -0,0 +1,61 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 36
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %S "S"
+ OpMemberName %S 0 "x"
+ OpMemberName %S 1 "a"
+ OpMemberName %S 2 "y"
+ OpName %wg "wg"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 1 Offset 4
+ OpMemberDecorate %S 2 Offset 8
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %int = OpTypeInt 32 1
+ %S = OpTypeStruct %int %uint %uint
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+ %wg = OpVariable %_ptr_Workgroup_S Workgroup
+ %void = OpTypeVoid
+ %8 = OpTypeFunction %void %uint
+ %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+ %16 = OpConstantNull %int
+ %uint_2 = OpConstant %uint 2
+ %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %23 = OpConstantNull %uint
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+ %uint_264 = OpConstant %uint 264
+ %31 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %8
+%local_invocation_index = OpFunctionParameter %uint
+ %12 = OpLabel
+ %15 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+ OpStore %15 %16
+ %22 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+ OpAtomicStore %22 %uint_2 %uint_0 %23
+ %25 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_2
+ OpStore %25 %23
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %30 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+ OpAtomicStore %30 %uint_2 %uint_0 %uint_1
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %31
+ %33 = OpLabel
+ %35 = OpLoad %uint %local_invocation_index_1
+ %34 = OpFunctionCall %void %compute_main_inner %35
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.wgsl
new file mode 100644
index 0000000..164ad73
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+struct S {
+ x : i32,
+ a : atomic<u32>,
+ y : u32,
+}
+
+var<workgroup> wg : S;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ atomicStore(&(wg.a), 1u);
+}
diff --git a/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.glsl b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.glsl
new file mode 100644
index 0000000..10f2103
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.glsl
@@ -0,0 +1,46 @@
+#version 310 es
+
+struct S0 {
+ int x;
+ uint a;
+ int y;
+ int z;
+};
+
+struct S1 {
+ int x;
+ S0 a;
+ int y;
+ int z;
+};
+
+struct S2 {
+ int x;
+ int y;
+ int z;
+ S1 a;
+};
+
+shared S2 wg;
+void compute_main(uint local_invocation_index) {
+ {
+ wg.x = 0;
+ wg.y = 0;
+ wg.z = 0;
+ wg.a.x = 0;
+ wg.a.a.x = 0;
+ atomicExchange(wg.a.a.a, 0u);
+ wg.a.a.y = 0;
+ wg.a.a.z = 0;
+ wg.a.y = 0;
+ wg.a.z = 0;
+ }
+ barrier();
+ atomicExchange(wg.a.a.a, 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.hlsl
new file mode 100644
index 0000000..4ceb2ff
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.hlsl
@@ -0,0 +1,49 @@
+struct S0 {
+ int x;
+ uint a;
+ int y;
+ int z;
+};
+struct S1 {
+ int x;
+ S0 a;
+ int y;
+ int z;
+};
+struct S2 {
+ int x;
+ int y;
+ int z;
+ S1 a;
+};
+
+groupshared S2 wg;
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ wg.x = 0;
+ wg.y = 0;
+ wg.z = 0;
+ wg.a.x = 0;
+ wg.a.a.x = 0;
+ uint atomic_result = 0u;
+ InterlockedExchange(wg.a.a.a, 0u, atomic_result);
+ wg.a.a.y = 0;
+ wg.a.a.z = 0;
+ wg.a.y = 0;
+ wg.a.z = 0;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ uint atomic_result_1 = 0u;
+ InterlockedExchange(wg.a.a.a, 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.msl
new file mode 100644
index 0000000..dbe7cc0
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.msl
@@ -0,0 +1,47 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S0 {
+ int x;
+ atomic_uint a;
+ int y;
+ int z;
+};
+
+struct S1 {
+ int x;
+ S0 a;
+ int y;
+ int z;
+};
+
+struct S2 {
+ int x;
+ int y;
+ int z;
+ S1 a;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup S2* const tint_symbol) {
+ {
+ (*(tint_symbol)).x = 0;
+ (*(tint_symbol)).y = 0;
+ (*(tint_symbol)).z = 0;
+ (*(tint_symbol)).a.x = 0;
+ (*(tint_symbol)).a.a.x = 0;
+ atomic_store_explicit(&((*(tint_symbol)).a.a.a), 0u, memory_order_relaxed);
+ (*(tint_symbol)).a.a.y = 0;
+ (*(tint_symbol)).a.a.z = 0;
+ (*(tint_symbol)).a.y = 0;
+ (*(tint_symbol)).a.z = 0;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomic_store_explicit(&((*(tint_symbol)).a.a.a), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup S2 tint_symbol_1;
+ compute_main_inner(local_invocation_index, &(tint_symbol_1));
+ return;
+}
+
diff --git a/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.spvasm
new file mode 100644
index 0000000..a8aaff3
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.spvasm
@@ -0,0 +1,97 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 45
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %S2 "S2"
+ OpMemberName %S2 0 "x"
+ OpMemberName %S2 1 "y"
+ OpMemberName %S2 2 "z"
+ OpMemberName %S2 3 "a"
+ OpName %S1 "S1"
+ OpMemberName %S1 0 "x"
+ OpMemberName %S1 1 "a"
+ OpName %S0 "S0"
+ OpMemberName %S0 0 "x"
+ OpMemberName %S0 1 "a"
+ OpMemberName %S0 2 "y"
+ OpMemberName %S0 3 "z"
+ OpMemberName %S1 2 "y"
+ OpMemberName %S1 3 "z"
+ OpName %wg "wg"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ OpMemberDecorate %S2 0 Offset 0
+ OpMemberDecorate %S2 1 Offset 4
+ OpMemberDecorate %S2 2 Offset 8
+ OpMemberDecorate %S2 3 Offset 12
+ OpMemberDecorate %S1 0 Offset 0
+ OpMemberDecorate %S1 1 Offset 4
+ OpMemberDecorate %S0 0 Offset 0
+ OpMemberDecorate %S0 1 Offset 4
+ OpMemberDecorate %S0 2 Offset 8
+ OpMemberDecorate %S0 3 Offset 12
+ OpMemberDecorate %S1 2 Offset 20
+ OpMemberDecorate %S1 3 Offset 24
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %int = OpTypeInt 32 1
+ %S0 = OpTypeStruct %int %uint %int %int
+ %S1 = OpTypeStruct %int %S0 %int %int
+ %S2 = OpTypeStruct %int %int %int %S1
+%_ptr_Workgroup_S2 = OpTypePointer Workgroup %S2
+ %wg = OpVariable %_ptr_Workgroup_S2 Workgroup
+ %void = OpTypeVoid
+ %10 = OpTypeFunction %void %uint
+ %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+ %18 = OpConstantNull %int
+ %uint_1 = OpConstant %uint 1
+ %uint_2 = OpConstant %uint 2
+ %uint_3 = OpConstant %uint 3
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %30 = OpConstantNull %uint
+ %uint_264 = OpConstant %uint 264
+ %40 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %10
+%local_invocation_index = OpFunctionParameter %uint
+ %14 = OpLabel
+ %17 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+ OpStore %17 %18
+ %20 = OpAccessChain %_ptr_Workgroup_int %wg %uint_1
+ OpStore %20 %18
+ %22 = OpAccessChain %_ptr_Workgroup_int %wg %uint_2
+ OpStore %22 %18
+ %24 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_0
+ OpStore %24 %18
+ %25 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_1 %uint_0
+ OpStore %25 %18
+ %29 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_3 %uint_1 %uint_1
+ OpAtomicStore %29 %uint_2 %uint_0 %30
+ %31 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_1 %uint_2
+ OpStore %31 %18
+ %32 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_1 %uint_3
+ OpStore %32 %18
+ %33 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_2
+ OpStore %33 %18
+ %34 = OpAccessChain %_ptr_Workgroup_int %wg %uint_3 %uint_3
+ OpStore %34 %18
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %39 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_3 %uint_1 %uint_1
+ OpAtomicStore %39 %uint_2 %uint_0 %uint_1
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %40
+ %42 = OpLabel
+ %44 = OpLoad %uint %local_invocation_index_1
+ %43 = OpFunctionCall %void %compute_main_inner %44
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.wgsl
new file mode 100644
index 0000000..6a9bdac
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.wgsl
@@ -0,0 +1,27 @@
+struct S0 {
+ x : i32,
+ a : atomic<u32>,
+ y : i32,
+ z : i32,
+}
+
+struct S1 {
+ x : i32,
+ a : S0,
+ y : i32,
+ z : i32,
+}
+
+struct S2 {
+ x : i32,
+ y : i32,
+ z : i32,
+ a : S1,
+}
+
+var<workgroup> wg : S2;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ atomicStore(&(wg.a.a.a), 1u);
+}
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
index 089c011..69d565e 100644
--- a/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.spvasm.expected.msl
@@ -1,23 +1,28 @@
#include <metal_stdlib>
using namespace metal;
-struct tint_array_wrapper {
- uint arr[10];
-};
-struct tint_array_wrapper_1 {
- atomic_uint arr[10];
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
};
struct S_atomic {
int x;
- tint_array_wrapper_1 a;
+ tint_array<atomic_uint, 10> a;
uint y;
};
struct S {
int x;
- tint_array_wrapper a;
+ tint_array<uint, 10> a;
uint y;
};
@@ -32,14 +37,14 @@
break;
}
uint const x_35 = idx;
- atomic_store_explicit(&((*(tint_symbol)).a.arr[x_35]), 0u, memory_order_relaxed);
+ atomic_store_explicit(&((*(tint_symbol)).a[x_35]), 0u, memory_order_relaxed);
{
uint const x_41 = idx;
idx = (x_41 + 1u);
}
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- atomic_store_explicit(&((*(tint_symbol)).a.arr[4]), 1u, memory_order_relaxed);
+ atomic_store_explicit(&((*(tint_symbol)).a[4]), 1u, memory_order_relaxed);
return;
}
@@ -56,7 +61,7 @@
}
for(uint idx_1 = local_invocation_index_1_param; (idx_1 < 10u); idx_1 = (idx_1 + 1u)) {
uint const i = idx_1;
- atomic_store_explicit(&((*(tint_symbol_3)).a.arr[i]), 0u, memory_order_relaxed);
+ atomic_store_explicit(&((*(tint_symbol_3)).a[i]), 0u, memory_order_relaxed);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
*(tint_symbol_4) = local_invocation_index_1_param;
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.glsl b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.glsl
new file mode 100644
index 0000000..287b69b
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.glsl
@@ -0,0 +1,29 @@
+#version 310 es
+
+struct S {
+ int x;
+ uint a[10];
+ uint y;
+};
+
+shared S wg;
+void compute_main(uint local_invocation_index) {
+ {
+ wg.x = 0;
+ wg.y = 0u;
+ }
+ {
+ for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
+ uint i = idx;
+ atomicExchange(wg.a[i], 0u);
+ }
+ }
+ barrier();
+ atomicExchange(wg.a[4], 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.hlsl
new file mode 100644
index 0000000..855c198
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.hlsl
@@ -0,0 +1,34 @@
+struct S {
+ int x;
+ uint a[10];
+ uint y;
+};
+
+groupshared S wg;
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ wg.x = 0;
+ wg.y = 0u;
+ }
+ {
+ [loop] for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
+ const uint i = idx;
+ uint atomic_result = 0u;
+ InterlockedExchange(wg.a[i], 0u, atomic_result);
+ }
+ }
+ GroupMemoryBarrierWithGroupSync();
+ uint atomic_result_1 = 0u;
+ InterlockedExchange(wg.a[4], 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl
new file mode 100644
index 0000000..00fa7d5
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl
@@ -0,0 +1,41 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct S {
+ int x;
+ tint_array<atomic_uint, 10> a;
+ uint y;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup S* const tint_symbol) {
+ {
+ (*(tint_symbol)).x = 0;
+ (*(tint_symbol)).y = 0u;
+ }
+ for(uint idx = local_invocation_index; (idx < 10u); idx = (idx + 1u)) {
+ uint const i = idx;
+ atomic_store_explicit(&((*(tint_symbol)).a[i]), 0u, memory_order_relaxed);
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomic_store_explicit(&((*(tint_symbol)).a[4]), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup S tint_symbol_1;
+ compute_main_inner(local_invocation_index, &(tint_symbol_1));
+ return;
+}
+
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.spvasm
new file mode 100644
index 0000000..9d06d13
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.spvasm
@@ -0,0 +1,91 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 54
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %S "S"
+ OpMemberName %S 0 "x"
+ OpMemberName %S 1 "a"
+ OpMemberName %S 2 "y"
+ OpName %wg "wg"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %idx "idx"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 1 Offset 4
+ OpDecorate %_arr_uint_uint_10 ArrayStride 4
+ OpMemberDecorate %S 2 Offset 44
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %int = OpTypeInt 32 1
+ %uint_10 = OpConstant %uint 10
+%_arr_uint_uint_10 = OpTypeArray %uint %uint_10
+ %S = OpTypeStruct %int %_arr_uint_uint_10 %uint
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+ %wg = OpVariable %_ptr_Workgroup_S Workgroup
+ %void = OpTypeVoid
+ %10 = OpTypeFunction %void %uint
+ %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+ %18 = OpConstantNull %int
+ %uint_2 = OpConstant %uint 2
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %22 = OpConstantNull %uint
+%_ptr_Function_uint = OpTypePointer Function %uint
+ %bool = OpTypeBool
+ %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+ %uint_264 = OpConstant %uint 264
+ %int_4 = OpConstant %int 4
+ %49 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %10
+%local_invocation_index = OpFunctionParameter %uint
+ %14 = OpLabel
+ %idx = OpVariable %_ptr_Function_uint Function %22
+ %17 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+ OpStore %17 %18
+ %21 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_2
+ OpStore %21 %22
+ OpStore %idx %local_invocation_index
+ OpBranch %25
+ %25 = OpLabel
+ OpLoopMerge %26 %27 None
+ OpBranch %28
+ %28 = OpLabel
+ %30 = OpLoad %uint %idx
+ %31 = OpULessThan %bool %30 %uint_10
+ %29 = OpLogicalNot %bool %31
+ OpSelectionMerge %33 None
+ OpBranchConditional %29 %34 %33
+ %34 = OpLabel
+ OpBranch %26
+ %33 = OpLabel
+ %35 = OpLoad %uint %idx
+ %40 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_1 %35
+ OpAtomicStore %40 %uint_2 %uint_0 %22
+ OpBranch %27
+ %27 = OpLabel
+ %41 = OpLoad %uint %idx
+ %42 = OpIAdd %uint %41 %uint_1
+ OpStore %idx %42
+ OpBranch %25
+ %26 = OpLabel
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %48 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_1 %int_4
+ OpAtomicStore %48 %uint_2 %uint_0 %uint_1
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %49
+ %51 = OpLabel
+ %53 = OpLoad %uint %local_invocation_index_1
+ %52 = OpFunctionCall %void %compute_main_inner %53
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.wgsl
new file mode 100644
index 0000000..793d8c8
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+struct S {
+ x : i32,
+ a : array<atomic<u32>, 10>,
+ y : u32,
+}
+
+var<workgroup> wg : S;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ atomicStore(&(wg.a[4]), 1u);
+}
diff --git a/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.glsl b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.glsl
new file mode 100644
index 0000000..39f13a5
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.glsl
@@ -0,0 +1,24 @@
+#version 310 es
+
+struct S {
+ int x;
+ uint a;
+ uint y;
+};
+
+shared S wg;
+void compute_main(uint local_invocation_index) {
+ {
+ wg.x = 0;
+ atomicExchange(wg.a, 0u);
+ wg.y = 0u;
+ }
+ barrier();
+ atomicExchange(wg.a, 1u);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ compute_main(gl_LocalInvocationIndex);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.hlsl b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.hlsl
new file mode 100644
index 0000000..fb93daa
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.hlsl
@@ -0,0 +1,29 @@
+struct S {
+ int x;
+ uint a;
+ uint y;
+};
+
+groupshared S wg;
+
+struct tint_symbol_1 {
+ uint local_invocation_index : SV_GroupIndex;
+};
+
+void compute_main_inner(uint local_invocation_index) {
+ {
+ wg.x = 0;
+ uint atomic_result = 0u;
+ InterlockedExchange(wg.a, 0u, atomic_result);
+ wg.y = 0u;
+ }
+ GroupMemoryBarrierWithGroupSync();
+ uint atomic_result_1 = 0u;
+ InterlockedExchange(wg.a, 1u, atomic_result_1);
+}
+
+[numthreads(1, 1, 1)]
+void compute_main(tint_symbol_1 tint_symbol) {
+ compute_main_inner(tint_symbol.local_invocation_index);
+ return;
+}
diff --git a/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.msl b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.msl
new file mode 100644
index 0000000..ccec811b
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.msl
@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+ int x;
+ atomic_uint a;
+ uint y;
+};
+
+void compute_main_inner(uint local_invocation_index, threadgroup S* const tint_symbol) {
+ {
+ (*(tint_symbol)).x = 0;
+ atomic_store_explicit(&((*(tint_symbol)).a), 0u, memory_order_relaxed);
+ (*(tint_symbol)).y = 0u;
+ }
+ threadgroup_barrier(mem_flags::mem_threadgroup);
+ atomic_store_explicit(&((*(tint_symbol)).a), 1u, memory_order_relaxed);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+ threadgroup S tint_symbol_1;
+ compute_main_inner(local_invocation_index, &(tint_symbol_1));
+ return;
+}
+
diff --git a/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.spvasm b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.spvasm
new file mode 100644
index 0000000..acad988
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.spvasm
@@ -0,0 +1,61 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 36
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main" %local_invocation_index_1
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpName %local_invocation_index_1 "local_invocation_index_1"
+ OpName %S "S"
+ OpMemberName %S 0 "x"
+ OpMemberName %S 1 "a"
+ OpMemberName %S 2 "y"
+ OpName %wg "wg"
+ OpName %compute_main_inner "compute_main_inner"
+ OpName %local_invocation_index "local_invocation_index"
+ OpName %compute_main "compute_main"
+ OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 1 Offset 4
+ OpMemberDecorate %S 2 Offset 8
+ %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+ %int = OpTypeInt 32 1
+ %S = OpTypeStruct %int %uint %uint
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+ %wg = OpVariable %_ptr_Workgroup_S Workgroup
+ %void = OpTypeVoid
+ %8 = OpTypeFunction %void %uint
+ %uint_0 = OpConstant %uint 0
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+ %16 = OpConstantNull %int
+ %uint_2 = OpConstant %uint 2
+ %uint_1 = OpConstant %uint 1
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+ %23 = OpConstantNull %uint
+%_ptr_Workgroup_uint_0 = OpTypePointer Workgroup %uint
+ %uint_264 = OpConstant %uint 264
+ %31 = OpTypeFunction %void
+%compute_main_inner = OpFunction %void None %8
+%local_invocation_index = OpFunctionParameter %uint
+ %12 = OpLabel
+ %15 = OpAccessChain %_ptr_Workgroup_int %wg %uint_0
+ OpStore %15 %16
+ %22 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+ OpAtomicStore %22 %uint_2 %uint_0 %23
+ %25 = OpAccessChain %_ptr_Workgroup_uint_0 %wg %uint_2
+ OpStore %25 %23
+ OpControlBarrier %uint_2 %uint_2 %uint_264
+ %30 = OpAccessChain %_ptr_Workgroup_uint %wg %uint_1
+ OpAtomicStore %30 %uint_2 %uint_0 %uint_1
+ OpReturn
+ OpFunctionEnd
+%compute_main = OpFunction %void None %31
+ %33 = OpLabel
+ %35 = OpLoad %uint %local_invocation_index_1
+ %34 = OpFunctionCall %void %compute_main_inner %35
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.wgsl b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.wgsl
new file mode 100644
index 0000000..5076b67
--- /dev/null
+++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.wgsl
@@ -0,0 +1,14 @@
+struct S {
+ x : i32,
+ a : atomic<u32>,
+ y : u32,
+}
+
+var<workgroup> wg : S;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ let p0 = &(wg);
+ let p1 = &((*(p0)).a);
+ atomicStore(p1, 1u);
+}
