CUDAアトミック関数

アトミック関数†

あるスレッドがグローバルメモリやシェアードメモリ上のデータを読み込み，修正し，書き込む(read-modify-write)という一連の処理を行うとき，その処理中にそのメモリ領域に他のスレッドが書き込みが行われないようにしたい場合があります．CUDAにはこれを保証するアトミック関数が用意されており，32bitまたは64bitワード長のread-modify-write操作中に他のスレッドからの干渉を防いでくれます．

CUDAのアトミック関数はその名前がatomicから始まっているのでわかりやすいと思います．また，signed integerとunsigned integerのみ対応しています(atomicExch()を除く)．各関数は引数として，intもしくはunsigned intのアドレスと値を取ります (atomicExch()はfloatも，atomicInc()とatomicDec()はunsigned intのみ)．例えば，atomicAdd()の場合，関数定義は以下です．

nonumber

address (グローバルまたはシェアードメモリ) の場所から 32bit もしくは 64bit ワードを読み込み(元の値を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,compare,valの3引数を取る．CASはCompare And Swapの略．
atomicAnd() : old & val (ビット演算)
atomicOr() : old | val (ビット演算)
atomicXor() : old ^ val (ビット演算)

返値はすべてoldです．

アトミック関数を用いるには最低でもcompute capability 1.1 が必要で，シェアードメモリで32bitワードの操作とグローバルメモリでの64bitワード操作のアトミック関数は compute capability 1.2以上，シェアードメモリで64bitワードを操作するアトミック関数は compute capability 2.0以上が必要です．

float版のアトミック関数はCUDA3.0+Fermiで対応するという話もありますが，それ以外の環境で使いたい場合は以下のような関数(Addの場合)を使うという手もあります (NVIDIA Forumsより転載)．

__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)
	{
		tmp0 = tmp1;
		i_val = __float_as_int(val + __int_as_float(tmp1));
	}
}

↑

2011年5月追記†

compute capability 2.0以上ではfloat版のatomicAddが使えるようになっています(少なくともCUDA3.2以上なら対応している様子)．また，Programming GuideにもatomicCASを用いてdoubleなどでアトミック関数を使うための方法が書かれています (ver4.0ならp.119参照)．

Menu

アトミック関数†

2011年5月追記†