ひびろぐ ver.2

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

Archive for the ‘GPU関連’

Pages: 1 2 3 ... 5 ►

割と久々に学会論文執筆中

Published 10月 19th, 2007 in 学会関連, GPU関連 | コメント（0） »

〆切は来週です．まだ全然書き終わってません．

いつもどおりの6P論文のつもりで書いていたら，実は参加するセッションが4P程度のShortPaperを要求していることに気がついた件について．


直前じゃなくて良かった．頑張って書き上げるぞ．

……研究室で寝てから．今週末は必死だな．

DavidKirk氏　講演中

Published 9月 11th, 2007 in GPU関連 | コメント（0） »

東大小柴ホールを満席にして（むしろ溢れてる）講演中。


tgbtの人は英語を聞き取る能力がカスですが、前提知識があるのでなんとかって感じ？　っつーか新しいことはあまり言ってないしな。


あどえすのカメラ機能はズームがへたれなので氏の写真を貼れない。デジカメとつなげる環境を構築しておくか……。

__device__な変数に初期値を与える方法を確認してみる

Published 9月 8th, 2007 in GPU関連 | コメント（0） »

PLAIN TEXT
CUDA:

  1. __device__ int flags[2]={0,0};


こういう書き方が出来るのはわかっている。

ここで、配列の長さに#defineした値なんかを使っていると困る。具体的には

PLAIN TEXT
CODE:

  1. __device__ int buf[N];


なんて場合。


というわけで、とりあえず__device__メモリの確保と初期化について少しコードを書いてみた。

PLAIN TEXT
CODE:

  1. // 値の設定方法

  2. // １：cudaMemcpyを利用する

  3. // ２：cudaMemsetを利用する

  4. #define TESTVERSION 1

  5.  

  6. #define N 128

  7.  

  8. // __device__変数のポインタををextern宣言しておく。

  9. // __global__関数経由で実体（Mallocされたメモリ）と関連付ける。

  10. // この方法で__global__関数以外の関数（__device__関数）からも

  11. // bufを利用することが可能になる。

  12. extern __device__ int* buf;

  13.  

  14. __global__ void gpu

  15. (int *mem0, int *_buf)

  16. {

  17.   int i;

  18.   buf = _buf;

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

  20.     mem0[i] = buf[i];

  21.   }

  22. }

  23.  

  24. void

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

  26. {

  27.   CUT_DEVICE_INIT();

  28.  

  29.   // メモリ確保

  30.   int* h_mem0;

  31.   int* d_mem0;

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

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

  34.   // 値を入れる必要性が無いのでコメントアウト

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

  36.  

  37.   // __device__メモリ(buf)の確保と初期化

  38.  

  39.   int* h_buf=NULL;

  40.   int* d_buf=NULL;

  41.  

  42.   // その１：cudaMemcpyを利用する

  43. #if TESTVERSION == 1

  44.   CUDA_SAFE_CALL(cudaMallocHost((void**)&h_buf, sizeof(int)*N)); // 普通にmallocしてもOK

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

  46.   // 値のセット

  47.   int n;

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

  49.     h_buf[n] = n;

  50.   }

  51.   // データのセット

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

  53. #endif

  54.  

  55.   // その２：cudaMemsetを利用する

  56. #if TESTVERSION == 2

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

  58.   CUDA_SAFE_CALL(cudaMemset(d_buf,0,sizeof(int)*N));

  59. #endif

  60.  

  61.   dim3 grid(1,1,1);

  62.   dim3 threads(1,1,1);

  63.  

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

  65.   CUDA_SAFE_CALL( cudaThreadSynchronize() );

  66.   CUT_CHECK_ERROR("Kernel execution failed");

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

  68.   CUDA_SAFE_CALL(cudaFree(d_mem0));

  69.   CUDA_SAFE_CALL(cudaFree(d_buf));

  70.   if(h_buf){CUDA_SAFE_CALL(cudaFreeHost(h_buf));}

  71.  

  72.   int i;

  73.   printf("mem0:");

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

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

  76.   }

  77.   printf("\n");

  78.  

  79.   free(h_mem0);

  80. }


まぁこんな感じ。

といいたいところだが、実は後者（#define TESTVERSION 2　の場合）に、cudaMemsetの第二引数を0じゃなくすると出力結果が無茶苦茶になってしまう。イージーミスか、バグか、仕様か、はてさて……？？？

行き当たりばったりのCUDAプログラミングその３＋：block間での同期を取る２

Published 9月 5th, 2007 in GPU関連 | コメント（2） »

今度はバリア同期っぽく。

global memory spaceにバリア同期ポイント到着フラグ（配列）を用意しておいて、全てのフラグが埋まるのを待つという単純な実装。

