代码之家 › 专栏 › 技术社区 › youpilat13 Ty Petrice

opencl 2.x-和归约函数

reduction macos-high-sierra opencl sum

youpilat13 Ty Petrice · 技术社区 · 7 年前

来自上一篇文章: 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执行最后的缩减,但理想情况下,我想直接用内核代码来实现 是的。

建议使用代码片段。

最后一点,我正在使用以下模型开发MacOS High Sierra 10.13.5:

opencl 2.x可以安装在我的macos硬件模型上吗?

当做

1 回复 | 直到 7 年前

colonel of truth 7 年前

应该避免使用原子函数,因为与并行缩减内核相比,原子函数会损害性能。您的内核看起来是在正确的轨道上,但是您需要记住,您必须多次调用它;不要在主机上执行最后的总和(除非您有来自上一次缩减的非常少量的数据)。也就是说,需要一直调用它,直到本地大小等于全局大小。无法对大量数据进行单一调用,因为无法在工作组之间进行同步。

此外,您需要小心设置适当的工作组大小(即本地大小),这取决于本地和全局内存吞吐量和延迟。不幸的是,据我所知,除了自我剖析代码之外,没有办法通过opencl来确定这一点,尽管这并不太难编写,因为ocl为您提供了jit编译。通过经验测试,我发现你应该在遭受过多的银行冲突(太大的本地大小)和全局内存延迟惩罚(太小的本地大小)之间找到一个最佳的位置。最好先做一个基准来确定您的缩减的最佳本地大小,然后使用该本地大小来进行将来的缩减。

编辑: 值得注意的是,将内核调用链接在一起的最佳方法是通过opencl事件。