Repository navigation
feat(mhc): use HIP Mega-MHC by default for supported DeepSeek V4.1 - #2568
Merged
Merged
Conversation
Contributor
🏷️ CI GuideRuns automatically on every eligible PR before approval:
Heavy model tests:
|
Reuse HIP Mega-MHC MXFP8 output for QKV projection while retaining BF16 hidden states for the compressor and indexer. Carry quantized inputs through compiled execution, DSpark and bounded replay. Enable fusion whenever attention selects HIP; keep FFN BF16. Extend compilation tests to cover fused outputs and optional inputs.
Preserve fused attention quantization with DP bounded replay stages. Handle fused inputs on placeholder ranks and cover both replay paths with regression tests.
Pass the optional fused FP8 hidden tensor through the capture-break attention copy and reuse it for QKV projection. Keep BF16 hidden for auxiliary readers and the attention output buffer. Refresh the core-forward digest after syncing the copy and add CPU regressions for fused/plain inputs and bounded-query forwarding.
zufayu
approved these changes
Oct 23, 2026
valarLip
pushed a commit
that referenced
this pull request
Oct 29, 2026
…2 on Model Runner V2 (#2818) * fix(plugin): build the reused TP communicator with aiter's own __init__ 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. * fix(deepseek_v2): set pcp_split only on an installed ATOM context #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. * fix(plugin): accept reduce_results in the vLLM GDN subclasses #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> * fix(plugin): carry V4.1's MX scale across the capture break as bytes 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> * fix(docker): backport vLLM's Kimi-K3 KDA autotune key fix 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> * fix(plugin): serve DeepSeek-V4 under vLLM's Model Runner V2 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> * ci(plugin): run the V4-Flash and GLM-5.2 cells on Model Runner V2 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> --------- Co-authored-by: XiaobingZhang <xiaobingzhangupc@gmail.com> Co-authored-by: perzhang <perzhang@amd.com> Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
1 of 6 tasks
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.
DeepSeek V4.1's delayed mHC seam uses AITER HIP Mega-MHC on supported gfx950 configurations, with fused quantization for attention inputs.
Implementation
ATOM_DSV41_MHC_BACKEND=autoselects HIP forhc_mult=4, H=4096/5120/6144/7168 and contiguous FP32 fn weights on gfx950 when the AITER API is available; otherwise it falls back to Triton. Explicithipandtritonoverrides remain available.Depends on ROCm/aiter#6520.
Performance
HIP Mega-MHC vs. Triton (2026-10-14)
This comparison measures HIP Mega-MHC with BF16 output, before fused attention quantization.
Measured ATOM
258f5c69c85dcd6ff88f40799b491db8f4613ee1with AITERe284c8ec380b9dcf09f2b72eadd1c9244cd58dd0. DeepSeek-V4.1-Flash on 4 x gfx950, TP4 without EP, DSpark with 5 speculative tokens, BF16 KV / FP4 index cache, compilation level 3 and FULL graphs. Concurrency 256, ISL/OSL 1024/1024, 512 warmup requests followed by 2560 measured requests per backend, random prompts with seed 42, unlimited request rate and ignore-eos. Mono and prefix caching disabled; max batched tokens 16384, max model length 4096, block size 16.Both backends completed all 512 warmup and 2560 measured requests without errors, with exactly 1024 input / 1024 output tokens per request. Each backend used a fresh ATOM server on the same GPUs; server arguments/environment differed only in ATOM_DSV41_MHC_BACKEND. The client used identical settings with request recording enabled for both runs. Startup and warmup are excluded; all measured requests, including the first wave, are included.
HIP's observed total-token throughput gain was 1.73%, with mean TPOT 1.77% lower. DSpark's observed acceptance, including warmup, was 20.99% / 21.13%, and mean accept length 2.0494 / 2.0564 for Triton / HIP. These are end-to-end serving results including actual speculative decoding behavior, with one measured run per backend.
Neither backend had a measured TTFT above 5 seconds; maxima were 3586.83 ms / 3457.63 ms for Triton / HIP.
Existing FlyDSL disk caches were reused, while the ATOM compile cache was isolated for each server. The client repeated one input for the 512 warmup requests. ATOM/AITER source remained unchanged across both runs.
Fused attention quantization (2026-10-21)
Remeasured after merging main at ATOM
cdb3af22cwith AITERe284c8ec3. DeepSeek-V4.1-Flash on 4 x gfx950, TP4 without EP / DP attention / TBO, DSpark 5, BF16 KV / FP4 index cache, level-3 FULL graphs and decoder bounded replay; mono and prefix caching disabled. Concurrency 256, ISL/OSL 1024/1024, seed 42 and ignore-eos.Four fresh-server runs in unfused → fused → fused → unfused order, each with a fresh ATOM compile cache, 512 warmup requests and 2560 measured requests. Profiling and detailed annotations were disabled. The unfused baseline uses the same commit with only attention
fuse_quantset toFalse, covering both target and draft; the saved compiled graphs confirm the intended path in each run.Pooled throughput is total tokens divided by total measured duration across the two runs; mean latencies average the equally sized runs. All 10,240 measured requests and 2,048 warmups succeeded, with exactly 1024 input / 1024 output tokens per request. Request records, warmup boundaries, token counts, throughput and TTFT were independently checked; source and native binaries were unchanged during the runs. With two runs per variant and actual MoE routing / DSpark acceptance not held fixed, the observed +0.59% gain is specific to this workload.
Earlier, separate steady-state traces on ATOM
f456bf565(before this main merge) confirmed 43 fewer quantization kernels per cycle (40 target + 3 draft), with 0.37–0.42 ms of net kernel time saved per rank. At matching batch/token counts and overlapping context lengths, the GPU cycle fell from 56.746 to 56.475 ms (-0.48%). Those trace timings were not rerun for the serving measurements above and measure GPU execution separately from serving throughput.Validation
git diff --checkpassed.