Продолжение. Начало здесь.
Итак, теперь мы попробуем оптимизировать код обработки изображения при помощи локальной памяти и рабочих групп. Будем, как и в случае работы с перемножением матриц, кешировать значения в локальный буфер. Код выглядит так:
// This kernel function convolves an image input_image[imgWidth, imgHeight]
// with a mask of size maskSize by caching submatrices from the input image
// in the device local memory.
__kernel void filterImageCached(
__global unsigned char* inImg,
__global unsigned char* outImg,
const int bytesPerPix,
const unsigned int maskSize,
__constant float* mask)
{
// Get work-item identifiers.
int x = get_global_id(0);
int y = get_global_id(1);
int lx = get_local_id(0);
int ly = get_local_id(1);
int imgW = get_global_size(0);
int imgH = get_global_size(1);
int offset = ((y * imgW) + x) * bytesPerPix;
// Declare submatrix used to cache the input image on local memory.
__local unsigned char sub[SUB_SIZE][SUB_SIZE][4];
sub[ly][lx][0] = inImg[offset];
sub[ly][lx][1] = inImg[offset + 1];
sub[ly][lx][2] = inImg[offset + 2];
if (bytesPerPix == 4) // if we have alfa-channel
sub[ly][lx][3] = inImg[offset + 3];
// Synchronize all work-items in this work-group.
barrier(CLK_LOCAL_MEM_FENCE);
// Check if the mask cannot be applied to the current pixel
if (x < maskSize / 2
|| y < maskSize / 2
|| x >= imgW - maskSize / 2
|| y >= imgH - maskSize / 2)
{
outImg[offset] = 0;
outImg[offset + 1] = 0;
outImg[offset + 2] = 0;
if (bytesPerPix == 4) // if we have alfa-channel
outImg[offset + 3] = inImg[offset + 3];
return;
}
// Apply mask based on the neighborhood of pixel inputImg.
int outSumB = 0;
int outSumG = 0;
int outSumR = 0;
for (size_t k = 0; k < maskSize; k++)
{
for (size_t l = 0; l < maskSize; l++)
{
// Calculate the current mask index.
size_t maskIdx = (maskSize - 1 - k) + (maskSize - 1 - l) * maskSize;
// Compute output pixel.
size_t maskLX = lx - maskSize / 2 + k;
size_t maskLY = ly - maskSize / 2 + l;
// Сheck if the current input pixel is in the local memory
if (maskLX >= 0 && maskLX < SUB_SIZE && maskLY >= 0 && maskLY < SUB_SIZE)
{
outSumB += sub[maskLY][maskLX][0] * mask[maskIdx];
outSumG += sub[maskLY][maskLX][1] * mask[maskIdx];
outSumR += sub[maskLY][maskLX][2] * mask[maskIdx];
}
else
{
// Read the current input pixel from the global memory
size_t maskX = x - maskSize / 2 + k;
size_t maskY = y - maskSize / 2 + l;
int offsetM = ((maskY * imgW) + maskX) * bytesPerPix;
outSumB += inImg[offsetM] * mask[maskIdx];
outSumG += inImg[offsetM + 1] * mask[maskIdx];
outSumR += inImg[offsetM + 2] * mask[maskIdx];
}
}
}
// Write output pixel.
outImg[offset] = MinMaxVal(0, 255, outSumB);
outImg[offset + 1] = MinMaxVal(0, 255, outSumG);
outImg[offset + 2] = MinMaxVal(0, 255, outSumR);
if (bytesPerPix == 4) // if we have alfa-channel
outImg[offset + 3] = inImg[offset + 3];
}
Вкратце работает так: создаем кеш размером SUB_SIZE на SUB_SIZE (в нашем случае 16 на 16) и записываем в него пиксели (строки 116-122). На краях рабочей группы озможны случаи, когда нам требуются пиксели, которые не закешированы и тогда мы извлекаем значения из глобальной памяти (строки 60-75). А дальше все как обычно: вычисляем сумму согласно коэффициентам фильтра и записываем ее в глобальный массив (строки 80-84).
Теперь самое интересное - результаты, которые довольно неожиданные. Использование кеша должно было улучшить производительность, по аналогии с прошлым разом, но получилось, что неоптимизированный метод (функция filterImage) работает быстрее, чем оптимизированный метод (функция filterImageCached), причем довольно заметно: 2.2-2.3мсек "неоптимизированного" против 2.7-2.8мсек "оптимизированного" кода для картинки 2048 на 1536. Я ожидал хотя бы незначительного, но прироста, а тут получилось значительное падение. Причины точно не ясны, но можно все же сделать несколько предположений. Например, аппаратный кеш видеокарты (у RTX 3060 он 3 мегабайта) вместе с быстрой памятью (192bit GDDR6) делает излишним дополнительные оптимизации в этом случае, но в то же время накладные расходы (как минимум дополнительные 4 сравнения в строке 60 и синхронизация тредов в строке 30) никуда не деваются. Тут, конечно, было бы интересно посмотреть на результаты с какой-то более слабой видеокартой c 64-битной шиной памяти - например GTX 1030. Так что не все так просто с оптимизациями и всегда надо очень тщательно проверять что и на каком оборудовании запускается.
Еще один нюанс с отладкой kernel-кода: оказывается, там можно использовать самый обычный printf(). Но, конечно, делать это надо с осторожностью, потому что обработка даже маленькой картинки 400 на 400 пикселей даст на выходе 16 тысяч трейсов в терминале, в которых можно потеряться.
Еще один важный нюанс при работе с OpenCL состоит в том, что при использовании рабочих групп непременно надо, чтобы размер входного массива был обязательно кратен размерам рабочей группы. Потому что я довольно долго не мог понять, почему неоптимизированная функция работает, а оптимизированная (и обработка разбита на группы тредов 16 на 16) не работает вообще. Оказалось, что enqueueNDRangeKernel возвращала мне значиние ошибки -54, что есть значение константы CL_INVALID_WORK_GROUP_SIZE. Об этом написали добрые люди на форуме:
Error -54 is CL_INVALID_WORK_GROUP_SIZE - values in your “globalWorkSize” array are not divisible with values in your “localWorkSize” array, and that’s it (remember that global work size is really total number of threads along each dimension, and not the size of the “block” of threads along corresponding dimension).
Так что неудивительно, что OpenCL отказался обрабатывать мою картинку размером 1012 на 737 пикселей. После того, как я сменил разрешение на 1024 на 768, все заработало. Проверяйте коды ошибок - это может быть очень полезно.
Тестовая платформа: Ryzen 7 3700X, 16GB RAM, RTX 3060 12GB GDDR6.
Весь код можно посмотреть здесь.
