cuda - 在 GPU 上用于双精度的 atomicAdd()

Question

我在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);
}

现在我想知道为什么该工具需要一个循环，而（假设！=旧）

Answer 1

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

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

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