floating-point - OpenCL 归约中 min() 和 max() 的中性元素

Question

我正在通过 OpenCL 对 GPU 上的 float[] 数组进行缩减(求最小值和最大值)。

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

但是 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)。

Answer 1

引用OpenCL标准:

HUGE_VALF evaluates to +infinity.

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

类似地，使用-INFINITY 来减少max。

如果您的数组包含无穷大，

MAX_FLOAT 将无法正确地作为中性元素运行。