ひびろぐ ver.3.1

アーカイブ

2007 年 7 月 19 日のアーカイブ

Twitter Updates for 2007-07-19

2007 年 7 月 19 日コメントはありません

CUDA????????UP??? #

CUDA?nqueen???? #

CUDA???N-QUEEN???? #

カテゴリー: Twitter タグ:

HelloWorld@CUDA　Runtime API版

2007 年 7 月 19 日 tgbt コメントはありません

というわけで再挑戦。

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

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

PLAIN TEXT

CU:

#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <math.h>
#include <cutil.h>
// サンプルの真似をしてGPUの計算を一応別ファイル化
#include "gpukernel.cu"
void
runTest(int argc, char** argv)
{
CUT_DEVICE_INIT();
int nSize = 10;
int n;
n = sizeof(float)*nSize;
// CPU側のデータを準備
float *h_sd1;
h_sd1 = (float*)malloc(n);
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側のメモリを準備
float* d_sd1;
CUDA_SAFE_CALL(cudaMalloc(<a href="#hs_ba1d0d500cc29d7d189a1a4b008b2cdd_footnote_1" id="hs_ba1d0d500cc29d7d189a1a4b008b2cdd_footnotelink_1" title="void**)&d_sd1, n">*1</a>);
CUDA_SAFE_CALL(cudaMemcpy(d_sd1, h_sd1, n, cudaMemcpyHostToDevice) );
// 演算結果取得のためのメモリを準備
float *h_sd2;
h_sd2 = (float*)malloc(n);
float* d_sd2;
CUDA_SAFE_CALL(cudaMalloc(<a href="#hs_ba1d0d500cc29d7d189a1a4b008b2cdd_footnote_2" id="hs_ba1d0d500cc29d7d189a1a4b008b2cdd_footnotelink_2" title="void**)&d_sd2, n">*2</a>);
// 実行のためのパラメタを準備
dim3 threads(nSize, 1, 1);
dim3 grid(1,1,1);
// 実行
test<<<grid, threads>>>(d_sd2, d_sd1, nSize, 2.0f);
// GPUの処理に問題が起きていないかの確認
CUT_CHECK_ERROR("Kernel execution failed");
// 演算結果の取得
CUDA_SAFE_CALL(cudaMemcpy(h_sd2, d_sd2, n, cudaMemcpyDeviceToHost) );
// 演算結果の確認
for(i=0; i<nSize; i++){
printf(" %f", h_sd2[i]);
}
printf("\n");
// クリーンアップ
free(h_sd1);
free(h_sd2);
CUDA_SAFE_CALL(cudaFree(d_sd1));
CUDA_SAFE_CALL(cudaFree(d_sd2));
}
int
main(int argc, char** argv)
{
runTest(argc, argv);
CUT_EXIT(argc, argv);
}

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

PLAIN TEXT

CU:

// 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;
}
}

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

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

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

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

*1: void**)&d_sd1, n

*2: void**)&d_sd2, n

カテゴリー: GPU・GPGPU関連タグ:

HelloWorld@CUDA　……あっ！

2007 年 7 月 19 日 tgbt コメントはありません

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

前エントリで書いたプログラムは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で十分なのかなぁ。まだ良くわかりません。

カテゴリー: GPU・GPGPU関連タグ:

HelloWorld@CUDA

2007 年 7 月 19 日 tgbt コメントはありません

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:

// 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側。

PLAIN TEXT

C++:

#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(<a href="#hs_4dc06060b7c05752dc70bb61bb33f4af_footnote_1" id="hs_4dc06060b7c05752dc70bb61bb33f4af_footnotelink_1" title="void *)h_sd2, d_sd2, sizeof(float)*nSize ">*1</a>);
// 実行結果のチェック
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もあるので、必要に応じて利用しましょう。

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

PLAIN TEXT

CODE:

$ ./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(*2);

から

CU_SAFE_CALL(cuMemcpyDtoH(*3);

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

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

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

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

*1: void *)h_sd2, d_sd2, sizeof(float)*nSize

*2: void *)h_sd2, d_sd2, sizeof(float)*nSize

*3: void *)h_sd1, d_sd2, sizeof(float)*nSize

カテゴリー: GPU・GPGPU関連タグ:

Page 1 of 11

カレンダー

2007年7月
日	月	火	水	木	金	土
« 6月		8月 »
1	2	3	4	5	6	7
8	9	10	11	12	13	14
15	16	17	18	19	20	21
22	23	24	25	26	27	28
29	30	31

-->

アーカイブ

Twitter Updates for 2007-07-19

HelloWorld@CUDA　Runtime API版

HelloWorld@CUDA　……あっ！

HelloWorld@CUDA

What I'm Doing...

使用中の外部サービスなど

はてぶ

カレンダー

最近の投稿

最近のコメント

アーカイブ

カテゴリー

Categories

Blogroll

メタ情報

ひびろぐ ver.3.1

アーカイブ

Twitter Updates for 2007-07-19

HelloWorld@CUDA Runtime API版

HelloWorld@CUDA ……あっ！

HelloWorld@CUDA

What I'm Doing...

使用中の外部サービスなど

はてぶ

カレンダー

最近の投稿

最近のコメント

アーカイブ

カテゴリー

Categories

Blogroll

メタ情報

HelloWorld@CUDA　Runtime API版

HelloWorld@CUDA　……あっ！