在 GPU 上用于双精度的 atomicAdd()答案

【问题标题】：atomicAdd() for double on GPU在 GPU 上用于双精度的 atomicAdd()
【发布时间】：2013-04-11 05:51:14
【问题描述】：

我在GPU上做一个项目，我必须使用atomicAdd()进行double，因为cuda不支持double，所以我使用下面的代码，这是NVIDIA提供的。

__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 atomic

【解决方案1】：

基本上是因为实现需要加载，而不能以原子方式执行。比较和交换操作是一个原子版本的

(*address == assumed) ? (assumed + val) : *address

不能保证*address 处的值不会在从*address 加载值的周期与atomicCAS 调用用于存储更新值的周期之间发生变化。如果发生这种情况，*address 的值将不会更新。因此，循环确保重复这两个操作，直到在读取和比较和交换操作之间*address 的值没有变化，这意味着更新发生了。

【讨论】：

谢谢！你的意思是加载操作，assumed=old，不是原子的，所以原子函数 old=atomicCAS(address_as_ull,assumed,__double_as_long(val+__longlong_as_double(assumed))) 启动了， old 的值可能会改变，add() 应该再做一次，保证添加当前值。
没有。 old 是线程局部变量。它的值不会改变，除非本地线程改变它。在线程控制之外，唯一可以更改的值是*address。在运行过程中被其他线程更改时，必须重复atomicCAS调用，否则不会发生更新。
我知道了，循环的功能保证*address_as_ull的当前值的变化是由当前线程完成的。