1#pragma clang diagnostic ignored "-Wmissing-prototypes"
2#pragma clang diagnostic ignored "-Wmissing-braces"
3
4#include <metal_stdlib>
5#include <simd/simd.h>
6
7using namespace metal;
8
9template<typename T, size_t Num>
10struct spvUnsafeArray
11{
12    T elements[Num ? Num : 1];
13    
14    thread T& operator [] (size_t pos) thread
15    {
16        return elements[pos];
17    }
18    constexpr const thread T& operator [] (size_t pos) const thread
19    {
20        return elements[pos];
21    }
22    
23    device T& operator [] (size_t pos) device
24    {
25        return elements[pos];
26    }
27    constexpr const device T& operator [] (size_t pos) const device
28    {
29        return elements[pos];
30    }
31    
32    constexpr const constant T& operator [] (size_t pos) const constant
33    {
34        return elements[pos];
35    }
36    
37    threadgroup T& operator [] (size_t pos) threadgroup
38    {
39        return elements[pos];
40    }
41    constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
42    {
43        return elements[pos];
44    }
45};
46
47struct Data
48{
49    float a;
50    float b;
51};
52
53constant float X_tmp [[function_constant(0)]];
54constant float X = is_function_constant_defined(X_tmp) ? X_tmp : 4.0;
55
56struct Data_1
57{
58    float a;
59    float b;
60};
61
62struct SSBO
63{
64    Data_1 outdata[1];
65};
66
67constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(2u, 1u, 1u);
68
69kernel void main0(device SSBO& _53 [[buffer(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
70{
71    spvUnsafeArray<Data, 2> _25 = spvUnsafeArray<Data, 2>({ Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } });
72    
73    spvUnsafeArray<Data, 2> _31 = spvUnsafeArray<Data, 2>({ Data{ X, 2.0 }, Data{ 3.0, 5.0 } });
74    spvUnsafeArray<Data, 2> data2;
75    data2 = _31;
76    _53.outdata[gl_WorkGroupID.x].a = _25[gl_LocalInvocationID.x].a + data2[gl_LocalInvocationID.x].a;
77    _53.outdata[gl_WorkGroupID.x].b = _25[gl_LocalInvocationID.x].b + data2[gl_LocalInvocationID.x].b;
78}
79
80