§Контекст главы

Глава 9 книги Fregly посвящена методам повышения эффективности CUDA-ядер за счёт роста арифметической интенсивности — отношения числа операций с плавающей запятой к объёму данных, переданных из памяти. Материал охватывает как концептуальные основы (модель Roofline), так и конкретные приёмы: многоуровневое тайлирование, слияние ядер, структурированную разреженность, понижение точности, Tensor Cores, TMEM, TMA, CUTLASS и тонкую настройку PTX. Ниже последовательно изложены ключевые положения главы, а также приведены уточнения по актуальным данным о платформе и проверка прогнозов.

IАрифметическая интенсивность и модель Roofline

Арифметическая интенсивность (её также называют операционной интенсивностью) — это метрика, показывающая, сколько операций с плавающей запятой выполняется на каждый байт данных, переданных из памяти. Измеряется в FLOPS на байт:

I = FLOPs / Bytes [FLOP/B]

Новые поколения GPU наращивают вычислительную мощность заметно быстрее, чем пропускную способность памяти. Этот разрыв делает повышение арифметической интенсивности всё более критичным. Например, на NVIDIA H100 пиковая производительность BF16 составляет ~990 TFLOPS, а пропускная способность HBM — ~3,35 ТБ/с. Точка перегиба (ridge point) — отношение этих величин — равна примерно 296 FLOP/байт. На Blackwell B200 с поддержкой NVFP4 пиковая производительность достигает 10 петаFLOPS, а пропускная способность HBM3e — около 8 ТБ/с, что сдвигает точку перегиба ещё выше. Чем выше интенсивность, тем больше вычислений приходится на каждый загруженный байт — а значит, тем полнее используется вычислительный потенциал GPU.

Арифметическая интенсивность — ключевая метрика в модели производительности Roofline. Эта модель визуально отображает производительность ядра (FLOPs/sec) в зависимости от арифметической интенсивности (FLOPs/byte). На графике видны два аппаратных «потолка»:

• Наклонная линия из начала координат — ограничение по пропускной способности памяти: P = B_peak × I.
• Горизонтальная линия — ограничение по вычислительной мощности: P = F_peak.

Достижимая производительность — минимум из этих двух:

P_attainable = min(F_peak, B_peak × I)

Это позволяет сразу понять, является ли ядро memory-bound или compute-bound. Ядро с интенсивностью ниже точки перегиба I* = F_peak / B_peak находится в memory-bound режиме — его скорость упирается в пропускную способность памяти. Ядро с интенсивностью выше I* — compute-bound, оно использует ALU и Tensor Cores близко к их пиковым возможностям.

На практике графики Roofline можно строить с помощью Nsight Compute, который включает режим анализа Roofline. С его помощью можно проверить, было ли ядро изначально ограничено памятью или вычислениями, а затем отслеживать улучшения по мере оптимизаций. Анализ Speed of Light в Nsight Compute дополняет Roofline, показывая конкретные причины простоев — «Memory Throttle», промахи кэша, зависимости конвейера.

Цель всегда одна: повышать арифметическую интенсивность там, где это возможно. Добиться этого можно несколькими способами:

• Тайлинг с повторным использованием данных — загрузка подматрицы в shared memory, где каждый элемент участвует в десятках или сотнях операций.
• Кэширование в L1/shared memory — удержание промежуточных результатов на кристалле.
• Слияние нескольких ядер в одно (kernel fusion) — чтобы промежуточные результаты не записывались в глобальную память.
• Использование типов данных пониженной точности (FP16, FP8, FP4) — сокращение объёма передаваемых данных и задействование Tensor Cores.
• Увеличение размера батча — амортизация загрузки весов по большему числу вычислений.
• Пересчёт вместо хранения — если пересчёт в регистрах дешевле загрузки из HBM.

Современные компиляторные фреймворки, такие как TorchInductor в PyTorch, автоматически выполняют часть этих оптимизаций: удерживают вычисления на GPU, сокращают трафик к внешней памяти и повышают эффективную арифметическую интенсивность.

Стоит помнить, что не каждая задача позволяет легко повысить арифметическую интенсивность — это ограничено характеристиками алгоритма. Поэлементные операции (ReLU, dropout, bias-add) имеют интенсивность порядка 0,125–0,25 FLOP/байт и катастрофически memory-bound по отдельности. Матричные умножения, напротив, имеют интенсивность 2MNK / (MN + NK + MK), которая растёт с размером матриц и легко превышает точку перегиба.

IIМногоуровневое микротайлирование и программная предвыборка

Всякий раз, когда код перестраивается так, что каждый элемент загружается один раз и используется десятки или сотни раз — как при тайлинге — соотношение FLOPs на байт умножается на коэффициент повторного использования. Например, при типичном матричном умножении C = A × B с тайлом T × T глобальные чтения сокращаются с O(N) на результат до O(N/T) на тайл плюс O(T) повторного использования внутри shared memory. Для тайла 128 × 128 с размером K-измерения 64 арифметическая интенсивность составляет:

I = (128 × 128 × 64 × 2) / ((128 × 64 + 64 × 128) × 2) ≈ 64 FLOP/байт

Это на несколько порядков выше, чем у наивной поэлементной операции.

Помимо простого тайлинга в shared memory, можно ещё больше повысить интенсивность и усилить ILP с помощью многоуровневого тайлинга. При таком подходе после размещения тайла в shared memory каждый поток загружает микротайлы в регистры, используя векторные типы вроде float4 и <half2>. Таким образом, повторные операции выполняются целиком в регистрах.

Необходимо обеспечивать коалесцирование глобальных чтений при заполнении shared memory и выравнивать/свиззлить данные в shared memory, чтобы избежать конфликтов банков. Конфликты банков shared memory могут снизить эффективную пропускную способность в 2–32 раза, поэтому свиззлинг (перестановка индексов) — стандартный приём для bank-free доступа.

