#pragma clang diagnostic ignored "-Wmissing-prototypes"

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

using namespace metal;

struct foo
{
    int a[128];
    uint b;
    float2 c;
};

struct bar
{
    int d;
};

struct baz
{
    int e[128];
};

static inline __attribute__((always_inline))
device int* select_buffer(device foo& buf, device baz& buf2, constant bar& cb)
{
    return (cb.d != 0) ? &buf.a[0u] : &buf2.e[0u];
}

static inline __attribute__((always_inline))
device int* select_buffer_null(device foo& buf, constant bar& cb)
{
    return (cb.d != 0) ? &buf.a[0u] : nullptr;
}

static inline __attribute__((always_inline))
threadgroup int* select_tgsm(constant bar& cb, threadgroup int (&tgsm)[128])
{
    return (cb.d != 0) ? &tgsm[0u] : nullptr;
}

kernel void main0(device foo& buf [[buffer(0)]], constant bar& cb [[buffer(1)]], device baz& buf2 [[buffer(2)]])
{
    threadgroup int tgsm[128];
    device int* sbuf = select_buffer(buf, buf2, cb);
    device int* sbuf2 = select_buffer_null(buf, cb);
    threadgroup int* stgsm = select_tgsm(cb, tgsm);
    threadgroup int* cur = stgsm;
    device int* _78;
    _78 = &buf.a[0u];
    threadgroup int* _81;
    int _82;
    for (;;)
    {
        _81 = cur;
        _82 = *_78;
        if (_82 != 0)
        {
            int _86 = *_81;
            int _87 = _82 + _86;
            *_78 = _87;
            *_81 = _87;
            cur = &_81[1u];
            _78 = &_78[1u];
            continue;
        }
        else
        {
            break;
        }
    }
}