JCuda - 初期化およびデバイスメモリの使い方
JCuda を使用するにあたって、ライブラリの初期化およびコードの読み込みが必要です。
また、デバイスメモリは GPU 側のみに存在する特別なメモリであるため、特別な方法で取得・初期化する必要があります。
この記事ではライブラリの初期化、デバイスメモリの確保まで説明します。
なお、JCuda を利用したサンプルプログラムが GitHub にありますので、併せてご参照ください。
JCuda の初期化
JCuda の初期化については２つポイントがあります。１つはライブラリ自身の初期化で、もう１つは CUDA を実行する GPU の選択です。
JCuda ライブラリを使用するには、以下のライブラリをインポートする必要があります。
import static jcuda.driver.JCudaDriver.*;
import jcuda.*;
import jcuda.driver.*;
JCuda ライブラリの初期化コードを以下に示します。
// 例外を有効化する
JCudaDriver.setExceptionEnabled(true);
// CUDA ドライバを初期化
cuInit(0);
// コンテキストを作成
CUcontext pctx = new CUcontext();
// Device(id=0: つまりデフォルト) の stream=0 をコンテキストに結びつける
CUdevice dev = new CUdevice();
cuDeviceGet(dev, 0);
cuCtxCreate(pctx, 0, dev);
// PTX ファイルを読み込む
String ptxfilename = preparePtxFile(CUFILENAME);
CUmodule module = new CUmodule();
cuModuleLoad(module, ptxfilename);
ここで preparePtxFile は .cu ファイルから .ptx ファイルを作成するための便利関数です。
実装については GitHub を参照してください。
ここまで来ると、module を使って .cu ファイル上 __global__ 宣言したエントリを取得することができます。
CUfunction entry = new CUfunction();
cuModuleGetFunction(entry, module, KERNELNAME);
// 以降、entry を使用する。他の変数は解放して構わない。
デバイスメモリの作成、初期化、および取得
デバイスメモリはバイト単位で管理されます。そのため、メモリの型に合わせたバイト数を与えるグローバルスタティックメンバー Sizeof.* が提供されています。
| 型 | 対象クラス | バイト数 |
|---|
| BYTE | byte | 1 |
| SHORT | short | 2 |
| INT | int | 4 |
| FLOAT | float | 4 |
| LONG | long | 8 |
| DOUBLE | double | 8 |
| POINTER | Pointer | 8 |
デバイスメモリを確保するには CUdeviceptr インスタンスを作成し、それを引数として cuMemAlloc() 関数を使用します。例えば float 型のメモリを SIZE 個確保するには
CUdeviceptr devMemFloat = new CUdeviceptr();
cuMemAlloc(devMemFloat, Sizeof.FLOAT * SIZE);
とします。ホストメモリは通常の Java のメモリ確保方法です。ホストメモリとデバイスメモリの間の転送は
// ホストメモリの確保
float[] hostMem = new float[SIZE];
// ホストメモリからデバイスメモリへの転送
cuMemcpyHtoD(devMemFloat, Pointer.to(hostmem), Sizeof.FLOAT * SIZE);
// デバイスメモリからホストメモリへの転送
cuMemcpyDtoH(Pointer.to(hostmem), devMemFloat, Sizeof.FLOAT * SIZE);
となります。
ここで Pointer.to() はホストメモリ配列の先頭アドレスを与える Pointerクラスのスタティックメソッドです。
カーネルの呼び出し
カーネルには以下のパラメータがあります。
- Grid 次元数
- Thread 次元数
- 引数ポインタ
- 共有メモリサイズ
- 使用ストリーム
Grid 次元数
CUDA の block 数の大きさを決定します。
３次元のパラメータで、使用しない次元は 1 を指定します。
Thread 次元数
CUDA の thread 数の大きさを決定します。
３次元のパラメータで、使用しない次元は 1 を指定します。
例えば１次元配列で大きさが SIZE だとした場合、NTHREAD を Thread 次元数として
Grid 次元数は (SIZE + NTHREAD - 1) / NTHREAD となります。
NTHREAD には上限があり、またどのような値がよいかは実験的に決める必要があります。
引数ポインタ
__global__ 関数に与える引数を指定します。JCuda は CUDA driver モードで動作するため、引数並びへのポインタが入力値として必要になります。
例えば、入力データ、入力サイズ、出力データをそれぞれ input, size, output と定義していたとすると
Pointer kp = Pointer.to(Pointer.to(input), Pointer.to(new int[]{size}), Pointer.to(output);
とします。２つの点に注意してください。
- Pointer.to(args, ...) は Pointer インスタンスしか引数に取らないため、デバイスメモリも Pointer.to() で Pointer 型に変換してください
- 非配列変数を直接指定することはできないため、new cls[]{初期値} として値を与えます。
この場合でも .cu プログラムでは配列ではなく、通常の引数が使用可能です。
上記の kp に対応する .cu 側のシグニチャは
extern "C" void kernel(float *input, int size, float *output);
となります。
共有メモリサイズ
shared memory の大きさをバイト単位で与えます。共有メモリを使用しない場合は 0 を設定しておきます。
使用ストリーム
カーネルを実行する CUDA stream を指定します。デフォルトの stream を使用する場合は null を指定します。
カーネルの呼び出し
カーネルの呼び出しは cuLaunchKernel で行います。
cuLaunchKernel(function,
(SIZE + NTHREAD - 1) / NTHREAD, 1, 1, // Grid
NTHREAD, 1, 1, // Block
0, null, // 共有メモリサイズ、使用 stream
kp, null // カーネルパラメータ、拡張パラメータ（常に null）
);
Copyright (c) 2017-2019 by TeqStock.tokyo