fa5b206d97
On MSL, the compiler refuses to allow access chains into a normal vector type. What happens in practice instead is a read-modify-write where a vector type is loaded, modified and written back. The workaround is to convert a vector into a pointer-to-scalar before the access chain continues to add the scalar index.
32 lines
882 B
Plaintext
32 lines
882 B
Plaintext
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using namespace metal;
|
|
|
|
struct SSBO1
|
|
{
|
|
short4 outputs[1];
|
|
};
|
|
|
|
struct SSBO0
|
|
{
|
|
int4 inputs[1];
|
|
};
|
|
|
|
struct UBO
|
|
{
|
|
half4 const0;
|
|
};
|
|
|
|
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
|
|
|
kernel void main0(device SSBO1& _21 [[buffer(0)]], device SSBO0& _29 [[buffer(1)]], constant UBO& _40 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
|
{
|
|
uint ident = gl_GlobalInvocationID.x;
|
|
short2 _47 = as_type<short2>(((device int*)&_29.inputs[ident])[0u]) + as_type<short2>(_40.const0.xy);
|
|
_21.outputs[ident] = short4(_47.x, _47.y, _21.outputs[ident].z, _21.outputs[ident].w);
|
|
short2 _66 = short2(as_type<ushort2>(uint(((device int*)&_29.inputs[ident])[1u])) - as_type<ushort2>(_40.const0.zw));
|
|
_21.outputs[ident] = short4(_21.outputs[ident].x, _21.outputs[ident].y, _66.x, _66.y);
|
|
}
|
|
|