atomicAdd（）在GPU上加倍

我正在GPU上做一個項目，我必須使用atomicAdd（）來加倍，因爲cuda不支持加倍，所以我使用下面的代碼，這是NVIDIA提供的。atomicAdd（）在GPU上加倍

__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); 
}

現在我想知道爲什麼工具需要一個循環，而（假定！=舊）

來源

2013-04-18 taoyuanjl

基本上是因爲實施需要一個負載，不能原子方式執行。在比較並交換操作的

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

原子的版本也不能保證在address的值將不是值從address加載的週期，而且atomicCAS呼叫週期之變化用於存儲更新的值。如果發生這種情況，address的值將不會更新。因此，循環確保重複執行這兩個操作，直到在讀取操作和比較與交換操作之間address的值沒有變化，這意味着發生了更新。

來源

2013-04-18 10:15:08 talonmies

感謝！你的意思是負荷運轉，假定=老了，是不是原子，因此原子函數舊= atomicCAS（address_as_ull，假設，__ double_as_long（VAL + __ longlong_as_double（假設）））推出後，舊的可能值被修改，並且add（）應該再次執行，保證當前值被添加。 – taoyuanjl

編號'old'是一個線程局部變量。它的值不會改變，除非本地線程改變它。在線程控制之外，唯一可以改變的值是'* address'。當在操作過程中被另一個線程更改時，必須重複'atomicCAS'調用，否則不會發生更新。 – talonmies

我知道了，循環函數保證* address_as_ull的當前變化值是由當前線程完成的。 – taoyuanjl

atomicAdd（）在GPU上加倍

回答

相關問題