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 три основных движка для внимания, и каждый ведёт себя по‑своему.
|
Движок |
Описание |
Что на GB10 |
|---|---|---|
|
FlashInfer |
NVIDIA‑специфичный CUDA, самый быстрый на массовом Blackwell |
FP8-attention сломан: |
|
CUTLASS / FA3 |
Альтернативный быстрый путь, есть поддержка NVFP4 |
Частично. |
|
TRITON_ATTN |
OpenAI Triton, компилируется JIT под GPU |
Работает везде, на 10–30% медленнее FlashInfer, но не падает |
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.xVLLM_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)
|
Квантизация |
Bytes/token (active) |
Roofline tok/s |
Реально замерено |
|---|---|---|---|
|
BF16 |
6 GB |
45 |
23 (Gemma 4 252K) |
|
FP8 |
3 GB |
91 |
52 (FP8 65K) |
|
NVFP4 |
1.5 GB |
182 |
40 (стандартная сборка) |
|
NVFP4 + speculative |
~0.85 GB эффективно |
~320 |
70 (AEON-7) |
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 работает.
Разброс по типам задач огромный:
|
Тип промпта |
tok/s |
|---|---|
|
Математика (T=0) |
117 |
|
Python‑код (T=0) |
103–114 |
|
Длинное рассуждение |
70 |
|
Бизнес‑чат на русском |
45–57 |
|
Креативное письмо |
33–50 |
Тот самый разброс 41–127 tok/s из README AEON-7. На структурированных задачах (математика, код, SQL) drafter попадает много раз подряд, спекулятивные накладки амортизируются. На открытом письме drafter промахивается, скорость откатывается к базовым сорока.
Сравнительная таблица
|
|
Gemma 4 BF16 252K |
FP8 65K |
FP8 260K |
NVFP4 260K |
NVFP4+MTP 131K |
AEON-7 DFlash 260K |
|---|---|---|---|---|---|---|
|
Один поток (среднее) |
23.3 |
52.5 |
51.2 |
40.9 |
54.1 |
69.7 |
|
Один поток (пик) |
— |
— |
— |
41 |
59 |
107 |
|
TTFT тёплый |
100ms |
~50ms |
61ms |
~70ms |
~65ms |
~100ms |
|
Параллельно 4 |
70.0 |
107.6 |
114.9 |
109 |
113.1 |
152.8 |
|
Параллельно 16 |
179.7 |
273.7 |
283.2 |
369 |
400.7 |
266.8 |
|
Параллельно 32 |
297.8 |
— |
349.2 |
448 |
498.6 |
350.1 |
|
Тяжёлый 2K |
22.5 |
— |
48.7 |
39.5 |
48.8 |
73.4 |
|
Максимум контекста |
252K |
65K |
260K |
260K |
131K |
260K |
|
Память на модель |
49 GiB |
35 |
35 |
22 |
22 |
22 |
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.
Если есть свои замеры или другие рабочие рецепты — приходите в комментарии. Тема сложная, информация разбросана по форумам, и проверенные руками цифры весьма ценный ресурс.
Источники
Предыдущие части серии:
ссылка на оригинал статьи https://habr.com/ru/articles/1033342/