1#include <metal_stdlib> 2#include <simd/simd.h> 3#ifdef __clang__ 4#pragma clang diagnostic ignored "-Wall" 5#endif 6using namespace metal; 7struct S { 8 atomic_uint structMemberAtomic; 9 array<atomic_uint, 2> structMemberAtomicArray; 10}; 11struct NestedS { 12 S nestedStructWithAtomicMember; 13}; 14struct Inputs { 15}; 16struct ssbo { 17 atomic_uint ssboAtomic; 18 array<atomic_uint, 2> ssboAtomicArray; 19 S ssboStructWithAtomicMember; 20 array<S, 2> ssboStructWithAtomicMemberArray; 21 NestedS ssboNestedStructWithAtomicMember; 22}; 23struct Globals { 24 device ssbo* _anonInterface0; 25}; 26struct Threadgroups { 27 atomic_uint wgAtomic; 28 array<atomic_uint, 2> wgAtomicArray; 29 NestedS wgNestedStructWithAtomicMember; 30}; 31kernel void computeMain(device ssbo& _anonInterface0 [[buffer(0)]]) { 32 Globals _globals{&_anonInterface0}; 33 (void)_globals; 34 threadgroup Threadgroups _threadgroups{{}, {}, {}}; 35 (void)_threadgroups; 36 Inputs _in = { }; 37 atomic_fetch_add_explicit(&_threadgroups.wgAtomicArray[1], atomic_load_explicit(&_threadgroups.wgAtomic, memory_order_relaxed), memory_order_relaxed); 38 atomic_fetch_add_explicit(&_threadgroups.wgAtomicArray[0], atomic_load_explicit(&_threadgroups.wgAtomicArray[1], memory_order_relaxed), memory_order_relaxed); 39 atomic_fetch_add_explicit(&_threadgroups.wgNestedStructWithAtomicMember.nestedStructWithAtomicMember.structMemberAtomic, 1u, memory_order_relaxed); 40 return; 41} 42