CUDA多块全局锁原子操作失效问题及__threadfence()作用咨询

环境配置 #

• GPU：Quadro RTX 4000
• CUDA版本：11.7

核函数与测试结果 #

实现1：floatAddLockExch（使用atomicExch释放锁） #

__global__ void floatAddLockExch(float* addr, volatile int* lock) {
    bool goon = true;
    while (goon) {
        if (atomicCAS((int*)lock, 0, 1) == 0) {
            *addr += 1; // 临界区任务

            int lockValue = atomicExch((int*)lock, 0);
            if (lockValue != 1) {
                printf("Error in <%d, %d>\n", blockIdx.x, threadIdx.x);
            }
//
//          *lock = 0;
//          __threadfence();

            goon = false;
        }
    }
}

测试结果：

• 以<<<1, 1024>>>启动时，*addr输出为1024；
• 以<<<2, 1024>>>启动时，输出为1025，且无"Error..."输出。

实现2：floatAddLockFence（使用__threadfence()） #

__global__ void floatAddLockFence(float* addr, volatile int* lock) {
    bool goon = true;
    while (goon) {
        if (atomicCAS((int*)lock, 0, 1) == 0) {
            *addr += 1; // 临界区任务

//          int lockValue = atomicExch((int*)lock, 0);
//          if (lockValue != 1) {
//              printf("Error in <%d, %d>\n", blockIdx.x, threadIdx.x);
//          }

            *lock = 0;
            __threadfence();

            goon = false;
        }
    }
}

测试结果：

• 以<<<1, 1024>>>启动时，*addr输出为1024；
• 以<<<2, 1024>>>启动时，输出为2048。

疑问 #

按预期，全局变量lock的原子操作应跨所有块原子执行，为何floatAddLockExch在多块场景失效？__threadfence()是如何解决该问题的？

1. floatAddLockExch多块场景失效的原因 #

CUDA的SM（流式多处理器）自带一级缓存，不同SM上的线程块访问全局内存时，会将数据缓存到本地。在floatAddLockExch中，虽然atomicCAS和atomicExch是原子操作，但临界区的*addr +=1操作结果可能还未同步到全局内存，其他SM上的线程块就通过原子操作获取了锁，读取到的addr是旧值，导致重复累加。

更关键的是：atomicExch释放锁后，当前线程块对addr的修改可能还停留在SM的缓存中，没有刷回全局内存。其他SM上的线程块拿到锁后，从全局内存或自身缓存读取addr，此时读到的是未更新的值，所以多次累加只生效了少数次数，最终结果远小于预期的2048。

另外，volatile修饰lock仅保证对lock的读写不被编译器优化，但无法保证跨SM的缓存一致性——不同SM的缓存不会自动同步，除非有内存栅栏或原子操作强制同步。

2. __threadfence()的作用 #

__threadfence()是CUDA的内存栅栏函数，它会强制当前线程块中所有未完成的内存操作（包括对addr的写操作）都刷回全局内存，并且确保这些操作对其他SM上的线程可见后，才执行栅栏之后的指令。

在floatAddLockFence中，先执行*lock =0释放锁，再调用__threadfence()：

• __threadfence()会确保临界区对addr的修改已完全刷入全局内存，同时lock=0的写操作也对所有SM可见；
• 这样其他SM上的线程块通过atomicCAS获取锁时，既能看到lock已释放，也能读取到addr的最新值，从而正确执行累加操作。

需要注意的是：atomicExch仅保证对lock的操作是原子的，不会自动同步临界区的其他内存操作；而__threadfence()专门解决跨SM的内存可见性问题，确保临界区修改对所有线程可见后再释放锁。

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