9.13. 実行プログラムのビルド

戻る		次へ

9.13.1. clBuildProgram

clBuildProgram関数はprogramと関連付けられたOpenCLコンテキスト内の全てもしくは特定のデバイス用の実行可能プログラムを、プログラムソースもしくはバイナリからビルド（コンパイルとリンク）します。

プログラムソースもしくはバイナリを使用して実行可能プログラムをビルド可能。clBuildProgramはclCreateProgramWithSourceまたは clCreateProgramWithBinaryを呼び出して生成したprogramに対して使用します。

プログラムがclCreateProgramWithBinaryを使用して生成された場合、プログラムバイナリは実行バイナリとなる（ライブラリ等ではない）。

実行バイナリの情報は clGetProgramInfo(program, CL_PROGRAM_BINARIES, …) を使用して取得でき、clCreateProgramWithBinaryを使用して新しいプログラムオブジェクトを生成する際に指定できる。

注記

詳しくは「表：clBuildProgram」（表B.67「表：clBuildProgram」）を参照ください。

以下が関数の定義です。

int org.jocl.CL.clBuildProgram(
    cl_program program, //(1)
    int num_devices, //(2)
    cl_device_id[] device_list, //(3)
    String options, //(4)
    BuildProgramFunction pfn_notify, //(5)
    Object user_data) //(6)

(1)	有効なプログラムオブジェクトを指定。
(2)	device_list引数に渡したデバイスの個数を指定。
(3)	programと関連付けられたデバイスのリストへのポインタを指定。
(4)	実行可能プログラムをビルドする際に適用するビルドオプションを指定するNULL終端文字列へのポインタを指定。表B.70「コンパイラ・ビルドオプション」を参照のこと。
(5)	通知ルーチンへの関数ポインタ。通知ルーチンはアプリケーションが登録できるコールバック関数であり、実行可能プログラムがビルドされた際に成功失敗に関わらず呼び出されます。
(6)	pfn_notifyを呼び出す際、user_data引数として渡します。NULL値を指定可能。

9.13.2. ビルドオプション

表B.70「コンパイラ・ビルドオプション」を参照のこと。

定数ループ

Intel GPUや他のGPUは、ループの処理には適した設計が行なわれていません。そのためループ処理では遅延が発生するものとされており、可能な限り避けるべきとの推奨をIntelはしています。

しかしカーネル内でループが避けられないケースは多くあるため、遅延をさける方法として、定数ループが推奨されています。

定数ループではループ回数が事前に知ることができる前提となります。例えば、以下のようにループサイズをカーネル引数として設定する方法があります。

__kernel void test_loop(
    float* input,
    float* result,
    int loop_size) { //(1)

    size_t gid = get_global_id(0);
    float p_result = 0.0f;
    float p_input = input[gid]; //(2)

    for(uint i=0; i < loop_size; i+=8) { //(3)

    }

    result = p_result; //(4)
}

(1)	loop_size変数をカーネル引数としてホストプログラムより渡されます。
(2)	プライベートメモリにコピーします。
(3)	loop_size変数を使ってループ回数が一定のものであることをコンパイラに知らせます。
(4)	計算結果をプライベートメモリにコピーします。

この場合はプライベートメモリを使用することで、メモリレイテンシを低減できます。その他にも、ループの部分アンローリングを使って絶対数を減らすることで、処理の高速化を実現しています。

定数ループのマクロ化

定数ループはマクロでも設定可能です。マクロをカーネルに受け渡す方法は、カーネル引数ではなく、ビルドオプションで指定する方法で行います。

例えば「LOOP_SIZE」というマクロを宣言するには、以下のようなビルドオプションを指定します。

-DLOOP_SIZE=8

前の項目で使ったコードは以下のように調整するだけです。

for(uint i=0; i < LOOP_SIZE; i++) {

}

ループのインデックスはGPUは符号無し整数のインクリメントに最適化されており、uintを使うことが推奨されています。