#pragma clang diagnostic ignored "-Wmissing-prototypes" #pragma clang diagnostic ignored "-Wmissing-braces" #include #include using namespace metal; template 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 Data { float a; float b; }; constant float X_tmp [[function_constant(0)]]; constant float X = is_function_constant_defined(X_tmp) ? X_tmp : 4.0; struct Data_1 { float a; float b; }; struct SSBO { Data_1 outdata[1]; }; constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(2u, 1u, 1u); constant spvUnsafeArray _25 = spvUnsafeArray({ Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } }); static inline __attribute__((always_inline)) Data combine(thread const Data& a, thread const Data& b) { return Data{ a.a + b.a, a.b + b.b }; } kernel void main0(device SSBO& _53 [[buffer(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]]) { spvUnsafeArray data = spvUnsafeArray({ Data{ 1.0, 2.0 }, Data{ 3.0, 4.0 } }); spvUnsafeArray _31 = spvUnsafeArray({ Data{ X, 2.0 }, Data{ 3.0, 5.0 } }); spvUnsafeArray data2 = _31; Data param = data[gl_LocalInvocationID.x]; Data param_1 = data2[gl_LocalInvocationID.x]; Data _73 = combine(param, param_1); _53.outdata[gl_WorkGroupID.x].a = _73.a; _53.outdata[gl_WorkGroupID.x].b = _73.b; }