cuTensorMapEncodeTiled illegal memory access

Что видно

Две сигнатуры одного и того же NVFP4-пути на GB10. На драйвере 595.58.03 NVFP4 ломался с illegal memory access в cuTensorMapEncodeTiled. В отчёте это записано именно так, словами, а не строкой из лога; точного текста сообщения рантайма на странице нет.

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

[AutoTuner]: Skipped 10 unsupported tactic(s) for trtllm::fused_moe::gemm2

Где мы это поймали

• Железо: один DGX Spark, GB10, Blackwell SM_121.
• Рантайм: vLLM mainline времён 0.15 с PyTorch 2.10, FlashInfer 0.6.8, CUDA 13.0, DGX OS 7.5.0. Драйвер 580.142 запинен; illegal memory access пришёл с 595.58.03.
• Модель: MoE класса Qwen3.6-35B-A3B в NVFP4, в конфигурации отчёта под длинный контекст.
• Когда: замеры мая 2026-го; отчёт датирован 2026-08-21.

Причина

До первопричины illegal memory access отчёт не докапывается, только связывает его с версией драйвера. Зафиксировано, что делал каждый драйвер новее: на 590.48.01 текла unified-память, причём ни в AnonPages, ни в Slab это не было видно; 595.58.03 ломал NVFP4 ошибкой cuTensorMapEncodeTiled. В Grafana при этом всё может выглядеть нормально.

Медленный путь — не та цель компиляции. В mainline-сборке, которую мы запинили, ядра NVFP4 были собраны под compute_120f, а нативные NVFP4-инструкции есть только в compute_120a и compute_121a. На SM_121 квантованные веса распаковывались программно: битовые операции в шейдере при простаивающих тензорных ядрах, а автотюнер отбрасывал большинство тактик fused-MoE GEMM как неподдерживаемые.

Отдельно отчёт называет FP4-пути CUTLASS/FA3, которые отвергают чип с порога: в них зашита эвристика, требующая нулевой минорной версии compute capability, и GB10 под неё не проходит по определению.

Лечение

Драйвер: запинить 580.142. На нашей коробке этим лечение и исчерпывается; ни один из двух драйверов новее, которые пробовали в отчёте, не выжил: на 590.48.01 текла unified-память, 595.58.03 ломал NVFP4.

Нехватка ядер: отчёт документирует обход, а не лечение. Флаги MoE-бэкенда выставьте руками:

VLLM_USE_FLASHINFER_MOE_FP4=0
VLLM_USE_FLASHINFER_MOE_FP8=1

FP4 выключен, потому что автотюнер отбрасывал его тактики; FP8 на SM_121 с FlashInfer 0.6.x работал.

Конфигурация, которая вытянула длинное окно с NVFP4-чекпоинтом, — vLLM, собранный из исходников с семью локальными патчами, FlashInfer 0.6.8 плюс PR #2520 и #2702, чекпоинт AEON-7/Qwen3.6-35B-A3B-heretic-NVFP4 и драфтер DFlash; флаги запуска — в отчёте. Нигде там не сказано, что эти патчи пересобирают ядра под compute_121a, так что не принимайте быстрый путь за лечение. У быстрого пути своя ловушка: файнтюн heretic ломает tool calling. Под агентов и MCP статья, которую цитирует отчёт, рекомендовала ванильный Qwen/Qwen3.6-35B-A3B-FP8 на стоковом vLLM.

Открыто ли апстрим?

На ошибку драйвера тикет не назван; в отчёте она закрыта пином версии. Нехватку ядер, которая стоит за строкой AutoTuner, отслеживает issue flashinfer #3170, аудит поддержки SM_121 для DGX Spark; на момент выхода статьи часть пунктов там оставалась открытой. Сначала сверьтесь с текущим mainline.

Доказательства

• Страница-источник: vLLM на DGX Spark с контекстом 256K, уровень доказательств lab_single_run.
• Статья, которую цитирует отчёт: лонгрид лаборатории на Хабре.
• Апстрим-трекер, названный на странице, только по нехватке ядер: issue flashinfer #3170.
• Почему коробка при этом может выглядеть нормально: почему мониторинг GB10 врёт.