OpenCL: CMakeLists.txt проекта, "размывающий" (нижних частот) и "контурный" (верхних частот) фильтры изображений

Работа с изображениями - достаточно типичная задача для видекорарт и OpenCL. Ну точно более типичная, чем перемножение матриц. Я долго искал наглядный пример, чтобы с изображением и чтобы все работало и без лишних телодвижений типа установки каких-то непонятных или очень больших библиотек и в итоге так ничего не нашел, но потом наткнулся на пример работы с bmp-файлом, который подходит мне идеально: небольшой размер библиотеки (по сути все укладывается водин небольшой заголовочный файл - "BMP.h") и не требует никаких дополнительных действий. А еще я решил изучить, насколько удобно использовать cmake, встроенный в Visual Studio 2019 и править конфигурацию через CMakeLists.txt. Оказывается, что довольно удобно - удобнее, чем например руками каждый раз отдельно создавать Release и Debug конфигурации и для каждой из них создавать x86 и x64 варианты (итого 4 конфигурации как минимум). Так вот, у нас есть вот такой CMakeLists.txt:

1. #
2. cmake_minimum_required (VERSION 3.8)
3.  
4. # create project
5. project ("OclFilterImg")
6. # Add source to this project's executable.
7. add_executable (OclFilterImg "OpenCL-filter-img.cpp" "BMP.h")
8.  
9. # OpencCL headers
10. target_include_directories(OclFilterImg PRIVATE "$ENV{CUDA_PATH}/include")
11.  
12. # OpenCL library
13. set(opencl_lib_folder "$ENV{CUDA_PATH}/lib")
14.  
15. # check 32 or 64 bits
16. if(CMAKE_SIZEOF_VOID_P EQUAL 8)
17. # 64 bits
18. set(opencl_lib_folder "${opencl_lib_folder}/x64")
19. elseif(CMAKE_SIZEOF_VOID_P EQUAL 4)
20. # 32 bits
21. set(opencl_lib_folder "${opencl_lib_folder}/Win32")
22. endif()
23.  
24. # link OpenCL library
25. target_link_libraries(OclFilterImg ${opencl_lib_folder}/OpenCL.lib)
26.  

В таком виде оно корректно работоает под Visual Studio 2019 (если создавать с помощью визарда, встроенного в VS2019). Выглядит достаточно лаконично, а выбор между x86 и x64 делается в строке 16, где по размеру указателя на void определяется, какая у нас сборка: x86 или x64. Очень удобно!

Теперь собственно про саму фильтрацию изображения. На входе у нас будет вот такой bmp-файл:

Нам надо наложить на него сначала один фильтр, а потом на результат наложить еще один фильтр и потом еще один. Каждый пиксел это 3 байта RGB цвета + альфа-канал прозрачности, итого 32 бита на пиксел. Причем альфа-канала может и не быть, тогда получается 24 бита на пиксел. Библиотека поддерживает оба варианта и для удобства хранит картинку в виде массива байтов (unsigned char). Именно так мы и передаем изображение в код ядра, где уже через высоту и ширину картинки и количества байтов на пиксел ( 3 или 4 байта) вычисляем позицию пикселя. Первый фильтр у нас делает из цветного изображения монохромное путем усреднения цвета по всем трем RGB-компонентам:

1. // converts an RBG image to grayscale
2. __kernel void rgbToGray(
3. __global unsigned char* inImg,
4. __global unsigned char* outImg,
5. const int bytesPerPix)
6. {
7. // Get work-item identifiers.
8. int x = get_global_id(0);
9. int y = get_global_id(1);
10. int imgW = get_global_size(0);
11.  
12. int offset = ((y * imgW) + x) * bytesPerPix;
13. unsigned char B = inImg[offset];
14. unsigned char G = inImg[offset + 1];
15. unsigned char R = inImg[offset + 2];
16. unsigned char gray = (B + G + R) / 3;
17.  
18. outImg[offset] = gray;
19. outImg[offset + 1] = gray;
20. outImg[offset + 2] = gray;
21. if (bytesPerPix == 4) // if we have alfa-channel
22. outImg[offset + 3] = inImg[offset + 3];
23. }

