cuda - CUDA 中的原子操作

Question

以下程序使用“Cuda By Example”中的原子锁实现，但运行该程序会使我的机器冻结。有人可以告诉我我的程序有什么问题吗？非常感谢

逸飞

#include <stdio.h>


__global__ void test()
{
    __shared__ int i, mutex;

    if (threadIdx.x == 0) {
       i = 0;
       mutex = 0;
    }
    __syncthreads();

    while( atomicCAS(&mutex, 0, 1) != 0);
    i++;
    printf("thread %d: %d\n", threadIdx.x, i);
    atomicExch(&mutex,0);
}

Answer 1

这是一个理论。我希望你熟悉扭曲的概念。在 while 循环中，warp 中的所有线程都将进入 while 循环。只有一个会退出，其余线程将驻留在 while 循环中。这将引入一个发散分支，使退出 while 循环的线程停止，直到分支再次收敛。因为这个线程是唯一可以释放互斥锁的线程，所以这永远不会发生，因为它等待其他线程确实收敛。