Repository navigation
Add an option to use dummy weights - #33
Merged
Merged
Conversation
hongxiayang
pushed a commit
to hongxiayang/vllm
that referenced
this pull request
Feb 13, 2024
tianyil1
pushed a commit
to tianyil1/vllm
that referenced
this pull request
Jun 5, 2024
* Bucketing/Warmup WIP * Cleanup * Revert "Fix model_output_idx on HPU (vllm-project#27)" This reverts commit 90dfa92. * Rework selected_token_indices fix to also work with block_size padding * Simple prompt attention POC * Remove cumsum * MQA/GQA support for simple prompt_attention * Cleanup * Fix typo * Restore profiling runs
dllehr-amd
pushed a commit
to dllehr-amd/vllm
that referenced
this pull request
Jul 22, 2024
…ernel tuning script for rocm. Merge pull request vllm-project#33 - tuned moe configs v2
bigPYJ1151
pushed a commit
to bigPYJ1151/vllm
that referenced
this pull request
Jul 31, 2024
Enable jit for com ops
Closed
1 task done
1 task done
1 task done
1 task done
zyongye
pushed a commit
to zyongye/vllm
that referenced
this pull request
Aug 5, 2025
zyongye
pushed a commit
to zyongye/vllm
that referenced
this pull request
Aug 6, 2025
heheda12345
added a commit
to heheda12345/vllm
that referenced
this pull request
Sep 29, 2025
…oject#26) * indexer medatata to separate prefill and decode * deep_gemm prefill kernel * decode kernel, can run for single batch * bug fixing insert decode k into kv before gemm * don't use tilelang quant function * faster non-looping torch for kv cache insertion * add chunked prefill impl * change quant kernel back to tilelang for promotion * fix format (vllm-project#31) Signed-off-by: Chen Zhang <zhangch99@outlook.com> * update unit tests * Fp8 indexer prefill (vllm-project#33) * init Signed-off-by: Chen Zhang <zhangch99@outlook.com> * can run --------- Signed-off-by: Chen Zhang <zhangch99@outlook.com> * remove debug comment Signed-off-by: Chen Zhang <zhangch99@outlook.com> * cleanup * further cleanup --------- Signed-off-by: Chen Zhang <zhangch99@outlook.com> Co-authored-by: mgoin <mgoin64@gmail.com> Co-authored-by: Chen Zhang <zhangch99@outlook.com>
1 task done
yma11
pushed a commit
to yma11/vllm
that referenced
this pull request
Dec 8, 2025
* Support cpu kv-cache offload on XPU platform Signed-off-by: chzhang <chaojun.zhang@intel.com> * Support cpu kv-cache offload on XPU platform Signed-off-by: chzhang <chaojun.zhang@intel.com> --------- Signed-off-by: chzhang <chaojun.zhang@intel.com>
GuoRen868
pushed a commit
to GuoRen868/vllm
that referenced
this pull request
Jan 27, 2026
AFDConnector, ubatch refactor
1 task done
1 task done
amd-xavierwang
pushed a commit
to amd-xavierwang/vllm
that referenced
this pull request
Apr 8, 2026
jxmorris12
added a commit
to jxmorris12/vllm
that referenced
this pull request
Apr 21, 2026
* fixes * bugfixes * make cartridge writes atomic * fix some compilation logic to match upstream
malaiwah
added a commit
to malaiwah/vllm-fairness
that referenced
this pull request
Aug 11, 2026
Verify the "align with r34" premise against real data rather than assuming it. Findings: - r34's vLLM base is dev/gilded-gnosis@e2666d9a, unchanged since 2026-08-07. Our branch is already an exact descendant, so no rebase is required to align with r34. - r34 is not a branch state. Its vLLM tree 4d006a4 is an isolated composition of the base plus 21 PR heads and 404s on GitHub. - The lock field disposition:"merged" does not mean merged; all 21 vLLM PRs are still open and vllm-project#145 is closed-unmerged yet shipped. - malaiwah/vllm-voipmonitor forks vllm-project/vllm directly, so PRs default to the wrong base and must be retargeted by hand. - sparkinfer and b12x are one repo mid-rename; exllamav3 comes from brandonmmusic-max, not the org. - The r34 merge chain is blocked at B12X vllm-project#144, which conflicts. PR target: local-inference-lab/vllm, base dev/gilded-gnosis. Sources: GitHub REST API, rtx6kpro issue vllm-project#33 and models/glm5.2_v20.md, blackwell-llm-docker release locks, skopeo inspect of the r34 manifest. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Ant9REBoa9Tpm5xHUkuW9x
1 task done
curvedinf
added a commit
to curvedinf/int8-vllm
that referenced
this pull request
Aug 30, 2026
…edges)
Three defects, convicted by a concurrent-load bisect (sequential repro
legs almost never fired; the user's agent workload runs parallel streams):
1. offloading/scheduler.py: the eagle catch-all flagged EVERY KV group
(incl. mamba) as volatile draft attention under method=dflash —
is_eagle_group is only ever set for DeepSeek V4 upstream, so the
fallback fired for all 15 groups (boot line "[0..14]"). Now only the
drafter's sliding-window group is flagged ("[14]"). Port of
qwen38-27b-rtx3090 issue vllm-project#33, plus their uniform-block chunk-geometry
warning.
2. kv_cache_manager.get_computed_blocks_for_connector: divergent
per-group prefix hits were kept whenever external tokens existed,
resuming at full-attention's computed count over mamba's shorter
state coverage -> dirty mamba state page on resume -> the garble.
Now always reconciled to the boundary every group agrees on.
3. GDN spec kernels: bound the (num_accepted-1) state lookups
(vLLM PR vllm-project#50021): causal_conv1d_update indexed the rolling buffer
unbounded; fused_sigmoid_gating could read into the next request's
state-indices row. Stale counts now degrade to a zeroed step instead
of cross-request state corruption or wild-address kernel spins
(the engine-wedge signature).
Validation: concurrent 3-stream 20k-in/4k-out rounds went from 3-5
acceptance-collapse windows/round + 2 full engine wedges to 0 over
5 rounds (13 streams) with the CPU tier actively storing; plus 4/4
clean on the reported 20k-in/4k-out single-shot shape. The tier stays
ON (OFFLOAD=1 default restored).
Also lands: pinned --kv-cache-memory 20.2GiB default (+4.9% arena,
1,031,145 tokens), GPU_PINNED_MIN_XFER_SIZE=67108864 (load 38.4->36.8s),
systemd-safe diagnostic flag files in serve_recipe_qwen38.sh
(OFFLOAD/PREFIXCACHE/SPECOFF/NS/UA/KVMEM under logs/serve_recipe_qwen38/),
repro + collapse-detection tooling, ledger rows G1-G8. The
hybrid-kv-groups bucket-padding port was evaluated and measured-REJECTED
(arena 1,031,145 -> 945,602; per-group mamba spec-window accounting),
reverted with a do-not-retry note.
Assisted-by: Kimi Code
cpuchip
pushed a commit
to cpuchip/vllm
that referenced
this pull request
Sep 14, 2026
OffloadingConnector under dflash: stop eagle-flagging every KV group, and warn on the chunk-geometry that silently zeroes cross-request reuse (vllm-project#33). Two findings from tracing the "KVarN stores 41 GB, loads exactly 0 bytes" report (issue vllm-project#33) on this card, with the connector instrumented: 1. Upstream only ever sets is_eagle_group for DeepSeek V4, and the connector's fallback marks EVERY group as draft attention whenever speculative_config.use_eagle() is true -- which includes method=dflash and mtp. Every group then has its trailing chunk excluded from store while decoding ("EAGLE/MTP draft attention groups [0..8] detected" at boot). For dflash the draft model's KV is exactly its sliding-window layers, so the fallback now flags SW-spec groups only; boot line reads "[8]". Under other eagle-family methods the upstream fallback is unchanged. 2. The actual zero-loads mechanism, reproduced and traced here (KVarN CTX=huge + TieringOffloadingSpec, 4 GiB CPU tier): the CPU pool allocates UNIFORM blocks sized for the largest group's chunk (2,176 tokens -> ~14.6 MiB blocks, 293 of them in 4.28 GiB). Under KVarN the promoted drafter SW group's chunks are 128 tokens -- each one still occupies a full block, so a single 23k-token request consumes ~264 blocks (~184 of them for tiny SW chunks) and EVICTS THE ENTIRE TIER. Store succeeds, complete_store succeeds, and by the time the same document returns, its blocks are gone: every lookup MISS, 0 bytes ever loaded, exactly the report. On bf16 the SW group shares the large block size (the same 4096-B/token coincidence hybrid-sw-block-promote.patch documents), so the geometry never explodes and offloading uplifts -- which is why the reporter saw the split along KV dtype. The real fix (per-group block sizing, or excluding never-hittable SW groups from both store and lookup) is upstream-shaped; until then the config builder computes the waste factor and WARNS at boot with the arithmetic (group, chunk tokens, blocks-per-chunk ratio, and the cpu_bytes_to_use multiplier that would compensate), so the next person gets the diagnosis in their log instead of a day of silence. Verified on this card: boot line shows eagle groups [8]; the warning fires for group 8 (128 vs 2,176 tokens, 17x) on KVarN + connector and stays silent on configs with uniform chunk sizes. Source: syv-ai/qwen38-27b-rtx3090 patches/offload-dflash-eagle-groups.patch (at d4c4cc2). Applied as-is from the patch file; the preamble above is the patch's own. Signed-off-by: Michael Stufflebeam <cpuchip@gmail.com>
cpuchip
pushed a commit
to cpuchip/vllm
that referenced
this pull request
Sep 23, 2026
OffloadingConnector under dflash: stop eagle-flagging every KV group, and warn on the chunk-geometry that silently zeroes cross-request reuse (vllm-project#33). Two findings from tracing the "KVarN stores 41 GB, loads exactly 0 bytes" report (issue vllm-project#33) on this card, with the connector instrumented: 1. Upstream only ever sets is_eagle_group for DeepSeek V4, and the connector's fallback marks EVERY group as draft attention whenever speculative_config.use_eagle() is true -- which includes method=dflash and mtp. Every group then has its trailing chunk excluded from store while decoding ("EAGLE/MTP draft attention groups [0..8] detected" at boot). For dflash the draft model's KV is exactly its sliding-window layers, so the fallback now flags SW-spec groups only; boot line reads "[8]". Under other eagle-family methods the upstream fallback is unchanged. 2. The actual zero-loads mechanism, reproduced and traced here (KVarN CTX=huge + TieringOffloadingSpec, 4 GiB CPU tier): the CPU pool allocates UNIFORM blocks sized for the largest group's chunk (2,176 tokens -> ~14.6 MiB blocks, 293 of them in 4.28 GiB). Under KVarN the promoted drafter SW group's chunks are 128 tokens -- each one still occupies a full block, so a single 23k-token request consumes ~264 blocks (~184 of them for tiny SW chunks) and EVICTS THE ENTIRE TIER. Store succeeds, complete_store succeeds, and by the time the same document returns, its blocks are gone: every lookup MISS, 0 bytes ever loaded, exactly the report. On bf16 the SW group shares the large block size (the same 4096-B/token coincidence hybrid-sw-block-promote.patch documents), so the geometry never explodes and offloading uplifts -- which is why the reporter saw the split along KV dtype. The real fix (per-group block sizing, or excluding never-hittable SW groups from both store and lookup) is upstream-shaped; until then the config builder computes the waste factor and WARNS at boot with the arithmetic (group, chunk tokens, blocks-per-chunk ratio, and the cpu_bytes_to_use multiplier that would compensate), so the next person gets the diagnosis in their log instead of a day of silence. Verified on this card: boot line shows eagle groups [8]; the warning fires for group 8 (128 vs 2,176 tokens, 17x) on KVarN + connector and stays silent on configs with uniform chunk sizes. Source: syv-ai/qwen38-27b-rtx3090 patches/offload-dflash-eagle-groups.patch (at d4c4cc2). Applied as-is from the patch file; the preamble above is the patch's own. Signed-off-by: Michael Stufflebeam <cpuchip@gmail.com> Re-cut onto v0.30.0 from the qwen38/main-track resolution (e5341a1); the topic is unchanged on the 0.29 line since main-track was cut.
cpuchip
pushed a commit
to cpuchip/vllm
that referenced
this pull request
Sep 23, 2026
OffloadingConnector under dflash: stop eagle-flagging every KV group, and warn on the chunk-geometry that silently zeroes cross-request reuse (vllm-project#33). Two findings from tracing the "KVarN stores 41 GB, loads exactly 0 bytes" report (issue vllm-project#33) on this card, with the connector instrumented: 1. Upstream only ever sets is_eagle_group for DeepSeek V4, and the connector's fallback marks EVERY group as draft attention whenever speculative_config.use_eagle() is true -- which includes method=dflash and mtp. Every group then has its trailing chunk excluded from store while decoding ("EAGLE/MTP draft attention groups [0..8] detected" at boot). For dflash the draft model's KV is exactly its sliding-window layers, so the fallback now flags SW-spec groups only; boot line reads "[8]". Under other eagle-family methods the upstream fallback is unchanged. 2. The actual zero-loads mechanism, reproduced and traced here (KVarN CTX=huge + TieringOffloadingSpec, 4 GiB CPU tier): the CPU pool allocates UNIFORM blocks sized for the largest group's chunk (2,176 tokens -> ~14.6 MiB blocks, 293 of them in 4.28 GiB). Under KVarN the promoted drafter SW group's chunks are 128 tokens -- each one still occupies a full block, so a single 23k-token request consumes ~264 blocks (~184 of them for tiny SW chunks) and EVICTS THE ENTIRE TIER. Store succeeds, complete_store succeeds, and by the time the same document returns, its blocks are gone: every lookup MISS, 0 bytes ever loaded, exactly the report. On bf16 the SW group shares the large block size (the same 4096-B/token coincidence hybrid-sw-block-promote.patch documents), so the geometry never explodes and offloading uplifts -- which is why the reporter saw the split along KV dtype. The real fix (per-group block sizing, or excluding never-hittable SW groups from both store and lookup) is upstream-shaped; until then the config builder computes the waste factor and WARNS at boot with the arithmetic (group, chunk tokens, blocks-per-chunk ratio, and the cpu_bytes_to_use multiplier that would compensate), so the next person gets the diagnosis in their log instead of a day of silence. Verified on this card: boot line shows eagle groups [8]; the warning fires for group 8 (128 vs 2,176 tokens, 17x) on KVarN + connector and stays silent on configs with uniform chunk sizes. Source: syv-ai/qwen38-27b-rtx3090 patches/offload-dflash-eagle-groups.patch (at d4c4cc2). Applied as-is from the patch file; the preamble above is the patch's own. Signed-off-by: Michael Stufflebeam <cpuchip@gmail.com> Re-cut onto v0.30.0 from the qwen38/main-track resolution (e5341a1); the topic is unchanged on the 0.29 line since main-track was cut.
nonameentername
pushed a commit
to nonameentername/vllm
that referenced
this pull request
Oct 7, 2026
OffloadingConnector under dflash: stop eagle-flagging every KV group, and warn on the chunk-geometry that silently zeroes cross-request reuse (vllm-project#33). Two findings from tracing the "KVarN stores 41 GB, loads exactly 0 bytes" report (issue vllm-project#33) on this card, with the connector instrumented: 1. Upstream only ever sets is_eagle_group for DeepSeek V4, and the connector's fallback marks EVERY group as draft attention whenever speculative_config.use_eagle() is true -- which includes method=dflash and mtp. Every group then has its trailing chunk excluded from store while decoding ("EAGLE/MTP draft attention groups [0..8] detected" at boot). For dflash the draft model's KV is exactly its sliding-window layers, so the fallback now flags SW-spec groups only; boot line reads "[8]". Under other eagle-family methods the upstream fallback is unchanged. 2. The actual zero-loads mechanism, reproduced and traced here (KVarN CTX=huge + TieringOffloadingSpec, 4 GiB CPU tier): the CPU pool allocates UNIFORM blocks sized for the largest group's chunk (2,176 tokens -> ~14.6 MiB blocks, 293 of them in 4.28 GiB). Under KVarN the promoted drafter SW group's chunks are 128 tokens -- each one still occupies a full block, so a single 23k-token request consumes ~264 blocks (~184 of them for tiny SW chunks) and EVICTS THE ENTIRE TIER. Store succeeds, complete_store succeeds, and by the time the same document returns, its blocks are gone: every lookup MISS, 0 bytes ever loaded, exactly the report. On bf16 the SW group shares the large block size (the same 4096-B/token coincidence hybrid-sw-block-promote.patch documents), so the geometry never explodes and offloading uplifts -- which is why the reporter saw the split along KV dtype. The real fix (per-group block sizing, or excluding never-hittable SW groups from both store and lookup) is upstream-shaped; until then the config builder computes the waste factor and WARNS at boot with the arithmetic (group, chunk tokens, blocks-per-chunk ratio, and the cpu_bytes_to_use multiplier that would compensate), so the next person gets the diagnosis in their log instead of a day of silence. Verified on this card: boot line shows eagle groups [8]; the warning fires for group 8 (128 vs 2,176 tokens, 17x) on KVarN + connector and stays silent on configs with uniform chunk sizes. Source: syv-ai/qwen38-27b-rtx3090 patches/offload-dflash-eagle-groups.patch (at d4c4cc2). Applied as-is from the patch file; the preamble above is the patch's own. Signed-off-by: Michael Stufflebeam <cpuchip@gmail.com> Re-cut onto v0.30.0 from the qwen38/main-track resolution (e5341a1); the topic is unchanged on the 0.29 line since main-track was cut.
cpuchip
pushed a commit
to cpuchip/vllm
that referenced
this pull request
Oct 9, 2026
OffloadingConnector under dflash: stop eagle-flagging every KV group, and warn on the chunk-geometry that silently zeroes cross-request reuse (vllm-project#33). Two findings from tracing the "KVarN stores 41 GB, loads exactly 0 bytes" report (issue vllm-project#33) on this card, with the connector instrumented: 1. Upstream only ever sets is_eagle_group for DeepSeek V4, and the connector's fallback marks EVERY group as draft attention whenever speculative_config.use_eagle() is true -- which includes method=dflash and mtp. Every group then has its trailing chunk excluded from store while decoding ("EAGLE/MTP draft attention groups [0..8] detected" at boot). For dflash the draft model's KV is exactly its sliding-window layers, so the fallback now flags SW-spec groups only; boot line reads "[8]". Under other eagle-family methods the upstream fallback is unchanged. 2. The actual zero-loads mechanism, reproduced and traced here (KVarN CTX=huge + TieringOffloadingSpec, 4 GiB CPU tier): the CPU pool allocates UNIFORM blocks sized for the largest group's chunk (2,176 tokens -> ~14.6 MiB blocks, 293 of them in 4.28 GiB). Under KVarN the promoted drafter SW group's chunks are 128 tokens -- each one still occupies a full block, so a single 23k-token request consumes ~264 blocks (~184 of them for tiny SW chunks) and EVICTS THE ENTIRE TIER. Store succeeds, complete_store succeeds, and by the time the same document returns, its blocks are gone: every lookup MISS, 0 bytes ever loaded, exactly the report. On bf16 the SW group shares the large block size (the same 4096-B/token coincidence hybrid-sw-block-promote.patch documents), so the geometry never explodes and offloading uplifts -- which is why the reporter saw the split along KV dtype. The real fix (per-group block sizing, or excluding never-hittable SW groups from both store and lookup) is upstream-shaped; until then the config builder computes the waste factor and WARNS at boot with the arithmetic (group, chunk tokens, blocks-per-chunk ratio, and the cpu_bytes_to_use multiplier that would compensate), so the next person gets the diagnosis in their log instead of a day of silence. Verified on this card: boot line shows eagle groups [8]; the warning fires for group 8 (128 vs 2,176 tokens, 17x) on KVarN + connector and stays silent on configs with uniform chunk sizes. Source: syv-ai/qwen38-27b-rtx3090 patches/offload-dflash-eagle-groups.patch (at d4c4cc2). Applied as-is from the patch file; the preamble above is the patch's own. Signed-off-by: Michael Stufflebeam <cpuchip@gmail.com> Re-cut onto v0.30.0 from the qwen38/main-track resolution (e5341a1); the topic is unchanged on the 0.29 line since main-track was cut.
zyongye
added a commit
to zyongye/vllm
that referenced
this pull request
Oct 12, 2026
…parse indexer Pin vllm-project/DeepGEMM dev 00918bc (upstream 26/09/30 release via vllm-project/DeepGEMM#30, the vendoring import fix vllm-project#31, and the paged sparse page-stride relaxation vllm-project#33) and vendor the new deep_gemm/locality_domain.py. vllm-project#33 drops the sparse indexer's page-stride requirement from 512 to 256 bytes, which is what the indexer's packed 64-token pages (8448 bytes FP8, 4352 bytes MXFP4) satisfy. The MQA logits API changed: clean_logits/logits_dtype are gone, max_seqlen_k is required, prefill logits are compressed to each row's K span, and SM100 takes only MXFP8/MXFP4 with BF16 weights and returns BF16 logits. On SM100 the FP8 indexer K writers store their existing power-of-two scale as four packed UE8M0 bytes (a lossless MXFP8 encoding), FP8 Q passes unit MX scales with its scale still folded into the weights, and BF16 logits are upcast row-aware for the FP32 top-k kernels. check_wheel_deepgemm.py now imports the vendored package. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
zyongye
added a commit
to zyongye/vllm
that referenced
this pull request
Oct 12, 2026
…parse indexer Pin vllm-project/DeepGEMM dev 00918bc (upstream 26/09/30 release via vllm-project/DeepGEMM#30, the vendoring import fix vllm-project#31, and the paged sparse page-stride relaxation vllm-project#33) and vendor the new deep_gemm/locality_domain.py. vllm-project#33 drops the sparse indexer's page-stride requirement from 512 to 256 bytes, which is what the indexer's packed 64-token pages (8448 bytes FP8, 4352 bytes MXFP4) satisfy. The MQA logits API changed: clean_logits/logits_dtype are gone, max_seqlen_k is required, prefill logits are compressed to each row's K span, and SM100 takes only MXFP8/MXFP4 with BF16 weights and returns BF16 logits. On SM100 the FP8 indexer K writers store their existing power-of-two scale as four packed UE8M0 bytes (a lossless MXFP8 encoding), FP8 Q passes unit MX scales with its scale still folded into the weights, and BF16 logits are upcast row-aware for the FP32 top-k kernels. check_wheel_deepgemm.py now imports the vendored package. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
cpuchip
pushed a commit
to cpuchip/vllm
that referenced
this pull request
Oct 19, 2026
OffloadingConnector under dflash: stop eagle-flagging every KV group, and warn on the chunk-geometry that silently zeroes cross-request reuse (vllm-project#33). Two findings from tracing the "KVarN stores 41 GB, loads exactly 0 bytes" report (issue vllm-project#33) on this card, with the connector instrumented: 1. Upstream only ever sets is_eagle_group for DeepSeek V4, and the connector's fallback marks EVERY group as draft attention whenever speculative_config.use_eagle() is true -- which includes method=dflash and mtp. Every group then has its trailing chunk excluded from store while decoding ("EAGLE/MTP draft attention groups [0..8] detected" at boot). For dflash the draft model's KV is exactly its sliding-window layers, so the fallback now flags SW-spec groups only; boot line reads "[8]". Under other eagle-family methods the upstream fallback is unchanged. 2. The actual zero-loads mechanism, reproduced and traced here (KVarN CTX=huge + TieringOffloadingSpec, 4 GiB CPU tier): the CPU pool allocates UNIFORM blocks sized for the largest group's chunk (2,176 tokens -> ~14.6 MiB blocks, 293 of them in 4.28 GiB). Under KVarN the promoted drafter SW group's chunks are 128 tokens -- each one still occupies a full block, so a single 23k-token request consumes ~264 blocks (~184 of them for tiny SW chunks) and EVICTS THE ENTIRE TIER. Store succeeds, complete_store succeeds, and by the time the same document returns, its blocks are gone: every lookup MISS, 0 bytes ever loaded, exactly the report. On bf16 the SW group shares the large block size (the same 4096-B/token coincidence hybrid-sw-block-promote.patch documents), so the geometry never explodes and offloading uplifts -- which is why the reporter saw the split along KV dtype. The real fix (per-group block sizing, or excluding never-hittable SW groups from both store and lookup) is upstream-shaped; until then the config builder computes the waste factor and WARNS at boot with the arithmetic (group, chunk tokens, blocks-per-chunk ratio, and the cpu_bytes_to_use multiplier that would compensate), so the next person gets the diagnosis in their log instead of a day of silence. Verified on this card: boot line shows eagle groups [8]; the warning fires for group 8 (128 vs 2,176 tokens, 17x) on KVarN + connector and stays silent on configs with uniform chunk sizes. Source: syv-ai/qwen38-27b-rtx3090 patches/offload-dflash-eagle-groups.patch (at d4c4cc2). Applied as-is from the patch file; the preamble above is the patch's own. Signed-off-by: Michael Stufflebeam <cpuchip@gmail.com> Re-cut onto v0.30.0 from the qwen38/main-track resolution (e5341a1); the topic is unchanged on the 0.29 line since main-track was cut.
cpuchip
pushed a commit
to cpuchip/vllm
that referenced
this pull request
Oct 20, 2026
OffloadingConnector under dflash: stop eagle-flagging every KV group, and warn on the chunk-geometry that silently zeroes cross-request reuse (vllm-project#33). Two findings from tracing the "KVarN stores 41 GB, loads exactly 0 bytes" report (issue vllm-project#33) on this card, with the connector instrumented: 1. Upstream only ever sets is_eagle_group for DeepSeek V4, and the connector's fallback marks EVERY group as draft attention whenever speculative_config.use_eagle() is true -- which includes method=dflash and mtp. Every group then has its trailing chunk excluded from store while decoding ("EAGLE/MTP draft attention groups [0..8] detected" at boot). For dflash the draft model's KV is exactly its sliding-window layers, so the fallback now flags SW-spec groups only; boot line reads "[8]". Under other eagle-family methods the upstream fallback is unchanged. 2. The actual zero-loads mechanism, reproduced and traced here (KVarN CTX=huge + TieringOffloadingSpec, 4 GiB CPU tier): the CPU pool allocates UNIFORM blocks sized for the largest group's chunk (2,176 tokens -> ~14.6 MiB blocks, 293 of them in 4.28 GiB). Under KVarN the promoted drafter SW group's chunks are 128 tokens -- each one still occupies a full block, so a single 23k-token request consumes ~264 blocks (~184 of them for tiny SW chunks) and EVICTS THE ENTIRE TIER. Store succeeds, complete_store succeeds, and by the time the same document returns, its blocks are gone: every lookup MISS, 0 bytes ever loaded, exactly the report. On bf16 the SW group shares the large block size (the same 4096-B/token coincidence hybrid-sw-block-promote.patch documents), so the geometry never explodes and offloading uplifts -- which is why the reporter saw the split along KV dtype. The real fix (per-group block sizing, or excluding never-hittable SW groups from both store and lookup) is upstream-shaped; until then the config builder computes the waste factor and WARNS at boot with the arithmetic (group, chunk tokens, blocks-per-chunk ratio, and the cpu_bytes_to_use multiplier that would compensate), so the next person gets the diagnosis in their log instead of a day of silence. Verified on this card: boot line shows eagle groups [8]; the warning fires for group 8 (128 vs 2,176 tokens, 17x) on KVarN + connector and stays silent on configs with uniform chunk sizes. Source: syv-ai/qwen38-27b-rtx3090 patches/offload-dflash-eagle-groups.patch (at d4c4cc2). Applied as-is from the patch file; the preamble above is the patch's own. Signed-off-by: Michael Stufflebeam <cpuchip@gmail.com> Re-cut onto v0.30.0 from the qwen38/main-track resolution (e5341a1); the topic is unchanged on the 0.29 line since main-track was cut.
cpuchip
pushed a commit
to cpuchip/vllm
that referenced
this pull request
Oct 20, 2026
OffloadingConnector under dflash: stop eagle-flagging every KV group, and warn on the chunk-geometry that silently zeroes cross-request reuse (vllm-project#33). Two findings from tracing the "KVarN stores 41 GB, loads exactly 0 bytes" report (issue vllm-project#33) on this card, with the connector instrumented: 1. Upstream only ever sets is_eagle_group for DeepSeek V4, and the connector's fallback marks EVERY group as draft attention whenever speculative_config.use_eagle() is true -- which includes method=dflash and mtp. Every group then has its trailing chunk excluded from store while decoding ("EAGLE/MTP draft attention groups [0..8] detected" at boot). For dflash the draft model's KV is exactly its sliding-window layers, so the fallback now flags SW-spec groups only; boot line reads "[8]". Under other eagle-family methods the upstream fallback is unchanged. 2. The actual zero-loads mechanism, reproduced and traced here (KVarN CTX=huge + TieringOffloadingSpec, 4 GiB CPU tier): the CPU pool allocates UNIFORM blocks sized for the largest group's chunk (2,176 tokens -> ~14.6 MiB blocks, 293 of them in 4.28 GiB). Under KVarN the promoted drafter SW group's chunks are 128 tokens -- each one still occupies a full block, so a single 23k-token request consumes ~264 blocks (~184 of them for tiny SW chunks) and EVICTS THE ENTIRE TIER. Store succeeds, complete_store succeeds, and by the time the same document returns, its blocks are gone: every lookup MISS, 0 bytes ever loaded, exactly the report. On bf16 the SW group shares the large block size (the same 4096-B/token coincidence hybrid-sw-block-promote.patch documents), so the geometry never explodes and offloading uplifts -- which is why the reporter saw the split along KV dtype. The real fix (per-group block sizing, or excluding never-hittable SW groups from both store and lookup) is upstream-shaped; until then the config builder computes the waste factor and WARNS at boot with the arithmetic (group, chunk tokens, blocks-per-chunk ratio, and the cpu_bytes_to_use multiplier that would compensate), so the next person gets the diagnosis in their log instead of a day of silence. Verified on this card: boot line shows eagle groups [8]; the warning fires for group 8 (128 vs 2,176 tokens, 17x) on KVarN + connector and stays silent on configs with uniform chunk sizes. Source: syv-ai/qwen38-27b-rtx3090 patches/offload-dflash-eagle-groups.patch (at d4c4cc2). Applied as-is from the patch file; the preamble above is the patch's own. Signed-off-by: Michael Stufflebeam <cpuchip@gmail.com> Re-cut onto v0.30.0 from the qwen38/main-track resolution (e5341a1); the topic is unchanged on the 0.29 line since main-track was cut.
cpuchip
pushed a commit
to cpuchip/vllm
that referenced
this pull request
Nov 5, 2026
OffloadingConnector under dflash: stop eagle-flagging every KV group, and warn on the chunk-geometry that silently zeroes cross-request reuse (vllm-project#33). Two findings from tracing the "KVarN stores 41 GB, loads exactly 0 bytes" report (issue vllm-project#33) on this card, with the connector instrumented: 1. Upstream only ever sets is_eagle_group for DeepSeek V4, and the connector's fallback marks EVERY group as draft attention whenever speculative_config.use_eagle() is true -- which includes method=dflash and mtp. Every group then has its trailing chunk excluded from store while decoding ("EAGLE/MTP draft attention groups [0..8] detected" at boot). For dflash the draft model's KV is exactly its sliding-window layers, so the fallback now flags SW-spec groups only; boot line reads "[8]". Under other eagle-family methods the upstream fallback is unchanged. 2. The actual zero-loads mechanism, reproduced and traced here (KVarN CTX=huge + TieringOffloadingSpec, 4 GiB CPU tier): the CPU pool allocates UNIFORM blocks sized for the largest group's chunk (2,176 tokens -> ~14.6 MiB blocks, 293 of them in 4.28 GiB). Under KVarN the promoted drafter SW group's chunks are 128 tokens -- each one still occupies a full block, so a single 23k-token request consumes ~264 blocks (~184 of them for tiny SW chunks) and EVICTS THE ENTIRE TIER. Store succeeds, complete_store succeeds, and by the time the same document returns, its blocks are gone: every lookup MISS, 0 bytes ever loaded, exactly the report. On bf16 the SW group shares the large block size (the same 4096-B/token coincidence hybrid-sw-block-promote.patch documents), so the geometry never explodes and offloading uplifts -- which is why the reporter saw the split along KV dtype. The real fix (per-group block sizing, or excluding never-hittable SW groups from both store and lookup) is upstream-shaped; until then the config builder computes the waste factor and WARNS at boot with the arithmetic (group, chunk tokens, blocks-per-chunk ratio, and the cpu_bytes_to_use multiplier that would compensate), so the next person gets the diagnosis in their log instead of a day of silence. Verified on this card: boot line shows eagle groups [8]; the warning fires for group 8 (128 vs 2,176 tokens, 17x) on KVarN + connector and stays silent on configs with uniform chunk sizes. Source: syv-ai/qwen38-27b-rtx3090 patches/offload-dflash-eagle-groups.patch (at d4c4cc2). Applied as-is from the patch file; the preamble above is the patch's own. Signed-off-by: Michael Stufflebeam <cpuchip@gmail.com> Re-cut onto v0.30.0 from the qwen38/main-track resolution (e5341a1); the topic is unchanged on the 0.29 line since main-track was cut.
triangleXIV
pushed a commit
to tokmo-infra/vllm
that referenced
this pull request
Nov 9, 2026
…A cache (vllm-project#33) The platform set prefix_match_unit = 64 for the merged MLA/indexer cache, which then had to be synced from the workers to the engine core. Leave it to the user (--prefix-match-unit 64) instead. Claude-Session: https://claude.ai/code/session_01LpGEdMazRvsGmbqGyTvfFF Signed-off-by: zjy0516 <riverclouds.zhu@qq.com> Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
No description provided.