在 OpenCL/CUDA 中处理边界条件

Question

给定一个 3D 统一网格，我想设置边框单元格的值相对于网格内最近邻单元格的值。例如，给定一个 10x10x10 的网格，对于坐标 (0, 8, 8) 处的体素，我想设置一个值如下：val(0, 8, 8)=a*val(1 ,8,8)。

因为 a 可以是 any 实数，我认为在这种情况下不能使用纹理 + 采样器。此外，该方法也适用于普通缓冲区。

此外，由于边界体素坐标可以是网格角、边或面的一部分，因此存在 26 (= 8 + 12 + 6) 种查找最近邻的不同选择（例如，如果坐标是在 (0,0,0) 处，其在网格内最近的邻居将是 (1, 1, 1)）。所以有很多潜在的分支。

在 OpenCL/CUDA 中是否有 "elegant" 方法来完成此操作？另外，是否建议使用单独的内核处理边界？

Answer 1

在 OpenCL 中，您可以使用 Image3d 来处理 3d 网格。可以使用采样器和特定地址模式实现边界处理：

CLK_ADDRESS_REPEAT - out-of-range image coordinates are wrapped to the valid range. This address mode can only be used with normalized coordinates. If normalized coordinates are not used, this addressing mode may generate image coordinates that are undefined.

CLK_ADDRESS_CLAMP_TO_EDGE - out-of-range image coordinates are clamped to the extent.

CLK_ADDRESS_CLAMP32 - out-of-range image coordinates will return a border color. The border color is (0.0f, 0.0f, 0.0f, 0.0f) if image channel order is CL_A, CL_INTENSITY, CL_RA, CL_ARGB, CL_BGRA or CL_RGBA and is (0.0f, 0.0f, 0.0f, 1.0f) if image channel order is CL_R, CL_RG, CL_RGB or CL_LUMINANCE.

CLK_ADDRESS_NONE - for this address mode the programmer guarantees that the image coordinates used to sample elements of the image refer to a location inside the image; otherwise the results are undefined.

此外，您可以定义插值的过滤模式（最近邻或线性）。

这符合您的需求吗？否则，请向我们提供有关您的数据及其边界要求的更多详细信息。

Answer 2

在CUDA中处理边界最常用的方法是检查所有可能的边界条件并据此采取行动，即：

如果"this element"超出范围，那么return（这在CUDA中非常有用，你可能会启动比严格需要更多的线程，所以额外的线程必须按顺序提前退出以避免写入越界内存）。
如果"this element"是at/near左边框（最小x）则对左边框做特殊操作。
右、上、下（以及 3D 中的前后）边框相同。

幸运的是，在大多数情况下你可以使用max/min来简化这些操作，这样你就避免了太多的if。我喜欢使用这种形式的表达式：

source_pixel_x = max(0, min(thread_2D_pos.x + j, MAX_X));
source_pixel_y = ... // you get the idea

这些表达式的结果总是介于 0 和某个 MAX 之间，因此将 out_of_bounds 源像素限制在边界像素上。

编辑：正如 DarkZeros 评论的那样，使用 clamp() 函数更容易（并且更不容易出错）。它不仅检查最小值和最大值，还允许像 float3 这样的向量类型并分别限制每个维度。参见：clamp

这是我做的一个例子，一个二维高斯模糊：

__global__
void gaussian_blur(const unsigned char* const inputChannel,
                   unsigned char* const outputChannel,
                   int numRows, int numCols,
                   const float* const filter, const int filterWidth)
{
  const int2 thread_2D_pos = make_int2( blockIdx.x * blockDim.x + threadIdx.x,
                                        blockIdx.y * blockDim.y + threadIdx.y);
  const int thread_1D_pos = thread_2D_pos.y * numCols + thread_2D_pos.x;

  if (thread_2D_pos.x >= numCols || thread_2D_pos.y >= numRows)
  {
      return;  // "this output pixel" is out-of-bounds. Do not compute
  }

  int j, k, jn, kn, filterIndex = 0;
  float value = 0.0;
  int2 pixel_2D_pos;
  int pixel_1D_pos;

  // Now we'll process input pixels.
  // Note the use of max(0, min(thread_2D_pos.x + j, numCols-1)),
  // which is a way to clamp the coordinates to the borders.
  for(k = -filterWidth/2; k <= filterWidth/2; ++k)
  {
      pixel_2D_pos.y = max(0, min(thread_2D_pos.y + k, numRows-1));
      for(j = -filterWidth/2; j <= filterWidth/2; ++j,++filterIndex)
      {
          pixel_2D_pos.x = max(0, min(thread_2D_pos.x + j, numCols-1));
          pixel_1D_pos =  pixel_2D_pos.y * numCols + pixel_2D_pos.x;

          value += ((float)(inputChannel[pixel_1D_pos])) * filter[filterIndex];
      }
  }

    outputChannel[thread_1D_pos] = (unsigned char)value;
}

在 OpenCL/CUDA 中处理边界条件

Handling Boundary Conditions in OpenCL/CUDA

cuda

gpgpu

opencl