C++でOpenCL(高速化編)

目指すところ

OpenCLは楽しいのですが、厳しい高速化を狙っていくと難しい側面も出てきます。そこで、C++でOpenCLシリーズ最終回の今回は、高速化のためにやってみたアルゴリズム工夫のうち、効果のあったものを紹介してみたいと思います。
サンプルコードの中ではデバイス、コンテキスト、コマンドキューなどあまり本質的でない要素は宣言が省略されています。その辺が知りたい方はC++でOpenCL(環境構築編)とか、C++でOpenCL(使ってみよう編)を見てください。

高速化のために

OpenCLなどの並列コンピューティングを用いて高速化するうえで問題となる箇所がいくつかありますが、主なものとしては
・メモリ転送の問題
・メモリアクセスの問題
・カーネル呼び出しの問題
などがあります。これらの問題を解決あるいは改善できる方法について紹介します。

実行時間計測

高速化手法とは違いますが、実行時間を詳しく計測することはボトルネック調査において大事なことです。OpenCLでは、コマンドキューにエンキューされたコマンドの処理にかかった時間を取得することができます。コマンドのエンキュー時にイベントを登録することができます。これは本来、コマンドの終了を待機したりコマンドキューを同期したりするために使用するものですが、このイベントを利用すると実行時間を計測することもできます。

cl_event handler;
clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, workSize, NULL, 0, NULL, &handler);
clWaitForEvents(1, &handler);
cl_ulong start, end;
clGetEventProfilingInfo(handler, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &start, NULL);
clGetEventProfilingInfo(handler, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &end, NULL);
clReleaseEvent(handler);

実行時間はナノ秒単位で取得することができます。また、タイマの精度はデバイス情報から取得できます。

ゼロコピー

メモリ転送の高速化といえばやっぱりゼロコピーです。OpenCLではCPUやオンボードGPUといったデバイスを使用することができます。これらのデバイスはホストと同じ物理メモリを使用しているため、メモリ転送することなく共有することが可能です。
ゼロコピーを行うためには、clCreateBufferのcl_mem_flagsにCL_MEM_USE_HOST_PTRかCL_MEM_ALLOC_HOST_PTRを指定する必要があります。CL_MEM_USE_HOST_PTRはすでに存在するホスト側のポインタをメモリオブジェクトとして使用します。CL_MEM_ALLOC_HOST_PTRではホスト側からアクセス可能な領域にメモリ確保を行います。
これらの方法で確保されたメモリ領域にホストからアクセスするためには、メモリのマップおよびアンマップの処理を行う必要があります。

cl_mem mem = clCreateBuffer(context, flag | CL_MEM_ALLOC_HOST_PTR, memsize, NULL, NULL);
void* ptr = clEnqueueMapBuffer(commandQueue, mem, CL_FALSE, flag, offset, size, 0, NULL, NULL, NULL);
//ptrに対して読み書き
clEnqueueUnmapMemObject(commandQueue, mem, ptr, 0, NULL, NULL);

Shared Virtual Memory

OpenCL 2.0から利用可能になった、ホストとデバイスで同じポインタを使用することができる機能です。SVMの利用には少しだけ癖があります。
まず、メモリ確保するとcl_mem型のオブジェクトではなくポインタが返されます。

void* ptr = clSVMAlloc(context, CL_MEM_READ_WRITE, size, 0);

このポインタを用いてホスト側から直接データを書き込むことができます。ただし、高度なSVMがサポートされていないデバイスではマップおよびアンマップの処理が必要になることがあります。

clEnqueueSVMMap(commandQueue, CL_TRUE, CL_MAP_WRITE, ptr, size, 0, NULL, NULL);
// 高度なSVMがサポートされていない場合、ここでしかポインタにアクセスできない
clEnqueueSVMUnmap(commandQueue, ptr, 0, NULL, NULL);

自分の所有するデバイスがSVMをどのレベルまでサポートしているかは、次のコードで調べることができます。

cl_device_svm_capabilities caps;
clGetDeviceInfo(device, CL_DEVICE_SVM_CAPABILITIES, sizeof(cl_device_svm_capabilities), &caps, NULL);
int svmCoarse     = 0!=(caps & CL_DEVICE_SVM_COARSE_GRAIN_BUFFER);
int svmFineBuffer = 0!=(caps & CL_DEVICE_SVM_FINE_GRAIN_BUFFER);
int svmFineSystem = 0!=(caps & CL_DEVICE_SVM_FINE_GRAIN_SYSTEM);

最後に、SVMをカーネルの引数に設定するときにはcl_mem用のものとは違う関数を用います。後のカーネル実行はいつも通りです。

clSetKernelArgSVMPointer(kernel, index, ptr);

cl_memと比べると利用できるプラットフォーム/デバイスも少なくコードの変更も必要ですが、便利で動作も速いので是非使っていきたいですね。

レジスタの活用

カーネル実行の高速化において意外に効果があるのが、レジスタを有効活用することです。OpenCLには多様なメモリ空間がありますが、その特徴はだいたい下の表のようになっています。

