35

为什么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 的一部分？

4

1 回答 1

41

编辑：从 CUDA 8 开始，双精度atomicAdd()在 CUDA 中实现，并在 SM_6X (Pascal) GPU 中提供硬件支持。

atomicAdd目前，硬件中不支持任何 CUDA 设备double。正如您所指出的，它可以根据atomicCAS64 位整数来实现，但这样做会带来不小的性能成本。

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

旁白：我认为这个问题不应该以“不具建设性”的方式结束。我认为这是一个完全有效的问题，+1。

于 2012-09-27T23:59:23.287 回答