默认 LVLLM_MOE_NUMA_ENABLED=0,不设置就等同原生 vLLM 纯 GPU 推理;从源码编译需要 GCC ≥ 13,且要设 VLLM_VERSION_OVERRIDE;混合架构主机需设 FLASHINFER_CUDA_ARCH_LIST 覆盖所有 rank,异构 TP 还要加 --disable-custom-all-reduce。
Lvllm
guqiong96Apache-2.0活跃维护基于 vLLM 的 NUMA 扩展推理引擎,充分利用 CPU 和内存资源降低显存占用,支持 MOE 大模型混合推理与 GPU/NUMA 并行。
LvLLM is a special NUMA extension of vllm that makes full use of CPU and memory resources, reduces GPU memory requirements, and features an efficient GPU parallel and NUMA parallel architecture, supporting hybrid inference for MOE large models.
pip install https://github.com/guqiong96/Lvllm/releases/download/lvllm-v2.5.2/lvllm-2.5.2-cp312-cp312-manylinux_2_34_x86_64.whl需要 x86 AVX2+ 的 CPU 和 NVIDIA GPU(SM75+)。README 未给出单卡最低显存,但参考机为 2× AMD EPYC 7642、1 TiB 内存,DeepSeek-V4.x 峰值常驻约 590 GB(建议 ≥640 GB),Qwen3.8 峰值约 335 GB。
vLLM 的 NUMA 扩展,把 MoE 专家分散到 CPU 内存与显存,降低显存门槛
基于 vLLM 的 NUMA 扩展推理引擎,充分利用 CPU 和内存资源降低显存占用,支持 MOE 大模型混合推理与 GPU/NUMA 并行。
核心能力
- 用 lk_moe 把 MoE 专家分布在显存与系统内存,CPU+GPU 混合调度
- 按 NUMA 感知调度专家,支持线程绑定到 CPU 核心或 NUMA 节点
- GPU prefill 与混合 decode 重叠,按 batch 大小决定在 GPU 还是 CPU 上 prefill
- 可指定常驻显存的专家层,逼近 100% 显存利用率
- 为 SM80/86/89 补齐支持并修复 SM120 路径,覆盖 DeepSeek-V4、Qwen3.8、GLM-5.3
容易踩的坑
- lvllm-v2.5.2最新正式版2026-09-21 · 2026-09-21 12:47 UTC
Base: lvllm-v2.5.1 · patch
What's in
- DeepSeek-V4.x failed engine startup on SM89 (Ada) — the CuTeDSL dequant-gather /
indexer-q kernel picks keyed on "cutlass installed + native fp8" (SM89+), but those
kernels emit inline PTX (
cvt.rn.bf16.f16,mul.bf16x2) with an sm_90 ISA floor, so ptxas aborted every Ada host with "requires .target sm_90 or higher". Addscutedsl_kernels_supported()(fp8 gate AND capability major ≥ 9, per device, fail closed) and applies it at all six pick sites; SM89 now falls back to the Triton path already validated on SM80/86. Kernel selection is unchanged on every previously validated config (SM86 / SM90+ / mixed-arch TP4); SM89 fallback is static-verified only (no Ada card on the build host — first Ada bring-up should run greedy + short GSM8K smoke). - Qwen3.8-Flash-Next FP8 builds —
FP8_PB_WOlayers now serve with the 2-Dweight_scale_invconvention (upstream #54126 contract), making the W4A16-4o6-FP8 checkpoint work end to end (SM86 TP2 MTP decode 61.8–65.9 t/s, accept 3.2–3.5, HTML smoke green). - qwen38 serve scripts point at the W4A16-4o6-FP8 checkpoint. https://huggingface.co/lovedheart/Qwen3.8-Flash-Next-NVFP4-W4A16-4-Over-6-FP8
Commits since v2.5.1:
cd532d25c5(cutedsl SM89 gate) ·badd007357+4c4a6fc0bf(FP8_PB_WO) ·33ded64cb2(release align, tag tip).Install
pip install https://github.com/guqiong96/Lvllm/releases/download/lvllm-v2.5.2/lvllm-2.5.2-cp312-cp312-manylinux_2_34_x86_64.whlAll dependencies resolve automatically, including the FlashInfer wheel attached to this release (
flashinfer_python-0.7.0-py3-none-any.whl, same bytes as the lvllm-v2.5.0/v2.5.1 assets; contains the SM120 sparse-MLA fix #5075, not on PyPI). First serve JIT-compiles the sparse-MLA kernels (nvcc required), results cached under~/.cache/flashinfer/. On mixed-arch hosts addFLASHINFER_CUDA_ARCH_LIST="12.0f".Assets
File Size sha256 lvllm-2.5.2-cp312-cp312-manylinux_2_34_x86_64.whl357,081,001 B b443ae3f33c728c21da004f5cc7258761b1e105434d4c8c45eb2291ab958f852flashinfer_python-0.7.0-py3-none-any.whlsame bytes as lvllm-v2.5.1 asset (sha prefix c3de3c93)— Build environment
python 3.12.11 · torch 2.13.0+cu130 · CUDA 13.2 · gcc-toolset-14 ·
TORCH_CUDA_ARCH_LIST="7.5 8.0 8.6 8.9 9.0 10.0 12.0"·VLLM_VERSION_OVERRIDE=2.5.2· tip33ded64cb2(= taglvllm-v2.5.2)Verify after install
python -c "import vllm; print(vllm.__version__)" # 2.5.2 cuobjdump --list-elf $(python -c "import vllm,os;print(os.path.join(os.path.dirname(vllm.__file__),'_C_stable_libtorch.abi3.so'))") | grep -o 'sm_[0-9]*' | sort -u # expect: sm_100 sm_120 sm_75 sm_80 sm_86 sm_89 sm_90 - DeepSeek-V4.x failed engine startup on SM89 (Ada) — the CuTeDSL dequant-gather /
indexer-q kernel picks keyed on "cutlass installed + native fp8" (SM89+), but those
kernels emit inline PTX (
- lvllm-v2.5.12026-09-19 · 2026-09-19 22:11 UTC
查看更新说明
Base: lvllm-v2.5.0 · Contents: full post-v2.5.0 tree up to release tip
2f47920d7f(Qwen3.8 NVFP4 validation + SM8x tiled prefill attention + KV-startup line + scheduler fixes) · Tested: 4-GPU mixed TP4 (2×3090+2×5060Ti) + pure-sm86 soak, config semantics re-passInstall (Python 3.12, Linux x86_64)
pip install https://github.com/guqiong96/Lvllm/releases/download/lvllm-v2.5.1/lvllm-2.5.1-cp312-cp312-manylinux_2_34_x86_64.whlDependencies (incl. the bundled FlashInfer wheel with the SM120 sparse-MLA fix) resolve automatically; first serve JIT-compiles the sparse-MLA kernels with nvcc (a few minutes), then caches under
~/.cache/flashinfer/. On a mixed-arch host exportFLASHINFER_CUDA_ARCH_LIST="<all ranks>"(e.g."8.6 12.0f").Source: tag
lvllm-v2.5.1(branchmain).Wheel artifacts
File Size SHA256 Notes lvllm-2.5.1-cp312-cp312-manylinux_2_34_x86_64.whl305 MiB (319,942,258 B) eaeb9cec593063149005f91c3dd44a260589faa2354515e5370c935e3f49dfb6recommended (auditwheel-repaired; bundles libcudart/libnvrtc/libssl/libcrypto/libpcre2) flashinfer_python-0.7.0-py3-none-any.whl37 MiB (38,985,337 B) c3de3c93764e0d570bf7a3f870d745f2679ff7e639aa0d60a706219ef607c9adbyte-identical to the v2.5.0 asset (JIT route; no flashinfer-cubin) Build environment: tip
2f47920d7f· Python 3.12.11 · torch 2.13.0+cu130 · nvcc CUDA 13.2 · gcc-toolset-14 (GCC 14.2.1) ·TORCH_CUDA_ARCH_LIST="7.5 8.0 8.6 8.9 10.0 12.0"·VLLM_VERSION_OVERRIDE=2.5.1.Note: these built wheels' METADATA pins flashinfer to the
lvllm-v2.5.0release URL (byte-identical asset, up forever);requirements/cuda.txton this tag already points at this release, so source/editable installs are self-contained too.Verify after install:
python -c "import vllm; print(vllm.__version__)" # -> 2.5.1 cuobjdump --list-elf $(python -c "import vllm,os;print(os.path.dirname(vllm.__file__))")/_C_stable_libtorch.abi3.so | grep -oE 'sm_[0-9]+' | sort -u # expected: sm_75 sm_80 sm_86 sm_89 sm_90 sm_100 sm_120(
_moe_C_stable_libtorch= sm_75/80/86/89/100/120,_flashkda_C= sm_100/120; sm_90 is not in the MoE arch list by design of the build recipe.)
lvllm-v2.5.1 (2026-09-20)
Base: lvllm-v2.5.0 · Release Type: patch (crash fixes + prefill speed + scheduling)
Behavior changes (read before upgrading)
VLLM_DSV4_UNPACK_INDEXERis now ON by default. The indexer-pool carve-out below is active unless you setVLLM_DSV4_UNPACK_INDEXER=0(escape hatch keeps the old packed layout). Cost: ~5 % more bytes per pool block (≈ −11 % block count); this is the fix for the mixed-arch Xid31 crash, so opting out is only meaningful on single-architecture hosts.- YAML config
key: falsenow actually takes effect.--configfiles expanded booleans to--no-*flags only when the flag was visible on the top-level parser; serve options live on theservesubparser, so every*: falseline (e.g.async_scheduling: false) was silently dropped until now. Anything that relied on a droppedfalseline keeps its default only by accident — check your YAML.
Fixes
-
DeepSeek-V4.x mixed-arch TP prefix-reuse crash (Xid31
FAULT_PDE) — under block-outermost interleaved KV packing, indexer pages packed with a ≈1 MBstride0faulted on Blackwell page tables on SM120 ranks once prefix reuse reached the same block depth. The unpacked indexer region (now default;VLLM_DSV4_UNPACK_INDEXER=0opts out) carves indexer layers out of the interleaved pool into a dense tail region (natural per-pagestride0); MLA/compressor groups keep interleaving. Validated on the exact crash seed under both guarded (blocking) and production (FULL-graph) timings, plus GSM8K 100/100 on SM120 TP2. (PR #111) -
Hybrid-model scheduling livelocks on tight KV pools — three scheduler-side fixes: the full-sequence admission gate now splits recycling-aware
- lvllm-v2.5.02026-09-16 · 2026-09-16 02:11 UTC
查看更新说明
Release Notes — lvllm-v2.5.0
Base Version: vllm commit 71888f507a (
upstream/main, 2026-09-14) + lk_moe Release Type: Feature release (Lvllm)Models and supported architectures
Model SM80 SM86 SM89 SM90 SM100 SM120 spec-decode DeepSeek-V4.1-Flash ✅ new ✅ new ✅ new ✅ native ✅ native ✅ fixed here ✅ dspark DeepSeek-V4-Flash (0731) ✅ new ✅ new ✅ new ✅ native ✅ native ✅ native ✅ dspark Qwen3.8-Flash-Next ✅ new ✅ new ✅ new ✅ native ✅ native ✅ new ✅ MTP GLM-5.3-Flash ✅ new ✅ new ✅ new ✅ native ✅ native ✅ fixed here ✅ MTP - native — done by upstream vLLM.
- new — support added by this Lvllm release.
- fixed here — upstream path corrected by this Lvllm release.
Installation
pip install lvllm-2.5.0-*.whlAll dependencies resolve automatically, including the FlashInfer wheel attached to this release (
flashinfer_python-0.7.0-py3-none-any.whl, contains the SM120 sparse-MLA fix #5075; not on PyPI). First serve JIT-compiles the sparse-MLA kernels (nvcc required, a few minutes), then results are cached under~/.cache/flashinfer/. On machines with mixed architectures addFLASHINFER_CUDA_ARCH_LIST="12.0f".Benchmark reference (2× RTX 3090, TP2)
DeepSeek-V4.1-Flash
Item Reference GPUs 2× RTX 3090 (SM86) CPU 2× AMD EPYC 7642 48-Core (96c/192t total, NPS4 ⇒ 8 NUMA nodes), LK_THREADS=48Host RAM machine 1 TiB DDR4-3200 (16-channel, 8 per socket); requirement ≥ 640 GB (peak observed ≈ 590 GB resident, MoE/engram host-resident) Launch:
commands/dsv41_serve_tp2_3090_dspark.shMeasured (single request, greedy):
Setup Plain decode dspark 2× RTX 5060 Ti (SM120, TP2) 25 t/s 32–37 t/s 2× RTX 3090 (SM86, TP2) 27 t/s 26–40 t/s 2× RTX 3090 + 2× RTX 5060 Ti (mixed, TP4) 22 t/s 24–34 t/s DeepSeek-V4-Flash (0731)
Same hardware as V4.1. Launch:
commands/dsv4_0731_serve_tp2_5060ti.sh(dspark variant…_dspark.sh; SM86/TP4 variants also undercommands/).Measured (single request, greedy):
Setup Plain decode dspark 2× RTX 5060 Ti (SM120, TP2) 28 t/s up to 44.8 t/s 2× RTX 3090 (SM86, TP2) 30.6–31.1 t/s 36–43 t/s 2× RTX 3090 + 2× RTX 5060 Ti (mixed, TP4) 30.5–33.6 t/s 38–43.5 t/s Qwen3.8-Flash-Next
Item Reference GPUs 2× RTX 3090 + 2× RTX 5060 Ti (mixed SM86/SM120), or 2× same-arch CPU 2× AMD EPYC 7642 48-Core (96c/192t total, NPS4 ⇒ 8 NUMA nodes), LK_THREADS=48Host RAM machine 1 TiB DDR4-3200 (16-channel, 8 per socket); PLE/FP8 table host-resident (peak observed ≈ 335 GB resident) Launch:
commands/qwen38_serve_tp4.sh(plain) /commands/qwen38_serve_tp4_mtp.sh(MTP).Measured (single request, greedy):
Setup Plain decode MTP 2× RTX 3090 (SM86, TP2) 45 t/s — 2× RTX 3090 + 2× RTX 5060 Ti (mixed, TP4) 40.3–40.6 t/s up to 78 t/s (accept ≈ 3.3) GLM-5.3-Flash
Item Reference GPUs 2× RTX 3090 + 2× RTX 5060 Ti (mixed SM86/SM120), or 2× same-arch CPU 2× AMD EPYC 7642 48-Core (96c/192t total, NPS4 ⇒ 8 NUMA nodes), LK_THREADS=48Host RAM machine 1 TiB DDR4-3200 (16-channel, 8 per socket); model ~190 GiB / 33 shards (nv-community NVFP4), MoE host-resident Launch:
commands/glm53_serve_tp2_5060ti_plain.sh(SM120 plain) /commands/glm53_serve_tp4.sh(mixed TP4; `--kv-cache-dtype - lvllm-v2.4.02026-09-10 · 2026-09-10 04:57 UTC
查看更新说明
Release Notes — v0.2.4.0
Base Version: vLLM [3116c5d] (upstream) + lk_moe v2.4.2 Release Type: Feature integration release
Run command Qwen3.8-Flash-Next ( SM120+) [epyc 7642 x 2, 16 x ddr4 3200, 5060Ti x 2]
normal decode: 35t/s :
CUDA_DEVICE_ORDER="PCI_BUS_ID" \ CUDA_VISIBLE_DEVICES="1,2" \ VLLM_PLE_CPU_OFFLOAD="1" \ LVLLM_MOE_NUMA_ENABLED="1" \ LK_THREADS="48" \ OMP_NUM_THREADS="1" \ LK_THREAD_BINDING="CPU_CORE" \ LVLLM_GPU_PREFETCH_WINDOW="1" \ LVLLM_GPU_PREFILL_MIN_BATCH_SIZE="1024" \ LVLLM_ENABLE_NUMA_INTERLEAVE="1" \ LK_POWER_SAVING="1" \ VLLM_USE_V2_MODEL_RUNNER="1" \ FLASHINFER_DISABLE_VERSION_CHECK="1" \ vllm serve \ --model "/home/guqiong/Models/Qwen3.8-Flash-Next-NVFP4" \ --host "0.0.0.0" \ --port 8070 \ --tensor-parallel-size 2 \ --max-model-len 66000 \ --gpu-memory-utilization 0.95 \ --trust-remote-code \ --tokenizer-mode "auto" \ --served-model-name "Qwen3.8-Flash-Next-NVFP4" \ --compilation-config '{"cudagraph_mode":"FULL_DECODE_ONLY","mode":"VLLM_COMPILE"}' \ --enable-prefix-caching \ --enable-chunked-prefill \ --max-num-batched-tokens 4096 \ --dtype "bfloat16" \ --max-num-seqs 2 \ --enable-auto-tool-choice \ --reasoning-parser "qwen3" \ --tool-call-parser "qwen3_xml" \ --speculative-config '{"method":"mtp","num_speculative_tokens":3}' \ --default-chat-template-kwargs '{"enable_thinking": false}' \ --disable-custom-all-reduceNote: do not add the VLLM_PLE_CPU_OFFLOAD parameter when running other models.
- lvllm-v2.3.112026-08-17 · 2026-08-17 13:03 UTC
查看更新说明
refactor: optimize LK_THREADS calculation and use cfg.use_gpu_prefill in lk-moe v2.3.3 fix https://github.com/guqiong96/Lsglang/issues/16#issuecomment-5297172594: add NUMA interleave override when LVLLM_MOE_NUMA_ENABLED=1 and LVLLM_ENABLE_NUMA_INTERLEAVE=1
⚠️ Important Note Set OMP_NUM_THREADS=1, otherwise model file shard loading will be 10x slower.
⚠️ Important Note Concurrent tilelang JIT compilation corrupts shared cache, crashing vllm. Workaround: rm ~/.cache/vllm/ -rf If vLLM crashes with OOM due to FlashInfer Autotuner during startup, add --no-enable-flashinfer-autotune to the launch command to disable it.
⚠️ Important Note
pip install https://github.com/guqiong96/Lvllm/releases/download/lvllm-v2.3.11/lvllm-2.3.11-cp312-cp312-manylinux_2_34_x86_64.whl # check the latest version at: https://github.com/guqiong96/Lvllm/releases pip install https://github.com/guqiong96/Lvllm/releases/download/lvllm-v2.3.10/flashinfer_cubin-0.6.16.post3-py3-none-any.whlCommand for 2× RTX 5060 Ti
LVLLM_MOE_NUMA_ENABLED=1 \ LK_THREADS=48 \ OMP_NUM_THREADS=1 \ LK_THREAD_BINDING=CPU_CORE \ LVLLM_GPU_PREFETCH_WINDOW=1 \ LVLLM_GPU_PREFILL_MIN_BATCH_SIZE=1024 \ LK_POWER_SAVING=1 \ FLASHINFER_DISABLE_VERSION_CHECK=1 \ vllm serve /home/guqiong/Downloads/DeepSeek-V4-Flash-0731 \ --host 0.0.0.0 \ --port 8070 \ --tensor-parallel-size 2 \ --max-model-len 32000 \ --gpu-memory-utilization 0.95 \ --trust-remote-code \ --served-model-name DeepSeek-V4-Flash-0731 \ --compilation-config '{"cudagraph_mode": "FULL_DECODE_ONLY", "mode": "VLLM_COMPILE"}' \ --enable-prefix-caching \ --enable-chunked-prefill \ --max-num-batched-tokens 4096 \ --dtype bfloat16 \ --max-num-seqs 2 \ --enable-auto-tool-choice \ --kv-cache-dtype fp8_ds_mla \ --tokenizer-mode deepseek_v4 \ --tool-call-parser deepseek_v4 \ --reasoning-parser deepseek_v4 \ --speculative-config '{"method":"dspark","num_speculative_tokens":5,"draft_sample_method":"probabilistic"}' \ --disable-custom-all-reduceRTX Pro 6000 Optimization
Parameter 5060 Ti ×2 RTX Pro 6000 Benefit CUDA_VISIBLE_DEVICES1,20Single GPU --tensor-parallel-size21No TP needed --max-model-len32000128000Extended context --max-num-batched-tokens40968192Faster Prefill LVLLM_GPU_RESIDENT_MOE_LAYERS- "0-13,43-45"Faster Prefill & Decode, 43-45 dSpark draft model on GPU For detailed optimization parameters, see: https://github.com/guqiong96/Lvllm#optimization
Docker: https://hub.docker.com/u/garenleeasa by @garen-lee-asa Note: Add --privileged to the container run command, which grants full access including NUMA topology visibility and SYS_NICE capability in one shot.
用户点评
登录用户的真实使用记录,可按版本和硬件筛选还没有人点评。第一个说出真实体验的人最有价值。
LvLLM
LvLLM = vLLM + lk_moe, plus SM80/86/89 adaptation and SM120 tuning/fixes for popular new models (DeepSeek-V4/V4.1, Qwen3.8, GLM-5.3, …).
- lk_moe is the CPU+GPU hybrid, NUMA-aware MoE engine (pip-installable); vLLM provides the GPU path. LvLLM is lk_moe's integration into vLLM plus the low-arch (SM80/86/89) bring-up and SM120 fixes for these models.
- Fully optional: with
LVLLM_MOE_NUMA_ENABLED=0it behaves exactly like stock vLLM.
Model support
| Model | SM80 | SM86 | SM89 | SM90 | SM100 | SM120 | spec-decode |
|---|---|---|---|---|---|---|---|
| DeepSeek-V4.1-Flash | ✅ new | ✅ new | ✅ new | ✅ native | ✅ native | ✅ fixed | ✅ dspark |
| DeepSeek-V4-Flash (0731) | ✅ new | ✅ new | ✅ new | ✅ native | ✅ native | ✅ native | ✅ dspark |
| Qwen3.8-Flash-Next | ✅ new | ✅ new | ✅ new | ✅ native | ✅ native | ✅ new | ✅ MTP |
| GLM-5.3-Flash | ✅ new | ✅ new | ✅ new | ✅ native | ✅ native | ✅ fixed | ✅ MTP |
native = upstream vLLM · new = support added by this Lvllm release · fixed = upstream path
corrected by this Lvllm release. Full per-model hardware tables, benchmarks and CLI:
RELEASE_NOTES.md.
Previously verified (lk_moe hybrid) models
Original MOE models verified on vLLM's Qwen3 / GLM / MiniMax lines:
| Family | Models |
|---|---|
| Qwen3 | 3.6 / 3.5-35B-A3B, 3.5-122B-A10B, 3.5-397B-A17B, Coder-Next, Next-80B-A3B-Instruct, Coder-30B-A3B-Instruct, VL-30B-A3B-Instruct |
| GLM | GLM-5.2-NVFP4 [sm120], GLM-4.7(-Flash) / 4.6V |
| MiniMax | M3 / M2.7 / M2.5 / M2.1 |
| Others | gemma-4-26B-A4B-it, NVIDIA-Nemotron-3-Super-120B-A12B-BF16, Ornith-1.0-35B-FP8, Kimi k2.6 / k2.5 |
Unlisted original MOE models from the Qwen3 / GLM / MiniMax lines are theoretically supported, pending testing.
Benchmarks
Single request, greedy decode (t/s). Each cell = plain / spec t/s; spec = dspark for
DeepSeek-V4.x, MTP for Qwen3.8. — = not applicable. Full hardware tables: RELEASE_NOTES.md.
| Model | SM120 TP2 (2× 5060 Ti) | SM86 TP2 (2× 3090) | Mixed TP4 (3090×2 + 5060 Ti×2) |
|---|---|---|---|
| DeepSeek-V4.1-Flash | 25 / 32–37 | 27 / 26–40 | 22 / 24–34 |
| DeepSeek-V4-Flash (0731) | 28 / up to 44.8 | 30.6–31.1 / 36–43 | 30.5–33.6 / 38–43.5 |
| Qwen3.8-Flash-Next | — | 45 / — | 40.3–40.6 / up to 78 |
| GLM-5.3-Flash | 20.8 / — | 22–23 / — | 22.4 / — |
Reference hardware
All numbers above were measured on one host:
| Item | Spec |
|---|---|
| CPU | 2× AMD EPYC 7642 48-Core (96c/192t total, NPS4 ⇒ 8 NUMA nodes), LK_THREADS=48 |
| Host RAM | 1 TiB DDR4-3200 (16-channel, 8 per socket) |
| RAM headroom | DeepSeek-V4.x peak ≈ 590 GB resident (≥ 640 GB advised); Qwen3.8 peak ≈ 335 GB |
Why lk_moe
lk_moe spans a MoE model across VRAM + system memory and schedules experts across CPU + GPU with NUMA awareness — reaching ~100% VRAM utilization and overlapping GPU prefill with hybrid decode.
| Role (per MoE layer) | Meaning | Env |
|---|---|---|
| master switch | 0 = stock vLLM pure-GPU, 1 = hybrid | LVLLM_MOE_NUMA_ENABLED |
| GPU-prefill layer | big batches on GPU, small on CPU | LVLLM_GPU_PREFILL_MIN_BATCH_SIZE + LVLLM_GPU_PREFETCH_WINDOW |
| GPU-resident layer | weights stay in VRAM | LVLLM_GPU_RESIDENT_MOE_LAYERS |
Requires x86 AVX2+ and an NVIDIA GPU (SM75+).
Launch
Ready-made serve scripts (per model × topology) live in commands/:
commands/
dsv4_0731_serve_tp2_{3090,5060ti}[_dspark].sh dsv4_0731_serve_tp4[_dspark].sh # DeepSeek-V4-Flash 0731
dsv41_serve_tp2_3090_dspark.sh # DeepSeek-V4.1-Flash
qwen38_serve_tp2_{3090,5060ti}[_mtp].sh qwen38_serve_tp4[_mtp].sh # Qwen3.8-Flash-Next
glm53_serve_tp2_{3090,5060ti}[_plain].sh glm53_serve_tp4.sh # GLM-5.3-Flash (--kv-cache-dtype bfloat16 on sm8x)
Each script is a complete, self-contained vllm serve … (env + args). Pick by model and GPU
topology, adjust the model path, and run:
bash commands/dsv4_0731_serve_tp4_dspark.sh
Configuration
| Env var | Default | Description |
|---|---|---|
LVLLM_MOE_NUMA_ENABLED | 0 | hybrid on/off (0 = stock vLLM) |
LK_THREADS | — | threads = physical cores ÷ #GPUs |
LK_THREAD_BINDING | CPU_CORE | CPU_CORE (best) / NUMA_NODE |
LVLLM_GPU_PREFETCH_WINDOW | — | prefetch window, typically 1 |
LVLLM_GPU_PREFILL_MIN_BATCH_SIZE | — | GPU prefill starts at input ≥ value; 0 = off |
LVLLM_GPU_RESIDENT_MOE_LAYERS | none | expert layers in VRAM, e.g. 0, 0-1,9 |
LVLLM_ENABLE_NUMA_INTERLEAVE | 1 | avoid NUMA node OOM |
LK_POWER_SAVING | 0 | 1 = CPU power saving |
VLLM_DSV4_UNPACK_INDEXER | 0 | DeepSeek-V4.x: carve indexer pages out of the interleaved KV pool (mixed-arch TP prefix-reuse crash workaround) |
On a mixed-arch host add FLASHINFER_CUDA_ARCH_LIST="<all ranks>" (e.g. "8.6 12.0f").
Optimization tips
- Output performance:
--compilation-config '{"mode":"VLLM_COMPILE","cudagraph_mode":"FULL_DECODE_ONLY"}'. - Enable GPU prefill:
LVLLM_GPU_PREFETCH_WINDOW=1,LVLLM_GPU_PREFILL_MIN_BATCH_SIZE=4096,--max-num-batched-tokens 32000; disable with…MIN_BATCH_SIZE=0+--max-num-batched-tokens 4096. - Heterogeneous TP (>2 PCIe GPUs): add
--disable-custom-all-reduce. - BIOS NUMA: AMD EPYC NPS4 / Intel SNC4, node count a multiple of GPU count.
Install
pip install https://github.com/guqiong96/Lvllm/releases/download/lvllm-v2.5.2/lvllm-2.5.2-cp312-cp312-manylinux_2_34_x86_64.whl
# deps (incl. the bundled FlashInfer wheel with the SM120 sparse-MLA fix) resolve automatically
From source (needs GCC ≥ 13 for <format>):
git clone https://github.com/guqiong96/Lvllm.git && cd Lvllm
for d in /opt/rh/gcc-toolset-{14,13}/enable; do [ -f "$d" ] && source "$d" && break; done # RHEL/Rocky
VLLM_VERSION_OVERRIDE="2.5.1" CMAKE_BUILD_TYPE=Release pip install -e . --no-build-isolation
Patches
- In this repo: only the lk_moe integration —
patches/01_lk_moe__71888f507a.patch, a single diff of the current tree over upstream base71888f507a.git checkout 71888f507a && git apply patches/01_lk_moe__71888f507a.patch - The per-model / PR patches are not shipped here; they live in
~/Downloads/opencode/patches/.