xref: /aosp_15_r20/external/skia/tests/sksl/compute/AtomicOperationsOverArrayAndStruct.metal (revision c8dee2aa9b3f27cf6c858bd81872bdeb2c07ed17)
1*c8dee2aaSAndroid Build Coastguard Worker#include <metal_stdlib>
2*c8dee2aaSAndroid Build Coastguard Worker#include <simd/simd.h>
3*c8dee2aaSAndroid Build Coastguard Worker#ifdef __clang__
4*c8dee2aaSAndroid Build Coastguard Worker#pragma clang diagnostic ignored "-Wall"
5*c8dee2aaSAndroid Build Coastguard Worker#endif
6*c8dee2aaSAndroid Build Coastguard Workerusing namespace metal;
7*c8dee2aaSAndroid Build Coastguard Workerstruct GlobalCounts {
8*c8dee2aaSAndroid Build Coastguard Worker    atomic_uint firstHalfCount;
9*c8dee2aaSAndroid Build Coastguard Worker    atomic_uint secondHalfCount;
10*c8dee2aaSAndroid Build Coastguard Worker};
11*c8dee2aaSAndroid Build Coastguard Workerstruct Inputs {
12*c8dee2aaSAndroid Build Coastguard Worker    uint3 sk_LocalInvocationID;
13*c8dee2aaSAndroid Build Coastguard Worker};
14*c8dee2aaSAndroid Build Coastguard Workerstruct ssbo {
15*c8dee2aaSAndroid Build Coastguard Worker    GlobalCounts globalCounts;
16*c8dee2aaSAndroid Build Coastguard Worker};
17*c8dee2aaSAndroid Build Coastguard Workerstruct Globals {
18*c8dee2aaSAndroid Build Coastguard Worker    device ssbo* _anonInterface0;
19*c8dee2aaSAndroid Build Coastguard Worker};
20*c8dee2aaSAndroid Build Coastguard Workerstruct Threadgroups {
21*c8dee2aaSAndroid Build Coastguard Worker    array<atomic_uint, 2> localCounts;
22*c8dee2aaSAndroid Build Coastguard Worker};
23*c8dee2aaSAndroid Build Coastguard Workerkernel void computeMain(uint3 sk_LocalInvocationID [[thread_position_in_threadgroup]], device ssbo& _anonInterface0 [[buffer(0)]]) {
24*c8dee2aaSAndroid Build Coastguard Worker    Globals _globals{&_anonInterface0};
25*c8dee2aaSAndroid Build Coastguard Worker    (void)_globals;
26*c8dee2aaSAndroid Build Coastguard Worker    threadgroup Threadgroups _threadgroups{{}};
27*c8dee2aaSAndroid Build Coastguard Worker    (void)_threadgroups;
28*c8dee2aaSAndroid Build Coastguard Worker    Inputs _in = { sk_LocalInvocationID };
29*c8dee2aaSAndroid Build Coastguard Worker    if (_in.sk_LocalInvocationID.x == 0u) {
30*c8dee2aaSAndroid Build Coastguard Worker        atomic_store_explicit(&_threadgroups.localCounts[0], 0u, memory_order_relaxed);
31*c8dee2aaSAndroid Build Coastguard Worker        atomic_store_explicit(&_threadgroups.localCounts[1], 0u, memory_order_relaxed);
32*c8dee2aaSAndroid Build Coastguard Worker    }
33*c8dee2aaSAndroid Build Coastguard Worker    threadgroup_barrier(mem_flags::mem_threadgroup);
34*c8dee2aaSAndroid Build Coastguard Worker    uint idx = uint(_in.sk_LocalInvocationID.x < 128u ? 0 : 1);
35*c8dee2aaSAndroid Build Coastguard Worker    atomic_fetch_add_explicit(&_threadgroups.localCounts[idx], 1u, memory_order_relaxed);
36*c8dee2aaSAndroid Build Coastguard Worker    threadgroup_barrier(mem_flags::mem_threadgroup);
37*c8dee2aaSAndroid Build Coastguard Worker    if (_in.sk_LocalInvocationID.x == 0u) {
38*c8dee2aaSAndroid Build Coastguard Worker        atomic_fetch_add_explicit(&_globals._anonInterface0->globalCounts.firstHalfCount, atomic_load_explicit(&_threadgroups.localCounts[0], memory_order_relaxed), memory_order_relaxed);
39*c8dee2aaSAndroid Build Coastguard Worker        atomic_fetch_add_explicit(&_globals._anonInterface0->globalCounts.secondHalfCount, atomic_load_explicit(&_threadgroups.localCounts[1], memory_order_relaxed), memory_order_relaxed);
40*c8dee2aaSAndroid Build Coastguard Worker    }
41*c8dee2aaSAndroid Build Coastguard Worker    return;
42*c8dee2aaSAndroid Build Coastguard Worker}
43