Разработка высокопроизводительных приложений для GPU долгое время оставалась искусством, доступным лишь избранным. NVIDIA CUDA остается фундаментом вычислений, ускоренных графическими процессорами, питая собой всё — от научных симуляций до обучения гигантских языковых моделей. Однако написание корректного, поддерживаемого и быстрого CUDA-кода сопряжено с серьезными трудностями. Ошибки памяти часто прячутся на виду, узкие места производительности невидимы без правильной инструментации, а самописные алгоритмы для GPU редко могут сравниться с эффективностью оптимизированных библиотек.
К счастью, современный инструментарий NVIDIA значительно созрел. Многие из этих проблем теперь имеют простые и элегантные решения. В этой статье мы подробно разберем инструменты, которые NVIDIA предлагает для отладки, бенчмаркинга и улучшения вашего кода. Мы пройдем путь от базового, но ошибочного примера до высокооптимизированного пайплайна, внося минимальные изменения на каждом шаге. Наша цель — сделать код безопаснее, проще в поддержке и значительно быстрее.
01Исходная точка: Пайплайн обработки изображений
Для демонстрации мы возьмем классическую задачу обработки изображений. Представьте себе поток входных изображений в формате RGB. Наша задача — сначала перенести данные с центрального процессора (CPU) на графический процессор (GPU), а затем преобразовать эти изображения из цветного пространства RGB в оттенки серого (grayscale). После конвертации в оттенки серого, для каждого блока (тайла) размером 32x32 пикселя необходимо вычислить медиану. Это делается путем сортировки пикселей внутри тайла и выбора среднего значения. Наконец, вычисленные медианы каждого тайла копируются обратно на CPU. Этот процесс кажется простым, но именно в нем кроются типичные ошибки новичков и узкие места профессионалов. Ниже представлен полный исходный код стартовой версии. Обратите внимание на использование ручного управления памятью, классического синтаксиса запуска ядер и отсутствия современных абстракций.
#define CUDA_CHECK_ERROR(call) do { \
cudaError_t err = call; \
if (err != cudaSuccess) { \
std::cerr << "CUDA error in " << __FILE__ << " at line " << __LINE__ << ": " \
<< cudaGetErrorString(err) << std::endl; \
std::exit(EXIT_FAILURE); \
} \
} while(0)
// Алиас для пикселя изображения
using pixel_t = uint8_t;
// Ядро, преобразующее красное, зеленое и синее изображения в одно серое
__global__ void computeRGBToGray(
const pixel_t* d_image_r,
const pixel_t* d_image_g,
const pixel_t* d_image_b,
pixel_t* d_image_gray,
int width,
int height) {
// Вычисляем глобальный индекс потока в сетке
const int x = threadIdx.x + blockIdx.x * blockDim.x;
const int y = threadIdx.y + blockIdx.y * blockDim.y;
// Проверка границ, выбираем только потоки внутри изображения
if (x < width && y < height) {
// Вычисляем индекс потока в изображении
const int i = x + y * width;
// Конвертируем из RGB в оттенки серого и сохраняем результат в глобальной памяти
d_image_gray[i] = static_cast(
0.299f * d_image_r[i] +
0.587f * d_image_g[i] +
0.114f * d_image_b[i]);
}
}
// Ядро, вычисляющее медиану каждого тайла в изображении в оттенках серого
template
__global__ void computeMedian(
pixel_t *d_image_gray,
pixel_t *d_median,
int width,
int height) {
// Вычисляем глобальный индекс потока в сетке
const int x = threadIdx.x + blockIdx.x * blockDim.x;
const int y = threadIdx.y + blockIdx.y * blockDim.y;
// Проверка границ
if (!(x < width && y < height))
return;
// Выделяем разделяемую память, в которой будем хранить тайл
__shared__ pixel_t tile[TILE_WIDTH * TILE_WIDTH];
// Вычисляем индекс потока в изображении
const int index = x + y * width;
// Загружаем значение оттенка серого тайла из глобальной памяти в разделяемую память
tile[index] = d_image_gray[index];
// Синхронизируем, чтобы убедиться, что все потоки загрузили свои данные
__syncthreads();
// Сортируем массив тайла с помощью однопоточной сортировки пузырьком
if (threadIdx.x == 0 && threadIdx.y == 0) {
for (int i = 0; i < TILE_WIDTH * TILE_WIDTH; ++i)
for (int j = i + 1; j < TILE_WIDTH * TILE_WIDTH; ++j)
if (tile[i] > tile[j])
cuda::std::swap(tile[i], tile[j]);
}
// Каждый блок потоков сохраняет медиану, найденную в среднем индексе после сортировки,
// в глобальный массив медиан
const int medianIndex = (TILE_WIDTH * TILE_WIDTH) / 2;
d_median[blockIdx.x + blockIdx.y * gridDim.x] = tile[medianIndex];
}
int main() {
// Определяем все константы для примера
constexpr auto TILE_WIDTH = 32;
constexpr auto HISTO_SIZE = 256;
constexpr auto NB_TILE_X = 250;
constexpr auto NB_TILE_Y = NB_TILE_X;
constexpr auto IMAGE_LENGTH = TILE_WIDTH * NB_TILE_X;
constexpr auto IMAGE_SIZE = IMAGE_LENGTH * IMAGE_LENGTH;
constexpr auto NB_IMAGES = 3;
constexpr auto INIT_VALUE = 0;
// Выделяем память CPU для хранения медиан тайлов изображений, а также для красных, зеленых, синих и серых изображений
std::vector> h_images_r(NB_IMAGES, std::vector(IMAGE_SIZE, INIT_VALUE));
std::vector> h_images_g(NB_IMAGES, std::vector(IMAGE_SIZE, INIT_VALUE));
std::vector> h_images_b(NB_IMAGES, std::vector(IMAGE_SIZE, INIT_VALUE));
std::vector> h_images_gray(NB_IMAGES, std::vector(IMAGE_SIZE, INIT_VALUE));
std::vector> h_medians(NB_IMAGES, std::vector(NB_TILE_X * NB_TILE_Y));
// Запускаем пайплайн обработки изображений для каждого изображения параллельно
#pragma omp pa
Источник: NVIDIA Developer ↗