// language: metal1.0
#include <metal_stdlib>
#include <simd/simd.h>

using metal::uint;

struct type_1 {
    uint inner[512];
};
struct type_3 {
    metal::atomic_int inner[8];
};
struct type_4 {
    type_3 inner[8];
};
struct WStruct {
    type_1 arr;
    metal::atomic_int atom;
    type_4 atom_arr;
};

kernel void main_(
  metal::uint3 __local_invocation_id [[thread_position_in_threadgroup]]
, threadgroup WStruct& w_mem
, device type_1& output [[buffer(0)]]
) {
    if (metal::all(__local_invocation_id == metal::uint3(0u))) {
        w_mem.arr = {};
        metal::atomic_store_explicit(&w_mem.atom, 0, metal::memory_order_relaxed);
        for (int __i0 = 0; __i0 < 8; __i0++) {
            for (int __i1 = 0; __i1 < 8; __i1++) {
                metal::atomic_store_explicit(&w_mem.atom_arr.inner[__i0].inner[__i1], 0, metal::memory_order_relaxed);
            }
        }
    }
    metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
    type_1 _e3 = w_mem.arr;
    output = _e3;
    return;
}
