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>
This commit is contained in:
Ben Clayton
Dawn LUCI CQ
parent
7289bca018
commit
bd8449f37d
|
|
@ -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 @@ 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<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 @@ 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<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;
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
||||
|
|
@ -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
|
||||
|
|
@ -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);
|
||||
}
|
||||
|
|
@ -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 @@ 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<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 @@ 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<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;
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
||||
|
|
@ -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
|
||||
|
|
@ -0,0 +1,6 @@
|
|||
var<workgroup> wg : array<atomic<u32>, 4>;
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn compute_main() {
|
||||
atomicStore(&(wg[1]), 1u);
|
||||
}
|
||||
|
|
@ -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 @@ 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<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 @@ 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<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;
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
||||
|
|
@ -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
|
||||
|
|
@ -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);
|
||||
}
|
||||
|
|
@ -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 @@ 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<S_atomic, 10>* 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<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 @@ 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<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;
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
||||
|
|
@ -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
|
||||
|
|
@ -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);
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
||||
|
|
@ -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
|
||||
|
|
@ -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);
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
||||
|
|
@ -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
|
||||
|
|
@ -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);
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
||||
|
|
@ -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
|
||||
|
|
@ -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);
|
||||
}
|
||||
|
|
@ -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 @@ 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;
|
||||
|
|
|
|||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
||||
|
|
@ -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
|
||||
|
|
@ -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);
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
|
@ -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;
|
||||
}
|
||||
|
||||
|
|
@ -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
|
||||
|
|
@ -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);
|
||||
}
|
||||
Loading…
Reference in New Issue