Средний

14 мин

13K

NVIDIA продаёт спарку с лозунгом «один петафлоп на FP4». Я купил коробку, поставил vLLM, запустил инференс и получил 40 токенов в секунду на 35B MoE‑модели. После маркетинговых слайдов цифра выглядит грустно.

Объяснение простое. NVFP4 в основной ветке vLLM и FlashInfer физически сломан на SM_121 — варианте Blackwell, который установлен в GB10. Ядра собраны под compute_120f, а нативные NVFP4-инструкции есть только в compute_120a и compute_121a. На SM_121 распаковка квантованных весов идёт через программные битовые манипуляции в шейдере, без участия тензорных ядер.

Сообщество вытащило стек руками: нашло обходные пути, собрало рабочие конфигурации. Я прогнал на своём Spark шесть разных конфигураций vLLM — от стокового BF16 до форка с DFlash speculative decoding — и замерил каждую одинаковым тестом. В этой статье разбираю, что в итоге работает и что выбирать под разные задачи.

Если коротко

• GB10 — неоднозначная платформа. На ней крутятся DGX Spark (NVIDIA reference), ASUS Ascent GX10, Lenovo ThinkStation PGX, HP ZGX Nano G1n, MSI EdgeXpert и решения от Dell, Acer и других OEM‑партнёров. Позиционируется всё это как «AI workstation для разработчиков». Но реальные ресурсы NVIDIA — инженеры, оптимизации, ранние сборки — идут в датацентровые B100/B200/B300 и Hopper. SM_121 в большинстве open‑source ядер (FlashInfer, CUTLASS, TensorRT‑LLM) либо не поддерживается, либо требует обходных путей. Поддержки от NVIDIA на форумах мало, ETA фиксов чаще всего нет.

• На моделях класса 35B параметров DGX Spark упирается в пропускную способность памяти. У него 273 GB/s LPDDR5x против 3.4 TB/s HBM3 у H100. В двенадцать раз медленнее, и инференс LLM эту разницу честно собирает.

• NVFP4 в основной ветке vLLM и FlashInfer на SM_121 не работает как должен. Issue flashinfer #3170 на момент написания статьи открыт, в нём 17 пунктов, прогресс по семи критическим я не смог подтвердить как закрытый.

• Стоковая конфигурация выдаёт 40–54 токена в секунду на одном потоке. Маркетинг «1 PFLOPS» на этом фоне выглядит ехидно. Не первый раз ловлю себя на знаменитой реакции Торвальдса в адрес NVIDIA.

• Конфигурация AEON-7 с DFlash пробивает потолок до 70 токенов в среднем и 107 на пике. Это форк vLLM из исходников с семью патчами и drafter из z‑lab. Важная оговорка: чекпойнт heretic — abliterated (без safety filtering), и на нём сломан function calling. Для обычного чата и RAG это не мешает, для Dify‑агентов и MCP — стоп‑фактор, тогда нужен чистый Qwen3.6-FP8.

• Контекст 260K на скорость почти не влияет. KV‑cache и раскладка памяти не упираются в потолок.

• MoE обыгрывает плотные модели на данном железе. На bandwidth‑bound GB10 важны активные параметры, а не общий размер.

• Драйвер залочен на 580.x. На 590+ есть две критических регрессии. Без обновлений можно жить долго и счастливо.

• vLLM 0.15+ на PyTorch 2.10 имеет отдельную проблему с FULL CUDA graph capture mode — лечится cudagraph_mode=piecewise или сборкой из исходников. Это не про драйвер.

• Если вы покупали DGX Spark под маркетинговый «один петафлоп на FP4» — забудьте слайды. На 35B‑моделях реальные цифры скромнее. Под классический сервинг и продакшн с function calling — Qwen3.6-FP8 на стандартной vLLM. Под максимум скорости в чате и RAG (без tool calling) — community‑форк AEON-7 с DFlash. Конкретные числа и ограничения каждого пути — ниже.

Зачем это вообще понадобилось