На современных GPU эти шаги внутреннего тайлинга часто покрываются с помощью API фрагментов MMA. Аппаратное обеспечение перемещает данные между shared memory и Tensor Memory (TMEM) неявно, используя инструкции Tensor Core. На Blackwell инструкции tcgen05 неявно переносят данные между shared memory и TMEM, используя отдельное адресное пространство TMEM. TMEM — это выделенный буфер SRAM на 256 КБ на каждый SM (512 столбцов × 128 строк × 32 бита), который существенно снижает давление на регистровый файл и уменьшает задержку промаха кэша примерно на 58% (с ~1000 циклов на H200 до ~420 циклов на B200).

Тесно связанный приём — программная предвыборка, часто реализуемая как двойная буферизация. Вместо ожидания завершения вычислений над текущим тайлом можно заранее выдать асинхронные загрузки следующего тайла в shared memory. Это перекрывает передачу DRAM → SMEM с текущими арифметическими операциями.

При использовании унифицированной памяти на суперчипе CPU-GPU, таком как Grace Blackwell, можно задействовать cudaMemPrefetchAsync(), чтобы намекнуть, что тайл скоро понадобится. Это подсказывает среде выполнения мигрировать страницы по NVLink-C2C. Однако предвыборка — это лишь подсказка, а не гарантия. Опытные пользователи часто предпочитают явный cudaMemcpy или выделение закреплённой памяти, чтобы полностью избежать накладных расходов на миграцию страниц.

IIIВнутриядерный конвейеринг и двойная буферизация

Традиционно GPU полагаются на многопоточность на уровне варпов для скрытия задержек. Пока один варп ждёт загрузки из памяти, другие варпы продолжают вычисления. Внутриядерный конвейеринг идёт дальше, перекрывая память и вычисления внутри одного варпа или ядра.

Ключевые вызовы CUDA Pipeline API:

• pipe.producer_acquire() — резервирует следующую стадию конвейера для записи.
• pipe.producer_commit() — сигнализирует, что ранее выданные асинхронные операции для этой стадии готовы к потреблению.
• pipe.consumer_wait() — ожидает завершения ранее зафиксированных операций для этой стадии.
• pipe.consumer_release() — освобождает текущую стадию для повторного использования.

Две стадии конвейера перекрывают загрузки из глобальной памяти с вычислениями. На первой стадии (Stage 0) один варп в потоковом блоке выдаёт асинхронную предвыборку следующего тайла в разделяемую память. Предвыборка порождает кооперативные копии cuda::memcpy_async, которые понижаются до per-thread cp.async в разделяемую память.

Пока Stage 0 производит (загружает) данные в одном варпе, остальные варпы потокового блока потребляют (вычисляют) загруженные данные во второй стадии (Stage 1). Эта простая реализация «производитель–потребитель» скрывает задержку DRAM за счёт текущих вычислений.

Ядро выделяет одну непрерывную область разделяемой памяти для тайлов A и B через extern __shared__ float shared_mem[]. Оно делит этот буфер на четыре подобласти — две для A и две для B — с помощью арифметики указателей. Это позволяет реализовать настоящую двойную буферизацию без дополнительных динамических выделений.

Перед входом в основной цикл ядро асинхронно предзагружает первые STAGES тайлов, используя асинхронные копии, привязанные к конвейеру: producer_acquire() → memcpy_async() → producer_commit(). Варпы-потребители используют consumer_wait() и consumer_release(), чтобы вычисления начинались точно тогда, когда предзагруженный тайл готов.

Одновременно другие варпы в блоке вызывают pipe.consumer_wait(). Это эффективно приостанавливает только те потоки, которые зависят от данных в текущем буфере, до тех пор, пока производитель не зафиксирует его. После завершения ожидания каждый поток вызывает device-функцию computeTile(...) для выполнения своего точечного произведения TILE_SIZE x TILE_SIZE. Результаты инкрементально накапливаются в регистрах.

Завершив вычисления на текущем буфере, варпы вызывают pipe.consumer_release(), чтобы освободить эту стадию для повторного использования в последующих итерациях. Этот тонкий release предотвращает блокировку всего блока и максимизирует перекрытие между вычислениями и передачей данных.

Эта простая схема с двойной буферизацией примерно в 2 раза быстрее наивного тайлинга, как показано в таблице, где сравнивается производительность наивной реализации тайлинга и оптимизированной реализации конвейера с двойной буферизацией.

Примечание: числовые значения приведены для иллюстрации концепций.

IVСпециализация варпов и модель «производитель–потребитель»

Такой тип специализации позволяет каждому набору варпов иметь собственные последовательности инструкций. Таким образом, инструкции выдаются и выполняются непрерывно, не прерываясь другими типами операций. В частности, специализация варпов позволяет назначить один набор «производящих» (producer) варпов для асинхронной предзагрузки тайлов через cuda::memcpy_async. Затем все остальные «потребляющие» (consumer) варпы выполняют вычисления.

Здесь четыре варпа назначены на роль производителей, а остальные восемь — на роль потребителей. Как и в большинстве паттернов «производитель–потребитель», можно назначить разное количество варпов для производителя и потребителя.

Поскольку каждый варп имеет собственный планировщик, GPU может выдать инструкцию загрузки, математическую инструкцию и инструкцию записи — всё в одном цикле от разных варпов, использующих разные планировщики. Так, SM с несколькими планировщиками может выдать инструкцию памяти от варпа 0, математическую инструкцию от варпа 1 и так далее — за один цикл.

Это фактически создаёт сценарий множественной выдачи на уровне потокового блока между варпами. Это невозможно при двойной буферизации в одном варпе, поскольку планировщик одного варпа может выдать только одну инструкцию за цикл.

Интересный паттерн для специализации варпов — использование трёх различных типов варпов: «загрузчик», «вычислитель» и «сохранитель». Варп-загрузчик помещает тайлы в очередь конвейера. Варп-вычислитель запускает вычислительное ядро на каждом тайле. Варп-сохранитель записывает результаты.

Этот конвейер со специализацией варпов устраняет простои, которые не может устранить последовательный цикл загрузки-вычисления с двойной буферизацией в одном варпе. Эффективное перекрытие передачи данных и вычислений при специализации варпов повышает загрузку GPU — особенно для длительных циклов и persistent-ядер.

