OpenCL 缩减中 min() 和 max() 的中性元素答案

【问题标题】：Neutral element for min() and max() in OpenCL reductionOpenCL 缩减中 min() 和 max() 的中性元素
【发布时间】：2013-01-06 23:05:55
【问题描述】：

我正在通过 OpenCL 在 GPU 上对 float[] 数组进行缩减（找到最小值和最大值）。

我正在将global 内存中的一些元素加载到每个工作组的local 内存中。当全局大小不是工作组大小的倍数时，我填充全局大小，使其成为全局大小的倍数。数组末尾之后的工作项将归约的中性元素放入local 内存中。

但是对于max()，那个中性元素应该是什么——最大的功能？ The OpenCL documentation 将 MAXFLOAT、HUGE_VALF 和 INFINITY 提供为非常大的正（或无符号）float 值。例如，将中性元素设为-INFINITY 是否有意义？

现在我使用HUGE_VALF 作为min() 的中性元素，但文档还说HUGE_VALF 被用作错误值，所以这可能是个坏主意。

归约核（代码）：

#define NEUTRAL_ELEMENT HUGE_VALF
#define REDUCTION_OP min

__kernel void reduce(__global float* weights,
                     __local float* weights_cached
                    )
{
  unsigned int id = get_global_id(0);

  // Load data
  if (id < {{ point_count }}) {
    weights_cached[get_local_id(0)] = weights[id];
  } else {
    weights_cached[get_local_id(0)] = NEUTRAL_ELEMENT;
  }

  barrier(CLK_LOCAL_MEM_FENCE);

  // Reduce
  for(unsigned int stride = get_local_size(0) / 2; stride >= 1; stride /= 2) {
    if (get_local_id(0) < stride) {
      weights_cached[get_local_id(0)] = REDUCTION_OP(weights_cached[get_local_id(0)], weights_cached[get_local_id(0) + stride]);
    barrier(CLK_LOCAL_MEM_FENCE);
  }

  // Save
  weights[get_group_id(0)] = weights_cached[0];
}

编辑： 我实际上最终使用了fmin() 和fmax() 以及NAN 作为中性元素——这基本上可以保证根据OpenCL documentation 工作，因为总是返回数值（NAN 只返回如果给出了两个NAN 值）。

【问题讨论】：

标签： floating-point opencl

【解决方案1】：

引用 OpenCL 标准：

HUGE_VALF 的计算结果为 +infinity。

所以使用HUGE_VALF 和INFINITY 之间没有真正的区别（隐含意图除外）；对于 min 减少，两者都可以正常工作。就清晰度而言，我稍微偏爱INFINITY，因为HUGE_VALF 在概念上是为边缘情况返回而设计的，但事实并非如此。

同样，使用-INFINITY 减少max。

如果您的数组包含无穷大，MAX_FLOAT 将无法作为中性元素正确运行。

【讨论】：

由于这个问题显示了一些与 -INFINITY 相关的不确定性，我们可能要强调这是一个常规支持的东西。您可以将一元 - 运算符应用于 INFINITY 并获得 -infinity 的值，它应该按照其数学含义的预期工作。
@EricPostpischil：对； -infinity 的存在由 OpenCL 标准保证（特别是要求 float 是 IEEE-754 binary32 格式）。