メモリ空間	レイテンシ	サイズ	特徴
グローバル	高	大	最も基本的なメモリで、引数のうち配列とかはここに入る
ローカル	低	小	ワークグループ内で共有できるメモリで、ホストからはアクセスできない
コンスタント	低	小	デバイスから書き込みできず、キャッシュが効くとレジスタ並みに速い
レジスタ	低	小	最も速いメモリで、カーネル内で宣言したローカル変数や、引数のうち値などはここに入る
ピクチャ	低	大	アクセスが速く容量も大きいが、特殊な命令でアクセスする必要がある

この中でも、レジスタを上手く利用することで高速化が期待できます。具体的には次の行為が効果あります。
・グローバルメモリの同じ領域に何度もアクセスする場合は、一度レジスタに値をコピーする
・何度も同じ計算を行う場合は計算結果をレジスタに入れておく

例として1次元配列A、2次元配列BとCを用いて以下のような計算をするカーネルを考えます。

//サンプルコード
__kernel void function(__global float* A, __global float* B, __global float* C, int2 size)
{
    int y = get_global_id(0);

    A[y] = 0;
    for (int i = 0; i < size.x; i++)
    {
        A[y] += B[size.x * y + i] * C[size.x * y + i];
    }
}

こうすることで高速化が期待できる。

//サンプルコード
__kernel void function(__global float* A, __global float* B, __global float* C, int2 size)
{
    int y = get_global_id(0);

    int idx = size.x * y;
    float tmp = 0;
    for (int i = 0; i < size.x; i++)
    {
        tmp += B[idx + i] * C[idx + i];
    }
    A[y] = tmp;
}

なお、for文の初期値を工夫すればもっと簡潔に書けるじゃんという指摘があると思いますが、まったくその通りでございます。あくまで例としてお楽しみください...。

if文禁止令

これはいろいろなところに書いてあると思いますが、OpenCLなどのSIMD型の計算フレームワークを用いる場合、すべてのスレッドで同じ処理を行う必要があります。したがって、スレッドによって処理が変わる可能性のあるif文は実行できません。しかしif文がないと不便であるため、すべてのスレッドでif文が成立した場合とifが成立しなかった場合の両方のコードを実行し、本来実行されないはずのコードを実行しているときは変数に変更が反映されないように制御するという特殊な方法でif文を実現しています。そのため、

・if...else...の中で重い処理を書くと、実際には両方とも実行するためとても遅い
・else ifを用いて条件分岐を増やすと、その分だけ遅くなる

といったことが起きます。そもそもGPUコアはCPUコアほどリッチではないため、if文は絶対に使わないという心構えでいくといいでしょう。（なお現場はこの限りではない...）

ホストとデバイスの並列化

OpenCLの関数は基本的に非同期実行なので、デバイス側でカーネルを実行している最中にホスト側で次の計算を行うことができます。たとえばメモリを2組用意して、片方にホストが書き込んでいる間にもう片方を使ってデバイスが計算するような構成にすると、かなり高速化することができます。もしホスト側の処理時間とデバイス側の処理時間が等しくなるように切り分けることができれば、単純に考えて2倍の高速化を達成できますよね。
潤沢なメモリと複数計算資源を使った物理で殴りに行くスタイルです。

ヘテロジニアス(マルチデバイス)

OpenCLでは、複数のデバイス(CPUとGPU、複数GPUなど)を用いたヘテロジニアスコンピューティングが可能です。なんか難しそうな雰囲気があって敬遠していたのですが、使ってみるとわりと簡単でした。

まず、使用したいデバイスのリストを作成します。

// プラットフォーム取得
cl_uint platformNumber = 0;
cl_platform_id platformIds[8];
clGetPlatformIDs(8, platformIds, &platformNumber);
cl_platform_id platform = platformIds[PLATFORM_INDEX];

// デバイス取得
cl_uint deviceNumber = 0;
cl_device_id deviceIds[8];
clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 8, deviceIds, &deviceNumber);
devices = (cl_device_id *)malloc(sizeof(cl_device_id) * 2);
devices[0] = deviceIds[DEVICE_INDEX_1];
devices[1] = deviceIds[DEVICE_INDEX_2];

次に、デバイスリストに紐づいたコンテキストを作成します。

context = clCreateContext(NULL, 2, devices, NULL, NULL, NULL);

このコンテキストを用いることで、すべてのデバイスで使えるメモリ、すべてのデバイスで使えるカーネルが取得できます。

cl_program prg = clCreateProgramWithSource(context, 1, (const char**)&sourceString, (const size_t*)&sourceSize, NULL);
clBuildProgram(prg, 2, devices, NULL, NULL, NULL);

最後に、それぞれのデバイス用のコマンドキューを作成します。

commandQueues = (cl_command_queue *)malloc(sizeof(cl_command_queue) * 2);
commandQueues[0] = clCreateCommandQueue(context, devices[0], 0, NULL);
commandQueues[1] = clCreateCommandQueue(context, devices[1], 0, NULL);

これで、メモリオブジェクトやカーネルオブジェクトを共有しながら複数デバイスで計算することができます。同じカーネルを同時に起動してもあまり高速化できませんでしたが、

・メモリ転送とカーネル実行でデバイスを分ける
・同時に実行できる別々のカーネルを割り当てる

などをすると高速化が期待できます。

続かない

C++でOpenCLシリーズは今回で最終回です。次回からはC#でOpenCLをやります。