ひびろぐ ver.2

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

Archive for 10月 10th, 2007


CUDAコンパイラ（ｒｙ　補足

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

変数idを__shared__にすると思い通りに動くよ！


　い　み　が　。

CUDAコンパイラの最適化が腐っている　のか？

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

アレな状況に陥ったので報告。本家forumで議論してきたほうがいいかもシレン。

PLAIN TEXT
CUDA:

  1. __global__ void

  2. test(int *value)

  3. {

  4.   int id;

  5.   id = blockIdx.x;

  6.   if(id==0){

  7.     value[0] = 1;

  8.   }else if(id==1){

  9.     value[1] = 2;

  10.   }

  11. }


grid(2,1,1)な実行後のvalueは何になるか。

答えは

value[0] = 0

value[1] = 0

……いや、それはおかしいだろう。

実は、例えば以下のように書くと想定どおりの答えが得られる。

PLAIN TEXT
CUDA:

  1. __global__ void

  2. test(int *value)

  3. {

  4.   int id;

  5.   id = blockIdx.x;

  6.   if(id==0){

  7.     value[0] = 1;

  8.   }else if(id==1){

  9.     value[id] = 2;

  10.   }

  11. }


これで答えは

value[0] = 1

value[1] = 2

となる。

PLAIN TEXT
CUDA:

  1. __global__ void

  2. test(int *value)

  3. {

  4.   int id;

  5.   id = blockIdx.x;

  6.   if(id==0){

  7.     value[id] = 1;

  8.   }else if(id==1){

  9.     value[1] = 2;

  10.   }

  11. }


でもいいや。

PLAIN TEXT
CUDA:

  1. __global__ void

  2. test(int *value)

  3. {

  4.   int id;

  5.   id = blockIdx.x;

  6.   if(id==0){

  7.     value[0] = id;

  8.   }else if(id==1){

  9.     value[1] = 2;

  10.   }

  11. }


なんてのもありらしい。

あくまで予想なんだけど、ifブロックの中で変数idが使われていないという理由で、コンパイラがid変数を最適化の課程で殺しているんじゃないだろうか。だとしたらコンパイラのバグだよなぁ……。

やっぱ本家にゴルァかな？勘違いかもしれないので一日熟成させよう。

block間での同期を取る3

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

このへんの続き。

前回は何故か__syncthreads()を使ったりと妙にキモイことをやっていたので、追実験。

atomic関数の存在も念頭に置いてみる。

まぁこんなケースを考えることにする。

PLAIN TEXT
CUDA:

  1. #define BLOCK_X 8      // grid(8,1,1) を意味する

  2.  

  3. __device__ int waitflag[8] = {0,0,0,0,0,0,0,0};

  4.  

  5. __global__ void gpu

  6. (int *mem0, int *mem1, unsigned int *_buf)

  7. {

  8.   int id = blockIdx.x;

  9.   // 実行時間をずらすための処理

  10.   if(id==3){

  11.     mem1[0]=0;

  12.     mem1[1]=0;

  13.     mem1[2]=0;

  14.   }

  15.   // 各blockで_buf[0]を叩く

  16.   atomicAdd(&_buf[0], id);

  17.  

  18.   // このへんで同期を取れないと、計算結果がblock毎に狂う可能性がある。

  19.   // 具体的には、id3以外のblockでは

  20.   // id3のatomicAddが終わる前に_buf[0]を読んでしまう可能性が高い。

  21.   // 恐らく、mem0は 26 26 26 29 26 26 26 26 になる（各要素ともデフォルトが1なので）。

  22.  

  23.   // 各blockからみえる計算結果を、

  24.   // global変数に書き戻して確認

  25.   mem0[id] = _buf[0];

  26.  

  27.   return;

  28. }


んで同期をどうやるか。

とりあえず考えたのが以下の実装。

PLAIN TEXT
CUDA:

  1. atomicAdd(&waitflag[0], 1);

  2.   while(waitflag[0]<BLOCK_X);


atomicAddで加算の保証が出来るので、あとは全blockが加算するのを待てばいいんじゃないかと。

うん。失敗した。無限ループっぽい。

直感的にはこれで行けそうなんだけど、何かひっかかるらしい。

PLAIN TEXT
CUDA:

  1. atomicAdd(&waitflag[0], 1);

  2.   while(waitflag[0]<BLOCK_X)__syncthreads();


という感じに__syncthreads();を入れると正しく動作する。

