blob: c3d8f74dcedaedd0e4b688881a50004b90c37eff [file] [log] [blame]
#include <metal_stdlib>
#include <simd/simd.h>
#ifdef __clang__
#pragma clang diagnostic ignored "-Wall"
#endif
using namespace metal;
struct GlobalCounts {
atomic_uint firstHalfCount;
atomic_uint secondHalfCount;
};
struct Inputs {
uint3 sk_LocalInvocationID;
};
struct ssbo {
GlobalCounts globalCounts;
};
struct Globals {
device ssbo* _anonInterface0;
};
struct Threadgroups {
array<atomic_uint, 2> localCounts;
};
kernel void computeMain(uint3 sk_LocalInvocationID [[thread_position_in_threadgroup]], device ssbo& _anonInterface0 [[buffer(0)]]) {
Globals _globals{&_anonInterface0};
(void)_globals;
threadgroup Threadgroups _threadgroups{{}};
(void)_threadgroups;
Inputs _in = { sk_LocalInvocationID };
if (_in.sk_LocalInvocationID.x == 0u) {
atomic_store_explicit(&_threadgroups.localCounts[0], 0u, memory_order_relaxed);
atomic_store_explicit(&_threadgroups.localCounts[1], 0u, memory_order_relaxed);
}
threadgroup_barrier(mem_flags::mem_threadgroup);
uint idx = uint(_in.sk_LocalInvocationID.x < 128u ? 0 : 1);
atomic_fetch_add_explicit(&_threadgroups.localCounts[idx], 1u, memory_order_relaxed);
threadgroup_barrier(mem_flags::mem_threadgroup);
if (_in.sk_LocalInvocationID.x == 0u) {
atomic_fetch_add_explicit(&_globals._anonInterface0->globalCounts.firstHalfCount, atomic_load_explicit(&_threadgroups.localCounts[0], memory_order_relaxed), memory_order_relaxed);
atomic_fetch_add_explicit(&_globals._anonInterface0->globalCounts.secondHalfCount, atomic_load_explicit(&_threadgroups.localCounts[1], memory_order_relaxed), memory_order_relaxed);
}
return;
}