1#include <metal_stdlib> 2#include <simd/simd.h> 3#ifdef __clang__ 4#pragma clang diagnostic ignored "-Wall" 5#endif 6using namespace metal; 7struct Inputs { 8 uint3 sk_LocalInvocationID; 9}; 10struct ssbo { 11 atomic_uint globalCounter; 12}; 13struct Globals { 14 device ssbo* _anonInterface0; 15}; 16struct Threadgroups { 17 atomic_uint localCounter; 18}; 19kernel void computeMain(uint3 sk_LocalInvocationID [[thread_position_in_threadgroup]], device ssbo& _anonInterface0 [[buffer(0)]]) { 20 Globals _globals{&_anonInterface0}; 21 (void)_globals; 22 threadgroup Threadgroups _threadgroups{{}}; 23 (void)_threadgroups; 24 Inputs _in = { sk_LocalInvocationID }; 25 if (_in.sk_LocalInvocationID.x == 0u) { 26 atomic_store_explicit(&_threadgroups.localCounter, 0u, memory_order_relaxed); 27 } 28 threadgroup_barrier(mem_flags::mem_threadgroup); 29 atomic_fetch_add_explicit(&_threadgroups.localCounter, 1u, memory_order_relaxed); 30 threadgroup_barrier(mem_flags::mem_threadgroup); 31 if (_in.sk_LocalInvocationID.x == 0u) { 32 atomic_fetch_add_explicit(&_globals._anonInterface0->globalCounter, atomic_load_explicit(&_threadgroups.localCounter, memory_order_relaxed), memory_order_relaxed); 33 } 34 return; 35} 36