min（）和max（）中减少OpenCL的中性元素 [英] Neutral element for min() and max() in OpenCL reduction

问题描述

我正在通过OpenCL在GPU上缩小（寻找最小和最大的） float [] 数组。

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

应该是中性元素为 max（） - 最大函数？
OpenCL文档给出 MAXFLOAT ， HUGE_VALF 和 INFINITY 或无符号） float 值。
举例来说，将中性元素设置为 -INFINITY 是否合理？

我使用 HUGE_VALF 作为 min（）的中性元素，但是文档也会说 HUGE_VALF 被用作一个错误值，所以也许这是一个坏主意。

还原内核（Code）：

#define NEUTRAL_ELEMENT HUGE_VALF
#define REDUCTION_OP min

__kernel void reduce __global float *权重，
__local float * weights_cached
）
{
unsigned int id = get_global_id（0）;

//载入数据
if（id< {{point_count}}）{
weights_cached [get_local_id（0）] =权重[id];
} else {
weights_cached [get_local_id（0）] = NEUTRAL_ELEMENT;
}

barrier（CLK_LOCAL_MEM_FENCE）;

//减少
（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]）;
障碍（CLK_LOCAL_MEM_FENCE）;
}

//保存
权重[get_group_id（0）] =权重_缓存[0];





$ b我实际上已经结束使用 fmin（）和 fmax（）和 NAN 作为中性元素 - 这基本上保证按照 OpenCL文档作为数值总是返回（ NAN 仅在两个 NAN

解决方案

引用OpenCL标准：

HUGE_VALF的计算结果为+ infinity。

所以没有实际的区别（除了隐含的意图）使用 HUGE_VALF 和 INFINITY ;要么对于 min 减少正常工作。在清晰度方面，我略微偏好 INFINITY ，因为 HUGE_VALF 在概念上用于边缘情况下的返回，这不是。

同样，对 max 使用 -INFINITY 如果你的数组包含无穷大，




$ b $> $ MAX_FLOAT 将不能正确表现为中性元素。

I'm doing a reduction (finding the minimum and maximum) of a float[] array on a GPU through OpenCL.

I'm loading the some elements from global memory into local memory for each workgroup. When the global size isn't a multiple of the workgroup size, I pad the global size, such that it becomes a multiple of the global size. Work-items past the end of the array put the neutral element of the reduction into local memory.

But what should that neutral element be for max() -- the maximum function? The OpenCL documentation gives MAXFLOAT, HUGE_VALF and INFINITY as very large positive (or unsigned) float values. Does it makes sense to have the neutral element to be -INFINITY for example?

Right now I'm using HUGE_VALF as the neutral element for min(), but the docs also say that HUGE_VALF is used as an error value, so maybe that's a bad idea.

Reduction kernel (Code):

#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];
}

Edit: I actually ended up using fmin() and fmax() together with NAN as the neutral element -- this is basically guaranteed to work according to the OpenCL documentation as the numerical value will always be returned (NAN is only returned if two NAN values are given).

解决方案

Quoting the OpenCL standard:

HUGE_VALF evaluates to +infinity.

So there's no real difference (except for implied intent) between using HUGE_VALF and INFINITY; either will work correctly for a min reduction. In terms of clarity, I have a slight preference for INFINITY, as HUGE_VALF is conceptually intended for edge-case returns, which this isn't.

Similarly, use -INFINITY for a max reduction.

MAX_FLOAT will not behave correctly as a neutral element if your array contains infinities.

这篇关于min（）和max（）中减少OpenCL的中性元素的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！