Ещё один паттерн специализации варпов — модификация трёхролевого конвейера. Он назначает набор варпов на загрузчик памяти, но затем использует два набора потребительских варпов, которые «пинг-понгуют» между ролями вычислителя и сохранителя. Эта трёхролевая архитектура реализована в CUTLASS как gemm_tma_warpspecialized_pingpong.

На практике специализация варпов чрезвычайно эффективна для выжимания последних крупиц производительности. Фактически, FlashAttention v3 приписывает свои ускорения частично конвейерам со специализацией варпов, которые перекрывают GEMM и вычисления softmax — вместе с передачей данных — чтобы держать все аппаратные блоки занятыми.

VКластеры блоков потоков и распределённая разделяемая память

На современных GPU можно расширить идею тайлинга с повторным использованием, используя кластеры блоков потоков из Cooperative Groups. Они позволяют нескольким блокам потоков совместно использовать данные через распределённую shared memory (DSMEM).

Кластеры блоков потоков напрямую повышают арифметическую интенсивность. Например, кластер из четырёх блоков потоков может совместно загрузить один тайл, используя функцию multicast в Tensor Memory Accelerator (TMA). Каждый тайл разделяется между четырьмя блоками так, что трафик глобальной памяти для этого тайла амортизируется по всему кластеру. Тайлы загружаются только один раз и переиспользуются всеми четырьмя блоками. Это может сократить трафик глобальной памяти до 4× в кластере 2×2, уменьшая знаменатель — количество байт, перемещаемых из глобальной памяти.

Blackwell поддерживает кластеры блоков потоков размером до 16 блоков, если включить непортируемый размер кластера сверх стандартного лимита в 8 CTA. Для этого нужно установить атрибут cudaFuncAttributeNonPortableClusterSizeAllowed для ядра. Большие кластеры могут повысить повторное использование, но способны снизить occupancy, поэтому профилируйте перед включением 16. Кроме того, кластеры ко-планируются в GPC; если размер кластера не делит число SM на GPC, часть SM простаивает.

Важно отметить, что DSMEM позволяет CTA внутри кластера читать общую shared memory, но такие загрузки пакетизируются и их пропускная способность заметно ниже локального SMEM (около 128 байт/такт), поэтому не следует относиться к удалённой DSMEM как к локальной shared memory. Для реальной меж-CTA пропускной способности применяется cp.async.bulk.

Кластеры потоковых блоков открывают новые стратегии распараллеливания нагрузок, которые ранее требовали коммуникации через глобальную память или множества запусков ядер. Например, представьте большое матричное умножение, в котором выходной тайл слишком велик для обработки одним потоковым блоком из-за ограничений разделяемой памяти.

С кластерами потоковых блоков и DSMEM два блока могут образовать кластер и напрямую разделить общую область разделяемой памяти на кристалле, чтобы бесшовно объединить свои результаты с помощью аппаратно-поддерживаемых примитивов. Аппаратное обеспечение DSMEM позволяет одному SM выполнять загрузки/сохранения/атомарные операции в разделяемую память другого SM через быструю сеть.

VIСлияние ядер (kernel fusion)

Ещё один способ повысить арифметическую интенсивность — объединить несколько операций или итераций цикла в одну операцию. Сливая несколько ядер вместе, данные, загруженные из памяти, можно использовать для нескольких вычислений и итераций, прежде чем записать обратно.

Поэлементные операции (их также называют точечными) применяют простое вычисление независимо к каждому элементу тензора. Слияние таких операций устраняет лишний трафик памяти, удерживая промежуточные значения на кристалле. Это увеличивает объём работы на каждый байт, прочитанный из глобальной памяти, — повышая арифметическую интенсивность.

Например, наивная реализация запускает два ядра. Первое читает x и записывает y в глобальную память. Второе читает y и записывает z:

y = sin(x);
z = sqrt(y);

Здесь каждый элемент обрабатывается дважды: один раз после sin(x) и один раз после sqrt(y). Арифметическая интенсивность каждого ядра очень низкая, поскольку на каждый элемент приходится лишь одна математическая функция на операцию загрузки/сохранения. В отличие от этого, слитое ядро выполняет тот же набор операций за один проход:

z[i] = sqrt(sin(x[i]));

Каждый x[i] загружается один раз, затем и sin, и sqrt применяются в регистрах, и только финальный z[i] записывается в память. Поскольку промежуточный y никогда не уходит в глобальную память, эффективные FLOPS на байт резко возрастают, приближая операцию к вычислительному потолку.

Вертикальное слияние объединяет последовательные операции над одними и теми же данными (как в примере выше), а горизонтальное слияние объединяет параллельные операции по данным (например, bias-add и активация после матричного умножения, выполненные в эпилоге GEMM). Оба типа слияния повышают арифметическую интенсивность.

Пример: L2-нормализация

Рассмотрим конкретный пример. Предположим, нужно L2-нормализовать каждую строку длины hidden в двумерном тензоре x (форма [batch, hidden]). Для каждой строки b вычисляется норма norm_b = sqrt(sum_i x[b,i]*x[b,i] + eps), после чего записывается y[b,i] = x[b,i] / norm_b для всех i.

Наивная реализация состояла бы из трёх отдельных ядер: возведение в квадрат, редукция, деление. Каждое ядро выполняет примерно 1 FLOP на элемент и требует 12 байт трафика на элемент. Итого арифметическая интенсивность 3/36 ≈ 0,083 FLOP/байт.

Вместо этого можно объединить все три шага в одно ядро:

__global__ void fusedL2Norm(const float* __restrict__ x,
                            float* __restrict__ y,
                            int hidden) {
  extern __shared__ float sdata[];
  const int batch = blockIdx.x;
  const int tid = threadIdx.x;
  const float* batch_ptr = x + size_t(batch) * hidden;
  float local = 0.f;
  for (int i = tid; i < hidden; i += blockDim.x) {
    float v = batch_ptr[i];
    local = fmaf(v, v, local);
  }
  sdata[tid] = local;
  __syncthreads();
  for (int offset = blockDim.x >> 1; offset > 0; offset >>= 1) {
    if (tid < offset) sdata[tid] += sdata[tid + offset];
    __syncthreads();
  }
  float inv = rsqrtf(sdata[0]);
  float* out_batch = y + size_t(batch) * hidden;
  for (int i = tid; i < hidden; i += blockDim.x) {
    out_batch[i] = batch_ptr[i] * inv;
  }
}

