GPU оптимизация

В 2012 году AlexNet выиграл ImageNet с точностью на 10% выше всех остальных. Секрет был не только в архитектуре сети, но и в том, что Алекс Крижевский и Илья Суцкевер вручную написали CUDA-kernel'ы для свёрток, тщательно учитывая coalescing, shared memory и occupancy на двух GTX 580. Без этих низкоуровневых оптимизаций обучение заняло бы недели вместо шести дней, и дедлайн соревнования бы не сошёлся. Тот же принцип сегодня лежит в основе FlashAttention, cuBLAS и любой high-performance ML-библиотеки: алгоритм - это полдела, остальное - выжимание железа.

• FlashAttention (Tri Dao, 2022) - переписанный attention с tiling в shared memory: 2-4x быстрее стандартного PyTorch на длинных последовательностях
• cuBLAS GEMM - годами вылизываемые tiled-kernel'ы достигают 95%+ пика тензорных ядер
• RAPIDS cuDF и cuML - GPU-аналитика, где обработка миллиардов строк держится на coalesced reads и shared-memory hash tables

Coalescing: warp читает 128 байт одной транзакцией

GPU global memory медленная: ~400-800 циклов латентности против ~20 у shared memory. Спасает её только пропускная способность - сотни GB/s. Чтобы её получить, нужен coalesced access: 32 нити warp'а должны обращаться к смежным 4-байтовым словам в одном 128-байтовом сегменте памяти. Тогда контроллер собирает все 32 запроса в одну транзакцию. Если же нити warp'а ходят 'в разнобой' (stride > 1, случайные индексы), каждая транзакция обслуживает один-два запроса, и эффективная пропускная способность падает в десятки раз. Это объясняет, почему row-major vs column-major выбирается под конкретный паттерн доступа: матрица 1024x1024 float, читаемая по строкам - coalesced, по столбцам - stride 1024 и катастрофа.

Memory transactions на современных NVIDIA GPU: L1/Texture cache работает с 32-байтовыми секторами, L2 - с 32-байтовыми, global - с 32-байтовыми сегментами (4 сегмента на cache line 128 байт). Команда ncu --metrics smsp__sass_thread_inst_executed_op_global_ld покажет, сколько секторов реально загружено на инструкцию. Идеал - 4 сектора на 32 нити (1.0 sector per thread). Stride-2 даст 8 секторов, stride-32 даст 32 сектора - в 8 раз больше работы.

Occupancy: сколько warp'ов помещается на SM

GPU прячет латентность памяти через параллелизм: пока один warp ждёт данные из global memory, SM (Streaming Multiprocessor) переключается на другой готовый warp. Чем больше активных warp'ов на SM, тем лучше прячется латентность. Отношение реально активных warp'ов к максимально возможным называется occupancy. На NVIDIA A100: 64 warp'а максимум на SM, 2048 нитей. Лимиты, ограничивающие occupancy: регистры на нить (255 max), shared memory на block, число блоков на SM (32). Если kernel использует 64 регистра на нить, на SM поместится только 65536/64 = 1024 нити = 32 warp'а = 50% occupancy. Иногда уменьшение регистров через __launch_bounds__ или пересборка с -maxrregcount поднимает occupancy и общую производительность.

Occupancy - не самоцель. 100% occupancy при memory-bound kernel может уступать 50% occupancy с лучшим coalescing. Volkov показал на GTC 2010, что high-ILP kernel'ы могут быть быстрее с 25% occupancy: каждой нити дать больше работы и регистров, использовать instruction-level parallelism вместо thread-level. CUDA Occupancy Calculator и cudaOccupancyMaxActiveBlocksPerMultiprocessor помогают подобрать block size.

Shared memory: programmable cache внутри SM

На каждом SM есть быстрая on-chip память, разделяемая всеми нитями блока: shared memory (или scratchpad на других архитектурах). Латентность - 20-30 циклов вместо 400-800 у global, пропускная способность - терабайты в секунду. В отличие от обычного кэша shared memory полностью под контролем программиста: данные туда явно загружаются, и каждая нить блока знает, что лежит в каком слоте. Классический паттерн - tiling: вместо того чтобы каждая нить читала свой блок данных из global, нити блока кооперативно загружают tile (например, 32x32 элемента) в shared, синхронизируются через __syncthreads(), и затем все совместно обрабатывают tile с быстрым доступом. Для матричного умножения это даёт 5-10x ускорение по сравнению с наивной реализацией.

