Настоящий материал представляет собой обзор восьмой главы книги Chris Fregly, «AI Systems Performance Engineering: Optimizing Model Training and Inference Workloads with GPUs, CUDA, and PyTorch», O'Reilly Media, 2026 (Copyright © 2026 Flux Capacitor, LLC).
Прямые цитаты из книги выделены отдельными блоками с пометкой «Цитата из книги». Остальной текст — систематизированное изложение содержания главы, сопоставление с актуальными данными, проверка прогнозов по открытым источникам, а также дополнительные замечания, выходящие за пределы текста главы. Цитаты приводятся в ограниченном объёме в целях обзора и критического анализа.
Восьмая глава книги Fregly посвящена продвинутым методам сокрытия задержек и повышения пропускной способности GPU-ядер. Она продолжает предыдущие главы, где были разобраны базовые оптимизации памяти, и переходит к систематическому анализу того, как профилировать, диагностировать и устранять узкие места, связанные с занятостью, эффективностью варпов и параллелизмом на уровне инструкций.
«Современные GPU-ускоренные нагрузки доводят аппаратное обеспечение до пределов. Многокристальные GPU, подобные Blackwell, соединяют несколько reticle-limited кристаллов через 10 ТБ/с NV-HBI и увеличивают L2 до 126 МБ. Эти архитектурные решения материально меняют компромиссы между памятью и вычислениями, а также точки оптимума занятости. Это делает профилирование и оптимизацию ещё более критичными, чем когда-либо».
В главе формулируется ключевой тезис: производительность GPU-ядра определяется не только количеством запущенных потоков, но и тем, насколько эффективно эти потоки используют доступные ресурсы. Глава разбивает проблему на три взаимодополняющих направления: настройка занятости (occupancy tuning), повышение эффективности варпов (warp efficiency) и раскрытие параллелизма на уровне инструкций (ILP).
Отдельно подчёркивается, что эти три направления не заменяют, а дополняют друг друга. Высокая занятость без ILP даёт много варпов, которые ждут память. Высокий ILP без занятости не может скрыть задержки, потому что мало варпов. Эффективные варпы без достаточного количества инструкций в полёте не могут заполнить конвейеры. Только совместное применение всех трёх подходов позволяет приблизиться к пиковой производительности.
Глава 8 — это систематический разбор трёх взаимодополняющих направлений оптимизации GPU-ядер: занятость, эффективность варпов, ILP. Каждое направление решает свою задачу и не заменяет остальные.
Архитектура Blackwell с 126 МБ L2 и 10 ТБ/с NV-HBI сдвигает баланс между памятью и вычислениями. Методики, работавшие на Ampere и Hopper, требуют пересмотра.
«Прежде чем оптимизировать, мы должны сначала идентифицировать узкие места в нашем коде, чтобы определить, какой аппаратный или программный ресурс ограничивает нашу производительность. Современные GPU NVIDIA сложны, и замедление может происходить из многих источников: пропускная способность памяти, задержка памяти, пропускная способность инструкций, накладные расходы на синхронизацию, недостаточный параллелизм, задержки передачи данных между хостом и устройством и многое другое».
Глава описывает двухэтапный подход к профилированию. На первом этапе используется Nsight Systems (nsys) для получения системной временной шкалы: потоки CPU, запуски ядер GPU, передачи данных, трассировки Python. Это позволяет найти горячие ядра и понять общую картину. На втором этапе используется Nsight Compute (ncu) для детального анализа каждого горячего ядра: достигнутая занятость, причины остановок варпов, пропускная способность памяти, загрузка исполнительных блоков.
Отдельно подчёркивается важность аннотаций NVTX (NVIDIA Tools Extension). Они позволяют размечать фазы выполнения (forward, backward, optimizer step) и делают временную шкалу читаемой. В PyTorch они доступны через torch.cuda.nvtx, а также через torch.profiler.record_function().
Приводится типичный сценарий: на временной шкале Nsight Systems видны промежутки между запусками ядер, а GPU простаивает. Это не всегда означает проблему в CUDA-коде — часто причина в DataLoader: слишком мало num_workers, отсутствие двойной буферизации, предобработка на CPU вместо GPU.
Единая временная шкала: потоки CPU, ядра GPU, диапазоны NVTX и transfer-движки
Экспорт trace из PyTorch profiler: op-уровень, kernels, memcpy
Если на временной шкале Nsight Systems видны промежутки между запусками ядер, а GPU простаивает — проблема не всегда в CUDA-коде. Часто причина в DataLoader: увеличьте num_workers, включите двойную буферизацию, перенесите предобработку на GPU.
Профилирование — двухэтапный процесс: Nsight Systems для общей картины и поиска горячих ядер, Nsight Compute для детального анализа. NVTX-аннотации делают временную шкалу читаемой.
PyTorch profiler (на базе Kineto) даёт op-уровневую разбивку и поддерживает экспорт в Chrome Tracing / Perfetto для совместного анализа.
Прежде чем оптимизировать ядро, убедитесь, что проблема именно в нём, а не в конвейере данных. Простаивающий GPU часто означает неэффективный DataLoader, а не плохое ядро.
«GPU могут останавливаться по четырём принципиально разным причинам: недоиспользование, ограничение задержкой, ограничение памятью и ограничение вычислениями. Эти режимы связаны, но важно понимать каждый из них независимо, чтобы выбирать правильные оптимизации. Часто исправление одного узкого места вскрывает следующее».
Глава вводит четыре режима, в которых может находиться GPU-ядро. Понимание того, в каком режиме работает конкретное ядро, определяет стратегию оптимизации.
Подчёркивается, что при итеративной оптимизации ядро проходит через режимы последовательно: недоиспользование → ограничение задержкой → ограничение памятью → ограничение вычислениями. Каждое исправление вскрывает следующее узкое место — это нормальный процесс.
Memory roof (наклонная) и Compute roof (горизонтальная). Точка ядра показывает: memory-bound или compute-bound
При итеративной оптимизации ядро проходит через режимы последовательно: недоиспользование → ограничение задержкой → ограничение памятью → ограничение вычислениями. Каждое исправление вскрывает следующее узкое место — это нормальный процесс оптимизации.
Каждое GPU-ядро находится в одном из четырёх режимов. Определение режима — первый шаг к правильной стратегии оптимизации.
Roofline-диаграмма в Nsight Compute: положение точки ядра относительно наклонной (memory) и горизонтальной (compute) линий. Дополнительно — анализ причин остановок варпов.
Не оптимизируйте вслепую. Сначала определите режим. Для memory-bound — тайлинг и слияние. Для latency-bound — занятость и ILP. Для compute-bound — пониженная точность и Tensor Cores.
«Одним из наиболее информативных представлений в Nsight Compute является разбивка статистики состояния варпов (Warp State Statistics) для ядра. Это иногда называют причинами остановок варпов. Анализируя эти причины, вы часто можете определить ограничивающий фактор».
Глава систематизирует основные категории остановок варпов и их интерпретацию. Ключевые категории:
Отдельно разбирается пример с Exec Dependency: код float u = x * x; float temp = u + 1.0f; float v = y * y; float sum = temp + v; создаёт ложную цепочку зависимостей, из-за которой умножение y * y не может быть выдано, пока не завершится u + 1.0f. Исправление — перестановка независимых операций подряд.
| Причина | Значение | Оптимизация |
|---|---|---|
| Long Scoreboard | Варп ожидает загрузки из глобальной DRAM | Увеличить занятость, асинхронная предвыборка, TMA |
| Short Scoreboard | Ожидание обмена между shared memory и регистрами | Оптимизировать шаблоны доступа к shared memory |
| Memory Throttle | Конвейеры загрузки/сохранения насыщены | Уменьшить трафик, улучшить объединение |
| Exec Dependency | Ожидание результатов предыдущих инструкций | Перестановка независимых операций, увеличение ILP |
| Compute Unit Busy | Математические блоки насыщены | Пониженная точность, Tensor Cores |
| Not Selected | Варп готов, но ждёт свободного слота | Обычно не проблема при здоровой занятости |
Варп ждёт данные из глобальной памяти (DRAM/HBM). Высокая latency → scoreboard «long»
Варп ждёт короткие по latency операции: разделяемая память ↔ регистры
Nsight Compute даёт разбивку причин остановок варпов. Long Scoreboard → память, Exec Dependency → ILP, Compute Unit Busy → вычисления.
Ложные цепочки зависимостей (как в примере с u = x*x; temp = u+1; v = y*y;) создают Exec Dependency. Исправление — перестановка независимых операций.
Не запоминайте все причины. Смотрите на доминирующую категорию: memory-related → оптимизация памяти; Exec Dependency → ILP; Compute Unit Busy → пониженная точность.
«Занятость — это отношение активных варпов на SM к максимуму (64 варпа / 2048 потоков на SM для Blackwell B200). Низкая достигнутая занятость — индикатор недоиспользования, но высокая занятость не гарантирует производительность».
«Занятость — это средство для сокрытия задержки, а не самоцель».
Глава подробно разбирает понятие достигнутой занятости и её ограничителей. Ключевые ограничители:
Для настройки занятости используются __launch_bounds__ (указывает максимальное число потоков на блок и минимальное число блоков на SM) и CUDA Occupancy API (cudaOccupancyMaxPotentialBlockSize). Компилятор при этом ограничивает использование регистров и разворачивание, что может вызвать лёгкий разлив в локальную память.
Отдельно подчёркивается: переход с 10% на 50% занятости может дать двукратное ускорение, а переход с 50% на 100% — лишь 5–15%. После 50–75% занятости другие факторы (ILP, паттерны доступа к памяти) начинают доминировать.
Баланс между ресурсами SM и параллелизмом. Цель — скрыть latency
| Сценарий | Оптимальная занятость | Почему |
|---|---|---|
| Ядро, ограниченное памятью | 60–80% | Нужно много варпов для сокрытия задержек DRAM |
| Ядро, ограниченное вычислениями | 25–50% | Tensor Cores поддерживают конвейер без массы варпов |
| Ядро с высоким ILP | 25% | Каждый варп обеспечивает 4 независимые операции |
| Ядро, ограниченное задержкой | 50–75% | Баланс между числом варпов и ресурсами на поток |
Пример использования __launch_bounds__:
__global__ __launch_bounds__(256, 8)
void kernel_optimized(...) {
// компилятор ужимается в ~32 регистра на поток
// → 8 блоков × 256 потоков = 2048 потоков на SM
// → занятость 100%
}Занятость — средство для сокрытия задержки, а не самоцель. После 50–75% другие факторы начинают доминировать.
Registers per thread (лимит 255 на Blackwell), shared memory per block, thread count. Для диагностики используются Nsight Compute Occupancy и CUDA Occupancy API.
__launch_bounds__ и CUDA Occupancy API позволяют найти оптимальную конфигурацию запуска. Но следите за разливом регистров в локальную память.
«В SIMT-выполнении варп из 32 потоков выполняет инструкции синхронно. Если потоки внутри варпа выбирают разные пути if/else, варп выполняет обе ветви последовательно. При разделении 50/50 эффективность падает до 50%».
Глава описывает проблему расхождения варпов (warp divergence) и методы её минимизации. Ключевые техники:
if (warpId % 2 == 0) { /* Task A */ } else { /* Task B */ }. Все потоки варпа выбирают одну ветвь.__ballot_sync, __any_sync, __all_sync позволяют варпу коллективно решать, какие lanes нуждаются в «специальной» работе.float cond = x > 0 ? 1.0f : 0.0f; y = cond * f(x) + (1.0f - cond) * g(x);.Также разбираются shuffle-интринсики (__shfl_sync, __shfl_down_sync) для внутриварпового обмена данными через регистры без обращений к shared memory. Пример warp reduce sum:
float warp_reduce_sum(float val) {
for (int offset = 16; offset > 0; offset >>= 1) {
val += __shfl_down_sync(0xFFFFFFFF, val, offset);
}
return val;
}Слева: if/else расходится → варп выполняет обе ветки. Справа: однородный путь — все lanes работают
Расхождение варпов снижает эффективность до 50% при разделении 50/50. Предикация, варп-однородные ветвления и shuffle-интринсики — основные методы борьбы.
Shuffle-интринсики позволяют обмениваться данными через регистры без shared memory и без __syncthreads(). Это быстрее и не создаёт bank conflicts.
PyTorch-операции вроде torch.where и torch.maximum компилируются в предикацию, обеспечивая варп-когерентное выполнение без расхождения.
«ILP — это количество независимых инструкций, которые один поток может выдать подряд, не дожидаясь завершения предыдущих. Высокий ILP позволяет одному варпу поддерживать занятость исполнительных блоков, даже когда отдельные инструкции имеют задержки».
Глава подробно разбирает ILP и его взаимодействие с занятостью. Ключевые положения:
Пример разрыва цепочки зависимостей:
// Без ILP: последовательная цепочка
float u = x * x;
float temp = u + 1.0f;
float v = y * y;
float sum = temp + v;
// С ILP: независимые операции подряд
float u = x * x;
float v = y * y;
float temp = u + 1.0f;
float sum = temp + v;Независимые инструкции перекрываются: пока загрузка x ждёт память, y*y уже выполняется
| ILP | Потоков на SM для ~100% загрузки | Занятость |
|---|---|---|
| 1 (без ILP) | ~1536 (48 варпов) | 75% |
| 2 | ~1024 (32 варпа) | 50% |
| 3 | ~768 (24 варпа) | 37.5% |
| 4 | ~512 (16 варпов) | 25% |
4 независимых warp-планировщика. Каждый такт может выпустить до 2 инструкций: одна math + одна memory
ILP позволяет одному варпу поддерживать занятость исполнительных блоков. С ILP=4 достаточно 25% занятости для насыщения SM.
Унифицированные ядра CUDA выполняют либо FP32, либо INT32 в одном такте. Смешанные потоки INT+FP больше не получают выгоды от dual-issue.
Перестановка независимых операций подряд, разворачивание циклов, увеличение числа независимых аккумуляторов. Оптимум — 2–4 операции на поток.
«Классический паттерн двойной буферизации: пока вычисления идут над текущим тайлом, в фоне загружается следующий. На Hopper и Blackwell этому помогает TMA — Tensor Memory Accelerator, аппаратный движок для массовых многомерных копий из глобальной памяти в разделяемую».
Глава описывает TMA и его взаимодействие с CUDA Pipeline API. Ключевые моменты:
tcgen05.mma, tcgen05.ld и tcgen05.st обеспечивают высокопроизводительные пути данных пятого поколения на Blackwell. Архитектура TMEM и инструкции tcgen05 подробно разобраны в обзоре главы 6.Пример двойной буферизации с TMA и CUDA Pipeline API:
__shared__ float tile[2][TILE_SIZE];
tile[0][threadIdx.x] = global_data[blockIdx.x * TILE_SIZE + threadIdx.x];
__syncthreads();
for (int t = 0; t < numTiles - 1; ++t) {
tile[(t + 1) % 2][threadIdx.x] =
global_data[(blockIdx.x + t + 1) * TILE_SIZE + threadIdx.x];
process(tile[t % 2][threadIdx.x]);
__syncthreads();
}Более современный вариант с использованием CUDA Pipeline API и TMA:
#include <cuda/pipeline>
#include <cooperative_groups.h>
namespace cg = cooperative_groups;
__global__ void kernelWithTMA(const float* __restrict__ global_ptr, int nTiles) {
__shared__ float tile0[TILE_SIZE];
__shared__ float tile1[TILE_SIZE];
float* tiles[2] = { tile0, tile1 };
__shared__ cuda::pipeline_shared_state<cuda::thread_scope_block, 2> state;
auto pipe = cuda::make_pipeline(cg::this_thread_block(), &state);
pipe.producer_acquire();
cuda::memcpy_async(tiles[0], global_ptr + 0 * TILE_SIZE,
cuda::aligned_size_t<32>(TILE_SIZE * sizeof(float)), pipe);
pipe.producer_commit();
for (int t = 1; t < nTiles; ++t) {
pipe.consumer_wait();
processTile(tiles[(t - 1) & 1]);
pipe.consumer_release();
pipe.producer_acquire();
cuda::memcpy_async(tiles[t & 1], global_ptr + t * TILE_SIZE,
cuda::aligned_size_t<32>(TILE_SIZE * sizeof(float)), pipe);
pipe.producer_commit();
}
pipe.consumer_wait();
processTile(tiles[(nTiles - 1) & 1]);
pipe.consumer_release();
}TMA — аппаратный движок для массовых многомерных копий. Он освобождает LD/ST блоки SM для полезной работы. Двойная буферизация с TMA скрывает задержку DRAM.
Пинг-понг между двумя буферами: пока TMA загружает следующий тайл, SM вычисляет на предыдущем. Это требует явного управления через CUDA Pipeline API.
Использование cuda::memcpy_async с cuda::pipeline на Hopper/Blackwell автоматически задействует TMA, если выравнивание и размеры позволяют.
«Для пользователей PyTorch, пишущих высокоуровневый код, обычно не нужно управлять занятостью напрямую — CUDA-ядра и библиотеки PyTorch обрабатывают параметры запуска под капотом. Однако понимание занятости может быть ценным. Если вы пишете кастомные CUDA-ядра как расширения PyTorch, применяются те же принципы».
Глава описывает несколько оптимизаций, доступных без написания CUDA:
model = torch.compile(model, mode="max-autotune") — генерирует объединённые ядра GPU, сливая множество небольших операций в одно более крупное ядро. Это увеличивает работу на запуск ядра и улучшает занятость.optimizer = torch.optim.AdamW(..., fused=True) — объединённая реализация выполняет все шаги оптимизатора в одном ядре, уменьшая накладные расходы на запуск и трафик памяти.y = torch.where(x > 0, x, torch.zeros_like(x)) — векторизованные операции компилируются в предикацию в стиле SIMD, обеспечивая варп-когерентное выполнение без расхождения.Отдельно подчёркивается, что torch.compile не гарантирует оптимальную занятость или ILP — для этого может потребоваться профилирование и дополнительная настройка. Также отмечается, что PyTorch compiler не всегда автоматически использует все возможности — иногда требуется регистрация кастомных ядер через Triton.
Для пользователей PyTorch многие оптимизации доступны без написания CUDA. torch.compile генерирует объединённые ядра, fused оптимизаторы уменьшают накладные расходы, torch.where обеспечивает предикацию.
torch.compile не гарантирует оптимальную занятость или ILP. Для достижения максимальной производительности может потребоваться профилирование и регистрация кастомных ядер через Triton.
Начинайте с torch.compile и fused оптимизаторов. Если производительность недостаточна — профилируйте и пишите кастомные ядра для горячих участков.
«Длительная почти пиковая загрузка может вызвать срабатывание ограничителей мощности на современных GPU для поддержания их в здоровом состоянии. Обязательно учитывайте это при профилировании, бенчмаркинге и настройке. Самый простой способ увидеть, ограничены ли вы по мощности — использовать nvidia-smi с запросом потребляемой мощности, текущих частот и причин троттлинга в реальном времени».
Глава описывает мониторинг энергопотребления как часть процесса оптимизации. Команда:
nvidia-smi \
--query-gpu=power.draw,clocks.current.sm,clocks.current.memory,\
clocks_event_reasons.active \
--format=csv -l 1выводит новую строку каждую секунду с текущим энергопотреблением, частотами графического ядра и памяти, а также причинами троттлинга. Это позволяет точно определить, когда GPU достигает предела мощности и снижает частоты.
Примечание для старых кластеров: на драйверах NVIDIA версии старше 535 то же поле называется clocks_throttle_reasons.active. Если запрос возвращает ошибку «Field not found» — переключитесь на старое имя поля, оно по-прежнему поддерживается на устаревших драйверах.
Программный доступ — через NVML: nvmlDeviceGetPowerUsage() и nvmlDeviceGetEnforcedPowerLimit(). Эти API идеальны для кастомных скриптов или интеграции с системами мониторинга.
При длительной почти пиковой нагрузке GPU может упираться в лимит мощности и снижать тактовые частоты. Мониторинг энергопотребления — часть процесса оптимизации.
nvidia-smi с запросом power.draw, clocks.current.sm, clocks_event_reasons.active (на драйверах <535 — clocks_throttle_reasons.active). Программный доступ через NVML.
Если GPU упирается в лимит мощности, рассмотрите настройку power limit или underclocking для повышения эффективности на ватт.
«В этой главе вы увидели, как выявлять и устранять узкие места GPU-ядер, перемещая работу из медленной глобальной памяти в более быстрые на кристалле ресурсы и вычислительные блоки. Следуя циклу профилирования, диагностики, оптимизации и повторного профилирования, вы можете трансформировать ядра из недоиспользованных или ограниченных памятью в вычислительно-насыщенные, высокопроизводительные процедуры».
Глава завершается чек-листом, который систематизирует процесс оптимизации:
nsys, найдите горячие ядра, углубитесь с ncu.__launch_bounds__ — если регистры ограничивают занятость.torch.where.Оптимизация GPU-ядер — систематический процесс, а не набор трюков. Чек-лист из 12 пунктов покрывает полный цикл.
Профилирование → диагностика → оптимизация → повторное профилирование. Каждое исправление вскрывает следующее узкое место.
Следуйте чек-листу последовательно. Не оптимизируйте вслепую. Проверяйте точность после каждого изменения.
Книга Fregly писалась в 2025 году, когда Blackwell только появлялся на рынке. Ниже представлено сопоставление прогнозов главы 8 с открытыми данными и уточнение цифр.
«Важно настроить автоматизированные тесты производительности для выявления регрессий на ранних этапах цикла разработки».
Подтверждено. Современные фреймворки (PyTorch, TensorRT-LLM) включают встроенные профайлеры и CI-пайплайны с бенчмарками. PyTorch performance HUD отслеживает регрессии в nightly-сборках. MLPerf стал де-факто стандартом воспроизводимых измерений. Инструменты вроде Nsight Systems и Nsight Compute обновляются для каждой новой архитектуры GPU.
«ILP и занятость взаимно компенсируют друг друга. С ILP=4 достаточно 25% занятости для насыщения SM; без ILP требуется 75%».
Подтверждено. Таблица зависимости ILP от необходимого числа потоков (1→1536, 2→1024, 3→768, 4→512) согласуется с данными NVIDIA и архитектурными руководствами. На Blackwell с его 64 варпами на SM и 256 КБ регистрового файла эта зависимость сохраняется. Единственное уточнение: на Blackwell унифицированные ядра CUDA выполняют либо FP32, либо INT32 в одном такте, что снижает эффективность смешанных потоков инструкций INT+FP по сравнению с более ранними архитектурами.
«TMA способен выполнять 1D–5D массовые копии с произвольными шагами между глобальной и разделяемой памятью без блокировки конвейера инструкций SM».
Подтверждено. TMA (Tensor Memory Accelerator) появился в Hopper и развит в Blackwell. CUDA Pipeline API (cuda::memcpy_async, cuda::pipeline) автоматически задействует TMA на поддерживаемых архитектурах. На Blackwell TMA работает в связке с TMEM для обслуживания Tensor Cores пятого поколения (см. обзор главы 6).
«torch.compile генерирует объединённые ядра GPU, сливая множество небольших операций в одно более крупное ядро».
Подтверждено. torch.compile с TorchInductor генерирует объединённые ядра, снижая накладные расходы на запуск. Однако, как отмечается в главе, компилятор не гарантирует оптимальную занятость или ILP. Для максимальной производительности может потребоваться регистрация кастомных ядер через Triton или ручная настройка.
«Длительная почти пиковая загрузка может вызвать срабатывание ограничителей мощности на современных GPU».
Частично подтверждено. На B200 TDP составляет 1000 Вт, на B300 — 1400 Вт. В стойке NVL72 — 130 кВт. Для большинства inference-нагрузок (которые обычно memory-bound, а не compute-bound) ограничение мощности не является основным фактором. Однако для плотных вычислений (например, обучение) ограничение мощности может снижать частоты. На Blackwell Ultra и Rubin вопрос охлаждения и питания становится более острым.
Глава 8 описывает фундаментальные принципы оптимизации, но оставляет за скобками несколько аспектов, которые проявились в реальной эксплуатации Blackwell.
Влияние двухкристальной архитектуры Blackwell на оптимизацию. Глава не упоминает, что B200 состоит из двух кристаллов, соединённых NV-HBI на 10 ТБ/с. Для программиста это выглядит как один GPU, но на уровне производительности памяти могут возникать аномалии, связанные с неравномерностью загрузки кристаллов. В редких случаях для диагностики может потребоваться проверка баланса между кристаллами.
Взаимодействие с NVLink и NVSwitch. Глава фокусируется на внутриядерной оптимизации, но для multi-GPU workloads критично понимание топологии. Оптимизация одного ядра не поможет, если данные передаются через медленный интерконнект: NVLink 5 на B200 даёт 1,8 ТБ/с на GPU, но при выходе за пределы NVL72 пропускная способность падает до 100–200 ГБ/с через InfiniBand. Топология NVLink/NVSwitch и пропускная способность интерконнектов подробно разобраны в обзоре главы 4.
Специфика MoE-моделей. Для MoE-моделей (как DeepSeek-V3) оптимизация attention и MLP слоёв требует учёта all-to-all коммуникаций между экспертами. В главе это не рассматривается, но для реальных inference-нагрузок это часто доминирующий фактор.
Ограничения FP4. Глава упоминает FP4 как множитель производительности, но не рассматривает ограничения: FP4 требует калибровки, не все модели tolerate такую точность, и для критичных слоёв (ранние слои, attention softmax, layer norms) требуется более высокая точность. Transformer Engine решает эту задачу за счёт автоматического выбора точности per-layer, но требует валидации на конкретной модели.
Специфика inference vs training. Глава описывает общие принципы, но не различает оптимизацию для training и inference. Для training доминируют compute-bound ядра (GEMM), для inference — memory-bound (attention, decode). Соответственно, оптимальные стратегии различаются: для training важнее ILP и Tensor Cores, для inference — пропускная способность памяти и низкая задержка.
clocks_throttle_reasons.active).Восьмая глава книги Fregly — систематический разбор трёх взаимодополняющих направлений оптимизации GPU-ядер: настройки занятости, эффективности варпов и параллелизма на уровне инструкций. Ключевая идея: оптимизация — это не набор трюков, а систематический процесс профилирования, диагностики и повторного профилирования. Для инженера это означает, что понимание фундаментальных принципов (иерархия памяти, ILP, занятость, расхождение варпов) важнее, чем знание конкретных цифр текущего поколения GPU.
Проверка прогнозов показывает: книга оказалась точной в своих предсказаниях. Методики, описанные в главе, применимы к Blackwell и, с учётом архитектурных изменений (унифицированные ядра CUDA, TMA, TMEM), остаются актуальными. Главный вывод остаётся неизменным: инвестиции в фундаментальные принципы обеспечивают адаптацию к любому поколению железа.
Источники: Chris Fregly, «AI Systems Performance Engineering» (O'Reilly, 2026), Глава 8 — цитаты приводятся в ограниченном объёме в целях обзора и критического анализа. Документация NVIDIA Nsight Systems, Nsight Compute, CUDA Toolkit, Blackwell Architecture; MLPerf Training/Inference v5.0; PyTorch Documentation; NVIDIA Technical Blog.
This material is a review of Chapter 8 of Chris Fregly, «AI Systems Performance Engineering: Optimizing Model Training and Inference Workloads with GPUs, CUDA, and PyTorch», O'Reilly Media, 2026 (Copyright © 2026 Flux Capacitor, LLC).
Direct quotes from the book are marked with separate blocks labeled «Quote from the book». The rest of the text is a systematized exposition of the chapter's content, a comparison with current data, verification of forecasts against open sources, and additional observations that go beyond the chapter's text. Quotes are provided in limited volume for review and critical analysis purposes.
Chapter 8 of Fregly's book is dedicated to advanced techniques for hiding latency and improving GPU kernel throughput. It continues previous chapters, where basic memory optimizations were discussed, and moves to a systematic analysis of how to profile, diagnose, and eliminate bottlenecks related to occupancy, warp efficiency, and instruction-level parallelism.
«Modern GPU-accelerated workloads are pushing hardware to its limits. Multi-die GPUs like Blackwell connect multiple reticle-limited dies with a 10 TB/s NV-HBI link and increase L2 to 126 MB. These hardware design choices materially change memory-vs-compute tradeoffs and occupancy sweet spots. This makes profiling and optimization even more critical than ever».
The chapter formulates a key thesis: GPU kernel performance is determined not only by the number of launched threads but by how efficiently those threads use available resources. The chapter breaks the problem into three complementary directions: occupancy tuning, warp efficiency improvement, and instruction-level parallelism (ILP).
The chapter separately emphasizes that these three directions do not replace but complement each other. High occupancy without ILP yields many warps that wait on memory. High ILP without occupancy cannot hide latency because there are too few warps. Efficient warps without sufficient instructions in flight cannot fill the pipelines. Only the combined application of all three approaches can approach peak performance.
Chapter 8 is a systematic breakdown of three complementary GPU kernel optimization directions: occupancy, warp efficiency, ILP. Each direction addresses its own task and does not replace the others.
The Blackwell architecture with 126 MB L2 and 10 TB/s NV-HBI shifts the balance between memory and compute. Techniques that worked on Ampere and Hopper require revision.
«Before optimizing, we must first identify the bottlenecks in our code to determine which hardware or software resource is limiting our performance. Modern NVIDIA GPUs are complex, and a slowdown could come from many sources, including memory bandwidth, memory latency, instruction throughput, synchronization overheads, insufficient parallelism, host-device transfer delays, and more».
The chapter describes a two-stage approach to profiling. In the first stage, Nsight Systems (nsys) is used to obtain a system timeline: CPU threads, GPU kernel launches, data transfers, Python traces. This helps find hot kernels and understand the big picture. In the second stage, Nsight Compute (ncu) is used for detailed analysis of each hot kernel: achieved occupancy, warp stall reasons, memory throughput, execution unit utilization.
The chapter separately emphasizes the importance of NVTX (NVIDIA Tools Extension) annotations. They allow marking execution phases (forward, backward, optimizer step) and make the timeline readable. In PyTorch they are available through torch.cuda.nvtx, as well as through torch.profiler.record_function().
A typical scenario is described: the Nsight Systems timeline shows gaps between kernel launches, and the GPU idles. This does not always mean a problem in CUDA code — often the cause is in the DataLoader: too few num_workers, lack of double buffering, CPU-side preprocessing instead of GPU.
A single timeline: CPU threads, GPU kernels, NVTX ranges, and transfer engines
Exporting a PyTorch profiler trace: op-level, kernels, memcpy
If the Nsight Systems timeline shows gaps between kernel launches while the GPU idles, the problem is not always in CUDA code. Often the cause is in the DataLoader: increase num_workers, enable double buffering, move preprocessing to the GPU.
Profiling is a two-stage process: Nsight Systems for the big picture and finding hot kernels, Nsight Compute for detailed analysis. NVTX annotations make the timeline readable.
PyTorch profiler (based on Kineto) provides op-level breakdown and supports export to Chrome Tracing / Perfetto for collaborative analysis.
Before optimizing a kernel, make sure the problem is in the kernel and not in the data pipeline. An idle GPU often means an inefficient DataLoader, not a bad kernel.
«GPUs can stall for four fundamentally different reasons: underutilization, latency bound, memory bound, and compute bound. These regimes are related, but it's important to understand each of them independently in order to choose the right optimizations to pursue. Often, fixing one bottleneck will reveal another».
The chapter introduces four regimes in which a GPU kernel may reside. Understanding which regime a specific kernel is in determines the optimization strategy.
It is emphasized that with iterative optimization, a kernel moves through regimes sequentially: underutilization → latency-bound → memory-bound → compute-bound. Each fix reveals the next bottleneck — this is a normal process.
Memory roof (sloped) and Compute roof (horizontal). The kernel dot shows: memory-bound or compute-bound
With iterative optimization, a kernel moves through regimes sequentially: underutilization → latency-bound → memory-bound → compute-bound. Each fix reveals the next bottleneck — this is normal.
Every GPU kernel is in one of four regimes. Determining the regime is the first step toward the right optimization strategy.
Roofline chart in Nsight Compute: kernel dot position relative to the sloped (memory) and horizontal (compute) lines. Additionally — analysis of warp stall reasons.
Don't optimize blindly. First determine the regime. For memory-bound — tiling and fusion. For latency-bound — occupancy and ILP. For compute-bound — reduced precision and Tensor Cores.
«One of the most insightful views in Nsight Compute is the Warp State Statistics breakdown for a kernel. This is sometimes called the warp stall reasons. By examining these stall reasons, you can often pinpoint the limiting factor».
The chapter systematizes the main categories of warp stalls and their interpretation. Key categories:
The chapter separately analyzes the Exec Dependency example: the code float u = x * x; float temp = u + 1.0f; float v = y * y; float sum = temp + v; creates a false dependency chain, so y * y cannot be issued until u + 1.0f completes. The fix — reorder independent operations to be adjacent.
| Reason | Meaning | Optimization |
|---|---|---|
| Long Scoreboard | Warp waiting for a global DRAM load | Increase occupancy, async prefetch, TMA |
| Short Scoreboard | Waiting on shared memory ↔ register transfers | Optimize shared memory access patterns |
| Memory Throttle | Load/store pipelines saturated | Reduce traffic, improve coalescing |
| Exec Dependency | Waiting on previous instruction results | Reorder independent operations, increase ILP |
| Compute Unit Busy | Math units saturated | Reduced precision, Tensor Cores |
| Not Selected | Warp ready but waiting for a free slot | Usually not a problem at healthy occupancy |
Warp waits for data from global memory (DRAM/HBM). High latency → «long» scoreboard
Warp waits for low-latency operations: shared memory ↔ registers
Nsight Compute provides a breakdown of warp stall reasons. Long Scoreboard → memory, Exec Dependency → ILP, Compute Unit Busy → compute.
False dependency chains (like in the example with u = x*x; temp = u+1; v = y*y;) create Exec Dependency. Fix — reorder independent operations.
Don't memorize all reasons. Look at the dominant category: memory-related → memory optimization; Exec Dependency → ILP; Compute Unit Busy → reduced precision.
«Occupancy is the ratio of active warps per SM to the maximum (64 warps / 2048 threads per SM for Blackwell B200). Low achieved occupancy is an indicator of underutilization, but high occupancy does not guarantee performance».
«Occupancy is a means to an end (hiding latency), not the end goal itself».
The chapter examines in detail the concept of achieved occupancy and its limiters. Key limiters:
To tune occupancy, __launch_bounds__ (specifies max threads per block and min blocks per SM) and the CUDA Occupancy API (cudaOccupancyMaxPotentialBlockSize) are used. The compiler limits register usage and unrolling, which may cause slight local memory spill.
It is separately emphasized: going from 10% to 50% occupancy can yield a 2× speedup, while going from 50% to 100% yields only 5–15%. After 50–75% occupancy, other factors (ILP, memory access patterns) begin to dominate.
The balance between SM resources and parallelism. Goal — hide latency
| Scenario | Optimal occupancy | Why |
|---|---|---|
| Memory-bound kernel | 60–80% | Needs many warps to hide DRAM latency |
| Compute-bound kernel | 25–50% | Tensor Cores keep the pipeline full without many warps |
| High-ILP kernel | 25% | Each warp provides 4 independent operations |
| Latency-bound kernel | 50–75% | Balance between warp count and per-thread resources |
Example usage of __launch_bounds__:
__global__ __launch_bounds__(256, 8)
void kernel_optimized(...) {
// compiler squeezes into ~32 registers per thread
// → 8 blocks × 256 threads = 2048 threads per SM
// → 100% occupancy
}Occupancy is a means to hide latency, not the goal itself. After 50–75%, other factors begin to dominate.
Registers per thread (limit 255 on Blackwell), shared memory per block, thread count. For diagnostics, Nsight Compute Occupancy and CUDA Occupancy API are used.
__launch_bounds__ and CUDA Occupancy API allow finding the optimal launch configuration. But watch out for register spilling to local memory.
«In SIMT execution, a warp of 32 threads executes instructions in lockstep. If threads within a warp take different if/else paths, the warp executes both branches sequentially. With a 50/50 split, warp efficiency drops to 50%».
The chapter describes the warp divergence problem and methods to minimize it. Key techniques:
if (warpId % 2 == 0) { /* Task A */ } else { /* Task B */ }. All threads in a warp take the same branch.__ballot_sync, __any_sync, __all_sync allow a warp to collectively decide which lanes need «special» work.float cond = x > 0 ? 1.0f : 0.0f; y = cond * f(x) + (1.0f - cond) * g(x);.Shuffle intrinsics (__shfl_sync, __shfl_down_sync) are also analyzed for intra-warp data exchange through registers without shared memory accesses. Example of warp reduce sum:
float warp_reduce_sum(float val) {
for (int offset = 16; offset > 0; offset >>= 1) {
val += __shfl_down_sync(0xFFFFFFFF, val, offset);
}
return val;
}Left: if/else diverges → warp runs both branches. Right: uniform path — all lanes work
Warp divergence reduces efficiency to 50% at a 50/50 split. Predication, warp-uniform branching and shuffle intrinsics are the main countermeasures.
Shuffle intrinsics allow exchanging data through registers without shared memory and without __syncthreads(). This is faster and creates no bank conflicts.
PyTorch operations like torch.where and torch.maximum compile into predication, providing warp-coherent execution without divergence.
«ILP is the number of independent instructions a single thread can issue in a row without waiting for previous ones to complete. High ILP lets a single warp keep execution units busy even when individual instructions have latency».
The chapter examines in detail ILP and its interaction with occupancy. Key points:
Example of breaking a dependency chain:
// No ILP: sequential dependency chain
float u = x * x;
float temp = u + 1.0f;
float v = y * y;
float sum = temp + v;
// With ILP: independent operations back-to-back
float u = x * x;
float v = y * y;
float temp = u + 1.0f;
float sum = temp + v;Independent instructions overlap: while load x waits for memory, y*y is already executing
| ILP | Threads per SM for ~100% issue | Occupancy |
|---|---|---|
| 1 (no ILP) | ~1536 (48 warps) | 75% |
| 2 | ~1024 (32 warps) | 50% |
| 3 | ~768 (24 warps) | 37.5% |
| 4 | ~512 (16 warps) | 25% |
4 independent warp schedulers. Each cycle can issue up to 2 instructions: one math + one memory
ILP allows a single warp to keep execution units busy. With ILP=4, 25% occupancy is enough to saturate the SM.
Unified CUDA cores execute either FP32 or INT32 in a single cycle. Mixed INT+FP streams no longer benefit from dual-issue.
Reorder independent operations back-to-back, unroll loops, increase the number of independent accumulators. Optimum — 2–4 operations per thread.
«The classic double-buffering pattern: while computation proceeds on the current tile, the next one is loaded in the background. On Hopper and Blackwell, TMA — Tensor Memory Accelerator — helps here: a hardware engine for bulk multidimensional copies from global to shared memory».
The chapter describes TMA and its interaction with the CUDA Pipeline API. Key points:
tcgen05.mma, tcgen05.ld and tcgen05.st provide high-performance data paths of the fifth generation on Blackwell. TMEM architecture and tcgen05 instructions are covered in detail in the Chapter 6 review.Example of double buffering with TMA and CUDA Pipeline API:
__shared__ float tile[2][TILE_SIZE];
tile[0][threadIdx.x] = global_data[blockIdx.x * TILE_SIZE + threadIdx.x];
__syncthreads();
for (int t = 0; t < numTiles - 1; ++t) {
tile[(t + 1) % 2][threadIdx.x] =
global_data[(blockIdx.x + t + 1) * TILE_SIZE + threadIdx.x];
process(tile[t % 2][threadIdx.x]);
__syncthreads();
}A more modern variant using the CUDA Pipeline API and TMA:
#include <cuda/pipeline>
#include <cooperative_groups.h>
namespace cg = cooperative_groups;
__global__ void kernelWithTMA(const float* __restrict__ global_ptr, int nTiles) {
__shared__ float tile0[TILE_SIZE];
__shared__ float tile1[TILE_SIZE];
float* tiles[2] = { tile0, tile1 };
__shared__ cuda::pipeline_shared_state<cuda::thread_scope_block, 2> state;
auto pipe = cuda::make_pipeline(cg::this_thread_block(), &state);
pipe.producer_acquire();
cuda::memcpy_async(tiles[0], global_ptr + 0 * TILE_SIZE,
cuda::aligned_size_t<32>(TILE_SIZE * sizeof(float)), pipe);
pipe.producer_commit();
for (int t = 1; t < nTiles; ++t) {
pipe.consumer_wait();
processTile(tiles[(t - 1) & 1]);
pipe.consumer_release();
pipe.producer_acquire();
cuda::memcpy_async(tiles[t & 1], global_ptr + t * TILE_SIZE,
cuda::aligned_size_t<32>(TILE_SIZE * sizeof(float)), pipe);
pipe.producer_commit();
}
pipe.consumer_wait();
processTile(tiles[(nTiles - 1) & 1]);
pipe.consumer_release();
}TMA is a hardware engine for bulk multidimensional copies. It frees the SM's LD/ST units for useful work. Double buffering with TMA hides DRAM latency.
Ping-pong between two buffers: while TMA loads the next tile, the SM computes on the previous one. This requires explicit management through the CUDA Pipeline API.
Using cuda::memcpy_async with cuda::pipeline on Hopper/Blackwell automatically engages TMA if alignment and sizes permit.
«For PyTorch users writing high-level code, you don't usually manage occupancy directly — PyTorch's CUDA kernels and libraries handle launch parameters under the hood. However, understanding occupancy can still be valuable. If you write custom CUDA kernels as PyTorch extensions, the same principles apply».
The chapter describes several optimizations available without writing CUDA:
model = torch.compile(model, mode="max-autotune") — generates fused GPU kernels, merging many small operations into one larger kernel. This increases work per kernel launch and improves occupancy.optimizer = torch.optim.AdamW(..., fused=True) — the fused implementation performs all optimizer steps in a single kernel, reducing launch overhead and memory traffic.y = torch.where(x > 0, x, torch.zeros_like(x)) — vectorized operations compile into SIMD-style predication, providing warp-coherent execution without divergence.It is separately emphasized that torch.compile does not guarantee optimal occupancy or ILP — this may require profiling and additional tuning. It is also noted that the PyTorch compiler does not always automatically use all capabilities — sometimes registering custom kernels through Triton is required.
For PyTorch users, many optimizations are available without writing CUDA. torch.compile generates fused kernels, fused optimizers reduce overhead, torch.where provides predication.
torch.compile does not guarantee optimal occupancy or ILP. Achieving maximum performance may require profiling and registering custom Triton kernels.
Start with torch.compile and fused optimizers. If performance is insufficient — profile and write custom kernels for hot spots.
«Prolonged near-peak compute utilization can sometimes invoke power management compute limiters on modern GPUs to keep them healthy. Be sure to take this into account when you're profiling, benchmarking, and tuning. The easiest way to see if you're being power-limited is to use nvidia-smi with a query for power draw, current clocks, and throttle reasons in real time».
The chapter describes power monitoring as part of the optimization process. The command:
nvidia-smi \
--query-gpu=power.draw,clocks.current.sm,clocks.current.memory,\
clocks_event_reasons.active \
--format=csv -l 1outputs a new line every second with current power draw, graphics and memory clocks, as well as throttle reasons. This allows precisely determining when the GPU hits its power cap and downclocks.
Note for legacy clusters: on NVIDIA drivers older than 535 the same field is named clocks_throttle_reasons.active. If the query returns a «Field not found» error, switch to the old field name; it is still supported on legacy drivers.
Programmatic access is available via NVML: nvmlDeviceGetPowerUsage() and nvmlDeviceGetEnforcedPowerLimit(). These APIs are ideal for custom scripts or integration with monitoring systems.
Under sustained near-peak load, the GPU may hit its power limit and reduce clocks. Power monitoring is part of the optimization process.
nvidia-smi with a query for power.draw, clocks.current.sm, clocks_event_reasons.active (on drivers <535 — clocks_throttle_reasons.active). Programmatic access via NVML.
If the GPU hits its power limit, consider tuning power limit or underclocking to improve efficiency per watt.
«In this chapter, you saw how to uncover and eliminate GPU kernel bottlenecks by moving work from slow global memory into faster on-chip resources and compute units. By following a cycle of profiling, diagnosing, optimizing, and reprofiling, you can transform kernels from underutilized or memory-bound into compute-saturated, high-throughput routines».
The chapter concludes with a checklist that systematizes the optimization process:
nsys, find hot kernels, drill down with ncu.__launch_bounds__ — if registers limit occupancy.torch.where.GPU kernel optimization is a systematic process, not a bag of tricks. The 12-point checklist covers the full cycle.
Profiling → diagnosing → optimizing → reprofiling. Each fix reveals the next bottleneck.
Follow the checklist sequentially. Don't optimize blindly. Verify accuracy after each change.
Fregly's book was written in 2025, when Blackwell was just entering the market. Below is a comparison of Chapter 8's forecasts with open data and clarification of figures.
«It is important to set up automated performance tests to detect regressions early in the development cycle».
Confirmed. Modern frameworks (PyTorch, TensorRT-LLM) include built-in profilers and CI pipelines with benchmarks. PyTorch performance HUD tracks regressions in nightly builds. MLPerf became the de facto standard for reproducible measurement. Tools like Nsight Systems and Nsight Compute are updated for each new GPU architecture.
«ILP and occupancy mutually compensate each other. With ILP=4, 25% occupancy is enough to saturate the SM; without ILP, 75% is required».
Confirmed. The table of ILP dependence on the required number of threads (1→1536, 2→1024, 3→768, 4→512) is consistent with NVIDIA data and architecture guides. On Blackwell with its 64 warps per SM and 256 KB register file, this dependence holds. The only clarification: on Blackwell, unified CUDA cores execute either FP32 or INT32 in one cycle, which reduces the efficiency of mixed INT+FP instruction streams compared to earlier architectures.
«TMA can perform 1D–5D bulk copies with arbitrary strides between global and shared memory without blocking the SM instruction pipeline».
Confirmed. TMA (Tensor Memory Accelerator) appeared in Hopper and was developed in Blackwell. The CUDA Pipeline API (cuda::memcpy_async, cuda::pipeline) automatically engages TMA on supported architectures. On Blackwell, TMA works together with TMEM to serve fifth-generation Tensor Cores (see the Chapter 6 review).
«torch.compile generates fused GPU kernels, merging many small operations into one larger kernel».
Confirmed. torch.compile with TorchInductor generates fused kernels, reducing launch overhead. However, as noted in the chapter, the compiler does not guarantee optimal occupancy or ILP. Achieving maximum performance may require registering custom Triton kernels or manual tuning.
«Prolonged near-peak compute utilization can sometimes invoke power management compute limiters on modern GPUs».
Partially confirmed. On B200, TDP is 1000 W; on B300 — 1400 W. In an NVL72 rack — 130 kW. For most inference workloads (which are usually memory-bound rather than compute-bound), power limiting is not the primary factor. However, for dense computations (e.g., training), power limiting can reduce clocks. On Blackwell Ultra and Rubin, cooling and power become more acute.
Chapter 8 describes fundamental optimization principles but leaves out several aspects that emerged in real-world Blackwell operation.
Impact of dual-die Blackwell architecture on optimization. The chapter does not mention that B200 consists of two dies connected by NV-HBI at 10 TB/s. For a programmer, this looks like one GPU, but at the memory performance level, anomalies related to die load imbalance may occur. In rare cases, diagnosing may require checking balance between dies.
Interaction with NVLink and NVSwitch. The chapter focuses on intra-kernel optimization, but for multi-GPU workloads, understanding topology is critical. Optimizing one kernel won't help if data is transferred over a slow interconnect: NVLink 5 on B200 delivers 1.8 TB/s per GPU, but beyond an NVL72 rack, throughput drops to 100–200 GB/s over InfiniBand. NVLink/NVSwitch topology and interconnect bandwidth are covered in detail in the Chapter 4 review.
MoE model specifics. For MoE models (like DeepSeek-V3), optimizing attention and MLP layers requires accounting for all-to-all communications between experts. The chapter does not cover this, but for real inference workloads, this is often the dominant factor.
FP4 limitations. The chapter mentions FP4 as a performance multiplier but does not examine limitations: FP4 requires calibration, not all models tolerate such precision, and for critical layers (early layers, attention softmax, layer norms) higher precision is required. Transformer Engine solves this through automatic per-layer precision selection but requires validation on the specific model.
Inference vs training specifics. The chapter describes general principles but does not distinguish optimization for training vs inference. For training, compute-bound kernels (GEMM) dominate; for inference — memory-bound (attention, decode). Accordingly, optimal strategies differ: for training, ILP and Tensor Cores are more important; for inference — memory bandwidth and low latency.
clocks_throttle_reasons.active).Chapter 8 of Fregly's book is a systematic breakdown of three complementary GPU kernel optimization directions: occupancy tuning, warp efficiency, and instruction-level parallelism. The key idea: optimization is not a bag of tricks but a systematic process of profiling, diagnosing, and reprofiling. For an engineer, this means understanding fundamental principles (memory hierarchy, ILP, occupancy, warp divergence) is more important than knowing specific numbers of the current GPU generation.
Forecast verification shows: the book proved accurate in its predictions. The techniques described in the chapter apply to Blackwell and, accounting for architectural changes (unified CUDA cores, TMA, TMEM), remain relevant. The main takeaway remains unchanged: investments in fundamental principles ensure adaptation to any generation of hardware.
Sources: Chris Fregly, «AI Systems Performance Engineering» (O'Reilly, 2026), Chapter 8 — quotes provided in limited volume for review and critical analysis purposes. NVIDIA Nsight Systems, Nsight Compute, CUDA Toolkit, Blackwell Architecture documentation; MLPerf Training/Inference v5.0; PyTorch Documentation; NVIDIA Technical Blog.
