OpenCL 浮点总和减少

Question

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

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

有办法吗？

我在 sum 上有一个竞争条件。

Answer 1

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

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

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

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

Answer 2

我知道这是一篇非常古老的帖子，但从我尝试过的所有内容来看，Bruce 的答案不起作用，而 Adam 的答案由于全局内存使用和内核执行开销而效率低下。

Jordan 对 Bruce 的回答的评论是正确的，即该算法在元素数量不均匀的每次迭代中都会崩溃。然而，它本质上与在多个搜索结果中可以找到的代码相同。

这几天我摸不着头脑，部分原因是我选择的语言不是基于 C/C++ 的，而且在 GPU 上调试也很棘手，如果不是不可能的话。但最终，我找到了一个有效的答案。

这是布鲁斯的答案和亚当的答案的组合。它将源从全局内存复制到本地，然后通过重复将上半部分折叠到底部来减少，直到没有数据留下。

结果是一个缓冲区包含与使用的工作组相同数量的项目（因此可以分解非常大的减少），必须由 CPU 求和，否则从另一个内核调用并执行此操作GPU 上的最后一步。

这部分有点过头了，但我相信，这段代码还可以通过从本地内存中顺序读取来避免存储库切换问题。 ** 希望知道的人对此进行确认。

注意：如果您的数据从偏移量零开始，则可以从源中省略全局“AOffset”参数。只需将其从内核原型和用作数组索引一部分的第四行代码中删除...

__kernel void Sum(__global float * A, __global float *output, ulong AOffset, __local float * target ) {
        const size_t globalId = get_global_id(0);
        const size_t localId = get_local_id(0);
        target[localId] = A[globalId+AOffset];

        barrier(CLK_LOCAL_MEM_FENCE);
        size_t blockSize = get_local_size(0);
        size_t halfBlockSize = blockSize / 2;
        while (halfBlockSize>0) {
            if (localId<halfBlockSize) {
                target[localId] += target[localId + halfBlockSize];
                if ((halfBlockSize*2)<blockSize) { // uneven block division
                    if (localId==0) { // when localID==0
                        target[localId] += target[localId + (blockSize-1)];
                    }
                }
            }
            barrier(CLK_LOCAL_MEM_FENCE);
            blockSize = halfBlockSize;
            halfBlockSize = blockSize / 2;
        }
        if (localId==0) {
            output[get_group_id(0)] = target[0];
        }
    }

https://pastebin.com/xN4yQ28N

Answer 3

如果您支持 OpenCL C 2.0 功能，您可以使用新的 work_group_reduce_add() 函数在单个工作组内减少总和

Answer 4

一种简单而快速的减少数据的方法是将数据的上半部分重复折叠到下半部分。

例如，请使用以下极其简单的 CL 代码：

__kernel void foldKernel(__global float *arVal, int offset) {
    int gid = get_global_id(0);
    arVal[gid] = arVal[gid]+arVal[gid+offset];
}

使用以下 Java/JOCL 主机代码（或将其移植到 C++ 等）：

    int t = totalDataSize;
    while (t > 1) {
        int m = t / 2;
        int n = (t + 1) / 2;
        clSetKernelArg(kernelFold, 0, Sizeof.cl_mem, Pointer.to(arVal));
        clSetKernelArg(kernelFold, 1, Sizeof.cl_int, Pointer.to(new int[]{n}));
        cl_event evFold = new cl_event();
        clEnqueueNDRangeKernel(commandQueue, kernelFold, 1, null, new long[]{m}, null, 0, null, evFold);
        clWaitForEvents(1, new cl_event[]{evFold});
        t = n;
    }

主机代码循环 log2(n) 次，因此即使是巨大的数组也能快速完成。 “m”和“n”的小提琴是处理非二次幂数组。

• OpenCL 很容易为任何 GPU 平台很好地并行化（即快速）。
• 内存不足，因为它可以就地工作
• 有效地处理非二次幂数据大小
• 灵活，例如您可以将内核更改为“min”而不是“+”