1
1Dデータセットはセグメントに分割され、各ワークアイテムは1つのセグメントを処理します。それはセグメントから要素の数を読みましたか?要素の数はあらかじめ分かっておらず、セグメントごとに異なります。OpenCLの部分和の計算
例えば:すべてのセグメントが処理された後
+----+----+----+----+----+----+----+----+----+ <-- segments
A BCD E FG HIJK L M N <-- elements in this segment
彼らはelements
連続に出力メモリを書き込む必要があり、そう
A B C D E F G H I J K L M N
一方からの元素の絶対出力位置等セグメントは、前のセグメントの要素数に依存します。 E
は、セグメントが1要素(A)を含み、セグメント2が3要素を含むので位置4にある。
OpenCLのカーネルは、共有/ローカルメモリバッファに各セグメントの要素の数を書き込み、この(擬似コード)のように動作
kernel void k(
constant uchar* input,
global int* output,
local int* segment_element_counts
) {
int segment = get_local_id(0);
int count = count_elements(&input[segment * segment_size]);
segment_element_counts[segment] = count;
barrier(CLK_LOCAL_MEM_FENCE);
ptrdiff_t position = 0;
for(int previous_segment = 0; previous_segment < segment; ++previous_segment)
position += segment_element_counts[previous_segment];
global int* output_ptr = &output[position];
read_elements(&input[segment * segment_size], output_ptr);
}
したがって、各作業項目は、部分和を計算しなければなりませんより大きいidを持つ作業項目がより多くの反復を行うループを使用します。
これを実装するより効率的な方法がありますか(各作業項目は、そのインデックスまでのシーケンスの部分和を計算します)、OpenCL 1.2ではありますか? OpenCL 2はこれにwork_group_scan_inclusive_add
を提供するようです。このようなあなたはLOG2にN部分(接頭辞)の合計を行うことができます