定義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

與數學加法不同，由於所涉及的舍入步驟，浮點加法不是關聯的。在需要原子操作的情況下，操作順序不是確定性的。所以非確定性舍入誤差是不可避免的。