ひびろぐ ver.2

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

Archive for the ‘GPU関連’

Pages: ◄ 1 2 3 4 ... 5 ►

続：CUDAのメモリ操作には何かが足りない

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

そうか、symbolを使えばconstantメモリにデータを格納できるのか。symbol系の転送命令ならちゃんとoffsetあるわ。

CUDAのメモリ操作には何かが足りない

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

オフセット指定が無い。

つまり、GPU上に大きなメモリを確保した場合は、それを送受信する際にCPU側にも同じサイズのメモリを確保する必要がある、と。


……何か間違ってないか？


よくよく考えてみると、OpenGLやDirectXでGPUとデータをやりとりする場合はテクスチャを確保して適当にデータを送受信していた。ここではオフセット調整ができた。CUDAにもテクスチャがあるので関数リファレンスを参照してみると、オフセットの調整が出来そうに見える。なるほど、こっちを使えばいいのか。


ここまで読んだところで、CUDAにはLocalMemory,SharedMemory,GlobalMemory,ConstantMemory,TextureMemoryという多数のメモリが存在しているわけだが、どれが何メガ使えるのかを把握していなかったことに気がついた。うーむ、まだまだだな。

glDeleteFramebuffersEXTで（ｒｙ

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

データサイズ間違えてた。そりゃセグメントエラーもするわな。orz

glDeleteFramebuffersEXTでセグメントエラーする件

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

おかしいなあ。

何かを忘れているんだろうか。

CUDAのatomic命令が使えない？

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

NVIDIA Forums -> CUDA 1.0 FAQ

なんてこった、G80系ミドルレンジ以下はCUDAのマイナーバージョンが上がって機能が増えていたのか。

並列処理屋さんとしてはatomic系の命令が欲しいんだけどねえ。誤魔化す手段はそれなりにあるわけだが。

CUDAの挙動を追いかけ中：共有メモリマルチスレッドプログラミングの実験みたいなもの

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

共有メモリに同時多数アクセスをしたらどうなるのか。

PLAIN TEXT
CUDA:

  1. __global__ void test

  2. (

  3.  int size

  4.  , int* nums

  5.  )

  6. {

  7.   int id = threadIdx.x + threadIdx.y*10;

  8.   __shared__ int data[100];

  9.   data[0] = id;

  10.   nums[id] = data[0];

  11. }


適当なサイズの共有メモリを確保して、いわゆるスレッド番号みたいなものを書き込んで、取得してみる。

同一の共有メモリにみんなで書き込んでいるため、値がどうなるかわかったもんじゃない。

共有メモリマルチスレッドプログラミングのバグの温床。

のはずなんだけど、値が狂わない。おかしいなあ？


CPUプログラミングで言うとことのスレッドであるCUDAのブロックは互いに同期を取れないらしいんだけど、これって一般的な並列プログラミングの足枷になるよなあ……。協調処理がしにくすぎる。アプリケーションとの相性になるよなぁ。

CUDAの並列実行モデル？についてのメモ

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

自分のための備忘録。というか自己解釈。間違ってる部分もあるかも。

GPUへ投げる``ジョブ''の単位は``Grid''。


``Grid''は複数の``Block''で構成される。

1つの``Block''は1つの``Multiprocessor''上で実行される。

（``Block''が更に分割されて``Multiprocessor''に割り当てられることは無い。）

既存のCPUプログラミングにおけるスレッド的な位置づけだが、``Block''間の同期は取れない。

``Multiprocessor''の数が並列実行可能な``Block''の数に対応。

``Block''の数が``Multiprocessor''の数を超えた場合は、

適当な``Multiprocessor''に複数の``Block''が割り当てられ、

タイムスライシングで平行実行される。


``Block''は複数の``Thread''で構成される。

``Thread''は既存のCPUプログラミングで言うところのSIMD的な位置づけ。

感覚的にはむしろ、OpenMPのループ並列に近く感じられる。

``Block''内で同時実行可能な``Thread''の数は8(=Multiprocessorを構成するプロセッサの数)。

warp size=32はどこに関係するんだろう？



うーむ、HWのモデルとメモリのモデルとがうまく脳内にマッピングしきれない。

とりあえず実行時間ベースでの調査に進むか。

CUDAでN-QUEENやってみた：修正版

Published 7月 23rd, 2007 in GPU関連 | コメント（0） »

前回のが思いっきりミスしてたので修正。

CUDAのスレッドってのはいわゆるSIMD。同じ処理を同時にやるだけ。探索で意味があるのは同時に別の処理を行える並列処理。前回のヤツだとうまく並列高速化ができていない。

っつーわけで修正。

