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.

lvllm-v2.5.2发布于 2026-09-21 12:47 UTC475star42fork同步于 2026-11-10 23:57 UTC

怎么上手

由 README 提炼,执行前请确认
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
硬件要求按 README 说明

需要 x86 AVX2+ 的 CPU 和 NVIDIA GPU(SM75+)。README 未给出单卡最低显存,但参考机为 2× AMD EPYC 7642、1 TiB 内存,DeepSeek-V4.x 峰值常驻约 590 GB(建议 ≥640 GB),Qwen3.8 峰值约 335 GB。

支持 Apple Silicon
下载:源码 zip·GitHub Releases(安装包)本站不镜像安装包

项目说明

由 README 自动翻译提炼

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_MOE_NUMA_ENABLED=0,不设置就等同原生 vLLM 纯 GPU 推理;从源码编译需要 GCC ≥ 13,且要设 VLLM_VERSION_OVERRIDE;混合架构主机需设 FLASHINFER_CUDA_ARCH_LIST 覆盖所有 rank,异构 TP 还要加 --disable-custom-all-reduce。

内容由模型根据 README 生成,可能不完整或不准确,以官方文档为准。看原文 README

最近发版

5 个版本lvllm-v2.5.2
  1. 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". Adds cutedsl_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_WO layers now serve with the 2-D weight_scale_inv convention (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.whl
    

    All 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 add FLASHINFER_CUDA_ARCH_LIST="12.0f".

    Assets

    FileSizesha256
    lvllm-2.5.2-cp312-cp312-manylinux_2_34_x86_64.whl357,081,001 Bb443ae3f33c728c21da004f5cc7258761b1e105434d4c8c45eb2291ab958f852
    flashinfer_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 · tip 33ded64cb2 (= tag lvllm-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
    
  2. 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-pass

    Install (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.whl
    

    Dependencies (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 export FLASHINFER_CUDA_ARCH_LIST="<all ranks>" (e.g. "8.6 12.0f").

    Source: tag lvllm-v2.5.1 (branch main).

    Wheel artifacts

    FileSizeSHA256Notes
    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.0 release URL (byte-identical asset, up forever); requirements/cuda.txt on 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_INDEXER is now ON by default. The indexer-pool carve-out below is active unless you set VLLM_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: false now actually takes effect. --config files expanded booleans to --no-* flags only when the flag was visible on the top-level parser; serve options live on the serve subparser, so every *: false line (e.g. async_scheduling: false) was silently dropped until now. Anything that relied on a dropped false line 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 MB stride0 faulted 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=0 opts out) carves indexer layers out of the interleaved pool into a dense tail region (natural per-page stride0); 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

  3. 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

    ModelSM80SM86SM89SM90SM100SM120spec-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-*.whl
    

    All 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 add FLASHINFER_CUDA_ARCH_LIST="12.0f".

    Benchmark reference (2× RTX 3090, TP2)

    DeepSeek-V4.1-Flash

    ItemReference
    GPUs2× RTX 3090 (SM86)
    CPU2× AMD EPYC 7642 48-Core (96c/192t total, NPS4 ⇒ 8 NUMA nodes), LK_THREADS=48
    Host RAMmachine 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.sh

    Measured (single request, greedy):

    SetupPlain decodedspark
    2× RTX 5060 Ti (SM120, TP2)25 t/s32–37 t/s
    2× RTX 3090 (SM86, TP2)27 t/s26–40 t/s
    2× RTX 3090 + 2× RTX 5060 Ti (mixed, TP4)22 t/s24–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 under commands/).

    Measured (single request, greedy):

    SetupPlain decodedspark
    2× RTX 5060 Ti (SM120, TP2)28 t/sup to 44.8 t/s
    2× RTX 3090 (SM86, TP2)30.6–31.1 t/s36–43 t/s
    2× RTX 3090 + 2× RTX 5060 Ti (mixed, TP4)30.5–33.6 t/s38–43.5 t/s

    Qwen3.8-Flash-Next

    ItemReference
    GPUs2× RTX 3090 + 2× RTX 5060 Ti (mixed SM86/SM120), or 2× same-arch
    CPU2× AMD EPYC 7642 48-Core (96c/192t total, NPS4 ⇒ 8 NUMA nodes), LK_THREADS=48
    Host RAMmachine 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):

    SetupPlain decodeMTP
    2× RTX 3090 (SM86, TP2)45 t/s—
    2× RTX 3090 + 2× RTX 5060 Ti (mixed, TP4)40.3–40.6 t/sup to 78 t/s (accept ≈ 3.3)

    GLM-5.3-Flash

    ItemReference
    GPUs2× RTX 3090 + 2× RTX 5060 Ti (mixed SM86/SM120), or 2× same-arch
    CPU2× AMD EPYC 7642 48-Core (96c/192t total, NPS4 ⇒ 8 NUMA nodes), LK_THREADS=48
    Host RAMmachine 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

  4. 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-reduce
     
    

    Note: do not add the VLLM_PLE_CPU_OFFLOAD parameter when running other models.

  5. 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.whl
    

    Command 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-reduce
    

    RTX Pro 6000 Optimization

    Parameter5060 Ti ×2RTX Pro 6000Benefit
    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.

