{{ message }}
Repository navigation
[Perf][Attention] Reduce GLM sparse MLA preparation overhead - #57458
Merged
vllm-bot merged 8 commits intoSep 22, 2026
Merged
Conversation
GirasoleY
force-pushed
the
codex/glm53-sparse-mla-prep
branch
3 times, most recently
from
September 18, 2026 01:51
1d499d7 to
2221877
Compare
Reuse a contiguous query when its rotary tail is empty. Preserve the concatenation fallback for nonempty rotary dimensions and strided queries. Cover both custom-op dispatch modes and FP8 value preservation. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Summer Yang <girasoleyang@gmail.com>
Use one padded program per query for counted sparse capacities up to 4096, including GLM's 2176 entries. Initialize padding and store counts in the remap kernel, removing two fills and atomic accumulation. Preserve the larger multi-tile fallback. Name the device entrypoint sparse_mla_index_remap_kernel for trace attribution. Extend remap and DCP reference tests for 2176 and the 4224-entry fallback. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Summer Yang <girasoleyang@gmail.com>
GirasoleY
force-pushed
the
codex/glm53-sparse-mla-prep
branch
from
September 18, 2026 01:55
2221877 to
6728a18
Compare
GirasoleY
marked this pull request as ready for review
September 18, 2026 02:08
GirasoleY
requested review from
AndreasKaratzas,
LucasWilkinson,
MatthewBonanni,
WoosukKwon,
mgoin,
pavanimajety,
tlrmchlsmth,
yewentao256 and
zyongye
as code owners
September 18, 2026 02:08
![claude[bot]](https://avatars.githubusercontent.com/in/1236702?s=60&v=4){kind=small}
ZJY0516
reviewed
Sep 18, 2026
| indices_ptr, | ||
| counts_ptr, | ||
| safe_lengths_ptr, | ||
| NUM_TOKENS: tl.constexpr, |
Member
There was a problem hiding this comment.
This will trigger JIT recompilation
Contributor
Author
There was a problem hiding this comment.
Fuse dummy-entry preparation and safe-length calculation while preserving raw valid counts. Resident TP skips this work and all empty-query masking. DCP masks only empty-query LSE to -inf; its existing merge suppresses zero-weight outputs, including NaNs. Preserve PyTorch output masking for HiSparse without LSE. Validate native BF16/FP8 attention and DCP merge behavior, safe lengths, strided storage, remapping, and query packing with focused regression cases. Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Summer Yang <girasoleyang@gmail.com>
GirasoleY
force-pushed
the
codex/glm53-sparse-mla-prep
branch
from
September 18, 2026 03:20
6728a18 to
736f29a
Compare
Contributor
Author
|
/ci run |
This was referenced Sep 18, 2026
22 of 39 tasks
Member
|
/ci run --allow-stale |
ZJY0516
enabled auto-merge (squash)
September 21, 2026 08:23
3 tasks
4 tasks
1 task done
michael-gpu
pushed a commit
to practical-inference/vllm-ampere-extended
that referenced
this pull request
Oct 20, 2026
…oject#57458) Signed-off-by: Summer Yang <girasoleyang@gmail.com> Co-authored-by: Codex <noreply@openai.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> Co-authored-by: Jiangyun Zhu <riverclouds.zhu@qq.com>
mgehre-amd
pushed a commit
to ROCm/vllm
that referenced
this pull request
Oct 20, 2026
* [DSV4.1] Add encoder cuda graph support for deepseek-v4.1-flash (vllm-project#56625) Signed-off-by: Isotr0py <Isotr0py@outlook.com> * [Feature][Humming] Humming feature integration (vllm-project#56685) Signed-off-by: jinzhen.ljz <jinzhen.ljz@antgroup.com> Signed-off-by: mgoin <mgoin64@gmail.com> Co-authored-by: Codex <noreply@openai.com> Co-authored-by: mgoin <mgoin64@gmail.com> * [Misc] Rename --enable-mamba-fine-grained-prefix-cache (vllm-project#53945 follow-up) (vllm-project#57382) Signed-off-by: Roi Koren <roik@nvidia.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> * [Refactor] Remove dead kernel code (vllm-project#57621) Signed-off-by: yewentao256 <zhyanwentao@126.com> * [ROCm][CI] Query HIP device memory for test GPU teardown waits. (vllm-project#57450) Signed-off-by: Sheral Kumar <shekumar@amd.com> Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com> Co-authored-by: Andreas Karatzas <Andreas.Karatzas@amd.com> Co-authored-by: Codex <noreply@openai.com> Co-authored-by: Andreas Karatzas <akaratza@amd.com> * [KVConnector][MoRIIO] Transfer hybrid mamba/KDA recurrent state in READ mode (vllm-project#51052) Signed-off-by: Yichao Zhu <Yichao.Zhu@amd.com> * [XPU][UT]Bugfix when the process can't see all the world_size meet accuracy issue. (vllm-project#57779) Signed-off-by: yisheng <yi.sheng@intel.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com> * [Perf][DSV4.1] Overlap mHC coefficients for small TP batches (vllm-project#57603) Signed-off-by: Woosuk Kwon <woosuk@inferact.ai> Co-authored-by: Codex <noreply@openai.com> * [Bugfix][Spec Decode] Stop dummy draft decode steps from writing KV through stale block-table rows (vllm-project#56734) Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai> Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com> * [Docs] Add ECMooncakeConnector usage example for EPD (vllm-project#57742) Signed-off-by: jiangkuaixue123 <jiangxiaozhou111@163.com> Co-authored-by: Codex <noreply@openai.com> * [Bugfix][KV Connector] Propagate cache reset failure during sleep (vllm-project#54581) Signed-off-by: CrazyCoder <18235787078@163.com> Signed-off-by: aoshen02 <aoshen@inferact.ai> Co-authored-by: aoshen02 <aoshen@inferact.ai> Co-authored-by: Codex <noreply@openai.com> * [Doc] Refresh Kthena integration guide (vllm-project#57314) Signed-off-by: Zhou Zihang <z@mcac.cc> Co-authored-by: Jee Jee Li <pandaleefree@gmail.com> * [ROCm][DSv4.1] Keep the Engram tables in host memory on ROCm (vllm-project#57491) * [Profiler] Unify platform-aware torch profiling (vllm-project#57460) * [BUGFIX][HY4] Record indexer completion event for full CUDA graph capture (vllm-project#57811) Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai> * [Build] Fix DeepGEMM CUDA 12.9 release builds (vllm-project#57554) Signed-off-by: khluu <khluu000@gmail.com> Co-authored-by: OpenAI Codex <noreply@openai.com> Co-authored-by: Jee Jee Li <pandaleefree@gmail.com> * [Perf][DSV4.1] Fuse TP all-reduce with mHC input preparation (vllm-project#57643) Signed-off-by: Woosuk Kwon <woosuk@inferact.ai> Co-authored-by: Codex <noreply@openai.com> * [Bugfix] DiffusionGemma: fix multimodal support (vllm-project#57589) Signed-off-by: Matt Mastracci <matthew@mastracci.com> Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com> Co-authored-by: Lucas Wilkinson <LucasWilkinson@users.noreply.github.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> * [Core][KDA] Generalize Mamba prefill checkpoint builder and exporter (vllm-project#57783) Signed-off-by: chaunceyjiang <chaunceyjiang@gmail.com> Signed-off-by: zjy0516 <riverclouds.zhu@qq.com> Co-authored-by: zjy0516 <riverclouds.zhu@qq.com> * [XPU][CI] fallback to 7-args of moe_align_block_size (vllm-project#57855) Signed-off-by: zhenwei-intel <zhenwei.liu@intel.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com> * [ROCm][Tests] Avoid pooling cleanup waits against live shared engines (vllm-project#57864) Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com> Co-authored-by: OpenAI Codex <noreply@openai.com> * [MRV2][XPU] use xpu sample kernel in mrv2 sampler (vllm-project#57277) Signed-off-by: zhenwei-intel <zhenwei.liu@intel.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com> * [Tests] Release HF embedding weights after prompt extraction (vllm-project#57868) Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com> Co-authored-by: OpenAI Codex <noreply@openai.com> * [Bugfix] Fix Whisper engine crash on audio clips longer than 30s (vllm-project#57769) Signed-off-by: Tianyao Wu <rayroy31@gmail.com> * [Bugfix][Model] Mistral3: align placeholder grid with public processor (vllm-project#53758) Signed-off-by: Oliver Holworthy <1216955+oliverholworthy@users.noreply.github.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> * [Bugfix] Fix xgrammar feature gate bypass when JSON Schema type is a list (vllm-project#48416) Signed-off-by: john <6712432+CaiJohn@users.noreply.github.com> Co-authored-by: john <6712432+CaiJohn@users.noreply.github.com> Co-authored-by: Codex <noreply@openai.com> * [Perf][Attention] Avoid redundant sparse attention metadata operations (vllm-project#57885) Signed-off-by: Woosuk Kwon <woosuk@inferact.ai> Co-authored-by: OpenAI Codex <noreply@openai.com> * [Bugfix][DSV4.1] Restrict mHC overlap to full CUDA graphs (vllm-project#57874) Signed-off-by: Woosuk Kwon <woosuk@inferact.ai> Co-authored-by: OpenAI Codex <noreply@openai.com> * [Perf][ROCm] Add a ROCm path for Hy4 and compile the backbone (vllm-project#57526) Signed-off-by: Aakif Nawaz <aakif.nawaz@amd.com> * [Model][Engram] Share host tables across co-located DP replicas by default (vllm-project#57651) Signed-off-by: Juntian Liu <Juntianl777@gmail.com> Signed-off-by: zjy0516 <riverclouds.zhu@qq.com> Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com> Co-authored-by: zjy0516 <riverclouds.zhu@qq.com> * [CI][ROCm] Make Python-only installation failures blocking (vllm-project#57870) Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com> Co-authored-by: OpenAI Codex <noreply@openai.com> * [AMD][Minimax-M3][perf] Enable packed LBHNC AITER QK-norm fusion for MiniMax-M3 on ROCm (vllm-project#54535) Signed-off-by: Wei-Ting Liao <wei-ting.liao@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Shanshan Shen <467638484@qq.com> * [CI][ROCm] Mirror remaining portable suites without duplicate coverage (vllm-project#57877) Signed-off-by: Andreas Karatzas <akaratza@amd.com> Co-authored-by: OpenAI Codex <noreply@openai.com> * [CI][ROCm] Increase timeout for MI300 Distributed DP Extended (vllm-project#57903) Signed-off-by: Djordje Ramic <djoramic@amd.com> * [AuxOutput] Add block-keyed storage for routed-expert outputs (vllm-project#45635) Signed-off-by: aoshen02 <aoshen@inferact.ai> Signed-off-by: Nick Hill <nickhill123@gmail.com> Co-authored-by: aoshen02 <aoshen@inferact.ai> Co-authored-by: OpenAI Codex <codex@openai.com> Co-authored-by: Roger Wang <hey@rogerw.io> Co-authored-by: Nick Hill <nickhill123@gmail.com> Co-authored-by: Codex <noreply@openai.com> * [Bugfix][ROCm][DSv4.1] Disable SWA bounded replay on ROCm (vllm-project#57906) Signed-off-by: Fangzhou Ai <fangzhou.ai@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> * [Test] Organize structured output utility tests (vllm-project#57545) Signed-off-by: john <6712432+CaiJohn@users.noreply.github.com> Co-authored-by: john <6712432+CaiJohn@users.noreply.github.com> * [CI][ROCm] Add ten AMD parity groups (vllm-project#57876) Signed-off-by: Andreas Karatzas <akaratza@amd.com> Co-authored-by: OpenAI Codex <noreply@openai.com> * [Security] Prefer fresh multimodal payloads over a stale receiver cache (vllm-project#57833) Signed-off-by: Juan Pérez de Algaba <jperezde@redhat.com> Co-authored-by: Cursor <cursoragent@cursor.com> * [CI][XPU] Deselect Ray UT in XPU V1 test Job (vllm-project#57889) Signed-off-by: Chaojun Zhang <chaojun.zhang@intel.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com> * [Bugfix][Multimodal] Tolerate malformed EXIF in Molmo 2 image preprocessing (vllm-project#57234) Signed-off-by: hotragn <hotragn.pettugani_2024@woxsen.edu.in> Co-authored-by: Claude Opus 5 <noreply@anthropic.com> * [Frontend] Add streaming parity tests and docs for derender (vllm-project#57922) Signed-off-by: Martin Hickey <martin.hickey@ie.ibm.com> * [ROCm][Tests] Reduce host memory when downcasting FP32 HF references (vllm-project#57869) Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com> Co-authored-by: OpenAI Codex <noreply@openai.com> * [MM] Move `supports_multimodal_inputs` and cache out of registry (vllm-project#57913) Signed-off-by: DarkLight1337 <tlleungac@connect.ust.hk> * [Pooling] Fix normalization of chunked long-text embeddings (vllm-project#57498) Signed-off-by: Taneem Ibrahim <taneem.ibrahim@gmail.com> * [Bugfix][Frontend] Accept diarized transcription responses in run-batch (vllm-project#57948) Signed-off-by: Sergio Figueras <sergio@yourecm.com> * [ROCm][DSv4] Enable DSpark adaptive verification (vllm-project#52362) Signed-off-by: Tuukka Sarvi <tuukka.sarvi@amd.com> Signed-off-by: larryli2-amd <larryli2@amd.com> Co-authored-by: larryli2-amd <larryli2@amd.com> * [Bugfix] Fix external LB DP rank handling when replicas share nodes (vllm-project#53743) Signed-off-by: Tony Lin <tony.lin@intel.com> * [docs] Fix legacy hf CLI references (vllm) (vllm-project#57958) Signed-off-by: Wauplin <lucainp@gmail.com> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com> * [Bugfix][NIXL] Fix DCP pulls across MLA cache regions (vllm-project#57389) Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com> Co-authored-by: OpenAI Codex <noreply@openai.com> * [ROCm] Refactor tuned gemms (vllm-project#55001) Signed-off-by: Andy Friedrich <afriedri@amd.com> Signed-off-by: afriedri <afriedri@amd.com> Co-authored-by: Douglas Lehr <91553416+dllehr-amd@users.noreply.github.com> Co-authored-by: Shanshan Shen <467638484@qq.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> * [Bugfix] unskip InternViT test for transformers v5 compatibility (vllm-project#55767) Signed-off-by: sahil <sahil@example.com> Co-authored-by: Cyrus Leung <tlleungac@connect.ust.hk> * [MM] Move get_dummy_processor_inputs into MM processor (vllm-project#57967) Signed-off-by: DarkLight1337 <tlleungac@connect.ust.hk> * [ROCm][CI] Use ROCm backend for DeepSeek V4.1 ViT test (vllm-project#57931) Signed-off-by: Djordje Ramic <djoramic@amd.com> Co-authored-by: Andreas Karatzas <akaratza@amd.com> * [Feature] Add first-class KV hints request envelope for programmatic KV management (vllm-project#53423) Signed-off-by: Karen Chung <karenc@nvidia.com> * [Docs] Add an Engram feature page explaining Engram usage in vLLM (vllm-project#57910) Signed-off-by: NickLucche <nicolo.lucchesi@mistral.ai> * [ROCm][Perf] Route the fused shared-expert gate GEMM through the platform dispatcher (vllm-project#54185) Signed-off-by: Mikko Tukiainen <Mikko.Tukiainen@amd.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> * [Engram] Drop redundant VLLM_PLE_CPU_OFFLOAD env var (vllm-project#57937) Signed-off-by: NickLucche <nicolo.lucchesi@mistral.ai> * [Bugfix][MoE] Reject hash routing for unsupported monolithic backends (vllm-project#57867) Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com> Co-authored-by: OpenAI Codex <noreply@openai.com> * [Bugfix][ROCm] Reject unsupported EP for monolithic AITER MXFP4 MoE (vllm-project#57866) Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com> Co-authored-by: OpenAI Codex <noreply@openai.com> * [Docker] Use zstd for CI images and offer a Docker Hub variant (vllm-project#55608) Signed-off-by: Nils Matteson <nilsmatteson@icloud.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> * [Kimi-K3][AMD] Return KDA and MLA projection outputs directly (vllm-project#50592) Signed-off-by: Liuyinfeng01 <yinfeliu@amd.com> Co-authored-by: Liuyinfeng01 <199041580+LiuYinfeng01@users.noreply.github.com> Co-authored-by: Andreas Karatzas <akaratza@amd.com> * [CI][Build] Harden triton-cpu sleef submodule fetch in CPU image build (vllm-project#57871) Signed-off-by: kimi-no-na-wa <kimi-no-na-wa@users.noreply.github.com> Co-authored-by: kimi-no-na-wa <kimi-no-na-wa@users.noreply.github.com> * [Bugfix][Kernel] Skip the fused silu-mul block-quant fast path when a swiglu clamp is set (vllm-project#57984) Signed-off-by: Garrett Goon <garrett@primeintellect.ai> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com> * [Frontend] Add reusable TP1 initialized-engine snapshots (vllm-project#51360) Signed-off-by: Nils Matteson <nilsmatteson@icloud.com> Signed-off-by: Taneem Ibrahim <taneem.ibrahim@gmail.com> Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com> Co-authored-by: Taneem Ibrahim <taneem.ibrahim@gmail.com> Co-authored-by: elehayym <52448798+Yuzu23@users.noreply.github.com> Co-authored-by: Andreas Karatzas <akaratza@amd.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> * [Pooling] MRV2 pooling shutdown model ref (vllm-project#57737) Signed-off-by: Taneem Ibrahim <taneem.ibrahim@gmail.com> * [CI] Split (H200) LM Eval Large Models into per-model jobs (vllm-project#57965) Signed-off-by: Thang Nguyen <thangnguyenvn647@gmail.com> Co-authored-by: Kimi Code <noreply@moonshot.cn> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com> * [Scheduler] Soften Long Prefill Tokens Threshhold (vllm-project#57951) Signed-off-by: Robert Shaw <robertgshaw2@gmail.com> Signed-off-by: Robert Shaw <robshaw@redhat.com> Co-authored-by: Robert Shaw <robshaw@redhat.com> * [Perf][Attention] Reduce GLM sparse MLA preparation overhead (vllm-project#57458) Signed-off-by: Summer Yang <girasoleyang@gmail.com> Co-authored-by: Codex <noreply@openai.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> Co-authored-by: Jiangyun Zhu <riverclouds.zhu@qq.com> * [Bugfix] Annotate MTP draft KV cache groups positionally on the hybrid grouping path (vllm-project#55390) Signed-off-by: Navjot Singh <navjot.singh@shopify.com> Co-authored-by: Codex <noreply@openai.com> * [Bugfix][GDN] Fix stateless first-chunk classification (vllm-project#51565) Signed-off-by: taking-lying-flat <1615405@qq.com> Signed-off-by: zjy0516 <riverclouds.zhu@qq.com> Co-authored-by: zjy0516 <riverclouds.zhu@qq.com> * [CI] Add pre-commit check that new tests are tethered to Buildkite jobs (vllm-project#54867) Signed-off-by: Turner Jabbour <doubleujabbour@gmail.com> Signed-off-by: Turner <doubleujabbour@gmail.com> Co-authored-by: Claude Sonnet 5 <noreply@anthropic.com> Co-authored-by: Kevin H. Luu <khluu000@gmail.com> Co-authored-by: Nick Hill <nickhill123@gmail.com> * [XPU] Fix Nemotron FP8 LM-eval config: drop CUDA-only moe_backend and wire to new Buildkite job (vllm-project#49685) Signed-off-by: Chaojun Zhang <chaojun.zhang@intel.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> * [Docker] Expose bundled vllm-rs on PATH (vllm-project#57606) Signed-off-by: Alec Flowers <aflowers@nvidia.com> Signed-off-by: Alec <35311602+alec-flowers@users.noreply.github.com> Co-authored-by: Codex <noreply@openai.com> Co-authored-by: Bugen Zhao <i@bugenzhao.com> * [Fast Start] Cache the MTP draft model in a separate daemon group (vllm-project#57312) Signed-off-by: liusy58 <mg21330037@smail.nju.edu.cn> Signed-off-by: Isotr0py <Isotr0py@outlook.com> Co-authored-by: Isotr0py <Isotr0py@outlook.com> * [Rust Frontend] Add MiMo V2.5 parser support (vllm-project#57933) Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Bugen Zhao <i@bugenzhao.com> * [Bugfix][Spec Decode] Cap DFlash/DSpark profiling query batch (vllm-project#56448) Signed-off-by: wangyicong <wangyicong@bytedance.com> * [RL][Sleep] Retain frozen weights across level-2 sleep (vllm-project#57891) Signed-off-by: aoshen02 <aoshen02@users.noreply.github.com> Co-authored-by: aoshen02 <aoshen02@users.noreply.github.com> Co-authored-by: Codex <noreply@openai.com> * [ROCm][Bugfix] Explicitly reject FSE=1 with DPA+ETP deployment for DeepSeek-V4 (vllm-project#57919) Signed-off-by: Shanshan Shen <87969357+shen-shanshan@users.noreply.github.com> * [Spec Decode] Enable async scheduling for DFlash (vllm-project#58065) * [Frontend][Rust] Add mm-processor benchmark for Rust frontend (vllm-project#51922) Signed-off-by: Karthik Gangula <gangula-karthik@users.noreply.github.com> Signed-off-by: gangula-karthik <gkarthik923@gmail.com> Signed-off-by: Isotr0py <Isotr0py@outlook.com> Co-authored-by: Karthik Gangula <gangula-karthik@users.noreply.github.com> Co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com> Co-authored-by: Isotr0py <Isotr0py@outlook.com> * [Rust Frontend] Introduce parser-owned output grammar interfaces (vllm-project#55269) Signed-off-by: Bugen Zhao <i@bugenzhao.com> Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com> * [Security] Reject min_tokens that exceeds the filled max_tokens default (vllm-project#57731) Signed-off-by: Juan Pérez de Algaba <jperezde@redhat.com> * [Bugfix][Frontend] Validate mixed prompt embedding mask lengths (vllm-project#57006) Signed-off-by: 子华 <huaxi.shx@alibaba-inc.com> Co-authored-by: Codex <noreply@openai.com> * [Bugfix][Qwen2.5-VL] Honor video fps for temporal M-RoPE (vllm-project#47736) Signed-off-by: Ting Sun <suntcrick@gmail.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> * [Rust Frontend] Build full-output grammars from initialized reasoning parsers (vllm-project#57340) Signed-off-by: Bugen Zhao <i@bugenzhao.com> Co-authored-by: Codex <noreply@openai.com> * [ROCm][Perf] Avoid extra reshape kernel in Qwen GDN output norm (vllm-project#47842) Signed-off-by: Mikko Tukiainen <Mikko.Tukiainen@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> * [Kernel] Add opt-in load-time MXFP4 dequantization (vllm-project#50814) Signed-off-by: Liuyinfeng01 <yinfeliu@amd.com> Co-authored-by: Shanshan Shen <467638484@qq.com> * [Rust Frontend] Separate multimodal instrumentation from request timing (vllm-project#58084) Signed-off-by: Bugen Zhao <i@bugenzhao.com> * [Kernel][DSV4.1] Fuse MXFP8 wo_b GEMM with sequence-parallel reduce-scatter (vllm-project#57428) Signed-off-by: Thien Tran <gau.nernst@yahoo.com.sg> Signed-off-by: Canlin <canlinguosdu@gmail.com> Co-authored-by: Thien Tran <gau.nernst@yahoo.com.sg> Co-authored-by: OpenAI Codex <codex@openai.com> Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com> Co-authored-by: Jee Jee Li <pandaleefree@gmail.com> * [CI] Emit a kernel symbol map from the csrc build (opt-in, for test selection) (vllm-project#58097) Signed-off-by: khluu <khluu000@gmail.com> Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com> * [perf] wire FA and FlashMLA for sm90 GLM5Next NoPE SparseMLA (vllm-project#55385) Signed-off-by: JaredforReal <w13431838023@gmail.com> Co-authored-by: Kimi Code <noreply@moonshot.cn> Co-authored-by: Leoyzen <leoyzen@gmail.com> Co-authored-by: Lucas Wilkinson <LucasWilkinson@users.noreply.github.com> * [CPU] Add device-memory-utilization CLI alias (vllm-project#56547) Signed-off-by: louie-tsai <louie.tsai@intel.com> Signed-off-by: Louie Tsai <louie.tsai@intel.com> Co-authored-by: Harry Mellor <19981378+hmellor@users.noreply.github.com> * [Rust Frontend] Construct model-owned vision processors through specs (vllm-project#58109) Co-authored-by: Codex <noreply@openai.com> Signed-off-by: Bugen Zhao <i@bugenzhao.com> * [BugFix][Core] Make the structured-output grammar poll non-blocking (vllm-project#55931) Signed-off-by: ubwzwd <ubwzwd@gmail.com> Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com> Co-authored-by: Artem Perevedentsev <aperevedents@nvidia.com> * [XPU][CI]Remove model_runner_v2 test from Intel GPU CI (vllm-project#58050) Signed-off-by: zengxian <xiangdong.zeng@intel.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com> * [Bugfix][Structured Outputs] Reject empty `structural_tag` at request validation (vllm-project#47450) Signed-off-by: linnea-lin-00638949 <15521435947@163.com> Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: Artem Perevedentsev <aperevedents@nvidia.com> * [Build] Fix CUDA 12 KV connector dependency selection (vllm-project#57945) Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai> * [Bugfix][GLM-5.3-Flash] Run the dense MLP layers on the sequence-parallel shard (vllm-project#58061) Signed-off-by: Jared Wen <w13431838023@gmail.com> Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> * [Bugfix][Structured Output] Disallow MRV1 + PP>1 + async sched + structured output (vllm-project#56250) Signed-off-by: Artem Perevedentsev <aperevedents@nvidia.com> Co-authored-by: CNE Pierre FICHEPOIL <pierre-1.fichepoil@gendarmerie.interieur.gouv.fr> * [Feat][XPU] VLLM_BATCH_INVARIANT support for Dense/MoE models (vllm-project#55881) Signed-off-by: Tony Lin <tony.lin@intel.com> Co-authored-by: Flora Feng <4florafeng@gmail.com> * [SpecDecode] Restore residual-logits comments in _resample_kernel (vllm-project#58166) Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com> Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com> * [Bugfix] Narrow AuxOutput KV restrictions to known PD connectors (vllm-project#58150) Signed-off-by: aoshen02 <aoshen@inferact.ai> * [Bugfix] prioritize architecture capability before DeepGEMM availability check (vllm-project#58073) Signed-off-by: Tony Lin <tony.lin@intel.com> * [Bugfix][Attention] Avoid NaN in the Triton softcap for large attention logits (vllm-project#56579) Signed-off-by: Kushal Dabbe <72650064+kushaldabbe@users.noreply.github.com> Co-authored-by: opencode <noreply@opencode.ai> Co-authored-by: Misha Goin <mgoin64@gmail.com> * [Bugfix][Engram] Fall back when /dev/shm is absent before sharing tables (vllm-project#57914) Signed-off-by: Juntian Liu <juntianl@inferact.ai> Signed-off-by: Juntian Liu <Juntianl777@gmail.com> Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> * [Bugfix] hadacore_transform: respect inplace parameter to fix garbage outputs with QuIP transforms (vllm-project#43462) Signed-off-by: Gilles Turpin <turpingilles15@gmail.com> Co-authored-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com> * [Bugfix][ROCm] Dispatch the QuantFP8 CUDA fallback on the class (vllm-project#58136) Signed-off-by: Stefan Koncarevic <Stefan.Koncarevic@amd.com> Co-authored-by: Andreas Karatzas <akaratza@amd.com> * [ROCm][DSv4.1][Perf] Fuse the inverse RoPE into the sparse decode reduce (vllm-project#57435) Signed-off-by: Fangzhou Ai <fangzhou.ai@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> * [Fast Start] Support data parallelism in the weight cache daemon (vllm-project#57386) Signed-off-by: liusy58 <mg21330037@smail.nju.edu.cn> Signed-off-by: Isotr0py <Isotr0py@outlook.com> Co-authored-by: Isotr0py <Isotr0py@outlook.com> * [Bugfix] batch_invariant: keep non-AllReduce collectives enabled on NCCL >= 2.31 (vllm-project#58179) Signed-off-by: Guanxin Li <38149783+guanxingithub@users.noreply.github.com> Co-authored-by: Flora Feng <4florafeng@gmail.com> * [MRV2] Release weight offloader on shutdown (vllm-project#57834) Signed-off-by: Taneem Ibrahim <taneem.ibrahim@gmail.com> * [Bugfix][V1] Honor enable_jit_warmup for V2 kernel warmup (vllm-project#55146) Co-authored-by: mgoin <mgoin64@gmail.com> * [ROCm][CI] Add GELU activation for AiterExperts in the modular-kernel coverage (vllm-project#58030) Signed-off-by: Divakar Verma <divakar.verma@amd.com> Co-authored-by: Andreas Karatzas <akaratza@amd.com> * [Bugfix][V1] Read ModelState max_model_len from model config (vllm-project#58149) Signed-off-by: Chenglun Hu <chenglunhu@gmail.com> Signed-off-by: mgoin <mgoin64@gmail.com> Co-authored-by: mgoin <mgoin64@gmail.com> * [Bugfix][Model][Spec Decode] Defer disposable GLM MTP head (vllm-project#55442) Signed-off-by: Luca Motz <luca.motz@icloud.com> * [Refactor] Remove dead code multiple places (vllm-project#58002) Signed-off-by: yewentao256 <zhyanwentao@126.com> * [Core] structured generation mode for DiffusionGemma model (Jev-like) (vllm-project#57250) Signed-off-by: Matt Mastracci <matthew@mastracci.com> Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com> Signed-off-by: Razorback16 <razorback16@protonmail.com> Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com> Co-authored-by: Lucas Wilkinson <lwilkins@redhat.com> Co-authored-by: Codex <noreply@openai.com> Co-authored-by: Razorback16 <razorback16@protonmail.com> * [Kernel] Remove AllSpark INT8 W8A16 GEMM backend (vllm-project#58001) Signed-off-by: mgoin <mgoin64@gmail.com> * [CI] Build the torch-nightly image on Ubuntu 24.04 (vllm-project#58204) * [Bugfix] Backport Inductor custom-op pattern matching fix (vllm-project#58189) Signed-off-by: mgoin <mgoin64@gmail.com> * [Core] Disable JIT warmup in eager mode (vllm-project#58197) Signed-off-by: mgoin <mgoin64@gmail.com> * [CI] Split LM Eval TurboQuant KV Cache into per-config jobs (vllm-project#57113) Signed-off-by: Thang Nguyen <thangnguyenvn647@gmail.com> Co-authored-by: Kimi <noreply@moonshot.ai> * [CI] Shard (H200 MIG 18GB) Spec Decode Draft Model across whole-directory replicas (vllm-project#58193) Signed-off-by: Thang Nguyen <thangnguyenvn647@gmail.com> Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com> * [Perf] Remove CPU-GPU sync in heterogeneous vocabulary speculative decoding (vllm-project#57396) * [ROCm][Build][The Rock] Bump Triton version to 3.8.x tip-of-tree with source build in The Rock image (vllm-project#58006) Signed-off-by: Randall Smith <Randall.Smith@amd.com> * [Bugfix][ROCm] Use the platform FP8 range in the concat MLA q test (vllm-project#58153) Co-authored-by: Andreas Karatzas <akaratza@amd.com> * [ROCm][Model][Bugfix] Enable GLM-5.2-MXFP4 on the deepseek_v32 path and fix sparse attention correctness (vllm-project#51915) Signed-off-by: Jack Hu <Jack.Hu@amd.com> Signed-off-by: Jack Hu <jack.hu@amd.com> Signed-off-by: Douglas Lehr <Doug.Lehr@amd.com> Co-authored-by: James E T Smith <jamesETsmith@users.noreply.github.com> Co-authored-by: Douglas Lehr <91553416+dllehr-amd@users.noreply.github.com> Co-authored-by: Douglas Lehr <Doug.Lehr@amd.com> * [ROCm] Use silu_and_mul_with_clamp's torch._C op (vllm-project#52052) Signed-off-by: Tres Popp <tres.popp@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Shanshan Shen <467638484@qq.com> Co-authored-by: Claude Sonnet 5 <noreply@anthropic.com> Co-authored-by: Andreas Karatzas <akaratza@amd.com> * [Bugfix] Skip VllmConfig re-validation for with_hf_config submodel views (vllm-project#58212) Signed-off-by: Nick Hill <nickhill123@gmail.com> Co-authored-by: Kimi Code <noreply@moonshot.ai> Co-authored-by: Roger Wang <rogerw@inferact.ai> * [EPD] Support metadata-only audio inputs (vllm-project#57887) Signed-off-by: Tianyu Guo <guoty@inferact.ai> Co-authored-by: Cyrus Leung <tlleungac@connect.ust.hk> * [ROCm][DSv4][Perf] Fuse the inverse RoPE into the sparse decode reduce (vllm-project#57451) Signed-off-by: Fangzhou Ai <fangzhou.ai@amd.com> * [ROCm][Compile] Support BF16 AsyncTP fusion (vllm-project#58098) Signed-off-by: Andreas Karatzas <akaratza@amd.com> * [CI][Bugfix] Update IPC test caller for vllm-project#57312's _apply_entries signature (vllm-project#58107) Signed-off-by: kimi-no-na-wa <kimi-no-na-wa@users.noreply.github.com> Co-authored-by: kimi-no-na-wa <kimi-no-na-wa@users.noreply.github.com> Co-authored-by: Kevin H. Luu <khluu000@gmail.com> * [ROCm][CI] Stage G gating (vllm-project#50922) Signed-off-by: Andreas Karatzas <akaratza@amd.com> Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com> Co-authored-by: OpenAI Codex <codex@openai.com> * [Bugfix] Set worker runtime threads before profiling and compilation (vllm-project#55891) Signed-off-by: Andreas Karatzas <andreas.karatzas@protonmail.com> Co-authored-by: Codex <noreply@openai.com> Co-authored-by: Nick Hill <nickhill123@gmail.com> * [XPU] Wire up SYCL apply_rotary_emb kernel in ApplyRotaryEmb (vllm-project#55721) Signed-off-by: Michal Ganczarenko <michal.ganczarenko@intel.com> Signed-off-by: Michał Ganczarenko <michal.ganczarenko@intel.com> Co-authored-by: Claude <noreply@anthropic.com> * [Compilation] Fix QuTLASS compilation with PyTorch 2.13 (vllm-project#58173) Signed-off-by: yewentao256 <zhyanwentao@126.com> * [Spec decode] Support variable-length decode for Kimi-K3 adaptive ver (vllm-project#52988) Signed-off-by: Albert Cheng <albecheng@nvidia.com> Signed-off-by: Albert Cheng (Engrg-Hardware 1) <albecheng@login-bia01.bia.clusters.nvidia.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> Co-authored-by: Benjamin Chislett <bchislett@nvidia.com> * [XPU][CI] Deselect tests/v1/spec_decode/test_mtp.py::test_glm_mtp_defers_lm_head (vllm-project#58237) Signed-off-by: zengxian <xiangdong.zeng@intel.com> * [MoE] Use GateLinear for all MoE models (vllm-project#58234) Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai> * [Bugfix][Frontend] Keep length finish_reason for max_tokens-truncated streaming tool calls (vllm-project#46303) Signed-off-by: Ting Sun <suntcrick@gmail.com> * [ROCm][Perf] Use wvSplitK for single-output GEMMs (vllm-project#53283) Signed-off-by: tangzzycc <3081129260@qq.com> * [Tests] Select V2 for diffusion scheduler unit tests (vllm-project#58272) Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com> Co-authored-by: OpenAI Codex <codex@openai.com> * [Quantization] Select per-token NVFP4 MoE backends explicitly (vllm-project#57176) Signed-off-by: aoshen02 <aoshen@inferact.ai> Signed-off-by: S1ro1 <matej.sirovatka@gmail.com> Signed-off-by: aoshen02 <aoshen02@users.noreply.github.com> Co-authored-by: aoshen02 <aoshen@inferact.ai> Co-authored-by: OpenAI Codex <noreply@openai.com> Co-authored-by: aoshen02 <aoshen02@users.noreply.github.com> Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com> * [ROCm][CI] Mirror the three TurboQuant evaluation groups on MI355 (vllm-project#58282) Signed-off-by: Andreas Karatzas <akaratza@amd.com> Co-authored-by: OpenAI Codex <noreply@openai.com> * [ROCm][CI] Add MI355 dense NVFP4 and MoRI kernel mirrors (vllm-project#58281) Signed-off-by: Andreas Karatzas <akaratza@amd.com> Co-authored-by: OpenAI Codex <noreply@openai.com> * [Mooncake] Address review nits from vllm-project#56855 (vllm-project#57174) Signed-off-by: NickLucche <nicolo.lucchesi@mistral.ai> Co-authored-by: Yifan Qiao <17067717+ivanium@users.noreply.github.com> * [Refactor][Quantization] Make FP8 and MLA weight transforms reusable pure functions (vllm-project#57732) Signed-off-by: aoshen02 <aoshen@inferact.ai> Co-authored-by: Yongye Zhu <zyy1102000@gmail.com> * [ROCm][Compile] Fuse AITER static FP8 attention output (vllm-project#58099) Signed-off-by: Andreas Karatzas <akaratza@amd.com> Co-authored-by: OpenAI Codex <noreply@openai.com> * [ROCm][Bugfix] Register MRV2 sampler JIT warmups (vllm-project#58092) Signed-off-by: Andreas Karatzas <akaratza@amd.com> Co-authored-by: OpenAI Codex <codex@openai.com> * [Perf][Attention] Avoid CPU-GPU sync in DCP sequence lengths (vllm-project#58169) Signed-off-by: zjy0516 <riverclouds.zhu@qq.com> Co-authored-by: OpenAI Codex <noreply@openai.com> * [CI][Bugfix] Extend groupwise rms_norm scale tolerance to CUDA (vllm-project#58252) Signed-off-by: kimi-no-na-wa <kimi-no-na-wa@users.noreply.github.com> Co-authored-by: kimi-no-na-wa <kimi-no-na-wa@users.noreply.github.com> * Disable pre-commit check-test-tethering Prevent this check from running in our branch because we are not testing our branch with Buildkite. Signed-off-by: Dan Eble <Dan.Eble@amd.com> --------- Signed-off-by: Isotr0py <Isotr0py@outlook.com> Signed-off-by: jinzhen.ljz <jinzhen.ljz@antgroup.com> Signed-off-by: mgoin <mgoin64@gmail.com> Signed-off-by: Roi Koren <roik@nvidia.com> Signed-off-by: yewentao256 <zhyanwentao@126.com> Signed-off-by: Sheral Kumar <shekumar@amd.com> Signed-off-by: Andreas Karatzas <Andreas.Karatzas@amd.com> Signed-off-by: Yichao Zhu <Yichao.Zhu@amd.com> Signed-off-by: yisheng <yi.sheng@intel.com> Signed-off-by: Woosuk Kwon <woosuk@inferact.ai> Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai> Signed-off-by: jiangkuaixue123 <jiangxiaozhou111@163.com> Signed-off-by: CrazyCoder <18235787078@163.com> Signed-off-by: aoshen02 <aoshen@inferact.ai> Signed-off-by: Zhou Zihang <z@mcac.cc> Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai> Signed-off-by: khluu <khluu000@gmail.com> Signed-off-by: Matt Mastracci <matthew@mastracci.com> Signed-off-by: chaunceyjiang <chaunceyjiang@gmail.com> Signed-off-by: zjy0516 <riverclouds.zhu@qq.com> Signed-off-by: zhenwei-intel <zhenwei.liu@intel.com> Signed-off-by: Tianyao Wu <rayroy31@gmail.com> Signed-off-by: Oliver Holworthy <1216955+oliverholworthy@users.noreply.github.com> Signed-off-by: john <6712432+CaiJohn@users.noreply.github.com> Signed-off-by: Aakif Nawaz <aakif.nawaz@amd.com> Signed-off-by: Juntian Liu <Juntianl777@gmail.com> Signed-off-by: Wei-Ting Liao <wei-ting.liao@amd.com> Signed-off-by: Andreas Karatzas <akaratza@amd.com> Signed-off-by: Djordje Ramic <djoramic@amd.com> Signed-off-by: Nick Hill <nickhill123@gmail.com> Signed-off-by: Fangzhou Ai <fangzhou.ai@amd.com> Signed-off-by: Juan Pérez de Algaba <jperezde@redhat.com> Signed-off-by: Chaojun Zhang <chaojun.zhang@intel.com> Signed-off-by: hotragn <hotragn.pettugani_2024@woxsen.edu.in> Signed-off-by: Martin Hickey <martin.hickey@ie.ibm.com> Signed-off-by: DarkLight1337 <tlleungac@connect.ust.hk> Signed-off-by: Taneem Ibrahim <taneem.ibrahim@gmail.com> Signed-off-by: Sergio Figueras <sergio@yourecm.com> Signed-off-by: Tuukka Sarvi <tuukka.sarvi@amd.com> Signed-off-by: larryli2-amd <larryli2@amd.com> Signed-off-by: Tony Lin <tony.lin@intel.com> Signed-off-by: Wauplin <lucainp@gmail.com> Signed-off-by: Lucas Wilkinson <lwilkinson@neuralmagic.com> Signed-off-by: Andy Friedrich <afriedri@amd.com> Signed-off-by: afriedri <afriedri@amd.com> Signed-off-by: sahil <sahil@example.com> Signed-off-by: Karen Chung <karenc@nvidia.com> Signed-off-by: NickLucche <nicolo.lucchesi@mistral.ai> Signed-off-by: Mikko Tukiainen <Mikko.Tukiainen@amd.com> Signed-off-by: Nils Matteson <nilsmatteson@icloud.com> Signed-off-by: Liuyinfeng01 <yinfeliu@amd.com> Signed-off-by: kimi-no-na-wa <kimi-no-na-wa@users.noreply.github.com> Signed-off-by: Garrett Goon <garrett@primeintellect.ai> Signed-off-by: Thang Nguyen <thangnguyenvn647@gmail.com> Signed-off-by: Robert Shaw <robertgshaw2@gmail.com> Signed-off-by: Robert Shaw <robshaw@redhat.com> Signed-off-by: Summer Yang <girasoleyang@gmail.com> Signed-off-by: Navjot Singh <navjot.singh@shopify.com> Signed-off-by: taking-lying-flat <1615405@qq.com> Signed-off-by: Turner Jabbour <doubleujabbour@gmail.com> Signed-off-by: Turner <doubleujabbour@gmail.com> Signed-off-by: Alec Flowers <aflowers@nvidia.com> Signed-off-by: Alec <35311602+alec-flowers@users.noreply.github.com> Signed-off-by: liusy58 <mg21330037@smail.nju.edu.cn> Signed-off-by: Bugen Zhao <i@bugenzhao.com> Signed-off-by: wangyicong <wangyicong@bytedance.com> Signed-off-by: aoshen02 <aoshen02@users.noreply.github.com> Signed-off-by: Shanshan Shen <87969357+shen-shanshan@users.noreply.github.com> Signed-off-by: Karthik Gangula <gangula-karthik@users.noreply.github.com> Signed-off-by: gangula-karthik <gkarthik923@gmail.com> Signed-off-by: 子华 <huaxi.shx@alibaba-inc.com> Signed-off-by: Ting Sun <suntcrick@gmail.com> Signed-off-by: Thien Tran <gau.nernst@yahoo.com.sg> Signed-off-by: Canlin <canlinguosdu@gmail.com> Signed-off-by: JaredforReal <w13431838023@gmail.com> Signed-off-by: louie-tsai <louie.tsai@intel.com> Signed-off-by: Louie Tsai <louie.tsai@intel.com> Signed-off-by: ubwzwd <ubwzwd@gmail.com> Signed-off-by: zengxian <xiangdong.zeng@intel.com> Signed-off-by: linnea-lin-00638949 <15521435947@163.com> Signed-off-by: Jared Wen <w13431838023@gmail.com> Signed-off-by: Artem Perevedentsev <aperevedents@nvidia.com> Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com> Signed-off-by: Kushal Dabbe <72650064+kushaldabbe@users.noreply.github.com> Signed-off-by: Juntian Liu <juntianl@inferact.ai> Signed-off-by: Gilles Turpin <turpingilles15@gmail.com> Signed-off-by: Stefan Koncarevic <Stefan.Koncarevic@amd.com> Signed-off-by: Guanxin Li <38149783+guanxingithub@users.noreply.github.com> Signed-off-by: Divakar Verma <divakar.verma@amd.com> Signed-off-by: Chenglun Hu <chenglunhu@gmail.com> Signed-off-by: Luca Motz <luca.motz@icloud.com> Signed-off-by: Razorback16 <razorback16@protonmail.com> Signed-off-by: Randall Smith <Randall.Smith@amd.com> Signed-off-by: Jack Hu <Jack.Hu@amd.com> Signed-off-by: Jack Hu <jack.hu@amd.com> Signed-off-by: Douglas Lehr <Doug.Lehr@amd.com> Signed-off-by: Tres Popp <tres.popp@amd.com> Signed-off-by: Tianyu Guo <guoty@inferact.ai> Signed-off-by: Andreas Karatzas <andreas.karatzas@protonmail.com> Signed-off-by: Michal Ganczarenko <michal.ganczarenko@intel.com> Signed-off-by: Michał Ganczarenko <michal.ganczarenko@intel.com> Signed-off-by: Albert Cheng <albecheng@nvidia.com> Signed-off-by: Albert Cheng (Engrg-Hardware 1) <albecheng@login-bia01.bia.clusters.nvidia.com> Signed-off-by: tangzzycc <3081129260@qq.com> Signed-off-by: S1ro1 <matej.sirovatka@gmail.com> Signed-off-by: Dan Eble <Dan.Eble@amd.com> Co-authored-by: Isotr0py <Isotr0py@outlook.com> Co-authored-by: Jinzhen Lin <jinzhen.ljz@antgroup.com> Co-authored-by: Codex <noreply@openai.com> Co-authored-by: mgoin <mgoin64@gmail.com> Co-authored-by: roikoren755 <26850796+roikoren755@users.noreply.github.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> Co-authored-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com> Co-authored-by: Sheral Kumar <shekumar@amd.com> Co-authored-by: Andreas Karatzas <Andreas.Karatzas@amd.com> Co-authored-by: Andreas Karatzas <akaratza@amd.com> Co-authored-by: YukioZzz <Yichao.Zhu@amd.com> Co-authored-by: YiSheng5 <yi.sheng@intel.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com> Co-authored-by: Woosuk Kwon <woosuk@inferact.ai> Co-authored-by: Yifan Qiao <yifanqiao@inferact.ai> Co-authored-by: Claude Fable 5.1 <noreply@anthropic.com> Co-authored-by: jiangkuaixue123 <jiangxiaozhou111@163.com> Co-authored-by: CrazyCoder <18235787078@163.com> Co-authored-by: aoshen02 <aoshen@inferact.ai> Co-authored-by: Zhou Zihang <mail@zzhgo.com> Co-authored-by: Jee Jee Li <pandaleefree@gmail.com> Co-authored-by: JohnQinAMD <yanyuan.qin@amd.com> Co-authored-by: czhu-cohere <conway.zhu@cohere.com> Co-authored-by: Kevin H. Luu <khluu000@gmail.com> Co-authored-by: Matt Mastracci <matthew@mastracci.com> Co-authored-by: Lucas Wilkinson <LucasWilkinson@users.noreply.github.com> Co-authored-by: Chauncey <chaunceyjiang@gmail.com> Co-authored-by: zjy0516 <riverclouds.zhu@qq.com> Co-authored-by: liuzhenwei <zhenweiliu@habana.ai> Co-authored-by: Tianyao Wu <54675599+twu3202@users.noreply.github.com> Co-authored-by: Oliver Holworthy <1216955+oliverholworthy@users.noreply.github.com> Co-authored-by: CaiJohn <CaiJohn@users.noreply.github.com> Co-authored-by: john <6712432+CaiJohn@users.noreply.github.com> Co-authored-by: akii96 <aakif.nawaz@amd.com> Co-authored-by: Juntian Liu <Juntianl777@gmail.com> Co-authored-by: weitliao <wei-ting.liao@amd.com> Co-authored-by: Cursor <cursoragent@cursor.com> Co-authored-by: Shanshan Shen <467638484@qq.com> Co-authored-by: djramic <djoramic@amd.com> Co-authored-by: Hongxin Xu <70438206+xhx1022@users.noreply.github.com> Co-authored-by: OpenAI Codex <codex@openai.com> Co-authored-by: Roger Wang <hey@rogerw.io> Co-authored-by: Nick Hill <nickhill123@gmail.com> Co-authored-by: Fangzhou Ai <31551580+Fangzhou-Ai@users.noreply.github.com> Co-authored-by: Juan Pérez de Algaba <124347725+jperezdealgaba@users.noreply.github.com> Co-authored-by: Chaojun Zhang <chaojun.zhang@intel.com> Co-authored-by: Hotragn Pettugani <103170876+Hotragn@users.noreply.github.com> Co-authored-by: Martin Hickey <martin.hickey@ie.ibm.com> Co-authored-by: Cyrus Leung <tlleungac@connect.ust.hk> Co-authored-by: Taneem Ibrahim <taneem.ibrahim@gmail.com> Co-authored-by: sfigs <sergio@yourecm.com> Co-authored-by: Tuukka Sarvi <tuukka.sarvi@amd.com> Co-authored-by: larryli2-amd <larryli2@amd.com> Co-authored-by: Tony Lin <tony.lin@intel.com> Co-authored-by: Lucain <lucainp@gmail.com> Co-authored-by: afriedri <afriedri@amd.com> Co-authored-by: Douglas Lehr <91553416+dllehr-amd@users.noreply.github.com> Co-authored-by: Sahil Patel <91423311+Sip4818@users.noreply.github.com> Co-authored-by: Karen Chung <karenc@nvidia.com> Co-authored-by: Nicolò Lucchesi <nicolo.lucchesi@mistral.ai> Co-authored-by: Mikko Tukiainen <mikko.tukiainen@amd.com> Co-authored-by: Nils Matteson <nilsmatteson@icloud.com> Co-authored-by: yinfengLiu <yinfeliu@amd.com> Co-authored-by: Liuyinfeng01 <199041580+LiuYinfeng01@users.noreply.github.com> Co-authored-by: vllm-agent <claw@inferact.ai> Co-authored-by: kimi-no-na-wa <kimi-no-na-wa@users.noreply.github.com> Co-authored-by: Garrett Goon <44747910+garrett361@users.noreply.github.com> Co-authored-by: elehayym <52448798+Yuzu23@users.noreply.github.com> Co-authored-by: Thang Nguyen <69278249+Thangnguyenvn98@users.noreply.github.com> Co-authored-by: Kimi Code <noreply@moonshot.cn> Co-authored-by: Robert Shaw <114415538+robertgshaw2-redhat@users.noreply.github.com> Co-authored-by: Robert Shaw <robshaw@redhat.com> Co-authored-by: Summer Yang <girasoleyang@gmail.com> Co-authored-by: Navjot Singh <navjot.singh@uwaterloo.ca> Co-authored-by: cherry77-cloud <1615405@qq.com> Co-authored-by: Turner Jabbour <doubleujabbour@gmail.com> Co-authored-by: Copilot <223556219+Copilot@users.noreply.github.com> Co-authored-by: Alec <35311602+alec-flowers@users.noreply.github.com> Co-authored-by: Bugen Zhao <i@bugenzhao.com> Co-authored-by: siyu <mg21330037@smail.nju.edu.cn> Co-authored-by: Yicong Wang <wangyicong@bytedance.com> Co-authored-by: aoshen02 <aoshen02@users.noreply.github.com> Co-authored-by: Jeff (Junze) Ma <93145857+majunze2001@users.noreply.github.com> Co-authored-by: karthik <56480632+gangula-karthik@users.noreply.github.com> Co-authored-by: Karthik Gangula <gangula-karthik@users.noreply.github.com> Co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com> Co-authored-by: shaohuaxi <huaxi.shx@alibaba-inc.com> Co-authored-by: Ting SUN <suntcrick@gmail.com> Co-authored-by: Canlin Guo <canlinguosdu@gmail.com> Co-authored-by: Thien Tran <gau.nernst@yahoo.com.sg> Co-authored-by: Jared Wen <w13431838023@gmail.com> Co-authored-by: Leoyzen <leoyzen@gmail.com> Co-authored-by: Louie Tsai <louie.tsai@intel.com> Co-authored-by: Harry Mellor <19981378+hmellor@users.noreply.github.com> Co-authored-by: ubwzwd <ubwzwd@gmail.com> Co-authored-by: Artem Perevedentsev <aperevedents@nvidia.com> Co-authored-by: xiangdong <40376367+zxd1997066@users.noreply.github.com> Co-authored-by: linyafeng <15521435947@163.com> Co-authored-by: CNE Pierre FICHEPOIL <pierre-1.fichepoil@gendarmerie.interieur.gouv.fr> Co-authored-by: Flora Feng <4florafeng@gmail.com> Co-authored-by: Kushal <72650064+kushaldabbe@users.noreply.github.com> Co-authored-by: opencode <noreply@opencode.ai> Co-authored-by: Gilles Turpin <turpingilles15@gmail.com> Co-authored-by: stefankoncarevic <Stefan.Koncarevic@amd.com> Co-authored-by: Guanxin Li <38149783+guanxingithub@users.noreply.github.com> Co-authored-by: pengyihang <1017861497@qq.com> Co-authored-by: Divakar Verma <137818590+divakar-amd@users.noreply.github.com> Co-authored-by: hcl <chenglunhu@gmail.com> Co-authored-by: lucamotz <luca.motz@icloud.com> Co-authored-by: Lucas Wilkinson <lwilkins@redhat.com> Co-authored-by: Razorback16 <razorback16@protonmail.com> Co-authored-by: Andrey Talman <atalman@fb.com> Co-authored-by: Kimi <noreply@moonshot.ai> Co-authored-by: Michael Lapshin <55516685+MichaelLapshin@users.noreply.github.com> Co-authored-by: rasmith <Randall.Smith@amd.com> Co-authored-by: Jack Hu <jack.hu@amd.com> Co-authored-by: James E T Smith <jamesETsmith@users.noreply.github.com> Co-authored-by: Douglas Lehr <Doug.Lehr@amd.com> Co-authored-by: Tres <tpopp@users.noreply.github.com> Co-authored-by: Roger Wang <rogerw@inferact.ai> Co-authored-by: Tianyu Guo <guoty@inferact.ai> Co-authored-by: Michał Ganczarenko <michal.ganczarenko@intel.com> Co-authored-by: Albert Cheng <albecheng@nvidia.com> Co-authored-by: Benjamin Chislett <bchislett@nvidia.com> Co-authored-by: 00CC <83797617+tangzzycc@users.noreply.github.com> Co-authored-by: Matej Sirovatka <54212263+S1ro1@users.noreply.github.com> Co-authored-by: Yifan Qiao <17067717+ivanium@users.noreply.github.com> Co-authored-by: Yongye Zhu <zyy1102000@gmail.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.
Purpose
Reduce query packing and sparse-index preparation overhead around FlashInfer NoPE MLA attention.
The three commits:
torch.catbefore FP8 quantization. Preserve concatenation for nonempty rotary dimensions and strided queries.sparse_mla_index_remap_kernelfor trace attribution.Duplicate-work checks found no matching open PR. Related PRs #54258 (separate ROCm remap kernel), #38354 (DeepSeek pre-attention GEMM/RoPE fusion), and #57161 (GLM kpool/indexer scheduling) address different kernels or stages. No existing issue is being claimed or closed.
Correctness
Prior TP validation on the experiment stack used GLM-5.3-Flash, TP8+EP8, all custom ops enabled, eager execution, no prefix cache, and exactly 32768 input tokens with either 1 or 16 generated tokens. Across 12 repeatability requests, all 8976 sparse-layer calls across eight ranks had positive valid counts (minimum 1). A direct full-shape FP8 attention comparison at 32768 queries × 8 heads × 512 dimensions, capacity 2176, matched the previous implementation bit for bit.
Performance
Controlled end-to-end comparison on four GB300 GPUs, TP4+EP4, eager,
custom_ops=["all"], no prefix cache, no speculative decoding, and FP8 E4M3 KV cache. The actual upstream PR model stayed loaded while an idle-worker extension switched only the three preparation changes to their PR-base implementations. All four ranks verified the implementation selection. Compatible prebuilt extensions were reused.Both implementations were warmed, then measured in four alternating baseline/PR and PR/baseline rounds with five matched samples per workload per round (20 pairs each). Each request had exactly 32768 input tokens, with a 32768-token prefill scheduling limit. Prefill generated one token; decode generated 128 tokens. These are client-observed serving latencies.
All 20 pairs improved for each workload. Per-round median reductions ranged from 1.78–2.27% for prefill and 1.21–1.98% for decode. These measurements apply to eager TP4 execution; CUDA-graph decode and other parallel configurations were not measured in this comparison.
Earlier isolated measurements on the same GB300 GPU, using CUDA graph replay and the median of five samples with 40 repetitions/sample:
Shape: 32768 queries × 8 heads × 512 dimensions, capacity 2176. These are preparation microbenchmarks from the experiment stack. Resident TP now skips safe-length preparation entirely. Earlier tracing also confirmed removal of all 11 redundant query concatenations. The final upstream revision has not been reprofiled; the controlled end-to-end results above are separate from these older microbenchmarks. Measurements of the removed output-zeroing kernel are excluded.