とりあえずWordPressにCUDAカテゴリを追加。
atomic命令はいわゆるRead-Modify-Writeの命令。
って言ってもわからない人が結構いそうな罠。
要するに、１命令でメモリの読み込みと変更と書き込みを行う命令。複数のスレッドやプロセスが協調して動作するときに、同じメモリアドレスを同じタイミングで書き換えることによる問題（後述）が生じるのを防ぐための命令。
起こりえる問題ってのは、例えば100回加算したはずなのに50回しか反映されていないみたいな問題。複数のスレッドが非同期に読み書きをするとよく起きるので困る。笑えない。
つーわけで、Read-Modify-Writeは並列プログラミングにおける必須アイテム。当然並列実行できないので、使いすぎるとパフォーマンスが落ちる諸刃の剣。それでも素人はとりあえず使いまくった方がバグが出ないという別の意味で諸刃の剣。
CUDA1.0仕様にはatomic命令がないんだけど、CUDA1.1仕様では使えるというので対応GPUを調達。ささっと実験プログラムを組んでみた。
以下実験コード。
- #define NSIZE 1
- #define THREAD_X 10
- #define BLOCK_X 10
- // -------- -------- -------- --------
- __global__ void
- test(int *value)
- {
- #ifdef ATOMICS
- atomicAdd(&value[0], 1);
- #else
- value[0] ++;
- #endif
- }
- // -------- -------- -------- --------
- void
- runTest(int argc, char** argv)
- {
- CUT_DEVICE_INIT();
- int n;
- n = sizeof(int)*NSIZE;
- // CPU側のデータを準備
- int *h_sd1;
- h_sd1 = (int*)malloc(n);
- int i;
- for(i=0; i<NSIZE; i++){
- h_sd1[i] = 0;
- }
- // 値の確認
- for(i=0; i<NSIZE; i++){
- printf(" %d", h_sd1[i]);
- }
- printf("\n");
- // GPU側のメモリを準備
- int* d_sd1;
- CUDA_SAFE_CALL(cudaMalloc((void**)&d_sd1, n));
- CUDA_SAFE_CALL(cudaMemcpy(d_sd1, h_sd1, n, cudaMemcpyHostToDevice) );
- // 実行のためのパラメタを準備
- dim3 threads(THREAD_X, 1, 1);
- dim3 grid(BLOCK_X,1,1);
- // 実行
- test<<<grid, threads>>>(d_sd1);
- // GPUの処理に問題が起きていないかの確認
- CUT_CHECK_ERROR("Kernel execution failed");
- // 演算結果の取得
- CUDA_SAFE_CALL(cudaMemcpy(h_sd1, d_sd1, n, cudaMemcpyDeviceToHost) );
- // 演算結果の確認
- for(i=0; i<NSIZE; i++){
- printf(" %d", h_sd1[i]);
- }
- printf("\n");
- // クリーンアップ
- free(h_sd1);
- CUDA_SAFE_CALL(cudaFree(d_sd1));
- }
- int
- main(int argc, char** argv)
- {
- runTest(argc, argv);
- CUT_EXIT(argc, argv);
- }
なんか無駄のあるコードだけど、だいたいこんなもん。
atomic命令はglobal memoryに対して有効だということで、手っ取り早く__global__関数の引数をターゲットにしてみた。
これでスレッドだろうがブロックだろうが跨ってRead-Modify-Writeが出来る模様。簡単。
私の環境（GeForce8600GTS）では、#define ATOMICSした状態だと100が出力されたのに対して、そうでない場合は1が出力された。どの環境でもatomic命令を使っていないときに1になるかはわからない（少なくとも100にはならないだろうけど）ので注意。
ちなみに、CUDA1.1向けのバイナリ生成時には -arch sm_11 が必要。忘れないように。