46 lines
1.3 KiB
Plaintext
46 lines
1.3 KiB
Plaintext
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
|
|
|
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using namespace metal;
|
|
|
|
// Implementation of the GLSL mod() function, which is slightly different than Metal fmod()
|
|
template<typename Tx, typename Ty>
|
|
inline Tx mod(Tx x, Ty y)
|
|
{
|
|
return x - y * floor(x / y);
|
|
}
|
|
|
|
// Implementation of signed integer mod accurate to SPIR-V specification
|
|
template<typename Tx, typename Ty>
|
|
inline Tx spvSMod(Tx x, Ty y)
|
|
{
|
|
Tx remainder = x - y * (x / y);
|
|
return select(Tx(remainder + y), remainder, remainder == 0 || (x >= 0) == (y >= 0));
|
|
}
|
|
|
|
struct SSBO
|
|
{
|
|
float4 in_data[1];
|
|
};
|
|
|
|
struct SSBO2
|
|
{
|
|
float4 out_data[1];
|
|
};
|
|
|
|
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
|
|
|
kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _33 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
|
{
|
|
uint ident = gl_GlobalInvocationID.x;
|
|
float4 v = mod(_23.in_data[ident], _33.out_data[ident]);
|
|
_33.out_data[ident] = v;
|
|
uint4 vu = as_type<uint4>(_23.in_data[ident]) % as_type<uint4>(_33.out_data[ident]);
|
|
_33.out_data[ident] = as_type<float4>(vu);
|
|
int4 vi = spvSMod(as_type<int4>(_23.in_data[ident]), as_type<int4>(_33.out_data[ident]));
|
|
_33.out_data[ident] = as_type<float4>(vi);
|
|
}
|
|
|