Cuda：使用单列常量值对图像（线性化二维数组）进行操作

Question

我正在处理很长的图像，通常有几十万像素。高度通常在 500-1000 像素范围内。该过程涉及逐列修改图像。因此，例如，我有一列常量值需要从图像中的每一列中减去。

目前我将图像分割成更小的块，将它们放入线性化的二维数组中。然后我从与较小块大小相同的常数值列中创建一个线性化的二维数组。然后执行（图像数组 - 常量数组）操作，直到处理完整图像。

是不是应该将常量列复制到设备中，然后逐列操作？或者我应该尝试使 "constant array" 尽可能大，然后执行减法。我不是在寻找 100% 的优化，甚至不是接近于此，而是想知道正确的方法是什么。

如何优化这个过程？任何用于了解有关此类处理的更多信息的资源都将不胜感激。

Answer 1

常量内存最大为 64KB，因此假设您的像素为 4 个字节或更少，那么您应该能够处理高达约 16K 像素的图像高度，并且仍然将整个 "constant column" 放在常量中内存。

之后，你就不需要处理东西了"column by column"。当每个线程都从常量内存中请求相同的值时，常量内存针对访问进行了优化，这完美地描述了您的情况。

因此，您的线程代码可以非常简单：

#define MAX_COL_SIZE 1024
__constant__ float const_column[MAX_COL_SIZE];

__global__ void img_col_kernel(float *in, float *out, int num_cols, int col_size){

  int idx = threadIdx.x + blockDim.x*blockIdx.x;
  if (idx < num_cols)
    for (int i = 0; i < col_size; i++)
      out[idx+i*num_cols] = in[idx+i*num_cols] - const_column[i];
}

（在浏览器中编码，未经测试）

在调用 img_col_kernel 之前，使用 cudaMemcpyToSymbol 在主机代码中设置 const_column。使用一维网格调用内核，包括等于或大于图像宽度 (num_cols) 的线程总数。将 "linearized 2D" 指针传递给您的输入和输出图像到内核（in 和 out）。上面的内核应该运行相当快，并且基本上受宽度为 1000 或更大的图像的内存带宽的限制。对于小图像，您可能希望通过将图像垂直分成 4 块来增加线程数，并使用 4 倍多的线程（和 4 个常量内存区域）进行操作。

Cuda：使用单列常量值对图像（线性化二维数组）进行操作

Cuda: Operating on images (linearized 2d arrays) with a single column of constant values

cuda

image-processing