vLLM-ը DGX Spark-ի վրա 256K կոնտեքստով. կոնֆիգուրացիաների որսը և NVFP4 ուղին, որը կոտրված էր mainline-ում

Սա լաբորատորիայի ռուսերեն լոնգրիդի անգլերեն ուղեկցի հայերեն թարգմանությունն է. բնօրինակը՝ Habr-ում. ինչպես ստիպել 35B դասի MoE-ին սպասարկել 256K կոնտեքստի պատուհան մեկ DGX Spark-ի վրա, և ինչ կոտրված գտանք ճանապարհին։ Այդ հոդվածն է այստեղ ամեն թվի առաջնային աղբյուրը — ներքևի թվերը նրանից արված մեջբերումներ են, ոչ թե լաբորատորիայի ռեեստրի claim-եր։

Նախ կարդացեք տարբերակների ամրացումը։ Այս էջի ամեն ինչ նկարագրում է շարժվող թիրախի մեկ պահ. vLLM 0.15 շարքի mainline՝ PyTorch 2.10-ով, FlashInfer 0.6.8, դրայվեր 580.142, CUDA 13.0, DGX OS 7.5.0, չափված 2026-ի մայիսին։ Mainline-ը շարժվում է ամեն շաբաթ. ինչ կոտրված էր այն ժամանակ, կարող է արդեն ուղղված լինել, իսկ ինչ աշխատում էր՝ հետընթաց ապրած։ Վերաբերվեք այս էջին որպես թվագրված վկայության՝ իր ամրացումներով, և կրկնելուց առաջ ստուգեք upstream-ի հղված issue-ները։

Մեքենան և նպատակը

Մեկ DGX Spark. GB10, Blackwell SM_121, 128 ԳիԲ unified LPDDR5x՝ 273 ԳԲ/վ թողունակությամբ ։ Նպատակը իսկապես երկար պատուհան էր — 262 144 թոքեն — Qwen3.6-35B-A3B դասի MoE-ի համար, սարքաշարի վրա, որի compute ճարտարապետության համար բաց կոդով kernel-ների մեծ մասը երբեք չէր կոմպիլացվել։

Կոնֆիգուրացիաների որսը

Ինչը հրաժարվեց աշխատել, հանդիպման հերթականությամբ.

• FlashInfer-ի FP8 attention-ը ուղղակի մերժեց չիպը — kernel-ը սպասարկում է sm120, և SM_121-ը fallback-ի փոխարեն ստանում էր անօրինական սինխրոնիզացիայի սխալներ։
• CUTLASS/FA3 FP4 ուղիները կրում են էվրիստիկա, որ պահանջում է զրո minor compute capability, ինչը GB10-ին բացառում է ըստ կառուցվածքի։
• Triton fallback-ը աշխատեց և մնաց կայուն՝ FlashInfer-ի դիմաց մոտավորապես 10–30% գնով ։
• MoE բэкենդի դրոշները պետք էր ձեռքով դնել. VLLM_USE_FLASHINFER_MOE_FP4=0, որովհետև autotuner-ը տասնյոթ տակտիկայից տասը բաց էր թողնում որպես չսպասարկվող , մինչդեռ VLLM_USE_FLASHINFER_MOE_FP8=1-ը SM_121-ի վրա FlashInfer 0.6.x-ով աշխատում էր։
• CUDA-գրաֆի ամբողջական capture-ը կախվում էր vLLM 0.15 շարքի սարքումներում. workaround-ը cudagraph_mode=piecewise-ն է։
• Դրայվերի ամրացումը 580.142-ն է։ 590.48.01 դրայվերը հոսեցնում էր unified հիշողությունը, մինչև 128-ից 80 ԳԲ սպառվում էր՝ AnonPages-ում ու Slab-ում ոչինչ չերևալով , իսկ 595.58.03-ը կոտրում էր NVFP4-ը cuTensorMapEncodeTiled անօրինական հիշողության դիմումով։ Եթե այս պահին ձեր Grafana-ն նորմալ տեսք ունի, տեսեք ինչու է GB10-ի մոնիտորինգը ստում։

Ինչու NVFP4-ը կոտրված էր mainline-ի մեր փորձարկած սարքումում

Ցավոտ հարցումը, որ հավանաբար ձեզ այստեղ բերեց։ Մեր ամրացրած mainline-սարքումում NVFP4 kernel-ները կոմպիլացված էին compute_120f-ի համար, մինչդեռ նատիվ NVFP4 հրահանգները գոյություն ունեն միայն compute_120a-ում և compute_121a-ում։ SM_121-ի վրա դա նշանակում էր, որ քվանտացված կշիռները բացվում էին ծրագրային ճանապարհով՝ բիթային մանիպուլյացիաներով shader-ում, պարապ տենզորային միջուկներով։ Լոգի մատնահետքը՝ [AutoTuner]: Skipped 10 unsupported tactic(s) for trtllm::fused_moe::gemm2։