Особо стоит обратить внимание на строку 10, где мы получаем фактическую ширину картинки функцией get_global_size(0). Этот параметр мы фактически сами задаем при постановке в очередь исполнения кода ядра:

Первый параметр тут ширина картинки, а второй соответственно высота. Этот код по сути говорит, что надо вызвать для каждого пиксела картинки width * height по одному треду. Результирующая картинка имеет все те же RGB-поля с усредненными значениями. В результате мы получим вот такую картинку:

Теперь эту картинку мы пропустим через "размывающий" фильтр (в коде называется "фильтром нижних частот"), который модифицирует каждый пиксел с учетом окружающиъх его пикселей и усредняет его значение, используя вот такие коэффициенты фильтра 5 на 5:

1. // Create a low-pass filter mask.
2. const int lpMaskSize = 5;
3. float lpMask[lpMaskSize][lpMaskSize] =
4. {
5. {.04,.04,.04,.04,.04},
6. {.04,.04,.04,.04,.04},
7. {.04,.04,.04,.04,.04},
8. {.04,.04,.04,.04,.04},
9. {.04,.04,.04,.04,.04},
10. };

 Вот код ядра:

1. int MinMaxVal(int min, int max, int val)
2. {
3. if (val < min)
4. {
5. return min;
6. }
7. else if (val > max)
8. {
9. return max;
10. }
11. return val;
12. }
13.  
14. // This kernel function convolves an image input_image[imgWidth, imgHeight]
15. // with a mask of size maskSize.
16. __kernel void filterImage(
17. __global unsigned char* inImg,
18. __global unsigned char* outImg,
19. const int bytesPerPix,
20. const unsigned int maskSize,
21. __constant float* mask)
22. {
23. // Get work-item identifiers.
24. int x = get_global_id(0);
25. int y = get_global_id(1);
26. int imgW = get_global_size(0);
27. int imgH = get_global_size(1);
28. int offset = ((y * imgW) + x) * bytesPerPix;
29.  
30. // Check if the mask cannot be applied to the current pixel
31. if (x < maskSize / 2
32. || y < maskSize / 2
33. || x >= imgW - maskSize / 2
34. || y >= imgH - maskSize / 2)
35. {
36. outImg[offset] = 0;
37. outImg[offset + 1] = 0;
38. outImg[offset + 2] = 0;
39. if (bytesPerPix == 4) // if we have alfa-channel
40. outImg[offset + 3] = inImg[offset + 3];
41. return;
42. }
43.  
44. // Apply mask based on the neighborhood of pixel inputImg.
45. int outSumB = 0;
46. int outSumG = 0;
47. int outSumR = 0;
48. for (size_t k = 0; k < maskSize; k++)
49. {
50. for (size_t l = 0; l < maskSize; l++)
51. {
52. // Calculate the current mask index.
53. size_t maskIdx = (maskSize - 1 - k) + (maskSize - 1 - l) * maskSize;
54. // Compute output pixel.
55. size_t xM = x - maskSize / 2 + k;
56. size_t yM = y - maskSize / 2 + l;
57. int offsetM = ((yM * imgW) + xM) * bytesPerPix;
58. outSumB += inImg[offsetM] * mask[maskIdx];
59. outSumG += inImg[offsetM + 1] * mask[maskIdx];
60. outSumR += inImg[offsetM + 2] * mask[maskIdx];
61. }
62. }
63.  
64. // Write output pixel.
65. outImg[offset] = MinMaxVal(0, 255, outSumB);
66. outImg[offset + 1] = MinMaxVal(0, 255, outSumG);
67. outImg[offset + 2] = MinMaxVal(0, 255, outSumR);
68. if (bytesPerPix == 4) // if we have alfa-channel
69. outImg[offset + 3] = inImg[offset + 3];
70. }

Результат работы "размывающего" фильтра получается вот таким:

Теперь попробуем наложить "контурный" фильтр. Он использует тот же код ядра, что и предыдущий, но кэффициенты фильтра вот такие:

1. // Create a high-pass filter mask.
2. const int hpMaskSize = 5;
3. float hpMask[hpMaskSize][hpMaskSize] =
4. {
5. {-1,-1,-1,-1,-1},
6. {-1,-1,-1,-1,-1},
7. {-1,-1,24,-1,-1},
8. {-1,-1,-1,-1,-1},
9. {-1,-1,-1,-1,-1},
10. };