Shared memory разбита на банки (32 на современных GPU). Если 32 нити warp'а попадают каждая в свой банк - идеальный bandwidth. Если две нити в один банк - bank conflict, обращения сериализуются. Stride-2 при чтении float-массива из shared - частая ловушка: банки 0, 2, 4, ..., 30, 0, 2, ... - 2-way conflict. Padding на 1 элемент (33 столбца вместо 32) часто решает проблему. ncu --metrics l1tex__data_bank_conflicts покажет реальный счёт.

Профилирование: Nsight как рентген kernel'а

Оптимизация без профилирования - угадайка. NVIDIA даёт два инструмента. Nsight Systems (nsys) - timeline-профайлер на уровне процесса: показывает CPU-GPU взаимодействие, копирования памяти, запуски kernel'ов и idle gap'ы. Полезно для поиска проблем верхнего уровня: 'почему GPU стоит 30% времени', 'какие kernel'ы доминируют'. Nsight Compute (ncu) - микроархитектурный профайлер одного kernel: roofline, memory throughput, instruction mix, warp stall reasons, occupancy, bank conflicts. Стандартный workflow: nsys собирает timeline, выявляет 2-3 доминирующих kernel'а, ncu прогоняет их в детальном режиме (--set full), Speed of Light отчёт показывает узкое место (memory-bound, compute-bound, latency-bound).

Roofline analysis - удобная визуализация: по оси X - arithmetic intensity (FLOP/byte), по Y - производительность (TFLOPS). Линия 'roof' - физический предел железа. Memory-bound kernel'ы лежат на наклонной части под линией пропускной способности, compute-bound - на горизонтальной части под пиком FLOPS. Цель оптимизации - подтянуть точку kernel'а к roof'у. Tiling и shared memory повышают intensity (переиспользование данных), сдвигая kernel в compute-bound область.

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

• Coalesced memory access - 32 нити warp'а в одну транзакцию: AoS-to-SoA, row-major выбор под паттерн доступа, проверка через ncu sectors per request
• Occupancy - сколько warp'ов на SM активны параллельно: лимиты по регистрам, shared memory и блокам; 100% не всегда оптимум, ILP может быть важнее
• Shared memory - programmable cache внутри SM с латентностью 20-30 циклов и контролем bank conflicts; tiling переводит memory-bound kernel в compute-bound
• Профилирование через Nsight Systems (timeline) и Nsight Compute (микроархитектура): Speed of Light, roofline, warp stall reasons - оптимизация без замеров это угадайка

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

GPU оптимизация - это сплав низкоуровневого знания железа и микроархитектурных метрик. От AlexNet 2012 до FlashAttention 2022 - одни и те же три рычага: coalescing, shared memory, профайлер.

• CUDA: основы программирования GPU — Базовые абстракции thread/warp/block, поверх которых строятся все оптимизации
• OpenCL и Vulkan Compute — Те же принципы coalescing и shared memory применимы и в OpenCL/Vulkan - меняется только синтаксис
• Кэш и иерархия памяти — Shared memory - явный аналог L1 cache; coalescing - аналог spatial locality в CPU world

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

• Если kernel memory-bound, какие три действия в порядке убывания эффекта стоит попробовать?
• Когда AoS лучше SoA и наоборот - какие признаки задачи диктуют выбор?
• Почему в реальных проектах сначала смотрят Nsight Systems, а не сразу Nsight Compute?

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

• par-12 — OpenCL/Vulkan - основа перед GPU-оптимизацией
• par-14 — GPU-оптимизация открывает путь к Spark и distributed
• dl-12 — GPU-обучение нейронных сетей - ключевое применение CUDA оптимизации
• opt-14 — Распределённая оптимизация строится на GPU-кластерах
• alg-01-big-o — Анализ GPU-алгоритмов требует понимания сложности
• arch-09-cache