28 lines
696 B
Plaintext
28 lines
696 B
Plaintext
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
|
|
|
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using namespace metal;
|
|
|
|
// 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 myBlock
|
|
{
|
|
int a;
|
|
float b[1];
|
|
};
|
|
|
|
kernel void main0(device myBlock& myStorage [[buffer(0)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
|
|
{
|
|
myStorage.a = spvSMod(myStorage.a + 1, 256);
|
|
myStorage.b[gl_LocalInvocationID.x] += 0.0199999995529651641845703125;
|
|
}
|
|
|