Результат работы такого "контурного фильтра" (в коде он называется "фильтром верхних частот") получается вот такой:

Интересный эффект, не правда ли? Если посмотреть на коэффициенты этого фильтра, то все можно обьяснить: если пиксел окружен такими же точно пикселями, то суммарное значение после прохождения такого фильтра будет равно нулю (24 пикселя с отрицательными коэффициентами плюс один этот пиксел с коэффициентом +24 дают в сумме ноль). А если есть вокруг отличные от этого пикселя, то значение будет больше нуля. То есть такой фильтр по сути будет давать значения больше нуля только на границах, например, какой-то заполненной геометрической фигуры.

Еще важно показать, как мы передаем результат работы предыдущего фильтра на вход следующего:

1. void gpuProcess(const BMP& bmpIn, BMP& bmpOut)
2. {
3. const auto imgWidth = bmpIn.bmp_info_header.width;
4. const auto imgHeight = bmpIn.bmp_info_header.height;
5. const uint32_t bytesPP = bmpIn.bmp_info_header.bit_count / 8;
6. cl::Buffer inImg(context, CL_MEM_READ_WRITE | CL_MEM_HOST_READ_ONLY | CL_MEM_COPY_HOST_PTR, imgWidth * imgHeight * bytesPP, (void*) bmpIn.data.data());
7. cl::Buffer grayImg(context, CL_MEM_READ_WRITE | CL_MEM_HOST_READ_ONLY, imgWidth * imgHeight * bytesPP);
8. cl::Buffer lpfImg(context, CL_MEM_READ_WRITE | CL_MEM_HOST_READ_ONLY, imgWidth * imgHeight * bytesPP);
9. cl::Buffer lpMaskBuf(context, CL_MEM_READ_ONLY | CL_MEM_HOST_NO_ACCESS | CL_MEM_COPY_HOST_PTR, lpMaskSize * lpMaskSize * sizeof(lpMask[0]), lpMask);
10. cl::Buffer hpfImg(context, CL_MEM_READ_WRITE | CL_MEM_HOST_READ_ONLY, imgWidth * imgHeight * bytesPP);
11. cl::Buffer hpMaskBuf(context, CL_MEM_READ_ONLY | CL_MEM_HOST_NO_ACCESS | CL_MEM_COPY_HOST_PTR, hpMaskSize * hpMaskSize * sizeof(hpMask[0]), hpMask);
12.  
13. cl::Kernel grayKernel(program, "rgbToGray");
14. grayKernel.setArg(0, inImg);
15. grayKernel.setArg(1, grayImg);
16. grayKernel.setArg(2, bytesPP);
17.  
18. cl::Kernel lpfKernel(program, "filterImage");
19. lpfKernel.setArg(0, grayImg);
20. lpfKernel.setArg(1, lpfImg);
21. lpfKernel.setArg(2, bytesPP);
22. lpfKernel.setArg(3, lpMaskSize);
23. lpfKernel.setArg(4, lpMaskBuf);
24.  
25. cl::Kernel hpfKernel(program, "filterImage");
26. hpfKernel.setArg(0, lpfImg);
27. hpfKernel.setArg(1, hpfImg);
28. hpfKernel.setArg(2, bytesPP);
29. hpfKernel.setArg(3, hpMaskSize);
30. hpfKernel.setArg(4, hpMaskBuf);
31.  
32. cl::CommandQueue queue(context, device);
33. queue.enqueueNDRangeKernel(grayKernel, cl::NullRange, cl::NDRange(imgWidth, imgHeight));
34. queue.enqueueNDRangeKernel(lpfKernel, cl::NullRange, cl::NDRange(imgWidth, imgHeight));
35. queue.enqueueNDRangeKernel(hpfKernel, cl::NullRange, cl::NDRange(imgWidth, imgHeight));
36. queue.enqueueReadBuffer(hpfImg, CL_TRUE, 0, bmpOut.data.size() * sizeof(bmpOut.data[0]), &bmpOut.data[0]);
37. }

Весь код можно найти здесь.