Настоящий материал представляет собой обзор девятой главы книги 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 посвящена повышению эффективности CUDA-ядер и арифметической интенсивности. В ней рассматривается, как превратить ограниченное памятью ядро в ограниченное вычислениями за счёт повторного использования данных, слияния операций, понижения точности и задействования Tensor Cores. Ниже последовательно изложены ключевые положения главы, а также приведены уточнения по актуальным данным о платформе.
«Даже если удалось полностью скрыть задержки за счёт массивного параллелизма и высокого ILP, производительность ядра всё ещё может упираться в то, сколько полезной работы оно выполняет на каждый байт, прочитанный из памяти».
«Новые поколения GPU наращивают вычислительную мощность заметно быстрее, чем пропускную способность памяти. Этот разрыв делает повышение арифметической интенсивности всё более критичным».
В главе формулировка «арифметическая интенсивность» используется не как абстрактная метрика, а как инструмент диагностики. Если ядро находится ниже точки перегиба на Roofline, оно ограничено памятью, и любые оптимизации вычислений не дадут эффекта. Если выше — оно ограничено вычислениями, и дальнейший рост возможен только за счёт более эффективного использования ALU и Tensor Cores.
Практическое следствие, отмечаемое в главе: цель всегда одна — повышать арифметическую интенсивность там, где это возможно, выполняя больше вычислений на каждый байт, переданный в глобальную память и из неё. Добиться этого можно несколькими способами: тайлинг, кэширование, слияние ядер, понижение точности, увеличение размера батча и пересчёт вместо хранения.
Отдельно подчёркивается, что не каждая задача позволяет легко повысить арифметическую интенсивность — это ограничено характеристиками алгоритма. Поэлементные операции имеют интенсивность порядка 0,125–0,25 FLOP/байт и катастрофически ограничены памятью по отдельности. Матричные умножения, напротив, имеют интенсивность, которая растёт с размером матриц и легко превышает точку перегиба.
«Арифметическая интенсивность, также называемая операционной интенсивностью, измеряет, сколько операций с плавающей запятой выполняется на каждый байт данных, переданных из памяти. Измеряется в FLOPS на байт».
«Арифметическая интенсивность — ключевая метрика в модели производительности Roofline. Эта модель визуально отображает производительность ядра (FLOPs/sec) в зависимости от арифметической интенсивности (FLOPs/byte). На графике видны два аппаратных потолка: наклонная линия из начала координат — ограничение по пропускной способности памяти, и горизонтальная линия — ограничение по вычислительной мощности».
В главе модель Roofline описывается как инструмент, позволяющий сразу понять, является ли ядро ограниченным памятью или вычислениями. Ядро с интенсивностью ниже точки перегиба находится в режиме ограничения памятью — его скорость упирается в пропускную способность памяти, потому что GPU большую часть времени ждёт данные, а не считает. Ядро с интенсивностью выше точки перегиба — ограничено вычислениями, оно использует ALU и Tensor Cores близко к их пиковым возможностям.
Приводится конкретный пример: на NVIDIA H100 пиковая производительность BF16 составляет ~990 TFLOPS, а пропускная способность HBM — ~3,35 ТБ/с. Точка перегиба — отношение этих величин — равна примерно 296 FLOP/байт. На Blackwell B200 с поддержкой NVFP4 пиковая производительность достигает 9 петаFLOPS (плотных) в FP4, а пропускная способность HBM3e — около 8 ТБ/с, что сдвигает точку перегиба ещё выше.
Практическое следствие: на практике графики Roofline можно строить с помощью Nsight Compute, который включает режим анализа Roofline. С его помощью можно проверить, было ли ядро изначально ограничено памятью или вычислениями, а затем отслеживать улучшения по мере оптимизаций.
GFLOP/s в зависимости от арифметической интенсивности (FLOPs/byte)
Арифметическая интенсивность — отношение FLOPs к байтам. Модель Roofline визуализирует два потолка: пропускную способность памяти и вычислительную мощность. Точка перегиба разделяет режимы ограничения памятью и вычислениями.
H100: ~990 TFLOPS BF16, ~3,35 ТБ/с HBM, точка перегиба ~296 FLOP/байт. B200: 9 PFLOPS (плотных) в FP4, 18 PFLOPS с разреженностью, ~8 ТБ/с HBM3e. Точка перегиба сдвигается вверх с каждым поколением.
Nsight Compute позволяет строить Roofline-графики и отслеживать, как оптимизации сдвигают ядро вправо — из режима ограничения памятью в режим ограничения вычислениями.
«Тайлинг и повторное использование данных — эффективный способ повысить арифметическую интенсивность. Загрузка небольшой подматрицы (тайла) из A и B в shared memory позволяет каждому байту, прочитанному из глобальной памяти, участвовать во множестве операций умножения-сложения на скорости SRAM».
«Всякий раз, когда код перестраивается так, что каждый элемент загружается один раз и используется десятки или сотни раз, соотношение FLOPs на байт умножается на коэффициент повторного использования».
В главе тайлинг описывается как базовый приём повышения арифметической интенсивности. При матричном умножении C = A × B с тайлом T × T глобальные чтения сокращаются с O(N) на результат до O(N/T) на тайл плюс O(T) повторного использования внутри shared memory. Для тайла 128 × 128 с размером K-измерения 64 арифметическая интенсивность составляет примерно 64 FLOP/байт — на порядки выше, чем у наивной поэлементной операции.
Помимо простого тайлинга в shared memory, в главе рассматривается многоуровневый тайлинг. При таком подходе после размещения тайла в shared memory каждый поток загружает микротайлы в регистры, используя векторные типы вроде float4 и <half2>. Такая иерархия повторного использования сокращает рабочий набор на каждом уровне и минимизирует трафик к внешней памяти.
Отдельно оговаривается, что необходимо обеспечивать коалесцирование глобальных чтений при заполнении shared memory и выравнивать/свиззлить данные в shared memory, чтобы избежать конфликтов банков. Конфликты банков shared memory могут снизить эффективную пропускную способность в 2–32 раза, поэтому свиззлинг — стандартный приём для доступа без конфликтов банков.
Практическое следствие: на современных GPU эти шаги внутреннего тайлинга часто покрываются с помощью API фрагментов MMA. Аппаратное обеспечение перемещает данные между shared memory и Tensor Memory неявно, используя инструкции Tensor Core.
DRAM → SMEM → Регистры: данные движутся вниз с повторным использованием
Глобальная память → разделяемая память напрямую, опционально минуя L1
Тайлинг повышает арифметическую интенсивность за счёт повторного использования данных в разделяемой памяти. Многоуровневый тайлинг добавляет повторное использование в регистрах.
Конфликты банков разделяемой памяти могут снизить эффективную пропускную способность в 2–32 раза. Свиззлинг или паддинг обязательны.
Программная предвыборка (двойная буферизация) перекрывает DRAM → SMEM с вычислениями. cudaMemPrefetchAsync используется как подсказка, но не как гарантия.
«На современных GPU можно расширить идею тайлинга с повторным использованием, используя кластеры блоков потоков из Cooperative Groups. Они позволяют нескольким блокам потоков совместно использовать данные через распределённую shared memory (DSMEM)».
«Кластер из четырёх блоков потоков может совместно загрузить один тайл, используя функцию multicast в Tensor Memory Accelerator (TMA). Каждый тайл разделяется между четырьмя блоками так, что трафик глобальной памяти для этого тайла амортизируется по всему кластеру».
В главе кластеры блоков потоков описываются как способ прямого повышения арифметической интенсивности. Тайлы загружаются только один раз и переиспользуются всеми блоками кластера. Это может сократить трафик глобальной памяти до 4× в кластере 2×2, уменьшая знаменатель — количество байт, перемещаемых из глобальной памяти.
Blackwell поддерживает кластеры блоков потоков размером до 16 блоков, если включить непортируемый размер кластера сверх стандартного лимита в 8 CTA. Для этого нужно установить атрибут cudaFuncAttributeNonPortableClusterSizeAllowed для ядра. Большие кластеры могут повысить повторное использование, но способны снизить occupancy, поэтому профилирование перед включением 16 обязательно.
Отдельно оговаривается, что DSMEM позволяет CTA внутри кластера читать общую shared memory, но такие загрузки пакетизируются и их пропускная способность заметно ниже локального SMEM. Для реальной меж-CTA пропускной способности следует использовать cp.async.bulk.
Практическое следствие: кластеры ко-планируются в GPC. Если размер кластера не делит число SM на GPC, часть SM простаивает. Blackwell вводит preferred/fallback размеры кластеров, чтобы смягчить эту квантованность.
Блоки потоков (CTA) внутри одного кластера совместно используют распределённую on-chip SRAM
Одна DMA-передача вещает один и тот же тайл в shared memory четырёх CTA
Кластеры блоков потоков + DSMEM + TMA multicast сокращают трафик глобальной памяти до 4× в кластере 2×2. Это прямой способ повышения арифметической интенсивности.
DSMEM-загрузки пакетизируются и медленнее локального SMEM. Кластеры ко-планируются в GPC; размер кластера влияет на occupancy.
Для реальной меж-CTA пропускной способности используется cp.async.bulk. Blackwell поддерживает кластеры до 16 CTA при включении непортируемого размера.
«Ещё один способ повысить арифметическую интенсивность — объединить несколько операций или итераций цикла в одну операцию. Сливая несколько ядер вместе, данные, загруженные из памяти, можно использовать для нескольких вычислений и итераций, прежде чем записать обратно».
«Поэлементные операции применяют простое вычисление независимо к каждому элементу тензора. Слияние таких операций устраняет лишний трафик памяти, удерживая промежуточные значения на кристалле».
В главе слияние ядер описывается как способ избежать записи промежуточных результатов в глобальную память. Приводятся два типа слияния: вертикальное — объединение последовательных операций над одними и теми же данными (например, sin и sqrt в одном ядре), и горизонтальное — объединение параллельных операций по данным (например, bias-add и активация после матричного умножения в эпилоге GEMM).
Приводится конкретный пример с L2-нормализацией. Наивная реализация состоит из трёх отдельных ядер: возведение в квадрат, редукция, деление. Каждое ядро выполняет примерно 1 FLOP на элемент и требует 12 байт трафика на элемент. Итого арифметическая интенсивность 3/36 ≈ 0,083 FLOP/байт. В слитом ядре каждый поток обходит свой срез дважды: один раз для накопления суммы квадратов и один раз для записи нормализованных выходов. Глобальный трафик — около 12 байт на элемент, на каждый элемент приходится примерно 3 FLOP. Арифметическая интенсивность слитого ядра: 3 / 12 ≈ 0,25 FLOP/байт — в 3 раза выше.
Отдельно оговаривается, что слишком агрессивное слияние может увеличить давление на регистры и снизить occupancy. Если использование регистров становится чрезмерным и начинается spilling в локальную память, выгода от слияния может быть перекрыта дополнительным трафиком памяти. Всегда профилирование слитых ядер с помощью Nsight Compute обязательно.
Практическое следствие: современные фреймворки глубокого обучения могут автоматически сливать и разворачивать операции через JIT-компиляторы и оптимизаторы графов. Например, torch.compile в PyTorch, а конкретно TorchInductor, может автоматически сливать последовательности поэлементных операций.
Наивная реализация (два ядра) и слитая реализация (одно ядро)
Слияние ядер — вертикальное и горизонтальное — повышает арифметическую интенсивность за счёт удержания промежуточных результатов в регистрах.
Чрезмерное слияние увеличивает давление на регистры и может вызвать spilling в локальную память. Всегда требуется профилирование.
torch.compile и TorchInductor автоматически сливают последовательности поэлементных операций. Для нестандартных паттернов может потребоваться ручное слияние.
«На современных GPU разреженность 2:4 ускоряется аппаратно с помощью Sparse Tensor Cores и cuSPARSELt. Схема 2:4 означает, что ровно два из каждых четырёх последовательных весов ненулевые».
«Обрезая половину весов в шаблоне 2:4, каждая загрузка памяти теперь доставляет вдвое больше ненулевых значений, которые реально участвуют в умножении. Это немедленно сокращает глобальные чтения памяти и повышает эффективную арифметическую интенсивность почти в 2 раза».
В главе структурированная разреженность описывается как способ повышения эффективной арифметической интенсивности. Поскольку трафик памяти для весов сокращается вдвое, а вычисления выполняются только над ненулевыми элементами, Sparse Tensor Cores могут приближаться к 2× ускорению по сравнению с плотными аналогами.
Отдельно оговаривается, что реальное ускорение сильно зависит от структуры матриц, плотности паттернов и overhead на хранение индексов. Для малых матриц разрежённость может даже замедлить выполнение. Ускорение предполагает, что точность модели сохраняется, — на практике важно дообучить или тщательно откалибровать модель после прунинга.
Практическое следствие: в PyTorch используется to_sparse_semi_structured() для конвертации обученных плотных модулей в формат 2:4 перед развёртыванием разреженных GEMM на Sparse Tensor Cores. Структурированная разреженность обычно применяется к задачам инференса; во время обучения градиенты не получают выгоды от разрежённости 2:4.
Перед применением разреженности важно сначала реализовать фундаментальные оптимизации: коалесцировать все глобальные загрузки, переиспользовать данные через тайлинг и объединять поэлементные операции.
Ровно 2 из каждых 4 последовательных весов ненулевые → 2× пропускная способность Tensor Core
Структурированная разреженность 2:4 повышает эффективную арифметическую интенсивность почти в 2 раза за счёт сокращения трафика весов.
Реальное ускорение зависит от структуры матриц и overhead на индексы. Для малых матриц возможно замедление. Применяется в основном на инференсе.
В PyTorch используется to_sparse_semi_structured(). Перед применением необходимо реализовать базовые оптимизации: коалесцирование, тайлинг, слияние.
«Вместо хранения или загрузки предварительно вычисленных значений рассмотрите возможность их пересчёта по требованию, если узким местом является пропускная способность памяти. Например, повторное вычисление x*x в регистрах часто быстрее, чем загрузка ранее вычисленного x² из глобальной памяти».
«Многие движки инференса LLM используют этот приём для экономии памяти. Вместо хранения больших тензоров активаций в HBM и их чтения обратно позже они могут пересчитывать определённые слои и активации на лету».
В главе пересчёт описывается как способ снижения трафика памяти за счёт увеличения вычислений. Классический пример — FlashAttention, который пересчитывает статистики softmax в прямом проходе вместо хранения полной матрицы внимания N × N в HBM. Это обменивает несколько дополнительных FLOPS на драматическое сокращение доступа к глобальной памяти, превращая O(N²) доступов к HBM в O(N² / T), где T — размер тайла.
Отдельно подчёркивается, что решение о том, пересчитывать или хранить, зависит от соотношения вычислительной мощности и пропускной способности памяти на конкретном GPU: если ядро ограничено памятью, пересчёт почти всегда выгоден; если ограничено вычислениями — хранение может быть предпочтительнее.
Практическое следствие: пересчёт улучшает эффективные FLOPS/байт и позволяет вместить большие модели в меньший объём памяти. Activation checkpointing — частный случай этого подхода, применяемый в обучении.
Пересчёт вместо хранения — обмен вычислений на пропускную способность памяти. FlashAttention и activation checkpointing — практические примеры.
Решение зависит от соотношения compute-to-memory-bandwidth на конкретном GPU. Для ядер, ограниченных памятью, пересчёт почти всегда выгоден.
«В PyTorch многие из описанных идей применяются автоматически. Компилятор PyTorch (torch.compile) может автоматически сливать цепочки поэлементных операций и даже некоторые редукции».
«Вложенные тензоры, или ragged tensors, позволяют представлять батчи входных данных переменной длины без паддинга. Каждый вложенный тензор упаковывает последовательности переменной длины в единый эффективный буфер. NestedTensor предоставляет обычный интерфейс тензора, но устраняет ненужный zero-padding».
В главе PyTorch описывается как среда, в которой многие оптимизации арифметической интенсивности применяются автоматически. Поскольку под капотом используются cuDNN и cuBLAS, PyTorch выполняет тайлинг и использует shared memory при матричных операциях. scaled_dot_product_attention (SDPA) может направлять вычисления к бэкендам FlashAttention, memory-efficient или cuDNN в зависимости от форм тензоров и типов данных.
Отдельно рассматриваются вложенные тензоры как способ представления батчей переменной длины без паддинга. При использовании вложенных тензоров операции attention и batch-matrix multiplies извлекают из памяти только необходимые данные, что приближает ядро к режиму ограничения вычислениями на модели Roofline.
Практическое следствие: хотя PyTorch может распознавать и компилировать большинство операций, некоторые нестандартные операции или пользовательские CUDA-операции могут не сливаться. В таких случаях может потребоваться ручная оптимизация.
PyTorch автоматически применяет тайлинг, слияние и выбор оптимального бэкенда attention. Вложенные тензоры устраняют паддинг, повышая эффективность загрузки.
Для нестандартных операций может потребоваться ручное слияние или написание кастомных ядер. torch.compile обрабатывает большинство стандартных паттернов.
«Современные NVIDIA GPU реализуют вычисления пониженной точности, такие как TF32, FP16, FP8, FP4 и INT8, в Tensor Cores. Каждый SM в Blackwell имеет 256 КБ внутрикристальной памяти TMEM, выделенной для данных Tensor Core».
«Ключевое архитектурное отличие Blackwell от Hopper: на Hopper результаты MMA были жёстко привязаны к регистровому файлу варп-группы, а на Blackwell аккумуляторы живут в TMEM, что декопулирует владение результатами от отдельных потоков».
В главе Tensor Cores описываются как ключевой инструмент повышения арифметической интенсивности. Понижение точности увеличивает количество операций на байт: FP8 даёт 2×, FP4 — до 4×. Память сокращается пропорционально: FP4 вдвое уменьшает объём на параметр по сравнению с FP8.
Архитектурное отличие Blackwell от Hopper описывается как фундаментальное. Инструкции tcgen05 являются CTA-scoped — один поток выдаёт инструкцию от имени всего CTA, в отличие от warpgroup-scoped wgmma на Hopper. Это даёт большую гибкость в планировании: одни потоки могут оркестрировать вычисления, другие — перемещать данные.
Практическое следствие: при правильном использовании эти возможности могут превратить когда-то ограниченное памятью ядро с интенсивными тензорными операциями в полностью ограниченное вычислениями, повысив арифметическую интенсивность в 2, 4 или даже 8 раз.
Входные данные пониженной точности, накопление в повышенной точности внутри Tensor Cores
Tensor Cores и TMEM — основа повышения арифметической интенсивности. Blackwell использует CTA-scoped инструкции tcgen05 с аккумуляторами в TMEM.
Понижение точности увеличивает количество операций на байт: FP8 — 2×, FP4 — до 4×. TMEM снижает давление на регистровый файл.
«В основе высокопроизводительных тензорных вычислений лежит TMEM — буфер SRAM на 256 КБ на каждый SM. Программисты не выделяют и не управляют TMEM явно — TMEM обрабатывается аппаратным обеспечением или библиотеками при использовании операций Tensor Core».
«Для выполнения передачи данных из глобальной памяти в shared memory используйте TMA или cuda::memcpy_async с конвейером из CUDA Pipeline API».
В главе TMEM описывается как выделенный буфер SRAM на 256 КБ на SM, который Transformer Engine и Tensor Cores используют для хранения результатов вместо регистров. Аппаратное обеспечение управляет перемещением данных между shared memory и TMEM неявно, используя инструкции Tensor Core.
Приводится пример двухступенчатого конвейера с двойной буферизацией, в котором cuda::memcpy_async загружает следующий тайл в shared memory параллельно с вычислениями Tensor Cores. Отдельно оговаривается, что на Blackwell UMMA имеет ограничение: операнд A может поступать из TMEM или SMEM, операнд B — только из SMEM. Комбинация A из TMEM и B из TMEM не поддерживается.
Практическое следствие: для передачи данных из глобальной памяти в shared memory используется TMA или cuda::memcpy_async. Это перекрывает перемещение данных с вычислениями Tensor Core и помогает скрыть задержку памяти.
Аккумуляторы живут в TMEM (256 КБ/SM), а не в регистровом файле
TMEM — выделенный буфер 256 КБ/SM для аккумуляторов Tensor Core. TMA и cuda::memcpy_async перекрывают передачу данных с вычислениями.
UMMA на Blackwell: A может быть из TMEM или SMEM, B — только из SMEM. Комбинация A из TMEM и B из TMEM не поддерживается.
Для передачи данных из HBM в shared memory используется TMA или cuda::memcpy_async. Двойная буферизация скрывает задержку памяти.
«Хотя Tensor Cores изначально проектировались для FP16, они также поддерживают TF32, который находится между FP32 и FP16. TF32 использует 8-битный порядок, как FP32, и 10-битную мантиссу, как FP16».
«Помимо TF32, автоматическая смешанная точность (AMP) в PyTorch может выбирать оптимальную точность (FP16 или BF16) для каждой операции и накапливать результаты в FP32 для стабильности».
В главе TF32 описывается как способ получить ускорение без изменения кода. В PyTorch достаточно установить torch.set_float32_matmul_precision('high'), и высокоуровневые операции автоматически выполняются как ядра TF32 Tensor Core.
Автоматическая смешанная точность (AMP) в PyTorch может выбирать оптимальную точность для каждой операции и накапливать результаты в FP32 для стабильности. BF16 помогает избежать проблем с переполнением, присущих FP16. Под капотом TorchInductor автоматически сливает преобразования точности, чтобы гарантировать: большие GEMM выполняются на Tensor Cores в FP16 или TF32, накопление остаётся в FP32, небольшие чувствительные ядра выполняются в FP32.
Практическое следствие: при использовании структурированной разреженности или экстремально низкой точности (FP8/FP4) необходимо убедиться, что размер батча или гранулярность тайла достаточно велики, чтобы TMEM и Tensor Cores оставались полностью загруженными. Малые батчи несут накладные расходы.
TF32 включается одной строкой. AMP автоматически выбирает точность на операцию и накапливает в FP32.
Для FP8/FP4 необходим достаточный размер батча. Малые батчи не амортизируют накладные расходы на преобразования форматов и обработку индексов.
«BF16/FP16 поддерживается уже несколько поколений GPU, но Tensor Cores на современных GPU часто могут поддерживать более 90% пиковой пропускной способности BF16/FP16, что примерно в 4 раза превышает пиковую пропускную способность FP32».
«Чтобы ещё больше повысить пропускную способность, можно использовать FP8. Уменьшая 16-битные веса на 50% до 8 бит, вы вдвое сокращаете трафик памяти — и вдвое увеличиваете количество весов, загружаемых за одну транзакцию HBM».
В главе понижение точности рассматривается как способ повышения арифметической интенсивности. Каждое снижение точности удваивает или учетверяет количество операций на байт. FP8 с накоплением FP32 или TF32 достигает в 2–3 раза больших TFLOPS, чем BF16/FP16.
Для решения проблем точности при очень низкой точности Transformer Engine поддерживает FP8, а также 4-битный формат NVFP4 с микро-масштабированием. NVFP4 применяет двухуровневое масштабирование, комбинируя масштабирование на уровне микроблока (каждые 32 элемента вдоль K-измерения имеют общий масштабный коэффициент в формате E8M0) и масштаб более высокого уровня.
Blackwell B200 обеспечивает 9 петаFLOPS (плотных) в NVFP4 (18 петаFLOPS с разреженностью), тогда как пиковая производительность FP32 составляет около 80 тераFLOPS (плотных) — то есть примерно в 112 раз больше FP32. Blackwell B300 (Ultra) — 15 петаFLOPS (плотных) в NVFP4.
Практическое следствие: когда TMEM/TMA перекрывают память и вычисления, эти форматы низкой точности превращают ранее ограниченные памятью ядра в полностью ограниченные вычислениями. Разные слои трансформера имеют разную чувствительность к понижению точности; рекомендуется калибровка с перебором точности по слоям.
Понижение точности приводит к кратному увеличению пропускной способности
FP8 даёт 2×, FP4 — до 4× пропускной способности. NVFP4 с микро-масштабированием сохраняет точность за счёт двухуровневого масштабирования.
Разные слои имеют разную чувствительность к понижению точности. Рекомендуется калибровка с перебором точности по слоям и измерением метрик качества.
«Задачи инференса LLM обычно могут допускать квантование до пониженной точности INT8, поддерживаемое современными GPU с помощью инструкций DP4A на обычных CUDA-ядрах и инструкций целочисленного матричного умножения-накопления на Tensor Cores».
«Поскольку трафик весов для INT8 составляет всего один байт на элемент вместо четырёх байт для FP32, трафик памяти для весов сокращается на 75%».
В главе INT8 описывается как способ повышения эффективности инференса. На уровне инструкций DP4A выполняет четыре операции умножения-накопления INT8 (MAC) за инструкцию по сравнению с одной операцией FMA FP32 за инструкцию. Каждый GPU может обрабатывать примерно в 4 раза больше данных в секунду из памяти при использовании весов INT8.
На момент написания поддержка INT8-квантования в PyTorch обеспечивается через TorchAO и вендорные бэкенды. Квантованные модули работают с использованием выделенных ядер INT8. Использование cuBLASLt или CUTLASS для низкоуровневого INT8 GEMM обеспечивает задействование Tensor Core.
Практическое следствие: INT8 — эффективный способ повышения пропускной способности инференса, особенно для моделей, где точность FP32 избыточна. Комбинация INT8 с разреженностью 2:4 может дать дополнительное ускорение.
INT8 сокращает трафик весов на 75% и позволяет обрабатывать в 4 раза больше данных в секунду. DP4A выполняет 4 MAC за инструкцию.
Поддержка INT8 в PyTorch обеспечивается через TorchAO и вендорные бэкенды. cuBLASLt и CUTLASS для INT8 GEMM задействуют Tensor Core.
«Современные NVIDIA GPU включают Transformer Engine, который сочетает аппаратную поддержку Tensor Core для форматов пониженной точности с программной средой выполнения для масштабирования и приведения типов».
«TMEM — это буфер SRAM на 256 КБ на SM, который Transformer Engine и Tensor Cores используют для хранения результатов вместо регистров. На практике явное выделение TMEM не требуется — всё обрабатывается аппаратным обеспечением».
В главе рабочий процесс Transformer Engine описывается как последовательность шагов: ядро выдаёт вызов MMA или запускает GEMM из CUTLASS; прошивка TE организует копирование TMA весов и активаций из HBM в shared memory; инструкции Tensor Core неявно перемещают операнды между SMEM и TMEM; частичные суммы накапливаются в TMEM с более высокой точностью; результаты записываются из TMEM в shared или глобальную память.
В идеале веса находятся в FP8 или FP4, а активации приводятся к FP8/FP4, когда это возможно. Операции Tensor Core MMA выполняются с низкой точностью, например FP8 × FP8 с накоплением более высокой точности или FP16 × FP16 с накоплением FP32.
Отдельно оговаривается ограничение UMMA на Blackwell: операнд A может поступать из TMEM или SMEM, операнд B — только из SMEM. Комбинация A из TMEM и B из TMEM не поддерживается.
Практическое следствие: если создаётся собственный цикл по тайлам, можно перекрыть перемещение данных с вычислениями Tensor Core с помощью cuda::memcpy_async и CUDA Pipeline API. Пока Tensor Cores обрабатывают текущий тайл, cuda::memcpy_async потоково передаёт следующий тайл из HBM в shared memory.
Transformer Engine автоматически управляет точностью и масштабированием. TMEM обрабатывается аппаратно без явного выделения.
UMMA на Blackwell: A — из TMEM или SMEM, B — только из SMEM. Комбинация A из TMEM и B из TMEM не поддерживается.
cuda::memcpy_async и CUDA Pipeline API перекрывают перемещение данных с вычислениями Tensor Core, скрывая задержку памяти.
«Один из самых простых способов задействовать описанные оптимизации — использовать библиотеку NVIDIA CUTLASS. С CUTLASS вы пишете один шаблонный вызов, и он автоматически применяет множество продвинутых оптимизаций: тайлинг в shared memory, асинхронные передачи памяти, двойную буферизацию с помощью буфера TMEM, warp specialization, кластеры блоков потоков и TMA multicast».
В главе CUTLASS описывается как способ получить производительность ручного MMA-ядра без недель низкоуровневой настройки. Приводится сравнение: ручное MMA-ядро и CUTLASS GEMM показывают ~98% утилизации Tensor Core, при этом CUTLASS требует всего несколько строк шаблонного кода.
Шаблоны CUTLASS поддерживают типы операндов FP4, FP8, FP16, TF32 и INT8. Для сборки под конкретную архитектуру необходимо указать цель с суффиксом a (например, sm100a для Blackwell), поскольку инструкции tcgen05 относятся к категории architecture-accelerated features без гарантий прямой совместимости. Код, скомпилированный для SM100, несовместим с GeForce RTX 50 (SM120).
Практическое следствие: оптимизированные библиотеки, такие как cuBLAS, построены на CUTLASS. Высокоуровневые библиотеки, такие как PyTorch, вызывают эти оптимизированные библиотеки для многих ядер. По мере того как всё больше операторов в PyTorch используют CUTLASS под капотом, можно задействовать те же оптимизации без написания CUDA C++.
| Метрика | Ручное MMA-ядро | CUTLASS GEMM |
|---|---|---|
| Утилизация Tensor Core | ~98% | ~98% |
| Регистров на поток | ~52 | ~60 (немного выше) |
| Разделяемая память на блок (CTA) | ~2 КБ | ~4 КБ |
| Трудоёмкость разработки | Высокая | Низкая (конфигурация шаблона) |
CUTLASS даёт ~98% утилизации Tensor Core при минимальных усилиях. Автоматически применяет тайлинг, асинхронные копии, двойную буферизацию, warp specialization и TMA multicast.
Инструкции tcgen05 требуют сборки под конкретную архитектуру с суффиксом a (sm100a для Blackwell). Код для SM100 несовместим с SM120.
cuBLAS и PyTorch используют CUTLASS под капотом. Для стандартных GEMM и fused attention рекомендуется использовать CUTLASS вместо ручного написания ядер.
«Для тех, кто готов выйти за пределы C++ и погрузиться в низкоуровневые микрооптимизации, CUDA позволяет использовать встраиваемый код Parallel Thread Execution (PTX) и SASS для извлечения последних крупиц производительности».
«PTX — это низкоуровневая виртуальная машина параллельного выполнения потоков. SASS — это низкоуровневый язык ассемблера, который фактически выполняется нативно на оборудовании NVIDIA GPU».
В главе встраиваемый PTX описывается как инструмент последней инстанции. Типичные сценарии: использование инструкции GPU, для которой нет прямого интринсика CUDA; применение модификаторов загрузки памяти (подсказок кэша) к конкретным доступам; вставка барьеров памяти или синхронизации в точных точках; ручная перестановка инструкций для избежания простоев конвейера; чтение специальных регистров.
Приводятся примеры: предвыборка в L2 через cp.async.bulk.prefetch.L2.global; управление поведением кэша через ld.global.cg.f32; доступ к специальным регистрам через mov.u32 %0, %smid. Отмечается, что некоторые разработчики использовали %smid в постоянных ядрах, чтобы только один блок на SM выполнял определённую работу.
Практическое следствие: встраиваемый PTX/SASS обеспечивает максимальный контроль ценой сложности. Рекомендуется использовать в последнюю очередь, когда нужна аппаратная функция, недоступная в CUDA C++, или когда точно определён небольшой фрагмент кода, где планирование компилятора можно улучшить.
| Версия | Простои warp (память) | Issue IPC | Время ядра | Ускорение |
|---|---|---|---|---|
| Оптимизированный C++ (расписание компилятора) | 35% циклов | 1.5 | 1.07 мс | 1.0× базовое |
| Настроенный вручную PTX (ручное планирование и подсказки) | 20% циклов | 1.6 | 1.00 мс | 1.07× |
Встраиваемый PTX/SASS — инструмент последней инстанции. Даёт 5–10% ускорения на уже оптимизированных ядрах, но требует высокой квалификации.
SASS меняется между поколениями GPU. PTX более стабилен, но не гарантирует неизменность производительности. Код требует обновления при смене архитектуры.
Использовать только после исчерпания всех высокоуровневых оптимизаций. Подходит для критических циклов, выполняемых миллионы раз.
«Хорошо известный пример пользовательского PTX — из библиотеки DeepEP для экспертного параллелизма DeepSeek. Эта библиотека использовала специальную инструкцию PTX ld.global.nc.l1::no_allocate.l2::256b для оптимизации доступа к глобальной памяти».
«Обходя L1, загрузки могут потоково передавать большие блоки данных непосредственно в L2 без вытеснения часто используемых данных, находящихся в L1».
В главе приводится разбор инструкции ld.global.nc.l1::no_allocate.l2::256b. Компоненты: ld.global — загрузка из глобальной памяти; .nc — некогерентная загрузка, использующая текстурный кэш; .l1::no_allocate — не выделять данные в кэше L1; .l2::256b — извлекать по 256 байт за раз в L2.
Отмечается, что эта инструкция не является частью официальной спецификации PTX ISA (хотя позже было обнаружено, что модификаторы документированы). Инженеры DeepSeek обнаружили её вне документации для точной настройки поведения кэша на их варианте Hopper GPU H800, ограниченном экспортными правилами США.
DeepEP также выделяет 20 SM из 132 на H800 исключительно для сетевого общения, чтобы L1-кэш этих SM оставался тёплым для сетевого кода. Это снижает вычислительную производительность на ~15%, но значительно повышает пропускную способность связи.
Практическое следствие: использование PTX таким образом несёт риски: инструкция не гарантирует стабильность между поколениями GPU. Настройка DeepEP включает флаг времени сборки DISABLE_AGGRESSIVE_PTX_INSTRS=1, чтобы отключить эти агрессивные инструкции при возникновении проблем совместимости.
Инструкция ld.global.nc.l1::no_allocate.l2::256b не гарантирует стабильность между поколениями GPU. Код, полагающийся на неё, может сломаться или дать некорректные результаты на будущих архитектурах. DeepSeek включает флаг DISABLE_AGGRESSIVE_PTX_INSTRS=1 для отключения.
DeepSeek использовала недокументированную PTX-инструкцию для оптимизации доступа к памяти. Обход L1 и загрузка в L2 по 256 байт сохраняет горячие данные в L1.
Инструкция не гарантирует стабильность между поколениями GPU. Требуется флаг DISABLE_AGGRESSIVE_PTX_INSTRS=1 для отключения при проблемах совместимости.
Использовать с осторожностью, тестировать на каждой новой архитектуре. Предпочитать документированные альтернативы (Triton eviction_policy, cache_modifier).
Девятая глава книги Fregly содержит прогнозы и рекомендации, которые можно сопоставить с текущим состоянием технологий. Ниже представлено сопоставление с открытыми данными.
«Каждое снижение точности удваивает или учетверяет количество операций на байт и, следовательно, повышает арифметическую интенсивность».
Подтверждён. NVIDIA Blackwell B200 обеспечивает 9 петаFLOPS (плотных) в NVFP4 (18 петаFLOPS с разреженностью) — примерно в 112 раз больше, чем FP32 (80 TFLOPS). Transformer Engine автоматически управляет точностью per-layer. PyTorch, TensorRT-LLM и vLLM используют FP8 и FP4 как стандарт для инференса.
«CUTLASS даёт ~98% утилизации Tensor Core при минимальных усилиях».
Подтверждён. CUTLASS используется cuBLAS под капотом. PyTorch, TensorRT-LLM и другие фреймворки используют CUTLASS для критических операций. Появились Python-обёртки (cuTile, CuTe DSL), снижающие порог входа.
«Структурированная разреженность обычно применяется к задачам инференса. Во время обучения градиенты не получают выгоды от разрежённости 2:4».
Подтверждён. Sparse Tensor Cores поддерживаются в cuSPARSELt и PyTorch через to_sparse_semi_structured(). NVIDIA продолжает развивать поддержку 2:4 в новых поколениях GPU. Для обучения разрежённость по-прежнему не даёт выгоды.
«FP4 (NVFP4) Tensor Cores удваивают пропускную способность FP8 и используются для инференса с микро-масштабированием для поддержания точности».
Частично подтверждён. FP4 и NVFP4 поддерживаются аппаратно на Blackwell, но требуют калибровки. Для многих моделей FP8 остаётся более безопасным выбором. FP4 активнее применяется для KV-кэша и весов, где tolerance к точности выше.
«TMEM — это выделенный буфер SRAM на 256 КБ на SM, который Transformer Engine и Tensor Cores используют для хранения результатов вместо регистров».
Подтверждён. TMEM на Blackwell (256 КБ/SM) позволяет аккумуляторам Tensor Core жить вне регистрового файла, снижая давление на регистры и уменьшая задержку промаха кэша примерно на 58% по сравнению с Hopper.
Глава описывает техники повышения арифметической интенсивности, но оставляет за скобками несколько аспектов, которые проявились в реальной эксплуатации.
Глава фокусируется на производительности, но не рассматривает энергоэффективность. FP4 и FP8 не только повышают пропускную способность, но и снижают энергопотребление на операцию. Для крупномасштабных развёртываний это критично: GB200 NVL72 потребляет ~130 кВт, и каждое удвоение эффективности напрямую снижает стоимость эксплуатации.
Глава упоминает, что FP4 требует калибровки, но не описывает методологию валидации. На практике необходим набор репрезентативных промптов и измерение метрик качества (perplexity, accuracy, win-rate) для каждого варианта точности. Без этого невозможно гарантировать, что ускорение не куплено ценой деградации модели.
Глава рассматривает CUTLASS и Transformer Engine, но не описывает интеграцию с inference-движками. На практике TensorRT-LLM и vLLM используют собственные оптимизированные ядра, которые могут отличаться от CUTLASS. Для инженера важно понимать, какие именно ядра используются в его стеке, и как их профилировать.
Глава описывает использование DeepSeek недокументированной PTX-инструкции, но не даёт рекомендаций по управлению рисками. На практике такие инструкции могут сломаться на новой архитектуре или драйвере. Необходимо иметь план Б и тестировать на каждой новой конфигурации.
Девятая глава книги Fregly — систематический разбор техник повышения арифметической интенсивности CUDA-ядер. Ключевая идея: превратить ограниченное памятью ядро в ограниченное вычислениями за счёт повторного использования данных, слияния операций, понижения точности и задействования Tensor Cores. Для инженера это означает, что понимание иерархии памяти, коммуникационных паттернов и принципов параллелизма важнее, чем знание конкретных цифр текущего поколения.
Проверка прогнозов показывает: книга оказалась точной в своих предсказаниях. Tensor Cores, CUTLASS, разреженность и TMEM стали стандартом. Единственное, что изменилось быстрее ожидаемого — доминирование inference над training и скорость внедрения FP4. Главный вывод остаётся неизменным: инвестиции в фундаментальные принципы обеспечивают адаптацию к любому поколению железа.
Источники: Chris Fregly, «AI Systems Performance Engineering» (O'Reilly, 2026), Глава 9 — цитаты приводятся в ограниченном объёме в целях обзора и критического анализа. Документация NVIDIA CUDA, CUTLASS, Transformer Engine, TensorRT-LLM; публикации DeepSeek; MLPerf Training/Inference v5.0; PyTorch documentation.
This material is a review of Chapter 9 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 about the platform, 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 9 of Fregly's book is dedicated to improving CUDA kernel efficiency and arithmetic intensity. It examines how to turn a memory-bound kernel into a compute-bound one through data reuse, operation fusion, precision reduction, and Tensor Core utilization. The key points of the chapter are presented below, along with clarifications based on current data about the platform.
«Even if you've fully hidden latency through massive parallelism and high ILP, kernel performance may still be limited by how much useful work it performs for every byte read from memory».
«New GPU generations increase compute power noticeably faster than memory bandwidth. This gap makes improving arithmetic intensity increasingly critical».
In the chapter, the term «arithmetic intensity» is used not as an abstract metric but as a diagnostic tool. If a kernel is below the ridge point on the Roofline, it is memory-bound, and any compute optimizations will have no effect. If above — it is compute-bound, and further growth is possible only through more efficient use of ALUs and Tensor Cores.
The practical consequence noted in the chapter: the goal is always the same — increase arithmetic intensity where possible by performing more compute per byte transferred to and from global memory. This can be achieved in several ways: tiling, caching, kernel fusion, precision reduction, larger batch size, and recompute instead of store.
It is separately emphasized that not every problem allows you to easily increase arithmetic intensity — that's limited by the algorithm's characteristics. Element-wise operations have intensity around 0.125–0.25 FLOP/byte and are catastrophically memory-bound individually. Matrix multiplications, in contrast, have intensity that grows with matrix size and easily exceeds the ridge point.
«Arithmetic intensity, also called operational intensity, measures how many floating-point operations are performed per byte of data transferred from memory. It is measured in FLOPS per byte».
«Arithmetic intensity is a key metric in the Roofline performance model. This model visually displays kernel performance (FLOPs/sec) as a function of arithmetic intensity (FLOPs/byte). The plot shows two hardware ceilings: a slanted line from the origin — memory bandwidth limit, and a horizontal line — compute power limit».
In the chapter, the Roofline model is described as a tool that immediately tells you whether a kernel is memory-bound or compute-bound. A kernel with intensity below the ridge point is in the memory-bound regime — its speed is limited by memory bandwidth because the GPU spends most of its time waiting for data rather than computing. A kernel with intensity above the ridge point is compute-bound, using ALU and Tensor Cores close to their peak capabilities.
A concrete example is given: on NVIDIA H100, peak BF16 performance is ~990 TFLOPS, while HBM bandwidth is ~3.35 TB/s. The ridge point — the ratio of these values — is about 296 FLOP/byte. On Blackwell B200 with NVFP4 support, peak performance reaches 9 petaFLOPS (dense) in FP4, while HBM3e bandwidth is around 8 TB/s, pushing the ridge point even higher.
Practical consequence: Roofline plots can be built with Nsight Compute, which includes a Roofline analysis mode. It lets you check whether a kernel was originally limited by memory or compute, and then track improvements as you optimize.
GFLOP/s vs Arithmetic Intensity (FLOPs/byte)
Arithmetic intensity is the ratio of FLOPs to bytes. The Roofline model visualizes two ceilings: memory bandwidth and compute power. The ridge point separates memory-bound and compute-bound regimes.
H100: ~990 TFLOPS BF16, ~3.35 TB/s HBM, ridge point ~296 FLOP/byte. B200: 9 PFLOPS (dense) in FP4, 18 PFLOPS with sparsity, ~8 TB/s HBM3e. The ridge point shifts upward with each generation.
Nsight Compute allows building Roofline plots and tracking how optimizations shift the kernel rightward — from memory-bound to compute-bound.
«Tiling and data reuse are effective ways to increase arithmetic intensity. Loading a small submatrix (tile) from A and B into shared memory lets every byte read from global memory participate in many multiply-add operations at SRAM speed».
«Whenever code is restructured so that each element is loaded once and used dozens or hundreds of times, the FLOPs-per-byte ratio is multiplied by the reuse factor».
In the chapter, tiling is described as a basic technique for increasing arithmetic intensity. For matrix multiply C = A × B with a T × T tile, global reads are reduced from O(N) per result to O(N/T) per tile plus O(T) reuse inside shared memory. For a 128 × 128 tile with K-dimension size 64, the arithmetic intensity is about 64 FLOP/byte — orders of magnitude higher than a naive element-wise operation.
Beyond simple shared-memory tiling, the chapter examines multilevel tiling. In this approach, after a tile is placed in shared memory, each thread loads microtiles into registers using vector types like float4 and <half2>. This hierarchy of reuse shrinks the working set at each level and minimizes external memory traffic.
It is separately stipulated that one must ensure coalesced global reads when filling shared memory and align/swizzle data in shared memory to avoid bank conflicts. Bank conflicts in shared memory can reduce effective bandwidth by 2–32×, so swizzling is a standard technique for bank-free access.
Practical consequence: on modern GPUs, these inner tiling steps are often covered by MMA fragment APIs. Hardware moves data between shared memory and Tensor Memory implicitly using Tensor Core instructions.
DRAM → SMEM → Registers, data flows downward with reuse
Global → shared memory directly, optionally skipping L1
Tiling increases arithmetic intensity through data reuse in shared memory. Multilevel tiling adds register-level reuse.
Shared memory bank conflicts can reduce effective bandwidth by 2–32×. Swizzling or padding is mandatory.
Software prefetching (double buffering) overlaps DRAM → SMEM with compute. cudaMemPrefetchAsync is a hint, not a guarantee.
«On modern GPUs, you can extend the reuse-tiling idea by using thread block clusters from Cooperative Groups. They let multiple thread blocks share data through distributed shared memory (DSMEM)».
«A cluster of four thread blocks can jointly load one tile using the multicast feature of the Tensor Memory Accelerator (TMA). Each tile is shared among the four blocks so that global memory traffic for that tile is amortized across the cluster».
In the chapter, thread block clusters are described as a way to directly increase arithmetic intensity. Tiles are loaded only once and reused by all blocks in the cluster. This can cut global memory traffic by up to 4× in a 2×2 cluster, reducing the denominator — the number of bytes moved from global memory.
Blackwell supports thread block clusters of up to 16 blocks if you enable the non-portable cluster size beyond the standard 8 CTA limit. This requires setting the cudaFuncAttributeNonPortableClusterSizeAllowed attribute for the kernel. Larger clusters can increase reuse but may reduce occupancy, so profiling before enabling 16 is mandatory.
It is separately stipulated that DSMEM lets CTAs in a cluster read shared memory of other CTAs, but those loads are packetized and their bandwidth is noticeably lower than local SMEM. For real inter-CTA throughput, use cp.async.bulk.
Practical consequence: clusters are co-scheduled within a GPC. If the cluster size doesn't divide the number of SMs per GPC, some SMs idle. Blackwell introduces preferred/fallback cluster sizes to soften this quantization.
Thread blocks (CTAs) in one cluster share distributed on-chip SRAM
One DMA transfer broadcasts the same tile into four CTAs' shared memory
Thread block clusters + DSMEM + TMA multicast reduce global memory traffic up to 4× in a 2×2 cluster. This directly increases arithmetic intensity.
DSMEM loads are packetized and slower than local SMEM. Clusters are co-scheduled within a GPC; cluster size affects occupancy.
For real inter-CTA throughput, use cp.async.bulk. Blackwell supports clusters up to 16 CTAs when the non-portable size is enabled.
«Another way to increase arithmetic intensity is to merge multiple operations or loop iterations into a single operation. By fusing several kernels together, data loaded from memory can be used for multiple computations and iterations before writing back».
«Element-wise operations apply a simple computation independently to each tensor element. Fusing such operations eliminates excess memory traffic by keeping intermediate values on-chip».
In the chapter, kernel fusion is described as a way to avoid writing intermediate results to global memory. Two types of fusion are presented: vertical — combining sequential operations on the same data (e.g., sin and sqrt in one kernel), and horizontal — combining parallel operations over data (e.g., bias-add and activation after matrix multiplication in the GEMM epilogue).
A concrete example with L2 normalization is given. A naive implementation consists of three separate kernels: square, reduce, divide. Each kernel performs about 1 FLOP per element and requires 12 bytes of traffic per element. That gives intensity 3/36 ≈ 0.083 FLOP/byte. In the fused kernel, each thread walks its slice twice: once for accumulating the sum of squares and once for writing normalized outputs. Global traffic is about 12 bytes per element, with about 3 FLOP per element. Fused kernel arithmetic intensity: 3 / 12 ≈ 0.25 FLOP/byte — 3× higher.
It is separately stipulated that overly aggressive fusion can increase register pressure and reduce occupancy. If register usage becomes excessive and spilling to local memory begins, the benefit of fusion may be offset by extra memory traffic. Always profiling fused kernels with Nsight Compute is mandatory.
Practical consequence: modern deep learning frameworks can automatically fuse and unroll operations through JIT compilers and graph optimizers. For example, torch.compile in PyTorch, specifically TorchInductor, can automatically fuse sequences of element-wise operations.
Naive implementation (two kernels) and fused implementation (one kernel)
Kernel fusion — vertical and horizontal — increases arithmetic intensity by keeping intermediate results in registers.
Over-fusion increases register pressure and may cause spilling to local memory. Profiling is always required.
torch.compile and TorchInductor automatically fuse sequences of element-wise operations. For non-standard patterns, manual fusion may be needed.
«On modern GPUs, 2:4 sparsity is accelerated in hardware by Sparse Tensor Cores and cuSPARSELt. The 2:4 pattern means exactly two of every four consecutive weights are non-zero».
«By pruning half the weights in a 2:4 pattern, each memory load now delivers twice as many non-zero values that actually participate in multiplication. This immediately cuts global memory reads and increases effective arithmetic intensity by almost 2×».
In the chapter, structured sparsity is described as a way to increase effective arithmetic intensity. Since weight memory traffic is halved and computation is performed only on non-zero elements, Sparse Tensor Cores can approach 2× speedup over dense counterparts.
It is separately stipulated that real speedup heavily depends on matrix structure, pattern density, and index storage overhead. For small matrices, sparsity may even slow execution. Speedup assumes model accuracy is preserved — in practice, it is important to fine-tune or carefully calibrate the model after pruning.
Practical consequence: in PyTorch, to_sparse_semi_structured() is used to convert trained dense modules to the 2:4 format before deploying sparse GEMMs on Sparse Tensor Cores. Structured sparsity is usually applied to inference; during training, gradients do not benefit from 2:4 sparsity.
Before applying sparsity, it is important to first implement fundamental optimizations: coalesce all global loads, reuse data via tiling, and combine element-wise operations.
Exactly 2 of every 4 consecutive weights are non-zero → 2× Tensor Core throughput
2:4 structured sparsity increases effective arithmetic intensity by almost 2× through reduced weight traffic.
Real speedup depends on matrix structure and index overhead. For small matrices, slowdown is possible. Applied mainly to inference.
PyTorch uses to_sparse_semi_structured(). Before applying, implement basic optimizations: coalescing, tiling, fusion.
«Instead of storing or loading precomputed values, consider recomputing them on demand if memory bandwidth is the bottleneck. For example, recomputing x*x in registers is often faster than loading a previously computed x² from global memory».
«Many LLM inference engines use this technique to save memory. Instead of storing large activation tensors in HBM and reading them back later, they can recompute certain layers and activations on the fly».
In the chapter, recompute is described as a way to reduce memory traffic at the cost of increased computation. The classic example is FlashAttention, which recomputes softmax statistics in the forward pass instead of storing the full N × N attention matrix in HBM. This trades a few extra FLOPS for a dramatic reduction in global memory access, turning O(N²) HBM accesses into O(N² / T), where T is the tile size.
It is separately emphasized that the decision to recompute or store depends on the compute-to-memory-bandwidth ratio on the specific GPU: if a kernel is memory-bound, recompute is almost always beneficial; if compute-bound, storing may be preferable.
Practical consequence: recompute improves effective FLOPS/byte and allows larger models to fit in less memory. Activation checkpointing is a special case of this approach used in training.
Recompute instead of store is a trade of computation for memory bandwidth. FlashAttention and activation checkpointing are practical examples.
The decision depends on the compute-to-memory-bandwidth ratio on the specific GPU. For memory-bound kernels, recompute is almost always beneficial.
«In PyTorch, many of these ideas are applied automatically. The PyTorch compiler (torch.compile) can automatically fuse chains of element-wise operations and even some reductions».
«Nested tensors, or ragged tensors, allow representing batches of variable-length inputs without padding. Each nested tensor packs variable-length sequences into a single efficient buffer. NestedTensor provides a normal tensor interface but eliminates unnecessary zero-padding».
In the chapter, PyTorch is described as an environment where many arithmetic intensity optimizations are applied automatically. Because it uses cuDNN and cuBLAS under the hood, PyTorch performs tiling and uses shared memory during matrix operations. scaled_dot_product_attention (SDPA) can route computation to FlashAttention, memory-efficient, or cuDNN backends depending on tensor shapes and data types.
Nested tensors are separately examined as a way to represent variable-length batches without padding. When using nested tensors, attention and batch matrix multiplies fetch only the necessary data from memory, bringing the kernel closer to compute-bound on the Roofline.
Practical consequence: while PyTorch can recognize and compile most operations, some non-standard operations or custom CUDA ops may not fuse. In such cases, manual optimization may be needed.
PyTorch automatically applies tiling, fusion, and optimal attention backend selection. Nested tensors eliminate padding, improving load efficiency.
For non-standard operations, manual fusion or custom kernels may be needed. torch.compile handles most standard patterns.
«Modern NVIDIA GPUs implement reduced-precision computation such as TF32, FP16, FP8, FP4 and INT8 in Tensor Cores. Each SM in Blackwell has 256 KB of on-chip Tensor Memory (TMEM) dedicated to Tensor Core data».
«The key architectural difference between Blackwell and Hopper: on Hopper, MMA results were tightly coupled to the warpgroup register file, while on Blackwell accumulators live in TMEM, decoupling result ownership from individual threads».
In the chapter, Tensor Cores are described as the key tool for increasing arithmetic intensity. Lowering precision increases operations per byte: FP8 gives 2×, FP4 up to 4×. Memory shrinks proportionally: FP4 halves bytes per parameter versus FP8.
The architectural difference between Blackwell and Hopper is described as fundamental. tcgen05 instructions are CTA-scoped — one thread issues an instruction on behalf of the entire CTA, unlike warpgroup-scoped wgmma on Hopper. This gives more scheduling flexibility: some threads can orchestrate computation, others can move data.
Practical consequence: used properly, these capabilities can turn a once memory-bound tensor-intensive kernel into a fully compute-bound one, boosting arithmetic intensity by 2×, 4×, or even 8×.
Low-precision inputs, higher-precision accumulation inside Tensor Cores
Tensor Cores and TMEM are the foundation for increasing arithmetic intensity. Blackwell uses CTA-scoped tcgen05 instructions with accumulators in TMEM.
Lowering precision increases operations per byte: FP8 — 2×, FP4 — up to 4×. TMEM reduces register file pressure.
«At the heart of high-performance tensor computation lies TMEM — a 256 KB SRAM buffer per SM. Programmers do not explicitly allocate or manage TMEM — TMEM is handled by hardware or libraries when using Tensor Core operations».
«To transfer data from global memory to shared memory, use TMA or cuda::memcpy_async with the CUDA Pipeline API».
In the chapter, TMEM is described as a dedicated 256 KB SRAM buffer per SM that Transformer Engine and Tensor Cores use for storing results instead of registers. Hardware manages the movement of data between shared memory and TMEM implicitly using Tensor Core instructions.
An example of a two-stage double-buffered pipeline is given, in which cuda::memcpy_async loads the next tile into shared memory in parallel with Tensor Core computation. It is separately stipulated that on Blackwell, UMMA has a constraint: operand A may come from TMEM or SMEM, operand B only from SMEM. A from TMEM and B from TMEM is not supported.
Practical consequence: for data transfer from global memory to shared memory, TMA or cuda::memcpy_async is used. This overlaps data movement with Tensor Core compute and helps hide memory latency.
Accumulators live in TMEM (256 KB/SM) instead of the register file
TMEM is a dedicated 256 KB/SM buffer for Tensor Core accumulators. TMA and cuda::memcpy_async overlap data transfer with compute.
UMMA on Blackwell: A may be from TMEM or SMEM, B only from SMEM. A from TMEM and B from TMEM is not supported.
For data transfer from HBM to shared memory, TMA or cuda::memcpy_async is used. Double buffering hides memory latency.
«Although Tensor Cores were originally designed for FP16, they also support TF32, which sits between FP32 and FP16. TF32 uses an 8-bit exponent like FP32 and a 10-bit mantissa like FP16».
«Beyond TF32, automatic mixed precision (AMP) in PyTorch can choose the optimal precision (FP16 or BF16) for each operation and accumulate results in FP32 for stability».
In the chapter, TF32 is described as a way to gain speed without changing code. In PyTorch, it is enough to set torch.set_float32_matmul_precision('high'), and high-level operations automatically execute as TF32 Tensor Core kernels.
Automatic mixed precision (AMP) in PyTorch can choose the optimal precision for each operation and accumulate results in FP32 for stability. BF16 helps avoid overflow issues inherent to FP16. Under the hood, TorchInductor automatically fuses precision conversions to ensure: large GEMMs run on Tensor Cores in FP16 or TF32, accumulation remains in FP32, and small sensitive kernels run in FP32.
Practical consequence: when using structured sparsity or extremely low precision (FP8/FP4), make sure the batch size or tile granularity is large enough that TMEM and Tensor Cores stay fully loaded. Small batches carry overhead.
TF32 is enabled with one line. AMP automatically selects precision per-operation and accumulates in FP32.
FP8/FP4 require sufficient batch size. Small batches do not amortize overhead from format conversions and index handling.
«BF16/FP16 has been supported for several GPU generations, but Tensor Cores on modern GPUs often deliver more than 90% of peak BF16/FP16 throughput, about 4× FP32 peak».
«To push throughput higher, use FP8. By shrinking 16-bit weights 50% down to 8 bits, you halve memory traffic — and double the number of weights loaded per HBM transaction».
In the chapter, precision reduction is considered a way to increase arithmetic intensity. Every precision reduction doubles or quadruples the number of operations per byte. FP8 with FP32 or TF32 accumulation reaches 2–3× higher TFLOPS than BF16/FP16.
To address precision issues at very low precision, Transformer Engine supports FP8 as well as NVIDIA's 4-bit NVFP4 format with micro-scaling. NVFP4 applies two-level scaling, combining microblock-level scaling (every 32 elements along the K-dimension share a scale factor in E8M0 format) with a higher-level scale.
Blackwell B200 delivers 9 petaFLOPS (dense) in NVFP4 (18 petaFLOPS with sparsity), while FP32 peak is around 80 teraFLOPS (dense) — that is about 112× more than FP32. Blackwell B300 (Ultra) — 15 petaFLOPS (dense) in NVFP4.
Practical consequence: when TMEM/TMA overlap memory and compute, these low-precision formats turn previously memory-bound kernels into fully compute-bound. Different transformer layers have different sensitivity to precision reduction; calibration by sweeping precision across layers is recommended.
Lower precision leads to multiplicative throughput gain
FP8 gives 2×, FP4 up to 4× throughput. NVFP4 with microscaling preserves accuracy through two-level scaling.
Different layers have different sensitivity to precision reduction. Calibration by sweeping precision across layers and measuring quality metrics is recommended.
«LLM inference tasks can typically tolerate quantization to reduced INT8 precision, supported by modern GPUs with DP4A instructions on regular CUDA cores and integer matrix-multiply-accumulate instructions on Tensor Cores».
«Because INT8 weight traffic is just one byte per element instead of four bytes for FP32, memory traffic for weights is cut by 75%».
In the chapter, INT8 is described as a way to improve inference efficiency. At the instruction level, DP4A performs four INT8 multiply-accumulate (MAC) operations per instruction, versus one FP32 FMA per instruction. Each GPU can process roughly 4× more data per second from memory when using INT8 weights.
At the time of writing, INT8 quantization support in PyTorch is provided through TorchAO and vendor backends. Quantized modules run using dedicated INT8 kernels. Using cuBLASLt or CUTLASS for low-level INT8 GEMM ensures Tensor Core engagement.
Practical consequence: INT8 is an effective way to increase inference throughput, especially for models where FP32 precision is redundant. Combining INT8 with 2:4 sparsity can yield additional speedup.
INT8 cuts weight traffic by 75% and allows processing 4× more data per second. DP4A performs 4 MAC per instruction.
INT8 support in PyTorch is provided through TorchAO and vendor backends. cuBLASLt and CUTLASS for INT8 GEMM engage Tensor Cores.
«Modern NVIDIA GPUs include Transformer Engine, which combines hardware Tensor Core support for reduced-precision formats with a software runtime for scaling and casting».
«TMEM is a 256 KB SRAM buffer per SM that Transformer Engine and Tensor Cores use for storing results instead of registers. In practice, you never explicitly allocate TMEM — it's handled by hardware».
In the chapter, the Transformer Engine workflow is described as a sequence of steps: the kernel issues an MMA call or launches a CUTLASS GEMM; TE firmware orchestrates TMA copies of weights and activations from HBM to shared memory; Tensor Core instructions implicitly move operands between SMEM and TMEM; partial sums accumulate in TMEM at higher precision; results are written from TMEM to shared or global memory.
Ideally, weights are in FP8 or FP4, and activations are cast to FP8/FP4 where possible. Tensor Core MMA operations run at low precision, e.g. FP8 × FP8 with higher-precision accumulation, or FP16 × FP16 with FP32 accumulation.
It is separately stipulated that UMMA on Blackwell has a constraint: operand A may come from TMEM or SMEM, operand B only from SMEM. A from TMEM and B from TMEM is not supported.
Practical consequence: if building a custom tile loop, one can overlap data movement with Tensor Core compute using cuda::memcpy_async and the CUDA Pipeline API. While Tensor Cores process the current tile, cuda::memcpy_async streams the next tile from HBM to shared memory.
Transformer Engine automatically manages precision and scaling. TMEM is handled by hardware without explicit allocation.
UMMA on Blackwell: A from TMEM or SMEM, B only from SMEM. A from TMEM and B from TMEM is not supported.
cuda::memcpy_async and CUDA Pipeline API overlap data movement with Tensor Core compute, hiding memory latency.
«One of the easiest ways to engage the optimizations described above is to use the NVIDIA CUTLASS library. With CUTLASS you write a single template call, and it automatically applies many advanced optimizations: shared-memory tiling, asynchronous memory transfers, double buffering using the 256 KB TMEM buffer per SM, warp specialization, thread block clusters, and TMA multicast».
In the chapter, CUTLASS is described as a way to get manual MMA kernel performance without weeks of low-level tuning. A comparison is given: manual MMA kernel and CUTLASS GEMM both show ~98% Tensor Core utilization, while CUTLASS requires only a few lines of template code.
CUTLASS templates support FP4, FP8, FP16, TF32, and INT8 operand types. To build for a specific architecture, you must specify the target with an a suffix (e.g., sm100a for Blackwell), because tcgen05 instructions fall under architecture-accelerated features without forward compatibility guarantees. Code compiled for SM100 is incompatible with GeForce RTX 50 (SM120).
Practical consequence: optimized libraries like cuBLAS are built on CUTLASS. High-level libraries like PyTorch call these optimized libraries for many kernels. As more PyTorch operators use CUTLASS under the hood, you can benefit from the same optimizations without writing a single line of CUDA C++.
| Metric | Manual MMA kernel | CUTLASS GEMM |
|---|---|---|
| Tensor Core utilization | ~98% | ~98% |
| Registers per thread | ~52 | ~60 (slightly higher) |
| Shared memory per CTA | ~2 KB | ~4 KB |
| Development effort | High | Low (template configuration) |
CUTLASS delivers ~98% Tensor Core utilization with minimal effort. Automatically applies tiling, async copies, double buffering, warp specialization, and TMA multicast.
tcgen05 instructions require building for a specific architecture with an a suffix (sm100a for Blackwell). SM100 code is incompatible with SM120.
cuBLAS and PyTorch use CUTLASS under the hood. For standard GEMM and fused attention, using CUTLASS instead of hand-writing kernels is recommended.
«For those willing to go beyond C++ and dive into low-level micro-optimizations, CUDA allows using inline Parallel Thread Execution (PTX) code and SASS to extract the last grains of performance».
«PTX is a low-level virtual machine for parallel thread execution. SASS is the low-level assembly language that actually executes natively on NVIDIA GPU hardware».
In the chapter, inline PTX is described as a tool of last resort. Typical use cases: using a GPU instruction with no direct CUDA intrinsic; applying memory load modifiers (cache hints) to specific accesses; inserting memory or synchronization barriers at precise points; manually reordering instructions to avoid pipeline stalls; reading special registers.
Examples are given: L2 prefetch via cp.async.bulk.prefetch.L2.global; controlling cache behavior via ld.global.cg.f32; accessing special registers via mov.u32 %0, %smid. It is noted that some developers have used %smid in persistent kernels so that only one block per SM performs a specific job.
Practical consequence: inline PTX/SASS provides maximum control at the cost of complexity. It is recommended to use it last, when you need a hardware feature unavailable in CUDA C++ — or when you've precisely identified a small code fragment where the compiler schedule can be improved.
| Version | Warp stalls (memory) | Issue IPC | Kernel time | Speedup |
|---|---|---|---|---|
| Optimized C++ (compiler schedule) | 35% cycles | 1.5 | 1.07 ms | 1.0× baseline |
| Hand-tuned PTX (manual scheduling and hints) | 20% cycles | 1.6 | 1.00 ms | 1.07× |
Inline PTX/SASS is a tool of last resort. Gives 5–10% speedup on already optimized kernels but requires high expertise.
SASS changes between GPU generations. PTX is more stable but doesn't guarantee performance stability. Code needs updating with each architecture change.
Use only after exhausting all high-level optimizations. Suitable for critical loops running millions of times.
«A well-known example of custom PTX is from the DeepEP library for DeepSeek's expert parallelism. This library used a special PTX instruction ld.global.nc.l1::no_allocate.l2::256b to optimize global memory access».
«By bypassing L1, loads can stream large blocks of data directly into L2 without evicting frequently used data residing in L1».
The chapter provides a breakdown of the ld.global.nc.l1::no_allocate.l2::256b instruction. Components: ld.global — load from global memory; .nc — non-coherent load using the texture cache; .l1::no_allocate — do not allocate data in L1 cache; .l2::256b — fetch 256 bytes at a time into L2.
It is noted that this instruction is not part of NVIDIA's official PTX ISA specification (though it was later discovered that the modifiers are documented). DeepSeek engineers found it "off the books" to fine-tune cache behavior on their Hopper H800 GPU variant, restricted by US export rules.
DeepEP also dedicates 20 SMs out of 132 on H800 exclusively for network communication, so the L1 cache of those SMs stays warm for network code. This reduces compute performance by ~15% but significantly improves communication throughput.
Practical consequence: using PTX this way carries risks: the instruction is not guaranteed stable across GPU generations. DeepSeek's DeepEP build includes a build-time flag DISABLE_AGGRESSIVE_PTX_INSTRS=1 to disable these aggressive instructions when compatibility problems arise.
The instruction ld.global.nc.l1::no_allocate.l2::256b is not guaranteed stable across GPU generations. Code that relies on it may break or produce incorrect results on future architectures. DeepSeek includes DISABLE_AGGRESSIVE_PTX_INSTRS=1 to disable it.
DeepSeek used an undocumented PTX instruction to optimize memory access. Bypassing L1 and loading into L2 in 256-byte chunks preserves hot data in L1.
The instruction is not guaranteed stable across GPU generations. The DISABLE_AGGRESSIVE_PTX_INSTRS=1 flag is required to disable it on compatibility issues.
Use with caution, test on each new architecture. Prefer documented alternatives (Triton eviction_policy, cache_modifier).
Chapter 9 of Fregly's book contains forecasts and recommendations that can be compared against the current state of technology. A comparison with open data is provided below.
«Every precision reduction doubles or quadruples the number of operations per byte and thus increases arithmetic intensity».
Confirmed. NVIDIA Blackwell B200 delivers 9 petaFLOPS (dense) in NVFP4 (18 petaFLOPS with sparsity) — approximately 112× more than FP32 (80 TFLOPS). Transformer Engine automatically manages per-layer precision. PyTorch, TensorRT-LLM and vLLM use FP8 and FP4 as the standard for inference.
«CUTLASS delivers ~98% Tensor Core utilization with minimal effort».
Confirmed. CUTLASS is used by cuBLAS under the hood. PyTorch, TensorRT-LLM and other frameworks use CUTLASS for critical operations. Python wrappers (cuTile, CuTe DSL) have appeared, lowering the barrier to entry.
«Structured sparsity is usually applied to inference. During training, gradients do not benefit from 2:4 sparsity».
Confirmed. Sparse Tensor Cores are supported in cuSPARSELt and PyTorch via to_sparse_semi_structured(). NVIDIA continues to develop 2:4 support in new GPU generations. For training, sparsity still does not deliver benefit.
«FP4 (NVFP4) Tensor Cores double FP8 throughput and are used for inference with microscaling to preserve accuracy».
Partially confirmed. FP4 and NVFP4 are hardware-supported on Blackwell but require calibration. For many models, FP8 remains the safer choice. FP4 is more actively used for KV cache and weights where precision tolerance is higher.
«TMEM is a 256 KB SRAM buffer per SM that Transformer Engine and Tensor Cores use for storing results instead of registers».
Confirmed. TMEM on Blackwell (256 KB/SM) allows Tensor Core accumulators to live outside the register file, reducing register pressure and cutting cache miss latency by about 58% compared to Hopper.
The chapter describes techniques for improving arithmetic intensity but leaves out several aspects that emerged in real-world operation.
The chapter focuses on performance but does not consider energy efficiency. FP4 and FP8 not only increase throughput but also reduce energy per operation. For large-scale deployments, this is critical: GB200 NVL72 consumes ~130 kW, and every doubling of efficiency directly reduces operating costs.
The chapter mentions that FP4 requires calibration but does not describe the validation methodology. In practice, a set of representative prompts is needed, along with measurement of quality metrics (perplexity, accuracy, win-rate) for each precision variant. Without this, it is impossible to guarantee that the speedup is not bought at the cost of model degradation.
The chapter examines CUTLASS and Transformer Engine but does not describe integration with inference engines. In practice, TensorRT-LLM and vLLM use their own optimized kernels that may differ from CUTLASS. For an engineer, it is important to understand exactly which kernels are used in their stack and how to profile them.
The chapter describes DeepSeek's use of an undocumented PTX instruction but does not provide risk management recommendations. In practice, such instructions may break on a new architecture or driver. A backup plan is required, and testing on each new configuration is essential.
Chapter 9 of Fregly's book is a systematic examination of techniques for improving CUDA kernel arithmetic intensity. The key idea: turn a memory-bound kernel into a compute-bound one through data reuse, operation fusion, precision reduction, and Tensor Core utilization. For an engineer, this means understanding memory hierarchy, communication patterns, and parallelism principles is more important than knowing specific numbers of the current generation.
Forecast verification shows: the book proved accurate in its predictions. Tensor Cores, CUTLASS, sparsity and TMEM have become standard. The only thing that changed faster than expected was the dominance of inference over training and the speed of FP4 adoption. 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 9 — quotes provided in limited volume for review and critical analysis purposes. NVIDIA CUDA, CUTLASS, Transformer Engine, TensorRT-LLM documentation; DeepSeek publications; MLPerf Training/Inference v5.0; PyTorch documentation.
