Repository navigation
Support official vLLM 0.29.0 - #23
Conversation
There was a problem hiding this comment.
Review summary
Minor changes requested.
I checked every runtime-interface change in this PR against the upstream v0.29.0 source (98dff2a) and they all hold: V2 prepare_inputs / dispatch signatures and call conventions, enable_batch_sharded_sampling + compute_logits_local, prompt_logprobs_worker.in_progress_prompt_logprobs, ParallelLMHead.tie_weights, the AutoWeightsLoader signature change (tied aliases are now auto-skipped), WeightsMapper None drop semantics, MoERunner.gate replacing is_internal_router, the FusedMoERouter.select_experts patch boundary, GraniteAttention.use_rope, the Kimi K3 SP/aux rewrite (a faithful mirror of vllm/models/kimi_k3/nvidia/model.py, including the shape[0] == M heuristic), the OLMo3 removal, and the engine APIs used by the smoke harness. The V1 wrapper signatures are unchanged in 0.29. compileall, bash -n, and git diff --check pass on the head.
On the residual-hook rewrite: the observed values are bit-identical to the old hidden_states + residual (single fp32 to BF16 rounding in both the CUDA kernel and the native path), and DMI's ring producer launches on the current CUDA stream, so the in-place aliasing of the returned residual cannot race the capture.
Findings (details inline):
- Medium - the GPU cells do not value-check the hooks whose semantics changed;
resid_midis not in the tested selection. - Medium - the
# Adapted from ... in official vLLM 0.27.1.headers are now stale in all copied files, andtest_models_source_contract.pyenforces the stale string. - Medium - the leftover OLMo3 module cannot import on 0.29 and is still lazily registered as an oracle target.
- Low - the two new V2 guards fail open via
getattrdefaults; a real-class attribute contract test would harden them. - Low - the smoke harness relies on
EngineCoreProc-onlypause_scheduler('wait'), the privateLLM.request_counter, a log-string grep, and a host-specificenv -u LD_LIBRARY_PATHbaked into the release gate. - Low - README points at the feature branch and adds
enforce_eagerto the quickstart although graph mode is qualified; the release gate pins DMI by raw commit. - Note (pre-existing) - the
MoERunner._fse_fuse_gatecombined-logits path is not guarded; no DMI model triggers it in 0.29. - Nits -
tests/test_request_order.py:488still setsis_internal_routeron a fake runner; the Kimi PP assertion lacks upstream's message.
Not run on my side: the portable pytest suite (no vLLM 0.29 environment here), the GPU cells, and the wheel build. The reported 668/4/11 and the six bitwise-equal cells are taken as stated.
| hidden_states = self.input_layernorm(hidden_states) | ||
| else: | ||
| hidden_states, residual = self.input_layernorm(hidden_states, residual) | ||
| self.hook_resid_pre(residual) |
There was a problem hiding this comment.
The rewrite to observe the fused-norm residual is sound: both the CUDA fused_add_rms_norm kernel and the native path round the fp32 sum once to BF16, so this equals the old hidden_states + residual, and DMI's ring producer launches on the current stream (native/csrc/ring/ring_torch_op.cpp, getCurrentCUDAStream), so the later in-place update of this buffer by post_attention_layernorm cannot race the copy.
What is missing is value-level evidence. check_storage in tests/v029_smoke.py only asserts torch.isfinite for resid_pre, and resid_mid is not in the tested hook selection at all, so the change most likely to alter what researchers store is verified only by the identity/order stubs in test_v029_port_contract.py. Suggest one GPU run with resid_pre,resid_mid,resid_final selected that compares the stored tensors against the old hidden_states + residual expression (bitwise in eager, tight tolerance under graphs), or an explicit note in docs/v029-port.md that captured residuals have not been numerically compared on 0.29. Same comment applies to qwen2.py and llama.py.
| @@ -59,7 +59,6 @@ | |||
| "minimax_m2.py": "vllm/model_executor/models/minimax_m2.py", | |||
| "mistral.py": "vllm/model_executor/models/mistral.py", | |||
There was a problem hiding this comment.
These copies have now been re-audited against 0.29.0, and several were rewritten in this PR (Qwen2/Qwen3/Llama residual taps and LM-head tying, Qwen3-MoE/GLM routing, Kimi SP, Granite RoPE, GPT-2 loader), but every model and oracle file still carries # Adapted from ... in official vLLM 0.27.1.. This test asserts that exact 0.27.1 string at lines 131 and 229, so the suite currently enforces the stale provenance in exactly the files that changed. Suggest bumping the header in all copied files and the two constants here to 0.29.0 (and the docstring at tests/test_model_contracts.py:1).
|
|
||
| from __future__ import annotations | ||
| Legacy source remains for reference; it is not a supported registry target. | ||
| """ |
There was a problem hiding this comment.
Failing closed on the registry is the right call, but the leftover source is no longer inert: src/dmi_vllm_integration/models/olmo3.py:63 imports vllm.model_executor.models.olmo3, which does not exist in 0.29, so the module raises ImportError. tests/oracles/__init__.py:54 still lazily registers DMIOlmo3CompareForCausalLM, tests/compare_worker.py:35 still remaps Olmo3ForCausalLM to it, and MODEL_CLASSES in test_models_source_contract.py still lists olmo3.py. Lazy registration hides the broken import until a compare run selects it. Suggest deleting the module, oracle, registry entry, and remap (git history keeps them), or moving them under an explicitly excluded legacy/ path.
| "speculative decoding" | ||
| ) | ||
| if ( | ||
| getattr(parallel, "enable_batch_sharded_sampling", False) |
There was a problem hiding this comment.
Attribute name is correct for 0.29 (ParallelConfig.enable_batch_sharded_sampling, typed bool | None), and the guard is necessary because the runner calls compute_logits_local (gpu/model_runner.py:1416) and bypasses the DMI final-logits hook.
Hardening suggestion: this getattr(..., False) and the nested getattr(..., "in_progress_prompt_logprobs", None) near line 2039 both fail open if upstream renames the attribute, and both unit tests use SimpleNamespace fakes so they would not notice. Since the package pins an exact vLLM version, a small contract test asserting hasattr(ParallelConfig, "enable_batch_sharded_sampling") and hasattr(PromptLogprobsWorker, "in_progress_prompt_logprobs") on the real classes would turn a silent no-op into a loud failure.
| ) | ||
| if captures_router_logits and any( | ||
| runner.is_monolithic and runner.is_internal_router | ||
| runner.is_monolithic and runner.gate is not None |
There was a problem hiding this comment.
Verified against 0.29: MoERunner.gate (runner/moe_runner.py:270) is the right replacement for the removed is_internal_router, and MoERunner is still re-exported from fused_moe/layer.py.
Not a regression, just a boundary the audit does not list: when a model hands the runner both gate and shared_expert_gate, MoERunner._fse_fuse_gate is true and the runner passes combined logits (num_experts + 1 columns) to select_experts (moe_runner.py:898-900), so the observer would capture a router-logits tensor with an extra column. This path existed in 0.27.1 as well, and in 0.29 only qwen3_next.py uses it, so no DMI model is exposed today. Worth a one-line rejection on runner._fse_fuse_gate or a documented boundary in docs/vllm_contract.md.
| if [[ "$execution" == graph ]]; then graph_args=(--graph); fi | ||
| capture_id="dmi-$cell-${run_root##*/}" | ||
| for mode in stock monitored; do | ||
| env -u LD_LIBRARY_PATH \ |
There was a problem hiding this comment.
env -u LD_LIBRARY_PATH is a workaround for one host's inherited path (per docs/v029-port.md), but baked into the shared script and into gpu-release-gate.yml it can break runners that legitimately need LD_LIBRARY_PATH for CUDA. Suggest making it opt-in (for example DMI_V029_CLEAR_LD_LIBRARY_PATH=1). Same thought for the grep -q 'Directly load AOT compilation' check on line 46: correct for 0.29 (compilation/decorators.py:314), but it is a log-string contract and should be recorded as one.
| checkout, including its native-backend build. This port was tested with DMI | ||
| `c222f18c4db55f6f2d7136b1d6608c0b540b273a` (API v1). Build the native backend | ||
| against the **same PyTorch/CUDA as vLLM 0.29.0**, not a 0.27 environment. | ||
| From the `vllm-0.29-support` integration checkout, run: |
There was a problem hiding this comment.
After merge this should point at main or the release tag rather than the feature branch.
| llm = LLM( | ||
| model="Qwen/Qwen3-0.6B", | ||
| worker_cls="dmi_vllm_integration.worker.DMXGPUWorker", | ||
| enforce_eager=True, |
There was a problem hiding this comment.
This README and docs/v029-port.md both say default compilation plus CUDA graphs are qualified, but the quickstart here (and --enforce-eager in the serve example below) reads as a requirement. Suggest dropping it, or stating plainly that graph mode is supported and eager is optional.
| with: | ||
| repository: ProjectDMX/DMI | ||
| ref: v1.1.0 | ||
| ref: c222f18c4db55f6f2d7136b1d6608c0b540b273a |
There was a problem hiding this comment.
Pinning a raw commit is reproducible but hard to read next to ci.yml's v1.1.0 / main matrix and pyproject's DMI>=1.1.0,<2.0. c222f18 is DMI 1.2.0 (per pyproject.toml in that checkout); tagging it and pinning the tag would make the release gate self-describing.
| assert hidden_states is not None | ||
| assert residual is not None | ||
| if not get_pp_group().is_last_rank: | ||
| assert not self.use_sequence_parallel |
There was a problem hiding this comment.
Nit: upstream's equivalent assertion carries a message ("Currently, SP is not supported with PP"); worth keeping it here for the same reason.
|
Addressed the review findings in the follow-up commit:
Validation: 680 passed, 4 skipped, 11 deselected in the portable gate. The three router-state skips pass in clean isolated test processes; the remaining optional layer-range module requires a newer DMI API than the pinned core. All 21 production lazy targets import without the native stub. Wheel and sdist build from a clean temporary source snapshot (no OLMo3 in the wheel); Committed GPU receipt: |
XbzOnGit
left a comment
XbzOnGit
left a comment
There was a problem hiding this comment.
Request changes: revert the residual-hook placement, keep this PR to the version bump
I ran this branch on hardware (3x RTX 4090, CUDA 13.0, DMI c222f18, torch 2.13.0+cu130, vLLM 0.29.0 from PyPI) against main at 420bafb on vLLM 0.27.1 in a second environment on the same torch. The forced parts of the port are correct and I verified them: all six smoke cells reproduce your receipts (744 rows each, 1,416 / 1,368 for the residual cell), request order at six topologies, online storage bitwise at TP1/TP2/PP2, DMILLM, Qwen2-MoE routing at TP2 (69,920 / 69,920 rows), compile-cache reload, V2 parity, and the CPU suite at 680 / 4 / 11. Thank you for that work.
One elective change does not hold up under measurement and I am asking for it to be reverted before merge: observing the fused norm's returned residual instead of computing hidden_states + residual for hook_resid_pre, hook_resid_mid, and hook_resid_final in qwen2.py, qwen3.py, llama.py.
1. The old placement perturbs the model's numerics less
Gate step 09 (Qwen3-0.6B, TP2, compiled, greedy) fails on this branch: request 3 flips from ' purple\nThis is the' (198) to ' purple\n\nThis is the' (271). I bisected it with one engine per process, recording top-5 logprobs at every step (4 prompts, 12 steps, 60 records), each monitored run compared against the same version's stock run.
| Version | Hooks in the compiled TP2 graph | Records differing from stock / 60 | Max logprob shift | Greedy token at the tie |
|---|---|---|---|---|
| 0.27.1 | residual only (old placement) | 15 | 0.095 | 198, tie position untouched |
| 0.27.1 | all hooks | 60 | 0.368 | 198 (exact tie) |
| 0.29 | residual only (new placement) | 60 | 0.867 | 198 (exact tie) |
| 0.29 | q,k,v,z |
60 | 0.802 | 198 (exact tie) |
| 0.29 | residual + q,k,v,z |
60 | 0.546 | 271, flipped |
| 0.29 | all hooks (2 shared-cache + 3 fresh-cache compiles + 1 null mode, all bitwise identical) | 60 | 0.546 | 271, flipped |
| 0.29 | final_logits only |
0 | 0 | 198 |
| 0.29 | ln1, attn_out, ln2, mlp_in, mlp_out |
0 | 0 | 198 |
The effect is deterministic across fresh compiles and identical in null mode, so it is the presence of the hook ops in the compiled graph, not the copies. The FX graphs in the compile cache have identical model ops in every variant; only the ring.producer_prefix nodes differ. The residual hooks are one of the two ingredients of the flip; the other is q/k/v/z, which predates this PR. At TP1 under compilation both placements are bitwise clean, and in eager TP2 tokens match, so the difference is confined to compiled TP>1, which is exactly where the old placement is the safer one: it reads the norm's inputs, which are already materialized buffers, and writes a private temporary, so it cannot change how inductor schedules the fused norm. The new placement adds an opaque reader to the fused region's output.
The 15 records that differ on 0.27.1 with the old placement carry the same values 0.29 stock produces at those positions, i.e. ordinary rounding-order noise of the same size as the upstream version bump itself (stock 0.27.1 vs stock 0.29: max 0.095). The one control still missing is the old placement compiled under 0.29; it is queued and I will attach it here. On everything measured so far the old placement is the less perturbing one.
2. "Authoritative" buys nothing measurable, and the recorded value is from a more perturbed forward
Your own follow-up oracle shows the two placements record bit-identical values: 1,368 / 1,368 residual rows equal to the old hidden_states + residual expression. I reproduced that exactly. Both the CUDA kernel and the native decomposition round the fp32 sum once, so there is no configuration in this PR's matrix where the fused norm's residual differs from the recomputed one. The case where it would matter (a fused all-reduce + norm kernel producing the residual from fp32 partial sums) is not enabled and not tested here.
Meanwhile the recorded tensor is only "authoritative" for the monitored run, and under compiled TP2 that run is now further from the unmonitored model than before. A residual stream that is bit-exact to a forward the user did not ask for is less accurate for research use, not more.
3. The performance gain is not measurable
Interleaved A/B, same GPU, 2 warmup + 5 timed iterations, tests/benchmark.py:
| Row | Mode | 0.27.1 (old placement) | 0.29 (new placement) |
|---|---|---|---|
| Qwen3-0.6B, graph, TP1, residual hooks only | ring_active | 174.6 ms (std 1.3) | 176.8 ms (std 2.9) |
| Qwen3-0.6B, graph, TP1, residual hooks only | ring_db | 191.4 ms (std 10.6) | 177.2 ms (std 1.1) |
| Qwen3-0.6B, graph, TP1, all hooks | ring_active | 287.2 ms (std 45.7) | 286.9 ms (std 46.1) |
| Qwen3-0.6B, graph, TP1, 32x128, all hooks | ring_active | 806.4 ms (std 46.3) | 812.6 ms (std 44.4) |
| Llama-3.1-8B, graph, TP1, all hooks | ring_active | 1243.7 ms (std 2.6) | 1233.8 ms (std 3.1) |
The saved add is one BF16 elementwise op per hook site. Under compilation inductor can fuse it into the norm's reduction kernel as an extra store, which is why it costs nothing to keep; under eager it is one extra kernel on a 0.6B model, below the noise floor.
4. A version bump should be a version bump
Diffing upstream v0.27.1 against v0.29.0 at every boundary this PR touches, the forced changes are the V2 prepare_inputs/dispatch signatures, is_internal_router to gate, the AutoWeightsLoader keyword removals, OLMo3, Granite use_rope, the Kimi K3 SP mirror, the pins, and the prompt-logprobs attribute fix. That is on the order of a hundred source lines and all of it is right. The residual-hook rewrite, the hook_v view cleanup, Llama compute_logits_local, and the tie_weights alignment are not needed for 0.29 to work; the loader change stands on its own. Please move them to a follow-up so that the port can be reviewed and bisected as a port, and so that a numerics change to three model copies gets its own TP2 compiled evidence before it lands.
Concretely
- Restore the explicit-add placement at the six hook sites in
qwen2.py,qwen3.py,llama.py; keep every loader and LM-head fix in those files. Adjusttests/test_v029_port_contract.pyaccordingly; your residual oracle already computes the old expression and continues to apply. - Split the other elective items into a separate PR.
- Independently of placement, gate step 09 compares greedy tokens of two differently compiled TP2 engines and passed on 0.27.1 only because the tie broke toward the lower id. It needs a logprob tolerance or an eager-mode comparison at TP>1, otherwise it will keep flipping.
| hidden_states = self.input_layernorm(hidden_states) | ||
| else: | ||
| hidden_states, residual = self.input_layernorm(hidden_states, residual) | ||
| self.hook_resid_pre(residual) |
There was a problem hiding this comment.
Please restore the previous placement here and at hook_resid_mid / hook_resid_final (and the same three sites in qwen2.py and llama.py): compute hidden_states + residual before the norm and pass that temporary to the hook. Values are bit-identical either way (your 1,368-row oracle), the saved add is not measurable, and under compiled TP2 this placement adds an opaque reader on the fused norm's output and perturbs the model's logits (60/60 records, up to 0.87 nats vs stock) where the old placement perturbs 15/60, 0.095. Details and tables in the review body.
|
Addressed the scope correction in
All three local GPUs were occupied, so I did not launch a GPU rerun or stop any process. The PR description and audit now clearly separate the old GPU receipts from corrected-head qualification. The same-version compiled TP2 control and bounded V1/V2 rerun remain pending; I am not claiming the rollback fixes the observed divergence. The compiled release gate and exact comparison thresholds are unchanged, so the outstanding step-09 issue is not hidden by a tolerance increase or an eager substitution. No review threads were auto-resolved and no optimization follow-up PR was opened. |
Summary
Port the out-of-tree DMI integration from official vLLM 0.27.1 to 0.29.0, keeping separate V1 and V2 workers. This PR changes DMI-vLLM-Integration only; it does not change the DMI root/gitlink, patch upstream vLLM, or publish a release tag.
2026-09-18 review update: reverted the nonessential residual-hook rewrite, Qwen3 V-view cleanup, added Llama
compute_logits_local, and optional weight-tying alignment. The corrected head preserves the pre-port observation and constructor behavior. GPU requalification of this correction is pending; historical GPU receipts are not results for the corrected head.Required adaptations retained
prepare_inputs(..., batch_req_state, ...)argument and graph dispatch's optionalmax_query_len; preserve packed ordering and capacity safeguards.AutoWeightsLoaderkeyword arguments and adapt GPT2's weight mapper. Preserve existing LM-head construction and weight aliases.Exact basis and dependencies
v0.27.1/6e448d0ea9bf3d88d898b65449ca6dc2aec170acv0.29.0wheel / source tag98dff2a81d747d1dba01a47f939f48c3526d4206420bafba42580291f1551ed129b5c55d5ac5799ec222f18c4db55f6f2d7136b1d6608c0b540b273a(1.2.0 / API v1)Corrected-head validation
420bafb: in the seven affected production model files, only necessary Qwen3/Llama loader, Qwen3-MoE routing and Granite RoPE method changes remain. Restored constructors, residual forwards and Qwen3 attention views match the base. Llama's already-post-normresid_midis unchanged.git diff --checkand smoke-wrapper shell syntax pass.Historical GPU evidence and pending checks
The earlier implementation (through
5ce33c9) passed the six bounded Qwen3-0.6B TP1 cells and the eager residual-value diagnostic: 744 storage rows per ordinary cell, 1,416 total / 1,368 residual rows in the diagnostic. Earlier wheel/import evidence and exact cell definitions are retained indocs/v029-port.md.The reviewer reported a compiled TP2 token divergence on that implementation. The same-version old-placement control was still pending in that review. Reverting the rewrite is not proof that this divergence is fixed.
Before claiming corrected-head GPU transparency, rerun the bounded V1/V2 workload and the same-version compiled TP2 control. Existing exact token/logit thresholds and the compiled release gate are unchanged; no tolerance was widened and no eager substitution hides this failure. No optimization follow-up PR is introduced.
Reproduce bounded TP1 cells with matching runtime/native, cached Qwen3-0.6B weights and ClickHouse:
Import/CPU success does not qualify other models or topologies. See
docs/v029-port.md,docs/v029-boundaries.tsv, anddocs/vllm_contract.mdfor historical evidence, limitations and outstanding GPU checks. This PR is not a blanket support or release qualification claim.