В этом слитом ядре каждый поток обходит свой срез x[b,*] дважды: один раз для накопления суммы квадратов и один раз для записи нормализованных выходов. Глобальный трафик — около 12 байт на элемент. На каждый элемент ядро выполняет примерно 1 умножение + 1 сложение во время редукции и 1 умножение при нормализации — итого ~3 FLOP. sqrt и rsqrt амортизируются по всей строке.

Арифметическая интенсивность слитого ядра: 3 / 12 ≈ 0,25 FLOP/байт — в 3 раза выше, чем у наивного варианта.

Дополнительные микрооптимизации: замена деления на умножение на rsqrtf; выполнение одного rsqrtf в одном потоке и трансляция скалярного результата другим потокам через shared memory.

Слияние не только повышает арифметическую интенсивность и сдвигает ядро в сторону compute-bound на Roofline, но и экономит пропускную способность памяти. Однако слишком агрессивное слияние может увеличить давление на регистры и снизить occupancy. Если использование регистров становится чрезмерным и начинается spilling в локальную память, выгода от слияния может быть перекрыта дополнительным трафиком памяти.

Когда слияние вредно: если объединяемые операции имеют радикально разные оптимальные размеры тайлов, разные требования к памяти или если одна операция compute-bound, а другая memory-bound с совершенно другим шаблоном доступа, их слияние может ухудшить производительность.

VIIСтруктурированная разреженность

На современных GPU разреженность 2:4 ускоряется аппаратно с помощью Sparse Tensor Cores и cuSPARSELt. Схема 2:4 означает, что ровно два из каждых четырёх последовательных весов ненулевые. Создание такой разреженности иногда называют прунингом.

Обрезая половину весов в шаблоне 2:4, каждая загрузка памяти теперь доставляет вдвое больше ненулевых значений, которые реально участвуют в умножении. Это немедленно сокращает глобальные чтения памяти и повышает эффективную арифметическую интенсивность почти в 2 раза.

Структурированная разреженность применяется после обучения модели. Модель обрезается и оптимизируется для инференса. В PyTorch применяется to_sparse_semi_structured() для конвертации обученных плотных модулей в формат 2:4 перед развёртыванием разреженных GEMM на Sparse Tensor Cores.

После конвертации модели будут вызываться оптимизированные разреженные GEMM-ядра, работающие на Sparse Tensor Cores. Sparse Tensor Cores могут приближаться к 2× ускорению по сравнению с плотными аналогами для многих задач инференса — особенно при подаче больших батчей входных данных, так как это амортизирует накладные расходы на запуск ядер и обработку индексов сжатого представления.

Ограничения разреженности 2:4:

• Реальное ускорение сильно зависит от структуры матриц, плотности паттернов и overhead на хранение индексов. Для малых матриц разрежённость может даже замедлить выполнение.
• Ускорение предполагает, что точность модели сохраняется. На практике важно дообучить или тщательно откалибровать модель после прунинга.
• Структурированная разреженность обычно применяется к задачам инференса. Во время обучения градиенты не получают выгоды от разрежённости 2:4.
• Перед применением разреженности важно сначала реализовать фундаментальные оптимизации: коалесцировать все глобальные загрузки, переиспользовать данные через тайлинг и объединять поэлементные операции.

VIIIКомпромисс между пересчётом и хранением

Вместо хранения или загрузки предварительно вычисленных значений (например, x²) возможен их пересчёт по требованию, если узким местом является пропускная способность памяти. Например, повторное вычисление x*x в регистрах часто быстрее, чем загрузка ранее вычисленного x² из глобальной памяти.

Многие движки инференса LLM используют этот приём для экономии памяти. Вместо хранения больших тензоров активаций в HBM и их чтения обратно позже они могут пересчитывать определённые слои и активации на лету. Это аналогично activation checkpointing в контексте обучения модели.

Классический пример — FlashAttention, который пересчитывает статистики softmax в прямом проходе вместо хранения полной матрицы внимания N × N в HBM. Это обменивает несколько дополнительных FLOPS на драматическое сокращение доступа к глобальной памяти, превращая O(N²) доступов к HBM в O(N² / T), где T — размер тайла.

Пересчёт улучшает эффективные FLOPS/байт и позволяет вместить большие модели в меньший объём памяти. Решение о том, пересчитывать или хранить, зависит от соотношения вычислительной мощности и пропускной способности памяти на конкретном GPU: если ядро memory-bound, пересчёт почти всегда выгоден; если compute-bound — хранение может быть предпочтительнее.

IXPyTorch и арифметическая интенсивность

В PyTorch многие из описанных идей применяются автоматически. Компилятор PyTorch (torch.compile) может автоматически сливать цепочки поэлементных операций и даже некоторые редукции. Он использует оптимизации на уровне графа выполнения, чтобы удерживать данные на GPU и максимально их переиспользовать.

Поскольку под капотом используются оптимизированные библиотеки, такие как cuDNN и cuBLAS, PyTorch выполняет тайлинг и использует shared memory при выполнении матричных операций через torch.matmul. Кроме того, scaled_dot_product_attention (SDPA) в PyTorch может направлять вычисления к бэкендам FlashAttention, memory-efficient или cuDNN в зависимости от форм тензоров и типов данных:

torch.nn.attention.sdpa_kernel(SDPBackend.FLASH_ATTENTION)

Важно отметить, что хотя PyTorch может распознавать и компилировать большинство операций, некоторые нестандартные операции или пользовательские CUDA-операции могут не сливаться. В таких случаях может потребоваться ручная оптимизация.

Вложенные тензоры (nested tensors), или ragged tensors, позволяют представлять батчи входных данных переменной длины без паддинга. Каждый вложенный тензор упаковывает последовательности переменной длины в единый эффективный буфер. NestedTensor предоставляет обычный интерфейс тензора, но устраняет ненужный zero-padding. Таким образом, глобальные загрузки памяти становятся более эффективными, потому что каждый загруженный байт полезен в вычислении.

