# cuTensorMapEncodeTiled illegal memory access

> 595.58.03 դրայվերը GB10-ի վրա կոտրեց NVFP4-ը cuTensorMapEncodeTiled-ի հիշողության անթույլատրելի դիմումով. ամրագրեք 580.142-ն։ Նույն mainline հավաքումում compute_120f-ի kernel-ները NVFP4 կշիռները ծրագրայնորեն էին բացում. VLLM_USE_FLASHINFER_MOE_FP4=0-ն շրջանցում է, ոչ թե ուղղում։

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

## Ինչ եք տեսնում

Նույն NVFP4 ուղին GB10-ի վրա երկու հետք է թողնում։ 595.58.03 դրայվերի վրա
NVFP4-ը կոտրվեց `cuTensorMapEncodeTiled`-ի հիշողության անթույլատրելի
դիմումով։ Հաշվետվությունը դա գրանցում է հենց այդ բառերով, ոչ թե որպես
լոգից վերցված տող; runtime-ի ճշգրիտ հաղորդագրությունն էջում չկա։

Ամրագրված դրայվերի վրա NVFP4-ը բեռնվեց և հարցումներ սպասարկեց, բայց լոգում
այս տողը կար, և դեկոդը նույն կոնտեքստում FP8-ից դանդաղ ստացվեց։ Կիսով չափ
պակաս բայթ, ավելի շատ ժամանակ.

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

## Որտեղ ենք հանդիպել

- Սարքաշար. մեկ DGX Spark, GB10, Blackwell SM_121։
- Runtime. vLLM 0.15 շարքի mainline՝ PyTorch 2.10-ով, FlashInfer 0.6.8-ով,
  CUDA 13.0-ով, DGX OS 7.5.0-ով։ Ամրագրված դրայվերը 580.142-ն է; հիշողության
  անթույլատրելի դիմումը 595.58.03-ից էր։
- Մոդել. Qwen3.6-35B-A3B դասի MoE NVFP4-ով՝ հաշվետվության երկար կոնտեքստի
  կոնֆիգուրացիայում։
- Երբ. չափվել է 2026-ի մայիսին; հաշվետվությունը թվագրված է 2026-08-21-ով։

## Պատճառ

Հաշվետվությունը հիշողության անթույլատրելի դիմումի արմատական պատճառը չի
պարզում; այն դա դրայվերի տարբերակի հետ է կապում։ Այնտեղ գրանցված է, թե
ինչ արեց ավելի նոր դրայվերներից յուրաքանչյուրը. 590.48.01-ը unified
հիշողության արտահոսք էր տալիս, ընդ որում AnonPages-ում ու Slab-ում ոչինչ
չէր երևում, իսկ 595.58.03-ը կոտրում էր NVFP4-ը `cuTensorMapEncodeTiled`
սխալով։ Grafana-ն այդ ընթացքում կարող է նորմալ տեսք ունենալ։

Դանդաղ ուղու հետևում կոմպիլացիայի թիրախի բացն է։ Մեր ամրագրած mainline
հավաքումում NVFP4 kernel-ները կոմպիլացված էին `compute_120f`-ի համար,
մինչդեռ նատիվ NVFP4 հրահանգները կան միայն `compute_120a`-ում և
`compute_121a`-ում։ SM_121-ի վրա քվանտավորված կշիռներն ապափաթեթավորվում
էին ծրագրայնորեն՝ shader-ում բիթային մանիպուլյացիաներով, մինչ տենզորային
միջուկները պարապ էին, իսկ autotuner-ը fused-MoE GEMM տակտիկաների մեծ մասը
բաց էր թողնում որպես չաջակցվող։

Առանձին՝ հաշվետվությունը CUTLASS/FA3 FP4 ուղիները նշում է որպես չիպն
առհասարակ մերժող. դրանցում կա էվրիստիկա, որը compute capability-ի
զրոյական minor համար է պահանջում, ինչն ինքնին բացառում է GB10-ը։

## Լուծում

Դրայվեր. ամրագրեք 580.142-ն։ Մեր տուփի վրա սա ամբողջ լուծումն է;
հաշվետվության փորձարկած երկու ավելի նոր դրայվերներից ոչ մեկը չդիմացավ.
590.48.01-ը unified հիշողության արտահոսք էր տալիս, 595.58.03-ը կոտրում էր
NVFP4-ը։

Kernel-ի բացը. հաշվետվությունը նկարագրում է շրջանցող լուծում, ոչ թե
ուղղում։ MoE բեքենդի դրոշները ձեռքով սահմանեք.

```
VLLM_USE_FLASHINFER_MOE_FP4=0
VLLM_USE_FLASHINFER_MOE_FP8=1
```

FP4-ն անջատված է, որովհետև autotuner-ը դրա տակտիկաները բաց էր թողնում;
FP8-ը SM_121-ի վրա FlashInfer 0.6.x-ով աշխատում էր։

Կոնֆիգուրացիան, որ NVFP4 checkpoint-ով երկար պատուհանը պահեց, հետևյալն էր.
ելակոդից հավաքված vLLM՝ յոթ տեղական պատչով, FlashInfer 0.6.8՝ #2520 և
#2702 PR-ներով, `AEON-7/Qwen3.6-35B-A3B-heretic-NVFP4` checkpoint-ը և
DFlash drafter-ը; գործարկման դրոշները հաշվետվությունում են։ Այնտեղ ոչ մի
տեղ չի ասվում, որ այդ պատչերը kernel-ները վերակոմպիլացնում են
`compute_121a`-ի համար, ուստի արագ ուղին որպես ուղղում մի՛ ընկալեք։ Արագ
ուղին իր թակարդն ունի. heretic fine-tune-ը կոտրում է tool calling-ը։
Ագենտների և MCP-ի սպասարկման համար հաշվետվության մեջբերած հոդվածը խորհուրդ
էր տալիս սովորական `Qwen/Qwen3.6-35B-A3B-FP8`-ը՝ չփոփոխված vLLM-ի վրա։

## Բա՞ց է դեռ upstream-ում

Դրայվերի սխալի համար ոչ մի տիկետ նշված չէ; այն գրանցված է որպես ամրագրում։
flashinfer-ի issue #3170-ը՝ DGX Spark SM_121-ի աջակցության աուդիտը, հետևում
է kernel-ի այն բացին, որի հետքն AutoTuner-ի տողն է, և հոդվածի հրապարակման
պահին այնտեղ դեռ բաց կետեր կային; նախ ստուգեք ընթացիկ mainline-ի վրա։

## Ապացույցներ

- Սկզբնաղբյուր էջ. [vLLM-ը DGX Spark-ի վրա 256K կոնտեքստով](https://agmind.ai/hy/reports/dgx-spark-256k-vllm/), ապացույցի մակարդակ՝ `lab_single_run`։
- Հոդվածը, որ այն մեջբերում է. լաբորատորիայի երկար նյութը [Habr-ում](https://habr.com/ru/articles/1033342/)։
- Էջում նշված upstream issue-ն՝ միայն kernel-ի բացի համար. [flashinfer issue #3170](https://github.com/flashinfer-ai/flashinfer/issues/3170)։
- Ինչու տուփն այդ ընթացքում կարող է նորմալ տեսք ունենալ. [ինչու է GB10-ի մոնիտորինգը ստում](https://agmind.ai/hy/reports/dgx-spark-gpu-monitoring/)։

---

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