ひびろぐ ver.2

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

Archive for the ‘GPU関連’

Pages: ◄ 1 2 3 4 5 ►

HelloWorld@CUDA　……あっ！

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

プログラミングガイドを読み直していたら気がついた。

前エントリで書いたプログラムはDriver APIってヤツを使って書きました。

CUDAにはDriver API以外にもRuntime APIというものがあります。

これらのAPIは排他的に利用します。

Driver APIはlow-levelなAPIです。Runtime APIはhigher-levelのAPIで、Driver APIの上に構築されています。

しまった、とりあえず動く簡単なサンプルを書くならRuntime APIの方が適していそうじゃないか。


というわけで、次はRuntime APIで書くとです……。



究極的には性能を求めてDriver APIかもしれないけど、手軽にGPUを活用するという意味ではRuntime APIで十分なのかなぁ。まだ良くわかりません。

HelloWorld@CUDA

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

CUDAで超単純なサンプルを書いたので、実行方法などと共に晒しておきます。

題してHelloWorld@CUDA。HelloWorldと表示するのが根本的に無理なので、簡単な配列計算だけどね。

ちなみに、とりあえず動くのは確認できた程度のレベルなので、解釈が色々と間違っているかもしれません。詳しい人の突っ込み希望（笑えない



実行環境はクアッドコアXeonのQuadroFX4600。要するに先日エログロ画像を晒した実験機。

環境構築については、nvidiaのサイトからCUDA SDK version 1.0 for LinuxをDL、マニュアルどおりインストール。

runファイルをshで走らせるだけですよ。

続いてプログラムの作成。

CPU(host)側とGPU(device)側で別にしておきます。別にすることが＊必要＊なのかは不明。


とりあえず、事前にnvidiaのサイトから適当なサンプルを入手し、cutilなどの便利アイテムを入手しておきましょう。


それではプログラムの晒し上げ。まずはGPU側から。

PLAIN TEXT
CUDA:

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

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

  3. __global__ void test(float* fOut, float* fIn, int nSize, float f)

  4. {

  5.     int i;

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

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

  8.     }

  9. }


こんな感じ。cuda.cuとして保存。

内容の説明は要らないよね。誰がどう見ても単純な配列計算。

見慣れないものとして、__global__という関数の装飾が付いています。

CUDAではメモリや関数に対するいくつかの装飾が利用できます。というか利用しなくてはなりません。

__global__のついた関数は、hostから呼び出されてdevice上で実行される関数を意味します。

再帰実行できません。staticな変数を内部に持てません。可変引数を取れません。void型専用です。

256byteを超える引数を受け取れません。

nvccコンパイラで--cudaオプションを利用してコンパイルすると、.cubinというモジュールファイルが生成されます。

実行時はこのcubinファイルをリンクして使います。


続いてCPU側。

