Оптимізація CUDA-конвеєра обробки зображень: від багів до бенчмарків

  • PublicPublic
  • 02 Sep, 2026
Оптимізація CUDA-конвеєра обробки зображень: від багів до бенчмарків

У цьому практичному прикладі CUDA обробляється потік червоних, зелених і синіх зображень: дані спочатку копіюються з CPU на GPU, переводяться з RGB у відтінки сірого, потім кожна плитка розміром 32 на 32 пікселі сортується, щоб можна було обчислити її медіану, і наприкінці медіани копіюються назад на хост. У вихідній реалізації визначено два ядра — computeRGBToGray для перетворення кольору та computeMedian<TILE_WIDTH, HISTO_SIZE> для медіани по кожній плитці, — а main виділяє пам'ять, запускає ядра та звільняє ресурси всередині паралельного циклу OpenMP по трьох зображеннях.

У ядрі медіани приховано тонкий, але серйозний дефект. Запис у спільну пам'ять tile[index] = d_image_gray[index] помилково використовує глобальний індекс для запису в спільну пам'ять, яка обмежена одним блоком потоків. Запуск бінарника видає помилку «illegal memory access», і compute-sanitizer точно вказує причину: запис за межі __shared__ на один байт у рядку 55, зафіксована для потоку (0,3,0) у блоці (20,0,0).

Figure 1. The first part of the image-processing pipeline. First, copy the RGB images from host to device, then convert them from RGB to grayscale
Figure 1. The first part of the image-processing pipeline. First, copy the RGB images from host to device, then convert them from RGB to grayscale

Щоб не допускати подібних помилок з індексацією, NVIDIA представила в CCCL новий API запуску. Ядра запускаються через cuda::make_config, cuda::block_dims, cuda::grid_dims та cuda::launch, потім ядро приймає конфігурацію першим параметром і використовує cuda::gpu_thread.index(cuda::grid, config) та cuda::gpu_thread.index(cuda::block, config), щоб отримувати глобальний та блочний індекси окремо. Сирі вказівники також можна замінити на cuda::std::span, cuda::std::mdspan та cuda::shared_memory_mdspan, які в налагоджувальному режимі спрацьовують асертами при доступі за межі масиву.

Figure 2. The second part of the image-processing pipeline: Divide each grayscale image into 32-by-32-pixel tiles, sort the pixels in each tile to select its median, then copy the median values from d
Figure 2. The second part of the image-processing pipeline: Divide each grayscale image into 32-by-32-pixel tiles, sort the pixels in each tile to select its median, then copy the median values from d

Коли код стає безпомилковим, його заміряють за допомогою Nsight Systems. Розмітка коду діапазонами NVTX — такими як nvtx3::scoped_range та nvtxRangePushA/nvtxRangePop — робить таймлайн зручнішим для читання та показує, що ядро обчислення медіани домінує у профілі: воно займає близько 2,1 секунди на зображення і відповідає за 98,5% часу GPU, тоді як операції з пам'яттю — лише за 1,5%.

Figure 3. The initial Nsight Systems timeline with NVTX annotations. Median-computation kernels dominate the profile, accounting for nearly all GPU activity and most of the 6.8-second image-processing
Figure 3. The initial Nsight Systems timeline with NVTX annotations. Median-computation kernels dominate the profile, accounting for nearly all GPU activity and most of the 6.8-second image-processing

З боку CPU весь етап обчислень над зображенням займає 6,8 секунди, і майже весь цей час іде на ядра медіани. Такий таймлайн дає розробнику ясну мету для наступних кроків оптимізації: обчислення медіани — найдорожча частина конвеєра, і наступні етапи розбору зосереджені саме на її покращенні.