143 lines
4.3 KiB
Plaintext
143 lines
4.3 KiB
Plaintext
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
|
|
|
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using namespace metal;
|
|
|
|
struct SSBO_A
|
|
{
|
|
float data[1];
|
|
};
|
|
|
|
struct UBO_C
|
|
{
|
|
float4 data[1024];
|
|
};
|
|
|
|
struct Registers
|
|
{
|
|
float reg;
|
|
};
|
|
|
|
struct SSBO_B
|
|
{
|
|
uint2 data[1];
|
|
};
|
|
|
|
struct UBO_D
|
|
{
|
|
uint4 data[1024];
|
|
};
|
|
|
|
struct SSBO_BRO
|
|
{
|
|
uint2 data[1];
|
|
};
|
|
|
|
struct SSBO_As
|
|
{
|
|
float data[1];
|
|
};
|
|
|
|
struct UBO_Cs
|
|
{
|
|
float4 data[1024];
|
|
};
|
|
|
|
struct SSBO_Bs
|
|
{
|
|
uint2 data[1024];
|
|
};
|
|
|
|
struct UBO_Ds
|
|
{
|
|
uint4 data[1024];
|
|
};
|
|
|
|
struct SSBO_BsRO
|
|
{
|
|
uint2 data[1024];
|
|
};
|
|
|
|
struct SSBO_E
|
|
{
|
|
float data[1];
|
|
};
|
|
|
|
struct UBO_G
|
|
{
|
|
float4 data[1024];
|
|
};
|
|
|
|
struct SSBO_F
|
|
{
|
|
uint2 data[1];
|
|
};
|
|
|
|
struct UBO_H
|
|
{
|
|
uint4 data[1024];
|
|
};
|
|
|
|
struct SSBO_I
|
|
{
|
|
uint2 data[1];
|
|
};
|
|
|
|
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
|
|
|
|
struct spvDescriptorSetBuffer0
|
|
{
|
|
device SSBO_A* ssbo_a [[id(0)]];
|
|
constant UBO_C* ubo_c [[id(1)]];
|
|
device SSBO_As* ssbo_as [[id(2)]][4];
|
|
constant UBO_Cs* ubo_cs [[id(6)]][4];
|
|
};
|
|
|
|
static inline __attribute__((always_inline))
|
|
void func0(device SSBO_A& ssbo_a, thread uint3& gl_GlobalInvocationID, constant UBO_C& ubo_c, thread uint3& gl_WorkGroupID, constant Registers& v_42, device SSBO_B& ssbo_b, constant UBO_D& ubo_d, const device SSBO_BRO& ssbo_b_readonly)
|
|
{
|
|
ssbo_a.data[gl_GlobalInvocationID.x] = ubo_c.data[gl_WorkGroupID.x].x + v_42.reg;
|
|
ssbo_b.data[gl_GlobalInvocationID.x] = ubo_d.data[gl_WorkGroupID.y].xy + ssbo_b_readonly.data[gl_GlobalInvocationID.x];
|
|
}
|
|
|
|
static inline __attribute__((always_inline))
|
|
void func1(thread uint3& gl_GlobalInvocationID, thread uint3& gl_WorkGroupID, device SSBO_As* const device (&ssbo_as)[4], constant UBO_Cs* const device (&ubo_cs)[4])
|
|
{
|
|
ssbo_as[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x] = ubo_cs[gl_WorkGroupID.x]->data[0].x;
|
|
}
|
|
|
|
static inline __attribute__((always_inline))
|
|
void func2(thread uint3& gl_GlobalInvocationID, thread uint3& gl_WorkGroupID, device SSBO_Bs* const device (&ssbo_bs)[4], constant UBO_Ds* const device (&ubo_ds)[4], const device SSBO_BsRO* const device (&ssbo_bs_readonly)[4])
|
|
{
|
|
ssbo_bs[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x] = ubo_ds[gl_WorkGroupID.x]->data[0].xy + ssbo_bs_readonly[gl_WorkGroupID.x]->data[gl_GlobalInvocationID.x];
|
|
}
|
|
|
|
static inline __attribute__((always_inline))
|
|
void func3(thread uint3& gl_GlobalInvocationID, thread uint3& gl_WorkGroupID, device SSBO_E& ssbo_e, constant UBO_G& ubo_g, device SSBO_F& ssbo_f, constant UBO_H& ubo_h, const device SSBO_I& ssbo_i)
|
|
{
|
|
ssbo_e.data[gl_GlobalInvocationID.x] = ubo_g.data[gl_WorkGroupID.x].x;
|
|
ssbo_f.data[gl_GlobalInvocationID.x] = ubo_h.data[gl_WorkGroupID.y].xy + ssbo_i.data[gl_GlobalInvocationID.x];
|
|
}
|
|
|
|
kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant Registers& v_42 [[buffer(1)]], device void* spvBufferAliasSet2Binding0 [[buffer(2)]], constant void* spvBufferAliasSet2Binding1 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
|
|
{
|
|
device auto& ssbo_e = *(device SSBO_E*)spvBufferAliasSet2Binding0;
|
|
constant auto& ubo_g = *(constant UBO_G*)spvBufferAliasSet2Binding1;
|
|
device auto& ssbo_f = *(device SSBO_F*)spvBufferAliasSet2Binding0;
|
|
constant auto& ubo_h = *(constant UBO_H*)spvBufferAliasSet2Binding1;
|
|
const device auto& ssbo_i = *(const device SSBO_I*)spvBufferAliasSet2Binding0;
|
|
device auto& ssbo_b = (device SSBO_B&)(*spvDescriptorSet0.ssbo_a);
|
|
constant auto& ubo_d = (constant UBO_D&)(*spvDescriptorSet0.ubo_c);
|
|
const device auto& ssbo_b_readonly = (const device SSBO_BRO&)(*spvDescriptorSet0.ssbo_a);
|
|
const device auto& ssbo_bs = (device SSBO_Bs* const device (&)[4])spvDescriptorSet0.ssbo_as;
|
|
const device auto& ubo_ds = (constant UBO_Ds* const device (&)[4])spvDescriptorSet0.ubo_cs;
|
|
const device auto& ssbo_bs_readonly = (const device SSBO_BsRO* const device (&)[4])spvDescriptorSet0.ssbo_as;
|
|
func0((*spvDescriptorSet0.ssbo_a), gl_GlobalInvocationID, (*spvDescriptorSet0.ubo_c), gl_WorkGroupID, v_42, ssbo_b, ubo_d, ssbo_b_readonly);
|
|
func1(gl_GlobalInvocationID, gl_WorkGroupID, spvDescriptorSet0.ssbo_as, spvDescriptorSet0.ubo_cs);
|
|
func2(gl_GlobalInvocationID, gl_WorkGroupID, ssbo_bs, ubo_ds, ssbo_bs_readonly);
|
|
func3(gl_GlobalInvocationID, gl_WorkGroupID, ssbo_e, ubo_g, ssbo_f, ubo_h, ssbo_i);
|
|
}
|
|
|