CUDA多线程修改内存元素遇竞态，atomicAdd编译报错求助

问题背景 #

• 硬件：GeForce GTX 1080 Ti（Pascal架构，计算能力6.1）
• CUDA版本：12.4，驱动版本：550.54.14
• 多线程并发修改设备内存中的double数组，存在竞态条件，导致实际输出与预期不符
• 尝试使用atomicAdd(&x[idx], 1.0)解决竞态，但编译器提示“无匹配的atomicAdd重载函数”

核心原因 #

GTX 1080 Ti属于Pascal架构（计算能力6.1），而CUDA原生的atomicAdd对double类型的支持仅从**Volta架构（计算能力7.0）**开始提供。旧架构不支持直接调用double版本的atomicAdd，因此编译报错。

解决方案 #

方案1：手动实现double原子加法（兼容旧架构） #

利用CUDA的atomicCAS（原子比较并交换）操作，结合类型转换手动实现double的原子加法，代码如下：

__device__ void increment(double *x, int start, int span) {
  for (int idx = start; idx < start + span; idx++) {
    unsigned long long int* addr = reinterpret_cast<unsigned long long int*>(&x[idx]);
    unsigned long long int old_val = *addr;
    unsigned long long int assumed_val;
    do {
        assumed_val = old_val;
        // 将double转成unsigned long long，完成加法后再转回去
        double new_val = __longlong_as_double(assumed_val) + 1.0;
        old_val = atomicCAS(addr, assumed_val, __double_as_longlong(new_val));
    } while (assumed_val != old_val); // 循环直到CAS操作成功
  }
}

方案2：改用float类型（精度允许时） #

如果业务场景可以接受float的精度损失，直接将数组类型改为float，即可使用原生atomicAdd：

// 修改数组类型为float
float *x;
cudaMallocManaged(&x, N * sizeof(float));

// 初始化
for (int i = 0; i < N; i++) {
  x[i] = 0.0f;
}

// 原子加法
atomicAdd(&x[idx], 1.0f);

修改后验证代码（方案1） #

完整可运行代码：

#include <iostream>
#include <math.h>

__device__ void increment(double *x, int start, int span) {
  for (int idx = start; idx < start + span; idx++) {
    unsigned long long int* addr = reinterpret_cast<unsigned long long int*>(&x[idx]);
    unsigned long long int old_val = *addr;
    unsigned long long int assumed_val;
    do {
        assumed_val = old_val;
        double new_val = __longlong_as_double(assumed_val) + 1.0;
        old_val = atomicCAS(addr, assumed_val, __double_as_longlong(new_val));
    } while (assumed_val != old_val);
  }
}

__global__ void test01(double *x, int N) {
  int idx = threadIdx.x + blockIdx.x * blockDim.x;

  if (idx >= 10)
    return;

  increment(x, idx, 10);
}

int main(void) {
  int N = 20;
  double *x;

  cudaMallocManaged(&x, N * sizeof(double));

  for (int i = 0; i < N; i++) {
    x[i] = 0.0;
  }

  int blockSize = 256;
  int numBlocks = (10 + blockSize - 1) / blockSize;
  test01<<<numBlocks, blockSize>>>(x, N);

  cudaDeviceSynchronize();

  for (int i = 0; i < N; i++)
    printf("%0.0f ", x[i]);
  printf("\n");

  cudaFree(x);

  return 0;
}

编译运行后将得到预期输出：1 2 3 4 5 6 7 8 9 10 9 8 7 6 5 4 3 2 1 0

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