Правильный способ использования памяти cuda __shared__ для фильтрации изображений ⇐ C++

[Anonymous]

Я пишу код Cuda C++ для фильтрации изображений. Cuda разделяет данные изображения на блоки для параллельной обработки. Для обычной попиксельной обработки, конечно, это быстро. Однако при фильтрации изображений для каждого пикселя нам нужны соседние пиксели, чтобы свернуть его с помощью маски фильтра (ядра фильтра).

Теперь для тех пикселей входного изображения, которые находятся на границе блока cuda, соседние пиксели будут находиться в соседнем блоке, и для этого требуется связь между различными блоками процессора, что делает скорость процесса резко падает!

Насколько я понял, решение для оптимизации такого случая состоит в том, чтобы использовать общую память и хранить все пиксели, необходимые для обработки блока, в этом общем массиве.

Я рассмотрел дополненный общий массив, чтобы сохранить соседние пиксели каждого блока и продолжить фильтрацию.

Вот фрагмент кода, который я написал:

/// в заголовке: #define MASK_WIDTH 5 #define TILE_SIZE 8 /// ... __global__ void local_filt ( const беззнаковый символ* inputImage, беззнаковый символ* выходное изображение, const int * filterKernel, внутренняя высота, ширина интервала ) { __shared__ плитка беззнакового символа[TILE_SIZE + MASK_WIDTH - 1][TILE_SIZE + MASK_WIDTH - 1]; int tx = threadIdx.x; int ty = threadIdx.y; int bx = blockIdx.x * TILE_SIZE; int by = blockIdx.y * TILE_SIZE; int row = by + ty; int col = bx + tx; int cx = MASK_WIDTH/2; int cy = MASK_WIDTH/2; __синхронные потоки(); if (строка