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