用户点评

登录用户的真实使用记录,可按版本和硬件筛选
共 0 条

    还没有人点评。第一个说出真实体验的人最有价值。

    README 原文

    未翻译

    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=0 it behaves exactly like stock vLLM.

    Model support

    ModelSM80SM86SM89SM90SM100SM120spec-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:

    FamilyModels
    Qwen33.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
    GLMGLM-5.2-NVFP4 [sm120], GLM-4.7(-Flash) / 4.6V
    MiniMaxM3 / M2.7 / M2.5 / M2.1
    Othersgemma-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.

    ModelSM120 TP2 (2× 5060 Ti)SM86 TP2 (2× 3090)Mixed TP4 (3090×2 + 5060 Ti×2)
    DeepSeek-V4.1-Flash25 / 32–3727 / 26–4022 / 24–34
    DeepSeek-V4-Flash (0731)28 / up to 44.830.6–31.1 / 36–4330.5–33.6 / 38–43.5
    Qwen3.8-Flash-Next—45 / —40.3–40.6 / up to 78
    GLM-5.3-Flash20.8 / —22–23 / —22.4 / —

    Reference hardware

    All numbers above were measured on one host:

    ItemSpec
    CPU2× AMD EPYC 7642 48-Core (96c/192t total, NPS4 ⇒ 8 NUMA nodes), LK_THREADS=48
    Host RAM1 TiB DDR4-3200 (16-channel, 8 per socket)
    RAM headroomDeepSeek-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)MeaningEnv
    master switch0 = stock vLLM pure-GPU, 1 = hybridLVLLM_MOE_NUMA_ENABLED
    GPU-prefill layerbig batches on GPU, small on CPULVLLM_GPU_PREFILL_MIN_BATCH_SIZE + LVLLM_GPU_PREFETCH_WINDOW
    GPU-resident layerweights stay in VRAMLVLLM_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 varDefaultDescription
    LVLLM_MOE_NUMA_ENABLED0hybrid on/off (0 = stock vLLM)
    LK_THREADS—threads = physical cores ÷ #GPUs
    LK_THREAD_BINDINGCPU_CORECPU_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_LAYERSnoneexpert layers in VRAM, e.g. 0, 0-1,9
    LVLLM_ENABLE_NUMA_INTERLEAVE1avoid NUMA node OOM
    LK_POWER_SAVING01 = CPU power saving
    VLLM_DSV4_UNPACK_INDEXER0DeepSeek-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 base 71888f507a.
      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/.

    Release / history

    See https://github.com/guqiong96/Lvllm/releases.