Настоящий материал представляет собой обзор главы 18 книги 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).
Прямые цитаты из книги выделены отдельными блоками с пометкой «Цитата из книги». Остальной текст — систематизированное изложение содержания главы, сопоставление с актуальными данными из открытых источников (arXiv, ASPLOS 2025, документация PyTorch и NVIDIA), проверка прогнозов, а также технические детали, выходящие за пределы исходного текста. Цитаты приводятся в ограниченном объёме в целях обзора и критического анализа.
Эта глава продолжает тему дезагрегации инференса и углубляется в тонкую настройку фаз prefill и decode. Рассматриваются низкоуровневые приёмы: от «мегаядер» для одного шага декодирования до интеллектуального управления KV-кэшем, его совместного использования между GPU, быстрой передачи состояния промпта, адаптивного планирования ресурсов и динамической маршрутизации между prefill- и decode-воркерами. Также обсуждаются аппаратные и программные инновации, открывающие новые уровни производительности и эффективности.
Фаза декодирования распределена во времени и часто ограничена пропускной способностью памяти (memory-bound). В отличие от prefill, где обрабатывается сразу много токенов и GPU загружен вычислениями, на каждом шаге декодирования генерируется лишь один токен — арифметическая интенсивность операций мала, и узким местом становится чтение весов и KV-кэша из памяти. Это мотивировало исследователей и практиков сделать декодирование максимально быстрым и заточенным под конкретное оборудование.
Три заметные инновации в этой области: FlashMLA (DeepSeek), ThunderMLA (Stanford) и FlexDecoding (PyTorch). Все они нацелены на повышение эффективности многоголового внимания трансформера при декодировании в сценариях с переменной длиной последовательностей, типичных для LLM-нагрузок.
Flash Multi-Latent Attention, или FlashMLA — оптимизированное ядро декодирования, представленное DeepSeek. Оно фокусируется на шаге декодирования одного токена. FlashMLA ускоряет декодирование за счёт объединения операций и более эффективного использования иерархии памяти GPU.
FlashMLA для инференса — то же, что FlashAttention для обучения. Он снижает накладные расходы на доступ к памяти и задержку. С FlashMLA можно добиться значительного сокращения латентности фазы декодирования по сравнению со стандартными ядрами. FlashMLA повышает арифметическую интенсивность, объединяя несколько операций внимания в одну.
FlashMLA особенно эффективен для моделей с архитектурой Multi-Latent Attention (MLA), используемой в семействе DeepSeek-V3/R1. В MLA KV-кэш хранится не в виде полноразмерных векторов ключей и значений, а в сжатом латентном пространстве, что кардинально уменьшает объём данных, читаемых из памяти на каждом шаге декодирования. Это даёт преимущество по сравнению с групповыми вариантами внимания (GQA) и multi-query attention (MQA), особенно на GPU архитектуры Hopper (H100), где пропускная способность HBM является ограничивающим фактором.
Согласно опубликованным данным, плотное MLA-ядро декодирования достигает до 3000 ГБ/с в memory-bound конфигурации и 580 TFLOPS в compute-bound конфигурации на H800 SXM5 с CUDA 12.6[reference:0]. В DeepSeek-V3.2 было представлено FP8 KV-кэш для поддержки контекста до 128K токенов, а токен-уровневое разреженное MLA-ядро декодирования использует FP8 KV-кэш с умножением матриц в bfloat16.
Roofline-график: MQA → GQA → MLA (Hopper H100)
Развивая идеи FlashMLA, исследователи из Стэнфорда представили ThunderMLA — полностью слитое «мегаядро» внимания для декодирования. Это «мегаядро» снижает накладные расходы на запуск и хвостовые эффекты, объединяя несколько запусков ядер в один, а также консолидируя промежуточные записи в память. ThunderMLA сообщает о 20–35% ускорении пропускной способности декодирования по сравнению с FlashMLA на различных нагрузках.
Ключевая идея ThunderMLA в том, что при декодировании последовательностей разной длины использование тонкого планирования и слитых операций позволяет избежать хвостового эффекта — ситуации, когда одни последовательности завершаются раньше, а другие оставляют GPU частично простаивающим. ThunderMLA поддерживает загрузку GPU, динамически упаковывая и обрабатывая оставшиеся потоки.
FlexAttention от PyTorch позволяет JIT-компилировать слитые ядра для произвольных паттернов разреженности внимания. Под капотом TorchInductor + OpenAI Triton генерируют слитое ядро, вычисляющее только разрешённые пары query-key для данного паттерна. FlexDecoding — это backend декодирования для torch.nn.attention.flex_attention, который компилирует специализированное ядро для фазы декодирования (Q_len=1) и переиспользует его на нескольких шагах декодирования.
Практические рекомендации по FlexDecoding:
torch.compile(mode="max-autotune") для стабильного, latency-критичного декодирования.torch.set_float32_matmul_precision("high") для включения fallback на TF32.Ключевая особенность FlexAttention и FlexDecoding — поддержка nested jagged-layout tensors (NJT). Они позволяют ragged-батчинг последовательностей переменной длины (обычное дело в LLM-нагрузках) во время декодирования.
Три последовательности переменной длины → один плоский буфер с offsets
| Критерий | FlashMLA | ThunderMLA | FlexDecoding |
|---|---|---|---|
| Модель | DeepSeek MLA | Любая с MLA | Любая |
| Ускорение | Значительное | +20–35% vs FlashMLA | Зависит от паттерна |
| Гибкость | Низкая | Средняя | Высокая (Python) |
| Кастомные маски | Нет | Нет | Да |
| Требует CUDA | Да (библиотека) | Да (библиотека) | Нет (Triton JIT) |
| Лучший сценарий | DeepSeek в продакшене | Максимум пропускной способности | MoE, нестандартные маски |
Дезагрегация требует рассматривать KV-кэш как первоклассный разделяемый ресурс всего кластера. Поскольку KV-кэши теперь могут жить дольше и перемещаться между узлами, высокопроизводительные системы инференса улучшили способы хранения и совместного использования KV-кэша.
Вместо того чтобы каждый GPU хранил KV только для запросов, которые он обслуживает в данный момент, дезагрегированный пул KV-кэша отделяет хранение KV от отдельных GPU. Данные распределяются по всей памяти кластера. Пул также может выгружаться в память CPU (включая унифицированную память CPU-GPU платформ Grace Blackwell и Vera Rubin) и в постоянное хранилище, такое как NVMe SSD.
Пул KV-кэша позволяет выгружать старые KV-блоки в DRAM CPU или NVMe, переиспользовать префиксы между запросами и восстанавливаться после отказов. vLLM реализует автоматическое префиксное кэширование, используя глобальную хеш-таблицу KV-«страниц» через механизм PagedAttention. Каждый уникальный 16-токенный блок контекста имеет хеш. Если новому запросу нужен префикс, совпадающий с сохранённым блоком, он может напрямую скопировать эти KV-тензоры вместо перевычисления.
SGLang’s RadixAttention использует сжатое префиксное дерево (radix tree) над целыми последовательностями токенов. Этот тип префиксного кэша сворачивает последовательности токенов в отдельные рёбра дерева для экономии места. Каждый узел в дереве указывает на KV-тензор, хранящийся как непрерывная GPU-страница, которая содержит KV-кэш префикса.
Приведённый расчёт справедлив для полноголового MHA. Многие современные 70B-модели (например, Llama 3) используют Grouped-Query Attention (GQA) с 8 KV-головами вместо 32, что уменьшает объём KV-кэша в 4 раза — до ~82 ГБ для тех же 250 000 токенов в FP16. Архитектуры MLA сжимают KV ещё сильнее.
KV-кэш живёт вне отдельного GPU — разделяется по кластеру через многоуровневую иерархию
Выгодно переиспользовать кэшированные KV-данные между запросами, которые разделяют общий префикс. Этот сценарий возникает часто: многоходовые беседы, общие системные промпты, прикреплённые документы. Вместо перевычисления выходов внимания трансформера для префикса каждого запроса система может сохранить KV-выходы для префикса и переиспользовать их напрямую.
vLLM реализует автоматическое префиксное кэширование, используя глобальную хеш-таблицу KV-«страниц» через механизм PagedAttention. Каждый уникальный 16-токенный блок контекста имеет хеш. Если новому запросу нужен префикс, совпадающий с сохранённым блоком, он может напрямую скопировать эти KV-тензоры вместо перевычисления.
LMCache — это открытый слой KV-кэширования, который располагается между vLLM (или SGLang) и иерархией хранения. Несколько экземпляров vLLM могут совместно использовать один пул KV-кэша. LMCache использует NIXL в качестве основы для передачи KV. Согласно arXiv:2510.09665, комбинация LMCache с vLLM обеспечивает до 15-кратного увеличения пропускной способности на таких нагрузках, как многораундовые вопросы-ответы и анализ документов, а также как минимум 2-кратное снижение задержки в различных сценариях[reference:1][reference:2].
Логические KV-блоки распределяются в физический пул без лишних копий
KV-кэш, хранящий ключи и значения для всех прошлых токенов, может стать огромным для многих одновременных потоков декодирования, поскольку каждый поток использует память примерно пропорционально num_layers × 2 × sequence_length × d_head. FlashMLA странично организует KV-кэш и выделяет его фиксированными блоками, чтобы для активных последовательностей происходили непрерывные доступы к памяти.
POD-Attention реорганизует вычисление внимания для уменьшения трафика HBM. Она использует планирование thread-block (CTA) с учётом SM. Это реализует привязку операций во время выполнения для динамического назначения каждой CTA, работающей на SM, либо задачи prefill, либо задачи decode. Согласно статье на ASPLOS 2025, POD-Attention ускоряет вычисление внимания до 59% (в среднем 28%) по сравнению с использованием независимо оптимизированных ядер внимания prefill и decode из FlashAttention и FlashInfer[reference:3]. Это первое GPU-ядро, которое эффективно вычисляет внимание для гибридных батчей, ускоряя вычисление внимания до 59% (в среднем 28%).
Prefill и decode задачи размещаются на одних SM — минимизируется движение памяти
Система NVIDIA Grace Blackwell GB200 NVL72 — стоечная платформа с 36 процессорами Grace и 72 GPU Blackwell — позволяет создать единый логический decode-блок с десятками терабайт памяти для KV-кэша. Согласно официальным спецификациям, GB200 NVL72 предоставляет 13,5 ТБ HBM3e для GPU и ~17 ТБ LPDDR5X для 36 процессоров Grace, что в сумме даёт около 30,5 ТБ быстрой памяти[reference:4][reference:5]. Однако для latency-критичного декодирования активные ключи и значения всё равно должны жить в HBM GPU, а память Grace CPU используется как кэш более низкого уровня.
Ключевое требование дезагрегированного инференса — быстрая и эффективная передача KV-кэша от prefill-воркера к decode-воркеру. Если эта передача не быстра, любое время, сэкономленное за счёт параллелизации prefill и decode, может быть потеряно на ожидание перемещения данных.
Выход prefill состоит в основном из KV-кэша для всех токенов промпта. Размер KV-кэша примерно равен 2 × L × N × (h × d) × bytes_per_val, где множитель 2 — для ключей и значений. Наивный подход — копирование KV в память CPU, затем отправка по TCP — может занимать сотни миллисекунд для больших промптов. Цель — сократить время передачи до нескольких миллисекунд.
Современные дезагрегированные системы используют техники zero-copy передачи GPU-to-GPU. На практике это включает использование remote direct memory access (RDMA) по высокоскоростным фабрикам: InfiniBand для передачи между стойками/узлами или NVLink/NVSwitch для прямых записей в память GPU внутри одного узла. Эти методы отправляют данные напрямую между GPU без копирования через память CPU.
Высокопроизводительная библиотека NVIDIA для передачи GPU-to-GPU при инференсе называется NVIDIA Inference Xfer Library (NIXL). NIXL обеспечивает архитектуру плагинов (NVLink, фабрики UCX, GPUDirect Storage) для zero-copy перемещения данных GPU↔GPU и GPU↔хранилище.
Общие стратегии быстрой передачи KV:
PagedAttention от vLLM хранит KV-кэш фиксированными блоками, обычно по 16 токенов на блок. Наивная отправка тысяч маленьких KV-страниц по RDMA повлечёт чрезмерные накладные расходы. LMCache решает эту неэффективность, объединяя KV-страницы в большие непрерывные буферы перед передачей. Например, если отправка KV-кэша на 7500 токенов в виде 470 маленьких передач занимает 20 мс, объединение их в более крупные блоки (страницы по 128 токенов) сокращает время передачи до 8 мс.
Пример конфигурации канала передачи NIXL для режима prefill-decode LMCache:
# lmcache-prefiller-config.yaml
enable_pd: true
transfer_channel: "nixl"
pd_role: "sender"
pd_proxy_host: "decode-host"
pd_proxy_port: 7500
pd_buffer_size: 1073741824
pd_buffer_device: "cuda"
# lmcache-decoder-config.yaml
enable_pd: true
transfer_channel: "nixl"
pd_role: "receiver"
pd_peer_host: "0.0.0.0"
pd_peer_init_port: 7300
pd_peer_alloc_port: 7400
pd_buffer_size: 1073741824
pd_buffer_device: "cuda"
nixl_backends: [UCX]NVIDIA Dynamo реализует глобальную очередь планирования, в которую decode-воркеры помещают новые задачи промптов, а prefill-воркеры их потребляют. Decode-узел ставит в очередь запрос на обработку промпта. Prefill-узел подхватывает этот запрос и, когда закончит, точно знает, какому decode-узлу отправить результаты, поскольку запрос несёт origin или reply-to ID. KV затем передаётся напрямую в GPU decode-воркера с помощью NIXL RDMA.
Decode тянет промпты из очереди · Prefill отправляет KV через NIXL
В реализации vLLM + LMCache используется более децентрализованный подход. Decode- и prefill-процессы устанавливают прямой канал с использованием pipe или буфера для KV каждого запроса. Под капотом это может использовать одноранговое TCP или RDMA-соединение, согласованное в начале запроса.
| Критерий | Глобальная очередь (Dynamo) | Прямые каналы (vLLM+LMCache) |
|---|---|---|
| Балансировка нагрузки | Надёжная многопользовательская | Требует согласования |
| Failover | Простой | Сложнее |
| Хвостовая латентность | Небольшая задержка очереди | Минимум джиттера |
| Лучший сценарий | Мультитенантность | Жёсткие SLO |
Одно из мощных преимуществ дезагрегации — свобода выбирать разное оборудование и разную конфигурацию параллелизма, которая лучше всего подходит для нужд prefill- и decode-кластеров по отдельности.
Фаза prefill выигрывает от GPU с высокой вычислительной пропускной способностью. Фаза decode, напротив, выигрывает от большой ёмкости памяти и пропускной способности памяти. Это открывает возможность использования разных поколений GPU: новейшие высокопроизводительные GPU для prefill-кластера, более старые или экономичные GPU для decode-кластера.
В исследовании Splitwise использование оборудования, специфичного для фазы, привело к конфигурации, достигающей в 1,4 раза более высокой пропускной способности при на 20% меньшей стоимости по сравнению с гомогенным базовым уровнем. Конкретно: 4× H100 для prefill и 4× A100 для decode дали ~2,35× RPS по сравнению с 8-GPU гомогенной системой при той же стоимости/мощности. Ещё одна система — HexGen-2 — рассматривает распределение дезагрегированного инференса на гетерогенных GPU как задачу оптимизации. На моделях вроде Llama 2 70B HexGen-2 показывает до 2× улучшения пропускной способности при той же ценовой точке.
Оптимальная стратегия параллелизма для prefill может не совпадать с таковой для decode. Фаза prefill — большой прямой проход через N токенов промпта, выигрывает от высокой степени параллелизации. Фаза decode, напротив, последовательна и latency-чувствительна на каждом шаге. Использование слишком многих GPU для одного decode может навредить латентности на выходной токен (TPOT).
Prefill — pipeline parallel (PP) · Decode — tensor parallel (TP)
Некоторые inference-движки позволяют использовать разную точность между фазами. Например, prefill в FP8 для ускорения, decode в FP16 для лучшей точности генерации. Обычно обе фазы выполняются в одной точности, чтобы KV-кэш был совместим. Однако можно применить преобразование: квантовать KV и сжать с FP16 до INT8/FP8/FP4 перед отправкой, затем преобразовать обратно на принимающей стороне.
При экстремальных масштабах стоит оценить, могут ли CPU снять нагрузку с GPU. Современные CPU гораздо медленнее GPU для вычислений нейронных сетей, но у них есть большой объём RAM и отсутствие конкуренции за пропускную способность памяти GPU. Практическое применение — обработка фоновых или низкоприоритетных задач prefill. Разгрузка CPU увеличила бы TTFT, поэтому она не используется для latency-чувствительных запросов.
Раннее отклонение означает: если система предсказывает, что не сможет обслужить запрос в рамках целевой латентности, она быстро отказывает, вернув ошибку или ответив «попробуйте позже». Это сохраняет goodput, гарантируя, что обслуживаемые запросы соответствуют гарантиям SLO. В Mooncake стратегия раннего отклонения оценивает входящие запросы на основе оценочной нагрузки как на prefill-, так и на decode-кластерах.
Без admission control система уходит в оверлоад · с early rejection — стабильна
Некоторые запросы могут иметь уровни приоритета: платящие клиенты — выше, чем бесплатные; интерактивные — выше, чем офлайн-пакетные. Планировщик может приоритезировать определённые запросы, чтобы всегда соответствовать SLO. Дезагрегация может помочь, выделив часть prefill- и decode-воркеров для высокоприоритетных задач.
В дезагрегированной системе, если decode-узел падает в середине генерации, пул KV-кэша может обеспечить восстановление на другом узле, поскольку KV был сохранён в пуле. Если KV не был сохранён, системе придётся пересчитать prefill на другом узле. Именно поэтому системы вроде vLLM периодически делают контрольные точки или копируют KV-кэш в пул. Дезагрегация обеспечивает чёткие точки контроля: можно наблюдать метрики prefill и decode отдельно и применять admission control именно там, где это необходимо.
При чисто статической конфигурации, если микс нагрузки меняется, фиксированное соотношение станет субоптимальным. Адаптивные механизмы планирования направлены на непрерывную настройку системы, чтобы ни одна фаза не стала узким местом.
Исследовательский прототип TetriInfer работает на двух уровнях детализации. Во-первых, на уровне отдельного запроса он назначает входящие запросы конкретным prefill- и decode-инстансам на основе текущей нагрузки. Во-вторых, он отслеживает утилизацию ресурсов по всему кластеру и прогнозирует, где могут сформироваться узкие места, проактивно смещая работу для предотвращения hotspots. TetriInfer разделяет промпты на чанки фиксированного размера, дезагрегирует prefill и decode и использует двухуровневый алгоритм планирования с прогнозируемым использованием ресурсов, чтобы избежать hotspots планирования decode.
Предсказание узких мест и балансировка
Arrow непрерывно анализирует нагрузку, измеряя скорость входных токенов versus скорость выходных токенов и backlog в каждом пуле воркеров. Затем он динамически корректирует распределение воркеров. Arrow рассматривает количество prefill- и decode-воркеров как настраиваемый параметр, который можно корректировать почти в реальном времени. Arrow использует stateless-инстансы, каждый из которых может обрабатывать как prefill-, так и decode-запросы без выделенных ролей. Система использует SLO-осведомлённый алгоритм планирования, где глобальный планировщик динамически корректирует диспетчеризацию запросов и распределение инстансов на основе настроек SLO и нагрузки на инстансы. Arrow демонстрирует до 5,62× более высокую скорость обслуживания запросов по сравнению с современными PD-дезагрегированными системами.
Непрерывный анализ потоков токенов → динамическое перераспределение инстансов
Mooncake представил admission control на основе прогнозирования — раннее отклонение. Если система прогнозирует, что у неё недостаточно ёмкости для обработки нового запроса в рамках SLO, она отклонит этот запрос. Адаптация может происходить на стороне предложения (добавление ресурсов) и на стороне спроса (throttling или отклонение запросов).
Сложные системы инференса используют прогнозные модели вроде ARIMA и другие техники прогнозирования для предвидения всплесков и сдвигов трафика с использованием паттернов времени суток. Если прогнозируется всплеск длинных выходов в 9 вечера, планировщик может упреждающе выделить больше ёмкости decode как раз вовремя.
Конкретные способы балансировки нагрузки: эластичные инстансы (Kubernetes HPA), «instance flip» (TetriInfer), statelessness для эластичности (Arrow). Быстрое переключение ролей может привести к осцилляции и thrashing, поэтому обычно вводится минимальное время, в течение которого роль должна сохраняться.
«Дезагрегация prefill и decode — не просто архитектурный паттерн, а фундамент для всего спектра оптимизаций».
Подтверждено. К 2026 году дезагрегация prefill и decode стала стандартной практикой для высокомасштабного инференса LLM. vLLM, SGLang, NVIDIA Dynamo и TensorRT-LLM поддерживают её нативно. Открытые проекты вроде llm-d обеспечивают Kubernetes-оркестрацию для дезагрегированного обслуживания.
«NIXL обеспечивает high-throughput, low-latency point-to-point communication для крупномасштабного распределённого инференса».
Подтверждено. NIXL интегрирован в vLLM, SGLang и NVIDIA Dynamo как основной транспорт для KV-кэша. LMCache использует NIXL в качестве основы для передачи KV.
«Adaptive scheduling and SLO-aware control are necessary to maintain latency guarantees under varying load».
Подтверждено. TetriInfer, Arrow и Mooncake демонстрируют эффективность адаптивного планирования. Kubernetes HPA и кастомные планировщики внедряют SLO-осведомлённые политики.
«POD-Attention — первое GPU-ядро, эффективно вычисляющее внимание для гибридных батчей».
Частично подтверждено. POD-Attention показала ускорение до 59% (в среднем 28%) на ASPLOS 2025. Однако её массовое внедрение в продакшене пока ограничено — требуются дополнительные исследования и интеграция в популярные inference-движки.
Ограничения динамических форм. Глава описывает динамические формы и CUDA Graphs, но не рассматривает компромиссы. CUDA Graphs требуют статических форм, и при включении dynamic=True они не используются. Для переменных длин последовательностей рекомендуется mode="default" или mode="max-autotune-no-cudagraphs".
Сложность отладки компилятора. Глава описывает инструменты отладки, но не подчёркивает сложность диагностики проблем компилятора. Graph breaks, рекомпиляции, проблемы с числовой корректностью — всё это требует глубокого понимания внутреннего устройства компилятора. Инструменты вроде TORCH_LOGS и torch._dynamo.explain() помогают, но кривая обучения остаётся высокой.
Взаимодействие с TensorRT и другими бэкендами. Глава фокусируется на TorchInductor и XLA, но не рассматривает взаимодействие с TensorRT. В production-средах TensorRT часто используется для инференса, и torch.compile может интегрироваться с ним. Это важный аспект для production-развёртываний, который остаётся за пределами главы.
Экосистема Triton. Глава описывает Triton, но не упоминает экосистему вокруг него. Triton активно развивается, появляются новые возможности: TLX (Triton Language Extensions) для низкоуровневого контроля, поддержка Blackwell с инструкциями tcgen05, MXFP4/MXFP8. Сообщество GPU MODE активно способствует развитию Triton.
MLPerf Logging. Для стандартизированного логирования производительности в распределённых нагрузках используется библиотека mlperf_logging. Она позволяет записывать ключевые события (начало обучения, конец эпохи, финальная точность) в формате, совместимом с MLPerf.
Основной источник: Chris Fregly, «AI Systems Performance Engineering: Optimizing Model Training and Inference Workloads with GPUs, CUDA, and PyTorch», O'Reilly Media, 2026, глава 18.
FlashMLA: DeepSeek Open Source Week, февраль 2025. Официальные бенчмарки: H800 SXM5, CUDA 12.6, 3000 ГБ/с (memory-bound), 580 TFLOPS (compute-bound). GitHub: github.com/deepseek-ai/FlashMLA.
ThunderMLA: Stanford University, исследование полностью слитого мегаядра внимания для декодирования. Отчёт о 20–35% ускорении относительно FlashMLA.
FlexDecoding: Документация PyTorch по torch.nn.attention.flex_attention и FlexDecoding. Поддержка nested jagged tensors и PagedAttention.
POD-Attention: «POD-Attention: Unlocking Full Prefill-Decode Overlap for Faster LLM Inference», ASPLOS 2025. Ускорение вычисления внимания до 59% (в среднем 28%) относительно FlashAttention и FlashInfer.
LMCache: «LMCache: An Efficient KV Cache Layer for Enterprise-Scale LLM Inference», arXiv:2510.09665, октябрь 2025. До 15-кратного увеличения пропускной способности с vLLM, как минимум 2-кратное снижение задержки.
NVIDIA GB200 NVL72: Официальные спецификации NVIDIA. 13,5 ТБ HBM3e (72 GPU × 192 ГБ), 17 ТБ LPDDR5X (36 Grace CPU × 480 ГБ), ~30,5 ТБ суммарной быстрой памяти.
NIXL: NVIDIA Inference Xfer Library. Документация NVIDIA. Интеграция с vLLM, SGLang, NVIDIA Dynamo.
NVIDIA Dynamo: Документация NVIDIA. Глобальная очередь планирования, NIXL-передача KV, интеграция с Kubernetes.
TetriInfer: «TetriInfer: Decode Scheduling and Resource Management for LLM Inference», исследовательский прототип. Двухуровневый планировщик с прогнозируемым использованием ресурсов.
Arrow: «Arrow: Adaptive Instance Scaling for Disaggregated LLM Inference». Stateless-инстансы, SLO-осведомлённый алгоритм планирования. До 5,62× более высокая скорость обслуживания запросов.
Mooncake: «Mooncake: A KVCache-Centric Architecture for Cloud-Based LLM Serving». Prediction-driven admission control.
Splitwise: «Splitwise: Efficient Generative LLM Inference Using Phase Splitting». Гетерогенное оборудование для prefill и decode. 1,4× пропускной способности при 20% меньшей стоимости; до 2,35× RPS на 4× H100 + 4× A100.
HexGen-2: «HexGen-2: Heterogeneous LLM Inference Deployment». До 2× пропускной способности на Llama 2 70B при той же ценовой точке.
MLPerf: MLPerf Training/Inference v5.0, 2025. Стандартизированные бенчмарки для сравнения производительности.
Ультрамасштабный инференс LLM требует целостного подхода, включающего как высокоуровневое адаптивное управление ресурсами, так и низкоуровневые оптимизации ядер и памяти. Комбинируя техники, представленные в этой главе, высокооптимизированное развёртывание инференса может достичь максимальной пропускной способности на современном оборудовании при соблюдении строгих гарантий латентности.
По мере дальнейшей эволюции оборудования и GPU и усложнения программных фреймворков эти оптимизации будут накапливаться и становиться ещё более мощными. Дезагрегация prefill и decode — не просто архитектурный паттерн, а фундамент для всего спектра оптимизаций: от выбора ядра до планирования кластера.
Каждый уровень выигрывает от независимости фаз, и именно их сочетание — от слитых мегаядер до адаптивного масштабирования — делает ультрамасштабный инференс экономически жизнеспособным.
This material is a review of Chapter 18 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 from open sources (arXiv, ASPLOS 2025, PyTorch and NVIDIA documentation), forecast verification, and technical details that go beyond the original text. Quotes are provided in limited volume for review and critical analysis purposes.
This chapter continues the theme of inference disaggregation and dives deep into fine-tuning the prefill and decode phases. It covers low-level techniques: from decoding megakernels to intelligent KV cache management, cross-GPU sharing, fast prompt state transfer, adaptive resource scheduling, and dynamic routing between prefill and decode workers. Hardware and software innovations that unlock new levels of performance and efficiency are also discussed.
The decoding phase is spread over time and is often memory-bound. Unlike prefill, which processes many tokens at once and keeps the GPU compute-saturated, each decoding step generates only one token — arithmetic intensity is low, and the bottleneck becomes reading weights and the KV cache from memory. This motivated researchers and practitioners to make decoding as fast and hardware-specific as possible.
Three notable innovations in this area: FlashMLA (DeepSeek), ThunderMLA (Stanford), and FlexDecoding (PyTorch). All target efficient multi-head transformer attention during decoding in variable-length scenarios typical of LLM workloads.
Flash Multi-Latent Attention (FlashMLA) is an optimized decoding kernel from DeepSeek. It focuses on the single-token decoding step. FlashMLA accelerates decoding by combining operations and using the GPU memory hierarchy more efficiently. FlashMLA is to inference what FlashAttention is to training. It reduces memory access overhead and latency. FlashMLA raises arithmetic intensity by combining several attention operations into one fused kernel launch.
FlashMLA is especially effective for Multi-Latent Attention (MLA) architectures used in the DeepSeek-V3/R1 family. In MLA, the KV cache is stored in a compressed latent space rather than full-size key/value vectors, dramatically reducing data read per decoding step. This gives an advantage over GQA and MQA, particularly on Hopper (H100), where HBM bandwidth is the limiting factor.
According to published data, the dense MLA decoding kernel achieves up to 3000 GB/s in a memory-bound configuration and 580 TFLOPS in a compute-bound configuration on H800 SXM5 with CUDA 12.6[reference:6]. DeepSeek-V3.2 introduced an FP8 KV cache to support contexts up to 128K tokens, and the token-level sparse MLA decoding kernel uses an FP8 KV cache while performing matrix multiplication in bfloat16.
Roofline plot: MQA → GQA → MLA (Hopper H100)
Building on FlashMLA, Stanford researchers introduced ThunderMLA — a fully fused attention "megakernel" for decoding. It cuts launch overhead and tail effects by merging multiple kernel launches into one. ThunderMLA reports 20–35% higher decoding throughput than FlashMLA across various workloads. The key idea is that when decoding variable-length sequences, fine-grained scheduling and fused operations avoid the tail effect where some sequences finish early while others leave the GPU partially idle.
PyTorch's FlexAttention JIT-compiles fused kernels for arbitrary attention sparsity patterns. Under the hood, TorchInductor + OpenAI Triton generate a fused kernel that computes only the allowed query-key pairs for the given pattern. FlexDecoding is the decoding backend for torch.nn.attention.flex_attention; it compiles a specialized decoding kernel (Q_len=1) and reuses it across decoding steps.
Practical FlexDecoding guidance:
torch.compile(mode="max-autotune") for stable, latency-critical decoding.torch.set_float32_matmul_precision("high") to enable the TF32 fallback.A key feature of FlexAttention and FlexDecoding is support for nested jagged-layout tensors (NJT). They enable ragged batching of variable-length sequences (common in LLM workloads) during decoding.
Three variable-length sequences → one flat buffer with offsets
| Criterion | FlashMLA | ThunderMLA | FlexDecoding |
|---|---|---|---|
| Model | DeepSeek MLA | Any with MLA | Any |
| Speedup | Significant | +20–35% vs FlashMLA | Pattern-dependent |
| Flexibility | Low | Medium | High (Python) |
| Custom masks | No | No | Yes |
| Needs CUDA | Yes (library) | Yes (library) | No (Triton JIT) |
| Best scenario | DeepSeek in production | Maximum decode throughput | MoE, custom masks |
Disaggregation requires treating the KV cache as a first-class, cluster-wide shared resource. Because KV caches can now outlive and move between nodes, high-performance inference systems have improved how the KV cache is stored and shared.
Instead of each GPU holding KV only for the requests it currently serves, a disaggregated KV cache pool decouples KV storage from individual GPUs. Data is distributed across the cluster's memory. The pool can also be offloaded to CPU memory (including unified CPU-GPU memory on Grace Blackwell and Vera Rubin platforms) and to persistent storage such as NVMe SSD.
The KV pool allows offloading old KV blocks to CPU DRAM or NVMe, reusing prefixes across requests, and recovering from failures. vLLM implements automatic prefix caching using a global hash table of KV "pages" through PagedAttention. Every unique 16-token context block has a hash. If a new request needs a prefix that matches a stored block, it can directly copy those KV tensors instead of recomputing.
SGLang's RadixAttention uses a compressed prefix tree (radix tree) over entire token sequences. This type of prefix cache collapses token sequences into single edges in the tree to save space. Each node in the tree points to a KV tensor stored as a contiguous GPU page that holds the prefix's KV cache.
This calculation is for full MHA. Many modern 70B models (e.g., Llama 3) use Grouped-Query Attention (GQA) with 8 KV heads instead of 32, cutting KV size 4× — down to ~82 GB for the same 250,000 tokens in FP16. MLA architectures compress KV even further.
KV cache lives outside a single GPU — shared across the cluster through a multi-tier hierarchy
It is beneficial to reuse cached KV data between requests that share a common prefix. This scenario is frequent: multi-turn conversations, shared system prompts, attached documents. Instead of recomputing transformer attention outputs for each request's prefix, the system can store the prefix KV outputs and reuse them directly.
vLLM implements automatic prefix caching using a global hash table of KV "pages" through PagedAttention. Every unique 16-token context block has a hash. If a new request needs a prefix that matches a stored block, it can directly copy those KV tensors instead of recomputing.
LMCache is an open-source key-value caching layer that sits between vLLM (or SGLang) and your storage hierarchy. Multiple vLLM instances can share one KV cache pool. LMCache uses NIXL as the underlying KV transmission. According to arXiv:2510.09665, combining LMCache with vLLM achieves up to 15× improvement in throughput across workloads such as multi-round question answering and document analysis, and at least 2× lower latency across diverse settings[reference:7][reference:8].
Logical KV blocks are mapped into a physical pool without extra copies
The KV cache, storing keys and values for all past tokens, can become huge for many concurrent decoding streams, since each stream uses memory roughly proportional to num_layers × 2 × sequence_length × d_head. FlashMLA organizes the KV cache in pages and allocates cache in fixed blocks so active sequences get contiguous memory accesses.
POD-Attention reorganizes attention computation to reduce HBM traffic. It uses SM-aware thread-block (CTA) scheduling, binding operations at runtime to dynamically assign each CTA on an SM either prefill or decode work. According to the ASPLOS 2025 paper, POD-Attention speeds up attention computation by up to 59% (mean 28%) compared to using independently optimized prefill and decode attention kernels from FlashAttention and FlashInfer[reference:9]. This is the first GPU kernel that efficiently computes attention for hybrid batches, speeding up attention computation by up to 59% (mean 28%).
Prefill and decode tasks are placed on the same SMs to minimize memory movement
NVIDIA Grace Blackwell GB200 NVL72 — a rack platform with 36 Grace CPUs and 72 Blackwell GPUs — allows a single logical decode block with tens of terabytes of memory for the KV cache. According to official specifications, the GB200 NVL72 provides 13.5 TB HBM3e for GPUs and ~17 TB LPDDR5X for the 36 Grace CPUs, totaling around 30.5 TB of fast memory[reference:10][reference:11]. However, for latency-critical decoding, active keys and values must still reside in GPU HBM, using Grace CPU memory as a lower-tier cache.
A key requirement of disaggregated inference is fast, efficient KV cache transfer from prefill worker to decode worker. If the transfer is not fast, any time saved from parallelizing prefill and decode can be lost waiting for data movement.
The prefill output is mainly the KV cache for all prompt tokens. The KV cache size is approximately 2 × L × N × (h × d) × bytes_per_val, where factor 2 is for keys and values. A naive approach — copying KV to CPU memory and sending over TCP — can take hundreds of milliseconds for large prompts. The goal is to reduce transfer time to a few milliseconds.
Modern disaggregated systems use zero-copy GPU-to-GPU transfer techniques. In practice this involves remote direct memory access (RDMA) over high-speed fabrics: InfiniBand for cross-rack / cross-node transfers, or NVLink/NVSwitch for direct writes to GPU memory within a node. These methods send data directly between GPUs without copying through CPU memory.
NVIDIA's high-performance library for GPU-to-GPU transfer in inference is the NVIDIA Inference Xfer Library (NIXL). NIXL provides a plugin architecture (NVLink, UCX fabrics, GPUDirect Storage) for zero-copy GPU↔GPU and GPU↔storage data movement.
Common fast KV transfer strategies:
vLLM's PagedAttention stores the KV cache in fixed blocks, typically 16 tokens per block. Naively sending thousands of small KV pages over RDMA incurs excessive overhead. LMCache addresses this inefficiency by packing KV pages into large contiguous buffers before transfer. For example, if sending a 7500-token KV cache as 470 small transfers takes 20 ms, packing into larger blocks (pages of 128 tokens) reduces the transfer time to 8 ms.
Example NIXL transfer channel configuration for LMCache prefill-decode mode:
# lmcache-prefiller-config.yaml
enable_pd: true
transfer_channel: "nixl"
pd_role: "sender"
pd_proxy_host: "decode-host"
pd_proxy_port: 7500
pd_buffer_size: 1073741824
pd_buffer_device: "cuda"
# lmcache-decoder-config.yaml
enable_pd: true
transfer_channel: "nixl"
pd_role: "receiver"
pd_peer_host: "0.0.0.0"
pd_peer_init_port: 7300
pd_peer_alloc_port: 7400
pd_buffer_size: 1073741824
pd_buffer_device: "cuda"
nixl_backends: [UCX]NVIDIA Dynamo implements a global scheduling queue into which decode workers post new prompt tasks and prefill workers consume them. A decode node enqueues a prompt processing request. A prefill node picks it up and, when finished, knows exactly which decode node to send results to, because the request carries an origin or reply-to ID. The KV is then transferred directly to the decode worker GPU via NIXL RDMA.
Decode pulls prompts from the queue · Prefill pushes KV via NIXL
vLLM + LMCache uses a more decentralized approach. Decode and prefill processes establish a direct channel using a pipe or buffer per request's KV. Under the hood, this can use peer-to-peer TCP or an RDMA connection agreed at request start.
| Criterion | Global queue (Dynamo) | Direct channels (vLLM+LMCache) |
|---|---|---|
| Load balancing | Robust multi-tenant | Requires negotiation |
| Failover | Simple | Harder |
| Tail latency | Small queue delay | Minimal jitter |
| Best scenario | Multi-tenant | Tight SLOs |
One of the powerful advantages of disaggregation is the freedom to choose different hardware and parallelism configurations that best suit the prefill and decode clusters separately.
Prefill benefits from GPUs with high compute throughput. Decode benefits from large memory capacity and memory bandwidth. This opens the possibility of using different GPU generations: newest high-performance GPUs for the prefill cluster, older or cost-effective GPUs for the decode cluster.
In the Splitwise study, phase-specific hardware achieved 1.4× higher throughput at 20% lower cost compared to a homogeneous baseline. Specifically: 4× H100 for prefill and 4× A100 for decode delivered ~2.35× RPS compared to an 8-GPU homogeneous system at the same cost/power. Another system, HexGen-2, treats disaggregated inference distribution on heterogeneous GPUs as an optimization problem. On models like Llama 2 70B, HexGen-2 shows up to 2× throughput improvement at the same price point.
The optimal parallelism strategy for prefill may not match decode. Prefill is a large forward pass through N prompt tokens and benefits from high parallelism. Decode is sequential and latency-sensitive at each step. Using too many GPUs for one decode can hurt output-token latency (TPOT).
Prefill — pipeline parallel (PP) · Decode — tensor parallel (TP)
Some inference engines allow using different precision between phases. For example, prefill in FP8 for speed, decode in FP16 for better generation quality. Usually both phases run at the same precision so the KV cache is compatible. However, conversion can be applied: quantize KV and compress from FP16 to INT8/FP8/FP4 before sending, then convert back on the receiving side.
At extreme scales it is worth evaluating whether CPUs can offload some work from GPUs. Modern CPUs are far slower than GPUs for neural network compute, but they have large RAM and no competition for GPU memory bandwidth. A practical application is background or low-priority prefill tasks. CPU offload would increase TTFT, so it is not used for latency-critical requests.
Early rejection means: if the system predicts it cannot serve a request within the target latency, it fails fast, returning an error or "try later". This preserves goodput by guaranteeing that served requests meet SLO guarantees. In Mooncake, the early rejection strategy evaluates incoming requests based on estimated load on both prefill and decode clusters.
Without admission control the system goes into overload · with early rejection — stable
Some requests may have priority levels: paying customers above free ones; interactive above offline batch. The scheduler can prioritize certain requests to always meet SLOs. Disaggregation can help by allocating a portion of prefill and decode workers for high-priority tasks.
In a disaggregated system, if a decode node crashes mid-generation, the KV cache pool enables recovery on another node, since the KV was stored in the pool. If KV was not stored, the system must recompute prefill on another node. This is why systems like vLLM periodically checkpoint or copy the KV cache to the pool. Disaggregation provides clear control points: prefill and decode metrics can be observed separately, and admission control can be applied exactly where needed.
With a purely static configuration, if the load mix changes, the fixed ratio becomes suboptimal. Adaptive scheduling mechanisms continuously tune the system so no phase becomes the bottleneck.
The TetriInfer research prototype works at two levels of granularity. First, at the request level, it assigns incoming requests to specific prefill and decode instances based on current load. Second, it tracks cluster-wide resource utilization and predicts where bottlenecks may form, proactively shifting work to prevent hotspots. TetriInfer partitions prompts into fixed-size chunks, disaggregates prefill and decode, and uses a two-level scheduling algorithm augmented with predicted resource usage to avoid decode scheduling hotspots.
Predicting bottlenecks and balancing
Arrow continuously analyzes load by measuring incoming token rate versus outgoing token rate and backlog in each worker pool. It then dynamically adjusts worker distribution. Arrow treats prefill and decode worker counts as a tunable parameter that can be adjusted almost in real time. Arrow uses stateless instances where each instance can process both prefill and decode requests without dedicated roles. The system features an SLO-aware scheduling algorithm where a global scheduler dynamically adjusts request dispatching and instance allocation based on SLO settings and instance load. Arrow achieves up to 5.62× higher request serving rates compared to state-of-the-art PD-disaggregated serving systems.
Continuous token stream analysis → dynamic instance reallocation
Mooncake introduced prediction-driven admission control — early rejection. If the system predicts it lacks capacity to process a new request within SLO, it rejects that request. Adaptation can happen on the supply side (adding resources) and on the demand side (throttling or rejecting requests).
Sophisticated inference systems use predictive models such as ARIMA and other forecasting techniques to anticipate traffic spikes and shifts using time-of-day patterns. If a spike of long outputs is predicted at 9 PM, the scheduler can proactively allocate more decode capacity just in time.
Specific load-balancing approaches: elastic instances (Kubernetes HPA), "instance flip" (TetriInfer), statelessness for elasticity (Arrow). Fast role switching can cause oscillation and thrashing, so a minimum role retention time is usually enforced.
«Prefill-decode disaggregation is not just an architectural pattern — it is the foundation for the entire spectrum of optimizations».
Confirmed. By 2026, prefill-decode disaggregation became standard practice for high-scale LLM inference. vLLM, SGLang, NVIDIA Dynamo, and TensorRT-LLM support it natively. Open-source projects like llm-d provide Kubernetes orchestration for disaggregated serving.
«NIXL provides high-throughput, low-latency point-to-point communication for large-scale distributed inference».
Confirmed. NIXL is integrated into vLLM, SGLang, and NVIDIA Dynamo as the primary transport for KV cache. LMCache uses NIXL as the underlying KV transmission.
«Adaptive scheduling and SLO-aware control are necessary to maintain latency guarantees under varying load».
Confirmed. TetriInfer, Arrow, and Mooncake demonstrate the effectiveness of adaptive scheduling. Kubernetes HPA and custom schedulers implement SLO-aware policies.
«POD-Attention is the first GPU kernel that efficiently computes attention for hybrid batches».
Partially confirmed. POD-Attention showed up to 59% speedup (mean 28%) at ASPLOS 2025. However, its mass adoption in production is still limited — further research and integration into popular inference engines are needed.
Dynamic shape limitations. The chapter describes dynamic shapes and CUDA Graphs but does not cover the trade-offs. CUDA Graphs require static shapes, and when dynamic=True is enabled, they are not used. For variable-length sequences, mode="default" or mode="max-autotune-no-cudagraphs" is recommended.
Complexity of compiler debugging. The chapter describes debugging tools but does not emphasize the complexity of diagnosing compiler issues. Graph breaks, recompilations, numerical correctness issues — all require deep understanding of compiler internals. Tools like TORCH_LOGS and torch._dynamo.explain() help, but the learning curve remains high.
Interaction with TensorRT and other backends. The chapter focuses on TorchInductor and XLA but does not consider interaction with TensorRT. In production environments, TensorRT is often used for inference, and torch.compile can integrate with it. This is an important aspect for production deployments that remains outside the chapter's scope.
Triton ecosystem. The chapter describes Triton but does not mention the ecosystem around it. Triton is actively evolving, with new capabilities emerging: TLX (Triton Language Extensions) for low-level control, Blackwell support with tcgen05 instructions, MXFP4/MXFP8. The GPU MODE community actively contributes to Triton's development.
MLPerf Logging. For standardized performance logging in distributed workloads, the mlperf_logging library is used. It allows recording key events (training start, epoch end, final accuracy) in a format compatible with MLPerf.
Primary source: Chris Fregly, «AI Systems Performance Engineering: Optimizing Model Training and Inference Workloads with GPUs, CUDA, and PyTorch», O'Reilly Media, 2026, Chapter 18.
FlashMLA: DeepSeek Open Source Week, February 2025. Official benchmarks: H800 SXM5, CUDA 12.6, 3000 GB/s (memory-bound), 580 TFLOPS (compute-bound). GitHub: github.com/deepseek-ai/FlashMLA.
ThunderMLA: Stanford University research on fully fused attention megakernel for decoding. Report of 20–35% speedup over FlashMLA.
FlexDecoding: PyTorch documentation on torch.nn.attention.flex_attention and FlexDecoding. Support for nested jagged tensors and PagedAttention.
POD-Attention: «POD-Attention: Unlocking Full Prefill-Decode Overlap for Faster LLM Inference», ASPLOS 2025. Attention computation speedup up to 59% (mean 28%) over FlashAttention and FlashInfer.
LMCache: «LMCache: An Efficient KV Cache Layer for Enterprise-Scale LLM Inference», arXiv:2510.09665, October 2025. Up to 15× throughput improvement with vLLM, at least 2× lower latency.
NVIDIA GB200 NVL72: Official NVIDIA specifications. 13.5 TB HBM3e (72 GPUs × 192 GB), 17 TB LPDDR5X (36 Grace CPUs × 480 GB), ~30.5 TB total fast memory.
NIXL: NVIDIA Inference Xfer Library. NVIDIA documentation. Integration with vLLM, SGLang, NVIDIA Dynamo.
NVIDIA Dynamo: NVIDIA documentation. Global scheduling queue, NIXL KV transfer, Kubernetes integration.
TetriInfer: «TetriInfer: Decode Scheduling and Resource Management for LLM Inference», research prototype. Two-level scheduler with predicted resource usage.
Arrow: «Arrow: Adaptive Instance Scaling for Disaggregated LLM Inference». Stateless instances, SLO-aware scheduling algorithm. Up to 5.62× higher request serving rates.
Mooncake: «Mooncake: A KVCache-Centric Architecture for Cloud-Based LLM Serving». Prediction-driven admission control.
Splitwise: «Splitwise: Efficient Generative LLM Inference Using Phase Splitting». Heterogeneous hardware for prefill and decode. 1.4× throughput at 20% lower cost; up to 2.35× RPS on 4× H100 + 4× A100.
HexGen-2: «HexGen-2: Heterogeneous LLM Inference Deployment». Up to 2× throughput improvement on Llama 2 70B at the same price point.
MLPerf: MLPerf Training/Inference v5.0, 2025. Standardized benchmarks for performance comparison.
Ultra-scale LLM inference requires a holistic approach that includes both high-level adaptive resource management and low-level kernel and memory optimizations. By combining the techniques presented in this chapter, a highly optimized inference deployment can achieve maximum throughput on modern hardware while meeting strict latency guarantees.
As hardware and GPUs continue to evolve and software frameworks grow more sophisticated, these optimizations will compound and become even more powerful. Prefill and decode disaggregation is not just an architectural pattern — it is the foundation for the entire spectrum of optimizations, from kernel choice to cluster scheduling.
Each layer benefits from phase independence, and it is their combination — from fused megakernels to adaptive scaling — that makes ultra-scale inference economically viable.
