CUDA,互斥和atomicCAS()

Question

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

下面的设备代码在1个块和几个线程上运行.

__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和line之后

*mutex = 0;

没有执行.其他线程永远保持循环.我尝试了这个循环的许多变体,比如while(){};,do {} while();, temp variable =*mutex inside loop,甚至是if(){}和goto的变体.但结果是一样的.

代码的主机部分:

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

我使用Visual Studio 2012和CUDA 5.5.

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

提前致谢.

---

UPD: 在今年春天的文凭项目工作一段时间后,我找到了关于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个组,通常称为warp.同一warp中的线程将执行完整锁定步骤中的指令.如果控制语句例如if或while导致32个线程中的某些线程与其余线程分开,则剩余线程将等待(即休眠)以使发散线程完成.[1]

回到有问题的循环,线程0变为非活动状态,因为线程1,2,...,31仍然停留在while循环中.所以线程0永远不会到达该行*mutex = 0,其他31个线程永远循环.

一个潜在的解决方案是让共享资源的本地副本中的问题,让32个线程修改副本,然后选择一个线程来"推"变回共享资源.甲__shared__变量是在这种情况下理想的:它会由属于同一块的线程而不是其他块共享.我们可以使用__syncthreads()成员线程来精细控制这个变量的访问.

[1] CUDA最佳实践指南 - 分支和分歧

避免在同一个warp中使用不同的执行路径.

任何流控制指令(if,switch,do,for,while)都会通过使相同warp的线程发散而显着影响指令吞吐量; 也就是说,遵循不同的执行路径.如果发生这种情况,必须序列化不同的执行路径,因为warp的所有线程共享一个程序计数器; 这会增加为此warp执行的指令总数.当所有不同的执行路径完成后,线程会聚回到相同的执行路径.