CUDA AtomicCAS死锁问题：并行递增数组元素时程序挂起

问题描述 #

有一个初始值全为0的matrix数组，希望根据indices数组中存储的索引，对matrix的部分元素执行+1操作。由于部分元素需要多次递增，尝试为matrix的每个元素设置一个互斥信号量数组，但运行代码时程序挂起，出现死锁。

最终需求是使用CUDA绘制可重叠的连续画笔笔触，因此需要并行访问画布的同一像素。

示例代码 #

#include <iostream>
using namespace std;

__global__ void add_kernel(int* matrix, int* indices, int* d_semaphores, int nof_indices)
{
    int index = threadIdx.x + blockIdx.x * blockDim.x; // thread id
    int ind = indices[index]; // indices of target array A to increment    

    if (index < nof_indices) {
        while (atomicCAS(&d_semaphores[ind], 0, 1) != 0);
        matrix[ind] += 1;
        atomicExch(&d_semaphores[ind], 0);
        __syncthreads();
    }
}

int main()
{
    int nof_indices = 6; // length of an array B
    int indices[6] = { 0,1,2,3,4,1 }; // array B; stores indices of an array A which to increment
    int canvas[10]; // array A
    int semaphores[10]; // mutex array with individual mutexes for each of array A elements

    int* d_canvas;
    int* d_indices;
    int* d_semaphores;

    memset(canvas, 0, sizeof(canvas)); // set all array A elements to 0
    memset(semaphores, 0, sizeof(semaphores)); // set all array A elements to 0    

    cudaMalloc(&d_canvas, sizeof(canvas));
    cudaMalloc(&d_semaphores, sizeof(semaphores));
    cudaMalloc(&d_indices, sizeof(indices));

    cudaMemcpy(d_canvas, &canvas, sizeof(canvas), cudaMemcpyHostToDevice);
    cudaMemcpy(d_indices, &indices, sizeof(indices), cudaMemcpyHostToDevice);
    cudaMemcpy(d_semaphores, &semaphores, sizeof(semaphores), cudaMemcpyHostToDevice);

    add_kernel <<<1, 6>>> (d_canvas, d_indices, d_semaphores, nof_indices);

    cudaMemcpy(&canvas, d_canvas, sizeof(canvas), cudaMemcpyHostToDevice);

    for (int it = 0; it < nof_indices; it++) {
        cout << canvas[it] << endl;
    }

    cudaFree(d_canvas);
    cudaFree(d_indices);
    cudaFree(d_semaphores);

    return 0;
}

本示例中，matrix的预期结果应为{1, 2 ,1 ,1,1,0}，但仅当以<<< 6,1 >>>的核函数维度运行时才能得到正确结果。使用环境为CUDA 12.1和GeForce RTX 3060显卡。（仅当每个线程块的线程数设为1时程序正常，但这并非想要的实现方式）

问题分析与解决方案 #

死锁原因 #

1. __syncthreads()的错误使用：__syncthreads()要求线程块内所有线程都到达同步点才能继续执行。当线程块中有线程因index >= nof_indices提前退出，或有线程卡在自旋锁循环中时，其他线程会一直等待同步，导致整个线程块挂起，引发死锁。
2. 自旋锁的低效与风险：手动实现的自旋锁在CUDA中容易浪费线程资源，多线程竞争同一锁时，可能出现线程饥饿或死锁，尤其是多线程块场景下。

正确实现方式 #

对同一内存地址多次递增这类需求，CUDA提供原生原子操作，无需手动实现互斥锁，高效且安全。直接使用atomicAdd函数即可完成线程安全的递增操作：

修改后的核函数

__global__ void add_kernel(int* matrix, int* indices, int nof_indices)
{
    int index = threadIdx.x + blockIdx.x * blockDim.x;

    if (index < nof_indices) {
        int ind = indices[index];
        atomicAdd(&matrix[ind], 1); // 原子递增操作，线程安全
    }
}

修改后的main函数

移除信号量相关的内存分配与拷贝代码：

int main()
{
    int nof_indices = 6;
    int indices[6] = { 0,1,2,3,4,1 };
    int canvas[10];

    int* d_canvas;
    int* d_indices;

    memset(canvas, 0, sizeof(canvas));

    cudaMalloc(&d_canvas, sizeof(canvas));
    cudaMalloc(&d_indices, sizeof(indices));

    cudaMemcpy(d_canvas, canvas, sizeof(canvas), cudaMemcpyHostToDevice);
    cudaMemcpy(d_indices, indices, sizeof(indices), cudaMemcpyHostToDevice);

    add_kernel <<<1, 6>>> (d_canvas, d_indices, nof_indices);

    cudaMemcpy(canvas, d_canvas, sizeof(canvas), cudaMemcpyDeviceToHost);

    for (int it = 0; it < 10; it++) { // 遍历整个canvas数组查看完整结果
        cout << canvas[it] << endl;
    }

    cudaFree(d_canvas);
    cudaFree(d_indices);

    return 0;
}

方案优势 #

• atomicAdd是硬件级原子操作，确保同一时间只有一个线程对指定内存地址执行写操作，完全避免竞态条件。
• 无需手动管理锁，代码更简洁，执行效率远高于自旋锁（硬件原子操作延迟远低于软件自旋等待）。
• 支持任意线程块配置（如<<<1,6>>>或<<<2,3>>>等），不会出现死锁或挂起问题。

画笔笔触场景扩展 #

如果后续需要更复杂的像素操作（如基于当前像素值的计算），可使用atomicCAS实现原子读-改-写操作，或使用CUDA协作组进行精细线程同步，但简单计数/递增场景下，atomicAdd是最优选择。

内容的提问来源于stack exchange，提问作者sergei