Вложенные тензоры полезны для LLM с последовательностями переменной длины. При использовании вложенных тензоров такие операции, как attention и batch-matrix multiplies, извлекают из памяти только необходимые данные. Это приближает ядро к режиму compute-bound на модели Roofline и помогает сократить простои, связанные с ограничением памяти.

XСмешанная точность и использование Tensor Cores

Современные NVIDIA GPU реализуют вычисления пониженной точности, такие как TF32, FP16, FP8, FP4 и INT8, в Tensor Cores. Каждый SM в Blackwell имеет 256 КБ внутрикристальной памяти TMEM, выделенной для данных Tensor Core. Также имеется специализированный блок TMA, который асинхронно копирует тайлы между глобальной памятью и shared memory. Инструкции Tensor Core (например, tcgen05.mma) затем неявно перемещают операнды и аккумуляторы между shared memory и TMEM.

Ключевое архитектурное отличие Blackwell от Hopper: на Hopper результаты MMA были жёстко привязаны к регистровому файлу варп-группы, а на Blackwell аккумуляторы живут в TMEM, что декопулирует владение результатами от отдельных потоков. Инструкции tcgen05 являются CTA-scoped — один поток выдаёт инструкцию от имени всего CTA, в отличие от warpgroup-scoped wgmma на Hopper. Это даёт большую гибкость в планировании: одни потоки могут оркестрировать вычисления, другие — перемещать данные, и весь CTA действует как исполнительная единица.

При правильном использовании эти возможности могут превратить когда-то ограниченное памятью ядро с интенсивными тензорными операциями в полностью compute-bound, повысив арифметическую интенсивность (FLOPS на байт) в 2, 4 или даже 8 раз.

XIПодача данных в Tensor Cores через TMEM и TMA

В основе высокопроизводительных тензорных вычислений лежит TMEM — буфер SRAM на 256 КБ на каждый SM. Программисты не выделяют и не управляют TMEM явно — TMEM обрабатывается аппаратным обеспечением или библиотеками при использовании операций Tensor Core.

Под капотом Blackwell использует инструкции tcgen05.mma, которые работают с TMEM для хранения операндов и аккумуляторов. CUTLASS и библиотечные ядра управляют необходимым выделением и использованием через конфигурацию ядра и ассемблер PTX. Transformer Engine использует TMEM для промежуточных результатов, что снижает зависимости MMA от регистров.

Высокоуровневые API, такие как CUTLASS, автоматически обрабатывают всю эту сложность. При запуске операции Tensor Core MMA с использованием интринсиков CUDA MMA или GEMM из CUTLASS реализация управляет перемещением операндов через shared memory и TMEM. TMA потоково передаёт тайлы между глобальной памятью и shared memory, а инструкции Tensor Core неявно перемещают операнды между shared memory и TMEM.

Для выполнения передачи данных из глобальной памяти в shared memory применяется TMA или cuda::memcpy_async с конвейером из CUDA Pipeline API (<cuda/pipeline>). Пример двухступенчатого конвейера с двойной буферизацией:

#include <cuda/pipeline>
#include <cooperative_groups.h>
namespace cg = cooperative_groups;
extern "C" __global__
void stage_ab_tiles(const float* __restrict__ globalA,
                    const float* __restrict__ globalB,
                    float* __restrict__ outC,
                    int tile_elems,
                    int num_tiles) {
  assert((tile_elems % (32 * 4)) == 0);
  extern __shared__ float smem[];
  auto block = cg::this_thread_block();
  float* A0 = smem + 0 * tile_elems;
  float* A1 = smem + 1 * tile_elems;
  float* B0 = smem + 2 * tile_elems;
  float* B1 = smem + 3 * tile_elems;
  constexpr auto scope = cuda::thread_scope_block;
  constexpr int stages = 2;
  __shared__ cuda::pipeline_shared_state<scope, stages> pstate;
  auto pipe = cuda::make_pipeline(block, &pstate);
  pipe.producer_acquire();
  cuda::memcpy_async(block, A0, globalA + 0 * tile_elems,
                     cuda::aligned_size_t<32>{tile_elems * sizeof(float)}, pipe);
  cuda::memcpy_async(block, B0, globalB + 0 * tile_elems,
                     cuda::aligned_size_t<32>{tile_elems * sizeof(float)}, pipe);
  pipe.producer_commit();
  for (int t = 1; t < num_tiles; ++t) {
    pipe.producer_acquire();
    cuda::memcpy_async(block, (t & 1) ? A1 : A0,
                       globalA + t * tile_elems,
                       cuda::aligned_size_t<32>{tile_elems*sizeof(float)}, pipe);
    cuda::memcpy_async(block, (t & 1) ? B1 : B0,
                       globalB + t * tile_elems,
                       cuda::aligned_size_t<32>{tile_elems*sizeof(float)}, pipe);
    pipe.producer_commit();
    pipe.consumer_wait();
    float* prevA = (t & 1) ? A0 : A1;
    float* prevB = (t & 1) ? B0 : B1;
    pipe.consumer_release();
  }
  pipe.consumer_wait();
  int last = (num_tiles - 1) & 1;
  float* lastA = last ? A1 : A0;
  float* lastB = last ? B1 : B0;
  pipe.consumer_release();
}

При запуске этого ядра установите размер динамической shared memory равным 4 * tile_elems * sizeof(float), чтобы выделить A0, A1, B0 и B1. Этот шаблон двойной буферизации гарантирует, что как только один тайл находится в shared memory, Tensor Cores могут начать его обработку. Тем временем cuda::memcpy_async загружает следующий тайл в shared memory параллельно.

Выбор между TMA и cp.async (LDGSTS): TMA оптимизирован для больших регулярных блоков и достигает пиковой пропускной способности ~7,2 ТБ/с, но требует большого объёма данных в pipeline (typically >64 КиБ in-flight) для насыщения. cp.async (LDGSTS) лучше подходит для нерегулярных паттернов и меньших объёмов (<32 КиБ in-flight), достигая насыщения ~6,6 ТБ/с при ~32 КиБ in-flight. Выбирайте путь копирования по шаблону доступа, а не по идеологии.

