cuda - 上传共享内存中的数据用于卷积核

Question

我在理解批加载时遇到了一些困难，正如评论中提到的那样。为了计算像素中的卷积，大小为 5 的掩码必须以该特定像素为中心。图像被划分为瓦片。应用卷积掩码后的这些瓦片是最终输出的瓦片，其大小为 TILE_WIDTH*TILE_WIDTH .对于属于输出图块边界的像素，当这个图块属于图像的边界时，掩码必须从相邻图块借用一些像素。否则，这些借用值将被赋值为零。这两个步骤描述在

if (srcY >= 0 && srcY < height && srcX >= 0 && srcX < width)
    N_ds[destY][destX] = I[src];
else
    N_ds[destY][destX] = 0;

出于这个原因，共享内存数组有 TILE_WIDTH + Mask_width - 1每边的尺寸。我不清楚代码的以下部分。

destY和 destX指数。
将输出索引除以输入平铺宽度是什么意思？

srcY添加 srcX指数。
为什么destY和 destX指数参与srcY添加 srcX指数？
srcY = blockIdx.y * TILE_WIDTH + destY - Mask_radius;srcX = blockIdx.x * TILE_WIDTH + destX - Mask_radius;

为什么在第二次加载中我们使用偏移量 TILE_WIDTH * TILE_WIDTH ?

一般来说，有两个加载的直观解释是什么？

所有这些问题都可以根据下图给出一个直观的例子吗？

谢谢!

编辑:图片已添加。绿色是输出图块，红色是 mask 以 114 索引为中心。很明显，蒙版借用了不同瓷砖的元素。
最后，该图像指的是一个 channel 。

示例:根据下图，我尝试编写了一个示例。输出磁贴有 blockIdx.x=1和 blockIdx.y=1基于此 destY=0和 destX=0 .还， srcY = 1*6+0-3=3 , srcX = 3和 src = (3*18+3)*3+0=171 .根据计算和图像示例，我们没有匹配项。在第一个共享内存中，应该存储的值是具有全局索引 57 的值。 .上述计算有什么问题？有人可以帮忙吗？

#define Mask_width  5
#define Mask_radius Mask_width/2
#define TILE_WIDTH 16
#define w (TILE_WIDTH + Mask_width - 1)
#define clamp(x) (min(max((x), 0.0), 1.0))

__global__ void convolution(float *I, const float* __restrict__ M, float *P,
                            int channels, int width, int height) {
   __shared__ float N_ds[w][w];
   int k;
   for (k = 0; k < channels; k++) {
      // First batch loading
      int dest = threadIdx.y * TILE_WIDTH + threadIdx.x,
         destY = dest / w, destX = dest % w,
         srcY = blockIdx.y * TILE_WIDTH + destY - Mask_radius,
         srcX = blockIdx.x * TILE_WIDTH + destX - Mask_radius,
         src = (srcY * width + srcX) * channels + k;
      if (srcY >= 0 && srcY < height && srcX >= 0 && srcX < width)
         N_ds[destY][destX] = I[src];
      else
         N_ds[destY][destX] = 0;

      // Second batch loading
      dest = threadIdx.y * TILE_WIDTH + threadIdx.x + TILE_WIDTH * TILE_WIDTH;
      destY = dest / w, destX = dest % w;
      srcY = blockIdx.y * TILE_WIDTH + destY - Mask_radius;
      srcX = blockIdx.x * TILE_WIDTH + destX - Mask_radius;
      src = (srcY * width + srcX) * channels + k;
      if (destY < w) {
         if (srcY >= 0 && srcY < height && srcX >= 0 && srcX < width)
            N_ds[destY][destX] = I[src];
         else
            N_ds[destY][destX] = 0;
      }
      __syncthreads();

      float accum = 0;
      int y, x;
      for (y = 0; y < Mask_width; y++)
         for (x = 0; x < Mask_width; x++)
            accum += N_ds[threadIdx.y + y][threadIdx.x + x] * M[y * Mask_width + x];
      y = blockIdx.y * TILE_WIDTH + threadIdx.y;
      x = blockIdx.x * TILE_WIDTH + threadIdx.x;
      if (y < height && x < width)
         P[(y * width + x) * channels + k] = clamp(accum);
      __syncthreads();
   }
}

Answer 1

您的问题在概念上与我在 StackOverflow 上的第一个问题类似:Moving a (BS_X+1)(BS_Y+1) global memory matrix by BS_XBS_Y threads .

您面临以下问题:每个线程块大小 TILE_WIDTHxTILE_WIDTH应该填充大小为 (TILE_WIDTH + Mask_width - 1)x(TILE_WIDTH + Mask_width - 1) 的共享内存区域 .

4) Generally, what is the intuitive explanation of having two loadings?

自共享内存区 (TILE_WIDTH + Mask_width - 1)x(TILE_WIDTH + Mask_width - 1)大于块大小 TILE_WIDTHxTILE_WIDTH并假设它小于 2xTILE_WIDTHxTILE_WIDTH ，那么每个线程最多应该将两个元素从全局内存移动到共享内存。这就是为什么你有一个两阶段加载的原因。

1) The destY and destX index. Dividing the output index by the input tile width what does it means?

这涉及指定加载 TILE_WIDTHxTILE_WIDTH 的第一个加载阶段。元素来自全局内存并填充共享内存区域的最上部。

所以，操作

dest = threadIdx.y * TILE_WIDTH + threadIdx.x;

展平通用线程的 2D 坐标，同时

destX = dest % w;
destY = dest / w; 

进行相反的操作，因为它计算通用线程相对于共享内存区域的二维坐标。

2) The srcY add srcX index. Why destY and destX index take part in srcY add srcX index?

srcY = blockIdx.y * TILE_WIDTH + destY - Mask_radius;

srcX = blockIdx.x * TILE_WIDTH + destX - Mask_radius;

(blockIdx.x * TILE_WIDTH, blockIdx.y * TILE_WIDTH)如果块大小和共享内存大小相同，则将是全局内存位置的坐标。由于您也是从邻居图块“借用”内存值，因此您必须将上述坐标移动 (destX - Mask_radius, destY - Mask_radius) .

3) Why in the second loading we use the offset TILE_WIDTH * TILE_WIDTH?

你有这个偏移量是因为在第一个内存阶段你已经填充了“第一个” TILE_WIDTHxTILE_WIDTH共享内存的位置。

编辑

下图说明了扁平化线程索引 dest之间的对应关系和共享内存位置。在图片中，蓝色框代表通用磁贴的元素，而红色框代表相邻磁贴的元素。蓝色和红色框的并集对应于整体共享内存位置。如您所见，所有 256一个线程块的线程参与填充绿线上方的共享内存的上半部分，而只有 145参与填充绿线下方的共享内存的下部。现在你应该也明白了 TILE_WIDTH x TILE_WIDTH抵消。

请注意，您最多只有 2由于参数的特定选择，每个线程的内存负载。例如，如果您有 TILE_WIDTH = 8 ，那么一个线程块中的线程数为 64 ，而共享内存大小为 12x12=144 ，这意味着每个线程至少负责执行 2共享内存写入自 144/64=2.25 .