13.18. NDRange

Prev		Next

データ並列プログラミングをOpenCLで実装するにはNDRange、つまりN次元（N-dimensional）の範囲を定義してカーネルのインスタンスを組成させます。他の言葉で表現するならば、Nを1にする場合、1次元の範囲のデータに対して並列処理を求めたいということを定義します。

１次元と定義した（線形）ベクトルのデータ（又は単にC言語でいう配列）数が例えば512個で、さらに全ての配列要素に対して並列・並行に処理をするならば、ワークアイテムは最低でも512個以上必要となります。

Nが2以上となるのは、画像データを取り扱う時です。この場合は、256x2、128x4、64x8、32x16といったように各次元のワークアイテムサイズを削減します。

13.18.1. NDRange：１次元

１次元については「図表：ワークグループ（１次元）」（Figure 13.5, “図表：ワークグループ（１次元）”）に記述した構成となります。

ワークグループをn個にすれば、演算する空間がn個に分割されます。各ワークグループの要素を指定するにはローカルIDを使うことで、ワークアイテムを指定しアクセスできます。

ワークグループ数が０個の場合は、全ての配列要素に対して並列・並行に処理をするためにOpenCL-C言語で以下のように、配列のインデックス（添字）をグローバルIDとして取得します。

Note

詳しくは「表：ワークアイテム・グループ関数」（Table B.102, “表：ワークアイテム・グループ関数”）を参照ください。

size_t id = get_global_id(0);
output[id] = input[id] * 2;

ここではget_global_id関数を使ってグローバルIDを取得します。

0 と引数を指定していますが 1 次元の NDRange に相当します。

各 id は OpenCL実装ライブラリが内部で割り当てるスレッド番号とでも考えてください。

コードの内容としては input 配列の各要素の 2 倍を output 配列の各要素に代入しています。

ローカルIDを使っていないためワークグループはこの例では１つと推定できます。

13.18.2. NDRange：２次元

2次元は以下の図のように2次元（8x8ピクセル）の映像を並列処理することに使うことができます。図では、ワークグループが４つの、4x4ピクセルの2次元のNDRangeとなります。左上が、workgroup-0、右上がworkgroup-1、右下がworkgroup-3となります。

Figure 13.6. 図表：NDRange2次元

例えば2x128、3x64、8x32の2次元行列データがあるとします。上図二示した例でいえば、行8x列8の2次元イメージデータがをしなくてはいけないので、Nは2とします。

2次元イメージのインデックス（添字）はOpenCL-C言語では以下のように取得します。OpenCL-C言語については後半の章で解説します。

Note

詳しくは「表：ワークアイテム・グループ関数」（Table B.102, “表：ワークアイテム・グループ関数”）を参照ください。

size_t x = get_global_id(0);
size_t y = get_global_id(1);

変数xとyは(x,y)座標として使用することができます。ワークグループと関係なく、グローバルに割り振られた座標を示します。以下の表は2次元（8 x 8）の場合のグローバルIDの展開した結果となります。

Table 13.16. グローバルID（8 x 8）

0,0	1,0	2,0	3,0	4,0	5,0	6,0	7,0
0,1	1,1	2,1	3,1	4,1	5,1	6,1	7,1
以下続く

ローカルIDを２次元（4 x 4）とする場合は下図のようにインデックス空間を表すことができます。

Table 13.17. ローカルID（4 x 4）：ワークグループ個数（4）

0,0	1,0	2,0	3,0	0,0	1,0	2,0	3,0
0,1	1,1	2,1	3,1	0,1	1,1	2,1	3,1
以下続く

Note

詳しくは「表：ワークアイテム・グループ関数」（Table B.102, “表：ワークアイテム・グループ関数”）を参照ください。

NDRangeカーネル.

