cuda::atomic功能解析及代码异常输出问题排查

问题代码 #

#include <atomic>
#include <cuda/atomic>
#include <stdio.h>


#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}


__global__ void atomic_test()
{
    cuda::atomic<int, cuda::thread_scope_block> x{0};
    x.fetch_add(1, cuda::memory_order_seq_cst);
    __syncthreads();
    int y = x.load(cuda::memory_order_acquire);
    printf("(%d %d) - Value of x is %d\n", blockIdx.x, threadIdx.x, y); 
}

int main()
{
    atomic_test<<<2, 32>>>();
    gpuErrchk( cudaDeviceSynchronize() );
    return 0;
}

问题现象 #

预期线程块内每个线程读取到的x值相同（应为32），但实际输出中仅线程31输出32，其余线程输出0。

错误原因 #

核心问题是**cuda::atomic变量的作用域错误**：

• 核函数中直接声明的x是线程局部变量，每个线程都拥有独立的x实例，存储在线程私有内存（寄存器/局部内存）中，线程之间无法互相访问对方的x。
• __syncthreads()仅同步线程块内的执行顺序，无法让线程看到其他线程的私有变量。
• 输出异常的表现（仅线程31显示32）是编译器优化或寄存器分配的副作用：大部分线程的私有x因未被有效引用被优化掉，只有线程31的x保留了修改后的值，但这本质上还是每个线程操作自己的变量，并非共享。

修正方案 #

要让cuda::thread_scope_block类型的原子变量被块内所有线程共享，必须将其声明为共享内存变量（用__shared__修饰），并正确初始化：

#include <atomic>
#include <cuda/atomic>
#include <stdio.h>


#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}


__global__ void atomic_test()
{
    // 声明为共享内存变量，块内所有线程共享同一个实例
    __shared__ cuda::atomic<int, cuda::thread_scope_block> x;
    
    // 仅由线程0完成初始化，避免多线程初始化冲突
    if (threadIdx.x == 0) {
        x.store(0, cuda::memory_order_relaxed);
    }
    __syncthreads(); // 等待所有线程看到初始化后的x值

    x.fetch_add(1, cuda::memory_order_seq_cst);
    __syncthreads(); // 等待所有线程完成加法操作

    int y = x.load(cuda::memory_order_acquire);
    printf("(%d %d) - Value of x is %d\n", blockIdx.x, threadIdx.x, y); 
}

int main()
{
    atomic_test<<<2, 32>>>();
    gpuErrchk( cudaDeviceSynchronize() );
    return 0;
}

修正要点 #

1. 共享内存修饰：用__shared__标记x，将其存储到线程块的共享内存区域，确保块内线程访问同一变量。
2. 安全初始化：共享内存变量不会自动初始化，需由单个线程（如线程0）完成初始化，再通过__syncthreads()同步，避免其他线程读取未初始化的值。
3. 同步保证：两次__syncthreads()分别确保初始化完成和所有加法操作完成，保证后续读取到正确的最终值。

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