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