OpenCL浮点数减少

Question

我想对我的内核代码(1维数据)应用reduce:

__local float sum = 0;
int i;
for(i = 0; i < length; i++)
  sum += //some operation depending on i here;

我想要有n个线程(n =长度),最后有1个线程来计算总和,而不是只有1个线程执行此操作.

在伪代码中,我希望能够写出这样的东西:

int i = get_global_id(0);
__local float sum = 0;
sum += //some operation depending on i here;
barrier(CLK_LOCAL_MEM_FENCE);
if(i == 0)
  res = sum;

有办法吗？

我总和有竞争条件.

Answer 1

为了帮助您入门,您可以执行以下示例(请参阅Scarpino).在这里,我们还通过使用OpenCL float4数据类型来利用向量处理.

请记住,下面的内核会返回一些部分总和:每个本地工作组一个,返回主机.这意味着您必须通过将所有部分金额加回主机来执行最终总和.这是因为(至少在OpenCL 1.2中)没有屏障功能可以同步不同工作组中的工作项.

如果不希望对主机上的部分和求和,可以通过启动多个内核来解决这个问题.这引入了一些内核调用开销,但在某些应用程序中,额外的惩罚是可接受的或无关紧要的.要使用下面的示例执行此操作,您需要修改主机代码以重复调用内核,然后包含逻辑以在输出向量的数量低于本地大小后停止执行内核(详细信息留给您或查看Scarpino参考).

编辑:为输出添加了额外的内核参数.添加了点积乘以浮点4向量.

__kernel void reduction_vector(__global float4* data,__local float4* partial_sums, __global float* output) 
{
    int lid = get_local_id(0);
    int group_size = get_local_size(0);
    partial_sums[lid] = data[get_global_id(0)];
    barrier(CLK_LOCAL_MEM_FENCE);

    for(int i = group_size/2; i>0; i >>= 1) {
        if(lid < i) {
            partial_sums[lid] += partial_sums[lid + i];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    if(lid == 0) {
        output[get_group_id(0)] = dot(partial_sums[0], (float4)(1.0f));
    }
}