#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;
};

static inline __attribute__((always_inline))
device foo* select_buffer(device foo& a, constant bar& cb)
{
    return (cb.d != 0) ? &a : nullptr;
}

static inline __attribute__((always_inline))
thread uint3* select_input(thread uint3& gl_GlobalInvocationID, thread uint3& gl_LocalInvocationID, constant bar& cb)
{
    return (cb.d != 0) ? &gl_GlobalInvocationID : &gl_LocalInvocationID;
}

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]])
{
    device foo* _46 = select_buffer(buf, cb);
    device foo* _45 = _46;
    device foo* _48 = _45;
    device int* _52;
    device int* _55;
    _52 = &_48->a[0u];
    _55 = &buf.a[0u];
    int _57;
    int _58;
    for (;;)
    {
        _57 = *_52;
        _58 = *_55;
        if (_57 != _58)
        {
            int _66 = (_57 + _58) + int((*select_input(gl_GlobalInvocationID, gl_LocalInvocationID, cb)).x);
            *_52 = _66;
            *_55 = _66;
            _52 = &_52[1u];
            _55 = &_55[1u];
            continue;
        }
        else
        {
            break;
        }
    }
}