OpenCL 浮点数求和减少

我想对我的这段内核代码（一维数据）应用减少：

__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;

有办法吗？

我对总和有竞争条件。

为了让您开始，您可以执行如下示例所示的操作（见斯卡皮诺）。在这里，我们还通过使用 OpenCL float4 数据类型来利用矢量处理。

请记住，下面的内核返回许多部分总和：每个本地工作组一个，返回到主机。这意味着您必须通过将所有部分总和返回到主机来执行最终总和。这是因为（至少对于 OpenCL 1.2）不存在同步不同工作组中的工作项的屏障函数。

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

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