PLAIN TEXT
C++:

  1. #include <stdlib.h>

  2. #include <stdio.h>

  3. #include <string.h>

  4. #include <math.h>

  5.  

  6. #include <cuda.h>

  7.  

  8. #include <cutil.h>

  9.  

  10. void runTest(int argc, char** argv);

  11. static CUresult initCUDA(char*, CUfunction *p);

  12.  

  13. CUdevice cuDevice;

  14. CUcontext cuContext;

  15. CUmodule cuModule;

  16.  

  17. // initCUDAはCUDAを利用するための準備処理

  18. // cubinファイルとの関連付けを行っている

  19. // ほぼ全てサンプルからのコピペです

  20. static CUresult

  21. initCUDA(char* executablePath, CUfunction *p)

  22. {

  23.   CUfunction cuFunction = 0;

  24.   char* module_path;

  25.  

  26.   CUT_DEVICE_INIT_DRV(cuDevice);

  27.  

  28.   CUresult status = cuCtxCreate( &cuContext, 0, cuDevice );

  29.   if ( CUDA_SUCCESS != status ){

  30.     printf("cuCtxCreate failed.\n");

  31.     goto Error;

  32.   }

  33.  

  34.   // 必要に応じて、利用するcubinファイル名に書き換えること

  35.   module_path = cutFindFilePath("cuda.cubin", executablePath);

  36.   if (module_path == 0) {

  37.     status = CUDA_ERROR_NOT_FOUND;

  38.     printf("cutFindFilePath failed.\n");

  39.     goto Error;

  40.   }

  41.  

  42.   status = cuModuleLoad(&cuModule, module_path);

  43.   cutFree(module_path);

  44.   if ( CUDA_SUCCESS != status ) {

  45.     printf("cuModuleLoad failed.\n");

  46.     goto Error;

  47.   }

  48.  

  49.   // 必要に応じて、呼び出す関数の名前に書き換えること

  50.   status = cuModuleGetFunction( &cuFunction, cuModule, "test" );

  51.   if ( CUDA_SUCCESS != status ){

  52.     printf("cuModuleGetFunction failed.\n");

  53.     goto Error;

  54.   }

  55.   *p = cuFunction;

  56.  

  57.   // 複数の関数を利用したい場合は

  58.   // cumodulleGetFunctionと*p=cuFunction=pの組み合わせが複数必要

  59.   // 実験して確かめてないけど

  60.  

  61.   return CUDA_SUCCESS;

  62.  

  63.  Error:

  64.   cuCtxDetach(cuContext);

  65.   return status;

  66. }

  67.  

  68. // 実際にGPUを叩く部分

  69. // サンプルを参考に作成

  70. void

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

  72. {

  73.   int nSize = 10;

  74.   CUfunction cumain = NULL;

  75.  

  76.   CU_SAFE_CALL(initCUDA(argv[0], &cumain ));

  77.   // _DEBUGをつけてmakeしている場合のみ、initCUDAでエラーが起きているとここで終了する

  78.   // _DEBGUをつけていない場合はスルーなので注意

  79.  

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

  81.   float *h_sd1;

  82.   h_sd1 = (float*)malloc(sizeof(float)*nSize);

  83.   int i;

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

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

  86.   }

  87.   // 値の確認

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

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

  90.   }

  91.   printf("\n");

  92.  

  93.   // GPU側のデータを準備

  94.   CUdeviceptr d_sd1;

  95.   CU_SAFE_CALL(cuMemAlloc( &d_sd1, sizeof(float)*nSize ));

  96.   // CPU側のデータをGPUに転送

  97.   CU_SAFE_CALL(cuMemcpyHtoD( d_sd1, h_sd1, sizeof(float)*nSize ));

  98.  

  99.   // CPUとGPUそれぞれにもう１セットのデータを用意

  100.   // 計算結果の受け取りに利用する

  101.   float *h_sd2;

  102.   h_sd2 = (float*)malloc(sizeof(float)*nSize);

  103.   CUdeviceptr d_sd2;

  104.   CU_SAFE_CALL(cuMemAlloc( &d_sd2, sizeof(float)*nSize ));

  105.  

  106.   // GPU側の内部実行単位ごとに使うメモリのサイズを指定している、はず

  107.   CU_SAFE_CALL(cuFuncSetBlockShape( cumain, nSize, 1, 1 ));

  108.   CU_SAFE_CALL(cuFuncSetSharedSize( cumain, sizeof(float)*nSize ) );

  109.   // GPUの関数の引数に対して設定を行っていると思われる

  110.   // 要するに、test関数の引数の順番にあわせて変数の割り当てをしている

  111.   CU_SAFE_CALL(cuParamSeti( cumain, 0, d_sd2 ));    // 引数の0byteめはd_sd2

  112.   CU_SAFE_CALL(cuParamSeti( cumain, 4, d_sd1 ));    // 引数の4byteめはd_sd1

  113.   CU_SAFE_CALL(cuParamSeti( cumain, 8, nSize ));    // 引数の8byteめはnSize

  114.   CU_SAFE_CALL(cuParamSetf( cumain, 12, 2.0 ));  // 引数の12byteめは2.0(float)

  115.   CU_SAFE_CALL(cuParamSetSize( cumain, 16 ));      // 引数全体のサイズ

  116.   // ここでGPUの処理を実行

  117.   // パラメタによってGPU内部で並列実行されると思われるが……

  118.   CU_SAFE_CALL(cuLaunchGrid( cumain, 1,1));

  119.  

  120.   // 実行結果の取得

  121.   CU_SAFE_CALL(cuMemcpyDtoH((void *)h_sd2, d_sd2, sizeof(float)*nSize ));

  122.  

  123.   // 実行結果のチェック

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

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

  126.   }

  127.   printf("\n");

  128.  

  129.   // 使い終わったメモリの破棄

  130.   free(h_sd1);

  131.   free(h_sd2);

  132.   CU_SAFE_CALL(cuMemFree(d_sd1));

  133.   CU_SAFE_CALL(cuMemFree(d_sd2));

  134.   CU_SAFE_CALL_NO_SYNC(cuCtxDetach(cuContext));

  135. }

  136.  

  137. int

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

  139. {

  140.   runTest(argc, argv);

  141.  

  142.   CUT_EXIT(argc, argv);

  143. }


