为什么没有为双打实现atomicAdd？

Question

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

从CUDA编程指南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)

同样的页面继续为我的项目中刚刚开始使用的下面的双打提供了一个小型的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开始,atomicAdd()在CUDA中实现了双精度,并在SM_6X(Pascal)GPU中提供硬件支持.

目前,还没有CUDA设备支持atomicAdd的double硬件.正如您所指出的,它可以atomicCAS在64位整数方面实现,但是有一个非平凡的性能成本.

因此,CUDA软件团队选择将正确的实现文档记录为开发人员的选项,而不是将其作为CUDA标准库的一部分.这样开发人员就不会在不知情的情况下选择他们不理解的性能成本.

旁白:我不认为这个问题应该被视为"没有建设性".我认为这是一个非常有效的问题,+ 1.