__kernel void ndrange_id(__global write_only float *output) {

   size_t global_id_0 = get_global_id(0);
   size_t global_id_1 = get_global_id(1);
   size_t global_id_2 = get_global_id(2);

   printf("GlobalIDs (%d:%d:%d)\n", global_id_0, global_id_1, global_id_2);

   size_t local_id_0 = get_local_id(0);
   size_t local_id_1 = get_local_id(1);
   size_t local_id_2 = get_local_id(2);

   printf("LocalIDs (%d:%d:%d)\n", local_id_0, local_id_1, local_id_2);

}

出力.

GlobalIDs (0:0:0)
GlobalIDs (1:0:0)
GlobalIDs (2:0:0)
GlobalIDs (3:0:0)
GlobalIDs (0:1:0)
GlobalIDs (1:1:0)
GlobalIDs (2:1:0)
GlobalIDs (3:1:0)
GlobalIDs (0:2:0)
GlobalIDs (1:2:0)
GlobalIDs (2:2:0)
GlobalIDs (3:2:0)
GlobalIDs (0:3:0)
GlobalIDs (1:3:0)
GlobalIDs (2:3:0)
GlobalIDs (3:3:0)
GlobalIDs (0:4:0)
GlobalIDs (1:4:0)
GlobalIDs (2:4:0)
GlobalIDs (3:4:0)
GlobalIDs (0:5:0)
GlobalIDs (1:5:0)
GlobalIDs (2:5:0)
GlobalIDs (3:5:0)
GlobalIDs (0:6:0)
GlobalIDs (1:6:0)
GlobalIDs (2:6:0)
GlobalIDs (3:6:0)
GlobalIDs (0:7:0)
GlobalIDs (1:7:0)
GlobalIDs (2:7:0)
GlobalIDs (3:7:0)
GlobalIDs (4:0:0)
GlobalIDs (5:0:0)
GlobalIDs (6:0:0)
GlobalIDs (7:0:0)
GlobalIDs (4:1:0)
GlobalIDs (5:1:0)
GlobalIDs (6:1:0)
GlobalIDs (7:1:0)
GlobalIDs (4:2:0)
GlobalIDs (5:2:0)
GlobalIDs (6:2:0)
GlobalIDs (7:2:0)
GlobalIDs (4:3:0)
GlobalIDs (5:3:0)
GlobalIDs (6:3:0)
GlobalIDs (7:3:0)
GlobalIDs (4:4:0)
GlobalIDs (5:4:0)
GlobalIDs (6:4:0)
GlobalIDs (7:4:0)
GlobalIDs (4:5:0)
GlobalIDs (5:5:0)
GlobalIDs (6:5:0)
GlobalIDs (7:5:0)
GlobalIDs (4:6:0)
GlobalIDs (5:6:0)
GlobalIDs (6:6:0)
GlobalIDs (7:6:0)
GlobalIDs (4:7:0)
GlobalIDs (5:7:0)
GlobalIDs (6:7:0)
GlobalIDs (7:7:0)
LocalIDs (0:0:0)
LocalIDs (1:0:0)
LocalIDs (2:0:0)
LocalIDs (3:0:0)
LocalIDs (0:1:0)
LocalIDs (1:1:0)
LocalIDs (2:1:0)
LocalIDs (3:1:0)
LocalIDs (0:2:0)
LocalIDs (1:2:0)
LocalIDs (2:2:0)
LocalIDs (3:2:0)
LocalIDs (0:3:0)
LocalIDs (1:3:0)
LocalIDs (2:3:0)
LocalIDs (3:3:0)
LocalIDs (0:0:0)
LocalIDs (1:0:0)
LocalIDs (2:0:0)
LocalIDs (3:0:0)
LocalIDs (0:1:0)
LocalIDs (1:1:0)
LocalIDs (2:1:0)
LocalIDs (3:1:0)
LocalIDs (0:2:0)
LocalIDs (1:2:0)
LocalIDs (2:2:0)
LocalIDs (3:2:0)
LocalIDs (0:3:0)
LocalIDs (1:3:0)
LocalIDs (2:3:0)
LocalIDs (3:3:0)
LocalIDs (0:0:0)
LocalIDs (1:0:0)
LocalIDs (2:0:0)
LocalIDs (3:0:0)
LocalIDs (0:1:0)
LocalIDs (1:1:0)
LocalIDs (2:1:0)
LocalIDs (3:1:0)
LocalIDs (0:2:0)
LocalIDs (1:2:0)
LocalIDs (2:2:0)
LocalIDs (3:2:0)
LocalIDs (0:3:0)
LocalIDs (1:3:0)
LocalIDs (2:3:0)
LocalIDs (3:3:0)
LocalIDs (0:0:0)
LocalIDs (1:0:0)
LocalIDs (2:0:0)
LocalIDs (3:0:0)
LocalIDs (0:1:0)
LocalIDs (1:1:0)
LocalIDs (2:1:0)
LocalIDs (3:1:0)
LocalIDs (0:2:0)
LocalIDs (1:2:0)
LocalIDs (2:2:0)
LocalIDs (3:2:0)
LocalIDs (0:3:0)
LocalIDs (1:3:0)
LocalIDs (2:3:0)
LocalIDs (3:3:0)

