From bd8449f37d312f633c7409d77d5a9419a9748dac Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Sat, 25 Jun 2022 00:55:59 +0000 Subject: [PATCH] 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 Commit-Queue: Ben Clayton Reviewed-by: Austin Eng --- .../array/aliased_arrays.spvasm.expected.msl | 47 ++++----- .../array/aliased_arrays.wgsl.expected.glsl | 21 ++++ .../array/aliased_arrays.wgsl.expected.hlsl | 26 +++++ .../array/aliased_arrays.wgsl.expected.msl | 33 +++++++ .../array/aliased_arrays.wgsl.expected.spvasm | 88 +++++++++++++++++ .../array/aliased_arrays.wgsl.expected.wgsl | 14 +++ .../array/array.spvasm.expected.msl | 31 +++--- .../array/array.wgsl.expected.glsl | 19 ++++ .../array/array.wgsl.expected.hlsl | 24 +++++ .../atomicStore/array/array.wgsl.expected.msl | 31 ++++++ .../array/array.wgsl.expected.spvasm | 76 +++++++++++++++ .../array/array.wgsl.expected.wgsl | 6 ++ .../array/arrays.spvasm.expected.msl | 47 ++++----- .../array/arrays.wgsl.expected.glsl | 21 ++++ .../array/arrays.wgsl.expected.hlsl | 26 +++++ .../array/arrays.wgsl.expected.msl | 33 +++++++ .../array/arrays.wgsl.expected.spvasm | 88 +++++++++++++++++ .../array/arrays.wgsl.expected.wgsl | 6 ++ .../array_of_struct.spvasm.expected.msl | 43 ++++---- .../struct/array_of_struct.wgsl.expected.glsl | 27 ++++++ .../struct/array_of_struct.wgsl.expected.hlsl | 32 ++++++ .../struct/array_of_struct.wgsl.expected.msl | 39 ++++++++ .../array_of_struct.wgsl.expected.spvasm | 91 +++++++++++++++++ .../struct/array_of_struct.wgsl.expected.wgsl | 12 +++ .../flat_multiple_atomics.wgsl.expected.glsl | 25 +++++ .../flat_multiple_atomics.wgsl.expected.hlsl | 32 ++++++ .../flat_multiple_atomics.wgsl.expected.msl | 26 +++++ ...flat_multiple_atomics.wgsl.expected.spvasm | 62 ++++++++++++ .../flat_multiple_atomics.wgsl.expected.wgsl | 13 +++ .../flat_single_atomic.wgsl.expected.glsl | 24 +++++ .../flat_single_atomic.wgsl.expected.hlsl | 29 ++++++ .../flat_single_atomic.wgsl.expected.msl | 25 +++++ .../flat_single_atomic.wgsl.expected.spvasm | 61 ++++++++++++ .../flat_single_atomic.wgsl.expected.wgsl | 12 +++ .../struct/nested.wgsl.expected.glsl | 46 +++++++++ .../struct/nested.wgsl.expected.hlsl | 49 ++++++++++ .../struct/nested.wgsl.expected.msl | 47 +++++++++ .../struct/nested.wgsl.expected.spvasm | 97 +++++++++++++++++++ .../struct/nested.wgsl.expected.wgsl | 27 ++++++ .../struct_of_array.spvasm.expected.msl | 25 +++-- .../struct/struct_of_array.wgsl.expected.glsl | 29 ++++++ .../struct/struct_of_array.wgsl.expected.hlsl | 34 +++++++ .../struct/struct_of_array.wgsl.expected.msl | 41 ++++++++ .../struct_of_array.wgsl.expected.spvasm | 91 +++++++++++++++++ .../struct/struct_of_array.wgsl.expected.wgsl | 12 +++ .../struct/via_ptr_let.wgsl.expected.glsl | 24 +++++ .../struct/via_ptr_let.wgsl.expected.hlsl | 29 ++++++ .../struct/via_ptr_let.wgsl.expected.msl | 25 +++++ .../struct/via_ptr_let.wgsl.expected.spvasm | 61 ++++++++++++ .../struct/via_ptr_let.wgsl.expected.wgsl | 14 +++ 50 files changed, 1741 insertions(+), 100 deletions(-) create mode 100644 test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.glsl create mode 100644 test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.hlsl create mode 100644 test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl create mode 100644 test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.spvasm create mode 100644 test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.wgsl create mode 100644 test/tint/builtins/atomicStore/array/array.wgsl.expected.glsl create mode 100644 test/tint/builtins/atomicStore/array/array.wgsl.expected.hlsl create mode 100644 test/tint/builtins/atomicStore/array/array.wgsl.expected.msl create mode 100644 test/tint/builtins/atomicStore/array/array.wgsl.expected.spvasm create mode 100644 test/tint/builtins/atomicStore/array/array.wgsl.expected.wgsl create mode 100644 test/tint/builtins/atomicStore/array/arrays.wgsl.expected.glsl create mode 100644 test/tint/builtins/atomicStore/array/arrays.wgsl.expected.hlsl create mode 100644 test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl create mode 100644 test/tint/builtins/atomicStore/array/arrays.wgsl.expected.spvasm create mode 100644 test/tint/builtins/atomicStore/array/arrays.wgsl.expected.wgsl create mode 100644 test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.glsl create mode 100644 test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.hlsl create mode 100644 test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl create mode 100644 test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.spvasm create mode 100644 test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.wgsl create mode 100644 test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.glsl create mode 100644 test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.hlsl create mode 100644 test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.msl create mode 100644 test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.spvasm create mode 100644 test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.wgsl create mode 100644 test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.glsl create mode 100644 test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.hlsl create mode 100644 test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.msl create mode 100644 test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.spvasm create mode 100644 test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.wgsl create mode 100644 test/tint/builtins/atomicStore/struct/nested.wgsl.expected.glsl create mode 100644 test/tint/builtins/atomicStore/struct/nested.wgsl.expected.hlsl create mode 100644 test/tint/builtins/atomicStore/struct/nested.wgsl.expected.msl create mode 100644 test/tint/builtins/atomicStore/struct/nested.wgsl.expected.spvasm create mode 100644 test/tint/builtins/atomicStore/struct/nested.wgsl.expected.wgsl create mode 100644 test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.glsl create mode 100644 test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.hlsl create mode 100644 test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl create mode 100644 test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.spvasm create mode 100644 test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.wgsl create mode 100644 test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.glsl create mode 100644 test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.hlsl create mode 100644 test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.msl create mode 100644 test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.spvasm create mode 100644 test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.wgsl 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 896696e8b0..943a64cf5c 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 using namespace metal; -struct tint_array_wrapper { - uint arr[1]; + +template +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, 2>, 3>* const tint_symbol) { uint idx = 0u; idx = local_invocation_index; while (true) { @@ -36,29 +25,29 @@ void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrap 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, 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, 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 @@ void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_ } 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, 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 0000000000..598fde91f4 --- /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 0000000000..6f993d5073 --- /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 0000000000..51e5245e4c --- /dev/null +++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; + +template +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, 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, 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 0000000000..0dbd964435 --- /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 0000000000..0a5ac8a00e --- /dev/null +++ b/test/tint/builtins/atomicStore/array/aliased_arrays.wgsl.expected.wgsl @@ -0,0 +1,14 @@ +type A0 = atomic; + +type A1 = array; + +type A2 = array; + +type A3 = array; + +var 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 d2ab1ecebf..8ef9d3975b 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 using namespace metal; -struct tint_array_wrapper { - uint arr[4]; + +template +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* const tint_symbol) { uint idx = 0u; idx = local_invocation_index; while (true) { @@ -18,27 +23,27 @@ void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrap 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* 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* 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 @@ void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_ } 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 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 0000000000..653bae6c55 --- /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 0000000000..d2ff575cd5 --- /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 0000000000..5f20d271bb --- /dev/null +++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.msl @@ -0,0 +1,31 @@ +#include + +using namespace metal; + +template +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* 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 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 0000000000..f41ddf706c --- /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 0000000000..f6812d4d30 --- /dev/null +++ b/test/tint/builtins/atomicStore/array/array.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +var wg : array, 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 896696e8b0..943a64cf5c 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 using namespace metal; -struct tint_array_wrapper { - uint arr[1]; + +template +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, 2>, 3>* const tint_symbol) { uint idx = 0u; idx = local_invocation_index; while (true) { @@ -36,29 +25,29 @@ void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrap 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, 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, 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 @@ void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_ } 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, 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 0000000000..598fde91f4 --- /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 0000000000..6f993d5073 --- /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 0000000000..51e5245e4c --- /dev/null +++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.msl @@ -0,0 +1,33 @@ +#include + +using namespace metal; + +template +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, 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, 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 0000000000..0dbd964435 --- /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 0000000000..3ca5597704 --- /dev/null +++ b/test/tint/builtins/atomicStore/array/arrays.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +var wg : array, 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 5360fda067..f4fd526a21 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 using namespace metal; + +template +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 @@ struct S { 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* const tint_symbol) { uint idx = 0u; idx = local_invocation_index; while (true) { @@ -30,31 +35,31 @@ void compute_main_inner(uint local_invocation_index, threadgroup tint_array_wrap 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* 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* 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 @@ void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup tint_ } 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 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 0000000000..bd0a3ba6bb --- /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 0000000000..4642386bf0 --- /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 0000000000..978edb198d --- /dev/null +++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.msl @@ -0,0 +1,39 @@ +#include + +using namespace metal; + +template +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* 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 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 0000000000..b9e0fed8f4 --- /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 0000000000..1def9bdcfb --- /dev/null +++ b/test/tint/builtins/atomicStore/struct/array_of_struct.wgsl.expected.wgsl @@ -0,0 +1,12 @@ +struct S { + x : i32, + a : atomic, + y : u32, +} + +var wg : array; + +@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 0000000000..d3a4b832ac --- /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 0000000000..35b428feee --- /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 0000000000..3ea40fa37e --- /dev/null +++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.msl @@ -0,0 +1,26 @@ +#include + +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 0000000000..f61269c96a --- /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 0000000000..1d25d8adbd --- /dev/null +++ b/test/tint/builtins/atomicStore/struct/flat_multiple_atomics.wgsl.expected.wgsl @@ -0,0 +1,13 @@ +struct S { + x : i32, + a : atomic, + b : atomic, +} + +var 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 0000000000..39f13a5880 --- /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 0000000000..fb93daa369 --- /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 0000000000..ccec811b27 --- /dev/null +++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.msl @@ -0,0 +1,25 @@ +#include + +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 0000000000..acad988c5f --- /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 0000000000..164ad73445 --- /dev/null +++ b/test/tint/builtins/atomicStore/struct/flat_single_atomic.wgsl.expected.wgsl @@ -0,0 +1,12 @@ +struct S { + x : i32, + a : atomic, + y : u32, +} + +var 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 0000000000..10f210371a --- /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 0000000000..4ceb2ff2de --- /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 0000000000..dbe7cc0854 --- /dev/null +++ b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.msl @@ -0,0 +1,47 @@ +#include + +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 0000000000..a8aaff3d02 --- /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 0000000000..6a9bdac1d6 --- /dev/null +++ b/test/tint/builtins/atomicStore/struct/nested.wgsl.expected.wgsl @@ -0,0 +1,27 @@ +struct S0 { + x : i32, + a : atomic, + 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 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 089c01164e..69d565e3d5 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 using namespace metal; -struct tint_array_wrapper { - uint arr[10]; -}; -struct tint_array_wrapper_1 { - atomic_uint arr[10]; +template +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 a; uint y; }; struct S { int x; - tint_array_wrapper a; + tint_array a; uint y; }; @@ -32,14 +37,14 @@ void compute_main_inner(uint local_invocation_index, threadgroup S_atomic* const 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 @@ void compute_main_inner_1(uint local_invocation_index_1_param, threadgroup S_ato } 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 0000000000..287b69b79b --- /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 0000000000..855c198a0f --- /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 0000000000..00fa7d560b --- /dev/null +++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.msl @@ -0,0 +1,41 @@ +#include + +using namespace metal; + +template +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 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 0000000000..9d06d130a0 --- /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 0000000000..793d8c8c66 --- /dev/null +++ b/test/tint/builtins/atomicStore/struct/struct_of_array.wgsl.expected.wgsl @@ -0,0 +1,12 @@ +struct S { + x : i32, + a : array, 10>, + y : u32, +} + +var 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 0000000000..39f13a5880 --- /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 0000000000..fb93daa369 --- /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 0000000000..ccec811b27 --- /dev/null +++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.msl @@ -0,0 +1,25 @@ +#include + +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 0000000000..acad988c5f --- /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 0000000000..5076b67e18 --- /dev/null +++ b/test/tint/builtins/atomicStore/struct/via_ptr_let.wgsl.expected.wgsl @@ -0,0 +1,14 @@ +struct S { + x : i32, + a : atomic, + y : u32, +} + +var wg : S; + +@compute @workgroup_size(1) +fn compute_main() { + let p0 = &(wg); + let p1 = &((*(p0)).a); + atomicStore(p1, 1u); +}