定义atomicAdd函数在CUDA中不起作用

Question

由于CUDA 2.0x没有atomicAdd()用于double的函数，因此我atomicAddd()根据此问题定义了“ atomicAdd（）”函数，

为什么对double而言没有实现atomicAdd？

这是设备功能的代码：

__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

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