13.18.3. NDRange：３次元

NDRangeのN=3はOpenGLやDirectXなどで使う3次元画像データの処理に適しています。

Figure 13.7. 図表：NDRange3次元

以下の表では、64個のカーネルインスタンス（ワークアイテム）に対して、8個のワークグループとして、各ワークグループに割り当てたワークアイテムは8としました。

「NDRangeカーネル」(NDRangeカーネル)を実行した場合の出力です。

グローバルID（4 x 4 x 4）、ローカルID（2 x 2 x 2）.

LocalIDs (0:0:0)
LocalIDs (1:0:0)
LocalIDs (0:1:0)
LocalIDs (1:1:0)
LocalIDs (0:0:1)
LocalIDs (1:0:1)
LocalIDs (0:1:1)
LocalIDs (1:1:1)
LocalIDs (0:0:0)
LocalIDs (1:0:0)
LocalIDs (0:1:0)
LocalIDs (1:1:0)
LocalIDs (0:0:1)
LocalIDs (1:0:1)
LocalIDs (0:1:1)
LocalIDs (1:1:1)
LocalIDs (0:0:0)
LocalIDs (1:0:0)
LocalIDs (0:1:0)
LocalIDs (1:1:0)
LocalIDs (0:0:1)
LocalIDs (1:0:1)
LocalIDs (0:1:1)
LocalIDs (1:1:1)
LocalIDs (0:0:0)
LocalIDs (1:0:0)
LocalIDs (0:1:0)
LocalIDs (1:1:0)
LocalIDs (0:0:1)
LocalIDs (1:0:1)
LocalIDs (0:1:1)
LocalIDs (1:1:1)
LocalIDs (0:0:0)
LocalIDs (1:0:0)
LocalIDs (0:1:0)
LocalIDs (1:1:0)
LocalIDs (0:0:1)
LocalIDs (1:0:1)
LocalIDs (0:1:1)
LocalIDs (1:1:1)
LocalIDs (0:0:0)
LocalIDs (1:0:0)
LocalIDs (0:1:0)
LocalIDs (1:1:0)
LocalIDs (0:0:1)
LocalIDs (1:0:1)
LocalIDs (0:1:1)
LocalIDs (1:1:1)
LocalIDs (0:0:0)
LocalIDs (1:0:0)
LocalIDs (0:1:0)
LocalIDs (1:1:0)
LocalIDs (0:0:1)
LocalIDs (1:0:1)
LocalIDs (0:1:1)
LocalIDs (1:1:1)
LocalIDs (0:0:0)
LocalIDs (1:0:0)
LocalIDs (0:1:0)
LocalIDs (1:1:0)
LocalIDs (0:0:1)
LocalIDs (1:0:1)
LocalIDs (0:1:1)
LocalIDs (1:1:1)