У меня возникли некоторые трудности, чтобы понять пакетную загрузку, как указано в комментариях. Чтобы вычислить свертку в пикселе, маска, размер которой равен 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Вообще, каково интуитивное объяснение наличия двух загрузок?Может ли за всеми этими вопросами последовать интуитивный пример, основанный на изображении ниже?Спасибо!

РЕДАКТИРОВАТЬ: Изображение добавлено. Зеленым цветом обозначены выходные плитки, а красным цветом - маска с центром в индексе 114. Очевидно, что маска заимствует элементы из разных плиток. Наконец, это изображение относится к одному каналу.

Пример: Основываясь на изображении ниже, я попытался написать пример. Выходной тайл имеетblockIdx.x=1 а такжеblockIdx.y=1 на основании этогоdestY=0 а такжеdestX=0, Также, ,srcY = 1*6+0-3=3srcX = 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();
   }
}