GPUカーネルのthreadIdxをblockIdxに変更して、呼び出し側はthreadを1,1,1、gridをN,1,1に変更。

  • 逐次実行N=12　9800msec

  • 偽並列実行N=12　2300ms

  • 真並列実行N=12　870ms

クソ速いな。

ちなみに実験機はQuadroFX4600、CUDAで言うところの「マルチプロセッサ」、要するに演算器のまとまりが12組ある。これは8800GTSと同数で、8800GTX/Ultraは16。そのためか、実装によってはN=13で実行するとうまくいかない。このへんの数の制限みたいなのはまだちょっと脳内で混乱が生じているので、ちゃんと把握しないといかんな。


それにしても、こんなに簡単に10倍以上の性能を拝めるのは、OpenMPプログラムをプロセッサ数16以上というややレアなスーパーコンで動かすときくらいのものだと思うのですよ。素晴らしいなぁ。

CUDAでN-QUEENやってみた

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

主に、CUDAでプログラムを組むとこんな感じだよーってのをチェックするため。こんなもの高速化しても論文になりませ……ならないと思います。多分。


とりあえず問題サイズと同じだけの並列化*もやってちょうど200行くらい。少なくともシェーダを叩くよりはやりやすいと思う。メモリの扱いだけ気をつければ、単純なCで書ける。まあそのメモリの扱いがめんどいんだけどさ。

（* : N=8の時は8並列。要するに、IDがｎの人は１列目のクイーンをｎ行目に置くと仮定してネ、っていう並列化。）


問題サイズを12より大きくするとこけるのはなんだろう。とりあえずでっちあげただけなのでちゃんと調査してません。もしかしたらスレッド（CUDAの実行モデルで言うところの一番小さい単位。スレッドを束にしたブロックとかいう単位もある。まぁいわゆるスレッド並列化のスレッドと考えていいよ。）の何らかの制限を突破してるかもね。メモリ量とか時間とか。シェーダみたいにループの回数制限とかあったっけかなぁ？

ちなみに実行時間を測定したところ、CPUの非並列非再帰版より明らかに遅かった。更に細かく並列化して高速化する余地というかアイディアはありまくりなので、この後もうちょっと頑張って高速化するかもね。しないかもね。


一応フルソース載せておきます。参考になるかは知らない。とっとと帰ってゼミ資料でも捏造して寝ます。CUDAたんと戯れていたら東急ストアの特売日逃したよチクショウ！

