CUDAアトミック関数
をテンプレートにして作成
[
トップ
|
新規
|
一覧
|
検索
|
最終更新
|
ヘルプ
]
開始行:
*アトミック関数 [#oade1c36]
あるスレッドがグローバルメモリやシェアードメモリ上のデー...
CUDAのアトミック関数はその名前がatomicから始まっているの...
また，signed integerとunsigned integerのみ対応しています(...
各関数は引数として，intもしくはunsigned intのアドレスと値...
(atomicExch()はfloatも，atomicInc()とatomicDec()はunsigne...
例えば，atomicAdd()の場合，関数定義は以下です．
#code(C,nonumber){{
int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address, unsigned i...
}}
address (グローバルまたはシェアードメモリ) の場所から 32b...
読み込み(元の値をoldとする)，その値にvalを足し(old+val)，...
atomicAdd()を含めて以下の11のアトミック関数が用意されてい...
(CUDA2.3)．
-atomicAdd() : old+val
-atomicSub() : old-val
-atomicExch() : old=val
-atomicMin() : min(old, val)
-atomicMax() : max(old, val)
-atomicInc() : (old >= val) ? 0 : (old+1)
-atomicDec() : ( old == 0 | old > val) ? val : (old-1)
-atomicCAS() : (old == compare) ? val : old → address,com...
-atomicAnd() : old & val (ビット演算)
-atomicOr() : old | val (ビット演算)
-atomicXor() : old ^ val (ビット演算)
返値はすべてoldです．
アトミック関数を用いるには最低でもcompute capability 1.1 ...
シェアードメモリで32bitワードの操作とグローバルメモリでの...
シェアードメモリで64bitワードを操作するアトミック関数は c...
float版のアトミック関数はCUDA3.0+Fermiで対応するという話...
それ以外の環境で使いたい場合は以下のような関数(Addの場合)...
([[NVIDIA Forums:http://forums.nvidia.com/index.php?act=S...
#code(C){{
__device__
inline void atomicFloatAdd(float *address, float val)
{
int i_val = __float_as_int(val);
int tmp0 = 0;
int tmp1;
while( (tmp1 = atomicCAS((int *)address, tmp0, i_val)) !...
{
tmp0 = tmp1;
i_val = __float_as_int(val + __int_as_float(tmp1));
}
}
}}
***2011年5月追記 [#r5389465]
compute capability 2.0以上ではfloat版のatomicAddが使える...
また，Programming GuideにもatomicCASを用いてdoubleなどで...
(ver4.0ならp.119参照)．
終了行:
*アトミック関数 [#oade1c36]
あるスレッドがグローバルメモリやシェアードメモリ上のデー...
CUDAのアトミック関数はその名前がatomicから始まっているの...
また，signed integerとunsigned integerのみ対応しています(...
各関数は引数として，intもしくはunsigned intのアドレスと値...
(atomicExch()はfloatも，atomicInc()とatomicDec()はunsigne...
例えば，atomicAdd()の場合，関数定義は以下です．
#code(C,nonumber){{
int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address, unsigned i...
}}
address (グローバルまたはシェアードメモリ) の場所から 32b...
読み込み(元の値をoldとする)，その値にvalを足し(old+val)，...
atomicAdd()を含めて以下の11のアトミック関数が用意されてい...
(CUDA2.3)．
-atomicAdd() : old+val
-atomicSub() : old-val
-atomicExch() : old=val
-atomicMin() : min(old, val)
-atomicMax() : max(old, val)
-atomicInc() : (old >= val) ? 0 : (old+1)
-atomicDec() : ( old == 0 | old > val) ? val : (old-1)
-atomicCAS() : (old == compare) ? val : old → address,com...
-atomicAnd() : old & val (ビット演算)
-atomicOr() : old | val (ビット演算)
-atomicXor() : old ^ val (ビット演算)
返値はすべてoldです．
アトミック関数を用いるには最低でもcompute capability 1.1 ...
シェアードメモリで32bitワードの操作とグローバルメモリでの...
シェアードメモリで64bitワードを操作するアトミック関数は c...
float版のアトミック関数はCUDA3.0+Fermiで対応するという話...
それ以外の環境で使いたい場合は以下のような関数(Addの場合)...
([[NVIDIA Forums:http://forums.nvidia.com/index.php?act=S...
#code(C){{
__device__
inline void atomicFloatAdd(float *address, float val)
{
int i_val = __float_as_int(val);
int tmp0 = 0;
int tmp1;
while( (tmp1 = atomicCAS((int *)address, tmp0, i_val)) !...
{
tmp0 = tmp1;
i_val = __float_as_int(val + __int_as_float(tmp1));
}
}
}}
***2011年5月追記 [#r5389465]
compute capability 2.0以上ではfloat版のatomicAddが使える...
また，Programming GuideにもatomicCASを用いてdoubleなどで...
(ver4.0ならp.119参照)．
ページ名:
![[PukiWiki]](image/pukiwiki.png "[PukiWiki]")
