cuTensorMapEncodeTiled illegal memory access

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

Նույն 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 կոնտեքստով, ապացույցի մակարդակ՝ lab_single_run։
• Հոդվածը, որ այն մեջբերում է. լաբորատորիայի երկար նյութը Habr-ում։
• Էջում նշված upstream issue-ն՝ միայն kernel-ի բացի համար. flashinfer issue #3170։
• Ինչու տուփն այդ ընթացքում կարող է նորմալ տեսք ունենալ. ինչու է GB10-ի մոնիտորինգը ստում։