CUDA atomicAdd用于双精度定义错误

Question

在以前的CUDA版本中,atomicAdd没有实现双精度,所以通常在这里实现这一点.使用新的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函数.

CUDA内置函数(例如atomicAdd)是实现定义的,可以在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

这个宏命名的体系结构标识宏在此处记录:

5.7.4.虚拟架构识别宏

在__CUDA_ARCH__编译compute_xy的每个nvcc编译阶段1期间,为体系结构标识宏分配三位数值字符串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,则不声明该方法.