From 0c16eee3f1ff777298cc894c3eeb85f3880c6d6a Mon Sep 17 00:00:00 2001 From: Kam Basra Date: Mon, 5 Oct 2026 09:00:07 +0700 Subject: [PATCH 1/6] [Bugfix][KV Offload] Reset SimpleCPU eager-store placement state on MRV2 resume (#57816) Signed-off-by: Kam Basra Co-authored-by: OpenAI Codex --- tests/v1/simple_kv_offload/test_scheduler.py | 47 ++++++++++++++++++++ vllm/v1/simple_kv_offload/manager.py | 13 ++++-- 2 files changed, 56 insertions(+), 4 deletions(-) diff --git a/tests/v1/simple_kv_offload/test_scheduler.py b/tests/v1/simple_kv_offload/test_scheduler.py index 0f2a98e4a6..5e299d2788 100644 --- a/tests/v1/simple_kv_offload/test_scheduler.py +++ b/tests/v1/simple_kv_offload/test_scheduler.py @@ -1552,6 +1552,53 @@ def test_chunked_prefill_reads_live_block_ids() -> None: assert len(meta2.store_gpu_blocks) == 2 +@pytest.mark.parametrize("use_v2_model_runner", [False, True]) +@pytest.mark.parametrize("same_output", [False, True]) +def test_eager_store_resets_block_cursor_after_preemption( + use_v2_model_runner: bool, + same_output: bool, +) -> None: + """A resumed request stores its confirmed tail after preemption.""" + fix = make_scheduler(num_cpu_blocks=8, num_gpu_blocks=8, lazy=False) + sched = fix.scheduler + req = make_request(num_blocks=3) + + initial_blocks = _alloc_and_register(fix, req, num_blocks=2) + sched.update_state_after_alloc(req, initial_blocks, num_external_tokens=0) + initial_output = make_scheduler_output( + {req.request_id: 2 * BLOCK_SIZE}, + new_reqs={req.request_id: initial_blocks.get_block_ids()}, + ) + initial_meta = sched.build_connector_meta(initial_output) + assert len(initial_meta.store_gpu_blocks) == 2 + simulate_store_completion(sched, initial_meta.store_event) + + if not same_output: + preemption_output = make_scheduler_output({}) + preemption_output.preempted_req_ids = {req.request_id} + sched.build_connector_meta(preemption_output) + fix.gpu_block_pool.free_blocks(initial_blocks.blocks[0]) + + resumed_blocks = _alloc_and_register(fix, req, num_blocks=3) + sched.update_state_after_alloc(req, resumed_blocks, num_external_tokens=0) + if use_v2_model_runner: + resumed_output = make_scheduler_output( + {req.request_id: 3 * BLOCK_SIZE}, + new_reqs={req.request_id: resumed_blocks.get_block_ids()}, + ) + else: + resumed_output = make_scheduler_output( + {req.request_id: 3 * BLOCK_SIZE}, + cached_req_new_blocks={req.request_id: resumed_blocks.get_block_ids()}, + ) + resumed_output.scheduled_cached_reqs.resumed_req_ids.add(req.request_id) + if same_output: + resumed_output.preempted_req_ids = {req.request_id} + + resumed_meta = sched.build_connector_meta(resumed_output) + assert resumed_meta.store_gpu_blocks == [resumed_blocks.blocks[0][2].block_id] + + # --------------------------------------------------------------------------- # Test 10: Partial GPU prefix hit + CPU load + new compute blocks # --------------------------------------------------------------------------- diff --git a/vllm/v1/simple_kv_offload/manager.py b/vllm/v1/simple_kv_offload/manager.py index 3d57e73566..9d95940a11 100644 --- a/vllm/v1/simple_kv_offload/manager.py +++ b/vllm/v1/simple_kv_offload/manager.py @@ -738,14 +738,19 @@ class SimpleCPUOffloadScheduler: if scheduled_for_req: req_ids.append(req_id) - for req_id, new_block_id_groups, preempted in yield_req_data(scheduler_output): + # Preemption frees the request's blocks. It resumes with a full new + # block table (MRV1 resumed-cached, MRV2 NewRequestData). + for req_id in preempted_req_ids: + state = self._reqs_to_store.get(req_id) + if state is not None: + state.block_ids = tuple([] for _ in range(num_groups)) + state.num_stored_blocks = [0] * num_groups + + for req_id, new_block_id_groups, _ in yield_req_data(scheduler_output): state = self._reqs_to_store.get(req_id) if state is None or state.finished: continue - if preempted: - state.block_ids = tuple([] for _ in range(num_groups)) - state.num_stored_blocks = [0] * num_groups if new_block_id_groups: for g in range(min(num_groups, len(new_block_id_groups))): if new_block_id_groups[g] is not None: From ac8c2130a8a6ecc4940f6c5433cb2ba15e983aa7 Mon Sep 17 00:00:00 2001 From: Micah Williamson Date: Sun, 4 Oct 2026 23:19:19 -0500 Subject: [PATCH 2/6] [ROCm] Bump AITER to 0.1.24.post1 (#59794) Signed-off-by: Micah Williamson Co-authored-by: Fangzhou Ai <31551580+Fangzhou-Ai@users.noreply.github.com> --- docker/Dockerfile.rock_base | 2 +- docker/Dockerfile.rocm_base | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/docker/Dockerfile.rock_base b/docker/Dockerfile.rock_base index 3ee203a347..40100a1bbf 100644 --- a/docker/Dockerfile.rock_base +++ b/docker/Dockerfile.rock_base @@ -8,7 +8,7 @@ ARG TRITON_VERSION=3.8.0+git669b31ac ARG TRITON_REPO="https://github.com/ROCm/triton.git" ARG FA_BRANCH="0e60e394" ARG FA_REPO="https://github.com/Dao-AILab/flash-attention.git" -ARG AITER_BRANCH="v0.1.23" +ARG AITER_BRANCH="v0.1.24.post1" ARG AITER_REPO="https://github.com/ROCm/aiter.git" ARG MORI_BRANCH="v1.2.3.post1" ARG MORI_REPO="https://github.com/ROCm/mori.git" diff --git a/docker/Dockerfile.rocm_base b/docker/Dockerfile.rocm_base index 15e3e2468e..bd8099dd15 100644 --- a/docker/Dockerfile.rocm_base +++ b/docker/Dockerfile.rocm_base @@ -9,7 +9,7 @@ ARG PYTORCH_AUDIO_BRANCH="v2.11.0" ARG PYTORCH_AUDIO_REPO="https://github.com/pytorch/audio.git" ARG FA_BRANCH="0e60e394" ARG FA_REPO="https://github.com/Dao-AILab/flash-attention.git" -ARG AITER_BRANCH="v0.1.23" +ARG AITER_BRANCH="v0.1.24.post1" ARG AITER_REPO="https://github.com/ROCm/aiter.git" ARG MORI_BRANCH="v1.2.3.post1" ARG MORI_REPO="https://github.com/ROCm/mori.git" From e8a53a55e6ffe92645d8183066aa72e228ba7181 Mon Sep 17 00:00:00 2001 From: Gawr Gura 0721 <121242164+LiRunGuo@users.noreply.github.com> Date: Mon, 5 Oct 2026 12:52:39 +0800 Subject: [PATCH 3/6] [Bugfix][Attention] Support int4_per_token_head on non-power-of-two head sizes (#56198) Signed-off-by: Gawr Gura 0721 Co-authored-by: Gawr Gura 0721 Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> Co-authored-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com> --- .../attention/test_attention_selector.py | 25 +++++++++++++++++++ vllm/v1/attention/ops/int4_per_token_head.py | 24 +++++++++++++----- 2 files changed, 43 insertions(+), 6 deletions(-) diff --git a/tests/kernels/attention/test_attention_selector.py b/tests/kernels/attention/test_attention_selector.py index 24f02836e9..c6f4094eb6 100644 --- a/tests/kernels/attention/test_attention_selector.py +++ b/tests/kernels/attention/test_attention_selector.py @@ -455,6 +455,31 @@ def test_per_head_quant_scales_backend_selection( assert backend_name in str(exc_info.value) +@pytest.mark.skipif( + CudaPlatform is None, reason="CUDA platform is required for this test" +) +@pytest.mark.parametrize("head_size", [64, 80, 96, 128, 192, 256]) +def test_int4_per_token_head_head_size_selection(head_size: int): + """INT4 per-token-head KV cache is selectable at any even head size. + + The Hadamard rotation the INT4 write path applies is now defined for + non-power-of-two rows as well (it runs block-diagonally, see + ``fast_hadamard_transform``), so a model such as Phi-3 (head_size=96) must + keep selecting TRITON_ATTN instead of being rejected. + """ + vllm_config = VllmConfig(cache_config=CacheConfig(block_size=64)) + with ( + set_current_vllm_config(vllm_config), + patch("vllm.platforms.current_platform", CudaPlatform()), + ): + backend = get_attn_backend( + head_size=head_size, + dtype=torch.float16, + kv_cache_dtype="int4_per_token_head", + ) + assert backend.get_name() == "TRITON_ATTN" + + @pytest.mark.parametrize( "backend_name,use_non_causal,should_succeed", [ diff --git a/vllm/v1/attention/ops/int4_per_token_head.py b/vllm/v1/attention/ops/int4_per_token_head.py index 322d32bb81..3567f49d95 100644 --- a/vllm/v1/attention/ops/int4_per_token_head.py +++ b/vllm/v1/attention/ops/int4_per_token_head.py @@ -18,6 +18,7 @@ import torch from vllm.platforms import current_platform from vllm.triton_utils import tl, triton +from vllm.utils.math_utils import largest_power_of_2_divisor from vllm.v1.attention.ops.triton_attention_helpers import ( apply_alibi_to_score, apply_softcap, @@ -908,14 +909,15 @@ def unified_attention_int4( ) -> None: """Paged attention over the INT4 packed cache, writing into *out*. - The forward RHT has norm ``sqrt(head_size)``, so ``softmax_scale`` is - divided by ``head_size`` and the inverse RHT divides the output by - ``head_size`` as well. + The forward RHT has norm ``sqrt(b)`` for ``b`` the largest power-of-two + divisor of ``head_size`` (``b == head_size`` for power-of-two sizes), so + ``softmax_scale`` is divided by ``b`` and the inverse RHT divides the output + by ``b`` as well. """ q_orig_dtype = q.dtype q = single_rht(q.float()).to(q_orig_dtype) - head_size = q.shape[2] - softmax_scale = softmax_scale / head_size + rht_norm = largest_power_of_2_divisor(q.shape[2]) + softmax_scale = softmax_scale / rht_norm _launch_packed_attn( q=q, @@ -945,7 +947,7 @@ def unified_attention_int4( packing_factor=_INT4_PACKING_FACTOR, ) - out_f = single_rht(out.float(), inverse=True) / head_size + out_f = single_rht(out.float(), inverse=True) / rht_norm out.copy_(out_f.to(q_orig_dtype)) @@ -1148,6 +1150,16 @@ def single_rht(x: torch.Tensor, inverse: bool = False) -> torch.Tensor: before asymmetric quantization. """ d = x.shape[-1] + block = largest_power_of_2_divisor(d) + if block < d: + # H is only defined for power-of-two rows, so apply it + # block-diagonally: H_b inside each of the d / b blocks, for b the + # largest power-of-two divisor of d. That keeps H @ H.T = b * I + # without padding the row, hence the norm is sqrt(b), not sqrt(d). + lead = x.shape[:-1] + return single_rht(x.reshape(*lead, d // block, block), inverse).reshape( + *lead, d + ) d1 = _get_rht_signs(d, 0, x.device) if inverse: return fast_hadamard_transform(x) * d1 From 92044241a02f05de51420d654daf669de20b4691 Mon Sep 17 00:00:00 2001 From: Thiago Mota Martins Date: Mon, 5 Oct 2026 02:28:33 -0300 Subject: [PATCH 4/6] [Bugfix][ROCm] Preserve config during GPU memory profiling (#58014) Co-authored-by: Fangzhou Ai <31551580+Fangzhou-Ai@users.noreply.github.com> --- tests/v1/worker/test_gpu_worker.py | 40 ++++++++++++++++++++++++++++++ vllm/v1/worker/gpu_worker.py | 6 ++++- 2 files changed, 45 insertions(+), 1 deletion(-) diff --git a/tests/v1/worker/test_gpu_worker.py b/tests/v1/worker/test_gpu_worker.py index b63681deca..90adce1874 100644 --- a/tests/v1/worker/test_gpu_worker.py +++ b/tests/v1/worker/test_gpu_worker.py @@ -152,6 +152,46 @@ def test_scoped_max_split_ignores_async_allocator(monkeypatch): set_settings.assert_not_called() +@pytest.mark.parametrize("kv_cache_memory_bytes", [None, 1]) +def test_memory_profile_bounds_decode_logits_rows(monkeypatch, kv_cache_memory_bytes): + """ROCm decode workspace uses max_num_seqs during startup profiling.""" + from vllm.config import get_current_vllm_config_or_none + from vllm.v1.attention.ops.rocm_aiter_mla_sparse import _max_decode_logits_rows + + config = SimpleNamespace( + scheduler_config=SimpleNamespace(max_num_seqs=16), + speculative_config=SimpleNamespace(num_speculative_tokens=5), + ) + + class Profiled(Exception): + pass + + def profile_run(*, randomize_inputs=False): + assert randomize_inputs is True + # Without the config context, the helper falls back to all 32768 + # batched tokens, which would reserve 128 GiB at 1M context. + assert _max_decode_logits_rows(32768) == 16 * (1 + 5) + raise Profiled + + worker = SimpleNamespace( + vllm_config=config, + randomize_dummy_inputs=True, + cache_config=SimpleNamespace(kv_cache_memory_bytes=kv_cache_memory_bytes), + model_runner=SimpleNamespace(profile_run=profile_run, model_memory_usage=0), + init_snapshot=SimpleNamespace(free_memory=2), + _scoped_allocator_max_split=lambda **kwargs: nullcontext(), + ) + monkeypatch.setattr(gpu_worker, "maybe_apply_startup_plan", lambda _: None) + monkeypatch.setattr( + gpu_worker, "memory_profiling", lambda *args, **kwargs: nullcontext() + ) + + assert get_current_vllm_config_or_none() is None + with pytest.raises(Profiled): + gpu_worker.Worker.determine_available_memory(worker) + assert get_current_vllm_config_or_none() is None + + # Startup-plan persistence (vllm/v1/worker/startup_plan.py), applied and # saved by Worker.determine_available_memory / compile_or_warm_up_model. diff --git a/vllm/v1/worker/gpu_worker.py b/vllm/v1/worker/gpu_worker.py index 55d56919a4..3f416b94c7 100644 --- a/vllm/v1/worker/gpu_worker.py +++ b/vllm/v1/worker/gpu_worker.py @@ -615,7 +615,10 @@ class Worker(WorkerBase): if kv_cache_memory_bytes := self.cache_config.kv_cache_memory_bytes: # still need a profile run which compiles the model for # max_num_batched_tokens - self.model_runner.profile_run(randomize_inputs=self.randomize_dummy_inputs) + with set_current_vllm_config(self.vllm_config): + self.model_runner.profile_run( + randomize_inputs=self.randomize_dummy_inputs + ) msg = ( f"Initial free memory {format_gib(self.init_snapshot.free_memory)} " @@ -651,6 +654,7 @@ class Worker(WorkerBase): # limit are never split, so they stay releasable. Exits before # memory_profiling measures, restoring the original limit. self._scoped_allocator_max_split(max_split_size_mb=20), + set_current_vllm_config(self.vllm_config), ): self.model_runner.profile_run(randomize_inputs=self.randomize_dummy_inputs) From 042ab0305cc4215e2c6fb215f1d6f2e65dd7d4ac Mon Sep 17 00:00:00 2001 From: Thien Tran Date: Mon, 5 Oct 2026 15:03:29 +0800 Subject: [PATCH 5/6] [Perf][Qwen4Exp] Merge QSA QKVG and indexer QK projections (#59533) Signed-off-by: Thien Tran Co-authored-by: Codex Co-authored-by: namgyu-youn --- tests/kernels/test_bf16_skinny_gemm.py | 18 ++- tests/models/qwen4_exp/test_qsa_prepare.py | 12 +- .../qwen4_exp/nvidia/low_latency_gemm.py | 137 +++++------------- vllm/models/qwen4_exp/nvidia/model.py | 15 +- .../qwen4_exp/nvidia/ops/qsa_prepare.py | 16 +- vllm/models/qwen4_exp/nvidia/qsa.py | 102 ++++++++++++- 6 files changed, 181 insertions(+), 119 deletions(-) diff --git a/tests/kernels/test_bf16_skinny_gemm.py b/tests/kernels/test_bf16_skinny_gemm.py index 7cdaf2b495..7ca207c8e4 100644 --- a/tests/kernels/test_bf16_skinny_gemm.py +++ b/tests/kernels/test_bf16_skinny_gemm.py @@ -713,8 +713,8 @@ def test_low_latency_table_capability_routing( def test_qwen4_exp_hopper_plans_are_valid() -> None: plans = qwen4_exp_gemm.QWEN4_EXP_SM90_GEMM_PLANS - assert len(plans) == 9 - assert sum(map(len, plans.values())) == 31 + assert len(plans) == 7 + assert sum(map(len, plans.values())) == 24 assert (320, 10240) in plans assert (10240, 320) not in plans for (n, k), shape_plans in plans.items(): @@ -728,10 +728,11 @@ def test_qwen4_exp_hopper_plans_are_valid() -> None: @pytest.mark.parametrize( "capability,expected_plans", [ - ((10, 3), qwen4_exp_gemm.QWEN4_EXP_GEMM_PLANS), + ((10, 3), qwen4_exp_gemm.QWEN4_EXP_SM103_GEMM_PLANS), ((10, 0), qwen4_exp_gemm.QWEN4_EXP_SM100_GEMM_PLANS), ((9, 0), qwen4_exp_gemm.QWEN4_EXP_SM90_GEMM_PLANS), ((12, 1), qwen4_exp_gemm.QWEN4_EXP_SM121_GEMM_PLANS), + ((12, 0), {}), ((8, 0), {}), ], ) @@ -742,11 +743,16 @@ def test_qwen4_exp_gemm_capability_routing( ) -> None: monkeypatch.setattr( qwen4_exp_gemm.current_platform, - "is_device_capability", - lambda target: capability == target, + "get_device_capability", + lambda: capability, ) - assert qwen4_exp_gemm._gemm_plans() == expected_plans + assert ( + qwen4_exp_gemm.QWEN4_EXP_GEMM_PLANS_BY_CAPABILITY.get( + qwen4_exp_gemm.current_platform.get_device_capability(), {} + ) + == expected_plans + ) def test_installation_is_shape_specific_and_unquantized( diff --git a/tests/models/qwen4_exp/test_qsa_prepare.py b/tests/models/qwen4_exp/test_qsa_prepare.py index 14b836c132..e56a50b6c9 100644 --- a/tests/models/qwen4_exp/test_qsa_prepare.py +++ b/tests/models/qwen4_exp/test_qsa_prepare.py @@ -74,13 +74,14 @@ def _make_main_inputs(num_tokens: int, fp8_cache: bool) -> dict: slots = torch.randperm(num_blocks * MAIN_PAGE, device="cuda")[:num_tokens] slots[-1] = -1 norm_weights = torch.randn(2, MAIN_D, dtype=torch.bfloat16, device="cuda") * 0.2 + qkv_width = 2 * (MAIN_HQ + MAIN_HK) * MAIN_D return dict( main_qkv=torch.randn( num_tokens, - 2 * (MAIN_HQ + MAIN_HK) * MAIN_D, + qkv_width + (HQ + 1) * D, dtype=torch.bfloat16, device="cuda", - ), + )[:, :qkv_width], main_q_norm_weight=norm_weights[0], main_k_norm_weight=norm_weights[1], main_eps=EPS, @@ -316,9 +317,12 @@ def test_qsa_fused_prepare_matches_unfused( ) unfused_compressed = fused_compressed.clone() + # TP=4 QKVG width makes the indexer slice row-strided, as in the merged + # projection. The exact width is unimportant; the larger stride is tested. + main_qkvg_width = 3584 projected_qk = torch.randn( - num_tokens, (HQ + 1) * D, dtype=torch.bfloat16, device=device - ) + num_tokens, main_qkvg_width + (HQ + 1) * D, dtype=torch.bfloat16, device=device + )[:, main_qkvg_width:] q_weight = torch.randn(D, dtype=torch.bfloat16, device=device) * 0.2 k_weight = torch.randn(D, dtype=torch.bfloat16, device=device) * 0.2 diff --git a/vllm/models/qwen4_exp/nvidia/low_latency_gemm.py b/vllm/models/qwen4_exp/nvidia/low_latency_gemm.py index fa01c75698..80bd4d2ddc 100644 --- a/vllm/models/qwen4_exp/nvidia/low_latency_gemm.py +++ b/vllm/models/qwen4_exp/nvidia/low_latency_gemm.py @@ -23,7 +23,7 @@ from vllm.model_executor.layers.vocab_parallel_embedding import ( from vllm.platforms import current_platform from vllm.utils.torch_utils import direct_register_custom_op -QWEN4_EXP_GEMM_PLANS: dict[tuple[int, int], dict[int, SkinnyGemmConfig]] = { +QWEN4_EXP_SM103_GEMM_PLANS: dict[tuple[int, int], dict[int, SkinnyGemmConfig]] = { # GDN fused QKVZ projection, TP=4. (4096, 2560): { 1: SkinnyGemmConfig(1, 64, 4, k_unroll=4), @@ -43,18 +43,11 @@ QWEN4_EXP_GEMM_PLANS: dict[tuple[int, int], dict[int, SkinnyGemmConfig]] = { 8: SkinnyGemmConfig(8, 128, 1, k_unroll=4, vector_width=4), 16: SkinnyGemmConfig(16, 128, 1, k_unroll=4, vector_width=4), }, - # QSA fused QKV/gate projection, TP=4. - (3584, 2560): { - 1: SkinnyGemmConfig(1, 128, 4, k_unroll=4, vector_width=4), - 2: SkinnyGemmConfig(2, 64, 2, k_unroll=2), - 4: SkinnyGemmConfig(4, 64, 2, k_unroll=2), - }, - # QSA indexer Q/K projection, replicated in a TP=4 deployment. - (640, 2560): { - 1: SkinnyGemmConfig(1, 128, 1, k_unroll=4, vector_width=4), - 2: SkinnyGemmConfig(2, 128, 1, k_unroll=4, vector_width=4), - 4: SkinnyGemmConfig(4, 128, 1, k_unroll=4, vector_width=4), - 8: SkinnyGemmConfig(8, 128, 1, k_unroll=4, vector_width=4), + # QSA fused QKV/gate + replicated indexer Q/K, TP=4 (GB300). + (4224, 2560): { + 1: SkinnyGemmConfig(1, 128, 4, vector_width=4, static_k=2560), + 2: SkinnyGemmConfig(2, 128, 2, vector_width=4, static_k=2560), + 4: SkinnyGemmConfig(4, 64, 4, vector_width=4, static_k=2560), }, # Shared-expert fused gate/up projection, TP=4. (320, 2560): { @@ -69,13 +62,6 @@ QWEN4_EXP_GEMM_PLANS: dict[tuple[int, int], dict[int, SkinnyGemmConfig]] = { 1: SkinnyGemmConfig(1, 64, 4, k_unroll=2), 2: SkinnyGemmConfig(2, 32, 4, k_unroll=2), }, - # HC merged down/injection projection, replicated in a TP=4 deployment. - (336, 10240): { - 1: SkinnyGemmConfig(1, 128, 1, static_k=10240), - 2: SkinnyGemmConfig(2, 128, 1, static_k=10240), - 4: SkinnyGemmConfig(4, 128, 2, static_k=10240), - 8: SkinnyGemmConfig(8, 128, 1, k_unroll=4), - }, } # H200 plans selected by exhaustive CUDA graph replay measurements over @@ -88,6 +74,12 @@ QWEN4_EXP_SM90_GEMM_PLANS: dict[tuple[int, int], dict[int, SkinnyGemmConfig]] = 1: SkinnyGemmConfig(1, 128, 2, vector_width=4, static_k=2560), 2: SkinnyGemmConfig(2, 64, 4, vector_width=4, static_k=2560), }, + # QSA fused QKV/gate + replicated indexer Q/K, TP=4. + (4224, 2560): { + 1: SkinnyGemmConfig(1, 128, 2, k_unroll=6, vector_width=2, static_k=2560), + 2: SkinnyGemmConfig(2, 128, 2, k_unroll=3, vector_width=2, static_k=2560), + 4: SkinnyGemmConfig(4, 64, 4, k_unroll=2, vector_width=4), + }, # GDN and QSA output projections, TP=4. (2560, 1536): { 1: SkinnyGemmConfig(1, 128, 4, vector_width=2, static_k=1536), @@ -102,18 +94,6 @@ QWEN4_EXP_SM90_GEMM_PLANS: dict[tuple[int, int], dict[int, SkinnyGemmConfig]] = 8: SkinnyGemmConfig(8, 128, 1, vector_width=4, static_k=2560), 16: SkinnyGemmConfig(16, 128, 1, vector_width=4, static_k=2560), }, - # QSA fused QKV/gate projection, TP=4. - (3584, 2560): { - 1: SkinnyGemmConfig(1, 128, 4, vector_width=2, static_k=2560), - 2: SkinnyGemmConfig(2, 64, 4, k_unroll=5), - }, - # QSA indexer Q/K projection, replicated in a TP=4 deployment. - (640, 2560): { - 1: SkinnyGemmConfig(1, 256, 2, vector_width=2, static_k=2560), - 2: SkinnyGemmConfig(2, 128, 2, vector_width=4, static_k=2560), - 4: SkinnyGemmConfig(4, 128, 1, vector_width=2, static_k=2560), - 8: SkinnyGemmConfig(8, 128, 2, vector_width=4, static_k=2560), - }, # Shared-expert fused gate/up projection, TP=4. (320, 2560): { 1: SkinnyGemmConfig(1, 64, 2, vector_width=4, static_k=2560), @@ -127,13 +107,6 @@ QWEN4_EXP_SM90_GEMM_PLANS: dict[tuple[int, int], dict[int, SkinnyGemmConfig]] = 1: SkinnyGemmConfig(1, 64, 2, vector_width=2, static_k=2560), 2: SkinnyGemmConfig(2, 64, 2, vector_width=2, static_k=2560), }, - # HC merged down/injection projection, replicated in a TP=4 deployment. - (336, 10240): { - 1: SkinnyGemmConfig(1, 256, 1, k_unroll=5), - 2: SkinnyGemmConfig(2, 256, 3, static_k=10240), - 4: SkinnyGemmConfig(4, 256, 3, static_k=10240), - 8: SkinnyGemmConfig(8, 256, 3, static_k=10240), - }, # Final HC down projection, replicated in a TP=4 deployment. (320, 10240): { 1: SkinnyGemmConfig(1, 256, 1, static_k=10240), @@ -169,30 +142,17 @@ QWEN4_EXP_SM100_GEMM_PLANS: dict[tuple[int, int], dict[int, SkinnyGemmConfig]] = 4: SkinnyGemmConfig(4, 256, 1, k_unroll=2), 8: SkinnyGemmConfig(8, 256, 1), }, - # HC merged down/injection projection, replicated in a TP=4 deployment. - (336, 10240): { - 1: SkinnyGemmConfig(1, 256, 1, k_unroll=4, static_k=10240), - 2: SkinnyGemmConfig(2, 256, 1, k_unroll=4, static_k=10240), - 8: SkinnyGemmConfig(8, 256, 3, static_k=10240), - }, - # QSA indexer Q/K projection, replicated in a TP=4 deployment. - (640, 2560): { - 1: SkinnyGemmConfig(1, 128, 2, k_unroll=4, vector_width=4, static_k=2560), - 2: SkinnyGemmConfig(2, 128, 1, k_unroll=4, vector_width=4, static_k=2560), - 4: SkinnyGemmConfig(4, 128, 1, k_unroll=4, vector_width=4, static_k=2560), - 8: SkinnyGemmConfig(8, 128, 1, k_unroll=4, vector_width=4, static_k=2560), - }, # GDN and QSA output projections, TP=4. (2560, 1536): { 1: SkinnyGemmConfig(1, 128, 2, k_unroll=2, vector_width=4, static_k=1536), 2: SkinnyGemmConfig(2, 64, 2, k_unroll=2, static_k=1536), 4: SkinnyGemmConfig(4, 64, 2, k_unroll=2, static_k=1536), }, - # QSA fused QKV/gate projection, TP=4. - (3584, 2560): { - 1: SkinnyGemmConfig(1, 128, 4, k_unroll=2, vector_width=4, static_k=2560), - 2: SkinnyGemmConfig(2, 128, 4, k_unroll=2, vector_width=4, static_k=2560), - 4: SkinnyGemmConfig(4, 64, 2, k_unroll=4, vector_width=4, static_k=2560), + # QSA fused QKV/gate + replicated indexer Q/K, TP=4. + (4224, 2560): { + 1: SkinnyGemmConfig(1, 128, 4, vector_width=4, static_k=2560), + 2: SkinnyGemmConfig(2, 128, 2, vector_width=4, static_k=2560), + 4: SkinnyGemmConfig(4, 64, 4, vector_width=4, static_k=2560), }, # GDN fused QKVZ projection, TP=4. (4096, 2560): { @@ -207,7 +167,7 @@ QWEN4_EXP_SM100_GEMM_PLANS: dict[tuple[int, int], dict[int, SkinnyGemmConfig]] = } -# DGX Spark (GB10) plans, TP=1. +# DGX Spark (GB10) plans, TP=1 unless noted. QWEN4_EXP_SM121_GEMM_PLANS: dict[tuple[int, int], dict[int, SkinnyGemmConfig]] = { # Shared-expert gate. (1, 2560): { @@ -225,14 +185,6 @@ QWEN4_EXP_SM121_GEMM_PLANS: dict[tuple[int, int], dict[int, SkinnyGemmConfig]] = 8: SkinnyGemmConfig(8, 128, 4, k_unroll=2, vector_width=4, static_k=2560), 16: SkinnyGemmConfig(16, 64, 1), }, - # HC merged down/injection projection. - (336, 10240): { - 1: SkinnyGemmConfig(1, 256, 1, vector_width=4, static_k=10240), - 2: SkinnyGemmConfig(2, 256, 1, k_unroll=4, vector_width=4, static_k=10240), - 4: SkinnyGemmConfig(4, 128, 1, k_unroll=4, static_k=10240), - 8: SkinnyGemmConfig(8, 128, 4, k_unroll=4, static_k=10240), - 16: SkinnyGemmConfig(16, 128, 4, k_unroll=4, static_k=10240), - }, # Router. (512, 2560): { 1: SkinnyGemmConfig(1, 256, 1, k_unroll=4, vector_width=2, static_k=2560), @@ -241,7 +193,7 @@ QWEN4_EXP_SM121_GEMM_PLANS: dict[tuple[int, int], dict[int, SkinnyGemmConfig]] = 8: SkinnyGemmConfig(8, 64, 2, static_k=2560), 16: SkinnyGemmConfig(16, 64, 2, k_unroll=2, static_k=2560), }, - # QSA indexer Q/K projection. + # Shared-expert fused gate/up, TP=2; retain after indexer fusion. (640, 2560): { 1: SkinnyGemmConfig(1, 256, 1, vector_width=2, static_k=2560), 2: SkinnyGemmConfig(2, 256, 1, vector_width=2, static_k=2560), @@ -280,13 +232,13 @@ QWEN4_EXP_SM121_GEMM_PLANS: dict[tuple[int, int], dict[int, SkinnyGemmConfig]] = 4: SkinnyGemmConfig(4, 32, 1, k_unroll=2, vector_width=2, static_k=320), 8: SkinnyGemmConfig(8, 32, 1, vector_width=2, static_k=320), }, - # QSA fused QKV/gate projection. - (13312, 2560): { + # QSA fused QKV/gate + replicated indexer Q/K projection. + (13952, 2560): { 1: SkinnyGemmConfig(1, 32, 1, vector_width=4, static_k=2560), - 2: SkinnyGemmConfig(2, 32, 1, vector_width=4, static_k=2560), + 2: SkinnyGemmConfig(2, 32, 1, k_unroll=2, vector_width=4, static_k=2560), 4: SkinnyGemmConfig(4, 32, 1, vector_width=4, static_k=2560), 8: SkinnyGemmConfig(8, 64, 1, vector_width=2, static_k=2560), - 16: SkinnyGemmConfig(16, 32, 1, vector_width=4, static_k=2560), + 16: SkinnyGemmConfig(16, 32, 1, k_unroll=2, vector_width=4, static_k=2560), }, # GDN fused QKVZ projection. (16384, 2560): { @@ -307,32 +259,14 @@ QWEN4_EXP_SM121_GEMM_PLANS: dict[tuple[int, int], dict[int, SkinnyGemmConfig]] = } -def _is_sm121() -> bool: - return current_platform.is_device_capability((12, 1)) - - -def _is_sm100() -> bool: - return current_platform.is_device_capability((10, 0)) - - -def _is_sm103() -> bool: - return current_platform.is_device_capability((10, 3)) - - -def _is_sm90() -> bool: - return current_platform.is_device_capability((9, 0)) - - -def _gemm_plans() -> dict[tuple[int, int], dict[int, SkinnyGemmConfig]]: - if _is_sm103(): - return QWEN4_EXP_GEMM_PLANS - if _is_sm100(): - return QWEN4_EXP_SM100_GEMM_PLANS - if _is_sm90(): - return QWEN4_EXP_SM90_GEMM_PLANS - if _is_sm121(): - return QWEN4_EXP_SM121_GEMM_PLANS - return {} +QWEN4_EXP_GEMM_PLANS_BY_CAPABILITY: dict[ + tuple[int, int] | None, dict[tuple[int, int], dict[int, SkinnyGemmConfig]] +] = { + (10, 3): QWEN4_EXP_SM103_GEMM_PLANS, + (10, 0): QWEN4_EXP_SM100_GEMM_PLANS, + (9, 0): QWEN4_EXP_SM90_GEMM_PLANS, + (12, 1): QWEN4_EXP_SM121_GEMM_PLANS, +} def _is_packed_row_major(tensor: torch.Tensor) -> bool: @@ -376,7 +310,10 @@ class Qwen4ExpLowLatencyEmbeddingMethod( def _qwen4_exp_low_latency_gemm(x: torch.Tensor, weight: torch.Tensor) -> torch.Tensor: - plan = _gemm_plans().get((weight.shape[0], weight.shape[1])) + plans = QWEN4_EXP_GEMM_PLANS_BY_CAPABILITY.get( + current_platform.get_device_capability(), {} + ) + plan = plans.get((weight.shape[0], weight.shape[1])) config = None if plan is None else plan.get(x.shape[0]) if ( config is not None @@ -404,7 +341,9 @@ def enable_qwen4_exp_low_latency_gemm( module: nn.Module, dtype: torch.dtype, ) -> None: - plans = _gemm_plans() + plans = QWEN4_EXP_GEMM_PLANS_BY_CAPABILITY.get( + current_platform.get_device_capability(), {} + ) if dtype != torch.bfloat16 or not plans: return if not shape_dynamic_skinny_gemm.is_available(): diff --git a/vllm/models/qwen4_exp/nvidia/model.py b/vllm/models/qwen4_exp/nvidia/model.py index 73fe2cc468..2de883cc47 100644 --- a/vllm/models/qwen4_exp/nvidia/model.py +++ b/vllm/models/qwen4_exp/nvidia/model.py @@ -622,7 +622,7 @@ class Qwen4ExpForCausalLM( IsHybrid, ): packed_modules_mapping = { - "qkv_proj": ["q_proj", "k_proj", "v_proj"], + "qkv_proj": ["q_proj", "k_proj", "v_proj", "indexer.index_qk_proj"], "gate_up_proj": ["gate_proj", "up_proj"], "kv_proj": ["key_proj", "value_proj"], "in_proj_qkvz": ["in_proj_qkv", "in_proj_z"], @@ -641,6 +641,12 @@ class Qwen4ExpForCausalLM( def __init__(self, *, vllm_config: VllmConfig, prefix: str = "") -> None: super().__init__() config: Qwen4ExpTextConfig = vllm_config.model_config.hf_text_config + if vllm_config.lora_config is not None: + # LoRA does not support the merged QKV/indexer projection, so its + # packed mapping must keep the indexer separate. + self.packed_modules_mapping = self.packed_modules_mapping | { + "qkv_proj": ["q_proj", "k_proj", "v_proj"], + } self.vllm_config = vllm_config self.model_config = vllm_config.model_config self.quant_config = vllm_config.quant_config @@ -866,6 +872,7 @@ class Qwen4ExpForConditionalGeneration( requires_raw_input_tokens = True packed_modules_mapping = Qwen3_5ForConditionalGeneration.packed_modules_mapping | { + "qkv_proj": Qwen4ExpForCausalLM.packed_modules_mapping["qkv_proj"], "kv_proj": ["key_proj", "value_proj"], "input_mix_weight_down_block_inject": [ "input_mix_weight_down", @@ -883,6 +890,12 @@ class Qwen4ExpForConditionalGeneration( def __init__(self, *, vllm_config: VllmConfig, prefix: str = "model") -> None: nn.Module.__init__(self) config: Qwen4ExpConfig = vllm_config.model_config.hf_config + if vllm_config.lora_config is not None: + # LoRA does not support the merged QKV/indexer projection, so its + # packed mapping must keep the indexer separate. + self.packed_modules_mapping = self.packed_modules_mapping | { + "qkv_proj": ["q_proj", "k_proj", "v_proj"], + } quant_config = vllm_config.quant_config multimodal_config = vllm_config.model_config.multimodal_config if multimodal_config is None: diff --git a/vllm/models/qwen4_exp/nvidia/ops/qsa_prepare.py b/vllm/models/qwen4_exp/nvidia/ops/qsa_prepare.py index 6e97a89781..a26b699657 100644 --- a/vllm/models/qwen4_exp/nvidia/ops/qsa_prepare.py +++ b/vllm/models/qwen4_exp/nvidia/ops/qsa_prepare.py @@ -154,6 +154,7 @@ def _qsa_prepare_kernel( MROPE_H: tl.constexpr, MROPE_W: tl.constexpr, main_qkv_ptr, + main_qkv_stride_token, main_q_norm_weight_ptr, main_k_norm_weight_ptr, main_eps, @@ -178,12 +179,12 @@ def _qsa_prepare_kernel( # indexer work. RoPE covers the first D // 2 dims, the width of the # shared cos/sin table. main_pid = pid - num_index_work - token = main_pid // (MAIN_HQ + MAIN_HK) + token = (main_pid // (MAIN_HQ + MAIN_HK)).to(tl.int64) head = main_pid % (MAIN_HQ + MAIN_HK) HALF: tl.constexpr = D // 4 dims = tl.arange(0, MAIN_D) rot = tl.arange(0, HALF) - row = main_qkv_ptr + token * (2 * (MAIN_HQ + MAIN_HK) * MAIN_D) + row = main_qkv_ptr + token * main_qkv_stride_token is_k = head >= MAIN_HQ kv_head = head - MAIN_HQ if is_k: @@ -241,7 +242,7 @@ def _qsa_prepare_kernel( num_head_tiles: tl.constexpr = tl.cdiv(HQ, TILE_H_Q) token_tile = q_pid // num_head_tiles head_tile = q_pid % num_head_tiles - tokens = token_tile * TILE_T_Q + tl.arange(0, TILE_T_Q) + tokens = (token_tile * TILE_T_Q + tl.arange(0, TILE_T_Q)).to(tl.int64) heads = head_tile * TILE_H_Q + tl.arange(0, TILE_H_Q) valid_tokens = tokens < num_tokens valid_heads = heads < HQ @@ -342,7 +343,8 @@ def _qsa_prepare_kernel( & (source_tokens < num_tokens) ) current_base = ( - k_ptr + tl.maximum(source_tokens, 0)[:, None] * k_stride_token + k_ptr + + tl.maximum(source_tokens, 0).to(tl.int64)[:, None] * k_stride_token ) cached_base = ( state_cache_ptr @@ -449,7 +451,7 @@ def _qsa_prepare_kernel( # One CTA per request commits only the suffix retained by the ring. num_state_rows = tl.minimum(query_len, STATE_SIZE) for state_offset in tl.range(0, num_state_rows): - token = query_end - num_state_rows + state_offset + token = (query_end - num_state_rows + state_offset).to(tl.int64) valid_token = ( (token >= query_start) & (token < query_end) & (token < num_tokens) ) @@ -567,7 +569,8 @@ def qsa_prepare( section = mrope_section if mrope_section is not None else (0, 0, 0) assert len(section) == 3 qkv_width = 2 * (num_main_q_heads + num_main_kv_heads) * main_head_dim - assert main_qkv.shape == (num_tokens, qkv_width) and main_qkv.is_contiguous() + assert main_qkv.shape == (num_tokens, qkv_width) + assert main_qkv.stride(-1) == 1 assert main_slot_mapping.shape == (num_tokens,) if num_tokens <= 4096: @@ -622,6 +625,7 @@ def qsa_prepare( MROPE_H=section[1], MROPE_W=section[2], main_qkv_ptr=main_qkv, + main_qkv_stride_token=main_qkv.stride(0), main_q_norm_weight_ptr=main_q_norm_weight, main_k_norm_weight_ptr=main_k_norm_weight, main_eps=main_eps, diff --git a/vllm/models/qwen4_exp/nvidia/qsa.py b/vllm/models/qwen4_exp/nvidia/qsa.py index d00cd186b0..6877189768 100644 --- a/vllm/models/qwen4_exp/nvidia/qsa.py +++ b/vllm/models/qwen4_exp/nvidia/qsa.py @@ -4,6 +4,7 @@ from __future__ import annotations +from collections.abc import Callable, Iterable from typing import ClassVar, cast import torch @@ -20,10 +21,15 @@ from vllm.model_executor.layers.attention.attention import ( ) from vllm.model_executor.layers.attention_layer_base import AttentionLayerBase from vllm.model_executor.layers.layernorm import GemmaRMSNorm -from vllm.model_executor.layers.linear import QKVParallelLinear, RowParallelLinear +from vllm.model_executor.layers.linear import ( + MergedColumnParallelLinear, + QKVParallelLinear, + RowParallelLinear, +) from vllm.model_executor.layers.quantization import QuantizationConfig from vllm.model_executor.layers.rotary_embedding import MRotaryEmbedding, get_rope from vllm.model_executor.models.qwen3_next import Qwen3NextAttention +from vllm.model_executor.models.utils import AutoWeightsLoader, WeightsMapper from vllm.platforms import current_platform from vllm.platforms.interface import DeviceCapability from vllm.utils.torch_utils import ( @@ -53,6 +59,72 @@ from . import model from .indexer_qsa import QSAIndexer +class Qwen4ExpQSAQKVIndexerLinear(MergedColumnParallelLinear): + """Pack sharded Q/gate, K, V and replicated indexer Q/K in one GEMM.""" + + def __init__( + self, + qkv_proj: QKVParallelLinear, + index_size: int, + quant_config: QuantizationConfig | None, + ) -> None: + self.num_kv_head_replicas = qkv_proj.num_kv_head_replicas + super().__init__( + input_size=qkv_proj.input_size, + output_sizes=[*qkv_proj.output_sizes, index_size * qkv_proj.tp_size], + bias=False, + params_dtype=qkv_proj.params_dtype, + quant_config=quant_config, + prefix=qkv_proj.prefix, + ) + + def _load_shard( + self, + loader: Callable[..., None], + param: nn.Parameter, + loaded_weight: torch.Tensor, + loaded_shard_id: int | tuple[int, ...] | None, + ) -> None: + tp_rank = self.tp_rank + param_tp_rank = getattr(param, "tp_rank", None) + if loaded_shard_id == 3: + shard_rank = 0 + elif loaded_shard_id in (1, 2): + shard_rank = tp_rank // self.num_kv_head_replicas + else: + shard_rank = tp_rank + self.tp_rank = shard_rank + if param_tp_rank is not None: + param.tp_rank = shard_rank + try: + loader(param, loaded_weight, loaded_shard_id) + finally: + self.tp_rank = tp_rank + if param_tp_rank is not None: + param.tp_rank = param_tp_rank + + def weight_loader(self, param, loaded_weight, loaded_shard_id=None) -> None: + self._load_shard(super().weight_loader, param, loaded_weight, loaded_shard_id) + + def weight_loader_v2(self, param, loaded_weight, loaded_shard_id=None) -> None: + self._load_shard( + super().weight_loader_v2, param, loaded_weight, loaded_shard_id + ) + + def load_weights( + self, weights: Iterable[tuple[str, torch.Tensor]] + ) -> Iterable[str]: + def remap_shards(): + for name, weight in weights: + shard_id = getattr(weight, "shard_id", None) + if isinstance(shard_id, str): + weight = weight.detach() + weight.shard_id = {"q": 0, "k": 1, "v": 2}[shard_id] + yield name, weight + + return super().load_weights(remap_shards()) + + class Qwen4ExpQSAMetadataBuilder(FlashAttentionMetadataBuilder): """Flash metadata supporting uniform decode and target-verify graphs.""" @@ -411,6 +483,16 @@ class Qwen4ExpQSAAttention(Qwen3NextAttention, AttentionLayerBase): self.use_fused_qsa_prepare = ( self.use_fused_qk_norm_rope_gate and self.indexer.use_fused_pre_indexer ) + self.fuse_indexer_projection = vllm_config.lora_config is None + if self.fuse_indexer_projection: + self.index_qk_size = self.indexer.index_qk_proj.output_size + self.qkv_proj = Qwen4ExpQSAQKVIndexerLinear( + self.qkv_proj, + self.index_qk_size, + model.without_modelopt_fp4(quant_config), + ) + del self.indexer.index_qk_proj + max_tokens = vllm_config.scheduler_config.max_num_batched_tokens # PACKED selection buffer: the trailing column holds each row's # valid-entry count (written by the expand kernel) — never a token @@ -432,6 +514,16 @@ class Qwen4ExpQSAAttention(Qwen3NextAttention, AttentionLayerBase): raise ValueError(f"Duplicate layer name: {self.layer_name}") static_context[self.layer_name] = self + def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]: + mapper = None + if self.fuse_indexer_projection: + mapper = WeightsMapper( + orig_to_new_stacked={ + "indexer.index_qk_proj.": ("qkv_proj.", 3), + } + ) + return AutoWeightsLoader(self).load_weights(weights, mapper=mapper) + def get_attn_backend(self) -> type[AttentionBackend]: return self.attn_backend @@ -518,6 +610,12 @@ class Qwen4ExpQSAAttention(Qwen3NextAttention, AttentionLayerBase): hidden_states: torch.Tensor, ) -> torch.Tensor: qkv, _ = self.qkv_proj(hidden_states) + if self.fuse_indexer_projection: + qkv, projected_qk = qkv.split( + [2 * self.q_size + 2 * self.kv_size, self.index_qk_size], dim=-1 + ) + else: + projected_qk, _ = self.indexer.index_qk_proj(hidden_states) num_tokens = hidden_states.shape[0] if not self.use_fused_qsa_prepare: q, k, v, gate = self._project_qkv_gate(qkv, positions) @@ -529,8 +627,6 @@ class Qwen4ExpQSAAttention(Qwen3NextAttention, AttentionLayerBase): # Norm/RoPE/gate and the K/V cache write happen inside _run_qsa. query = key = value = gate = None attn_output = qkv.new_empty(num_tokens, self.num_heads, self.head_dim) - # Keep the index projection outside the eager break. - projected_qk, _ = self.indexer.index_qk_proj(hidden_states) self._run_qsa( projected_qk, positions, From 34051ad7145104e73ba3a7f5b39506ad0949bb8d Mon Sep 17 00:00:00 2001 From: Harshal Janjani Date: Mon, 5 Oct 2026 11:51:15 +0400 Subject: [PATCH 6/6] feat: Add video support for the Transformers backend (#57441) Signed-off-by: Harshal Janjani Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com> Co-authored-by: Harry Mellor <19981378+hmellor@users.noreply.github.com> Co-authored-by: Claude Opus 5 (1M context) --- docs/models/supported_models.md | 4 +- .../generation/test_transformers_audio.py | 16 +- .../generation/test_transformers_video.py | 100 ++++ .../processing/test_transformers_video.py | 277 +++++++++++ tests/models/transformers/test_backend.py | 19 +- vllm/config/model.py | 24 +- .../models/transformers/multimodal.py | 432 +++++++++++++++--- .../models/transformers/utils.py | 10 +- vllm/transformers_utils/config.py | 8 +- 9 files changed, 775 insertions(+), 115 deletions(-) create mode 100644 tests/models/multimodal/generation/test_transformers_video.py create mode 100644 tests/models/multimodal/processing/test_transformers_video.py diff --git a/docs/models/supported_models.md b/docs/models/supported_models.md index e461379bcb..a55559b0d6 100644 --- a/docs/models/supported_models.md +++ b/docs/models/supported_models.md @@ -19,12 +19,10 @@ vLLM also supports model implementations that are available in Transformers. We Currently, the Transformers modeling backend works for the following: -- Modalities: embedding models, language models, vision-language models* and audio-language models +- Modalities: embedding models, language models, vision-language models and audio-language models - Architectures: encoder-only, decoder-only, mixture-of-experts - Attention types: full attention and/or sliding attention -_*Vision-language models currently accept only image inputs. Support for video inputs will be added in a future release._ - If the Transformers model implementation follows all the steps in [writing a custom model](#writing-custom-models) then, when used with the Transformers modeling backend, it will be compatible with the following features of vLLM: - All the features listed in the [compatibility matrix](../features/README.md#feature-x-feature) diff --git a/tests/models/multimodal/generation/test_transformers_audio.py b/tests/models/multimodal/generation/test_transformers_audio.py index 9195579877..230cbc910c 100644 --- a/tests/models/multimodal/generation/test_transformers_audio.py +++ b/tests/models/multimodal/generation/test_transformers_audio.py @@ -69,13 +69,8 @@ AUDIO_MODEL_SETTINGS: dict[str, dict[str, Any]] = { } -@pytest.mark.parametrize("model_id", list(AUDIO_MODEL_SETTINGS)) -def test_transformers_audio_generation( - hf_runner: type[HfRunner], - vllm_runner: type[VllmRunner], - monkeypatch, - model_id: str, -): +@pytest.fixture(autouse=True) +def use_spawn_for_audio_models(monkeypatch): """Single-process workaround for V1 fork safety deadlock issue (vllm-project/vllm/issues/17676). Running multiple audio models together under pytest can cause (possibly flaky) hangs, so they are grouped under @@ -87,6 +82,13 @@ def test_transformers_audio_generation( disable_envs_cache() monkeypatch.setenv("VLLM_WORKER_MULTIPROC_METHOD", "spawn") + +@pytest.mark.parametrize("model_id", list(AUDIO_MODEL_SETTINGS)) +def test_transformers_audio_generation( + hf_runner: type[HfRunner], + vllm_runner: type[VllmRunner], + model_id: str, +): settings = AUDIO_MODEL_SETTINGS[model_id] audio_lora_path = settings.get("audio_lora_path") diff --git a/tests/models/multimodal/generation/test_transformers_video.py b/tests/models/multimodal/generation/test_transformers_video.py new file mode 100644 index 0000000000..bbb31b7be8 --- /dev/null +++ b/tests/models/multimodal/generation/test_transformers_video.py @@ -0,0 +1,100 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project + +from typing import Any + +import pytest +from transformers import AutoModelForImageTextToText + +from vllm.assets.video import VideoAsset +from vllm.envs import disable_envs_cache + +from ....conftest import HfRunner, VllmRunner +from ...utils import check_logprobs_close +from .vlm_utils import model_utils + +VIDEO_ASSET = VideoAsset("baby_reading", num_frames=8) + +VIDEO_MODEL_SETTINGS: dict[str, dict[str, Any]] = { + "Qwen/Qwen3-VL-2B-Instruct": { + "prompt": ( + "<|im_start|>user\n" + "<|vision_start|><|video_pad|><|vision_end|>" + "Describe this video.<|im_end|>\n" + "<|im_start|>assistant\n" + ), + }, + "llava-hf/llava-onevision-qwen2-0.5b-ov-hf": { + "prompt": ( + "<|im_start|>user