atomicAdd()用于GPU上的double

Question

我正在做一个关于GPU的项目，我必须使用atomicAdd() for 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);
}

现在我想知道为什么实现需要循环，而(假设!=old)

Answer 1

主要是因为实现需要加载，这不能自动执行。比较与交换操作是

的原子版本。

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

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