双打定义错误的CUDA atomicAdd

Question

在以前的 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 内置函数相同。

Answer 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，则不会声明该方法。