vello/tests/shader/gen/linkedlist.msl

#pragma clang diagnostic ignored "-Wunused-variable"

#include <metal_stdlib>
#include <simd/simd.h>
#include <metal_atomic>

using namespace metal;

struct MemBuf
{
    uint mem[1];
};

constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(256u, 1u, 1u);

kernel void main0(device MemBuf& _56 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
    uint rng = gl_GlobalInvocationID.x + 1u;
    for (uint i = 0u; i < 100u; i++)
    {
        rng ^= (rng << uint(13));
        rng ^= (rng >> uint(17));
        rng ^= (rng << uint(5));
        uint bucket = rng % 65536u;
        if (bucket != 0u)
        {
            uint _61 = atomic_fetch_add_explicit((device atomic_uint*)&_56.mem[0], 2u, memory_order_relaxed);
            uint alloc = _61 + 65536u;
            uint _67 = atomic_exchange_explicit((device atomic_uint*)&_56.mem[bucket], alloc, memory_order_relaxed);
            uint old = _67;
            _56.mem[alloc] = old;
            _56.mem[alloc + 1u] = gl_GlobalInvocationID.x;
        }
    }
}