爲什麼atomicAdd未實現雙打？

爲什麼hasnt atomicAdd()作爲CUDA 4.0或更高版本的一部分明確實施了雙打？爲什麼atomicAdd未實現雙打？

從的附錄F頁97開始，下列版本的 atomicAdd已經實現。

int atomicAdd(int* address, int val); 
unsigned int atomicAdd(unsigned int* address, 
         unsigned int val); 
unsigned long long int atomicAdd(unsigned long long int* address, 
           unsigned long long int val); 
float atomicAdd(float* address, float val)

在同一頁接着給出一個小的實現atomicAdd的雙打如下我已經用在我的項目剛剛開始。

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

爲什麼不把上面的代碼定義爲CUDA的一部分？

來源

2012-09-27 smilingbuddha

可能致使其每個用戶知道它的實現，因爲它不是一個內置的指令和重試邏輯可以承受活鎖（因爲沒有公平的保證，一個線程可以得到停滯不前的只要有其他線程更新相同的變量）。 – tera

編輯：正如CUDA 8，雙精度atomicAdd()在CUDA實現與SM_6X（帕斯卡）的GPU的硬件支持。

~~目前，沒有CUDA設備在double硬件中支持atomicAdd。~~正如您所指出的那樣，它可以在64位整數上以atomicCAS的形式實現，但是這對性能成本來說是不平凡的。

因此，CUDA軟件團隊選擇將正確的實現文檔記錄爲開發人員的選項，而不是將其作爲CUDA標準庫的一部分。這樣開發人員不會在不知不覺中選擇性能成本。另外：我不認爲這個問題應該被視爲「不具有建設性」。我認爲這是一個完全有效的問題，+1。

來源

2012-09-27 23:59:23 harrism

是的，但從技術上講，你是少數人之一，可能會回答這個問題。雖然我已經說過爲什麼我認爲這樣做很有意義，但只有您可以說這是CUDA團隊如此選擇它的原因。 ;-)無論如何，我不是一個倒下了這個問題的人。 – tera

還有誰閱讀並回答SO CUDA疑問多個NVIDIA人們（尤其是在我們的開發者論壇下降），而事實使得這樣有效的問題。你可以發表你的評論作爲答案，這將是正確的，我會投票贊成。順便說一句，我沒有假設你低調;我指的是一個投票來結束這個問題。 – harrism

我同意，這是一個完全有效的問題，CUDA頭文件可以在軟件中實現雙原子。雖然它制定的方式引發了紅光對某些人來說，我認爲決定應該恢復！ – pszilard

爲什麼atomicAdd未實現雙打？

回答

相關問題