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）
);