atomicAdd()for 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);
}

现在我想知道为什么工具需要一个循环,而(假设!=旧)

Answer 1

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

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

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