atomicAdd（）for double for GPU

Question

私はGPUでプロジェクトをやっていますが、cudaはdoubleをサポートしていないため、atomicAdd（）をdoubleに使う必要があります。 。

__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の値は更新されません。したがって、ループは、読み取りと比較およびスワップ操作の間の値の変更がなくなるまで（つまり、更新が行われたことを意味する）、2つの操作が繰り返されることを保証します。