ひびろぐ ver.2

Google検索したときの見た目を考慮してブログ名を変更してみた。馬鹿かと阿呆かと。

Archive for the ‘CUDA’

Pages: ◄ 1 2

条件分岐がシリアライズされるthread同期とかいう罠

Published 10月 9th, 2007 in CUDA | コメント（0） »

CUDAいじってて嵌ったのでメモ。

PLAIN TEXT
CUDA:

  1. // grid(1,1,1) thread(2,1,1) で実行、_buf[0]の初期値は1

  2.   if(threadIdx.x==0){

  3.     atomicAdd(&_buf[0],10);

  4.     __syncthreads();

  5.     mem0[0] = _buf[0];

  6.   }else if(threadIdx.x==1){

  7.     atomicAdd(&_buf[0],100);

  8.     __syncthreads();

  9.     mem1[0] = _buf[0];

  10.   }

  11.   return;


一見すると_buf[0]もmem0[0]もmem1[0]も111になりそうなコード。

でもmem1[0]は101になっちゃう。

何を間違えていたかというと、CUDAのthread内では条件分岐系が本当の分岐じゃないというか、マスクで潰されるだけで処理されるというか、そんな感じなのをちゃんと理解していなかったというお話。

今回の場合、ifの分岐は本当に分岐するわけではなく、マスクを生成して全スレッドが同じ処理をするわけだ。

コンパイル結果の都合だろうけど、今回はelse ifの方が先に実行されるというバイナリ（アセンブラ？）が出力されることになったらしく、

PLAIN TEXT
CUDA:

  1. atomicAdd(&_buf[0],100);

  2.     __syncthreads();

  3.     mem1[0] = _buf[0];


を各スレッドが実行。ただしthread0は演算結果を保持しない。この時点でmem1[0]=_buf[0]は101。

続いて

PLAIN TEXT
CUDA:

  1. atomicAdd(&_buf[0],10);

  2.     __syncthreads();

  3.     mem0[0] = _buf[0];


を各スレッドが実行。ただしthread1は演算結果を保持しない。この時点でmem0[0]=_buf[0]は111。


まぁこんな感じ。

なるほど納得。

まぁ今回の場合は

PLAIN TEXT
CUDA:

  1. if(threadIdx.x==0){

  2.     atomicAdd(&_buf[0],10);

  3.   }else if(threadIdx.x==1){

  4.     atomicAdd(&_buf[0],100);

  5.   }

  6.   __syncthreads();

  7.   if(threadIdx.x==0){

  8.     mem0[0] = _buf[0];

  9.   }else if(threadIdx.x==1){

  10.     mem1[0] = _buf[0];

  11.   }

  12.   return;


って書けば良かったわけだな。


なんかソース内の空白がおかしくなった気がするけど、気にしない方向で。

CUDAのatomic命令を試してみた

Published 10月 1st, 2007 in CUDA | コメント（0） »

とりあえずWordPressにCUDAカテゴリを追加。


atomic命令はいわゆるRead-Modify-Writeの命令。

って言ってもわからない人が結構いそうな罠。

要するに、１命令でメモリの読み込みと変更と書き込みを行う命令。複数のスレッドやプロセスが協調して動作するときに、同じメモリアドレスを同じタイミングで書き換えることによる問題（後述）が生じるのを防ぐための命令。

起こりえる問題ってのは、例えば100回加算したはずなのに50回しか反映されていないみたいな問題。複数のスレッドが非同期に読み書きをするとよく起きるので困る。笑えない。

つーわけで、Read-Modify-Writeは並列プログラミングにおける必須アイテム。当然並列実行できないので、使いすぎるとパフォーマンスが落ちる諸刃の剣。それでも素人はとりあえず使いまくった方がバグが出ないという別の意味で諸刃の剣。

CUDA1.0仕様にはatomic命令がないんだけど、CUDA1.1仕様では使えるというので対応GPUを調達。ささっと実験プログラムを組んでみた。

以下実験コード。

PLAIN TEXT
CUDA:

  1. #define NSIZE   1

  2. #define THREAD_X        10

  3. #define BLOCK_X         10

  4.  

  5. // -------- -------- -------- --------

  6. __global__ void

  7. test(int *value)

  8. {

  9. #ifdef ATOMICS

  10.   atomicAdd(&value[0], 1);

  11. #else

  12.   value[0] ++;

  13. #endif

  14. }

  15. // -------- -------- -------- --------

  16.  

  17. void

  18. runTest(int argc, char** argv)

  19. {

  20.   CUT_DEVICE_INIT();

  21.  

  22.   int n;

  23.   n = sizeof(int)*NSIZE;

  24.  

  25.   // CPU側のデータを準備

  26.   int *h_sd1;

  27.   h_sd1 = (int*)malloc(n);

  28.   int i;

  29.   for(i=0; i<NSIZE; i++){

  30.     h_sd1[i] = 0;

  31.   }

  32.   // 値の確認

  33.   for(i=0; i<NSIZE; i++){

  34.     printf(" %d", h_sd1[i]);

  35.   }

  36.   printf("\n");

  37.  

  38.   // GPU側のメモリを準備

  39.   int* d_sd1;

  40.   CUDA_SAFE_CALL(cudaMalloc((void**)&d_sd1, n));

  41.   CUDA_SAFE_CALL(cudaMemcpy(d_sd1, h_sd1, n, cudaMemcpyHostToDevice) );

  42.  

  43.   // 実行のためのパラメタを準備

  44.   dim3 threads(THREAD_X, 1, 1);

  45.   dim3 grid(BLOCK_X,1,1);

  46.  

  47.   // 実行

  48.   test<<<grid, threads>>>(d_sd1);

  49.  

  50.   // GPUの処理に問題が起きていないかの確認

  51.   CUT_CHECK_ERROR("Kernel execution failed");

  52.  

  53.   // 演算結果の取得

  54.   CUDA_SAFE_CALL(cudaMemcpy(h_sd1, d_sd1, n, cudaMemcpyDeviceToHost) );

  55.  

  56.   // 演算結果の確認

  57.   for(i=0; i<NSIZE; i++){

  58.     printf(" %d", h_sd1[i]);

  59.   }

  60.   printf("\n");

  61.  

  62.   // クリーンアップ

  63.   free(h_sd1);

  64.   CUDA_SAFE_CALL(cudaFree(d_sd1));

  65. }

  66.  

  67.  

  68. int

  69. main(int argc, char** argv)

  70. {

  71.   runTest(argc, argv);

  72.  

  73.   CUT_EXIT(argc, argv);

  74. }


なんか無駄のあるコードだけど、だいたいこんなもん。

atomic命令はglobal memoryに対して有効だということで、手っ取り早く__global__関数の引数をターゲットにしてみた。

これでスレッドだろうがブロックだろうが跨ってRead-Modify-Writeが出来る模様。簡単。

私の環境（GeForce8600GTS）では、#define ATOMICSした状態だと100が出力されたのに対して、そうでない場合は1が出力された。どの環境でもatomic命令を使っていないときに1になるかはわからない（少なくとも100にはならないだろうけど）ので注意。


ちなみに、CUDA1.1向けのバイナリ生成時には -arch sm_11 が必要。忘れないように。

Pages: ◄ 1 2