OpenCL 2.x - 减和函数

Question

从上一篇文章：strategy-for-doing-final-reduction，我想知道OpenCL 2.x 提供的最后一个功能（不是上面这篇文章的主题的 1.x），尤其是关于允许执行数组缩减的原子函数（在我的情况下是总和缩减）。

有人告诉我 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 执行此最终缩减，但理想情况下，我想直接使用内核代码来执行此操作。

欢迎提供代码 sn-p 的建议。

最后一点，我正在使用以下型号开发 MacOS High Sierra 10.13.5：

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

Answer 1

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

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

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