#pragma clang diagnostic ignored "-Wmissing-prototypes"
#pragma clang diagnostic ignored "-Wmissing-braces"

#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

template<typename T, size_t Num>
struct spvUnsafeArray
{
    T elements[Num ? Num : 1];
    
    thread T& operator [] (size_t pos) thread
    {
        return elements[pos];
    }
    constexpr const thread T& operator [] (size_t pos) const thread
    {
        return elements[pos];
    }
    
    device T& operator [] (size_t pos) device
    {
        return elements[pos];
    }
    constexpr const device T& operator [] (size_t pos) const device
    {
        return elements[pos];
    }
    
    constexpr const constant T& operator [] (size_t pos) const constant
    {
        return elements[pos];
    }
    
    threadgroup T& operator [] (size_t pos) threadgroup
    {
        return elements[pos];
    }
    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
    {
        return elements[pos];
    }
};

struct foo
{
    spvUnsafeArray<int, 128> a;
    uint b;
    float2 c;
};

struct bar
{
    int d;
};

kernel void main0(device foo& buf [[buffer(0)]], constant bar& cb [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
{
    bool _70 = cb.d != 0;
    device foo* _71 = _70 ? &buf : nullptr;
    device foo* _67 = _71;
    device foo* _45 = _71;
    thread uint3* _77 = _70 ? &gl_GlobalInvocationID : &gl_LocalInvocationID;
    thread uint3* _73 = _77;
    device int* _52;
    device int* _55;
    _52 = &_71->a[0u];
    _55 = &buf.a[0u];
    int _57;
    int _58;
    for (;;)
    {
        _57 = *_52;
        _58 = *_55;
        if (_57 != _58)
        {
            int _66 = (_57 + _58) + int((*_77).x);
            *_52 = _66;
            *_55 = _66;
            _52 = &_52[1u];
            _55 = &_55[1u];
            continue;
        }
        else
        {
            break;
        }
    }
}