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

【问题标题】：Define atomicAdd function doesn't work in CUDA定义 atomicAdd 函数在 CUDA 中不起作用
【发布时间】：2016-11-24 06:27:02
【问题描述】：

由于 CUDA 2.0x 没有 atomicAdd() double 函数，所以我根据这个问题将 '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之间的差异很小，不知道是什么问题。

【问题讨论】：

标签： cuda

【解决方案1】：

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

【讨论】：

现在就明白了。谢谢。