How to implement a lock and unlock sequence in a metal shader?

user_19346490 0
Viewed 64

How should I implement a lock/unlock sequence with Compare and Swap using a Metal compute shader.

I’ve tested this sample code but it does not seem to work. For some reason, the threads are not detecting that the lock was released.

Here is a brief explanation of the code below:

The depthFlag is an array of atomic_bools. In this simple example, I simply try to do a lock by comparing the content of depthFlag[1]. I then go ahead and do my operation and once the operation is done, I do an unlock.

As stated above, only one thread is able to do the locking/work/unlocking but the rest of the threads get stuck in the while loop. They NEVER leave. I expect another thread to detect the unlock and go through the sequence.

What am I doing wrong? My knowledge on CAS is limited, so I appreciate any tips.

kernel void testFunction(device float *depthBuffer[[buffer(4)]], device atomic_bool *depthFlag [[buffer(5)]], uint index[[thread_position_in_grid]]){
    
    //lock
    bool expected=false;
while(!atomic_compare_exchange_weak_explicit(&depthFlag[1],&expected,true,memory_order_relaxed,memory_order_relaxed)){
        //wait
        expected=false;

    }

    //Do my operation here

        
    //unlock
    atomic_store_explicit(&depthFlag[1], false, memory_order_relaxed);
        
//barrier
}
compute-shader metal compare-and-swap
1 Answers

You essentially can't use the locking programming model for GPU concurrency. For one, the relaxed memory order model (the only one available) is not suitable for this; for another, you can't guarantee that other threads will make progress between your atomic operations. Your code must always be able to make progress, regardless of what the other threads are doing.

My recommendation is that you use something like the following model instead:

  • Read atomic value to check if another thread has already completed the operation in question.
  • If no other thread has done it yet, perform the operation. (But don't cause any side effects, i.e. don't write to device memory.)
  • Perform an atomic operation to indicate your thread has completed the operation while checking whether another thread got there first. (e.g. compare-and-swap a boolean, but increasing a counter also works)
  • If another thread got there first, don't perform side effects.
  • If your thread "won" and no other thread registered completion, perform your operation's side effects, e.g. do whatever you need to do to write out the result etc.

This works well if there's not much competition, and if the result does not vary depending on which thread performs the operation.

The occasional discarded work should not matter. If there is significant competition, use thread groups; within a thread group, the threads can coordinate which thread will perform which operation. You may still end up with wasted computation from competition between groups. If this is a problem, you may need to change your approach more fundamentally.

If the results of the operation are not deterministic, and the threads all need to proceed using the same result, you will need to change your approach. For example, split your kernels up so any computation which depends on the result of the operation in question runs in a sequentially queued kernel.

user_486600
Related