CUDA原子操作（CAS、Exch）循环挂死问题求助

问题原因 #

你的代码挂死的核心原因是自旋锁实现不符合GPU线程调度特性：

• GPU以线程束（Warp）为调度单位，当一个线程束内多数线程进入自旋循环时，会持续占用SM计算资源，导致持有锁的线程无法获得足够执行时间片完成锁释放操作。
• 自旋循环未插入线程暂停指令，线程无限制重复执行原子操作，进一步加剧资源竞争与调度阻塞。

修正后的代码 #

主机端代码 #

int state = 0;
int * mutex;
cudaMalloc(&mutex, sizeof(int) );
cudaMemcpy(mutex, &state , sizeof(int), cudaMemcpyHostToDevice);
    
mykernel<<< 1, 9>>>(mutex);
cudaDeviceSynchronize(); // 等待核函数执行完成
cudaFree(mutex);

设备端核函数 #

__global__ void mykernel(int * mutex){
    int ti = threadIdx.x + blockDim.x * blockIdx.x;

    printf("TESTATOM0 -- THREAD%d MUTEX:%d\n", ti, *mutex);
    
    // 优化后的自旋锁逻辑
    while (atomicCAS(mutex, 0, 1) != 0) {
        __pause(); // 让出资源，允许其他线程调度
    }
    
    // 临界区操作
    for(int i = 0; i < 5; i++){
        printf("BLOCKED -- THREAD%d MUTEX:%d\n", ti, *mutex);
    }
    
    // 内存屏障确保临界区操作全局可见
    __threadfence();
    atomicExch(mutex, 0);
    
    for(int i = 0; i < 5; i++){
        printf("UNBLOCKED -- THREAD%d MUTEX:%d\n", ti, *mutex);
    }
}

关键修改说明 #

1. __pause()指令：
   告诉GPU硬件当前线程处于自旋等待状态，暂时让出计算资源，让持有锁的线程有机会执行并释放锁，避免资源独占导致的死锁。

2. __threadfence()内存屏障：
   确保临界区的所有内存操作（如printf缓冲写入）在锁被释放前已完成并对其他线程可见，避免其他线程提前获取锁后看到不一致的状态。

3. cudaDeviceSynchronize()同步：
   主机端等待核函数完全执行完毕，避免程序提前退出导致的输出不完整或异常。

额外注意事项 #

• GPU自旋锁仅适合短期临界区操作，长期阻塞会严重降低性能，优先使用CUDA原生同步原语（如共享内存同步、流同步等）替代自定义自旋锁。
• 当线程数超过线程束大小（32）时，线程束间调度复杂度提升，自旋锁性能会进一步下降，需谨慎使用。

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