想像より長い。

ある程度コメントを入れておきましたが、基本的にはサンプルをぱくりました。

cuda.cppとして保存


説明するより、とにかく実行してみよう。


プログラムのコンパイルにはnvccを利用。

まずは

nvcc --cubin cuda.cu

で.cuファイルを.cubinファイルに変換。

cuFindFilePathがdataディレクトリを参照してcubinファイルを探すようなので、

変換したファイルはdataディレクトリを作成して投げ込みましょう。

つづいて.cppファイルもコンパイル。

基本的には

nvcc cuda.cpp

でOKですが、例によってインクルードなりライブラリなりの設定が必要です。

特にcutilの参照が必要なこと、更にGLEWが必要なことに注意。

サンプルを落としてあれば、

展開して得られるNVIDIA_CUDA_SDKディレクトリの中のcommon/incにcutil関係のヘッダファイルが、libの中にlibcutil.aがあります。

common/libにlibGLEW.aもあるので、必要に応じて利用しましょう。


作成した実行ファイルを走らせてみます。

PLAIN TEXT
CODE:

  1. $ ./test

  2.  8.300000 8.600000 7.700000 1.500000 9.300000 3.500000 8.600000 9.200000 4.900000 2.100000

  3.  16.600000 17.200001 15.400000 3.000000 18.600000 7.000000 17.200001 18.400000 9.800000 4.200000

  4.  

  5. Press ENTER to exit...


１行目が元データ、２行目が計算結果となっています。

ちゃんと２倍になっていますね。


さて、重要なことは以下のいくつかの点。

CUDA実行のための手続き。

initCUDA関数にまとまっています。

コンテキストの作成(cuCtxCreate)、cubinファイルの読み込みおよび実行する関数の確認(cutFindFilePath,cuModulelLoad,cuModuleGetFunction)が必要です。

データの管理。device側のメモリをcuMemAllocで確保したり、host-device間のコピーをcuMemcpy*で行ったりします。

実行。cuFuncSetBlockShapeとcuFuncSetSharedSizeは実行単位で利用するメモリサイズの指定のはずなのですが、実はちゃんと把握できてませんゴメンナサイ。とりあえず動くけど、これはあっていない気がする。

cuParamSeti/cuParamSetfは__global__関数の引数設定だと思います。何バイト目にどの変数を割り当てればよいかを指定するわけです。

cuParamSetSizeが全体のサイズです。

実行はcuLaunchGridで、この関数のパラメタによってdevice内の並列実行が行われると思うのですが、まだそこまでチェックしてません。



まあこんな感じです。

ちょっとソースコードは長いけど、定型処理が多いので難しくはありません。

今回はcuFunc*,cuLaunch*系を利用した並列実行の方法をちゃんと把握できていません。

これがわかれば高速な計算が望めるんですが……。それから、device側での関数呼び出しなどを全然行っていません。

というわけで、現在これらについて実験をしようとしているところ。把握し次第またエントリ化する予定ということで。

ちなみに、cuのソースを

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

から

fIn[i] *= f;

に書き換えて、

cppの方も

CU_SAFE_CALL(cuMemcpyDtoH((void *)h_sd2, d_sd2, sizeof(float)*nSize ));

から