PLAIN TEXT
CUDA:

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

  2. /*

  3.   block制御のテスト

  4.   バリア同期のようなものを作成し、待ち合わせを行う

  5. */

  6.  

  7. #include <stdlib.h>

  8. #include <stdio.h>

  9. #include <string.h>

  10. #include <math.h>

  11. #include <unistd.h>

  12.  

  13. #include <cutil.h>

  14.  

  15. #define  N      4

  16.  

  17. #define USE_BARRIER

  18.  

  19. __device__ int flags[N]={0,0,0,0}; // 待ち合わせのためのフラグ

  20. __device__ int tmp[N]={0,0,0,0}; // 計算途中のデータ

  21.  

  22. __global__ void gpu

  23. (int* answer)

  24. {

  25.   int id = blockIdx.x;

  26.   int id2;

  27.  

  28.   id2 = 4 - (id+1);

  29.   tmp[id] = id+1;

  30.  

  31.   // tmp[] = {1,2,3,4}

  32.  

  33.   // blockによって実行時間にばらつきを出すためのダミー演算

  34.   // id=0およびid=1のtmpは二倍になります

  35.   if(id<2){

  36.     int n = 1;

  37.     n *= 2;

  38.     n *= (id+1);

  39.     if(id==1){

  40.       n /= 2;

  41.     }

  42.     tmp[id] *= n;

  43.   }

  44.   // tmp[] = {2,4,3,4}

  45.  

  46.   // 擬似バリア同期

  47.   // 全てのフラグが埋まるのを待つことで、バリア同期のような感じに

  48. #ifdef USE_BARRIER

  49.   int sum = 0;

  50.   flags[id] = 1;

  51.   while(sum <4){

  52.     __syncthreads();

  53.     sum = flags[0]+flags[1]+flags[2]+flags[3];

  54.   }

  55. #endif

  56.  

  57.   answer[id] = tmp[id2] * (id+2);

  58.   // answer[] =

  59.   // 同期ミス = 8,9, 8, 5 (16と10の片方のみ狂う可能性もあり)

  60.   // 同期成功 = 8,9,16,10

  61. }

  62. void

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

  64. {

  65.   CUT_DEVICE_INIT();

  66.   int answer[N];

  67.   int i;

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

  69.     answer[i] = 0;

  70.   }

  71.   int* d_a;

  72.   CUDA_SAFE_CALL(cudaMalloc((void**)&d_a, sizeof(int)*N));

  73.   CUDA_SAFE_CALL(cudaMemcpy(d_a, &answer, sizeof(int)*N, cudaMemcpyHostToDevice) );

  74.   dim3 grid(N,1,1);

  75.   dim3 threads(1,1,1);

  76.   gpu<<<grid, threads>>>(d_a);

  77.   CUDA_SAFE_CALL( cudaThreadSynchronize() );

  78.   CUT_CHECK_ERROR("Kernel execution failed");

  79.   CUDA_SAFE_CALL(cudaMemcpy(&answer, d_a, sizeof(int)*N, cudaMemcpyDeviceToHost) );

  80.   CUDA_SAFE_CALL(cudaFree(d_a));

  81.  

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

  83.     printf(" %d\n", answer[i]);

  84.   }

  85.   printf("\n");

  86. }

  87.  

  88. int

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

  90. {

  91.   runTest(argc, argv);

  92.  

  93.   CUT_EXIT(argc, argv);

  94. }


それっぽく動いた。

へー。

行き当たりばったりのCUDAプログラミングその３：block間での同期を取る

Published 9月 5th, 2007 in GPU関連 | コメント（0） »

ＣＵＤＡプログラミングガイドより意訳：block間での同期は取れないよ

流石に困るので、なんとかして同期っぽいことができないか試してみた。

考え方は以下：

global memory space内の変数をフラグとして利用する。以上。


適当に実装してみたが、とりあえずそれっぽく動いてくれた。

まぁ同期とはちょっと違うんだけど。

PLAIN TEXT
CUDA:

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

  2. /*

  3.   block制御のテスト

  4.   同期のようなものを作成して逐次処理を行う

  5. */

  6.  

  7. #include <stdlib.h>

  8. #include <stdio.h>

  9. #include <string.h>

  10. #include <math.h>

  11. #include <unistd.h>

  12.  

  13. #include <cutil.h>

  14.  

  15. #define  N      4

  16.  

  17. __device__ int  counter=0; // 逐次処理のためのカウンタ

  18.  

  19. __global__ void gpu

  20. (int* answer)

  21. {

  22.   int id = blockIdx.x;

  23.  

  24.   // カウンタを監視して自分の順番を待つ

  25.   while(counter != id){

  26.     __syncthreads();

  27.   }

  28.   *answer = (*answer*id) + (id+1);

  29.   counter += 1;

  30. }

  31.  

  32. void

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

  34. {

  35.   CUT_DEVICE_INIT();

  36.   int answer=0;

  37.   int* d_a;

  38.   CUDA_SAFE_CALL(cudaMalloc((void**)&d_a, sizeof(int)));

  39.   CUDA_SAFE_CALL(cudaMemcpy(d_a, &answer, sizeof(int), cudaMemcpyHostToDevice) );

  40.   dim3 grid(N,1,1);

  41.   dim3 threads(1,1,1);

  42.   gpu<<<grid, threads>>>(d_a);

  43.   CUDA_SAFE_CALL( cudaThreadSynchronize() );

  44.   CUT_CHECK_ERROR("Kernel execution failed");

  45.   CUDA_SAFE_CALL(cudaMemcpy(&answer, d_a, sizeof(int), cudaMemcpyDeviceToHost) );

  46.   CUDA_SAFE_CALL(cudaFree(d_a));

  47.  

  48.   printf("answer %d\n", answer);

  49. }

  50.  

  51. int

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

  53. {

  54.   runTest(argc, argv);

  55.  

  56.   CUT_EXIT(argc, argv);

  57. }


