Myujth

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;