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

Prev		Next

16.8.1. build

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

プログラムソースもしくはバイナリを使用して実行可能プログラムをビルド可能が可能です。buildは生成したprogramオブジェクトに対して使用します。

Note

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

以下が関数の定義です。

build(
    options=[], #(1)
    devices=None, #(2)
    cache_dir=None) #(3)

(1)	programと関連付けられたデバイスのリストを指定。
(2)	実行可能プログラムをビルドする際に適用するビルドオプションを指定するNULL終端文字列を指定。Table B.70, “コンパイラ・ビルドオプション”を参照のこと。
(3)	ビルドしたバイナリのキャッシュを有効化。かつディレクトリを指定。

ビルド関数については以下のように使用します。

program = cl.Program(ctx, """
    __kernel void square(__global int* data) {

    }
    """).build(options=["-DSIZE=256"], cache_dir="cache")

カーネル関数のソース文字列をビルドしますが、引数でSIZE変数を256に設定し、キャッシュディレクトリをcacheフォルダに指定します。cacheフォルダにはバイナリファイルが保管されます。

16.8.2. ビルドオプション

Table 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を使うことが推奨されています。