XIITF32 и автоматическая смешанная точность

Хотя Tensor Cores изначально проектировались для FP16, они также поддерживают TF32, который находится между FP32 и FP16. TF32 использует 8-битный порядок, как FP32, и 10-битную мантиссу, как FP16. TF32 выполняется на Tensor Cores с существенно более высокой пропускной способностью, чем FP32 на CUDA-ядрах, сохраняя при этом диапазон порядка FP32. В PyTorch включить TF32 так же просто, как установить:

import torch
torch.set_float32_matmul_precision('high')

После установки этих флагов высокоуровневые операции, такие как torch.matmul и torch.nn.Linear, автоматически выполняются как ядра TF32 Tensor Core, а не в FP32 на стандартных CUDA-ядрах.

Помимо TF32, автоматическая смешанная точность (AMP) в PyTorch может выбирать оптимальную точность (FP16 или BF16) для каждой операции и накапливать результаты в FP32 для стабильности. BF16 помогает избежать проблем с переполнением, присущих FP16:

with torch.amp.autocast("cuda", dtype=torch.bfloat16):
    output = model(input)

Под капотом TorchInductor автоматически сливает эти преобразования точности, чтобы гарантировать: большие операции GEMM выполняются на Tensor Cores в FP16 или TF32, накопление остаётся в FP32 для числовой стабильности, небольшие «чувствительные» ядра, такие как layer normalization и softmax, выполняются в FP32, а GradScaler предотвращает underflow при обучении с FP16. BF16 имеет диапазон порядка FP32, поэтому GradScaler обычно не нужен при обучении с BF16.

При использовании структурированной разреженности или экстремально низкой точности (FP8/FP4) размер батча или гранулярность тайла должны быть достаточно велики, чтобы TMEM и Tensor Cores оставались полностью загруженными. Малые батчи несут накладные расходы, включая преобразования форматов, обработку разреженных индексов, нерегулярные шаблоны памяти. Батч размером 128 или 256 полностью задействует конвейер TMEM и обеспечит производительность, близкую к пиковой.

XIIIBF16/FP16, FP8 и FP4 пониженной точности

BF16/FP16 поддерживается уже несколько поколений GPU, но Tensor Cores на современных GPU часто могут поддерживать более 90% пиковой пропускной способности BF16/FP16, что примерно в 4 раза превышает пиковую пропускную способность FP32.

Обучение с FP16 использует более узкий 5-битный порядок, чем FP32, поэтому очень малые значения градиентов могут исчезать в ноль (underflow) без применения масштабирования потерь (loss scaling). BF16 соответствует 8-битному диапазону порядка FP32 и изначально избегает underflow — масштабирование потерь требуется редко. BF16 обычно предпочтительнее для обучения на современных GPU.

Чтобы ещё больше повысить пропускную способность, можно использовать FP8. Уменьшая 16-битные веса на 50% до 8 бит, вы вдвое сокращаете трафик памяти — и вдвое увеличиваете количество весов, загружаемых за одну транзакцию HBM. На практике матричные умножения FP8 с накоплением FP32 или TF32 достигают в 2–3 раза больших TFLOPS, чем BF16/FP16.

Для решения проблем точности при очень низкой точности Transformer Engine поддерживает FP8, а также 4-битный формат NVFP4 от NVIDIA с микро-масштабированием. NVFP4 применяет двухуровневое масштабирование, комбинируя масштабирование на уровне микроблока (каждые 32 элемента вдоль K-измерения имеют общий масштабный коэффициент в формате E8M0) и масштаб более высокого уровня. Blackwell B200 обеспечивает 10 петаFLOPS (плотных) в NVFP4, тогда как пиковая производительность FP32 составляет около 80 тераFLOPS (плотных). Blackwell B300 (Ultra) — 15 петаFLOPS (плотных) в NVFP4.

Чувствительность слоёв к квантованию: на практике разные слои трансформера имеют разную чувствительность к понижению точности. Слои attention (особенно QK-умножение) и FFN обычно более устойчивы к FP8/FP4, чем normalization и softmax, которые требуют FP32 для стабильности. Рекомендуется калибровка модели с перебором точности по слоям и измерением метрик качества (perplexity, accuracy) для каждого варианта.

Поскольку точность настолько низка, 256 КБ TMEM на SM могут вместить большие тайлы FP4 (например, 256×256), что дополнительно увеличивает внутрикристальное повторное использование и повышает производительность.

Каждое снижение точности удваивает или учетверяет количество операций на байт и, следовательно, повышает арифметическую интенсивность. Когда TMEM/TMA перекрывают память и вычисления, эти форматы низкой точности превращают ранее memory-bound ядра в полностью compute-bound, полностью задействуя вычислительные движки Tensor Core.

XIVINT8 и инструкции DP4A для инференса

Задачи инференса LLM обычно могут допускать квантование до пониженной точности INT8, поддерживаемое современными GPU с помощью инструкций DP4A (SIMD dot-product) на обычных CUDA-ядрах и инструкций целочисленного матричного умножения-накопления (MMA) на Tensor Cores. На уровне инструкций DP4A выполняет четыре операции умножения-накопления INT8 (MAC) за инструкцию по сравнению с одной операцией FMA FP32 за инструкцию.

Поскольку трафик весов для INT8 составляет всего один байт на элемент вместо четырёх байт для FP32, трафик памяти для весов сокращается на 75%. Задачи инференса INT8 могут значительно превосходить FP32 благодаря более высокой пиковой пропускной способности INT8 Tensor Core и сниженному трафику памяти — каждый GPU может обрабатывать примерно в 4 раза больше данных в секунду из памяти при использовании весов INT8.

На момент написания поддержка INT8-квантования в PyTorch обеспечивается через TorchAO и вендорные бэкенды. Квантованные модули работают с использованием выделенных ядер INT8. Использование cuBLASLt или CUTLASS для низкоуровневого INT8 GEMM обеспечивает задействование Tensor Core.

XVTransformer Engine и TMEM в деталях

