一个一维数据集被分成多个段,每个工作项处理一个段。它从段中读取一定数量的元素。这个数量事先不知道,并且对于每个段来说都不同。
例如:
+----+----+----+----+----+----+----+----+----+ <-- segments
A BCD E FG HIJK L M N <-- elements in this segment
在处理完所有的段落后,应该将
元素
写入连续的输出内存中,例如:A B C D E F G H I J K L M N
因此,一个片段中元素的绝对输出位置取决于前面片段中元素的数量。 E
在位置4处,因为片段1包含1个元素(A),而片段2包含3个元素。
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
来解决这个问题。