multithreading - OpenCL float 减少

Question

我想对我的这部分内核代码(1维数据)进行归约:

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

我不想只有1个线程来执行此操作，我想拥有n个线程(n =长度)，最后有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

要开始使用，可以执行以下示例(see Scarpino)。在这里，我们还通过使用OpenCL float4数据类型来利用矢量处理。

请记住，下面的内核返回许多部分和:每个本地工作组一个，然后返回主机。这意味着您必须通过将所有部分总和加回到主机上来执行最终总和。这是因为(至少在OpenCL 1.2中)没有屏障功能来同步不同工作组中的工作项目。

如果不希望将主机上的部分总和相加，则可以通过启动多个内核来解决此问题。这引入了一些内核调用开销，但是在某些应用程序中，额外的损失是可以接受的或可以忽略的。要在下面的示例中执行此操作，您将需要修改主机代码以重复调用内核，然后在输出矢量的数量低于本地大小之后添加逻辑以停止执行内核(详细信息留给您或检查Scarpino reference) 。

编辑:为输出添加了额外的内核参数。添加了点积，将浮点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));
    }
}