43 lines
1.3 KiB
Plaintext
43 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]])
|
|
{
|
|
float4 _29 = _23.in_data[gl_GlobalInvocationID.x];
|
|
_33.out_data[gl_GlobalInvocationID.x] = mod(_29, _33.out_data[gl_GlobalInvocationID.x]);
|
|
_33.out_data[gl_GlobalInvocationID.x] = as_type<float4>(as_type<uint4>(_29) % as_type<uint4>(_33.out_data[gl_GlobalInvocationID.x]));
|
|
_33.out_data[gl_GlobalInvocationID.x] = as_type<float4>(spvSMod(as_type<int4>(_29), as_type<int4>(_33.out_data[gl_GlobalInvocationID.x])));
|
|
}
|
|
|