cuda - 在 CUDA 中实现临界区

Question

我正在尝试使用原子指令在 CUDA 中实现关键部分，但遇到了一些麻烦。我已经创建了测试程序来显示问题:

#include <cuda_runtime.h>
#include <cutil_inline.h>
#include <stdio.h>

__global__ void k_testLocking(unsigned int* locks, int n) {
    int id = threadIdx.x % n;
    while (atomicExch(&(locks[id]), 1u) != 0u) {} //lock
    //critical section would go here
    atomicExch(&(locks[id]),0u); //unlock
}

int main(int argc, char** argv) {
    //initialize the locks array on the GPU to (0...0)
    unsigned int* locks;
    unsigned int zeros[10]; for (int i = 0; i < 10; i++) {zeros[i] = 0u;}
    cutilSafeCall(cudaMalloc((void**)&locks, sizeof(unsigned int)*10));
    cutilSafeCall(cudaMemcpy(locks, zeros, sizeof(unsigned int)*10, cudaMemcpyHostToDevice));

    //Run the kernel:
    k_testLocking<<<dim3(1), dim3(256)>>>(locks, 10);

    //Check the error messages:
    cudaError_t error = cudaGetLastError();
    cutilSafeCall(cudaFree(locks));
    if (cudaSuccess != error) {
        printf("error 1: CUDA ERROR (%d) {%s}\n", error, cudaGetErrorString(error));
        exit(-1);
    }
    return 0;
}

不幸的是，这段代码将我的机器硬卡住了几秒钟，最后退出，打印出以下消息:

fcudaSafeCall() Runtime API error in file <XXX.cu>, line XXX : the launch timed out and was terminated.

这意味着其中一个 while 循环没有返回，但这似乎应该有效。

温馨提示 atomicExch(unsigned int* address, unsigned int val)原子地将存储在地址中的内存位置的值设置为 val并返回 old值(value)。所以我的锁定机制背后的想法是它最初是 0u ，所以一个线程应该通过 while循环和所有其他线程应该等待 while循环，因为他们会读 locks[id]如 1u .然后当线程完成临界区时，它将锁重置为 0u所以另一个线程可以进入。

我错过了什么？

顺便说一下，我正在编译:

nvcc -arch sm_11 -Ipath/to/cuda/C/common/inc XXX.cu

Answer 1

好的，我想通了，这是另一个 cuda 范式的痛苦。

任何优秀的 cuda 程序员都知道(请注意，我不记得这让我成为一个糟糕的 cuda 程序员，我认为)warp 中的所有线程都必须执行相同的代码。如果不是因为这个事实，我编写的代码将完美运行。然而，事实上，同一个经线中可能有两个线程访问同一个锁。如果其中之一获得了锁，它只会忘记执行循环，但在其经纱中的所有其他线程完成循环之前，它无法继续执行循环。不幸的是，另一个线程永远不会完成，因为它正在等待第一个解锁。

这是一个可以毫无错误地完成任务的内核:

__global__ void k_testLocking(unsigned int* locks, int n) {
    int id = threadIdx.x % n;
    bool leaveLoop = false;
    while (!leaveLoop) {
        if (atomicExch(&(locks[id]), 1u) == 0u) {
            //critical section
            leaveLoop = true;
            atomicExch(&(locks[id]),0u);
        }
    } 
}