GPU Computing: CUDA основы

2012 год: AlexNet выигрывает ImageNet с отрывом в 10 процентных пунктов от второго места. Ключевая техническая деталь в статье: обучение заняло 6 дней на двух GTX 580 GPU. На CPU того времени это заняло бы месяцы. С этого момента GPU стал главным инструментом ML - и понимание CUDA перестало быть экзотикой, превратившись в базовый навык ML-инженера.

• **PyTorch и TensorFlow** реализованы на CUDA: каждая операция - это CUDA kernel. При написании кастомных операторов через torch.autograd.Function необходимо писать CUDA kernel напрямую для достижения максимальной производительности
• **FlashAttention** - CUDA kernel для attention в трансформерах, сокращающий потребление памяти с O(n^2) до O(n) через тайлинг в Shared Memory. Используется в GPT-4, Claude и большинстве современных LLM
• **NVIDIA cuDNN** содержит высокооптимизированные CUDA kernels для свёрток, нормализации, активаций - результат тысяч часов ручной оптимизации под конкретные архитектуры GPU

CUDA kernels: функции для тысяч потоков

NVIDIA GeForce RTX 4090 содержит 16384 CUDA ядра. Для сравнения - топовый CPU имеет 24-32 ядра. Секрет не в том, что каждое GPU ядро мощнее - наоборот, оно проще и медленнее. Секрет в массивном параллелизме: одна инструкция выполняется на тысячах элементов данных одновременно (SIMT - Single Instruction Multiple Threads). CUDA kernel - это C-функция с атрибутом __global__, которая при вызове запускается на тысячах потоков параллельно. Каждый поток знает свой уникальный threadIdx и blockIdx - через них он определяет, с каким элементом данных работать.

CUDA работает только на GPU от NVIDIA. AMD предлагает HIP (гетерогенный интерфейс переносимости) с совместимым API. Apple Silicon использует Metal compute shaders. Для кросс-платформенного кода: OpenCL или Vulkan Compute (следующий урок). PyTorch и TensorFlow скрывают CUDA за Python API, но понимание CUDA необходимо для написания кастомных операторов.

Иерархия: threads, blocks, grid

CUDA организует параллелизм в трёхуровневую иерархию: Thread - один вычислительный поток с собственными регистрами. Block - группа до 1024 потоков, разделяющих Shared Memory и способных синхронизироваться через __syncthreads(). Grid - набор блоков, выполняющих один kernel. Блоки не синхронизируются между собой - это фундаментальное ограничение, позволяющее GPU независимо планировать блоки на разные Streaming Multiprocessors (SM). Размер блока 128-256 потоков часто оптимален: кратность 32 (размер warp) обязательна.

CUDA поддерживает 3D индексацию: threadIdx.x/y/z и blockIdx.x/y/z. Это удобно для обработки изображений (2D) или объёмных данных (3D). Максимальный размер блока: 1024 потока суммарно (например, 32x32 = 1024 для 2D). Максимальный grid размер: 2^31-1 блоков по оси X, 65535 по Y и Z.

Иерархия памяти GPU: от регистров до DRAM

Пропускная способность памяти - главное узкое место GPU вычислений. RTX 4090 имеет 1 ТБ/с пропускной способности GDDR6X (в 10 раз больше, чем у CPU DDR5), но при неправильном паттерне доступа реальная пропускная способность падает в 10-32 раза. Иерархия: Registers (8192 на SM, ~1 пикосекунда) -> Shared Memory (48-96 КБ на SM, ~4 нс) -> L1/L2 Cache (автоматически) -> Global Memory (16-24 ГБ, ~200-400 нс). Ключевой паттерн: данные загружаются из Global Memory в Shared Memory один раз, обрабатываются несколько раз.

Coalesced access - критически важный паттерн: потоки warp должны обращаться к последовательным адресам в Global Memory. При 128-байтовой cache line и 32 потоках в warp: доступ к float (4 байта) по последовательным адресам дает одну транзакцию. Случайный доступ - 32 транзакции. Это разница в 32x производительности.

Warp: единица SIMT исполнения

GPU не выполняет 16384 потоков независимо - под капотом они сгруппированы в warps по 32 потока. Все 32 потока warp выполняют одинаковую инструкцию в один момент времени (SIMT). Проблема возникает при условных ветвлениях: if (idx % 2 == 0). Потоки с чётными idx идут по одной ветке, нечётные - по другой. GPU последовательно выполняет обе ветки, маскируя потоки из другой ветки. Это warp divergence - потенциально двукратное падение производительности для простого if/else.

Warp scheduler GPU скрывает latency через переключение между warps: пока warp ждёт данных из Global Memory (~400 нс = ~800 тактов), другой warp готов к исполнению и занимает SM. GPU нужно достаточное число активных warps для скрытия latency - это Occupancy. Высокий occupancy (>50%) обычно важнее абсолютной минимизации кода kernel.

Ключевые идеи

• **CUDA kernel** - функция, выполняемая тысячами потоков параллельно. Каждый поток идентифицирует свои данные через threadIdx и blockIdx. Защита от выхода за границы обязательна
• **Иерархия памяти** определяет производительность: Shared Memory (48КБ, быстрая) для межпоточного обмена внутри блока; coalesced access к Global Memory; регистры для промежуточных значений
• **Warp = 32 потока** - реальная единица исполнения. Warp divergence при if/else снижает производительность; Occupancy (число активных warps) скрывает latency памяти

Связанные темы

CUDA GPU computing фундаментально меняет подход к параллельным вычислениям:

• OpenCL и Vulkan Compute — Кросс-платформенные альтернативы CUDA: OpenCL для compute, Vulkan Compute для graphics pipeline. Похожая иерархия work-items/work-groups
• GPU Optimization — Продвинутые техники: kernel fusion, persistent kernels, CUDA streams для перекрытия compute и memory transfers

Вопросы для размышления

• FlashAttention использует Shared Memory тайлинг для attention матрицы. Как бы рассуждали о компромиссе между размером тайла и occupancy при проектировании такого kernel?
• CUDA программа работает в 100 раз быстрее на GPU, чем на CPU для умножения матриц, но в 2 раза медленнее для последовательного алгоритма. Почему, и что это говорит о природе задач, подходящих для GPU?
• Как изменится архитектура CUDA программы при переходе с одного GPU на несколько GPU (multi-GPU)? Какие новые проблемы синхронизации возникнут?

Связанные уроки

• arch-09-cache
• la-04-matrix-ops