Ndtyjky

I've been trying to put together a falling sand type game in Metal, that uses a compute shader to do the work. My initial work looked ok, until I ran into the following problem:

In one update, both particle 1A and 1B could flow into space T.

Sequential

In a sequential loop, I could understand how I would simply move 1A into T, then when it came to evaluate 1B, I could see that the space was already filled.

Parallel

My problem is that with the GPU, The code to check 1A and 1B are both running on different threads, so in one update loop, they can both move to position T, essentially killing one of them.

My question is then, can I make a parallel solution to this, or can it only be solved sequentially?

Here's my calculation code, running a GPU thread for every pixel in a 640 by 480 grid:

struct Particle {

    float4 color;

    int live;

    int particleType;

};





kernel void calculateFrame(

                           device Particle *src [[buffer(1)]],

                           device Particle *dst [[buffer(2)]],

                           uint2 threadgroup_position_in_grid   [[ threadgroup_position_in_grid ]],

                           uint2 thread_position_in_threadgroup [[ thread_position_in_threadgroup ]],

                           uint2 threads_per_threadgroup        [[ threads_per_threadgroup ]]) {



    uint2 gid = threadgroup_position_in_grid;

    int index = (gid * 640) + gid;

    Particle particle = src[index];



    Particle bottom = src[toLinear( gid + uint2( 0, 1))]; //bottom

    Particle bottomLeft = src[toLinear( gid + uint2(-1, 1))]; //bottom left

    Particle bottomRight = src[toLinear( gid + uint2( 1, 1))]; //bottom right



    if (particle.live == 1 && particle.particleType == SAND) { // move down



        if (bottom.live == 0) {



            dst[index].live = 0;



            dst[toLinear( gid + uint2( 0, 1))].particleType = particle.particleType;

            dst[toLinear( gid + uint2( 0, 1))].color = particle.color;

            dst[toLinear( gid + uint2( 0, 1))].live = 1;



        }

        else if (bottomLeft.live == 0) { // move down to the left if clear

            dst[index].live = 0;



            dst[toLinear( gid + uint2( -1, 1))].particleType = particle.particleType;

            dst[toLinear( gid + uint2( -1, 1))].color = red;

            dst[toLinear( gid + uint2( -1, 1))].live = 1;





        }

        else if (bottomRight.live == 0) { // move down to the right if clear



            dst[index].live = 0;



            dst[toLinear( gid + uint2( 1, 1))].particleType = particle.particleType;

            dst[toLinear( gid + uint2( 1, 1))].color = red;

            dst[toLinear( gid + uint2( 1, 1))].live = 1;

        }

        else { // cant move

            dst[index] = src[index];

        }



    }



    if (particle.particleType == WALL) {

        dst[index] = particle;

        dst[index].live = 1;

    }



    if (particle.particleType == ERASER) {

        dst[index].color = blue;

        dst[index].live = 0;

    }

}