Индивидуальная траектория 04 · 15 конспектов · две параллельные линии

Многопоточность и алгоритмы на GPU

Траектория о том, как заставить алгоритм выполняться не в одном потоке, а в тысячах — на ядрах процессора и на видеокарте. Основная линия — CUDA на C++: устройство видеокарты как вычислителя, компиляция, ядра, память, финальные проекты от сортировки до фрактала. Сопутствующая линия — те же идеи из Python с компилятором Numba: сначала многопоточные алгоритмы на CPU (там же — главный вопрос траектории: какие алгоритмы вообще можно распараллелить, а какие нет), затем ядра CUDA без единой строчки C++. Линии пересекаются: модель «грид → блоки → треды» одна и та же, и финальный проект обеих линий — одна задача, решённая двумя инструментами.

Линия C · основная · последовательно

CUDA на C++

Конспекты читаются по порядку. Весь код — CUDA C++, компилятор nvcc. По материалам лаборатории суперкомпьютерных и квантовых вычислений ДВФУ.

Инструменты и модельC01–C03
  1. C01 Рабочее место: установка CUDA и компиляция .cu CUDA Toolkit и nvcc; что лежит в .cu-файле и как один компилятор собирает код для двух процессоров; флаги компиляции; Google Colab для тех, у кого нет NVIDIA; «Hello, CUDA!» как дымовой тест.
  2. C02 Программная модель CUDA: грид, блоки, треды Почему GPU быстрее и когда нет; словарь CUDA одной аналогией; спецификаторы __global__/__device__/__host__; запуск ядра, threadIdx/blockIdx/dim3; интерактивный грид с варпами.
  3. C03 Память GPU: шесть этажей Две разные памяти и мост cudaMalloc/cudaMemcpy/cudaFree; регистры, локальная, разделяемая, глобальная, константная, текстурная — с исправленной таблицей; банки shared; тренажёр «куда положить данные».
Ядра и данныеC04–C06
  1. C04 Первые ядра: от сложения чисел к векторам Полное сложение векторов; формула глобального индекса и защита if (i < n); округление вверх; когда цикл можно раздать тредам — map против рекуррентности; тренажёр индексации. ⇄ стык с P02: критерий независимости итераций
  2. C05 Общение тредов: скалярное произведение и редукция __shared__ и __syncthreads; наивная сумма → древовидная за log₂ n шагов; atomicAdd между блоками; правило «барьер вне if»; пошаговый тренажёр дерева.
  3. C06 Меряем время: события CUDA и пропускная способность cudaEvent вместо системных часов; SAXPY и метрика ГБ/с с долей от пика; честный замер с копированиями и когда GPU проигрывает CPU; калькулятор пропускной способности.
Алгоритмы и проектыC07–C09
  1. C07 Перемножение матриц: наивное и тайловое Двумерный грид; наивное ядро и его 2N чтений DRAM; тайлы в разделяемой памяти — трафик в 16 раз меньше; два обязательных барьера; интерактивная карта тайлов.
  2. C08 Финальные проекты: битонная сортировка и Мандельброт Сортирующая сеть Бэтчера и магия i ^ j; фрактал по треду на пиксель с записью PPM без библиотек; галерея студенческих проектов; тренажёр сети на живых числах. ⇄ стык с P06: та же задача на Numba
  3. C09 Разреженные матрицы: форматы и SpMV Матрицы из пустоты: COO и CSR по косточкам (измерено: память ×645, скорость ×215); ядро «строка на тред» и его дисбаланс; «варп на строку», cuSPARSE и scipy.sparse; тренажёр — редактор матрицы с живым CSR.
Линия P · сопутствующая · параллельно линии C

Python: многопоточность и Numba

Те же идеи из Python. P01–P03 не требуют видеокарты вообще; P04–P06 — CUDA без C++. Можно начинать после C02.

Многопоточные алгоритмы · CPUP01–P03
  1. P01 Numba: компилятор для numpy-кода @njit и ускорение в 59 раз одной строкой; цена первого вызова; режим nopython и TypingError; что ускоряется, а что нет; тренажёр «ускорит или нет».
  2. P02 Из последовательного алгоритма — в многопоточный Паттерны map/reduce/scan/рекуррентность с замерами (prange: ×19.5 на 24 ядрах); что компилятор не распараллелит; гонка вживую: 833 тысячи вместо 10 миллионов; parallel_diagnostics; GIL.
  3. P03 Ufunc за минуту: @vectorize и @stencil Своя универсальная функция: один проход вместо трёх numpy-проходов (84 → 10 мс); @guvectorize; стенсилы и «Жизнь» Конвея одним ядром — живой автомат на canvas.
Алгоритмы на видеокарте · GPUP04–P06
  1. P04 CUDA без C++: ядра на Python @cuda.jit и kernel[blocks, threads]; cuda.grid как формула индекса; словарь соответствий C++ ↔ Numba; симулятор CUDASIM для машин без видеокарты; тренажёр-переводчик.
  2. P05 Данные на видеокарте из Python to_device / copy_to_host / device_array; данные живут на устройстве, а не копируются вокруг каждого ядра; редукция через cuda.shared.array — перевод C05; тренажёр-счётчик копирований PCIe.
  3. P06 Проект: Мандельброт на Numba и итоги Тот же фрактал, что в C08, — через @cuda.jit и PIL, сверен с эталоном до пикселя; случайные числа xoroshiro и Монте-Карло π; когда C++, когда Numba; живой Мандельброт с зумом.