为什么atomicAdd未实现双打？

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

从的附录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的一部分？

来源

2012-09-27 smilingbuddha

可能致使其每个用户知道它的实现，因为它不是一个内置的指令和重试逻辑可以承受活锁（因为没有公平的保证，一个线程可以得到停滞不前的只要有其他线程更新相同的变量）。 – tera

编辑：正如CUDA 8，双精度atomicAdd()在CUDA实现与SM_6X（帕斯卡）的GPU的硬件支持。

~~目前，没有CUDA设备在double硬件中支持atomicAdd。~~正如您所指出的那样，它可以在64位整数上以atomicCAS的形式实现，但是这对性能成本来说是不平凡的。

因此，CUDA软件团队选择将正确的实现文档记录为开发人员的选项，而不是将其作为CUDA标准库的一部分。这样开发人员不会在不知不觉中选择性能成本。另外：我不认为这个问题应该被视为“不具有建设性”。我认为这是一个完全有效的问题，+1。

来源

2012-09-27 23:59:23 harrism

是的，但从技术上讲，你是少数人之一，可能会回答这个问题。虽然我已经说过为什么我认为这样做很有意义，但只有您可以说这是CUDA团队如此选择它的原因。 ;-)无论如何，我不是一个倒下了这个问题的人。 – tera

还有谁阅读并回答SO CUDA疑问多个NVIDIA人们（尤其是在我们的开发者论坛下降），而事实使得这样有效的问题。你可以发表你的评论作为答案，这将是正确的，我会投票赞成。顺便说一句，我没有假设你低调;我指的是一个投票来结束这个问题。 – harrism

我同意，这是一个完全有效的问题，CUDA头文件可以在软件中实现双原子。虽然它制定的方式引发了红光对某些人来说，我认为决定应该恢复！ – pszilard

为什么atomicAdd未实现双打？

回答

相关问题