Չափված հետևանքը. կիսով չափ հիշողության հետք ունեցող քվանտացումը դեկոդում էր FP8-ից դանդաղ — 40,9՝ 51,2 tok/s-ի դիմաց մեկ հոսքով, 260K կոնֆիգուրացիայում ։ Կիսով չափ բայթ, ավելի շատ ժամանակ. բացազերծման ծախսը կերավ թողունակության շահը։

Ճշգրիտ լինենք ժամանակաձևի հարցում. NVFP4-ը կոտրված էր այդ սարքումում, այդ չիպի վրա, այդ ամսաթվին։ Սա կոմպիլացիայի թիրախի բաց է, ոչ թե սարքաշարի օրենք — upstream-ի հետագծողը flashinfer-ի issue #3170-ն էր, «DGX Spark SM_121 Audit»-ը՝ տասնյոթ բաց կետով հոդվածի հրապարակման պահին ։ Որևէ բան որոշելուց առաջ ստուգեք ընթացիկ mainline-ի դեմ։

Ինչն աշխատեց

260K պահած և ամենաարագ կոնֆիգուրացիան ամենևին ստոկային չէր. vLLM՝ սորսից՝ յոթ լոկալ պատչով, FlashInfer 0.6.8՝ գումարած PR #2520 և #2702, community AEON-7/Qwen3.6-35B-A3B-heretic-NVFP4 checkpoint-ը և DFlash block-diffusion drafter-ը՝ տասնհինգ սպեկուլյատիվ թոքենով.

--max-model-len 262144
--max-num-seqs 128
--max-num-batched-tokens 65536
--gpu-memory-utilization 0.85
--enable-chunked-prefill --enable-prefix-caching
--speculative-config '{"method":"dflash","num_speculative_tokens":15}'
--attention-backend flash_attn

Արագ ուղու թակարդը. heretic fine-tune-ը կոտրում է tool calling-ը (նրա issue #4-ը, այն պահին չուղղված)։ Agent և MCP սպասարկման համար հոդվածի հանձնարարականը սովորական Qwen/Qwen3.6-35B-A3B-FP8-ն էր ստոկ vLLM-ի վրա — ավելի դանդաղ, բայց ֆունկցիաների կանչերը վերադառնում են ամբողջական։

Չափված վարքը 256K-ում

Մեջբերված հոդվածի մեկ մեքենայի գործարկումներից .

Տաք time-to-first-token-ը մոտ 100 մվ էր՝ ներառյալ drafter-ի կոմպիլացիան ։ Ամենաօգտակար անակնկալը հենց պատուհանի էժանությունն է. max-model-len-ի աճը 65K-ից 260K արժեցավ մեկ հոսքի դեկոդի մոտ երկուսուկես տոկոս , մինչդեռ մեկ օգտագործողի ակտիվ կոնտեքստը հազվադեպ էր անցնում 260K պատուհանի 30K-ն ։ Սպեկուլյատիվ թվերը ուժեղ տատանվում են ըստ բեռի — DFlash-ի ընդունումը ընդհանուր 17.8% էր, բայց զրո դիրքում 78%՝ մաթեմատիկայի ու կոդի վրա , ինչը դեկոդը փռում էր մաթեմատիկայի 117 tok/s-ից մինչև բիզնես-չաթի 45–57 tok/s ։ Իսկ առաստաղը հիշողությունն է, ոչ թե մարքեթինգը. 273 ԳԲ/վ-ի դեպքում հոդվածի roofline-ը արդյունավետ սպեկուլյատիվ դեկոդը դնում է մոտ 320 tok/s ՝ անկախ նրանից, թե ինչ է ակնարկում տուփի վրայի FP4 պետաֆլոպի թիվը։

Ազնվության բաժին

• Այստեղ ամեն ինչ lab_single_run է. մեկ լաբորատորիա, մեկ մեքենա, միայնակ անցումներ։ Այս թվերից ոչ մեկը ռեեստրի claim չէ — դրանք մեր իսկ հոդվածից արված մեջբերումներ են՝ իր վերապահումներով։
• Վերևի տարբերակների ամրացումը ամբողջ իմաստն է։ «Կոտրված» վճիռն առանց իր build-ի մատնահետքի աղմուկ է; նրա հետ միասին՝ վերարտադրելի bug report։
• Երկնոդանոց սպասարկման բաղադրատոմսը՝ իր սեփական թակարդներով, առանձին էջ է, իսկ թե ինչու են Spark-ի հրապարակված թվերը իրար հակասում՝ առանձին քննություն։