定义atomicAdd功能在CUDA

Question

不起作用作为CUDA 2.0倍不具有双重功能atomicAdd()，然后我根据这个问题定义“atomicAdd（）”功能atomicAddd()，

Why has atomicAdd not been implemented for doubles?

这里是设备功能的代码：

__device__ double atomicAddd(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 test(double *dev_like, double *dev_sum){ 
    __shared__ double lik; 
    // some code to compute lik; 
    // copy lik back to global dev_lik; 
    dev_like[blockIdx.x] = lik; 

    // add lik to dev_sum 
    if(threadIdx.x == 0){ 
     atomicAddd(dev_sum, loglik); 
    } 

} 

后，我复制dev_lik回主机并将其添加到sum，我也复制dev_sum回主机sum1。我的理解是sum应该与sum1相同，这里是我的主机代码来打印它们。

for (int m = 0; m < 100; ++m){ 
     if(sum[m] == sum1[m]){ 
      std::cout << "True" << std::endl; 
     } 
     else{ 
      std::cout << "False" << "\t" << std::setprecision(20) << sum[m] << "\t" << std::setprecision(20) << sum1[m] << std::endl; 
     } 
    } 

，我得到的结果如下：

True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
False -1564.0205173292260952 -1564.0205173292256404 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
False -1563.4011523293495429 -1563.4011523293493156 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 
True 

一些结果表明False但sum和sum1之间的差别非常小，不知道是什么问题。

Answer 1

与数学加法不同，由于所涉及的舍入步骤，浮点加法不是关联的。在需要原子操作的情况下，操作顺序不是确定性的。所以非确定性舍入误差是不可避免的。