Compare commits

..
Author SHA1 Message Date
Tyler Michael SmithandClaude 5c846ccbda [CI] Inline env vars lost by CONTINUE_ON_FAILURE subshell wrapping
Postmerge/nightly/daily builds set CONTINUE_ON_FAILURE=1, which wraps
each YAML command in `(cmd) || CI_OVERALL_STATUS=1`. A standalone
`export VAR=val` executes inside the subshell and the variable is
immediately lost — subsequent commands never see it.

This caused real failures (model_executor fastsafetensors crash from
missing VLLM_WORKER_MULTIPROC_METHOD=spawn) and silent coverage gaps
(rust_frontend tests silently running the Python frontend because
VLLM_USE_RUST_FRONTEND=1 was never set).

Fix: inline env vars as command prefixes (`VAR=val command`) instead
of standalone `export` lines across all 11 affected CI YAML files.

Co-authored-by: Claude <noreply@anthropic.com>
Signed-off-by: Tyler Michael Smith <tyler@tylermsmith.com>
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
2026-06-22 17:53:18 -04:00
Taneem IbrahimGitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
f2069b005b [Pooling] Validate non-negative rerank top_n (#46119)
Signed-off-by: Taneem Ibrahim <taneem.ibrahim@gmail.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-06-22 11:40:47 -04:00
Martin HickeyandGitHub ccd49f6821 [MyPy] Fix mypy for vllm/lora (#41722)
Signed-off-by: Martin Hickey <martin.hickey@ie.ibm.com>
2026-06-22 10:57:09 -04:00
Li, JiangandGitHub 1c7bc18318 [Bugfix][CPU] Fix CPU model runner v2 (#46365)
Signed-off-by: jiang1.li <jiang1.li@intel.com>
2026-06-22 22:52:05 +08:00
AlexHuangandGitHub 9a938df64e [Test][KV Offloading] Add unit tests for OffloadingSpecFactory and SecondaryTierFactory (#46355)
Signed-off-by: Alex <alex.tech.lab@outlook.com>
2026-06-22 17:45:04 +03:00
Liangliang MaandGitHub 3da4a1b124 [XPU] add awq format for INCXPULinear (#43404)
Signed-off-by: Ma, Liangliang <liangliang.ma@intel.com>
2026-06-22 22:29:13 +08:00
6871738777 [Doc] Document pull request limit (#46376)
Signed-off-by: simon-mo <simon.mo@hey.com>
Co-authored-by: OpenAI Codex <codex@openai.com>
2026-06-22 14:04:56 +00:00
Yifan QiaoandGitHub aa4990a9a2 [Attention] Re-enable cross-layer KV cache layout for MLA via stride-aware kernels (#45111)
Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai>
2026-06-22 06:57:02 -07:00
a4610da0c6 [docs] link security docs from AGENTS (#46373)
Add a security-review routing sentence to AGENTS.md that points agents to SECURITY.md, docs/usage/security.md, and docs/contributing/vulnerability_management.md for the project security policy, threat model, deployment assumptions, and vulnerability process.

Co-authored-by: OpenAI Codex <codex@openai.com>
2026-06-22 06:28:25 -07:00
liuzhenweiandGitHub 09cdcf34aa [XPU] update nixl to v1.2.0 (#46327)
Signed-off-by: zhenwei-intel <zhenwei.liu@intel.com>
2026-06-22 20:55:06 +08:00
wcyGitHubLi, Jiang <jiang1.li@intel.com>
d2c671c29b [CPU][RISC-V] Add RVV micro GEMM for WNA16 (#44324)
Signed-off-by: wcy <233313160abc@gmail.com>
Co-authored-by: Li, Jiang <jiang1.li@intel.com>
2026-06-22 12:53:54 +00:00
xiangdongandGitHub b5a2adec4b [XPU][CI]Skip v1/spec_decode/test_speculators_correctness.py in intel GPU nightly (#46356)
Signed-off-by: zengxian <xiangdong.zeng@intel.com>
2026-06-22 19:30:41 +08:00
78739e3bda [Bugfix] Reject matryoshka embedding dimensions above hidden size (#46313)
Signed-off-by: EazyReal <8047065+EazyReal@users.noreply.github.com>
Co-authored-by: EazyReal <8047065+EazyReal@users.noreply.github.com>
2026-06-22 10:16:35 +00:00
Tuukka SarviandGitHub 89accad2cc [ROCm][DSV4] Disable TileLang MHC dispatch on gfx942 (#45931)
Signed-off-by: Tuukka Sarvi <tuukka.sarvi@amd.com>
2026-06-22 09:26:54 +00:00
3c8e49596c [Model] ColQwen3.5: fix retrieval correctness (bias + bidirectional) (#46108)
Signed-off-by: Athrael Soju <athrael.soju@gmail.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
2026-06-22 17:25:54 +08:00
Weiwei SunGitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>Jiangyun Zhu
cec2ec1176 [Bugfix] Avoid racy accepted counts in async spec decode (#45100)
Signed-off-by: Weiwei Sun <68775773+sunnweiwei@users.noreply.github.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
Co-authored-by: Jiangyun Zhu <riverclouds.zhu@qq.com>
2026-06-22 08:53:16 +00:00
liuzhenweiandGitHub 435f82d61a [Bugfix] Fix Llama4ForCausalLM initialization test failure (#46341)
Signed-off-by: zhenwei-intel <zhenwei.liu@intel.com>
2026-06-22 08:40:43 +00:00
Roger WangandGitHub 1c4b51b990 Temporarily skip M3 on CI (#46352)
Signed-off-by: Roger Wang <hey@rogerw.io>
2026-06-22 01:35:31 -07:00
2e2c47928b [Doc] Update MiniMax-M3 (#45940)
Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai>
Signed-off-by: Roger Wang <hey@rogerw.io>
Co-authored-by: Jiangyun Zhu <riverclouds.zhu@qq.com>
Co-authored-by: Roger Wang <hey@rogerw.io>
2026-06-22 01:23:27 -07:00
Chao-Ju ChenGitHubBugen Zhaomergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
80abe0de7d [Rust Frontend] Support thinking_token_budget for chat and completions (#46137)
Co-authored-by: Bugen Zhao <i@bugenzhao.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
Signed-off-by: RickyChen / 陳昭儒 <ricky.chen@infinirc.com>
Signed-off-by: Bugen Zhao <i@bugenzhao.com>
2026-06-22 16:00:02 +08:00
a9f7b2d41c [feature][kv_offload] Self-describing KV events for OffloadingConnector (#43468)
Signed-off-by: Change72 <changg@nvidia.com>
Co-authored-by: Claude <noreply@anthropic.com>
2026-06-22 07:27:46 +00:00
Tiezhen WANGGitHubClaudemergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
d14e551a53 [Model] Remove MiniMaxText01, MiniMaxVL01, MiniMaxForCausalLM (#45993)
Signed-off-by: Xianbao QIAN <xianbao.qian@gmail.com>
Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-06-22 15:20:46 +08:00
68567ef2df [CPUOffloadingManager] Maintain evictable list in LRUCachePolicy (#46216)
Signed-off-by: <>
Co-authored-by: Varun Sundar Rabindranath <varun-sundar-rabindranath@h100-01.nemg-001.lab.rdu2.dc.redhat.com>
2026-06-22 06:54:44 +00:00
6bc6f2d86d [1/N][Core] add partial prefix cache primitives (#45939)
Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>
Co-authored-by: Yifan Qiao <yifanqiao@inferact.ai>
2026-06-21 23:43:10 -07:00
wang.yuqiandGitHub 1eb2cc961e [Frontend] Refactor ServingTokenization entrypoint. (#46022)
Signed-off-by: wang.yuqi <yuqi.wang@daocloud.io>
2026-06-22 06:27:58 +00:00
ReidGitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
31124749d1 [Bugfix] [Rust Frontend] Fix stop string truncation with repeated matches (#46113)
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
Signed-off-by: reidliu41 <reid201711@gmail.com>
2026-06-22 14:11:29 +08:00
Ma JianandGitHub 9037498c22 [DSV4][XPU] Pass gemm1_clamp_limit to XpuFusedMoe (#44517)
Signed-off-by: Ma Jian <jian1.ma@intel.com>
2026-06-22 12:57:10 +08:00
db32b53e30 [SpecDecode] Support DFlash with FlashInfer (#43081)
Signed-off-by: gss <2783977641@qq.com>
Co-authored-by: gss <2783977641@qq.com>
2026-06-22 04:55:30 +00:00
xiangdongandGitHub b529bfd6c5 [XPU][CI] Add agent_tags for Intel GPU CI (#45768)
Signed-off-by: zengxian <xiangdong.zeng@intel.com>
2026-06-22 10:33:17 +08:00
Micah WilliamsonandGitHub f3df7a7231 [ROCm][CI] Enable kv_connector unit tests on ROCm (#45955)
Signed-off-by: Micah Williamson <micah.williamson@amd.com>
2026-06-22 05:08:44 +03:00
485bbe1c6f [CI] Fix missing tp_size attribute on RoutedExperts (#46163)
Signed-off-by: Felix Marty <Felix.Marty@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
2026-06-21 18:46:49 -06:00
MattandGitHub a19ff2218a [Hardware][AMD][CI] Fix Spec Decode Eagle test group (#46018)
Signed-off-by: Matthew Wong <Matthew.Wong2@amd.com>
2026-06-21 17:40:02 -05:00
MattandGitHub 4f0d0049a0 [Hardware][AMD][CI] Fix Kernels Attention test groups (#46080)
Signed-off-by: Matthew Wong <Matthew.Wong2@amd.com>
2026-06-21 17:10:51 -05:00
13b83d77ad [ROCm][CI] skip test_double_aiter_rms_quant_fusion (#45967)
Signed-off-by: charlifu <charlifu@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
2026-06-21 16:53:11 -05:00
MattandGitHub 50241602fd [Hardware][AMD][CI] Fix gfx942 Kernels MoE test group (#46298)
Signed-off-by: Matthew Wong <Matthew.Wong2@amd.com>
2026-06-21 16:45:37 -05:00
Ting SUNandGitHub 12fe2a9aac [Bugfix][Qwen3-VL] Fix multi-video crash with list-valued fps/num_frames (#46305)
Signed-off-by: Ting Sun <suntcrick@gmail.com>
2026-06-21 14:31:23 -07:00
Benjamin ChislettandGitHub 89bd2c14d3 [Spec Decode] Add Qwen3 architecture support for EAGLE3 (#43132)
Signed-off-by: Benjamin Chislett <bchislett@nvidia.com>
2026-06-21 13:55:26 -07:00
ZedongLiuandGitHub 9c450b1027 [Kernel][Bugfix] Fix INT8 per-token-head KV cache rounding in Triton reshape-and-cache (#45361)
Signed-off-by: ZedongLiu <113341356+Zedong-Liu@users.noreply.github.com>
2026-06-21 15:59:40 -04:00
RanranGitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>Isotr0py
635c38338a [Multimodal] Add Qwen2-VL/Qwen2.5-VL processor-mapped video loader (#45555)
Signed-off-by: Ranran <hzz5361@psu.edu>
Signed-off-by: Ranran Haoran Zhang <ranzhang@redhat.com>
Signed-off-by: Isotr0py <mozf@mail2.sysu.edu.cn>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
Co-authored-by: Isotr0py <mozf@mail2.sysu.edu.cn>
2026-06-21 18:56:50 +00:00
c441ad1c07 [KV Offloading] Add labeled metrics support (#45957)
Signed-off-by: srinivas_oo7 <sklinkedin0120@gmail.com>
Co-authored-by: srinivas_oo7 <sklinkedin0120@gmail.com>
2026-06-21 18:04:01 +00:00
Jee Jee LiandGitHub 745bba5ea8 [Model]Fix MiniMaxM2ForCausalLM perf regression (#45935)
Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai>
2026-06-22 00:28:52 +08:00
2cac89f9da [Spec Decode] Support mixed KV page sizes for DFlash (#45181)
Signed-off-by: Alex Steiner <asteiner@nvidia.com>
Signed-off-by: Giancarlo Delfin <gdelfin@inferact.ai>
Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
Co-authored-by: Giancarlo Delfin <gdelfin@inferact.ai>
Co-authored-by: Yifan Qiao <yifanqiao@inferact.ai>
2026-06-21 22:45:14 +08:00
3e6e33526d [Disagg] return routed_experts on streaming generate responses (#44638)
Signed-off-by: aoshen02 <aoshen@inferact.ai>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Co-authored-by: Roger Wang <hey@rogerw.io>
2026-06-21 07:37:10 -07:00
junkang1991GitHubHongxia YangTan Pin SiangvllmellmChun FangTianDi101functionstackxtjtanaamergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
b91b7726e0 [ROCm][P/D] Support MiniMax-M3 mixed KV layouts in MoRIIO READ mode (#46039)
Signed-off-by: Jun Kang Chow <junkangchow@gmail.com>
Signed-off-by: tjtanaa <tunjian.tan@embeddedllm.com>
Co-authored-by: Hongxia Yang <hongxia.yang@amd.com>
Co-authored-by: Tan Pin Siang <tanpinsiang@gmail.com>
Co-authored-by: vllmellm <vllm.ellm@embeddedllm.com>
Co-authored-by: Chun Fang <chun.fang@amd.com>
Co-authored-by: TianDi101 <ditian12@amd.com>
Co-authored-by: functionstackx <47992694+functionstackx@users.noreply.github.com>
Co-authored-by: tjtanaa <tunjian.tan@embeddedllm.com>
Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
2026-06-21 12:55:19 +00:00
Palaiologos1453andGitHub d3ad8e8bcd [Bugfix] Defer offload reads while transfers are pending (#46231)
Signed-off-by: test test <2260891073@qq.com>
2026-06-21 14:30:13 +03:00
b80ce9dd2f [CI][test] Replace InternVL2-1B with InternVL3-1B in test_pipeline_parallel.py (#46241)
Signed-off-by: wentian-byte <192079369+wentian-byte@users.noreply.github.com>
Co-authored-by: wentian-byte <192079369+wentian-byte@users.noreply.github.com>
2026-06-21 15:11:19 +08:00
b5495cc5f9 Fix memory pointer overflow in Mamba state buffers (#44665)
Signed-off-by: Shifani Rajabose <shifani.rajabose@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>
2026-06-21 14:00:50 +08:00
199 changed files with 7249 additions and 5403 deletions
@@ -21,6 +21,10 @@ steps:
timeout_in_minutes: 30
optional: true
device: intel_gpu
agent_tags:
label: production
gpu: 2+
mem: 24+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -38,6 +42,10 @@ steps:
timeout_in_minutes: 30
optional: true
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -55,6 +63,10 @@ steps:
timeout_in_minutes: 30
optional: true
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -5,6 +5,10 @@ steps:
- label: XPU Sleep Mode
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
+4
View File
@@ -5,6 +5,10 @@ steps:
- label: Engine (1 GPU)
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -6,6 +6,10 @@ steps:
key: eplb-algorithm
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
+4
View File
@@ -5,6 +5,10 @@ steps:
- label: vLLM IR Tests
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
+24
View File
@@ -5,6 +5,10 @@ steps:
- label: LoRA Runtime + Utils
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 24+
no_plugin: true
working_dir: "."
env:
@@ -34,6 +38,10 @@ steps:
- label: LoRA Fused/MoE Kernels
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -54,6 +62,10 @@ steps:
- label: LoRA Punica Kernels
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -74,6 +86,10 @@ steps:
- label: LoRA Punica FP8/XPU Ops
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -94,6 +110,10 @@ steps:
- label: LoRA Models
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 2+
mem: 24+
no_plugin: true
working_dir: "."
env:
@@ -117,6 +137,10 @@ steps:
- label: LoRA Multimodal
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
+24
View File
@@ -5,6 +5,10 @@ steps:
- label: V1 Core + KV + Metrics
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -31,6 +35,10 @@ steps:
- label: V1 Sample + Logits
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -71,6 +79,10 @@ steps:
- label: XPU CPU Offload
timeout_in_minutes: 60
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -95,6 +107,10 @@ steps:
key: regression
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -126,6 +142,10 @@ steps:
timeout_in_minutes: 30
num_devices: 2
device: intel_gpu
agent_tags:
label: production
gpu: 2+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -157,6 +177,10 @@ steps:
key: async-engine-inputs-utils-worker
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 24+
no_plugin: true
working_dir: "."
env:
@@ -5,6 +5,10 @@ steps:
- label: Model Runner V2 Core Tests (Intel)
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 2+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -30,6 +34,10 @@ steps:
- label: Model Runner V2 Examples (Intel)
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 24+
no_plugin: true
working_dir: "."
env:
@@ -6,6 +6,10 @@ steps:
key: multi-modal-models-standard-1-qwen2
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -27,6 +31,10 @@ steps:
key: multi-modal-models-standard-2-qwen3-gemma
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -47,6 +55,10 @@ steps:
key: multi-modal-models-standard-3-llava-qwen2-vl
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 24+
no_plugin: true
working_dir: "."
env:
@@ -68,6 +80,10 @@ steps:
key: multi-modal-models-standard-4-other-whisper
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -88,6 +104,10 @@ steps:
key: multi-modal-processor
timeout_in_minutes: 45
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
+16
View File
@@ -19,6 +19,10 @@ steps:
- image-build-xpu
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 2+
mem: 24+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -49,6 +53,10 @@ steps:
- image-build-xpu
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -74,6 +82,10 @@ steps:
- image-build-xpu
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -93,6 +105,10 @@ steps:
- image-build-xpu
timeout_in_minutes: 30
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -34,7 +34,7 @@ case "${test_suite}" in
pytest -v -s v1/worker --ignore=v1/worker/test_gpu_model_runner.py --ignore=v1/worker/test_worker_memory_snapshot.py
pytest -v -s v1/structured_output
pytest -v -s v1/test_serial_utils.py
pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py
pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py --ignore=v1/spec_decode/test_speculators_correctness.py
pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --ignore=v1/kv_connector/unit/test_lmcache_integration.py --ignore=v1/kv_connector/unit/test_hf3fs_client.py --ignore=v1/kv_connector/unit/test_hf3fs_connector.py --ignore=v1/kv_connector/unit/test_hf3fs_metadata_server.py --ignore=v1/kv_connector/unit/test_offloading_connector.py
;;
server)
@@ -4,6 +4,11 @@
set -euo pipefail
if python3 -c "import torch; raise SystemExit(0 if torch.version.hip is not None else 1)"; then
uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
exit 0
fi
REQUIREMENTS_FILE="${KV_CONNECTORS_REQUIREMENTS:-/vllm-workspace/requirements/kv_connectors.txt}"
uv pip install --system -r "${REQUIREMENTS_FILE}"
+11 -9
View File
@@ -1594,9 +1594,10 @@ steps:
#---------------------------------------------------------- mi300 · kernels ----------------------------------------------------------#
- label: Kernels Attention Test %N # TBD
timeout_in_minutes: 180
timeout_in_minutes: 55
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
parallelism: 2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -1627,10 +1628,11 @@ steps:
- pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py kernels/test_concat_mla_q.py kernels/test_top_k_per_row.py
- label: Kernels MoE Test %N # TBD
timeout_in_minutes: 180
timeout_in_minutes: 50
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
parallelism: 4
optional: true
parallelism: 5
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- csrc/quantization/cutlass_w8a8/moe/
@@ -2120,9 +2122,10 @@ steps:
- pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
- label: Spec Decode Eagle # TBD
timeout_in_minutes: 180
timeout_in_minutes: 45
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/v1/spec_decode/
@@ -3040,7 +3043,7 @@ steps:
#---------------------------------------------------------- mi355 · kernels ----------------------------------------------------------#
- label: Kernels (B200-MI355) # TBD
timeout_in_minutes: 180
timeout_in_minutes: 15
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
working_dir: "/vllm-workspace/"
@@ -3064,11 +3067,10 @@ steps:
- pytest -v -s tests/kernels/attention/test_attention_selector.py
- label: Kernels Attention Test %N # TBD
timeout_in_minutes: 180
timeout_in_minutes: 60
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
parallelism: 2
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- csrc/attention/
@@ -3082,10 +3084,10 @@ steps:
- pytest -v -s kernels/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
- label: Kernels MoE Test %N # TBD
timeout_in_minutes: 180
timeout_in_minutes: 50
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
parallelism: 4
parallelism: 5
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- csrc/quantization/cutlass_w8a8/moe/
+3 -4
View File
@@ -12,7 +12,6 @@ steps:
- tests/basic_correctness/test_cpu_offload
- tests/basic_correctness/test_mem.py
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s basic_correctness/test_mem.py
- pytest -v -s basic_correctness/test_basic_correctness.py
- pytest -v -s basic_correctness/test_cpu_offload.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s basic_correctness/test_mem.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s basic_correctness/test_basic_correctness.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s basic_correctness/test_cpu_offload.py
+5 -10
View File
@@ -14,8 +14,7 @@ steps:
- vllm/v1/cudagraph_dispatcher.py
- tests/compile/correctness_e2e/test_sequence_parallel.py
commands:
- export VLLM_TEST_CLEAN_GPU_MEMORY=1
- pytest -v -s tests/compile/correctness_e2e/test_sequence_parallel.py
- VLLM_TEST_CLEAN_GPU_MEMORY=1 pytest -v -s tests/compile/correctness_e2e/test_sequence_parallel.py
- label: Sequence Parallel Correctness Tests (2xH100)
key: sequence-parallel-correctness-tests-2xh100
@@ -25,8 +24,7 @@ steps:
optional: true
num_devices: 2
commands:
- export VLLM_TEST_CLEAN_GPU_MEMORY=1
- pytest -v -s tests/compile/correctness_e2e/test_sequence_parallel.py
- VLLM_TEST_CLEAN_GPU_MEMORY=1 pytest -v -s tests/compile/correctness_e2e/test_sequence_parallel.py
- label: AsyncTP Correctness Tests (2xH100)
key: asynctp-correctness-tests-2xh100
@@ -36,8 +34,7 @@ steps:
optional: true
num_devices: 2
commands:
- export VLLM_TEST_CLEAN_GPU_MEMORY=1
- pytest -v -s tests/compile/correctness_e2e/test_async_tp.py
- VLLM_TEST_CLEAN_GPU_MEMORY=1 pytest -v -s tests/compile/correctness_e2e/test_async_tp.py
- label: AsyncTP Correctness Tests (B200)
key: asynctp-correctness-tests-b200
@@ -47,8 +44,7 @@ steps:
optional: true
num_devices: 2
commands:
- export VLLM_TEST_CLEAN_GPU_MEMORY=1
- pytest -v -s tests/compile/correctness_e2e/test_async_tp.py
- VLLM_TEST_CLEAN_GPU_MEMORY=1 pytest -v -s tests/compile/correctness_e2e/test_async_tp.py
- label: Distributed Compile Unit Tests (2xH100)
key: distributed-compile-unit-tests-2xh100
@@ -61,8 +57,7 @@ steps:
- vllm/model_executor/layers
- tests/compile/passes/distributed/
commands:
- export VLLM_TEST_CLEAN_GPU_MEMORY=1
- pytest -s -v tests/compile/passes/distributed
- VLLM_TEST_CLEAN_GPU_MEMORY=1 pytest -s -v tests/compile/passes/distributed
- label: Fusion and Compile Unit Tests (2xB200)
key: fusion-and-compile-unit-tests-2xb200
+37 -46
View File
@@ -32,11 +32,10 @@ steps:
- tests/entrypoints/openai/test_multi_api_servers.py
commands:
# https://github.com/NVIDIA/nccl/issues/1838
- export NCCL_CUMEM_HOST_ENABLE=0
- TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
- TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
- TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
- DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py
- NCCL_CUMEM_HOST_ENABLE=0 TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
- NCCL_CUMEM_HOST_ENABLE=0 TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
- NCCL_CUMEM_HOST_ENABLE=0 TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
- NCCL_CUMEM_HOST_ENABLE=0 DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py
- label: Distributed Compile + RPC Tests (2 GPUs)
key: distributed-compile-rpc-tests-2-gpus
@@ -56,10 +55,9 @@ steps:
- tests/entrypoints/llm/test_collective_rpc.py
commands:
# https://github.com/NVIDIA/nccl/issues/1838
- export NCCL_CUMEM_HOST_ENABLE=0
- pytest -v -s entrypoints/llm/test_collective_rpc.py
- pytest -v -s ./compile/fullgraph/test_basic_correctness.py
- pytest -v -s ./compile/test_wrapper.py
- NCCL_CUMEM_HOST_ENABLE=0 pytest -v -s entrypoints/llm/test_collective_rpc.py
- NCCL_CUMEM_HOST_ENABLE=0 pytest -v -s ./compile/fullgraph/test_basic_correctness.py
- NCCL_CUMEM_HOST_ENABLE=0 pytest -v -s ./compile/test_wrapper.py
- label: Distributed Torchrun + Shutdown Tests (2 GPUs)
key: distributed-torchrun-shutdown-tests-2-gpus
@@ -78,11 +76,10 @@ steps:
- tests/v1/worker/test_worker_memory_snapshot.py
commands:
# https://github.com/NVIDIA/nccl/issues/1838
- export NCCL_CUMEM_HOST_ENABLE=0
- VLLM_TEST_SAME_HOST=1 torchrun --nproc-per-node=4 distributed/test_same_node.py | grep 'Same node test passed'
- VLLM_TEST_SAME_HOST=1 VLLM_TEST_WITH_DEFAULT_DEVICE_SET=1 torchrun --nproc-per-node=4 distributed/test_same_node.py | grep 'Same node test passed'
- CUDA_VISIBLE_DEVICES=0,1 pytest -v -s v1/shutdown
- pytest -v -s v1/worker/test_worker_memory_snapshot.py
- NCCL_CUMEM_HOST_ENABLE=0 VLLM_TEST_SAME_HOST=1 torchrun --nproc-per-node=4 distributed/test_same_node.py | grep 'Same node test passed'
- NCCL_CUMEM_HOST_ENABLE=0 VLLM_TEST_SAME_HOST=1 VLLM_TEST_WITH_DEFAULT_DEVICE_SET=1 torchrun --nproc-per-node=4 distributed/test_same_node.py | grep 'Same node test passed'
- NCCL_CUMEM_HOST_ENABLE=0 CUDA_VISIBLE_DEVICES=0,1 pytest -v -s v1/shutdown
- NCCL_CUMEM_HOST_ENABLE=0 pytest -v -s v1/worker/test_worker_memory_snapshot.py
- label: Distributed Torchrun + Examples (4 GPUs)
key: distributed-torchrun-examples-4-gpus
@@ -97,24 +94,23 @@ steps:
- tests/examples/features/data_parallel/data_parallel_offline.py
commands:
# https://github.com/NVIDIA/nccl/issues/1838
- export NCCL_CUMEM_HOST_ENABLE=0
# test with torchrun tp=2 and external_dp=2
- torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example.py
- NCCL_CUMEM_HOST_ENABLE=0 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example.py
# test with torchrun tp=2 and pp=2
- PP_SIZE=2 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example.py
- NCCL_CUMEM_HOST_ENABLE=0 PP_SIZE=2 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example.py
# test with torchrun tp=4 and dp=1
- TP_SIZE=4 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example_moe.py
- NCCL_CUMEM_HOST_ENABLE=0 TP_SIZE=4 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example_moe.py
# test with torchrun tp=2, pp=2 and dp=1
- PP_SIZE=2 TP_SIZE=2 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example_moe.py
- NCCL_CUMEM_HOST_ENABLE=0 PP_SIZE=2 TP_SIZE=2 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example_moe.py
# test with torchrun tp=1 and dp=4 with ep
- DP_SIZE=4 ENABLE_EP=1 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example_moe.py
- NCCL_CUMEM_HOST_ENABLE=0 DP_SIZE=4 ENABLE_EP=1 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example_moe.py
# test with torchrun tp=2 and dp=2 with ep
- TP_SIZE=2 DP_SIZE=2 ENABLE_EP=1 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example_moe.py
- NCCL_CUMEM_HOST_ENABLE=0 TP_SIZE=2 DP_SIZE=2 ENABLE_EP=1 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example_moe.py
# test with internal dp
- python3 examples/features/data_parallel/data_parallel_offline.py --enforce-eager
- NCCL_CUMEM_HOST_ENABLE=0 python3 examples/features/data_parallel/data_parallel_offline.py --enforce-eager
# rlhf examples
- VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_nccl.py
- VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_ipc.py
- NCCL_CUMEM_HOST_ENABLE=0 VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_nccl.py
- NCCL_CUMEM_HOST_ENABLE=0 VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_ipc.py
- label: Distributed DP Tests (4 GPUs)
key: distributed-dp-tests-4-gpus
@@ -128,14 +124,13 @@ steps:
- tests/distributed/test_utils
commands:
# https://github.com/NVIDIA/nccl/issues/1838
- export NCCL_CUMEM_HOST_ENABLE=0
- TP_SIZE=2 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
- TP_SIZE=2 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
- TP_SIZE=2 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
- TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_internal_lb_dp.py
- TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_hybrid_lb_dp.py
- pytest -v -s v1/engine/test_engine_core_client.py::test_kv_cache_events_dp
- pytest -v -s distributed/test_utils.py
- NCCL_CUMEM_HOST_ENABLE=0 TP_SIZE=2 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
- NCCL_CUMEM_HOST_ENABLE=0 TP_SIZE=2 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
- NCCL_CUMEM_HOST_ENABLE=0 TP_SIZE=2 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
- NCCL_CUMEM_HOST_ENABLE=0 TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_internal_lb_dp.py
- NCCL_CUMEM_HOST_ENABLE=0 TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_hybrid_lb_dp.py
- NCCL_CUMEM_HOST_ENABLE=0 pytest -v -s v1/engine/test_engine_core_client.py::test_kv_cache_events_dp
- NCCL_CUMEM_HOST_ENABLE=0 pytest -v -s distributed/test_utils.py
- label: Distributed Compile + Comm (4 GPUs)
key: distributed-compile-comm-4-gpus
@@ -151,13 +146,12 @@ steps:
- tests/distributed/test_multiproc_executor.py
commands:
# https://github.com/NVIDIA/nccl/issues/1838
- export NCCL_CUMEM_HOST_ENABLE=0
- pytest -v -s compile/fullgraph/test_basic_correctness.py
- pytest -v -s distributed/test_pynccl.py
- pytest -v -s distributed/test_events.py
- pytest -v -s distributed/test_symm_mem_allreduce.py
- NCCL_CUMEM_HOST_ENABLE=0 pytest -v -s compile/fullgraph/test_basic_correctness.py
- NCCL_CUMEM_HOST_ENABLE=0 pytest -v -s distributed/test_pynccl.py
- NCCL_CUMEM_HOST_ENABLE=0 pytest -v -s distributed/test_events.py
- NCCL_CUMEM_HOST_ENABLE=0 pytest -v -s distributed/test_symm_mem_allreduce.py
# test multi-node TP with multiproc executor (simulated on single node)
- pytest -v -s distributed/test_multiproc_executor.py::test_multiproc_executor_multi_node
- NCCL_CUMEM_HOST_ENABLE=0 pytest -v -s distributed/test_multiproc_executor.py::test_multiproc_executor_multi_node
- label: Distributed Tests (8 GPUs)(H100)
key: distributed-tests-8-gpus-h100
@@ -176,9 +170,8 @@ steps:
commands:
# https://github.com/NVIDIA/nccl/issues/1838
- export NCCL_CUMEM_HOST_ENABLE=0
# test with torchrun tp=2 and dp=4 with ep
- torchrun --nproc-per-node=8 ../examples/features/torchrun/torchrun_dp_example_offline.py --tp-size=2 --pp-size=1 --dp-size=4 --enable-ep
- NCCL_CUMEM_HOST_ENABLE=0 torchrun --nproc-per-node=8 ../examples/features/torchrun/torchrun_dp_example_offline.py --tp-size=2 --pp-size=1 --dp-size=4 --enable-ep
- label: Distributed Tests (4 GPUs)(A100)
key: distributed-tests-4-gpus-a100
@@ -271,9 +264,7 @@ steps:
- tests/distributed/test_pipeline_parallel.py
- tests/basic_correctness/test_basic_correctness.py
commands:
- export VLLM_USE_RAY_V2_EXECUTOR_BACKEND=1
- export NCCL_CUMEM_HOST_ENABLE=0
- pytest -v -s distributed/test_ray_v2_executor.py
- pytest -v -s distributed/test_ray_v2_executor_e2e.py
- pytest -v -s distributed/test_pipeline_parallel.py -k "ray"
- TARGET_TEST_SUITE=L4 pytest -v -s basic_correctness/test_basic_correctness.py -k "ray"
- VLLM_USE_RAY_V2_EXECUTOR_BACKEND=1 NCCL_CUMEM_HOST_ENABLE=0 pytest -v -s distributed/test_ray_v2_executor.py
- VLLM_USE_RAY_V2_EXECUTOR_BACKEND=1 NCCL_CUMEM_HOST_ENABLE=0 pytest -v -s distributed/test_ray_v2_executor_e2e.py
- VLLM_USE_RAY_V2_EXECUTOR_BACKEND=1 NCCL_CUMEM_HOST_ENABLE=0 pytest -v -s distributed/test_pipeline_parallel.py -k "ray"
- VLLM_USE_RAY_V2_EXECUTOR_BACKEND=1 NCCL_CUMEM_HOST_ENABLE=0 TARGET_TEST_SUITE=L4 pytest -v -s basic_correctness/test_basic_correctness.py -k "ray"
+11 -18
View File
@@ -22,10 +22,9 @@ steps:
- vllm/
- tests/entrypoints/llm
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/llm --ignore=entrypoints/llm/test_generate.py --ignore=entrypoints/llm/test_collective_rpc.py --ignore=entrypoints/llm/offline_mode
- pytest -v -s entrypoints/llm/test_generate.py # it needs a clean process
- pytest -v -s entrypoints/llm/offline_mode # Needs to avoid interference with other tests
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/llm --ignore=entrypoints/llm/test_generate.py --ignore=entrypoints/llm/test_collective_rpc.py --ignore=entrypoints/llm/offline_mode
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/llm/test_generate.py # it needs a clean process
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/llm/offline_mode # Needs to avoid interference with other tests
mirror:
amd:
device: mi325_1
@@ -41,9 +40,8 @@ steps:
- vllm/
- tests/entrypoints/serve
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
- VLLM_WORKER_MULTIPROC_METHOD=spawn PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
mirror:
amd:
device: mi325_1
@@ -59,8 +57,7 @@ steps:
- tests/entrypoints/openai
- tests/entrypoints/test_chat_utils
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/openai --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/correctness
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/openai --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/correctness
mirror:
amd:
device: mi325_1
@@ -77,9 +74,8 @@ steps:
- tests/entrypoints/openai
- tests/entrypoints/test_chat_utils
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/openai/chat_completion
- pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/openai/chat_completion
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
mirror:
amd:
device: mi325_1
@@ -128,8 +124,7 @@ steps:
- vllm/
- tests/entrypoints/speech_to_text
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/speech_to_text
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/speech_to_text
- label: Entrypoints Integration (Multimodal)
device: h200_35gb
@@ -140,8 +135,7 @@ steps:
- vllm/
- tests/entrypoints/multimodal
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/multimodal
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/multimodal
- label: Entrypoints Integration (Pooling)
key: entrypoints-integration-pooling
@@ -151,8 +145,7 @@ steps:
- vllm/
- tests/entrypoints/pooling
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/pooling
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/pooling
- label: OpenAI API Correctness
key: openai-api-correctness
+30
View File
@@ -74,6 +74,20 @@ steps:
commands:
- pytest -v -s kernels/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
parallelism: 2
mirror:
amd:
device: mi325_1
timeout_in_minutes: 55
depends_on:
- image-build-amd
source_file_dependencies:
- csrc/attention/
- vllm/v1/attention
- vllm/model_executor/layers/attention
- tests/kernels/attention
- vllm/_aiter_ops.py
- vllm/envs.py
- vllm/platforms/rocm.py
- label: Kernels Attention DiffKV Test (H100)
key: kernels-attention-diffkv-test-h100
@@ -128,6 +142,22 @@ steps:
- pytest -v -s kernels/moe --ignore=kernels/moe/test_modular_oai_triton_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
- pytest -v -s kernels/moe/test_modular_oai_triton_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
parallelism: 5
mirror:
amd:
device: mi325_1
timeout_in_minutes: 50
source_file_dependencies:
- csrc/quantization/cutlass_w8a8/moe/
- csrc/moe/
- tests/kernels/moe
- vllm/model_executor/layers/fused_moe/
- vllm/distributed/device_communicators/
- vllm/envs.py
- vllm/config
- vllm/_aiter_ops.py
- vllm/platforms/rocm.py
depends_on:
- image-build-amd
- label: Kernels Mamba Test
key: kernels-mamba-test
+2 -5
View File
@@ -50,8 +50,7 @@ steps:
- csrc/
- vllm/model_executor/layers/quantization
commands:
- export VLLM_USE_DEEP_GEMM=0 # We found Triton is faster than DeepGEMM for H100
- pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-hopper.txt --tp-size=4
- VLLM_USE_DEEP_GEMM=0 pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-hopper.txt --tp-size=4 # Triton is faster than DeepGEMM for H100
- label: LM Eval Small Models (B200)
key: lm-eval-small-models-b200
@@ -108,9 +107,7 @@ steps:
depends_on:
- image-build-amd
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- export PYTORCH_ROCM_ARCH=gfx942 # Limit Quark compilation to save time
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx.txt
- VLLM_WORKER_MULTIPROC_METHOD=spawn PYTORCH_ROCM_ARCH=gfx942 pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx.txt # Limit Quark compilation to save time
- label: MoE Refactor Integration Test (H100 - TEMPORARY)
key: moe-refactor-integration-test-h100-temporary
+8 -8
View File
@@ -36,14 +36,14 @@ steps:
commands:
# FIXIT: find out which code initialize cuda before running the test
# before the fix, we need to use spawn to test it
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
#
# Alot of these tests are on the edge of OOMing
- export PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True
#
# There is some Tensor Parallelism related processing logic in LoRA that
# requires multi-GPU testing for validation.
- pytest -v -s -x lora/test_chatglm3_tp.py
- pytest -v -s -x lora/test_llama_tp.py
- pytest -v -s -x lora/test_qwen3_with_multi_loras.py
- pytest -v -s -x lora/test_olmoe_tp.py
- pytest -v -s -x lora/test_gptoss_tp.py
- pytest -v -s -x lora/test_qwen35_densemodel_lora.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True pytest -v -s -x lora/test_chatglm3_tp.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True pytest -v -s -x lora/test_llama_tp.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True pytest -v -s -x lora/test_qwen3_with_multi_loras.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True pytest -v -s -x lora/test_olmoe_tp.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True pytest -v -s -x lora/test_gptoss_tp.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True pytest -v -s -x lora/test_qwen35_densemodel_lora.py
+35 -38
View File
@@ -18,9 +18,8 @@ steps:
- vllm/v1/
- tests/v1/spec_decode
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
# TODO: create another `optional` test group for slow tests
- pytest -v -s -m 'not slow_test' v1/spec_decode
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s -m 'not slow_test' v1/spec_decode
mirror:
amd:
device: mi300_1
@@ -50,12 +49,11 @@ steps:
- tests/v1/test_request.py
- tests/v1/test_outputs.py
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s v1/sample
- pytest -v -s v1/logits_processors
- pytest -v -s v1/test_oracle.py
- pytest -v -s v1/test_request.py
- pytest -v -s v1/test_outputs.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/sample
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/logits_processors
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/test_oracle.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/test_request.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/test_outputs.py
mirror:
amd:
device: mi325_1
@@ -93,18 +91,23 @@ steps:
- tests/entrypoints/openai/correctness/test_lmeval.py
commands:
- 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
- pytest -v -s v1/executor
- pytest -v -s v1/kv_offload
- pytest -v -s v1/simple_kv_offload
- pytest -v -s v1/worker
- pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
- pytest -v -s -m 'not cpu_test' v1/metrics
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s -m 'not cpu_test' v1/core
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/executor
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/kv_offload
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/simple_kv_offload
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/worker
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s -m 'not cpu_test' v1/metrics
# Integration test for streaming correctness (requires special branch).
- pip install -U git+https://github.com/robertgshaw2-redhat/lm-evaluation-harness.git@streaming-api
- pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
mirror:
amd:
device: mi325_1
timeout_in_minutes: 60
depends_on:
- image-build-amd
- label: V1 Others (CPU)
key: v1-others-cpu
@@ -147,8 +150,7 @@ steps:
- vllm/distributed/kv_transfer/kv_connector/v1/example_hidden_states_connector.py
- tests/v1/kv_connector/extract_hidden_states_integration
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s v1/kv_connector/extract_hidden_states_integration
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/kv_connector/extract_hidden_states_integration
- label: Extract Hidden States Integration (2 GPUs)
key: extract-hidden-states-integration-2-gpus
@@ -161,8 +163,7 @@ steps:
- vllm/distributed/kv_transfer/kv_connector/v1/example_hidden_states_connector.py
- tests/v1/kv_connector/extract_hidden_states_integration
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s -m 'distributed' v1/kv_connector/extract_hidden_states_integration
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s -m 'distributed' v1/kv_connector/extract_hidden_states_integration
- label: Regression
key: regression
@@ -354,10 +355,9 @@ steps:
- vllm/model_executor/layers
- tests/v1/determinism/
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pip install pytest-timeout pytest-forked
- pytest -v -s v1/determinism/test_batch_invariance.py
- VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[TRITON_MLA]
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/determinism/test_batch_invariance.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[TRITON_MLA]
- label: Batch Invariance (H100)
key: batch-invariance-h100
@@ -368,12 +368,11 @@ steps:
- vllm/model_executor/layers
- tests/v1/determinism/
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pip install pytest-timeout pytest-forked
- pytest -v -s v1/determinism/test_batch_invariance.py
- pytest -v -s v1/determinism/test_rms_norm_batch_invariant.py
- VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[TRITON_MLA]
- VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[FLASH_ATTN]
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/determinism/test_batch_invariance.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/determinism/test_rms_norm_batch_invariant.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[TRITON_MLA]
- VLLM_WORKER_MULTIPROC_METHOD=spawn VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[FLASH_ATTN]
- label: Batch Invariance (B200)
key: batch-invariance-b200
@@ -384,14 +383,13 @@ steps:
- vllm/model_executor/layers
- tests/v1/determinism/
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pip install pytest-timeout pytest-forked
- pytest -v -s v1/determinism/test_batch_invariance.py
- pytest -v -s v1/determinism/test_rms_norm_batch_invariant.py
- VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[TRITON_MLA]
- VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[FLASH_ATTN]
- pytest -v -s v1/determinism/test_nvfp4_batch_invariant.py
- pytest -v -s v1/determinism/test_nvfp4_batch_invariant_scaled_mm.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/determinism/test_batch_invariance.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/determinism/test_rms_norm_batch_invariant.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[TRITON_MLA]
- VLLM_WORKER_MULTIPROC_METHOD=spawn VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[FLASH_ATTN]
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/determinism/test_nvfp4_batch_invariant.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/determinism/test_nvfp4_batch_invariant_scaled_mm.py
- label: Acceptance Length Test (Large Models) # optional
device: h200_35gb
@@ -406,5 +404,4 @@ steps:
- vllm/model_executor/models/mlp_speculator.py
- tests/v1/spec_decode/test_acceptance_length.py
commands:
- export VLLM_ALLOW_INSECURE_SERIALIZATION=1
- pytest -v -s v1/spec_decode/test_acceptance_length.py -m slow_test
- VLLM_ALLOW_INSECURE_SERIALIZATION=1 pytest -v -s v1/spec_decode/test_acceptance_length.py -m slow_test
+5 -6
View File
@@ -13,13 +13,12 @@ steps:
- tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
commands:
- apt-get update && apt-get install -y curl libsodium23
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
# Dump tracebacks of all threads if a test hangs, so a wedged GPU/CUDA
# init surfaces a stack instead of silently stalling.
- export PYTHONFAULTHANDLER=1
# Per-test watchdog: a single hung test (e.g. stuck during engine/CUDA
# init) fails fast with a traceback instead of running until the global
# build timeout. The `thread` method also handles hangs inside C/CUDA
# calls that the signal method cannot interrupt.
- pytest -v -s model_executor -m '(not slow_test)' --timeout=900 --timeout-method=thread
- pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py --timeout=900 --timeout-method=thread
#
# Env vars are inlined because CONTINUE_ON_FAILURE wraps each command
# in a subshell, so a standalone `export` would be lost.
- VLLM_WORKER_MULTIPROC_METHOD=spawn PYTHONFAULTHANDLER=1 pytest -v -s model_executor -m '(not slow_test)' --timeout=900 --timeout-method=thread
- VLLM_WORKER_MULTIPROC_METHOD=spawn PYTHONFAULTHANDLER=1 pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py --timeout=900 --timeout-method=thread
+29 -35
View File
@@ -16,15 +16,14 @@ steps:
- tests/entrypoints/llm/test_struct_output_generate.py
commands:
- set -x
- export VLLM_USE_V2_MODEL_RUNNER=1
- pytest -v -s v1/engine/test_llm_engine.py -k "not test_engine_metrics"
- VLLM_USE_V2_MODEL_RUNNER=1 pytest -v -s v1/engine/test_llm_engine.py -k "not test_engine_metrics"
# This requires eager until we sort out CG correctness issues.
# TODO: remove ENFORCE_EAGER here after https://github.com/vllm-project/vllm/pull/32936 is merged.
- ENFORCE_EAGER=1 pytest -v -s v1/e2e/general/test_async_scheduling.py -k "not ngram"
- pytest -v -s v1/e2e/general/test_context_length.py
- pytest -v -s v1/e2e/general/test_min_tokens.py
- VLLM_USE_V2_MODEL_RUNNER=1 ENFORCE_EAGER=1 pytest -v -s v1/e2e/general/test_async_scheduling.py -k "not ngram"
- VLLM_USE_V2_MODEL_RUNNER=1 pytest -v -s v1/e2e/general/test_context_length.py
- VLLM_USE_V2_MODEL_RUNNER=1 pytest -v -s v1/e2e/general/test_min_tokens.py
# Temporary hack filter to exclude ngram spec decoding based tests.
- pytest -v -s entrypoints/llm/test_struct_output_generate.py -k "xgrammar and not speculative_config6 and not speculative_config7 and not speculative_config8 and not speculative_config0"
- VLLM_USE_V2_MODEL_RUNNER=1 pytest -v -s entrypoints/llm/test_struct_output_generate.py -k "xgrammar and not speculative_config6 and not speculative_config7 and not speculative_config8 and not speculative_config0"
- label: Model Runner V2 Examples
device: h200_35gb
@@ -42,26 +41,25 @@ steps:
- examples/features/tensorize_vllm_model.py
commands:
- set -x
- export VLLM_USE_V2_MODEL_RUNNER=1
- pip install tensorizer # for tensorizer test
- python3 basic/offline_inference/chat.py # for basic
- python3 basic/offline_inference/generate.py --model facebook/opt-125m
#- python3 basic/offline_inference/generate.py --model meta-llama/Llama-2-13b-chat-hf --cpu-offload-gb 10 # TODO
#- python3 basic/offline_inference/embed.py # TODO
- VLLM_USE_V2_MODEL_RUNNER=1 python3 basic/offline_inference/chat.py # for basic
- VLLM_USE_V2_MODEL_RUNNER=1 python3 basic/offline_inference/generate.py --model facebook/opt-125m
#- VLLM_USE_V2_MODEL_RUNNER=1 python3 basic/offline_inference/generate.py --model meta-llama/Llama-2-13b-chat-hf --cpu-offload-gb 10 # TODO
#- VLLM_USE_V2_MODEL_RUNNER=1 python3 basic/offline_inference/embed.py # TODO
# for multi-modal models
- python3 generate/multimodal/audio_language_offline.py --seed 0
- python3 generate/multimodal/vision_language_offline.py --seed 0
- python3 generate/multimodal/vision_language_multi_image_offline.py --seed 0
- python3 generate/multimodal/encoder_decoder_multimodal_offline.py --model-type whisper --seed 0
- VLLM_USE_V2_MODEL_RUNNER=1 python3 generate/multimodal/audio_language_offline.py --seed 0
- VLLM_USE_V2_MODEL_RUNNER=1 python3 generate/multimodal/vision_language_offline.py --seed 0
- VLLM_USE_V2_MODEL_RUNNER=1 python3 generate/multimodal/vision_language_multi_image_offline.py --seed 0
- VLLM_USE_V2_MODEL_RUNNER=1 python3 generate/multimodal/encoder_decoder_multimodal_offline.py --model-type whisper --seed 0
# for pooling models
- python3 pooling/embed/vision_embedding_offline.py --seed 0
- VLLM_USE_V2_MODEL_RUNNER=1 python3 pooling/embed/vision_embedding_offline.py --seed 0
# for features demo
- python3 features/automatic_prefix_caching/prefix_caching_offline.py
- python3 deployment/llm_engine_example.py
- python3 features/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 features/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
- python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
- VLLM_USE_V2_MODEL_RUNNER=1 python3 features/automatic_prefix_caching/prefix_caching_offline.py
- VLLM_USE_V2_MODEL_RUNNER=1 python3 deployment/llm_engine_example.py
- VLLM_USE_V2_MODEL_RUNNER=1 python3 features/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && VLLM_USE_V2_MODEL_RUNNER=1 python3 features/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
- VLLM_USE_V2_MODEL_RUNNER=1 python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
# https://github.com/vllm-project/vllm/pull/26682 uses slightly more memory in PyTorch 2.9+ causing this test to OOM in 1xL4 GPU
- python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
- VLLM_USE_V2_MODEL_RUNNER=1 python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
- label: Model Runner V2 Distributed (2 GPUs)
key: model-runner-v2-distributed-2-gpus
@@ -76,13 +74,11 @@ steps:
- tests/v1/distributed/test_eagle_dp.py
commands:
- set -x
- export VLLM_USE_V2_MODEL_RUNNER=1
# The "and not True" here is a hacky way to exclude the prompt_embeds cases which aren't yet supported.
- TARGET_TEST_SUITE=L4 pytest -v -s basic_correctness/test_basic_correctness.py -m 'distributed(num_gpus=2)' -k "not ray and not True"
- VLLM_USE_V2_MODEL_RUNNER=1 TARGET_TEST_SUITE=L4 pytest -v -s basic_correctness/test_basic_correctness.py -m 'distributed(num_gpus=2)' -k "not ray and not True"
# https://github.com/NVIDIA/nccl/issues/1838
- export NCCL_CUMEM_HOST_ENABLE=0
- TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py -k "not ray"
- TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
- VLLM_USE_V2_MODEL_RUNNER=1 NCCL_CUMEM_HOST_ENABLE=0 TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py -k "not ray"
- VLLM_USE_V2_MODEL_RUNNER=1 NCCL_CUMEM_HOST_ENABLE=0 TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
- label: Model Runner V2 Pipeline Parallelism (4 GPUs)
key: model-runner-v2-pipeline-parallelism-4-gpus
@@ -97,10 +93,9 @@ steps:
- tests/v1/distributed/test_pp_dp_v2.py
commands:
- set -x
- export VLLM_USE_V2_MODEL_RUNNER=1
- pytest -v -s distributed/test_pipeline_parallel.py -k "not ray and not Jamba"
- pytest -v -s distributed/test_pp_cudagraph.py -k "not ray"
- pytest -v -s v1/distributed/test_pp_dp_v2.py
- VLLM_USE_V2_MODEL_RUNNER=1 pytest -v -s distributed/test_pipeline_parallel.py -k "not ray and not Jamba"
- VLLM_USE_V2_MODEL_RUNNER=1 pytest -v -s distributed/test_pp_cudagraph.py -k "not ray"
- VLLM_USE_V2_MODEL_RUNNER=1 pytest -v -s v1/distributed/test_pp_dp_v2.py
- label: Model Runner V2 Spec Decode
device: h200_35gb
@@ -115,8 +110,7 @@ steps:
- tests/v1/e2e/spec_decode/test_spec_decode.py
commands:
- set -x
- export VLLM_USE_V2_MODEL_RUNNER=1
- pytest -v -s v1/spec_decode/test_max_len.py -k "eagle or mtp"
- pytest -v -s v1/spec_decode/test_rejection_sampler_utils.py
- pytest -v -s v1/spec_decode/test_synthetic_rejection_sampler_utils.py
- pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "eagle or mtp"
- VLLM_USE_V2_MODEL_RUNNER=1 pytest -v -s v1/spec_decode/test_max_len.py -k "eagle or mtp"
- VLLM_USE_V2_MODEL_RUNNER=1 pytest -v -s v1/spec_decode/test_rejection_sampler_utils.py
- VLLM_USE_V2_MODEL_RUNNER=1 pytest -v -s v1/spec_decode/test_synthetic_rejection_sampler_utils.py
- VLLM_USE_V2_MODEL_RUNNER=1 pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "eagle or mtp"
+6
View File
@@ -107,6 +107,12 @@ steps:
- tests/compile/passes
commands:
- pytest -s -v compile/passes --ignore compile/passes/distributed
mirror:
amd:
device: mi300_1
timeout_in_minutes: 180
depends_on:
- image-build-amd
- label: PyTorch Fullgraph Smoke Test
key: pytorch-fullgraph-smoke-test
+18 -29
View File
@@ -23,17 +23,15 @@ steps:
# - tests/entrypoints/openai/test_uds.py
- tests/v1/sample/test_logprobs_e2e.py
commands:
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s benchmarks/test_serve_cli.py -k "not insecure and not (test_bench_serve and not test_bench_serve_chat)"
- pytest -v -s entrypoints/openai/chat_completion/test_chat_completion.py
# - pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not invalid"
- VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s benchmarks/test_serve_cli.py -k "not insecure and not (test_bench_serve and not test_bench_serve_chat)"
- VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/openai/chat_completion/test_chat_completion.py
# - VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not invalid"
# - pytest -v -s entrypoints/openai/completion/test_prompt_validation.py -k "not prompt_embeds"
- pytest -v -s entrypoints/openai/completion/test_shutdown.py -k "not engine_failure and not test_abort_timeout_exits_quickly"
# - pytest -v -s entrypoints/openai/test_return_token_ids.py
# - pytest -v -s entrypoints/openai/test_uds.py
- pytest -v -s v1/sample/test_logprobs_e2e.py -k "test_prompt_logprobs_e2e_server"
# - VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/openai/completion/test_prompt_validation.py -k "not prompt_embeds"
- VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/openai/completion/test_shutdown.py -k "not engine_failure and not test_abort_timeout_exits_quickly"
# - VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/openai/test_return_token_ids.py
# - VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/openai/test_uds.py
- VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s v1/sample/test_logprobs_e2e.py -k "test_prompt_logprobs_e2e_server"
- label: Rust Frontend Serve/Admin Coverage
timeout_in_minutes: 60
@@ -51,13 +49,11 @@ steps:
- tests/entrypoints/serve/instrumentator/test_metrics.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/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/dev/test_sleep.py
# - VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
- VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not show_version and not server_load"
- VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn 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"
- VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn 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"
# - VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s entrypoints/serve/dev/test_sleep.py
- label: Rust Frontend Core Correctness
timeout_in_minutes: 30
@@ -69,9 +65,7 @@ steps:
- tests/utils.py
- tests/entrypoints/openai/correctness/test_lmeval.py
commands:
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
- VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
- label: Rust Frontend Tool Use
timeout_in_minutes: 60
@@ -83,9 +77,7 @@ steps:
- tests/utils.py
- tests/tool_use/
commands:
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s tool_use --ignore=tool_use/mistral --models llama3.2 -k "not test_response_format_with_tool_choice_required and not test_parallel_tool_calls_false and not test_tool_call_and_choice"
- VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s tool_use --ignore=tool_use/mistral --models llama3.2 -k "not test_response_format_with_tool_choice_required and not test_parallel_tool_calls_false and not test_tool_call_and_choice"
- label: Rust Frontend Distributed
timeout_in_minutes: 30
@@ -103,9 +95,6 @@ steps:
- tests/v1/distributed/test_hybrid_lb_dp.py
- tests/v1/distributed/test_internal_lb_dp.py
commands:
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- export NCCL_CUMEM_HOST_ENABLE=0
- TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_internal_lb_dp.py -k "not 4 and not server_info"
- TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py -k "not 4 and not server_info"
- TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_hybrid_lb_dp.py -k "not 4 and not server_info"
- VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn NCCL_CUMEM_HOST_ENABLE=0 TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_internal_lb_dp.py -k "not 4 and not server_info"
- VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn NCCL_CUMEM_HOST_ENABLE=0 TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py -k "not 4 and not server_info"
- VLLM_USE_RUST_FRONTEND=1 VLLM_WORKER_MULTIPROC_METHOD=spawn NCCL_CUMEM_HOST_ENABLE=0 TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_hybrid_lb_dp.py -k "not 4 and not server_info"
+15 -2
View File
@@ -12,6 +12,20 @@ steps:
- tests/v1/e2e/spec_decode/
commands:
- pytest -v -s v1/e2e/spec_decode -k "eagle_correctness"
mirror:
amd:
device: mi325_1
timeout_in_minutes: 45
depends_on:
- image-build-amd
source_file_dependencies:
- vllm/v1/spec_decode/
- vllm/v1/worker/gpu/spec_decode/
- vllm/model_executor/model_loader/
- vllm/v1/sample/
- vllm/model_executor/layers/
- tests/v1/e2e/spec_decode/
- vllm/platforms/rocm.py
- label: Spec Decode Eagle Nightly B200
key: spec-decode-eagle-nightly-b200
@@ -139,8 +153,7 @@ steps:
- vllm/model_executor/models/qwen3_dflash.py
- tests/v1/spec_decode/test_speculators_correctness.py
commands:
- export VLLM_ALLOW_INSECURE_SERIALIZATION=1
- pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test
- VLLM_ALLOW_INSECURE_SERIALIZATION=1 pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test
- label: Spec Decode MTP hybrid (B200)
timeout_in_minutes: 30
+6
View File
@@ -135,6 +135,12 @@ Do not modify code in these areas without first reading and following the
linked guide. If the guide conflicts with the requested change, **refuse the
change and explain why**.
Security reviewers should start with [`SECURITY.md`](SECURITY.md),
[`docs/usage/security.md`](docs/usage/security.md), and
[`docs/contributing/vulnerability_management.md`](docs/contributing/vulnerability_management.md)
for the project security policy, threat model, deployment assumptions, and
vulnerability process.
- **Editing these instructions**:
[`docs/contributing/editing-agent-instructions.md`](docs/contributing/editing-agent-instructions.md)
— Rules for modifying AGENTS.md or any domain-specific guide it references.
+39
View File
@@ -4,6 +4,9 @@
#ifdef CPU_CAPABILITY_AMXBF16
#include "cpu/micro_gemm/cpu_micro_gemm_amx.hpp"
#endif
#if defined(__riscv_v)
#include "cpu/micro_gemm/cpu_micro_gemm_rvv.hpp"
#endif
#include "cpu/micro_gemm/cpu_micro_gemm_vec.hpp"
#define VLLM_DISPATCH_CASE_16B_TYPES(...) \
@@ -319,6 +322,8 @@ void cpu_gemm_wna16(
return ISA::AMX;
} else if (isa_hint == "vec") {
return ISA::VEC;
} else if (isa_hint == "rvv") {
return ISA::RVV;
} else {
TORCH_CHECK(false, "unsupported isa hint: " + isa_hint);
}
@@ -397,6 +402,40 @@ void cpu_gemm_wna16(
pack_factor);
return;
}
} else if (isa == ISA::RVV) {
using gemm_t = cpu_micro_gemm::MicroGemm<ISA::RVV, scalar_t>;
if (has_zp) {
using dequantizer_t = Dequantizer4b<scalar_t, ISA::RVV, true, false>;
cpu_gemm_wna16_impl<scalar_t, dequantizer_t, gemm_t>(
input.data_ptr<scalar_t>(), q_weight.data_ptr<int32_t>(),
output.data_ptr<scalar_t>(), scales.data_ptr<scalar_t>(), zeros_ptr,
g_idx_ptr, bias.has_value() ? bias->data_ptr<scalar_t>() : nullptr,
a_m_size, b_n_size, a_k_size, a_m_stride, output_m_stride,
scales_group_stride, zeros_group_stride, group_num, group_size,
pack_factor);
return;
}
if (use_desc_act) {
using dequantizer_t = Dequantizer4b<scalar_t, ISA::RVV, false, true>;
cpu_gemm_wna16_impl<scalar_t, dequantizer_t, gemm_t>(
input.data_ptr<scalar_t>(), q_weight.data_ptr<int32_t>(),
output.data_ptr<scalar_t>(), scales.data_ptr<scalar_t>(), zeros_ptr,
g_idx_ptr, bias.has_value() ? bias->data_ptr<scalar_t>() : nullptr,
a_m_size, b_n_size, a_k_size, a_m_stride, output_m_stride,
scales_group_stride, zeros_group_stride, group_num, group_size,
pack_factor);
return;
} else {
using dequantizer_t = Dequantizer4b<scalar_t, ISA::RVV, false, false>;
cpu_gemm_wna16_impl<scalar_t, dequantizer_t, gemm_t>(
input.data_ptr<scalar_t>(), q_weight.data_ptr<int32_t>(),
output.data_ptr<scalar_t>(), scales.data_ptr<scalar_t>(), zeros_ptr,
g_idx_ptr, bias.has_value() ? bias->data_ptr<scalar_t>() : nullptr,
a_m_size, b_n_size, a_k_size, a_m_stride, output_m_stride,
scales_group_stride, zeros_group_stride, group_num, group_size,
pack_factor);
return;
}
}
});
}
+228
View File
@@ -0,0 +1,228 @@
#ifndef CPU_MICRO_GEMM_RVV_HPP
#define CPU_MICRO_GEMM_RVV_HPP
#include "cpu/micro_gemm/cpu_micro_gemm_impl.hpp"
#if defined(__riscv_v)
namespace cpu_micro_gemm {
namespace {
constexpr int32_t RVV_MGEMM_N8 = 8;
constexpr int32_t RVV_MGEMM_B_GROUP_STRIDE = 16;
template <typename scalar_t>
FORCE_INLINE fixed_fp32x8_t load_row8_b_as_f32(const scalar_t* ptr);
template <>
FORCE_INLINE fixed_fp32x8_t load_row8_b_as_f32<float>(const float* ptr) {
return RVVI(__riscv_vle32_v_f32, LMUL_256)(ptr, RVV_MGEMM_N8);
}
template <>
FORCE_INLINE fixed_fp32x8_t
load_row8_b_as_f32<c10::Half>(const c10::Half* ptr) {
#if defined(__riscv_zvfh)
fixed_fp16x8_t vec = RVVI(__riscv_vle16_v_f16, LMUL_128)(
reinterpret_cast<const _Float16*>(ptr), RVV_MGEMM_N8);
return RVVI(__riscv_vfwcvt_f_f_v_f32, LMUL_256)(vec, RVV_MGEMM_N8);
#else
alignas(32) float values[RVV_MGEMM_N8];
for (int32_t i = 0; i < RVV_MGEMM_N8; ++i) {
values[i] = static_cast<float>(ptr[i]);
}
return RVVI(__riscv_vle32_v_f32, LMUL_256)(values, RVV_MGEMM_N8);
#endif
}
template <>
FORCE_INLINE fixed_fp32x8_t
load_row8_b_as_f32<c10::BFloat16>(const c10::BFloat16* ptr) {
#if defined(__riscv_zvfbfmin)
fixed_u16x8_t raw = RVVI(__riscv_vle16_v_u16, LMUL_128)(
reinterpret_cast<const uint16_t*>(ptr), RVV_MGEMM_N8);
fixed_bf16x8_t vec =
RVVI4(__riscv_vreinterpret_v_u16, LMUL_128, _bf16, LMUL_128)(raw);
return RVVI(__riscv_vfwcvtbf16_f_f_v_f32, LMUL_256)(vec, RVV_MGEMM_N8);
#else
fixed_u16x8_t raw = RVVI(__riscv_vle16_v_u16, LMUL_128)(
reinterpret_cast<const uint16_t*>(ptr), RVV_MGEMM_N8);
auto wide = RVVI(__riscv_vzext_vf2_u32, LMUL_256)(raw, RVV_MGEMM_N8);
auto shifted = RVVI(__riscv_vsll_vx_u32, LMUL_256)(wide, 16, RVV_MGEMM_N8);
return RVVI4(__riscv_vreinterpret_v_u32, LMUL_256, _f32, LMUL_256)(shifted);
#endif
}
// Mx8 RVV kernel. B points at one 8-channel half of a 16-channel packed group,
// with rows separated by RVV_MGEMM_B_GROUP_STRIDE scalar elements.
template <int32_t M, typename scalar_t>
FORCE_INLINE void gemm_micro_rvv_fma_mx8_ku4(const scalar_t* __restrict__ a_ptr,
const scalar_t* __restrict__ b_ptr,
float* __restrict__ c_ptr,
const int64_t lda,
const int64_t ldc, const int32_t k,
const bool accum_c) {
static_assert(0 < M && M <= 8);
#define RVV_ROWS_APPLY(OP) OP(0) OP(1) OP(2) OP(3) OP(4) OP(5) OP(6) OP(7)
#define RVV_IF_M(i) if constexpr (M > (i))
#define RVV_DECL_A(i) const scalar_t* __restrict__ a##i = a_ptr + (i) * lda;
RVV_ROWS_APPLY(RVV_DECL_A)
#undef RVV_DECL_A
#define RVV_DECL_ACC(i) fixed_fp32x8_t acc##i;
RVV_ROWS_APPLY(RVV_DECL_ACC)
#undef RVV_DECL_ACC
#define RVV_INIT_ACC(i) \
RVV_IF_M(i) { \
if (accum_c) { \
acc##i = RVVI(__riscv_vle32_v_f32, LMUL_256)(c_ptr + (i) * ldc, \
RVV_MGEMM_N8); \
} else { \
acc##i = RVVI(__riscv_vfmv_v_f_f32, LMUL_256)(0.0f, RVV_MGEMM_N8); \
} \
}
RVV_ROWS_APPLY(RVV_INIT_ACC)
#undef RVV_INIT_ACC
int32_t k_idx = 0;
for (; k_idx + 3 < k; k_idx += 4) {
#define RVV_FMA_ROW(i, K_OFFSET) \
RVV_IF_M(i) { \
acc##i = RVVI(__riscv_vfmacc_vf_f32, LMUL_256)( \
acc##i, static_cast<float>(*(a##i + k_idx + (K_OFFSET))), b, \
RVV_MGEMM_N8); \
}
#define RVV_STEP_K(K_OFFSET) \
{ \
fixed_fp32x8_t b = load_row8_b_as_f32<scalar_t>( \
b_ptr + (k_idx + (K_OFFSET)) * RVV_MGEMM_B_GROUP_STRIDE); \
RVV_FMA_ROW(0, K_OFFSET) \
RVV_FMA_ROW(1, K_OFFSET) \
RVV_FMA_ROW(2, K_OFFSET) \
RVV_FMA_ROW(3, K_OFFSET) \
RVV_FMA_ROW(4, K_OFFSET) \
RVV_FMA_ROW(5, K_OFFSET) \
RVV_FMA_ROW(6, K_OFFSET) \
RVV_FMA_ROW(7, K_OFFSET) \
}
RVV_STEP_K(0)
RVV_STEP_K(1)
RVV_STEP_K(2)
RVV_STEP_K(3)
#undef RVV_STEP_K
#undef RVV_FMA_ROW
}
for (; k_idx < k; ++k_idx) {
fixed_fp32x8_t b =
load_row8_b_as_f32<scalar_t>(b_ptr + k_idx * RVV_MGEMM_B_GROUP_STRIDE);
#define RVV_TAIL_ROW(i) \
RVV_IF_M(i) { \
acc##i = RVVI(__riscv_vfmacc_vf_f32, LMUL_256)( \
acc##i, static_cast<float>(*(a##i + k_idx)), b, RVV_MGEMM_N8); \
}
RVV_ROWS_APPLY(RVV_TAIL_ROW)
#undef RVV_TAIL_ROW
}
#define RVV_STORE_ROW(i) \
RVV_IF_M(i) { \
RVVI(__riscv_vse32_v_f32, LMUL_256)(c_ptr + (i) * ldc, acc##i, \
RVV_MGEMM_N8); \
}
RVV_ROWS_APPLY(RVV_STORE_ROW)
#undef RVV_STORE_ROW
#undef RVV_ROWS_APPLY
#undef RVV_IF_M
}
template <int32_t M, typename scalar_t>
FORCE_INLINE void gemm_micro_rvv_mx32_ku4(DEFINE_CPU_MICRO_GEMM_PARAMS) {
static_assert(0 < M && M <= 8);
scalar_t* __restrict__ curr_b_0 = b_ptr;
scalar_t* __restrict__ curr_b_1 = b_ptr + b_n_group_stride;
gemm_micro_rvv_fma_mx8_ku4<M>(a_ptr, curr_b_0, c_ptr, lda, ldc, k, accum_c);
gemm_micro_rvv_fma_mx8_ku4<M>(a_ptr, curr_b_0 + RVV_MGEMM_N8,
c_ptr + RVV_MGEMM_N8, lda, ldc, k, accum_c);
gemm_micro_rvv_fma_mx8_ku4<M>(a_ptr, curr_b_1, c_ptr + 16, lda, ldc, k,
accum_c);
gemm_micro_rvv_fma_mx8_ku4<M>(a_ptr, curr_b_1 + RVV_MGEMM_N8, c_ptr + 24, lda,
ldc, k, accum_c);
}
class TileGemmRVV {
public:
template <typename scalar_t>
FORCE_INLINE static void gemm(DEFINE_CPU_MICRO_GEMM_PARAMS) {
switch (m) {
case 1:
gemm_micro_rvv_mx32_ku4<1>(CPU_MICRO_GEMM_PARAMS);
break;
case 2:
gemm_micro_rvv_mx32_ku4<2>(CPU_MICRO_GEMM_PARAMS);
break;
case 3:
gemm_micro_rvv_mx32_ku4<3>(CPU_MICRO_GEMM_PARAMS);
break;
case 4:
gemm_micro_rvv_mx32_ku4<4>(CPU_MICRO_GEMM_PARAMS);
break;
case 5:
gemm_micro_rvv_mx32_ku4<5>(CPU_MICRO_GEMM_PARAMS);
break;
case 6:
gemm_micro_rvv_mx32_ku4<6>(CPU_MICRO_GEMM_PARAMS);
break;
case 7:
gemm_micro_rvv_mx32_ku4<7>(CPU_MICRO_GEMM_PARAMS);
break;
case 8:
gemm_micro_rvv_mx32_ku4<8>(CPU_MICRO_GEMM_PARAMS);
break;
}
}
};
} // namespace
template <typename scalar_t>
class MicroGemm<cpu_utils::ISA::RVV, scalar_t> {
public:
static constexpr int32_t MaxMSize = 8;
static constexpr int32_t NSize = 32;
public:
void gemm(DEFINE_CPU_MICRO_GEMM_PARAMS) {
TileGemmRVV::gemm<scalar_t>(CPU_MICRO_GEMM_PARAMS);
}
static void pack_weight(const scalar_t* __restrict__ weight,
scalar_t* __restrict__ packed_weight,
const int32_t output_size, const int32_t input_size) {
TORCH_CHECK_EQ(output_size % 16, 0);
for (int32_t o_idx = 0; o_idx < output_size; ++o_idx) {
const scalar_t* __restrict__ curr_weight = weight + o_idx * input_size;
scalar_t* __restrict__ curr_packed_weight =
packed_weight + (o_idx / 16) * (16 * input_size) + o_idx % 16;
for (int32_t i_idx = 0; i_idx < input_size; ++i_idx) {
*curr_packed_weight = *curr_weight;
curr_packed_weight += 16;
++curr_weight;
}
}
}
};
} // namespace cpu_micro_gemm
#endif // defined(__riscv_v)
#endif // CPU_MICRO_GEMM_RVV_HPP
+3 -1
View File
@@ -8,13 +8,15 @@
#include "cpu/cpu_types.hpp"
namespace cpu_utils {
enum class ISA { AMX, VEC };
enum class ISA { AMX, VEC, RVV };
inline ISA get_isa(const std::string& isa) {
if (isa == "amx") {
return ISA::AMX;
} else if (isa == "vec") {
return ISA::VEC;
} else if (isa == "rvv") {
return ISA::RVV;
} else {
TORCH_CHECK(false, "Invalid isa type: " + isa);
}
@@ -136,8 +136,12 @@ typename T::Fmha::Arguments args_from_options(
StrideQ stride_Q_pe = cute::make_tuple(
static_cast<int64_t>(q_pe.stride(1)), _1{}, static_cast<int64_t>(q_pe.stride(0)));
// Read the token and page strides from the cache tensor instead of assuming
// packed pages, so strided views (e.g. per-layer views into a cross-layer
// block-major cache) are addressed correctly.
StrideK stride_C = cute::make_tuple(
static_cast<int64_t>(0 + D_latent + D_rope), _1{}, static_cast<int64_t>(page_size * (D_latent + D_rope)));
static_cast<int64_t>(kv_c_and_k_pe_cache.stride(1)), _1{},
static_cast<int64_t>(kv_c_and_k_pe_cache.stride(0)));
StrideLSE stride_PT = cute::make_stride(_1{}, page_count_per_seq);
StrideLSE stride_LSE = cute::make_tuple(_1{}, 0 + H);
StrideO stride_O = cute::make_tuple(static_cast<int64_t>(0 + D_latent), _1{}, static_cast<int64_t>(0 + H * D_latent));
+6 -7
View File
@@ -549,7 +549,7 @@ __global__ void indexer_k_quant_and_cache_kernel(
const int head_dim, // dimension of each head
const int quant_block_size, // quantization block size
const int cache_block_size, // cache block size
const int cache_stride, // stride for each token in kv_cache
const int64_t cache_block_stride, // stride for each block in kv_cache
const bool use_ue8m0 // use ue8m0 scale format
) {
@@ -590,16 +590,15 @@ __global__ void indexer_k_quant_and_cache_kernel(
scale = exp2f(ceilf(log2f(scale)));
}
const int64_t dst_offset = block_idx * cache_block_size * cache_stride +
block_offset * head_dim + head_dim_idx;
const int64_t dst_offset =
block_idx * cache_block_stride + block_offset * head_dim + head_dim_idx;
for (int i = 0; i < VEC_SIZE; i++) {
kv_cache[dst_offset + i] =
fp8::scaled_convert<cache_t, scalar_t, kv_dt>(k_val_ptr[i], scale);
}
if (threadIdx.x == 0) {
const int64_t dst_scale_idx =
block_idx * cache_block_size * cache_stride +
cache_block_size * head_dim +
block_idx * cache_block_stride + cache_block_size * head_dim +
(block_offset * head_dim + head_dim_idx) * 4 / quant_block_size;
reinterpret_cast<float*>(kv_cache)[dst_scale_idx / 4] = scale;
}
@@ -1452,7 +1451,7 @@ void cp_gather_and_upconvert_fp8_kv_cache(
reinterpret_cast<KV_T*>(k.data_ptr()), \
reinterpret_cast<CACHE_T*>(kv_cache.data_ptr()), \
slot_mapping.const_data_ptr<int64_t>(), head_dim, quant_block_size, \
cache_block_size, cache_stride, use_ue8m0);
cache_block_size, cache_block_stride, use_ue8m0);
void indexer_k_quant_and_cache(
torch::stable::Tensor& k, // [num_tokens, head_dim]
@@ -1463,7 +1462,7 @@ void indexer_k_quant_and_cache(
int num_tokens = k.size(0);
int head_dim = k.size(1);
int cache_block_size = kv_cache.size(1);
int cache_stride = kv_cache.size(2);
int64_t cache_block_stride = kv_cache.stride(0);
bool use_ue8m0 = scale_fmt == "ue8m0";
STD_TORCH_CHECK(k.device() == kv_cache.device(),
+10 -10
View File
@@ -132,7 +132,7 @@ CMD ["/bin/bash"]
FROM vllm-base AS ucx-nixl-build
ARG UCX_VERSION=v1.21.0-rc2
ARG NIXL_VERSION=0.10.1
ARG NIXL_VERSION=v1.2.0
# Build-time only: compiler, autotools, and verbs dev headers
RUN apt-get update -y && apt-get install -y --no-install-recommends \
@@ -149,25 +149,25 @@ RUN apt-get update -y && apt-get install -y --no-install-recommends \
# patchelf (installed via uv) is used by the NIXL wheel build to rewrite
# RPATH entries, making the wheel portable across stages.
RUN --mount=type=cache,target=/root/.cache/uv \
git clone https://github.com/openucx/ucx /tmp/ucx_source && \
cd /tmp/ucx_source && git checkout "${UCX_VERSION}" && \
git clone --depth 1 --branch "${UCX_VERSION}" https://github.com/openucx/ucx /tmp/ucx_source && \
cd /tmp/ucx_source && \
bash autogen.sh && \
./configure --prefix=/tmp/ucx_install --with-ze=yes --enable-examples --enable-mt && \
make CFLAGS="-Wno-error=incompatible-pointer-types" -j8 && make install && \
git clone https://github.com/ai-dynamo/nixl /tmp/nixl_source && \
cd /tmp/nixl_source && git checkout "${NIXL_VERSION}" && \
make CFLAGS="-Wno-error=incompatible-pointer-types" -j"$(nproc)" && make install && \
git clone --depth 1 --branch "${NIXL_VERSION}" https://github.com/ai-dynamo/nixl /tmp/nixl_source && \
cd /tmp/nixl_source && \
uv pip install --upgrade meson pybind11 patchelf && \
uv pip install -r requirements.txt && \
PKG_CONFIG_PATH=/tmp/ucx_install/lib/pkgconfig \
LD_LIBRARY_PATH=/tmp/ucx_install/lib \
python -m pip wheel --no-deps . -w /tmp/nixl_wheels/ && \
find /tmp/ucx_install -type f \( -name '*.a' -o -name '*.la' \) -delete && \
rm -rf /tmp/ucx_install/include /tmp/ucx_install/share /tmp/ucx_install/etc /tmp/ucx_install/lib/cmake /tmp/ucx_install/bin && \
rm -rf /tmp/ucx_source /tmp/nixl_source
rm -rf /tmp/ucx_install/{include,share,etc,bin} /tmp/ucx_install/lib/cmake \
/tmp/ucx_source /tmp/nixl_source
FROM vllm-base AS vllm-openai
ARG NIXL_VERSION=0.10.1
ARG NIXL_VERSION=v1.2.0
# Copy compiled UCX runtime libraries and the pre-built NIXL wheel.
# No compiler or autotools are installed in this stage.
@@ -192,7 +192,7 @@ RUN --mount=type=cache,target=/root/.cache/uv \
librdmacm1t64 \
&& rm -rf /var/lib/apt/lists/* \
&& uv pip install --no-deps /tmp/nixl_wheels/nixl*.whl \
&& uv pip install nixl==${NIXL_VERSION} \
&& uv pip install nixl==${NIXL_VERSION} && uv pip uninstall nixl-cu13 \
&& rm -rf /tmp/nixl_wheels
RUN --mount=type=cache,target=/root/.cache/uv \
+8 -2
View File
@@ -304,9 +304,15 @@ review process:
resources. The reviewer will add `ready` label to the PR when the PR is
ready to merge or a full CI run is needed.
### Escalating Stalled Contributions
### Pull Request Limits and Escalation
If you have an important contribution that has not yet received maintainer attention, please email us at:
vLLM uses GitHub's [pull request limit](https://github.blog/open-source/maintainers/how-pull-request-limits-are-cutting-down-the-noise/)
for contributors without write access. The current cap is 6 open PRs. If this
blocks well-intentioned critical work, contact a committer to request bypass
list access.
If you need an expedited review for an important contribution, please email us
at:
<pr-review-request@vllm.ai>
+2 -2
View File
@@ -136,7 +136,7 @@ The model should also be added to the `MODELS_CONFIG_MAP` dictionary in [vllm/mo
For case (2), we recommend using as a reference the implementation of [`JambaForCausalLM`](../../../vllm/model_executor/models/jamba.py) (for an example of a model that uses Mamba-1 and attention together) or [`NemotronHForCausalLM`](../../../vllm/model_executor/models/nemotron_h.py) (for an example of a model that uses Mamba-2 and attention together).
These models should follow the same instructions as case (1), but they should inherit protocol `IsHybrid` (instead of `IsAttentionFree`) and it is *not* necessary to add them to the `MODELS_CONFIG_MAP` (their runtime defaults will be inferred from the protocol).
For case (3), we recommend looking at the implementation of [`MiniMaxText01ForCausalLM`](../../../vllm/model_executor/models/minimax_text_01.py) or [`Lfm2ForCausalLM`](../../../vllm/model_executor/models/lfm2.py) as a reference, which use custom "mamba-like" layers `MiniMaxText01LinearAttention` and `ShortConv` respectively.
For case (3), we recommend looking at the implementation of [`Lfm2ForCausalLM`](../../../vllm/model_executor/models/lfm2.py) as a reference, which uses a custom "mamba-like" layer `ShortConv`.
Please follow the same guidelines as case (2) for implementing these models.
We use "mamba-like" to refer to layers that possess a state that is updated in-place, rather than being appended-to (like KV cache for attention).
For implementing new custom mamba-like layers, one should inherit from `MambaBase` and implement the methods `get_state_dtype`, `get_state_shape` to calculate the data types and state shapes at runtime, as well as `mamba_type` and `get_attn_backend`.
@@ -144,5 +144,5 @@ It is also necessary to implement the "attention meta-data" class which handles
Please see [`LinearAttentionMetadata`](../../../vllm/v1/attention/backends/linear_attn.py) or [`ShortConvAttentionMetadata`](../../../vllm/v1/attention/backends/short_conv_attn.py) for examples of this.
It is also worth noting that we should update `MambaAttentionBackendEnum` in [`registry.py`](../../../vllm/v1/attention/backends/registry.py) when adding a new mamba backend.
Finally, if one wants to support torch compile and CUDA graphs, it necessary to wrap the call to the mamba-like layer inside a custom op and register it.
Please see the calls to `direct_register_custom_op` in [vllm/model_executor/models/minimax_text_01.py](../../../vllm/model_executor/models/minimax_text_01.py) or [vllm/model_executor/layers/mamba/short_conv.py](../../../vllm/model_executor/layers/mamba/short_conv.py) for examples of this.
Please see the calls to `direct_register_custom_op` in [vllm/model_executor/layers/mamba/linear/minimax_linear_attn.py](../../../vllm/model_executor/layers/mamba/linear/minimax_linear_attn.py) or [vllm/model_executor/layers/mamba/short_conv.py](../../../vllm/model_executor/layers/mamba/short_conv.py) for examples of this.
The new custom op should then be added to the list `_attention_ops` in [vllm/config/compilation.py](../../../vllm/config/compilation.py) to ensure that piecewise CUDA graphs works as intended.
+2 -2
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, 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 |
| `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 |
+1
View File
@@ -74,6 +74,7 @@ vllm serve <model> \
| `max_tracker_size` | no | `64000` | single-tier | Max entries in the lookup tracker. |
| `secondary_tiers` | no | `[]` | multi-tier | List of secondary tier configs (see below). |
| `offload_prompt_only` | no | `true` | both | If `true`, only prompt (prefill) blocks are offloaded; decode blocks are skipped. |
| `self_describing_kv_events` | no | `false` | single-tier | Opt-in. When `true` *and* KV cache events are enabled (`--kv-events-config` with `enable_kv_cache_events`), the connector emits self-describing block-granular `BlockStored`/`BlockRemoved` payloads (constituent block hashes, whole-chunk `token_ids`, per-block `block_size`, parent hash, LoRA + group/cache-spec metadata) instead of the placeholder fallback, so external KV-event consumers can index offloaded blocks. Inert unless events are enabled. Currently rejected by `TieringOffloadingSpec`. Full-attention groups only; sliding-window/SSM groups keep the placeholder fallback. In chunk mode (`block_size` > GPU block size), overlapping chunks re-announce shared per-block hashes, so consumers must reference-count (deduplicate) repeated store/remove announcements. |
| `spec_module_path` | no | — | both | Python import path for a custom `OffloadingSpec` not in the built-in registry. Required only when `spec_name` is not built-in (advanced). |
## Secondary Tiers
-9
View File
@@ -321,15 +321,6 @@ For Qwen2.5, the chat template in tokenizer_config.json has already included sup
Flags: `--tool-call-parser hermes`
### MiniMax Models (`minimax_m1`)
Supported models:
* `MiniMaxAi/MiniMax-M1-40k` (use with [examples/tool_chat_template_minimax_m1.jinja](../../examples/tool_chat_template_minimax_m1.jinja))
* `MiniMaxAi/MiniMax-M1-80k` (use with [examples/tool_chat_template_minimax_m1.jinja](../../examples/tool_chat_template_minimax_m1.jinja))
Flags: `--tool-call-parser minimax --chat-template examples/tool_chat_template_minimax_m1.jinja`
### DeepSeek-V3 Models (`deepseek_v3`)
Supported models:
+1 -1
View File
@@ -61,7 +61,7 @@ Models of any architecture can be converted into embedding models using `--conve
| `ColModernVBertForRetrieval` | ColModernVBERT | T / I | `ModernVBERT/colmodernvbert-merged` | | |
| `ColPaliForRetrieval` | ColPali | T / I | `vidore/colpali-v1.3-hf` | | |
| `ColQwen3` | Qwen3-VL | T / I | `TomoroAI/tomoro-colqwen3-embed-4b`, `TomoroAI/tomoro-colqwen3-embed-8b` | | |
| `ColQwen3_5` | ColQwen3.5 | T + I + V | `athrael-soju/colqwen3.5-4.5B-v3` | | |
| `ColQwen3_5` | ColQwen3.5 | T + I + V | `athrael-soju/colqwen3.5-4.5B-v3`, `vultr/VultronRetrieverPrime-Qwen3.5-8B` | | |
| `OpsColQwen3Model` | Qwen3-VL | T / I | `OpenSearch-AI/Ops-Colqwen3-4B`, `OpenSearch-AI/Ops-Colqwen3-8B` | | |
| `Qwen3VLNemotronEmbedModel` | Qwen3-VL | T / I | `nvidia/nemotron-colembed-vl-4b-v2`, `nvidia/nemotron-colembed-vl-8b-v2` | ✅︎ | ✅︎ |
| `*ForConditionalGeneration`<sup>C</sup>, `*ForCausalLM`<sup>C</sup>, etc. | Generative models | \* | N/A | \* | \* |
+1 -3
View File
@@ -441,7 +441,6 @@ th {
| `MiMoV2ForCausalLM` | MiMoV2Pro | `XiaomiMiMo/MiMo-V2.5-Pro`, etc. | | ✅︎ |
| `MiniCPMForCausalLM` | MiniCPM | `openbmb/MiniCPM-2B-sft-bf16`, `openbmb/MiniCPM-2B-dpo-bf16`, `openbmb/MiniCPM-S-1B-sft`, etc. | ✅︎ | ✅︎ |
| `MiniCPM3ForCausalLM` | MiniCPM3 | `openbmb/MiniCPM3-4B`, etc. | ✅︎ | ✅︎ |
| `MiniMaxForCausalLM` | MiniMax-Text | `MiniMaxAI/MiniMax-Text-01-hf`, etc. | | |
| `MiniMaxM2ForCausalLM` | MiniMax-M2, MiniMax-M2.1 | `MiniMaxAI/MiniMax-M2`, etc. | ✅︎ | ✅︎ |
| `MistralForCausalLM` | Ministral-3, Mistral, Mistral-Instruct | `mistralai/Ministral-3-3B-Instruct-2512`, `mistralai/Mistral-7B-v0.1`, `mistralai/Mistral-7B-Instruct-v0.1`, etc. | ✅︎ | ✅︎ |
| `MistralLarge3ForCausalLM` | Mistral-Large-3-675B-Base-2512, Mistral-Large-3-675B-Instruct-2512 | `mistralai/Mistral-Large-3-675B-Base-2512`, `mistralai/Mistral-Large-3-675B-Instruct-2512`, etc. | ✅︎ | ✅︎ |
@@ -487,8 +486,6 @@ th {
| `TeleChat2ForCausalLM` | TeleChat2 | `Tele-AI/TeleChat2-3B`, `Tele-AI/TeleChat2-7B`, `Tele-AI/TeleChat2-35B`, etc. | ✅︎ | ✅︎ |
| `TeleChat3ForCausalLM` | TeleChat3 | `Tele-AI/TeleChat3-36B-Thinking`, `Tele-AI/TeleChat3-Coder-36B-Thinking`, etc. | ✅︎ | ✅︎ |
| `TeleFLMForCausalLM` | TeleFLM | `CofeAI/FLM-2-52B-Instruct-2407`, `CofeAI/Tele-FLM`, etc. | ✅︎ | ✅︎ |
| `MiniMaxM1ForCausalLM` | MiniMax-Text | `MiniMaxAI/MiniMax-M1-40k`, `MiniMaxAI/MiniMax-M1-80k`, etc. | | |
| `MiniMaxText01ForCausalLM` | MiniMax-Text | `MiniMaxAI/MiniMax-Text-01`, etc. | | |
| `Zamba2ForCausalLM` | Zamba2 | `Zyphra/Zamba2-7B-instruct`, `Zyphra/Zamba2-2.7B-instruct`, `Zyphra/Zamba2-1.2B-instruct`, etc. | | |
!!! note
@@ -595,6 +592,7 @@ These models primarily accept the [`LLM.generate`](./generative_models.md#llmgen
| `MiMoV2OmniForCausalLM` | MiMo-V2.5-Omni | T + I<sup>E+</sup> + V<sup>E+</sup> + A<sup>+</sup> | `XiaomiMiMo/MiMo-V2.5-Omni` | | ✅︎ |
| `MiniCPMO` | MiniCPM-O | T + I<sup>E+</sup> + V<sup>E+</sup> + A<sup>E+</sup> | `openbmb/MiniCPM-o-2_6`, etc. | ✅︎ | ✅︎ |
| `MiniCPMV` | MiniCPM-V | T + I<sup>E+</sup> + V<sup>E+</sup> | `openbmb/MiniCPM-V-2` (see note), `openbmb/MiniCPM-Llama3-V-2_5`, `openbmb/MiniCPM-V-2_6`, `openbmb/MiniCPM-V-4`, `openbmb/MiniCPM-V-4_5`, etc. | ✅︎ | |
| `MiniMaxM3SparseForConditionalGeneration` | MiniMax-M3 | T + I<sup>+</sup> + V<sup>+</sup> | `MiniMaxAI/MiniMax-M3`, `MiniMaxAI/MiniMax-M3-MXFP8`, etc. | | |
| `MiniMaxVL01ForConditionalGeneration` | MiniMax-VL | T + I<sup>E+</sup> | `MiniMaxAI/MiniMax-VL-01`, etc. | | ✅︎ |
| `Mistral3ForConditionalGeneration` | Mistral3 (HF Transformers) | T + I<sup>+</sup> | `mistralai/Mistral-Small-3.1-24B-Instruct-2503`, etc. | ✅︎ | ✅︎ |
| `MolmoForCausalLM` | Molmo | T + I<sup>+</sup> | `allenai/Molmo-7B-D-0924`, `allenai/Molmo-7B-O-0924`, etc. | ✅︎ | ✅︎ |
+1 -1
View File
@@ -128,7 +128,7 @@ Models that use Mamba-2 and Mamba-1 layers (e.g., `Mamba2ForCausalLM`, `MambaFor
Hybrid models that combine Mamba-2 and Mamba-1 layers with standard attention layers are also supported (e.g., `BambaForCausalLM`,
`Zamba2ForCausalLM`, `NemotronHForCausalLM`, `FalconH1ForCausalLM` and `GraniteMoeHybridForCausalLM`, `JambaForCausalLM`, `Plamo2ForCausalLM`).
Hybrid models with mechanisms different to Mamba are also supported (e.g, `MiniMaxText01ForCausalLM`, `MiniMaxM1ForCausalLM`, `Lfm2ForCausalLM`).
Hybrid models with mechanisms different to Mamba are also supported (e.g, `Lfm2ForCausalLM`).
Please note that prefix caching is not yet supported for any of the above models.
@@ -1481,39 +1481,6 @@ def run_minicpmv(questions: list[str], modality: str) -> ModelRequestData:
return run_minicpmv_base(questions, modality, "openbmb/MiniCPM-V-2_6")
def run_minimax_vl_01(questions: list[str], modality: str) -> ModelRequestData:
assert modality == "image"
model_name = "MiniMaxAI/MiniMax-VL-01"
engine_args = EngineArgs(
model=model_name,
max_num_seqs=2,
limit_mm_per_prompt={modality: 1},
trust_remote_code=True,
tensor_parallel_size=8,
)
tokenizer = AutoTokenizer.from_pretrained(model_name)
messages = [
[
{
"role": "user",
"content": [{"type": "image"}, {"type": "text", "text": question}],
}
]
for question in questions
]
prompts = tokenizer.apply_chat_template(
messages, add_generation_prompt=True, tokenize=False
)
return ModelRequestData(
engine_args=engine_args,
prompts=prompts,
)
# Mistral-3 HF-format
def run_mistral3(questions: list[str], modality: str) -> ModelRequestData:
assert modality == "image"
@@ -2485,7 +2452,6 @@ model_example_map = {
"mantis": run_mantis,
"minicpmo": run_minicpmo,
"minicpmv": run_minicpmv,
"minimax_vl_01": run_minimax_vl_01,
"mistral3": run_mistral3,
"molmo": run_molmo,
"molmo2": run_molmo2,
@@ -7,11 +7,27 @@ ColQwen3.5 is a multi-modal ColBERT-style model based on Qwen3.5.
It produces per-token embeddings and uses MaxSim scoring for retrieval
and reranking. Supports both text and image inputs.
Works for any ColQwen3.5 checkpoint, e.g. `athrael-soju/colqwen3.5-4.5B-v3`
or `vultr/VultronRetrieverPrime-Qwen3.5-8B`.
Start the server with:
vllm serve athrael-soju/colqwen3.5-4.5B --max-model-len 4096
vllm serve athrael-soju/colqwen3.5-4.5B-v3 --max-model-len 4096 \
--mm-processor-kwargs '{"min_pixels": 65536, "max_pixels": 1835008}'
Then run this script:
python colqwen3_5_rerank_online.py
Parity note (matching the native colpali ColQwen3_5Processor pipeline):
- Visual-token budget: ColQwen3_5Processor uses max_num_visual_tokens=1792,
i.e. max_pixels = 1792 * (patch_size*merge_size)^2 = 1792 * 32^2 = 1835008
(with min_pixels = shortest_edge = 65536). Pass these via --mm-processor-kwargs
as above; the default budget gives fewer visual tokens and lower retrieval ndcg.
- When you build prompts yourself (token_embed), reproduce the processor exactly:
image (document): wrap in the instruction template
"<|im_start|>user\n<|vision_start|><|image_pad|><|vision_end|>"
"Describe the image.<|im_end|><|endoftext|>"
query: append the augmentation suffix <text> + "<|endoftext|>" * 10
Omitting these reproduces a silent ~2.5 ndcg@10 drop vs the native pipeline.
"""
import requests
@@ -1,91 +0,0 @@
{{ '<begin_of_document>' -}}
{%- if custom_tools is defined %}
{%- set tools = custom_tools %}
{%- endif %}
{%- if not tools is defined %}
{%- set tools = none %}
{%- endif %}
{#- Extract system message #}
{% set ns = namespace(system_prompt='') -%}
{%- if messages[0]['role'] == 'system' %}
{%- if messages[0]['content'] is string %}
{%- set ns.system_prompt = messages[0]['content']|trim %}
{%- else %}
{%- set ns.system_prompt = messages[0]['content'][0]['text']|trim %}
{%- endif %}
{%- set messages = messages[1:] %}
{%- else %}
{%- if tools is not none %}
{%- set ns.system_prompt = "You are a helpful assistant created by Minimax based on MiniMax-M1 model." %}
{%- else %}
{%- set ns.system_prompt = "You are a helpful assistant created by Minimax based on MiniMax-M1 model." %}
{%- endif %}
{%- endif %}
{#- System message #}
{%- if ns.system_prompt != '' %}
{{ '<beginning_of_sentence>system ai_setting=assistant\n' + ns.system_prompt + '<end_of_sentence>\n' -}}
{%- endif %}
{#- Tools configuration #}
{%- if tools is not none %}
{{ '<beginning_of_sentence>system tool_setting=tools\nYou are provided with these tools:\n<tools>\n' -}}
{%- for tool in tools %}
{{ tool | tojson ~ '\n' -}}
{%- endfor %}
{{ '</tools>\n\nIf you need to call tools, please respond with <tool_calls></tool_calls> XML tags, and provide tool-name and json-object of arguments, following the format below:\n<tool_calls>\n{"name": <tool-name>, "arguments": <args-json-object>}\n...\n</tool_calls><end_of_sentence>\n' -}}
{%- endif %}
{#- Process messages #}
{%- for message in messages %}
{%- if not (message.role == 'ipython' or message.role == 'tool' or 'tool_calls' in message) %}
{%- if message['role'] == 'user' %}
{{ '<beginning_of_sentence>user name=user\n' -}}
{%- if message['content'] is string %}
{{ message['content']|trim -}}
{%- else %}
{%- for content in message['content'] %}
{%- if content['type'] == 'text' %}
{{ content['text']|trim -}}
{%- endif %}
{%- endfor %}
{%- endif %}
{{ '<end_of_sentence>\n' -}}
{%- elif message['role'] == 'assistant' %}
{{ '<beginning_of_sentence>ai name=assistant\n' -}}
{%- if message['content'] is string %}
{{ message['content']|trim -}}
{%- else %}
{%- for content in message['content'] | selectattr('type', 'equalto', 'text') %}
{{ content['text']|trim -}}
{%- endfor %}
{%- endif %}
{{ '<end_of_sentence>\n' -}}
{%- endif %}
{%- elif 'tool_calls' in message %}
{{ '<beginning_of_sentence>ai name=assistant\n<tool_calls>\n' -}}
{%- for tool_call in message.tool_calls %}
{{ '{"name": "' + tool_call.function.name + '", "arguments": ' + tool_call.function.arguments | tojson + '}\n' -}}
{%- endfor %}
{{ '</tool_calls><end_of_sentence>\n' -}}
{%- elif message.role == "tool" or message.role == "ipython" %}
{{ '<beginning_of_sentence>tool name=tools\n' -}}
{%- if message.content is string %}
{{ 'tool result: ' + message.content + '\n\n' -}}
{%- else %}
{%- for content in message['content'] %}
{%- if content['type'] == 'text' %}
{{ 'tool result: ' + content['text'] + '\n\n' -}}
{%- elif content.get('name') %}
{{ 'tool name: ' + content['name'] + '\ntool result: ' + content['text'] + '\n\n' -}}
{%- endif %}
{%- endfor %}
{%- endif %}
{{ '<end_of_sentence>\n' -}}
{%- endif %}
{%- endfor %}
{%- if add_generation_prompt %}
{{ '<beginning_of_sentence>ai name=assistant\n' -}}
{%- endif %}
-1
View File
@@ -386,7 +386,6 @@ mod tests {
tool_chat_template_llama3.2_pythonic.jinja => String
tool_chat_template_llama4_json.jinja => OpenAi
tool_chat_template_llama4_pythonic.jinja => OpenAi
tool_chat_template_minimax_m1.jinja => OpenAi
tool_chat_template_mistral.jinja => String
tool_chat_template_mistral3.jinja => OpenAi
tool_chat_template_mistral_parallel.jinja => String
@@ -1,91 +0,0 @@
{{ '<begin_of_document>' -}}
{%- if custom_tools is defined %}
{%- set tools = custom_tools %}
{%- endif %}
{%- if not tools is defined %}
{%- set tools = none %}
{%- endif %}
{#- Extract system message #}
{% set ns = namespace(system_prompt='') -%}
{%- if messages[0]['role'] == 'system' %}
{%- if messages[0]['content'] is string %}
{%- set ns.system_prompt = messages[0]['content']|trim %}
{%- else %}
{%- set ns.system_prompt = messages[0]['content'][0]['text']|trim %}
{%- endif %}
{%- set messages = messages[1:] %}
{%- else %}
{%- if tools is not none %}
{%- set ns.system_prompt = "You are a helpful assistant created by Minimax based on MiniMax-M1 model." %}
{%- else %}
{%- set ns.system_prompt = "You are a helpful assistant created by Minimax based on MiniMax-M1 model." %}
{%- endif %}
{%- endif %}
{#- System message #}
{%- if ns.system_prompt != '' %}
{{ '<beginning_of_sentence>system ai_setting=assistant\n' + ns.system_prompt + '<end_of_sentence>\n' -}}
{%- endif %}
{#- Tools configuration #}
{%- if tools is not none %}
{{ '<beginning_of_sentence>system tool_setting=tools\nYou are provided with these tools:\n<tools>\n' -}}
{%- for tool in tools %}
{{ tool | tojson ~ '\n' -}}
{%- endfor %}
{{ '</tools>\n\nIf you need to call tools, please respond with <tool_calls></tool_calls> XML tags, and provide tool-name and json-object of arguments, following the format below:\n<tool_calls>\n{"name": <tool-name>, "arguments": <args-json-object>}\n...\n</tool_calls><end_of_sentence>\n' -}}
{%- endif %}
{#- Process messages #}
{%- for message in messages %}
{%- if not (message.role == 'ipython' or message.role == 'tool' or 'tool_calls' in message) %}
{%- if message['role'] == 'user' %}
{{ '<beginning_of_sentence>user name=user\n' -}}
{%- if message['content'] is string %}
{{ message['content']|trim -}}
{%- else %}
{%- for content in message['content'] %}
{%- if content['type'] == 'text' %}
{{ content['text']|trim -}}
{%- endif %}
{%- endfor %}
{%- endif %}
{{ '<end_of_sentence>\n' -}}
{%- elif message['role'] == 'assistant' %}
{{ '<beginning_of_sentence>ai name=assistant\n' -}}
{%- if message['content'] is string %}
{{ message['content']|trim -}}
{%- else %}
{%- for content in message['content'] | selectattr('type', 'equalto', 'text') %}
{{ content['text']|trim -}}
{%- endfor %}
{%- endif %}
{{ '<end_of_sentence>\n' -}}
{%- endif %}
{%- elif 'tool_calls' in message %}
{{ '<beginning_of_sentence>ai name=assistant\n<tool_calls>\n' -}}
{%- for tool_call in message.tool_calls %}
{{ '{"name": "' + tool_call.function.name + '", "arguments": ' + tool_call.function.arguments | tojson + '}\n' -}}
{%- endfor %}
{{ '</tool_calls><end_of_sentence>\n' -}}
{%- elif message.role == "tool" or message.role == "ipython" %}
{{ '<beginning_of_sentence>tool name=tools\n' -}}
{%- if message.content is string %}
{{ 'tool result: ' + message.content + '\n\n' -}}
{%- else %}
{%- for content in message['content'] %}
{%- if content['type'] == 'text' %}
{{ 'tool result: ' + content['text'] + '\n\n' -}}
{%- elif content.get('name') %}
{{ 'tool name: ' + content['name'] + '\ntool result: ' + content['text'] + '\n\n' -}}
{%- endif %}
{%- endfor %}
{%- endif %}
{{ '<end_of_sentence>\n' -}}
{%- endif %}
{%- endfor %}
{%- if add_generation_prompt %}
{{ '<beginning_of_sentence>ai name=assistant\n' -}}
{%- endif %}
@@ -277,6 +277,12 @@ pub struct EngineCoreSamplingParams {
pub max_tokens: u32,
/// Minimum number of tokens to generate before EOS or stop-token handling.
pub min_tokens: u32,
/// Maximum number of reasoning ("thinking") tokens to emit before the
/// reasoning section is force-closed. `None` means unlimited; the
/// user-facing `-1` sentinel is normalized to `None` by the frontend before
/// reaching this DTO, so only non-negative values are sent. Enforced
/// engine-side (and only when a reasoning parser is configured).
pub thinking_token_budget: Option<u64>,
/// Number of log probabilities to return per generated token.
///
/// `None` disables sample logprobs. `-1` requests the full vocabulary.
@@ -345,6 +351,7 @@ impl EngineCoreSamplingParams {
seed: None,
max_tokens: 65536,
min_tokens: 0,
thinking_token_budget: None,
logprobs: None,
prompt_logprobs: None,
min_p: 0.0,
@@ -150,6 +150,7 @@ fn sample_request_with_id(request_id: &str) -> EngineCoreRequest {
top_k: 8,
max_tokens: 32,
min_tokens: 1,
thinking_token_budget: Some(256),
stop_token_ids: vec![151643],
eos_token_id: Some(151645),
all_stop_token_ids: BTreeSet::from([151643, 151645]),
@@ -2502,6 +2503,7 @@ fn python_msgpack_fixtures_match_rust_encoding() {
seed: None,
max_tokens: 16,
min_tokens: 0,
thinking_token_budget: None,
logprobs: None,
prompt_logprobs: None,
min_p: 0.0,
@@ -39,6 +39,7 @@ class EngineCoreSamplingParams(msgspec.Struct, dict=True, omit_defaults=True):
seed: int | None = None
max_tokens: int = 16
min_tokens: int = 0
thinking_token_budget: int | None = None
min_p: float = 0.0
frequency_penalty: float = 0.0
presence_penalty: float = 0.0
@@ -122,6 +123,7 @@ request = EngineCoreRequest(
seed=None,
max_tokens=32,
min_tokens=1,
thinking_token_budget=256,
min_p=0.0,
frequency_penalty=0.0,
presence_penalty=0.0,
+13
View File
@@ -103,6 +103,7 @@ fn is_request_validation_error(error: &vllm_text::Error) -> bool {
| vllm_text::Error::EmptyPromptTokenIds { .. }
| vllm_text::Error::Logprobs(_)
| vllm_text::Error::OutOfVocab(_)
| vllm_text::Error::InvalidThinkingTokenBudget
// An empty tokenized prompt detected later, at request prepare
// time, surfaces through the transparent Llm wrapper.
| vllm_text::Error::Llm(vllm_llm::Error::EmptyPromptTokenIds { .. })
@@ -127,6 +128,18 @@ mod tests {
assert!(response.error.message.contains("9000"));
}
#[test]
fn invalid_thinking_token_budget_maps_to_invalid_request() {
let api_error = text_submit_error(
"failed to submit completion request",
vllm_text::Error::InvalidThinkingTokenBudget,
);
assert_eq!(api_error.status_code(), StatusCode::BAD_REQUEST);
let response = api_error.to_error_response();
assert_eq!(response.error.error_type, "invalid_request_error");
assert!(response.error.message.contains("thinking_token_budget"));
}
#[test]
fn chat_wrapped_prompt_too_long_maps_to_invalid_request() {
let error = vllm_chat::Error::Text(vllm_text::Error::PromptTooLong {
@@ -150,6 +150,33 @@ mod tests {
);
}
#[test]
fn prepare_generate_request_forwards_thinking_token_budget() {
let request: GenerateRequest = serde_json::from_value(json!({
"model": "Qwen/Qwen1.5-0.5B-Chat",
"token_ids": [11, 22, 33],
"sampling_params": {
"thinking_token_budget": 64
}
}))
.expect("parse request");
let prepared = prepare_generate_request(
request,
&served(&["Qwen/Qwen1.5-0.5B-Chat"]),
ResolvedRequestContext::default(),
)
.expect("prepare");
// The raw inference route shares `vllm_text::SamplingParams`, so the
// field is carried through to lowering exactly like the OpenAI routes
// (normalization/validation then happens in `lower_sampling_params`).
assert_eq!(
prepared.text_request.sampling_params.thinking_token_budget,
Some(64)
);
}
#[test]
fn prepare_generate_request_gates_continuous_usage_on_include_usage() {
let request: GenerateRequest = serde_json::from_value(json!({
@@ -115,6 +115,7 @@ pub(super) fn prepare_chat_request(
seed: request.seed,
max_tokens: request.max_completion_tokens,
min_tokens: request.min_tokens,
thinking_token_budget: request.thinking_token_budget,
logprobs: request.logprobs.then_some(top_logprobs),
prompt_logprobs,
min_p: request.min_p,
@@ -613,6 +614,31 @@ mod tests {
assert_eq!(prepared.chat_request.sampling_params, expected);
}
#[test]
fn prepare_chat_request_passes_through_thinking_token_budget() {
let prepare = |budget: Option<i64>| {
prepare_chat_request(
ChatCompletionRequest {
thinking_token_budget: budget,
..base_request()
},
&served(&["Qwen/Qwen1.5-0.5B-Chat"]),
ResolvedRequestContext::default(),
)
.expect("request is valid")
.chat_request
.sampling_params
.thinking_token_budget
};
// The convert layer forwards the raw value verbatim (including the `-1`
// "unlimited" sentinel); normalization/validation happens during
// lowering (see `vllm_text::lower`).
assert_eq!(prepare(Some(64)), Some(64));
assert_eq!(prepare(Some(-1)), Some(-1));
assert_eq!(prepare(None), None);
}
#[test]
fn prepare_chat_request_accepts_developer_messages() {
let request = ChatCompletionRequest {
@@ -165,8 +165,10 @@ pub struct ChatCompletionRequest {
pub bad_words: Option<Vec<String>>,
// -------- Extra vLLM Parameters --------
/// Token budget for reasoning/thinking
pub thinking_token_budget: Option<u32>,
/// Token budget for reasoning/thinking. Accepts a non-negative integer, or
/// `-1` for unlimited (mirroring the Python frontend, which normalizes `-1`
/// to "no budget").
pub thinking_token_budget: Option<i64>,
/// Whether to include reasoning content in the response
#[serde(default = "default_true")]
@@ -108,11 +108,6 @@ pub(super) fn validate_request_compat(
"truncate_prompt_tokens",
"truncate_prompt_tokens is not supported.",
)?;
reject_non_default(
request.thinking_token_budget.as_ref(),
"thinking_token_budget",
"thinking_token_budget is not supported.",
)?;
reject_non_default(
request.media_io_kwargs.as_ref(),
"media_io_kwargs",
@@ -108,6 +108,7 @@ pub(super) fn prepare_completion_request(
seed: request.seed,
max_tokens,
min_tokens: request.min_tokens,
thinking_token_budget: request.thinking_token_budget,
logprobs,
prompt_logprobs,
min_p: request.min_p,
@@ -266,6 +267,34 @@ mod tests {
assert!(!prepared.text_request.decode_options.skip_special_tokens);
}
#[test]
fn prepare_completion_request_passes_through_thinking_token_budget() {
let prepare = |budget: serde_json::Value| {
let request: CompletionRequest = serde_json::from_value(json!({
"model": "Qwen/Qwen1.5-0.5B-Chat",
"prompt": "hello",
"thinking_token_budget": budget,
}))
.expect("parse request");
prepare_completion_request(
request,
&served(&["Qwen/Qwen1.5-0.5B-Chat"]),
ResolvedRequestContext::default(),
)
.expect("prepare")
.text_request
.sampling_params
.thinking_token_budget
};
// The convert layer forwards the raw value verbatim (including the `-1`
// "unlimited" sentinel); normalization/validation happens during
// lowering (see `vllm_text::lower`).
assert_eq!(prepare(json!(64)), Some(64));
assert_eq!(prepare(json!(-1)), Some(-1));
assert_eq!(prepare(json!(null)), None);
}
#[test]
fn prepare_completion_request_maps_stream_usage_and_token_format_options() {
let request: CompletionRequest = serde_json::from_value(json!({
@@ -146,6 +146,11 @@ pub struct CompletionRequest {
/// Additional kwargs for structured outputs
pub structured_outputs: Option<Value>,
/// Token budget for reasoning/thinking. Accepts a non-negative integer, or
/// `-1` for unlimited (mirroring the Python frontend, which normalizes `-1`
/// to "no budget").
pub thinking_token_budget: Option<i64>,
/// Request scheduling priority (lower means earlier; default 0)
pub priority: Option<i32>,
+2
View File
@@ -20,6 +20,8 @@ pub enum Error {
Logprobs(#[from] LogprobsError),
#[error(transparent)]
OutOfVocab(#[from] OutOfVocabError),
#[error("`thinking_token_budget` must be a non-negative integer or -1 for unlimited.")]
InvalidThinkingTokenBudget,
#[error("text request stream `{request_id}` closed before terminal output")]
StreamClosedBeforeTerminalOutput { request_id: String },
#[error(transparent)]
+54
View File
@@ -87,6 +87,7 @@ pub fn lower_sampling_params(
seed,
max_tokens,
min_tokens,
thinking_token_budget,
logprobs,
prompt_logprobs,
min_p,
@@ -128,6 +129,7 @@ pub fn lower_sampling_params(
prompt_len,
)?;
let min_tokens = min_tokens.unwrap_or(0);
let thinking_token_budget = normalize_thinking_token_budget(thinking_token_budget)?;
let frequency_penalty = frequency_penalty.unwrap_or(0.0);
let presence_penalty = presence_penalty.unwrap_or(0.0);
@@ -149,6 +151,7 @@ pub fn lower_sampling_params(
seed,
max_tokens,
min_tokens,
thinking_token_budget,
logprobs,
prompt_logprobs,
min_p,
@@ -170,6 +173,21 @@ pub fn lower_sampling_params(
Ok(params)
}
/// Normalize the user-facing `thinking_token_budget` into the engine value.
///
/// Mirrors Python's `validate_thinking_token_budget`
/// (<https://github.com/vllm-project/vllm/blob/ecf9d83520eb217401b47d8a5451a27c5231b8c2/vllm/sampling_params.py#L35-L55>):
/// `None` and the `-1` "unlimited" sentinel both map to `None`; any other
/// negative value is rejected; non-negative values pass through unchanged. Like
/// Python's `int`, no upper bound is imposed.
fn normalize_thinking_token_budget(value: Option<i64>) -> Result<Option<u64>> {
match value {
None | Some(-1) => Ok(None),
Some(budget) if budget >= 0 => Ok(Some(budget as u64)),
Some(_) => Err(Error::InvalidThinkingTokenBudget),
}
}
/// Convert bad-word strings into token-ID sequences, following the Python vLLM
/// logic in `SamplingParams.update_from_tokenizer()`.
///
@@ -366,6 +384,36 @@ mod tests {
)
}
#[test]
fn lower_sampling_params_normalizes_thinking_token_budget() {
let lower = |budget: Option<i64>| {
lower_sampling_params_with_limits(
SamplingParams {
thinking_token_budget: budget,
..SamplingParams::default()
},
sample_sampling_limits(),
)
};
// Non-negative budgets (including 0) pass through unchanged.
assert_eq!(lower(Some(256)).unwrap().thinking_token_budget, Some(256));
assert_eq!(lower(Some(0)).unwrap().thinking_token_budget, Some(0));
// `None` and the `-1` "unlimited" sentinel both disable the budget.
assert_eq!(lower(None).unwrap().thinking_token_budget, None);
assert_eq!(lower(Some(-1)).unwrap().thinking_token_budget, None);
// No upper bound is imposed, matching Python's `int`.
assert_eq!(
lower(Some(i64::from(u32::MAX) + 1)).unwrap().thinking_token_budget,
Some(u64::from(u32::MAX) + 1)
);
// Other negatives are rejected.
assert!(matches!(
lower(Some(-2)),
Err(Error::InvalidThinkingTokenBudget)
));
}
#[test]
fn lower_text_request_applies_python_style_eos_hints() {
let prepared = lower_text_request(
@@ -386,6 +434,7 @@ mod tests {
seed: None,
max_tokens: 999997,
min_tokens: 0,
thinking_token_budget: None,
logprobs: None,
prompt_logprobs: None,
min_p: 0.0,
@@ -437,6 +486,7 @@ mod tests {
seed: None,
max_tokens: 999997,
min_tokens: 0,
thinking_token_budget: None,
logprobs: None,
prompt_logprobs: None,
min_p: 0.0,
@@ -567,6 +617,7 @@ mod tests {
seed: None,
max_tokens: 40957,
min_tokens: 0,
thinking_token_budget: None,
logprobs: None,
prompt_logprobs: None,
min_p: 0.0,
@@ -628,6 +679,7 @@ mod tests {
seed: None,
max_tokens: 999997,
min_tokens: 0,
thinking_token_budget: None,
logprobs: None,
prompt_logprobs: None,
min_p: 0.0,
@@ -697,6 +749,7 @@ mod tests {
seed: None,
max_tokens: 32,
min_tokens: 2,
thinking_token_budget: None,
logprobs: None,
prompt_logprobs: None,
min_p: 0.1,
@@ -929,6 +982,7 @@ mod tests {
seed: None,
max_tokens: 128,
min_tokens: 0,
thinking_token_budget: None,
logprobs: None,
prompt_logprobs: None,
min_p: 0.1,
+8 -1
View File
@@ -309,7 +309,7 @@ fn matches_stop_string(stops: &[String], output: &str, new_bytes: usize) -> Opti
.find_map(|(ss_idx, (ss, len, start_off))| {
output[start_off..]
.windows(len)
.rposition(|w| w == ss)
.position(|w| w == ss)
.map(|pos| (ss_idx, start_off + pos))
})
}
@@ -562,6 +562,13 @@ mod tests {
assert_eq!(result, Some((0, 4)));
}
#[test]
fn stop_string_matches_leftmost_with_multiple_new_bytes() {
let stops = vec!["\n".to_string()];
let result = matches_stop_string(&stops, "Answer\n\n", 2);
assert_eq!(result, Some((0, 6)));
}
#[test]
fn stop_string_matches_at_beginning() {
let stops = vec!["say".to_string()];
+7
View File
@@ -56,6 +56,12 @@ pub struct SamplingParams {
pub max_tokens: Option<u32>,
/// Minimum number of tokens to generate before EOS or stop-token handling.
pub min_tokens: Option<u32>,
/// Maximum number of reasoning ("thinking") tokens to emit before the
/// reasoning section is force-closed. `None` or the user-facing `-1`
/// "unlimited" sentinel both disable the budget. The raw value is carried
/// here; `-1` is normalized to `None` (and other negatives rejected) during
/// lowering (see `lower_sampling_params`).
pub thinking_token_budget: Option<i64>,
/// Number of log probabilities to return per generated token.
///
/// `None` disables sample logprobs. `-1` requests the full vocabulary.
@@ -116,6 +122,7 @@ impl Default for SamplingParams {
seed: None,
max_tokens: None,
min_tokens: None,
thinking_token_budget: None,
logprobs: None,
prompt_logprobs: None,
min_p: None,
@@ -22,7 +22,7 @@ import torch
import vllm.config
from tests.compile.backend import TestBackend
from vllm._aiter_ops import is_aiter_found_and_supported, rocm_aiter_ops
from vllm._aiter_ops import rocm_aiter_ops
from vllm.compilation.passes.utility.noop_elimination import NoOpEliminationPass
from vllm.compilation.passes.utility.post_cleanup import PostCleanupPass
from vllm.config import (
@@ -83,9 +83,8 @@ class _ViewDoubleQuantModel(torch.nn.Module):
[_NoViewDoubleQuantModel, _ViewDoubleQuantModel],
ids=["no_view", "with_view"],
)
@pytest.mark.skipif(
not is_aiter_found_and_supported(),
reason="Only test on ROCm with AITER installed and supported",
@pytest.mark.skip(
reason="Skipping for now because pytorch compiler removes one the two quant ops"
)
def test_double_aiter_rms_fp8_group_quant_fusion(
model_cls: type[torch.nn.Module],
+2 -2
View File
@@ -175,7 +175,7 @@ MULTIMODAL_MODELS = {
"facebook/chameleon-7b": PPTestSettings.fast(),
"adept/fuyu-8b": PPTestSettings.fast(),
"zai-org/glm-4v-9b": PPTestSettings.fast(),
"OpenGVLab/InternVL2-1B": PPTestSettings.fast(),
"OpenGVLab/InternVL3-1B": PPTestSettings.fast(),
"llava-hf/llava-1.5-7b-hf": PPTestSettings.fast(),
"llava-hf/llava-v1.6-mistral-7b-hf": PPTestSettings.fast(),
"llava-hf/LLaVA-NeXT-Video-7B-hf": PPTestSettings.fast(),
@@ -203,7 +203,7 @@ TEST_MODELS = [
"intfloat/e5-mistral-7b-instruct",
"BAAI/bge-multilingual-gemma2",
# [MULTIMODAL GENERATION]
"OpenGVLab/InternVL2-1B",
"OpenGVLab/InternVL3-1B",
"microsoft/Phi-3.5-vision-instruct",
"fixie-ai/ultravox-v0_5-llama-3_2-1b",
# [LANGUAGE GENERATION - HYBRID ARCH]
@@ -25,7 +25,7 @@ def server():
"--runner",
"pooling",
"--max-model-len",
"5000",
"16384",
"--enforce-eager",
"--limit-mm-per-prompt",
json.dumps({"video": MAXIMUM_VIDEOS}),
@@ -143,4 +143,4 @@ def test_chat_video_url_request(server: RemoteOpenAIServer, model_name: str):
assert output.model == model_name
assert len(output.data) == 1
assert len(output.data[0].probs) == 2
assert output.usage.prompt_tokens == 4807
assert output.usage.prompt_tokens == 8993
@@ -15,7 +15,7 @@ from vllm.entrypoints.serve.tokenize.protocol import (
TokenizeChatRequest,
TokenizeCompletionRequest,
)
from vllm.entrypoints.serve.tokenize.serving import OpenAIServingTokenization
from vllm.entrypoints.serve.tokenize.serving import ServingTokenization
from vllm.v1.engine.async_llm import AsyncLLM
MODEL_NAME = "openai-community/gpt2"
@@ -58,7 +58,7 @@ class MockModelConfig:
return self.diff_sampling_param or {}
def _build_serving_tokenization(engine: AsyncLLM) -> OpenAIServingTokenization:
def _build_serving_tokenization(engine: AsyncLLM) -> ServingTokenization:
models = OpenAIServingModels(
engine_client=engine,
base_model_paths=BASE_MODEL_PATHS,
@@ -71,8 +71,7 @@ def _build_serving_tokenization(engine: AsyncLLM) -> OpenAIServingTokenization:
chat_template=None,
chat_template_content_format="auto",
)
return OpenAIServingTokenization(
engine,
return ServingTokenization(
models,
openai_serving_render=serving_render,
request_logger=None,
@@ -15,16 +15,14 @@ from vllm.config import (
from vllm.platforms import current_platform
from vllm.platforms.cpu import CpuPlatform
# CudaPlatform and RocmPlatform import their respective compiled C extensions
# at module level, raising ModuleNotFoundError on incompatible builds.
try:
if current_platform.is_cuda():
from vllm.platforms.cuda import CudaPlatform
except (ImportError, ModuleNotFoundError):
else:
CudaPlatform = None
try:
if current_platform.is_rocm():
from vllm.platforms.rocm import RocmPlatform
except (ImportError, ModuleNotFoundError):
else:
RocmPlatform = None
from vllm.v1.attention.backends.registry import AttentionBackendEnum
@@ -434,9 +432,15 @@ def test_per_head_quant_scales_backend_selection(
[
("FLASH_ATTN", True, True), # FlashAttn supports non-causal
("FLASH_ATTN", False, True), # FlashAttn also works with causal
("FLASHINFER", True, False), # FlashInfer does not support non-causal
("FLASHINFER", False, True), # FlashInfer works with causal
],
]
+ (
[
("FLASHINFER", True, True), # FlashInfer supports non-causal
("FLASHINFER", False, True), # FlashInfer works with causal
]
if CudaPlatform is not None
else []
),
)
def test_non_causal_backend_selection(
backend_name: str, use_non_causal: bool, should_succeed: bool
@@ -459,11 +463,12 @@ def test_non_causal_backend_selection(
attention_config=attention_config, cache_config=cache_config
)
if CudaPlatform is None:
pytest.skip("CudaPlatform not available")
platform = CudaPlatform or RocmPlatform
if platform is None:
pytest.skip("CudaPlatform and RocmPlatform are not available")
with (
set_current_vllm_config(vllm_config),
patch("vllm.platforms.current_platform", CudaPlatform()),
patch("vllm.platforms.current_platform", platform()),
):
if should_succeed:
backend = get_attn_backend(
@@ -212,3 +212,69 @@ def test_cutlass_mla_decode(
print(
f"{t:.3f} ms, {FLOPS / 10**9 / t:.0f} TFLOPS,", f"{bytes / 10**6 / t:.0f} GB/s"
)
@pytest.mark.skipif(
not current_platform.has_device_capability(100),
reason=CUTLASS_MLA_UNSUPPORTED_REASON,
)
@torch.inference_mode()
def test_cutlass_mla_decode_cross_layer_view():
"""The kernel must read the cache's page-dim stride instead of assuming
pages are packed back-to-back. A per-layer view into a cross-layer
(block-major) cache has stride(0) inflated by num_layers; outputs must
match a contiguous cache holding the same data exactly."""
device = torch.device("cuda:0")
torch.set_default_dtype(torch.bfloat16)
torch.set_default_device(device)
torch.manual_seed(42)
b, mean_sk, d, dv, block_size = 4, 512, 576, 512, 64
num_layers, layer_idx = 3, 1
scale = math.sqrt(d) ** (-1)
num_pages = b * (mean_sk // block_size)
cache_seqlens = torch.full((b,), mean_sk, dtype=torch.int32)
block_table = torch.arange(num_pages, dtype=torch.int32).view(
b, mean_sk // block_size
)
kv_contig = torch.randn(num_pages, block_size, d)
# Neighbor layers hold random data so packed-pages addressing reads
# garbage rather than zeros.
kv_cross_layer = torch.randn(num_pages, num_layers, block_size, d)
kv_view = kv_cross_layer[:, layer_idx]
kv_view.copy_(kv_contig)
assert kv_view.stride(0) == num_layers * block_size * d
q_nope = torch.randn(b, 128, dv)
q_pe = torch.randn(b, 128, d - dv)
sm_count = num_compute_units(device.index)
workspace_size = ops.sm100_cutlass_mla_get_workspace_size(
mean_sk, b, sm_count, num_kv_splits=1
)
workspace = torch.empty(workspace_size, dtype=torch.uint8)
def run(cache):
out = torch.empty(b, 128, dv)
lse = torch.empty(b, 128, dtype=torch.float32)
ops.sm100_cutlass_mla_decode(
out,
lse,
q_nope,
q_pe,
cache,
cache_seqlens,
block_table,
workspace,
scale,
1,
)
return out, lse
out_contig, lse_contig = run(kv_contig)
out_view, lse_view = run(kv_view)
# Same data and same compute order; only addressing differs.
assert torch.equal(out_contig, out_view)
assert torch.equal(lse_contig, lse_view)
@@ -0,0 +1,566 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
"""Bit-exact kernel equivalence for MLA decode/write kernels on the
cross-layer (block-major) KV cache layout.
The cross-layer layout carves each layer's per-block page out of a single
unified slot, so the per-layer view has an inflated ``stride(0)`` (the full
unified slot) and a non-zero storage offset. These tests confirm the MLA
kernels behind the backends that opt in to the layout (FlashMLA dense,
FlashInfer MLA dense, FlashMLA fp8 sparse, plus the ``concat_and_cache_mla``
write) honor that strided view bit-identically to a contiguous per-layer
cache, and that writes do not bleed into neighbouring layers' segments.
"""
import pytest
import torch
pytestmark = pytest.mark.skipif(
not torch.cuda.is_available(), reason="MLA cache kernels require CUDA"
)
def test_concat_and_cache_mla_into_unified_slot_view():
"""concat_and_cache_mla must write correctly into a per-layer view whose
block stride is the full unified slot (block-major), with zero bleed into
the other layers' segments of the same slot."""
from vllm import _custom_ops as ops
torch.manual_seed(0)
dev = "cuda"
kv_lora_rank = 512
pe = 64
entry = kv_lora_rank + pe
page = 64
num_blocks = 32
ntok = 200
kv_c = torch.randn(ntok, kv_lora_rank, device=dev, dtype=torch.bfloat16)
k_pe = torch.randn(ntok, pe, device=dev, dtype=torch.bfloat16)
slot = torch.randperm(num_blocks * page, device=dev, dtype=torch.int64)[:ntok]
scale = torch.tensor(1.0, device=dev)
def write(cache):
ops.concat_and_cache_mla(kv_c, k_pe, cache, slot, "auto", scale)
# Contiguous per-layer reference: (num_blocks, page, entry).
ref = torch.zeros(num_blocks, page, entry, device=dev, dtype=torch.bfloat16)
write(ref)
# Unified slot holding three layer pages per block. Carve the middle
# layer's view (non-zero offset, block stride == full unified slot).
layer_page_elems = page * entry
n_layers = 3
unified_slot_elems = n_layers * layer_page_elems
big = torch.zeros(num_blocks, unified_slot_elems, device=dev, dtype=torch.bfloat16)
flat = big.view(-1)
offset = layer_page_elems # middle layer
view = torch.as_strided(
flat,
size=(num_blocks, page, entry),
stride=(unified_slot_elems, entry, 1),
storage_offset=offset,
)
assert not view.is_contiguous()
assert view.stride(0) == unified_slot_elems
write(view)
# Bit-exact equivalence and zero bleed into the neighbour segments.
max_diff = (ref.float() - view.float()).abs().max().item()
assert max_diff == 0.0, f"max|Δ| = {max_diff}"
neighbour_lo = torch.as_strided(
flat, (num_blocks, layer_page_elems), (unified_slot_elems, 1), 0
)
neighbour_hi = torch.as_strided(
flat,
(num_blocks, layer_page_elems),
(unified_slot_elems, 1),
2 * layer_page_elems,
)
assert neighbour_lo.abs().max().item() == 0.0
assert neighbour_hi.abs().max().item() == 0.0
def test_flashmla_dense_decode_unified_slot_view():
"""FlashMLA dense decode (FLASHMLA backend, e.g. Kimi-K2-style dense MLA
on Hopper) must read a unified-slot block-major view bit-identically to a
contiguous per-layer cache."""
import vllm.v1.attention.ops.flashmla as fm
ok, reason = fm.is_flashmla_dense_supported()
if not ok:
pytest.skip(reason)
torch.manual_seed(0)
dev = "cuda"
dt = torch.bfloat16
head_dim = 576
hdv = 512
h_q = 128
page = 64
num_blocks = 64
bs = 4
n_layers = 3
layer = 1
q = torch.randn(bs, 1, h_q, head_dim, device=dev, dtype=dt) * 0.1
kv_data = torch.randn(num_blocks, page, 1, head_dim, device=dev, dtype=dt) * 0.1
# (A) contiguous per-layer reference.
cache_contiguous = kv_data.clone().contiguous()
# (B) unified slot: view one layer -> inflated stride(0), non-zero offset.
unified = (
torch.randn(num_blocks, n_layers, page, 1, head_dim, device=dev, dtype=dt) * 0.1
)
unified[:, layer].copy_(kv_data)
cache_view = unified[:, layer]
assert not cache_view.is_contiguous()
assert cache_view.stride(0) == n_layers * page * 1 * head_dim
max_blk = num_blocks // bs
block_table = torch.arange(num_blocks, device=dev, dtype=torch.int32).view(
bs, max_blk
)
cache_seqlens = torch.full((bs,), max_blk * page, device=dev, dtype=torch.int32)
def run(kc):
meta, num_splits = fm.get_mla_metadata()
out, _ = fm.flash_mla_with_kvcache(
q=q,
k_cache=kc,
block_table=block_table,
cache_seqlens=cache_seqlens,
head_dim_v=hdv,
tile_scheduler_metadata=meta,
num_splits=num_splits,
softmax_scale=head_dim**-0.5,
causal=True,
)
return out.clone().float()
out_ref = run(cache_contiguous)
out_view = run(cache_view)
assert torch.isfinite(out_ref).all()
assert out_ref.abs().max().item() > 0.0
assert (out_ref - out_view).abs().max().item() == 0.0
def test_flashinfer_mla_dense_decode_unified_slot_view():
"""FlashInfer MLA dense decode must read a unified-slot block-major view
(inflated stride(0), non-zero storage offset) bit-identically to a
contiguous per-layer cache."""
try:
from flashinfer.decode import trtllm_batch_decode_with_kv_cache_mla
except ImportError:
pytest.skip("flashinfer is not available")
from vllm.platforms import current_platform
if not current_platform.is_device_capability_family(100):
pytest.skip("FlashInfer trtllm-gen MLA requires sm100")
torch.manual_seed(0)
dev = "cuda"
dt = torch.bfloat16
kv_lora_rank = 512
qk_rope_head_dim = 64
qk_nope_head_dim = 128
head_dim = kv_lora_rank + qk_rope_head_dim # 576
num_qo_heads = 128
page = 64
num_blocks = 64
bs = 4
n_layers = 3 # >1 so the per-layer view's block stride is inflated.
layer = 1
q = torch.randn(bs, 1, num_qo_heads, head_dim, device=dev, dtype=dt)
kv_data = torch.randn(num_blocks, 1, page, head_dim, device=dev, dtype=dt)
# (A) contiguous per-layer reference.
kv_contiguous = kv_data.clone().contiguous()
# (B) unified slot: block b of every layer packed together; view one layer
# -> stride(0) is n_layers x larger and storage offset is non-zero.
unified = torch.randn(num_blocks, n_layers, 1, page, head_dim, device=dev, dtype=dt)
unified[:, layer].copy_(kv_data)
kv_view = unified[:, layer]
assert not kv_view.is_contiguous()
assert kv_view.stride(0) == n_layers * 1 * page * head_dim
max_blk = num_blocks // bs
block_tables = torch.arange(num_blocks, device=dev, dtype=torch.int32).view(
bs, max_blk
)
seq_lens = torch.full((bs,), max_blk * page, device=dev, dtype=torch.int32)
ws = torch.empty(128 * 1024 * 1024, dtype=torch.int8, device=dev)
scale = head_dim**-0.5
def run(kv):
return trtllm_batch_decode_with_kv_cache_mla(
query=q,
kv_cache=kv,
workspace_buffer=ws,
qk_nope_head_dim=qk_nope_head_dim,
kv_lora_rank=kv_lora_rank,
qk_rope_head_dim=qk_rope_head_dim,
block_tables=block_tables,
seq_lens=seq_lens,
max_seq_len=int(seq_lens.max().item()),
bmm1_scale=scale,
bmm2_scale=1.0,
).clone()
out_ref = run(kv_contiguous).float()
out_view = run(kv_view).float()
assert torch.isfinite(out_ref).all()
assert (out_ref - out_view).abs().max().item() == 0.0
def test_flashmla_fp8_sparse_decode_unified_slot_view():
"""FlashMLA fp8 sparse decode (DeepSeek V3.2/V4 DSA path) must read a
unified-slot block-major view bit-identically to a contiguous fp8_ds_mla
cache, with finite nonzero output."""
import vllm.v1.attention.ops.flashmla as fm
ok, reason = fm.is_flashmla_sparse_supported()
if not ok:
pytest.skip(reason)
torch.manual_seed(0)
dev = "cuda"
entry = 656 # fp8_ds_mla bytes per token
page = 64
num_blocks = 32
h_q = 128
head_dim = 576
hdv = 512
batch = 2
topk = 128
n_layers = 3
layer = 1
q = torch.randn(batch, 1, h_q, head_dim, device=dev, dtype=torch.bfloat16) * 0.1
# Structurally valid fp8 ds_mla payload: 512B fp8 + 16B f32 scales + 128B
# bf16 rope (random bytes corrupt the scale region and yield NaNs).
nope = (torch.randn(num_blocks, page, 1, 512, device=dev) * 0.1).to(
torch.float8_e4m3fn
)
scales = torch.ones(num_blocks, page, 1, 4, device=dev, dtype=torch.float32)
rope = (torch.randn(num_blocks, page, 1, 64, device=dev) * 0.1).to(torch.bfloat16)
payload = torch.cat(
[
nope.view(torch.uint8).view(num_blocks, page, 1, 512),
scales.view(torch.uint8).view(num_blocks, page, 1, 16),
rope.view(torch.uint8).view(num_blocks, page, 1, 128),
],
dim=-1,
).contiguous()
assert payload.shape[-1] == entry and payload.dtype == torch.uint8
# (A) contiguous reference.
cache_contiguous = payload.clone().contiguous()
# (B) unified slot: view one layer -> inflated stride(0), non-zero offset.
unified = torch.randint(
0, 256, (num_blocks, n_layers, page, 1, entry), device=dev, dtype=torch.uint8
)
unified[:, layer].copy_(payload)
cache_view = unified[:, layer]
assert not cache_view.is_contiguous()
assert cache_view.stride(0) == n_layers * page * 1 * entry
# Sparse indices: each batch uses its own disjoint blocks.
blocks_per_batch = num_blocks // batch
idx = torch.full((batch, 1, topk), -1, device=dev, dtype=torch.int32)
for b in range(batch):
slots: list[int] = []
for blk in range(b * blocks_per_batch, (b + 1) * blocks_per_batch):
slots.extend(blk * page + off for off in range(page))
slots_t = torch.tensor(slots[:topk], device=dev, dtype=torch.int32)
idx[b, 0, : slots_t.numel()] = slots_t
def run(kc):
meta, num_splits = fm.get_mla_metadata()
out, _ = fm.flash_mla_with_kvcache(
q=q,
k_cache=kc,
block_table=None,
cache_seqlens=None,
head_dim_v=hdv,
tile_scheduler_metadata=meta,
is_fp8_kvcache=True,
indices=idx,
softmax_scale=head_dim**-0.5,
)
return out.clone().float()
out_ref = run(cache_contiguous)
out_view = run(cache_view)
assert torch.isfinite(out_ref).all()
assert out_ref.abs().max().item() > 0.0
assert (out_ref - out_view).abs().max().item() == 0.0
def test_indexer_k_quant_and_cache_into_unified_slot_view():
"""indexer_k_quant_and_cache (DeepSeek V3.2/V4 DSA indexer K write) must
write correctly into a per-layer view whose block stride is the full
unified slot, with zero bleed into the other layers' segments."""
from vllm import _custom_ops as ops
torch.manual_seed(0)
dev = "cuda"
head_dim = 128
quant_block_size = 128
block_size = 64
num_blocks = 16
ntok = 100
# Indexer cache layout per token: head_dim fp8 bytes followed by
# head_dim * 4 / quant_block_size scale bytes.
cache_stride = head_dim + head_dim * 4 // quant_block_size
k = torch.randn(ntok, head_dim, device=dev, dtype=torch.bfloat16)
slot = torch.randperm(num_blocks * block_size, device=dev, dtype=torch.int64)[:ntok]
def write(cache):
ops.indexer_k_quant_and_cache(k, cache, slot, quant_block_size, "ue8m0")
# Contiguous per-layer reference.
ref = torch.zeros(
num_blocks, block_size, cache_stride, device=dev, dtype=torch.uint8
)
write(ref)
# Unified slot holding three layer pages per block; carve the middle one.
n_layers = 3
layer = 1
unified = torch.zeros(
num_blocks, n_layers, block_size, cache_stride, device=dev, dtype=torch.uint8
)
view = unified[:, layer]
assert not view.is_contiguous()
assert view.stride(0) == n_layers * block_size * cache_stride
write(view)
assert torch.equal(ref, view.contiguous())
# Zero bleed into the neighbour layers' segments.
assert unified[:, 0].abs().max().item() == 0
assert unified[:, 2].abs().max().item() == 0
def test_flashattn_mla_dense_decode_unified_slot_view():
"""FA3 decode (FLASH_ATTN_MLA backend) must read a unified-slot
block-major view bit-identically to a contiguous per-layer cache."""
try:
from vllm.vllm_flash_attn import flash_attn_varlen_func
except ImportError:
pytest.skip("vllm_flash_attn is not available")
from vllm.v1.attention.backends.fa_utils import flash_attn_supports_mla
if not flash_attn_supports_mla():
pytest.skip("FA3 MLA requires a Hopper device")
torch.manual_seed(0)
dev = "cuda"
dt = torch.bfloat16
kv_lora_rank = 512
rope_dim = 64
entry = kv_lora_rank + rope_dim # 576
h_q = 16
page = 64
num_blocks = 64
bs = 4
n_layers = 3
layer = 1
q_pe = torch.randn(bs, h_q, rope_dim, device=dev, dtype=dt) * 0.1
q_nope = torch.randn(bs, h_q, kv_lora_rank, device=dev, dtype=dt) * 0.1
kv_data = torch.randn(num_blocks, page, entry, device=dev, dtype=dt) * 0.1
# (A) contiguous per-layer reference.
cache_contiguous = kv_data.clone().contiguous()
# (B) unified slot: view one layer -> inflated stride(0), non-zero offset.
unified = torch.randn(num_blocks, n_layers, page, entry, device=dev, dtype=dt) * 0.1
unified[:, layer].copy_(kv_data)
cache_view = unified[:, layer]
assert not cache_view.is_contiguous()
assert cache_view.stride(0) == n_layers * page * entry
max_blk = num_blocks // bs
block_table = torch.arange(num_blocks, device=dev, dtype=torch.int32).view(
bs, max_blk
)
seq_lens = torch.full((bs,), max_blk * page, device=dev, dtype=torch.int32)
cu_seqlens_q = torch.arange(bs + 1, device=dev, dtype=torch.int32)
def run(cache):
kv_c_cache = cache[..., :kv_lora_rank]
k_pe_cache = cache[..., kv_lora_rank:]
out = flash_attn_varlen_func(
q=q_pe,
k=k_pe_cache.unsqueeze(-2), # Add head dim of 1
v=kv_c_cache.unsqueeze(-2), # Add head dim of 1
q_v=q_nope,
max_seqlen_q=1,
cu_seqlens_q=cu_seqlens_q,
max_seqlen_k=int(seq_lens.max().item()),
seqused_k=seq_lens,
block_table=block_table,
softmax_scale=entry**-0.5,
causal=True,
fa_version=3,
)
return out.clone().float()
out_ref = run(cache_contiguous)
out_view = run(cache_view)
assert torch.isfinite(out_ref).all()
assert out_ref.abs().max().item() > 0.0
assert (out_ref - out_view).abs().max().item() == 0.0
def test_flashmla_dense_fp8_decode_unified_slot_view():
"""FlashMLA dense fp8 decode (FLASHMLA backend with quantized KV cache)
must read a unified-slot block-major view bit-identically to a contiguous
per-layer fp8 cache."""
import vllm.v1.attention.ops.flashmla as fm
ok, reason = fm.is_flashmla_dense_supported()
if not ok:
pytest.skip(reason)
torch.manual_seed(0)
dev = "cuda"
head_dim = 576
hdv = 512
h_q = 128
page = 64
num_blocks = 64
bs = 4
n_layers = 3
layer = 1
q = torch.randn(bs, 1, h_q, head_dim, device=dev, dtype=torch.bfloat16) * 0.1
kv_data = (torch.randn(num_blocks, page, head_dim, device=dev) * 0.1).to(
torch.float8_e4m3fn
)
# (A) contiguous per-layer reference.
cache_contiguous = kv_data.clone().contiguous()
# (B) unified slot: view one layer -> inflated stride(0), non-zero offset.
unified = (torch.randn(num_blocks, n_layers, page, head_dim, device=dev) * 0.1).to(
torch.float8_e4m3fn
)
unified[:, layer].copy_(kv_data)
cache_view = unified[:, layer]
assert not cache_view.is_contiguous()
assert cache_view.stride(0) == n_layers * page * head_dim
max_blk = num_blocks // bs
block_table = torch.arange(num_blocks, device=dev, dtype=torch.int32).view(
bs, max_blk
)
cache_seqlens = torch.full((bs,), max_blk * page, device=dev, dtype=torch.int32)
descale = torch.ones(1, device=dev, dtype=torch.float32)
def run(kc):
tile_md, num_splits = fm.get_mla_metadata_dense_fp8(cache_seqlens, h_q, 1)
out, _ = fm.flash_mla_with_kvcache_fp8(
q=q,
k_cache=kc.unsqueeze(-2), # Add head dim of 1
block_table=block_table,
cache_seqlens=cache_seqlens,
head_dim_v=hdv,
tile_scheduler_metadata=tile_md,
num_splits=num_splits,
softmax_scale=head_dim**-0.5,
causal=True,
descale_q=descale,
descale_k=descale,
)
return out.clone().float()
out_ref = run(cache_contiguous)
out_view = run(cache_view)
assert torch.isfinite(out_ref).all()
assert out_ref.abs().max().item() > 0.0
assert (out_ref - out_view).abs().max().item() == 0.0
def test_flashinfer_mla_dense_fp8_decode_unified_slot_view():
"""FlashInfer MLA dense decode with an fp8 KV cache must read a
unified-slot block-major view bit-identically to a contiguous per-layer
cache."""
try:
from flashinfer.decode import trtllm_batch_decode_with_kv_cache_mla
except ImportError:
pytest.skip("flashinfer is not available")
from vllm.platforms import current_platform
if not current_platform.is_device_capability_family(100):
pytest.skip("FlashInfer trtllm-gen MLA requires sm100")
torch.manual_seed(0)
dev = "cuda"
kv_lora_rank = 512
qk_rope_head_dim = 64
qk_nope_head_dim = 128
head_dim = kv_lora_rank + qk_rope_head_dim # 576
num_qo_heads = 128
page = 64
num_blocks = 64
bs = 4
n_layers = 3
layer = 1
# With a quantized KV cache the decode query is quantized to fp8 as well
# (trtllm-gen has no bf16-query x fp8-cache decode kernel).
q = (torch.randn(bs, 1, num_qo_heads, head_dim, device=dev) * 0.1).to(
torch.float8_e4m3fn
)
kv_data = (torch.randn(num_blocks, 1, page, head_dim, device=dev) * 0.1).to(
torch.float8_e4m3fn
)
# (A) contiguous per-layer reference.
kv_contiguous = kv_data.clone().contiguous()
# (B) unified slot: view one layer -> inflated stride(0), non-zero offset.
unified = (
torch.randn(num_blocks, n_layers, 1, page, head_dim, device=dev) * 0.1
).to(torch.float8_e4m3fn)
unified[:, layer].copy_(kv_data)
kv_view = unified[:, layer]
assert not kv_view.is_contiguous()
assert kv_view.stride(0) == n_layers * 1 * page * head_dim
max_blk = num_blocks // bs
block_tables = torch.arange(num_blocks, device=dev, dtype=torch.int32).view(
bs, max_blk
)
seq_lens = torch.full((bs,), max_blk * page, device=dev, dtype=torch.int32)
ws = torch.empty(128 * 1024 * 1024, dtype=torch.int8, device=dev)
scale = head_dim**-0.5
def run(kv):
return trtllm_batch_decode_with_kv_cache_mla(
query=q,
kv_cache=kv,
workspace_buffer=ws,
qk_nope_head_dim=qk_nope_head_dim,
kv_lora_rank=kv_lora_rank,
qk_rope_head_dim=qk_rope_head_dim,
block_tables=block_tables,
seq_lens=seq_lens,
max_seq_len=int(seq_lens.max().item()),
bmm1_scale=scale,
bmm2_scale=1.0,
).clone()
out_ref = run(kv_contiguous).float()
out_view = run(kv_view).float()
assert torch.isfinite(out_ref).all()
assert (out_ref - out_view).abs().max().item() == 0.0
+17 -9
View File
@@ -5,10 +5,12 @@ import math
import random
import time
from collections.abc import Callable
from contextlib import nullcontext
import pytest
import torch
import torch.nn.functional as F
from torch.nn.attention import SDPBackend, sdpa_kernel
from vllm.platforms import current_platform
from vllm.utils.torch_utils import STR_DTYPE_TO_TORCH_DTYPE, set_random_seed
@@ -557,15 +559,21 @@ def test_contexted_kv_attention_alibi(
query_len, seq_len, alibi_slopes, device, dtype
)
# Compute attention
out = F.scaled_dot_product_attention(
q_sdpa,
k_sdpa,
v_sdpa,
attn_mask=alibi_mask,
dropout_p=0.0,
scale=scale,
)
# Compute attention. On ROCm we force use of the Math SDPA backend rather than
# the Flash or Mem-Efficient backends for increased numerical accuracy
if current_platform.is_rocm():
sdpa_context = sdpa_kernel(SDPBackend.MATH)
else:
sdpa_context = nullcontext()
with sdpa_context:
out = F.scaled_dot_product_attention(
q_sdpa,
k_sdpa,
v_sdpa,
attn_mask=alibi_mask,
dropout_p=0.0,
scale=scale,
)
# Reshape output back to [query_len, num_heads, head_size]
out = out.view(num_heads, query_len, head_size).permute(1, 0, 2)
@@ -90,7 +90,9 @@ def _ref_sparse_prefill_ragged(
return out.to(torch.bfloat16)
def _pack_fp8_ds_mla_cache(kv: torch.Tensor, block_size: int) -> torch.Tensor:
def _pack_fp8_ds_mla_cache(
kv: torch.Tensor, block_size: int, is_extra: bool = False
) -> torch.Tensor:
assert kv.shape[-1] == HEAD_DIM
num_tokens = kv.shape[0]
num_blocks = (num_tokens + block_size - 1) // block_size
@@ -101,7 +103,9 @@ def _pack_fp8_ds_mla_cache(kv: torch.Tensor, block_size: int) -> torch.Tensor:
)
cache_flat = cache.view(torch.uint8).flatten()
kv_nope_fp8 = (
kv[:, :NOPE_HEAD_DIM].to(current_platform.fp8_dtype()).view(torch.uint8)
kv[:, :NOPE_HEAD_DIM]
.to(torch.float8_e4m3fn if is_extra else current_platform.fp8_dtype())
.view(torch.uint8)
)
kv_rope_u8 = kv[:, NOPE_HEAD_DIM:].contiguous().view(torch.uint8)
@@ -120,7 +124,7 @@ def _pack_fp8_ds_mla_cache(kv: torch.Tensor, block_size: int) -> torch.Tensor:
def _read_fp8_ds_mla_cache(
cache: torch.Tensor, slot: int, block_size: int
cache: torch.Tensor, slot: int, block_size: int, is_extra: bool = False
) -> torch.Tensor:
cache_flat = cache.view(torch.uint8).flatten()
block_idx = slot // block_size
@@ -129,7 +133,9 @@ def _read_fp8_ds_mla_cache(
token_base = block_base + pos * 576
nope_u8 = cache_flat[token_base : token_base + NOPE_HEAD_DIM]
nope = nope_u8.view(current_platform.fp8_dtype()).to(torch.float32)
nope = nope_u8.view(
torch.float8_e4m3fn if is_extra else current_platform.fp8_dtype()
).to(torch.float32)
rope_u8 = cache_flat[
token_base + NOPE_HEAD_DIM : token_base + NOPE_HEAD_DIM + ROPE_HEAD_DIM * 2
]
@@ -157,7 +163,9 @@ def _ref_sparse_decode_ragged(
]
if extra_cache is not None and extra_rows is not None:
row_kv.extend(
_read_fp8_ds_mla_cache(extra_cache, int(slot), block_size)
_read_fp8_ds_mla_cache(
extra_cache, int(slot), block_size, is_extra=True
)
for slot in extra_rows[query_idx]
)
@@ -326,7 +334,7 @@ def test_sparse_attn_decode_ragged_kernel() -> None:
main_kv = torch.randn(6, HEAD_DIM, dtype=torch.bfloat16, device=device) * 0.125
extra_kv = torch.randn(5, HEAD_DIM, dtype=torch.bfloat16, device=device) * 0.125
main_cache = _pack_fp8_ds_mla_cache(main_kv, block_size)
extra_cache = _pack_fp8_ds_mla_cache(extra_kv, block_size)
extra_cache = _pack_fp8_ds_mla_cache(extra_kv, block_size, is_extra=True)
main_indices = torch.tensor([0, 2, 4, 1], dtype=torch.int32, device=device)
main_indptr = torch.tensor([0, 2, 4], dtype=torch.int32, device=device)
extra_indices = torch.tensor([1, 3, 0], dtype=torch.int32, device=device)
@@ -477,7 +485,7 @@ def test_sparse_attn_decode_split_k_kernel(
rows = [[1, 3, 0, 5, 2, 4], [3, 0, 6]]
extra_kv = torch.randn(7, HEAD_DIM, dtype=torch.bfloat16, device=device) * 0.125
extra_rows = rows
extra_cache = _pack_fp8_ds_mla_cache(extra_kv, block_size)
extra_cache = _pack_fp8_ds_mla_cache(extra_kv, block_size, is_extra=True)
extra_indices, extra_indptr = _ragged_from_rows(rows, device)
attn_sink = (
@@ -231,3 +231,95 @@ def test_decode_attention_fp8(B, L, H_Q, H_KV, D_QK, D_V, CACHE_SIZE, PAGE_SIZE)
# FP8 tolerances match test_mla_backends.py test_backend_correctness.
torch.testing.assert_close(o_ref, o_fp8, atol=5e-1, rtol=1e-2)
@pytest.mark.parametrize(
"H_Q,H_KV,D_QK,D_V,is_mla",
[
(16, 1, 576, 512, True), # MLA path (grouped kernel, v = trans(k))
(32, 8, 128, 128, False), # GQA path (grouped kernel)
(32, 32, 128, 128, False), # MHA path (normal kernel)
],
)
@pytest.mark.parametrize("PAGE_SIZE", [16])
def test_decode_attention_cross_layer_view(H_Q, H_KV, D_QK, D_V, is_mla, PAGE_SIZE):
"""The kernel must honor the cache's page-dim stride, not assume pages are
packed back-to-back. A per-layer view into a cross-layer (block-major)
cache has stride(0) inflated by num_layers; outputs must match a
contiguous cache holding the same data exactly."""
B = 3
seq_len = 1027
CACHE_SIZE = 16384
NUM_LAYERS = 3
LAYER_IDX = 1
dtype = torch.bfloat16
sm_scale = 1.0 / (D_QK**0.5)
num_kv_splits = 8
num_pages = CACHE_SIZE // PAGE_SIZE
num_pages_per_batch = cdiv(seq_len, PAGE_SIZE)
req_to_page = torch.randint(
0, num_pages, (B, num_pages_per_batch), device=DEVICE_TYPE
)
q = torch.randn(B, H_Q, D_QK, dtype=dtype, device=DEVICE_TYPE)
b_seq_len = torch.full((B,), seq_len, device=DEVICE_TYPE)
# Reference: contiguous paged cache.
k_ref = torch.randn(
num_pages, PAGE_SIZE, H_KV, D_QK, dtype=dtype, device=DEVICE_TYPE
)
if is_mla:
v_ref = k_ref[..., :D_V]
else:
v_ref = torch.randn(
num_pages, PAGE_SIZE, H_KV, D_V, dtype=dtype, device=DEVICE_TYPE
)
# Cross-layer cache: all layers' pages for a block are adjacent. The
# per-layer view has the same shape as the contiguous cache but
# stride(0) is NUM_LAYERS x larger. Neighbor layers hold random data so
# any packed-pages addressing reads garbage rather than zeros.
k_xl = torch.randn(
num_pages, NUM_LAYERS, PAGE_SIZE, H_KV, D_QK, dtype=dtype, device=DEVICE_TYPE
)
k_view = k_xl[:, LAYER_IDX]
k_view.copy_(k_ref)
assert k_view.stride(0) == NUM_LAYERS * PAGE_SIZE * H_KV * D_QK
if is_mla:
v_view = k_view[..., :D_V]
else:
v_xl = torch.randn(
num_pages, NUM_LAYERS, PAGE_SIZE, H_KV, D_V, dtype=dtype, device=DEVICE_TYPE
)
v_view = v_xl[:, LAYER_IDX]
v_view.copy_(v_ref)
def run(k_buffer, v_buffer):
o = torch.zeros(B, H_Q, D_V, dtype=dtype, device=DEVICE_TYPE)
lse = torch.zeros(B, H_Q, dtype=dtype, device=DEVICE_TYPE)
attn_logits = torch.empty(
(B, H_Q, num_kv_splits, D_V + 1), dtype=torch.float32, device=DEVICE_TYPE
)
decode_attention_fwd(
q,
k_buffer,
v_buffer,
o,
lse,
req_to_page,
b_seq_len,
attn_logits,
num_kv_splits,
sm_scale,
PAGE_SIZE,
is_mla=is_mla,
)
return o, lse
o_ref, lse_ref = run(k_ref, v_ref)
o_xl, lse_xl = run(k_view, v_view)
# Same data and same compute order; only addressing differs.
assert torch.equal(o_ref, o_xl)
assert torch.equal(lse_ref, lse_xl)
@@ -18,11 +18,7 @@ HEAD_SIZES = [128, 256]
BLOCK_SIZES = [16]
DTYPES = [torch.bfloat16]
QDTYPES = (
[None, torch.float8_e4m3fn]
if not current_platform.is_rocm()
else [None, torch.float8_e4m3fnuz]
)
QDTYPES = [None, current_platform.fp8_dtype()]
FP8_DTYPE = current_platform.fp8_dtype()
# one value large enough to test overflow in index calculation.
+59 -3
View File
@@ -10,8 +10,12 @@ from torch.multiprocessing import spawn
from tests.kernels.utils import opcheck
from tests.utils import ensure_current_vllm_config, init_test_distributed_environment
from vllm.distributed import cleanup_dist_env_and_memory
from vllm.model_executor.layers.minimax_rms_norm import MiniMaxText01RMSNormTP
from vllm.model_executor.layers.minimax_rms_norm import (
MiniMaxText01RMSNormTP,
rms_norm_tp,
)
from vllm.platforms import current_platform
from vllm.triton_utils import HAS_TRITON
from vllm.utils.network_utils import get_open_port
from vllm.utils.torch_utils import set_random_seed
@@ -54,8 +58,19 @@ def _worker_forward_qk(
torch.manual_seed(seed + 1000 + local_rank)
qkv = torch.randn(num_tokens, hq + hk + hk, dtype=dtype, device="cuda")
q_ref, k_ref, v_ref = qkv.clone().split([hq, hk, hk], dim=-1)
ref_q, ref_k = MiniMaxText01RMSNormTP.forward_qk(q_norm, k_norm, q_ref, k_ref)
# Reference: eager all-reduce path. ``forward_qk`` no longer all-reduces
# the variance (it is the tp==1 / already-reduced building block), so the
# multi-rank reference must use the eager path that performs the global
# variance all-reduce, matching the fused kernel below.
ref_q, ref_k = rms_norm_tp._minimax_qk_norm_tp_eager(
qkv.clone(),
q_norm.weight,
k_norm.weight,
hq,
hk,
world_size,
eps,
)
# Set up Lamport workspace.
from vllm.distributed.parallel_state import get_tp_group
@@ -150,3 +165,44 @@ def test_minimax_reduce_rms_qk(
nprocs=world_size,
join=True,
)
@pytest.mark.skipif(
not current_platform.is_cuda() or not HAS_TRITON,
reason="CUDA and Triton required",
)
@pytest.mark.parametrize("num_tokens", [1, 7, 128, 333, 2049])
@pytest.mark.parametrize("hidden_dims", [(3072, 512), (768, 256), (3000, 500)])
@pytest.mark.parametrize("dtype", [torch.bfloat16, torch.float16])
@pytest.mark.parametrize("tp_world", [1, 4, 8])
@pytest.mark.parametrize("eps", [1e-6])
@pytest.mark.parametrize("seed", [42])
def test_minimax_qk_norm_triton_fallback(
monkeypatch, num_tokens, hidden_dims, dtype, tp_world, eps, seed
):
"""Single-GPU check: Triton fallback kernels vs the pure-torch reference.
The all-reduce is a TP communication barrier, so it is monkeypatched to
identity here; both the Triton path and the reference see the same
(patched) reduction. This validates the kernel math and the folded
``/ tp_world`` scaling without needing multiple ranks -- ``hidden_dims``
are the per-rank q/k segment widths.
"""
monkeypatch.setattr(rms_norm_tp, "_all_reduce_variance", lambda v: v)
q_size, kv_size = hidden_dims
device = "cuda"
torch.manual_seed(seed)
qkv = torch.randn(num_tokens, q_size + 2 * kv_size, dtype=dtype, device=device)
q_weight = torch.randn(q_size, dtype=dtype, device=device)
k_weight = torch.randn(kv_size, dtype=dtype, device=device)
q_triton, k_triton = rms_norm_tp._minimax_qk_norm_tp_fallback(
qkv, q_weight, k_weight, q_size, kv_size, 0, tp_world, eps
)
q_ref, k_ref = rms_norm_tp._minimax_qk_norm_tp_eager(
qkv, q_weight, k_weight, q_size, kv_size, tp_world, eps
)
torch.testing.assert_close(q_triton, q_ref, atol=3e-2, rtol=3e-2)
torch.testing.assert_close(k_triton, k_ref, atol=3e-2, rtol=3e-2)
+8 -8
View File
@@ -9,6 +9,7 @@ import pytest
import torch
from packaging import version
from vllm._aiter_ops import is_aiter_found
from vllm.platforms import current_platform
from vllm.utils.flashinfer import has_flashinfer
@@ -31,17 +32,15 @@ HOPPER_MXFP4_BF16_AVAILABLE = (
# ROCm platform and dependencies
ROCM_AVAILABLE = current_platform.is_rocm()
ROCM_TRITON_KERNELS_AVAILABLE = False
ROCM_AITER_AVAILABLE = False
ROCM_AITER_AVAILABLE = is_aiter_found()
ROCM_GFX950 = False
if ROCM_AVAILABLE:
from vllm._aiter_ops import rocm_aiter_ops
from vllm.platforms.rocm import on_gfx950
from vllm.utils.import_utils import has_triton_kernels
ROCM_TRITON_KERNELS_AVAILABLE = has_triton_kernels()
ROCM_GFX950 = on_gfx950()
ROCM_AITER_AVAILABLE = rocm_aiter_ops.is_enabled()
if ROCM_AITER_AVAILABLE:
from aiter.ops.triton.moe.quant_moe import upcast_from_mxfp
@@ -83,7 +82,7 @@ def enable_pickle(monkeypatch):
[
ModelCase("fxmarty/qwen_1.5-moe-a2.7b-mxfp4", tp=2),
ModelCase("fxmarty/deepseek_r1_3_layers_mxfp4", tp=8),
ModelCase("fxmarty/Llama-4-Scout-17B-16E-Instruct-2-layers-mxfp4", tp=1),
ModelCase("mawong-amd/Llama-4-Scout-17B-16E-Instruct-2-layers-mxfp4", tp=1),
ModelCase("fxmarty/Llama-3.1-70B-Instruct-2-layers-mxfp6", tp=1),
ModelCase("fxmarty/Llama-3.1-70B-Instruct-2-layers-mxfp6", tp=4),
],
@@ -102,6 +101,7 @@ def test_mxfp4_loading_and_execution_moe(vllm_runner, model_case: ModelCase):
tensor_parallel_size=model_case.tp,
load_format="dummy",
compilation_config={"cudagraph_capture_sizes": [16]},
gpu_memory_utilization=0.8, # mxfp6 models use more scratch space
) as llm:
# Disabled as check_model is broken: https://github.com/vllm-project/vllm/pull/18465#issuecomment-3329880562
# def check_model(model):
@@ -1267,7 +1267,7 @@ def test_rocm_mxfp4_moe_oracle(
This test validates that the oracle functions work end-to-end:
- select_mxfp4_moe_backend() selects a valid backend
- convert_to_mxfp4_moe_kernel_format() converts weights without error
- convert_gpt_oss_weight_to_mxfp4_moe_kernel_format() converts weights without error
- make_mxfp4_moe_quant_config() builds a valid quant config
- make_mxfp4_moe_kernel() creates a kernel that runs without error
- The kernel output is within accuracy tolerance of reference
@@ -1287,7 +1287,7 @@ def test_rocm_mxfp4_moe_oracle(
from vllm.model_executor.layers.fused_moe.oracle.mxfp4 import (
Mxfp4MoeBackend,
backend_to_kernel_cls,
convert_to_mxfp4_moe_kernel_format,
convert_gpt_oss_weight_to_mxfp4_moe_kernel_format,
make_mxfp4_moe_kernel,
make_mxfp4_moe_quant_config,
)
@@ -1387,7 +1387,7 @@ def test_rocm_mxfp4_moe_oracle(
# Convert weights using oracle
w13_conv, w2_conv, w13_scale_conv, w2_scale_conv, w13_bias_conv, w2_bias_conv = (
convert_to_mxfp4_moe_kernel_format(
convert_gpt_oss_weight_to_mxfp4_moe_kernel_format(
mxfp4_backend=backend,
layer=layer, # type: ignore[arg-type]
w13_weight=w13_quant,
@@ -1423,7 +1423,7 @@ def test_rocm_mxfp4_moe_oracle(
mxfp4_backend=backend,
experts_cls=experts_cls,
routing_tables=None,
shared_experts=None,
layer=None,
)
# Create inputs
+11 -11
View File
@@ -8,8 +8,8 @@ from vllm.model_executor.kernels.mhc.tilelang import (
_tilelang_hc_prenorm_gemm,
_torch_hc_prenorm_gemm,
)
from vllm.model_executor.layers.mhc import HAS_TILELANG_MHC
from vllm.platforms import current_platform
from vllm.utils.import_utils import has_tilelang
from vllm.utils.torch_utils import set_random_seed
DEVICE = current_platform.device_type
@@ -97,8 +97,8 @@ def hc_head_ref(
@pytest.mark.skipif(
not (current_platform.is_cuda_alike() and has_tilelang()),
reason="CUDA or ROCm and tilelang required",
not HAS_TILELANG_MHC,
reason="TileLang MHC support required",
)
@pytest.mark.parametrize("num_tokens", [1, 4, 8, 128])
@pytest.mark.parametrize("hidden_size", [4096, 7168])
@@ -150,8 +150,8 @@ def test_mhc_pre_tilelang(num_tokens, hidden_size, hc_mult):
@pytest.mark.skipif(
not (current_platform.is_cuda_alike() and has_tilelang()),
reason="CUDA or ROCm and tilelang required",
not HAS_TILELANG_MHC,
reason="TileLang MHC support required",
)
@pytest.mark.parametrize(
("num_tokens", "hidden_size"),
@@ -190,8 +190,8 @@ def test_hc_prenorm_gemm_tilelang(num_tokens, hidden_size):
@pytest.mark.skipif(
not (current_platform.is_cuda_alike() and has_tilelang()),
reason="CUDA or ROCm and tilelang required",
not HAS_TILELANG_MHC,
reason="TileLang MHC support required",
)
@pytest.mark.parametrize("num_tokens", [1, 4, 8, 128])
@pytest.mark.parametrize("hidden_size", [4096, 7168])
@@ -217,8 +217,8 @@ def test_mhc_post_tilelang(num_tokens, hidden_size, hc_mult):
@pytest.mark.skipif(
not (current_platform.is_cuda_alike() and has_tilelang()),
reason="CUDA or ROCm and tilelang required",
not HAS_TILELANG_MHC,
reason="TileLang MHC support required",
)
@pytest.mark.parametrize("num_tokens", [1, 4, 8, 128])
@pytest.mark.parametrize("hidden_size", [4096, 7168])
@@ -324,8 +324,8 @@ def test_hc_head_triton(num_tokens, hidden_size, hc_mult):
@pytest.mark.skipif(
not (current_platform.is_cuda_alike() and has_tilelang()),
reason="CUDA or ROCm and tilelang required",
not HAS_TILELANG_MHC,
reason="TileLang MHC support required",
)
@pytest.mark.parametrize("num_tokens", [1, 4, 8, 128])
@pytest.mark.parametrize("hidden_size", [4096, 7168])
+21 -8
View File
@@ -124,6 +124,7 @@ def test_replace_submodules(default_vllm_config, dist_init, dummy_model):
max_lora_rank=8, max_cpu_loras=8, max_loras=8, lora_dtype=DEFAULT_DTYPE
),
torch.device(DEVICES[0]),
default_vllm_config,
)
model = manager.model
assert isinstance(model.get_submodule("dense1"), ColumnParallelLinearWithLoRA)
@@ -152,6 +153,7 @@ def test_wrap_replicated_linear_subclasses(default_vllm_config, dist_init, dummy
max_lora_rank=8, max_cpu_loras=8, max_loras=8, lora_dtype=DEFAULT_DTYPE
),
torch.device(DEVICES[0]),
default_vllm_config,
)
assert isinstance(
@@ -172,6 +174,7 @@ def test_wrap_gate_linear(default_vllm_config, dist_init, dummy_model):
max_lora_rank=8, max_cpu_loras=8, max_loras=8, lora_dtype=DEFAULT_DTYPE
),
torch.device(DEVICES[0]),
default_vllm_config,
)
assert isinstance(
@@ -219,6 +222,7 @@ def test_dedup_shared_module_across_paths(default_vllm_config, dist_init, dummy_
max_lora_rank=8, max_cpu_loras=8, max_loras=8, lora_dtype=DEFAULT_DTYPE
),
torch.device(DEVICES[0]),
default_vllm_config,
)
canonical = manager.model.get_submodule("moe.gate")
@@ -263,6 +267,7 @@ def test_lm_head_exempt_from_dedup(default_vllm_config, dist_init, dummy_model):
max_lora_rank=8, max_cpu_loras=8, max_loras=8, lora_dtype=DEFAULT_DTYPE
),
torch.device(DEVICES[0]),
default_vllm_config,
)
# lm_head's special handling still ran: logits_processor got wrapped
@@ -293,6 +298,7 @@ def test_skip_unsupported_matched_modules(default_vllm_config, dist_init, dummy_
max_lora_rank=8, max_cpu_loras=8, max_loras=8, lora_dtype=DEFAULT_DTYPE
),
torch.device(DEVICES[0]),
default_vllm_config,
)
# Should not crash and should keep unsupported matched modules unchanged.
@@ -325,6 +331,7 @@ def test_target_modules_fail_closed_on_unsupported_matched_modules(
target_modules=["dense1"],
),
torch.device(DEVICES[0]),
default_vllm_config,
)
@@ -374,6 +381,7 @@ def test_lora_model_manager(default_vllm_config, dist_init, dummy_model, device)
max_lora_rank=8, max_cpu_loras=3, max_loras=2, lora_dtype=DEFAULT_DTYPE
),
device=device,
vllm_config=default_vllm_config,
)
assert all(x is None for x in manager.lora_index_to_id)
assert manager.add_adapter(model_lora1)
@@ -442,6 +450,7 @@ def test_lora_lru_cache_model_manager(
max_lora_rank=8, max_cpu_loras=3, max_loras=2, lora_dtype=DEFAULT_DTYPE
),
device=device,
vllm_config=default_vllm_config,
)
assert all(x is None for x in manager.lora_index_to_id)
assert manager.add_adapter(model_lora1)
@@ -535,6 +544,7 @@ def test_lru_lora_model_manager(default_vllm_config, dist_init, dummy_model, dev
max_lora_rank=8, max_cpu_loras=2, max_loras=2, lora_dtype=DEFAULT_DTYPE
),
device=device,
vllm_config=default_vllm_config,
)
assert all(x is None for x in manager.lora_index_to_id)
@@ -642,9 +652,7 @@ def test_lru_lora_model_manager(default_vllm_config, dist_init, dummy_model, dev
@pytest.mark.parametrize("device", DEVICES)
def test_lru_cache_worker_adapter_manager(
default_vllm_config, dist_init, dummy_model, device, tmp_path
):
def test_lru_cache_worker_adapter_manager(dist_init, dummy_model, device, tmp_path):
lora_config = LoRAConfig(
max_lora_rank=8, max_cpu_loras=4, max_loras=4, lora_dtype=DEFAULT_DTYPE
)
@@ -670,7 +678,7 @@ def test_lru_cache_worker_adapter_manager(
worker_adapter_manager.max_num_seqs = 4
worker_adapter_manager.max_num_batched_tokens = 2
worker_adapter_manager.create_lora_manager(dummy_model)
worker_adapter_manager.create_lora_manager(dummy_model, vllm_config)
mapping = LoRAMapping([], [])
worker_adapter_manager.set_active_adapters(
@@ -758,9 +766,7 @@ def test_lru_cache_worker_adapter_manager(
@pytest.mark.parametrize("device", DEVICES)
def test_worker_adapter_manager(
default_vllm_config, dist_init, dummy_model_gate_up, device, tmp_path
):
def test_worker_adapter_manager(dist_init, dummy_model_gate_up, device, tmp_path):
# Should remove every LoRA not specified in the request.
lora_config = LoRAConfig(
max_lora_rank=8, max_cpu_loras=4, max_loras=4, lora_dtype=DEFAULT_DTYPE
@@ -774,7 +780,7 @@ def test_worker_adapter_manager(
worker_adapter_manager = WorkerLoRAManager(vllm_config, device, EMBEDDING_MODULES)
worker_adapter_manager.vocab_size = dummy_model_gate_up.unpadded_vocab_size
worker_adapter_manager.create_lora_manager(dummy_model_gate_up)
worker_adapter_manager.create_lora_manager(dummy_model_gate_up, vllm_config)
dummy_lora_files = f"{tmp_path}/lora_adapter"
os.makedirs(dummy_lora_files, exist_ok=True)
@@ -894,6 +900,7 @@ def test_packed_loras(default_vllm_config, dist_init, dummy_model_gate_up, devic
max_lora_rank=8, max_cpu_loras=2, max_loras=2, lora_dtype=DEFAULT_DTYPE
),
device=device,
vllm_config=default_vllm_config,
)
model = manager.model
@@ -944,6 +951,7 @@ def _test_target_modules(
device: str,
expected_lora: list[tuple[str, type]],
expected_no_lora: list[tuple[str, type]],
vllm_config,
):
"""Create a LoRAModelManager and assert which modules have LoRA applied."""
LoRAModelManager(
@@ -959,6 +967,7 @@ def _test_target_modules(
target_modules=target_modules,
),
device=device,
vllm_config=vllm_config,
)
for module_path, lora_cls in expected_lora:
assert isinstance(model.get_submodule(module_path), lora_cls)
@@ -981,6 +990,7 @@ def test_target_modules_config(default_vllm_config, dist_init, dummy_model, devi
("dense2", RowParallelLinearWithLoRA),
("layer1.dense2", RowParallelLinearWithLoRA),
],
vllm_config=default_vllm_config,
)
@@ -998,6 +1008,7 @@ def test_target_modules_multiple(default_vllm_config, dist_init, dummy_model, de
("layer1.dense2", RowParallelLinearWithLoRA),
],
expected_no_lora=[],
vllm_config=default_vllm_config,
)
@@ -1017,6 +1028,7 @@ def test_target_modules_none_uses_all(
("layer1.dense2", RowParallelLinearWithLoRA),
],
expected_no_lora=[],
vllm_config=default_vllm_config,
)
@@ -1036,4 +1048,5 @@ def test_target_modules_match_packed_runtime_modules(
("layer1.dense1", ColumnParallelLinearWithLoRA),
("layer1.dense2", RowParallelLinearWithLoRA),
],
vllm_config=default_vllm_config,
)
@@ -810,29 +810,6 @@ VLM_TEST_SETTINGS = {
hf_output_post_proc=model_utils.minicpmv_trunc_hf_output,
patch_hf_runner=model_utils.minicpmv_26_patch_hf_runner,
),
"minimax_vl_01": VLMTestInfo(
models=["MiniMaxAI/MiniMax-VL-01"],
prompt_formatter=lambda img_prompt: f"<beginning_of_sentence>user: {img_prompt} assistant:<end_of_sentence>", # noqa: E501
img_idx_to_prompt=lambda _: "<image>",
test_type=(VLMTestType.IMAGE, VLMTestType.MULTI_IMAGE),
max_model_len=8192,
max_num_seqs=4,
dtype="bfloat16",
hf_output_post_proc=model_utils.minimax_vl_01_hf_output,
patch_hf_runner=model_utils.minimax_vl_01_patch_hf_runner,
auto_cls=AutoModelForImageTextToText,
marks=[
large_gpu_mark(min_gb=80),
# TODO: [ROCm] Fix pickle issue with ROCm spawn and tp>1
pytest.mark.skipif(
current_platform.is_rocm(),
reason=(
"ROCm: Model too large for single GPU; "
"multi-GPU blocked by HF _LazyConfigMapping pickle issue with spawn"
),
),
],
),
"molmo": VLMTestInfo(
models=["allenai/Molmo-7B-D-0924"],
test_type=(VLMTestType.IMAGE, VLMTestType.MULTI_IMAGE),
@@ -245,13 +245,6 @@ def minicpmv_trunc_hf_output(hf_output: RunnerOutput, model: str) -> RunnerOutpu
return output_ids, output_str, out_logprobs
def minimax_vl_01_hf_output(hf_output: RunnerOutput, model: str) -> RunnerOutput:
output_ids, output_str, out_logprobs = hf_output
if output_str.endswith("<end_of_sentence>"):
output_str = output_str.split("<end_of_sentence>")[0]
return output_ids, output_str, out_logprobs
def ultravox_trunc_hf_output(hf_output: RunnerOutput, model: str) -> RunnerOutput:
output_ids, output_str, out_logprobs = hf_output
@@ -1023,17 +1016,6 @@ def minicpmv_26_patch_hf_runner(hf_model: HfRunner) -> HfRunner:
return hf_model
def minimax_vl_01_patch_hf_runner(hf_model: HfRunner) -> HfRunner:
orig_generate = hf_model.model.generate
def _generate(self, *args, image_sizes=None, **kwargs):
return orig_generate(*args, decode_text=False, **kwargs)
hf_model.model.generate = types.MethodType(_generate, hf_model.model)
return hf_model
def molmo_patch_hf_runner(hf_model: HfRunner) -> HfRunner:
"""Patches and returns an instance of the HfRunner to use for Molmo."""
hf_processor = hf_model.processor
@@ -152,3 +152,21 @@ def test_colqwen3_5_relevance_ordering(
dtype: str,
) -> None:
_run_relevance_test(vllm_runner, model, dtype=dtype)
def test_colqwen3_5_config_enables_bidirectional_attention() -> None:
"""ColQwen3.5 retrieval must be served BIDIRECTIONAL (is_causal=False) so the
full_attention layers build with AttentionType.ENCODER_ONLY. This guards the
silent-causal regression (no GPU / model load needed)."""
from types import SimpleNamespace
from vllm.model_executor.models.config import (
MODELS_CONFIG_MAP,
ColQwen3_5Config,
)
assert MODELS_CONFIG_MAP["ColQwen3_5"] is ColQwen3_5Config
model_config = SimpleNamespace(hf_config=SimpleNamespace())
ColQwen3_5Config.verify_and_update_model_config(model_config)
assert model_config.hf_config.is_causal is False
@@ -1,113 +0,0 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
import pytest
from PIL import Image
from vllm.multimodal import MULTIMODAL_REGISTRY
from vllm.multimodal.parse import ImageSize
from vllm.multimodal.processing import BaseMultiModalProcessor
from ....conftest import ImageTestAssets
from ...utils import build_model_context
@pytest.mark.parametrize("model_id", ["MiniMaxAI/MiniMax-VL-01"])
@pytest.mark.parametrize("num_imgs", [1, 2])
def test_processor_override(
image_assets: ImageTestAssets,
model_id: str,
num_imgs: int,
):
ctx = build_model_context(
model_id,
mm_processor_kwargs=None,
limit_mm_per_prompt={"image": num_imgs},
)
processor = MULTIMODAL_REGISTRY.create_processor(ctx.model_config)
prompt = "<image>" * num_imgs
image = Image.new("RGB", size=(364, 364))
mm_data = {"image": [image] * num_imgs}
processed_inputs = processor(
prompt,
mm_items=processor.info.parse_mm_data(mm_data),
hf_processor_mm_kwargs={},
)
image_placeholders = processed_inputs["mm_placeholders"]["image"]
assert len(image_placeholders) == num_imgs
def _validate_image_prompt_replacements_one(
processor: BaseMultiModalProcessor,
num_imgs: int,
failed_size_excs: list[tuple[ImageSize, Exception]],
image_size: ImageSize,
) -> None:
prompt = "<image>" * num_imgs
image = Image.new("RGB", size=image_size)
mm_data = {"image": [image] * num_imgs}
try:
processed_inputs = processor(
prompt,
mm_items=processor.info.parse_mm_data(mm_data),
hf_processor_mm_kwargs={},
)
image_placeholders = processed_inputs["mm_placeholders"]["image"]
assert len(image_placeholders) == num_imgs
except Exception as exc:
failed_size_excs.append((image_size, exc))
def _test_image_prompt_replacements(
processor,
*,
num_imgs: int,
image_sizes: list[ImageSize],
) -> None:
failed_size_excs = list[tuple[ImageSize, Exception]]()
for size in image_sizes:
_validate_image_prompt_replacements_one(
processor, num_imgs, failed_size_excs, size
)
if failed_size_excs:
msg = "Found failing image sizes:" + "\n========\n".join(
f"[{size}]\n{exc}" for size, exc in failed_size_excs
)
raise AssertionError(msg)
@pytest.mark.parametrize("model_id", ["MiniMaxAI/MiniMax-VL-01"])
@pytest.mark.parametrize("num_imgs", [1, 2])
def test_processor_prompt_replacements_regression(model_id, num_imgs):
ctx = build_model_context(
model_id,
mm_processor_kwargs=None,
limit_mm_per_prompt={"image": num_imgs},
)
processor = MULTIMODAL_REGISTRY.create_processor(ctx.model_config)
image_ratios = [
(171, 152),
(184, 161),
(198, 176),
(333, 296),
(369, 328),
(488, 183),
(2560, 1669),
]
image_sizes = [
size for w, h in image_ratios for size in [ImageSize(w, h), ImageSize(h, w)]
]
_test_image_prompt_replacements(
processor,
num_imgs=num_imgs,
image_sizes=image_sizes,
)
@@ -138,3 +138,49 @@ def test_processor_multi_video(
assert video_phs[i].offset >= prev_end, (
f"Placeholder {i} overlaps with placeholder {i - 1}"
)
@pytest.mark.parametrize("model_id", [MODEL_ID])
@pytest.mark.parametrize(
"hf_mm_kwargs",
[{"num_frames": [8, 16]}, {"fps": [2.0, 4.0]}],
)
def test_processor_multi_video_list_kwargs(
model_id: str,
hf_mm_kwargs: dict[str, Any],
) -> None:
"""Regression test: a multi-video request with list-valued per-video
``mm_processor_kwargs`` (one ``fps``/``num_frames`` per video) must not
crash.
Before the fix, ``_call_hf_processor`` copied the whole kwargs to every
video without slicing, so ``_get_video_second_idx`` received the list
where a scalar was expected and raised ``TypeError``.
"""
ctx = build_model_context(
model_id,
limit_mm_per_prompt={"image": 0, "video": 2},
)
processor = MULTIMODAL_REGISTRY.create_processor(ctx.model_config)
prompt = (
"<|vision_start|><|video_pad|><|vision_end|>"
"<|vision_start|><|video_pad|><|vision_end|>"
)
mm_data = {
"video": [
_build_video_mm_data(num_frames=16)["video"][0],
_build_video_mm_data(num_frames=32)["video"][0],
]
}
processed = processor(
prompt,
mm_items=processor.info.parse_mm_data(mm_data),
hf_processor_mm_kwargs=hf_mm_kwargs,
)
video_phs = processed["mm_placeholders"].get("video", [])
assert len(video_phs) == 2, (
f"Expected exactly 2 video placeholders, got {len(video_phs)}"
)
+19 -13
View File
@@ -421,15 +421,6 @@ _TEXT_GENERATION_EXAMPLE_MODELS = {
},
trust_remote_code=True,
),
"MiniMaxForCausalLM": _HfExamplesInfo("MiniMaxAI/MiniMax-Text-01-hf"),
"MiniMaxText01ForCausalLM": _HfExamplesInfo(
"MiniMaxAI/MiniMax-Text-01",
trust_remote_code=True,
revision="a59aa9cbc53b9fb8742ca4e9e1531b9802b6fdc3",
),
"MiniMaxM1ForCausalLM": _HfExamplesInfo(
"MiniMaxAI/MiniMax-M1-40k", trust_remote_code=True
),
"MiniMaxM2ForCausalLM": _HfExamplesInfo(
"MiniMaxAI/MiniMax-M2",
trust_remote_code=True,
@@ -1113,10 +1104,6 @@ _MULTIMODAL_EXAMPLE_MODELS = {
"openbmb/MiniCPM-V-4_6",
min_transformers_version="5.7.0",
),
"MiniMaxVL01ForConditionalGeneration": _HfExamplesInfo(
"MiniMaxAI/MiniMax-VL-01",
trust_remote_code=True,
),
"MiniMaxM3SparseForConditionalGeneration": _HfExamplesInfo(
"MiniMaxAI/MiniMax-M3",
trust_remote_code=True,
@@ -1530,6 +1517,16 @@ _SPECULATIVE_DECODING_EXAMPLE_MODELS = {
"Qwen/Qwen3-VL-8B-Instruct",
speculative_model="taobao-mnn/Qwen3-VL-8B-Instruct-Eagle3",
),
"Eagle3Qwen3ForCausalLM": _HfExamplesInfo(
"Qwen/Qwen3-8B",
trust_remote_code=True,
speculative_model=(
"inference-optimization/"
"Qwen3-8B-from-Qwen3-8B_regen-speculators.eagle3-qwen3arch-ckpt1"
),
tokenizer="Qwen/Qwen3-8B",
use_original_num_layers=True,
),
# [PEagle]
"PEagleDraftModel": _HfExamplesInfo(
"Qwen/Qwen3-8B",
@@ -1545,6 +1542,15 @@ _SPECULATIVE_DECODING_EXAMPLE_MODELS = {
tokenizer="Qwen/Qwen3-8B",
use_original_num_layers=True,
),
"PeagleQwen3ForCausalLM": _HfExamplesInfo(
"Qwen/Qwen3-8B",
trust_remote_code=True,
speculative_model=(
"inference-optimization/Qwen3-8B-speculators.peagle-qwen3arch-ckpt4"
),
tokenizer="Qwen/Qwen3-8B",
use_original_num_layers=True,
),
# [MTP]
"DeepSeekMTPModel": _HfExamplesInfo(
"luccafong/deepseek_mtp_main_random",
-5
View File
@@ -98,11 +98,6 @@ def can_initialize(
vllm_config.validate_block_size()
return scheduler_kv_cache_config
if model_arch == "MiniMaxVL01ForConditionalGeneration":
pytest.skip(
"pickle error when loading `transformers.models.auto.CONFIG_MAPPING`"
)
if model_arch == "MoonshotKimiaForCausalLM":
pytest.skip(
"Kimi-Audio requires SpeechToTextConfig "
+7 -1
View File
@@ -507,7 +507,13 @@ def dummy_hf_overrides(
# Only set MoE related config when the model has MoE layers.
# Otherwise all models detected as MoE by _get_transformers_backend_cls.
if model_arch_config.num_experts > 0:
num_experts_per_tok = 1 if model_arch == "Llama4ForConditionalGeneration" else 2
num_experts_per_tok = 2
if model_arch in (
"Llama4ForConditionalGeneration",
"Llama4ForCausalLM",
"EagleLlama4ForCausalLM",
):
num_experts_per_tok = 1
update_dict.update(
{
"num_experts": num_experts,
+25 -2
View File
@@ -15,6 +15,7 @@ from vllm.multimodal.video import (
DynamicVideoBackend,
GLM46VVideoBackend,
Molmo2VideoBackend,
Qwen2VLVideoBackend,
Qwen3VLVideoBackend,
VideoLoader,
VideoSourceMetadata,
@@ -70,11 +71,12 @@ def test_video_loader_type_doesnt_exist():
@pytest.mark.parametrize(
"model_repo, expected_loader_cls",
"model_repo, expected_loader_cls, hf_sample_kwargs",
[
pytest.param(
"allenai/Molmo2-4B",
Molmo2VideoBackend,
None,
marks=pytest.mark.skip(
reason="Video processor not aligned, investigate later.",
),
@@ -83,23 +85,44 @@ def test_video_loader_type_doesnt_exist():
pytest.param(
"zai-org/GLM-4.1V-9B-Thinking",
DynamicVideoBackend,
None,
id="glm4v",
),
pytest.param(
"zai-org/GLM-4.6V-Flash",
GLM46VVideoBackend,
None,
id="glm46v",
),
pytest.param(
"Qwen/Qwen3-VL-4B-Instruct",
Qwen3VLVideoBackend,
None,
id="qwen3vl",
),
# Qwen2-VL/Qwen2.5-VL ship no ``video_processor_type`` in their
# preprocessor config, so resolution relies on the model_type ->
# video processor fallback in get_video_processor_cls_name_from_config.
# They also ship no default fps/num_frames, so the HF sampler needs an
# explicit target rate; pass fps=2 to match the loader default.
pytest.param(
"Qwen/Qwen2-VL-7B-Instruct",
Qwen2VLVideoBackend,
{"fps": 2},
id="qwen2vl",
),
pytest.param(
"Qwen/Qwen2.5-VL-7B-Instruct",
Qwen2VLVideoBackend,
{"fps": 2},
id="qwen2_5_vl",
),
],
)
def test_video_processor_from_model_repo(
model_repo: str,
expected_loader_cls: type,
hf_sample_kwargs: dict[str, int | float] | None,
):
"""Test that a model repo resolves to the correct video loader backend.
@@ -143,7 +166,7 @@ def test_video_processor_from_model_repo(
fps=vllm_meta["fps"],
duration=vllm_meta["duration"],
)
hf_indices = processor.sample_frames(hf_metadata)
hf_indices = processor.sample_frames(hf_metadata, **(hf_sample_kwargs or {}))
vllm_indices = np.array(vllm_meta["frames_indices"])
np.testing.assert_array_equal(
hf_indices,
+2 -2
View File
@@ -43,8 +43,8 @@ MODELS = [
pytest.param(
"Intel/Qwen2-0.5B-Instruct-int4-sym-AutoRound",
marks=pytest.mark.skipif(
not current_platform.is_cuda(),
reason="AWQ AutoRound model only supports CUDA backend for now.",
not (current_platform.is_cuda() or current_platform.is_xpu()),
reason="AWQ AutoRound model only supports CUDA/XPU backend for now.",
),
id="auto_round:auto_awq",
),
+53 -7
View File
@@ -61,8 +61,8 @@ class QuantConfig:
quant_max: float
quant_min: float
kv_quant_mode: KVQuantMode
# INT8 Triton stores truncate; FP8 hardware casts round.
uses_trunc: bool
# INT8 rounds explicitly; FP8 relies on dtype cast rounding.
rounds_before_store: bool
INT8_CONFIG = QuantConfig(
@@ -71,7 +71,7 @@ INT8_CONFIG = QuantConfig(
quant_max=127.0,
quant_min=-128.0,
kv_quant_mode=KVQuantMode.INT8_PER_TOKEN_HEAD,
uses_trunc=True,
rounds_before_store=True,
)
FP8_CONFIG = QuantConfig(
cache_dtype=FP8_DTYPE,
@@ -79,7 +79,7 @@ FP8_CONFIG = QuantConfig(
quant_max=FP8_MAX,
quant_min=FP8_MIN,
kv_quant_mode=KVQuantMode.FP8_PER_TOKEN_HEAD,
uses_trunc=False,
rounds_before_store=False,
)
QUANT_CONFIGS = [INT8_CONFIG, FP8_CONFIG]
@@ -104,7 +104,7 @@ def _quantize_per_token_head_ref(
absmax = data.float().abs().amax(dim=2) # [num_tokens, num_heads]
scales = (absmax / cfg.quant_max).clamp(min=1e-6)
scaled = data.float() * (1.0 / scales[:, :, None])
if cfg.uses_trunc:
if cfg.rounds_before_store:
q = scaled.round().clamp(cfg.quant_min, cfg.quant_max).to(cfg.cache_dtype)
else:
q = scaled.clamp(cfg.quant_min, cfg.quant_max).to(cfg.cache_dtype)
@@ -255,7 +255,7 @@ def test_per_token_head_round_trip_accuracy(
):
"""Verify per-token-head round-trip: kernel dequant matches reference.
INT8: Triton truncates on float->int8 store.
INT8: round-to-nearest before int8 store.
FP8: hardware cast (clamp then cast).
"""
from vllm.v1.attention.ops.triton_reshape_and_cache_flash import (
@@ -315,6 +315,52 @@ def test_per_token_head_round_trip_accuracy(
)
@torch.inference_mode()
def test_int8_per_token_head_raw_cache_matches_round_reference():
"""INT8 cache writes should match round-to-nearest quantization exactly."""
from vllm.v1.attention.ops.triton_reshape_and_cache_flash import (
triton_reshape_and_cache_flash_per_token_head_quant,
)
torch.set_default_device(DEVICE_TYPE)
head_size = 8
block_size = 4
key = torch.tensor(
[[[-127.0, -2.6, -2.4, -1.6, -1.4, -0.6, -0.4, 127.0]]],
dtype=torch.bfloat16,
)
value = -key
key_cache = torch.zeros(1, block_size, 1, head_size, dtype=torch.int8)
value_cache = torch.zeros_like(key_cache)
k_scale_cache = torch.ones(1, block_size, 1, dtype=torch.float32)
v_scale_cache = torch.ones_like(k_scale_cache)
slot_mapping = torch.tensor([2], dtype=torch.long)
triton_reshape_and_cache_flash_per_token_head_quant(
key,
value,
key_cache,
value_cache,
k_scale_cache,
v_scale_cache,
slot_mapping,
)
ref_k_quant, ref_k_scales = _quantize_per_token_head_ref(key, INT8_CONFIG)
ref_v_quant, ref_v_scales = _quantize_per_token_head_ref(value, INT8_CONFIG)
slot = slot_mapping.item()
blk = slot // block_size
off = slot % block_size
assert torch.equal(key_cache[blk, off], ref_k_quant[0])
assert torch.equal(value_cache[blk, off], ref_v_quant[0])
torch.testing.assert_close(k_scale_cache[blk, off], ref_k_scales[0])
torch.testing.assert_close(v_scale_cache[blk, off], ref_v_scales[0])
# ===========================================================================
# 4. Negative slot mapping (padding tokens should be skipped)
# ===========================================================================
@@ -461,7 +507,7 @@ def test_triton_unified_attention_per_token_head_scale(
scaled_k = key_cache_bf16.float() / k_scale_cache[:, :, :, None]
scaled_v = value_cache_bf16.float() / v_scale_cache[:, :, :, None]
if qcfg.uses_trunc:
if qcfg.rounds_before_store:
key_cache_q = (
scaled_k.round().clamp(qcfg.quant_min, qcfg.quant_max).to(qcfg.cache_dtype)
)
@@ -83,20 +83,6 @@ class MiniMaxM3Tokenizer:
return "".join(tokens)
class SplitMiniMaxM3Tokenizer(MiniMaxM3Tokenizer):
"""Tokenizer that exposes marker vocab entries but encodes them as text."""
def tokenize(self, text: str) -> list[str]:
return list(text)
class RuntimeSplitMiniMaxM3Tokenizer(MiniMaxM3Tokenizer):
"""Tokenizer whose runtime output splits markers despite atomic encodes."""
def encode_runtime(self, text: str) -> list[int]:
return [self._add_token(token) for token in list(text)]
def make_parser(
chat_template_kwargs: dict[str, str] | None = None,
) -> tuple[MiniMaxM3ReasoningParser, MiniMaxM3Tokenizer]:
@@ -119,8 +105,7 @@ def run_streaming(
reasoning_end_states: list[bool] = []
for chunk in chunks:
encode_runtime = getattr(tokenizer, "encode_runtime", tokenizer.encode)
delta_token_ids = encode_runtime(chunk)
delta_token_ids = tokenizer.encode(chunk, add_special_tokens=False)
current_text = previous_text + chunk
current_token_ids = previous_token_ids + delta_token_ids
delta = parser.extract_reasoning_streaming(
@@ -189,14 +174,14 @@ def test_nonstreaming_drops_leading_end_tag():
assert content == "answer"
def test_nonstreaming_end_tag_in_content_state_is_dropped():
def test_nonstreaming_non_leading_end_tag_is_content():
parser, _ = make_parser()
request = ChatCompletionRequest(messages=[], model="test-model")
reasoning, content = parser.extract_reasoning("XXX</mm:think>YYY", request)
assert reasoning is None
assert content == "XXXYYY"
assert content == "XXX</mm:think>YYY"
def test_nonstreaming_enabled_mode_starts_in_reasoning():
@@ -261,7 +246,7 @@ def test_streaming_drops_leading_end_tag():
assert end_states == [True, True]
def test_streaming_end_tag_in_content_state_is_dropped():
def test_streaming_non_leading_end_tag_is_content():
parser, tokenizer = make_parser()
reasoning, content, end_states = run_streaming(
@@ -271,7 +256,7 @@ def test_streaming_end_tag_in_content_state_is_dropped():
)
assert reasoning is None
assert content == "XXXYYY"
assert content == "XXX</mm:think>YYY"
assert end_states == [True]
@@ -303,110 +288,6 @@ def test_streaming_plain_content_ends_reasoning_phase():
assert end_states == [True, True]
def test_streaming_split_marker_tokens_are_not_returned():
tokenizer = RuntimeSplitMiniMaxM3Tokenizer()
parser = MiniMaxM3ReasoningParser(tokenizer)
reasoning, content, end_states = run_streaming(
parser,
tokenizer,
["<mm:think>", "Reasoning", " content", "</mm:think>", "content"],
)
assert reasoning == "Reasoning content"
assert content == "content"
assert end_states == [False, False, False, True, True]
def test_streaming_split_marker_text_drives_end_state():
tokenizer = RuntimeSplitMiniMaxM3Tokenizer()
parser = MiniMaxM3ReasoningParser(tokenizer)
previous_text = ""
previous_token_ids: list[int] = []
for chunk in ["<mm:think>", "Reasoning", " content", "</mm:think>"]:
delta_token_ids = tokenizer.encode_runtime(chunk)
current_text = previous_text + chunk
current_token_ids = previous_token_ids + delta_token_ids
parser.extract_reasoning_streaming(
previous_text=previous_text,
current_text=current_text,
delta_text=chunk,
previous_token_ids=previous_token_ids,
current_token_ids=current_token_ids,
delta_token_ids=delta_token_ids,
)
previous_text = current_text
previous_token_ids = current_token_ids
assert parser.is_reasoning_end_streaming(previous_token_ids, []) is True
def test_streaming_split_marker_tokens_enabled_mode():
tokenizer = RuntimeSplitMiniMaxM3Tokenizer()
parser = MiniMaxM3ReasoningParser(
tokenizer, chat_template_kwargs={"thinking_mode": "enabled"}
)
reasoning, content, end_states = run_streaming(
parser,
tokenizer,
["Reasoning", " content", "</mm:think>", "content"],
)
assert reasoning == "Reasoning content"
assert content == "content"
assert end_states == [False, False, True, True]
def test_streaming_split_marker_text_across_deltas():
tokenizer = RuntimeSplitMiniMaxM3Tokenizer()
parser = MiniMaxM3ReasoningParser(tokenizer)
reasoning, content, end_states = run_streaming(
parser,
tokenizer,
["<mm:", "think>", "Reasoning", " content", "</mm:", "think>", "content"],
)
assert reasoning == "Reasoning content"
assert content == "content"
assert end_states == [False, False, False, False, False, True, True]
def test_streaming_split_leading_end_marker_text_across_deltas():
tokenizer = RuntimeSplitMiniMaxM3Tokenizer()
parser = MiniMaxM3ReasoningParser(tokenizer)
reasoning, content, end_states = run_streaming(
parser,
tokenizer,
["</mm:", "think>", "content"],
)
assert reasoning is None
assert content == "content"
assert end_states == [False, True, True]
def test_token_id_helpers_with_split_marker_tokens():
tokenizer = SplitMiniMaxM3Tokenizer()
parser = MiniMaxM3ReasoningParser(tokenizer)
output_ids = tokenizer.encode(
"<mm:think>abc</mm:think>def", add_special_tokens=False
)
open_reasoning_ids = tokenizer.encode("<mm:think>abc", add_special_tokens=False)
content_ids = tokenizer.encode("plain", add_special_tokens=False)
assert parser.is_reasoning_end(output_ids)
assert not parser.is_reasoning_end(open_reasoning_ids)
assert not parser.is_reasoning_end(content_ids)
assert tokenizer.decode(parser.extract_content_ids(output_ids)) == "def"
assert parser.extract_content_ids(open_reasoning_ids) == []
assert parser.extract_content_ids(content_ids) == content_ids
assert parser.count_reasoning_tokens(output_ids) == len(tokenizer.encode("abc"))
def test_token_id_helpers():
parser, tokenizer = make_parser()
output_ids = tokenizer.encode(
+23
View File
@@ -74,6 +74,29 @@ def test_embed_dimensions(model_info: EmbedModelInfo):
pooling_params.verify(model_config)
@dataclass()
class MockMatryoshkaModelConfig:
pooler_config: PoolerConfig
is_matryoshka: bool = True
matryoshka_dimensions: list[int] | None = None
served_model_name: str = "mock-matryoshka-model"
embedding_size: int = 32
def test_embed_dimensions_matryoshka_without_list_upper_bound():
task = "embed"
model_config = MockMatryoshkaModelConfig(
pooler_config=PoolerConfig(seq_pooling_type="CLS"),
matryoshka_dimensions=None,
embedding_size=32,
)
PoolingParams(task=task, dimensions=16).verify(model_config)
with pytest.raises(ValueError):
PoolingParams(task=task, dimensions=64).verify(model_config)
@pytest.mark.parametrize("task", ["classify"])
def test_classify(task):
model_config = MockModelConfig(pooler_config=PoolerConfig(seq_pooling_type="CLS"))
File diff suppressed because it is too large Load Diff
@@ -6,10 +6,12 @@ from types import SimpleNamespace
import pytest
from vllm.model_executor.layers.mamba.linear.minimax_linear_attn import (
MiniMaxText01LinearAttention,
)
from vllm.model_executor.layers.mamba.mamba_mixer import MambaMixer
from vllm.model_executor.layers.mamba.mamba_mixer2 import MambaMixer2
from vllm.model_executor.layers.mamba.short_conv import ShortConv
from vllm.model_executor.models.minimax_text_01 import MiniMaxText01LinearAttention
from vllm.v1.attention.backends.linear_attn import LinearAttentionBackend
from vllm.v1.attention.backends.mamba1_attn import Mamba1AttentionBackend
from vllm.v1.attention.backends.mamba2_attn import Mamba2AttentionBackend
@@ -16,6 +16,7 @@ from tests.v1.attention.utils import (
create_vllm_config,
)
from vllm.config import SpeculativeConfig
from vllm.config.compilation import CUDAGraphMode
from vllm.v1.attention.backends.gdn_attn import (
GDNAttentionMetadata,
GDNAttentionMetadataBuilder,
@@ -123,9 +124,15 @@ GDN_BUILD_TEST_CASES = {
def _create_gdn_builder(
num_speculative_tokens: int = 0,
full_cuda_graph: bool = False,
) -> GDNAttentionMetadataBuilder:
"""Create a GDNAttentionMetadataBuilder with minimal config."""
vllm_config = create_vllm_config(block_size=BLOCK_SIZE)
vllm_config = create_vllm_config(
model_name="Qwen/Qwen3.5-0.8B",
block_size=BLOCK_SIZE,
)
if full_cuda_graph:
vllm_config.compilation_config.cudagraph_mode = CUDAGraphMode.FULL_AND_PIECEWISE
if num_speculative_tokens > 0:
vllm_config.speculative_config = SpeculativeConfig(
method="ngram",
@@ -189,3 +196,28 @@ def test_has_initial_state_after_reclassification():
assert meta.has_initial_state is not None
# req0 has context_lens = 65 - 1 = 64 > 0, so has_initial_state[0] = True
assert meta.has_initial_state[0].item() is True
def test_full_cudagraph_spec_metadata_uses_request_count():
"""FULL cudagraph token padding must not pad request-indexed metadata."""
num_speculative_tokens = 3
builder = _create_gdn_builder(
num_speculative_tokens=num_speculative_tokens,
full_cuda_graph=True,
)
batch = BatchSpec(seq_lens=[80, 96], query_lens=[4, 4])
meta = _build(builder, batch, num_decode_draft_tokens=[3, 3])
assert meta.num_spec_decodes == batch.batch_size
assert meta.num_spec_decode_tokens == batch.compute_num_tokens()
assert meta.spec_state_indices_tensor is not None
assert meta.spec_state_indices_tensor.shape == (
batch.batch_size,
num_speculative_tokens + 1,
)
assert meta.spec_sequence_masks is not None
assert meta.spec_sequence_masks.shape == (batch.batch_size,)
assert meta.spec_query_start_loc is not None
assert meta.spec_query_start_loc.shape == (batch.batch_size + 1,)
assert meta.num_accepted_tokens is not None
assert meta.num_accepted_tokens.shape == (batch.batch_size,)
@@ -0,0 +1,460 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
from collections.abc import Callable
import pytest
import vllm.v1.core.kv_cache_utils as kv_cache_utils
from vllm.distributed.kv_events import BlockRemoved, BlockStored
from vllm.sampling_params import SamplingParams
from vllm.utils.hashing import sha256
from vllm.v1.core.block_pool import BlockPool
from vllm.v1.core.kv_cache_utils import (
BlockHash,
BlockHashListWithBlockSize,
KVCacheBlock,
get_request_block_hasher,
hash_block_tokens,
init_none_hash,
)
from vllm.v1.request import Request
pytestmark = pytest.mark.cpu_test
@pytest.fixture(autouse=True)
def _auto_init_hash_fn():
init_none_hash(sha256)
def make_request(
request_id: str,
prompt_token_ids: list[int],
hash_block_size: int,
hash_fn: Callable,
) -> Request:
sampling_params = SamplingParams(max_tokens=17)
sampling_params.update_from_generation_config({}, eos_token_id=100)
return Request(
request_id=request_id,
prompt_token_ids=prompt_token_ids,
sampling_params=sampling_params,
pooling_params=None,
block_hasher=get_request_block_hasher(hash_block_size, hash_fn),
)
def boundary_hash(req: Request, hash_block_size: int, num_tokens: int) -> BlockHash:
# Every boundary at a hash_block_size multiple is just the fine-grained
# chain hash ending there.
return req.block_hashes[num_tokens // hash_block_size - 1]
def cache_full_block_and_partial_tail(
token_ids: list[int],
*,
enable_kv_cache_events: bool = False,
) -> tuple[BlockPool, Request, list[KVCacheBlock], BlockHash]:
hash_block_size = 2
block_size = 6
kv_cache_group_id = 0
req = make_request("0", token_ids, hash_block_size, sha256)
pool = BlockPool(
num_gpu_blocks=3,
enable_caching=True,
hash_block_size=hash_block_size,
enable_kv_cache_events=enable_kv_cache_events,
)
blocks = pool.get_new_blocks(2)
pool.cache_full_blocks(
request=req,
blocks=blocks,
num_cached_blocks=0,
num_full_blocks=1,
block_size=block_size,
kv_cache_group_id=kv_cache_group_id,
)
partial_hash = boundary_hash(req, hash_block_size, len(token_ids))
assert pool.cache_partial_block(
request=req,
block=blocks[1],
num_tokens=len(token_ids),
kv_cache_group_id=kv_cache_group_id,
block_size=block_size,
)
return pool, req, blocks, partial_hash
def test_boundary_hashes_reuse_fine_grained_chain():
hash_block_size = 2
block_size = 6
token_ids = [0, 0, 1, 1, 2, 2, 3, 3, 4, 4]
req = make_request("0", token_ids, hash_block_size, sha256)
coarse = BlockHashListWithBlockSize(req.block_hashes, hash_block_size, block_size)
# The block_size=6 full-block hash is the fine hash at the 6-token boundary,
# not a concatenation of the three fine hashes inside the block.
assert coarse[0] == req.block_hashes[6 // hash_block_size - 1]
assert coarse[0] != BlockHash(
req.block_hashes[0] + req.block_hashes[1] + req.block_hashes[2]
)
# A partial tail at 10 tokens is the fine hash at the 10-token boundary,
# which chains over the entire prefix.
tail_hash = boundary_hash(req, hash_block_size, 10)
assert tail_hash == req.block_hashes[4]
assert tail_hash == hash_block_tokens(sha256, req.block_hashes[3], token_ids[8:10])
def test_cache_partial_block_kv_cache_events():
hash_block_size = 4
block_size = 12
kv_cache_group_id = 2
pool = BlockPool(
num_gpu_blocks=2,
enable_caching=True,
hash_block_size=hash_block_size,
enable_kv_cache_events=True,
)
req = make_request(
"req_partial_events",
prompt_token_ids=list(range(hash_block_size * 2)),
hash_block_size=hash_block_size,
hash_fn=sha256,
)
block = pool.get_new_blocks(1)[0]
partial_entry_hash = pool.cache_partial_block(
request=req,
block=block,
num_tokens=hash_block_size * 2,
kv_cache_group_id=kv_cache_group_id,
block_size=block_size,
)
events = pool.take_events()
assert len(events) == 1
stored_event = events[0]
assert isinstance(stored_event, BlockStored)
assert partial_entry_hash is not None
assert stored_event.block_hashes == [
kv_cache_utils.maybe_convert_block_hash(req.block_hashes[1])
]
assert stored_event.parent_block_hash == kv_cache_utils.maybe_convert_block_hash(
req.block_hashes[0]
)
assert stored_event.token_ids == req.all_token_ids[hash_block_size:]
assert stored_event.block_size == 4
assert stored_event.group_idx == kv_cache_group_id
duplicate_entry_hash = pool.cache_partial_block(
request=req,
block=block,
num_tokens=hash_block_size * 2,
kv_cache_group_id=kv_cache_group_id,
block_size=block_size,
)
assert duplicate_entry_hash == partial_entry_hash
assert pool.take_events() == []
pool.free_blocks([block])
pool.get_new_blocks(1)
events = pool.take_events()
assert len(events) == 1
removed_event = events[0]
assert isinstance(removed_event, BlockRemoved)
assert removed_event.block_hashes == stored_event.block_hashes
assert removed_event.group_idx == kv_cache_group_id
def test_partial_block_replacement_emits_remove_then_store_events():
hash_block_size = 2
block_size = 6
kv_cache_group_id = 0
req = make_request("0", [0, 0, 1, 1, 2, 2, 3, 3], hash_block_size, sha256)
pool = BlockPool(
num_gpu_blocks=3,
enable_caching=True,
hash_block_size=hash_block_size,
enable_kv_cache_events=True,
)
blocks = pool.get_new_blocks(2)
pool.cache_full_blocks(
request=req,
blocks=blocks,
num_cached_blocks=0,
num_full_blocks=1,
block_size=block_size,
kv_cache_group_id=kv_cache_group_id,
)
partial_hash_8 = boundary_hash(req, hash_block_size, 8)
assert pool.cache_partial_block(
request=req,
block=blocks[1],
num_tokens=8,
kv_cache_group_id=kv_cache_group_id,
block_size=block_size,
)
assert pool.get_cached_block(partial_hash_8, [kv_cache_group_id]) == [blocks[1]]
pool.take_events()
req.append_output_token_ids([4, 4])
partial_hash_10 = boundary_hash(req, hash_block_size, 10)
assert pool.cache_partial_block(
request=req,
block=blocks[1],
num_tokens=10,
kv_cache_group_id=kv_cache_group_id,
block_size=block_size,
)
events = pool.take_events()
assert len(events) == 2
removed_event, stored_event = events
assert isinstance(removed_event, BlockRemoved)
assert removed_event.block_hashes == [
kv_cache_utils.maybe_convert_block_hash(partial_hash_8)
]
assert removed_event.group_idx == kv_cache_group_id
assert isinstance(stored_event, BlockStored)
assert stored_event.block_hashes == [
kv_cache_utils.maybe_convert_block_hash(partial_hash_10)
]
assert stored_event.parent_block_hash == kv_cache_utils.maybe_convert_block_hash(
boundary_hash(req, hash_block_size, 8)
)
assert stored_event.token_ids == req.all_token_ids[8:10]
assert stored_event.block_size == hash_block_size
assert stored_event.group_idx == kv_cache_group_id
assert pool.get_cached_block(partial_hash_8, [kv_cache_group_id]) is None
assert pool.get_cached_block(partial_hash_10, [kv_cache_group_id]) == [blocks[1]]
def test_later_request_hits_cached_partial_tail():
hash_block_size = 2
block_size = 6
kv_cache_group_id = 0
cached_token_ids = [0, 0, 1, 1, 2, 2, 3, 3, 4, 4]
req = make_request("0", cached_token_ids, hash_block_size, sha256)
pool = BlockPool(
num_gpu_blocks=3,
enable_caching=True,
hash_block_size=hash_block_size,
)
blocks = pool.get_new_blocks(2)
pool.cache_full_blocks(
request=req,
blocks=blocks,
num_cached_blocks=0,
num_full_blocks=1,
block_size=block_size,
kv_cache_group_id=kv_cache_group_id,
)
partial_hash_10 = boundary_hash(req, hash_block_size, 10)
assert pool.cache_partial_block(
request=req,
block=blocks[1],
num_tokens=10,
kv_cache_group_id=kv_cache_group_id,
block_size=block_size,
)
replay = make_request("1", cached_token_ids, hash_block_size, sha256)
replay_hash_10 = boundary_hash(replay, hash_block_size, 10)
assert replay_hash_10 == partial_hash_10
assert pool.get_cached_block(replay_hash_10, [kv_cache_group_id]) == [blocks[1]]
extended = make_request("2", cached_token_ids + [10], hash_block_size, sha256)
extended_hash_10 = boundary_hash(extended, hash_block_size, 10)
assert extended_hash_10 == partial_hash_10
assert pool.get_cached_block(extended_hash_10, [kv_cache_group_id]) == [blocks[1]]
def test_cache_partial_block_uses_fine_grained_boundary_hash():
hash_block_size = 2
block_size = 6
kv_cache_group_id = 0
token_ids = [0, 0, 1, 1, 2, 2, 3, 3, 4, 4]
req = make_request("0", token_ids, hash_block_size, sha256)
pool = BlockPool(
num_gpu_blocks=3,
enable_caching=True,
hash_block_size=hash_block_size,
)
blocks = pool.get_new_blocks(2)
pool.cache_full_blocks(
request=req,
blocks=blocks,
num_cached_blocks=0,
num_full_blocks=1,
block_size=block_size,
kv_cache_group_id=kv_cache_group_id,
)
partial_entry_hash = pool.cache_partial_block(
request=req,
block=blocks[1],
num_tokens=10,
kv_cache_group_id=kv_cache_group_id,
block_size=block_size,
)
# The partial entry is keyed by the fine-grained hash at the 10-token
# boundary, regardless of the owning group's block_size.
expected = boundary_hash(req, hash_block_size, 10)
assert partial_entry_hash == kv_cache_utils.make_block_hash_with_group_id(
expected, kv_cache_group_id
)
assert pool.get_cached_block(expected, [kv_cache_group_id]) == [blocks[1]]
def test_cache_partial_block_requires_hash_boundary():
hash_block_size = 2
block_size = 4
req = make_request("0", [0, 0, 1, 1], hash_block_size, sha256)
pool = BlockPool(
num_gpu_blocks=2,
enable_caching=True,
hash_block_size=hash_block_size,
)
block = pool.get_new_blocks(1)[0]
with pytest.raises(AssertionError):
pool.cache_partial_block(
request=req,
block=block,
num_tokens=3,
kv_cache_group_id=0,
block_size=block_size,
)
def test_cache_partial_block_duplicate_checks_all_blocks_for_hash():
hash_block_size = 2
block_size = 4
kv_cache_group_id = 0
req = make_request("0", [0, 0, 1, 1], hash_block_size, sha256)
pool = BlockPool(
num_gpu_blocks=4,
enable_caching=True,
hash_block_size=hash_block_size,
)
blocks = pool.get_new_blocks(2)
first_entry_hash = pool.cache_partial_block(
request=req,
block=blocks[0],
num_tokens=2,
kv_cache_group_id=kv_cache_group_id,
block_size=block_size,
)
second_entry_hash = pool.cache_partial_block(
request=req,
block=blocks[1],
num_tokens=2,
kv_cache_group_id=kv_cache_group_id,
block_size=block_size,
)
assert first_entry_hash == second_entry_hash
duplicate_entry_hash = pool.cache_partial_block(
request=req,
block=blocks[1],
num_tokens=2,
kv_cache_group_id=kv_cache_group_id,
block_size=block_size,
)
assert duplicate_entry_hash == second_entry_hash
assert pool.cached_block_hashes_by_block == {}
def test_reset_prefix_cache_clears_partial_entry_metadata():
pool, req, blocks, partial_hash_10 = cache_full_block_and_partial_tail(
[0, 0, 1, 1, 2, 2, 3, 3, 4, 4]
)
full_hash = BlockHashListWithBlockSize(req.block_hashes, 2, 6)[0]
assert pool.get_cached_block(full_hash, [0]) == [blocks[0]]
assert pool.get_cached_block(partial_hash_10, [0]) == [blocks[1]]
pool.free_blocks(blocks)
assert pool.reset_prefix_cache()
assert pool.get_cached_block(full_hash, [0]) is None
assert pool.get_cached_block(partial_hash_10, [0]) is None
assert pool.cached_block_hashes_by_block == {}
def test_evict_cached_block_removes_full_hash_and_partial_entry():
pool, req, blocks, partial_hash_10 = cache_full_block_and_partial_tail(
[0, 0, 1, 1, 2, 2, 3, 3, 4, 4]
)
full_hash = BlockHashListWithBlockSize(req.block_hashes, 2, 6)[0]
assert pool.get_cached_block(full_hash, [0]) == [blocks[0]]
assert pool.get_cached_block(partial_hash_10, [0]) == [blocks[1]]
pool.evict_blocks({blocks[0].block_id, blocks[1].block_id})
assert pool.get_cached_block(full_hash, [0]) is None
assert pool.get_cached_block(partial_hash_10, [0]) is None
assert pool.cached_block_hashes_by_block == {}
def test_partial_block_promotes_to_direct_full_block_hash():
hash_block_size = 2
block_size = 6
kv_cache_group_id = 0
token_ids = [0, 0, 1, 1, 2, 2, 3, 3, 4, 4]
req = make_request("0", token_ids, hash_block_size, sha256)
pool = BlockPool(
num_gpu_blocks=3,
enable_caching=True,
hash_block_size=hash_block_size,
)
blocks = pool.get_new_blocks(2)
pool.cache_full_blocks(
request=req,
blocks=blocks,
num_cached_blocks=0,
num_full_blocks=1,
block_size=block_size,
kv_cache_group_id=kv_cache_group_id,
)
partial_hash_10 = boundary_hash(req, hash_block_size, 10)
assert pool.cache_partial_block(
request=req,
block=blocks[1],
num_tokens=10,
kv_cache_group_id=kv_cache_group_id,
block_size=block_size,
)
assert pool.get_cached_block(partial_hash_10, [kv_cache_group_id]) == [blocks[1]]
req.append_output_token_ids([5, 5])
full_hashes = BlockHashListWithBlockSize(
req.block_hashes, hash_block_size, block_size
)
promoted_full_hash = full_hashes[1]
# The promoted full-block hash is the fine hash at the 12-token boundary,
# not a concatenation of the fine hashes inside the block.
assert promoted_full_hash == req.block_hashes[12 // hash_block_size - 1]
assert promoted_full_hash != BlockHash(
req.block_hashes[3] + req.block_hashes[4] + req.block_hashes[5]
)
pool.cache_full_blocks(
request=req,
blocks=blocks,
num_cached_blocks=1,
num_full_blocks=2,
block_size=block_size,
kv_cache_group_id=kv_cache_group_id,
)
assert pool.get_cached_block(promoted_full_hash, [kv_cache_group_id]) == [blocks[1]]
assert pool.get_cached_block(partial_hash_10, [kv_cache_group_id]) is None
+103 -8
View File
@@ -117,6 +117,7 @@ def new_kv_cache_spec(
page_size_padded=None,
sliding_window=None,
attention_chunk_size=None,
indexes_kv_by_block_stride=False,
):
return FullAttentionSpec(
block_size=block_size,
@@ -126,6 +127,7 @@ def new_kv_cache_spec(
page_size_padded=page_size_padded,
sliding_window=sliding_window,
attention_chunk_size=attention_chunk_size,
indexes_kv_by_block_stride=indexes_kv_by_block_stride,
)
@@ -136,6 +138,7 @@ def new_sliding_window_spec(
dtype=torch.float32,
page_size_padded=None,
sliding_window=1,
indexes_kv_by_block_stride=False,
):
return SlidingWindowSpec(
block_size=block_size,
@@ -144,6 +147,7 @@ def new_sliding_window_spec(
dtype=dtype,
page_size_padded=page_size_padded,
sliding_window=sliding_window,
indexes_kv_by_block_stride=indexes_kv_by_block_stride,
)
@@ -221,7 +225,7 @@ def test_kv_cache_block():
# Test block hash setting and resetting
block_hash = make_block_hash_with_group_id(BlockHash(b"abc"), 0)
block.block_hash = block_hash
block.set_block_hash(block_hash)
assert block.block_hash == block_hash
block.reset_hash()
@@ -1799,16 +1803,38 @@ def test_get_kv_cache_config_one_worker():
],
)
# different hidden size that cannot be aligned by using different block size
# different hidden size that cannot be aligned by using different block size,
# but can be aligned by padding the smaller physical page.
swa_spec = new_sliding_window_spec(head_size=96, indexes_kv_by_block_stride=True)
kv_cache_specs_hybrid = {
"layer_1": new_kv_cache_spec(head_size=64),
"layer_2": new_sliding_window_spec(head_size=96),
"layer_1": new_kv_cache_spec(head_size=64, indexes_kv_by_block_stride=True),
"layer_2": swa_spec,
}
with pytest.raises(NotImplementedError):
get_kv_cache_configs(
vllm_config, [kv_cache_specs_hybrid], [mem_per_block_per_layer * 2 * 32]
)[0]
kv_cache_config_hybrid = get_kv_cache_configs(
vllm_config, [kv_cache_specs_hybrid], [mem_per_block_per_layer * 2 * 32]
)[0]
padded_page_size = swa_spec.page_size_bytes
assert kv_cache_config_hybrid == KVCacheConfig(
num_blocks=42,
kv_cache_tensors=[
KVCacheTensor(size=padded_page_size * 42, shared_by=["layer_1", "layer_2"]),
],
kv_cache_groups=[
KVCacheGroupSpec(
["layer_1"],
new_kv_cache_spec(
head_size=64,
page_size_padded=padded_page_size,
indexes_kv_by_block_stride=True,
),
),
KVCacheGroupSpec(
["layer_2"],
new_sliding_window_spec(head_size=96, indexes_kv_by_block_stride=True),
),
],
)
# Test num_gpu_blocks_override
vllm_config.cache_config.num_gpu_blocks_override = 16
@@ -2322,6 +2348,75 @@ def test_check_enough_kv_cache_memory_respects_num_gpu_blocks_override():
get_kv_cache_configs(vllm_config, [kv_cache_specs], [large_available_memory])
def test_unify_kv_cache_page_size_uses_padding_for_non_divisible_sizes():
"""DFlash drafters can have a smaller head size than the target model.
For example, MiMo uses 192-dim target KV heads while its DFlash draft uses
128-dim KV heads. The resulting page sizes are 3:2 rather than an integer
block-size multiple, so the smaller page must be padded instead.
"""
# Both layers' backends opt into the padded-page strided view (e.g.
# FlashAttention / its DiffKV subclass), so padding is allowed.
target_spec = new_kv_cache_spec(
block_size=16,
num_kv_heads=1,
head_size=192,
dtype=torch.bfloat16,
indexes_kv_by_block_stride=True,
)
draft_spec = new_sliding_window_spec(
block_size=16,
num_kv_heads=1,
head_size=128,
dtype=torch.bfloat16,
sliding_window=1024,
indexes_kv_by_block_stride=True,
)
unified_specs = kv_cache_utils.unify_kv_cache_spec_page_size(
{
"target_attn": target_spec,
"draft_attn": draft_spec,
}
)
assert unified_specs["target_attn"] == target_spec
unified_draft_spec = unified_specs["draft_attn"]
assert unified_draft_spec.block_size == draft_spec.block_size
assert unified_draft_spec.real_page_size_bytes == draft_spec.real_page_size_bytes
assert unified_draft_spec.page_size_padded == target_spec.page_size_bytes
assert unified_draft_spec.page_size_bytes == target_spec.page_size_bytes
def test_unify_kv_cache_page_size_padding_requires_backend_support():
"""Padding is gated on the backend declaring ``indexes_kv_by_block_stride``.
A backend that does not support the strided padded-page view must raise
rather than silently padding (and misreading KV at runtime).
"""
target_spec = new_kv_cache_spec(
block_size=16,
num_kv_heads=1,
head_size=192,
dtype=torch.bfloat16,
indexes_kv_by_block_stride=True,
)
# The non-divisible draft layer needs padding but its backend does not
# support the strided padded-page view -> must raise, not silently pad.
draft_spec = new_sliding_window_spec(
block_size=16,
num_kv_heads=1,
head_size=128,
dtype=torch.bfloat16,
sliding_window=1024,
indexes_kv_by_block_stride=False,
)
specs = {"target_attn": target_spec, "draft_attn": draft_spec}
with pytest.raises(NotImplementedError):
kv_cache_utils.unify_kv_cache_spec_page_size(specs)
def test_unify_hybrid_kv_cache_specs():
# 1. has_full_attention and has_sliding_window
before_spec_1 = new_kv_cache_spec()
+1 -1
View File
@@ -2003,7 +2003,7 @@ def test_maybe_evict_cached_block():
assert len(pool.blocks) == len(block_hashes)
# Manually add all blocks to cached_blocks
for block, block_hash in zip(pool.blocks, block_hashes):
block.block_hash = block_hash
block.set_block_hash(block_hash)
pool.cached_block_hash_to_block.insert(block_hash, block)
block0, block1, block2, block3 = pool.blocks
+1 -1
View File
@@ -425,7 +425,7 @@ def _run_eagle_correctness(
if "deepseek" in model_setup[1].lower():
m.setenv("VLLM_ROCM_USE_AITER", "1")
m.delenv("VLLM_MLA_DISABLE", raising=False)
attention_config = {"backend": "TRITON_MLA"}
attention_config = {"backend": "ROCM_AITER_MLA"}
else:
m.setenv("VLLM_ROCM_USE_AITER", "1")
@@ -0,0 +1,355 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
from unittest.mock import MagicMock
import pytest
import torch
from tests.v1.kv_connector.unit.utils import create_vllm_config
from vllm.config import KVEventsConfig, KVTransferConfig
from vllm.distributed.kv_events import BlockRemoved, BlockStored
from vllm.distributed.kv_transfer.kv_connector.v1.offloading.events import (
OffloadingEventGroupSpec,
OffloadingEventsTracker,
)
from vllm.distributed.kv_transfer.kv_connector.v1.offloading.scheduler import (
GroupOffloadConfig,
)
from vllm.v1.core.kv_cache_utils import BlockHash, maybe_convert_block_hash
from vllm.v1.kv_cache_interface import (
FullAttentionSpec,
KVCacheConfig,
KVCacheGroupSpec,
KVCacheSpecKind,
)
from vllm.v1.kv_offload.base import (
OffloadingEvent,
OffloadingKVEventsConfig,
OffloadKey,
make_offload_key,
)
from vllm.v1.kv_offload.cpu.common import CPULoadStoreSpec
from vllm.v1.kv_offload.tiering.spec import TieringOffloadingSpec
_CPU_MEDIUM = CPULoadStoreSpec.medium()
_FULL_ATTENTION_EVENT_SPEC = OffloadingEventGroupSpec(
kv_cache_spec_kind=KVCacheSpecKind.FULL_ATTENTION.value,
kv_cache_spec_sliding_window=None,
)
def _tracker(
*,
enable_kv_cache_events: bool = True,
self_describing_kv_events: bool = True,
) -> OffloadingEventsTracker:
return OffloadingEventsTracker(
OffloadingKVEventsConfig(
enable_kv_cache_events=enable_kv_cache_events,
self_describing_kv_events=self_describing_kv_events,
)
)
def _hash(i: int) -> BlockHash:
return BlockHash(str(i).encode())
def _wire_hash(block_hash: BlockHash):
return maybe_convert_block_hash(block_hash)
def _request(*, block_hashes: list[BlockHash], token_count: int):
req = MagicMock()
req.block_hashes = block_hashes
req.all_token_ids = list(range(1, token_count + 1))
req.lora_request = None
return req
def _group_config(
*,
group_idx: int = 0,
block_size: int = 4,
block_size_factor: int = 1,
sliding_window_size_in_blocks: int | None = None,
) -> GroupOffloadConfig:
return GroupOffloadConfig(
group_idx=group_idx,
gpu_block_size=block_size,
offloaded_block_size=block_size * block_size_factor,
hash_block_size_factor=block_size_factor,
sliding_window_size_in_blocks=sliding_window_size_in_blocks,
kv_event_group_spec=_FULL_ATTENTION_EVENT_SPEC,
)
def _record_chunks(
tracker: OffloadingEventsTracker,
req,
group_config: GroupOffloadConfig,
num_chunks: int,
) -> list[OffloadKey]:
keys: list[OffloadKey] = []
hbf = group_config.hash_block_size_factor
for chunk_idx in range(num_chunks):
tail_hash = req.block_hashes[(chunk_idx + 1) * hbf - 1]
assert tail_hash is not None
key = make_offload_key(tail_hash, group_config.group_idx)
tracker.record_store(req, group_config, chunk_idx, key)
keys.append(key)
return keys
def _stored_event(keys: list[OffloadKey]) -> OffloadingEvent:
return OffloadingEvent(keys=keys, medium=_CPU_MEDIUM, removed=False)
def _removed_event(keys: list[OffloadKey]) -> OffloadingEvent:
return OffloadingEvent(keys=keys, medium=_CPU_MEDIUM, removed=True)
def test_take_events_publishes_routable_block_stored():
block_size = 4
tracker = _tracker()
group_config = _group_config(block_size=block_size)
req = _request(
block_hashes=[_hash(i) for i in range(6)],
token_count=block_size * 6,
)
keys = _record_chunks(tracker, req, group_config, num_chunks=6)
batch1 = list(tracker.take_events([_stored_event(keys[:3])]))
assert len(batch1) == 3
for i, event in enumerate(batch1):
assert isinstance(event, BlockStored)
assert event.medium == _CPU_MEDIUM
assert event.block_hashes == [_wire_hash(_hash(i))]
assert event.block_size == block_size
assert event.token_ids == list(
range(i * block_size + 1, (i + 1) * block_size + 1)
)
if i == 0:
assert event.parent_block_hash is None
else:
assert event.parent_block_hash == _wire_hash(_hash(i - 1))
assert event.lora_id is None
assert event.lora_name is None
assert event.extra_keys is None
assert event.group_idx == 0
assert event.kv_cache_spec_kind == KVCacheSpecKind.FULL_ATTENTION.value
assert event.kv_cache_spec_sliding_window is None
batch2 = list(tracker.take_events([_stored_event(keys[3:])]))
assert len(batch2) == 3
assert batch2[0].parent_block_hash == batch1[-1].block_hashes[-1]
assert len(tracker._pending_event_metadata) == 6
def test_take_events_factor_gt_1_chunk_store_and_remove():
block_size = 4
block_size_factor = 3
tracker = _tracker()
group_config = _group_config(
block_size=block_size, block_size_factor=block_size_factor
)
req = _request(
block_hashes=[_hash(i) for i in range(6)],
token_count=block_size * block_size_factor * 2,
)
keys = _record_chunks(tracker, req, group_config, num_chunks=2)
stored = list(tracker.take_events([_stored_event(keys)]))
assert len(stored) == 2
expected_hashes = []
for chunk_idx, event in enumerate(stored):
assert isinstance(event, BlockStored)
expected_chunk_hashes = [
_wire_hash(_hash(i))
for i in range(
chunk_idx * block_size_factor,
(chunk_idx + 1) * block_size_factor,
)
]
assert event.block_hashes == expected_chunk_hashes
assert event.block_size == block_size
assert len(event.token_ids) == block_size * block_size_factor
if chunk_idx == 0:
assert event.parent_block_hash is None
else:
assert event.parent_block_hash == _wire_hash(_hash(block_size_factor - 1))
expected_hashes.extend(expected_chunk_hashes)
assert len(tracker._pending_event_metadata) == 2
removed = list(tracker.take_events([_removed_event(keys)]))
assert len(removed) == 1
assert isinstance(removed[0], BlockRemoved)
assert removed[0].block_hashes == expected_hashes
assert removed[0].medium == _CPU_MEDIUM
assert removed[0].group_idx == 0
assert not tracker._pending_event_metadata
def test_take_events_factor_gt_1_store_is_order_independent():
block_size_factor = 3
tracker = _tracker()
group_config = _group_config(block_size_factor=block_size_factor)
req = _request(
block_hashes=[_hash(i) for i in range(6)],
token_count=4 * block_size_factor * 2,
)
keys = _record_chunks(tracker, req, group_config, num_chunks=2)
unknown_key = make_offload_key(_hash(12345), 0)
events = list(tracker.take_events([_stored_event([keys[1], unknown_key, keys[0]])]))
assert len(events) == 3
chunk1, placeholder, chunk0 = events
assert [len(event.block_hashes) for event in events] == [3, 1, 3]
assert placeholder.block_size == 0
assert placeholder.token_ids == []
assert chunk0.parent_block_hash is None
assert chunk1.parent_block_hash == chunk0.block_hashes[-1]
def test_take_events_opt_out_keeps_placeholders():
tracker = _tracker(self_describing_kv_events=False)
group_config = _group_config()
req = _request(block_hashes=[_hash(i) for i in range(3)], token_count=12)
keys = _record_chunks(tracker, req, group_config, num_chunks=3)
assert not tracker.self_describing_enabled
assert not tracker._pending_event_metadata
events = list(
tracker.take_events(
[
_stored_event(keys),
_removed_event(keys),
]
)
)
assert len(events) == 4
for event in events[:3]:
assert isinstance(event, BlockStored)
assert event.block_size == 0
assert event.token_ids == []
assert event.parent_block_hash is None
assert isinstance(events[3], BlockRemoved)
assert len(events[3].block_hashes) == 3
def test_record_store_skips_sliding_window_group():
tracker = _tracker()
group_config = _group_config(sliding_window_size_in_blocks=2)
req = _request(block_hashes=[_hash(i) for i in range(3)], token_count=12)
keys = _record_chunks(tracker, req, group_config, num_chunks=3)
assert not tracker._pending_event_metadata
events = list(tracker.take_events([_stored_event(keys[:1])]))
assert len(events) == 1
assert isinstance(events[0], BlockStored)
assert events[0].block_size == 0
def test_take_events_groups_removed_hashes_by_kv_group():
tracker = _tracker()
group0_config = _group_config(group_idx=0, block_size_factor=2)
group1_config = _group_config(group_idx=1, block_size_factor=2)
req0 = _request(block_hashes=[_hash(0), _hash(1)], token_count=8)
req1 = _request(block_hashes=[_hash(10), _hash(11)], token_count=8)
key0 = _record_chunks(tracker, req0, group0_config, num_chunks=1)[0]
key1 = _record_chunks(tracker, req1, group1_config, num_chunks=1)[0]
removed = list(tracker.take_events([_removed_event([key0, key1])]))
assert len(removed) == 2
by_group = {event.group_idx: event.block_hashes for event in removed}
assert by_group == {
0: [_wire_hash(_hash(0)), _wire_hash(_hash(1))],
1: [_wire_hash(_hash(10)), _wire_hash(_hash(11))],
}
def test_take_events_supports_restore_after_eviction():
block_size = 4
tracker = _tracker()
group_config = _group_config(block_size=block_size)
req = _request(block_hashes=[_hash(0)], token_count=block_size)
key = _record_chunks(tracker, req, group_config, num_chunks=1)[0]
first_store = list(tracker.take_events([_stored_event([key])]))
assert len(first_store) == 1
assert isinstance(first_store[0], BlockStored)
assert first_store[0].token_ids == [1, 2, 3, 4]
removed = list(tracker.take_events([_removed_event([key])]))
assert len(removed) == 1
assert isinstance(removed[0], BlockRemoved)
assert not tracker._pending_event_metadata
req.all_token_ids = [5, 6, 7, 8]
tracker.record_store(req, group_config, offload_block_idx=0, offload_key=key)
second_store = list(tracker.take_events([_stored_event([key])]))
assert len(second_store) == 1
assert isinstance(second_store[0], BlockStored)
assert second_store[0].token_ids == [5, 6, 7, 8]
def test_reset_cache_clears_side_table():
tracker = _tracker()
group_config = _group_config()
req = _request(block_hashes=[_hash(i) for i in range(3)], token_count=12)
_record_chunks(tracker, req, group_config, num_chunks=3)
assert tracker._pending_event_metadata
tracker.reset()
assert not tracker._pending_event_metadata
def test_tiering_rejects_self_describing_kv_events():
vllm_config = create_vllm_config(
block_size=4,
max_num_batched_tokens=16,
disable_hybrid_kv_cache_manager=False,
)
vllm_config.kv_transfer_config = KVTransferConfig(
kv_connector="OffloadingConnector",
kv_role="kv_both",
kv_connector_extra_config={
"spec_name": "TieringOffloadingSpec",
"cpu_bytes_to_use": 1 << 20,
"self_describing_kv_events": True,
"secondary_tiers": [{"type": "example"}],
},
)
vllm_config.kv_events_config = KVEventsConfig(
enable_kv_cache_events=True,
publisher="null",
)
kv_cache_config = KVCacheConfig(
num_blocks=0,
kv_cache_tensors=[],
kv_cache_groups=[
KVCacheGroupSpec(
["layer"],
FullAttentionSpec(
block_size=4,
num_kv_heads=1,
head_size=1,
dtype=torch.float32,
),
)
],
)
with pytest.raises(ValueError, match="TieringOffloadingSpec"):
TieringOffloadingSpec(vllm_config, kv_cache_config)
@@ -22,7 +22,7 @@ from vllm.v1.kv_offload.base import (
OffloadingGaugeMetadata,
OffloadingHistogramMetadata,
)
from vllm.v1.kv_offload.cpu.spec import CPUOffloadingSpec
from vllm.v1.kv_offload.factory import OffloadingSpecFactory
LOAD_BYTES = _TransferMetricName.LOAD_BYTES
LOAD_TIME = _TransferMetricName.LOAD_TIME
@@ -33,6 +33,8 @@ STORE_SIZE = _TransferMetricName.STORE_SIZE
STORES_SKIPPED = "vllm:kv_offload_stores_skipped"
PENDING_STORES = "vllm:kv_offload_pending_stores"
LOOKUP_LATENCY = "vllm:kv_offload_lookup_latency_seconds"
MY_COUNTER = "my_counter"
MY_LABEL = "my_label"
class _FakeMetric:
@@ -67,6 +69,20 @@ class _FakeVllmConfig:
)
def _spec_cls_with_metric_definitions(
metric_definitions: dict[str, Any],
) -> type:
"""Build a fake offloading spec class reporting the given metric
definitions, so tests don't need to patch the real CPU spec."""
class _FakeOffloadingSpec:
@staticmethod
def build_metric_definitions(extra_config):
return metric_definitions
return _FakeOffloadingSpec
def _metric_metadata():
return {
LOAD_BYTES: OffloadingCounterMetadata(
@@ -96,9 +112,17 @@ def _metric_metadata():
LOOKUP_LATENCY: OffloadingHistogramMetadata(
documentation="lookup latency",
),
MY_COUNTER: OffloadingCounterMetadata(
documentation="counter with a label",
labelnames=(MY_LABEL,),
),
}
def _unlabeled(values: dict[str, Any], metric_name: str) -> Any:
return values[metric_name][()]
def test_build_kv_connector_stats_with_none():
"""Test that build_kv_connector_stats returns empty stats when given None."""
stats = OffloadingConnector.build_kv_connector_stats(data=None)
@@ -131,13 +155,13 @@ def test_build_kv_connector_stats_reconstructs_offload_stats():
STORES_SKIPPED: _MetricType.COUNTER,
},
_StatsKey.DATA: {
LOAD_BYTES: 24,
LOAD_TIME: 1.5,
LOAD_SIZE: [16, 8],
STORE_BYTES: 3,
STORE_TIME: 0.3,
STORE_SIZE: [1, 2],
STORES_SKIPPED: 5,
LOAD_BYTES: {(): 24},
LOAD_TIME: {(): 1.5},
LOAD_SIZE: {(): [16, 8]},
STORE_BYTES: {(): 3},
STORE_TIME: {(): 0.3},
STORE_SIZE: {(): [1, 2]},
STORES_SKIPPED: {(): 5},
},
}
@@ -145,22 +169,28 @@ def test_build_kv_connector_stats_reconstructs_offload_stats():
assert isinstance(stats, OffloadingConnectorStats)
values = stats.data[_StatsKey.DATA]
assert values[LOAD_BYTES] == 24
assert values[LOAD_TIME] == 1.5
assert values[LOAD_SIZE] == [16, 8]
assert values[STORE_BYTES] == 3
assert values[STORE_TIME] == 0.3
assert values[STORE_SIZE] == [1, 2]
assert values[STORES_SKIPPED] == 5
assert _unlabeled(values, LOAD_BYTES) == 24
assert _unlabeled(values, LOAD_TIME) == 1.5
assert _unlabeled(values, LOAD_SIZE) == [16, 8]
assert _unlabeled(values, STORE_BYTES) == 3
assert _unlabeled(values, STORE_TIME) == 0.3
assert _unlabeled(values, STORE_SIZE) == [1, 2]
assert _unlabeled(values, STORES_SKIPPED) == 5
def _make_stats_data(
metric_data: dict[str, Any],
metric_metadata: dict[str, Any],
) -> dict[str, Any]:
"""Build a structured data dict from flat metric data and metadata."""
"""Build a structured data dict from flat metric data and metadata.
Values for unlabeled metrics may be passed flat (wrapped here under the
empty label tuple); values for labeled metrics must already be passed as
a ``{labelvalues: value}`` map.
"""
metric_types = {}
for key in metric_data:
data = {}
for key, value in metric_data.items():
md = metric_metadata[key]
if isinstance(md, OffloadingCounterMetadata):
metric_types[key] = _MetricType.COUNTER
@@ -168,9 +198,10 @@ def _make_stats_data(
metric_types[key] = _MetricType.GAUGE
elif isinstance(md, OffloadingHistogramMetadata):
metric_types[key] = _MetricType.HISTOGRAM
data[key] = value if md.labelnames else {(): value}
return {
_StatsKey.TYPES: metric_types,
_StatsKey.DATA: metric_data,
_StatsKey.DATA: data,
}
@@ -215,34 +246,106 @@ def test_aggregate_same_connector():
assert result is stats1 # Should return self
values = result.data[_StatsKey.DATA]
assert values[LOAD_BYTES] == 34
assert values[LOAD_TIME] == 2.6
assert values[LOAD_SIZE] == [16, 8, 3, 7]
assert values[STORE_BYTES] == 19
assert values[STORE_TIME] == 2.3
assert values[STORE_SIZE] == [1, 2, 16]
assert values[STORES_SKIPPED] == 4
assert values[PENDING_STORES] == 1
assert values[LOOKUP_LATENCY] == [0.1, 0.2, 0.3]
assert _unlabeled(values, LOAD_BYTES) == 34
assert _unlabeled(values, LOAD_TIME) == 2.6
assert _unlabeled(values, LOAD_SIZE) == [16, 8, 3, 7]
assert _unlabeled(values, STORE_BYTES) == 19
assert _unlabeled(values, STORE_TIME) == 2.3
assert _unlabeled(values, STORE_SIZE) == [1, 2, 16]
assert _unlabeled(values, STORES_SKIPPED) == 4
assert _unlabeled(values, PENDING_STORES) == 1
assert _unlabeled(values, LOOKUP_LATENCY) == [0.1, 0.2, 0.3]
def test_aggregate_labeled_metrics():
metadata = _metric_metadata()
stats1 = OffloadingConnectorStats(
data=_make_stats_data(
{
MY_COUNTER: {
("a",): 10,
("b",): 3,
},
},
metadata,
),
)
stats2 = OffloadingConnectorStats(
data=_make_stats_data(
{
MY_COUNTER: {
("a",): 7,
("c",): 5,
},
},
metadata,
),
)
stats1.aggregate(stats2)
values = stats1.data[_StatsKey.DATA][MY_COUNTER]
assert values[("a",)] == 17
assert values[("b",)] == 3
assert values[("c",)] == 5
def test_aggregate_labeled_metric_missing_from_self():
"""Aggregating a labeled metric that self doesn't have at all yet."""
metadata = _metric_metadata()
stats1 = OffloadingConnectorStats()
stats2 = OffloadingConnectorStats(
data=_make_stats_data(
{
MY_COUNTER: {
("a",): 7,
("b",): 5,
},
},
metadata,
),
)
stats1.aggregate(stats2)
values = stats1.data[_StatsKey.DATA][MY_COUNTER]
assert values[("a",)] == 7
assert values[("b",)] == 5
assert stats1.data[_StatsKey.TYPES][MY_COUNTER] == _MetricType.COUNTER
def test_helper_methods_accept_labeled_metrics():
stats = OffloadingConnectorStats()
stats.increase_counter(MY_COUNTER, 3, ("a",))
stats.increase_counter(MY_COUNTER, 4, ("a",))
stats.set_gauge(PENDING_STORES, 2, ("b",))
stats.observe_histogram(LOOKUP_LATENCY, 0.1, ("b",))
stats.observe_histogram(LOOKUP_LATENCY, 0.2, ("b",))
values = stats.data[_StatsKey.DATA]
assert values[MY_COUNTER][("a",)] == 7
assert values[PENDING_STORES][("b",)] == 2
assert values[LOOKUP_LATENCY][("b",)] == [0.1, 0.2]
def test_aggregate_merges_types():
stats1 = OffloadingConnectorStats(
data={
_StatsKey.TYPES: {LOAD_BYTES: _MetricType.COUNTER},
_StatsKey.DATA: {LOAD_BYTES: 1},
_StatsKey.DATA: {LOAD_BYTES: {(): 1}},
},
)
stats2 = OffloadingConnectorStats(
data={
_StatsKey.TYPES: {PENDING_STORES: _MetricType.GAUGE},
_StatsKey.DATA: {PENDING_STORES: 2},
_StatsKey.DATA: {PENDING_STORES: {(): 2}},
},
)
result = stats1.aggregate(stats2)
assert result.data[_StatsKey.DATA][PENDING_STORES] == 2
assert _unlabeled(result.data[_StatsKey.DATA], PENDING_STORES) == 2
assert result.data[_StatsKey.TYPES][PENDING_STORES] == _MetricType.GAUGE
@@ -283,6 +386,26 @@ def test_reduce():
assert reduced[f"{LOOKUP_LATENCY}_sum"] == sum([0.1, 0.2, 0.3])
def test_reduce_labeled_metrics():
metadata = _metric_metadata()
stats = OffloadingConnectorStats(
data=_make_stats_data(
{
MY_COUNTER: {
("a",): 17,
("b",): 3,
},
},
metadata,
),
)
reduced = stats.reduce()
assert reduced[f"{MY_COUNTER}:{('a',)}"] == 17
assert reduced[f"{MY_COUNTER}:{('b',)}"] == 3
def test_reset():
"""Test that reset() resets all connector stats."""
metadata = _metric_metadata()
@@ -326,11 +449,11 @@ def test_prom_metrics_observes_manager_counter():
prom_metrics.observe(
{
_StatsKey.TYPES: {STORES_SKIPPED: _MetricType.COUNTER},
_StatsKey.DATA: {STORES_SKIPPED: 7},
_StatsKey.DATA: {STORES_SKIPPED: {(): 7}},
}
)
counter = prom_metrics.offloading_metrics[(0, STORES_SKIPPED)]
counter = prom_metrics.offloading_metrics[(0, STORES_SKIPPED, ())]
assert counter.increments == [7]
counter_def = prom_metrics._offloading_metric_defs[STORES_SKIPPED]
assert counter_def.kwargs["name"] == "vllm:kv_offload_stores_skipped"
@@ -360,22 +483,22 @@ def test_prom_metrics_observes_flat_transfer_metrics_and_legacy_metrics():
STORE_SIZE: _MetricType.HISTOGRAM,
},
_StatsKey.DATA: {
LOAD_BYTES: 24,
LOAD_TIME: 1.5,
LOAD_SIZE: [16, 8],
STORE_BYTES: 3,
STORE_TIME: 0.3,
STORE_SIZE: [1, 2],
LOAD_BYTES: {(): 24},
LOAD_TIME: {(): 1.5},
LOAD_SIZE: {(): [16, 8]},
STORE_BYTES: {(): 3},
STORE_TIME: {(): 0.3},
STORE_SIZE: {(): [1, 2]},
},
}
)
assert prom_metrics.offloading_metrics[(0, LOAD_BYTES)].increments == [24]
assert prom_metrics.offloading_metrics[(0, LOAD_TIME)].increments == [1.5]
assert prom_metrics.offloading_metrics[(0, LOAD_SIZE)].observed == [16, 8]
assert prom_metrics.offloading_metrics[(0, STORE_BYTES)].increments == [3]
assert prom_metrics.offloading_metrics[(0, STORE_TIME)].increments == [0.3]
assert prom_metrics.offloading_metrics[(0, STORE_SIZE)].observed == [1, 2]
assert prom_metrics.offloading_metrics[(0, LOAD_BYTES, ())].increments == [24]
assert prom_metrics.offloading_metrics[(0, LOAD_TIME, ())].increments == [1.5]
assert prom_metrics.offloading_metrics[(0, LOAD_SIZE, ())].observed == [16, 8]
assert prom_metrics.offloading_metrics[(0, STORE_BYTES, ())].increments == [3]
assert prom_metrics.offloading_metrics[(0, STORE_TIME, ())].increments == [0.3]
assert prom_metrics.offloading_metrics[(0, STORE_SIZE, ())].observed == [1, 2]
assert prom_metrics.counter_kv_bytes[(0, "CPU_to_GPU")].increments == [24]
assert prom_metrics.counter_kv_transfer_time[(0, "CPU_to_GPU")].increments == [1.5]
@@ -396,7 +519,9 @@ def test_prom_metrics_observes_manager_gauge_and_histogram():
),
}
with patch.object(
CPUOffloadingSpec, "build_metric_definitions", return_value=metric_definitions
OffloadingSpecFactory,
"get_spec_cls",
return_value=_spec_cls_with_metric_definitions(metric_definitions),
):
prom_metrics = OffloadPromMetrics(
vllm_config=_FakeVllmConfig(store_threshold=0), # type: ignore[arg-type]
@@ -416,20 +541,91 @@ def test_prom_metrics_observes_manager_gauge_and_histogram():
LOOKUP_LATENCY: _MetricType.HISTOGRAM,
},
_StatsKey.DATA: {
PENDING_STORES: 5,
LOOKUP_LATENCY: [0.2, 0.4],
PENDING_STORES: {(): 5},
LOOKUP_LATENCY: {(): [0.2, 0.4]},
},
}
)
gauge = prom_metrics.offloading_metrics[(0, PENDING_STORES)]
histogram = prom_metrics.offloading_metrics[(0, LOOKUP_LATENCY)]
gauge = prom_metrics.offloading_metrics[(0, PENDING_STORES, ())]
histogram = prom_metrics.offloading_metrics[(0, LOOKUP_LATENCY, ())]
assert gauge.set_values == [5]
assert histogram.observed == [0.2, 0.4]
histogram_def = prom_metrics._offloading_metric_defs[LOOKUP_LATENCY]
assert histogram_def.kwargs["buckets"] == (0.1, 1.0)
def test_prom_metrics_lazily_observes_labeled_metric():
metric_definitions = {
MY_COUNTER: OffloadingCounterMetadata(
documentation="counter with a label",
labelnames=(MY_LABEL,),
),
}
with patch.object(
OffloadingSpecFactory,
"get_spec_cls",
return_value=_spec_cls_with_metric_definitions(metric_definitions),
):
prom_metrics = OffloadPromMetrics(
vllm_config=_FakeVllmConfig(store_threshold=0), # type: ignore[arg-type]
metric_types={
Gauge: _FakeMetric,
Counter: _FakeMetric,
Histogram: _FakeMetric,
},
labelnames=["model_name", "engine"],
per_engine_labelvalues={0: ["model", "0"]},
)
assert (0, MY_COUNTER, ("a",)) not in prom_metrics.offloading_metrics
prom_metrics.observe(
{
_StatsKey.TYPES: {MY_COUNTER: _MetricType.COUNTER},
_StatsKey.DATA: {MY_COUNTER: {("a",): 7}},
}
)
counter = prom_metrics.offloading_metrics[(0, MY_COUNTER, ("a",))]
assert counter.increments == [7]
assert counter.labelvalues == ("model", "0", "a")
counter_def = prom_metrics._offloading_metric_defs[MY_COUNTER]
assert counter_def.kwargs["labelnames"] == ["model_name", "engine", MY_LABEL]
def test_prom_metrics_rejects_wrong_label_count():
metric_definitions = {
MY_COUNTER: OffloadingCounterMetadata(
documentation="counter with a label",
labelnames=(MY_LABEL,),
),
}
with patch.object(
OffloadingSpecFactory,
"get_spec_cls",
return_value=_spec_cls_with_metric_definitions(metric_definitions),
):
prom_metrics = OffloadPromMetrics(
vllm_config=_FakeVllmConfig(store_threshold=0), # type: ignore[arg-type]
metric_types={
Gauge: _FakeMetric,
Counter: _FakeMetric,
Histogram: _FakeMetric,
},
labelnames=["model_name", "engine"],
per_engine_labelvalues={0: ["model", "0"]},
)
with pytest.raises(AssertionError, match="expects 1 labels"):
prom_metrics.observe(
{
_StatsKey.TYPES: {MY_COUNTER: _MetricType.COUNTER},
_StatsKey.DATA: {MY_COUNTER: {("a", "extra"): 7}},
}
)
def test_prom_metrics_uses_configured_manager_metrics():
prom_metrics = OffloadPromMetrics(
vllm_config=_FakeVllmConfig(store_threshold=0), # type: ignore[arg-type]
@@ -458,9 +654,9 @@ def test_aggregate_into_empty_stats():
PENDING_STORES: _MetricType.GAUGE,
},
_StatsKey.DATA: {
LOAD_BYTES: 42,
LOAD_SIZE: [10, 20],
PENDING_STORES: 3,
LOAD_BYTES: {(): 42},
LOAD_SIZE: {(): [10, 20]},
PENDING_STORES: {(): 3},
},
},
)
@@ -469,9 +665,9 @@ def test_aggregate_into_empty_stats():
assert result is empty
values = result.data[_StatsKey.DATA]
assert values[LOAD_BYTES] == 42
assert values[LOAD_SIZE] == [10, 20]
assert values[PENDING_STORES] == 3
assert _unlabeled(values, LOAD_BYTES) == 42
assert _unlabeled(values, LOAD_SIZE) == [10, 20]
assert _unlabeled(values, PENDING_STORES) == 3
def test_prom_metrics_multi_engine_routing():
@@ -490,14 +686,13 @@ def test_prom_metrics_multi_engine_routing():
prom_metrics.observe(
{
_StatsKey.TYPES: {LOAD_BYTES: _MetricType.COUNTER},
_StatsKey.DATA: {LOAD_BYTES: 100},
_StatsKey.DATA: {LOAD_BYTES: {(): 100}},
},
engine_idx=1,
)
engine0 = prom_metrics.offloading_metrics[(0, LOAD_BYTES)]
engine1 = prom_metrics.offloading_metrics[(1, LOAD_BYTES)]
assert engine0.increments == []
assert (0, LOAD_BYTES, ()) not in prom_metrics.offloading_metrics
engine1 = prom_metrics.offloading_metrics[(1, LOAD_BYTES, ())]
assert engine1.increments == [100]
@@ -518,6 +713,6 @@ def test_prom_metrics_rejects_undeclared_metric():
prom_metrics.observe(
{
_StatsKey.TYPES: {"unknown:metric": _MetricType.COUNTER},
_StatsKey.DATA: {"unknown:metric": 1},
_StatsKey.DATA: {"unknown:metric": {(): 1}},
}
)
@@ -1,6 +1,6 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
from collections.abc import Iterable
from types import SimpleNamespace
from unittest.mock import MagicMock
import pytest
@@ -11,19 +11,16 @@ from tests.v1.kv_connector.unit.offloading_connector.utils import (
to_keys,
)
from tests.v1.kv_connector.unit.utils import EOS_TOKEN_ID
from vllm.distributed.kv_events import BlockRemoved, BlockStored
from vllm.distributed.kv_transfer.kv_connector.v1.offloading.scheduler import (
OffloadingConnectorScheduler,
RequestOffloadState,
)
from vllm.v1.core.kv_cache_utils import BlockHash
from vllm.v1.kv_cache_interface import (
FullAttentionSpec,
KVCacheGroupSpec,
SlidingWindowSpec,
)
from vllm.v1.kv_offload.base import (
OffloadingEvent,
OffloadingManager,
OffloadPolicy,
ReqContext,
@@ -145,31 +142,6 @@ def test_offloading_connector(request_runner, async_scheduling: bool):
runner.connector_scheduler._maximal_prefix_lookup = lambda key, req_context: 1
runner.run(decoded_tokens=[EOS_TOKEN_ID], expected_loaded=(3, 4, 5))
# test take_events
def to_hashes(int_hashes: list[int]) -> list[BlockHash]:
return [BlockHash(str(i).encode()) for i in int_hashes]
def take_events() -> Iterable[OffloadingEvent]:
yield OffloadingEvent(keys=to_keys([1, 2, 3]), medium="A", removed=False)
yield OffloadingEvent(keys=to_keys([4, 5, 6]), medium="B", removed=True)
runner.manager.take_events.side_effect = take_events
events = list(runner.scheduler_connector.take_events())
assert len(events) == 2
event = events[0]
assert isinstance(event, BlockStored)
assert event.block_hashes == to_hashes([1, 2, 3])
assert event.block_size == 0
assert event.medium == "A"
assert event.token_ids == []
assert event.parent_block_hash is None
assert event.lora_id is None
assert event.lora_name is None
event = events[1]
assert isinstance(event, BlockRemoved)
assert event.block_hashes == to_hashes([4, 5, 6])
assert event.medium == "B"
@pytest.mark.parametrize("async_scheduling", [True, False])
def test_request_preemption(request_runner, async_scheduling: bool):
@@ -1278,11 +1250,11 @@ def test_reset_cache_finalizes_finished_request_with_pending_store(
)
finalized: list[str] = []
runner.manager.on_request_finished.side_effect = (
lambda req_context: finalized.append(req_context.req_id)
runner.manager.on_request_finished.side_effect = lambda req_context: (
finalized.append(req_context.req_id)
)
runner.manager.prepare_store.side_effect = (
lambda keys, req_context: generate_store_output(keys)
runner.manager.prepare_store.side_effect = lambda keys, req_context: (
generate_store_output(keys)
)
# Decode a couple of blocks and keep every transfer in flight, so the
@@ -1314,6 +1286,100 @@ def test_reset_cache_finalizes_finished_request_with_pending_store(
assert req_id not in cs._req_status
def test_pending_transfer_defers_prefix_lookup():
"""A request with an in-flight store must not issue a load on re-admission.
With async scheduling, a preempted request's store can be flushed by the
worker before the scheduler consumes its completion. If the request is
re-admitted in that window, the connector should defer it instead of
looking up offloaded blocks and later asserting when a load is queued while
the store job is still tracked.
"""
scheduler = object.__new__(OffloadingConnectorScheduler)
scheduler.manager = MagicMock(spec=OffloadingManager)
request = SimpleNamespace(request_id="req-0")
group_state = SimpleNamespace(block_ids=[1, 2, 3])
req_status = SimpleNamespace(
group_states=[group_state],
transfer_jobs={123},
)
scheduler._req_status = {request.request_id: req_status}
matched_tokens, is_async = scheduler.get_num_new_matched_tokens(
request,
num_computed_tokens=0,
)
assert matched_tokens is None
assert is_async is False
assert group_state.block_ids == []
scheduler.manager.lookup.assert_not_called()
def test_async_preempt_readmit_before_transfer_output_is_deferred(request_runner):
"""A preempted request can be scheduled again before flush output is read.
EngineCore.step_with_batch_queue() may schedule a new batch while a prior
preemption batch is still queued. The store completion from jobs_to_flush is
only cleared when that queued output reaches update_from_output(), so the
re-admission path must defer while the scheduler still tracks the store.
"""
block_size = 4
block_size_factor = 3
offloaded_block_size = block_size * block_size_factor
runner = request_runner(
block_size=block_size,
num_gpu_blocks=100,
async_scheduling=True,
block_size_factor=block_size_factor,
)
free_block_queue = runner.scheduler.kv_cache_manager.block_pool.free_block_queue
num_free_blocks_empty = free_block_queue.num_free_blocks
req_id = "0"
runner.new_request(token_ids=[0] * offloaded_block_size * 2)
runner.manager.prepare_store.side_effect = lambda keys, req_context: (
generate_store_output(keys)
)
runner.run(decoded_tokens=[0], complete_transfers=False)
runner.run(
decoded_tokens=[0] * (2 * offloaded_block_size - block_size),
complete_transfers=False,
)
req_status = runner.connector_scheduler._req_status[req_id]
pending_store_jobs = set(req_status.transfer_jobs)
assert pending_store_jobs
assert all(
runner.connector_scheduler._jobs[jid].is_store for jid in pending_store_jobs
)
free_block_queue.num_free_blocks = 0
preempt_output = runner.scheduler.schedule()
assert preempt_output.preempted_req_ids == {req_id}
assert preempt_output.kv_connector_metadata is not None
assert pending_store_jobs <= preempt_output.kv_connector_metadata.jobs_to_flush
assert req_status.transfer_jobs == pending_store_jobs
# Simulate the async batch-queue window: schedule again before the
# preemption batch's ModelRunnerOutput is consumed by update_from_output().
free_block_queue.num_free_blocks = num_free_blocks_empty
assert runner.scheduler.reset_prefix_cache()
runner.connector_scheduler._maximal_prefix_lookup = lambda key, req_context: len(
key
)
readmit_output = runner.scheduler.schedule()
assert readmit_output.num_scheduled_tokens == {}
assert readmit_output.kv_connector_metadata is not None
assert readmit_output.kv_connector_metadata.load_jobs == {}
assert req_status.transfer_jobs == pending_store_jobs
@pytest.mark.parametrize("async_scheduling", [True, False])
def test_swa_alignment_skip(request_runner, async_scheduling: bool):
"""SWA blocks unreachable by the load path are skipped during store.
@@ -14,7 +14,12 @@ from tests.v1.kv_connector.unit.utils import (
create_vllm_config,
)
from vllm import SamplingParams
from vllm.config import KVTransferConfig, VllmConfig, set_current_vllm_config
from vllm.config import (
KVEventsConfig,
KVTransferConfig,
VllmConfig,
set_current_vllm_config,
)
from vllm.distributed.kv_transfer.kv_connector.v1 import KVConnectorRole
from vllm.distributed.kv_transfer.kv_connector.v1.offloading.common import (
OffloadingConnectorMetadata,
@@ -198,6 +203,9 @@ class RequestRunner:
"spec_module_path": "tests.v1.kv_connector.unit.offloading_connector.utils", # noqa: E501
# Preserve legacy behavior for tests; new opt-in tests override.
"offload_prompt_only": False,
# Exercise the self-describing KV events path by default;
# opt-out tests override this to cover the legacy placeholders.
"self_describing_kv_events": True,
}
if block_size_factor > 1:
extra_config["block_size"] = block_size * block_size_factor
@@ -209,6 +217,13 @@ class RequestRunner:
kv_role="kv_both",
kv_connector_extra_config=extra_config,
)
vllm_config.kv_events_config = KVEventsConfig(
# Enable so the offloading events tracker is active, but use the
# null publisher: these tests drain take_events directly and a
# real ZMQ publisher would bind a port per test.
enable_kv_cache_events=True,
publisher="null",
)
if kv_cache_groups is None:
kv_cache_groups = [
@@ -1,10 +1,13 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
import pytest
def test_mla_backend_rejects_cross_layer_kv_cache():
"""MLA backends return identity permutation (layers dim first)
to signal cross-layer KV cache is unsupported."""
def test_mla_common_backend_rejects_cross_layer_kv_cache():
"""MLACommonBackend defaults to the identity permutation (layers dim
first) so MLA backends whose decode kernels are not verified to honor
the cache's block-dim stride stay opted out of cross-layer KV cache."""
from vllm.model_executor.layers.attention.mla_attention import (
MLACommonBackend,
)
@@ -19,6 +22,35 @@ def test_mla_backend_rejects_cross_layer_kv_cache():
) == (0, 1, 2)
@pytest.mark.parametrize(
"backend_path",
[
"vllm.v1.attention.backends.mla.triton_mla.TritonMLABackend",
"vllm.v1.attention.backends.mla.cutlass_mla.CutlassMLABackend",
"vllm.v1.attention.backends.mla.flashattn_mla.FlashAttnMLABackend",
"vllm.v1.attention.backends.mla.flashmla.FlashMLABackend",
"vllm.v1.attention.backends.mla.flashinfer_mla.FlashInferMLABackend",
],
)
def test_verified_mla_backends_support_cross_layer_kv_cache(backend_path):
"""Backends whose decode kernels honor the cache's block-dim stride opt
in to the cross-layer layout with a non-identity permutation placing
num_blocks first in physical layout."""
module_path, name = backend_path.rsplit(".", 1)
backend = getattr(
pytest.importorskip(module_path, reason="backend deps unavailable"), name
)
stride_order = backend.get_kv_cache_stride_order(include_num_layers_dimension=True)
assert stride_order == (1, 0, 2, 3)
assert stride_order[0] != 0 # num_blocks first => cross-layer supported
assert backend.get_kv_cache_stride_order(include_num_layers_dimension=False) == (
0,
1,
2,
)
def test_deepseek_v32_indexer_rejects_cross_layer_kv_cache():
"""DeepseekV32Indexer returns identity permutation (layers dim first)
to signal cross-layer KV cache is unsupported."""

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