原子的乘法和除法？

Question

有atomicAdd和atomicSub，但似乎atomicMul和atomicDiv不存在！可能嗎？我需要執行以下代碼：

atomicMul(&accumulation[index],value) 

我該怎麼辦？

Answer 1

好的，我解決了。但我無法理解atomicMul是如何工作的，我不知道如何爲浮點數編寫它。

#include <stdio.h> 
#include <cuda_runtime.h> 

__device__ double atomicMul(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); 
}  
__global__ void try_atomicMul(double* d_a, double* d_out) 
{ 
    atomicMul(d_out,d_a[threadIdx.x]); 
} 
int main() 
{ 
    double h_a[]={5,6,7,8}, h_out=1; 
    double *d_a, *d_out; 

cudaMalloc((void **)&d_a, 4 * sizeof(double)); 
cudaMalloc((void **)&d_out,sizeof(double)); 

cudaMemcpy(d_a, h_a, 4 * sizeof(double),cudaMemcpyHostToDevice); 
cudaMemcpy(d_out, &h_out, sizeof(double),cudaMemcpyHostToDevice); 

dim3 blockDim(4); 
dim3 gridDim(1); 

    try_atomicMul<<<gridDim, blockDim>>>(d_a,d_out); 
cudaMemcpy(&h_out, d_out, sizeof(double), cudaMemcpyDeviceToHost); 

printf("%f \n",h_out); 
cudaFree(d_a); 
return 0; 
} 

Answer 2

我會根據我對atomicCAS的理解來補充horus的答案。我的答案可能是錯誤的細節，因爲我沒有看到atomicCAS函數，但只是閱讀有關它的文檔（atomicCAS,Atomic Functions）。隨時解決我的答案。

atomicMul如何工作

據筆者瞭解，中atomicCAS(int* address, int compare, int val)行爲是下面。

1. 複製*address爲old（即old = *address）
2. 商店(old == compare ? val : old)到*address。 （在這一點上，old和*address值可以根據如果條件符合與否有所不同。）
3. 返回old

瞭解關於它的行爲變得更好，當我們在atomicMul函數的定義一起來看一下。

unsigned long long int* address_as_ull = (unsigned long long int*)address; 
unsigned long long int oldValue = *address_as_ull, assumed; // Modified the name 'old' to 'oldValue' because it can be confused with 'old' inside the atomicCAS. 
do { 
    assumed = oldValue; 
    // other threads can access and modify value of *address_as_ull between upper and lower line. 
    oldValue = atomicCAS(address_as_ull, assumed, __double_as_longlong(val * 
         __longlong_as_double(assumed))); 
} while (assumed != oldValue); return __longlong_as_double(oldValue); 

我們想要做的是讀取address值（其值是eqaul到address_as_ull），繁衍一定的參考價值，然後將它寫回。問題是其他線程可以在讀取，修改和寫入之間訪問和修改*address的值。

爲了確保沒有其他線程的攔截，我們檢查*address的值是否等於我們的assumed在那裏。假設assumed=oldValue和oldValue = atomicCAS(...)之後的其他線程的修改值爲*address。修改後的*address的值將被複制到atomicCAS內的變量old（參見上面的atomicCAS的行爲1）。 由於atomicCAS更新*address根據*address = (old == compare ? val : old)，*address將不會更改（old==*address）。

然後atomicCAS返回old它進入oldValue，這樣循環可以繼續，我們可以在下一次迭代時嘗試另一個鏡頭。當*address在讀取和寫入之間未被修改時，val被寫入*address，並且循環將結束。

怎麼寫呢浮法

簡短的回答：

__device__ float atomicMul(float* address, float val) 
{ 
    int* address_as_int = (int*)address; 
    int old = *address_as_int, assumed; 
    do { 
    assumed = old; 
    old = atomicCAS(address_as_int, assumed, __float_as_int(val * 
__float_as_int(assumed))); 
} while (assumed != old); return __int_as_float(old); 
} 

我沒有測試，所以可以有一些錯誤。修復我，如果我錯了。

它是如何工作： 出於某種原因，atomicCAS只支持整數類型。因此，我們應該手動將float/double類型變量轉換爲整數類型以輸入到該函數，然後將整數結果重新轉換爲float/double類型。我上面修改的是double到float和unsigned long long到int，因爲float的尺寸匹配到int。