在 cuda 内核函数上实现互斥锁恰好死锁

Question

我是cuda的新手，我尝试在内核函数中执行互斥锁。

我阅读了一些教程并编写了我的函数，但在某些情况下，发生了死锁。

这是我的代码，内核函数非常简单，可以计算由 main 函数启动的正在运行的线程数。

#include <iostream>
#include <cuda_runtime.h>

__global__ void countThreads(int* sum, int* mutex) {
    while(atomicCAS(mutex, 0, 1) != 0); // lock
    
    *sum += 1;
    __threadfence();

    atomicExch(mutex, 0); // unlock
}

int main() {
    int* mutex = nullptr;
    cudaMalloc(&mutex, sizeof(int));
    cudaMemset(&mutex, 0, sizeof(int));

    int* sum = nullptr;
    cudaMalloc(&sum, sizeof(int));
    cudaMemset(&mutex, 0, sizeof(int));

    int ret = 0;
    // pass, result is 1024
    countThreads<<<1024, 1>>>(sum, mutex);
    cudaMemcpy(&ret, sum, sizeof(int), cudaMemcpyDeviceToHost);
    std::cout << ret << std::endl; 
    
    // deadlock, why?
    countThreads<<<1, 2>>>(sum, mutex);
    cudaMemcpy(&ret, sum, sizeof(int), cudaMemcpyDeviceToHost);
    std::cout << ret << std::endl;

    return 0;
}

那么，谁能告诉我为什么在调用countThreads>>()时程序死锁，以及如何解决它？我想执行跨块互斥锁，但这可能不是一个好主意。非常感谢。

我试验了一段时间，发现如果在同一个块中使用线程，就会发生死锁，否则一切正常。

Answer 1

同一 warp 中的线程尝试协商锁定或互斥锁可能是最坏的情况。正确编程相当困难，并且行为可能会根据您运行的确切 GPU 而改变。

Here 是解释特定情况下死锁的确切原因所需的分析类型示例。由于您没有指出您正在编译或运行的 GPU 类型，因此无法根据您在此处显示的内容轻松完成此类分析。提供用于编译的 CUDA 版本也很重要。我目睹了从一代编译器到另一代编译器的代码更改，这可能会影响这一点。即使您提供了这些信息，我也不确定该分析是否真的值得，因为我认为正确编程的协商中的案例非常麻烦。 This question/answer 也可能感兴趣。

我对 CUDA 新人的一般建议（如您所说）是使用类似于here 描述的方法。简而言之，在线程块级别协商锁定（即每个块中的一个线程在其他块之间协商锁定）然后使用标准的可用块级协调方案管理块内的单例活动，例如__syncthreads()，和条件编码。

您可以通过在cuda 标签上搜索“锁定”、“临界区”等关键字来了解有关此主题的更多信息。

FWIW，无论如何，对我来说，您的代码在 Kepler 设备上确实死锁，而在 Volta 设备上不会死锁，正如 reference in the comments 所建议的那样。我不是试图传达任何关于您的代码是否无缺陷的声明，这只是一个观察。如果我将您的内核修改为如下所示：

__global__ void countThreads(int* sum, int* mutex) {

    int old = 1;
    while (old){
      old = atomicCAS(mutex, 0, 1);  // lock
      if (old == 0){
        *sum += 1;
        __threadfence();

        atomicExch(mutex, 0); // unlock
        }
      }
}

那么在我看来，无论是开普勒案例还是沃尔特案例都可以工作。我并不是要提出这个例子来暗示它是“正确的”，而是表明一些无害的代码修改可以将代码从死锁更改为非死锁情况，反之亦然。在我看来，最好避免这种脆弱性，当然是在 Volta 之前的情况下。

对于 volta 和前向案例，CUDA 11 和前向，您可能希望使用 libcu++ 库中的功能，例如 semaphore