В моём проекте встал вопрос: какая конфигурация vLLM выжимает максимум на контекстном окне 256K. Длинный контекст нужен для серьёзных вещей.

Юридические договоры целиком в ввод, без RAG. Договор поставки на сорок страниц с многоуровневой нумерацией — это около пятидесяти тысяч токенов на русском. Резать его на чанки часто хуже, чем отдать модели целиком. Анализ кодовой базы за один проход — средний микросервис на пятьдесят-сто файлов влезает в двести тысяч токенов с метаданными. Многошаговые агентские сценарии: история разговора растёт, агент держит в голове предыдущие шаги, инструменты, промежуточные результаты. На двадцати-тридцати шагах легко уходит за сто тысяч. Техническая документация в одном запросе — Confluence-выгрузка с прикреплёнными файлами или регламент компании со всеми приложениями.

Стандартная сборка AGmind крутила Gemma 4 26B-A4B-it в BF16 на 64K. Когда клиент захотел 256K, я пошёл в эксперименты. NVFP4 — квантизация, в два раза меньше памяти. MTP — speculative decoding. DFlash — block-diffusion drafter. Форки сообщества. Получилось шесть разных конфигураций, каждую прогнал одним и тем же тестом.

Железо

NVIDIA DGX Spark
├── GPU:        GB10 (Blackwell SM_121)
├── Память:     128 GiB единая (LPDDR5x)
├── Bandwidth:  273 GB/s
├── Драйвер:    580.142 (NVIDIA не сертифицирует более новые на Spark)
├── CUDA:       13.0
├── ОС:         DGX OS 7.5.0 (aarch64, Ubuntu 24.04)
└── Кластер:    2 ноды через QSFP 200G

Spark позиционируется как рабочая станция для разработчиков AI. Главное — единая память. Модель видит CPU и GPU как один пул на 128 GiB, ходить через PCIe не надо.

Минусов у платформы три. Главный — bandwidth: 273 GB/s против 3.4 TB/s у H100 HBM3, в двенадцать раз меньше. Декод LLM упирается именно в это. На MoE‑модели с тремя миллиардами активных параметров теоретический потолок при NVFP4 — около 180 токенов в секунду. На H100 в FP8 на той же модели потолок более тысячи; на датацентровых Blackwell B200 в NVFP4 — за пять тысяч.

Дальше — драйвер залочен на 580.x. На 590 и выше есть критические регрессии. Подробности были во второй части серии.

И отдельно — SM_121 это узкая разновидность Blackwell, на которой собирается только GB10. Под массовый SM_120 экосистема тестируется и собирается основательно. Под SM_121 ядра либо отсутствуют, либо собраны вторым приоритетом.

Зоопарк attention-движков

В vLLM три основных движка для внимания, и каждый ведёт себя по‑своему.

TRITON_ATTN — де‑факто запасной вариант. У меня в проекте закреплено правило: все FP8-нагрузки идут через VLLM_ATTENTION_BACKEND=TRITON_ATTN, если используется FlashInfer FP8.

Триггеры для перехода: ошибка kernel only supports sm120 в логах, краши illegal synchronization operations, апгрейд vLLM с пересборкой FlashInfer (старая версия работала, новая нет).

В моём эксперименте AEON-7 уже собран с flash_attn (библиотека FlashAttention от Tri Dao, не Triton). На GB10 она работает корректно, поэтому AEON-7 не упирается в проблемы FlashInfer.

Бэкенд для MoE

Для MoE‑моделей (Qwen3.6, Mixtral, GLM) есть отдельный выбор:

bash

VLLM_USE_FLASHINFER_MOE_FP8=1   # FP8 MoE — работает на SM_121 с flashinfer 0.6.x
VLLM_USE_FLASHINFER_MOE_FP4=0   # NVFP4 MoE — десять стратегий не поддерживаются
VLLM_NVFP4_GEMM_BACKEND=marlin  # запасной вариант на ядре Marlin INT4

