Repository navigation
fix(plugin): green the vLLM 0.31 plugin CI cells; V4-Flash and GLM-5.2 on Model Runner V2 - #2818
Merged
Merged
Conversation
ReusedPyNcclCudaCommunicator copied CudaCommunicator.__init__ by hand to avoid a second ncclCommInitRank, so every field aiter adds is missing on it. aiter now reads self.gluon_ar_plain in all_reduce, and every vLLM-plugin run dies at engine init with AttributeError. Build a plain CudaCommunicator instead, handing it vLLM's PyNccl communicator in place of the one its __init__ would create. aiter then sets every field it reads (Gluon all-reduce, ownership, all2all state) and RCCL still inits no duplicate communicator; a check fails loudly if aiter ever builds its own.
#2725 writes get_forward_context().context.pcp_split on every forward, but the vLLM plugin installs no ATOM context, so every DeepSeek-V2-family model on the plugin (GLM-5.2 in CI) dies on its first forward with 'NoneType' object has no attribute 'pcp_split'. The plugin runs no PCP and the readers short-circuit on PCP being off, so the flag is only needed where the context exists.
#2520 made the Qwen3.5 decoder layer pass reduce_results to its GatedDeltaNet, so the norm after it can fuse the TP all-reduce. The plugin swaps that class for Qwen3_5GatedDeltaNetVllm, whose __init__ spelled out the old signature, and model construction failed: TypeError: Qwen3_5GatedDeltaNetVllm.__init__() got an unexpected keyword argument 'reduce_results' Accept it and pass it through. allreduce_norm_fusion_enabled() keeps the fusion off in plugin mode, so the plugin layer passes True and computes what it did before #2520. Qwen3NextGatedDeltaNetVllm narrowed its parent's signature the same way; the Qwen3-Next layer does not pass the argument today, so it is aligned before it does. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
vLLM's eager_break_during_capture weak-refs every argument of the call it breaks out through torch.ops._C.weak_ref_tensor, a stable-ABI op that accepts the fp8 e4m3/e5m2 family but not float8_e8m0fnu. Since the HIP Mega-MHC norm became the default (#2568) it hands QKV an e8m0 MX scale, and graph capture of the V4.1 plugin cell failed: Not yet supported ScalarType 44, please file an issue ... _forward now passes an e8m0 scale as a uint8 view of the same storage, with its dtype as a non-tensor keyword, and _attend_block views it back. The weak ref still points at the caller's buffer on every replay. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
The base image moved from Triton 3.7.0 to 3.8.0 between the 10-13 and
10-25 nightlies. 3.8 checks autotune keys against the kernel's arguments
at decoration time, and the pinned vLLM (db9527a46) keys
chunk_kda_fwd_kernel_intra_sub_chunk on NC, which is not one. Importing
vllm.models.kimi_k3 raised, and model inspection failed for every
Kimi-K3 deployment:
ValueError: key contains names that are not kernel arguments: NC.
vllm-project/vllm#60446 (fc3931739) drops NC from the key, upstream on
10-09 but after the pinned commit. Apply it as a patch on top of the
pinned commit. A scan of every @triton.autotune in the image's vLLM finds
no other key naming a non-argument.
The PR job reuses the nightly vLLM image unless the vLLM commit or the
ATOM base recipe changes, so a change to the vLLM recipe or its patches
was never built by its own PR. Rebuild the vLLM image when either changes.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
vLLM runs DeepseekV4ForCausalLM on V1 on ROCm; with
VLLM_USE_V2_MODEL_RUNNER=1 the V4 cell died in V2's kernel warmup:
RuntimeError: DeepSeek-V4 request(s) resumed at chunk_start [2, 2, ...]
into a fresh slot, but no state checkpoint path is bound
The state-slot hooks were V1's alone. V2 has no _update_states; it drops
finished and preempted requests in finish_requests, which execute_model
calls first. Hook that method on V2 (the same body as V1's
_update_states hook) and initialize_kv_cache on every runner class, so
slots come back and the checkpoint groups get their ids. Without this,
slots would run out after max_num_seqs requests.
V2's kernel warmup also schedules real-looking requests. Their prefill
was built before the forward bound the proxy cache, so it took no slots
and the following decode found them fresh at chunk_start 2. The metadata
builder now binds from the proxy, as V4.1's does, with the arguments the
wrapper's forward would pass. V1 is unchanged: its dummy runs carry no
request ids, so they still take throwaway arange slots after the bind.
Prefix caching (state-checkpoint stores run from V1's _model_forward) and
the MTP draft (a V1 get_kv_cache_spec hook, and V2 sizing the drafter's
hidden buffer from the class) still hook V1 alone; under V2 they are
refused at model construction instead of served wrong.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
vLLM keeps DeepseekV4ForCausalLM and GlmMoeDsaForCausalLM on V1 on ROCm (ROCM_DEFAULT_MRV1_ARCHITECTURES) unless VLLM_USE_V2_MODEL_RUNNER is set, so no PR cell ran either model on V2; the other four cells already select V2 by default. Set it on both cells. GLM-5.2 needs no code change for V2; DeepSeek-V4 needs the previous commit. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Contributor
🏷️ CI GuideRuns automatically on every eligible PR before approval:
Heavy model tests:
|
valarLip
approved these changes
Oct 29, 2026
valarLip
pushed a commit
that referenced
this pull request
Oct 31, 2026
…LM image (#2835) #2818 fixed Kimi-K3's import failure under Triton 3.8 ValueError: key contains names that are not kernel arguments: NC. by applying vllm-project/vllm#60446 as a patch on the vLLM build. That only reaches a deployment once its vLLM image is rebuilt: the nightly vllm-latest built before #2818 merged still fails every PR's Kimi-K3 cell. Fix it where the plugin fixes vLLM. register_model wraps triton.autotune: for a kernel defined under vllm.models.kimi_k3 it drops the key names that are not kernel arguments, which is what #60446 does by hand, and passes every other kernel through unchanged. The decorator runs at import, so the wrapper has to precede any import of vllm.models.kimi_k3; vLLM's model-inspection subprocess, where the failure surfaced, runs load_general_plugins and so this hook first. A vLLM that carries #60446 names no such key and nothing is dropped; an INFO line names each dropped key. The image patch and its dockerfile steps are removed; the vLLM recipe is back to its state before #2818. Co-authored-by: perzhang <perzhang@amd.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.
Depends on #2816, whose two commits sit at the base of this branch. Ready for review.
With #2816 the plugin cells get past engine init; three still fail on 0.31.0. V4-Flash and GLM-5.2 never ran on Model Runner V2 in CI (vLLM keeps both on V1 on ROCm). This fixes the three and moves both to V2.
Dev record
reduce_results, since (Qwen3.5) Optimize decode: add fusion, remove copybuffer, opt _replayssm_fwd_kernel #2520), V4.1 graph capture (ScalarType 44, the e8m0 scale from feat(mhc): use HIP Mega-MHC by default for supported DeepSeek V4.1 #2568), Kimi-K3 model inspection (Triton 3.8 rejects vLLM'sNCautotune key).finish_requests, the V4 metadata builder binds from the proxy as V4.1's does, and V4 prefix caching and MTP on V2 are refused at startup. Check first: V1 now binds at the first build, not the first forward, of the same dummy run._kda_runscale crosses the same break; float32 under the CI's ptpc_fp8.Named result
MI355X,
vllm-latestimage (vLLM0.31.1.dev0+gdb9527a46), each cell's CI config, gsm8k flexible-extract:This PR's CI (40885171738): all six cells pass, all on V2. V4-Flash 0.9583, GLM-5.2 0.9287, M3 0.9378, Kimi-K3 0.9613, V4.1 0.9249, Qwen3.5 MXFP4 0.8400.
Gates
tests/pluginfails the same 18 on base and branch.Evidence
TypeError: Qwen3_5GatedDeltaNetVllm.__init__() got an unexpected keyword argument 'reduce_results'; V4.1RuntimeError: ... stableivalue_conversions.h:530, Not yet supported ScalarType 44fromweak_ref_tensorinbreakable_cudagraph._weak_ref_capture_arg; Kimi-K3ValueError: key contains names that are not kernel arguments: NC.import vllm.models.kimi_k3raises theNCerror unpatched and succeeds withdocker/patches/vllm-kimi-k3-kda-autotune-key.patchapplied. An AST scan of every@triton.autotunein the image's vLLM and in aiter finds no other key that names a non-argument. Base image Triton: 3.7.0 in the 10-13 nightly, 3.8.0 from the 10-25 nightly.RuntimeError: DeepSeek-V4 request(s) resumed at chunk_start [2, 2, ...] into a fresh slot, fromvllm/v1/worker/gpu/warmup.py_run_decode_step.Qwen/Qwen3.5-397B-A17B-FP8; the CI cell'samd/Qwen3.5-397B-A17B-MXFP4is not on the test host. In plugin modeallreduce_norm_fusion_enabled()is off, so the layer passesreduce_results=True, as before (Qwen3.5) Optimize decode: add fusion, remove copybuffer, opt _replayssm_fwd_kernel #2520.test_vllm_deepseek_v4_proxy_state_arena_layout.py(12,DeepSeek V4 proxy cache must be block-major contiguous),test_sglang_model_wrapper.py(5),test_rtpllm_model_wrapper.py(1).🤖 Generated with Claude Code