__syncthread()にそれらしい仕様はない気がするんだが……？

ところで、atomicAddは確実に読み出して加算ができる、ってのはいいんだけど、その処理の途中で別のblock/threadが値を読んでしまうことってあるんだろうか。いや、別に「0が読めるか1が読めるかわからないよ」ってのは構わないというか当然なんだけど、何故か全然違う値が読めたりすると危険だよなぁ……。

さて、胡散臭い__syncthreads()を使わずにちゃんとやるにはどうすればいいかなんだが、作業用の変数を追加して

PLAIN TEXT
CODE:

  1. __device__ int waitflag[8] = {0,0,0,0,0,0,0,0}; // 関数外宣言

  2.   int wait = 0; // 関数内宣言

  3.  

  4.   // 計算終了通知

  5.   waitflag[id] = 1;

  6.   // 全blockが計算終了に到達していることを確認する

  7.   while(wait <BLOCK_X){

  8.     wait = 0;

  9.     for(int i=0; i<BLOCK_X; i++){

  10.       wait += waitflag[i];

  11.     }

  12.   }


とやったらうまくいった。

安全を確保するならwaitflagの変更はatomicの方がいいのかもしれないけど、まぁ大丈夫だろう。

さて。以上の振る舞いから、今回のglobal memory spaceの変数の参照はコンパイラによって最適化されてしまっているのではないかという仮説が立てられる。

nvccの-ptxや-keepで中間コードが覗けるんだけど、これ見て判断するのはちょっとツライんだよなぁ。

最後。

仮設を踏まえて適当にうまく行く記述を考えてみた。

PLAIN TEXT
CUDA:

  1. atomicAdd(&waitflag[0], 1);

  2.   while(wait<BLOCK_X){wait=waitflag[0];}


明示的にglobal memory spaceからの再読み込みを行うような感じ。

実際、これで思い通りの動作は出来ている。

ちなみに、どの実装にしても複数回の実行にはちょっとだけ工夫が必要かもしれない。

作業用変数のクリアが必要だから。

とはいえ、作業用変数を多重化して、同期が終了するときに今使ったものとは別の作業用変数を初期化してやれば動きそうだから多分大丈夫だろう。たぶん。

以下、#ifで切ったプログラム全体。長いけど参考にベタッと貼っておく。