Современные NVIDIA GPU включают Transformer Engine, который сочетает аппаратную поддержку Tensor Core для форматов пониженной точности с программной средой выполнения для масштабирования и приведения типов. Ядра в cuBLASLt, cuDNN, CUTLASS или OpenAI Triton выполняют инструкции cp.async или передачи TMA в shared memory. Инструкции Tensor Core затем неявно перемещают операнды между shared memory и TMEM.

TMEM — это буфер SRAM на 256 КБ на SM, который Transformer Engine и Tensor Cores используют для хранения результатов (вместо регистров). На практике TMEM не выделяется явно — всё обрабатывается аппаратным обеспечением.

Рабочий процесс Transformer Engine:

1. Ядро выдаёт вызов MMA или запускает GEMM из CUTLASS.
2. Прошивка Transformer Engine организует копирование TMA (или cuda::memcpy_async) весов и активаций из HBM в shared memory.
3. Инструкции Tensor Core (например, tcgen05.mma) неявно перемещают операнды между SMEM и TMEM во время конвейера MMA.
4. Частичные суммы накапливаются в TMEM с более высокой точностью (BF16, FP16, FP32).
5. Результаты записываются из TMEM в shared или глобальную память коалесцированно.

В идеале веса находятся в FP8 или FP4, а активации приводятся к FP8/FP4, когда это возможно — иначе активации могут оставаться в формате FP16/FP32. Операции Tensor Core MMA выполняются с низкой точностью, например FP8 × FP8 с накоплением более высокой точности или FP16 × FP16 с накоплением FP32.

Если создаётся собственный цикл по тайлам, можно перекрыть перемещение данных с вычислениями Tensor Core с помощью cuda::memcpy_async и CUDA Pipeline API. Пока Tensor Cores обрабатывают текущий тайл, cuda::memcpy_async потоково передаёт следующий тайл из HBM в shared memory. Это перекрытие помогает скрыть задержку памяти и поддерживать занятость Tensor Cores.

Ограничение UMMA на Blackwell: операнд A может поступать из TMEM или SMEM, операнд B — только из SMEM. Комбинация A из TMEM и B из TMEM не поддерживается. Это ограничение влияет на структуру ядра и требует планирования размещения данных.

XVIИспользование CUTLASS для оптимальной арифметической интенсивности

Один из самых простых способов задействовать описанные оптимизации — использовать библиотеку NVIDIA CUTLASS. С CUTLASS вы пишете один шаблонный вызов, и он автоматически применяет множество продвинутых оптимизаций: тайлинг в shared memory, асинхронные передачи памяти, двойную буферизацию с помощью буфера TMEM на 256 КБ на SM, warp specialization, кластеры блоков потоков и TMA multicast с распределённой shared memory.

Например, предположим, нужно вычислить GEMM C = A * B с входами половинной точности и выходом половинной точности, накапливая в FP16 или FP32:

#include <cutlass/numeric_types.h>
#include <cutlass/gemm/device/gemm.h>
using Gemm = cutlass::gemm::device::Gemm<
    cutlass::half_t,
    cutlass::layout::RowMajor,
    cutlass::half_t,
    cutlass::layout::ColumnMajor,
    cutlass::half_t,
    cutlass::layout::RowMajor,
    float,
    cutlass::arch::OpClassTensorOp,
    cutlass::arch::Sm100
>;
Gemm gemm_op;
cutlass::Status status = gemm_op(
    { M, N, K },
    float(1.0f),
    A_d, lda,
    B_d, ldb,
    float(0.0f),
    C_d, ldc
);

При компиляции и запуске этого кода CUTLASS автоматически выбирает тайлы, размещает их в shared memory, генерирует асинхронные копии памяти, использует инструкции Tensor Core tcgen05.mma, накапливает частичные суммы с указанной точностью, записывает результаты из TMEM коалесцированно, и использует кластеры блоков потоков, когда это выгодно.

CUTLASS соответствует или превосходит производительность ручного MMA-ядра в пределах около 2%. Небольшие увеличения использования регистров и shared memory связаны с тем, что 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++.

XVIIВстраиваемый PTX и тонкая настройка SASS

Для тех, кто готов выйти за пределы C++ и погрузиться в низкоуровневые микрооптимизации, CUDA позволяет использовать встраиваемый код Parallel Thread Execution (PTX) и SASS (ассемблер NVIDIA) для извлечения последних крупиц производительности.

PTX — это низкоуровневая виртуальная машина параллельного выполнения потоков. Высокоуровневый компилятор (например, CUDA C++) генерирует инструкции PTX, которые транслируются в машинные инструкции целевой архитектуры. SASS — это низкоуровневый язык ассемблера, который фактически выполняется нативно на оборудовании NVIDIA GPU.

Типичные сценарии для встраиваемого PTX:

• Использование инструкции GPU, для которой нет прямого интринсика CUDA.
• Применение модификаторов загрузки памяти (подсказок кэша) к конкретным доступам.
• Вставка барьеров памяти или синхронизации в точных точках.
• Ручная перестановка инструкций для избежания простоев конвейера.
• Чтение специальных регистров (SM ID, warp lane ID).

Пример предвыборки в L2:

__global__ void PrefetchExample(const float *in, float *out) {
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  asm volatile("cp.async.bulk.prefetch.L2.global [%0], %1;"
               :: "l"(in + idx + 32), "n"(128));
  float x = in[idx];
  out[idx] = x;
}

Управление поведением кэша через встраиваемый PTX:

asm("ld.global.cg.f32 %0, [%1];" : "=f"(val) : "l"(ptr))

Загрузка float с модификатором .cg («cache global») означает, что мы хотим кэшировать данные в L2, но обойти кэш L1.

Доступ к специальным регистрам:

asm("mov.u32 %0, %smid;" : "=r"(smid))

Некоторые разработчики использовали %smid в постоянных ядрах (persistent kernels), чтобы только один блок на SM выполнял определённую работу.

После настройки простои памяти ядра уменьшились за счёт перекрытия загрузок, а подсказки кэша снизили часть задержки. Чистое улучшение — 7% по общему времени выполнения ядра.

