2018-05-02 08:01:08 +00:00
|
|
|
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
|
|
|
|
2018-05-02 07:38:41 +00:00
|
|
|
#include <metal_stdlib>
|
|
|
|
#include <simd/simd.h>
|
|
|
|
|
|
|
|
using namespace metal;
|
|
|
|
|
|
|
|
struct SSBO0
|
|
|
|
{
|
|
|
|
float4 as[1];
|
|
|
|
};
|
|
|
|
|
2018-09-27 09:10:22 +00:00
|
|
|
struct SSBO1
|
|
|
|
{
|
|
|
|
float4 bs[1];
|
|
|
|
};
|
|
|
|
|
2019-08-27 11:16:16 +00:00
|
|
|
template<typename T, uint A>
|
|
|
|
void spvArrayCopyFromConstantToStack1(thread T (&dst)[A], constant T (&src)[A])
|
2018-05-02 08:01:08 +00:00
|
|
|
{
|
2019-08-27 11:16:16 +00:00
|
|
|
for (uint i = 0; i < A; i++)
|
|
|
|
{
|
|
|
|
dst[i] = src[i];
|
|
|
|
}
|
2018-05-02 08:01:08 +00:00
|
|
|
}
|
|
|
|
|
2019-08-27 11:16:16 +00:00
|
|
|
template<typename T, uint A>
|
|
|
|
void spvArrayCopyFromConstantToThreadGroup1(threadgroup T (&dst)[A], constant T (&src)[A])
|
2018-05-02 08:01:08 +00:00
|
|
|
{
|
2019-08-27 11:16:16 +00:00
|
|
|
for (uint i = 0; i < A; i++)
|
|
|
|
{
|
|
|
|
dst[i] = src[i];
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
template<typename T, uint A>
|
|
|
|
void spvArrayCopyFromStackToStack1(thread T (&dst)[A], thread const T (&src)[A])
|
|
|
|
{
|
|
|
|
for (uint i = 0; i < A; i++)
|
|
|
|
{
|
|
|
|
dst[i] = src[i];
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
template<typename T, uint A>
|
|
|
|
void spvArrayCopyFromStackToThreadGroup1(threadgroup T (&dst)[A], thread const T (&src)[A])
|
|
|
|
{
|
|
|
|
for (uint i = 0; i < A; i++)
|
|
|
|
{
|
|
|
|
dst[i] = src[i];
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
template<typename T, uint A>
|
|
|
|
void spvArrayCopyFromThreadGroupToStack1(thread T (&dst)[A], threadgroup const T (&src)[A])
|
|
|
|
{
|
|
|
|
for (uint i = 0; i < A; i++)
|
|
|
|
{
|
|
|
|
dst[i] = src[i];
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
template<typename T, uint A>
|
|
|
|
void spvArrayCopyFromThreadGroupToThreadGroup1(threadgroup T (&dst)[A], threadgroup const T (&src)[A])
|
|
|
|
{
|
|
|
|
for (uint i = 0; i < A; i++)
|
|
|
|
{
|
|
|
|
dst[i] = src[i];
|
|
|
|
}
|
2018-05-02 08:01:08 +00:00
|
|
|
}
|
|
|
|
|
2018-09-27 09:10:22 +00:00
|
|
|
kernel void main0(device SSBO0& _16 [[buffer(0)]], device SSBO1& _32 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
|
2018-05-02 07:38:41 +00:00
|
|
|
{
|
2018-09-27 09:10:22 +00:00
|
|
|
float4 _37[2] = { _16.as[gl_GlobalInvocationID.x], _32.bs[gl_GlobalInvocationID.x] };
|
2018-05-02 08:01:08 +00:00
|
|
|
float4 values[2];
|
2019-08-27 11:16:16 +00:00
|
|
|
spvArrayCopyFromStackToStack1(values, _37);
|
2018-05-02 08:01:08 +00:00
|
|
|
_16.as[0] = values[gl_LocalInvocationIndex];
|
2018-09-27 09:10:22 +00:00
|
|
|
_32.bs[1] = float4(40.0);
|
2018-05-02 07:38:41 +00:00
|
|
|
}
|
|
|
|
|