2012 年 3 月仲里 (執筆)2012 年 4 月大野 (修正)sciGPGPU サンプルプログラム解説
ここでは sciGPGPU のホームページの Project Home(http://forge.scilab.org/index.php/p/sciCuda/) に載せてあるサンプルコードの解説を行います。下記のプログラムは全て sciGPGPU をインストールすることで scilab 上で実行可能です。

• 1) simple.sce//カーネル関数を用いたプログラム。scilab(CPU) 上で定義された 256 × 256 の行列 A を GPU 上で//確保したメモリ dB に代入し，それを scilab への行列 B へ引き渡すプログラム
abs path=get absolute file path(”simple.sce”);

bin = gpuBuild(abs path+”simple”);
A=rand(256,256);dA=gpuSetData(A);dB=gpuAlloc(256,256);
Kernel= gpuLoadFunction(bin,”someSimpleKernel”);
lst=list(dA,dB,int32(256*256));
gpuApplyFunction(Kernel,lst,128,1,256*256/128,1);
B=gpuGetData(dB);
gpuFree(dA);
gpuFree(dB);
clear
解説：
まず上記のコードを sciNotes 上で simple.sce として保存し、実行します。また、simple.sce のディレクトリ内に simple.cuという CUDA 用のコードを置いておく必要があります。これらはプロジェクトのホームページからダウンロードすることが可能です。

• abs path=get absolute file path(”simple.sce”);
get absolute file path(”.sce ファイル”) で.sce ファイルまでのディレクトリパスを abs path へ返します。しかしこれは現在開いているファイルのみを参照するため.sce ファイル内でしか使えないことに注意が必要です。
• bin = gpuBuild(abs path+”simple”);
gpuBuild(”ファイル名”) で”ファイル名.cu”ファイルをコンパイルします。.cu ファイルの中身はカーネル関数 (複数可) のみを記述します。bin には gpuLoadFunction の第一引数へ用います。
• dA=gpuSetData(A);
gpuSetData(A) で scilab 上で定義された行列 A を GPU 上 (global memory) へロードします。返り値 dA はロードされたGPU 上の行列へのポインタとなります。
• dB=gpuAlloc(256,256);
gpuAlloc(256,256) で GPU(global memory) 上に 256 × 256 分のメモリの確保を行います。返り値 dB はメモリへのポインタとなります。

• Kernel= gpuLoadFunction(bin,”someSimpleKernel”);
gpuLoadFunction(bin,”カーネル関数名”) で GPU へ使用したいカーネル関数をロードします。返り値は gpuApplyFunction の第一引数で用います• lst=list(dA,dB,int32(256*256));
scilab の関数 list を用いて GPU へロードしたカーネル関数に渡す引数を定義しています。dA、dB という順に渡していきます。ポインタではなく数値をカーネルへ直接渡したい場合は int32 型を用いる必要があります。

• gpuApplyFunction(Kernel,lst,128,1,256*256/128,1);
gpuApplyFunction で実際に GPU 上での演算を開始します。Kernel でロードしたカーネルにアクセスし、lst で引数を順次カーネルへ渡します。第 3、4 引数はそれぞれブロック内に配置するスレッドの縦 (height) と横 (width) の数、第 5、6 引数はグリッド内に配置するブロックの縦 (height) と横 (width) の数を指定します。このプログラムではブロック内にスレッドを 128 × 1 とグリッド内にブロックを 256 × 256/128、計 256 × 256 個のスレッドを使用しています。
• B=gpuGetData(dB);gpuGetData(dB) で、GPU での演算結果を格納している dB のデータを scilab 上での関数 B へロードします。
• gpuFree(dA)、gpuFree(dB);
使用し終わった GPU 上のメモリ dA、dB を解放します。
• simple.cu の中身

extern ”C”{
global void someSimpleKernel(double* src,double* dst, int numberOfElement){
int idx=threadIdx.x+blockIdx.x*blockDim.x;
if(idx < numberOfElement){
dst[idx]=src[idx];}}}
解説：
引数は lis で定義した順に参照されるので*src は*dA、*dst は*dB、numberOfElement は 256 × 256 となります。idx で各スレッドがそれぞれのビルトイン変数を獲得し、dA の idx 番目の要素を dB の idx 番目の要素に代入しています。

• 2)//sciGPGPU で実装される関数のみを使用したプログラム。
//scilab 上で定義した行列 A、B、C を GPU 上で乗算している。

stacksize(’max’);
A=rand(1000,1000);
B=rand(1000,1000);
C=rand(1000,1000);
dC=gpuSetData(C);//gpuMult(A,B,[opt1 opt2])//opt1 : 0 = Return the result on host.
//
1 = Return the result on device.
//
2 = Return the result on host and device.
//opt2 : 0 = Free memory allocated on device for two input.
////

1 = Keep the first argument on device.
2 = Keep the second argument on device.

//

12 = keep both on device.

[GPU dA dB]=gpuMult(A,B,[1 12]);
GPU1=gpuMult(dA,dC,[1 0]);
CPU=gpuMult(GPU,GPU1);
dA=gpuFree(dA);
dB=gpuFree(dB);
dC=gpuFree(dC)
GPU=gpuFree(GPU);
GPU1=gpuFree(GPU1);
clear
解説：

• [GPU dA dB]=gpuMult(A,B,[1 12]);scilab 上で定義された行列 A、B の乗算を GPU 上で演算します。演算結果は変数 GPU に格納されますが、opt1 で 1 を指定しているため演算結果は GPU のメモリに格納され、変数 GPU は演算結果を格納している行列へのポインタとなります。
また opt2 で 12 を指定しているため scilab で演算のために用いた行列 A、B は GPU 上にも dA、dB として保存されています。

• GPU1=gpuMult(dA,dC,[1 0]);GPU のメモリに格納されている行列 dA、dC の乗算を GPU 上で演算します。opt2 で 0 を指定しているため演算に用いたdA、dB にメモリは保存されません。
• CPU=gpuMult(GPU,GPU1);
GPU のメモリに格納されている行列 GPU、GPU1 の乗算を GPU 上で演算します。第 3 引数に何も指定していないため([0 0] と同義) 演算結果は scilab 上の変数 CPU に格納され、演算に用いられた GPU、GPU1 のメモリは保存されません。
しかしこれらのポインタは全て前述のプログラムですでに保存されているので解放する必要があります。