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