CU_SAFE_CALL(cuMemcpyDtoH((void *)h_sd1, d_sd2, sizeof(float)*nSize ));

に書き換えると、ちゃんと計算結果が取得できます。

シェーダのときみたいに計算結果は別にしろっていう制限があったりはしない模様。

パフォーマンスが変わるか？　とか、そういうのはわかりません。


せっかくだからサンプル一式添付しておきますね→CUDAのサンプル

意外なところにGPUの機能制限

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

と言っていいのかなんなのか。


研究室のメインPCにはX1600がささってるんだけど、FBO+AlphaBlendがうまく実行できない模様。

ふと気がついてQuadroFX4600に持っていったら普通に動く。

この制限はどこで線が引かれるんだろう？

RGBAでF32なフォーマットのとき限定かもしれないけど、その辺の詳細は未チェック。

y=-( ゜д゜)・∵; タ－ン

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

row-majorとcolumn-majorを間違えていた。


cublasってcolumn-majorだったのか……しかもマニュアルの１ページ目に書いてあったとか、これはもう死ぬしかないな。

っつーかmajor-modeを選択できないんだけど、これって逆の問題が出てきたらどうするんだ？AとBだけひっくり返して計算してから、Cを足せとか言われるのか？




というわけで、これから必死にベンチマーク。

ハッハッハァ！

CUDAの計算ライブラリが腐っている気がする件

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

cublasSgemmの行列Aと行列Bって、逆にしないと結果が合わなくないか？


という結論に至ったので、本家フォーラムでゴルアしてきた。

こんなんで性能評価できるんだろうか……。

GotoBLASの使い方がわからないっ

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

CUDAの実験をするついでに真面目に実験環境の性能評価をしてやろうとGotoBLASの導入を


しようとしてるんだけど。

何をincludeして、どういう風に関数呼び出しすれば動くんだ？

というわけで手詰まりに。

先日実験していたと言っていた某東工大のGPU仲間にヘルプメールを投げて死亡。

CUDA実験機をATLASのチューニングで陵辱中

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

CoreQuadの性能が見たいぜわくわく。


それにしても、GotoBLASのquickinstallの手軽さは異常だと思うんだ。

とりあえず性能評価するぞ！

共用マシンでCUDAを使うために

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

$ xhost +

$ chmod o+rwx /dev/nvidiactl

$ chmod o+rwx /dev/nvidia0

まあ、LANの中だし？

そろそろ真面目にCUDAに手を出そうと思ったので

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

最新のFAQに目を通している。

  • 最新のCUDAであればミドルレンジGPUでも使える。8500とか、8400Mとか

  • M/Bに物理的に搭載可能であれば、複数のGPUにも対応

  • GPU同士でのDMAデータ転送は不可能

  • X-serverの設定次第でリモート実行も問題なし

  • GPU計算の起動は非同期なので、直後にCPUで演算をすることが可能

  • 計算結果の取得と計算を並列に行うことは不可能（download dataとしか書かれていないので、uploadは不明）


んー、CPUとGPUの並列処理自体は普通に書けそうである。

とりあえずCUBLASとgotoblasで性能比較とかやっておくかねえ。

JAVAでGPGPUできた

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

JOGLからGLSL叩けたよ！

単純な配列加算しかやってないけどな！！！


JOGLはGL、GLU、GLUTあたりがそのまま叩けて、GLEW関連は最初っからぶっこまれてるっぽい。

割と使えるな。

とりあえずJAVA自体は最低限使えるようになったかもしれん。

不満なのはむしろ、eclipseのショートカット（ホットキー）がVisualStudioと全然違うこととか、MSDNのリファレンスって何気にサンプルとか充実してて強力だったんだよなあとかそういう方向。


今は「GPGPUの触り方」（詳細は当然言えませんよ）にフォーカスして研究を進めようとしているんだけど、JAVAって選択肢は割と悪くないのかもしれない。標準Cより強力なんだろ？くらいにしか把握してないけどさ。

ちなみに、GPUの機能が足りていないときにどうなるのかはまだチェックしていない。ちゃんと検出できるよねきっと。

Pages: ◄ 1 2 3 4 5 ►