Marlin — это GEMM‑ядро для INT4-весов с FP16/BF16-активациями. На SM_121 оно работает корректно, но веса должны быть заранее проквантизованы в формат, который Marlin понимает. Использовать его как запасной вариант имеет смысл, если NVFP4-путь не идёт.

Почему важно НЕ обновлять драйвер

DGX OS 7.5.0 пинает 580.142 как единственный сертифицированный драйвер. На GB10 на единой памяти есть две независимые регрессии в 590+, которых нет на дискретной видеопамяти.

Первая — утечка UMA в 590.48.01. После чистого выхода CUDA‑процесса около 80 из 128 GB «съедены», их не видно в AnonPages, Slab, PageTables. MemAvailable падает. Лечится echo 3 > /proc/sys/vm/drop_caches или выгрузкой модуля nvidia. Тред форума #359969 — NVIDIA подтвердила баг и явно сказала, что новые драйверы после 580.126.09 на Spark пока не поддерживаются.

Вторая — баг TMA в 595.58.03. cuTensorMapEncodeTiled ловит illegal memory access. Валит NVFP4 на Spark (vLLM #35519).

Нужно захолдить драйвер:

bash

sudo apt-mark hold nvidia-driver-580-open

Иначе unattended‑upgrades подтянет 590+ и стек посыпется.

Что в итоге сломано в основной ветке

NVFP4 откатывается на программную распаковку, в три раза медленнее нативного. TensorRT‑LLM на Spark требует workaround для PTXAS (sm_121a is not defined for option gpu-name, Triton issue #8539) — лечится передачей корректного TRITON_PTXAS_PATH через mpirun. FP8-attention в FlashInfer ломается на illegal sync, надо TRITON_ATTN. NVFP4 в MoE — десять стратегий из семнадцати не поддерживаются. Скомпилированных кубинов TRTLLM‑gen FMHA для SM_121 не существует. Marlin INT4 работает.

Сколько физически возможно: roofline

Простая математика для авторегрессивного декода на MoE 35B / 3B-active:

peak_tok_s = bandwidth / (bytes_per_param × active_params_per_token)

Roofline даёт верхнюю границу. Реальные числа всегда ниже. Причины: трафик KV-cache (приходится читать состояние прошлых токенов), нормализации слоёв, softmax, сэмплинг, накладные расходы на запуск ядра при каждом шаге декода. Для NVFP4 добавляется ещё и распаковка блоковых масштабов.

Эксперимент: шесть конфигураций

Все запускались на одной машине, одним и тем же бенч‑скриптом. Пять одиночных запусков с разными промптами (T=0.7, до 300 токенов на ответ), TTFT через стриминг, параллельно 4, 8, 16 и 32 запроса, тяжёлый промпт на 2K ввода и 400 ответа. Между конфигурациями я останавливал текущий контейнер vLLM и поднимал новый. Параллельно две модели не крутил — вторая нода кластера в это время держала отдельный production‑инстанс, перенастраивать её ради бенчей не имело смысла.

Gemma 4 26B-A4B-it BF16 на 252K

Это базовая конфигурация AGmind до экспериментов. BF16-веса, FP8 KV-cache (без него 252K не влезает: 252K × 580 KB = 146 GB).

bash

docker run vllm/vllm-openai:gemma4-cu130 \
  google/gemma-4-26B-A4B-it \
  --max-model-len 252000 \
  --gpu-memory-utilization 0.85 \
  --kv-cache-dtype fp8

Один поток в среднем — 23.3 tok/s, разброс меньше процента. TTFT на тёплом — 94-106 мс. Параллельно 32 запроса — 297.8 tok/s суммарно. Тяжёлый промпт 2K — 22.5 tok/s.

Speculative decoding нет, каждый токен идёт через целевую модель. Квантизации нет, BF16 на 49 GiB тяжело тащить через 273 GB/s. Стабильность пять из пяти, скорость удручает.

Qwen3.6-35B-A3B-FP8 на 65K

Свежая Qwen3.6 (релиз апрель 2026) — MoE 35B / 3B активных, нативный FP8 от команды Qwen. Образ vllm/vllm-openai:cu130-nightly.

bash

docker run vllm/vllm-openai:cu130-nightly \
  Qwen/Qwen3.6-35B-A3B-FP8 \
  --max-model-len 65536 \
  --gpu-memory-utilization 0.85 \
  --port 8000

Один поток — 52.5 tok/s. Это в два с половиной раза быстрее Gemma. Параллельно 32 я не запускал, упёрся в лимит concurrency на 65K.

Прирост относительно Gemma даёт сразу несколько вещей. FP8 — половина байта на параметр против BF16. Гибридное внимание Qwen3.6 (Gated DeltaNet и стандартный attention) даёт почти линейную сложность на длинном контексте. Активных параметров 3B против 4B у Gemma — тоже играет в плюс.

Qwen3.6-35B-A3B-FP8 на 260K

Та же модель, но контекст бахнул до 260K с явным FP8-кэшем.

Один поток — 51.2 tok/s, минус два с половиной процента от 65K. Ничтожно. Параллельно 32 — 349.2 tok/s. Тяжёлый 2K — 48.7 tok/s.

На FP8 контекст 65K против 260K на одиночный поток не влияет. KV-pool вырастает с 805K до 1.6M токенов, но на одного юзера активно используется не больше 30K. Узкое место в пропускной способности памяти не зависит от длины контекста.

RedHatAI/Qwen3.6-35B-A3B-NVFP4 на 260K

Та же модель в квантизации NVFP4: четырёхбитные веса, FP8 в блоковых масштабах. На бумаге вдвое меньше памяти, должно ускорить декод.

bash

docker run vllm/vllm-openai:cu130-nightly \
  RedHatAI/Qwen3.6-35B-A3B-NVFP4 \
  --max-model-len 260000 \
  --gpu-memory-utilization 0.85 \
  --moe-backend flashinfer_cutlass \
  --kv-cache-dtype fp8

Один поток — 40.9 tok/s. Медленнее, чем FP8. Параллельно 8 — 232 tok/s, FP8 выигрывает на concurrency. Параллельно 32 — 448 tok/s.

В логах vLLM значилось [AutoTuner]: Skipped 10 unsupported tactic(s) for trtllm::fused_moe::gemm2. Полез копать, нашёл flashinfer issue #3170 «DGX Spark / SM121 Audit». Там семнадцать пунктов, и среди них именно то, что описано выше: ядра под compute_120f, нативные NVFP4-инструкции только в compute_120a/121a, на SM_121 распаковка квантованных весов идёт через программные битовые манипуляции. CUTLASS-ядро dense_blockscaled_gemm_sm120.py:1591 строкой if sm_version != "sm_120" отвергает SM_121 целиком. Кубинов TRTLLM-gen FMHA для SM_121 не существует (issue #11799).

Итог: NVFP4 на DGX Spark в основной ветке подломлен на уровне ядер. Профит по памяти (вдвое меньше байтов) компенсируется накладными расходами программной распаковки.

NVFP4 + MTP speculative на 131K

Дальше я взял рецепт Стива Скаргалла и добавил MTP — Multi-Token Prediction. Это нативная фича Qwen3.6: один слой в самой модели предсказывает следующий токен параллельно с основным forward pass. Если предсказание совпадает — экономим один шаг декода.

docker run vllm/vllm-openai:cu130-nightly \
  RedHatAI/Qwen3.6-35B-A3B-NVFP4 \
  --max-model-len 131072 \
  --gpu-memory-utilization 0.87 \
  --moe-backend flashinfer_cutlass \
  --kv-cache-dtype fp8_e4m3 \
  --max-num-batched-tokens 32768 \
  --max-num-seqs 32 \
  --enable-chunked-prefill --enable-prefix-caching \
  --speculative-config '{"method":"mtp","num_speculative_tokens":1}'

На запуске поймал AssertionError: In Mamba cache align mode, block_size (2112) must be <= max_num_batched_tokens (2048). Лечится обязательным --max-num-batched-tokens 32768. Гибридное внимание Qwen3-Next с GDN и Mamba-слоями требует специфической конфигурации block_size.

Ещё одна засада: для Qwen3.6 строго num_speculative_tokens=1. MTP — это один слой в модели, и больше одного спекулятивного токена даёт ноль процентов acceptance (issue #36331).

Один поток — 54.1 tok/s, плюс шесть процентов к чистому FP8 и плюс тридцать два процента к NVFP4 без speculative. Параллельно 8 — 160 tok/s (плюс девятнадцать процентов к FP8). Параллельно 16 — 400.7 tok/s (плюс сорок два). Параллельно 32 — 498.6 tok/s (плюс сорок три, пик в эксперименте). MTP acceptance — около восьмидесяти процентов.

MTP заметно сильнее на параллельной нагрузке: больше места для амортизации спекулятивных накладок через batched verification.

AEON-7 DFlash на 260K — победитель

Этот форк я нашёл через NVIDIA developer forum. Репозиторий AEON-7/Qwen3.6-NVFP4-DFlash. Внутри:

— vLLM, собранный из исходников с HEAD, плюс семь патчей. Среди них register_qwen3_5_text.py, четыре патча на patch_kv_cache_utils.py, патч на mrope-fallback, патч на выравнивание CUDAGraph. — Build target TORCH_CUDA_ARCH_LIST="12.0+PTX" плюс переменная VLLM_TEST_FORCE_FP8_MARLIN=1. — FlashInfer 0.6.8 с PR #2520 и #2702 (ядра NVFP4 KV-cache decode для sm_120). — Модель AEON-7/Qwen3.6-35B-A3B-heretic-NVFP4 — abliterated fine-tune Qwen3.6 (без safety filtering, KL divergence 0.000492 от базовой модели, в README заявлено что capacity preserved). Лицензия Apache 2.0, мультимодальность сохранена, квантизация в NVFP4. Корпоративным заказчикам с compliance-требованиями по safety этот чекпойнт не подойдёт — берите оригинальный Qwen/Qwen3.6-35B-A3B-FP8 (скорость ниже, но политика родная). — Drafter z-lab/Qwen3.6-35B-A3B-DFlash — block-diffusion, дообучены пять слоёв, ест промежуточные скрытые состояния из слоёв 1, 16, 31, 46, 61 целевой модели и за один forward pass предсказывает блок из шестнадцати токенов.

bash

docker run \
  -v /opt/qwen36/qwen36-nvfp4:/models/qwen36 \
  -v /opt/qwen36/qwen36-dflash:/models/qwen36-dflash \
  ghcr.io/aeon-7/vllm-spark-omni-q36:v1.2 \
  bash -c "exec vllm serve /models/qwen36 \
    --max-model-len 262144 --max-num-seqs 128 \
    --max-num-batched-tokens 65536 \
    --gpu-memory-utilization 0.85 \
    --enable-chunked-prefill --enable-prefix-caching \
    --reasoning-parser qwen3 \
    --speculative-config '{\"method\":\"dflash\",\"model\":\"/models/qwen36-dflash\",\"num_speculative_tokens\":15}' \
    --attention-backend flash_attn"

Один поток в среднем — 69.7 tok/s. Это втрое больше базовой Gemma и на тридцать четыре процента больше чистого FP8. Пиковая скорость на одном потоке — 107 tok/s, ловится на математических и кодовых промптах. TTFT тёплый — около ста миллисекунд (с учётом компиляции drafter). Параллельно 4 — 152.8 tok/s суммарно (плюс тридцать четыре от FP8). Тяжёлый промпт 2K на 400 ответа — 73.4 tok/s (плюс пятьдесят процентов от FP8). Acceptance DFlash — 17.8% общий, 78% на нулевой позиции для математики и кода.

Известное ограничение, важное для agentic‑сценариев: на heretic‑чекпойнте сломан tool calling. Модель видит OpenAI‑совместимый tools payload в запросе, но возвращает пустой tool_calls: [] и finish_reason: length вместо вызова функции. Issue #4 на GitHub открыт без фикса. Для классического чата и RAG это не критично. Для Dify‑агентов с function calling, MCP‑интеграций и любых workflow с вызовом инструментов — берите вместо AEON-7 чистый Qwen/Qwen3.6-35B-A3B-FP8 без heretic‑fine‑tune. Скорость ниже (51 vs 70 tok/s), но tool calling работает.

Разброс по типам задач огромный:

Тот самый разброс 41–127 tok/s из README AEON-7. На структурированных задачах (математика, код, SQL) drafter попадает много раз подряд, спекулятивные накладки амортизируются. На открытом письме drafter промахивается, скорость откатывается к базовым сорока.

Сравнительная таблица

AEON-7 DFlash — лучший на одном потоке и при невысокой параллельности. Это типичный сценарий бизнес-чата. NVFP4+MTP — лучший при высокой параллельной нагрузке.

Что в итоге выбрать

Для бизнес‑чата, RAG и Dify я остановился на AEON-7 DFlash 260K. 70 tok/s на бизнес‑чате, 100+ на коде, полные 260K контекста. Думалка в Dify работает в поле reasoning_content, Dify его сворачивает, юзер видит чистый финальный ответ.

Для ИИ‑агентов с function calling и MCP‑плагинов AEON-7 не подходит: на heretic‑чекпойнте сломан tool calling (Issue #4). Тут нужен чистый Qwen3.6-FP8 260K — он чуть медленнее, зато tools работают.

Для пакетной обработки или API‑сервиса с большим количеством юзеров лучше брать NVFP4+MTP 131K. Скорость на одного юзера ниже, но суммарная пропускная способность выше 500 tok/s на 32 параллельных запросах.

Если нужна стабильность, а не максимум скорости — Qwen3.6-FP8 260K. 53 tok/s, без форков, на стандартной сборке vLLM. Никаких кастомных билдов и tool calling работает из коробки, политика безопасности — родная Qwen.

В сухом остатке

Маркетинговый «один петафлоп на FP4» на DGX Spark в реальности упирается в пропускную способность памяти и в то, что открытый стек до сих пор не доведён под SM_121 как следует. Сорок токенов в секунду на стандартной сборке — это текущий уровень готовности экосистемы.

Сообщество ушло вперёд: AEON-7, Стив Скаргалл, Avarok разогнали NVFP4 в открытом стеке через форки vLLM, патчи FlashInfer и speculative decoding. У кого Spark и нужен реальный инференс на 35B‑моделях с длинным контекстом — ставьте AEON-7.

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

Источники

• flashinfer #3170 — DGX Spark SM_121 Audit

• AEON-7/Qwen3.6-NVFP4-DFlash

• z-lab DFlash project (GitHub)

• z-lab DFlash paper page

• arXiv 2602.06036 — DFlash: Block Diffusion for Flash Speculative Decoding

• Steve Scargall NVFP4 recipe

• ai-muninn Gemma 4 + MTP recipe

• LMSYS DGX Spark in-depth review

• NVIDIA forum #359969 — UMA leak in 590.x

• NVIDIA forum #360181 — vLLM 0.15+ CUDAGraph FULL mode hang

• NVIDIA forum #366006 — TRT-LLM PTXAS workaround

• vLLM #36331 — MTP 0% acceptance at N=4

• vLLM #36821 — No sm_121 support on aarch64

• Cutlass #3096 — SM120 NVFP4 MoE garbage output

• Triton #8539 — sm_121a not defined for option gpu-name

Предыдущие части серии:

• Часть 1. Я собрал на DGX Spark приватный AI-сервер

• Часть 2. Мониторинг единой памяти, когда NVML и dcgm-exporter молчат

• Часть 3. Кириллица в LLM: почему русский язык в нейросетях стоит дороже