CUDA: вычисления на GPU из C++

CUDA — платформа NVIDIA для программирования графических процессоров на расширенном диалекте C++. Видеокарта содержит тысячи простых ядер, и задачи, которые раскладываются на независимые элементы — умножение матриц, свёртки нейросетей, обработка изображений, физические симуляции — на ней выполняются в десятки раз быстрее, чем на CPU. Именно CUDA лежит под PyTorch, TensorRT и большинством научных пакетов, поэтому C++-разработчик в ML, HPC или компьютерном зрении рано или поздно пишет ядра сам.

Как это работает

Программа на CUDA состоит из хост-кода, выполняемого на CPU, и ядер (kernels) — функций с квалификатором __global__, которые запускаются на GPU. Запуск ядра записывается синтаксисом kernel<<<blocks, threads>>>(args): указывается сетка из блоков и число потоков в блоке. Каждый поток вычисляет свой индекс через blockIdx, blockDim и threadIdx и обрабатывает свой элемент данных. Потоки внутри блока исполняются группами по 32 — варпами (warps) — и по SIMT-модели выполняют одну инструкцию над разными данными.

Память устроена иерархически. Глобальная память GPU велика, но медленна; разделяемая память (__shared__) внутри блока быстра, но ограничена десятками килобайт; регистры принадлежат потоку. Классическая оптимизация — загрузить фрагмент данных в разделяемую память, синхронизировать потоки вызовом __syncthreads() и работать с ним локально. Данные между хостом и устройством копируют через cudaMemcpy или используют unified memory, где перенос выполняет драйвер. Компилирует всё это nvcc: он отделяет device-код, генерирует PTX и машинный код для указанных архитектур (-arch=sm_86), а хост-часть передаёт обычному компилятору.

Поверх сырых ядер существуют библиотеки: cuBLAS для линейной алгебры, cuDNN для нейросетей, Thrust — аналог STL с thrust::device_vector и параллельными алгоритмами, CUB для блоковых примитивов.

Когда применять

CUDA оправдана, когда работа делится на тысячи одинаковых независимых операций и объём вычислений перекрывает стоимость копирования данных на устройство. Умножение матриц 4096×4096 — идеальный случай; рекурсивный обход дерева с непредсказуемыми ветвлениями — плохой. Прежде чем писать ядро, проверьте, нет ли готовой реализации в cuBLAS, Thrust или библиотеках фреймворков: собственные ядра пишут для нестандартных операций, слияния нескольких шагов в один проход и специфичных форматов данных.

Типичные ошибки

Самая дорогая ошибка — расходящиеся ветвления внутри варпа: если половина из 32 потоков идёт по if, а половина по else, обе ветки исполняются последовательно. Вторая — некоалесцированный доступ к глобальной памяти, когда соседние потоки читают далёкие адреса и пропускная способность падает на порядок. Третья — игнорирование кодов возврата cudaError_t; ошибка запуска ядра проявится позже в другом месте. Четвёртая — синхронное копирование маленьких буферов в цикле вместо потоков (streams) и асинхронных операций, из-за чего GPU простаивает. Для диагностики используйте compute-sanitizer и профилировщик Nsight Compute.

Связанные понятия

• GPU computing
• OpenCL
• Nsight / cuda-gdb
• TensorRT
• Cache-friendly код