双打定义错误的CUDA atomicAdd

Posted 2023-04-15

【中文标题】双打定义错误的CUDA atomicAdd

【英文标题】：CUDA atomicAdd for doubles definition error

【发布时间】：2016-09-30 17:20:42

【问题描述】：

在以前的 CUDA 版本中，没有为双精度实现 atomicAdd，因此通常像 here 那样实现这一点。使用新的 CUDA 8 RC，当我尝试编译包含此类功能的代码时遇到了麻烦。我猜这是因为在 Pascal 和 Compute Capability 6.0 中，添加了 atomicAdd 的原生双版本，但不知何故，以前的 Compute Capabilities 并没有正确忽略这一点。

下面的代码在以前的 CUDA 版本中可以正常编译和运行，但现在我得到了这个编译错误：

test.cu(3): error: function "atomicAdd(double *, double)" has already been defined

但是如果我删除我的实现，我会得到这个错误：

test.cu(33): error: no instance of overloaded function "atomicAdd" matches the argument list
            argument types are: (double *, double)

我应该补充一点，我只有在使用 -arch=sm_35 或类似代码编译时才会看到这个。如果我使用 -arch=sm_60 编译，我会得到预期的行为，即只有第一个错误，而在第二种情况下编译成功。

编辑：另外，它是特定于atomicAdd 的——如果我更改名称，它会很好用。

它看起来真的像一个编译器错误。其他人可以确认是这种情况吗？

示例代码：

__device__ double atomicAdd(double* address, double val)

    unsigned long long int* address_as_ull = (unsigned long long int*)address;
    unsigned long long int old = *address_as_ull, assumed;
    do 
        assumed = old;
        old = atomicCAS(address_as_ull, assumed,
                __double_as_longlong(val + __longlong_as_double(assumed)));
     while (assumed != old);
    return __longlong_as_double(old);


__global__ void kernel(double *a)

    double b=1.3;
    atomicAdd(a,b);


int main(int argc, char **argv)

    double *a;
    cudaMalloc(&a,sizeof(double));

    kernel<<<1,1>>>(a);

    cudaFree(a);
    return 0;

---

编辑：我从 Nvidia 那里得到了一个认识到这个问题的答案，以下是开发人员对此的看法：

CUDA 8.0 新支持的 sm_60 架构具有 本机 fp64 atomicAdd 函数。由于我们的局限性 工具链和CUDA语言，这个函数的声明需要 即使没有专门为代码编译，也存在 sm_60。这会导致您的代码出现问题，因为您还定义了一个 fp64 atomicAdd 函数。

诸如 atomicAdd 之类的 CUDA 内置函数是实现定义的 并且可以在 CUDA 版本之间更改。用户不应定义 与任何 CUDA 内置函数同名的函数。我们会 建议您将 atomicAdd 函数重命名为不是 与任何 CUDA 内置函数相同。

【问题讨论】：

对我来说看起来像是 CUDA 8 RC 中的一个错误。似乎本机双 atomicAdd() 仅适用于 sm_60，但也可以与 sm_35 一起使用。也许你可以通过重命名自己的版本来解决这个问题。

@Eric 是的，重命名解决了它。已编辑帖子以包含此内容。

【参考方案1】：

atomicAdd 的这种风格是为计算能力 6.0 引入的一种新方法。您可以使用宏定义来保护您之前实现的其他计算能力

#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 600
#else
<... place here your own pre-pascal atomicAdd definition ...>
#endif

这个名为架构识别宏的宏被记录在here：

5.7.4。虚拟架构识别宏

在为 compute_xy 编译的每个 nvcc 编译阶段 1 期间，架构标识宏 __CUDA_ARCH__ 被分配一个三位值字符串 xy0（以文字 0 结尾）。

此宏可用于 GPU 函数的实现，以确定当前正在为其编译的虚拟架构。主机代码（非 GPU 代码）不能依赖它。

我假设 NVIDIA 没有将它放在以前的 CC 中，以避免用户定义它和不迁移到 Compute Capability >= 6.x 的冲突。不过，我不认为这是一个 BUG，而是一种发布交付实践。

编辑：宏保护不完整（已修复） - 这里是一个完整的示例。

#if !defined(__CUDA_ARCH__) || __CUDA_ARCH__ >= 600
#else
__device__ double atomicAdd(double* a, double b)  return b; 
#endif

__device__ double s_global ;
__global__ void kernel ()  atomicAdd (&s_global, 1.0) ; 


int main (int argc, char* argv[])

        kernel<<<1,1>>> () ;
        return ::cudaDeviceSynchronize () ;

编译：

$> nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2016 NVIDIA Corporation
Built on Wed_May__4_21:01:56_CDT_2016
Cuda compilation tools, release 8.0, V8.0.26

命令行（均成功）：

$> nvcc main.cu -arch=sm_60
$> nvcc main.cu -arch=sm_35

您可能会发现为什么它适用于包含文件：sm_60_atomic_functions.h，如果__CUDA_ARCH__ 低于 600，则不会声明该方法。

【讨论】：

我可能不会使用相同的名称，因为这会导致上面的第一个错误，“函数...已经被定义”。给出一个完全不必要的错误并带有非常混乱的消息怎么不是错误？

@kalj，您可以保持相同的名称，但由__CUDA_ARCH__ 保护。如果您的声明受到此宏的保护，则不应出现上面列出的错误。此外，这将使您的代码具有一定的一致性和清晰度。诚然，无论是错误还是 API 支持选择，更多的是意见而不是技术声明。选择哪一个，但 NVIDIA 将获得最后的决定权。

@kalj，我的宏测试确实不完整（已修复）。我提供了一个在 linux Ubuntu 16.04 上使用 CUDA 8.0 RC 编译的完整示例。

如何防止自定义 pre-pascal atomicAdd(double, double) 在系统上隐藏 atomicAdd(float, float) CUDA 函数 (device_atomic_function.hpp) >=CC2.X &&

@FlorentDUGUET 这是一个命名空间的东西， atomicAdd(double, double) 应该在命名空间之外，否则你必须为自己提供浮动版本