CUDAで超単純なサンプルを書いたので、実行方法などと共に晒しておきます。
題してHelloWorld@CUDA。HelloWorldと表示するのが根本的に無理なので、簡単な配列計算だけどね。
ちなみに、とりあえず動くのは確認できた程度のレベルなので、解釈が色々と間違っているかもしれません。詳しい人の突っ込み希望（笑えない
実行環境はクアッドコアXeonのQuadroFX4600。要するに先日エログロ画像を晒した実験機。
環境構築については、nvidiaのサイトからCUDA SDK version 1.0 for LinuxをDL、マニュアルどおりインストール。
runファイルをshで走らせるだけですよ。
続いてプログラムの作成。
CPU(host)側とGPU(device)側で別にしておきます。別にすることが＊必要＊なのかは不明。
とりあえず、事前にnvidiaのサイトから適当なサンプルを入手し、cutilなどの便利アイテムを入手しておきましょう。
それではプログラムの晒し上げ。まずはGPU側から。
- // CPUから呼び出す関数
- // 配列に定数を掛けているだけの簡単なものです
- __global__ void test(float* fOut, float* fIn, int nSize, float f)
- {
- int i;
- for(i=0; i<nSize; i++){
- fOut[i] = fIn[i] * f;
- }
- }
こんな感じ。cuda.cuとして保存。
内容の説明は要らないよね。誰がどう見ても単純な配列計算。
見慣れないものとして、__global__という関数の装飾が付いています。
CUDAではメモリや関数に対するいくつかの装飾が利用できます。というか利用しなくてはなりません。
__global__のついた関数は、hostから呼び出されてdevice上で実行される関数を意味します。
再帰実行できません。staticな変数を内部に持てません。可変引数を取れません。void型専用です。
256byteを超える引数を受け取れません。
nvccコンパイラで--cudaオプションを利用してコンパイルすると、.cubinというモジュールファイルが生成されます。
実行時はこのcubinファイルをリンクして使います。
続いてCPU側。
- #include <stdlib.h>
- #include <stdio.h>
- #include <string.h>
- #include <math.h>
- #include <cuda.h>
- #include <cutil.h>
- void runTest(int argc, char** argv);
- static CUresult initCUDA(char*, CUfunction *p);
- CUdevice cuDevice;
- CUcontext cuContext;
- CUmodule cuModule;
- // initCUDAはCUDAを利用するための準備処理
- // cubinファイルとの関連付けを行っている
- // ほぼ全てサンプルからのコピペです
- static CUresult
- initCUDA(char* executablePath, CUfunction *p)
- {
- CUfunction cuFunction = 0;
- char* module_path;
- CUT_DEVICE_INIT_DRV(cuDevice);
- CUresult status = cuCtxCreate( &cuContext, 0, cuDevice );
- if ( CUDA_SUCCESS != status ){
- printf("cuCtxCreate failed.\n");
- goto Error;
- }
- // 必要に応じて、利用するcubinファイル名に書き換えること
- module_path = cutFindFilePath("cuda.cubin", executablePath);
- if (module_path == 0) {
- status = CUDA_ERROR_NOT_FOUND;
- printf("cutFindFilePath failed.\n");
- goto Error;
- }
- status = cuModuleLoad(&cuModule, module_path);
- cutFree(module_path);
- if ( CUDA_SUCCESS != status ) {
- printf("cuModuleLoad failed.\n");
- goto Error;
- }
- // 必要に応じて、呼び出す関数の名前に書き換えること
- status = cuModuleGetFunction( &cuFunction, cuModule, "test" );
- if ( CUDA_SUCCESS != status ){
- printf("cuModuleGetFunction failed.\n");
- goto Error;
- }
- *p = cuFunction;
- // 複数の関数を利用したい場合は
- // cumodulleGetFunctionと*p=cuFunction=pの組み合わせが複数必要
- // 実験して確かめてないけど
- return CUDA_SUCCESS;
- Error:
- cuCtxDetach(cuContext);
- return status;
- }
- // 実際にGPUを叩く部分
- // サンプルを参考に作成
- void
- runTest(int argc, char** argv)
- {
- int nSize = 10;
- CUfunction cumain = NULL;
- CU_SAFE_CALL(initCUDA(argv[0], &cumain ));
- // _DEBUGをつけてmakeしている場合のみ、initCUDAでエラーが起きているとここで終了する
- // _DEBGUをつけていない場合はスルーなので注意
- // CPU側のデータを準備
- float *h_sd1;
- h_sd1 = (float*)malloc(sizeof(float)*nSize);
- int i;
- for(i=0; i<nSize; i++){
- h_sd1[i] = (float)(rand()%100)/10.0f;
- }
- // 値の確認
- for(i=0; i<nSize; i++){
- printf(" %f", h_sd1[i]);
- }
- printf("\n");
- // GPU側のデータを準備
- CUdeviceptr d_sd1;
- CU_SAFE_CALL(cuMemAlloc( &d_sd1, sizeof(float)*nSize ));
- // CPU側のデータをGPUに転送
- CU_SAFE_CALL(cuMemcpyHtoD( d_sd1, h_sd1, sizeof(float)*nSize ));
- // CPUとGPUそれぞれにもう１セットのデータを用意
- // 計算結果の受け取りに利用する
- float *h_sd2;
- h_sd2 = (float*)malloc(sizeof(float)*nSize);
- CUdeviceptr d_sd2;
- CU_SAFE_CALL(cuMemAlloc( &d_sd2, sizeof(float)*nSize ));
- // GPU側の内部実行単位ごとに使うメモリのサイズを指定している、はず
- CU_SAFE_CALL(cuFuncSetBlockShape( cumain, nSize, 1, 1 ));
- CU_SAFE_CALL(cuFuncSetSharedSize( cumain, sizeof(float)*nSize ) );
- // GPUの関数の引数に対して設定を行っていると思われる
- // 要するに、test関数の引数の順番にあわせて変数の割り当てをしている
- CU_SAFE_CALL(cuParamSeti( cumain, 0, d_sd2 )); // 引数の0byteめはd_sd2
- CU_SAFE_CALL(cuParamSeti( cumain, 4, d_sd1 )); // 引数の4byteめはd_sd1
- CU_SAFE_CALL(cuParamSeti( cumain, 8, nSize )); // 引数の8byteめはnSize
- CU_SAFE_CALL(cuParamSetf( cumain, 12, 2.0 )); // 引数の12byteめは2.0(float)
- CU_SAFE_CALL(cuParamSetSize( cumain, 16 )); // 引数全体のサイズ
- // ここでGPUの処理を実行
- // パラメタによってGPU内部で並列実行されると思われるが……
- CU_SAFE_CALL(cuLaunchGrid( cumain, 1,1));
- // 実行結果の取得
- CU_SAFE_CALL(cuMemcpyDtoH((void *)h_sd2, d_sd2, sizeof(float)*nSize ));
- // 実行結果のチェック
- for(i=0; i<nSize; i++){
- printf(" %f", h_sd2[i]);
- }
- printf("\n");
- // 使い終わったメモリの破棄
- free(h_sd1);
- free(h_sd2);
- CU_SAFE_CALL(cuMemFree(d_sd1));
- CU_SAFE_CALL(cuMemFree(d_sd2));
- CU_SAFE_CALL_NO_SYNC(cuCtxDetach(cuContext));
- }
- int
- main(int argc, char** argv)
- {
- runTest(argc, argv);
- CUT_EXIT(argc, argv);
- }
想像より長い。
ある程度コメントを入れておきましたが、基本的にはサンプルをぱくりました。
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もあるので、必要に応じて利用しましょう。
作成した実行ファイルを走らせてみます。
- $ ./test
- 8.300000 8.600000 7.700000 1.500000 9.300000 3.500000 8.600000 9.200000 4.900000 2.100000
- 16.600000 17.200001 15.400000 3.000000 18.600000 7.000000 17.200001 18.400000 9.800000 4.200000
- 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のサンプル