原子の乗算と除算？

Question

atomicAddとatomicSubがありますが、atomicMulとatomicDivは存在しないようです！出来ますか？次のコードを実装する必要があります。

atomicMul(&accumulation[index],value) 

どうすればいいですか？

Answer 1

[OK]を解決しました。しかし、私は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)の動作は以下の通りです。 *addressからoldに

1. コピー*address（すなわちold = *address）
2. ストア(old == compare ? val : old)。 （この時点では、oldと*addressの値は、条件が一致したか否かによって異なる場合があります。）
3. 戻り、我々は一緒にatomicMul関数の定義を見たときにその行動についての理解oldが

良くなります。私たちが何をしたいのか

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から値を読み込む（その値はaddress_as_ullにeqaulである）、およびそれにいくつかの値を乗算してから戻ってそれを書きます。問題は、他のスレッドがread、modify、writeの間で値*addressにアクセスして変更できることです。

*addressの値がassumedと等しいかどうかを確認するために、他のスレッドの切片がないことを確認します。 assumed=oldValueとoldValue = atomicCAS(...)の後に他のスレッドが*addressの値を変更したとします。 *addressの変更された値は、内の変数oldにコピーされます（上記の動作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型に整数型にフロート/ double型の変数を変換する必要があります。私は上記変更したことintへfloatマッチのサイズのでintにfloatとunsigned long longからdoubleです。