cuda - atomicAdd()用于在GPU上加倍

Question

我正在GPU上做一个项目，我必须使用atomicAdd()来实现两倍，因为cuda不支持它来实现两倍，因此我使用了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处的值没有变化为止，这意味着已进行了更新。