カウンタを監視して自分の順番を待つ、ってところがあるおかげで正答である31が得られる。これがないと変な値が得られてしまう。

次はバリア同期を実装だな。

行き当たりばったりのCUDAプログラミングその２＋＋

Published 9月 5th, 2007 in GPU関連 | コメント（0） »

ifでもswitchでも同じ感じ。

PLAIN TEXT
CUDA:

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

  2. /*

  3.   block制御のテスト

  4.   if実験

  5. */

  6.  

  7. #include <stdlib.h>

  8. #include <stdio.h>

  9. #include <string.h>

  10. #include <math.h>

  11. #include <unistd.h>

  12.  

  13. #include <cutil.h>

  14.  

  15. #define  N      4

  16.  

  17. int  a[N];

  18. int  b[N];

  19.  

  20. #define PATTERN 0

  21. /*

  22.   0 : id          => OK

  23.   1 : 即値        => NG

  24.   2 : table[id]   => OK

  25.   3 : table[即値] => NG

  26. */

  27.  

  28. __constant__ int table[4] = {0,1,2,3};

  29. __global__ void gpu

  30. (int* d_a, int* d_b)

  31. {

  32.   int id = blockIdx.x;

  33. #if PATTERN == 0

  34.   if(id==0){

  35.     d_b[id] = 1;

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

  37.     d_b[id] = 2;

  38.   }else if(id==2){

  39.     d_b[id] = 3;

  40.   }else if(id==3){

  41.     d_b[id] = 4;

  42.   }

  43. #endif

  44. #if PATTERN == 1

  45.   if(id==0){

  46.     d_b[0] = 1;

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

  48.     d_b[1] = 2;

  49.   }else if(id==2){

  50.     d_b[2] = 3;

  51.   }else if(id==3){

  52.     d_b[3] = 4;

  53.   }

  54. #endif

  55. #if PATTERN == 2

  56.   if(id==0){

  57.     d_b[table[id]] = 1;

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

  59.     d_b[table[id]] = 2;

  60.   }else if(id==2){

  61.     d_b[table[id]] = 3;

  62.   }else if(id==3){

  63.     d_b[table[id]] = 4;

  64.   }

  65. #endif

  66. #if PATTERN == 3

  67.   if(id==0){

  68.     d_b[table[0]] = 1;

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

  70.     d_b[table[1]] = 2;

  71.   }else if(id==2){

  72.     d_b[table[2]] = 3;

  73.   }else if(id==3){

  74.     d_b[table[3]] = 4;

  75.   }

  76. #endif

  77. }

  78.  

  79. void

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

  81. {

  82.   int  i, n;

  83.   n = N;

  84.   for( i = 0 ; i <n ; i++ ){

  85.     a[i] = 0;

  86.     b[i] = 0;

  87.   }

  88.  

  89.   CUT_DEVICE_INIT();

  90.   int* d_a;

  91.   CUDA_SAFE_CALL(cudaMalloc((void**)&d_a, sizeof(int)*n));

  92.   CUDA_SAFE_CALL(cudaMemcpy(d_a, a, sizeof(int)*n, cudaMemcpyHostToDevice) );

  93.   int* d_b;

  94.   CUDA_SAFE_CALL(cudaMalloc((void**)&d_b, sizeof(int)*n));

  95.   CUDA_SAFE_CALL(cudaMemcpy(d_b, b, sizeof(int)*n, cudaMemcpyHostToDevice) );

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

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

  98.   gpu<<<grid, threads>>>(d_a, d_b);

  99.   CUDA_SAFE_CALL( cudaThreadSynchronize() );

  100.   CUT_CHECK_ERROR("Kernel execution failed");

  101.   CUDA_SAFE_CALL(cudaMemcpy(b, d_b, sizeof(int)*n, cudaMemcpyDeviceToHost) );

  102.   CUDA_SAFE_CALL(cudaFree(d_a));

  103.   CUDA_SAFE_CALL(cudaFree(d_b));

  104.  

  105.   for( i = 0 ; i <n ; i++ ){

  106.     printf(" %d", b[i]);

  107.   }

  108.   printf("\n");

  109. }

  110.  

  111. int

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

  113. {

  114.   runTest(argc, argv);

  115.  

  116.   CUT_EXIT(argc, argv);

  117. }


PLAIN TEXT
CUDA:

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

  2. /*

  3.   block制御のテスト

  4.   switch実験

  5. */

  6.  

  7. #include <stdlib.h>

  8. #include <stdio.h>

  9. #include <string.h>

  10. #include <math.h>

  11. #include <unistd.h>

  12.  

  13. #include <cutil.h>

  14.