Встраиваемый PTX/SASS — инструмент оптимизации последней инстанции. Обеспечивает максимальный контроль ценой сложности. Применяется в последнюю очередь, когда нужна аппаратная функция, недоступная в CUDA C++, — или когда точно определён небольшой фрагмент кода, где планирование компилятора можно улучшить.

XVIIIИспользование inline PTX в DeepSeek для оптимизации доступа к памяти

Разбор инструкции

Разберём компоненты инструкции ld.global.nc.l1::no_allocate.l2::256b:

• ld.global — загрузка из глобальной памяти.
• .nc — некогерентная (non-coherent) загрузка, использующая текстурный кэш, который обеспечивает более высокую пропускную способность для доступа только для чтения.
• .l1::no_allocate — не выделять данные в кэше L1, чтобы не вытеснять горячие данные, находящиеся там.
• .l2::256b — извлекать по 256 байт за раз в L2, увеличивая эффективную ширину транзакции.

Обходя L1, загрузки могут потоково передавать большие блоки данных непосредственно в L2 без вытеснения часто используемых данных, находящихся в L1.

Контекст и риски

Эта инструкция не является частью официальной спецификации PTX ISA от NVIDIA (хотя позже было обнаружено, что модификаторы .L1::no_allocate и .L2::256B документированы в PTX ISA). Инженеры DeepSeek обнаружили её «вне документации» для точной настройки поведения кэша на их варианте Hopper GPU H800, ограниченном экспортными правилами США.

Ключевая проблема, которую решает инструкция: стандартная некогерентная загрузка (ld.global.nc) может возвращать устаревшие данные, поскольку текстурный кэш не поддерживает когерентность с глобальной памятью. Добавление .l1::no_allocate обходит L1, заставляя каждую загрузку извлекать свежие данные из L2 или глобальной памяти.

DeepEP также выделяет 20 SM из 132 на H800 исключительно для сетевого общения, чтобы L1-кэш этих SM оставался «тёплым» для сетевого кода. Это снижает вычислительную производительность на ~15%, но значительно повышает пропускную способность связи.

XIXПрактический чек-лист оптимизации ядра

Прежде чем переходить к низкоуровневым оптимизациям, убедитесь, что фундамент на месте:

1. Профилируйте с помощью Nsight Systems для общего обзора и Nsight Compute для детального анализа ядра.
2. Постройте Roofline — определите, memory-bound или compute-bound ваше ядро.
3. Коалесцируйте все глобальные загрузки — это базовое требование.
4. Тайлируйте данные в shared memory для повторного использования.
5. Сливайте поэлементные операции, чтобы исключить промежуточные записи в HBM.
6. Увеличьте размер батча для амортизации загрузки весов.
7. Используйте пониженную точность (TF32, FP16, BF16, FP8, FP4) где возможно.
8. Задействуйте Tensor Cores через CUTLASS или PyTorch SDPA.
9. Примените структурированную разреженность 2:4 для инференса после калибровки.
10. Рассмотрите пересчёт вместо хранения для memory-bound ядер.
11. Проверьте occupancy — не слишком ли много регистров? Нет ли spilling?
12. Только потом переходите к встраиваемому PTX/SASS.

Следуя этому систематическому рабочему процессу профилирования, диагностики и применения целевых оптимизаций — от настройки occupancy и уточнений на уровне warp-а до ILP и высокой арифметической интенсивности через тайлинг, слияние и Tensor Cores — можно превратить memory-bound ядро GPU в compute-bound. Это может дать большое ускорение, а в сильно memory-bound задачах иногда и на порядок.

Сначала используется системный профилировщик (например, NVIDIA Nsight Systems), чтобы соотнести активность CPU, GPU-ядра и трафик NVLink/NVSwitch по нескольким GPU. Затем Nsight Compute для глубокого погружения в детали каждого ядра.

Насыщение интерконнектов и пропускной способности памяти произойдёт задолго до того, как ядра смогут полностью использовать вычислительные возможности множества взаимосвязанных GPU (например, 72 в случае NVL72). Поэтому повышение арифметической интенсивности ядер — ключ к эффективному масштабированию для много-GPU задач обучения и инференса.

XXПроверка прогнозов и актуальные данные

Книга Fregly писалась в 2025 году. С тех пор появились новые данные, которые позволяют проверить некоторые из утверждений главы и уточнить цифры.

XXIДополнительные замечания, выходящие за пределы главы

Глава описывает методы оптимизации ядер, но оставляет за скобками несколько аспектов, которые проявились в реальной практике.

Сложность отладки на уровне PTX/SASS

Экосистема вокруг Tensor Cores

Влияние разреженности на обучение

Ограничения SDPA и FlashAttention

Memory-bound vs compute-bound в контексте NVL72

XXIIДополнительные выводы

1. Арифметическая интенсивность — центральная метрика. Она определяет режим работы ядра и направление оптимизации. Модель Roofline даёт визуальную основу для принятия решений.
2. Тайлинг и слияние — базовые приёмы. Они дают наибольший эффект при минимальных усилиях. Начинать нужно с них.
3. Понижение точности — множитель производительности. FP8 и FP4 дают кратный рост без изменения алгоритма. Но требуется калибровка и проверка точности.
4. Tensor Cores и TMEM меняют модель программирования. На Blackwell аккумуляторы живут в TMEM, а инструкции CTA-scoped. Это даёт больше гибкости, но требует понимания новой архитектуры.
5. CUTLASS — оптимальный выбор для большинства случаев. Он даёт производительность ручного ядра в пределах ~2% без недель настройки.
6. PTX/SASS — инструмент последней инстанции. Применять только после исчерпания высокоуровневых оптимизаций.
7. Контекст NVL72 меняет акценты. Внутри стойки коммуникации настолько быстры, что оптимизация ядер может быть даже важнее, чем для традиционных кластеров.

Источники: Chris Fregly, «AI Systems Performance Engineering» (O'Reilly, 2026), Глава 9 — цитаты приводятся в ограниченном объёме в целях обзора и критического анализа. NVIDIA CUTLASS documentation; NVIDIA Tensor Memory Accelerator documentation; NVIDIA Transformer Engine documentation; PyTorch documentation; MLPerf Training/Inference v5.0; открытые публикации DeepSeek.