このへんの続き。
前回は何故か__syncthreads()を使ったりと妙にキモイことをやっていたので、追実験。
atomic関数の存在も念頭に置いてみる。
まぁこんなケースを考えることにする。
CUDA:
- #define BLOCK_X 8 // grid(8,1,1) を意味する
- __device__ int waitflag[8] = {0,0,0,0,0,0,0,0};
- __global__ void gpu
- (int *mem0, int *mem1, unsigned int *_buf)
- {
- int id = blockIdx.x;
- // 実行時間をずらすための処理
- if(id==3){
- mem1[0]=0;
- mem1[1]=0;
- mem1[2]=0;
- }
- // 各blockで_buf[0]を叩く
- atomicAdd(&_buf[0], id);
- // このへんで同期を取れないと、計算結果がblock毎に狂う可能性がある。
- // 具体的には、id3以外のblockでは
- // id3のatomicAddが終わる前に_buf[0]を読んでしまう可能性が高い。
- // 恐らく、mem0は 26 26 26 29 26 26 26 26 になる（各要素ともデフォルトが1なので）。
- // 各blockからみえる計算結果を、
- // global変数に書き戻して確認
- mem0[id] = _buf[0];
- return;
- }
んで同期をどうやるか。
とりあえず考えたのが以下の実装。
CUDA:
- atomicAdd(&waitflag[0], 1);
- while(waitflag[0]<BLOCK_X);
atomicAddで加算の保証が出来るので、あとは全blockが加算するのを待てばいいんじゃないかと。
うん。失敗した。無限ループっぽい。
直感的にはこれで行けそうなんだけど、何かひっかかるらしい。
CUDA:
- atomicAdd(&waitflag[0], 1);
- while(waitflag[0]<BLOCK_X)__syncthreads();
という感じに__syncthreads();を入れると正しく動作する。
__syncthread()にそれらしい仕様はない気がするんだが……？
ところで、atomicAddは確実に読み出して加算ができる、ってのはいいんだけど、その処理の途中で別のblock/threadが値を読んでしまうことってあるんだろうか。いや、別に「0が読めるか1が読めるかわからないよ」ってのは構わないというか当然なんだけど、何故か全然違う値が読めたりすると危険だよなぁ……。
さて、胡散臭い__syncthreads()を使わずにちゃんとやるにはどうすればいいかなんだが、作業用の変数を追加して
CODE:
- __device__ int waitflag[8] = {0,0,0,0,0,0,0,0}; // 関数外宣言
- int wait = 0; // 関数内宣言
- // 計算終了通知
- waitflag[id] = 1;
- // 全blockが計算終了に到達していることを確認する
- while(wait <BLOCK_X){
- wait = 0;
- for(int i=0; i<BLOCK_X; i++){
- wait += waitflag[i];
- }
- }
とやったらうまくいった。
安全を確保するならwaitflagの変更はatomicの方がいいのかもしれないけど、まぁ大丈夫だろう。
さて。以上の振る舞いから、今回のglobal memory spaceの変数の参照はコンパイラによって最適化されてしまっているのではないかという仮説が立てられる。
nvccの-ptxや-keepで中間コードが覗けるんだけど、これ見て判断するのはちょっとツライんだよなぁ。
最後。
仮設を踏まえて適当にうまく行く記述を考えてみた。
CUDA:
- atomicAdd(&waitflag[0], 1);
- while(wait<BLOCK_X){wait=waitflag[0];}
明示的にglobal memory spaceからの再読み込みを行うような感じ。
実際、これで思い通りの動作は出来ている。
ちなみに、どの実装にしても複数回の実行にはちょっとだけ工夫が必要かもしれない。
作業用変数のクリアが必要だから。
とはいえ、作業用変数を多重化して、同期が終了するときに今使ったものとは別の作業用変数を初期化してやれば動きそうだから多分大丈夫だろう。たぶん。
以下、#ifで切ったプログラム全体。長いけど参考にベタッと貼っておく。
CUDA:
- // -*- C++ -*-
- /*
- block間の同期を行ってみるテスト
- */
- #include <stdlib.h>
- #include <stdio.h>
- #include <string.h>
- #include <math.h>
- #include <unistd.h>
- #include <cutil.h>
- #define N 8
- #define BLOCK_X N
- #define THREAD_X 1
- //extern __device__ int* buf;
- __device__ int waitflag[8] = {0,0,0,0,0,0,0,0};
- __global__ void gpu
- (int *mem0, int *mem1, unsigned int *_buf)
- {
- // -------- -------- -------- --------
- int wait = 0;
- int id = blockIdx.x;
- // 実行時間をずらすための処理
- if(id==3){
- mem1[0]=0;
- mem1[1]=0;
- mem1[2]=0;
- }
- // 各blockで_buf[0]を叩く
- atomicAdd(&_buf[0], id);
- #if SYNC==3
- // 計算終了通知
- waitflag[id] = 1;
- // 全blockが計算終了に到達していることを確認する
- while(wait <BLOCK_X){
- wait = 0;
- for(int i=0; i<BLOCK_X; i++){
- wait += waitflag[i];
- }
- }
- #elif SYNC==4
- atomicAdd(&waitflag[0], 1);
- while(wait<BLOCK_X){wait=waitflag[0];}
- #elif SYNC==2
- atomicAdd(&waitflag[0], 1);
- while(waitflag[0]<BLOCK_X)__syncthreads();
- #else // NG
- atomicAdd(&waitflag[0], 1);
- while(waitflag[0]<BLOCK_X);//__syncthreads();
- #endif
- // ここから、各blockが計算結果を確実に取得できる
- // global変数に書き戻して確認
- mem0[id] = _buf[0];
- // 待ち合わせデータのチェックのため、
- // global変数に書き戻して確認
- mem1[id] = wait;
- return;
- // -------- -------- -------- --------
- }
- // -------- -------- -------- -------- -------- -------- -------- --------
- void
- runTest(int argc, char** argv)
- {
- CUT_DEVICE_INIT();
- int* h_mem0;
- int* h_mem1;
- int* d_mem0;
- int* d_mem1;
- h_mem0 = (int*)malloc(sizeof(int)*N);
- h_mem1 = (int*)malloc(sizeof(int)*N);
- CUDA_SAFE_CALL(cudaMalloc((void**)&d_mem0, sizeof(int)*N));
- CUDA_SAFE_CALL(cudaMalloc((void**)&d_mem1, sizeof(int)*N));
- //CUDA_SAFE_CALL(cudaMemcpy(d_mem0, h_mem0, sizeof(int)*N, cudaMemcpyHostToDevice) );
- //CUDA_SAFE_CALL(cudaMemcpy(d_mem1, h_mem1, sizeof(int)*N, cudaMemcpyHostToDevice) );
- unsigned int* h_buf;
- unsigned int* d_buf;
- CUDA_SAFE_CALL(cudaMallocHost((void**)&h_buf, sizeof(int)*N));
- int n;
- for(n=0; n<N; n++){
- h_buf[n] = 1;
- }
- CUDA_SAFE_CALL(cudaMalloc((void**)&d_buf, sizeof(int)*N));
- CUDA_SAFE_CALL(cudaMemcpy(d_buf, h_buf, sizeof(int)*N, cudaMemcpyHostToDevice) );
- //CUDA_SAFE_CALL(cudaMemset(buf,0,sizeof(int)*N));
- dim3 grid(BLOCK_X,1,1);
- dim3 threads(THREAD_X,1,1);
- gpu<<<grid, threads>>>(d_mem0, d_mem1, d_buf);
- CUDA_SAFE_CALL( cudaThreadSynchronize() );
- CUT_CHECK_ERROR("Kernel execution failed");
- CUDA_SAFE_CALL(cudaMemcpy(h_mem0, d_mem0, sizeof(int)*N, cudaMemcpyDeviceToHost) );
- CUDA_SAFE_CALL(cudaMemcpy(h_mem1, d_mem1, sizeof(int)*N, cudaMemcpyDeviceToHost) );
- CUDA_SAFE_CALL(cudaMemcpy(h_buf, d_buf, sizeof(int)*N, cudaMemcpyDeviceToHost) );
- CUDA_SAFE_CALL(cudaFree(d_mem0));
- CUDA_SAFE_CALL(cudaFree(d_mem1));
- CUDA_SAFE_CALL(cudaFree(d_buf));
- int i;
- printf("mem0:");
- for(i=0; i<N; i++){
- printf(" %d", h_mem0[i]);
- }
- printf("\n");
- printf("mem1:");
- for(i=0; i<N; i++){
- printf(" %d", h_mem1[i]);
- }
- printf("\n");
- printf("buf:");
- for(i=0; i<N; i++){
- printf(" %d", h_buf[i]);
- }
- printf("\n");
- CUDA_SAFE_CALL(cudaFreeHost(h_buf));
- free(h_mem0);
- free(h_mem1);
- }
- int
- main(int argc, char** argv)
- {
- runTest(argc, argv);
- CUT_EXIT(argc, argv);
- }