1.7. OpenCLカーネルプログラミング

ここまでカーネルのNDRangeや、グローバルID、ローカルID、グループIDといった名称が頻繁に出てきましたが、これらは全てカーネル内で記述されます。

カーネルはOpenCL-C言語で記述されます。

1.7.1. カーネル関数

カーネルは以下のように、 __kernel という修飾子が付き、さらに戻り値型は常にvoid型の関数を指します。以下がカーネル関数の例です。

__kernel void func(args) { //(1)

}

(1)

__kernel 修飾子を関数に付けます。

1.7.2. 引数とバッファオブジェクト

前の例では引数が args となっていますが、C 言語と同様に OpenCL-C では、型名と変数名を引数で指定します。

__kernel void func(__global float* data) { //(1)

}

(1)

グローバルメモリ領域にあるメモリオブジェクト(正確にはメモリオブジェクトから派生するバッファオブジェクト)をカーネル関数の引数として宣言します。

この例ではグローバルメモリを指定しているので、引数に渡されたdataはメモリオブジェクトになります。

メモリオブジェクトはカーネル内ではなく、カーネルを呼び出すホストアプリケーションで生成し、これをカーネルの引数として複製・転送します。

注記

メモリオブジェクトについては後ほど解説します。現時点ではカーネル関数の引数として渡すためのデータ型程度の認識で構いませんし、この段階でメモリオブジェクトという名前も覚える必要はありません。

1.7.3. カーネルカウンタ

データ並列プログラミングやSPMDの概念では、同一カーネルを実行すると解説してきましたが、同時にカーネルは固有のカウンタを持つという定義がありました。

これをカーネル関数で実現する場合には、まずは以下の関数を使ってカーネルのカウンタであるグローバルIDを取得します。

size_t get_global_id(uint dimindx)

引数のdimindxは次元1の場合は常に「0」の値となります。2次元の場合は、「0」か「1」のいずれかを指定して、取得したいグローバルIDがどの次元に属するかを設定します。関数が戻す値は、ワークアイテムに固有に割り振られる整数型(size_t)のグローバルIDです。

以下のカーネル関数では、グローバルIDをidという変数に代入して、idをグローバルメモリ領域に確保した配列の要素インデックス(添字)として使います。

__kernel void func(__global float* data) {
    size_t id = get_global_id(0); //(1)
    data[id] = data[id] * data[id]; //(2)
}

(1)

グローバルIDを取得

(2)

グローバルIDを配列要素の添字にして、添字がアクセスする要素を二乗にした値をdata変数に代入します。data変数は、ホストプログラムに複製・転送ができます。

Copyright 2018-2019, by Masaki Komatsu