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

Question

本人初学cuda，尝试在内核函数中进行mutex

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

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

#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<<<1, 2>>>() 时死锁了，如何解决？我想执行跨块互斥，但这可能不是一个好主意。非常感谢。

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

Answer 1

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

is an example of the type of analysis needed to explain the exact reason for the deadlock in a particular case. Such analysis is not readily done on what you have shown here because you have not indicated the type of GPU you are compiling for, or running on. It's also fairly important to provide the CUDA version you are using for compilation. I have witnessed code changes from one compiler generation to another, that may impact this. Even if you provided that information, I'm not sure the analysis is really worth-while, because I consider the negotiation-within-a-warp case to be extra troublesome to program correctly. 可能也有兴趣。

我对 CUDA 新手（如您所说）的一般建议是使用与描述的方法类似的方法 here。简而言之，在线程块级别协商锁定（即每个块中的一个线程在其他块之间协商锁定）然后使用标准的可用块级协调方案在块内管理单例 activity，例如__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 和 forward 情况，CUDA 11 及更高版本，您可能希望使用 libcu++ 库中的功能，例如 semaphore

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

Implementing of mutex on cuda kernel function happens to be deadlocked

cuda

mutex

deadlock