PLAIN TEXT
CUDA:

  1. #include <stdlib.h>

  2. #include <stdio.h>

  3. #include <string.h>

  4. #include <math.h>

  5. #include <unistd.h>

  6.  

  7. #include <cutil.h>

  8.  

  9. #define N 12

  10.  

  11. __global__ void test(

  12.   int size,

  13.   int* numAnswers,

  14.   int* E,

  15.   int* SE,

  16.   int* NW,

  17.   int* answer,

  18.   int* etc

  19. )

  20. {

  21.   int id = threadIdx.x;

  22.   if(id>size){

  23.     if(id<size*2){

  24.       etc[id] = -1;

  25.     }

  26.     return;

  27.   }

  28.   etc[id] = 1;

  29.   if(threadIdx.y>0){

  30.     return;

  31.   }

  32.   int offset1 = id*size;

  33.   int offset2 = id*(size*2-1);

  34.   int x, y;

  35.   x=0;

  36.   y=id;

  37.   answer[offset1+x] = y;

  38.   E[offset1+y] = 1;

  39.   SE[offset2+(size-1)-x+y] = 1;

  40.   NW[offset2+x+y] = 1;

  41.   x++;

  42.   y=0;

  43.  

  44.   while(1){

  45.     if(x==size){

  46.       // all OK

  47.       //printAnswer(d);

  48.       numAnswers[id]++;

  49.       x--;

  50.       y = answer[offset1+x];

  51.       //removeQueen(d,x,y);

  52.       E[offset1+y] = 0;

  53.       SE[offset2+(size-1)-x+y] = 0;

  54.       NW[offset2+x+y] = 0;

  55.       y++;

  56.     }else if(y==size){

  57.       x--;

  58.       if(x<1){

  59.     break;

  60.       }

  61.       y = answer[offset1+x];

  62.       //removeQueen(d,x,y);

  63.       E[offset1+y] = 0;

  64.       SE[offset2+(size-1)-x+y] = 0;

  65.       NW[offset2+x+y] = 0;

  66.       y++;

  67.     }else{

  68.  

  69.       int test=0;

  70.       if(E[offset1+y] || SE[offset2+(size-1)-x+y] || NW[offset2+x+y]){

  71.     test=0;

  72.       }else{

  73.     answer[offset1+x] = y;

  74.     E[offset1+y] = 1;

  75.     SE[offset2+(size-1)-x+y] = 1;

  76.     NW[offset2+x+y] = 1;

  77.     test=1;

  78.       }

  79.  

  80.       if(test==1){

  81.     x++;

  82.     y=0;

  83.       }else{

  84.     y++;

  85.       }

  86.     }

  87.   }

  88.   //etc[id] = offset;

  89. }

  90.  

  91. void

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

  93. {

  94.   CUT_DEVICE_INIT();

  95.  

  96.   printf("%d Byte\n", (1+N+N*N+(N*(N*2-1))+(N*(N*2-1))+N*N)*4);

  97.  

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

  99.   int size;

  100.   int numAnswers[N];

  101.   int E[N*N];

  102.   int SE[N*(N*2-1)];

  103.   int NW[N*(N*2-1)];

  104.   int answer[N*N];

  105.   int etc[N*2];

  106.  

  107.   int i;

  108.   size = N;

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

  110.     numAnswers[i] = 0;

  111.   }

  112.   for(i=0; i<N*2; i++){

  113.     etc[i] = 0;

  114.   }

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

  116.     E[i] = 0;

  117.     answer[i] = 0;

  118.   }

  119.   for(i=0; i<N*(N*2-1); i++){

  120.     SE[i] = 0;

  121.     NW[i] = 0;

  122.   }

  123.  

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

  125.   int* d_numAnswers;

  126.   int* d_E;

  127.   int* d_SE;

  128.   int* d_NW;

  129.   int* d_answer;

  130.   int* d_etc;

  131.  

  132.   CUDA_SAFE_CALL(cudaMalloc((void**)&d_numAnswers, sizeof(int)*N));

  133.   CUDA_SAFE_CALL(cudaMalloc((void**)&d_E, sizeof(int)*N*N));

  134.   CUDA_SAFE_CALL(cudaMalloc((void**)&d_SE, sizeof(int)*(N*2-1)*N));

  135.   CUDA_SAFE_CALL(cudaMalloc((void**)&d_NW, sizeof(int)*(N*2-1)*N));

  136.   CUDA_SAFE_CALL(cudaMalloc((void**)&d_answer, sizeof(int)*N*N));

  137.   CUDA_SAFE_CALL(cudaMalloc((void**)&d_etc, sizeof(int)*N*2));

  138.  

  139.   CUDA_SAFE_CALL(cudaMemcpy(d_numAnswers, &numAnswers, sizeof(int)*N, cudaMemcpyHostToDevice) );

  140.   CUDA_SAFE_CALL(cudaMemcpy(d_E, E, sizeof(int)*N*N, cudaMemcpyHostToDevice) );

  141.   CUDA_SAFE_CALL(cudaMemcpy(d_SE, SE, sizeof(int)*(N*2-1)*N, cudaMemcpyHostToDevice) );

  142.   CUDA_SAFE_CALL(cudaMemcpy(d_NW, NW, sizeof(int)*(N*2-1)*N, cudaMemcpyHostToDevice) );

  143.   CUDA_SAFE_CALL(cudaMemcpy(d_answer, &answer, sizeof(int)*N*N, cudaMemcpyHostToDevice) );

  144.   CUDA_SAFE_CALL(cudaMemcpy(d_etc, &etc, sizeof(int)*N*2, cudaMemcpyHostToDevice) );

  145.  

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

  147.   dim3 threads(N, 1, 1);

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

  149.  

  150.   // 実行

  151.   unsigned int hTimer;

  152.   CUT_SAFE_CALL( cutCreateTimer(&hTimer) );

  153.   CUT_SAFE_CALL( cutResetTimer(hTimer) );

  154.   CUT_SAFE_CALL( cutStartTimer(hTimer) );

  155.   test<<<grid, threads>>>(size, d_numAnswers, d_E, d_SE, d_NW, d_answer, d_etc);

  156.   CUDA_SAFE_CALL( cudaThreadSynchronize() );

  157.   CUT_SAFE_CALL( cutStopTimer(hTimer) );

  158.   double gpuTime = cutGetTimerValue(hTimer);

  159.   printf("Time: %f ms\n", gpuTime);

  160.  

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

  162.   CUT_CHECK_ERROR("Kernel execution failed");

  163.  

  164.   // 演算結果の取得

  165.   CUDA_SAFE_CALL(cudaMemcpy(&numAnswers, d_numAnswers, sizeof(int)*N, cudaMemcpyDeviceToHost) );

  166.   CUDA_SAFE_CALL(cudaMemcpy(&etc, d_etc, sizeof(int)*N, cudaMemcpyDeviceToHost) );

  167.  

  168.   // 演算結果の確認

  169.   int nAll=0;

  170.   for(i=0; i<N*2; i++){

  171.     printf(" %d", etc[i]);

  172.   }

  173.   printf("\n");

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

  175.     nAll += numAnswers[i];

  176.     printf(" %d\n", numAnswers[i]);

  177.   }

  178.   printf("sum %d\n", nAll);

  179.  

  180.   // クリーンアップ

  181.   //free();

  182.   CUDA_SAFE_CALL(cudaFree(d_numAnswers));

  183.   CUDA_SAFE_CALL(cudaFree(d_E));

  184.   CUDA_SAFE_CALL(cudaFree(d_SE));

  185.   CUDA_SAFE_CALL(cudaFree(d_NW));

  186.   CUDA_SAFE_CALL(cudaFree(d_answer));

  187.   CUDA_SAFE_CALL(cudaFree(d_etc));

  188. }

  189.  

  190.  

  191. int

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

  193. {

  194.   runTest(argc, argv);

  195.  

  196.   CUT_EXIT(argc, argv);

  197. }