PLAIN TEXT
CUDA:

  1. // -*- C++ -*-

  2. /*

  3.   block間の同期を行ってみるテスト

  4. */

  5.  

  6. #include <stdlib.h>

  7. #include <stdio.h>

  8. #include <string.h>

  9. #include <math.h>

  10. #include <unistd.h>

  11.  

  12. #include <cutil.h>

  13.  

  14. #define N 8

  15. #define BLOCK_X N

  16. #define THREAD_X 1

  17.  

  18. //extern __device__ int* buf;

  19. __device__ int waitflag[8] = {0,0,0,0,0,0,0,0};

  20.  

  21. __global__ void gpu

  22. (int *mem0, int *mem1, unsigned int *_buf)

  23. {

  24.   // -------- -------- -------- --------

  25.   int wait = 0;

  26.   int id = blockIdx.x;

  27.   // 実行時間をずらすための処理

  28.   if(id==3){

  29.     mem1[0]=0;

  30.     mem1[1]=0;

  31.     mem1[2]=0;

  32.   }

  33.   // 各blockで_buf[0]を叩く

  34.   atomicAdd(&_buf[0], id);

  35.  

  36. #if SYNC==3

  37.   // 計算終了通知

  38.   waitflag[id] = 1;

  39.   // 全blockが計算終了に到達していることを確認する

  40.   while(wait <BLOCK_X){

  41.     wait = 0;

  42.     for(int i=0; i<BLOCK_X; i++){

  43.       wait += waitflag[i];

  44.     }

  45.   }

  46. #elif SYNC==4

  47.   atomicAdd(&waitflag[0], 1);

  48.   while(wait<BLOCK_X){wait=waitflag[0];}

  49. #elif SYNC==2

  50.   atomicAdd(&waitflag[0], 1);

  51.   while(waitflag[0]<BLOCK_X)__syncthreads();

  52. #else // NG

  53.   atomicAdd(&waitflag[0], 1);

  54.   while(waitflag[0]<BLOCK_X);//__syncthreads();

  55. #endif

  56.   // ここから、各blockが計算結果を確実に取得できる

  57.   // global変数に書き戻して確認

  58.   mem0[id] = _buf[0];

  59.   // 待ち合わせデータのチェックのため、

  60.   // global変数に書き戻して確認

  61.   mem1[id] = wait;

  62.  

  63.   return;

  64.  

  65.   // -------- -------- -------- --------

  66. }

  67.  

  68. // -------- -------- -------- -------- -------- -------- -------- --------

  69.  

  70. void

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

  72. {

  73.   CUT_DEVICE_INIT();

  74.  

  75.   int* h_mem0;

  76.   int* h_mem1;

  77.   int* d_mem0;

  78.   int* d_mem1;

  79.   h_mem0 = (int*)malloc(sizeof(int)*N);

  80.   h_mem1 = (int*)malloc(sizeof(int)*N);

  81.   CUDA_SAFE_CALL(cudaMalloc((void**)&d_mem0, sizeof(int)*N));

  82.   CUDA_SAFE_CALL(cudaMalloc((void**)&d_mem1, sizeof(int)*N));

  83.   //CUDA_SAFE_CALL(cudaMemcpy(d_mem0, h_mem0, sizeof(int)*N, cudaMemcpyHostToDevice) );

  84.   //CUDA_SAFE_CALL(cudaMemcpy(d_mem1, h_mem1, sizeof(int)*N, cudaMemcpyHostToDevice) );

  85.   unsigned int* h_buf;

  86.   unsigned int* d_buf;

  87.   CUDA_SAFE_CALL(cudaMallocHost((void**)&h_buf, sizeof(int)*N));

  88.   int n;

  89.   for(n=0; n<N; n++){

  90.     h_buf[n] = 1;

  91.   }

  92.   CUDA_SAFE_CALL(cudaMalloc((void**)&d_buf, sizeof(int)*N));

  93.   CUDA_SAFE_CALL(cudaMemcpy(d_buf, h_buf, sizeof(int)*N, cudaMemcpyHostToDevice) );

  94.   //CUDA_SAFE_CALL(cudaMemset(buf,0,sizeof(int)*N));

  95.  

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

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

  98.  

  99.   gpu<<<grid, threads>>>(d_mem0, d_mem1, d_buf);

  100.   CUDA_SAFE_CALL( cudaThreadSynchronize() );

  101.   CUT_CHECK_ERROR("Kernel execution failed");

  102.   CUDA_SAFE_CALL(cudaMemcpy(h_mem0, d_mem0, sizeof(int)*N, cudaMemcpyDeviceToHost) );

  103.   CUDA_SAFE_CALL(cudaMemcpy(h_mem1, d_mem1, sizeof(int)*N, cudaMemcpyDeviceToHost) );

  104.   CUDA_SAFE_CALL(cudaMemcpy(h_buf, d_buf, sizeof(int)*N, cudaMemcpyDeviceToHost) );

  105.   CUDA_SAFE_CALL(cudaFree(d_mem0));

  106.   CUDA_SAFE_CALL(cudaFree(d_mem1));

  107.   CUDA_SAFE_CALL(cudaFree(d_buf));

  108.  

  109.   int i;

  110.   printf("mem0:");

  111.   for(i=0; i<N; i++){

  112.     printf(" %d", h_mem0[i]);

  113.   }

  114.   printf("\n");

  115.   printf("mem1:");

  116.   for(i=0; i<N; i++){

  117.     printf(" %d", h_mem1[i]);

  118.   }

  119.   printf("\n");

  120.   printf("buf:");

  121.   for(i=0; i<N; i++){

  122.     printf(" %d", h_buf[i]);

  123.   }

  124.   printf("\n");

  125.  

  126.   CUDA_SAFE_CALL(cudaFreeHost(h_buf));

  127.   free(h_mem0);

  128.   free(h_mem1);

  129. }

  130.  

  131.  

  132. int

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

  134. {

  135.   runTest(argc, argv);

  136.  

  137.   CUT_EXIT(argc, argv);

  138. }


Linuxでlatex/PDF

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

うまくいってなかったのが解決したかもしれないのでメモ。

っつーかptetex -- teTeX 用日本語パッチ集に従っただけ。

makeにこけたけど、flexとyaccとlibXaw-develとncurses-develあたりを追加してインストール完了。

platexとdvipdfmxでPDFの生成までは確認。

（……あれ？dvipdfmxっていつ入れたんだっけ……？）


何か気がついたことがあったらまた書く。