78 lines
1.7 KiB
Plaintext
78 lines
1.7 KiB
Plaintext
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
|
|
|
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using namespace metal;
|
|
|
|
template<typename T>
|
|
struct spvDescriptor
|
|
{
|
|
T value;
|
|
};
|
|
|
|
template<typename T>
|
|
struct spvDescriptorArray
|
|
{
|
|
spvDescriptorArray(const device spvDescriptor<T>* ptr) : ptr(&ptr->value)
|
|
{
|
|
}
|
|
const device T& operator [] (size_t i) const
|
|
{
|
|
return ptr[i];
|
|
}
|
|
const device T* ptr;
|
|
};
|
|
|
|
struct SSBO
|
|
{
|
|
float4 v[1];
|
|
};
|
|
|
|
struct UBO
|
|
{
|
|
float4 v[1024];
|
|
};
|
|
|
|
struct UBOs
|
|
{
|
|
float4 v;
|
|
};
|
|
|
|
struct SSBOIn
|
|
{
|
|
float4 v[1024];
|
|
};
|
|
|
|
struct SSBOIns
|
|
{
|
|
float4 v;
|
|
};
|
|
|
|
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
|
|
|
|
struct spvDescriptorSetBuffer0
|
|
{
|
|
device SSBO* o [[id(0)]];
|
|
constant UBO* v [[id(1)]];
|
|
spvDescriptor<constant UBOs *> vs [[id(2)]][1] /* unsized array hack */;
|
|
};
|
|
|
|
struct spvDescriptorSetBuffer1
|
|
{
|
|
device SSBOIn* w [[id(0)]];
|
|
spvDescriptor<device SSBOIns *> ws [[id(1)]][1] /* unsized array hack */;
|
|
};
|
|
|
|
kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
|
|
{
|
|
spvDescriptorArray<constant UBOs*> vs {spvDescriptorSet0.vs};
|
|
spvDescriptorArray<device SSBOIns*> ws {spvDescriptorSet1.ws};
|
|
|
|
(*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = (*spvDescriptorSet0.v).v[gl_WorkGroupID.x];
|
|
(*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = vs[gl_WorkGroupID.x]->v;
|
|
(*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = (*spvDescriptorSet1.w).v[gl_WorkGroupID.x];
|
|
(*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = ws[gl_WorkGroupID.x]->v;
|
|
}
|
|
|