CUDA 中的原子操作答案

【问题标题】：atomic operations in CUDACUDA 中的原子操作
【发布时间】：2012-02-13 14:18:50
【问题描述】：

以下程序使用“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);
}

【问题讨论】：

上面的程序只是我想要做的一个例子。我知道我可以在上述情况下使用 atomicAdd 来增加 i。但是，在我的实际案例中，关键部分更复杂

标签： cuda gpu-atomics

【解决方案1】：

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

【讨论】：