CUDA互斥锁为何有时需要__threadfence()？

先看你的测试内核代码：

__global__ void write(int *count, int *a, int length, int *lock){
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if(i == 0)
        *lock = 0;
    if(i < length){
        bool blocked = true;
        while(blocked) {
            if(0 == atomicCAS(lock, 0, 1)) {
                //*count = *count + 1;
                atomicAdd(count, 1);
                a[i] = *count;
                atomicExch(lock, 0);
                blocked = false;
            }
        }
    }
}

write<<<(length + 31) / 32, 32>>>(dev_count, dev_a, length, lock);

当你把atomicAdd(count, 1)替换为*count = *count + 1时，数组a会出现重复值；但在atomicExch前添加__threadfence()后问题解决。你后续在临界区前后都加了该函数，代码如下：

__global__ void func(int length, int *lock){
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if(i == 0)
        *lock = 0;
    if(i < length){
        bool blocked = true;
        while(blocked) {
            if(0 == atomicCAS(lock, 0, 1)) {
                __threadfence();
                /*
                    critical section
                */
                __threadfence();
                atomicExch(lock, 0);
                blocked = false;
            }
        }
    }
}

为什么需要__threadfence() #

• CUDA线程的内存访问存在缓存延迟和指令重排序：线程对全局内存的写入不会立刻同步到全局内存，而是先暂存在线程私有缓存或L1/L2缓存中；同时编译器、硬件为了提升效率，可能会调整内存操作的执行顺序。
• 用*count = *count + 1时，这是两个独立操作：先读取count的当前值，加1后再写回。虽然锁机制（atomicCAS+atomicExch）保证了临界区的互斥，但如果没有内存栅栏，线程A修改count和a[i]后，这些改动可能还停留在本地缓存，未同步到全局内存。此时线程B获取锁后，读取到的count还是旧值，自然会写出重复的a[i]。
• 而atomicAdd是硬件层面实现的原子操作，会保证读写改的全过程原子性，且操作结果会立即同步到全局内存，因此不会出现重复问题。手动实现的自增操作不具备这个特性，必须依赖内存栅栏来保证可见性。

__threadfence()的工作原理 #

• __threadfence()是CUDA的内存栅栏函数，核心作用是强制当前线程：
  1. 栅栏之前的所有全局内存写操作，必须完成并同步到全局内存（让其他线程可见）；
  2. 栅栏之后的内存操作，不会被重排到栅栏之前执行。
• 对应你的代码场景：
  • 临界区前的__threadfence()：确保当前线程获取锁后，其他线程之前释放锁时的所有内存修改都已同步到全局内存，避免读取到过期的缓存值。
  • 临界区后的__threadfence()：确保临界区内的所有写操作（比如修改count、a[i]）都已同步到全局内存，再释放锁。这样下一个获取锁的线程能读取到最新的正确值，不会出现重复。
• 注意：__threadfence()针对全局内存可见性，会阻塞当前线程，直到所有未完成的全局内存写操作都提交到内存系统，保证后续其他线程的内存访问能看到这些修改。

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