OpenCL 2.x - Sum Reduction 函数

Question

从之前的 post：，我想知道 OpenCL 2.x 提供的最后功能（不是 1.x 这是上一个 post 的主题），特别是关于允许执行数组缩减的原子函数（在我的例子中是总和缩减）。

有人告诉我 OpenCL 1.x 原子函数 (atom_add) 的性能很差，我可以检查一下，所以我正在寻找一种方法来获得 final reduction function（即对应于每个工作组的每个计算总和）。

我记得我目前使用的典型内核代码类型：

__kernel void sumGPU ( __global const double *input, 
                       __global double *partialSums,
               __local double *localSums)
 {
  uint local_id = get_local_id(0);
  uint group_size = get_local_size(0);

  // Copy from global memory to local memory
  localSums[local_id] = input[get_global_id(0)];

  // Loop for computing localSums
  for (uint stride = group_size/2; stride>0; stride /=2)
     {
      // Waiting for each 2x2 addition into given workgroup
      barrier(CLK_LOCAL_MEM_FENCE);

      // Divide WorkGroup into 2 parts and add elements 2 by 2
      // between local_id and local_id + stride
      if (local_id < stride)
        localSums[local_id] += localSums[local_id + stride];
     }

  // Write result into partialSums[nWorkGroups]
  if (local_id == 0)
    partialSums[get_group_id(0)] = localSums[0];
 }

如您所见，在内核代码执行结束时，我得到了包含所有部分和的数组 partialSums[number_of_workgroups]。

您能告诉我如何使用 OpenCL 2.x 提供的函数的最佳性能来执行此数组的第二次也是最后一次缩减吗？ 一个经典的解决方案是使用 CPU 执行最终归约，但理想情况下，我想直接使用内核代码 .

欢迎提出代码片段的建议。

最后一点，我正在使用以下模型在 MacOS High Sierra 10.13.5 上工作：

OpenCL 2.x 可以安装在我的硬件 MacOS 型号上吗？

Answer 1

应避免使用原子函数，因为与并行缩减内核相比，它们确实会损害性能。您的内核看起来走在正确的轨道上，但您需要记住，您必须多次调用它；不要在主机上执行最终求和（除非您从之前的缩减中获得非常少量的数据）。也就是说，您需要不断调用它，直到您的本地大小等于您的全局大小。无法对大量数据进行单次调用，因为无法在工作组之间进行同步。

此外，您要小心设置适当的工作组大小（即本地大小），这取决于本地和全局内存吞吐量和延迟。不幸的是，据我所知，除了自我分析代码之外，没有办法通过 OpenCL 确定这一点，尽管这并不难编写，因为 OCL 为您提供了 JIT 编译。通过实证测试，我发现您应该在遭受太多存储区冲突（本地大小太大）与全局内存延迟惩罚（本地大小太小）之间找到一个最佳平衡点。最好先做一个基准测试来确定你的缩减的最佳局部大小，然后使用该局部大小进行未来的缩减。

编辑： 同样值得注意的是，将内核调用链接在一起的最佳方式是通过 OpenCL 事件。

OpenCL 2.x - Sum Reduction 函数

OpenCL 2.x - Sum Reduction function

sum

reduction

opencl

macos-high-sierra