CUDA: How to atomically execute these instructions (4 addresses)

Question

I'm implementing an algorithm in Cuda that needs to perform the following steps:

Given an array x (in shared mem) and some device function f,

• Select a pair of indices (i,j) to x (randomly).
• Calculate y = f(x[i], x[i - 1], x[j], x[j + 1]).
• Based on y decide whether to exchange the positions of x[i] and x[j].

The problem is that the function f depends on 4 values in shared memory, all of which have to be guaranteed to remain unchanged until after the swap.

For a minute I figured this could be the body of a critical section, but I don't see how I could use a single lock-address to lock 4 variables. The main problem, I think, is that when some thread is working on (i,j), other threads are not allowed to work on any pair (k,l) where k or l are any of {i, i-1, j, j+1}.

EDIT

Right after posting, an idea came to mind... Would it be possible to have a cascade of locks? First lock x[i], if that succeeds lock x[i-1], etc for all 4 values. Only if the final lock succeeds, process the above mentioned steps. I'll go experiment and keep this question open to other suggestions.

Answer 1

CUDA is very lock-unfriendly and critical-section-unfriendly :) One of many reasons is that it operates in a 32-wide SIMD mode. This may cause unexpected deadlocks because of it. Consider for example:

__shared__ int crit;
crit = 0;
__syncthreads();
int old;
do {
    old = atomicCas(&crit, 0, 1);
} while (old==1);
//critical section
crit = 0;

The intetion is that threads actively wait in the do-while loop. Only one thread exists the loop at a time, perform action in the critical secton and then resets crit to 0. However, in CUDA, a warp scheduler will always give priority to 31 threads in the loop over the 1 thread that exits. Because warps operate in SIMD, the thread in the critical section never executes and you get an unexpected deadlock.

---

For that reason I would strongly suggest trying to avoid critical sections completely.

Now, I don't know the details of your algorithm. I assume that you have some "master" for/while loop and in each iteration you pick a random pair for a possible swap.

You say the collisions don't happen often. If it does, could you just choose to drop one of the conflicting pairs completely, instead of waiting for it to succeed?

If that is something you would accept, then just detecting the collision would be the problem, not the action that you take afterwards. To detect collisions you could for example:

• After each thread comes up with a pair candidate, sort the pair indices and then check the values held by the neighbours.

• Have a flag array f of the same size as x and atomicCas on it 4 times, similarly to what you suggested. If f is in shared memory, it should not be costly.

Now, when a thread sees it is in conflict, it does nothing. Just waits for all other threads to complete their work, __syncthreads, and then goes for the next iteration of the master for/while loop.

The difference from your proposed solution is that if you fail the lock, your thread just drops his work, instead of trying to wait.

Answer 2

It seems to me you are hugely overthinking this. If all the memory transactions must be serialised for the operation to be thread safe, then the simplest solution is to have one thread per block perform the operation. So something like

if (threadIdx.x == 0) // assume 1D block for simplicity
{
    y = f(x[i], x[i - 1], x[j], x[j + 1]);
    compare_and_swap(y, x[i], x[j];
}
__syncthreads();

will work fine because the array being operated on is in shared memory, so a guaranteed single thread per block is performing the operation, and there are no read-after-write hazards. In practice, this approach shouldn't be slower than hve a whole block of threads in contention for a lock, or a large number of serialised atomic memory transactions.