在 OpenCL/CUDA 中处理边界条件

Question

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

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

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

在 OpenCL/CUDA 中是否有一种“优雅”的方式来实现这一点？另外，是否建议使用单独的内核处理边界？

Answer 1

在 CUDA 中处理边界最常用的方法是检查所有可能的边界条件并采取相应措施，即：

• 如果“此元素”超出范围，则return（这在 CUDA 中非常有用，您可能会启动比严格必要的线程更多的线程，因此必须提前退出额外的线程以避免写入输出-of-bounds 内存）。
• 如果“此元素”位于/靠近左边框（最小 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

这是我作为练习做的一个例子，一个 2D 高斯模糊：

__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;
} 

Answer 2

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

• CLK_ADDRESS_REPEAT - 超出范围的图像坐标被包装到有效范围。此地址模式只能与归一化坐标一起使用。如果不使用归一化坐标，这种寻址模式可能会生成未定义的图像坐标。
• CLK_ADDRESS_CLAMP_TO_EDGE - 超出范围的图像坐标被限制在一定范围内。
• CLK_ADDRESS_CLAMP32 - 超出范围的图像坐标将返回边框颜色。如果图像通道顺序为 CL_A、CL_INTENSITY、CL_RA、CL_ARGB、CL_BGRA 或 CL_RGBA，则边框颜色为 (0.0f, 0.0f, 0.0f, 0.0f)，如果图像通道顺序为 (0.0f, 0.0f, 0.0f, 1.0f)通道顺序为 CL_R、CL_RG、CL_RGB 或 CL_LUMINANCE。
• CLK_ADDRESS_NONE - 对于这种地址模式，程序员保证用于对图像元素进行采样的图像坐标指的是图像内部的位置；否则结果是不确定的。

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

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