HelloWorld@CUDA　Runtime API版

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

というわけで再挑戦。

nvidiaのサンプルに倣って、gpuの処理を書いたgpukernel.cuをcuda.cuからincludeしてみた。

まずはCPUの処理を記述しているcuda.cu。

PLAIN TEXT
CU:

  1. #include <stdlib.h>

  2. #include <stdio.h>

  3. #include <string.h>

  4. #include <math.h>

  5.  

  6. #include <cutil.h>

  7.  

  8. // サンプルの真似をしてGPUの計算を一応別ファイル化

  9. #include "gpukernel.cu"

  10.  

  11. void

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

  13. {

  14.   CUT_DEVICE_INIT();

  15.  

  16.   int nSize = 10;

  17.   int n;

  18.   n = sizeof(float)*nSize;

  19.  

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

  21.   float *h_sd1;

  22.   h_sd1 = (float*)malloc(n);

  23.   int i;

  24.   for(i=0; i<nSize; i++){

  25.     h_sd1[i] = (float)(rand()%100)/10.0f;

  26.   }

  27.   // 値の確認

  28.   for(i=0; i<nSize; i++){

  29.     printf(" %f", h_sd1[i]);

  30.   }

  31.   printf("\n");

  32.  

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

  34.   float* d_sd1;

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

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

  37.  

  38.   // 演算結果取得のためのメモリを準備

  39.   float *h_sd2;

  40.   h_sd2 = (float*)malloc(n);

  41.   float* d_sd2;

  42.   CUDA_SAFE_CALL(cudaMalloc((void**)&d_sd2, n));

  43.  

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

  45.   dim3 threads(nSize, 1, 1);

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

  47.  

  48.   // 実行

  49.   test<<<grid, threads>>>(d_sd2, d_sd1, nSize, 2.0f);

  50.  

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

  52.   CUT_CHECK_ERROR("Kernel execution failed");

  53.  

  54.   // 演算結果の取得

  55.   CUDA_SAFE_CALL(cudaMemcpy(h_sd2, d_sd2, n, cudaMemcpyDeviceToHost) );

  56.  

  57.   // 演算結果の確認

  58.   for(i=0; i<nSize; i++){

  59.     printf(" %f", h_sd2[i]);

  60.   }

  61.   printf("\n");

  62.  

  63.   // クリーンアップ

  64.   free(h_sd1);

  65.   free(h_sd2);

  66.   CUDA_SAFE_CALL(cudaFree(d_sd1));

  67.   CUDA_SAFE_CALL(cudaFree(d_sd2));

  68. }

  69.  

  70.  

  71. int

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

  73. {

  74.   runTest(argc, argv);

  75.  

  76.   CUT_EXIT(argc, argv);

  77. }


続いてGPUの処理を書いたgpukernel.cu

PLAIN TEXT
CU:

  1. // CPUから呼び出す関数

  2. // 配列に定数を掛けているだけの簡単なものです

  3. __global__ void

  4. test(float* fOut, float* fIn, int nSize, float f)

  5. {

  6.     int i;

  7.     for(i=0; i<nSize; i++){

  8.         fOut[i] = fIn[i] * f;

  9.     }

  10. }


ああ、これは楽チンだ。めんどくさい処理が全部消えてくれました。

コンパイルもnvccにcuda.cuを喰わせれば終了。例によってlibcutil.aとlibGLUT.aのリンクが必要だけど。


gridとthreadsという例のアレがまだ把握できていないのは問題として、とりあえずこれならサクサク組めそうです。

さて、もうちょっと色々といじってみるか。

Pages: ◄ 1 2 3 4 ... 5 ►