将数据上传到共享内存中以用于卷积内核

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每一侧都有尺寸.我不清楚代码的以下部分.

1. 在destY和destX指数.将输出索引除以输入切片宽度意味着什么？

2. 在srcY添加srcX索引.为什么destY和destX索引参与srcY添加srcX索引？

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

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

3. 为什么在第二次加载时我们使用偏移TILE_WIDTH * TILE_WIDTH？
4. 一般来说,有两次加载的直观解释是什么？
5. 所有这些问题都可以根据下面的图像进行直观的示例吗？
6. 谢谢!

编辑:图片添加.在绿色中有输出瓦片,在红色中我们有掩模以114索引为中心.很明显,面具借用了不同瓷砖的元素.最后,该图像指的是一个通道.

示例:根据下图,我试图写一个例子.输出磁贴具有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上的第一个问题:通过BS_X BS_Y线程移动(BS_X + 1)(BS_Y + 1)全局存储矩阵.

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

4)一般来说,有两个载荷的直观解释是什么？

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

1)destY和destX索引.将输出索引除以输入切片宽度意味着什么？

这涉及第一个加载阶段,它被指定TILE_WIDTHxTILE_WIDTH从全局内存加载元素并填充共享内存区域的最上部分.

所以,操作

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

平坦化通用线程的2D坐标

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

进行逆运算,因为它计算通用线程相对于共享存储区的2D坐标.

2)srcY添加srcX索引.为什么destY和destX索引参与srcY添加srcX索引？

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

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

(blockIdx.x * TILE_WIDTH, blockIdx.y * TILE_WIDTH)如果块大小和共享内存大小相同,则将是全局内存位置的坐标.由于你也是从neighboor瓷砖"借用"记忆值,你必须将上面的坐标移动(destX - Mask_radius, destY - Mask_radius).

3)为什么在第二次加载时我们使用偏移量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.