Технология CUDA и параллельные алгоритмы

Перенести вычислительное ядро на GPU обычно мешает не синтаксис. Ядро пишется за вечер, запускается — и работает вдвое медленнее однопоточной версии на CPU. Дальше начинается настоящая работа: понять, во что упирается код, почему warp вместо одной транзакции памяти выпускает тридцать две, зачем __syncthreads() стоит там, где стоит, и почему занятость в сто процентов не сделала ядро быстрее.

Курс устроен вокруг этой работы. Он не пересказывает Programming Guide и не учит программировать: указатели, кэш, выравнивание и CMake считаются известными. Он вводит в CUDA как в технологию и, главное, ставит способ думать о задаче в терминах пропускной способности, дивергенции, занятости и стоимости коммуникации.

Что внутри

Модель исполнения и иерархия памяти. Коалесинг, раскладка данных, shared memory и банковые конфликты. Warp-примитивы, кооперативные группы, атомарные операции и модель памяти libcu++. Базовые параллельные паттерны: редукция, скан, гистограмма, компактификация, сортировка, отбор top-k, тайлинг. Потоки, события и CUDA Graphs. Разреженные структуры и нерегулярный параллелизм. Библиотеки NVIDIA — cuBLAS, cuRAND, cuFFT, cuSPARSE и CCCL. Профилирование в Nsight Systems и Nsight Compute, чтение SASS, сокрытие латентности. Продвинутая работа с памятью: пулы, асинхронные копирования, unified memory, VMM, IPC. Несколько GPU, шардирование и NCCL. Корректность: compute-sanitizer, юнит-тесты ядер, детерминизм, контейнеризация и мониторинг. Отдельная глава — проектирование параллельных алгоритмов: модель стоимости, выбор гранулярности, анти-паттерны.

Сквозной проект

VecLens — движок приближённого поиска ближайших соседей по эмбеддингам, то, что стоит под retrieval в RAG-системах, рекомендациях и дедупликации. Задача выбрана потому, что она одновременно упирается в память, в арифметику и в нерегулярный доступ, а ещё требует конвейера — на ней естественно ложится почти весь материал курса.

Проект растёт вместе с главами: от каркаса CMake и первого ядра расстояния к layout базы векторов, инвертированным спискам, обучению центроидов на cuBLAS, конвейеру на CUDA Graphs, шардированию по картам, тестовому контуру и JIT-специализации ядер. К концу собирается библиотека и CLI с командами генерации данных, обучения, построения индекса, поиска, измерения recall и бенчмарка.

Как читать код курса

Приоритет отдан исходникам и разбору чисел, а не описательному тексту. Каждый приём сопровождается ответом на вопрос, сколько он стоит и как измерить выигрыш. Показатели профилировщика подаются как порядок величины для названного класса карты, а не как универсальная константа, — на своей машине студент получает свои числа и учится их читать.

Чего в курсе нет

Обучения программированию с нуля. OpenCL, SYCL, HIP и OpenACC. Обучения нейросетей и теории машинного обучения. Программирования графики, кроме короткого обзора интеропа. Готовых ответов вида «ставьте блок 256» без объяснения, откуда это число.