ＣＵＤＡプログラミングガイドより意訳：block間での同期は取れないよ
流石に困るので、なんとかして同期っぽいことができないか試してみた。
考え方は以下：
global memory space内の変数をフラグとして利用する。以上。
適当に実装してみたが、とりあえずそれっぽく動いてくれた。
まぁ同期とはちょっと違うんだけど。
CUDA:
- // -*- C++ -*-
- /*
- block制御のテスト
- 同期のようなものを作成して逐次処理を行う
- */
- #include <stdlib.h>
- #include <stdio.h>
- #include <string.h>
- #include <math.h>
- #include <unistd.h>
- #include <cutil.h>
- #define N 4
- __device__ int counter=0; // 逐次処理のためのカウンタ
- __global__ void gpu
- (int* answer)
- {
- int id = blockIdx.x;
- // カウンタを監視して自分の順番を待つ
- while(counter != id){
- __syncthreads();
- }
- *answer = (*answer*id) + (id+1);
- counter += 1;
- }
- void
- runTest(int argc, char** argv)
- {
- CUT_DEVICE_INIT();
- int answer=0;
- int* d_a;
- CUDA_SAFE_CALL(cudaMalloc((void**)&d_a, sizeof(int)));
- CUDA_SAFE_CALL(cudaMemcpy(d_a, &answer, sizeof(int), cudaMemcpyHostToDevice) );
- dim3 grid(N,1,1);
- dim3 threads(1,1,1);
- gpu<<<grid, threads>>>(d_a);
- CUDA_SAFE_CALL( cudaThreadSynchronize() );
- CUT_CHECK_ERROR("Kernel execution failed");
- CUDA_SAFE_CALL(cudaMemcpy(&answer, d_a, sizeof(int), cudaMemcpyDeviceToHost) );
- CUDA_SAFE_CALL(cudaFree(d_a));
- printf("answer %d\n", answer);
- }
- int
- main(int argc, char** argv)
- {
- runTest(argc, argv);
- CUT_EXIT(argc, argv);
- }
カウンタを監視して自分の順番を待つ、ってところがあるおかげで正答である31が得られる。これがないと変な値が得られてしまう。
次はバリア同期を実装だな。