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.
24 lines
486 B
Plaintext
24 lines
486 B
Plaintext
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using namespace metal;
|
|
|
|
struct SSBO
|
|
{
|
|
float4 out_data[1];
|
|
};
|
|
|
|
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
|
|
|
|
kernel void main0(device SSBO& _27 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
|
{
|
|
float4 v;
|
|
v.x = 10.0;
|
|
v.y = 30.0;
|
|
v.z = 70.0;
|
|
v.w = 90.0;
|
|
_27.out_data[gl_GlobalInvocationID.x] = v;
|
|
((device float*)&_27.out_data[gl_GlobalInvocationID.x])[1u] = 20.0;
|
|
}
|
|
|