1#pragma clang diagnostic ignored "-Wmissing-prototypes"
2
3#include <metal_stdlib>
4#include <simd/simd.h>
5
6using namespace metal;
7
8struct foo
9{
10    int a[128];
11    uint b;
12    float2 c;
13};
14
15struct bar
16{
17    int d;
18};
19
20struct baz
21{
22    int e[128];
23};
24
25static inline __attribute__((always_inline))
26device int* select_buffer(device foo& buf, device baz& buf2, constant bar& cb)
27{
28    return (cb.d != 0) ? &buf.a[0u] : &buf2.e[0u];
29}
30
31static inline __attribute__((always_inline))
32device int* select_buffer_null(device foo& buf, constant bar& cb)
33{
34    return (cb.d != 0) ? &buf.a[0u] : nullptr;
35}
36
37static inline __attribute__((always_inline))
38threadgroup int* select_tgsm(constant bar& cb, threadgroup int (&tgsm)[128])
39{
40    return (cb.d != 0) ? &tgsm[0u] : nullptr;
41}
42
43kernel void main0(device foo& buf [[buffer(0)]], constant bar& cb [[buffer(1)]], device baz& buf2 [[buffer(2)]])
44{
45    threadgroup int tgsm[128];
46    device int* sbuf = select_buffer(buf, buf2, cb);
47    device int* sbuf2 = select_buffer_null(buf, cb);
48    threadgroup int* stgsm = select_tgsm(cb, tgsm);
49    threadgroup int* cur = stgsm;
50    device int* _73;
51    _73 = &buf.a[0u];
52    threadgroup int* _76;
53    int _77;
54    for (;;)
55    {
56        _76 = cur;
57        _77 = *_73;
58        if (_77 != 0)
59        {
60            int _81 = *_76;
61            int _82 = _77 + _81;
62            *_73 = _82;
63            *_76 = _82;
64            cur = &_76[1u];
65            _73 = &_73[1u];
66            continue;
67        }
68        else
69        {
70            break;
71        }
72    }
73}
74
75