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