cuda - 为什么对double而言未实现atomicAdd？

Question

为什么没有将double的atomicAdd()明确实现为CUDA 4.0或更高版本的一部分？

在CUDA programming guide 4.1的附录F第97页中，以下版本的
atomicAdd已实现。

int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address,
                       unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address,
                                 unsigned long long int val);
float atomicAdd(float* address, float val)

继续在同一页面上提供对double的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);
}

为什么不将上述代码定义为CUDA的一部分？

Answer 1

编辑:从CUDA 8开始， double atomicAdd()在CUDA中实现，SM_6X(Pascal)GPU中具有硬件支持。

当前，没有CUDA设备在硬件中支持atomicAdd的double。如您所述，可以使用atomicCAS来实现64位整数，但是为此付出了不小的代价。

因此，CUDA软件团队选择将正确的实现记录为开发人员的选择，而不是使其成为CUDA标准库的一部分。这样，开发人员不会在不知不觉中选择他们不了解的性能成本。

旁白:我认为这个问题不应以“非 build 性的”来解决。我认为这是一个非常有效的问题，+ 1。