# cuTensorMapEncodeTiled illegal memory access

> Драйвер 595.58.03 ломал NVFP4 на GB10 ошибкой illegal memory access в cuTensorMapEncodeTiled; запиньте драйвер 580.142. В той же mainline-сборке vLLM ядра под compute_120f распаковывали веса NVFP4 программно, а VLLM_USE_FLASHINFER_MOE_FP4=0 был обходом, а не лечением.

- Platform: DGX Spark (GB10)
- Runtime: vLLM
- Published: 2026-09-02
- Sources: https://agmind.ai/ru/reports/dgx-spark-256k-vllm/
- Upstream: https://github.com/flashinfer-ai/flashinfer/issues/3170
- Canonical: https://agmind.ai/ru/symptoms/vllm-nvfp4-illegal-memory-gb10/

## Что видно

Две сигнатуры одного и того же 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](https://agmind.ai/ru/reports/dgx-spark-256k-vllm/), уровень доказательств `lab_single_run`.
- Статья, которую цитирует отчёт: лонгрид лаборатории на [Хабре](https://habr.com/ru/articles/1033342/).
- Апстрим-трекер, названный на странице, только по нехватке ядер: [issue flashinfer #3170](https://github.com/flashinfer-ai/flashinfer/issues/3170).
- Почему коробка при этом может выглядеть нормально: [почему мониторинг GB10 врёт](https://agmind.ai/ru/reports/dgx-spark-gpu-monitoring/).

---

Machine-readable claim registry: https://agmind.ai/claims.json · llms.txt: https://agmind.ai/llms.txt