c++ - CUDA、互斥量和 atomicCAS()

Question

最近开始在CUDA上开发，遇到了atomicCAS()的问题。要在设备代码中对内存进行一些操作，我必须创建一个互斥量，以便只有一个线程可以在代码的关键部分使用内存。

下面的设备代码在 1 个 block 和多个线程上运行。

__global__ void cudaKernelGenerateRandomGraph(..., int* mutex)
{
    int i = threadIdx.x;
    ...

    do 
    {
        atomicCAS(mutex, 0, 1 + i);
    }
    while (*mutex != i + 1);

    //critical section
    //do some manipulations with objects in device memory

    *mutex = 0;

    ...
}

当第一个线程执行时

atomicCAS(mutex, 0, 1 + i);

mutex 为 1。在第一个线程将其状态从 Active 更改为 Inactive 之后，行

*mutex = 0;

未执行。其他线程永远处于循环中。我已经尝试了这个循环的许多变体，比如 while(){};、do{}while();，在循环内使用 temp variable = *mutex，甚至使用 if(){} 和 goto 的变体。但结果是一样的。

宿主部分代码:

...
int verticlesCount = 5;
int *mutex;
cudaMalloc((void **)&mutex, sizeof(int));
cudaMemset(mutex, 0, sizeof(int));
cudaKernelGenerateRandomGraph<<<1, verticlesCount>>>(..., mutex);

我使用带有 CUDA 5.5 的 Visual Studio 2012。

该设备是具有 1.2 计算能力的 NVidia GeForce GT 240。

提前致谢。

---

更新:今年 Spring 在我的文凭项目上工作了一段时间后，我找到了 cuda 上关键部分的解决方案。这是无锁和互斥机制的组合。这是工作代码。用它来插入原子动态调整大小的数组。

// *mutex should be 0 before calling this function
__global__ void kernelFunction(..., unsigned long long* mutex) 
{
    bool isSet = false; 
    do 
    {
        if (isSet = atomicCAS(mutex, 0, 1) == 0) 
        {
            // critical section goes here
        }
        if (isSet) 
        {
            mutex = 0;
        }
    } 
    while (!isSet);
}

Answer 1

有问题的循环

do 
{
    atomicCAS(mutex, 0, 1 + i);
}
while (*mutex != i + 1);

如果它在主机 (CPU) 端运行，可以正常工作；一旦线程 0 将 *mutex 设置为 1，其他线程将等待直到线程 0 将 *mutex 设置回 0。

但是，GPU 线程并不像 CPU 线程那样独立。 GPU 线程以 32 个为一组进行分组，通常称为 warps。 同一个 warp 中的线程将以完全锁步的方式执行指令。如果诸如 if 或 while 之类的控制语句导致 32 个线程中的某些线程与其余线程分道扬镳，其余线程将等待(即休眠) 为了完成不同的线程。 [1]

回到有问题的循环，线程 0 变为非事件状态，因为线程 1、2、...、31 仍停留在 while 循环中。所以线程 0 永远不会到达 *mutex = 0 行，而其他 31 个线程永远循环。

一个可能的解决方案是制作一个有问题的共享资源的本地拷贝，让 32 个线程修改拷贝，然后选择一个线程将更改“推送”回共享资源。 __shared__ 变量在这种情况下是理想的:它将由属于同一 block 但不属于其他 block 的线程共享。我们可以使用__syncthreads()来精细控制成员线程对该变量的访问。

[1] CUDA Best Practices Guide - Branching and Divergence

Avoid different execution paths within the same warp.

Any flow control instruction (if, switch, do, for, while) can significantly affect the instruction throughput by causing threads of the same warp to diverge; that is, to follow different execution paths. If this happens, the different execution paths must be serialized, since all of the threads of a warp share a program counter; this increases the total number of instructions executed for this warp. When all the different execution paths have completed, the threads converge back to the same execution path.