Compare commits

..
Author SHA1 Message Date
dependabot[bot]andGitHub f36d57ec34 Bump actions/setup-python from 6.1.0 to 6.2.0
Bumps [actions/setup-python](https://github.com/actions/setup-python) from 6.1.0 to 6.2.0.
- [Release notes](https://github.com/actions/setup-python/releases)
- [Commits](https://github.com/actions/setup-python/compare/83679a892e2d95755f2dac6acb0bfd1e9ac5d548...a309ff8b426b58ec0e2a45f0f869d46889d02405)

---
updated-dependencies:
- dependency-name: actions/setup-python
  dependency-version: 6.2.0
  dependency-type: direct:production
  update-type: version-update:semver-minor
...

Signed-off-by: dependabot[bot] <support@github.com>
2026-06-02 18:31:03 +00:00
dependabot[bot]GitHubdependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
b8b49e2395 Bump actions/github-script from 8.0.0 to 9.0.0 (#39667)
Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>
2026-06-02 11:26:57 -07:00
da107a59e5 [MRV2] Also enable MRV2 for Llama and Mistral dense models (#43458)
Signed-off-by: Nick Hill <nickhill123@gmail.com>
Signed-off-by: yewentao256 <zhyanwentao@126.com>
Co-authored-by: yewentao256 <zhyanwentao@126.com>
2026-06-02 11:18:46 -07:00
ChaunceyGitHubAleksandar YanakievAng Kah Min, Kelvin <syraxius@hotmail.com>
ed9a7526b6 [Anthropic] Support system role messages inside messages array (#44283)
Signed-off-by: chaunceyjiang <chaunceyjiang@gmail.com>
Co-authored-by: Aleksandar Yanakiev <alexander.yanakiev@discretestack.com>
Co-authored-by: Ang Kah Min, Kelvin <syraxius@hotmail.com>
2026-06-02 18:13:54 +00:00
2427094152 [Feature] Support EPLB for DeepSeek v4 Mega Moe (#43339)
Signed-off-by: wzhao18 <wzhao18.sz@gmail.com>
Co-authored-by: Wei Zhao (Engrg-Hardware 1) <weizha@login-lyris01.lyris.clusters.nvidia.com>
2026-06-02 10:56:44 -07:00
Kartavya sonarandGitHub fe32e7830b [Bugfix] flashinfer: fail fast when --kv-cache-dtype nvfp4 used on unsupported arch (#43669)
Signed-off-by: Kartavya Sonar <sonarkartavya@gmail.com>
2026-06-02 10:50:00 -07:00
afcb580715 [BugFix] Fix Humming MoE deploy error (#43100)
Signed-off-by: Alireza Dadgarnia <dadgarnia@Alirezas-MacBook-Pro-2.local>
Signed-off-by: Alireza Dadgarnia <49554709+adotdad@users.noreply.github.com>
Co-authored-by: Alireza Dadgarnia <dadgarnia@Alirezas-MacBook-Pro-2.local>
Co-authored-by: Jinzhen Lin <linjinzhen@hotmail.com>
2026-06-02 09:32:50 -07:00
3f3e2702c2 [XPU] Enable rms_norm/act quant fusions (#43963)
Signed-off-by: zhenwei-intel <zhenwei.liu@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>
2026-06-02 16:14:41 +00:00
Flora FengandGitHub 478b49ddec [Refactor] Remove dead code from parser infrastructure (#44279)
Signed-off-by: sfeng33 <4florafeng@gmail.com>
2026-06-02 12:08:27 -04:00
Nick HillandGitHub cab5c9a2a9 [Core] Move max_concurrent_batches to VllmConfig (#44274)
Signed-off-by: Nick Hill <nickhill123@gmail.com>
2026-06-02 08:57:25 -07:00
Brian DellabettaandGitHub 774e552397 [compressed-tensors] Asymmetric support for MoE WNA16 marlin (#44025)
Signed-off-by: Brian Dellabetta <bdellabe@redhat.com>
2026-06-02 08:51:45 -07:00
XiaoZandGitHub 53fa09d085 [Misc] Support local image encoding in benchmarks (#43843)
Signed-off-by: xiaoz <Sukra1@outlook.com>
2026-06-02 15:15:06 +00:00
Chris LeonardandGitHub 4d93bc35c9 Migrate header files to torch stable abi (#44013) 2026-06-02 08:09:52 -07:00
Bugen ZhaoandGitHub 586201ebdc [Rust Frontend] Cover different thinking modes in roundtrip tests (#44320)
Signed-off-by: Bugen Zhao <i@bugenzhao.com>
2026-06-02 07:51:25 -07:00
pschlan-amdandGitHub 88f172188b [ROCm] Fix AITER RMSNormQuantFusion for Kimi-Linear (#44308)
Signed-off-by: Patrick Schlangen <pschlan@amd.com>
2026-06-02 14:50:21 +00:00
Bugen ZhaoandGitHub 880fc032f4 [Rust Frontend] Support recursive tool parameter conversion (#44299)
Signed-off-by: Bugen Zhao <i@bugenzhao.com>
2026-06-02 07:45:35 -07:00
6314de8bad [XPU] [Bug] remove xpuw4a16 output size check (#44168)
Signed-off-by: Zhu, Zufang <zufang.zhu@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>
2026-06-02 22:26:20 +08:00
IdoAtadTDandGitHub c91a87f01a [BugFix] [GDN] Read linear_key_head_dim from hf_text_config for multimodal models (#43978)
Signed-off-by: IdoAtadTD <ido.atad@twodelta.com>
2026-06-02 17:17:55 +03:00
Matthew BonanniandGitHub ea0d045a05 [FlashAttention] Sync FA with upstream (#44065)
Signed-off-by: Matthew Bonanni <mbonanni@redhat.com>
2026-06-02 07:15:37 -07:00
0bdfd5eb84 [Bugfix] Vendor MiniCPMV/MiniCPMO processors to unblock Transformers v5 (#44282)
Signed-off-by: guanwei-wu <b08901019@ntu.edu.tw>
Signed-off-by: wjinxu <1299461899@qq.com>
Co-authored-by: guanwei-wu <b08901019@ntu.edu.tw>
Co-authored-by: Cursor <cursoragent@cursor.com>
2026-06-02 07:14:38 -07:00
0cbc48c4f9 Support ModelOpt MXFP8 non-gated MoE (#42958)
Signed-off-by: tbarnatan <tbarnatan@nvidia.com>
Co-authored-by: Michael Goin <mgoin64@gmail.com>
2026-06-02 13:56:03 +00:00
2fd0e52252 [Bugfix] Fix Gemma4 startup crash with recent transformers multimodal processor (#44232)
Signed-off-by: Luciano Martins <lucianommartins@users.noreply.github.com>
Co-authored-by: Luciano Martins <lucianommartins@users.noreply.github.com>
2026-06-02 13:42:40 +00:00
654bd2bca4 [Bugfix] Sync block_size from EngineCore to frontend for hybrid Mamba… (#42967)
Signed-off-by: Amit Gruner <agruner@crusoe.ai>
Co-authored-by: Amit Gruner <agruner@crusoe.ai>
Co-authored-by: Jiangyun Zhu <riverclouds.zhu@qq.com>
2026-06-02 13:41:00 +00:00
wang.yuqiandGitHub b623f7ea95 [Frontend] Consolidate dev entrypoints. (#44170)
Signed-off-by: wang.yuqi <yuqi.wang@daocloud.io>
2026-06-02 06:30:21 -07:00
Shreyas KulkarniandGitHub 0eeba5eec1 Fix DFlash prefix cache corruption due to missing lookahead block (#42971)
Signed-off-by: Shreyas Kulkarni <shreyas.gp269@gmail.com>
2026-06-02 12:06:33 +00:00
f69ede495b [XPU][Mamba] Triton-based selective scan forward op for XPU (#43421)
Signed-off-by: Marceli Fylcek <marceli.fylcek@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>
2026-06-02 03:50:26 -07:00
Ronen SchafferandGitHub 2a2b5ca791 [KV Offload] Add on_schedule_end() hook to separate step lifecycle from event draining (#44206)
Signed-off-by: Ronen Schaffer <ronen.schaffer@ibm.com>
2026-06-02 13:42:52 +03:00
Rukhaiya2004GitHubRukhaiyaAkash kaothalkarLi, Jiang <jiang1.li@intel.com>
689b0eeb9e [HARDWARE][POWER] Enable SHM communicator support for PowerPC (#43754)
Signed-off-by: Rukhaiya <rukhaiya@c643n08aix1-lp1.pok.stglabs.ibm.com>
Signed-off-by: Rukhaiya <bibirukhaiya123@gmail.com>
Co-authored-by: Rukhaiya <rukhaiya@c643n08aix1-lp1.pok.stglabs.ibm.com>
Co-authored-by: Akash kaothalkar <61960177+Akashcodes732@users.noreply.github.com>
Co-authored-by: Li, Jiang <jiang1.li@intel.com>
2026-06-02 18:06:32 +08:00
Isotr0pyandGitHub f8e9c56d15 [Multimodal] Automatically select registered video loader for VLM (#44126)
Signed-off-by: Isotr0py <mozf@mail2.sysu.edu.cn>
2026-06-02 09:09:47 +00:00
albertoandGitHub e30313220c [Parser] Migrate ResponsesParser to unified Parser interface (#42977)
Signed-off-by: Alberto Perdomo <aperdomo@redhat.com>
2026-06-02 08:50:05 +00:00
omerpaz95GitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
d247a9dc13 [EC Connector] Non blocking EC Connector lookup (#41627)
Signed-off-by: omerpaz95 <omerpaz95@gmail.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-06-02 08:48:25 +00:00
Yifan QiaoandGitHub 7c37096620 [Core][Refactor]: thread scheduler_block_size into KVCacheManager and KVCacheCoordinator (#44165)
Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai>
2026-06-02 01:14:44 -07:00
Maria GuevaraandGitHub b817b23f7b [Rust Frontend] add --enable-request-id-headers flag support. (#43883)
Signed-off-by: Maria Guevara <kawaiiplush14@gmail.com>
2026-06-02 16:08:37 +08:00
Ronen SchafferandGitHub 93da882e73 [kv_offload] Add @override decorators to subclass method implementations (#44177)
Signed-off-by: Ronen Schaffer <ronen.schaffer@ibm.com>
2026-06-02 08:07:47 +00:00
Fadi ArafehGitHubLi, Jiang <jiang1.li@intel.com>
0b25cf4419 [CPU][Perf] Enable fused kernels for GDN's gated delta rules (#43534)
Signed-off-by: Fadi Arafeh <fadi.arafeh@arm.com>
Co-authored-by: Li, Jiang <jiang1.li@intel.com>
2026-06-02 08:00:48 +00:00
Jiangyun ZhuandGitHub dcdfe66bfa [Perf] use triton moe backend on hopper by default (#44220)
Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>
2026-06-02 15:52:30 +08:00
Flora FengandGitHub 68dafcca75 [Refactor] Unify reasoning + tool-call parsing behind Parser.parse() (#44267)
Signed-off-by: sfeng33 <4florafeng@gmail.com>
2026-06-02 15:11:42 +08:00
zhrrrandGitHub 1edfd09ffd [Model Runner V2] Use actual batch max_seq_len for attn metadata (#43991)
Signed-off-by: zhuhaoran <zhuhaoran.zhr@alibaba-inc.com>
2026-06-02 06:07:56 +00:00
zhrrrandGitHub 8a9eb40808 [Model Runner V2] Support zeroing freshly allocated KV blocks for hybrid + fp8 KVCache (#43990)
Signed-off-by: zhuhaoran <zhuhaoran.zhr@alibaba-inc.com>
2026-06-02 05:56:53 +00:00
f91fb2fcf3 [Bugfix] Convert Gemma4-MM ViT linear layers to vllm native impl (#43798)
Signed-off-by: Isotr0py <mozf@mail2.sysu.edu.cn>
Co-authored-by: ZiTian Zhao <zitian.zhao@tencentmusic.com>
Co-authored-by: B-201 <Joy25810@foxmail.com>
2026-06-01 21:41:16 -07:00
JooHo LeeandGitHub a045c7425f [MM][CG] Profile encoder CUDA graph pool memory (#41714)
Signed-off-by: JooHo Lee <jooho414@gmail.com>
2026-06-02 12:27:34 +08:00
a3a5a5ece5 [XPU][Bugfix] Fix per_token_group_fp8_quant missing dummy args on XPU (#43930)
Signed-off-by: Chaojun,Zhang <chaojun.zhang@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>
2026-06-02 03:09:21 +00:00
Or OzeriandGitHub 480fadab1b [BugFix][kv_offload]: Prevent offloading stale sliding window blocks (#42959)
Signed-off-by: Or Ozeri <oro@il.ibm.com>
2026-06-02 05:59:48 +03:00
279d25f5cb [BugFix] Fix TypeError in MiniCPM-O audio feature unpadding (#38053)
Signed-off-by: Krishna Chaitanya Balusu <krishnabkc15@gmail.com>
Signed-off-by: wjinxu <1299461899@qq.com>
Signed-off-by: Kc Balusu <kcbalusu@users.noreply.github.com>
Co-authored-by: wjinxu <1299461899@qq.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cyrus Leung <tlleungac@connect.ust.hk>
Co-authored-by: Kc Balusu <kcbalusu@users.noreply.github.com>
2026-06-01 19:57:28 -07:00
Andreas KaratzasandGitHub 54d0c36fff [CI] Stabilize OpenAI schema fuzzing for malformed structural tags (#44131)
Signed-off-by: Andreas Karatzas <akaratza@amd.com>
2026-06-01 19:56:15 -07:00
Flora FengandGitHub 9affc17a05 [Refactor] Move unstreamed tool-arg flush from serving layer to parser (#44017)
Signed-off-by: sfeng33 <4florafeng@gmail.com>
2026-06-02 10:37:43 +08:00
AlecandGitHub 816cc73a9b [Bugfix][CI] Normalize NIXL connector CUDA wheel installs (#44266)
Signed-off-by: Alec Flowers <aflowers@nvidia.com>
2026-06-01 19:34:05 -07:00
Micah WilliamsonandGitHub 2588ec4f0a [ROCm] Upgrade AITER to v0.1.13.post1 (#44265)
Signed-off-by: Micah Williamson <micah.williamson@amd.com>
2026-06-02 01:48:59 +00:00
d68f0b220e [Bugfix][Mooncake] Release GPU pin on failed store in MooncakeStoreConnector (#43742)
Signed-off-by: Dao Le <Dao007forever@gmail.com>
Co-authored-by: Claude <noreply@anthropic.com>
2026-06-01 18:29:18 -07:00
Woosuk KwonandGitHub 517e74a964 [DSV4] Refactor RoPE initialization (#44262)
Signed-off-by: Woosuk Kwon <woosuk@inferact.ai>
2026-06-02 01:26:58 +00:00
JartXandGitHub 48c0d13e65 [ROCm][CI] Skip unbacked dynamic shapes tests on PyTorch < 2.11 (#44256)
Signed-off-by: JartX <sagformas@epdcenter.es>
2026-06-01 19:09:01 -05:00
Woosuk KwonandGitHub 8c3cc98cff [DSV4] Remove unncessary classes & functions (#44246)
Signed-off-by: Woosuk Kwon <woosuk@inferact.ai>
2026-06-01 14:43:00 -07:00
Nick HillandGitHub e4cbc4385d [Test][BugFix] Fix double-BOS in PD+specdec acceptance test (#44234)
Signed-off-by: Nick Hill <nickhill123@gmail.com>
2026-06-01 14:31:12 -07:00
Nick HillandGitHub 6f8b40a23f [BugFix][CI] Fix added _has_module tests (#44248)
Signed-off-by: Nick Hill <nickhill123@gmail.com>
2026-06-01 14:23:12 -07:00
266b9d9c64 [Frontend][Core] Add sparse NCCL weight transfer support for in-place updates (#40096)
Signed-off-by: Siddharth Bedekar <bedeksid@gmail.com>
Co-authored-by: OpenAI Codex <codex@openai.com>
2026-06-01 15:37:30 -04:00
182c67daf1 [Rust Frontend] Support streaming generate endpoint (#43779)
Signed-off-by: xunzhuo <xunzhuo@vllm-semantic-router.ai>
Co-authored-by: Bugen Zhao <i@bugenzhao.com>
2026-06-01 19:30:55 +00:00
fd9e91d7e4 [ROCm][CI] Fix and stabilize EAGLE3 acceptance tests (#41294)
Signed-off-by: Andreas Karatzas <akaratza@amd.com>
Signed-off-by: Micah Williamson <micah.williamson@amd.com>
Co-authored-by: Micah Williamson <micah.williamson@amd.com>
2026-06-01 12:40:01 -05:00
Yongye ZhuandGitHub 035733515f [Kernel][DSv4] Optimize sparse FP8 compressor kernels (#44161)
Signed-off-by: Yongye Zhu <zyy1102000@gmail.com>
2026-06-02 00:18:32 +08:00
023808c23d [Feature] Add support for JetBrains' Mellum v2 code generation model (#43992)
Signed-off-by: Madeesh Kannan <madeeswaran.kannan@jetbrains.com>
Co-authored-by: Robert Shaw <114415538+robertgshaw2-redhat@users.noreply.github.com>
2026-06-01 10:11:35 -04:00
Wentao YeGitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
985c97a6a8 [Perf] Optimize cutlass fp8 scaled mm bypassing padding, 20% kernel performance improvement (#43706)
Signed-off-by: yewentao256 <zhyanwentao@126.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-06-01 09:05:21 -04:00
Chaojun ZhangandGitHub bd0aecdc08 [XPU][CI] Fix test_audio_in_video flake by using module-scoped server fixture (#44146)
Signed-off-by: Chaojun Zhang <chaojun.zhang@intel.com>
2026-06-01 11:21:36 +00:00
8796838910 [Bugfix] fix wrong partial_rotary_factor calculation for bailing_moe model. (#43770)
Signed-off-by: zzt <zengzetang.zzt@antgroup.com>
Co-authored-by: Jiangyun Zhu <riverclouds.zhu@qq.com>
2026-06-01 02:42:49 -07:00
de21863419 [Rust Frontend] Add InternLM2 tool parser (#43481)
Signed-off-by: Will.hou <1205157517@qq.com>
Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: Bugen Zhao <i@bugenzhao.com>
2026-06-01 08:58:46 +00:00
wang.yuqiandGitHub 0910f7e0e1 [Frontend] Resettle generative scoring entrypoint. (#44153)
Signed-off-by: wang.yuqi <yuqi.wang@daocloud.io>
2026-06-01 07:54:59 +00:00
UranusandGitHub 1f6048abe5 fix: glm5.1 pp model loading (#42944)
Signed-off-by: UranusSeven <109661872+UranusSeven@users.noreply.github.com>
2026-06-01 15:14:47 +08:00
wcyGitHubLi, Jiang <jiang1.li@intel.com>
98f1279815 [CPU][RISC-V] Add missing RVV cpu_types helpers for WNA16 (#42730)
Signed-off-by: wcy <233313160abc@gmail.com>
Co-authored-by: Li, Jiang <jiang1.li@intel.com>
2026-06-01 14:56:41 +08:00
Isotr0pyandGitHub 1fd8bd02a4 [Docs] Replace broken video url in examples (#44159)
Signed-off-by: Isotr0py <mozf@mail2.sysu.edu.cn>
2026-06-01 06:01:10 +00:00
29d69332aa [BugFix] Fix _has_module to verify native deps via trial import (#44035)
Signed-off-by: esmeetu <jasonailu87@gmail.com>
Signed-off-by: Jeffrey Wang <jeffreywang@anyscale.com>
Signed-off-by: Nick Hill <nickhill123@gmail.com>
Co-authored-by: esmeetu <jasonailu87@gmail.com>
Co-authored-by: Nick Hill <nickhill123@gmail.com>
2026-05-31 22:06:33 -07:00
Lucas WilkinsonandGitHub 4721bb3aa4 [MRV2] Remove Eagle's dedicated CUDA graph pool (#44078)
Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
2026-05-31 22:00:33 -07:00
Umut PolatandGitHub f46e6be169 [Misc] Use VLLMValidationError consistently in chat completion and completion protocol validators (#36254)
Signed-off-by: umut-polat <52835619+umut-polat@users.noreply.github.com>
2026-06-01 04:04:11 +00:00
8b8546da1c docs: fix MLA attention docstring examples (#44118)
Co-authored-by: nightcityblade <nightcityblade@gmail.com>
2026-05-31 12:28:38 -07:00
362 changed files with 10195 additions and 26851 deletions
+3 -1
View File
@@ -16,6 +16,7 @@ steps:
- tests/kernels/test_onednn.py
- tests/kernels/test_awq_int4_to_int8.py
- tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
- tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
commands:
- |
bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
@@ -24,7 +25,8 @@ steps:
pytest -x -v -s tests/kernels/moe/test_cpu_quant_fused_moe.py
pytest -x -v -s tests/kernels/test_onednn.py
pytest -x -v -s tests/kernels/test_awq_int4_to_int8.py
pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py"
pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py"
- label: CPU-Compatibility Tests
depends_on: []
@@ -37,7 +37,8 @@ function cpu_tests() {
pytest -x -v -s tests/kernels/test_onednn.py
pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
pytest -x -v -s tests/kernels/core/test_cpu_activation.py
pytest -x -v -s tests/kernels/moe/test_moe.py -k test_cpu_fused_moe_basic"
pytest -x -v -s tests/kernels/moe/test_moe.py -k test_cpu_fused_moe_basic
pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py"
# skip tests requiring model downloads if HF_TOKEN is not set
# due to rate-limits
+39
View File
@@ -0,0 +1,39 @@
#!/bin/bash
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
set -euo pipefail
REQUIREMENTS_FILE="${KV_CONNECTORS_REQUIREMENTS:-/vllm-workspace/requirements/kv_connectors.txt}"
uv pip install --system -r "${REQUIREMENTS_FILE}"
NIXL_METADATA=$(python3 - <<'PY'
import importlib.metadata as metadata
import torch
cuda_version = torch.version.cuda
if cuda_version is None:
raise SystemExit("torch.version.cuda is not set")
print(cuda_version.split(".", 1)[0], metadata.version("nixl"))
PY
)
read -r CUDA_MAJOR NIXL_VERSION <<<"${NIXL_METADATA}"
# nixl>=1.1.0 can install multiple CUDA wheel variants. Keep only the variant
# matching this CI image so nixl_ep_cpp links against the available libcudart.
uv pip uninstall --system nixl-cu12 nixl-cu13 2>/dev/null || true
uv pip install --system --no-deps "nixl-cu${CUDA_MAJOR}==${NIXL_VERSION}"
python3 - <<'PY'
import importlib.metadata as metadata
for package_name in ("nixl", "nixl-cu12", "nixl-cu13"):
try:
version = metadata.version(package_name)
except metadata.PackageNotFoundError:
version = "not installed"
print(f"{package_name}: {version}")
PY
+15 -13
View File
@@ -1238,14 +1238,11 @@ steps:
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- tests/entrypoints/rpc
- tests/entrypoints/serve/instrumentator
- tests/tool_use
- tests/entrypoints/serve
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/serve/instrumentator
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
- pytest -v -s tool_use
- pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
- label: Entrypoints Integration (API Server openai - Part 1) # TBD
timeout_in_minutes: 180
@@ -1275,10 +1272,14 @@ steps:
- vllm/
- tests/entrypoints/openai
- tests/entrypoints/test_chat_utils
- tests/entrypoints/generate
- tests/tool_use
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
- pytest -v -s entrypoints/test_chat_utils.py
- pytest -v -s entrypoints/generate
- pytest -v -s tool_use
- label: Entrypoints Integration (API Server openai - Part 3) # TBD
timeout_in_minutes: 180
@@ -1368,7 +1369,7 @@ steps:
- vllm/platforms/rocm.py
commands:
- pytest -v -s entrypoints/openai/tool_parsers
- pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py --ignore=entrypoints/pooling
- pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/offline_mode --ignore=entrypoints/openai --ignore=entrypoints/serve --ignore=entrypoints/test_chat_utils.py --ignore=entrypoints/pooling --ignore=entrypoints/speech_to_text --ignore=tests/entrypoints/generate
- label: OpenAI API correctness # TBD
timeout_in_minutes: 180
@@ -2745,14 +2746,11 @@ steps:
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- tests/entrypoints/rpc
- tests/entrypoints/serve/instrumentator
- tests/tool_use
- tests/entrypoints/serve
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/serve/instrumentator
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
- pytest -v -s tool_use
- pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
- label: Entrypoints Integration (API Server openai - Part 1) # TBD
timeout_in_minutes: 180
@@ -2782,10 +2780,14 @@ steps:
- vllm/
- tests/entrypoints/openai
- tests/entrypoints/test_chat_utils
- tests/entrypoints/generate
- tests/tool_use
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
- pytest -v -s entrypoints/test_chat_utils.py
- pytest -v -s entrypoints/generate
- pytest -v -s tool_use
- label: Entrypoints Integration (API Server openai - Part 3) # TBD
timeout_in_minutes: 180
+9 -9
View File
@@ -11,7 +11,7 @@ steps:
- vllm/distributed/kv_transfer/kv_connector/v1/nixl/
- tests/v1/kv_connector/nixl_integration/
commands:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
- label: Distributed FlashInfer NixlConnector PD accuracy (4 GPUs)
key: distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus
@@ -22,7 +22,7 @@ steps:
- vllm/distributed/kv_transfer/kv_connector/v1/nixl/
- tests/v1/kv_connector/nixl_integration/
commands:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- FLASHINFER=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
- label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs)
@@ -34,7 +34,7 @@ steps:
- vllm/distributed/kv_transfer/kv_connector/v1/nixl/
- tests/v1/kv_connector/nixl_integration/
commands:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- DP_EP=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
- label: CrossLayer KV layout Distributed NixlConnector PD accuracy tests (4 GPUs)
@@ -46,7 +46,7 @@ steps:
- vllm/distributed/kv_transfer/kv_connector/v1/nixl/
- tests/v1/kv_connector/nixl_integration/
commands:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- CROSS_LAYERS_BLOCKS=True bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
- label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs)
@@ -58,7 +58,7 @@ steps:
- vllm/distributed/kv_transfer/kv_connector/v1/nixl/
- tests/v1/kv_connector/nixl_integration/
commands:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- HYBRID_SSM=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
- label: MultiConnector (Nixl+Offloading) PD accuracy (2 GPUs)
@@ -73,7 +73,7 @@ steps:
- vllm/distributed/kv_transfer/kv_connector/v1/offloading/
- tests/v1/kv_connector/nixl_integration/
commands:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- bash v1/kv_connector/nixl_integration/run_multi_connector_accuracy_test.sh
- label: NixlConnector PD + Spec Decode acceptance (2 GPUs)
@@ -87,7 +87,7 @@ steps:
- vllm/v1/worker/kv_connector_model_runner_mixin.py
- tests/v1/kv_connector/nixl_integration/
commands:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- bash v1/kv_connector/nixl_integration/config_sweep_spec_decode_test.sh
- label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs)
@@ -102,5 +102,5 @@ steps:
- vllm/distributed/kv_transfer/kv_connector/v1/offloading/
- tests/v1/kv_connector/nixl_integration/
commands:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
- bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh
+8 -8
View File
@@ -11,7 +11,7 @@ steps:
- tests/entrypoints/
commands:
- pytest -v -s entrypoints/openai/tool_parsers
- pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py --ignore=entrypoints/pooling --ignore=entrypoints/speech_to_text
- pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/offline_mode --ignore=entrypoints/openai --ignore=entrypoints/serve --ignore=entrypoints/test_chat_utils.py --ignore=entrypoints/pooling --ignore=entrypoints/speech_to_text --ignore=tests/entrypoints/generate
- label: Entrypoints Integration (LLM)
key: entrypoints-integration-llm
@@ -60,9 +60,13 @@ steps:
- vllm/
- tests/entrypoints/openai
- tests/entrypoints/test_chat_utils
- tests/entrypoints/generate
- tests/tool_use
commands:
- pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
- pytest -v -s entrypoints/test_chat_utils.py
- pytest -v -s entrypoints/generate
- pytest -v -s tool_use
mirror:
amd:
device: mi325_1
@@ -98,14 +102,11 @@ steps:
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- tests/entrypoints/rpc
- tests/entrypoints/serve/instrumentator
- tests/tool_use
- tests/entrypoints/serve
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/serve/instrumentator
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
- pytest -v -s tool_use
- pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
mirror:
amd:
device: mi325_1
@@ -153,6 +154,5 @@ steps:
source_file_dependencies:
- csrc/
- vllm/entrypoints/openai/
- vllm/model_executor/models/whisper.py
commands: # LMEval
- pytest -s entrypoints/openai/correctness/
+1 -1
View File
@@ -86,7 +86,7 @@ steps:
- tests/v1/metrics
- tests/entrypoints/openai/correctness/test_lmeval.py
commands:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
# split the test to avoid interference
- pytest -v -s -m 'not cpu_test' v1/core
+4 -4
View File
@@ -45,19 +45,19 @@ steps:
- vllm/entrypoints/serve/
- vllm/v1/engine/
- tests/utils.py
# - tests/entrypoints/rpc/test_collective_rpc.py
# - tests/entrypoints/serve/dev/rpc/test_collective_rpc.py
- tests/entrypoints/serve/disagg/test_serving_tokens.py
- tests/entrypoints/serve/instrumentator/test_basic.py
- tests/entrypoints/serve/instrumentator/test_metrics.py
# - tests/entrypoints/serve/instrumentator/test_sleep.py
# - tests/entrypoints/serve/dev/test_sleep.py
commands:
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
# - pytest -v -s entrypoints/rpc/test_collective_rpc.py
# - pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
- pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not show_version and not server_load"
- pytest -v -s entrypoints/serve/disagg/test_serving_tokens.py -k "not stream and not lora and not test_generate_logprobs and not stop_string_workflow"
- pytest -v -s entrypoints/serve/instrumentator/test_metrics.py -k "text and not show and not run_batch and not test_metrics_counts and not test_metrics_exist"
# - pytest -v -s entrypoints/serve/instrumentator/test_sleep.py
# - pytest -v -s entrypoints/serve/dev/test_sleep.py
- label: Rust Frontend Core Correctness
timeout_in_minutes: 30
+1 -1
View File
@@ -10,7 +10,7 @@ jobs:
runs-on: ubuntu-latest
steps:
- name: Add label
uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
with:
script: |
github.rest.issues.addLabels({
+3 -3
View File
@@ -14,7 +14,7 @@ jobs:
steps:
- name: Label issues based on keywords
id: label-step
uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
with:
script: |
// Configuration: Add new labels and keywords here
@@ -315,7 +315,7 @@ jobs:
- name: CC users for labeled issues
if: steps.label-step.outputs.labels_added != '[]'
uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
with:
script: |
// Configuration: Map labels to GitHub users to CC
@@ -392,7 +392,7 @@ jobs:
- name: Request missing ROCm info from issue author
if: contains(steps.label-step.outputs.labels_added, 'rocm') && contains(toJSON(github.event.issue.labels.*.name), 'bug')
uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
with:
script: |
const body = (context.payload.issue.body || '').toLowerCase();
+2 -2
View File
@@ -12,7 +12,7 @@ jobs:
runs-on: ubuntu-latest
steps:
- name: Update PR description
uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
with:
script: |
const { owner, repo } = context.repo;
@@ -55,7 +55,7 @@ jobs:
runs-on: ubuntu-latest
steps:
- name: Post welcome comment for first-time contributors
uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
with:
script: |
const { owner, repo } = context.repo;
+2 -2
View File
@@ -20,7 +20,7 @@ jobs:
runs-on: ubuntu-latest
steps:
- name: Check PR label and author merge count
uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
with:
script: |
const { data: pr } = await github.rest.pulls.get({
@@ -49,7 +49,7 @@ jobs:
runs-on: ubuntu-latest
steps:
- uses: actions/checkout@8e8c483db84b4bee98b60c0593521ed34d9990e8 # v6.0.1
- uses: actions/setup-python@83679a892e2d95755f2dac6acb0bfd1e9ac5d548 # v6.1.0
- uses: actions/setup-python@a309ff8b426b58ec0e2a45f0f869d46889d02405 # v6.2.0
with:
python-version: "3.12"
- run: echo "::add-matcher::.github/workflows/matchers/actionlint.json"
+1 -1
View File
@@ -21,7 +21,7 @@ repos:
rev: v21.1.2
hooks:
- id: clang-format
exclude: 'csrc/(moe/topk_softmax_kernels.cu|quantization/gguf/(ggml-common.h|dequantize.cuh|vecdotq.cuh|mmq.cuh|mmvq.cuh))|vllm/third_party/.*'
exclude: 'csrc/(moe/topk_softmax_kernels.cu|libtorch_stable/quantization/gguf/(ggml-common.h|dequantize.cuh|vecdotq.cuh|mmq.cuh|mmvq.cuh))|vllm/third_party/.*'
types_or: [c++, cuda]
args: [--style=file, --verbose]
- repo: https://github.com/DavidAnson/markdownlint-cli2
+1 -3
View File
@@ -315,8 +315,7 @@ set(VLLM_EXT_SRC
if(VLLM_GPU_LANG STREQUAL "CUDA")
list(APPEND VLLM_EXT_SRC
"csrc/minimax_reduce_rms_kernel.cu"
"csrc/minimax_m3_build_k2q_csr.cu")
"csrc/minimax_reduce_rms_kernel.cu")
SET(CUTLASS_ENABLE_HEADERS_ONLY ON CACHE BOOL "Enable only the header library")
@@ -638,7 +637,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
"csrc/libtorch_stable/quantization/gguf/gguf_kernel.cu"
"csrc/libtorch_stable/pos_encoding_kernels.cu"
"csrc/libtorch_stable/fused_qknorm_rope_kernel.cu"
"csrc/libtorch_stable/fused_minimax_m3_qknorm_rope_kv_insert_kernel.cu"
"csrc/libtorch_stable/layernorm_kernels.cu"
"csrc/libtorch_stable/layernorm_quant_kernels.cu"
"csrc/libtorch_stable/quantization/fused_kernels/fused_layernorm_dynamic_per_token_quant.cu"
+31 -1
View File
@@ -369,6 +369,18 @@ else()
add_compile_definitions(-DVLLM_NUMA_DISABLED)
endif()
# check if the pytorch wheel ships libopenblas.so.
set(VLLM_OPENBLAS_LIB "")
if (NOT ENABLE_X86_ISA)
file(GLOB _VLLM_TORCH_OPENBLAS_LIBS
"${TORCH_INSTALL_PREFIX}/lib/libopenblas*.so*")
# Note: we don't link openblas directly to _C extension, as it's available through libtorch.so
if (_VLLM_TORCH_OPENBLAS_LIBS)
list(GET _VLLM_TORCH_OPENBLAS_LIBS 0 VLLM_OPENBLAS_LIB)
message(STATUS "CPU OpenBLAS library: ${VLLM_OPENBLAS_LIB}")
endif()
endif()
#
# Generate CPU attention dispatch header
#
@@ -387,6 +399,7 @@ endif()
#
set(VLLM_EXT_SRC
"csrc/cpu/activation.cpp"
"csrc/cpu/sgl-kernels/fla.cpp"
"csrc/cpu/utils.cpp"
"csrc/cpu/spec_decode_utils.cpp"
"csrc/cpu/layernorm.cpp"
@@ -396,6 +409,13 @@ set(VLLM_EXT_SRC
"csrc/cpu/cpu_attn.cpp"
"csrc/cpu/torch_bindings.cpp")
if (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64" AND VLLM_RVV_VLEN AND
VLLM_RVV_VLEN GREATER 0 AND (RVV_FP16_FOUND OR RVV_BF16_FOUND))
set(VLLM_EXT_SRC
"csrc/cpu/cpu_wna16.cpp"
${VLLM_EXT_SRC})
endif()
if (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND)
set(VLLM_EXT_SRC
"csrc/cpu/shm.cpp"
@@ -403,6 +423,12 @@ if (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND)
${VLLM_EXT_SRC})
endif()
if (POWER9_FOUND OR POWER10_FOUND OR POWER11_FOUND)
set(VLLM_EXT_SRC
"csrc/cpu/shm.cpp"
${VLLM_EXT_SRC})
endif()
if(USE_ONEDNN)
set(VLLM_EXT_SRC
"csrc/cpu/dnnl_kernels.cpp"
@@ -411,7 +437,6 @@ endif()
if (ENABLE_X86_ISA)
set(VLLM_EXT_SRC_SGL
"csrc/cpu/sgl-kernels/fla.cpp"
"csrc/cpu/sgl-kernels/conv.cpp"
"csrc/cpu/sgl-kernels/gemm.cpp"
"csrc/cpu/sgl-kernels/gemm_int8.cpp"
@@ -423,6 +448,7 @@ if (ENABLE_X86_ISA)
"csrc/cpu/sgl-kernels/moe_fp8.cpp")
set(VLLM_EXT_SRC_AVX512
"csrc/cpu/sgl-kernels/fla.cpp"
"csrc/cpu/shm.cpp"
"csrc/cpu/cpu_wna16.cpp"
"csrc/cpu/cpu_fused_moe.cpp"
@@ -439,6 +465,7 @@ if (ENABLE_X86_ISA)
"csrc/moe/dynamic_4bit_int_moe_cpu.cpp")
set(VLLM_EXT_SRC_AVX2
"csrc/cpu/sgl-kernels/fla.cpp"
"csrc/cpu/utils.cpp"
"csrc/cpu/spec_decode_utils.cpp"
"csrc/cpu/cpu_attn.cpp"
@@ -512,6 +539,9 @@ else()
USE_SABI 3
WITH_SOABI
)
if (VLLM_OPENBLAS_LIB)
target_compile_definitions(_C PRIVATE VLLM_HAS_OPENBLAS)
endif()
endif()
message(STATUS "Enabling C extension.")
@@ -31,7 +31,7 @@ endif()
if(VLLM_FLASH_ATTN_SRC_DIR)
FetchContent_Declare(
vllm-flash-attn SOURCE_DIR
vllm-flash-attn SOURCE_DIR
${VLLM_FLASH_ATTN_SRC_DIR}
BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn
)
@@ -39,7 +39,7 @@ else()
FetchContent_Declare(
vllm-flash-attn
GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
GIT_TAG bce29425653ec0fbc579d329883030e832d15ada
GIT_TAG dd62dac706b1cf7895bd99b18c6cb7e7e117ee25
GIT_PROGRESS TRUE
# Don't share the vllm-flash-attn build between build types
BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn
+51
View File
@@ -94,6 +94,10 @@ struct FP16Vec16 : public Vec<FP16Vec16> {
: reg(RVVI(__riscv_vle16_v_f16, LMUL_256)(
static_cast<const _Float16*>(ptr), VEC_ELEM_NUM)) {};
explicit FP16Vec16(const c10::Half v)
: reg(RVVI4(__riscv_vreinterpret_v_u16, LMUL_256, _f16, LMUL_256)(
RVVI(__riscv_vmv_v_x_u16, LMUL_256)(v.x, VEC_ELEM_NUM))) {};
explicit FP16Vec16(const FP32Vec16& vec);
void save(void* ptr) const {
@@ -165,6 +169,9 @@ struct BF16Vec16 : public Vec<BF16Vec16> {
reinterpret_cast<const uint16_t*>(ptr), VEC_ELEM_NUM))) {};
explicit BF16Vec16(fixed_bf16x16_t data) : reg(data) {};
explicit BF16Vec16(const c10::BFloat16 v)
: reg(RVVI4(__riscv_vreinterpret_v_u16, LMUL_256, _bf16, LMUL_256)(
RVVI(__riscv_vmv_v_x_u16, LMUL_256)(v.x, VEC_ELEM_NUM))) {};
explicit BF16Vec16(const FP32Vec16&);
void save(void* ptr) const {
@@ -290,6 +297,9 @@ struct BF16Vec16 : public Vec<BF16Vec16> {
}
reg_fp32 = RVVI(__riscv_vle32_v_f32, LMUL_512)(tmp, 16);
}
explicit BF16Vec16(const c10::BFloat16 v)
: reg_fp32(RVVI(__riscv_vfmv_v_f_f32, LMUL_512)(static_cast<float>(v),
VEC_ELEM_NUM)) {}
explicit BF16Vec16(const FP32Vec16&);
void save(void* ptr) const {
float tmp[16];
@@ -629,6 +639,19 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
: reg(RVVI4(__riscv_vcreate_v_f32, LMUL_256, _f32, LMUL_512)(
data.reg, data.reg)) {};
explicit FP32Vec16(const FP32Vec16& data) : reg(data.reg) {};
explicit FP32Vec16(int64_t value, const FP32Vec16& lut) {
const uint64_t q_values = static_cast<uint64_t>(value);
auto packed = RVVI(__riscv_vmv_v_x_u64, LMUL_1024)(q_values, VEC_ELEM_NUM);
auto lane_ids = RVVI(__riscv_vid_v_u64, LMUL_1024)(VEC_ELEM_NUM);
auto shifts =
RVVI(__riscv_vsll_vx_u64, LMUL_1024)(lane_ids, 2, VEC_ELEM_NUM);
auto shifted =
RVVI(__riscv_vsrl_vv_u64, LMUL_1024)(packed, shifts, VEC_ELEM_NUM);
auto idx64 =
RVVI(__riscv_vand_vx_u64, LMUL_1024)(shifted, 0xF, VEC_ELEM_NUM);
auto idx32 = RVVI(__riscv_vnsrl_wx_u32, LMUL_512)(idx64, 0, VEC_ELEM_NUM);
reg = RVVI(__riscv_vrgather_vv_f32, LMUL_512)(lut.reg, idx32, VEC_ELEM_NUM);
}
explicit FP32Vec16(const FP16Vec16& v);
#ifdef __riscv_zvfbfmin
@@ -641,6 +664,10 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
explicit FP32Vec16(const BF16Vec16& v) : reg(v.reg_fp32) {};
#endif
// FP8 stub: dead code on RISC-V (fp8 KV cache is x86-only), needed for
// load_b_pair_vec template to compile on all platforms.
explicit FP32Vec16(const BF16Vec32&, int) : FP32Vec16() {}
FP32Vec16 operator+(const FP32Vec16& b) const {
return FP32Vec16(
RVVI(__riscv_vfadd_vv_f32, LMUL_512)(reg, b.reg, VEC_ELEM_NUM));
@@ -891,6 +918,30 @@ inline void fma(FP32Vec16& acc, const FP32Vec16& a, const FP32Vec16& b) {
acc = acc.fma(a, b);
}
template <typename VecT>
static void interleave_save_16b(const VecT& vec0, const VecT& vec1, void* ptr) {
alignas(64) uint16_t values0[VecT::VEC_ELEM_NUM];
alignas(64) uint16_t values1[VecT::VEC_ELEM_NUM];
vec0.save(values0);
vec1.save(values1);
auto* packed = reinterpret_cast<uint32_t*>(ptr);
for (int32_t i = 0; i < VecT::VEC_ELEM_NUM; ++i) {
packed[i] = static_cast<uint32_t>(values0[i]) |
(static_cast<uint32_t>(values1[i]) << 16);
}
}
static void interleave_save(const FP16Vec16& vec0, const FP16Vec16& vec1,
void* ptr) {
interleave_save_16b(vec0, vec1, ptr);
}
static void interleave_save(const BF16Vec16& vec0, const BF16Vec16& vec1,
void* ptr) {
interleave_save_16b(vec0, vec1, ptr);
}
#ifdef __riscv_zvfbfmin
template <>
inline void storeFP32<c10::BFloat16>(float v, c10::BFloat16* ptr) {
+106 -1
View File
@@ -89,6 +89,35 @@ struct BF16Vec8 : public Vec<BF16Vec8> {
}
};
struct FP16Vec16 : public Vec<FP16Vec16> {
constexpr static int VEC_ELEM_NUM = 16;
ss16x8x2_t reg;
explicit FP16Vec16(const void* ptr) {
reg.val[0] = (__vector signed short)vec_xl(0, (signed short*)ptr);
reg.val[1] = (__vector signed short)vec_xl(16, (signed short*)ptr);
}
explicit FP16Vec16(bool, const void* ptr) : FP16Vec16(ptr) {}
explicit FP16Vec16(const FP32Vec16&);
void save(void* ptr) const {
vec_xst(reg.val[0], 0, (signed short*)ptr);
vec_xst(reg.val[1], 16, (signed short*)ptr);
}
void save(void* ptr, int elem_num) const {
int num = std::max(0, std::min(elem_num, VEC_ELEM_NUM));
if (num <= 8) {
vec_xst_len(reg.val[0], (signed short*)ptr, num * 2);
} else {
vec_xst(reg.val[0], 0, (signed short*)ptr);
vec_xst_len(reg.val[1], (signed short*)ptr + 8, (num - 8) * 2);
}
}
};
struct BF16Vec16 : public Vec<BF16Vec16> {
constexpr static int VEC_ELEM_NUM = 16;
@@ -100,6 +129,8 @@ struct BF16Vec16 : public Vec<BF16Vec16> {
reg.val[1] = (__vector signed short)vec_xl(16, (signed short*)ptr);
}
explicit BF16Vec16(bool, const void* ptr) : BF16Vec16(ptr) {}
explicit BF16Vec16(const FP32Vec16&);
void save(void* ptr) const {
@@ -379,6 +410,8 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
reg.val[3] = vec_xl(48, ptr);
}
explicit FP32Vec16(bool, const float* ptr) : FP32Vec16(ptr) {}
explicit FP32Vec16(f32x4x4_t data) : reg(data) {}
explicit FP32Vec16(const FP32Vec16& data) {
@@ -402,6 +435,7 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
reg.val[3] = data.reg.val[1];
}
explicit FP32Vec16(const FP16Vec16& v);
explicit FP32Vec16(const BF16Vec16& v) {
reg.val[0] = (__vector float)vec_mergeh(zero, v.reg.val[0]);
reg.val[1] = (__vector float)vec_mergel(zero, v.reg.val[0]);
@@ -735,6 +769,40 @@ inline BF16Vec8::BF16Vec8(const FP32Vec8& v) {
#endif
}
inline FP16Vec16::FP16Vec16(const FP32Vec16& v) {
alignas(16) float temp_fp32[16];
alignas(16) c10::Half temp_fp16[16];
vec_xst(v.reg.val[0], 0, temp_fp32);
vec_xst(v.reg.val[1], 16, temp_fp32);
vec_xst(v.reg.val[2], 32, temp_fp32);
vec_xst(v.reg.val[3], 48, temp_fp32);
for (int i = 0; i < 16; i++) {
temp_fp16[i] = c10::Half(temp_fp32[i]);
}
reg.val[0] = (__vector signed short)vec_xl(0, (signed short*)temp_fp16);
reg.val[1] = (__vector signed short)vec_xl(16, (signed short*)temp_fp16);
}
inline FP32Vec16::FP32Vec16(const FP16Vec16& v) {
alignas(16) c10::Half temp_fp16[16];
alignas(16) float temp_fp32[16];
vec_xst(v.reg.val[0], 0, (signed short*)temp_fp16);
vec_xst(v.reg.val[1], 16, (signed short*)temp_fp16);
for (int i = 0; i < 16; i++) {
temp_fp32[i] = float(temp_fp16[i]);
}
reg.val[0] = vec_xl(0, temp_fp32);
reg.val[1] = vec_xl(16, temp_fp32);
reg.val[2] = vec_xl(32, temp_fp32);
reg.val[3] = vec_xl(48, temp_fp32);
}
inline BF16Vec16::BF16Vec16(const FP32Vec16& v) {
#ifdef _ARCH_PWR10
__vector signed short ret[4];
@@ -794,6 +862,43 @@ inline void prefetch(const void* addr) {
__asm__ __volatile__("dcbt 0, %0" : : "r"(addr) : "memory");
}
}; // namespace vec_op
struct INT8Vec64 {
__vector signed char data[4];
INT8Vec64() = default;
explicit INT8Vec64(const int8_t* ptr) {
data[0] = vec_xl(0, ptr);
data[1] = vec_xl(16, ptr);
data[2] = vec_xl(32, ptr);
data[3] = vec_xl(48, ptr);
}
explicit INT8Vec64(bool, const int8_t* ptr) : INT8Vec64(ptr) {}
void save(int8_t* ptr) const {
vec_xst(data[0], 0, ptr);
vec_xst(data[1], 16, ptr);
vec_xst(data[2], 32, ptr);
vec_xst(data[3], 48, ptr);
}
void save(int8_t* ptr, int elem_num) const {
if (elem_num <= 0) return;
int full_vecs = elem_num / 16;
for (int i = 0; i < full_vecs && i < 4; i++) {
vec_xst(data[i], i * 16, ptr);
}
int remaining = elem_num % 16;
if (remaining > 0 && full_vecs < 4) {
vec_xst_len(data[full_vecs], ptr + full_vecs * 16, remaining);
}
}
void nt_save(int8_t* ptr) const { save(ptr); }
};
} // namespace vec_op
#endif
+82
View File
@@ -0,0 +1,82 @@
// SPDX-License-Identifier: Apache-2.0
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
#include <ATen/native/CPUBlas.h>
// Unlike brgemm, PyTorch does not publicly expose at::native::cpublas::gemm
// If OpenBLS is available in the PyTorch wheel, we rely on it for fast
// bf16:bf16->fp32 GEMMs Otherwise, we fall back to PyTorch reference BLAS path.
#if defined(VLLM_HAS_OPENBLAS)
extern "C" void sbgemm_(char* transa, char* transb, int* m, int* n, int* k,
float* alpha, const at::BFloat16* a, int* lda,
const at::BFloat16* b, int* ldb, float* beta, float* c,
int* ldc);
extern "C" void sgemm_(char* transa, char* transb, int* m, int* n, int* k,
float* alpha, const float* a, int* lda, const float* b,
int* ldb, float* beta, float* c, int* ldc);
inline char blas_transpose(at::native::TransposeType trans) {
switch (trans) {
case at::native::TransposeType::NoTranspose:
return 'n';
case at::native::TransposeType::Transpose:
return 't';
case at::native::TransposeType::ConjTranspose:
return 'c';
}
return 'n';
}
inline void blas_gemm(at::native::TransposeType transa,
at::native::TransposeType transb, int64_t m, int64_t n,
int64_t k, float alpha, const at::BFloat16* a,
int64_t lda, const at::BFloat16* b, int64_t ldb,
float beta, float* c, int64_t ldc) {
char transa_ = blas_transpose(transa);
char transb_ = blas_transpose(transb);
int m_ = static_cast<int>(m);
int n_ = static_cast<int>(n);
int k_ = static_cast<int>(k);
int lda_ = static_cast<int>(lda);
int ldb_ = static_cast<int>(ldb);
int ldc_ = static_cast<int>(ldc);
sbgemm_(&transa_, &transb_, &m_, &n_, &k_, &alpha, a, &lda_, b, &ldb_, &beta,
c, &ldc_);
}
inline void blas_gemm(at::native::TransposeType transa,
at::native::TransposeType transb, int64_t m, int64_t n,
int64_t k, float alpha, const float* a, int64_t lda,
const float* b, int64_t ldb, float beta, float* c,
int64_t ldc) {
char transa_ = blas_transpose(transa);
char transb_ = blas_transpose(transb);
int m_ = static_cast<int>(m);
int n_ = static_cast<int>(n);
int k_ = static_cast<int>(k);
int lda_ = static_cast<int>(lda);
int ldb_ = static_cast<int>(ldb);
int ldc_ = static_cast<int>(ldc);
sgemm_(&transa_, &transb_, &m_, &n_, &k_, &alpha, a, &lda_, b, &ldb_, &beta,
c, &ldc_);
}
inline void blas_gemm(at::native::TransposeType, at::native::TransposeType,
int64_t, int64_t, int64_t, float, const at::Half*,
int64_t, const at::Half*, int64_t, float, float*,
int64_t) {
TORCH_CHECK(false, "CPU OpenBLAS hgemm is not available.");
}
#else
template <typename scalar_t>
inline void blas_gemm(at::native::TransposeType transa,
at::native::TransposeType transb, int64_t m, int64_t n,
int64_t k, float alpha, const scalar_t* a, int64_t lda,
const scalar_t* b, int64_t ldb, float beta, float* c,
int64_t ldc) {
auto gemm = at::native::cpublas::gemm_no_downcast_stub.DEFAULT;
gemm(c10::CppTypeToScalarType<scalar_t>::value, transa, transb, m, n, k,
at::Scalar(alpha), a, lda, b, ldb, at::Scalar(beta), c, ldc);
}
#endif
+278 -141
View File
@@ -301,25 +301,42 @@ void chunk_gated_delta_rule_kernel_impl(
// attn = k_beta @ key.transpose(-1, -2)
// attn: [B, HV, num_chunk, chunk_size, chunk_size]
// transpose and pack for key
pack_vnni<scalar_t>(
/* dst */ k_transpose,
/* src */ curr_k_pad,
/* N */ chunk_size,
/* K */ qk_head_size,
/* ld_src */ qk_head_size,
/* ld_dst */ chunk_size);
// k_beta @ key.transpose(-1, -2)
at::native::cpublas::brgemm(
/* M */ chunk_size,
/* N */ chunk_size,
/* K */ qk_head_size,
/* lda */ qk_head_size,
/* ldb */ chunk_size,
/* ldc */ chunk_size,
/* add_C */ false,
/* A */ curr_k_beta,
/* B */ k_transpose,
/* C */ curr_attn);
if constexpr (brgemm_supported()) {
pack_vnni<scalar_t>(
/* dst */ k_transpose,
/* src */ curr_k_pad,
/* N */ chunk_size,
/* K */ qk_head_size,
/* ld_src */ qk_head_size,
/* ld_dst */ chunk_size);
// k_beta @ key.transpose(-1, -2)
at::native::cpublas::brgemm(
/* M */ chunk_size,
/* N */ chunk_size,
/* K */ qk_head_size,
/* lda */ qk_head_size,
/* ldb */ chunk_size,
/* ldc */ chunk_size,
/* add_C */ false,
/* A */ curr_k_beta,
/* B */ k_transpose,
/* C */ curr_attn);
} else {
blas_gemm(
at::native::TransposeType::Transpose,
at::native::TransposeType::NoTranspose,
chunk_size,
chunk_size,
qk_head_size,
1.0f,
curr_k_pad,
qk_head_size,
curr_k_beta,
qk_head_size,
0.0f,
curr_attn,
chunk_size);
}
// attn = attn * decay_mask
for (int64_t m = 0; m < chunk_size; m++) {
at::vec::map2<float>(
@@ -413,25 +430,42 @@ void chunk_gated_delta_rule_kernel_impl(
// k_beta_g = k_beta * g: [B, HV, num_chunk, chunk_size, EK]
// k_cumdecay: [B, HV, num_chunk, chunk_size, EK]
// pack for value
pack_vnni2<scalar_t>(
/* dst */ v_pack,
/* src */ curr_v_beta,
/* N */ chunk_size,
/* K */ v_head_size,
/* ld_src */ v_head_size,
/* ld_dst */ v_head_size);
// value = attn @ v_beta
at::native::cpublas::brgemm(
/* M */ chunk_size,
/* N */ v_head_size,
/* K */ chunk_size,
/* lda */ chunk_size,
/* ldb */ v_head_size,
/* ldc */ v_head_size,
/* add_C */ false,
/* A */ curr_attn_reduced,
/* B */ v_pack,
/* C */ curr_value);
if constexpr (brgemm_supported()) {
pack_vnni2<scalar_t>(
/* dst */ v_pack,
/* src */ curr_v_beta,
/* N */ chunk_size,
/* K */ v_head_size,
/* ld_src */ v_head_size,
/* ld_dst */ v_head_size);
// value = attn @ v_beta
at::native::cpublas::brgemm(
/* M */ chunk_size,
/* N */ v_head_size,
/* K */ chunk_size,
/* lda */ chunk_size,
/* ldb */ v_head_size,
/* ldc */ v_head_size,
/* add_C */ false,
/* A */ curr_attn_reduced,
/* B */ v_pack,
/* C */ curr_value);
} else {
blas_gemm(
at::native::TransposeType::NoTranspose,
at::native::TransposeType::NoTranspose,
v_head_size,
chunk_size,
chunk_size,
1.0f,
curr_v_beta,
v_head_size,
curr_attn_reduced,
chunk_size,
0.0f,
curr_value,
v_head_size);
}
// k_beta_g = k_beta * g.exp().unsqueeze(-1)
for (int64_t j = 0; j < chunk_size; j++) {
int64_t i = 0;
@@ -445,25 +479,42 @@ void chunk_gated_delta_rule_kernel_impl(
}
}
// pack for k_beta_g
pack_vnni2<scalar_t>(
/* dst */ k_beta_g_pack,
/* src */ k_beta_g,
/* N */ chunk_size,
/* K */ qk_head_size,
/* ld_src */ qk_head_size,
/* ld_dst */ qk_head_size);
// k_cumdecay = attn @ k_beta_g
at::native::cpublas::brgemm(
/* M */ chunk_size,
/* N */ qk_head_size,
/* K */ chunk_size,
/* lda */ chunk_size,
/* ldb */ qk_head_size,
/* ldc */ qk_head_size,
/* add_C */ false,
/* A */ curr_attn_reduced,
/* B */ k_beta_g_pack,
/* C */ k_cumdecay);
if constexpr (brgemm_supported()) {
pack_vnni2<scalar_t>(
/* dst */ k_beta_g_pack,
/* src */ k_beta_g,
/* N */ chunk_size,
/* K */ qk_head_size,
/* ld_src */ qk_head_size,
/* ld_dst */ qk_head_size);
// k_cumdecay = attn @ k_beta_g
at::native::cpublas::brgemm(
/* M */ chunk_size,
/* N */ qk_head_size,
/* K */ chunk_size,
/* lda */ chunk_size,
/* ldb */ qk_head_size,
/* ldc */ qk_head_size,
/* add_C */ false,
/* A */ curr_attn_reduced,
/* B */ k_beta_g_pack,
/* C */ k_cumdecay);
} else {
blas_gemm(
at::native::TransposeType::NoTranspose,
at::native::TransposeType::NoTranspose,
qk_head_size,
chunk_size,
chunk_size,
1.0f,
k_beta_g,
qk_head_size,
curr_attn_reduced,
chunk_size,
0.0f,
k_cumdecay,
qk_head_size);
}
for (int i = 0; i < chunk_size; i++) {
at::vec::map<scalar_t>(
[](fVec x) { return x; },
@@ -551,25 +602,42 @@ void chunk_gated_delta_rule_kernel_impl(
// attn_i = (q_i @ k_i.transpose(-1, -2) * decay_mask[:, :, i]).masked_fill_(mask, 0)
// k_transpose_i = k_i.transpose(-1, -2)
pack_vnni<scalar_t>(
/* dst */ k_transpose_i,
/* src */ k_i,
/* N */ chunk_size,
/* K */ qk_head_size,
/* ld_src */ qk_head_size,
/* ld_dst */ chunk_size);
// attn_i = q_i @ k_transpose_i
at::native::cpublas::brgemm(
/* M */ chunk_size,
/* N */ chunk_size,
/* K */ qk_head_size,
/* lda */ qk_head_size,
/* ldb */ chunk_size,
/* ldc */ chunk_size,
/* add_C */ false,
/* A */ q_i,
/* B */ k_transpose_i,
/* C */ attn_i);
if constexpr (brgemm_supported()) {
pack_vnni<scalar_t>(
/* dst */ k_transpose_i,
/* src */ k_i,
/* N */ chunk_size,
/* K */ qk_head_size,
/* ld_src */ qk_head_size,
/* ld_dst */ chunk_size);
// attn_i = q_i @ k_transpose_i
at::native::cpublas::brgemm(
/* M */ chunk_size,
/* N */ chunk_size,
/* K */ qk_head_size,
/* lda */ qk_head_size,
/* ldb */ chunk_size,
/* ldc */ chunk_size,
/* add_C */ false,
/* A */ q_i,
/* B */ k_transpose_i,
/* C */ attn_i);
} else {
blas_gemm(
at::native::TransposeType::Transpose,
at::native::TransposeType::NoTranspose,
chunk_size,
chunk_size,
qk_head_size,
1.0f,
k_i,
qk_head_size,
q_i,
qk_head_size,
0.0f,
attn_i,
chunk_size);
}
// attn_i = attn_i * decay_mask_i
for (int64_t m = 0; m < chunk_size; m++) {
auto attn_i_m = attn_i + m * chunk_size;
@@ -609,28 +677,45 @@ void chunk_gated_delta_rule_kernel_impl(
}
// pack for curr_last_recurrent_state
pack_vnni2<scalar_t>(
/* dst */ curr_last_recurrent_state_pack_reduced,
/* src */ curr_last_recurrent_state_reduced,
/* N */ qk_head_size,
/* K */ v_head_size,
/* ld_src */ v_head_size,
/* ld_dst */ v_head_size);
if constexpr (brgemm_supported()) {
pack_vnni2<scalar_t>(
/* dst */ curr_last_recurrent_state_pack_reduced,
/* src */ curr_last_recurrent_state_reduced,
/* N */ qk_head_size,
/* K */ v_head_size,
/* ld_src */ v_head_size,
/* ld_dst */ v_head_size);
// v_prime = k_cumdecay_i @ curr_last_recurrent_state: [chunk_size, EV]
// k_cumdecay_i: [chunk_size, EK]
// curr_last_recurrent_state: [EK, EV]
at::native::cpublas::brgemm(
/* M */ chunk_size,
/* N */ v_head_size,
/* K */ qk_head_size,
/* lda */ qk_head_size,
/* ldb */ v_head_size,
/* ldc */ v_head_size,
/* add_C */ false,
/* A */ k_cumdecay_i_reduced,
/* B */ curr_last_recurrent_state_pack_reduced,
/* C */ v_prime);
// v_prime = k_cumdecay_i @ curr_last_recurrent_state: [chunk_size, EV]
// k_cumdecay_i: [chunk_size, EK]
// curr_last_recurrent_state: [EK, EV]
at::native::cpublas::brgemm(
/* M */ chunk_size,
/* N */ v_head_size,
/* K */ qk_head_size,
/* lda */ qk_head_size,
/* ldb */ v_head_size,
/* ldc */ v_head_size,
/* add_C */ false,
/* A */ k_cumdecay_i_reduced,
/* B */ curr_last_recurrent_state_pack_reduced,
/* C */ v_prime);
} else {
blas_gemm(
at::native::TransposeType::NoTranspose,
at::native::TransposeType::NoTranspose,
v_head_size,
chunk_size,
qk_head_size,
1.0f,
curr_last_recurrent_state_reduced,
v_head_size,
k_cumdecay_i_reduced,
qk_head_size,
0.0f,
v_prime,
v_head_size);
}
// v_new = v_prime = v_i - v_prime
// v_i: [chunk_size, EV]
@@ -663,41 +748,75 @@ void chunk_gated_delta_rule_kernel_impl(
}
// attn_inter = qg @ curr_last_recurrent_state: [chunk_size, EV]
// curr_last_recurrent_state: [EK, EV]
at::native::cpublas::brgemm(
/* M */ chunk_size,
/* N */ v_head_size,
/* K */ qk_head_size,
/* lda */ qk_head_size,
/* ldb */ v_head_size,
/* ldc */ v_head_size,
/* add_C */ false,
/* A */ qg,
/* B */ curr_last_recurrent_state_pack_reduced,
/* C */ attn_inter);
if constexpr (brgemm_supported()) {
at::native::cpublas::brgemm(
/* M */ chunk_size,
/* N */ v_head_size,
/* K */ qk_head_size,
/* lda */ qk_head_size,
/* ldb */ v_head_size,
/* ldc */ v_head_size,
/* add_C */ false,
/* A */ qg,
/* B */ curr_last_recurrent_state_pack_reduced,
/* C */ attn_inter);
} else {
blas_gemm(
at::native::TransposeType::NoTranspose,
at::native::TransposeType::NoTranspose,
v_head_size,
chunk_size,
qk_head_size,
1.0f,
curr_last_recurrent_state_reduced,
v_head_size,
qg,
qk_head_size,
0.0f,
attn_inter,
v_head_size);
}
// core_attn_out[:, :, i] = attn_inter + attn_i @ v_new
// pack for v_prime
pack_vnni2<scalar_t>(
/* dst */ v_prime_pack_reduced,
/* src */ v_prime_reduced,
/* N */ chunk_size,
/* K */ v_head_size,
/* ld_src */ v_head_size,
/* ld_dst */ v_head_size);
// attn_inter = attn_inter + attn_i @ v_new: [chunk_size, EV]
// attn_i: [chunk_size, chunk_size]
// v_new: [chunk_size, EV]
at::native::cpublas::brgemm(
/* M */ chunk_size,
/* N */ v_head_size,
/* K */ chunk_size,
/* lda */ chunk_size,
/* ldb */ v_head_size,
/* ldc */ v_head_size,
/* add_C */ true,
/* A */ attn_i_reduced,
/* B */ v_prime_pack_reduced,
/* C */ attn_inter);
if constexpr (brgemm_supported()) {
pack_vnni2<scalar_t>(
/* dst */ v_prime_pack_reduced,
/* src */ v_prime_reduced,
/* N */ chunk_size,
/* K */ v_head_size,
/* ld_src */ v_head_size,
/* ld_dst */ v_head_size);
// attn_inter = attn_inter + attn_i @ v_new: [chunk_size, EV]
// attn_i: [chunk_size, chunk_size]
// v_new: [chunk_size, EV]
at::native::cpublas::brgemm(
/* M */ chunk_size,
/* N */ v_head_size,
/* K */ chunk_size,
/* lda */ chunk_size,
/* ldb */ v_head_size,
/* ldc */ v_head_size,
/* add_C */ true,
/* A */ attn_i_reduced,
/* B */ v_prime_pack_reduced,
/* C */ attn_inter);
} else {
blas_gemm(
at::native::TransposeType::NoTranspose,
at::native::TransposeType::NoTranspose,
v_head_size,
chunk_size,
chunk_size,
1.0f,
v_prime_reduced,
v_head_size,
attn_i_reduced,
chunk_size,
1.0f,
attn_inter,
v_head_size);
}
// core_attn_out[:, :, i] = attn_inter
for (int64_t m = 0; m < chunk_size; m++) {
@@ -762,17 +881,34 @@ void chunk_gated_delta_rule_kernel_impl(
/* ld_dst */ chunk_size);
// kgv = kg.transpose(-1, -2) @ v_new
// v_new: [chunk_size, EV]
at::native::cpublas::brgemm(
/* M */ qk_head_size,
/* N */ v_head_size,
/* K */ chunk_size,
/* lda */ chunk_size,
/* ldb */ v_head_size,
/* ldc */ v_head_size,
/* add_C */ false,
/* A */ kg_transpose,
/* B */ v_prime_pack_reduced,
/* C */ kgv);
if constexpr (brgemm_supported()) {
at::native::cpublas::brgemm(
/* M */ qk_head_size,
/* N */ v_head_size,
/* K */ chunk_size,
/* lda */ chunk_size,
/* ldb */ v_head_size,
/* ldc */ v_head_size,
/* add_C */ false,
/* A */ kg_transpose,
/* B */ v_prime_pack_reduced,
/* C */ kgv);
} else {
blas_gemm(
at::native::TransposeType::NoTranspose,
at::native::TransposeType::NoTranspose,
v_head_size,
qk_head_size,
chunk_size,
1.0f,
v_prime_reduced,
v_head_size,
kg_transpose,
chunk_size,
0.0f,
kgv,
v_head_size);
}
// last_recurrent_state = 1) + 2)
for (int64_t m = 0; m < qk_head_size; m++) {
at::vec::map2<float>(
@@ -921,7 +1057,8 @@ void fused_sigmoid_gating_delta_rule_update_kernel_impl(
float k_scale = use_qk_l2norm_in_kernel ? qk_scale_buf[k_scale_offset] : 1.0f;
int64_t v_offset = si * v_strideS + bi * v_strideB + ni * v_strideH;
int64_t o_offset = ((bi * seq_len + si) * v_num_heads + ni) * v_head_dim;
float beta_val = 1 / (1 + std::exp(-b_ptr[ni]));
// See: https://github.com/sgl-project/sglang/pull/26634
float beta_val = 1 / (1 + std::exp(-b_ptr[bi * v_num_heads + ni]));
fVec beta_vec = fVec(beta_val);
int64_t dvi = 0;
for (; dvi <= v_head_dim - VecSize; dvi += VecSize) {
+18 -7
View File
@@ -4,9 +4,12 @@
// clang-format off
#pragma once
#include <ATen/native/CPUBlas.h>
#include "common.h"
#include "blas_gemm.h"
#if defined(__AVX512F__) && defined(__AVX512BF16__) && defined(__AMX_BF16__)
#define CPU_CAPABILITY_AVX512
#endif
// amx-bf16
#define TILE_M 16
@@ -21,31 +24,39 @@ constexpr int block_size_n() {
return 2 * TILE_N;
}
constexpr bool brgemm_supported() {
#if defined(CPU_CAPABILITY_AVX512)
return true;
#else
return false;
#endif
}
// define threshold using brgemm (intel AMX)
template <typename T>
inline bool can_use_brgemm(int M);
template <>
inline bool can_use_brgemm<at::BFloat16>(int M) {
return M > 4;
return brgemm_supported() && M > 4;
}
template <>
inline bool can_use_brgemm<at::Half>(int M) {
return true;
return brgemm_supported();
}
// this requires PyTorch 2.7 or above
template <>
inline bool can_use_brgemm<int8_t>(int M) {
return M > 4;
return brgemm_supported() && M > 4;
}
template <>
inline bool can_use_brgemm<uint8_t>(int M) {
return M > 4;
return brgemm_supported() && M > 4;
}
template <>
inline bool can_use_brgemm<at::Float8_e4m3fn>(int M) {
return M > 4;
return brgemm_supported() && M > 4;
}
// work around compiler internal error
+2
View File
@@ -11,7 +11,9 @@
#include <ATen/cpu/vec/functional.h>
#include <ATen/cpu/vec/vec.h>
#if defined(CPU_CAPABILITY_AVX512)
#include <immintrin.h>
#endif
namespace {
using namespace at::vec;
+10 -8
View File
@@ -5,7 +5,7 @@
#include <sys/stat.h>
#include <unistd.h>
#ifdef __aarch64__
#if defined(__aarch64__) || defined(__powerpc64__)
#include <atomic>
#endif
@@ -38,7 +38,7 @@ struct KernelVecType<c10::Half> {
};
struct ThreadSHMContext {
#ifdef __aarch64__
#if defined(__aarch64__) || defined(__powerpc64__)
// memory model is weaker on AArch64, so we use atomic variables for
// consumer (load-acquire) and producer (store-release) to make sure
// that a stamp cannot be ready before the corresponding data is ready.
@@ -75,7 +75,7 @@ struct ThreadSHMContext {
TORCH_CHECK(group_size <= MAX_SHM_RANK_NUM);
TORCH_CHECK((size_t)this % 64 == 0);
TORCH_CHECK((size_t)thread_shm_ptr % 64 == 0);
#ifdef __aarch64__
#if defined(__aarch64__) || defined(__powerpc64__)
_curr_thread_stamp[0].store(1, std::memory_order_relaxed);
_curr_thread_stamp[1].store(1, std::memory_order_relaxed);
_ready_thread_stamp[0].store(0, std::memory_order_relaxed);
@@ -124,7 +124,7 @@ struct ThreadSHMContext {
}
char get_curr_stamp(int idx) const {
#ifdef __aarch64__
#if defined(__aarch64__) || defined(__powerpc64__)
return _curr_thread_stamp[idx].load(std::memory_order_acquire);
#else
return _curr_thread_stamp[idx];
@@ -132,7 +132,7 @@ struct ThreadSHMContext {
}
char get_ready_stamp(int idx) const {
#ifdef __aarch64__
#if defined(__aarch64__) || defined(__powerpc64__)
return _ready_thread_stamp[idx].load(std::memory_order_acquire);
#else
return _ready_thread_stamp[idx];
@@ -140,7 +140,7 @@ struct ThreadSHMContext {
}
void next_stamp() {
#ifdef __aarch64__
#if defined(__aarch64__) || defined(__powerpc64__)
_curr_thread_stamp[local_stamp_buffer_idx].fetch_add(
1, std::memory_order_release);
#else
@@ -150,7 +150,7 @@ struct ThreadSHMContext {
}
void commit_ready_stamp() {
#ifdef __aarch64__
#if defined(__aarch64__) || defined(__powerpc64__)
_ready_thread_stamp[local_stamp_buffer_idx].store(
_curr_thread_stamp[local_stamp_buffer_idx].load(
std::memory_order_relaxed),
@@ -186,8 +186,10 @@ struct ThreadSHMContext {
break;
}
++_spinning_count;
#ifdef __aarch64__
#if defined(__aarch64__)
__asm__ __volatile__("yield");
#elif defined(__powerpc64__)
__asm__ __volatile__("or 1,1,1");
#else
_mm_pause();
#endif // __aarch64__
+22 -21
View File
@@ -378,7 +378,8 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
#endif
// SHM CCL
#if defined(__AVX512F__) || (defined(__aarch64__) && !defined(__APPLE__))
#if defined(__AVX512F__) || (defined(__aarch64__) && !defined(__APPLE__)) || \
defined(__powerpc64__)
ops.def(
"init_shm_manager(str name, int group_size, int rank, int thread_num) -> "
"int",
@@ -447,6 +448,25 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
"bool is_vnni) -> Tensor");
ops.impl("fp8_scaled_mm_cpu", torch::kCPU, &fp8_scaled_mm_cpu);
// Adapted from sglang: casual_conv1d kernels
ops.def("causal_conv1d_weight_pack(Tensor weight) -> Tensor");
ops.impl("causal_conv1d_weight_pack", torch::kCPU,
&causal_conv1d_weight_pack);
ops.def(
"causal_conv1d_fwd_cpu(Tensor x, Tensor weight, Tensor? bias, Tensor? "
"conv_states, Tensor? query_start_loc,"
"Tensor? cache_indices, Tensor? has_initial_state, bool silu_activation, "
"int pad_slot_id, bool is_vnni) -> "
"Tensor");
ops.impl("causal_conv1d_fwd_cpu", torch::kCPU, &causal_conv1d_fwd_cpu);
ops.def(
"causal_conv1d_update_cpu(Tensor x, Tensor(a!) conv_states, Tensor "
"weight, Tensor? bias, bool silu_activation,"
"Tensor? cache_seqlens, Tensor? conv_state_indices, int pad_slot_id, "
"bool is_vnni) -> Tensor");
ops.impl("causal_conv1d_update_cpu", torch::kCPU, &causal_conv1d_update_cpu);
#endif
// Adapted from sglang: GDN kernels
ops.def(
"chunk_gated_delta_rule_cpu(Tensor query, Tensor key, Tensor value, "
@@ -470,25 +490,6 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
"-> (Tensor, Tensor)");
ops.impl("fused_gdn_gating_cpu", torch::kCPU, &fused_gdn_gating_cpu);
// Adapted from sglang: casual_conv1d kernels
ops.def("causal_conv1d_weight_pack(Tensor weight) -> Tensor");
ops.impl("causal_conv1d_weight_pack", torch::kCPU,
&causal_conv1d_weight_pack);
ops.def(
"causal_conv1d_fwd_cpu(Tensor x, Tensor weight, Tensor? bias, Tensor? "
"conv_states, Tensor? query_start_loc,"
"Tensor? cache_indices, Tensor? has_initial_state, bool silu_activation, "
"int pad_slot_id, bool is_vnni) -> "
"Tensor");
ops.impl("causal_conv1d_fwd_cpu", torch::kCPU, &causal_conv1d_fwd_cpu);
ops.def(
"causal_conv1d_update_cpu(Tensor x, Tensor(a!) conv_states, Tensor "
"weight, Tensor? bias, bool silu_activation,"
"Tensor? cache_seqlens, Tensor? conv_state_indices, int pad_slot_id, "
"bool is_vnni) -> Tensor");
ops.impl("causal_conv1d_update_cpu", torch::kCPU, &causal_conv1d_update_cpu);
#endif
// CPU attention kernels
ops.def(
"get_scheduler_metadata(int num_req, int num_heads_q, int num_heads_kv, "
@@ -518,7 +519,7 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
ops.def("dynamic_per_token_scaled_fp8_quant() -> ()", placeholder_op);
// WNA16
#if defined(__AVX512F__)
#if defined(__AVX512F__) || defined(__riscv_v)
ops.def(
"cpu_gemm_wna16(Tensor input, Tensor q_weight, Tensor(a2!) output, "
"Tensor scales, Tensor? zeros, Tensor? g_idx, Tensor? bias, SymInt "
+41 -75
View File
@@ -10,20 +10,11 @@
namespace vllm {
// `alpha` and `beta` are applied to opposite operands:
// - alpha lives INSIDE the activation (the activated half): the gated
// activation computes act_half * sigmoid(alpha * act_half).
// - beta is added to the OTHER (non-activated) half before the multiply.
// So the result is always ACT(act_half, alpha) * (other_half + beta).
// Which half is which depends on `act_first` (see below). Defaults
// alpha=1.0, beta=0.0 reproduce the plain SwiGLU/GeGLU behavior.
template <typename scalar_t, scalar_t (*ACT_FN)(const scalar_t&, const float),
template <typename scalar_t, scalar_t (*ACT_FN)(const scalar_t&),
bool act_first, bool HAS_CLAMP>
__device__ __forceinline__ scalar_t compute(const scalar_t& x,
const scalar_t& y,
const float limit,
const float alpha,
const float beta) {
const float limit) {
if constexpr (act_first) {
scalar_t gate = x;
scalar_t up = y;
@@ -31,9 +22,7 @@ __device__ __forceinline__ scalar_t compute(const scalar_t& x,
gate = (scalar_t)fminf((float)gate, limit);
up = (scalar_t)fmaxf(fminf((float)up, limit), -limit);
}
// act_first: gate is the activated half -> alpha applies to gate;
// beta is added to up (the non-activated half).
return ACT_FN(gate, alpha) * (scalar_t)((float)up + beta);
return ACT_FN(gate) * up;
} else {
scalar_t gate = x;
scalar_t up = y;
@@ -41,66 +30,55 @@ __device__ __forceinline__ scalar_t compute(const scalar_t& x,
gate = (scalar_t)fmaxf(fminf((float)gate, limit), -limit);
up = (scalar_t)fminf((float)up, limit);
}
// !act_first: up is the activated half -> alpha applies to up;
// beta is added to gate (the non-activated half).
return (scalar_t)((float)gate + beta) * ACT_FN(up, alpha);
return gate * ACT_FN(up);
}
}
template <typename packed_t,
packed_t (*PACKED_ACT_FN)(const packed_t&, const float),
template <typename packed_t, packed_t (*PACKED_ACT_FN)(const packed_t&),
bool act_first, bool HAS_CLAMP>
__device__ __forceinline__ packed_t packed_compute(const packed_t& x,
const packed_t& y,
const float limit,
const float alpha,
const float beta) {
const float limit) {
if constexpr (act_first) {
packed_t gate = x;
packed_t up = y;
float2 u = cast_to_float2(up);
if constexpr (HAS_CLAMP) {
float2 g = cast_to_float2(gate);
float2 u = cast_to_float2(up);
g.x = fminf(g.x, limit);
g.y = fminf(g.y, limit);
u.x = fmaxf(fminf(u.x, limit), -limit);
u.y = fmaxf(fminf(u.y, limit), -limit);
gate = cast_to_packed<packed_t>(g);
up = cast_to_packed<packed_t>(u);
}
// act_first: gate is the activated half -> alpha applies to gate;
// beta is added to up (the non-activated half).
u.x += beta;
u.y += beta;
return packed_mul(PACKED_ACT_FN(gate, alpha), cast_to_packed<packed_t>(u));
return packed_mul(PACKED_ACT_FN(gate), up);
} else {
packed_t gate = x;
packed_t up = y;
float2 g = cast_to_float2(gate);
if constexpr (HAS_CLAMP) {
float2 g = cast_to_float2(gate);
float2 u = cast_to_float2(up);
g.x = fmaxf(fminf(g.x, limit), -limit);
g.y = fmaxf(fminf(g.y, limit), -limit);
u.x = fminf(u.x, limit);
u.y = fminf(u.y, limit);
gate = cast_to_packed<packed_t>(g);
up = cast_to_packed<packed_t>(u);
}
// !act_first: up is the activated half -> alpha applies to up;
// beta is added to gate (the non-activated half).
g.x += beta;
g.y += beta;
return packed_mul(cast_to_packed<packed_t>(g), PACKED_ACT_FN(up, alpha));
return packed_mul(gate, PACKED_ACT_FN(up));
}
}
// Activation and gating kernel template.
template <typename scalar_t, typename packed_t,
scalar_t (*ACT_FN)(const scalar_t&, const float),
packed_t (*PACKED_ACT_FN)(const packed_t&, const float),
bool act_first, bool use_vec, bool HAS_CLAMP, bool use_256b = false>
scalar_t (*ACT_FN)(const scalar_t&),
packed_t (*PACKED_ACT_FN)(const packed_t&), bool act_first,
bool use_vec, bool HAS_CLAMP, bool use_256b = false>
__global__ void act_and_mul_kernel(
scalar_t* __restrict__ out, // [..., d]
const scalar_t* __restrict__ input, // [..., 2, d]
const int d, const float limit, const float alpha, const float beta) {
const int d, const float limit) {
const scalar_t* x_ptr = input + blockIdx.x * 2 * d;
const scalar_t* y_ptr = x_ptr + d;
scalar_t* out_ptr = out + blockIdx.x * d;
@@ -127,7 +105,7 @@ __global__ void act_and_mul_kernel(
for (int j = 0; j < pvec_t::NUM_ELTS; j++) {
x.elts[j] =
packed_compute<packed_t, PACKED_ACT_FN, act_first, HAS_CLAMP>(
x.elts[j], y.elts[j], limit, alpha, beta);
x.elts[j], y.elts[j], limit);
}
if constexpr (use_256b) {
st256(x, &out_vec[i]);
@@ -140,34 +118,29 @@ __global__ void act_and_mul_kernel(
for (int64_t idx = threadIdx.x; idx < d; idx += blockDim.x) {
const scalar_t x = VLLM_LDG(&x_ptr[idx]);
const scalar_t y = VLLM_LDG(&y_ptr[idx]);
out_ptr[idx] = compute<scalar_t, ACT_FN, act_first, HAS_CLAMP>(
x, y, limit, alpha, beta);
out_ptr[idx] =
compute<scalar_t, ACT_FN, act_first, HAS_CLAMP>(x, y, limit);
}
}
}
// Gated activations take an `alpha` argument that scales the sigmoid input
// (`x * sigmoid(alpha * x)`). alpha defaults to 1.0 at all call sites, which
// is exactly SiLU; only the clamp path (silu_and_mul_with_clamp) passes a
// non-default alpha. Activations that do not use alpha simply ignore it.
template <typename T>
__device__ __forceinline__ T silu_kernel(const T& x, const float alpha) {
// x * sigmoid(alpha * x)
return (T)(((float)x) / (1.0f + expf((float)-x * alpha)));
__device__ __forceinline__ T silu_kernel(const T& x) {
// x * sigmoid(x)
return (T)(((float)x) / (1.0f + expf((float)-x)));
}
template <typename packed_t>
__device__ __forceinline__ packed_t packed_silu_kernel(const packed_t& val,
const float alpha) {
// x * sigmoid(alpha * x)
__device__ __forceinline__ packed_t packed_silu_kernel(const packed_t& val) {
// x * sigmoid(x)
float2 fval = cast_to_float2(val);
fval.x = fval.x / (1.0f + expf(-fval.x * alpha));
fval.y = fval.y / (1.0f + expf(-fval.y * alpha));
fval.x = fval.x / (1.0f + expf(-fval.x));
fval.y = fval.y / (1.0f + expf(-fval.y));
return cast_to_packed<packed_t>(fval);
}
template <typename T>
__device__ __forceinline__ T gelu_kernel(const T& x, const float /*alpha*/) {
__device__ __forceinline__ T gelu_kernel(const T& x) {
// Equivalent to PyTorch GELU with 'none' approximation.
// Refer to:
// https://github.com/pytorch/pytorch/blob/8ac9b20d4b090c213799e81acf48a55ea8d437d6/aten/src/ATen/native/cuda/ActivationGeluKernel.cu#L36-L38
@@ -177,8 +150,7 @@ __device__ __forceinline__ T gelu_kernel(const T& x, const float /*alpha*/) {
}
template <typename packed_t>
__device__ __forceinline__ packed_t packed_gelu_kernel(const packed_t& val,
const float /*alpha*/) {
__device__ __forceinline__ packed_t packed_gelu_kernel(const packed_t& val) {
// Equivalent to PyTorch GELU with 'none' approximation.
// Refer to:
// https://github.com/pytorch/pytorch/blob/8ac9b20d4b090c213799e81acf48a55ea8d437d6/aten/src/ATen/native/cuda/ActivationGeluKernel.cu#L36-L38
@@ -190,8 +162,7 @@ __device__ __forceinline__ packed_t packed_gelu_kernel(const packed_t& val,
}
template <typename T>
__device__ __forceinline__ T gelu_tanh_kernel(const T& x,
const float /*alpha*/) {
__device__ __forceinline__ T gelu_tanh_kernel(const T& x) {
// Equivalent to PyTorch GELU with 'tanh' approximation.
// Refer to:
// https://github.com/pytorch/pytorch/blob/8ac9b20d4b090c213799e81acf48a55ea8d437d6/aten/src/ATen/native/cuda/ActivationGeluKernel.cu#L25-L30
@@ -205,7 +176,7 @@ __device__ __forceinline__ T gelu_tanh_kernel(const T& x,
template <typename packed_t>
__device__ __forceinline__ packed_t
packed_gelu_tanh_kernel(const packed_t& val, const float /*alpha*/) {
packed_gelu_tanh_kernel(const packed_t& val) {
// Equivalent to PyTorch GELU with 'tanh' approximation.
// Refer to:
// https://github.com/pytorch/pytorch/blob/8ac9b20d4b090c213799e81acf48a55ea8d437d6/aten/src/ATen/native/cuda/ActivationGeluKernel.cu#L25-L30
@@ -231,7 +202,7 @@ packed_gelu_tanh_kernel(const packed_t& val, const float /*alpha*/) {
// clamped (max only) and up input is clamped (both sides) before the
// activation function is applied.
#define LAUNCH_ACTIVATION_GATE_KERNEL(KERNEL, PACKED_KERNEL, ACT_FIRST, \
HAS_CLAMP, LIMIT, ALPHA, BETA) \
HAS_CLAMP, LIMIT) \
auto dtype = input.scalar_type(); \
int d = input.size(-1) / 2; \
int64_t num_tokens = input.numel() / input.size(-1); \
@@ -259,7 +230,7 @@ packed_gelu_tanh_kernel(const packed_t& val, const float /*alpha*/) {
PACKED_KERNEL<typename vllm::PackedTypeConverter<scalar_t>::Type>, \
ACT_FIRST, true, HAS_CLAMP, true><<<grid, block, 0, stream>>>( \
out.mutable_data_ptr<scalar_t>(), \
input.const_data_ptr<scalar_t>(), d, LIMIT, ALPHA, BETA); \
input.const_data_ptr<scalar_t>(), d, LIMIT); \
}); \
} else { \
VLLM_STABLE_DISPATCH_FLOATING_TYPES(dtype, "act_and_mul_kernel", [&] { \
@@ -269,7 +240,7 @@ packed_gelu_tanh_kernel(const packed_t& val, const float /*alpha*/) {
PACKED_KERNEL<typename vllm::PackedTypeConverter<scalar_t>::Type>, \
ACT_FIRST, true, HAS_CLAMP, false><<<grid, block, 0, stream>>>( \
out.mutable_data_ptr<scalar_t>(), \
input.const_data_ptr<scalar_t>(), d, LIMIT, ALPHA, BETA); \
input.const_data_ptr<scalar_t>(), d, LIMIT); \
}); \
} \
} else { \
@@ -281,7 +252,7 @@ packed_gelu_tanh_kernel(const packed_t& val, const float /*alpha*/) {
PACKED_KERNEL<typename vllm::PackedTypeConverter<scalar_t>::Type>, \
ACT_FIRST, false, HAS_CLAMP><<<grid, block, 0, stream>>>( \
out.mutable_data_ptr<scalar_t>(), input.const_data_ptr<scalar_t>(), \
d, LIMIT, ALPHA, BETA); \
d, LIMIT); \
}); \
}
@@ -289,18 +260,14 @@ void silu_and_mul(torch::stable::Tensor& out, // [..., d]
torch::stable::Tensor& input) // [..., 2 * d]
{
LAUNCH_ACTIVATION_GATE_KERNEL(vllm::silu_kernel, vllm::packed_silu_kernel,
true, false, 0.0f, 1.0f, 0.0f);
true, false, 0.0f);
}
void silu_and_mul_clamp(torch::stable::Tensor& out, // [..., d]
torch::stable::Tensor& input, // [..., 2 * d]
double limit, double alpha, double beta) {
// out = (gate.clamp(max=limit) * sigmoid(alpha * gate.clamp(max=limit)))
// * (up.clamp(+-limit) + beta)
// alpha=1.0, beta=0.0 reduce this to silu(gate) * up.
double limit) {
LAUNCH_ACTIVATION_GATE_KERNEL(vllm::silu_kernel, vllm::packed_silu_kernel,
true, true, (float)limit, (float)alpha,
(float)beta);
true, true, (float)limit);
}
void mul_and_silu(torch::stable::Tensor& out, // [..., d]
@@ -309,22 +276,21 @@ void mul_and_silu(torch::stable::Tensor& out, // [..., d]
// The difference between mul_and_silu and silu_and_mul is that mul_and_silu
// applies the silu to the latter half of the input.
LAUNCH_ACTIVATION_GATE_KERNEL(vllm::silu_kernel, vllm::packed_silu_kernel,
false, false, 0.0f, 1.0f, 0.0f);
false, false, 0.0f);
}
void gelu_and_mul(torch::stable::Tensor& out, // [..., d]
torch::stable::Tensor& input) // [..., 2 * d]
{
LAUNCH_ACTIVATION_GATE_KERNEL(vllm::gelu_kernel, vllm::packed_gelu_kernel,
true, false, 0.0f, 1.0f, 0.0f);
true, false, 0.0f);
}
void gelu_tanh_and_mul(torch::stable::Tensor& out, // [..., d]
torch::stable::Tensor& input) // [..., 2 * d]
{
LAUNCH_ACTIVATION_GATE_KERNEL(vllm::gelu_tanh_kernel,
vllm::packed_gelu_tanh_kernel, true, false,
0.0f, 1.0f, 0.0f);
LAUNCH_ACTIVATION_GATE_KERNEL(
vllm::gelu_tanh_kernel, vllm::packed_gelu_tanh_kernel, true, false, 0.0f);
}
namespace vllm {
@@ -2,7 +2,7 @@
#include <torch/csrc/stable/tensor.h>
#include "cutlass_extensions/epilogue/broadcast_load_epilogue_c2x.hpp"
#include "broadcast_load_epilogue_c2x.hpp"
/*
This file defines custom epilogues for fusing channel scales, token scales,
@@ -1,615 +0,0 @@
/*
* SPDX-License-Identifier: Apache-2.0
* SPDX-FileCopyrightText: Copyright contributors to the vLLM project
*
* Horizontally-fused MiniMax-M3 attention pre-processing kernel.
*
* Replaces the per-token Python sequence in
* ``MiniMaxM3SparseAttention.forward`` / ``MiniMaxM3Attention.forward``:
*
* q = q_norm(q); k = k_norm(k); q, k = rotary_emb(pos, q, k)
* index_q = index_q_norm(index_q); index_k = index_k_norm(index_k)
* index_q, index_k = rotary_emb(pos, index_q, index_k)
* _insert_kv(k, v, index_k)
*
* All branches share head_dim=128 and the *same* partial-NeoX RoPE table
* (``rotary_dim`` rotated, the trailing dims pass through). The four norms
* are Gemma-style RMSNorm (``x * rsqrt(mean(x^2)+eps) * (1 + weight)``) with
* independent weights.
*
* Everything lives in a single fused ``qkv`` tensor. The sparse layer's
* fused projection (MinimaxM3QKVParallelLinearWithIndexer) emits, per token::
*
* [ q | k | v | index_q | index_k ] (the "5 results")
*
* while the dense layer emits just ``[ q | k | v ]``. The kernel reads the
* index branch straight out of that packed row -- no separate index tensors.
*
* One kernel, one grid; each warp owns one (token, head-slot) pair. Slot
* enumeration per token:
* [0, nq) Q heads -> norm(q_w) + RoPE, write
* qkv [nq, nq+nkv) K heads -> norm(k_w) + RoPE, write
* qkv
* (+ insert into key cache)
* [nq+nkv, nq+2*nkv) V heads -> insert into value cache
* IQ heads (niq) -> norm(iq_w) + RoPE, write iq
* IK (1) -> norm(ik_w) + RoPE
* (+ insert into index cache)
*
* The IQ/IK warps address the index_q/index_k sub-blocks *inside* qkv at the
* fixed physical offsets (nq+2*nkv)*128 and (nq+2*nkv+niq)*128.
*
* Dense vs sparse is a compile-time choice via the ``kIsSparse``/``kInsertKV``
* template bools (3 instantiations: dense <false,false>, sparse-profiling
* <true,false>, sparse-serving <true,true>), so the index slots, the V slots
* and the cache inserts fold away entirely on paths that don't use them. The
* dense layer passes no caches/index: norm+RoPE happens in place and the
* generic ``Attention`` layer owns the cache write.
*
* Q/K and (sparse) index_q/index_k are all rewritten in place inside the fused
* ``qkv`` tensor. Caches (bf16) are scatter-written by slot.
*/
#include <cmath>
#include <cuda_runtime.h>
#include <type_traits>
#include "torch_utils.h"
#include "../cuda_compat.h"
#include "../type_convert.cuh"
#include "dispatch_utils.h"
#ifndef FINAL_MASK
#ifdef USE_ROCM
#define FINAL_MASK 0xffffffffffffffffULL
#else
#define FINAL_MASK 0xffffffffu
#endif
#endif
namespace vllm {
namespace minimax_m3_fused_ops {
namespace {
inline int getSMVersion() {
auto* props = get_device_prop();
return props->major * 10 + props->minor;
}
} // namespace
// ────────────────────────────────────────────────────────────────────────────
// Constants (hard-coded for MiniMax-M3-preview).
// ────────────────────────────────────────────────────────────────────────────
constexpr int kHeadDim = 128;
constexpr int kNumLanes = 32;
constexpr int kElemsPerLane = kHeadDim / kNumLanes; // 4
// ────────────────────────────────────────────────────────────────────────────
// Helpers
// ────────────────────────────────────────────────────────────────────────────
__device__ __forceinline__ float warpReduceSum(float val) {
#pragma unroll
for (int mask = 16; mask > 0; mask >>= 1) {
val += __shfl_xor_sync(FINAL_MASK, val, mask, 32);
}
return val;
}
// Gemma RMSNorm over the full head (no-op when ``weight == nullptr``) followed
// by partial NeoX RoPE on the leading ``rotary_dim`` dims, all in fp32. Each
// lane owns ``kElemsPerLane`` contiguous dims [laneId*4, laneId*4+4).
template <typename scalar_t>
__device__ __forceinline__ void normAndRope(
float (&elems)[kElemsPerLane], int const laneId, float const eps,
scalar_t const* __restrict__ weight, // [kHeadDim] or nullptr (no norm)
bool const do_rope, int const rotary_dim,
scalar_t const* __restrict__ cos_ptr, // cos_sin_cache + pos*rotary_dim
bool const apply_norm) {
// ── Gemma RMSNorm: x * rsqrt(mean(x^2)+eps) * (1 + w) ──────────────────
if (apply_norm) {
float sumsq = 0.0f;
#pragma unroll
for (int i = 0; i < kElemsPerLane; i++) sumsq += elems[i] * elems[i];
sumsq = warpReduceSum(sumsq);
float const rms_rcp = rsqrtf(sumsq / static_cast<float>(kHeadDim) + eps);
#pragma unroll
for (int i = 0; i < kElemsPerLane; i++) {
int const dim = laneId * kElemsPerLane + i;
float const w = 1.0f + static_cast<float>(weight[dim]);
elems[i] = elems[i] * rms_rcp * w;
}
}
// ── Partial NeoX RoPE on dims [0, rotary_dim) ──────────────────────────
// half = rotary_dim/2. Pair (i, i+half) for i in [0, half). Lane L owns
// dims [4L, 4L+4); since half is a multiple of 4, a lane lies wholly in the
// first half (own=x[i]) or second half (own=x[i+half]); its partner lives
// ``half/4`` lanes away (XOR with that distance).
if (do_rope) {
int const half = rotary_dim / 2;
int const dim0 = laneId * kElemsPerLane;
bool const in_rope = dim0 < rotary_dim;
int const lane_xor = half / kElemsPerLane; // partner-lane distance
float partner[kElemsPerLane];
#pragma unroll
for (int i = 0; i < kElemsPerLane; i++) {
partner[i] = __shfl_xor_sync(FINAL_MASK, elems[i], lane_xor, 32);
}
if (in_rope) {
bool const first_half = dim0 < half;
int const i_base = first_half ? dim0 : (dim0 - half); // cos/sin index
scalar_t const* sin_ptr = cos_ptr + half;
#pragma unroll
for (int i = 0; i < kElemsPerLane; i++) {
float const c = static_cast<float>(cos_ptr[i_base + i]);
float const s = static_cast<float>(sin_ptr[i_base + i]);
if (first_half) {
elems[i] = elems[i] * c - partner[i] * s;
} else {
elems[i] = elems[i] * c + partner[i] * s;
}
}
}
}
}
// Load 4 contiguous bf16 -> 4 fp32 registers.
template <typename scalar_t>
__device__ __forceinline__ void loadElems(scalar_t const* __restrict__ src,
float (&elems)[kElemsPerLane]) {
using Converter = vllm::_typeConvert<scalar_t>;
uint2 v = *reinterpret_cast<uint2 const*>(src);
auto const* p =
reinterpret_cast<typename Converter::packed_hip_type const*>(&v);
#pragma unroll
for (int i = 0; i < kElemsPerLane / 2; i++) {
float2 f2 = Converter::convert(p[i]);
elems[2 * i] = f2.x;
elems[2 * i + 1] = f2.y;
}
}
// Store 4 fp32 registers -> 4 contiguous bf16.
template <typename scalar_t>
__device__ __forceinline__ void storeElems(
scalar_t* __restrict__ dst, float const (&elems)[kElemsPerLane]) {
using Converter = vllm::_typeConvert<scalar_t>;
uint2 v;
auto* p = reinterpret_cast<typename Converter::packed_hip_type*>(&v);
#pragma unroll
for (int i = 0; i < kElemsPerLane / 2; i++) {
p[i] = Converter::convert(make_float2(elems[2 * i], elems[2 * i + 1]));
}
*reinterpret_cast<uint2*>(dst) = v;
}
// ────────────────────────────────────────────────────────────────────────────
// Kernel
// ────────────────────────────────────────────────────────────────────────────
// Grid: 1D, ceil(num_tokens * slots_per_token / warps_per_block).
// Each warp = one (token, slot).
//
// `kIsSparse` and `kInsertKV` are compile-time template bools, so all the
// branch decisions that distinguish the dense layer from the sparse layer
// (index slots, KV/index inserts, V slots) fold away per instantiation.
// Three instantiations are built: dense <false,false>, sparse-profiling
// <true,false> and sparse-serving <true,true>. Slots per token:
// Q : nq (always — norm+RoPE)
// K : nkv (always — norm+RoPE; +K-cache insert)
// V : nkv only if kInsertKV (V-cache insert; no warps in dense)
// IQ: niq only if kIsSparse (norm+RoPE)
// IK: 1 only if kIsSparse (norm+RoPE; +index-cache insert)
template <typename scalar_t, bool kIsSparse, bool kInsertKV>
__global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
scalar_t* __restrict__ qkv, // [N, qkv_row] in/out (packs index if sparse)
scalar_t* __restrict__ q_out, // [N, nq*128] contiguous, or nullptr
scalar_t* __restrict__ index_q_out, // [N, niq*128] contiguous, or nullptr
scalar_t const* __restrict__ q_norm_w,
scalar_t const* __restrict__ k_norm_w,
scalar_t const* __restrict__ iq_norm_w,
scalar_t const* __restrict__ ik_norm_w,
scalar_t const* __restrict__ cos_sin_cache, // [max_pos, rotary_dim]
int64_t const* __restrict__ positions, // [N] i64
int64_t const* __restrict__ slot_mapping, // [N] i64 or nullptr
scalar_t* __restrict__ kv_cache, // [nb,2,bs,nkv,128] or nullptr
scalar_t* __restrict__ index_cache, // [nb*bs, 128] or nullptr
float const eps, int const rotary_dim, int const num_tokens, int const nq,
int const nkv, int const niq, int const block_size,
// kv_cache strides (in elements) for logical shape [nb, 2, bs, nkv, 128].
// The head_dim (last) dim is always innermost-contiguous (stride 1), so the
// NHD/HND layout choice is fully captured by these four strides: NHD keeps
// s_token < s_head, HND swaps them. dim_base addresses head_dim directly.
int64_t const kv_s_block, int64_t const kv_s_kv, int64_t const kv_s_token,
int64_t const kv_s_head) {
#if (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ < 800) && !defined(USE_ROCM)
// _typeConvert<BFloat16> is unavailable on pre-Ampere; the M3 kernel only
// runs with bf16/fp16 inputs in practice. Discard the bf16 body there.
if constexpr (std::is_same_v<scalar_t, c10::BFloat16>) {
return;
} else {
#endif
int const warpsPerBlock = blockDim.x / 32;
int const laneId = threadIdx.x % 32;
int const globalWarpIdx = blockIdx.x * warpsPerBlock + (threadIdx.x / 32);
// Slot layout (compile-time gated: dense has neither V nor index slots).
int const v_slots = kInsertKV ? nkv : 0;
int const idx_slots = kIsSparse ? niq + 1 : 0;
int const slots_per_token = nq + nkv + v_slots + idx_slots;
int const tokenIdx = globalWarpIdx / slots_per_token;
int const slot = globalWarpIdx % slots_per_token;
if (tokenIdx >= num_tokens) return;
// Slot boundaries.
int const k_begin = nq;
int const v_begin = nq + nkv; // valid only when kInsertKV
int const iq_begin = nq + nkv + v_slots; // index block start
int const ik_slot = iq_begin + niq; // valid only when kIsSparse
bool const isQ = slot < k_begin;
bool const isK = slot >= k_begin && slot < v_begin;
bool isV = false;
if constexpr (kInsertKV) isV = slot >= v_begin && slot < v_begin + nkv;
bool isIQ = false, isIK = false;
if constexpr (kIsSparse) {
isIQ = slot >= iq_begin && slot < ik_slot;
isIK = slot == ik_slot;
}
int const dim_base = laneId * kElemsPerLane;
// Physical row width of qkv: the dense layer packs [q|k|v]; the sparse
// layer additionally packs [index_q (niq heads) | index_k (1 head)].
int const qkv_row = (nq + 2 * nkv + (kIsSparse ? (niq + 1) : 0)) * kHeadDim;
// ── Resolve source pointer + per-branch parameters. ────────────────────
scalar_t* row_ptr = nullptr; // in-place output location
scalar_t const* norm_w = nullptr; // nullptr -> skip norm (V)
bool do_rope = true;
int head = 0; // kv head index for inserts
if (isQ) {
row_ptr =
qkv + static_cast<int64_t>(tokenIdx) * qkv_row + slot * kHeadDim;
norm_w = q_norm_w;
} else if (isK) {
head = slot - k_begin;
row_ptr =
qkv + static_cast<int64_t>(tokenIdx) * qkv_row + slot * kHeadDim;
norm_w = k_norm_w;
} else if (isV) {
// qkv V section starts at slot index (nq + nkv): slot * kHeadDim is the
// correct in-tensor offset.
head = slot - v_begin;
row_ptr =
qkv + static_cast<int64_t>(tokenIdx) * qkv_row + slot * kHeadDim;
norm_w = nullptr; // V: no norm, no rope
do_rope = false;
} else if (isIQ) {
// index_q sub-block lives at physical offset (nq+2*nkv)*128 in qkv.
int const ih = slot - iq_begin;
row_ptr = qkv + static_cast<int64_t>(tokenIdx) * qkv_row +
(nq + 2 * nkv + ih) * kHeadDim;
norm_w = iq_norm_w;
} else { // isIK -- single shared index key at (nq+2*nkv+niq)*128.
row_ptr = qkv + static_cast<int64_t>(tokenIdx) * qkv_row +
(nq + 2 * nkv + niq) * kHeadDim;
norm_w = ik_norm_w;
}
// Store destination. Q and index_q are gathered into dedicated contiguous
// output buffers (when provided) so the downstream SM100 sparse kernel's
// flat TMA descriptor can address them as [tokens*heads, head_dim]; this
// folds the de-interleaving into the store the kernel already does, instead
// of a separate q.contiguous() copy. Everything else stays in place.
scalar_t* store_ptr = row_ptr;
if (isQ && q_out != nullptr) {
store_ptr = q_out + static_cast<int64_t>(tokenIdx) * nq * kHeadDim +
slot * kHeadDim;
} else if (isIQ && index_q_out != nullptr) {
store_ptr = index_q_out +
static_cast<int64_t>(tokenIdx) * niq * kHeadDim +
(slot - iq_begin) * kHeadDim;
}
// PDL: wait for the predecessor kernel (the qkv-projection GEMM that
// produces ``qkv``) to finish before touching any global memory. No-op
// when PDL is not enabled on the launch. The CUDA runtime wrapper emits
// the griddepcontrol.wait PTX with the required memory clobber internally.
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
cudaGridDependencySynchronize();
#endif
// ── Load -> norm+rope (fp32) -> store back in place. ───────────────────
float elems[kElemsPerLane];
loadElems<scalar_t>(row_ptr + dim_base, elems);
if (!isV) {
int64_t const pos = positions[tokenIdx];
scalar_t const* cos_ptr = cos_sin_cache + pos * rotary_dim;
normAndRope<scalar_t>(elems, laneId, eps, norm_w, do_rope, rotary_dim,
cos_ptr, /*apply_norm=*/norm_w != nullptr);
storeElems<scalar_t>(store_ptr + dim_base, elems);
}
// ── Cache inserts (sparse serving only). ───────────────────────────────
if constexpr (kInsertKV) {
// Guard (not early-return) so every thread reaches the PDL trigger below.
int64_t const sm = (isK || isV || isIK) ? slot_mapping[tokenIdx] : -1;
if (sm >= 0) { // skip padded / unscheduled tokens
if (isIK) {
scalar_t* dst = index_cache + sm * kHeadDim + dim_base;
storeElems<scalar_t>(dst, elems);
} else if (isK || isV) {
// kv_cache logical shape [num_blocks, 2, block_size, nkv, head_dim].
// Paging is logical (block = sm/block_size, token = sm%block_size);
// the physical NHD/HND layout is honoured via the passed strides.
int64_t const b = sm / block_size;
int64_t const t = sm % block_size;
int const kv = isK ? 0 : 1;
int64_t const off =
b * kv_s_block + kv * kv_s_kv + t * kv_s_token + head * kv_s_head;
storeElems<scalar_t>(kv_cache + off + dim_base, elems);
}
}
}
// PDL: signal that this kernel is done so a dependent successor may launch
// early. No-op when PDL is not enabled on the launch.
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
cudaTriggerProgrammaticLaunchCompletion();
#endif
#if (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ < 800) && !defined(USE_ROCM)
}
#endif
}
// ────────────────────────────────────────────────────────────────────────────
// Launch wrapper
// ────────────────────────────────────────────────────────────────────────────
template <typename scalar_t>
void launchFusedMiniMaxM3(scalar_t* qkv, scalar_t* q_out, scalar_t* index_q_out,
scalar_t const* q_norm_w, scalar_t const* k_norm_w,
scalar_t const* iq_norm_w, scalar_t const* ik_norm_w,
scalar_t const* cos_sin_cache,
int64_t const* positions, int64_t const* slot_mapping,
scalar_t* kv_cache, scalar_t* index_cache,
float const eps, int const rotary_dim,
int const num_tokens, int const nq, int const nkv,
int const niq, int const block_size,
int64_t const kv_s_block, int64_t const kv_s_kv,
int64_t const kv_s_token, int64_t const kv_s_head,
bool const has_index, bool const insert_kv,
cudaStream_t stream) {
// Slot count must match the kernel's compile-time gating.
int const v_slots = insert_kv ? nkv : 0;
int const idx_slots = has_index ? niq + 1 : 0;
int const slots_per_token = nq + nkv + v_slots + idx_slots;
constexpr int kBlockSize = 256;
constexpr int kWarpsPerBlock = kBlockSize / 32;
int64_t const total_warps =
static_cast<int64_t>(num_tokens) * slots_per_token;
int const grid =
static_cast<int>((total_warps + kWarpsPerBlock - 1) / kWarpsPerBlock);
if (grid == 0) return;
#ifndef USE_ROCM
// PDL: enable programmatic stream serialization whenever the hardware
// supports it (SM90+). On pre-Hopper GPUs the attribute is unavailable, so
// leave numAttrs = 0 and launch as a regular kernel via cudaLaunchKernelEx.
static int const sm_version = getSMVersion();
cudaLaunchConfig_t config;
config.gridDim = dim3(grid);
config.blockDim = dim3(kBlockSize);
config.dynamicSmemBytes = 0;
config.stream = stream;
cudaLaunchAttribute attrs[1];
attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
attrs[0].val.programmaticStreamSerializationAllowed = 1;
config.attrs = attrs;
config.numAttrs = (sm_version >= 90) ? 1 : 0;
#define LAUNCH(IS_SPARSE, INSERT) \
cudaLaunchKernelEx( \
&config, \
fusedMiniMaxM3QNormRopeKVInsertKernel<scalar_t, IS_SPARSE, INSERT>, \
qkv, q_out, index_q_out, q_norm_w, k_norm_w, iq_norm_w, ik_norm_w, \
cos_sin_cache, positions, slot_mapping, kv_cache, index_cache, eps, \
rotary_dim, num_tokens, nq, nkv, niq, block_size, kv_s_block, kv_s_kv, \
kv_s_token, kv_s_head)
#else
// ROCm: standard kernel launch syntax (no PDL/stream serialization).
// clang-format off
#define LAUNCH(IS_SPARSE, INSERT) \
fusedMiniMaxM3QNormRopeKVInsertKernel<scalar_t, IS_SPARSE, INSERT> \
<<<grid, kBlockSize, 0, stream>>>( \
qkv, q_out, index_q_out, q_norm_w, k_norm_w, iq_norm_w, \
ik_norm_w, cos_sin_cache, positions, slot_mapping, kv_cache, \
index_cache, eps, rotary_dim, num_tokens, nq, nkv, niq, \
block_size, kv_s_block, kv_s_kv, kv_s_token, kv_s_head)
// clang-format on
#endif
if (has_index) {
if (insert_kv) {
LAUNCH(true, true); // sparse serving
} else {
LAUNCH(true, false); // sparse profiling
}
} else {
// Dense layer: never has an index branch and never inserts here (the
// generic Attention layer owns the KV insert).
LAUNCH(false, false);
}
#undef LAUNCH
}
} // namespace minimax_m3_fused_ops
} // namespace vllm
// ────────────────────────────────────────────────────────────────────────────
// Torch op wrapper
// ────────────────────────────────────────────────────────────────────────────
void fused_minimax_m3_qknorm_rope_kv_insert(
torch::stable::Tensor& qkv, // [N, qkv_row] (packs index if sparse)
torch::stable::Tensor const& q_norm_weight, // [128]
torch::stable::Tensor const& k_norm_weight, // [128]
torch::stable::Tensor const& cos_sin_cache, // [max_pos, rotary_dim]
torch::stable::Tensor const& positions, // [N] i64
int64_t num_heads, int64_t num_kv_heads, int64_t rotary_dim, double eps,
std::optional<torch::stable::Tensor> index_q_norm_weight, // [128]
std::optional<torch::stable::Tensor> index_k_norm_weight, // [128]
int64_t num_index_heads, // niq; 0 => dense
std::optional<torch::stable::Tensor> slot_mapping, // [N] i64
std::optional<torch::stable::Tensor> kv_cache, // [nb,2,bs,nkv,128]
std::optional<torch::stable::Tensor> index_cache, // [nb,bs,128]
int64_t block_size,
std::optional<torch::stable::Tensor> q_out, // [N, nq*128] contiguous
std::optional<torch::stable::Tensor>
index_q_out) { // [N, niq*128] contiguous
STD_TORCH_CHECK(qkv.is_cuda() && qkv.is_contiguous(),
"qkv must be contiguous CUDA");
STD_TORCH_CHECK(
positions.is_cuda() &&
positions.scalar_type() == torch::headeronly::ScalarType::Long,
"positions must be int64 CUDA");
STD_TORCH_CHECK(cos_sin_cache.is_cuda() && cos_sin_cache.is_contiguous(),
"cos_sin_cache must be contiguous CUDA");
STD_TORCH_CHECK(cos_sin_cache.scalar_type() == qkv.scalar_type(),
"cos_sin_cache dtype must match qkv");
STD_TORCH_CHECK(
cos_sin_cache.dim() == 2 && cos_sin_cache.size(1) == rotary_dim,
"cos_sin_cache shape [max_pos, rotary_dim]");
STD_TORCH_CHECK(q_norm_weight.scalar_type() == qkv.scalar_type() &&
k_norm_weight.scalar_type() == qkv.scalar_type(),
"q/k norm weight dtype must match qkv");
STD_TORCH_CHECK(
q_norm_weight.numel() == vllm::minimax_m3_fused_ops::kHeadDim &&
k_norm_weight.numel() == vllm::minimax_m3_fused_ops::kHeadDim,
"q/k norm weight must have 128 elements");
STD_TORCH_CHECK(rotary_dim > 0 && rotary_dim % 8 == 0 &&
rotary_dim <= vllm::minimax_m3_fused_ops::kHeadDim,
"rotary_dim must be a positive multiple of 8 and <= 128");
int const num_tokens = static_cast<int>(qkv.size(0));
int const nq = static_cast<int>(num_heads);
int const nkv = static_cast<int>(num_kv_heads);
int const niq = static_cast<int>(num_index_heads);
// The sparse layer packs the index branch ([index_q (niq heads) | index_k
// (1 head)]) right after [q|k|v] in the same row; the dense layer does not.
bool const has_index = niq > 0;
bool const insert_kv = kv_cache.has_value();
int const kHeadDim = vllm::minimax_m3_fused_ops::kHeadDim;
int const expected_row =
(nq + 2 * nkv + (has_index ? niq + 1 : 0)) * kHeadDim;
STD_TORCH_CHECK(qkv.size(1) == expected_row,
"qkv last dim must be (num_heads + 2*num_kv_heads"
" + num_index_heads + 1) * 128 for sparse, "
"(num_heads + 2*num_kv_heads) * 128 for dense");
// Only the sparse layer inserts here (dense lets the generic Attention layer
// own the KV write); there is no dense+insert kernel instantiation.
STD_TORCH_CHECK(
!insert_kv || has_index,
"insert mode (kv_cache) requires the index branch (sparse layer)");
if (has_index) {
STD_TORCH_CHECK(
index_q_norm_weight.has_value() && index_k_norm_weight.has_value(),
"index branch requires both index norm weights");
STD_TORCH_CHECK(index_q_norm_weight->scalar_type() == qkv.scalar_type() &&
index_k_norm_weight->scalar_type() == qkv.scalar_type(),
"index norm weights dtype must match qkv");
STD_TORCH_CHECK(index_q_norm_weight->numel() == kHeadDim &&
index_k_norm_weight->numel() == kHeadDim,
"index norm weights must have 128 elements");
}
// kv_cache strides (logical shape [nb, 2, bs, nkv, head_dim]). Read straight
// off the tensor so the kernel honours whatever physical layout the attention
// backend allocated (NHD: stride order (0,1,2,3,4); HND: (0,1,3,2,4)). No new
// op argument is needed -- the strides ride along with the tensor itself.
int64_t kv_s_block = 0, kv_s_kv = 0, kv_s_token = 0, kv_s_head = 0;
if (insert_kv) {
STD_TORCH_CHECK(
slot_mapping.has_value() &&
slot_mapping->scalar_type() == torch::headeronly::ScalarType::Long,
"insert mode requires int64 slot_mapping");
STD_TORCH_CHECK(kv_cache->scalar_type() == qkv.scalar_type(),
"kv_cache dtype must match qkv (bf16 cache only)");
STD_TORCH_CHECK(index_cache.has_value() &&
index_cache->scalar_type() == qkv.scalar_type(),
"insert mode requires matching index_cache");
STD_TORCH_CHECK(kv_cache->dim() == 5 && kv_cache->stride(4) == 1,
"kv_cache must be [nb,2,bs,nkv,head_dim] with contiguous "
"head_dim (stride(4)==1)");
kv_s_block = kv_cache->stride(0);
kv_s_kv = kv_cache->stride(1);
kv_s_token = kv_cache->stride(2);
kv_s_head = kv_cache->stride(3);
}
// Optional contiguous gather targets: when given, the normed/roped q (and
// index_q) are written here instead of in place, so callers avoid a separate
// .contiguous() copy. index_q_out only makes sense on the sparse path.
if (q_out.has_value()) {
STD_TORCH_CHECK(
q_out->is_cuda() && q_out->is_contiguous() &&
q_out->scalar_type() == qkv.scalar_type(),
"q_out must be a contiguous CUDA tensor matching qkv dtype");
STD_TORCH_CHECK(
q_out->numel() == static_cast<int64_t>(num_tokens) * nq * kHeadDim,
"q_out must have num_tokens * num_heads * 128 elements");
}
if (index_q_out.has_value()) {
STD_TORCH_CHECK(
has_index,
"index_q_out requires the index branch (num_index_heads > 0)");
STD_TORCH_CHECK(
index_q_out->is_cuda() && index_q_out->is_contiguous() &&
index_q_out->scalar_type() == qkv.scalar_type(),
"index_q_out must be a contiguous CUDA tensor matching qkv dtype");
STD_TORCH_CHECK(index_q_out->numel() ==
static_cast<int64_t>(num_tokens) * niq * kHeadDim,
"index_q_out must have num_tokens * num_index_heads * 128 "
"elements");
}
const torch::stable::accelerator::DeviceGuard device_guard(
qkv.get_device_index());
auto stream = get_current_cuda_stream(qkv.get_device_index());
VLLM_STABLE_DISPATCH_HALF_TYPES(
qkv.scalar_type(), "fused_minimax_m3_qknorm_rope_kv_insert", [&] {
using st = scalar_t;
vllm::minimax_m3_fused_ops::launchFusedMiniMaxM3<st>(
reinterpret_cast<st*>(qkv.data_ptr()),
q_out.has_value() ? reinterpret_cast<st*>(q_out->data_ptr())
: nullptr,
index_q_out.has_value()
? reinterpret_cast<st*>(index_q_out->data_ptr())
: nullptr,
reinterpret_cast<st const*>(q_norm_weight.data_ptr()),
reinterpret_cast<st const*>(k_norm_weight.data_ptr()),
has_index
? reinterpret_cast<st const*>(index_q_norm_weight->data_ptr())
: nullptr,
has_index
? reinterpret_cast<st const*>(index_k_norm_weight->data_ptr())
: nullptr,
reinterpret_cast<st const*>(cos_sin_cache.data_ptr()),
reinterpret_cast<int64_t const*>(positions.data_ptr()),
insert_kv
? reinterpret_cast<int64_t const*>(slot_mapping->data_ptr())
: nullptr,
insert_kv ? reinterpret_cast<st*>(kv_cache->data_ptr()) : nullptr,
(insert_kv && has_index)
? reinterpret_cast<st*>(index_cache->data_ptr())
: nullptr,
static_cast<float>(eps), static_cast<int>(rotary_dim), num_tokens,
nq, nkv, niq, static_cast<int>(block_size), kv_s_block, kv_s_kv,
kv_s_token, kv_s_head, has_index, insert_kv, stream);
});
}
@@ -20,7 +20,7 @@
#include "torch_utils.h"
#include "../async_util.cuh"
#include "async_util.cuh"
#include "../cuda_compat.h"
#include "../type_convert.cuh"
#include "dispatch_utils.h"
+1 -19
View File
@@ -231,23 +231,6 @@ void fused_qk_norm_rope(torch::stable::Tensor& qkv, int64_t num_heads_q,
torch::stable::Tensor& position_ids,
int64_t forced_token_heads_per_warp);
// Horizontally-fused MiniMax-M3 QK-norm + partial NeoX RoPE (+ optional KV /
// index-cache insert). Dense layer: norm+RoPE only; sparse layer: also packs
// the index branch and scatters k/v/index_k into their paged caches.
void fused_minimax_m3_qknorm_rope_kv_insert(
torch::stable::Tensor& qkv, torch::stable::Tensor const& q_norm_weight,
torch::stable::Tensor const& k_norm_weight,
torch::stable::Tensor const& cos_sin_cache,
torch::stable::Tensor const& positions, int64_t num_heads,
int64_t num_kv_heads, int64_t rotary_dim, double eps,
std::optional<torch::stable::Tensor> index_q_norm_weight,
std::optional<torch::stable::Tensor> index_k_norm_weight,
int64_t num_index_heads, std::optional<torch::stable::Tensor> slot_mapping,
std::optional<torch::stable::Tensor> kv_cache,
std::optional<torch::stable::Tensor> index_cache, int64_t block_size,
std::optional<torch::stable::Tensor> q_out,
std::optional<torch::stable::Tensor> index_q_out);
// Sampler kernels (shared CUDA/ROCm)
void apply_repetition_penalties_(
torch::stable::Tensor& logits, const torch::stable::Tensor& prompt_mask,
@@ -293,8 +276,7 @@ void selective_scan_fwd(
// Activation kernels (shared CUDA/ROCm)
void silu_and_mul(torch::stable::Tensor& out, torch::stable::Tensor& input);
void silu_and_mul_clamp(torch::stable::Tensor& out,
torch::stable::Tensor& input, double limit,
double alpha = 1.0, double beta = 0.0);
torch::stable::Tensor& input, double limit);
void mul_and_silu(torch::stable::Tensor& out, torch::stable::Tensor& input);
void gelu_and_mul(torch::stable::Tensor& out, torch::stable::Tensor& input);
void gelu_tanh_and_mul(torch::stable::Tensor& out,
@@ -25,7 +25,7 @@
#include <cuda_fp8.h>
#include "cuda_utils.h"
#include "launch_bounds_utils.h"
#include "libtorch_stable/launch_bounds_utils.h"
// Define before including nvfp4_utils.cuh so the header
// can use this macro during compilation.
@@ -34,7 +34,7 @@
static_assert(CVT_FP4_ELTS_PER_THREAD == 16,
"MXFP4 experts quant requires PACK16 mode (CUDA >= 12.9)");
#include "launch_bounds_utils.h"
#include "libtorch_stable/launch_bounds_utils.h"
namespace vllm {
@@ -26,7 +26,7 @@
#include "cuda_utils.h"
#include "nvfp4_utils.cuh"
#include "launch_bounds_utils.h"
#include "libtorch_stable/launch_bounds_utils.h"
namespace vllm {
@@ -26,7 +26,7 @@
#include "../../cuda_vec_utils.cuh"
#include "cuda_utils.h"
#include "launch_bounds_utils.h"
#include "libtorch_stable/launch_bounds_utils.h"
// Define before including nvfp4_utils.cuh so the header
// can use this macro during compilation.
@@ -7,14 +7,11 @@
#include <torch/csrc/stable/ops.h>
// NOTE: These headers are intentionally kept in csrc/quantization/gguf/ (not
// moved to libtorch_stable) to avoid unnecessary reformatting that would break
// git rename detection and pollute blame history.
#include "../../../quantization/gguf/ggml-common.h"
#include "../../../quantization/gguf/vecdotq.cuh"
#include "../../../quantization/gguf/dequantize.cuh"
#include "../../../quantization/gguf/mmvq.cuh"
#include "../../../quantization/gguf/mmq.cuh"
#include "ggml-common.h"
#include "vecdotq.cuh"
#include "dequantize.cuh"
#include "mmvq.cuh"
#include "mmq.cuh"
#include "moe.cuh"
#include "moe_vec.cuh"
+1 -1
View File
@@ -7,7 +7,7 @@
#include "torch_utils.h"
#ifndef USE_ROCM
#include "../persistent_topk.cuh"
#include "persistent_topk.cuh"
#endif
namespace {
+2 -17
View File
@@ -337,17 +337,6 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
"bool is_neox, Tensor position_ids, "
"int forced_token_heads_per_warp=-1) -> ()");
// Horizontally-fused MiniMax-M3 QK-norm + partial NeoX RoPE + KV-insert.
ops.def(
"fused_minimax_m3_qknorm_rope_kv_insert("
"Tensor! qkv, Tensor q_norm_weight, Tensor k_norm_weight, "
"Tensor cos_sin_cache, Tensor positions, int num_heads, "
"int num_kv_heads, int rotary_dim, float eps, "
"Tensor? index_q_norm_weight, Tensor? index_k_norm_weight, "
"int num_index_heads, "
"Tensor? slot_mapping, Tensor!? kv_cache, Tensor!? index_cache, "
"int block_size, Tensor!? q_out, Tensor!? index_q_out) -> ()");
// Apply repetition penalties to logits in-place.
ops.def(
"apply_repetition_penalties_(Tensor! logits, Tensor prompt_mask, "
@@ -375,11 +364,9 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
ops.def("mul_and_silu(Tensor! out, Tensor input) -> ()");
// SwiGLU activation with input clamping.
// alpha scales the sigmoid (gate * sigmoid(alpha * gate)); beta is added to
// the up half (up + beta). Defaults alpha=1.0, beta=0.0 give silu(gate)*up.
ops.def(
"silu_and_mul_with_clamp(Tensor! result, Tensor input, float limit, "
"float alpha=1.0, float beta=0.0) -> ()");
"silu_and_mul_with_clamp(Tensor! result, Tensor input, float limit) "
"-> ()");
// Activation function used in GeGLU with `none` approximation.
ops.def("gelu_and_mul(Tensor! out, Tensor input) -> ()");
@@ -584,8 +571,6 @@ STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
// Positional encoding kernels (shared CUDA/ROCm)
ops.impl("rotary_embedding", TORCH_BOX(&rotary_embedding));
ops.impl("fused_qk_norm_rope", TORCH_BOX(&fused_qk_norm_rope));
ops.impl("fused_minimax_m3_qknorm_rope_kv_insert",
TORCH_BOX(&fused_minimax_m3_qknorm_rope_kv_insert));
// Sampler kernels (shared CUDA/ROCm)
ops.impl("apply_repetition_penalties_",
-742
View File
@@ -1,742 +0,0 @@
// CUDA C++ q2k -> k2q CSR builder.
//
// Five-stage pipeline. q-ascending order within each CSR row is preserved
// by partitioning q across (CTA, warp_in_CTA) units; each unit owns a
// contiguous q-sub-range and reserves a contiguous slot range per row via
// a precomputed exclusive prefix scan.
//
// M: build_row_map -- round-robin packing of rows across batches
// H: histogram + tile_counts
// PR: row prefix -- single block per head, row_counts -> row_ptr
// PT: tile prefix -- multi-block, scan tile_counts along (c, w) axis
// S: scatter (sorted) -- per-warp slot range, q-sequential within warp
//
// Per-warp partitioning: each CTA has kWarps warps; warp w of CTA c owns
// q-range [c*q_per_cta + w*q_per_warp, c*q_per_cta + (w+1)*q_per_warp).
// tile_counts is shaped [G * kWarps, H, total_rows]; the "row" dimension
// of the prefix scan is the flattened (c * kWarps + w) index, scanned in
// lexicographic order so that warp-local slot ranges concatenate to the
// global q-sorted output.
#include <torch/all.h>
#include <ATen/cuda/CUDAContext.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <algorithm>
#define CHECK_CUDA(x) TORCH_CHECK((x).is_cuda(), #x " must be CUDA")
#define CHECK_CONTIGUOUS(x) \
TORCH_CHECK((x).is_contiguous(), #x " must be contiguous")
#define CHECK_INT(x) \
TORCH_CHECK((x).scalar_type() == at::kInt, #x " must be int32")
#define CHECK_INPUT(x) \
CHECK_CUDA(x); \
CHECK_CONTIGUOUS(x); \
CHECK_INT(x)
namespace {
constexpr int kWarpSize = 32;
__device__ __forceinline__ void advance_batch_only(int const* __restrict__ cu_q,
int B, int q_abs, int& bi) {
while (bi < B && cu_q[bi + 1] <= q_abs) ++bi;
}
// Atomic increment of a 16-bit half within a 32-bit SMEM word; returns the
// OLD 16-bit value (slot). Per-warp count must stay < 32768 so the low
// half does not carry into the high half.
// base_int32 : int32 pointer; element i holds rows 2*i (low) and 2*i+1
// (high).
__device__ __forceinline__ int atomic_inc_int16_packed(int* base_int32,
int row) {
int idx = row >> 1;
int shift = (row & 1) << 4; // 0 or 16
int delta = 1 << shift;
int old = atomicAdd(&base_int32[idx], delta);
return (old >> shift) & 0xFFFF;
}
// Read 16-bit half from packed int32 storage.
__device__ __forceinline__ int read_int16_packed(int const* base_int32,
int row) {
int v = base_int32[row >> 1];
int shift = (row & 1) << 4;
return (v >> shift) & 0xFFFF;
}
// ---------------------------------------------------------------------------
// M: round-robin row map.
// ---------------------------------------------------------------------------
template <int kBlockK>
__global__ void k2q_build_row_map_kernel(int const* __restrict__ cu_k,
int* __restrict__ row_map,
int* __restrict__ row_coords, int B,
int max_kv_blocks) {
int level = blockIdx.x;
if (level >= max_kv_blocks) return;
if (threadIdx.x != 0) return;
int rows_before = 0;
for (int b = 0; b < B; ++b) {
int rb = (cu_k[b + 1] - cu_k[b] + kBlockK - 1) / kBlockK;
rows_before += (rb < level ? rb : level);
}
int active_before = 0;
for (int b = 0; b < B; ++b) {
int rb = (cu_k[b + 1] - cu_k[b] + kBlockK - 1) / kBlockK;
if (rb > level) {
int row_linear = rows_before + active_before;
row_map[(size_t)b * max_kv_blocks + level] = row_linear;
if (row_coords != nullptr) {
row_coords[(size_t)row_linear * 2] = b;
row_coords[(size_t)row_linear * 2 + 1] = level;
}
++active_before;
} else {
row_map[(size_t)b * max_kv_blocks + level] = -1;
}
}
}
// ---------------------------------------------------------------------------
// H: per-warp histogram + tile_counts.
// kWarps warps per CTA, each owns q-sub-range = q_per_cta / kWarps.
// SMEM hist[kWarps, total_rows] int32 (stored as packed int16 cursor:
// 2 entries per int32 word). Each warp counts to its own row.
// At end-of-CTA, write tile_counts[c*kWarps + w, h, r] = smem_hist[w, r]
// and atomicAdd(row_counts[h, r], sum over w of smem_hist[w, r]).
// ---------------------------------------------------------------------------
template <int kTopK, int kBlockK, int kWarps>
__global__ void k2q_hist_kernel(int const* __restrict__ q2k,
int const* __restrict__ cu_q,
int const* __restrict__ row_map,
int* __restrict__ row_counts,
int* __restrict__ tile_counts, int H, int B,
int S_Q, int total_rows, int max_kv_blocks,
int q_per_cta, int q_per_warp) {
constexpr int kThreads = kWarps * kWarpSize;
extern __shared__ int smem_hist_int[];
int* smem_hist = smem_hist_int;
int tid = threadIdx.x;
int warp_id = tid >> 5;
int lane = tid & 31;
int c = blockIdx.x;
int q_start_cta = c * q_per_cta;
int q_end_cta = min(q_start_cta + q_per_cta, S_Q);
int q_start_warp = min(q_start_cta + warp_id * q_per_warp, q_end_cta);
int q_end_warp = min(q_start_warp + q_per_warp, q_end_cta);
constexpr int kInt4PerToken = kTopK / 4;
int packed_per_warp = (total_rows + 1) >> 1;
int* my_hist = smem_hist + warp_id * packed_per_warp;
for (int h = 0; h < H; ++h) {
for (int i = lane; i < packed_per_warp; i += kWarpSize) my_hist[i] = 0;
__syncthreads();
if (q_start_warp < q_end_warp) {
int bi = 0;
int qi = q_start_warp + lane;
advance_batch_only(cu_q, B, qi, bi);
int4 const* head_topk4 =
reinterpret_cast<int4 const*>(q2k + (size_t)h * S_Q * kTopK);
for (; qi < q_end_warp; qi += kWarpSize) {
advance_batch_only(cu_q, B, qi, bi);
int const* my_row_map = row_map + (size_t)bi * max_kv_blocks;
int4 buf[kInt4PerToken];
#pragma unroll
for (int v = 0; v < kInt4PerToken; ++v) {
buf[v] = head_topk4[(size_t)qi * kInt4PerToken + v];
}
#pragma unroll
for (int t = 0; t < kTopK; ++t) {
int kvb_local = reinterpret_cast<int const*>(buf)[t];
if (kvb_local >= 0 && kvb_local < max_kv_blocks) {
int row = my_row_map[kvb_local];
if (row >= 0 && row < total_rows) {
atomic_inc_int16_packed(my_hist, row);
}
}
}
}
}
__syncthreads();
int* head_row_counts = row_counts + (size_t)h * total_rows;
// Each warp writes its own slice of tile_counts (full int32) by
// unpacking int16 entries from SMEM.
int* my_tile =
tile_counts + ((size_t)(c * kWarps + warp_id) * H + h) * total_rows;
for (int i = lane; i < total_rows; i += kWarpSize) {
my_tile[i] = read_int16_packed(my_hist, i);
}
__syncthreads();
// Sum across warps (int32 accumulator), atomicAdd to row_counts.
for (int i = tid; i < total_rows; i += kThreads) {
int sum = 0;
#pragma unroll
for (int w = 0; w < kWarps; ++w) {
sum += read_int16_packed(smem_hist + w * packed_per_warp, i);
}
if (sum > 0) atomicAdd(&head_row_counts[i], sum);
}
if (h + 1 < H) __syncthreads();
}
}
// ---------------------------------------------------------------------------
// PR: row prefix. One block per head.
// ---------------------------------------------------------------------------
template <int kThreads>
__global__ void k2q_row_prefix_kernel(int const* __restrict__ row_counts,
int* __restrict__ row_ptr,
int const* __restrict__ row_coords,
int* __restrict__ scheduler_metadata,
int* __restrict__ work_count,
int total_rows, int target_q_per_cta,
int work_capacity) {
int h = blockIdx.x;
int tid = threadIdx.x;
__shared__ int scan_buf[kThreads];
int const* head_counts = row_counts + (size_t)h * total_rows;
int* head_rowptr = row_ptr + (size_t)h * (total_rows + 1);
int chunk = (total_rows + kThreads - 1) / kThreads;
int lo = tid * chunk;
int hi = min(lo + chunk, total_rows);
int local_sum = 0;
for (int i = lo; i < hi; ++i) local_sum += head_counts[i];
scan_buf[tid] = local_sum;
__syncthreads();
for (int off = 1; off < kThreads; off <<= 1) {
int add = (tid >= off) ? scan_buf[tid - off] : 0;
__syncthreads();
scan_buf[tid] += add;
__syncthreads();
}
int running = scan_buf[tid] - local_sum;
for (int i = lo; i < hi; ++i) {
int row_count = head_counts[i];
running += row_count;
head_rowptr[i + 1] = running;
if (scheduler_metadata != nullptr && work_count != nullptr &&
row_count > 0) {
int num_chunks = (row_count + target_q_per_cta - 1) / target_q_per_cta;
int base = atomicAdd(work_count, num_chunks);
int batch_idx = row_coords[(size_t)i * 2];
int kv_block_idx = row_coords[(size_t)i * 2 + 1];
for (int c = 0; c < num_chunks; ++c) {
int work_idx = base + c;
if (work_idx < work_capacity) {
int q_begin = c * target_q_per_cta;
int q_count = min(target_q_per_cta, row_count - q_begin);
int* meta = scheduler_metadata + (size_t)work_idx * 6;
meta[0] = h;
meta[1] = i;
meta[2] = q_begin;
meta[3] = q_count;
meta[4] = batch_idx;
meta[5] = kv_block_idx;
}
}
}
}
}
// ---------------------------------------------------------------------------
// PT_smem: SMEM-staged tile prefix scan.
// Each block handles kRowsPerBlock rows for one head h. Cooperative load
// of tile_counts[*, h, base_r..base_r+M) into SMEM (better coalescing
// than per-warp uncoalesced stride reads), then per-warp scan in SMEM,
// then cooperative store back. Fuses row_ptr into the base.
// ---------------------------------------------------------------------------
template <int kThreads, int kRowsPerBlock>
__global__ void k2q_tile_prefix_smem_kernel(int* __restrict__ tile_counts,
int const* __restrict__ row_ptr,
int H, int total_rows,
int G_total) {
static_assert(kRowsPerBlock > 0, "kRowsPerBlock must be positive");
extern __shared__ int smem_tprefix[];
// smem layout: smem[r_off][g] for r_off in [0, M), g in [0, G_total).
int tid = threadIdx.x;
int lane = tid & 31;
int warp_id = tid >> 5;
// Grid: H * blocks_per_h. Each block stays within a single head h
// and processes kRowsPerBlock contiguous rows starting at b_in_h *
// kRowsPerBlock. (Earlier flat-grid mapping `h = block_job /
// total_rows; base_r = block_job - h*total_rows` skipped rows when
// total_rows was not a multiple of kRowsPerBlock and H > 1, because
// the last partial block of head h-1 left blocks of head h starting
// at a non-zero row offset.)
int blocks_per_h = (total_rows + kRowsPerBlock - 1) / kRowsPerBlock;
int h = blockIdx.x / blocks_per_h;
int b_in_h = blockIdx.x - h * blocks_per_h;
if (h >= H) return;
int base_r = b_in_h * kRowsPerBlock;
if (base_r >= total_rows) return;
int actual_M = min(kRowsPerBlock, total_rows - base_r);
size_t stride_g = (size_t)H * total_rows;
int* base_ptr = tile_counts + (size_t)h * total_rows + base_r;
int total_elems = G_total * actual_M;
// Cooperative load. Pattern: thread tid -> (r_off=tid%M, g=tid/M),
// then strided. 32 lanes hit M r's × (32/M) g's, giving 32/M cache
// lines per warp (vs 32 in the naive stride-along-g pattern).
for (int i = tid; i < total_elems; i += kThreads) {
int r_off = i % actual_M;
int g = i / actual_M;
smem_tprefix[r_off * G_total + g] = base_ptr[g * stride_g + r_off];
}
__syncthreads();
// Per-warp scan: warp w scans row (base_r + w) if w < actual_M.
if (warp_id < actual_M) {
int abs_r = base_r + warp_id;
int rp = row_ptr[(size_t)h * (total_rows + 1) + abs_r];
int* my_smem = smem_tprefix + warp_id * G_total;
int running = rp;
for (int g0 = 0; g0 < G_total; g0 += kWarpSize) {
int g = g0 + lane;
int v = (g < G_total) ? my_smem[g] : 0;
int x = v;
#pragma unroll
for (int off = 1; off < kWarpSize; off <<= 1) {
int nbr = __shfl_up_sync(0xFFFFFFFF, x, off);
if (lane >= off) x += nbr;
}
int excl = running + x - v;
if (g < G_total) my_smem[g] = excl;
int chunk_sum = __shfl_sync(0xFFFFFFFF, x, 31);
running += chunk_sum;
}
}
__syncthreads();
// Cooperative store back.
for (int i = tid; i < total_elems; i += kThreads) {
int r_off = i % actual_M;
int g = i / actual_M;
base_ptr[g * stride_g + r_off] = smem_tprefix[r_off * G_total + g];
}
}
// ---------------------------------------------------------------------------
// S: scatter. kWarps warps per CTA, each owns q-sub-range. Per-warp SMEM
// cursor and per-warp tile_offset slot range. Within a warp, q's are
// processed sequentially; lanes 0..kTopK-1 handle the topK slots in
// lockstep. Across distinct q's in the same warp, the lockstep ordering
// guarantees q-monotonic atomicAdd on smem_cursor[r].
// ---------------------------------------------------------------------------
// kQPerIter * kTopK lanes are active per warp iter; remaining lanes idle.
// For kTopK=16, kQPerIter=2 uses all 32 lanes; for kTopK=8, kQPerIter=4.
// CORRECTNESS NOTE: relies on lane-ordered SMEM atomicAdd return values
// within a single warp instruction (verified on B200; tests pass).
//
// SMEM cursor stored as packed int16 (two cursors per int32). Per-warp
// row count must stay < 32768 (~q_per_warp * kTopK at max sink), which
// holds for all task.md sizes up to 1024K.
template <int kTopK, int kBlockK, int kWarps>
__global__ void k2q_scatter_kernel(
int const* __restrict__ q2k, int const* __restrict__ cu_q,
int const* __restrict__ row_map, int const* __restrict__ abs_base,
int* __restrict__ q_idx, int* __restrict__ qsplit_idx,
int* __restrict__ split_counts, int H, int B, int S_Q, int total_rows,
int max_kv_blocks, int q_per_cta, int q_per_warp, int max_seqlen_q) {
constexpr int kQPerIter = kWarpSize / kTopK > 0 ? kWarpSize / kTopK : 1;
extern __shared__ int smem_cursor_int[];
int* smem_cursor = smem_cursor_int;
int tid = threadIdx.x;
int warp_id = tid >> 5;
int lane = tid & 31;
int c = blockIdx.x;
int q_start_cta = c * q_per_cta;
int q_end_cta = min(q_start_cta + q_per_cta, S_Q);
int q_start_warp = min(q_start_cta + warp_id * q_per_warp, q_end_cta);
int q_end_warp = min(q_start_warp + q_per_warp, q_end_cta);
int q_in_iter = lane / kTopK;
int slot_in_q = lane % kTopK;
bool lane_active = (lane < kQPerIter * kTopK);
// Per-warp packed cursor: total_rows int16 entries -> ceil(total_rows/2)
// int32.
int packed_per_warp = (total_rows + 1) >> 1;
int* my_cursor = smem_cursor + warp_id * packed_per_warp;
for (int h = 0; h < H; ++h) {
for (int i = lane; i < packed_per_warp; i += kWarpSize) my_cursor[i] = 0;
__syncwarp();
if (q_start_warp < q_end_warp) {
int bi = 0;
advance_batch_only(cu_q, B, q_start_warp, bi);
int const* head_q2k = q2k + (size_t)h * S_Q * kTopK;
int const* my_abs_base =
abs_base + ((size_t)(c * kWarps + warp_id) * H + h) * total_rows;
int* head_qidx = q_idx + (size_t)h * S_Q * kTopK;
// (Hot-row register cache experiment showed no measurable
// benefit; relying on L1 to keep row 0 / row total_rows-1
// hot since they're hit every iteration in sink workloads.)
constexpr int kUnroll = 16;
int qi_base = q_start_warp;
for (; qi_base + kUnroll * kQPerIter <= q_end_warp;
qi_base += kUnroll * kQPerIter) {
int kvb[kUnroll];
int qloc[kUnroll];
int batch[kUnroll];
int const* rmap[kUnroll];
#pragma unroll
for (int u = 0; u < kUnroll; ++u) {
int qi_u = qi_base + u * kQPerIter + q_in_iter;
kvb[u] = -1;
qloc[u] = 0;
batch[u] = 0;
if (lane_active) {
advance_batch_only(cu_q, B, qi_u, bi);
qloc[u] = qi_u - cu_q[bi];
batch[u] = bi;
kvb[u] = head_q2k[(size_t)qi_u * kTopK + slot_in_q];
}
rmap[u] = row_map + (size_t)bi * max_kv_blocks;
}
int row[kUnroll];
#pragma unroll
for (int u = 0; u < kUnroll; ++u) {
row[u] = -1;
if (lane_active && kvb[u] >= 0 && kvb[u] < max_kv_blocks)
row[u] = rmap[u][kvb[u]];
}
// Pre-issue all kUnroll abs_base loads in parallel before
// the atomic chain so memory pipeline runs concurrently
// with SMEM atomic-adds.
int abs_v[kUnroll];
#pragma unroll
for (int u = 0; u < kUnroll; ++u) {
abs_v[u] =
(row[u] >= 0 && row[u] < total_rows) ? my_abs_base[row[u]] : 0;
}
#pragma unroll
for (int u = 0; u < kUnroll; ++u) {
int r = row[u];
bool valid_edge = r >= 0 && r < total_rows;
unsigned int valid_mask = __ballot_sync(0xFFFFFFFFu, valid_edge);
unsigned int group_mask =
(kTopK == 32) ? 0xFFFFFFFFu
: (((1u << kTopK) - 1u) << (q_in_iter * kTopK));
unsigned int lower_lane_mask = lane == 0 ? 0u : ((1u << lane) - 1u);
int split_slot = __popc(valid_mask & group_mask & lower_lane_mask);
int valid_count = __popc(valid_mask & group_mask);
if (split_counts != nullptr && slot_in_q == 0) {
split_counts[((size_t)batch[u] * max_seqlen_q + qloc[u]) * H + h] =
valid_count;
}
if (valid_edge) {
int slot = atomic_inc_int16_packed(my_cursor, r);
int out_pos = abs_v[u] + slot;
head_qidx[out_pos] = qloc[u];
if (qsplit_idx != nullptr) {
qsplit_idx[(size_t)h * S_Q * kTopK + out_pos] =
qloc[u] | ((split_slot & 0xFF) << 24);
}
}
}
}
// Tail: 1-3 iters left.
for (; qi_base < q_end_warp; qi_base += kQPerIter) {
int my_qi = qi_base + q_in_iter;
bool valid_q = (my_qi < q_end_warp) && lane_active;
int kvb_local = -1;
int q_local = 0;
int batch_local = 0;
if (valid_q) {
advance_batch_only(cu_q, B, my_qi, bi);
batch_local = bi;
q_local = my_qi - cu_q[bi];
kvb_local = head_q2k[(size_t)my_qi * kTopK + slot_in_q];
}
int const* my_row_map = row_map + (size_t)bi * max_kv_blocks;
int row = -1;
if (valid_q && kvb_local >= 0 && kvb_local < max_kv_blocks) {
row = my_row_map[kvb_local];
}
bool valid_edge = row >= 0 && row < total_rows;
unsigned int valid_mask = __ballot_sync(0xFFFFFFFFu, valid_edge);
unsigned int group_mask =
(kTopK == 32) ? 0xFFFFFFFFu
: (((1u << kTopK) - 1u) << (q_in_iter * kTopK));
unsigned int lower_lane_mask = lane == 0 ? 0u : ((1u << lane) - 1u);
int split_slot = __popc(valid_mask & group_mask & lower_lane_mask);
int valid_count = __popc(valid_mask & group_mask);
if (split_counts != nullptr && valid_q && slot_in_q == 0) {
split_counts[((size_t)batch_local * max_seqlen_q + q_local) * H + h] =
valid_count;
}
if (valid_edge) {
int slot = atomic_inc_int16_packed(my_cursor, row);
int out_pos = my_abs_base[row] + slot;
head_qidx[out_pos] = q_local;
if (qsplit_idx != nullptr) {
qsplit_idx[(size_t)h * S_Q * kTopK + out_pos] =
q_local | ((split_slot & 0xFF) << 24);
}
}
}
}
if (h + 1 < H) __syncthreads();
}
}
} // anonymous namespace
// ===========================================================================
// Host orchestration
// ===========================================================================
template <int kTopK, int kBlockK>
static void launch_pipeline(torch::Tensor q2k, torch::Tensor cu_q,
torch::Tensor cu_k, torch::Tensor row_ptr,
torch::Tensor q_idx, int total_rows,
int max_kv_blocks,
torch::Tensor scheduler_metadata = torch::Tensor(),
torch::Tensor work_count = torch::Tensor(),
torch::Tensor qsplit_idx = torch::Tensor(),
torch::Tensor split_counts = torch::Tensor(),
int target_q_per_cta = 1, int work_capacity = 0,
int max_seqlen_q = 0) {
int H = (int)q2k.size(0);
int S_Q = (int)q2k.size(1);
int topK = (int)q2k.size(2);
TORCH_CHECK(topK == kTopK, "topK runtime != template kTopK");
int B = (int)cu_q.size(0) - 1;
auto device = q2k.device();
cudaStream_t stream = at::cuda::getCurrentCUDAStream();
AT_CUDA_CHECK(cudaMemsetAsync(row_ptr.data_ptr<int>(), 0,
(size_t)H * (total_rows + 1) * sizeof(int),
stream));
AT_CUDA_CHECK(cudaMemsetAsync(q_idx.data_ptr<int>(), 0xFF,
(size_t)H * S_Q * kTopK * sizeof(int), stream));
auto opts = torch::TensorOptions().dtype(torch::kInt32).device(device);
auto row_counts = torch::zeros({H, total_rows}, opts);
auto row_map = torch::empty({B, max_kv_blocks}, opts);
bool emit_schedule = scheduler_metadata.defined();
auto row_coords =
emit_schedule ? torch::empty({total_rows, 2}, opts) : torch::Tensor();
int* scheduler_metadata_ptr =
emit_schedule ? scheduler_metadata.data_ptr<int>() : nullptr;
int* work_count_ptr = emit_schedule ? work_count.data_ptr<int>() : nullptr;
int* qsplit_idx_ptr = emit_schedule ? qsplit_idx.data_ptr<int>() : nullptr;
int* split_counts_ptr =
emit_schedule ? split_counts.data_ptr<int>() : nullptr;
int* row_coords_ptr = emit_schedule ? row_coords.data_ptr<int>() : nullptr;
if (emit_schedule) {
AT_CUDA_CHECK(cudaMemsetAsync(work_count_ptr, 0, sizeof(int), stream));
AT_CUDA_CHECK(cudaMemsetAsync(scheduler_metadata_ptr, 0,
(size_t)work_capacity * 6 * sizeof(int),
stream));
}
int dev = q2k.get_device();
int num_sms = 0;
AT_CUDA_CHECK(
cudaDeviceGetAttribute(&num_sms, cudaDevAttrMultiProcessorCount, dev));
// -- Pick kWarps per CTA based on SMEM budget for cursor/hist ---------
// SMEM per CTA = kWarps * total_rows * sizeof(int) (for both H and S).
// Want at least 2 CTAs/SM for memory parallelism. SM100 SMEM = 228KB.
// Pick the largest kWarps that fits two CTAs/SM, capped at 4.
// SMEM cursor packed as int16 (2 entries per int32 word):
int per_warp_smem = ((total_rows + 1) >> 1) * (int)sizeof(int);
int kWarps_pick = 4;
while (kWarps_pick > 1 && (kWarps_pick * per_warp_smem) * 2 > 228 * 1024) {
kWarps_pick >>= 1;
}
if (kWarps_pick < 1) kWarps_pick = 1;
// -- Pick G (CTAs) ----------------------------------------------------
// For each (kWarps, per_warp_smem) pair, the SMEM-bound occupancy is
// 228KB / (kWarps*per_warp_smem) CTAs/SM. We size G as
// num_sms * occupancy so a single resident wave covers all CTAs and
// the memory pipeline runs at peak.
int per_cta_smem_bytes = kWarps_pick * per_warp_smem;
int max_ctas_per_sm =
std::max(1, (228 * 1024) / std::max(1, per_cta_smem_bytes));
if (max_ctas_per_sm > 8) max_ctas_per_sm = 8;
constexpr int kMinQPerCta = 256;
// Cap target_g at num_sms * 3 — empirically this balances
// per-CTA work-size against parallelism. Higher caps regress
// mid-size cases due to row_counts atomicAdd contention and
// smaller q_per_cta. SMEM-bound configurations naturally cap
// lower if max_ctas_per_sm < 3.
int target_g = num_sms * std::min(max_ctas_per_sm, 3);
int max_g_for_q = (S_Q + kMinQPerCta - 1) / kMinQPerCta;
int G = std::min({target_g, max_g_for_q, S_Q});
if (G < 1) G = 1;
int q_per_cta = (S_Q + G - 1) / G;
G = (S_Q + q_per_cta - 1) / q_per_cta;
int q_per_warp = (q_per_cta + kWarps_pick - 1) / kWarps_pick;
int G_total = G * kWarps_pick;
auto tile_counts = torch::empty({G_total, H, total_rows}, opts);
// -- Compile-time switch on kWarps for the templated kernels ---------
auto rmap_fn = k2q_build_row_map_kernel<kBlockK>;
auto rprefix_fn = k2q_row_prefix_kernel<1024>;
constexpr int kPtRowsPerBlock = 8;
constexpr int kPtThreads = 256;
auto tprefix_smem_fn =
k2q_tile_prefix_smem_kernel<kPtThreads, kPtRowsPerBlock>;
if (max_kv_blocks > 0) {
rmap_fn<<<max_kv_blocks, 32, 0, stream>>>(cu_k.data_ptr<int>(),
row_map.data_ptr<int>(),
row_coords_ptr, B, max_kv_blocks);
}
auto launch_hist_scatter = [&](auto kWarps_const) {
constexpr int W = decltype(kWarps_const)::value;
size_t smem_bytes = (size_t)W * per_warp_smem;
auto hist_fn = k2q_hist_kernel<kTopK, kBlockK, W>;
auto scat_fn = k2q_scatter_kernel<kTopK, kBlockK, W>;
AT_CUDA_CHECK(cudaFuncSetAttribute(
hist_fn, cudaFuncAttributeMaxDynamicSharedMemorySize, (int)smem_bytes));
AT_CUDA_CHECK(cudaFuncSetAttribute(
scat_fn, cudaFuncAttributeMaxDynamicSharedMemorySize, (int)smem_bytes));
hist_fn<<<G, W * kWarpSize, smem_bytes, stream>>>(
q2k.data_ptr<int>(), cu_q.data_ptr<int>(), row_map.data_ptr<int>(),
row_counts.data_ptr<int>(), tile_counts.data_ptr<int>(), H, B, S_Q,
total_rows, max_kv_blocks, q_per_cta, q_per_warp);
rprefix_fn<<<H, 1024, 0, stream>>>(
row_counts.data_ptr<int>(), row_ptr.data_ptr<int>(),
emit_schedule ? row_coords.data_ptr<int>() : nullptr,
scheduler_metadata_ptr, work_count_ptr, total_rows, target_q_per_cta,
work_capacity);
// Grid is H * blocks_per_h so each block stays within a single
// head; flat (H*total_rows) grid would skip rows when total_rows
// is not a multiple of kPtRowsPerBlock.
int blocks_per_h = (total_rows + kPtRowsPerBlock - 1) / kPtRowsPerBlock;
int pt_grid = H * blocks_per_h;
if (pt_grid < 1) pt_grid = 1;
size_t pt_smem = (size_t)kPtRowsPerBlock * G_total * sizeof(int);
AT_CUDA_CHECK(cudaFuncSetAttribute(
tprefix_smem_fn, cudaFuncAttributeMaxDynamicSharedMemorySize,
(int)pt_smem));
tprefix_smem_fn<<<pt_grid, kPtThreads, pt_smem, stream>>>(
tile_counts.data_ptr<int>(), row_ptr.data_ptr<int>(), H, total_rows,
G_total);
scat_fn<<<G, W * kWarpSize, smem_bytes, stream>>>(
q2k.data_ptr<int>(), cu_q.data_ptr<int>(), row_map.data_ptr<int>(),
tile_counts.data_ptr<int>(), q_idx.data_ptr<int>(), qsplit_idx_ptr,
split_counts_ptr, H, B, S_Q, total_rows, max_kv_blocks, q_per_cta,
q_per_warp, max_seqlen_q);
};
if (kWarps_pick == 4) {
launch_hist_scatter(std::integral_constant<int, 4>{});
} else if (kWarps_pick == 2) {
launch_hist_scatter(std::integral_constant<int, 2>{});
} else {
launch_hist_scatter(std::integral_constant<int, 1>{});
}
}
void run_minimax_m3_build_k2q_csr_with_schedule(
torch::Tensor q2k, torch::Tensor cu_q, torch::Tensor cu_k,
torch::Tensor row_ptr, torch::Tensor q_idx,
torch::Tensor scheduler_metadata, torch::Tensor work_count,
torch::Tensor qsplit_idx, torch::Tensor split_counts, int64_t topk,
int64_t blk_kv, int64_t total_rows, int64_t max_kv_blocks,
int64_t target_q_per_cta, int64_t work_capacity, int64_t max_seqlen_q) {
CHECK_INPUT(q2k);
CHECK_INPUT(cu_q);
CHECK_INPUT(cu_k);
CHECK_INPUT(row_ptr);
CHECK_INPUT(q_idx);
CHECK_INPUT(scheduler_metadata);
CHECK_INPUT(work_count);
CHECK_INPUT(qsplit_idx);
CHECK_INPUT(split_counts);
TORCH_CHECK(blk_kv == 128, "build_k2q_csr only supports blk_kv == 128");
int H = (int)q2k.size(0);
int S_Q = (int)q2k.size(1);
int tr = (int)total_rows;
int mkv = (int)max_kv_blocks;
int target = (int)target_q_per_cta;
int capacity = (int)work_capacity;
int max_sq = (int)max_seqlen_q;
TORCH_CHECK(tr >= 0 && mkv >= 0 && target > 0 && capacity > 0 && max_sq >= 0,
"invalid schedule sizing arguments");
TORCH_CHECK(row_ptr.size(0) == H && row_ptr.size(1) == tr + 1,
"row_ptr shape mismatch");
TORCH_CHECK(q_idx.size(0) == H && q_idx.size(1) == (int64_t)S_Q * (int)topk,
"q_idx shape mismatch");
TORCH_CHECK(qsplit_idx.sizes() == q_idx.sizes(), "qsplit_idx shape mismatch");
TORCH_CHECK(
scheduler_metadata.size(0) == capacity && scheduler_metadata.size(1) == 6,
"scheduler_metadata shape mismatch");
TORCH_CHECK(work_count.numel() == 1,
"work_count must have one int32 element");
TORCH_CHECK(split_counts.dim() == 3 &&
split_counts.size(0) == cu_q.size(0) - 1 &&
split_counts.size(1) == max_sq && split_counts.size(2) == H,
"split_counts shape mismatch");
if (S_Q == 0 || tr == 0 || H == 0 || mkv == 0) {
cudaStream_t stream = at::cuda::getCurrentCUDAStream();
AT_CUDA_CHECK(cudaMemsetAsync(row_ptr.data_ptr<int>(), 0,
(size_t)H * (tr + 1) * sizeof(int), stream));
AT_CUDA_CHECK(cudaMemsetAsync(q_idx.data_ptr<int>(), 0xFF,
(size_t)H * S_Q * (int)topk * sizeof(int),
stream));
AT_CUDA_CHECK(
cudaMemsetAsync(work_count.data_ptr<int>(), 0, sizeof(int), stream));
if (split_counts.numel() > 0) {
AT_CUDA_CHECK(cudaMemsetAsync(split_counts.data_ptr<int>(), 0,
(size_t)split_counts.numel() * sizeof(int),
stream));
}
return;
}
if (topk == 16) {
launch_pipeline<16, 128>(q2k, cu_q, cu_k, row_ptr, q_idx, tr, mkv,
scheduler_metadata, work_count, qsplit_idx,
split_counts, target, capacity, max_sq);
} else if (topk == 8) {
launch_pipeline<8, 128>(q2k, cu_q, cu_k, row_ptr, q_idx, tr, mkv,
scheduler_metadata, work_count, qsplit_idx,
split_counts, target, capacity, max_sq);
} else if (topk == 32) {
launch_pipeline<32, 128>(q2k, cu_q, cu_k, row_ptr, q_idx, tr, mkv,
scheduler_metadata, work_count, qsplit_idx,
split_counts, target, capacity, max_sq);
} else if (topk == 4) {
launch_pipeline<4, 128>(q2k, cu_q, cu_k, row_ptr, q_idx, tr, mkv,
scheduler_metadata, work_count, qsplit_idx,
split_counts, target, capacity, max_sq);
} else {
TORCH_CHECK(false, "unsupported topK ", topk,
" (expected 4, 8, 16, or 32)");
}
}
+1 -10
View File
@@ -62,8 +62,7 @@ void rotary_embedding(torch::Tensor& positions, torch::Tensor& query,
void silu_and_mul(torch::Tensor& out, torch::Tensor& input);
void silu_and_mul_clamp(torch::Tensor& out, torch::Tensor& input, double limit,
double alpha = 1.0, double beta = 0.0);
void silu_and_mul_clamp(torch::Tensor& out, torch::Tensor& input, double limit);
void silu_and_mul_quant(torch::Tensor& out, torch::Tensor& input,
torch::Tensor& scale);
@@ -147,12 +146,4 @@ std::tuple<torch::Tensor, torch::Tensor> minimax_allreduce_rms_qk(
torch::Tensor const& norm_weight_k, torch::Tensor workspace,
int64_t const q_size, int64_t const kv_size, int64_t const rank,
int64_t const nranks, double const eps);
void run_minimax_m3_build_k2q_csr_with_schedule(
torch::Tensor q2k, torch::Tensor cu_q, torch::Tensor cu_k,
torch::Tensor row_ptr, torch::Tensor q_idx,
torch::Tensor scheduler_metadata, torch::Tensor work_count,
torch::Tensor qsplit_idx, torch::Tensor split_counts, int64_t topk,
int64_t blk_kv, int64_t total_rows, int64_t max_kv_blocks,
int64_t target_q_per_cta, int64_t work_capacity, int64_t max_seqlen_q);
#endif
-21
View File
@@ -187,27 +187,6 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
"float eps) -> (Tensor, Tensor)");
ops.impl("minimax_allreduce_rms_qk", torch::kCUDA, &minimax_allreduce_rms_qk);
ops.def(
"minimax_m3_build_k2q_csr_with_schedule("
"Tensor q2k,"
"Tensor cu_q,"
"Tensor cu_k,"
"Tensor! row_ptr,"
"Tensor! q_idx,"
"Tensor! scheduler_metadata,"
"Tensor! work_count,"
"Tensor! qsplit_idx,"
"Tensor! split_counts,"
"int topk,"
"int blk_kv,"
"int total_rows,"
"int max_kv_blocks,"
"int target_q_per_cta,"
"int work_capacity,"
"int max_seqlen_q) -> ()");
ops.impl("minimax_m3_build_k2q_csr_with_schedule", torch::kCUDA,
&run_minimax_m3_build_k2q_csr_with_schedule);
// conditionally compiled so impl in source file
#endif
}
+1 -1
View File
@@ -757,7 +757,7 @@ RUN --mount=type=cache,target=/opt/uv/cache \
# Install FlashInfer JIT cache (requires CUDA-version-specific index URL)
# https://docs.flashinfer.ai/installation.html
# From versions.json: .flashinfer.version
ARG FLASHINFER_VERSION=0.6.12
ARG FLASHINFER_VERSION=0.6.11.post2
RUN --mount=type=cache,target=/opt/uv/cache \
uv pip install --system flashinfer-jit-cache==${FLASHINFER_VERSION} \
--extra-index-url https://flashinfer.ai/whl/cu$(echo $CUDA_VERSION | cut -d. -f1,2 | tr -d '.')
+2 -2
View File
@@ -256,13 +256,13 @@ RUN pip install setuptools==75.6.0 packaging==23.2 ninja==1.11.1.3 build==1.2.2.
# build flashinfer for torch nightly from source around 10 mins
# release version: v0.6.12
# release version: v0.6.11.post2
# todo(elainewy): cache flashinfer build result for faster build
ENV CCACHE_DIR=/root/.cache/ccache
RUN --mount=type=cache,target=/root/.cache/ccache \
--mount=type=cache,target=/root/.cache/uv \
echo "git clone flashinfer..." \
&& git clone --depth 1 --branch v0.6.12 --recursive https://github.com/flashinfer-ai/flashinfer.git \
&& git clone --depth 1 --branch v0.6.11.post2 --recursive https://github.com/flashinfer-ai/flashinfer.git \
&& cd flashinfer \
&& git submodule update --init --recursive \
&& echo "finish git clone flashinfer..." \
+1 -1
View File
@@ -9,7 +9,7 @@ ARG PYTORCH_AUDIO_BRANCH="v2.9.0"
ARG PYTORCH_AUDIO_REPO="https://github.com/pytorch/audio.git"
ARG FA_BRANCH="0e60e394"
ARG FA_REPO="https://github.com/Dao-AILab/flash-attention.git"
ARG AITER_BRANCH="v0.1.13"
ARG AITER_BRANCH="v0.1.13.post1"
ARG AITER_REPO="https://github.com/ROCm/aiter.git"
ARG MORI_BRANCH="v1.1.0"
ARG MORI_REPO="https://github.com/ROCm/mori.git"
+1 -1
View File
@@ -68,7 +68,7 @@
"default": "true"
},
"FLASHINFER_VERSION": {
"default": "0.6.12"
"default": "0.6.11.post2"
},
"GDRCOPY_CUDA_VERSION": {
"default": "12.8"
+8 -2
View File
@@ -246,6 +246,12 @@ Every image listed in "image_files" is added to the request in the listed order
The "image" shorthand accepts the same values as "image_files". The "image_url" field accepts either an OpenAI-style object with a "url" field or a URL string.
By default, image references are sent to the serving endpoint as provided, with local image paths converted to `file://` URLs.
If the benchmark client should load local and HTTP(S) images before sending requests, pass `--custom-ensure-client-side-data` to encode them as base64 data URLs on the client side.
Existing `data:image/...` URLs are already self-contained and are kept unchanged.
```bash
# need a model with vision capability here
vllm serve Qwen/Qwen2-VL-7B-Instruct
@@ -253,13 +259,13 @@ vllm serve Qwen/Qwen2-VL-7B-Instruct
```bash
# run benchmarking script
vllm bench serve--save-result --save-detailed \
vllm bench serve --save-result --save-detailed \
--backend openai-chat \
--model Qwen/Qwen2-VL-7B-Instruct \
--endpoint /v1/chat/completions \
--dataset-name custom_image \
--dataset-path <path-to-your-image-data-jsonl> \
--allowed-local-media-path /path/to/image/folder
--custom-ensure-client-side-data
```
Note that we need to use the `openai-chat` backend and `/v1/chat/completions` endpoint for multimodal inputs.
+2 -14
View File
@@ -170,8 +170,8 @@ Priority is **1 = highest** (tried first).
| Backend | Version | Dtypes | KV Dtypes | Block Sizes | Head Sizes | Sink | Non-Causal | MM Prefix | DCP | Attention Types | Compute Cap. |
| ------- | ------- | ------ | --------- | ----------- | ---------- | ---- | ---------- | --------- | --- | --------------- | ------------ |
| `CPU_ATTN` | | fp16, bf16, fp32 | `auto`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | %16 | 32, 64, 80, 96, 112, 128, 160, 192, 224, 256, 512 | ❌ | ❌ | ❌ | ❌ | All | N/A |
| `FLASHINFER` | Native† | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | 16, 32, 64, 128, 256, 512, 1024 | 64, 128, 256, 512 | ❌ | ❌ | ❌ | ✅ | Decoder | 7.x-9.x |
| `FLASHINFER` | TRTLLM† | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2`, `nvfp4` | 16, 32, 64, 128, 256, 512, 1024 | 64, 128, 256, 512 | ✅ | ❌ | ❌ | ✅ | Decoder | 10.x |
| `FLASHINFER` | Native† | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | 16, 32, 64 | 64, 128, 256, 512 | ❌ | ❌ | ❌ | ✅ | Decoder | 7.x-9.x |
| `FLASHINFER` | TRTLLM† | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2`, `nvfp4` | 16, 32, 64 | 64, 128, 256, 512 | ✅ | ❌ | ❌ | ✅ | Decoder | 10.x |
| `FLASH_ATTN` | FA2* | fp16, bf16 | `auto`, `float16`, `bfloat16` | %16 | Any | ❌ | ✅ | ❌ | ✅ | All | ≥8.0 |
| `FLASH_ATTN` | FA3* | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | %16 | Any | ✅ | ✅ | ❌ | ✅ | All | 9.x |
| `FLASH_ATTN` | FA4* | fp16, bf16 | `auto`, `float16`, `bfloat16` | %16 | Any | ✅ | ✅ | ❌ | ✅ | All | ≥10.0 |
@@ -187,18 +187,6 @@ Priority is **1 = highest** (tried first).
>
> **\*** Specify the FlashAttention version via `--attention-config.flash_attn_version=2`, `3`, or `4`. Default is FA4 on SM100+ (Blackwell), FA3 on SM90 (Hopper), FA2 otherwise.
## MiniMax M3 Sparse Attention Backends
Block-sparse GQA backend used by MiniMax M3 sparse ("lightning indexer")
layers. It is wired in directly by the model and is not part of the
automatic priority lists above. A lightning indexer scores KV blocks, the
top-k blocks (plus fixed init/local blocks) are selected, and attention
attends only to those blocks; index keys live in a separate side cache.
| Backend | Dtypes | KV Dtypes | Block Sizes | Head Sizes | Sink | Non-Causal | MM Prefix | DCP | Attention Types | Compute Cap. |
| ------- | ------ | --------- | ----------- | ---------- | ---- | ---------- | --------- | --- | --------------- | ------------ |
| `MINIMAX_M3_SPARSE` | bf16, fp16 | `bfloat16` | 128 | 128 | ❌ | ❌ | ❌ | ❌ | Decoder | Any |
## MLA (Multi-head Latent Attention) Backends
MLA uses separate backends for prefill and decode phases.
+1 -1
View File
@@ -778,7 +778,7 @@ Then, you can use the OpenAI client as follows:
base_url=openai_api_base,
)
video_url = "http://commondatastorage.googleapis.com/gtv-videos-bucket/sample/ForBiggerFun.mp4"
video_url = "https://huggingface.co/datasets/raushan-testing-hf/videos-test/resolve/main/sample_demo_1.mp4"
## Use video url in the payload
chat_completion_from_url = client.chat.completions.create(
+1
View File
@@ -437,6 +437,7 @@ th {
| `LongcatFlashForCausalLM` | LongCat-Flash | `meituan-longcat/LongCat-Flash-Chat`, `meituan-longcat/LongCat-Flash-Chat-FP8` | ✅︎ | ✅︎ |
| `MambaForCausalLM` | Mamba | `state-spaces/mamba-130m-hf`, `state-spaces/mamba-790m-hf`, `state-spaces/mamba-2.8b-hf`, etc. | | ✅︎ |
| `Mamba2ForCausalLM` | Mamba2 | `mistralai/Mamba-Codestral-7B-v0.1`, etc. | | ✅︎ |
| `MellumForCausalLM` | Mellum 2 | `JetBrains/Mellum2-12B-A2.5B-Base`, etc. | | ✅︎ |
| `MiMoForCausalLM` | MiMo | `XiaomiMiMo/MiMo-7B-RL`, etc. | ✅︎ | ✅︎ |
| `MiMoV2FlashForCausalLM` | MiMoV2Flash | `XiaomiMiMo/MiMo-V2-Flash`, etc. | | ✅︎ |
| `MiMoV2ForCausalLM` | MiMoV2Pro | `XiaomiMiMo/MiMo-V2.5-Pro`, etc. | | ✅︎ |
+32 -2
View File
@@ -100,14 +100,44 @@ For further details on renderer APIs, please refer to [this page](renderer.md).
- `/version` - Version information
- `/load` - Server load metrics
## Sleep Mode APIs
## Server in development mode
When using the flag VLLM_SERVER_DEV_MODE=1, you enable development endpoints.
**SECURITY WARNING: These endpoints should NOT be used in production!**
### Cache Management APIs
- `/reset_prefix_cache` - Reset prefix cache (can disrupt service)
- `/reset_mm_cache` - Reset multimodal cache (can disrupt service)
- `/reset_encoder_cache` - Reset encoder cache (can disrupt service)
### Weight Transfer APIs (RL Training)
For further details on Weight Transfer, please refer to [this page](../../training/weight_transfer/README.md).
- `/pause` - Pause generation (causes denial of service)
- `/resume` - Resume generation
- `/is_paused` - Check if generation is paused
- `/init_weight_transfer_engine` - Initialize weight transfer engine for RLHF
- `/update_weights` - Update model weights (can alter model behavior)
- `/get_world_size` - Get distributed world size
### Collective RPC
- `/collective_rpc` - Execute arbitrary RPC methods on the engine (extremely dangerous)
### Server info
- `/server_info` - Get detailed server configuration
### Sleep Mode APIs
For further details on sleep mode, please refer to [this page](../../features/sleep_mode.md).
- `/sleep` - Put engine to sleep (causes denial of service)
- `/wake_up` - Wake engine from sleep
- `/is_sleeping` - Check if engine is sleeping
- `/collective_rpc` - Execute arbitrary RPC methods on the engine (extremely dangerous)
## Chat Template
+18 -3
View File
@@ -84,7 +84,10 @@ Both the trainer (`NCCLTrainerSendWeightsArgs`) and inference side (`NCCLWeightT
## Receiving Weights (Inference Side)
The inference side triggers weight reception using the four-phase protocol`init_weight_transfer_engine`, `start_weight_update`, `update_weights`, `finish_weight_update`. The init phase is shown [above](#initialization). The remaining three steps are:
The inference side triggers weight reception using the four-phase protocol:
`init_weight_transfer_engine`, `start_weight_update`, `update_weights`,
`finish_weight_update`. The init phase is shown [above](#initialization). The
remaining three steps are:
```python
from vllm.distributed.weight_transfer.base import WeightTransferUpdateRequest
@@ -108,12 +111,24 @@ llm.update_weights(
llm.finish_weight_update()
```
The `names`, `dtype_names`, and `shapes` lists describe each parameter. These must match the order in which the trainer iterates over its parameters.
The `names`, `dtype_names`, and `shapes` lists describe each parameter. These
must match the order in which the trainer iterates over its parameters.
`start_weight_update` must be called before `update_weights`, and `finish_weight_update` must be called after all weight chunks have been transferred. The `is_checkpoint_format` flag controls whether layerwise reload processing is applied (`True` for checkpoint-format weights, `False` for pre-processed kernel-format weights).
`start_weight_update` must be called before `update_weights`, and
`finish_weight_update` must be called after all weight chunks have been
transferred. The `is_checkpoint_format` flag controls whether layerwise reload
processing is applied (`True` for checkpoint-format weights, `False` for
pre-processed kernel-format weights).
Sparse NCCL patches still use `update_kind="sparse_flat"` inside
`update_info`, but they should be wrapped in
`start_weight_update(is_checkpoint_format=False)` because sparse patches apply
directly to runtime/kernel-format parameters. The current sparse MVP requires
`TP=1` and `PP=1`.
## Examples
- [RLHF with NCCL weight syncing (offline, Ray)](../../../examples/rl/rlhf_nccl.py) - Trainer on one GPU, 2x tensor-parallel vLLM engine on two others, with packed NCCL weight broadcast
- [RLHF with sparse NCCL weight syncing (offline, Ray)](../../../examples/rl/rlhf_sparse_nccl.py) - Dense-vs-sparse equivalence demo with a real model on a 2-GPU trainer/inference setup; sparse patches use `start_weight_update(is_checkpoint_format=False)` and currently require `TP=1` and `PP=1`
- [RLHF with async weight syncing (offline, Ray)](../../../examples/rl/rlhf_async_new_apis.py) - Async generation with mid-flight pause, weight sync, resume, and validation against a fresh model
- [RLHF with NCCL weight syncing (online serving, HTTP)](../../../examples/rl/rlhf_http_nccl.py) - Weight transfer with a running vLLM HTTP server using HTTP control plane and NCCL data plane
@@ -203,7 +203,7 @@ def run_multi_image(model: str, max_completion_tokens: int) -> None:
# Video input inference
def run_video(model: str, max_completion_tokens: int) -> None:
video_url = "http://commondatastorage.googleapis.com/gtv-videos-bucket/sample/ForBiggerFun.mp4"
video_url = "https://huggingface.co/datasets/raushan-testing-hf/videos-test/resolve/main/sample_demo_1.mp4"
video_base64 = encode_base64_content_from_url(video_url)
## Use video url in the payload
+526
View File
@@ -0,0 +1,526 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
"""
Demonstrates dense-vs-sparse NCCL weight syncing with a real model.
This example mirrors the validation story used for the sparse NCCL MVP:
both the dense update path and the sparse patch path start from the same real
checkpoint and apply the same deterministic trainer-side patch. The script then
checks that greedy 1-token outputs match between the dense and sparse vLLM
engines after the update.
The example performs the following steps:
* Load a training model on one GPU via a Ray actor.
* Launch a vLLM engine with the same real model on a second GPU.
* Verify trainer vs vLLM baseline agreement before any update.
* Apply a deterministic patch to ``model.embed_tokens.weight`` on the trainer.
* Run a dense NCCL update into a fresh vLLM engine and collect post-update
outputs.
* Reset the trainer back to the baseline checkpoint.
* Apply the same deterministic patch again.
* Run a sparse NCCL update into another fresh vLLM engine and collect
post-update outputs.
* Compare dense vs sparse baseline outputs, dense vs sparse post-update
outputs, estimated payload sizes, and trainer-side send times.
Current sparse weight transfer MVP limitations:
* ``TP=1`` and ``PP=1`` only
* sparse updates use runtime/kernel-format parameter names
* sparse updates are not composable with checkpoint-format or packed updates
This example assumes a single-node cluster with two GPUs.
"""
import hashlib
import os
import time
from collections.abc import Sequence
import ray
import torch
from ray.util.placement_group import placement_group
from ray.util.scheduling_strategies import PlacementGroupSchedulingStrategy
from transformers import AutoModelForCausalLM, AutoTokenizer
from vllm import LLM, SamplingParams
from vllm.config import WeightTransferConfig
from vllm.distributed.weight_transfer.base import SparseWeightPatch
from vllm.distributed.weight_transfer.nccl_engine import (
NCCLTrainerSendWeightsArgs,
NCCLWeightTransferEngine,
)
from vllm.utils.network_utils import get_ip, get_open_port
MODEL_NAME = "Qwen/Qwen2.5-0.5B-Instruct"
PATCHED_PARAM_NAME = "model.embed_tokens.weight"
MAX_PATCH_ROWS = 32
PROMPTS = [
"Hello, my name is",
"The president of the United States is",
"The capital of France is",
"The future of AI is",
]
SAMPLING_PARAMS = SamplingParams(temperature=0.0, max_tokens=1)
class MyLLM(LLM):
"""Configure the vLLM worker for Ray placement group execution."""
def __init__(self, *args, **kwargs):
os.environ["VLLM_RAY_BUNDLE_INDICES"] = "0"
super().__init__(*args, **kwargs)
@ray.remote(num_gpus=1)
class TrainModel:
"""Ray actor that owns the trainer-side model and deterministic patch state."""
def __init__(self, model_name: str):
self.model_name = model_name
self.tokenizer = AutoTokenizer.from_pretrained(model_name)
if self.tokenizer.pad_token_id is None:
self.tokenizer.pad_token = self.tokenizer.eos_token
self.model = None
self.patched_param = None
self.pending_sparse_patches: list[SparseWeightPatch] | None = None
self.model_update_group = None
self.master_address = get_ip()
self.port = get_open_port()
self.reset_model()
def reset_model(self) -> None:
self.model = AutoModelForCausalLM.from_pretrained(
self.model_name,
torch_dtype=torch.bfloat16,
).to("cuda:0")
self.model.eval()
try:
self.patched_param = self.model.get_parameter(PATCHED_PARAM_NAME)
except AttributeError as exc:
raise RuntimeError(
f"Expected trainer model to expose `{PATCHED_PARAM_NAME}`"
) from exc
self.pending_sparse_patches = None
def create_rendezvous(self) -> tuple[str, int]:
self.port = get_open_port()
return self.master_address, self.port
def init_weight_transfer_group(self, world_size: int) -> None:
self.model_update_group = NCCLWeightTransferEngine.trainer_init(
dict(
master_address=self.master_address,
master_port=self.port,
world_size=world_size,
)
)
def get_dense_update_info(self, packed: bool = False) -> tuple[dict, int]:
names = []
dtype_names = []
shapes = []
payload_bytes = 0
for name, param in self.model.named_parameters():
names.append(name)
dtype_names.append(str(param.dtype).split(".")[-1])
shapes.append(list(param.shape))
payload_bytes += param.numel() * param.element_size()
return (
dict(
names=names,
dtype_names=dtype_names,
shapes=shapes,
packed=packed,
),
payload_bytes,
)
@torch.inference_mode()
def generate(
self,
prompts: Sequence[str],
max_new_tokens: int = 1,
) -> list[dict[str, object]]:
generations = []
for prompt in prompts:
model_inputs = self.tokenizer(prompt, return_tensors="pt").to("cuda:0")
output = self.model.generate(
**model_inputs,
max_new_tokens=max_new_tokens,
do_sample=False,
pad_token_id=self.tokenizer.pad_token_id,
)
new_token_ids = output[0, model_inputs["input_ids"].shape[1] :].tolist()
generations.append(
{
"token_ids": new_token_ids,
"text": self.tokenizer.decode(
new_token_ids,
skip_special_tokens=False,
),
}
)
return generations
def prepare_sparse_patch(
self,
prompts: Sequence[str],
max_patch_rows: int = MAX_PATCH_ROWS,
) -> tuple[dict[str, object], list[int], str, int]:
selected_token_ids: list[int] = []
special_ids = set(self.tokenizer.all_special_ids)
for prompt in prompts:
token_ids = self.tokenizer(prompt, add_special_tokens=False)["input_ids"]
for token_id in token_ids:
if token_id in special_ids or token_id in selected_token_ids:
continue
selected_token_ids.append(token_id)
if len(selected_token_ids) == max_patch_rows:
break
if len(selected_token_ids) == max_patch_rows:
break
if not selected_token_ids:
raise ValueError("Could not derive any non-special token IDs to patch")
vocab_size = self.patched_param.shape[0]
next_token_id = selected_token_ids[-1]
while len(selected_token_ids) < max_patch_rows:
next_token_id = (next_token_id + 1) % vocab_size
if next_token_id in special_ids or next_token_id in selected_token_ids:
continue
selected_token_ids.append(next_token_id)
row_ids = torch.tensor(
selected_token_ids,
device=self.patched_param.device,
dtype=torch.long,
)
hidden_size = self.patched_param.shape[1]
column_offsets = torch.arange(
hidden_size,
device=self.patched_param.device,
dtype=torch.long,
)
with torch.no_grad():
# Rotate the selected embedding rows instead of zeroing them so the
# patch remains deterministic while avoiding a degenerate collapse
# to the same special token after the update.
replacement_rows = self.patched_param[row_ids].roll(shifts=1, dims=0)
self.patched_param[row_ids] = replacement_rows
flat_indices = (
row_ids.unsqueeze(1).mul(hidden_size).add(column_offsets).reshape(-1)
)
flat_values = self.patched_param[row_ids].reshape(-1).contiguous()
self.pending_sparse_patches = [
SparseWeightPatch(
name=PATCHED_PARAM_NAME,
indices=flat_indices.to(torch.int32),
values=flat_values,
)
]
patch_digest = hashlib.sha256(
self.pending_sparse_patches[0].indices.cpu().numpy().tobytes()
+ self.pending_sparse_patches[0]
.values.detach()
.float()
.cpu()
.numpy()
.tobytes()
).hexdigest()
sparse_payload_bytes = (
flat_indices.numel() * torch.tensor([], dtype=torch.int32).element_size()
+ flat_values.numel() * flat_values.element_size()
)
update_info = dict(
names=[PATCHED_PARAM_NAME],
dtype_names=[str(self.patched_param.dtype).split(".")[-1]],
shapes=[list(self.patched_param.shape)],
num_updates_list=[flat_indices.numel()],
update_kind="sparse_flat",
)
return update_info, selected_token_ids, patch_digest, sparse_payload_bytes
def broadcast_weights(self, packed: bool = False) -> float:
if self.model_update_group is None:
raise RuntimeError("Weight transfer group is not initialized")
trainer_args = NCCLTrainerSendWeightsArgs(
group=self.model_update_group,
packed=packed,
)
start = time.perf_counter()
NCCLWeightTransferEngine.trainer_send_weights(
iterator=self.model.named_parameters(),
trainer_args=trainer_args,
)
torch.accelerator.synchronize()
return (time.perf_counter() - start) * 1000.0
def broadcast_pending_sparse_patch(self) -> float:
if self.model_update_group is None:
raise RuntimeError("Weight transfer group is not initialized")
if self.pending_sparse_patches is None:
raise RuntimeError("Sparse patch has not been prepared")
start = time.perf_counter()
NCCLWeightTransferEngine.trainer_send_sparse_weights(
iter(self.pending_sparse_patches),
NCCLTrainerSendWeightsArgs(group=self.model_update_group),
)
torch.accelerator.synchronize()
self.pending_sparse_patches = None
return (time.perf_counter() - start) * 1000.0
def launch_llm(
scheduling_inference: PlacementGroupSchedulingStrategy,
):
return ray.remote(
num_cpus=0,
num_gpus=0,
scheduling_strategy=scheduling_inference,
)(MyLLM).remote(
model=MODEL_NAME,
enforce_eager=True,
tensor_parallel_size=1,
distributed_executor_backend="ray",
gpu_memory_utilization=0.7,
weight_transfer_config=WeightTransferConfig(backend="nccl"),
)
def collect_vllm_generations(llm_handle) -> list[dict[str, object]]:
outputs = ray.get(llm_handle.generate.remote(PROMPTS, SAMPLING_PARAMS))
generations = []
for output in outputs:
generations.append(
{
"token_ids": output.outputs[0].token_ids,
"text": output.outputs[0].text,
}
)
return generations
def token_sequences_match(
left: Sequence[dict[str, object]],
right: Sequence[dict[str, object]],
) -> bool:
return [item["token_ids"] for item in left] == [item["token_ids"] for item in right]
def print_generations(label: str, prompts: Sequence[str], generations) -> None:
print(f"\n{label}")
print("-" * 50)
for prompt, generation in zip(prompts, generations):
print(f"Prompt: {prompt!r}")
print(f"Token IDs: {generation['token_ids']}")
print(f"Text: {generation['text']!r}")
print("-" * 50)
def run_dense_phase(
train_model,
scheduling_inference: PlacementGroupSchedulingStrategy,
) -> dict[str, object]:
ray.get(train_model.reset_model.remote())
llm = launch_llm(scheduling_inference)
try:
dense_before = collect_vllm_generations(llm)
ray.get(llm.sleep.remote(level=0))
master_address, master_port = ray.get(train_model.create_rendezvous.remote())
world_size = ray.get(llm.get_world_size.remote()) + 1
inference_init = llm.init_weight_transfer_engine.remote(
dict(
init_info=dict(
master_address=master_address,
master_port=master_port,
rank_offset=1,
world_size=world_size,
)
)
)
trainer_init = train_model.init_weight_transfer_group.remote(world_size)
ray.get([trainer_init, inference_init])
ray.get(llm.start_weight_update.remote(is_checkpoint_format=True))
dense_update_info, dense_payload_bytes = ray.get(
train_model.get_dense_update_info.remote()
)
_, selected_token_ids, patch_digest, _ = ray.get(
train_model.prepare_sparse_patch.remote(PROMPTS)
)
inference_update = llm.update_weights.remote(
dict(update_info=dense_update_info)
)
dense_send_ms, _ = ray.get(
[
train_model.broadcast_weights.remote(packed=False),
inference_update,
]
)
ray.get(llm.finish_weight_update.remote())
ray.get(llm.wake_up.remote(tags=["scheduling"]))
dense_after = collect_vllm_generations(llm)
return {
"dense_before": dense_before,
"dense_after": dense_after,
"selected_token_ids": selected_token_ids,
"patch_digest": patch_digest,
"dense_payload_bytes": dense_payload_bytes,
"dense_send_ms": dense_send_ms,
}
finally:
ray.kill(llm)
def run_sparse_phase(
train_model,
scheduling_inference: PlacementGroupSchedulingStrategy,
) -> dict[str, object]:
ray.get(train_model.reset_model.remote())
llm = launch_llm(scheduling_inference)
try:
sparse_before = collect_vllm_generations(llm)
ray.get(llm.sleep.remote(level=0))
master_address, master_port = ray.get(train_model.create_rendezvous.remote())
world_size = ray.get(llm.get_world_size.remote()) + 1
inference_init = llm.init_weight_transfer_engine.remote(
dict(
init_info=dict(
master_address=master_address,
master_port=master_port,
rank_offset=1,
world_size=world_size,
)
)
)
trainer_init = train_model.init_weight_transfer_group.remote(world_size)
ray.get([trainer_init, inference_init])
ray.get(llm.start_weight_update.remote(is_checkpoint_format=False))
sparse_update_info, selected_token_ids, patch_digest, sparse_payload_bytes = (
ray.get(train_model.prepare_sparse_patch.remote(PROMPTS))
)
inference_update = llm.update_weights.remote(
dict(update_info=sparse_update_info)
)
sparse_send_ms, _ = ray.get(
[
train_model.broadcast_pending_sparse_patch.remote(),
inference_update,
]
)
ray.get(llm.finish_weight_update.remote())
ray.get(llm.wake_up.remote(tags=["scheduling"]))
sparse_after = collect_vllm_generations(llm)
return {
"sparse_before": sparse_before,
"sparse_after": sparse_after,
"selected_token_ids": selected_token_ids,
"patch_digest": patch_digest,
"sparse_payload_bytes": sparse_payload_bytes,
"sparse_send_ms": sparse_send_ms,
}
finally:
ray.kill(llm)
ray.init()
try:
train_model = TrainModel.remote(MODEL_NAME)
pg_inference = placement_group([{"GPU": 1, "CPU": 0}])
ray.get(pg_inference.ready())
scheduling_inference = PlacementGroupSchedulingStrategy(
placement_group=pg_inference,
placement_group_capture_child_tasks=True,
placement_group_bundle_index=0,
)
dense_results = run_dense_phase(train_model, scheduling_inference)
sparse_results = run_sparse_phase(train_model, scheduling_inference)
baseline_equal = token_sequences_match(
dense_results["dense_before"],
sparse_results["sparse_before"],
)
patch_selection_equal = (
dense_results["selected_token_ids"] == sparse_results["selected_token_ids"]
)
patch_digest_equal = dense_results["patch_digest"] == sparse_results["patch_digest"]
after_equal = token_sequences_match(
dense_results["dense_after"],
sparse_results["sparse_after"],
)
any_output_changed = any(
before["token_ids"] != after["token_ids"]
for before, after in zip(
dense_results["dense_before"],
dense_results["dense_after"],
)
)
dense_payload_mb = dense_results["dense_payload_bytes"] / (1024 * 1024)
sparse_payload_mb = sparse_results["sparse_payload_bytes"] / (1024 * 1024)
print_generations(
"Dense baseline outputs",
PROMPTS,
dense_results["dense_before"],
)
print_generations(
"Sparse baseline outputs", PROMPTS, sparse_results["sparse_before"]
)
print_generations(
"Dense outputs after update", PROMPTS, dense_results["dense_after"]
)
print_generations(
"Sparse outputs after update",
PROMPTS,
sparse_results["sparse_after"],
)
print(f"patched_token_ids = {dense_results['selected_token_ids']}")
print(f"patch_selection_equal = {patch_selection_equal}")
print(f"dense_patch_digest = {dense_results['patch_digest']}")
print(f"sparse_patch_digest = {sparse_results['patch_digest']}")
print(f"patch_digest_equal = {patch_digest_equal}")
print(f"baseline_equal = {baseline_equal}")
print(f"after_equal = {after_equal}")
print(f"any_output_changed = {any_output_changed}")
print(f"dense_payload_mb = {dense_payload_mb:.2f}")
print(f"sparse_payload_mb = {sparse_payload_mb:.2f}")
print(f"dense_send_ms = {dense_results['dense_send_ms']:.2f}")
print(f"sparse_send_ms = {sparse_results['sparse_send_ms']:.2f}")
if not baseline_equal:
raise RuntimeError(
"Dense and sparse phases did not start from the same baseline"
)
if not patch_selection_equal:
raise RuntimeError("Dense and sparse phases used different sparse patches")
if not patch_digest_equal:
raise RuntimeError("Dense and sparse phases produced different patch values")
if not after_equal:
raise RuntimeError("Dense and sparse updates produced different outputs")
if not any_output_changed:
raise RuntimeError("Patch did not change the observed outputs")
finally:
ray.shutdown()
-5
View File
@@ -162,12 +162,7 @@ dout = "dout"
Pn = "Pn"
arange = "arange"
thw = "thw"
# temporal position ids (parallels hpos/wpos in vision RoPE)
tpos = "tpos"
subtile = "subtile"
subtiles = "subtiles"
reord = "reord"
Ot = "Ot"
HSA = "HSA"
setp = "setp"
CPY = "CPY"
-1
View File
@@ -29,7 +29,6 @@ xgrammar >= 0.2.0, < 1.0.0; platform_machine == "x86_64" or platform_machine ==
typing_extensions >= 4.10
filelock >= 3.16.1 # need to contain https://github.com/tox-dev/filelock/pull/317
partial-json-parser # used for parsing partial JSON outputs
jsonschema >= 4.23.0 # required for MiniMax M3 tool schema validation
pyzmq >= 25.0.0
msgspec
gguf >= 0.17.0
+2 -2
View File
@@ -9,8 +9,8 @@ torchaudio==2.11.0
# These must be updated alongside torch
torchvision==0.26.0 # Required for phi3v processor. See https://github.com/pytorch/vision?tab=readme-ov-file#installation for corresponding version
# FlashInfer should be updated together with the Dockerfile
flashinfer-python==0.6.12
flashinfer-cubin==0.6.12
flashinfer-python==0.6.11.post2
flashinfer-cubin==0.6.11.post2
apache-tvm-ffi==0.1.9
tilelang==0.1.9
# Cap nvidia-cudnn-frontend (transitive dep of flashinfer) due to
-1
View File
@@ -360,7 +360,6 @@ jsonpointer==3.0.0
# via jsonschema
jsonschema==4.23.0
# via
# -c requirements/common.txt
# hypothesis-jsonschema
# mistral-common
# ray
-2
View File
@@ -440,8 +440,6 @@ jsonpointer==3.1.0
# via jsonschema
jsonschema==4.26.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# hypothesis-jsonschema
# mcp
# mistral-common
-1
View File
@@ -229,7 +229,6 @@ jsonlines==4.0.0
# via lm-eval
jsonschema==4.26.0
# via
# -c requirements/common.txt
# hypothesis-jsonschema
# mistral-common
# schemathesis
-87
View File
@@ -3458,75 +3458,6 @@ version = "0.1.29"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "e0c5ccf5294c6ccd63a74f1565028353830a9c2f5eb0c682c355c471726a6e3f"
[[package]]
name = "pyo3"
version = "0.28.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "91fd8e38a3b50ed1167fb981cd6fd60147e091784c427b8f7183a7ee32c31c12"
dependencies = [
"libc",
"once_cell",
"portable-atomic",
"pyo3-build-config",
"pyo3-ffi",
"pyo3-macros",
]
[[package]]
name = "pyo3-build-config"
version = "0.28.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "e368e7ddfdeb98c9bca7f8383be1648fd84ab466bf2bc015e94008db6d35611e"
dependencies = [
"target-lexicon",
]
[[package]]
name = "pyo3-ffi"
version = "0.28.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "7f29e10af80b1f7ccaf7f69eace800a03ecd13e883acfacc1e5d0988605f651e"
dependencies = [
"libc",
"pyo3-build-config",
]
[[package]]
name = "pyo3-macros"
version = "0.28.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "df6e520eff47c45997d2fc7dd8214b25dd1310918bbb2642156ef66a67f29813"
dependencies = [
"proc-macro2",
"pyo3-macros-backend",
"quote",
"syn 2.0.117",
]
[[package]]
name = "pyo3-macros-backend"
version = "0.28.3"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "c4cdc218d835738f81c2338f822078af45b4afdf8b2e33cbb5916f108b813acb"
dependencies = [
"heck",
"proc-macro2",
"pyo3-build-config",
"quote",
"syn 2.0.117",
]
[[package]]
name = "pythonize"
version = "0.28.0"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "0b79f670c9626c8b651c0581011b57b6ba6970bb69faf01a7c4c0cfc81c43f95"
dependencies = [
"pyo3",
"serde",
"serde_json",
]
[[package]]
name = "qoi"
version = "0.4.1"
@@ -4738,12 +4669,6 @@ dependencies = [
"libc",
]
[[package]]
name = "target-lexicon"
version = "0.13.5"
source = "registry+https://github.com/rust-lang/crates.io-index"
checksum = "adb6935a6f5c20170eeceb1a3835a49e12e19d792f6dd344ccc76a985ca5a6ca"
[[package]]
name = "task-local"
version = "0.1.1"
@@ -5697,7 +5622,6 @@ dependencies = [
"expect-test",
"futures",
"half",
"indexmap 2.13.0",
"itertools 0.14.0",
"llm-multimodal",
"minijinja",
@@ -5976,17 +5900,6 @@ dependencies = [
"winnow",
]
[[package]]
name = "vllm-tool-parser-py"
version = "0.1.0"
dependencies = [
"pyo3",
"pythonize",
"serde_json",
"thiserror-ext",
"vllm-tool-parser",
]
[[package]]
name = "walkdir"
version = "2.5.0"
+1 -5
View File
@@ -12,7 +12,6 @@ members = [
"src/text",
"src/tokenizer",
"src/tool-parser",
"src/tool-parser/python",
]
resolver = "3"
@@ -44,7 +43,6 @@ half = { version = "2.7.1", features = ["bytemuck"] }
hex = "0.4.3"
hf-hub = { version = "0.5.0", features = ["tokio"] }
http-body = "1.0.1"
indexmap = "2.13.0"
itertools = "0.14.0"
libc = "0.2.177"
llm-multimodal = { git = "https://github.com/vllm-project/llm-multimodal", rev = "5b558989844d1c7af3e43d0f604069ffd9c06320" }
@@ -61,8 +59,6 @@ prometheus-client = "0.24.0"
prometheus-client-derive-encode = "0.5.0"
prost = "0.14.3"
prost-types = "0.14.3"
pyo3 = "0.28.3"
pythonize = "0.28.0"
rand = "0.9.2"
reasoning-parser = "1.2.2"
reqwest = { version = "0.12.8", default-features = false, features = ["rustls-tls"] }
@@ -73,7 +69,7 @@ rustc-hash = "1.1.0"
serde = { version = "1.0.228", features = ["derive"] }
serde-json-fmt = "0.1.0"
serde_default = "0.2.0"
serde_json = { version = "1.0.145", features = ["preserve_order"] }
serde_json = { version = "1.0.145", features = ["arbitrary_precision", "preserve_order"] }
serde_repr = "0.1.20"
serde_tuple = "1.1.3"
serde_with = "3.18.0"
-1
View File
@@ -10,7 +10,6 @@ asynk-strim-attr.workspace = true
easy-ext.workspace = true
futures.workspace = true
half.workspace = true
indexmap.workspace = true
itertools.workspace = true
llm-multimodal.workspace = true
minijinja.workspace = true
+1 -1
View File
@@ -233,7 +233,7 @@ mod tests {
)
.unwrap_err();
expect_test::expect!["tool parser `definitely_missing_tool_parser` is not registered (choose from: deepseek_v3, deepseek_v31, deepseek_v32, deepseek_v4, gemma4, glm45, glm47, hermes, hy_v3, kimi_k2, llama3_json, llama4_json, minimax_m2, minimax_m3, mistral, qwen3_coder, qwen3_xml)"].assert_eq(&error.to_report_string());
expect_test::expect!["tool parser `definitely_missing_tool_parser` is not registered (choose from: deepseek_v3, deepseek_v31, deepseek_v32, deepseek_v4, gemma4, glm45, glm47, hermes, hy_v3, internlm, kimi_k2, llama3_json, llama4_json, minimax_m2, mistral, qwen3_coder, qwen3_xml)"].assert_eq(&error.to_report_string());
}
#[test]
+2 -7
View File
@@ -5,9 +5,8 @@ use std::sync::LazyLock;
pub use vllm_reasoning_parser::{
CohereCmdReasoningParser, DeepSeekR1ReasoningParser, DeepSeekV3ReasoningParser,
DeepSeekV4ReasoningParser, Gemma4ReasoningParser, Glm45ReasoningParser, KimiK2ReasoningParser,
KimiReasoningParser, MiniMaxM2ReasoningParser, MiniMaxM3ReasoningParser,
NemotronV3ReasoningParser, Qwen3ReasoningParser, ReasoningDelta, ReasoningError,
ReasoningParser, Step3ReasoningParser,
KimiReasoningParser, MiniMaxM2ReasoningParser, NemotronV3ReasoningParser, Qwen3ReasoningParser,
ReasoningDelta, ReasoningError, ReasoningParser, Step3ReasoningParser,
};
use vllm_tokenizer::DynTokenizer;
@@ -24,7 +23,6 @@ pub mod names {
pub const KIMI: &str = "kimi";
pub const KIMI_K2: &str = "kimi_k2";
pub const MINIMAX_M2: &str = "minimax_m2";
pub const MINIMAX_M3: &str = "minimax_m3";
pub const NEMOTRON_V3: &str = "nemotron_v3";
pub const QWEN3: &str = "qwen3";
pub const STEP3: &str = "step3";
@@ -61,7 +59,6 @@ impl ReasoningParserFactory {
.register_parser::<KimiReasoningParser>(names::KIMI)
.register_parser::<KimiK2ReasoningParser>(names::KIMI_K2)
.register_parser::<MiniMaxM2ReasoningParser>(names::MINIMAX_M2)
.register_parser::<MiniMaxM3ReasoningParser>(names::MINIMAX_M3)
.register_parser::<NemotronV3ReasoningParser>(names::NEMOTRON_V3)
.register_parser::<Qwen3ReasoningParser>(names::QWEN3)
.register_parser::<Step3ReasoningParser>(names::STEP3);
@@ -81,8 +78,6 @@ impl ReasoningParserFactory {
.register_pattern("kimi-k2", names::KIMI_K2)
.register_pattern("kimi", names::KIMI)
.register_pattern("step3", names::STEP3)
.register_pattern("minimax-m3", names::MINIMAX_M3)
.register_pattern("mm-m3", names::MINIMAX_M3)
.register_pattern("minimax", names::MINIMAX_M2)
.register_pattern("mm-m2", names::MINIMAX_M2)
.register_pattern("cohere", names::COHERE_CMD)
@@ -32,10 +32,8 @@ fn factory_contains_and_lists_registered_parsers() {
let factory = ReasoningParserFactory::new();
assert!(factory.contains(names::QWEN3));
assert!(factory.contains(names::DEEPSEEK_V4));
assert!(factory.contains(names::MINIMAX_M3));
assert!(factory.list().contains(&names::QWEN3.to_string()));
assert!(factory.list().contains(&names::DEEPSEEK_V4.to_string()));
assert!(factory.list().contains(&names::MINIMAX_M3.to_string()));
}
#[test]
@@ -51,19 +49,6 @@ fn factory_resolves_deepseek_v4_to_qwen3_alias() {
);
}
#[test]
fn factory_resolves_minimax_m3_before_generic_minimax() {
let factory = ReasoningParserFactory::new();
assert_eq!(
factory.resolve_name_for_model("MiniMaxAI/Minimax-M3-preview"),
Some(names::MINIMAX_M3)
);
assert_eq!(
factory.resolve_name_for_model("mm-m3"),
Some(names::MINIMAX_M3)
);
}
#[test]
fn factory_rejects_unknown_parser_names() {
let tokenizer = Arc::new(FakeTokenizer);
+11 -5
View File
@@ -5,7 +5,7 @@ use std::sync::LazyLock;
pub use vllm_tool_parser::{
DeepSeekV3ToolParser, DeepSeekV4ToolParser, DeepSeekV31ToolParser, DeepSeekV32ToolParser,
Gemma4ToolParser, Glm45MoeToolParser, Glm47MoeToolParser, HermesToolParser, HyV3ToolParser,
KimiK2ToolParser, Llama3JsonToolParser, MinimaxM2ToolParser, MinimaxM3ToolParser,
Internlm2ToolParser, KimiK2ToolParser, Llama3JsonToolParser, MinimaxM2ToolParser,
MistralToolParser, Qwen3CoderToolParser, Qwen3XmlToolParser, ToolCallDelta, ToolParser,
ToolParserError, ToolParserOutput,
};
@@ -24,11 +24,13 @@ pub mod names {
pub const GEMMA4: &str = "gemma4";
pub const HERMES: &str = "hermes";
pub const HY_V3: &str = "hy_v3";
// Matches the Python CLI name `--tool-call-parser internlm`, which Python
// also routes to `Internlm2ToolParser` despite the version-agnostic name.
pub const INTERNLM: &str = "internlm";
pub const KIMI_K2: &str = "kimi_k2";
pub const LLAMA3_JSON: &str = "llama3_json";
pub const LLAMA4_JSON: &str = "llama4_json";
pub const MINIMAX_M2: &str = "minimax_m2";
pub const MINIMAX_M3: &str = "minimax_m3";
pub const MISTRAL: &str = "mistral";
pub const QWEN3_CODER: &str = "qwen3_coder";
pub const QWEN3_XML: &str = "qwen3_xml";
@@ -63,11 +65,11 @@ impl ToolParserFactory {
.register_parser::<Gemma4ToolParser>(names::GEMMA4)
.register_parser::<HermesToolParser>(names::HERMES)
.register_parser::<HyV3ToolParser>(names::HY_V3)
.register_parser::<Internlm2ToolParser>(names::INTERNLM)
.register_parser::<KimiK2ToolParser>(names::KIMI_K2)
.register_parser::<Llama3JsonToolParser>(names::LLAMA3_JSON)
.register_parser::<Llama3JsonToolParser>(names::LLAMA4_JSON)
.register_parser::<MinimaxM2ToolParser>(names::MINIMAX_M2)
.register_parser::<MinimaxM3ToolParser>(names::MINIMAX_M3)
.register_parser::<MistralToolParser>(names::MISTRAL)
.register_parser::<Qwen3XmlToolParser>(names::QWEN3_XML)
.register_parser::<Qwen3CoderToolParser>(names::QWEN3_CODER);
@@ -82,6 +84,12 @@ impl ToolParserFactory {
.register_pattern("hermes", names::HERMES)
.register_pattern("hy3", names::HY_V3)
.register_pattern("hy_v3", names::HY_V3)
// Narrow to `internlm2` substring so it matches `internlm2-chat-7b`
// and `internlm2_5-7b-chat` but NOT `internlm-chat-7b` (InternLM v1,
// routes to Llama), `internlm3-*` (also Llama-architecture per
// vllm/model_executor/models/registry.py:146), or `Intern-S1` /
// `Intern-S1-Pro` (separate intern-s1 parser, see PR #40115).
.register_pattern("internlm2", names::INTERNLM)
.register_pattern("llama-4", names::LLAMA4_JSON)
.register_pattern("llama-3.2", names::LLAMA3_JSON)
.register_pattern("llama-3.1", names::LLAMA3_JSON)
@@ -98,8 +106,6 @@ impl ToolParserFactory {
.register_pattern("gemma4", names::GEMMA4)
.register_pattern("gemma-4", names::GEMMA4)
.register_pattern("kimi-k2", names::KIMI_K2)
.register_pattern("minimax-m3", names::MINIMAX_M3)
.register_pattern("mm-m3", names::MINIMAX_M3)
.register_pattern("minimax", names::MINIMAX_M2)
.register_pattern("mm-m2", names::MINIMAX_M2);
+29 -8
View File
@@ -153,14 +153,6 @@ fn factory_new_resolves_default_patterns() {
factory.resolve_name_for_model("tencent/Hy3-preview"),
Some(names::HY_V3)
);
assert_eq!(
factory.resolve_name_for_model("MiniMax/MiniMax-M3-Text"),
Some(names::MINIMAX_M3)
);
assert_eq!(
factory.resolve_name_for_model("org/mm-m3-base"),
Some(names::MINIMAX_M3)
);
assert_eq!(
factory.resolve_name_for_model("MiniMax/MiniMax-M2-01"),
Some(names::MINIMAX_M2)
@@ -169,4 +161,33 @@ fn factory_new_resolves_default_patterns() {
factory.resolve_name_for_model("org/mm-m2-base"),
Some(names::MINIMAX_M2)
);
// InternLM2 positive: both dashed and underscored versioned names route.
assert_eq!(
factory.resolve_name_for_model("internlm/internlm2-chat-7b"),
Some(names::INTERNLM)
);
assert_eq!(
factory.resolve_name_for_model("internlm/internlm2_5-7b-chat"),
Some(names::INTERNLM)
);
// Negative: other internlm-org models do NOT route to the InternLM2 parser,
// since they use unrelated prompt formats.
// - InternLM v1 (`internlm-chat-7b`) routes to Llama
// - InternLM3 (`internlm3-8b-instruct`) routes to Llama
// - Intern-S1 / Intern-S1-Pro have their own parser (Python PR #40115)
assert_eq!(
factory.resolve_name_for_model("internlm/internlm-chat-7b"),
None
);
assert_eq!(
factory.resolve_name_for_model("internlm/internlm3-8b-instruct"),
None
);
assert_eq!(factory.resolve_name_for_model("internlm/Intern-S1"), None);
assert_eq!(
factory.resolve_name_for_model("internlm/Intern-S1-Pro"),
None
);
}
+7 -33
View File
@@ -1,7 +1,7 @@
use std::collections::HashMap;
use serde::Serialize;
use serde_json::Value as JsonValue;
use serde_json::Value;
use thiserror_ext::AsReport as _;
use tracing::{info, trace, warn};
use vllm_text::Prompt;
@@ -13,7 +13,6 @@ use self::format::{
ChatTemplateContentFormat, ChatTemplateContentFormatOption as ContentFormatOption,
};
use self::template::{CompiledChatTemplate, TemplateContext};
use self::value::{TemplateValue, to_template_value};
use super::{ChatRenderer, RenderedPrompt};
use crate::error::Result;
use crate::request::{ChatContent, ChatContentPart, ChatMessage, ChatRequest};
@@ -25,7 +24,6 @@ mod error;
mod format;
mod template;
mod tojson;
mod value;
pub use template::{load_chat_template, resolve_chat_template};
@@ -40,7 +38,7 @@ pub struct MultimodalRenderInfo {
/// state.
pub struct HfChatRenderer {
default_template: Option<CompiledChatTemplate>,
default_template_kwargs: HashMap<String, JsonValue>,
default_template_kwargs: HashMap<String, Value>,
content_format: ContentFormatOption,
special_tokens: Option<HfSpecialTokens>,
multimodal: Option<MultimodalRenderInfo>,
@@ -50,7 +48,7 @@ impl HfChatRenderer {
/// Create a renderer from the given template string.
pub fn new(
template: Option<String>,
default_template_kwargs: HashMap<String, JsonValue>,
default_template_kwargs: HashMap<String, Value>,
content_format: ContentFormatOption,
) -> Result<Self> {
Ok(Self {
@@ -247,7 +245,7 @@ struct TemplateToolCall {
#[derive(Debug, Serialize)]
struct TemplateToolFunction {
name: String,
arguments: TemplateValue,
arguments: Value,
}
#[derive(Debug, Serialize)]
@@ -261,7 +259,7 @@ pub(super) struct TemplateTool {
struct TemplateToolDefinition {
name: String,
description: Option<String>,
parameters: TemplateValue,
parameters: Value,
strict: Option<bool>,
}
@@ -347,14 +345,13 @@ fn to_template_tool_calls(
let mut tool_calls = Vec::new();
for tool_call in content.tool_calls() {
let arguments = serde_json::from_str(&tool_call.arguments).map_err(|error| {
let arguments = serde_json::from_str::<Value>(&tool_call.arguments).map_err(|error| {
Error::ChatTemplate(format!(
"assistant tool call `{}` has invalid JSON arguments: {}",
tool_call.id,
error.as_report()
))
})?;
let arguments = to_template_value(arguments);
tool_calls.push(TemplateToolCall {
id: tool_call.id.clone(),
@@ -437,7 +434,7 @@ fn to_template_tools(tools: &[ChatTool]) -> Vec<TemplateTool> {
function: TemplateToolDefinition {
name: tool.name.clone(),
description: tool.description.clone(),
parameters: to_template_value(tool.parameters.clone()),
parameters: tool.parameters.clone(),
strict: tool.strict,
},
})
@@ -912,29 +909,6 @@ mod tests {
assert_eq!(rendered, "get_weather|Paris|call_1|Sunny");
}
#[test]
fn chat_template_tool_call_argument_items_method_is_not_shadowed_by_field() {
let request = sample_request(vec![ChatMessage::assistant_blocks(vec![
AssistantContentBlock::ToolCall(crate::AssistantToolCall {
id: "call_1".to_string(),
name: "add".to_string(),
arguments: r#"{"items":"operands","x":2,"y":1.0}"#.to_string(),
}),
])]);
let rendered = render(
Some(
"{%- set arguments = messages[0].tool_calls[0].function.arguments -%}
{%- for key, value in arguments.items() -%}{{ key }}={{ value }};{%- endfor -%}
|{{ arguments['items'] }}",
),
&request,
)
.unwrap();
assert_eq!(rendered, "items=operands;x=2;y=1.0;|operands");
}
#[test]
fn qwen35_template_renders_prefilled_reasoning_start_when_thinking_enabled() {
let mut request = sample_request(vec![ChatMessage::text(ChatRole::User, "hello")]);
+2 -18
View File
@@ -208,27 +208,11 @@ mod tests {
}
#[test]
fn tojson_uses_standard_serde_json_number_spelling() {
fn tojson_preserves_arbitrary_precision_number_spelling() {
let payload = serde_json::from_str(r#"{"x":2,"y":1.00}"#).unwrap();
let rendered = render("{{ payload|tojson }}", payload);
// TODO: we cannot preserve the original number precision by enabling `serde_json`'s
// `arbitrary_precision` feature, otherwise the following test
// `serialized_json_numbers_do_not_leak_serde_private_representation` will fail.
// See issue: https://github.com/mitsuhiko/minijinja/issues/641
assert_eq!(rendered, "{\"x\": 2, \"y\": 1.0}");
}
#[test]
fn serialized_json_numbers_do_not_leak_serde_private_representation() {
let payload: serde_json::Value = serde_json::from_str(r#"{"x":2,"y":1.00}"#).unwrap();
let rendered = render("{{ payload }}", payload);
// TODO: we cannot preserve the original number precision by enabling `serde_json`'s
// `arbitrary_precision` feature, otherwise this will fail.
// See issue: https://github.com/mitsuhiko/minijinja/issues/641
assert!(!rendered.contains("$serde_json::private::Number"));
assert_eq!(rendered, r#"{"x": 2, "y": 1.0}"#);
assert_eq!(rendered, "{\"x\": 2, \"y\": 1.00}");
}
#[test]
-77
View File
@@ -1,77 +0,0 @@
use std::sync::Arc;
use indexmap::IndexMap;
use minijinja::value::{Enumerator, Object, ObjectExt, ObjectRepr};
use minijinja::{Error as TemplateError, ErrorKind as TemplateErrorKind, State};
use serde::Serialize;
use serde_json::Value as JsonValue;
/// A wrapper around `minijinja::Value` that can be constructed with `to_template_value` and used
/// as a value in the chat template.
#[derive(Debug, Serialize)]
#[serde(transparent)]
pub(super) struct TemplateValue(minijinja::Value);
pub(super) fn to_template_value(value: JsonValue) -> TemplateValue {
TemplateValue(match value {
JsonValue::Array(values) => values
.into_iter()
.map(to_template_value)
.map(|value| value.0)
.collect::<minijinja::Value>(),
JsonValue::Object(values) => minijinja::Value::from_object(TemplateMap(
values
.into_iter()
.map(|(key, value)| (key, to_template_value(value).0))
.collect(),
)),
// For primitive values, directly convert them to `minijinja::Value` using `from_serialize`.
value => minijinja::Value::from_serialize(value),
})
}
/// A custom map type that always returns `UnknownMethod` for method calls, so that pycompat can
/// always handle dict methods through the unknown-method callback.
///
/// Use `IndexMap` to preserve the original key order when iterating.
///
/// MiniJinja's default map can resolve a same-named field before Python dict methods. HF templates
/// commonly call `dict.items()`, which would fail if the map had an `items` field.
/// See issue: https://github.com/mitsuhiko/minijinja/issues/903
#[derive(Debug)]
struct TemplateMap(IndexMap<String, minijinja::Value>);
impl Object for TemplateMap {
fn repr(self: &Arc<Self>) -> ObjectRepr {
ObjectRepr::Map
}
fn get_value(self: &Arc<Self>, key: &minijinja::Value) -> Option<minijinja::Value> {
self.0.get(key.as_str()?).cloned()
}
fn get_value_by_str(self: &Arc<Self>, key: &str) -> Option<minijinja::Value> {
self.0.get(key).cloned()
}
fn enumerate(self: &Arc<Self>) -> Enumerator {
self.mapped_rev_enumerator(|this| {
Box::new(this.0.keys().map(|key| minijinja::Value::from(key.as_str())))
})
}
fn enumerator_len(self: &Arc<Self>) -> Option<usize> {
Some(self.0.len())
}
fn call_method(
self: &Arc<Self>,
_state: &State<'_, '_>,
_method: &str,
_args: &[minijinja::Value],
) -> std::result::Result<minijinja::Value, TemplateError> {
// Always return `UnknownMethod` for method calls,
// so that pycompat can handle dict methods through the unknown-method callback.
Err(TemplateError::from(TemplateErrorKind::UnknownMethod))
}
}
+82 -33
View File
@@ -20,6 +20,7 @@ use vllm_chat::{
use vllm_text::{DecodedTextEvent, Finished, Prompt};
/// One model/parser configuration used to run the fixed roundtrip fixtures.
#[derive(Clone)]
struct RoundtripCase {
/// Hugging Face model id resolved through the production backend loader.
model_id: &'static str,
@@ -31,11 +32,45 @@ struct RoundtripCase {
tool_call_parser: ParserSelection,
/// Reasoning parser selection used by the output processor.
reasoning_parser: ParserSelection,
/// How this model's chat template handles thinking mode.
thinking_behavior: ThinkingBehavior,
/// JSON formatting expected after this model's template has materialized
/// tool-call arguments.
json_fmt: JsonFmt,
}
#[derive(Clone, Copy)]
enum ThinkingBehavior {
/// The chat template accepts explicit thinking on/off kwargs, and uses
/// `default` when the request does not specify either kwarg.
Toggleable { default: bool },
/// The chat template always behaves as `value` for this fixture.
Always { value: bool },
}
impl ThinkingBehavior {
fn default(self) -> bool {
match self {
Self::Toggleable { default } => default,
Self::Always { value } => value,
}
}
fn fixtures(self) -> Vec<Option<bool>> {
match self {
Self::Toggleable { .. } => vec![
Some(true), // explicitly enable thinking
Some(false), // explicitly disable thinking
None, // use default template behavior
],
Self::Always { value } => vec![
Some(value), // explicitly request the supported thinking behavior
None, // use default template behavior
],
}
}
}
impl RoundtripCase {
/// Qwen3 XML tool-call format with `qwen3` reasoning tags.
fn qwen3() -> Self {
@@ -44,6 +79,7 @@ impl RoundtripCase {
assistant_stop_suffix: "<|im_end|>\n",
tool_call_parser: ParserSelection::Auto,
reasoning_parser: ParserSelection::Auto,
thinking_behavior: ThinkingBehavior::Toggleable { default: true },
json_fmt: spaced_json_fmt(),
}
}
@@ -55,6 +91,7 @@ impl RoundtripCase {
assistant_stop_suffix: "<|im_end|>\n",
tool_call_parser: ParserSelection::Auto,
reasoning_parser: ParserSelection::Auto,
thinking_behavior: ThinkingBehavior::Toggleable { default: true },
json_fmt: compact_json_fmt(),
}
}
@@ -66,6 +103,7 @@ impl RoundtripCase {
assistant_stop_suffix: "[e~[\n",
tool_call_parser: ParserSelection::Auto,
reasoning_parser: ParserSelection::Auto,
thinking_behavior: ThinkingBehavior::Always { value: true },
json_fmt: compact_json_fmt(),
}
}
@@ -77,6 +115,7 @@ impl RoundtripCase {
assistant_stop_suffix: "<end▁of▁sentence>",
tool_call_parser: ParserSelection::Auto,
reasoning_parser: ParserSelection::Auto,
thinking_behavior: ThinkingBehavior::Toggleable { default: false },
json_fmt: compact_json_fmt(),
}
}
@@ -88,6 +127,7 @@ impl RoundtripCase {
assistant_stop_suffix: "",
tool_call_parser: ParserSelection::Auto,
reasoning_parser: ParserSelection::Auto,
thinking_behavior: ThinkingBehavior::Toggleable { default: true },
json_fmt: compact_json_fmt(),
}
}
@@ -100,6 +140,7 @@ impl RoundtripCase {
assistant_stop_suffix: "<|im_end|>",
tool_call_parser: ParserSelection::Auto,
reasoning_parser: ParserSelection::Auto,
thinking_behavior: ThinkingBehavior::Toggleable { default: true },
json_fmt: spaced_json_fmt(),
}
}
@@ -135,35 +176,44 @@ roundtrip_tests! {
/// Run the fixed reasoning+content fixture for one model/parser case.
async fn run_roundtrip_reasoning_and_content(case: RoundtripCase) -> Result<()> {
for thinking in case.thinking_behavior.fixtures() {
run_roundtrip_reasoning_and_content_inner(case.clone(), thinking).await?;
}
Ok(())
}
async fn run_roundtrip_reasoning_and_content_inner(
case: RoundtripCase,
thinking: Option<bool>,
) -> Result<()> {
let backends = load_roundtrip_backends(&case).await?;
let request = roundtrip_request(
"roundtrip-reasoning-content",
vec![ChatMessage::text(ChatRole::User, "What is 2 + 2?")],
Vec::new(),
thinking,
);
let expected_reasoning = "Need compute 2 + 2 directly.";
let expected_text = "The answer is 4.";
let effective_thinking = thinking.unwrap_or(case.thinking_behavior.default());
let result = run_roundtrip(
&case,
&backends,
&request,
AssistantMessage {
content: vec![
AssistantContentBlock::Reasoning {
text: expected_reasoning.to_string(),
},
AssistantContentBlock::Text {
text: expected_text.to_string(),
},
],
},
)
.await?;
let assistant = {
let mut content = Vec::new();
if effective_thinking {
content.push(AssistantContentBlock::Reasoning {
text: expected_reasoning.to_string(),
});
}
content.push(AssistantContentBlock::Text {
text: expected_text.to_string(),
});
AssistantMessage { content }
};
let result = run_roundtrip(&case, &backends, &request, assistant).await?;
assert_eq!(
result.parsed_message.reasoning().as_deref().map(str::trim),
Some(expected_reasoning)
effective_thinking.then_some(expected_reasoning)
);
assert_eq!(result.parsed_message.text().trim(), expected_text);
assert_eq!(result.parsed_message.tool_calls().count(), 0);
@@ -183,9 +233,10 @@ async fn run_roundtrip_tool_call_mix(case: RoundtripCase) -> Result<()> {
"roundtrip-reasoning-tools",
vec![ChatMessage::text(
ChatRole::User,
"Check Shanghai weather and add 1.0 plus 2.",
"Check Shanghai weather and add 1.00 plus 2.",
)],
test_tools(),
Some(true), // always enable thinking in this fixture
);
let expected_reasoning = "Need call the weather and add tools.";
let expected_text = "I will call the tools.";
@@ -210,10 +261,9 @@ async fn run_roundtrip_tool_call_mix(case: RoundtripCase) -> Result<()> {
AssistantContentBlock::ToolCall(AssistantToolCall {
id: "functions.add:1".to_string(),
name: "add".to_string(),
// Intentionally use a non-lexical order of keys to verify text-level
// fidelity of the roundtrip where JSON formatting remains stable. The
// `items` key also exercises templates that call `arguments.items()`.
arguments: r#"{"y":1.0,"x":2,"items":["left","right"]}"#.to_string(),
// Intentionally use a non-lexical order of keys and a different number
// formatting style to verify text-level fidelity of the roundtrip.
arguments: r#"{"y":1.00,"x":2}"#.to_string(),
}),
],
},
@@ -241,7 +291,7 @@ async fn run_roundtrip_tool_call_mix(case: RoundtripCase) -> Result<()> {
assert_eq!(tool_calls[1].name, "add");
assert_eq!(
tool_calls[1].arguments,
expected_arguments(&case, r#"{"y": 1.0, "x": 2, "items": ["left", "right"]}"#)?,
expected_arguments(&case, r#"{"y": 1.00, "x": 2}"#)?,
);
assert_eq!(
@@ -488,6 +538,7 @@ fn roundtrip_request(
request_id: impl Into<String>,
messages: Vec<ChatMessage>,
tools: Vec<ChatTool>,
thinking: Option<bool>,
) -> ChatRequest {
let mut request = ChatRequest {
request_id: request_id.into(),
@@ -501,10 +552,12 @@ fn roundtrip_request(
..ChatRequest::for_test()
};
// Enable thinking for some models so that rendering and parsing the reasoning block is
// exercised in the roundtrip.
for key in ["thinking", "enable_thinking"] {
request.chat_options.template_kwargs.insert(key.to_string(), true.into());
// Explicitly enable or disable thinking so that rendering and parsing the reasoning block is
// exercised or skipped in the roundtrip. If unspecified, use the default template behavior.
if let Some(thinking) = thinking {
for key in ["thinking", "enable_thinking"] {
request.chat_options.template_kwargs.insert(key.to_string(), thinking.into());
}
}
request
@@ -532,13 +585,9 @@ fn test_tools() -> Vec<ChatTool> {
"type": "object",
"properties": {
"y": { "type": "number" },
"x": { "type": "number" },
"items": {
"type": "array",
"items": { "type": "string" }
}
"x": { "type": "number" }
},
"required": ["y", "x", "items"]
"required": ["y", "x"]
}),
strict: None,
},
+11
View File
@@ -165,6 +165,15 @@ pub struct SharedRuntimeArgs {
#[serde(default)]
pub enable_log_requests: bool,
/// If specified, API server will add X-Request-Id header to responses.
#[arg(
long,
default_missing_value = "true",
num_args = 0..=1
)]
#[serde(default)]
pub enable_request_id_headers: bool,
/// Disable periodic logging of engine statistics (throughput, queue depth,
/// cache usage).
#[arg(long)]
@@ -238,6 +247,7 @@ impl SharedRuntimeArgs {
default_chat_template_kwargs: self.default_chat_template_kwargs,
chat_template_content_format: self.chat_template_content_format,
enable_log_requests: self.enable_log_requests,
enable_request_id_headers: self.enable_request_id_headers,
disable_log_stats: self.disable_log_stats,
grpc_port: self.grpc_port,
shutdown_timeout,
@@ -278,6 +288,7 @@ impl SharedRuntimeArgs {
default_chat_template_kwargs: self.default_chat_template_kwargs,
chat_template_content_format: self.chat_template_content_format,
enable_log_requests: self.enable_log_requests,
enable_request_id_headers: self.enable_request_id_headers,
disable_log_stats: self.disable_log_stats,
grpc_port: self.grpc_port,
shutdown_timeout,
+46
View File
@@ -43,6 +43,7 @@ fn serve_args_forward_python_flags_with_separator() {
default_chat_template_kwargs: None,
chat_template_content_format: Auto,
enable_log_requests: false,
enable_request_id_headers: false,
disable_log_stats: false,
served_model_name: [],
},
@@ -116,6 +117,46 @@ fn serve_args_accept_explicit_deepseek_v32_renderer() {
assert_eq!(args.runtime.renderer, RendererSelection::DeepSeekV32);
}
#[test]
fn serve_passes_enable_request_id_headers_into_config() {
let cli = Cli::try_parse_from([
"vllm-rs",
"serve",
"Qwen/Qwen3-0.6B",
"--enable-request-id-headers",
])
.unwrap();
let Command::Serve(args) = cli.command else {
panic!("expected serve args");
};
let config = args.to_frontend_config("tcp://127.0.0.1:62100".to_string());
assert!(config.enable_request_id_headers);
}
#[test]
fn frontend_args_json_passes_enable_request_id_headers_into_config() {
let cli = Cli::try_parse_from([
"vllm-rs",
"frontend",
"--listen-fd",
"3",
"--input-address",
"ipc:///tmp/input.sock",
"--output-address",
"ipc:///tmp/output.sock",
"--args-json",
r#"{"model_tag":"Qwen/Qwen3-0.6B","enable_request_id_headers":true}"#,
])
.unwrap();
let Command::Frontend(args) = cli.command else {
panic!("expected frontend args");
};
let config = args.into_config();
assert!(config.enable_request_id_headers);
}
#[test]
fn serve_args_reject_unknown_renderer_value() {
let error = Cli::try_parse_from([
@@ -218,6 +259,7 @@ fn frontend_args_accept_json() {
default_chat_template_kwargs: None,
chat_template_content_format: Auto,
enable_log_requests: false,
enable_request_id_headers: false,
disable_log_stats: false,
served_model_name: [],
},
@@ -616,6 +658,7 @@ fn serve_args_accept_handshake_aliases() {
default_chat_template_kwargs: None,
chat_template_content_format: Auto,
enable_log_requests: false,
enable_request_id_headers: false,
disable_log_stats: false,
served_model_name: [],
},
@@ -733,6 +776,7 @@ fn serve_frontend_config_uses_dp_address_as_advertised_host() {
default_chat_template_kwargs: None,
chat_template_content_format: Auto,
enable_log_requests: false,
enable_request_id_headers: false,
disable_log_stats: false,
grpc_port: None,
shutdown_timeout: 0ns,
@@ -795,6 +839,7 @@ fn serve_frontend_config_keeps_tcp_transport_for_non_local_only_topology() {
default_chat_template_kwargs: None,
chat_template_content_format: Auto,
enable_log_requests: false,
enable_request_id_headers: false,
disable_log_stats: false,
grpc_port: None,
shutdown_timeout: 0ns,
@@ -872,6 +917,7 @@ fn frontend_config_uses_external_coordinator_when_coordinator_address_is_present
default_chat_template_kwargs: None,
chat_template_content_format: Auto,
enable_log_requests: false,
enable_request_id_headers: false,
disable_log_stats: false,
grpc_port: None,
shutdown_timeout: 0ns,
-9
View File
@@ -620,15 +620,6 @@ pub struct ServerUnsupportedArgs {
#[arg(long)]
pub middleware: Option<Unsupported>,
/// If specified, API server will add X-Request-Id header to responses.
#[arg(
long,
visible_alias = "no-enable-request-id-headers",
default_missing_value = "true",
num_args = 0..=1
)]
pub enable_request_id_headers: Option<Unsupported>,
/// Disable FastAPI's OpenAPI schema, Swagger UI, and ReDoc endpoint.
#[arg(
long,
-2
View File
@@ -19,7 +19,6 @@ mod deepseek_r1;
mod delimited;
mod gemma4;
mod kimi;
mod minimax_m3;
mod qwen3;
use thiserror::Error;
@@ -30,7 +29,6 @@ pub use self::deepseek_r1::DeepSeekR1ReasoningParser;
pub(crate) use self::delimited::DelimitedReasoningParser;
pub use self::gemma4::Gemma4ReasoningParser;
pub use self::kimi::KimiReasoningParser;
pub use self::minimax_m3::MiniMaxM3ReasoningParser;
pub use self::qwen3::Qwen3ReasoningParser;
/// DeepSeek V3 currently shares the standard `<think>...</think>` parser.
@@ -1,43 +0,0 @@
use vllm_tokenizer::DynTokenizer;
use super::{DelimitedReasoningParser, ReasoningDelta, ReasoningParser, Result};
/// Reasoning parser for MiniMax M3 style outputs.
///
/// MiniMax M3 uses `<mm:think>...</mm:think>` delimiters. Its chat template may
/// prefill either delimiter depending on the requested thinking mode, so the
/// shared delimited parser derives the starting state from the rendered prompt.
pub struct MiniMaxM3ReasoningParser {
inner: DelimitedReasoningParser,
}
impl MiniMaxM3ReasoningParser {
/// Create a MiniMax M3 parser backed by the shared delimited state machine.
pub fn new(tokenizer: DynTokenizer) -> Result<Self> {
Ok(Self {
inner: DelimitedReasoningParser::new(tokenizer, "<mm:think>", "</mm:think>", false)?,
})
}
}
impl ReasoningParser for MiniMaxM3ReasoningParser {
fn create(tokenizer: DynTokenizer) -> Result<Box<dyn ReasoningParser>>
where
Self: Sized + 'static,
{
Ok(Box::new(Self::new(tokenizer)?))
}
fn initialize(&mut self, prompt_token_ids: &[u32]) -> Result<()> {
self.inner.initialize(prompt_token_ids);
Ok(())
}
fn push(&mut self, delta: &str) -> Result<ReasoningDelta> {
Ok(self.inner.push(delta))
}
fn finish(&mut self) -> Result<ReasoningDelta> {
Ok(self.inner.finish())
}
}
+1 -36
View File
@@ -3,8 +3,7 @@ use std::sync::Arc;
use vllm_tokenizer::Tokenizer;
use super::{
DeepSeekR1ReasoningParser, DelimitedReasoningParser, MiniMaxM3ReasoningParser,
Qwen3ReasoningParser, ReasoningParser,
DeepSeekR1ReasoningParser, DelimitedReasoningParser, Qwen3ReasoningParser, ReasoningParser,
};
struct FakeTokenizer;
@@ -33,8 +32,6 @@ impl Tokenizer for FakeTokenizer {
"<|END_THINKING|>" => Some(4),
"◁think▷" => Some(5),
"◁/think▷" => Some(6),
"<mm:think>" => Some(8),
"</mm:think>" => Some(9),
_ => None,
}
}
@@ -162,35 +159,3 @@ fn deepseek_r1_stops_scanning_at_last_special_token() {
assert_eq!(delta.reasoning.as_deref(), Some("reason"));
assert_eq!(delta.content.as_deref(), Some("answer"));
}
#[test]
fn minimax_m3_handles_explicit_think_delimiters() {
let tokenizer = Arc::new(FakeTokenizer);
let mut parser = MiniMaxM3ReasoningParser::new(tokenizer).unwrap();
let delta = parser.push("<mm:think>reason</mm:think>answer").unwrap();
assert_eq!(delta.reasoning.as_deref(), Some("reason"));
assert_eq!(delta.content.as_deref(), Some("answer"));
}
#[test]
fn minimax_m3_uses_prompt_prefilled_start_marker() {
let tokenizer = Arc::new(FakeTokenizer);
let mut parser = MiniMaxM3ReasoningParser::new(tokenizer).unwrap();
parser.initialize(&[8]).unwrap();
let delta = parser.push("reason</mm:think>answer").unwrap();
assert_eq!(delta.reasoning.as_deref(), Some("reason"));
assert_eq!(delta.content.as_deref(), Some("answer"));
}
#[test]
fn minimax_m3_uses_prompt_prefilled_end_marker() {
let tokenizer = Arc::new(FakeTokenizer);
let mut parser = MiniMaxM3ReasoningParser::new(tokenizer).unwrap();
parser.initialize(&[9]).unwrap();
let delta = parser.push("answer").unwrap();
assert_eq!(delta.reasoning, None);
assert_eq!(delta.content.as_deref(), Some("answer"));
}
@@ -68,6 +68,7 @@ async fn main() -> Result<()> {
default_chat_template_kwargs: None,
chat_template_content_format: ChatTemplateContentFormatOption::Auto,
enable_log_requests: false,
enable_request_id_headers: false,
disable_log_stats: false,
grpc_port: None,
shutdown_timeout: Duration::ZERO,
+2
View File
@@ -61,6 +61,8 @@ pub struct Config {
pub chat_template_content_format: ChatTemplateContentFormatOption,
/// Log a summary line for each completed request.
pub enable_log_requests: bool,
/// When `true`, set `X-Request-Id` on every HTTP response.
pub enable_request_id_headers: bool,
/// When `true`, suppress periodic stats logging (throughput, queue depth,
/// cache usage).
pub disable_log_stats: bool,
+3 -1
View File
@@ -85,7 +85,9 @@ async fn build_state(config: &Config) -> Result<Arc<AppState>> {
};
Ok(Arc::new(
AppState::new(served_model_names, chat).with_log_requests(config.enable_log_requests),
AppState::new(served_model_names, chat)
.with_log_requests(config.enable_log_requests)
.with_request_id_headers(config.enable_request_id_headers),
))
}
+2
View File
@@ -1,5 +1,7 @@
mod load;
mod metrics;
mod request_id;
pub use load::track_server_load;
pub use metrics::track_http_metrics;
pub use request_id::set_request_id_header;
@@ -0,0 +1,24 @@
use axum::extract::Request;
use axum::http::HeaderValue;
use axum::http::header::HeaderName;
use axum::middleware::Next;
use axum::response::Response;
use uuid::Uuid;
const X_REQUEST_ID: HeaderName = HeaderName::from_static("x-request-id");
/// Echo the request's `X-Request-Id` on the response, or generate a fresh
/// `uuid4` hex if the request did not provide one.
///
/// Original Python:
/// `vllm.entrypoints.openai.server_utils.XRequestIdMiddleware`.
pub async fn set_request_id_header(req: Request, next: Next) -> Response {
let incoming = req.headers().get(&X_REQUEST_ID).cloned();
let mut response = next.run(req).await;
let value = incoming.unwrap_or_else(|| {
HeaderValue::from_str(&Uuid::new_v4().simple().to_string())
.expect("uuid hex is valid header value")
});
response.headers_mut().insert(X_REQUEST_ID, value);
response
}
+9 -2
View File
@@ -56,11 +56,18 @@ fn build_router_with_dev_mode(state: Arc<AppState>, dev_mode_enabled: bool) -> R
.route("/is_sleeping", get(sleep::is_sleeping))
}
router
let enable_request_id_headers = state.enable_request_id_headers;
let mut router = router
.with_state(state.clone())
.layer(from_fn_with_state(state, middleware::track_server_load))
.layer(from_fn(middleware::track_http_metrics))
.layer(TraceLayer::new_for_http())
.layer(TraceLayer::new_for_http());
if enable_request_id_headers {
router = router.layer(from_fn(middleware::set_request_id_header));
}
router
}
#[cfg(test)]
@@ -3,23 +3,33 @@ mod types;
mod validate;
use std::collections::HashMap;
use std::convert::Infallible;
use std::result::Result;
use std::sync::Arc;
use asynk_strim_attr::{TryYielder, try_stream};
use axum::Json;
use axum::extract::State;
use axum::http::HeaderMap;
use axum::response::sse::{Event, Sse};
use axum::response::{IntoResponse, Response};
use futures::{Stream, StreamExt as _, pin_mut};
use thiserror_ext::AsReport as _;
use tracing::info;
use tracing::{error, info, trace};
use tracing_futures::Instrument as _;
use vllm_engine_core_client::protocol::logprobs::{Logprobs, PositionLogprobs};
use vllm_llm::{CollectedGenerateOutput, GenerateOutputStreamExt as _};
use vllm_llm::{
CollectedGenerateOutput, FinishReason, GenerateOutput, GenerateOutputStreamExt as _,
};
use self::convert::prepare_generate_request;
use self::types::{GenerateLogprob, GenerateRequest, GenerateResponse, GenerateResponseChoice};
use crate::error::{ApiError, server_error};
use self::types::{
GenerateLogprob, GenerateRequest, GenerateResponse, GenerateResponseChoice,
GenerateResponseStreamChoice, GenerateStreamResponse,
};
use crate::error::{ApiError, bail_server_error, server_error};
use crate::routes::openai::utils::logprobs::clamp_logprob;
use crate::routes::openai::utils::types::{ChatLogProbs, ChatLogProbsContent, TopLogProb};
use crate::routes::openai::utils::types::{ChatLogProbs, ChatLogProbsContent, TopLogProb, Usage};
use crate::routes::openai::utils::validated_json::ValidatedJson;
use crate::state::AppState;
use crate::utils::resolve_request_context;
@@ -46,6 +56,7 @@ pub async fn generate(
let log_request = state.enable_log_requests;
let include_logprobs = prepared.include_logprobs;
let include_prompt_logprobs = prepared.include_prompt_logprobs;
let stream = prepared.stream;
let raw_stream = match state
.chat
@@ -64,6 +75,20 @@ pub async fn generate(
}
};
if stream {
let chunk_stream = generate_chunk_stream(
raw_stream,
prepared.request_id,
log_request,
prepared.include_usage,
prepared.include_continuous_usage,
include_logprobs,
);
let sse_stream = generate_sse_stream(chunk_stream).instrument(request_span);
return Sse::new(sse_stream).into_response();
}
let collected = match raw_stream.collect_output().instrument(request_span.clone()).await {
Ok(collected) => collected,
Err(error) => {
@@ -98,6 +123,102 @@ pub async fn generate(
Json(response).into_response()
}
#[try_stream]
async fn generate_chunk_stream(
stream: impl Stream<Item = vllm_llm::Result<GenerateOutput>>,
request_id: String,
log_request: bool,
include_usage: bool,
include_continuous_usage: bool,
include_logprobs: bool,
mut y: TryYielder<GenerateStreamResponse, ApiError>,
) -> Result<(), ApiError> {
pin_mut!(stream);
let mut prompt_tokens: Option<u32> = None;
let mut output_tokens = 0_u32;
while let Some(next) = stream.next().await {
match next {
Ok(output) => {
if prompt_tokens.is_none() {
prompt_tokens =
output.prompt_info.as_ref().map(|info| info.prompt_token_ids.len() as u32);
}
let usage_prompt_tokens = prompt_tokens.unwrap_or_default();
let token_ids = output.token_ids;
output_tokens = output_tokens.saturating_add(token_ids.len() as u32);
let finish_reason = output.finish_reason;
if matches!(finish_reason.as_ref(), Some(FinishReason::Error)) {
bail_server_error!("Internal server error");
}
if let Some(finish_reason) = finish_reason.as_ref()
&& log_request
{
info!(
stream = true,
prompt_tokens = usage_prompt_tokens,
output_tokens,
finish_reason = finish_reason.as_str(),
"generate finished"
);
}
if token_ids.is_empty() && finish_reason.is_none() {
continue;
}
let logprobs = if include_logprobs && !token_ids.is_empty() {
let logprobs = output.logprobs.as_ref().ok_or_else(|| {
server_error!(
"raw generate stream requested logprobs but generation returned none"
)
})?;
Some(raw_logprobs_to_openai_chat(logprobs)?)
} else {
None
};
y.yield_ok(GenerateStreamResponse {
request_id: request_id.clone(),
choices: vec![GenerateResponseStreamChoice {
index: 0,
logprobs,
finish_reason: finish_reason.map(|reason| reason.as_str().to_string()),
token_ids,
}],
usage: include_continuous_usage
.then(|| Usage::from_counts(usage_prompt_tokens, output_tokens)),
})
.await;
}
Err(error) => {
error!(
error = %error.as_report(),
"raw generate stream failed"
);
bail_server_error!("{}", error.to_report_string());
}
}
}
if include_usage {
y.yield_ok(GenerateStreamResponse {
request_id,
choices: Vec::new(),
usage: Some(Usage::from_counts(
prompt_tokens.unwrap_or_default(),
output_tokens,
)),
})
.await;
}
Ok(())
}
fn collect_generate(
collected: CollectedGenerateOutput,
request_id: String,
@@ -213,3 +334,94 @@ fn position_to_logprob_map(position: &PositionLogprobs) -> HashMap<u32, Generate
fn format_token_id(token_id: u32) -> String {
format!("token_id:{token_id}")
}
/// Convert one raw-generate chunk stream into SSE events.
#[try_stream]
async fn generate_sse_stream(
stream: impl Stream<Item = Result<GenerateStreamResponse, ApiError>>,
mut y: TryYielder<Event, Infallible>,
) -> Result<(), Infallible> {
pin_mut!(stream);
while let Some(next) = stream.next().await {
match next {
Ok(chunk) => y.yield_ok(to_sse_event(&chunk)).await,
Err(error) => {
y.yield_ok(to_error_sse_event(&error)).await;
break;
}
}
}
y.yield_ok(done_sse_event()).await;
Ok(())
}
fn to_sse_event(chunk: &GenerateStreamResponse) -> Event {
let payload = serde_json::to_string(chunk).expect("generate chunk must serialize to JSON");
trace!(payload, "generate emitting chunk");
Event::default().data(payload)
}
fn to_error_sse_event(error: &ApiError) -> Event {
let payload = serde_json::to_string(&error.to_error_response())
.expect("ErrorResponse must serialize to JSON");
trace!(payload, "generate emitting error");
Event::default().data(payload)
}
fn done_sse_event() -> Event {
trace!("generate emitting done");
Event::default().data("[DONE]")
}
#[cfg(test)]
mod tests {
use std::sync::Arc;
use futures::{TryStreamExt as _, stream};
use vllm_llm::GeneratePromptInfo;
use super::*;
#[tokio::test]
async fn generate_chunk_stream_captures_late_prompt_info() {
let stream = stream::iter(vec![
Ok(GenerateOutput {
request_id: String::new(),
prompt_info: None,
token_ids: Vec::new(),
logprobs: None,
finish_reason: None,
kv_transfer_params: None,
}),
Ok(GenerateOutput {
request_id: String::new(),
prompt_info: Some(GeneratePromptInfo {
prompt_token_ids: Arc::from([11_u32, 22_u32]),
prompt_logprobs: None,
}),
token_ids: vec![33],
logprobs: None,
finish_reason: Some(FinishReason::stop_eos()),
kv_transfer_params: None,
}),
]);
let chunks: Vec<_> =
generate_chunk_stream(stream, "raw-stream".to_string(), false, true, true, false)
.try_collect()
.await
.expect("collect chunks");
assert_eq!(chunks.len(), 2);
assert_eq!(
chunks[0].usage.as_ref().expect("chunk usage").prompt_tokens,
2
);
assert_eq!(
chunks[1].usage.as_ref().expect("final usage").prompt_tokens,
2
);
}
}
@@ -10,6 +10,9 @@ use crate::utils::{ResolvedRequestContext, merge_kv_transfer_params};
pub struct PreparedRequest {
pub request_id: String,
pub text_request: TextRequest,
pub stream: bool,
pub include_usage: bool,
pub include_continuous_usage: bool,
pub include_logprobs: bool,
pub include_prompt_logprobs: bool,
}
@@ -23,6 +26,18 @@ pub fn prepare_generate_request(
) -> Result<PreparedRequest, ApiError> {
validate::validate_request_compat(&request, served_model_names)?;
let stream = request.stream;
let include_usage = request
.stream_options
.as_ref()
.and_then(|options| options.include_usage)
.unwrap_or(false);
let include_continuous_usage = include_usage
&& request
.stream_options
.as_ref()
.and_then(|options| options.continuous_usage_stats)
.unwrap_or(false);
let include_logprobs = request.sampling_params.logprobs.is_some();
let include_prompt_logprobs = request.sampling_params.prompt_logprobs.is_some();
let mut sampling_params = request.sampling_params;
@@ -47,6 +62,9 @@ pub fn prepare_generate_request(
Ok(PreparedRequest {
request_id: ctx.request_id,
text_request,
stream,
include_usage,
include_continuous_usage,
include_logprobs,
include_prompt_logprobs,
})
@@ -109,4 +127,28 @@ mod tests {
Some(json!({"connector": "x"}))
);
}
#[test]
fn prepare_generate_request_gates_continuous_usage_on_include_usage() {
let request: GenerateRequest = serde_json::from_value(json!({
"model": "Qwen/Qwen1.5-0.5B-Chat",
"token_ids": [11, 22],
"stream": true,
"stream_options": {
"continuous_usage_stats": true
},
"sampling_params": {}
}))
.expect("parse request");
let prepared = prepare_generate_request(
request,
&["Qwen/Qwen1.5-0.5B-Chat".to_string()],
ResolvedRequestContext::default(),
)
.expect("prepare");
assert!(!prepared.include_usage);
assert!(!prepared.include_continuous_usage);
}
}
@@ -5,7 +5,7 @@ use serde_json::{Map, Value};
use validator::Validate;
use vllm_text::SamplingParams;
use crate::routes::openai::utils::types::{ChatLogProbs, Normalizable};
use crate::routes::openai::utils::types::{ChatLogProbs, Normalizable, StreamOptions, Usage};
/// vLLM-compatible request type for the token-in/token-out generate API.
#[serde_with::skip_serializing_none]
@@ -17,6 +17,7 @@ pub struct GenerateRequest {
pub sampling_params: SamplingParams,
#[serde(default)]
pub stream: bool,
pub stream_options: Option<StreamOptions>,
pub cache_salt: Option<String>,
#[serde(default)]
pub priority: i32,
@@ -37,6 +38,25 @@ pub(super) struct GenerateResponseChoice {
pub token_ids: Vec<u32>,
}
/// Mirrors the Python vLLM `GenerateResponseStreamChoice` class.
#[serde_with::skip_serializing_none]
#[derive(Debug, Clone, Serialize)]
pub(super) struct GenerateResponseStreamChoice {
pub index: u32,
pub logprobs: Option<ChatLogProbs>,
pub finish_reason: Option<String>,
pub token_ids: Vec<u32>,
}
/// Mirrors the Python vLLM `GenerateStreamResponse` class.
#[serde_with::skip_serializing_none]
#[derive(Debug, Clone, Serialize)]
pub(super) struct GenerateStreamResponse {
pub request_id: String,
pub choices: Vec<GenerateResponseStreamChoice>,
pub usage: Option<Usage>,
}
/// Mirrors the Python vLLM `GenerateResponse` class.
#[serde_with::skip_serializing_none]
#[derive(Debug, Clone, Serialize)]
@@ -13,8 +13,11 @@ pub(super) fn validate_request_compat(
return Err(ApiError::model_not_found(model.clone()));
}
if request.stream {
bail_invalid_request!(param = "stream", "stream=true is not supported.");
if request.stream_options.is_some() && !request.stream {
bail_invalid_request!(
param = "stream_options",
"stream_options are only supported when stream=true."
);
}
if request.token_ids.is_empty() {
@@ -65,11 +68,24 @@ mod tests {
}
#[test]
fn validate_request_compat_rejects_streaming() {
fn validate_request_compat_accepts_streaming() {
let request = GenerateRequest {
stream: true,
..base_request()
};
assert!(validate_request_compat(&request, &served(&["Qwen/Qwen1.5-0.5B-Chat"])).is_ok());
}
#[test]
fn validate_request_compat_rejects_stream_options_without_streaming() {
let request: GenerateRequest = serde_json::from_value(json!({
"model": "Qwen/Qwen1.5-0.5B-Chat",
"token_ids": [11, 22],
"stream": false,
"stream_options": {"include_usage": true},
"sampling_params": {}
}))
.expect("parse request");
assert!(validate_request_compat(&request, &served(&["Qwen/Qwen1.5-0.5B-Chat"])).is_err());
}
+327 -6
View File
@@ -746,6 +746,20 @@ async fn test_app() -> axum::Router {
)))
}
async fn test_app_with_request_id_headers() -> (axum::Router, MockEngineTask) {
let (chat, engine_task) = test_models_with_engine_outputs_and_backend(
b"engine-openai-request-id",
default_stream_output_specs(),
Arc::new(FakeChatBackend::new()),
)
.await;
let app = build_router(Arc::new(
AppState::new(vec!["Qwen/Qwen1.5-0.5B-Chat".to_string()], chat)
.with_request_id_headers(true),
));
(app, engine_task)
}
async fn test_health_app_with_engine_script<F>(
script: F,
) -> (axum::Router, Arc<AppState>, MockEngineTask)
@@ -947,6 +961,18 @@ async fn health_status(app: &axum::Router) -> (StatusCode, Bytes) {
(status, body)
}
async fn health_response(app: &axum::Router, request_id: Option<&str>) -> axum::response::Response {
let mut builder = Request::builder().method("GET").uri("/health");
if let Some(request_id) = request_id {
builder = builder.header("X-Request-Id", request_id);
}
app.clone()
.call(builder.body(Body::empty()).expect("build request"))
.await
.expect("call app")
}
fn metric_value(rendered: &str, metric: &str, labels: Option<&str>) -> Option<f64> {
rendered.lines().find_map(|line| {
let rest = line.strip_prefix(metric)?;
@@ -994,6 +1020,43 @@ async fn list_models_returns_configured_model() {
assert_eq!(json["data"][0]["id"], "Qwen/Qwen1.5-0.5B-Chat");
}
#[tokio::test(flavor = "multi_thread", worker_threads = 2)]
#[serial]
async fn request_id_header_is_absent_by_default() {
let app = test_app().await;
let response = health_response(&app, None).await;
assert_eq!(response.status(), StatusCode::OK);
assert!(!response.headers().contains_key("x-request-id"));
}
#[tokio::test(flavor = "multi_thread", worker_threads = 2)]
#[serial]
async fn request_id_header_generates_uuid_hex_when_enabled() {
let (app, _engine_task) = test_app_with_request_id_headers().await;
let response = health_response(&app, None).await;
assert_eq!(response.status(), StatusCode::OK);
let request_id = response
.headers()
.get("x-request-id")
.expect("x-request-id header")
.to_str()
.expect("header is ascii");
assert_eq!(request_id.len(), 32);
assert!(request_id.chars().all(|ch| ch.is_ascii_hexdigit()));
}
#[tokio::test(flavor = "multi_thread", worker_threads = 2)]
#[serial]
async fn request_id_header_echoes_incoming_header_when_enabled() {
let (app, _engine_task) = test_app_with_request_id_headers().await;
let response = health_response(&app, Some("req-123")).await;
assert_eq!(response.status(), StatusCode::OK);
assert_eq!(response.headers().get("x-request-id").unwrap(), "req-123");
}
#[tokio::test(flavor = "multi_thread", worker_threads = 2)]
#[serial]
async fn version_returns_engine_vllm_version() {
@@ -2425,8 +2488,72 @@ async fn non_stream_raw_generate_returns_token_output_envelope() {
#[tokio::test(flavor = "multi_thread", worker_threads = 2)]
#[serial]
async fn raw_generate_rejects_streaming() {
let mut app = test_app().await;
async fn stream_raw_generate_returns_sse_chunks_and_usage() {
let ipc = IpcNamespace::new().expect("create ipc namespace");
let handshake_address = ipc.handshake_endpoint();
let engine_id = b"engine-raw-generate-stream".to_vec();
let engine_task = MockEngineTask::new(spawn_mock_engine_task(
handshake_address.clone(),
engine_id.clone(),
|dealer, push| {
boxed_test_future(async move {
let add = recv_engine_message(dealer).await;
let request: EngineCoreRequest =
rmp_serde::from_slice(&add[1]).expect("decode request");
assert_eq!(request.prompt_token_ids.as_deref(), Some(&[11, 22][..]));
assert_eq!(request.external_req_id.as_deref(), Some("raw-stream"));
send_outputs(
push,
EngineCoreOutputs {
engine_index: 0,
outputs: vec![
request_output_with_logprobs(
&request.request_id,
vec![33],
None,
None,
Some(sample_logprobs_for_token(33, 34)),
None,
),
request_output_with_logprobs(
&request.request_id,
vec![44],
Some(EngineCoreFinishReason::Stop),
None,
Some(sample_logprobs_for_token(44, 45)),
None,
),
],
scheduler_stats: None,
timestamp: 0.0,
utility_output: None,
finished_requests: None,
wave_complete: None,
start_wave: None,
},
)
.await;
})
},
));
let client = EngineCoreClient::connect(
EngineCoreClientConfig::new_single(handshake_address)
.with_model_name("test-model")
.with_local_input_output_addresses(
Some(ipc.input_endpoint()),
Some(ipc.output_endpoint()),
),
)
.await
.expect("connect client");
let chat = ChatLlm::from_shared_backend(Llm::new(client), Arc::new(FakeChatBackend::new()));
let mut app = build_router(Arc::new(AppState::new(
vec!["Qwen/Qwen1.5-0.5B-Chat".to_string()],
chat,
)));
let response = app
.call(
@@ -2437,9 +2564,17 @@ async fn raw_generate_rejects_streaming() {
.body(Body::from(
json!({
"model": "Qwen/Qwen1.5-0.5B-Chat",
"request_id": "raw-stream",
"token_ids": [11, 22],
"stream": true,
"sampling_params": {}
"stream_options": {
"include_usage": true,
"continuous_usage_stats": true
},
"sampling_params": {
"max_tokens": 2,
"logprobs": 1
}
})
.to_string(),
))
@@ -2448,10 +2583,196 @@ async fn raw_generate_rejects_streaming() {
.await
.expect("call app");
assert_eq!(response.status(), StatusCode::BAD_REQUEST);
assert_eq!(response.status(), StatusCode::OK);
assert_eq!(
response.headers().get("content-type").and_then(|value| value.to_str().ok()),
Some("text/event-stream")
);
let body = to_bytes(response.into_body(), usize::MAX).await.expect("read body");
let json: serde_json::Value = serde_json::from_slice(&body).expect("decode json");
assert_eq!(json["error"]["param"], "stream");
engine_task.await.expect("mock engine task");
let text = String::from_utf8(body.to_vec()).expect("utf8 body");
let payloads = sse_data_payloads(&text);
assert_eq!(payloads.len(), 4, "{text}");
let first: serde_json::Value = serde_json::from_str(payloads[0]).expect("first chunk json");
assert_eq!(first["request_id"], "raw-stream");
assert_eq!(first["choices"][0]["index"], 0);
assert_eq!(first["choices"][0]["token_ids"], json!([33]));
assert_eq!(
first["choices"][0]["logprobs"]["content"][0]["token"],
"token_id:33"
);
assert_eq!(first["usage"]["prompt_tokens"], 2);
assert_eq!(first["usage"]["completion_tokens"], 1);
let second: serde_json::Value = serde_json::from_str(payloads[1]).expect("second chunk json");
assert_eq!(second["choices"][0]["token_ids"], json!([44]));
assert_eq!(second["choices"][0]["finish_reason"], "stop");
assert_eq!(second["usage"]["completion_tokens"], 2);
let usage: serde_json::Value = serde_json::from_str(payloads[2]).expect("usage chunk json");
assert_eq!(usage["choices"], json!([]));
assert_eq!(usage["usage"]["prompt_tokens"], 2);
assert_eq!(usage["usage"]["completion_tokens"], 2);
assert_eq!(usage["usage"]["total_tokens"], 4);
assert_eq!(payloads[3], "[DONE]");
}
#[tokio::test(flavor = "multi_thread", worker_threads = 2)]
#[serial]
async fn stream_raw_generate_emits_final_usage_without_continuous_usage() {
let (mut app, engine_task) = test_app_with_stream_output_specs(vec![
(vec![33], None),
(vec![44], Some(EngineCoreFinishReason::Stop)),
])
.await;
let response = app
.call(
Request::builder()
.method("POST")
.uri("/inference/v1/generate")
.header("content-type", "application/json")
.body(Body::from(
json!({
"model": "Qwen/Qwen1.5-0.5B-Chat",
"request_id": "raw-stream-final-usage",
"token_ids": [11, 22],
"stream": true,
"stream_options": {
"include_usage": true
},
"sampling_params": {
"max_tokens": 2
}
})
.to_string(),
))
.expect("build request"),
)
.await
.expect("call app");
assert_eq!(response.status(), StatusCode::OK);
let body = to_bytes(response.into_body(), usize::MAX).await.expect("read body");
engine_task.await.expect("mock engine task");
let text = String::from_utf8(body.to_vec()).expect("utf8 body");
let payloads = sse_data_payloads(&text);
assert_eq!(payloads.len(), 4, "{text}");
let first: serde_json::Value = serde_json::from_str(payloads[0]).expect("first chunk json");
assert_eq!(first["choices"][0]["token_ids"], json!([33]));
assert!(first.get("usage").is_none());
let second: serde_json::Value = serde_json::from_str(payloads[1]).expect("second chunk json");
assert_eq!(second["choices"][0]["token_ids"], json!([44]));
assert_eq!(second["choices"][0]["finish_reason"], "stop");
assert!(second.get("usage").is_none());
let usage: serde_json::Value = serde_json::from_str(payloads[2]).expect("usage chunk json");
assert_eq!(usage["choices"], json!([]));
assert_eq!(usage["usage"]["prompt_tokens"], 2);
assert_eq!(usage["usage"]["completion_tokens"], 2);
assert_eq!(usage["usage"]["total_tokens"], 4);
assert_eq!(payloads[3], "[DONE]");
}
#[tokio::test(flavor = "multi_thread", worker_threads = 2)]
#[serial]
async fn stream_raw_generate_emits_empty_finish_chunk() {
let (mut app, engine_task) = test_app_with_stream_output_specs(vec![
(vec![33], None),
(vec![], Some(EngineCoreFinishReason::Stop)),
])
.await;
let response = app
.call(
Request::builder()
.method("POST")
.uri("/inference/v1/generate")
.header("content-type", "application/json")
.body(Body::from(
json!({
"model": "Qwen/Qwen1.5-0.5B-Chat",
"request_id": "raw-stream-empty-finish",
"token_ids": [11, 22],
"stream": true,
"sampling_params": {
"max_tokens": 2
}
})
.to_string(),
))
.expect("build request"),
)
.await
.expect("call app");
assert_eq!(response.status(), StatusCode::OK);
let body = to_bytes(response.into_body(), usize::MAX).await.expect("read body");
engine_task.await.expect("mock engine task");
let text = String::from_utf8(body.to_vec()).expect("utf8 body");
let payloads = sse_data_payloads(&text);
assert_eq!(payloads.len(), 3, "{text}");
let first: serde_json::Value = serde_json::from_str(payloads[0]).expect("first chunk json");
assert_eq!(first["choices"][0]["token_ids"], json!([33]));
assert!(first["choices"][0].get("finish_reason").is_none());
let second: serde_json::Value = serde_json::from_str(payloads[1]).expect("second chunk json");
assert_eq!(second["choices"][0]["token_ids"], json!([]));
assert_eq!(second["choices"][0]["finish_reason"], "stop");
assert_eq!(payloads[2], "[DONE]");
}
#[tokio::test(flavor = "multi_thread", worker_threads = 2)]
#[serial]
async fn stream_raw_generate_error_finish_returns_sse_error() {
let (mut app, engine_task) =
test_app_with_stream_output_specs(vec![(vec![], Some(EngineCoreFinishReason::Error))])
.await;
let response = app
.call(
Request::builder()
.method("POST")
.uri("/inference/v1/generate")
.header("content-type", "application/json")
.body(Body::from(
json!({
"model": "Qwen/Qwen1.5-0.5B-Chat",
"request_id": "raw-stream-error",
"token_ids": [11, 22],
"stream": true,
"stream_options": {
"include_usage": true
},
"sampling_params": {
"max_tokens": 2
}
})
.to_string(),
))
.expect("build request"),
)
.await
.expect("call app");
assert_eq!(response.status(), StatusCode::OK);
let body = to_bytes(response.into_body(), usize::MAX).await.expect("read body");
engine_task.await.expect("mock engine task");
let text = String::from_utf8(body.to_vec()).expect("utf8 body");
assert!(text.contains("\"type\":\"server_error\""), "{text}");
assert!(text.contains("Internal server error"), "{text}");
assert!(!text.contains("\"finish_reason\":\"error\""), "{text}");
assert!(!text.contains("\"usage\":"), "{text}");
assert!(text.trim_end().ends_with("data: [DONE]"), "{text}");
}
#[tokio::test(flavor = "multi_thread", worker_threads = 2)]
+9
View File
@@ -17,6 +17,8 @@ pub struct AppState {
pub chat: ChatLlm,
/// Whether to log a summary line for each completed request.
pub enable_log_requests: bool,
/// Whether to set X-Request-Id on every HTTP response.
pub enable_request_id_headers: bool,
/// Number of in-flight inference requests currently owned by this frontend.
server_load: AtomicU64,
}
@@ -39,6 +41,7 @@ impl AppState {
served_model_names,
chat,
enable_log_requests: false,
enable_request_id_headers: false,
server_load: AtomicU64::new(0),
}
}
@@ -49,6 +52,12 @@ impl AppState {
self
}
/// Enable X-Request-Id response headers.
pub fn with_request_id_headers(mut self, enabled: bool) -> Self {
self.enable_request_id_headers = enabled;
self
}
/// The primary model name echoed back in API responses (the first served
/// name).
pub fn primary_model_name(&self) -> &str {
-23
View File
@@ -1,23 +0,0 @@
[package]
name = "vllm-tool-parser-py"
version.workspace = true
edition.workspace = true
license.workspace = true
[lib]
name = "_rust_tool_parser"
crate-type = ["cdylib", "rlib"]
[features]
default = []
extension-module = ["pyo3/extension-module"]
[dependencies]
pyo3.workspace = true
pythonize = { workspace = true, features = ["serde_json"] }
serde_json.workspace = true
thiserror-ext.workspace = true
vllm-tool-parser.workspace = true
[lints]
workspace = true
-367
View File
@@ -1,367 +0,0 @@
//! Thin PyO3 bindings for `vllm_tool_parser`.
//!
//! This crate exposes the Rust tool parser trait and data shapes to Python
//! while keeping parser state, grammar, and schema-aware argument conversion in
//! Rust. Python callers should use this module as a typed bridge and keep any
//! vLLM protocol adaptation outside the binding.
use pyo3::exceptions::PyValueError;
use pyo3::prelude::*;
use pyo3::types::{PyAny, PyModule};
use pythonize::{depythonize, pythonize};
use serde_json::Value;
use thiserror_ext::AsReport as _;
use vllm_tool_parser::{Tool, ToolCallDelta, ToolParser, ToolParserOutput};
macro_rules! tool_parser_factory {
($($parser:ident),+ $(,)?) => {
fn create_tool_parser(
name: &str,
tools: &[Tool],
) -> PyResult<Box<dyn ToolParser>> {
match name {
$(
stringify!($parser) => {
<vllm_tool_parser::$parser as ToolParser>::create(tools)
}
)+
_ => {
return Err(PyValueError::new_err(format!(
"unsupported tool parser `{name}`"
)));
}
}
.map_err(|error| PyValueError::new_err(error.to_report_string()))
}
};
}
// Export a tool parser to Python by registering it here.
tool_parser_factory! {
DeepSeekV4ToolParser,
MinimaxM3ToolParser,
}
#[pyclass(name = "Tool", module = "vllm._rust_tool_parser", skip_from_py_object)]
#[derive(Clone)]
struct PyTool(Tool);
#[pymethods]
impl PyTool {
#[new]
#[pyo3(signature = (name, description, parameters, strict=None))]
fn new(
name: String,
description: Option<String>,
parameters: &Bound<'_, PyAny>,
strict: Option<bool>,
) -> PyResult<Self> {
let parameters = depythonize::<Value>(parameters).map_err(|error| {
PyValueError::new_err(format!(
"failed to convert tool parameters from Python to JSON: {error}"
))
})?;
Ok(Self(Tool {
name,
description,
parameters,
strict,
}))
}
#[getter]
fn name(&self) -> &str {
&self.0.name
}
#[getter]
fn description(&self) -> Option<&str> {
self.0.description.as_deref()
}
#[getter]
fn parameters(&self, py: Python<'_>) -> PyResult<Py<PyAny>> {
pythonize(py, &self.0.parameters).map(Bound::unbind).map_err(|error| {
PyValueError::new_err(format!(
"failed to convert tool parameters from JSON to Python: {error}"
))
})
}
#[getter]
fn strict(&self) -> Option<bool> {
self.0.strict
}
}
#[pyclass(
name = "ToolCallDelta",
module = "vllm._rust_tool_parser",
skip_from_py_object
)]
#[derive(Clone)]
struct PyToolCallDelta(ToolCallDelta);
#[pymethods]
impl PyToolCallDelta {
#[new]
#[pyo3(signature = (tool_index, name, arguments))]
fn new(tool_index: usize, name: Option<String>, arguments: String) -> Self {
Self(ToolCallDelta {
tool_index,
name,
arguments,
})
}
#[getter]
fn tool_index(&self) -> usize {
self.0.tool_index
}
#[getter]
fn name(&self) -> Option<&str> {
self.0.name.as_deref()
}
#[getter]
fn arguments(&self) -> &str {
&self.0.arguments
}
}
#[pyclass(
name = "ToolParserOutput",
module = "vllm._rust_tool_parser",
skip_from_py_object
)]
#[derive(Clone)]
struct PyToolParserOutput(ToolParserOutput);
#[pymethods]
impl PyToolParserOutput {
#[new]
#[pyo3(signature = (normal_text="", calls=None))]
fn new(py: Python<'_>, normal_text: &str, calls: Option<Vec<Py<PyToolCallDelta>>>) -> Self {
let calls =
calls.unwrap_or_default().iter().map(|call| call.borrow(py).0.clone()).collect();
Self(ToolParserOutput {
normal_text: normal_text.to_owned(),
calls,
})
}
#[getter]
fn normal_text(&self) -> &str {
&self.0.normal_text
}
#[getter]
fn calls(&self) -> Vec<PyToolCallDelta> {
self.0.calls.iter().cloned().map(PyToolCallDelta).collect()
}
fn append(&mut self, other: PyRef<'_, PyToolParserOutput>) {
self.0.append(other.0.clone());
}
fn coalesce_calls(&self) -> Self {
Self(self.0.clone().coalesce_calls())
}
}
#[pyclass(name = "ToolParser", module = "vllm._rust_tool_parser", unsendable)]
struct PyToolParser(Box<dyn ToolParser>);
impl PyToolParser {
fn parse_into_output(&mut self, chunk: &str, output: &mut PyToolParserOutput) -> PyResult<()> {
self.0
.parse_into(chunk, &mut output.0)
.map_err(|error| PyValueError::new_err(error.to_report_string()))
}
}
#[pymethods]
impl PyToolParser {
#[new]
fn new(py: Python<'_>, parser_name: &str, tools: Vec<Py<PyTool>>) -> PyResult<Self> {
let tools = tools.iter().map(|tool| tool.borrow(py).0.clone()).collect::<Vec<_>>();
create_tool_parser(parser_name, &tools).map(Self)
}
fn parse_into(
&mut self,
chunk: &str,
mut output: PyRefMut<'_, PyToolParserOutput>,
) -> PyResult<()> {
self.parse_into_output(chunk, &mut output)
}
fn finish(&mut self) -> PyResult<PyToolParserOutput> {
self.0
.finish()
.map(PyToolParserOutput)
.map_err(|error| PyValueError::new_err(error.to_report_string()))
}
fn reset(&mut self) -> String {
self.0.reset()
}
fn preserve_special_tokens(&self) -> bool {
self.0.preserve_special_tokens()
}
}
#[pymodule]
fn _rust_tool_parser(m: &Bound<'_, PyModule>) -> PyResult<()> {
m.add_class::<PyTool>()?;
m.add_class::<PyToolCallDelta>()?;
m.add_class::<PyToolParserOutput>()?;
m.add_class::<PyToolParser>()?;
Ok(())
}
#[cfg(test)]
mod tests {
use super::*;
use serde_json::json;
fn with_python<R>(f: impl for<'py> FnOnce(Python<'py>) -> R) -> R {
Python::initialize();
Python::attach(f)
}
fn tool_schema() -> Value {
json!({
"type": "object",
"properties": {
"user_id": {"type": "integer"},
"shipping": {
"type": "object",
"properties": {
"city": {"type": "string"},
"zip": {"type": "integer"}
}
}
}
})
}
fn build_call() -> String {
r#"<DSMLtool_calls>
<DSMLinvoke name="create_order">
<DSMLparameter name="user_id" string="false">42</DSMLparameter>
<DSMLparameter name="shipping" string="false">{"city":"Singapore","zip":18956}</DSMLparameter>
</DSMLinvoke>
</DSMLtool_calls>"#
.to_owned()
}
fn make_py_tool(py: Python<'_>) -> PyResult<Py<PyTool>> {
let parameters = pythonize(py, &tool_schema()).map_err(|error| {
PyValueError::new_err(format!(
"failed to convert test schema from JSON to Python: {error}"
))
})?;
Py::new(
py,
PyTool::new(
"create_order".to_owned(),
Some("Create an order".to_owned()),
&parameters,
None,
)?,
)
}
#[test]
fn tool_round_trips_typed_fields() {
with_python(|py| {
let tool = make_py_tool(py)?;
let borrowed = tool.borrow(py);
assert_eq!(borrowed.name(), "create_order");
assert_eq!(borrowed.description(), Some("Create an order"));
assert_eq!(borrowed.strict(), None);
let parameters = borrowed.parameters(py)?;
let parameters = depythonize::<Value>(parameters.bind(py))?;
assert_eq!(parameters, tool_schema());
PyResult::Ok(())
})
.unwrap();
}
#[test]
fn output_append_and_coalesce_calls() {
with_python(|py| {
let first = Py::new(
py,
PyToolCallDelta::new(0, Some("create_order".to_owned()), "{\"a\"".to_owned()),
)?;
let second = Py::new(py, PyToolCallDelta::new(0, None, ":1}".to_owned()))?;
let mut output = PyToolParserOutput::new(py, "text", Some(vec![first]));
let other = Py::new(py, PyToolParserOutput::new(py, "", Some(vec![second])))?;
output.append(other.borrow(py));
let coalesced = output.coalesce_calls();
assert_eq!(coalesced.normal_text(), "text");
let calls = coalesced.calls();
assert_eq!(calls.len(), 1);
assert_eq!(calls[0].tool_index(), 0);
assert_eq!(calls[0].name(), Some("create_order"));
assert_eq!(calls[0].arguments(), "{\"a\":1}");
PyResult::Ok(())
})
.unwrap();
}
#[test]
fn parser_parse_finish_and_preserve_special_tokens() {
with_python(|py| {
let tool = make_py_tool(py)?;
let mut parser = PyToolParser::new(py, "DeepSeekV4ToolParser", vec![tool])?;
assert!(parser.preserve_special_tokens());
let mut output = PyToolParserOutput::new(py, "", None);
parser.parse_into_output(&build_call(), &mut output)?;
let finish = Py::new(py, parser.finish()?)?;
output.append(finish.borrow(py));
let output = output.coalesce_calls();
assert_eq!(output.normal_text(), "");
let calls = output.calls();
assert_eq!(calls.len(), 1);
assert_eq!(calls[0].name(), Some("create_order"));
assert_eq!(
serde_json::from_str::<Value>(calls[0].arguments()).unwrap(),
json!({
"user_id": 42,
"shipping": {
"city": "Singapore",
"zip": 18956
}
})
);
assert_eq!(parser.reset(), "");
PyResult::Ok(())
})
.unwrap();
}
#[test]
fn parser_errors_for_unknown_name() {
with_python(|py| {
let tool = make_py_tool(py)?;
let error = match PyToolParser::new(py, "missing", vec![tool]) {
Ok(_) => panic!("missing parser name unexpectedly succeeded"),
Err(error) => error,
};
let message = format!("{error}");
assert!(message.contains("unsupported tool parser `missing`"));
PyResult::Ok(())
})
.unwrap();
}
}
+1 -1
View File
@@ -8,7 +8,7 @@ const HERMES_CONFIG: JsonToolCallConfig = JsonToolCallConfig {
marker_whitespace: JsonToolCallWhitespace::Optional,
delimiter: None,
name_key: "name",
arguments_key: "arguments",
arguments_key: &["arguments"],
};
/// Tool parser for Hermes XML-wrapped JSON tool calls.

Some files were not shown because too many files have changed in this diff Show More