From 7f4a1b7e246c5df89077c65831f349f698aa2b73 Mon Sep 17 00:00:00 2001 From: stefankoncarevic Date: Sat, 22 Aug 2026 03:26:19 +0200 Subject: [PATCH 001/235] [Bugfix][ROCM] Fix the MXFP8 block scale exponent (#53110) Signed-off-by: Stefan Koncarevic --- .../layers/quantization/utils/mxfp8_utils.py | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) diff --git a/vllm/model_executor/layers/quantization/utils/mxfp8_utils.py b/vllm/model_executor/layers/quantization/utils/mxfp8_utils.py index 1e45ea9ab03b..0c9324a5490a 100644 --- a/vllm/model_executor/layers/quantization/utils/mxfp8_utils.py +++ b/vllm/model_executor/layers/quantization/utils/mxfp8_utils.py @@ -107,6 +107,8 @@ def _kernel( ssm, ssk, BLOCK_M: tl.constexpr, + FP8_MAX: tl.constexpr, + TINY: tl.constexpr, ): pid_m = tl.program_id(0) pid_b = tl.program_id(1) # which 32-element block along K @@ -118,8 +120,11 @@ def _kernel( mask=m_mask[:, None], other=0.0, ).to(tl.float32) - amax = tl.maximum(tl.max(tl.abs(x), axis=1), 1e-30) # [BLOCK_M] - sb = tl.floor(tl.log2(amax)) + 127.0 + # Mirror _mxfp8_e4m3_quantize_torch: the scale has to put the block amax + # at the top of the e4m3 range rather than at 1.0, or small elements of + # the block end up in the subnormals. + amax = tl.maximum(tl.max(tl.abs(x), axis=1), TINY) # [BLOCK_M] + sb = tl.ceil(tl.log2(amax / FP8_MAX)) + 127.0 sb = tl.minimum(tl.maximum(sb, 0.0), 254.0) descale = tl.exp2(sb - 127.0) xq = (x / descale[:, None]).to(xq_ptr.dtype.element_ty) @@ -167,6 +172,8 @@ def _mxfp8_e4m3_quantize_triton( scales.stride(0), scales.stride(1), BLOCK_M=BLOCK_M, + FP8_MAX=float(torch.finfo(MXFP8_VALUE_DTYPE).max), + TINY=float(torch.finfo(torch.float32).tiny), ) return xq, scales From b2db227a7c4c5e55f85524a094b607e5d27408b4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=E2=84=8D=F0=9D=95=A0=F0=9D=95=9D=F0=9D=95=9D=F0=9D=95=A0?= =?UTF-8?q?=F0=9D=95=A8=20=F0=9D=95=84=F0=9D=95=92=F0=9D=95=9F?= Date: Sat, 22 Aug 2026 05:13:53 +0300 Subject: [PATCH 002/235] [Bugfix][R3] Unwrap UniformTypeKVCacheSpecs when selecting the routed-experts KV group (#53240) Signed-off-by: Hollow Man --- .../test_routed_experts_capture.py | 53 ++++++++++++++++--- tests/v1/core/test_kv_cache_utils.py | 39 ++++++++++++++ .../fused_moe/routed_experts_capturer.py | 4 +- vllm/v1/attention/backends/flashinfer.py | 9 +--- vllm/v1/kv_cache_interface.py | 37 ++++++++++--- 5 files changed, 119 insertions(+), 23 deletions(-) diff --git a/tests/model_executor/test_routed_experts_capture.py b/tests/model_executor/test_routed_experts_capture.py index 21b7ba4239b6..b33b62cb4a29 100644 --- a/tests/model_executor/test_routed_experts_capture.py +++ b/tests/model_executor/test_routed_experts_capture.py @@ -25,6 +25,9 @@ FullAttentionSpec, KVCacheConfig, KVCacheGroupSpec, + MLAAttentionSpec, + SlidingWindowMLASpec, + UniformTypeKVCacheSpecs, ) pytestmark = pytest.mark.cpu_test @@ -321,15 +324,17 @@ def test_routed_experts_capturer_dp_unexpected_batch_raises(): assert capturer.device_buffer[0, 0, 0].item() == -1 -def test_routed_experts_attention_group_is_shared_and_fail_closed(monkeypatch): - class FullAttentionSpec: - pass - - monkeypatch.setattr(f"{_REC_MODULE}.FullAttentionSpec", FullAttentionSpec) +def test_routed_experts_attention_group_is_shared_and_fail_closed(): + """Both sides key routing data by this gid, so it must skip non-full-attention + groups rather than defaulting to 0, and fail closed when none exists.""" + common = dict(num_kv_heads=1, head_size=1, dtype=torch.float32) config = SimpleNamespace( kv_cache_groups=[ - SimpleNamespace(kv_cache_spec=object()), - SimpleNamespace(kv_cache_spec=FullAttentionSpec()), + KVCacheGroupSpec( + ["swa_layer"], + SlidingWindowMLASpec(block_size=4, sliding_window=8, **common), + ), + KVCacheGroupSpec(["full_layer"], FullAttentionSpec(block_size=4, **common)), ] ) assert get_routed_experts_attn_gid(config) == 1 @@ -338,6 +343,40 @@ class FullAttentionSpec: get_routed_experts_attn_gid(SimpleNamespace(kv_cache_groups=[])) +def test_routed_experts_attention_group_unwraps_uniform_type_specs(): + """DeepSeek-V4-shaped groups wrap their specs in ``UniformTypeKVCacheSpecs``. + + The wrapper is not a ``FullAttentionSpec``, so a bare isinstance check finds + no group and fails closed on every worker. Unwrap semantics themselves are + covered by ``test_is_full_attention_spec_*`` in tests/v1/core. + """ + common = dict(num_kv_heads=1, head_size=1, dtype=torch.float32) + swa_spec = SlidingWindowMLASpec(block_size=4, sliding_window=8, **common) + mla_spec = MLAAttentionSpec(block_size=4, **common) + config = SimpleNamespace( + kv_cache_groups=[ + KVCacheGroupSpec( + ["swa_layer"], + UniformTypeKVCacheSpecs( + block_size=4, kv_cache_specs={"swa_layer": swa_spec} + ), + ), + KVCacheGroupSpec( + ["mla_layer"], + UniformTypeKVCacheSpecs( + block_size=4, kv_cache_specs={"mla_layer": mla_spec} + ), + ), + ] + ) + + assert get_routed_experts_attn_gid(config) == 1 + + swa_only = SimpleNamespace(kv_cache_groups=[config.kv_cache_groups[0]]) + with pytest.raises(ValueError, match="requires a full-attention KV cache group"): + get_routed_experts_attn_gid(swa_only) + + @pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required") def test_mrv2_async_output_returns_existing_routed_experts_field(): from vllm.v1.outputs import ModelRunnerOutput, RoutedExpertsTensors diff --git a/tests/v1/core/test_kv_cache_utils.py b/tests/v1/core/test_kv_cache_utils.py index 870c139ae0e9..39f18c6d4f7c 100644 --- a/tests/v1/core/test_kv_cache_utils.py +++ b/tests/v1/core/test_kv_cache_utils.py @@ -61,6 +61,8 @@ UniformTypeKVCacheSpecs, get_kv_cache_spec_kind, get_kv_cache_spec_sliding_window, + is_full_attention_spec, + iter_layer_specs, ) from vllm.v1.metrics.stats import CachingMetrics, PrefixCacheStats from vllm.v1.request import Request @@ -3179,3 +3181,40 @@ def test_check_enough_kv_cache_memory_reserves_null_block(): check_enough_kv_cache_memory( vllm_config, {"layer1": spec}, spec.page_size_bytes * 33 ) + + +def test_is_full_attention_spec_unwraps_uniform_type_specs(): + """``UniformTypeKVCacheSpecs`` is not a ``FullAttentionSpec``, so callers + scanning groups with a bare isinstance miss DeepSeek-V4-shaped configs.""" + common = dict(num_kv_heads=1, head_size=1, dtype=torch.float32) + full = FullAttentionSpec(block_size=4, **common) + mla = MLAAttentionSpec(block_size=4, **common) + swa = SlidingWindowMLASpec(block_size=4, sliding_window=8, **common) + + def wrap(**specs): + return UniformTypeKVCacheSpecs(block_size=4, kv_cache_specs=specs) + + assert is_full_attention_spec(full) + assert is_full_attention_spec(mla) + assert not is_full_attention_spec(swa) + + assert is_full_attention_spec(wrap(a=mla, b=mla)) + assert not is_full_attention_spec(wrap(a=swa, b=swa)) + + # Every layer must qualify: a group holding a recycling sliding-window layer + # has no stable slot layout. Grouping forbids mixing today, but the wrapper + # is also built directly (e.g. projecting groups onto a PP worker). + assert not is_full_attention_spec(wrap(a=mla, b=swa)) + assert not is_full_attention_spec(wrap()) + + +def test_iter_layer_specs_returns_group_members(): + common = dict(num_kv_heads=1, head_size=1, dtype=torch.float32) + full = FullAttentionSpec(block_size=4, **common) + mla = MLAAttentionSpec(block_size=4, **common) + + assert list(iter_layer_specs(full)) == [full] + wrapped = UniformTypeKVCacheSpecs( + block_size=4, kv_cache_specs={"a": full, "b": mla} + ) + assert list(iter_layer_specs(wrapped)) == [full, mla] diff --git a/vllm/model_executor/layers/fused_moe/routed_experts_capturer.py b/vllm/model_executor/layers/fused_moe/routed_experts_capturer.py index fa57ea383f83..edd3e3beaa19 100644 --- a/vllm/model_executor/layers/fused_moe/routed_experts_capturer.py +++ b/vllm/model_executor/layers/fused_moe/routed_experts_capturer.py @@ -17,7 +17,7 @@ from vllm.distributed.parallel_state import get_tp_group from vllm.forward_context import get_forward_context from vllm.platforms import current_platform -from vllm.v1.kv_cache_interface import FullAttentionSpec, KVCacheConfig +from vllm.v1.kv_cache_interface import KVCacheConfig, is_full_attention_spec from vllm.v1.outputs import RoutedExpertsTensors logger = logging.getLogger(__name__) @@ -307,7 +307,7 @@ def capture_fn( def get_routed_experts_attn_gid(kv_cache_config: KVCacheConfig) -> int: """Return the full-attention KV cache group used for routed experts.""" for gid, group in enumerate(kv_cache_config.kv_cache_groups): - if isinstance(group.kv_cache_spec, FullAttentionSpec): + if is_full_attention_spec(group.kv_cache_spec): return gid raise ValueError("Routed-experts capture requires a full-attention KV cache group.") diff --git a/vllm/v1/attention/backends/flashinfer.py b/vllm/v1/attention/backends/flashinfer.py index 792056f61d43..c9aea6cb8349 100755 --- a/vllm/v1/attention/backends/flashinfer.py +++ b/vllm/v1/attention/backends/flashinfer.py @@ -85,7 +85,7 @@ KVCacheLayout, KVCacheSpec, KVQuantMode, - UniformTypeKVCacheSpecs, + iter_layer_specs, ) from vllm.v1.utils import CpuGpuBuffer @@ -981,12 +981,7 @@ def get_cudagraph_support( if is_sm12x and vllm_config.parallel_config.decode_context_parallel_size > 1: return AttentionCGSupport.UNIFORM_SINGLE_TOKEN_DECODE - # For UniformTypeKVCacheSpecs, check all contained specs - kv_specs = ( - kv_cache_spec.kv_cache_specs.values() - if isinstance(kv_cache_spec, UniformTypeKVCacheSpecs) - else [kv_cache_spec] - ) + kv_specs = iter_layer_specs(kv_cache_spec) num_qo_heads = vllm_config.model_config.get_num_attention_heads( vllm_config.parallel_config ) diff --git a/vllm/v1/kv_cache_interface.py b/vllm/v1/kv_cache_interface.py index efc284bcbbe5..b660f4931709 100644 --- a/vllm/v1/kv_cache_interface.py +++ b/vllm/v1/kv_cache_interface.py @@ -1020,6 +1020,35 @@ def max_memory_usage_pages(self, vllm_config: VllmConfig) -> int: ) +def iter_layer_specs(kv_cache_spec: KVCacheSpec) -> Collection[KVCacheSpec]: + """The per-layer specs a KV cache group spec covers. + + ``UniformTypeKVCacheSpecs`` groups keep one spec per layer; every other + spec describes its group on its own. Returns the layer specs either way so + callers do not have to special-case the wrapper. + """ + if isinstance(kv_cache_spec, UniformTypeKVCacheSpecs): + return kv_cache_spec.kv_cache_specs.values() + return (kv_cache_spec,) + + +def is_full_attention_spec(kv_cache_spec: KVCacheSpec) -> bool: + """Whether a KV cache group spec is (or wraps) full attention. + + ``UniformTypeKVCacheSpecs`` is not itself a ``FullAttentionSpec``, so a bare + isinstance check misses groups that carry the wrapper -- DeepSeek-V4's MLA + layers, or any model taking the ``UniformTypeKVCacheSpecs.from_specs`` path. + + Every layer must be full attention: a group holding a recycling + (sliding-window) layer has no stable slot layout, so callers that key data + by slot cannot use it. + """ + layer_specs = iter_layer_specs(kv_cache_spec) + return len(layer_specs) > 0 and all( + isinstance(spec, FullAttentionSpec) for spec in layer_specs + ) + + def get_kv_cache_spec_kind(kv_cache_spec: KVCacheSpec) -> KVCacheSpecKind: if isinstance(kv_cache_spec, UniformTypeKVCacheSpecs): inner_kinds = { @@ -1144,15 +1173,9 @@ def has_mixed_precision_kv_cache(self) -> bool: """Whether attention groups store their KV cache at more than one precision.""" kv_cache_precisions: set[tuple[torch.dtype, KVQuantMode]] = set() for group in self.kv_cache_groups: - group_spec = group.kv_cache_spec - group_specs = ( - list(group_spec.kv_cache_specs.values()) - if isinstance(group_spec, UniformTypeKVCacheSpecs) - else [group_spec] - ) kv_cache_precisions.update( (spec.dtype, spec.kv_quant_mode) - for spec in group_specs + for spec in iter_layer_specs(group.kv_cache_spec) if isinstance(spec, AttentionSpec) ) return len(kv_cache_precisions) > 1 From e9d1398d9edfd90fcc1cf783805240e3effec013 Mon Sep 17 00:00:00 2001 From: Yongye Zhu Date: Fri, 21 Aug 2026 20:23:32 -0700 Subject: [PATCH 003/235] [Bugfix][Kimi K3] Enable deferred MoE finalization before weight loading (#53327) Signed-off-by: Yongye Zhu Co-authored-by: OpenAI Codex --- tests/models/kimi_k3/test_latent_moe_tail.py | 79 +++++++++++++++++++ .../kimi_k3/nvidia/latent_moe_runner.py | 20 ++--- 2 files changed, 90 insertions(+), 9 deletions(-) diff --git a/tests/models/kimi_k3/test_latent_moe_tail.py b/tests/models/kimi_k3/test_latent_moe_tail.py index f5447edbf54b..b956d765ef6d 100644 --- a/tests/models/kimi_k3/test_latent_moe_tail.py +++ b/tests/models/kimi_k3/test_latent_moe_tail.py @@ -1,6 +1,8 @@ # SPDX-License-Identifier: Apache-2.0 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project +from types import SimpleNamespace + import pytest import ray import torch @@ -13,8 +15,13 @@ multi_process_parallel, ) from vllm.distributed import get_tp_group +from vllm.model_executor.layers.fused_moe.experts.trtllm_mxfp4_moe import ( + TrtLlmMxfp4ExpertsMonolithic, +) from vllm.model_executor.layers.fused_moe.moe_output import UnfinalizedMoEOutput +from vllm.model_executor.layers.fused_moe.runner.moe_runner import MoERunner from vllm.model_executor.warmup.cutedsl_warmup import cutedsl_warmup +from vllm.models.kimi_k3.nvidia import latent_moe_runner from vllm.models.kimi_k3.nvidia.ops.latent_moe_tail import KimiK3LatentMoETailOp from vllm.platforms import current_platform @@ -24,6 +31,78 @@ TOP_K = 8 +def test_deferred_finalize_enabled_before_moe_kernel_setup( + monkeypatch: pytest.MonkeyPatch, +) -> None: + class FakeMoEConfig: + tp_size = 8 + dp_size = 1 + ep_size = 1 + pcp_size = 1 + is_sequence_parallel = False + hidden_dim = LATENT_SIZE + hidden_dim_unpadded = LATENT_SIZE + experts_per_token = 16 + defer_moe_finalize = False + defer_moe_finalize_max_num_tokens = -1 + + @property + def use_deferred_moe_finalize(self) -> bool: + return self.defer_moe_finalize + + moe_config = FakeMoEConfig() + quant_method = SimpleNamespace( + experts_cls=TrtLlmMxfp4ExpertsMonolithic, + moe_kernel=None, + ) + norm_weight = torch.empty(LATENT_SIZE, dtype=torch.bfloat16) + transform = SimpleNamespace( + norm=SimpleNamespace(weight=norm_weight, variance_epsilon=EPS), + up_proj=SimpleNamespace( + weight=SimpleNamespace(shape=(HIDDEN_SIZE, LATENT_SIZE)) + ), + ) + + def fake_runner_init(runner, *args, **kwargs) -> None: + runner.moe_config = moe_config + runner.routed_experts = SimpleNamespace(quant_method=quant_method) + runner._shared_experts = object() + runner.routed_output_transform = transform + + initialized_with: dict[str, object] = {} + tail_op = SimpleNamespace(contract=SimpleNamespace(max_num_tokens=128)) + + def fake_tail_initialize(**kwargs): + initialized_with.update(kwargs) + return tail_op + + monkeypatch.setattr(MoERunner, "__init__", fake_runner_init) + monkeypatch.setattr(latent_moe_runner.torch.cuda, "Event", lambda: object()) + monkeypatch.setattr( + latent_moe_runner, + "current_platform", + SimpleNamespace( + is_cuda=lambda: True, + is_device_capability_family=lambda capability: capability == 100, + ), + ) + monkeypatch.setattr( + latent_moe_runner, + "get_current_vllm_config", + lambda: SimpleNamespace( + parallel_config=SimpleNamespace(use_ubatching=False), + model_config=SimpleNamespace(enable_sleep_mode=False), + ), + ) + monkeypatch.setattr(KimiK3LatentMoETailOp, "initialize", fake_tail_initialize) + + latent_moe_runner.LatentMoERunner() + + assert moe_config.defer_moe_finalize + assert moe_config.defer_moe_finalize_max_num_tokens == 128 + assert initialized_with["experts_per_token"] == 16 + + def _make_deferred_routed_output( num_tokens: int, device: torch.device, diff --git a/vllm/models/kimi_k3/nvidia/latent_moe_runner.py b/vllm/models/kimi_k3/nvidia/latent_moe_runner.py index 905b8b16948f..d6c8eca32753 100644 --- a/vllm/models/kimi_k3/nvidia/latent_moe_runner.py +++ b/vllm/models/kimi_k3/nvidia/latent_moe_runner.py @@ -113,16 +113,12 @@ def __init__( ) experts_cls = getattr(self._quant_method, "experts_cls", None) - moe_kernel = self._quant_method.moe_kernel + # The kernel instance is built after weight loading, while the tail + # must register its CuTeDSL warmup units during runner construction. self.moe_config.defer_moe_finalize = ( - ( - experts_cls is TrtLlmMxfp4ExpertsMonolithic - or experts_cls is TrtLlmNvFp4ExpertsMonolithic - ) - and self.moe_config.hidden_dim == self.moe_config.hidden_dim_unpadded - and moe_kernel is not None - and moe_kernel.supports_deferred_moe_finalize() - ) + experts_cls is TrtLlmMxfp4ExpertsMonolithic + or experts_cls is TrtLlmNvFp4ExpertsMonolithic + ) and self.moe_config.hidden_dim == self.moe_config.hidden_dim_unpadded from vllm.models.kimi_k3.nvidia.ops.latent_moe_tail import ( KimiK3LatentMoETailOp, ) @@ -145,6 +141,12 @@ def __init__( if self.moe_config.use_deferred_moe_finalize else -1 ) + if self.moe_config.use_deferred_moe_finalize: + logger.info_once( + "K3 latent-MoE tail fusion with deferred top-k finalization " + "is enabled for up to %d tokens.", + op.contract.max_num_tokens, + ) else: self.moe_config.defer_moe_finalize = False self.moe_config.defer_moe_finalize_max_num_tokens = -1 From 7ca49fbe4bab019e55d57cdc4b7fd3d55c67c1a6 Mon Sep 17 00:00:00 2001 From: Tianyu Guo Date: Sat, 22 Aug 2026 13:11:44 +0800 Subject: [PATCH 004/235] [Refactor][Model Runner V2][Multimodal] Move the encoder-only path out of the shared runner (#53176) Signed-off-by: Tianyu Guo --- tests/v1/worker/test_gpu_warmup_blocks.py | 5 +- vllm/config/vllm.py | 12 +- vllm/v1/core/sched/scheduler.py | 4 +- vllm/v1/worker/gpu/model_runner.py | 36 +----- vllm/v1/worker/gpu/warmup.py | 2 +- vllm/v1/worker/gpu_worker.py | 11 +- vllm/v1/worker/mm_encoder_model_runner.py | 141 ++++++++++++++++++++++ 7 files changed, 169 insertions(+), 42 deletions(-) create mode 100644 vllm/v1/worker/mm_encoder_model_runner.py diff --git a/tests/v1/worker/test_gpu_warmup_blocks.py b/tests/v1/worker/test_gpu_warmup_blocks.py index 7b3c5b011bbd..11be99f2089f 100644 --- a/tests/v1/worker/test_gpu_warmup_blocks.py +++ b/tests/v1/worker/test_gpu_warmup_blocks.py @@ -77,7 +77,6 @@ def _make_runner( num_speculative_steps=num_spec_steps, decode_query_len=num_spec_steps + 1, is_pooling_model=False, - is_encoder_only=False, is_encoder_decoder=False, is_last_pp_rank=True, max_num_reqs=4, @@ -88,7 +87,9 @@ def _make_runner( kv_cache_config=SimpleNamespace( kv_cache_groups=kv_cache_groups, num_blocks=1024 ), - vllm_config=SimpleNamespace(num_lookahead_tokens=num_lookahead_tokens), + vllm_config=SimpleNamespace( + num_lookahead_tokens=num_lookahead_tokens, is_mm_encoder_only=False + ), kv_block_zeroer=None, kv_connector=SimpleNamespace(set_disabled=lambda disabled: None), ) diff --git a/vllm/config/vllm.py b/vllm/config/vllm.py index a3e027617685..acde48a931a2 100644 --- a/vllm/config/vllm.py +++ b/vllm/config/vllm.py @@ -572,7 +572,7 @@ def is_ec_producer_only(self) -> bool: ) @property - def is_encoder_only(self) -> bool: + def is_mm_encoder_only(self) -> bool: mm_config = ( self.model_config.multimodal_config if self.model_config is not None @@ -1782,6 +1782,16 @@ def has_blocked_weights(): # before the HMA check below, which inspects the connector class. self._post_init_kv_transfer_config() + if self.is_mm_encoder_only and self.cache_config.enable_prefix_caching: + # Such an instance publishes encoder embeddings and runs no language + # model, so it holds no KV cache for prefix caching to reuse and its + # coordinator would have no group to manage. + logger.info( + "Disabling prefix caching: this instance runs the " + "multi-modal encoder only." + ) + self.cache_config.enable_prefix_caching = False + # Hybrid KV cache manager (HMA) runtime rules: # - Explicit enable (--no-disable-kv-cache-manager): error if runtime # disables it diff --git a/vllm/v1/core/sched/scheduler.py b/vllm/v1/core/sched/scheduler.py index 26674154c648..b41727f273cc 100644 --- a/vllm/v1/core/sched/scheduler.py +++ b/vllm/v1/core/sched/scheduler.py @@ -102,7 +102,7 @@ def __init__( ) self.structured_output_manager = structured_output_manager self.is_encoder_decoder = vllm_config.model_config.is_encoder_decoder - self.is_encoder_only = vllm_config.is_encoder_only + self.is_mm_encoder_only = vllm_config.is_mm_encoder_only # include_finished_set controls whether a separate set of finished # request ids should be included in the EngineCoreOutputs returned @@ -1904,7 +1904,7 @@ def update_from_output( request.status = RequestStatus.FINISHED_STOPPED stopped = True elif ( - self.is_encoder_only + self.is_mm_encoder_only and request.num_computed_tokens >= request.num_prompt_tokens ): # An encoder instance runs the encoder and publishes the diff --git a/vllm/v1/worker/gpu/model_runner.py b/vllm/v1/worker/gpu/model_runner.py index 6d8a28933f53..07e080a5bb9f 100644 --- a/vllm/v1/worker/gpu/model_runner.py +++ b/vllm/v1/worker/gpu/model_runner.py @@ -69,7 +69,6 @@ ECConnectorOutput, ModelRunnerOutput, RoutedExpertsTensors, - make_empty_encoder_model_runner_output, ) from vllm.v1.worker.block_table import get_block_table_width from vllm.v1.worker.cp_utils import check_attention_cp_compatibility @@ -175,7 +174,6 @@ def __init__(self, vllm_config: VllmConfig, device: torch.device): self.device = device self.dtype = self.model_config.dtype - self.is_encoder_only = vllm_config.is_encoder_only self.kv_cache_dtype = self.dtype if self.cache_config.cache_dtype != "auto": # Quantized KV cache. @@ -499,8 +497,6 @@ def get_encoder_timing_stats(self) -> dict[str, dict[str, float | int]]: return encoder_runner.get_encoder_timing_stats() def get_kv_cache_spec(self): - if self.is_encoder_only: - return {} return get_kv_cache_spec(self.vllm_config) def initialize_kv_cache(self, kv_cache_config: KVCacheConfig) -> None: @@ -646,9 +642,6 @@ def _dummy_run( is_profile: bool = False, **kwargs, ) -> tuple[torch.Tensor | None, torch.Tensor | None]: - if self.is_encoder_only: - empty = torch.empty(0, device=self.device) - return empty, empty if skip_attn and not is_profile: raise ValueError( "skip_attn must only be True for initial memory profiling." @@ -809,12 +802,6 @@ def profile_run(self) -> None: dummy_mm_inputs, mm_budget ) - if self.is_encoder_only: - torch.accelerator.synchronize() - self.reset_encoder_cache() - gc.collect() - return - hidden_states, sample_hidden_states = self._dummy_run( self.max_num_tokens, skip_attn=True, is_profile=True ) @@ -851,9 +838,6 @@ def profile_cudagraph_memory(self) -> int: @torch.inference_mode() def capture_model(self) -> int: - if self.is_encoder_only: - return 0 - assert self.cudagraph_manager is not None capture_encoder = ( self.model_state.supports_mm_inputs @@ -1575,26 +1559,12 @@ def execute_model( with self.ec_connector.maybe_get_output( scheduler_output ) as ec_connector_output: - if self.is_encoder_only: - # Encode and publish, nothing else: this instance runs no - # language model, so the gather inside prepare_inputs_embeds - # would build an inputs_embeds nobody reads -- and it - # raises "Encoder cache miss" for any scheduled item this - # instance did not encode, taking the engine down with it. - self.model_state.execute_mm_encoder(scheduled_encoder_inputs) - else: - inputs_embeds = self.model_state.prepare_inputs_embeds( - scheduled_encoder_inputs, input_batch, self.req_states - ) + inputs_embeds = self.model_state.prepare_inputs_embeds( + scheduled_encoder_inputs, input_batch, self.req_states + ) if inputs_embeds is not None and not requires_raw_input_tokens(self.model): input_ids = None - if self.is_encoder_only: - return ModelRunnerOutput.with_ec_conn_output( - make_empty_encoder_model_runner_output(scheduler_output), - ec_connector_output, - ) - model_inputs = { "input_ids": input_ids, "positions": input_batch.positions, diff --git a/vllm/v1/worker/gpu/warmup.py b/vllm/v1/worker/gpu/warmup.py index 9d0085233743..e3c5e11e034d 100644 --- a/vllm/v1/worker/gpu/warmup.py +++ b/vllm/v1/worker/gpu/warmup.py @@ -209,7 +209,7 @@ def warmup_kernels( We must call the provided worker's execute_model for pipeline parallel coordination. """ - if model_runner.is_encoder_only: + if model_runner.vllm_config.is_mm_encoder_only: return num_spec_steps = model_runner.num_speculative_steps diff --git a/vllm/v1/worker/gpu_worker.py b/vllm/v1/worker/gpu_worker.py index c35e22643f52..05c03af75eac 100644 --- a/vllm/v1/worker/gpu_worker.py +++ b/vllm/v1/worker/gpu_worker.py @@ -423,9 +423,14 @@ def init_device(self): # Construct the model runner if self.use_v2_model_runner: - from vllm.v1.worker.gpu.model_runner import ( - GPUModelRunner as GPUModelRunnerV2, - ) + if self.vllm_config.is_mm_encoder_only: + from vllm.v1.worker.mm_encoder_model_runner import ( + MMEncoderModelRunner as GPUModelRunnerV2, + ) + else: + from vllm.v1.worker.gpu.model_runner import ( # type: ignore[assignment] + GPUModelRunner as GPUModelRunnerV2, + ) # HACK(woosuk): This is a temporary fix to avoid type errors. self.model_runner: GPUModelRunner = GPUModelRunnerV2( # type: ignore diff --git a/vllm/v1/worker/mm_encoder_model_runner.py b/vllm/v1/worker/mm_encoder_model_runner.py new file mode 100644 index 000000000000..2a68ab37a937 --- /dev/null +++ b/vllm/v1/worker/mm_encoder_model_runner.py @@ -0,0 +1,141 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Model runner for instances that only run the multi-modal encoder. + +An encoder-only instance -- `--mm-encoder-only`, or the producer side of +encoder-cache disaggregation -- encodes the multi-modal items and publishes the +embeddings. It runs no language model: no KV cache, no sampler, no CUDA graphs. +""" + +from collections.abc import Iterator +from contextlib import contextmanager +from typing import TYPE_CHECKING, Any + +import torch + +from vllm.config import VllmConfig +from vllm.config.compilation import CUDAGraphMode +from vllm.sequence import IntermediateTensors +from vllm.v1.kv_cache_interface import KVCacheSpec +from vllm.v1.outputs import ( + ModelRunnerOutput, + make_empty_encoder_model_runner_output, +) +from vllm.v1.worker.gpu.cudagraph_utils import BatchExecutionDescriptor +from vllm.v1.worker.gpu.mm.lora import set_active_mm_loras +from vllm.v1.worker.gpu.model_runner import GPUModelRunner + +if TYPE_CHECKING: + from vllm.v1.core.sched.output import SchedulerOutput + + +class MMEncoderModelRunner(GPUModelRunner): + """Encoder-only variant of the V2 GPU model runner.""" + + def __init__(self, vllm_config: VllmConfig, device: torch.device): + super().__init__(vllm_config, device) + assert self.supports_mm_inputs, ( + "An encoder-only instance must serve a multi-modal model." + ) + assert self.dp_size == 1, "An encoder-only instance does not support DP." + + depth = vllm_config.max_concurrent_batches + # Blocking (sleep) events: busy-polling the driver lock can make this + # rank a straggler under contention. + self._input_events = [torch.Event(blocking=True) for _ in range(depth)] + self._input_event_idx = 0 + + @contextmanager + def input_tensor_semaphore(self) -> Iterator[None]: + """Guard the host writes into the pooled input buffers. + + `UvaBufferPool` recycles a slot every `max_concurrent_batches` steps and + the device reads the pooled host buffers in place. A sampling step is + ordered by the device wait in `AsyncOutput.get_output()`; this one never + waits, so entering blocks until the step that last wrote these slots is + done with them and leaving records this step's last write. + """ + idx = self._input_event_idx + input_event = self._input_events[idx] + input_event.synchronize() + try: + yield + finally: + input_event.record() + self._input_event_idx = (idx + 1) % len(self._input_events) + + def get_kv_cache_spec(self) -> dict[str, KVCacheSpec]: + return {} + + def capture_model(self) -> int: + return 0 + + def _dummy_run( + self, *args: Any, **kwargs: Any + ) -> tuple[torch.Tensor, torch.Tensor]: + empty = torch.empty(0, device=self.device) + return empty, empty + + def _dummy_sampler_run(self, hidden_states: torch.Tensor) -> None: + return + + def _dummy_pooler_run(self, hidden_states: torch.Tensor) -> None: + return + + def _no_forward(self, scheduler_output: "SchedulerOutput") -> ModelRunnerOutput: + return self._merge_ec_connector_no_forward( + scheduler_output, self.kv_connector.no_forward(scheduler_output) + ) + + @torch.inference_mode() + def execute_model( + self, + scheduler_output: "SchedulerOutput", + intermediate_tensors: IntermediateTensors | None = None, + dummy_run: bool = False, + skip_attn_for_dummy_run: bool = False, + is_profile: bool = False, + context_len: int = 0, + ) -> ModelRunnerOutput: + assert not dummy_run, "An encoder-only instance runs no dummy batch." + + with self.input_tensor_semaphore(): + self.update_pp_decode_requests() + self.finish_requests(scheduler_output) + self.free_states(scheduler_output) + self.add_requests(scheduler_output) + self.update_requests(scheduler_output) + self.block_tables.apply_staged_writes() + if scheduler_output.total_num_scheduled_tokens == 0: + return self._no_forward(scheduler_output) + + batch_req_state, _ = self.gather_batch_req_state(scheduler_output, False) + assert batch_req_state is not None + # No CUDA graph, and no DP peer to agree a padded shape with. + batch_desc = BatchExecutionDescriptor( + cg_mode=CUDAGraphMode.NONE, + num_tokens=batch_req_state.num_tokens, + num_reqs=None, + ) + self.prepare_inputs(scheduler_output, batch_req_state, batch_desc) + + scheduled_encoder_inputs = scheduler_output.scheduled_encoder_inputs + if self.lora_config is not None: + set_active_mm_loras( + model=self.model, + lora_manager=self.lora_manager, + encoder_cache=self.encoder_cache, + req_id_to_index=self.req_states.req_id_to_index, + lora_state=self.lora_state, + scheduled_encoder_inputs=scheduled_encoder_inputs, + ) + + with self.ec_connector.maybe_get_output( + scheduler_output + ) as ec_connector_output: + self.model_state.execute_mm_encoder(scheduled_encoder_inputs) + + return ModelRunnerOutput.with_ec_conn_output( + make_empty_encoder_model_runner_output(scheduler_output), + ec_connector_output, + ) From 0b19ebcacd0e9578460206035c4b3575361c19c4 Mon Sep 17 00:00:00 2001 From: Cyrus Leung Date: Sat, 22 Aug 2026 16:49:37 +0800 Subject: [PATCH 005/235] [MM] Simplify _apply_hf_processor_main (#53275) Signed-off-by: DarkLight1337 --- docs/contributing/model/multimodal.md | 63 ++++-- docs/design/mm_processing.md | 53 +---- requirements/common.txt | 2 +- .../processing/test_audio_in_video.py | 2 +- .../processing/test_audioflamingo3.py | 4 +- .../multimodal/processing/test_common.py | 22 +- .../processing/test_llava_onevision2.py | 2 +- .../multimodal/processing/test_qwen3_vl.py | 4 +- vllm/model_executor/models/audioflamingo3.py | 47 ++-- vllm/model_executor/models/blip2.py | 18 -- vllm/model_executor/models/chameleon.py | 30 +-- vllm/model_executor/models/cheers.py | 9 - vllm/model_executor/models/cohere2_vision.py | 27 ++- vllm/model_executor/models/cohere_asr.py | 65 ++++-- vllm/model_executor/models/colmodernvbert.py | 23 +- vllm/model_executor/models/colpali.py | 40 ++-- vllm/model_executor/models/deepseek_ocr.py | 29 ++- vllm/model_executor/models/deepseek_ocr2.py | 30 ++- vllm/model_executor/models/deepseek_vl2.py | 28 +-- vllm/model_executor/models/ernie45_vl.py | 39 ++-- vllm/model_executor/models/fireredasr2.py | 45 ++-- vllm/model_executor/models/fireredlid.py | 45 ++-- vllm/model_executor/models/funasr.py | 62 ++++-- vllm/model_executor/models/funaudiochat.py | 29 ++- vllm/model_executor/models/gemma3_mm.py | 25 ++- vllm/model_executor/models/gemma3n_mm.py | 84 +++++--- vllm/model_executor/models/gemma4_mm.py | 66 +++--- vllm/model_executor/models/glm4_1v.py | 72 ++++--- vllm/model_executor/models/glmasr.py | 76 ++++--- vllm/model_executor/models/granite_speech.py | 28 +-- vllm/model_executor/models/h2ovl.py | 2 +- vllm/model_executor/models/hunyuan_vision.py | 22 -- .../models/hyperclovax_vision.py | 104 +++++---- .../models/hyperclovax_vision_v2.py | 28 ++- vllm/model_executor/models/idefics3.py | 45 ++-- vllm/model_executor/models/interns1.py | 76 ++++--- vllm/model_executor/models/internvl.py | 36 ++-- vllm/model_executor/models/jina_vl.py | 28 ++- vllm/model_executor/models/kanana_v.py | 24 ++- vllm/model_executor/models/keye.py | 11 +- vllm/model_executor/models/keye_vl1_5.py | 11 +- vllm/model_executor/models/kimi_audio.py | 26 ++- vllm/model_executor/models/kimi_k25.py | 10 - vllm/model_executor/models/kimi_vl.py | 3 + vllm/model_executor/models/lfm2_vl.py | 33 ++- vllm/model_executor/models/lightonocr.py | 37 ++-- vllm/model_executor/models/llava.py | 24 +-- vllm/model_executor/models/llava_next.py | 3 + vllm/model_executor/models/llava_onevision.py | 55 ++--- .../model_executor/models/llava_onevision2.py | 81 ++++--- vllm/model_executor/models/midashenglm.py | 29 +-- vllm/model_executor/models/mimo_v2_omni.py | 23 +- vllm/model_executor/models/minicpmo.py | 2 +- vllm/model_executor/models/minicpmv.py | 44 ++-- vllm/model_executor/models/mistral3.py | 24 +-- vllm/model_executor/models/mllama4.py | 26 +-- vllm/model_executor/models/molmo.py | 81 +++---- vllm/model_executor/models/molmo2.py | 73 ++++--- vllm/model_executor/models/moondream3.py | 10 - vllm/model_executor/models/moss_audio.py | 25 ++- .../models/moss_transcribe_diarize.py | 33 +-- vllm/model_executor/models/muse_glimmer.py | 22 +- .../model_executor/models/nano_nemotron_vl.py | 27 +-- vllm/model_executor/models/nemotron_parse.py | 16 -- vllm/model_executor/models/ovis.py | 31 +-- vllm/model_executor/models/ovis2_5.py | 36 ++-- vllm/model_executor/models/paddleocr_vl.py | 53 ++--- vllm/model_executor/models/paligemma.py | 17 -- vllm/model_executor/models/phi3v.py | 20 +- vllm/model_executor/models/phi4mm.py | 39 ++-- vllm/model_executor/models/phi4siglip.py | 28 ++- vllm/model_executor/models/pixtral.py | 42 ++-- .../models/qwen2_5_omni_thinker.py | 163 +++++++------- vllm/model_executor/models/qwen2_5_vl.py | 10 - vllm/model_executor/models/qwen2_audio.py | 33 ++- .../models/qwen3_omni_moe_thinker.py | 126 ++++++----- vllm/model_executor/models/qwen3_vl.py | 57 +++-- .../models/transformers/multimodal.py | 131 ++++++----- vllm/model_executor/models/ultravox.py | 91 +++++--- vllm/model_executor/models/unlimited_ocr.py | 2 +- vllm/model_executor/models/voxtral.py | 48 ++--- vllm/model_executor/models/whisper.py | 45 ++-- vllm/models/dots3_note/common/processor.py | 9 +- vllm/models/inkling/common/mm_preprocess.py | 32 +-- vllm/models/kimi_k3/common/mm_preprocess.py | 13 -- .../models/minimax_m3/common/mm_preprocess.py | 27 ++- vllm/multimodal/processing/context.py | 9 +- vllm/multimodal/processing/dummy_inputs.py | 5 +- vllm/multimodal/processing/processor.py | 204 +++++++----------- vllm/transformers_utils/processor.py | 44 ---- 90 files changed, 1810 insertions(+), 1668 deletions(-) diff --git a/docs/contributing/model/multimodal.md b/docs/contributing/model/multimodal.md index 3d16e0e91fbc..36830dd8ea3b 100644 --- a/docs/contributing/model/multimodal.md +++ b/docs/contributing/model/multimodal.md @@ -153,11 +153,13 @@ def get_supported_mm_limits(self) -> Mapping[str, int | None]: ## 3. Specify dummy inputs -Then, inherit [BaseDummyInputsBuilder][vllm.multimodal.processing.BaseDummyInputsBuilder] to construct dummy inputs for -HF processing. The processed outputs are also used for memory profiling. +Then, inherit [BaseDummyInputsBuilder][vllm.multimodal.processing.BaseDummyInputsBuilder] to construct the dummy inputs +that are used for memory profiling. Override the abstract methods [get_dummy_text][vllm.multimodal.processing.BaseDummyInputsBuilder.get_dummy_text] and [get_dummy_mm_data][vllm.multimodal.processing.BaseDummyInputsBuilder.get_dummy_mm_data] to construct dummy inputs. These dummy inputs should result in the worst-case memory usage of the model so that vLLM can reserve the correct amount of memory for it. +Besides profiling, the dummy text is passed to the HF processor by models whose HF processor requires text corresponding to the multi-modal items (see [Multi-modal fields](#multi-modal-fields)). If the HF processor can process multi-modal data without any text, the dummy text is not used during HF processing. + Assuming that the memory usage increases with the number of tokens, the dummy inputs can be constructed to maximize the number of output embeddings, which is the same number as placeholder feature tokens. === "Basic example: LLaVA" @@ -446,37 +448,64 @@ return a schema of the tensors outputted by the HF processor that are related to like in LLaVA, each image's features must be independent of the others (which is also required for prefix caching to work correctly). So, we un-pad each image back to its own size by overriding - [BaseMultiModalProcessor._call_hf_processor][vllm.multimodal.processing.BaseMultiModalProcessor._call_hf_processor]: + [BaseMultiModalProcessor._postprocess_hf_mm_data][vllm.multimodal.processing.BaseMultiModalProcessor._postprocess_hf_mm_data] + to post-process the outputs of the HF processor: ??? code ```python - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + # Mistral3Processor requires text corresponding to the images + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _postprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) + if not mm_data: + return processed_data - pixel_values = processed_outputs.get("pixel_values") + pixel_values = processed_data.get("pixel_values") if pixel_values is not None: # Avoid padding since we need the output for each image to be # independent of other images for the cache to work correctly - image_sizes = processed_outputs["image_sizes"] + image_sizes = processed_data["image_sizes"] assert len(pixel_values) == len(image_sizes) - processed_outputs["pixel_values"] = [ + processed_data["pixel_values"] = [ p[:, :h, :w] for p, (h, w) in zip(pixel_values, image_sizes) ] - return processed_outputs + return processed_data ``` + The default implementation of + [_apply_hf_processor_main][vllm.multimodal.processing.BaseMultiModalProcessor._apply_hf_processor_main] + calls the HF processor on the multi-modal data without passing any text. + If the HF processor instead requires text corresponding to the multi-modal items, + you should override + [_get_hf_processor_text][vllm.multimodal.processing.BaseMultiModalProcessor._get_hf_processor_text] + to return the dummy text from + [BaseDummyInputsBuilder.get_dummy_text][vllm.multimodal.processing.BaseDummyInputsBuilder.get_dummy_text] + like in the example above. If the HF processor expects the multi-modal data + under different keys than those provided by the multi-modal items + (e.g. `audio` instead of `audios`), or requires additional keyword arguments + (e.g. `sampling_rate`), you should override + [_preprocess_hf_mm_data][vllm.multimodal.processing.BaseMultiModalProcessor._preprocess_hf_mm_data]. + If you need to modify the output of the HF processor, + you should override + [_postprocess_hf_mm_data][vllm.multimodal.processing.BaseMultiModalProcessor._postprocess_hf_mm_data]. + If the HF processor applies additional transformations to the prompt tokens + that are not reflected in its multi-modal outputs (e.g. prepending a BOS token), + you should override + [_postprocess_prompt][vllm.multimodal.processing.BaseMultiModalProcessor._postprocess_prompt] + to replicate them. + For even more control over how the HF processor is called, you can override + [_apply_hf_processor_main][vllm.multimodal.processing.BaseMultiModalProcessor._apply_hf_processor_main] + directly. + Since `pixel_values` is now a list with one tensor per image, we can override [_get_mm_fields_config][vllm.multimodal.processing.BaseMultiModalProcessor._get_mm_fields_config] as follows: @@ -680,7 +709,7 @@ Examples: ### Handling prompt updates unrelated to multi-modal data -[_get_prompt_updates][vllm.multimodal.processing.BaseMultiModalProcessor._get_prompt_updates] assumes that each application of prompt update corresponds to one multi-modal item. If the HF processor performs additional processing regardless of how many multi-modal items there are, you should override [_apply_hf_processor_tokens_only][vllm.multimodal.processing.BaseMultiModalProcessor._apply_hf_processor_tokens_only] so that the processed token inputs are consistent with the result of applying the HF processor on text inputs. This is because token inputs bypass the HF processor according to [our design](../../design/mm_processing.md). +[_get_prompt_updates][vllm.multimodal.processing.BaseMultiModalProcessor._get_prompt_updates] assumes that each application of prompt update corresponds to one multi-modal item. If the HF processor performs additional processing on the prompt regardless of how many multi-modal items there are, you should override [_postprocess_prompt][vllm.multimodal.processing.BaseMultiModalProcessor._postprocess_prompt] so that the prompt token IDs are consistent with the result of applying the HF processor on text inputs. This is because token inputs bypass the HF processor according to [our design](../../design/mm_processing.md). Examples: @@ -690,7 +719,7 @@ Examples: ### Custom HF processor -Some models don't define an HF processor class on HF Hub. In that case, you can define a custom HF processor that has the same call signature as HF processors and pass it to [_call_hf_processor][vllm.multimodal.processing.BaseMultiModalProcessor._call_hf_processor]. +Some models don't define an HF processor class on HF Hub. In that case, you can define a custom HF processor that has the same call signature as HF processors and return it from [BaseProcessingInfo.get_hf_processor][vllm.multimodal.processing.BaseProcessingInfo.get_hf_processor]. It is then applied to the multi-modal data via [InputProcessingContext.call_hf_processor][vllm.multimodal.processing.InputProcessingContext.call_hf_processor] inside [_apply_hf_processor_main][vllm.multimodal.processing.BaseMultiModalProcessor._apply_hf_processor_main]. Examples: diff --git a/docs/design/mm_processing.md b/docs/design/mm_processing.md index b2b1ac307eea..a2112a2922ed 100644 --- a/docs/design/mm_processing.md +++ b/docs/design/mm_processing.md @@ -2,7 +2,15 @@ To enable various optimizations in vLLM such as [chunked prefill](../configuration/optimization.md#chunked-prefill) and [prefix caching](../features/automatic_prefix_caching.md), we use [BaseMultiModalProcessor][vllm.multimodal.processing.BaseMultiModalProcessor] to provide the correspondence between placeholder feature tokens (e.g. ``) and multi-modal inputs (e.g. the raw input image) based on the outputs of HF processor. -Here are the main features of [BaseMultiModalProcessor][vllm.multimodal.processing.BaseMultiModalProcessor]: +In vLLM's rendering pipeline (see [BaseRenderer][vllm.renderers.base.BaseRenderer]), tokenization is performed as a separate step before multi-modal processing. Therefore, `BaseMultiModalProcessor` needs to recreate the output of calling HF processor end-to-end, while not being able to see the original text. This is achieved through **Dummy Input Text** and **Prompt Update Detection**. + +## Dummy Input Text + +Since Transformers 5.10, `ProcessorMixin` now allows multi-modal inputs to be passed by themselves. However, certain subclasses (such as `ChameleonProcessor`) and older out-of-tree implementations may still define their own `__call__` method that assumes the presence of text with corresponding placeholder tokens. This causes a problem as we don't have the original text anymore to pass to these processors. + +To work around this, each model defines how to generate dummy text based on the number of multi-modal inputs, via [get_dummy_text][vllm.multimodal.processing.BaseDummyInputsBuilder.get_dummy_text], which its override of [_get_hf_processor_text][vllm.multimodal.processing.BaseMultiModalProcessor._get_hf_processor_text] returns so that [_apply_hf_processor_main][vllm.multimodal.processing.BaseMultiModalProcessor._apply_hf_processor_main] passes it to the HF processor together with the multi-modal inputs to obtain the processed multi-modal data. + +Similarly, since the multi-modal data extracted by vLLM may not match what a specific HF processor expects, [_apply_hf_processor_main][vllm.multimodal.processing.BaseMultiModalProcessor._apply_hf_processor_main] allows each model to adapt the inputs via [_preprocess_hf_mm_data][vllm.multimodal.processing.BaseMultiModalProcessor._preprocess_hf_mm_data] (e.g. renaming keys like `audios` to `audio` or injecting extra keyword arguments such as `sampling_rate`) and the outputs via [_postprocess_hf_mm_data][vllm.multimodal.processing.BaseMultiModalProcessor._postprocess_hf_mm_data], without having to reimplement the entire method. ## Prompt Update Detection @@ -13,46 +21,9 @@ One of the main responsibilities of HF processor is to update the prompt with pl The information about which tokens have been updated is key to finding the correspondence between placeholder feature tokens and multi-modal inputs. -In vLLM, this information is specified using [PromptUpdate][vllm.multimodal.processing.PromptUpdate] in [_get_prompt_updates][vllm.multimodal.processing.BaseMultiModalProcessor._get_prompt_updates]. We can automatically detect whether HF has updated the prompt by checking the existence of the updated tokens. - -## Tokenized Prompt Inputs - -To enable tokenization in a separate process, we support passing input token IDs alongside multi-modal data. - -### The problem - -Consider that HF processors follow these main steps: - -1. Tokenize the text -2. Process multi-modal inputs -3. Perform prompt updates - -And we require that: - -- For text + multi-modal inputs, apply all steps 1--3. -- For tokenized + multi-modal inputs, apply only steps 2--3. +Since we call HF processor without the input text, we have to perform this update by ourselves. In vLLM, we represent the necessary information using [PromptUpdate][vllm.multimodal.processing.PromptUpdate] in [_get_prompt_updates][vllm.multimodal.processing.BaseMultiModalProcessor._get_prompt_updates], and apply them via [_apply_prompt_updates][vllm.multimodal.processing.BaseMultiModalProcessor._apply_prompt_updates]. -How can we achieve this without rewriting HF processors? We can try to call the HF processor several times on different inputs: - -- For text + multi-modal inputs, simply call the HF processor directly. -- For tokenized + multi-modal inputs, call the processor only on the multi-modal inputs. - -While HF processors support text + multi-modal inputs natively, this is not so for tokenized + multi-modal inputs: an error is thrown if the number of input placeholder tokens do not correspond to the number of multi-modal inputs. - -Moreover, since the tokenized text has not passed through the HF processor, we have to apply Step 3 by ourselves to keep the output tokens and multi-modal data consistent with each other. - -### Dummy text - -We work around the first issue by requiring each model to define how to generate dummy text based on the number of multi-modal inputs, via [get_dummy_text][vllm.multimodal.processing.BaseDummyInputsBuilder.get_dummy_text]. This lets us generate dummy text corresponding to the multi-modal inputs and input them together to obtain the processed multi-modal data. - -### Automatic prompt updating - -We address the second issue by implementing model-agnostic code in -[_apply_prompt_updates][vllm.multimodal.processing.BaseMultiModalProcessor._apply_prompt_updates] to automatically update the prompt with feature placeholder tokens based on the specification outputted by [_get_prompt_updates][vllm.multimodal.processing.BaseMultiModalProcessor._get_prompt_updates]. - -### Summary - -With the help of dummy text and automatic prompt updating, our multi-modal processor can finally accept both text and token prompts with multi-modal data. The detailed logic is shown in [_apply_hf_processor_main][vllm.multimodal.processing.BaseMultiModalProcessor._apply_hf_processor_main]. +Some HF processors additionally transform the prompt itself regardless of the multi-modal inputs (such as `ChameleonProcessor` appending a sep token for chat mode). Since the prompt tokens likewise bypass the HF processor, such transformations are replicated via [_postprocess_prompt][vllm.multimodal.processing.BaseMultiModalProcessor._postprocess_prompt] before the prompt updates are located or applied. ## Processor Output Caching @@ -60,8 +31,6 @@ Some HF processors, such as the one for Qwen2-VL, are [very slow](https://github When new data is passed in, we first check which items are in the cache, and which ones are missing. The missing items are passed into the HF processor in a single batch and cached, before being merged with the existing items in the cache. -Since we only process the missing multi-modal data items, the number of input placeholder tokens no longer corresponds to the number of the multi-modal inputs, so they can't be passed alongside the text prompt to HF processor. Therefore, we process the text and multi-modal inputs separately, using [dummy text](#dummy-text) to avoid HF errors. Since this skips HF's prompt updating code, we apply [automatic prompt updating](#automatic-prompt-updating) afterwards to keep the output tokens and multi-modal data consistent with each other. - ## Speeding Up Multi‑Modal Data Processing ### Fused Normalisation on the Device diff --git a/requirements/common.txt b/requirements/common.txt index 385dd28dacb1..a1f8dbea3c2b 100644 --- a/requirements/common.txt +++ b/requirements/common.txt @@ -7,7 +7,7 @@ requests >= 2.26.0 tqdm blake3 py-cpuinfo -transformers >= 5.5.3 +transformers >= 5.10.4 huggingface_hub >= 1.28.0 tokenizers >= 0.21.1 # Required for fast incremental detokenization. safetensors >= 0.6.2 # MXFP4/MXFP6 dtype support (F8_E8M0, F4) added in 0.6.0: https://github.com/huggingface/safetensors/pull/611 diff --git a/tests/models/multimodal/processing/test_audio_in_video.py b/tests/models/multimodal/processing/test_audio_in_video.py index 67235cfc8841..df6eea0f3865 100644 --- a/tests/models/multimodal/processing/test_audio_in_video.py +++ b/tests/models/multimodal/processing/test_audio_in_video.py @@ -136,7 +136,7 @@ def test_use_audio_in_video_without_audio_track(model_id: str) -> None: # no audio track" case. (The audio limit above is 1, not 0, so that # `mm_counts["audio"]` is still declared for this request; a 0 limit # would make vLLM omit the "audio" key from `mm_counts` entirely, which - # trips an unrelated assertion in `_apply_hf_processor_mm_only`.) + # trips an unrelated assertion in `_apply_hf_processor_main`.) mm_data = {"video": [video]} with pytest.raises(ValueError, match="doesn't have audio track"): diff --git a/tests/models/multimodal/processing/test_audioflamingo3.py b/tests/models/multimodal/processing/test_audioflamingo3.py index c7dffa90de29..f5920a3be4b8 100644 --- a/tests/models/multimodal/processing/test_audioflamingo3.py +++ b/tests/models/multimodal/processing/test_audioflamingo3.py @@ -105,9 +105,9 @@ def test_audio_chunk_counting(mock_ctx): audio_2 = np.zeros(75 * sr) mm_data = {"audio": [audio_1, audio_2]} - prompt = "<|user|>Listen.<|end|>" - processed = processor._call_hf_processor(prompt, mm_data, {}) + mm_items = processor.info.parse_mm_data(mm_data, validate=False) + processed = processor._apply_hf_processor_main(mm_items, {}) chunk_counts = processed["chunk_counts"] diff --git a/tests/models/multimodal/processing/test_common.py b/tests/models/multimodal/processing/test_common.py index ce79ce243583..7755fc74a886 100644 --- a/tests/models/multimodal/processing/test_common.py +++ b/tests/models/multimodal/processing/test_common.py @@ -385,17 +385,19 @@ def test_processing_correctness( num_batches: int, simplify_rate: float, ): - if model_id == "google/gemma-3n-E2B-it": - pytest.skip("Fix later") - if model_id == "OpenGVLab/InternVL2-2B": - pytest.skip("Fix later") if model_id == "openvla/openvla-7b": pytest.skip( "OpenVLA uses a custom vLLM processor because its HF remote " "processor is incompatible with current Transformers." ) if model_id == "jinaai/jina-reranker-m0": - pytest.skip("Fix later") + pytest.skip( + "JinaVL reverses the multi-modal item order inside " + "_apply_hf_processor_main to match its document-before-query " + "score template, which breaks the positional correspondence " + "between multi-modal kwargs and hashes assumed by the " + "processing cache. Needs a fix that keeps both consistent." + ) if model_id == "mistralai/Voxtral-Mini-4B-Realtime-2602": pytest.skip( "Voxtral Realtime doesn't make use of any place-holder " @@ -403,22 +405,12 @@ def test_processing_correctness( "correctness test as is. Let's revisit adapting this " "test once more realtime models exist." ) - if model_id == "CohereLabs/cohere-transcribe-03-2026": - pytest.skip("Fix later") if model_id.startswith("OpenMOSS-Team/MOSS-Audio-"): pytest.skip( "MOSS-Audio uses a custom processor that dynamically expands " "audio placeholders from processed audio lengths. Its vLLM " "processor paths are covered by test_moss_audio.py." ) - # TODO: Remove when transformers 5.15.0 is released, which contains - # https://github.com/huggingface/transformers/pull/47483. - if model_id == "microsoft/VibeVoice-ASR-HF": - pytest.skip( - "VibeVoice ASR requires audio as a positional argument and hence " - "cannot pass the processing correctness test as is. Its generation " - "is covered by test_transformers_audio.py." - ) if model_id == "lmms-lab-encoder/LLaVA-OneVision-2-8B-Instruct": pytest.skip( "LLaVA-OneVision-2 video processing routes frames through custom " diff --git a/tests/models/multimodal/processing/test_llava_onevision2.py b/tests/models/multimodal/processing/test_llava_onevision2.py index 817d9a4f670c..5a84d049fd1d 100644 --- a/tests/models/multimodal/processing/test_llava_onevision2.py +++ b/tests/models/multimodal/processing/test_llava_onevision2.py @@ -202,7 +202,7 @@ def test_validate_relative_bare_path_blocked(): # --------------------------------------------------------------------------- # Frame backend marker -> PIL + timestamps (_frame_video_to_pil_and_timestamps) # -# Non-codec videos reach _call_hf_processor as a ``(frames_ndarray, metadata)`` +# Non-codec videos reach _apply_hf_processor_main as a ``(frames_ndarray, metadata)`` # tuple -- produced by the registered ``LlavaOnevision2VideoBackend`` for real # ``video_url`` inputs, or by the dummy-inputs builder during profiling # (``video_needs_metadata=True``). The helper materialises PIL frames and diff --git a/tests/models/multimodal/processing/test_qwen3_vl.py b/tests/models/multimodal/processing/test_qwen3_vl.py index a1ab94ca43af..2465a528efd0 100644 --- a/tests/models/multimodal/processing/test_qwen3_vl.py +++ b/tests/models/multimodal/processing/test_qwen3_vl.py @@ -103,7 +103,7 @@ def test_processor_multi_video( """Verify that multi-video processing produces correct placeholders. This exercises the token-level replacement path in - ``_call_hf_processor`` which avoids the quadratic text-level + ``_apply_hf_processor_main`` which avoids the quadratic text-level prompt expansion. """ ctx = build_model_context( @@ -153,7 +153,7 @@ def test_processor_multi_video_list_kwargs( ``mm_processor_kwargs`` (one ``fps``/``num_frames`` per video) must not crash. - Before the fix, ``_call_hf_processor`` copied the whole kwargs to every + Before the fix, ``_apply_hf_processor_main`` 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``. """ diff --git a/vllm/model_executor/models/audioflamingo3.py b/vllm/model_executor/models/audioflamingo3.py index c5583b669b2c..43e6212c7aee 100644 --- a/vllm/model_executor/models/audioflamingo3.py +++ b/vllm/model_executor/models/audioflamingo3.py @@ -377,35 +377,41 @@ def _parse_audio_data( class AudioFlamingo3MultiModalProcessor( BaseMultiModalProcessor[AudioFlamingo3ProcessingInfo] ): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _preprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, Any], - ) -> BatchFeature: - processor_mm_data = dict(mm_data) - audios = processor_mm_data.pop("audios", None) + hf_processor_mm_kwargs: Mapping[str, object], + ) -> tuple[Mapping[str, object], Mapping[str, object]]: + mm_data = dict(mm_data) + audios = mm_data.pop("audios", None) if audios is not None: - processor_mm_data["audio"] = audios + mm_data["audio"] = audios - outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=processor_mm_data, - mm_kwargs=mm_kwargs, - ) + return mm_data, hf_processor_mm_kwargs - if "input_features_mask" in outputs: - outputs["feature_attention_mask"] = outputs.pop("input_features_mask") + def _postprocess_hf_mm_data( + self, + mm_data: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, + ) -> BatchFeature: + if "input_features_mask" in processed_data: + processed_data["feature_attention_mask"] = processed_data.pop( + "input_features_mask" + ) - audio_data = processor_mm_data.get("audio") + audio_data = mm_data.get("audio") if audio_data is None: - return outputs + return processed_data audio_list = audio_data if isinstance(audio_data, list) else [audio_data] if len(audio_list) == 0: - return outputs + return processed_data - processor = self.info.get_hf_processor(**mm_kwargs) + processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) feature_extractor = processor.feature_extractor sampling_rate = feature_extractor.sampling_rate chunk_length = feature_extractor.chunk_length @@ -423,8 +429,9 @@ def _call_hf_processor( n_win = max_windows chunk_counts.append(n_win) - outputs["chunk_counts"] = torch.tensor(chunk_counts, dtype=torch.long) - return outputs + processed_data["chunk_counts"] = torch.tensor(chunk_counts, dtype=torch.long) + + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/blip2.py b/vllm/model_executor/models/blip2.py index 6b1be9860a55..cc7fcd335d98 100644 --- a/vllm/model_executor/models/blip2.py +++ b/vllm/model_executor/models/blip2.py @@ -467,24 +467,6 @@ def get_dummy_mm_data( class Blip2MultiModalProcessor(BaseMultiModalProcessor[Blip2ProcessingInfo]): - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - if not mm_data: - # HF processor always adds placeholders even when there's no image - tokenizer = self.info.get_tokenizer() - prompt_ids = tokenizer.encode(prompt) - return BatchFeature(dict(input_ids=[prompt_ids]), tensor_type="pt") - - return super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) - def _get_mm_fields_config( self, hf_inputs: BatchFeature, diff --git a/vllm/model_executor/models/chameleon.py b/vllm/model_executor/models/chameleon.py index af147b7f1a59..fa9cc13f0b5d 100644 --- a/vllm/model_executor/models/chameleon.py +++ b/vllm/model_executor/models/chameleon.py @@ -135,34 +135,16 @@ def get_dummy_mm_data( class ChameleonMultiModalProcessor(BaseMultiModalProcessor[ChameleonProcessingInfo]): - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - if not mm_data: - prompt_ids = self.info.get_tokenizer().encode(prompt) - prompt_ids = self._apply_hf_processor_tokens_only(prompt_ids) - return BatchFeature(dict(input_ids=[prompt_ids]), tensor_type="pt") - - return super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + # ChameleonProcessor requires text corresponding to the images + return self.dummy_inputs.get_dummy_text(mm_counts) - def _apply_hf_processor_tokens_only( - self, - prompt_tokens: list[int], - ) -> list[int]: + def _postprocess_prompt(self, prompt: list[int]) -> list[int]: # HF processor adds sep token for chat mode tokenizer = self.info.get_tokenizer() - vocab = tokenizer.get_vocab() - - sep_token_id = vocab[tokenizer.sep_token] # type: ignore + sep_token_id = tokenizer.get_vocab()[tokenizer.sep_token] # type: ignore - return prompt_tokens + [sep_token_id] + return [*prompt, sep_token_id] def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/cheers.py b/vllm/model_executor/models/cheers.py index b9e88e10558a..2b934f38675f 100644 --- a/vllm/model_executor/models/cheers.py +++ b/vllm/model_executor/models/cheers.py @@ -16,7 +16,6 @@ import torch.nn as nn import torch.nn.functional as F from einops import rearrange -from transformers import BatchFeature from vllm.config import VllmConfig from vllm.config.multimodal import BaseDummyOptions @@ -492,14 +491,6 @@ def get_dummy_mm_data( class CheersMultiModalProcessor(BaseMultiModalProcessor[CheersProcessingInfo]): """Multimodal processor for Cheers model.""" - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - return super()._call_hf_processor(prompt, mm_data, mm_kwargs) - def _get_prompt_updates( self, mm_items: MultiModalDataItems, diff --git a/vllm/model_executor/models/cohere2_vision.py b/vllm/model_executor/models/cohere2_vision.py index 147d462b23c6..6c13203b2739 100644 --- a/vllm/model_executor/models/cohere2_vision.py +++ b/vllm/model_executor/models/cohere2_vision.py @@ -222,24 +222,23 @@ def get_dummy_mm_data( class Cohere2VisionMultiModalProcessor( BaseMultiModalProcessor[Cohere2VisionProcessingInfo] ): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _postprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: - processed_outputs = super()._call_hf_processor( - prompt, - mm_data, - mm_kwargs, - ) + if not mm_data: + return processed_data - # Ensure num_patches is available for proper tensor splitting if ( - "num_patches" not in processed_outputs + "num_patches" not in processed_data and (images := mm_data.get("images")) is not None ): - hf_processor = self.info.get_hf_processor(**mm_kwargs) + hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) # Fallback calculation if HF processor didn't provide num_patches mm_items = self.info.parse_mm_data({"image": images}, validate=False) @@ -250,13 +249,13 @@ def _call_hf_processor( image_width=parsed_images.get_image_size(i).width, image_height=parsed_images.get_image_size(i).height, processor=hf_processor, - mm_kwargs=mm_kwargs, + mm_kwargs=hf_processor_mm_kwargs, ) for i in range(len(parsed_images)) ] - processed_outputs["num_patches"] = torch.tensor(num_patches) + processed_data["num_patches"] = torch.tensor(num_patches) - return processed_outputs + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/cohere_asr.py b/vllm/model_executor/models/cohere_asr.py index 9cbf9bb205c5..b4e72d5b3f68 100644 --- a/vllm/model_executor/models/cohere_asr.py +++ b/vllm/model_executor/models/cohere_asr.py @@ -8,7 +8,7 @@ import torch import torch.nn.functional as F from torch import nn -from transformers import PretrainedConfig +from transformers import BatchFeature, PretrainedConfig from vllm.compilation.decorators import support_torch_compile from vllm.config import CacheConfig, ModelConfig, SpeechToTextConfig, VllmConfig @@ -45,9 +45,12 @@ BaseDummyInputsBuilder, BaseProcessingInfo, EncDecMultiModalProcessor, + ProcessorInputs, PromptReplacement, PromptUpdate, + TimingContext, ) +from vllm.multimodal.processing.processor import MultiModalProcessingInfo from vllm.renderers import TokenizeParams from vllm.tokenizers import cached_tokenizer_from_config from vllm.transformers_utils.processors.cohere_asr import ( @@ -1941,27 +1944,51 @@ def create_encoder_prompt( ) -> list[int]: return [0] - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _preprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ): - if mm_data: - feature_extractor = self.info.get_feature_extractor(**mm_kwargs) - mm_data = dict(audio=dict(mm_data).pop("audios")) - mm_kwargs = dict( - **mm_kwargs, - sampling_rate=feature_extractor.sampling_rate, - ) - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, + hf_processor_mm_kwargs: Mapping[str, object], + ) -> tuple[Mapping[str, object], Mapping[str, object]]: + feature_extractor = self.info.get_feature_extractor(**hf_processor_mm_kwargs) + + mm_data = dict(mm_data) + mm_data["audio"] = mm_data.pop("audios") + + hf_processor_mm_kwargs = dict( + **hf_processor_mm_kwargs, + sampling_rate=feature_extractor.sampling_rate, ) - if "labels" in processed_outputs: - processed_outputs["input_ids"] = processed_outputs.pop("labels") - return processed_outputs + + return mm_data, hf_processor_mm_kwargs + + def _postprocess_hf_mm_data( + self, + mm_data: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, + ) -> BatchFeature: + if "labels" in processed_data: + processed_data["input_ids"] = processed_data.pop("labels") + + return processed_data + + def _cached_apply_hf_processor( + self, + inputs: ProcessorInputs, + timing_ctx: TimingContext, + ) -> MultiModalProcessingInfo: + # Dithering injects noise into the extracted features, so the + # feature extractor is not a pure function of its input. Since the + # processing cache assumes that processor outputs are invariant + # across calls, bypass the cache when dithering is active. + preproc = self.info.get_hf_config().preprocessor + if preproc.get("dither", 1e-05) > 0: + return self._apply_hf_processor(inputs, timing_ctx) + + return super()._cached_apply_hf_processor(inputs, timing_ctx) def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/colmodernvbert.py b/vllm/model_executor/models/colmodernvbert.py index ffa10452c529..89e55d40e275 100644 --- a/vllm/model_executor/models/colmodernvbert.py +++ b/vllm/model_executor/models/colmodernvbert.py @@ -154,16 +154,25 @@ def get_dummy_mm_data( class ColModernVBertMultiModalProcessor( BaseMultiModalProcessor[ColModernVBertProcessingInfo], ): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + tokenizer = self.info.get_tokenizer() assert isinstance(tokenizer, HfTokenizer) text_encoding = tokenizer( - prompt, + prompt_text, return_tensors="pt", ) result = BatchFeature(data=dict(text_encoding)) @@ -183,7 +192,9 @@ def _call_hf_processor( ) result.update(image_outputs) - return result + processed_data = result + processed_data.update(passthrough_data) + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/colpali.py b/vllm/model_executor/models/colpali.py index aaf3b307a1f5..f31d63c00632 100644 --- a/vllm/model_executor/models/colpali.py +++ b/vllm/model_executor/models/colpali.py @@ -28,6 +28,7 @@ from vllm.model_executor.layers.pooler.tokwise.heads import TokenPoolerHead from vllm.model_executor.model_loader.weight_utils import default_weight_loader from vllm.multimodal import MULTIMODAL_REGISTRY +from vllm.multimodal.parse import MultiModalDataItems from .interfaces import SupportsLateInteraction from .interfaces_base import default_pooling_type @@ -59,24 +60,31 @@ def get_hf_processor(self, **kwargs: object) -> PaliGemmaProcessor: class ColPaliMultiModalProcessor(PaliGemmaMultiModalProcessor): """Multimodal processor for ColPali.""" - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: - if mm_data: - # The ColPali tokenizer_config.json ships with a small default - # max_length (50) that truncates the 1024 image tokens inserted - # by PaliGemmaProcessor, causing a token-count mismatch. - # vLLM enforces its own max_model_len, so we disable HF - # truncation to keep all image + text tokens intact. - return self.info.ctx.call_hf_processor( - self.info.get_hf_processor(**mm_kwargs), - dict(text=prompt, **mm_data), - dict(**mm_kwargs, truncation=False), - ) - return super()._call_hf_processor(prompt, mm_data, mm_kwargs) + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + + # The ColPali tokenizer_config.json ships with a small default + # max_length (50) that truncates the 1024 image tokens inserted + # by PaliGemmaProcessor, causing a token-count mismatch. + # vLLM enforces its own max_model_len, so we disable HF + # truncation to keep all image + text tokens intact. + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + mm_data, + dict(**hf_processor_mm_kwargs, truncation=False), + ) + processed_data.update(passthrough_data) + return processed_data @default_pooling_type(seq_pooling_type="CLS", tok_pooling_type="ALL") diff --git a/vllm/model_executor/models/deepseek_ocr.py b/vllm/model_executor/models/deepseek_ocr.py index 5365b031f9f3..8365c9069e1f 100644 --- a/vllm/model_executor/models/deepseek_ocr.py +++ b/vllm/model_executor/models/deepseek_ocr.py @@ -287,27 +287,34 @@ def get_dummy_mm_data( class DeepseekOCRMultiModalProcessor( BaseMultiModalProcessor[DeepseekOCRProcessingInfo] ): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + if mm_data: - processed_outputs = self.info.ctx.call_hf_processor( - self.info.get_hf_processor(**mm_kwargs), - dict(prompt=prompt, **mm_data), - mm_kwargs, + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(prompt=prompt_text, **mm_data), + hf_processor_mm_kwargs, ) else: tokenizer = self.info.get_tokenizer() assert isinstance(tokenizer, HfTokenizer) - processed_outputs = tokenizer( - prompt, add_special_tokens=True, return_tensors="pt" + processed_data = tokenizer( + prompt_text, add_special_tokens=True, return_tensors="pt" ) - return processed_outputs + processed_data.update(passthrough_data) + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/deepseek_ocr2.py b/vllm/model_executor/models/deepseek_ocr2.py index d19af4713b6f..0e907a344a87 100644 --- a/vllm/model_executor/models/deepseek_ocr2.py +++ b/vllm/model_executor/models/deepseek_ocr2.py @@ -158,27 +158,35 @@ def get_dummy_mm_data( class DeepseekOCR2MultiModalProcessor( BaseMultiModalProcessor[DeepseekOCR2ProcessingInfo] ): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + if mm_data: - processed_outputs = self.info.ctx.call_hf_processor( - self.info.get_hf_processor(**mm_kwargs), - dict(prompt=prompt, **mm_data), - mm_kwargs, + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(prompt=prompt_text, **mm_data), + hf_processor_mm_kwargs, ) else: tokenizer = self.info.get_tokenizer() assert isinstance(tokenizer, HfTokenizer) - processed_outputs = tokenizer( - prompt, add_special_tokens=True, return_tensors="pt" + processed_data = tokenizer( + prompt_text, add_special_tokens=True, return_tensors="pt" ) - return processed_outputs + processed_data.update(passthrough_data) + + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/deepseek_vl2.py b/vllm/model_executor/models/deepseek_vl2.py index ac83b6028ad4..e77dd3d13f59 100644 --- a/vllm/model_executor/models/deepseek_vl2.py +++ b/vllm/model_executor/models/deepseek_vl2.py @@ -43,7 +43,6 @@ ) from vllm.sequence import IntermediateTensors from vllm.tokenizers import cached_tokenizer_from_config -from vllm.tokenizers.hf import HfTokenizer from vllm.transformers_utils.configs.deepseek_vl2 import ( DeepseekVLV2Config, MlpProjectorConfig, @@ -236,28 +235,23 @@ def get_dummy_mm_data( class DeepseekVL2MultiModalProcessor( BaseMultiModalProcessor[DeepseekVL2ProcessingInfo] ): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _postprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: if not mm_data: - tokenizer = self.info.get_tokenizer() - assert isinstance(tokenizer, HfTokenizer) - return tokenizer(prompt, add_special_tokens=True, return_tensors="pt") - - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) + return processed_data - processed_outputs["num_patches"] = ( - processed_outputs["images_spatial_crop"].prod(-1) + 1 + processed_data["num_patches"] = ( + processed_data["images_spatial_crop"].prod(-1) + 1 ) - return processed_outputs + return processed_data def _get_mm_fields_config( self, @@ -315,7 +309,7 @@ def _cached_apply_hf_processor( self, inputs: ProcessorInputs, timing_ctx: TimingContext, - ) -> tuple[list[int], MultiModalProcessingInfo, bool]: + ) -> MultiModalProcessingInfo: # The processor logic is different for len(images) <= 2 vs > 2 # Since the processing cache assumes that the processor output is # invariant of how many images are passed per prompt, we only diff --git a/vllm/model_executor/models/ernie45_vl.py b/vllm/model_executor/models/ernie45_vl.py index 4b495fae7aaa..a1de2ae54e50 100644 --- a/vllm/model_executor/models/ernie45_vl.py +++ b/vllm/model_executor/models/ernie45_vl.py @@ -1110,21 +1110,20 @@ def _pixel_values_norm( pixel_values = pixel_values.to(hf_config.dtype) return pixel_values - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: - # when the prompt is not empty but the multimodal data is empty, - # directly invoke the tokenizer. - if "images" not in mm_data and "videos" not in mm_data and prompt != "": - tokenizer = self.info.get_tokenizer() - prompt_ids = tokenizer.encode(prompt) - tokenizer_output = BatchFeature( - dict(input_ids=[prompt_ids]), tensor_type="pt" - ) - return tokenizer_output + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if "images" not in mm_data and "videos" not in mm_data: + return BatchFeature(dict(passthrough_data)) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) if "images" not in mm_data: mm_data["images"] = [] @@ -1132,7 +1131,7 @@ def _call_hf_processor( mm_data["videos"] = [] # Check if HF processor supports video metadata - hf_processor = self.info.get_hf_processor(**mm_kwargs) + hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) supports_video_metadata = getattr( hf_processor, "supports_video_metadata", False ) @@ -1149,8 +1148,10 @@ def _call_hf_processor( processor_output = self.info.ctx.call_hf_processor( hf_processor, - dict(text=[prompt], images=mm_data["images"], videos=mm_data["videos"]), - mm_kwargs, + dict( + text=[prompt_text], images=mm_data["images"], videos=mm_data["videos"] + ), + hf_processor_mm_kwargs, ) # Divide the processor_output into two modalities: image and video. @@ -1158,7 +1159,7 @@ def _call_hf_processor( pixel_values = processor_output["images"] if pixel_values is not None: processor_output["images"] = self._pixel_values_norm( - pixel_values, mm_kwargs + pixel_values, hf_processor_mm_kwargs ) for key in list(processor_output.keys()): if processor_output[key] is None: @@ -1184,7 +1185,9 @@ def _call_hf_processor( ] del processor_output["images"] - return processor_output + processed_data = processor_output + processed_data.update(passthrough_data) + return processed_data def _get_prompt_updates( self, diff --git a/vllm/model_executor/models/fireredasr2.py b/vllm/model_executor/models/fireredasr2.py index 670e44cb8642..f652eb7518db 100644 --- a/vllm/model_executor/models/fireredasr2.py +++ b/vllm/model_executor/models/fireredasr2.py @@ -228,27 +228,36 @@ def get_dummy_mm_data( class FireRedASR2MultiModalProcessor( BaseMultiModalProcessor[FireRedASR2ProcessingInfo] ): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _preprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - if mm_data: - feature_extractor = self.info.get_feature_extractor(**mm_kwargs) - mm_data = dict(audio=mm_data.pop("audios")) - mm_kwargs = dict( - **mm_kwargs, - sampling_rate=feature_extractor.sampling_rate, - ) - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, + hf_processor_mm_kwargs: Mapping[str, object], + ) -> tuple[Mapping[str, object], Mapping[str, object]]: + feature_extractor = self.info.get_feature_extractor(**hf_processor_mm_kwargs) + + mm_data = dict(mm_data) + mm_data["audio"] = mm_data.pop("audios") + + hf_processor_mm_kwargs = dict( + **hf_processor_mm_kwargs, + sampling_rate=feature_extractor.sampling_rate, ) - if "labels" in processed_outputs: - processed_outputs["input_ids"] = processed_outputs.pop("labels") - return processed_outputs + + return mm_data, hf_processor_mm_kwargs + + def _postprocess_hf_mm_data( + self, + mm_data: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, + ) -> BatchFeature: + if "labels" in processed_data: + processed_data["input_ids"] = processed_data.pop("labels") + + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/fireredlid.py b/vllm/model_executor/models/fireredlid.py index 32efe3dc0eac..2225d70f9e18 100644 --- a/vllm/model_executor/models/fireredlid.py +++ b/vllm/model_executor/models/fireredlid.py @@ -451,27 +451,36 @@ def create_encoder_prompt( # Dummy encoder prompt for profiling (encoder only processes audio). return [0] - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _preprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - if mm_data: - feature_extractor = self.info.get_feature_extractor(**mm_kwargs) - mm_data = dict(audio=mm_data.pop("audios")) - mm_kwargs = dict( - **mm_kwargs, - sampling_rate=feature_extractor.sampling_rate, - ) - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, + hf_processor_mm_kwargs: Mapping[str, object], + ) -> tuple[Mapping[str, object], Mapping[str, object]]: + feature_extractor = self.info.get_feature_extractor(**hf_processor_mm_kwargs) + + mm_data = dict(mm_data) + mm_data["audio"] = mm_data.pop("audios") + + hf_processor_mm_kwargs = dict( + **hf_processor_mm_kwargs, + sampling_rate=feature_extractor.sampling_rate, ) - if "labels" in processed_outputs: - processed_outputs["input_ids"] = processed_outputs.pop("labels") - return processed_outputs + + return mm_data, hf_processor_mm_kwargs + + def _postprocess_hf_mm_data( + self, + mm_data: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, + ) -> BatchFeature: + if "labels" in processed_data: + processed_data["input_ids"] = processed_data.pop("labels") + + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/funasr.py b/vllm/model_executor/models/funasr.py index 205cb5c0e629..e15cde2d7519 100644 --- a/vllm/model_executor/models/funasr.py +++ b/vllm/model_executor/models/funasr.py @@ -44,9 +44,12 @@ BaseDummyInputsBuilder, BaseMultiModalProcessor, BaseProcessingInfo, + ProcessorInputs, PromptReplacement, PromptUpdate, + TimingContext, ) +from vllm.multimodal.processing.processor import MultiModalProcessingInfo from vllm.transformers_utils.processor import cached_processor_from_config from vllm.transformers_utils.processors.funasr import FunASRFeatureExtractor from vllm.utils.tensor_schema import TensorSchema, TensorShape @@ -747,27 +750,50 @@ def get_dummy_mm_data( class FunASRMultiModalProcessor(BaseMultiModalProcessor[FunASRProcessingInfo]): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _preprocess_hf_mm_data( + self, + mm_data: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + ) -> tuple[Mapping[str, object], Mapping[str, object]]: + feature_extractor = self.info.get_feature_extractor(**hf_processor_mm_kwargs) + + mm_data = dict(mm_data) + mm_data["audio"] = mm_data.pop("audios") + + hf_processor_mm_kwargs = dict( + **hf_processor_mm_kwargs, + sampling_rate=feature_extractor.sampling_rate, + ) + + return mm_data, hf_processor_mm_kwargs + + def _postprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: - if mm_data: - feature_extractor = self.info.get_feature_extractor(**mm_kwargs) - mm_data = dict(audio=mm_data.pop("audios")) - mm_kwargs = dict( - **mm_kwargs, - sampling_rate=feature_extractor.sampling_rate, - ) - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) - if "labels" in processed_outputs: - processed_outputs["input_ids"] = processed_outputs.pop("labels") - return processed_outputs + if "labels" in processed_data: + processed_data["input_ids"] = processed_data.pop("labels") + + return processed_data + + def _cached_apply_hf_processor( + self, + inputs: ProcessorInputs, + timing_ctx: TimingContext, + ) -> MultiModalProcessingInfo: + # Dithering injects noise into the extracted features, so the + # feature extractor is not a pure function of its input. Since the + # processing cache assumes that processor outputs are invariant + # across calls, bypass the cache when dithering is active. + if self.info.get_feature_extractor().dither > 0: + return self._apply_hf_processor(inputs, timing_ctx) + + return super()._cached_apply_hf_processor(inputs, timing_ctx) def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/funaudiochat.py b/vllm/model_executor/models/funaudiochat.py index 255e691df3d4..8a0cf1364c6a 100644 --- a/vllm/model_executor/models/funaudiochat.py +++ b/vllm/model_executor/models/funaudiochat.py @@ -612,20 +612,31 @@ def get_dummy_mm_data( class FunAudioChatMultiModalProcessor( BaseMultiModalProcessor[FunAudioChatProcessingInfo] ): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + tokenizer = self.info.get_tokenizer() - input_ids = torch.tensor([tokenizer.encode(prompt)]) + input_ids = torch.tensor([tokenizer.encode(prompt_text)]) audios = mm_data.get("audios", []) if not audios: - return BatchFeature({"input_ids": input_ids}) + processed_data = BatchFeature({"input_ids": input_ids}) + processed_data.update(passthrough_data) + return processed_data - feature_extractor = self.info.get_feature_extractor(**mm_kwargs) + feature_extractor = self.info.get_feature_extractor(**hf_processor_mm_kwargs) sr = int(feature_extractor.sampling_rate) min_samples = int(getattr(feature_extractor, "n_fft", 400) or 400) @@ -676,7 +687,9 @@ def _call_hf_processor( "feature_exist_mask": torch.ones((len(wavs),), dtype=torch.bool), } - return BatchFeature({"input_ids": input_ids, **mm_inputs}) + processed_data = BatchFeature({"input_ids": input_ids, **mm_inputs}) + processed_data.update(passthrough_data) + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/gemma3_mm.py b/vllm/model_executor/models/gemma3_mm.py index 2176b023fad3..d698483cb806 100644 --- a/vllm/model_executor/models/gemma3_mm.py +++ b/vllm/model_executor/models/gemma3_mm.py @@ -261,39 +261,38 @@ def get_dummy_mm_data( class Gemma3MultiModalProcessor(BaseMultiModalProcessor[Gemma3ProcessingInfo]): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _postprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: - processed_outputs = super()._call_hf_processor( - prompt, - mm_data, - mm_kwargs, - ) + if not mm_data: + return processed_data - # HF processor pops the `num_crops` kwarg, which is needed by vLLM if (images := mm_data.get("images")) is not None: mm_items = self.info.parse_mm_data({"image": images}, validate=False) parsed_images = mm_items.get_items("image", ImageProcessorItems) image_sizes = [ parsed_images.get_image_size(i) for i in range(len(parsed_images)) ] - hf_processor = self.info.get_hf_processor(**mm_kwargs) + hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) num_crops = [ self.info.get_num_crops( image_width=size.width, image_height=size.height, processor=hf_processor, - mm_kwargs=mm_kwargs, + mm_kwargs=hf_processor_mm_kwargs, ) for size in image_sizes ] - processed_outputs["num_patches"] = torch.tensor(num_crops) + 1 + processed_data["num_patches"] = torch.tensor(num_crops) + 1 - return processed_outputs + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/gemma3n_mm.py b/vllm/model_executor/models/gemma3n_mm.py index 787642fa1eb1..a473d9f437fc 100644 --- a/vllm/model_executor/models/gemma3n_mm.py +++ b/vllm/model_executor/models/gemma3n_mm.py @@ -1,5 +1,6 @@ # SPDX-License-Identifier: Apache-2.0 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project +import inspect from collections.abc import Iterable, Mapping, Sequence from typing import Annotated, Any, Literal @@ -257,40 +258,69 @@ def get_dummy_mm_data( class Gemma3nMultiModalProcessor(BaseMultiModalProcessor[Gemma3nProcessingInfo]): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _preprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - # HF Transformers audio processor no longer accepts `audios` key. - # We pop `audios` and replace it with `audio` key to suppress - # the warning. + hf_processor_mm_kwargs: Mapping[str, object], + ) -> tuple[Mapping[str, object], Mapping[str, object]]: + mm_data = dict(mm_data) if "audios" in mm_data: mm_data["audio"] = mm_data.pop("audios") - processed_outputs = super()._call_hf_processor( - prompt, - mm_data, - mm_kwargs, - ) - if "input_features" in processed_outputs: - # Padding enables audio_tower to run in batched mode - processed_outputs["input_features_padded"] = processed_outputs[ - "input_features" - ] + return mm_data, hf_processor_mm_kwargs - # Unpad features here since we need the output of each item to be - # independent of other items for the cache to work correctly - unpadded_features = [ - f[mask] - for f, mask in zip( - processed_outputs["input_features"], - processed_outputs["input_features_mask"], - ) + def _postprocess_hf_mm_data( + self, + mm_data: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, + ) -> BatchFeature: + if "input_features" in processed_data: + # The feature extractor pads every clip to the longest one in + # the batch, so each clip's frame count depends on the other + # clips it was processed with. Trim each clip to the number of + # frames it would have if processed alone so that each item's + # output (and thus its cache entry) is independent of the other + # items in the batch. `batch_audio_features` re-pads the + # per-item features so audio_tower can still run batched. + num_frames = [ + self._get_num_audio_frames_alone(len(audio)) + for audio in mm_data["audio"] ] - processed_outputs["input_features"] = unpadded_features - return processed_outputs + processed_data["input_features_padded"] = [ + features[:n] + for features, n in zip(processed_data.pop("input_features"), num_frames) + ] + processed_data["input_features_mask"] = [ + mask[:n] + for mask, n in zip(processed_data["input_features_mask"], num_frames) + ] + + return processed_data + + def _get_num_audio_frames_alone(self, num_samples: int) -> int: + """ + Get the number of mel frames that the feature extractor produces + for a single audio clip of the given length, independent of batch + padding. + """ + feature_extractor = self.info.get_feature_extractor() + call_params = inspect.signature(type(feature_extractor).__call__).parameters + max_length = call_params["max_length"].default + pad_to_multiple_of = call_params["pad_to_multiple_of"].default + + num_samples = min(num_samples, max_length) + remainder = num_samples % pad_to_multiple_of + if remainder: + num_samples += pad_to_multiple_of - remainder + + frame_size = feature_extractor.frame_length + 1 + num_frames = (num_samples - frame_size) // feature_extractor.hop_length + 1 + + return max(num_frames, 0) def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/gemma4_mm.py b/vllm/model_executor/models/gemma4_mm.py index 15774a9ad2f3..29b9a08aefb8 100644 --- a/vllm/model_executor/models/gemma4_mm.py +++ b/vllm/model_executor/models/gemma4_mm.py @@ -560,13 +560,22 @@ def _get_dummy_videos( class Gemma4MultiModalProcessor(BaseMultiModalProcessor[Gemma4ProcessingInfo]): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: - merged_kwargs = self.info.ctx.get_merged_mm_kwargs(mm_kwargs) + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + + merged_kwargs = self.info.ctx.get_merged_mm_kwargs(hf_processor_mm_kwargs) val, is_top_level_max_soft_tokens = _get_max_soft_tokens(merged_kwargs) if val is not None and val not in _SUPPORTED_SOFT_TOKENS: @@ -611,15 +620,15 @@ def _call_hf_processor( timestamps = [idx / fps for idx in frame_indices] # Process frames as images with max_soft_tokens=70 - video_mm_kwargs = dict(mm_kwargs) + video_mm_kwargs = dict(hf_processor_mm_kwargs) video_mm_kwargs["max_soft_tokens"] = _VIDEO_MAX_SOFT_TOKENS dummy_prompt = ("\t" + processor.image_token) * len(frames) - frame_outputs = super()._call_hf_processor( - prompt=dummy_prompt, - mm_data={"images": frames}, - mm_kwargs=video_mm_kwargs, + frame_outputs = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**video_mm_kwargs), + dict(text=dummy_prompt, **{"images": frames}), + video_mm_kwargs, ) # Remap HF key name @@ -660,7 +669,7 @@ def _call_hf_processor( # split-replace bug where replacement text (which itself # contains <|video|> tokens) collides with later splits. vt = processor.video_token - parts = prompt.split(vt, len(video_replacements)) + parts = prompt_text.split(vt, len(video_replacements)) # NOTE: len(parts) <= len(video_replacements) + 1 parts_with_repl: list[str] = [] @@ -668,7 +677,7 @@ def _call_hf_processor( parts_with_repl.extend([part, repl]) parts_with_repl.extend(parts[len(video_replacements) :]) - prompt = "".join(parts_with_repl) + prompt_text = "".join(parts_with_repl) video_outputs = { "pixel_values_videos": torch.cat(all_video_pixel_values, dim=0), @@ -712,45 +721,46 @@ def _call_hf_processor( # NOTE: This requires a corresponding type annotation on the # HF side (Gemma4ProcessorKwargs.images_kwargs) so that # _merge_kwargs routes max_soft_tokens into images_kwargs. - patched_mm_kwargs = dict(mm_kwargs) + patched_mm_kwargs = dict(hf_processor_mm_kwargs) if val is not None and is_top_level_max_soft_tokens: patched_mm_kwargs["max_soft_tokens"] = val - processed_outputs = super()._call_hf_processor( - prompt, - mm_data, + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**patched_mm_kwargs), + dict(text=prompt_text, **mm_data), patched_mm_kwargs, ) # HF uses 'image_position_ids'; vLLM uses 'pixel_position_ids'. # Remap here to keep a single translation point. - if "image_position_ids" in processed_outputs: - processed_outputs["pixel_position_ids"] = processed_outputs.pop( + if "image_position_ids" in processed_data: + processed_data["pixel_position_ids"] = processed_data.pop( "image_position_ids" ) - if "input_features" in processed_outputs: + if "input_features" in processed_data: # Unpad per-item so each item's cache entry is # self-contained. The batched() field config in # _get_mm_fields_config stacks equal-length items and leaves # the rest as a list; _process_audio_input re-pads that list # to the batch's max length. - masks = processed_outputs["input_features_mask"] + masks = processed_data["input_features_mask"] unpadded_features = [ f[mask] for f, mask in zip( - processed_outputs["input_features"], + processed_data["input_features"], masks, ) ] unpadded_masks = [mask[mask] for mask in masks] - processed_outputs["input_features"] = unpadded_features - processed_outputs["input_features_padded"] = unpadded_features - processed_outputs["input_features_mask"] = unpadded_masks + processed_data["input_features"] = unpadded_features + processed_data["input_features_padded"] = unpadded_features + processed_data["input_features_mask"] = unpadded_masks + + processed_data.update(video_outputs) + processed_data.update(passthrough_data) - # Merge video outputs into the final result - combined_outputs = dict(processed_outputs, **video_outputs) - return BatchFeature(combined_outputs) + return processed_data def _get_mm_fields_config( self, @@ -815,7 +825,7 @@ def get_replacement_image(item_idx: int): image_size = images.get_image_size(item_idx) # Resolve the effective max_soft_tokens by merging # per-prompt kwargs with the config-level defaults, - # consistent with how _call_hf_processor resolves it. + # consistent with how _apply_hf_processor_main resolves it. # Without this merge, a missing per-prompt override # would fall back to vision_cfg.default_output_length # instead of the config's mm_processor_kwargs default. diff --git a/vllm/model_executor/models/glm4_1v.py b/vllm/model_executor/models/glm4_1v.py index 9bcca2f8bffc..a4690e9a47d7 100644 --- a/vllm/model_executor/models/glm4_1v.py +++ b/vllm/model_executor/models/glm4_1v.py @@ -1576,32 +1576,39 @@ def _get_direct_path_inputs( return prepared_data, prepared_kwargs - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: - mm_data = dict(mm_data) + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + if not mm_data: - tokenizer = self.info.get_tokenizer() - prompt_ids = tokenizer.encode(prompt, add_special_tokens=False) - return BatchFeature(dict(input_ids=[prompt_ids]), tensor_type="pt") + return BatchFeature(dict(passthrough_data)) - processor = self.info.get_hf_processor(**mm_kwargs) + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + + mm_data = dict(mm_data) + + processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) use_direct_path = ( not isinstance(processor, Glm4vProcessor) and TRANSFORMERS_WITH_GA ) if use_direct_path: prepared_data, prepared_kwargs = self._get_direct_path_inputs( - mm_data, mm_kwargs + mm_data, hf_processor_mm_kwargs ) - return super()._call_hf_processor( - prompt=prompt, - mm_data=prepared_data, - mm_kwargs=prepared_kwargs, + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**prepared_kwargs), + dict(text=prompt_text, **prepared_data), + prepared_kwargs, ) + processed_data.update(passthrough_data) + return processed_data if ( "videos" in mm_data @@ -1616,7 +1623,7 @@ def _call_hf_processor( video_array, metadata = item # don't update mm_kwargs inplace - video_mm_kwargs = dict(**mm_kwargs) + video_mm_kwargs = dict(**hf_processor_mm_kwargs) video_mm_kwargs["do_sample_frames"] = metadata.get( "do_sample_frames", True ) @@ -1625,10 +1632,13 @@ def _call_hf_processor( video_mm_data["videos"] = [[video_array]] video_mm_data["video_metadata"] = [[_to_video_metadata(metadata)]] - video_outputs = super()._call_hf_processor( - prompt="<|begin_of_video|><|video|><|end_of_video|>", - mm_data=video_mm_data, - mm_kwargs=video_mm_kwargs, + video_outputs = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**video_mm_kwargs), + dict( + text="<|begin_of_video|><|video|><|end_of_video|>", + **video_mm_data, + ), + video_mm_kwargs, ) input_ids = video_outputs.pop("input_ids") if swap_video_frame_tokens: @@ -1636,7 +1646,7 @@ def _call_hf_processor( processor.video_token_id ) video_placeholder = processor.tokenizer.batch_decode(input_ids)[0] - prompt = prompt.replace( + prompt_text = prompt_text.replace( "<|begin_of_video|><|video|><|end_of_video|>", video_placeholder, 1, @@ -1652,20 +1662,20 @@ def _call_hf_processor( video_outputs = dict() swap_video_frame_tokens = False - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(text=prompt_text, **mm_data), + hf_processor_mm_kwargs, ) if swap_video_frame_tokens: - input_ids = processed_outputs["input_ids"] + input_ids = processed_data["input_ids"] input_ids[input_ids == processor.video_token_id] = processor.image_token_id - processed_outputs["input_ids"] = input_ids - combined_outputs = dict( - processed_outputs, - **video_outputs, - ) - return BatchFeature(combined_outputs) + processed_data["input_ids"] = input_ids + + processed_data.update(video_outputs) + processed_data.update(passthrough_data) + + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/glmasr.py b/vllm/model_executor/models/glmasr.py index d90bbb5634da..fcf17c550ca5 100644 --- a/vllm/model_executor/models/glmasr.py +++ b/vllm/model_executor/models/glmasr.py @@ -750,46 +750,40 @@ def _calculate_chunk_counts( chunk_counts.append(min(n_chunks, max_windows)) return chunk_counts - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _preprocess_hf_mm_data( self, - prompt: str, - mm_data: dict[str, object], - mm_kwargs: Mapping[str, Any], - ) -> BatchFeature: - # Normalize input: handle deprecated key and list conversion. + mm_data: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + ) -> tuple[Mapping[str, object], Mapping[str, object]]: + mm_data = dict(mm_data) if "audios" in mm_data: mm_data["audio"] = mm_data.pop("audios") - audio = mm_data.get("audio", []) - audio_list = [audio] if audio and not isinstance(audio, list) else audio - - # Early return for text-only. - if not audio_list: - prompt_ids = self.info.get_tokenizer().encode(prompt) - prompt_ids = self._apply_hf_processor_tokens_only(prompt_ids) - return BatchFeature(dict(input_ids=[prompt_ids]), tensor_type="pt") - - # Handle sampling_rate - feature_extractor = self.info.get_feature_extractor(**mm_kwargs) - mm_kwargs = dict( - **mm_kwargs, + feature_extractor = self.info.get_feature_extractor(**hf_processor_mm_kwargs) + hf_processor_mm_kwargs = dict( + **hf_processor_mm_kwargs, sampling_rate=feature_extractor.sampling_rate, ) - # Call parent method - outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) + return mm_data, hf_processor_mm_kwargs - # Postprocess: rename mask and add chunk counts + def _postprocess_hf_mm_data( + self, + mm_data: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, + ) -> BatchFeature: # Handle different key names from different transformers versions - if "input_features_mask" in outputs: - outputs["feature_attention_mask"] = outputs.pop("input_features_mask") - elif "input_features_mask" not in outputs and "input_features" in outputs: + if "input_features_mask" in processed_data: + processed_data["feature_attention_mask"] = processed_data.pop( + "input_features_mask" + ) + elif "input_features" in processed_data: # If no mask is provided, create one from input_features - input_features = outputs["input_features"] + input_features = processed_data["input_features"] if isinstance(input_features, torch.Tensor): # Create a mask of all ones matching the sequence length mask = torch.ones( @@ -797,18 +791,22 @@ def _call_hf_processor( input_features.shape[-1], dtype=torch.long, ) - outputs["feature_attention_mask"] = mask + processed_data["feature_attention_mask"] = mask - # Get processor for chunk counts calculation - processor = self.info.get_hf_processor(**mm_kwargs) + audio = mm_data.get("audio", []) + audio_list = [audio] if audio and not isinstance(audio, list) else audio + if audio_list: + processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) - # Override chunk counts calculation with GLM-ASR specific logic - chunk_counts = self._calculate_chunk_counts( - audio_list, processor.feature_extractor, processor - ) - outputs["chunk_counts"] = torch.tensor(chunk_counts, dtype=torch.long) + # Override chunk counts calculation with GLM-ASR specific logic + chunk_counts = self._calculate_chunk_counts( + audio_list, processor.feature_extractor, processor + ) + processed_data["chunk_counts"] = torch.tensor( + chunk_counts, dtype=torch.long + ) - return outputs + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/granite_speech.py b/vllm/model_executor/models/granite_speech.py index 626750bd960d..9233e772d7d1 100644 --- a/vllm/model_executor/models/granite_speech.py +++ b/vllm/model_executor/models/granite_speech.py @@ -180,12 +180,14 @@ def get_replacement(item_idx: int): ) ] - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _preprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: + hf_processor_mm_kwargs: Mapping[str, object], + ) -> tuple[Mapping[str, object], Mapping[str, object]]: mm_data = dict(mm_data) audios = mm_data.pop("audios", []) @@ -193,21 +195,23 @@ def _call_hf_processor( # GraniteSpeechFeatureExtractor accepts "audio" mm_data["audio"] = audios - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) + return mm_data, hf_processor_mm_kwargs + def _postprocess_hf_mm_data( + self, + mm_data: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, + ) -> BatchFeature: if "audio" in mm_data: # Calculate the number of audio tokens per entry in the batch; # This is used to split the batch back out after padding. audio_token_index = self.info.get_hf_config().audio_token_index - processed_outputs["audio_embed_sizes"] = ( - processed_outputs["input_ids"] == audio_token_index + processed_data["audio_embed_sizes"] = ( + processed_data["input_ids"] == audio_token_index ).sum(-1) - return processed_outputs + return processed_data class GraniteSpeechDummyInputsBuilder( diff --git a/vllm/model_executor/models/h2ovl.py b/vllm/model_executor/models/h2ovl.py index 40240d3e4ee2..dc63276d8269 100644 --- a/vllm/model_executor/models/h2ovl.py +++ b/vllm/model_executor/models/h2ovl.py @@ -135,7 +135,7 @@ def _cached_apply_hf_processor( self, inputs: ProcessorInputs, timing_ctx: TimingContext, - ) -> tuple[list[int], MultiModalProcessingInfo, bool]: + ) -> MultiModalProcessingInfo: # The processor logic is different for len(images) <= 1 vs > 1 # Since the processing cache assumes that the processor output is # invariant of how many images are passed per prompt, we only diff --git a/vllm/model_executor/models/hunyuan_vision.py b/vllm/model_executor/models/hunyuan_vision.py index 4454ce47c1a6..ccc3f8fee9ac 100644 --- a/vllm/model_executor/models/hunyuan_vision.py +++ b/vllm/model_executor/models/hunyuan_vision.py @@ -731,28 +731,6 @@ def get_dummy_mm_data( class HunYuanVLMultiModalProcessor(BaseMultiModalProcessor[HunYuanVLProcessingInfo]): - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - hf_processor = self.info.get_hf_processor(**mm_kwargs) - # HunYuanVLProcessor requires image placeholders wrapped with start/end tokens. - if mm_data.get("images") is not None and prompt: - img_tok = hf_processor.image_token - wrapped = ( - f"{hf_processor.image_start_token}{img_tok}" - f"{hf_processor.image_end_token}" - ) - if img_tok in prompt and wrapped not in prompt: - prompt = prompt.replace(img_tok, wrapped) - return self.info.ctx.call_hf_processor( - hf_processor, - dict(text=prompt, **mm_data), - mm_kwargs, - ) - def _get_prompt_updates( self, mm_items: MultiModalDataItems, diff --git a/vllm/model_executor/models/hyperclovax_vision.py b/vllm/model_executor/models/hyperclovax_vision.py index e469c0202254..41387ad7e2e8 100644 --- a/vllm/model_executor/models/hyperclovax_vision.py +++ b/vllm/model_executor/models/hyperclovax_vision.py @@ -190,77 +190,71 @@ def get_dummy_mm_data( class HCXVisionMultiModalProcessor(BaseMultiModalProcessor[HCXVisionProcessingInfo]): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + for video_idx, video_arr in enumerate(mm_data.get("videos", [])): if video_arr.dtype != np.uint8: mm_data["videos"][video_idx] = video_arr.astype(np.uint8) - processed_outputs = self.info.ctx.call_hf_processor( - hf_processor=self.info.get_hf_processor(**mm_kwargs), + images = mm_data.get("images") + videos = mm_data.get("videos") + + # batchify input as a single item + processed_data = self.info.ctx.call_hf_processor( + hf_processor=self.info.get_hf_processor(**hf_processor_mm_kwargs), data=dict( - text=prompt, - images=None, - videos=None, + images=None if images is None else [images], + videos=None if videos is None else [videos], ), - ) # text-only - - if len(mm_data) > 0: - images = mm_data.get("images") - videos = mm_data.get("videos") - - # batchify input as a single item - _processed_outputs = self.info.ctx.call_hf_processor( - hf_processor=self.info.get_hf_processor(**mm_kwargs), - data=dict( - text=None, - images=None if images is None else [images], - videos=None if videos is None else [videos], - ), - ) # mm-only + ) - for k, v in _processed_outputs.items(): - if isinstance(v, list) and len(v) > 0: - assert len(v) == 1 - _processed_outputs[k] = v[0] + for k, v in processed_data.items(): + if isinstance(v, list) and len(v) > 0: + assert len(v) == 1 + processed_data[k] = v[0] - if images: - _processed_outputs["image_sizes_images"] = torch.tensor( - _processed_outputs["image_sizes_images"] - ) - _processed_outputs["vision_query_lengths_images"] = torch.tensor( - _processed_outputs["vision_query_lengths_images"] - ) + if images: + processed_data["image_sizes_images"] = torch.tensor( + processed_data["image_sizes_images"] + ) + processed_data["vision_query_lengths_images"] = torch.tensor( + processed_data["vision_query_lengths_images"] + ) - if videos: - _idx_per_video = [ - 0, - *accumulate( - get_num_combined_frames(len(video)) for video in videos - ), + if videos: + _idx_per_video = [ + 0, + *accumulate(get_num_combined_frames(len(video)) for video in videos), + ] + processed_data["pixel_values_videos"] = [ + processed_data["pixel_values_videos"][ + _idx_per_video[i] : _idx_per_video[i + 1] ] - _processed_outputs["pixel_values_videos"] = [ - _processed_outputs["pixel_values_videos"][ + for i in range(len(videos)) + ] + processed_data["vision_query_lengths_videos"] = [ + torch.tensor( + processed_data["vision_query_lengths_videos"][ _idx_per_video[i] : _idx_per_video[i + 1] ] - for i in range(len(videos)) - ] - _processed_outputs["vision_query_lengths_videos"] = [ - torch.tensor( - _processed_outputs["vision_query_lengths_videos"][ - _idx_per_video[i] : _idx_per_video[i + 1] - ] - ) - for i in range(len(videos)) - ] + ) + for i in range(len(videos)) + ] - processed_outputs.update(_processed_outputs) + processed_data.update(passthrough_data) - return processed_outputs + return processed_data def _get_prompt_updates( self, diff --git a/vllm/model_executor/models/hyperclovax_vision_v2.py b/vllm/model_executor/models/hyperclovax_vision_v2.py index e50faf885529..b1d64db9d90c 100644 --- a/vllm/model_executor/models/hyperclovax_vision_v2.py +++ b/vllm/model_executor/models/hyperclovax_vision_v2.py @@ -261,34 +261,44 @@ class HCXVisionV2MultiModalProcessor( ): """Multimodal processor for HyperCLOVAX V2 (32B Think model).""" - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + images = mm_data.get("images") videos = mm_data.get("videos") # Get the HF processor - hf_processor = self.info.get_hf_processor(**mm_kwargs) + hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) # Build data dict for HF processor (images/videos only) # NOTE: We pass the prompt as-is without token normalization. # Token expansion is handled by vLLM via _get_prompt_updates. data: dict[str, object] = dict( - text=prompt, + text=prompt_text, images=images, videos=videos, ) - processed_outputs = self.info.ctx.call_hf_processor( + processed_data = self.info.ctx.call_hf_processor( hf_processor=hf_processor, data=data, - kwargs=mm_kwargs, + kwargs=hf_processor_mm_kwargs, ) + processed_data.update(passthrough_data) - return processed_outputs + return processed_data def _get_prompt_updates( self, diff --git a/vllm/model_executor/models/idefics3.py b/vllm/model_executor/models/idefics3.py index 226bfbff79b6..fbf0e7a7c3f4 100644 --- a/vllm/model_executor/models/idefics3.py +++ b/vllm/model_executor/models/idefics3.py @@ -239,51 +239,58 @@ def get_dummy_mm_data( class Idefics3MultiModalProcessor(BaseMultiModalProcessor[Idefics3ProcessingInfo]): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: - # Text-only input not supported in composite processor - if not (images := mm_data.get("images", [])): - prompt_ids = self.info.get_tokenizer().encode(prompt) - prompt_ids = self._apply_hf_processor_tokens_only(prompt_ids) - return BatchFeature(dict(input_ids=[prompt_ids]), tensor_type="pt") + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) image_processor = self.info.get_hf_processor().image_processor if getattr(image_processor, "backend", "pil") == "pil": - mm_kwargs = {"input_data_format": "channels_last", **mm_kwargs} + hf_processor_mm_kwargs = { + "input_data_format": "channels_last", + **hf_processor_mm_kwargs, + } - processed_outputs = super()._call_hf_processor( - prompt, + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), mm_data, - mm_kwargs, + hf_processor_mm_kwargs, ) + images = mm_data.get("images", []) mm_items = self.info.parse_mm_data({"image": images}, validate=False) parsed_images = mm_items.get_items("image", ImageProcessorItems) image_sizes = [ parsed_images.get_image_size(i) for i in range(len(parsed_images)) ] - hf_processor = self.info.get_hf_processor(**mm_kwargs) + hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) num_patches = [ self.info.get_num_patches( image_width=size.width, image_height=size.height, processor=hf_processor, - mm_kwargs=mm_kwargs, + mm_kwargs=hf_processor_mm_kwargs, ) for size in image_sizes ] - processed_outputs["num_patches"] = torch.tensor(num_patches) + processed_data["num_patches"] = torch.tensor(num_patches) # Remove the extra batch dimension - processed_outputs["pixel_values"].squeeze_(0) - processed_outputs["pixel_attention_mask"].squeeze_(0) + processed_data["pixel_values"].squeeze_(0) + processed_data["pixel_attention_mask"].squeeze_(0) + + processed_data.update(passthrough_data) - return processed_outputs + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/interns1.py b/vllm/model_executor/models/interns1.py index bb3ba6a55620..2248911b2010 100644 --- a/vllm/model_executor/models/interns1.py +++ b/vllm/model_executor/models/interns1.py @@ -332,19 +332,28 @@ def get_dummy_mm_data( class InternS1MultiModalProcessor(BaseMultiModalProcessor[InternS1ProcessingInfo]): """Basic image-only MultiModalProcessor for InternS1-style models.""" - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + mm_data = dict(mm_data) videos = mm_data.pop("videos", []) images = mm_data.pop("images", []) assert isinstance(videos, list) assert isinstance(images, list) - hf_processor = self.info.get_hf_processor(**mm_kwargs) + hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) tokenizer = hf_processor.tokenizer video_token_id = tokenizer.encode( hf_processor.video_token, add_special_tokens=False @@ -352,23 +361,29 @@ def _call_hf_processor( assert len(video_token_id) == 1 video_token_id = video_token_id[0] - prompt = re.sub(hf_processor.image_token, "", prompt) - prompt = re.sub(hf_processor.video_token, "", prompt) + prompt_text = re.sub( + hf_processor.image_token, "", prompt_text + ) + prompt_text = re.sub( + hf_processor.video_token, "", prompt_text + ) image_outputs = {} if images: image_pixel_values = [] for image in images: - processed_outputs = super()._call_hf_processor( - prompt=hf_processor.image_token, - mm_data={"images": image}, - mm_kwargs=mm_kwargs, + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(text=hf_processor.image_token, **{"images": image}), + hf_processor_mm_kwargs, ) - image_pixel_values.append(processed_outputs.pop("pixel_values")) + image_pixel_values.append(processed_data.pop("pixel_values")) - input_ids = processed_outputs.pop("input_ids") + input_ids = processed_data.pop("input_ids") image_placeholder = tokenizer.batch_decode(input_ids)[0] - prompt = prompt.replace("", image_placeholder, 1) + prompt_text = prompt_text.replace( + "", image_placeholder, 1 + ) num_patches = [len(item) for item in image_pixel_values] image_outputs = { @@ -381,18 +396,20 @@ def _call_hf_processor( if videos: video_pixel_values = [] for video in videos: - processed_outputs = super()._call_hf_processor( - prompt=hf_processor.video_token, - mm_data={"videos": video}, - mm_kwargs=mm_kwargs, + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(text=hf_processor.video_token, **{"videos": video}), + hf_processor_mm_kwargs, ) - video_pixel_values.append(processed_outputs.pop("pixel_values")) + video_pixel_values.append(processed_data.pop("pixel_values")) - input_ids = processed_outputs.pop("input_ids") + input_ids = processed_data.pop("input_ids") input_ids[input_ids == hf_processor.image_token_id] = video_token_id video_placeholder = tokenizer.batch_decode(input_ids)[0] - prompt = prompt.replace("", video_placeholder, 1) + prompt_text = prompt_text.replace( + "", video_placeholder, 1 + ) num_frames = [len(item) for item in video_pixel_values] video_outputs = { @@ -401,11 +418,20 @@ def _call_hf_processor( "video_token_id": torch.tensor(video_token_id), } - prompt = re.sub("", hf_processor.image_token, prompt) - prompt = re.sub("", hf_processor.video_token, prompt) - text_outputs = tokenizer(prompt, return_tensors="pt") + prompt_text = re.sub( + "", hf_processor.image_token, prompt_text + ) + prompt_text = re.sub( + "", hf_processor.video_token, prompt_text + ) + text_outputs = tokenizer(prompt_text, return_tensors="pt") + + processed_data.update(text_outputs) + processed_data.update(image_outputs) + processed_data.update(video_outputs) + processed_data.update(passthrough_data) - return BatchFeature({**text_outputs, **image_outputs, **video_outputs}) + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/internvl.py b/vllm/model_executor/models/internvl.py index 3895533439bb..ae7b05a553c5 100644 --- a/vllm/model_executor/models/internvl.py +++ b/vllm/model_executor/models/internvl.py @@ -220,27 +220,21 @@ def get_dummy_mm_data( class BaseInternVLMultiModalProcessor(BaseMultiModalProcessor[_I]): """Basic image-only MultiModalProcessor for InternVL-style models.""" - def _call_hf_processor( + def _postprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) - - hf_processor = self.info.get_hf_processor(**mm_kwargs) + hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) image_token_id = hf_processor.ctx_image_token_id # Since there may be extra tokens in the feature placeholders, # we need to pass the image token ID to the model to select the # tokens to merge from the vision encoder outputs - processed_outputs["image_token_id"] = torch.tensor(image_token_id) + processed_data["image_token_id"] = torch.tensor(image_token_id) - return processed_outputs + return processed_data def _get_image_fields_config(self, hf_inputs: BatchFeature): image_num_patches = hf_inputs.get("image_num_patches", torch.empty(0)) @@ -455,19 +449,23 @@ class InternVLMultiModalProcessor( ): """InternVL MultiModalProcessor extended for video support""" - def _call_hf_processor( + def _postprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: - processed_outputs = super()._call_hf_processor(prompt, mm_data, mm_kwargs) + processed_data = super()._postprocess_hf_mm_data( + mm_data, + hf_processor_mm_kwargs, + processed_data, + ) - hf_processor = self.info.get_hf_processor(**mm_kwargs) + hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) if (video_token_id := hf_processor.ctx_video_token_id) is not None: - processed_outputs["video_token_id"] = torch.tensor(video_token_id) + processed_data["video_token_id"] = torch.tensor(video_token_id) - return processed_outputs + return processed_data def _get_video_fields_config(self, hf_inputs: BatchFeature): video_num_patches = hf_inputs.get("video_num_patches", torch.empty(0)) diff --git a/vllm/model_executor/models/jina_vl.py b/vllm/model_executor/models/jina_vl.py index c1cdf207e329..fdfc53def66f 100644 --- a/vllm/model_executor/models/jina_vl.py +++ b/vllm/model_executor/models/jina_vl.py @@ -12,6 +12,7 @@ from vllm.model_executor.layers.linear import ColumnParallelLinear, RowParallelLinear from vllm.model_executor.layers.pooler import DispatchPooler from vllm.multimodal import MULTIMODAL_REGISTRY +from vllm.multimodal.parse import MultiModalDataItems from vllm.sequence import IntermediateTensors from .interfaces import SupportsCrossEncoding, SupportsMultiModal, SupportsScoreTemplate @@ -54,19 +55,28 @@ def forward(self, x, **kwargs): class JinaVLMultiModalProcessor(Qwen2VLMultiModalProcessor): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: - # NOTE: We should reverse the order of the mm_data because the - # query prompt is placed after the document prompt in the score - # template for JinaVLForRanking model, but in mm_data they are - # stored in the opposite order (query first, then document). + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + for _, value in mm_data.items(): value.reverse() - return super()._call_hf_processor(prompt, mm_data, mm_kwargs) + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + mm_data, + hf_processor_mm_kwargs, + ) + processed_data.update(passthrough_data) + return processed_data @MULTIMODAL_REGISTRY.register_processor( diff --git a/vllm/model_executor/models/kanana_v.py b/vllm/model_executor/models/kanana_v.py index 88999149977e..8e55054d3892 100644 --- a/vllm/model_executor/models/kanana_v.py +++ b/vllm/model_executor/models/kanana_v.py @@ -461,17 +461,21 @@ class KananaVMultiModalProcessor(BaseMultiModalProcessor[KananaVProcessingInfo]) def media_token_id(self) -> int: return self.info.get_hf_config().text_config.eos_token_id + 1 - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: """Run the underlying HF processor on text and image data.""" - # Text-only input is handled as a special case here. + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + if not mm_data or not mm_data.get("images", []): - prompt_ids = self.info.get_tokenizer().encode(prompt) - return BatchFeature(dict(input_ids=[prompt_ids]), tensor_type="pt") + return BatchFeature(dict(passthrough_data)) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) # Images image_inputs = mm_data.get("images", []) @@ -493,7 +497,7 @@ def _call_hf_processor( tokenizer = self.info.get_tokenizer() media_token = tokenizer.convert_ids_to_tokens([self.media_token_id])[0] - prompt_replaced = prompt.replace("", media_token) + prompt_replaced = prompt_text.replace("", media_token) input_ids = tokenizer.encode(prompt_replaced) input_ids = torch.tensor(input_ids) @@ -527,7 +531,9 @@ def _call_hf_processor( image_token_thw=torch.tensor(image_meta["image_token_thw"]), pixel_sizes=torch.tensor(pixel_sizes), ) - return BatchFeature(combined_outputs, tensor_type="pt") + processed_data = BatchFeature(combined_outputs, tensor_type="pt") + processed_data.update(passthrough_data) + return processed_data def _get_prompt_updates( self, diff --git a/vllm/model_executor/models/keye.py b/vllm/model_executor/models/keye.py index 7a504e83b73f..bd7eb79f6f57 100644 --- a/vllm/model_executor/models/keye.py +++ b/vllm/model_executor/models/keye.py @@ -1162,15 +1162,8 @@ class KeyeDummyInputsBuilder(KeyeBaseDummyInputsBuilder[KeyeProcessingInfo]): class KeyeMultiModalProcessor(BaseMultiModalProcessor[KeyeProcessingInfo]): - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - # Override to use the text path instead of token path to use the - # video-specific logic in processing_keye.py - return super()._call_hf_processor(prompt, mm_data, mm_kwargs) + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) def _get_prompt_updates( self, diff --git a/vllm/model_executor/models/keye_vl1_5.py b/vllm/model_executor/models/keye_vl1_5.py index 80a94373cb4a..8625fd9c0188 100644 --- a/vllm/model_executor/models/keye_vl1_5.py +++ b/vllm/model_executor/models/keye_vl1_5.py @@ -378,15 +378,8 @@ def get_supported_mm_limits( class KeyeVL1_5MultiModalProcessor(BaseMultiModalProcessor[KeyeVL1_5ProcessingInfo]): - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - # Override to use the text path instead of token path to use the - # video-specific logic in processing_keye.py - return super()._call_hf_processor(prompt, mm_data, mm_kwargs) + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) def _get_prompt_updates( self, diff --git a/vllm/model_executor/models/kimi_audio.py b/vllm/model_executor/models/kimi_audio.py index 60534658ae6d..71cf389f98f7 100644 --- a/vllm/model_executor/models/kimi_audio.py +++ b/vllm/model_executor/models/kimi_audio.py @@ -37,6 +37,7 @@ DictEmbeddingItems, ModalityData, ModalityDataItems, + MultiModalDataItems, MultiModalDataParser, ) from vllm.multimodal.processing import ( @@ -226,13 +227,20 @@ def _parse_audio_data( class KimiAudioMultiModalProcessor(BaseMultiModalProcessor[KimiAudioProcessingInfo]): """vLLM multi-modal processor wrapper for Kimi-Audio.""" - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: """Call the HuggingFace processor.""" + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + # Convert mm_data format: {'audios': [...]} -> {'audio': ...} mm_data = dict(mm_data) audios = mm_data.pop("audios", []) @@ -252,11 +260,13 @@ def _call_hf_processor( mm_data["audio"] = audio_arrays # Use the context's call_hf_processor for proper handling - return self.info.ctx.call_hf_processor( - self.info.get_hf_processor(**mm_kwargs), - dict(text=prompt, **mm_data), - mm_kwargs, + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + mm_data, + hf_processor_mm_kwargs, ) + processed_data.update(passthrough_data) + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/kimi_k25.py b/vllm/model_executor/models/kimi_k25.py index 3054fcc5d89f..7f313f44ea84 100644 --- a/vllm/model_executor/models/kimi_k25.py +++ b/vllm/model_executor/models/kimi_k25.py @@ -258,16 +258,6 @@ def _get_mm_fields_config( grid_thws=MultiModalFieldConfig.batched("vision_chunk", keep_on_cpu=True), ) - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - # Override to use the text path instead of token path because vision chunk - # is not considered - return super()._call_hf_processor(prompt, mm_data, mm_kwargs) - def _get_prompt_updates( self, mm_items: MultiModalDataItems, diff --git a/vllm/model_executor/models/kimi_vl.py b/vllm/model_executor/models/kimi_vl.py index e5fa22de7e1a..fc7533b192a1 100644 --- a/vllm/model_executor/models/kimi_vl.py +++ b/vllm/model_executor/models/kimi_vl.py @@ -237,6 +237,9 @@ def get_dummy_mm_data( class KimiVLMultiModalProcessor(BaseMultiModalProcessor[KimiVLProcessingInfo]): + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + def _get_mm_fields_config( self, hf_inputs: BatchFeature, diff --git a/vllm/model_executor/models/lfm2_vl.py b/vllm/model_executor/models/lfm2_vl.py index 85286712feb6..e6ae23cca47e 100644 --- a/vllm/model_executor/models/lfm2_vl.py +++ b/vllm/model_executor/models/lfm2_vl.py @@ -393,45 +393,38 @@ def get_dummy_mm_data( class Lfm2VLMultiModalProcessor(BaseMultiModalProcessor[Lfm2VLProcessingInfo]): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _postprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: - # Text-only input not supported in composite processor - if not (images := mm_data.get("images", [])): - prompt_ids = self.info.get_tokenizer().encode( - prompt, add_special_tokens=False - ) - prompt_ids = self._apply_hf_processor_tokens_only(prompt_ids) - return BatchFeature(dict(input_ids=[prompt_ids]), tensor_type="pt") - - processed_outputs = super()._call_hf_processor( - prompt, - mm_data, - mm_kwargs, - ) + if not mm_data: + return processed_data + images = mm_data.get("images", []) mm_items = self.info.parse_mm_data({"image": images}, validate=False) parsed_images = mm_items.get_items("image", ImageProcessorItems) image_sizes = [ parsed_images.get_image_size(i) for i in range(len(parsed_images)) ] - hf_processor = self.info.get_hf_processor(**mm_kwargs) + hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) num_patches = [ self.info.get_num_patches( image_width=size.width, image_height=size.height, processor=hf_processor, - mm_kwargs=mm_kwargs, + mm_kwargs=hf_processor_mm_kwargs, ) for size in image_sizes ] - processed_outputs["num_patches"] = torch.tensor(num_patches) + processed_data["num_patches"] = torch.tensor(num_patches) - return processed_outputs + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/lightonocr.py b/vllm/model_executor/models/lightonocr.py index 83133e400f84..ccb33177fbca 100644 --- a/vllm/model_executor/models/lightonocr.py +++ b/vllm/model_executor/models/lightonocr.py @@ -39,20 +39,19 @@ class LightOnOCRMultiModalProcessor(BaseMultiModalProcessor[Mistral3ProcessingInfo]): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _postprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) + if not mm_data: + return processed_data - # NOTE: LightOnOCR does not use break/end tokens, so we remove them here. - input_ids = processed_outputs.get("input_ids") + input_ids = processed_data.get("input_ids") if input_ids is not None: processor = self.info.get_hf_processor() tokenizer = self.info.get_tokenizer() @@ -67,22 +66,22 @@ def _call_hf_processor( torch.tensor([break_id, end_id]), ) - processed_outputs["input_ids"] = input_ids[keep_mask].unsqueeze(0) - if "attention_mask" in processed_outputs: - processed_outputs["attention_mask"] = processed_outputs[ - "attention_mask" - ][keep_mask].unsqueeze(0) + processed_data["input_ids"] = input_ids[keep_mask].unsqueeze(0) + if "attention_mask" in processed_data: + processed_data["attention_mask"] = processed_data["attention_mask"][ + keep_mask + ].unsqueeze(0) # un-pad pixel_values per-image so caches remain independent. - pixel_values = processed_outputs.get("pixel_values") + pixel_values = processed_data.get("pixel_values") if pixel_values is not None: - image_sizes = processed_outputs["image_sizes"] + image_sizes = processed_data["image_sizes"] assert len(pixel_values) == len(image_sizes) - processed_outputs["pixel_values"] = [ + processed_data["pixel_values"] = [ p[:, :h, :w] for p, (h, w) in zip(pixel_values, image_sizes) ] - return processed_outputs + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/llava.py b/vllm/model_executor/models/llava.py index 3ddb60452e22..c3ec7a4daa6a 100644 --- a/vllm/model_executor/models/llava.py +++ b/vllm/model_executor/models/llava.py @@ -321,30 +321,30 @@ def get_hf_processor(self, **kwargs: object): class PixtralHFMultiModalProcessor(BaseMultiModalProcessor[PixtralHFProcessingInfo]): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _postprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) + if not mm_data: + return processed_data - pixel_values = processed_outputs.get("pixel_values") + pixel_values = processed_data.get("pixel_values") if pixel_values is not None: # Avoid padding since we need the output for each image to be # independent of other images for the cache to work correctly - image_sizes = processed_outputs["image_sizes"] + image_sizes = processed_data["image_sizes"] assert len(pixel_values) == len(image_sizes) - processed_outputs["pixel_values"] = [ + processed_data["pixel_values"] = [ p[:, :h, :w] for p, (h, w) in zip(pixel_values, image_sizes) ] - return processed_outputs + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/llava_next.py b/vllm/model_executor/models/llava_next.py index fec9d6af5ede..cf93da29243a 100644 --- a/vllm/model_executor/models/llava_next.py +++ b/vllm/model_executor/models/llava_next.py @@ -211,6 +211,9 @@ def _get_mm_fields_config( class LlavaNextMultiModalProcessor(BaseLlavaNextMultiModalProcessor[_I]): + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + def _get_mm_fields_config( self, hf_inputs: BatchFeature, diff --git a/vllm/model_executor/models/llava_onevision.py b/vllm/model_executor/models/llava_onevision.py index 7b8896f9e0d1..848372f1f606 100644 --- a/vllm/model_executor/models/llava_onevision.py +++ b/vllm/model_executor/models/llava_onevision.py @@ -321,22 +321,33 @@ def _get_mm_fields_config( pixel_values_videos=MultiModalFieldConfig.batched("video"), ) - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + mm_data = dict(mm_data) videos = mm_data.pop("videos", []) assert isinstance(videos, list) if not videos: - return super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(text=prompt_text, **mm_data), + hf_processor_mm_kwargs, ) + processed_data.update(passthrough_data) + return processed_data # LLaVA-OneVision processor doesn't support multiple videos # with different sizes when converting back to tensors @@ -346,19 +357,13 @@ def _call_hf_processor( image_token = processor.image_token video_token = processor.video_token - text_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data={}, - mm_kwargs=mm_kwargs, - ) - images = mm_data.pop("images", []) assert isinstance(images, list) if images: - processor_outputs = super()._call_hf_processor( - prompt=image_token * len(images), - mm_data={"images": images}, - mm_kwargs=mm_kwargs, + processor_outputs = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(text=image_token * len(images), **{"images": images}), + hf_processor_mm_kwargs, ) image_outputs = { k: v @@ -370,10 +375,10 @@ def _call_hf_processor( pixel_values_videos = [] for video in videos: - item_outputs = super()._call_hf_processor( - prompt=video_token, - mm_data={"videos": video}, - mm_kwargs=mm_kwargs, + item_outputs = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(text=video_token, **{"videos": video}), + hf_processor_mm_kwargs, ) pixel_values_videos.append(item_outputs["pixel_values_videos"][0]) @@ -381,11 +386,13 @@ def _call_hf_processor( video_outputs = {"pixel_values_videos": pixel_values_videos} combined_outputs = dict( - text_outputs, + input_ids=prompt_text, **image_outputs, **video_outputs, ) - return BatchFeature(combined_outputs) + processed_data = BatchFeature(combined_outputs) + processed_data.update(passthrough_data) + return processed_data def _get_prompt_updates( self, diff --git a/vllm/model_executor/models/llava_onevision2.py b/vllm/model_executor/models/llava_onevision2.py index be51b1f0f5f3..5dfb83412887 100644 --- a/vllm/model_executor/models/llava_onevision2.py +++ b/vllm/model_executor/models/llava_onevision2.py @@ -286,7 +286,7 @@ def _validate_video_sources(paths, model_config) -> list[str]: # participates in the standard decode-stage pipeline. # # * codec backend: NOT a loader. OV2's codec path needs the video path string -# to survive into ``_call_hf_processor``, where the HF processor builds the +# to survive into ``_apply_hf_processor_main``, where the HF processor builds the # codec canvas + smart_resize + patchify # (pixel_values/image_grid_thw/patch_positions). That transform is # path-level and inseparable; it cannot be reconstructed from pre-decoded @@ -302,7 +302,7 @@ def prepare_codec_video_input(video_path: str) -> tuple: Returns ``(dummy_ndarray, metadata)`` where the ndarray satisfies the parser's 4-D shape check and the metadata carries the actual path to - our ``_call_hf_processor``. Use as:: + our ``_apply_hf_processor_main``. Use as:: multi_modal_data = {"video": prepare_codec_video_input("foo.mp4")} @@ -1283,7 +1283,7 @@ def get_data_parser(self) -> MultiModalDataParser: # ``video_needs_metadata=True`` makes the parser preserve both the # ``(frames, metadata)`` tuples from the frame backend and the # ``(dummy, {marker: path})`` tuples from prepare_codec_video_input; - # both are dispatched by metadata content in ``_call_hf_processor``. + # both are dispatched by metadata content in ``_apply_hf_processor_main``. return LlavaOnevision2MultiModalDataParser( self.get_hf_config().vision_config.spatial_merge_size, video_needs_metadata=True, @@ -1547,18 +1547,23 @@ def _parse_image_data( class LlavaOnevision2MultiModalProcessor( BaseMultiModalProcessor[LlavaOnevision2ProcessingInfo] ): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: - # The wrapped OV2 processor is a bare custom class without the standard - # ProcessorMixin ``_merge_kwargs`` machinery, so vLLM's default path - # fails; overriding this method routes the base class to call us - # directly. - hf_processor = self.info.get_hf_processor(**mm_kwargs) - merged_kwargs = self.info.ctx.get_merged_mm_kwargs(mm_kwargs) + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + + hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) + merged_kwargs = self.info.ctx.get_merged_mm_kwargs(hf_processor_mm_kwargs) merged_kwargs.setdefault("return_tensors", "pt") call_kwargs = { k: v @@ -1627,21 +1632,25 @@ def _call_hf_processor( if len(codec_video_paths) > 1 else codec_video_paths[0] ) - # Route through the base ``_call_hf_processor`` so float-tensor + # Route through ``ctx.call_hf_processor`` so float-tensor # dtype postprocessing is applied automatically; inject # ``video_backend="codec"`` via mm_kwargs so the wrapped processor # dispatches to its codec branch. - output = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs={**mm_kwargs, "video_backend": "codec"}, + output = self.info.ctx.call_hf_processor( + self.info.get_hf_processor( + **{**hf_processor_mm_kwargs, "video_backend": "codec"} + ), + dict(text=prompt_text, **mm_data), + {**hf_processor_mm_kwargs, "video_backend": "codec"}, ) data = dict(output) - return BatchFeature( + processed_data = BatchFeature( self._rename_codec_outputs_to_video( data, codec_video_paths, hf_processor ) ) + processed_data.update(passthrough_data) + return processed_data # ---- Frame backend (registered LlavaOnevision2VideoBackend) ------ # Every non-codec video reaches here as a ``(frames_ndarray, metadata)`` @@ -1657,7 +1666,9 @@ def _call_hf_processor( # local-file gating is enforced by the connector before decoding. if videos_present: timestamp_decimals = int( - mm_kwargs.get("timestamp_decimals", _DEFAULT_TIMESTAMP_DECIMALS) + hf_processor_mm_kwargs.get( + "timestamp_decimals", _DEFAULT_TIMESTAMP_DECIMALS + ) ) per_video_frames: list[list[Image.Image]] = [] @@ -1671,7 +1682,7 @@ def _call_hf_processor( # ``<{t} seconds><|vision_start|><|image_pad|><|vision_end|>`` # blocks (matches the OV2 hf-chat reference exactly). new_prompt = _expand_video_markers_in_prompt( - prompt, + prompt_text, per_video_timestamps, timestamp_decimals=timestamp_decimals, ) @@ -1700,7 +1711,7 @@ def _call_hf_processor( row_is_video: list[bool] = [] vid_idx = 0 img_idx = 0 - for marker in marker_pattern.finditer(prompt): + for marker in marker_pattern.finditer(prompt_text): if marker.lastgroup == "video": frames = per_video_frames[vid_idx] vid_idx += 1 @@ -1723,15 +1734,15 @@ def _call_hf_processor( merged_mm_data.pop("videos", None) merged_mm_data["images"] = flat_frames - # Route through the base ``_call_hf_processor`` (applies float-tensor + # Route through ``ctx.call_hf_processor`` (applies float-tensor # dtype postprocessing automatically). The wrapped processor's image # branch ignores video/codec-only kwargs and does not forward extra # **kwargs to the image processor, so passing the full merged kwarg # set here is a no-op beyond return_tensors/padding. - output = super()._call_hf_processor( - prompt=new_prompt, - mm_data=merged_mm_data, - mm_kwargs=mm_kwargs, + output = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(text=new_prompt, **merged_mm_data), + hf_processor_mm_kwargs, ) data = dict(output) @@ -1788,17 +1799,21 @@ def _gather_rows(tensor, rows): (len(per_video_timestamps),), dtype=torch.long ) - return BatchFeature(data) + processed_data = BatchFeature(data) + processed_data.update(passthrough_data) + return processed_data # ---- Image-only / text-only call -------------------------------- - # No videos present: delegate to the base ``_call_hf_processor``, which + # No videos present: delegate to ``ctx.call_hf_processor``, which # runs the wrapped processor over the (possibly empty) image set and # applies float-tensor dtype postprocessing automatically. - return super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(text=prompt_text, **mm_data), + hf_processor_mm_kwargs, ) + processed_data.update(passthrough_data) + return processed_data def _rename_codec_outputs_to_video( self, diff --git a/vllm/model_executor/models/midashenglm.py b/vllm/model_executor/models/midashenglm.py index 8ddfd05f7294..9464c3f248a9 100644 --- a/vllm/model_executor/models/midashenglm.py +++ b/vllm/model_executor/models/midashenglm.py @@ -26,7 +26,7 @@ import collections import collections.abc from collections.abc import Callable, Iterable, Mapping, Sequence -from typing import Annotated, Any, TypeAlias, cast +from typing import Annotated, TypeAlias, cast import numpy as np import torch @@ -583,15 +583,17 @@ def get_dummy_mm_data( class MiDashengLMMultiModalProcessor( BaseMultiModalProcessor[MiDashengLMProcessingInfo] ): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _preprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, Any], - ) -> BatchFeature: + hf_processor_mm_kwargs: Mapping[str, object], + ) -> tuple[Mapping[str, object], Mapping[str, object]]: + mm_data = dict(mm_data) audios = mm_data.pop("audios", []) - # + Padding min_audio_len = self.info.get_min_audio_len() processed_audios = [ np.pad( @@ -608,20 +610,7 @@ def _call_hf_processor( if processed_audios: mm_data["audio"] = processed_audios - if not mm_data.get("audio", []): - prompt_ids = self.info.get_tokenizer().encode(prompt) - prompt_ids = self._apply_hf_processor_tokens_only(prompt_ids) - return BatchFeature(dict(input_ids=[prompt_ids]), tensor_type="pt") - - mm_kwargs = dict( - **mm_kwargs, - ) - - return super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) + return mm_data, hf_processor_mm_kwargs def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/mimo_v2_omni.py b/vllm/model_executor/models/mimo_v2_omni.py index 8f78409581f8..bd418d9c093c 100644 --- a/vllm/model_executor/models/mimo_v2_omni.py +++ b/vllm/model_executor/models/mimo_v2_omni.py @@ -912,16 +912,23 @@ def _get_mm_fields_config( fields["va_audio_features"] = MultiModalFieldConfig.batched("va_audio") return fields - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: """Convert numpy video arrays to (TCHW, timestamps) tuples for MiMo. Also remap 'audios' → 'audio' since MiMoOmniProcessor.__call__ uses the singular form. """ + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + # Remap audios → audio (MiMoOmniProcessor uses singular param name) if "audios" in mm_data: mm_data = {**mm_data, "audio": mm_data["audios"]} @@ -999,7 +1006,13 @@ def _call_hf_processor( mm_data = {**mm_data, "videos": converted} - return super()._call_hf_processor(prompt, mm_data, mm_kwargs) + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + mm_data, + hf_processor_mm_kwargs, + ) + processed_data.update(passthrough_data) + return processed_data def _get_prompt_updates( self, diff --git a/vllm/model_executor/models/minicpmo.py b/vllm/model_executor/models/minicpmo.py index 3dd3f8fd3b0c..527117bd0b76 100644 --- a/vllm/model_executor/models/minicpmo.py +++ b/vllm/model_executor/models/minicpmo.py @@ -448,7 +448,7 @@ def process_audios( if isinstance(parsed_audios, MiniCPMOAudioEmbeddingItems): audio_inputs = {} else: - audio_inputs = self._base_call_hf_processor( + audio_inputs = self._call_hf_processor_on_prompts( prompts=[self.info.audio_pattern] * len(parsed_audios), mm_data={"audios": [[audio] for audio in parsed_audios]}, mm_kwargs={**mm_kwargs, "chunk_input": True}, diff --git a/vllm/model_executor/models/minicpmv.py b/vllm/model_executor/models/minicpmv.py index 968490d4e9de..b638e1034db1 100644 --- a/vllm/model_executor/models/minicpmv.py +++ b/vllm/model_executor/models/minicpmv.py @@ -881,7 +881,7 @@ def process_images( if isinstance(parsed_images, MiniCPMVImageEmbeddingItems): image_inputs = {} else: - image_inputs = self._base_call_hf_processor( + image_inputs = self._call_hf_processor_on_prompts( prompts=[self.info.image_pattern] * len(parsed_images), mm_data={"images": [[image] for image in parsed_images]}, mm_kwargs=mm_kwargs, @@ -906,7 +906,7 @@ def process_videos( if isinstance(parsed_videos, MiniCPMVVideoEmbeddingItems): video_inputs = {} else: - video_inputs = self._base_call_hf_processor( + video_inputs = self._call_hf_processor_on_prompts( prompts=[ self.info.image_pattern * len(video) for video in parsed_videos ], @@ -983,7 +983,7 @@ def _apply_prompt_updates( return new_token_ids, placeholders - def _base_call_hf_processor( + def _call_hf_processor_on_prompts( self, prompts: list[str], mm_data: Mapping[str, Sequence[object]], @@ -993,19 +993,19 @@ def _base_call_hf_processor( ) -> dict[str, NestedTensors]: # This processor supports zipping prompt and mm_data together if self.info.get_model_version() in {(2, 6), (4, 0), (4, 5), (4, 6)}: - inputs = super()._call_hf_processor( - prompt=prompts, # type: ignore - mm_data=mm_data, - mm_kwargs=mm_kwargs, + inputs = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**mm_kwargs), + dict(text=prompts, **mm_data), + mm_kwargs, ) else: inputs = defaultdict[str, list[torch.Tensor]](list) for i, prompt in enumerate(prompts): - inputs_one = super()._call_hf_processor( - prompt=prompt, - mm_data={k: v[i] for k, v in mm_data.items()}, - mm_kwargs=mm_kwargs, + inputs_one = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**mm_kwargs), + dict(text=prompt, **{k: v[i] for k, v in mm_data.items()}), + mm_kwargs, ) for k, v in inputs_one.items(): @@ -1014,23 +1014,31 @@ def _base_call_hf_processor( return {k: inputs[k] for k in out_keys} - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + tokenizer = self.info.get_tokenizer() - input_ids = torch.tensor([tokenizer.encode(prompt)]) - mm_inputs = self.process_mm_inputs(mm_data, mm_kwargs) + input_ids = torch.tensor([tokenizer.encode(prompt_text)]) + mm_inputs = self.process_mm_inputs(mm_data, hf_processor_mm_kwargs) - return BatchFeature( + processed_data = BatchFeature( { "input_ids": input_ids, **mm_inputs, } ) + processed_data.update(passthrough_data) + return processed_data def _get_prompt_updates( self, diff --git a/vllm/model_executor/models/mistral3.py b/vllm/model_executor/models/mistral3.py index 22adf1c484af..11f214c65b1a 100644 --- a/vllm/model_executor/models/mistral3.py +++ b/vllm/model_executor/models/mistral3.py @@ -233,30 +233,30 @@ def get_dummy_mm_data( class Mistral3MultiModalProcessor(BaseMultiModalProcessor[Mistral3ProcessingInfo]): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _postprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) + if not mm_data: + return processed_data - pixel_values = processed_outputs.get("pixel_values") + pixel_values = processed_data.get("pixel_values") if pixel_values is not None: # Avoid padding since we need the output for each image to be # independent of other images for the cache to work correctly - image_sizes = processed_outputs["image_sizes"] + image_sizes = processed_data["image_sizes"] assert len(pixel_values) == len(image_sizes) - processed_outputs["pixel_values"] = [ + processed_data["pixel_values"] = [ p[:, :h, :w] for p, (h, w) in zip(pixel_values, image_sizes) ] - return processed_outputs + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/mllama4.py b/vllm/model_executor/models/mllama4.py index ca862596ef54..9860996bcf1c 100644 --- a/vllm/model_executor/models/mllama4.py +++ b/vllm/model_executor/models/mllama4.py @@ -589,23 +589,23 @@ def get_image_size_with_most_features(self) -> ImageSize: class Mllama4MultiModalProcessor(BaseMultiModalProcessor[Mllama4ProcessingInfo]): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _postprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) + if not mm_data: + return processed_data - processor = self.info.get_hf_processor(**mm_kwargs) + processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) image_processor = processor.image_processor vision_config = self.info.get_hf_config().vision_config - if processed_outputs.get("pixel_values") is not None: + if processed_data.get("pixel_values") is not None: assert "images" in mm_data, ( "images expected to be in mm_data when pixel_values is present" ) @@ -636,10 +636,10 @@ def _call_hf_processor( 1 if r_h * r_w == 1 else 1 + r_h * r_w for (r_h, r_w) in aspect_ratios ] - processed_outputs["aspect_ratios"] = torch.tensor(aspect_ratios) - processed_outputs["patches_per_image"] = torch.tensor(patches_per_image) + processed_data["aspect_ratios"] = torch.tensor(aspect_ratios) + processed_data["patches_per_image"] = torch.tensor(patches_per_image) - return processed_outputs + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/molmo.py b/vllm/model_executor/models/molmo.py index 2b963e339a92..b823d9c09056 100644 --- a/vllm/model_executor/models/molmo.py +++ b/vllm/model_executor/models/molmo.py @@ -1123,17 +1123,44 @@ def get_dummy_mm_data( class MolmoMultiModalProcessor(BaseMultiModalProcessor[MolmoProcessingInfo]): - def _call_hf_processor( + def _postprocess_prompt(self, prompt: list[int]) -> list[int]: + processor = self.info.get_hf_processor() + + # The chat template is already applied to the prompt tokens + # Use message_format="none" to avoid applying it again + # Prepend an empty space if `always_start_with_space` is True + tokens = processor.get_tokens_input( + self.info.get_tokenizer().decode(prompt), + message_format="none", + always_start_with_space=True, + ) + + # Prepend a BOS token id to the tokens + return self.info.ctx.call_hf_processor( + processor.process, + dict(tokens=tokens), + )["input_ids"].tolist() + + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: - hf_processor = self.info.get_hf_processor(**mm_kwargs) - processed_outputs = self.info.ctx.call_hf_processor( + mm_counts = mm_items.get_all_counts() + + valid_mm_items = mm_items.select({k for k, c in mm_counts.items() if c > 0}) + processor_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not processor_data: + return BatchFeature(dict(passthrough_data)) + + hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) + prompt_text = self.dummy_inputs.get_dummy_text(mm_counts) + + processed_data = self.info.ctx.call_hf_processor( hf_processor.process, - dict(text=prompt, **mm_data), - mm_kwargs, + dict(text=prompt_text, **processor_data), + hf_processor_mm_kwargs, ) tokenizer = hf_processor.tokenizer @@ -1141,17 +1168,16 @@ def _call_hf_processor( image_processor = hf_processor.image_processor - input_ids: torch.Tensor = processed_outputs.pop("input_ids") - processed_outputs["input_ids"] = input_ids.unsqueeze(0) + processed_data.pop("input_ids") - if (images := mm_data.get("images")) is not None: + if (images := processor_data.get("images")) is not None: mm_items = self.info.parse_mm_data({"image": images}, validate=False) parsed_images = mm_items.get_items("image", ImageProcessorItems) image_sizes = [ parsed_images.get_image_size(i) for i in range(len(parsed_images)) ] - feat_is_patch = processed_outputs["image_input_idx"] >= 0 + feat_is_patch = processed_data["image_input_idx"] >= 0 tilings = [ self.info.select_tiling( @@ -1165,35 +1191,12 @@ def _call_hf_processor( num_crops = torch.tensor(tilings).prod(-1) + 1 assert num_crops.sum() == len(feat_is_patch) - processed_outputs["num_crops"] = num_crops - processed_outputs["img_patch_id"] = image_patch_id - - return processed_outputs - - def _apply_hf_processor_tokens_only( - self, - prompt_tokens: list[int], - ) -> list[int]: - processor = self.info.get_hf_processor() + processed_data["num_crops"] = num_crops + processed_data["img_patch_id"] = image_patch_id - # The chat template is already applied to the prompt tokens - # Use message_format="none" to avoid applying it again - # Prepend an empty space if `always_start_with_space` is True - tokens = processor.get_tokens_input( - self.info.get_tokenizer().decode(prompt_tokens), - message_format="none", - always_start_with_space=True, - ) - - # Prepend a BOS token id to the tokens - processed_data = self.info.ctx.call_hf_processor( - processor.process, - dict(tokens=tokens), - ) - prompt_ids = processed_data.pop("input_ids").tolist() - print(prompt_ids, len(prompt_ids)) + processed_data.update(passthrough_data) - return prompt_ids + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/molmo2.py b/vllm/model_executor/models/molmo2.py index 861e31928bdc..d73160d68989 100644 --- a/vllm/model_executor/models/molmo2.py +++ b/vllm/model_executor/models/molmo2.py @@ -1929,30 +1929,33 @@ def _get_dummy_videos( class Molmo2MultiModalProcessor(BaseMultiModalProcessor[Molmo2ProcessingInfo]): - def _apply_hf_processor_tokens_only( - self, - prompt_tokens: list[int], - ) -> list[int]: + def _postprocess_prompt(self, prompt: list[int]) -> list[int]: processor = self.info.get_hf_processor() tokenizer = processor.tokenizer bos_token_id = tokenizer.bos_token_id or tokenizer.eos_token_id - if len(prompt_tokens) == 0 or prompt_tokens[0] != bos_token_id: + if len(prompt) == 0 or prompt[0] != bos_token_id: # Prepend the bos token to the prompt tokens - prompt_tokens = [bos_token_id] + prompt_tokens + prompt = [bos_token_id] + prompt - return prompt_tokens + return prompt - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: - mm_data = dict(mm_data) + mm_counts = mm_items.get_all_counts() + + valid_mm_items = mm_items.select({k for k, c in mm_counts.items() if c > 0}) + processor_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_counts) + + mm_data = dict(processor_data) hf_config = self.info.get_hf_config() - hf_processor = self.info.get_hf_processor(**mm_kwargs) + hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) def patched_call(text=None, images=None, videos=None, **kwargs) -> BatchFeature: res = hf_processor(text=text, images=images, videos=videos, **kwargs) @@ -1984,14 +1987,16 @@ def patched_call(text=None, images=None, videos=None, **kwargs) -> BatchFeature: # NOTE: metadata.frames_indices indicates # the sampled frames indices of pre-sampled videos, which is # used to calculate the timestamps. Make sure that - # do_sample_frames in mm_kwargs is false for presampled videos. + # do_sample_frames in hf_processor_mm_kwargs is false for + # presampled videos. - # NOTE: a copy of mm_kwargs is created to update do_sample_frames, - # otherwise mm_hash for the object will be incorrect. - video_mm_kwargs = dict(**mm_kwargs) + # NOTE: a copy of hf_processor_mm_kwargs is created to update + # do_sample_frames, otherwise mm_hash for the object will be + # incorrect. + video_mm_kwargs = dict(**hf_processor_mm_kwargs) if "do_sample_frames" not in video_mm_kwargs: # molmo_utils already has "do_sample_frames" in - # mm_kwargs, don't overwrite it. + # hf_processor_mm_kwargs, don't overwrite it. video_mm_kwargs["do_sample_frames"] = metadata.get( "do_sample_frames", False ) @@ -2015,7 +2020,7 @@ def patched_call(text=None, images=None, videos=None, **kwargs) -> BatchFeature: input_ids = input_ids[:, 1:] video_string = tokenizer.batch_decode(input_ids)[0] - prompt = prompt.replace(VIDEO_PROMPT, video_string, 1) + prompt_text = prompt_text.replace(VIDEO_PROMPT, video_string, 1) video_grids = video_outputs.pop("video_grids") assert video_grids[:, 0].sum() == len( @@ -2054,10 +2059,10 @@ def patched_call(text=None, images=None, videos=None, **kwargs) -> BatchFeature: else: all_video_outputs = dict() - processed_outputs = self.info.ctx.call_hf_processor( + processed_data = self.info.ctx.call_hf_processor( patched_call, - dict(text=prompt, **mm_data), - mm_kwargs, + dict(text=prompt_text, **mm_data), + hf_processor_mm_kwargs, ) if (images := mm_data.get("images")) is not None: @@ -2077,22 +2082,22 @@ def patched_call(text=None, images=None, videos=None, **kwargs) -> BatchFeature: for image_size in image_sizes ] num_crops = torch.tensor(tilings).prod(-1) + 1 - assert sum(num_crops) == len(processed_outputs["pixel_values"]) - assert sum(num_crops) == processed_outputs["image_num_crops"].sum().item() + assert sum(num_crops) == len(processed_data["pixel_values"]) + assert sum(num_crops) == processed_data["image_num_crops"].sum().item() - image_grids = processed_outputs.pop("image_grids") + image_grids = processed_data.pop("image_grids") image_num_pooled_patches = image_grids[:, :2].prod(dim=1) + image_grids[ :, 2: ].prod(dim=1) - processed_outputs["image_num_pooled_patches"] = image_num_pooled_patches - n_patches = processed_outputs["pixel_values"].shape[1] - processed_outputs["image_num_patches"] = ( - processed_outputs["image_num_crops"] * n_patches + processed_data["image_num_pooled_patches"] = image_num_pooled_patches + n_patches = processed_data["pixel_values"].shape[1] + processed_data["image_num_patches"] = ( + processed_data["image_num_crops"] * n_patches ) ( - processed_outputs["image_tokens"], - processed_outputs["num_image_tokens"], + processed_data["image_tokens"], + processed_data["num_image_tokens"], ) = build_flat_image_bool_length( image_grids, hf_config, @@ -2101,7 +2106,11 @@ def patched_call(text=None, images=None, videos=None, **kwargs) -> BatchFeature: use_single_crop_start_token=hf_processor.use_single_crop_start_token, ) - return BatchFeature({**processed_outputs, **all_video_outputs}) + processed_data.update(all_video_outputs) + processed_data.update(passthrough_data) + processed_data.pop("input_ids") + + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/moondream3.py b/vllm/model_executor/models/moondream3.py index 0464150bd49e..45833f6241bd 100644 --- a/vllm/model_executor/models/moondream3.py +++ b/vllm/model_executor/models/moondream3.py @@ -942,16 +942,6 @@ class Moondream3MultiModalProcessor(BaseMultiModalProcessor[Moondream3Processing image_placeholder: str = "" bos_image_placeholder: str = "<|endoftext|>" - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - # Moondream3's processor handles images directly rather than exposing a - # separate `image_processor`, so keep the cache path on text+MM calls. - return super()._call_hf_processor(prompt, mm_data, mm_kwargs) - @cached_property def bos_image_placeholder_tokens(self) -> list[int]: tokenizer = self.info.get_tokenizer() diff --git a/vllm/model_executor/models/moss_audio.py b/vllm/model_executor/models/moss_audio.py index ceed6245a1ca..7c4fa3b8f889 100644 --- a/vllm/model_executor/models/moss_audio.py +++ b/vllm/model_executor/models/moss_audio.py @@ -1335,23 +1335,32 @@ def get_dummy_mm_data( class MossAudioMultiModalProcessor(BaseMultiModalProcessor[MossAudioProcessingInfo]): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + mm_data = dict(mm_data) audios = mm_data.pop("audios", []) if audios: mm_data["audio"] = audios - mm_kwargs = dict(mm_kwargs) - processor_kwargs = _filter_moss_audio_processor_config(mm_kwargs) - return self.info.ctx.call_hf_processor( + hf_processor_mm_kwargs = dict(hf_processor_mm_kwargs) + processor_kwargs = _filter_moss_audio_processor_config(hf_processor_mm_kwargs) + processed_data = self.info.ctx.call_hf_processor( self.info.get_hf_processor(**processor_kwargs), - dict(text=prompt, **mm_data), + mm_data, {}, ) + processed_data.update(passthrough_data) + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/moss_transcribe_diarize.py b/vllm/model_executor/models/moss_transcribe_diarize.py index 421901d0ee8f..8122308af4af 100644 --- a/vllm/model_executor/models/moss_transcribe_diarize.py +++ b/vllm/model_executor/models/moss_transcribe_diarize.py @@ -461,24 +461,29 @@ def get_dummy_processor_inputs( class MossTranscribeDiarizeMultiModalProcessor( BaseMultiModalProcessor[MossTranscribeDiarizeProcessingInfo] ): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _preprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + ) -> tuple[Mapping[str, object], Mapping[str, object]]: + audios = _get_audios_from_mm_data(mm_data) + + return dict(audio=audios), hf_processor_mm_kwargs + + def _postprocess_hf_mm_data( + self, + mm_data: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: - tokenizer = self.info.get_tokenizer() audios = _get_audios_from_mm_data(mm_data) - if not audios: - input_ids = tokenizer.encode(prompt, add_special_tokens=False) - return BatchFeature({"input_ids": [input_ids]}, tensor_type="pt") - - processed = self.info.ctx.call_hf_processor( - self.info.get_hf_processor(**mm_kwargs), - dict(text=prompt, audio=audios), - mm_kwargs, - ) - return _add_vllm_audio_metadata(processed, len(audios)) + if audios: + processed_data = _add_vllm_audio_metadata(processed_data, len(audios)) + + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/muse_glimmer.py b/vllm/model_executor/models/muse_glimmer.py index fcdfb82333cd..a0281ec0aa4d 100644 --- a/vllm/model_executor/models/muse_glimmer.py +++ b/vllm/model_executor/models/muse_glimmer.py @@ -262,13 +262,19 @@ def get_dummy_mm_data( class MuseGlimmerMultiModalProcessor( BaseMultiModalProcessor[MuseGlimmerProcessingInfo] ): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: - processor = self.info.get_hf_processor(**mm_kwargs) + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + + processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) tokenizer = processor.tokenizer config = self.info.get_hf_config() images = mm_data.get("images", ()) @@ -276,7 +282,7 @@ def _call_hf_processor( if not isinstance(images, Sequence) or not isinstance(videos, Sequence): raise TypeError("MuseGlimmer multi-modal data must be a sequence") - prompt_ids = tokenizer.encode(prompt, add_special_tokens=False) + prompt_ids = tokenizer.encode(prompt_text, add_special_tokens=False) image_sentinel_id = tokenizer.convert_tokens_to_ids(IMAGE_TOKEN) video_sentinel_id = tokenizer.convert_tokens_to_ids(VIDEO_TOKEN) if prompt_ids.count(image_sentinel_id) != len(images): @@ -354,7 +360,9 @@ def _call_hf_processor( video_pixel_values=video_pixels, video_feature_sizes=torch.tensor(video_sizes), ) - return BatchFeature(data=data, tensor_type=None) + processed_data = BatchFeature(data=data, tensor_type=None) + processed_data.update(passthrough_data) + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/nano_nemotron_vl.py b/vllm/model_executor/models/nano_nemotron_vl.py index 679495e3ecf0..0e9eda5c82a4 100644 --- a/vllm/model_executor/models/nano_nemotron_vl.py +++ b/vllm/model_executor/models/nano_nemotron_vl.py @@ -361,18 +361,8 @@ def get_mm_max_tokens_per_item( class NanoNemotronVLMultiModalProcessor( BaseMultiModalProcessor[NanoNemotronVLProcessingInfo] ): - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - """ - Bypass `call_hf_processor_mm_only` by no-op overriding`_call_hf_processor`, - so it chooses this path: - `type(self)._call_hf_processor != BaseMultiModalProcessor._call_hf_processor` - """ - return super()._call_hf_processor(prompt, mm_data, mm_kwargs) + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) def _get_image_fields_config(self, hf_inputs: BatchFeature): if self.info.is_dynamic_tiler: @@ -406,7 +396,9 @@ def _get_video_fields_config(self, hf_inputs: BatchFeature): ) def _get_audio_fields_config(self, hf_inputs: BatchFeature): - audio_num_clips = torch.as_tensor(hf_inputs["audio_num_clips"]) + audio_num_clips = torch.as_tensor( + hf_inputs.get("audio_num_clips", torch.empty(0)) + ) return dict( input_audio_features=MultiModalFieldConfig.flat_from_sizes( @@ -733,11 +725,8 @@ def apply( # Bypass the cached path: the HF processor must receive the # prompt (with injected ) and the audio data # together so it can perform audio-token replacement natively. - ( - prompt_ids, - mm_info, - is_update_applied, - ) = self._apply_hf_processor(inputs, timing_ctx) + mm_info = self._apply_hf_processor(inputs, timing_ctx) + prompt_ids = self._postprocess_prompt(inputs.prompt) with timing_ctx.record("apply_prompt_updates"): prompt_ids, mm_placeholders = self._maybe_apply_prompt_updates( @@ -745,7 +734,7 @@ def apply( prompt_ids=prompt_ids, mm_kwargs=mm_info.kwargs, mm_prompt_updates=mm_info.prompt_updates, - is_update_applied=is_update_applied, + is_update_applied=self.hf_processor_applies_updates, ) mm_placeholder_ranges = { diff --git a/vllm/model_executor/models/nemotron_parse.py b/vllm/model_executor/models/nemotron_parse.py index 9ce8ba22463c..12c7b580484b 100644 --- a/vllm/model_executor/models/nemotron_parse.py +++ b/vllm/model_executor/models/nemotron_parse.py @@ -411,22 +411,6 @@ def create_encoder_prompt( ) -> list[int]: return [0] - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - if mm_data: - processed_outputs = super()._call_hf_processor(prompt, mm_data, mm_kwargs) - else: - hf_processor = self.info.get_hf_processor() - tokenizer = hf_processor.tokenizer - processed_outputs = tokenizer( - prompt, add_special_tokens=False, return_tensors="pt" - ) - return processed_outputs - def _get_mm_fields_config( self, hf_inputs: BatchFeature, diff --git a/vllm/model_executor/models/ovis.py b/vllm/model_executor/models/ovis.py index 4d023f51e3cb..d1a805f608fc 100644 --- a/vllm/model_executor/models/ovis.py +++ b/vllm/model_executor/models/ovis.py @@ -341,41 +341,30 @@ def image_indicators_to_visual_tokens( # -300 is image_atom token, filter them out return [vte_vocab_size + x + 300 for x in image_indicators if x < -300] - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _postprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: if not mm_data: - # Avoid warning from HF logger for text-only input - tokenizer = self.info.get_tokenizer() - prompt_ids = tokenizer.encode(prompt, add_special_tokens=False) - return BatchFeature(dict(input_ids=[prompt_ids]), tensor_type="pt") - - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) + return processed_data hf_processor = self.info.get_hf_processor() image_indicators = [ hf_processor.construct_image_indicators(grid) - for grid in processed_outputs["grids"] + for grid in processed_data["grids"] ] indicator_tokens = [ self.image_indicators_to_visual_tokens(indicator) for indicator in image_indicators ] - processed_outputs["indicator_tokens"] = torch.tensor(indicator_tokens) - return processed_outputs + processed_data["indicator_tokens"] = torch.tensor(indicator_tokens) - def _apply_hf_processor_tokens_only( - self, - prompt_tokens: list[int], - ) -> list[int]: - return prompt_tokens + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/ovis2_5.py b/vllm/model_executor/models/ovis2_5.py index 1ccf820e0fbf..c55f4805377a 100644 --- a/vllm/model_executor/models/ovis2_5.py +++ b/vllm/model_executor/models/ovis2_5.py @@ -332,53 +332,43 @@ def visual_indicators_to_visual_tokens( if x >= INDICATOR_IDS[0] ] - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _postprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: if not mm_data: - # Avoid warning from HF logger for text-only input - tokenizer = self.info.get_tokenizer() - prompt_ids = tokenizer.encode(prompt, add_special_tokens=False) - return BatchFeature(dict(input_ids=[prompt_ids]), tensor_type="pt") - - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) + return processed_data + hf_processor = self.info.get_hf_processor() if "videos" in mm_data: visual_indicators = [ hf_processor.construct_visual_indicators((1, 1, 1), True) - for grid in processed_outputs["video_grids"] + for grid in processed_data["video_grids"] ] indicator_tokens = [ self.visual_indicators_to_visual_tokens(indicator) for indicator in visual_indicators ] - processed_outputs["video_indicator_tokens"] = torch.tensor(indicator_tokens) + processed_data["video_indicator_tokens"] = torch.tensor(indicator_tokens) if "images" in mm_data: visual_indicators = [ hf_processor.construct_visual_indicators((1, 1, 1), False) - for grid in processed_outputs["grids"] + for grid in processed_data["grids"] ] indicator_tokens = [ self.visual_indicators_to_visual_tokens(indicator) for indicator in visual_indicators ] - processed_outputs["indicator_tokens"] = torch.tensor(indicator_tokens) - return processed_outputs + processed_data["indicator_tokens"] = torch.tensor(indicator_tokens) - def _apply_hf_processor_tokens_only( - self, - prompt_tokens: list[int], - ) -> list[int]: - return prompt_tokens + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/paddleocr_vl.py b/vllm/model_executor/models/paddleocr_vl.py index 61f1ed9d8324..a432a8d5f0d4 100644 --- a/vllm/model_executor/models/paddleocr_vl.py +++ b/vllm/model_executor/models/paddleocr_vl.py @@ -244,32 +244,37 @@ def get_dummy_mm_data( class PaddleOCRVLMultiModalProcessor( BaseMultiModalProcessor[PaddleOCRVLProcessingInfo] ): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: - if mm_data: - final_mm_kwargs = dict(mm_kwargs or {}) - final_mm_kwargs.setdefault("images_kwargs", {}) - # vLLM use PIL.Image, always set channel_last - final_mm_kwargs["input_data_format"] = ChannelDimension.LAST - processed_outputs = self.info.ctx.call_hf_processor( - self.info.get_hf_processor(**final_mm_kwargs), - dict(text=prompt, **mm_data), - mm_kwargs, - ) - num_patches_per_image = processed_outputs["image_grid_thw"].prod(-1) - processed_outputs["pixel_values"] = processed_outputs["pixel_values"].split( - num_patches_per_image.tolist() - ) - else: - tokenizer = self.info.get_tokenizer() - processed_outputs = tokenizer( - prompt, add_special_tokens=True, return_tensors="pt" - ) - return processed_outputs + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + + final_mm_kwargs = dict(hf_processor_mm_kwargs or {}) + final_mm_kwargs.setdefault("images_kwargs", {}) + # vLLM use PIL.Image, always set channel_last + final_mm_kwargs["input_data_format"] = ChannelDimension.LAST + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**final_mm_kwargs), + dict(text=prompt_text, **mm_data), + hf_processor_mm_kwargs, + ) + num_patches_per_image = processed_data["image_grid_thw"].prod(-1) + processed_data["pixel_values"] = processed_data["pixel_values"].split( + num_patches_per_image.tolist() + ) + processed_data.update(passthrough_data) + + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/paligemma.py b/vllm/model_executor/models/paligemma.py index 1e80e8d8186a..781e9a620d2f 100644 --- a/vllm/model_executor/models/paligemma.py +++ b/vllm/model_executor/models/paligemma.py @@ -152,23 +152,6 @@ def get_dummy_mm_data( class PaliGemmaMultiModalProcessor(BaseMultiModalProcessor[PaliGemmaProcessingInfo]): - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - tokenizer = self.info.get_tokenizer() - if not mm_data: - prompt_ids = tokenizer.encode(prompt, add_special_tokens=False) - return BatchFeature(dict(input_ids=[prompt_ids]), tensor_type="pt") - - return super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) - def _get_mm_fields_config( self, hf_inputs: BatchFeature, diff --git a/vllm/model_executor/models/phi3v.py b/vllm/model_executor/models/phi3v.py index aca1b9a6328c..5ed955db828a 100644 --- a/vllm/model_executor/models/phi3v.py +++ b/vllm/model_executor/models/phi3v.py @@ -396,19 +396,19 @@ def get_dummy_mm_data( class Phi3VMultiModalProcessor(BaseMultiModalProcessor[Phi3VProcessingInfo]): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _postprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, ) -> BatchFeature: - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) + if not mm_data: + return processed_data - input_ids = processed_outputs["input_ids"] + input_ids = processed_data["input_ids"] assert isinstance(input_ids, torch.Tensor) # Phi3v processor has inserted -1, -2 etc as placeholder in prompt_ids, @@ -416,7 +416,7 @@ def _call_hf_processor( # Therefore, we need to do an early replacement here input_ids.masked_fill_(input_ids < 0, _IMAGE_TOKEN_ID) - return processed_outputs + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/phi4mm.py b/vllm/model_executor/models/phi4mm.py index b18ff6bb0ffd..4dc5d408a4d4 100644 --- a/vllm/model_executor/models/phi4mm.py +++ b/vllm/model_executor/models/phi4mm.py @@ -868,43 +868,52 @@ def get_dummy_mm_data( class Phi4MMMultiModalProcessor(BaseMultiModalProcessor[Phi4MMProcessingInfo]): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + if not mm_data: - prompt_ids = self.info.get_tokenizer().encode(prompt) - prompt_ids = self._apply_hf_processor_tokens_only(prompt_ids) - return BatchFeature(dict(input_ids=[prompt_ids]), tensor_type="pt") + return BatchFeature(dict(passthrough_data)) - sr = self.info.get_feature_extractor(**mm_kwargs).sampling_rate + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + + sr = self.info.get_feature_extractor(**hf_processor_mm_kwargs).sampling_rate if audio_data := mm_data.get("audios", []): mm_data["audios"] = [(data, sr) for data in audio_data] - processed_outputs = super()._call_hf_processor(prompt, mm_data, mm_kwargs) + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(text=prompt_text, **mm_data), + hf_processor_mm_kwargs, + ) - hf_processor = self.info.get_hf_processor(**mm_kwargs) + hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) num_img_tokens = [ self.info.get_num_image_tokens( image_width=img_size[0], image_height=img_size[1], processor=hf_processor, ) - for img_size in processed_outputs["image_sizes"] + for img_size in processed_data["image_sizes"] ] - processed_outputs["num_img_tokens"] = num_img_tokens + processed_data["num_img_tokens"] = num_img_tokens - audio_features = processed_outputs["input_audio_embeds"] + audio_features = processed_data["input_audio_embeds"] feature_sizes = [ self.info.get_audio_num_frames(len(audio), sr) for audio in audio_data ] - processed_outputs["input_audio_embeds"] = [ + processed_data["input_audio_embeds"] = [ audio_features[idx, :size] for idx, size in enumerate(feature_sizes) ] + processed_data.update(passthrough_data) - return processed_outputs + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/phi4siglip.py b/vllm/model_executor/models/phi4siglip.py index cb43161486b7..6c6c6efa2dea 100644 --- a/vllm/model_executor/models/phi4siglip.py +++ b/vllm/model_executor/models/phi4siglip.py @@ -142,16 +142,22 @@ def get_dummy_mm_data( class Phi4SiglipMultiModalProcessor( BaseMultiModalProcessor[Phi4SiglipProcessingInfo], ): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: - processed = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + + processed = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(text=prompt_text, **mm_data), + hf_processor_mm_kwargs, ) # The HF processor's tokenizer_image_token() replaces the "" @@ -163,10 +169,12 @@ def _call_hf_processor( # NOTE: tokenizer.__call__() (not .encode()) must be used so that # added/special tokens like <|user|>, <|end|> are kept as single IDs. tokenizer = self.info.get_tokenizer() - new_ids = tokenizer(prompt).input_ids + new_ids = tokenizer(prompt_text).input_ids processed["input_ids"] = torch.tensor([new_ids]) - return processed + processed_data = processed + processed_data.update(passthrough_data) + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/pixtral.py b/vllm/model_executor/models/pixtral.py index 48b3310d30e9..a24d2484ab29 100644 --- a/vllm/model_executor/models/pixtral.py +++ b/vllm/model_executor/models/pixtral.py @@ -51,12 +51,10 @@ from vllm.multimodal.processing.processor import ( BaseMultiModalProcessor, BaseProcessingInfo, - MultiModalProcessingInfo, ProcessorInputs, PromptReplacement, PromptUpdate, PromptUpdateDetails, - TimingContext, ) from vllm.platforms import current_platform from vllm.sequence import IntermediateTensors @@ -213,6 +211,9 @@ def get_dummy_processor_inputs( class PixtralMultiModalProcessor(BaseMultiModalProcessor[PixtralProcessingInfo]): + # The tokens are already inserted by the chat template + hf_processor_applies_updates = True + def _get_mm_fields_config( self, hf_inputs: Mapping[str, NestedTensors], @@ -220,24 +221,35 @@ def _get_mm_fields_config( ) -> Mapping[str, MultiModalFieldConfig]: return dict(images=MultiModalFieldConfig.batched("image")) - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + outputs = self.info.ctx.call_hf_processor( - self.info.get_hf_processor(**mm_kwargs), - dict(text=prompt, **mm_data), + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(text=prompt_text, **mm_data), # Avoid padding issue - dict(**mm_kwargs, return_tensors=None), + dict(**hf_processor_mm_kwargs, return_tensors=None), ) # Missing batch dimension if is_list_of(outputs["input_ids"], int): outputs["input_ids"] = [outputs["input_ids"]] - return outputs + processed_data = outputs + processed_data.update(passthrough_data) + return processed_data def _get_prompt_updates( self, @@ -273,16 +285,6 @@ def get_replacement(item_idx: int): ), ] - def _cached_apply_hf_processor( - self, - inputs: ProcessorInputs, - timing_ctx: TimingContext, - ) -> tuple[list[int], MultiModalProcessingInfo, bool]: - prompt_ids, mm_info, _ = super()._cached_apply_hf_processor(inputs, timing_ctx) - - # NOTE: The tokens are already inserted by the chat template - return prompt_ids, mm_info, True - @MULTIMODAL_REGISTRY.register_processor( PixtralMultiModalProcessor, diff --git a/vllm/model_executor/models/qwen2_5_omni_thinker.py b/vllm/model_executor/models/qwen2_5_omni_thinker.py index 052b2823de26..1f335cfc2420 100644 --- a/vllm/model_executor/models/qwen2_5_omni_thinker.py +++ b/vllm/model_executor/models/qwen2_5_omni_thinker.py @@ -499,83 +499,6 @@ def get_dummy_mm_data( class Qwen2_5OmniThinkerMultiModalProcessor( BaseMultiModalProcessor[Qwen2_5OmniThinkerProcessingInfo] ): - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - mm_data = dict(mm_data) - audios = mm_data.pop("audios", []) - - # NOTE: WhisperFeatureExtractor cannot handle empty list of audios - if audios: - # NOTE: Qwen2.5-Omni processor accept "audio" - mm_data["audio"] = audios - mm_kwargs = dict( - **mm_kwargs, - ) - elif ( - mm_kwargs.get("use_audio_in_video") - and "audio" not in mm_data - and mm_data.get("videos") - ): - # A subclass (e.g. Qwen3-Omni) may have already popped "audios" - # and populated mm_data["audio"] itself before delegating here, - # in which case this isn't a genuine "no audio" case. Likewise, - # mm_data can be empty on a multimodal-processor-cache hit, where - # there's nothing to (re-)process this call and the real result - # comes from the cache — not a genuine "no audio" case either. - raise ValueError( - "Video doesn't have audio track with `audio_in_video=True`" - ) - - merged = self.info.ctx.get_merged_mm_kwargs(mm_kwargs) - if mm_data.get("videos") and ( - merged.keys() & {"size", "min_pixels", "max_pixels"} - ): - mm_kwargs = dict(mm_kwargs) - video_size = dict(self.info.get_hf_processor().video_processor.size) - size_override = merged.get("size") - if size_override is not None: - video_size = video_size | size_override - min_pixels = merged.get("min_pixels") - if min_pixels is not None: - video_size["shortest_edge"] = min_pixels - max_pixels = merged.get("max_pixels") - if max_pixels is not None: - video_size["longest_edge"] = max_pixels - mm_kwargs["size"] = video_size - - hf_inputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) - - input_features = hf_inputs.pop("input_features", None) - feature_attention_mask = hf_inputs.get("feature_attention_mask", None) - if "input_audio_features" not in hf_inputs and input_features is not None: - if feature_attention_mask is not None: - input_features = input_features.permute(0, 2, 1)[ - feature_attention_mask.bool() - ].permute(1, 0) - hf_inputs["input_audio_features"] = input_features - if ( - "audio_feature_lengths" not in hf_inputs - and feature_attention_mask is not None - ): - hf_inputs["audio_feature_lengths"] = feature_attention_mask.sum(-1) - - video_second_per_grid = hf_inputs.get("video_second_per_grid", None) - if video_second_per_grid is not None: - hf_inputs["second_per_grid_ts"] = video_second_per_grid - - use_audio_in_video = mm_kwargs.get("use_audio_in_video", False) - hf_inputs["use_audio_in_video"] = torch.tensor(use_audio_in_video) - - return hf_inputs - def _get_mm_fields_config( self, hf_inputs: BatchFeature, @@ -885,7 +808,7 @@ def get_replacement_qwen2_use_audio_in_video(item_idx: int): ), ] - def _apply_hf_processor_mm_only( + def _apply_hf_processor_main( self, mm_items: MultiModalDataItems, hf_processor_mm_kwargs: Mapping[str, object], @@ -903,13 +826,87 @@ def _apply_hf_processor_mm_only( ) mm_counts["audio"] -= mm_counts["video"] - _, mm_processed_data = self._apply_hf_processor_text_mm( - prompt_text=self.dummy_inputs.get_dummy_text(mm_counts), - mm_items=mm_items, - hf_processor_mm_kwargs=hf_processor_mm_kwargs, + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} ) + processor_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_counts) + + mm_data = dict(processor_data) + audios = mm_data.pop("audios", []) + + # NOTE: WhisperFeatureExtractor cannot handle empty list of audios + if audios: + # NOTE: Qwen2.5-Omni processor accept "audio" + mm_data["audio"] = audios + hf_processor_mm_kwargs = dict( + **hf_processor_mm_kwargs, + ) + elif ( + hf_processor_mm_kwargs.get("use_audio_in_video") + and "audio" not in mm_data + and mm_data.get("videos") + ): + # A subclass (e.g. Qwen3-Omni) may have already popped "audios" + # and populated mm_data["audio"] itself before delegating here, + # in which case this isn't a genuine "no audio" case. Likewise, + # mm_data can be empty on a multimodal-processor-cache hit, where + # there's nothing to (re-)process this call and the real result + # comes from the cache — not a genuine "no audio" case either. + raise ValueError( + "Video doesn't have audio track with `audio_in_video=True`" + ) + + merged = self.info.ctx.get_merged_mm_kwargs(hf_processor_mm_kwargs) + if mm_data.get("videos") and ( + merged.keys() & {"size", "min_pixels", "max_pixels"} + ): + hf_processor_mm_kwargs = dict(hf_processor_mm_kwargs) + video_size = dict(self.info.get_hf_processor().video_processor.size) + size_override = merged.get("size") + if size_override is not None: + video_size = video_size | size_override + min_pixels = merged.get("min_pixels") + if min_pixels is not None: + video_size["shortest_edge"] = min_pixels + max_pixels = merged.get("max_pixels") + if max_pixels is not None: + video_size["longest_edge"] = max_pixels + hf_processor_mm_kwargs["size"] = video_size + + hf_inputs = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(text=prompt_text, **mm_data), + hf_processor_mm_kwargs, + ) + + input_features = hf_inputs.pop("input_features", None) + feature_attention_mask = hf_inputs.get("feature_attention_mask", None) + if "input_audio_features" not in hf_inputs and input_features is not None: + if feature_attention_mask is not None: + input_features = input_features.permute(0, 2, 1)[ + feature_attention_mask.bool() + ].permute(1, 0) + hf_inputs["input_audio_features"] = input_features + if ( + "audio_feature_lengths" not in hf_inputs + and feature_attention_mask is not None + ): + hf_inputs["audio_feature_lengths"] = feature_attention_mask.sum(-1) + + video_second_per_grid = hf_inputs.get("video_second_per_grid", None) + if video_second_per_grid is not None: + hf_inputs["second_per_grid_ts"] = video_second_per_grid + + use_audio_in_video = hf_processor_mm_kwargs.get("use_audio_in_video", False) + hf_inputs["use_audio_in_video"] = torch.tensor(use_audio_in_video) + + processed_data = hf_inputs + processed_data.update(passthrough_data) + processed_data.pop("input_ids") - return mm_processed_data + return processed_data class Qwen2_5OmniConditionalGenerationMixin: diff --git a/vllm/model_executor/models/qwen2_5_vl.py b/vllm/model_executor/models/qwen2_5_vl.py index 396a892cabba..3e5deaad91b9 100644 --- a/vllm/model_executor/models/qwen2_5_vl.py +++ b/vllm/model_executor/models/qwen2_5_vl.py @@ -1154,16 +1154,6 @@ def _get_mm_fields_config( second_per_grid_ts=MultiModalFieldConfig.batched("video", keep_on_cpu=True), ) - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - # Override to use the text path instead of token path to use the - # video-specific logic in processing_qwen2_5_vl.py - return super()._call_hf_processor(prompt, mm_data, mm_kwargs) - def _get_prompt_updates( self, mm_items: MultiModalDataItems, diff --git a/vllm/model_executor/models/qwen2_audio.py b/vllm/model_executor/models/qwen2_audio.py index 695d4833557e..74341575e744 100644 --- a/vllm/model_executor/models/qwen2_audio.py +++ b/vllm/model_executor/models/qwen2_audio.py @@ -235,36 +235,27 @@ def get_dummy_mm_data( class Qwen2AudioMultiModalProcessor(BaseMultiModalProcessor[Qwen2AudioProcessingInfo]): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _preprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, Any], - ) -> BatchFeature: - # NOTE - we rename audios -> audio in mm data because transformers has - # deprecated audios for the qwen2audio processor and will remove - # support for it in transformers 4.54. + hf_processor_mm_kwargs: Mapping[str, object], + ) -> tuple[Mapping[str, object], Mapping[str, object]]: + feature_extractor = self.info.get_feature_extractor(**hf_processor_mm_kwargs) + + mm_data = dict(mm_data) audios = mm_data.pop("audios", []) if audios: mm_data["audio"] = audios - # Text-only input not supported in composite processor - if not mm_data.get("audio", []): - prompt_ids = self.info.get_tokenizer().encode(prompt) - prompt_ids = self._apply_hf_processor_tokens_only(prompt_ids) - return BatchFeature(dict(input_ids=[prompt_ids]), tensor_type="pt") - - feature_extractor = self.info.get_feature_extractor(**mm_kwargs) - mm_kwargs = dict( - **mm_kwargs, + hf_processor_mm_kwargs = dict( + **hf_processor_mm_kwargs, sampling_rate=feature_extractor.sampling_rate, ) - return super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, - ) + return mm_data, hf_processor_mm_kwargs def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/qwen3_omni_moe_thinker.py b/vllm/model_executor/models/qwen3_omni_moe_thinker.py index 622376f0b3f4..a3691742ffc5 100755 --- a/vllm/model_executor/models/qwen3_omni_moe_thinker.py +++ b/vllm/model_executor/models/qwen3_omni_moe_thinker.py @@ -1195,79 +1195,95 @@ def get_mm_max_tokens_per_item( class Qwen3OmniMoeThinkerMultiModalProcessor( Qwen2_5OmniThinkerMultiModalProcessor, ): - def _call_hf_processor( + def _get_hf_mm_data( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - mm_data = dict(mm_data) - audios = mm_data.pop("audios", []) - - def pad_to_hop_length(x: np.ndarray, hop_length: int) -> np.ndarray: - length = x.shape[-1] - if length % hop_length != 0: - pad_length = hop_length - (length % hop_length) - x = np.pad(x, (0, pad_length), mode="constant", constant_values=0) - return x - - # NOTE: WhisperFeatureExtractor cannot handle empty list of audios - feature_extractor = self.info.get_feature_extractor(**mm_kwargs) - hop_length = feature_extractor.hop_length + mm_items: MultiModalDataItems, + ) -> tuple[Mapping[str, object], Mapping[str, object]]: + processor_data, passthrough_data = super()._get_hf_mm_data(mm_items) + + audios = processor_data.get("audios") if audios: - # NOTE: Qwen3-Omni processor accept "audio" + feature_extractor = self.info.get_feature_extractor() + hop_length = feature_extractor.hop_length + + def pad_to_hop_length(x: np.ndarray) -> np.ndarray: + length = x.shape[-1] + if length % hop_length != 0: + pad_length = hop_length - (length % hop_length) + x = np.pad(x, (0, pad_length), mode="constant", constant_values=0) + return x + # To make sure the cache works with padding=True, we pre-padded - # the audio to multiple of hop_length. - mm_data["audio"] = [ - pad_to_hop_length(audio, hop_length) + # the audio to a multiple of hop_length. + processor_data = dict(processor_data) + processor_data["audios"] = [ + pad_to_hop_length(audio) if isinstance(audio, np.ndarray) - else (pad_to_hop_length(audio[0], hop_length), audio[1]) + else (pad_to_hop_length(audio[0]), audio[1]) for audio in audios ] + return processor_data, passthrough_data + + def _apply_hf_processor_main( + self, + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], + ) -> BatchFeature: + if mm_items.get_all_counts().get("audio", 0): # TODO(Isotr0py): Remove this patch after upstream fix PR # released and Transformers version update: # https://github.com/huggingface/transformers/pull/41473 - mm_kwargs = dict(mm_kwargs) - mm_kwargs["audio_kwargs"] = dict(mm_kwargs.get("audio_kwargs") or {}) - mm_kwargs["text_kwargs"] = dict(mm_kwargs.get("text_kwargs") or {}) - elif mm_kwargs.get("use_audio_in_video") and mm_data.get("videos"): - # mm_data can be empty on a multimodal-processor-cache hit, where - # there's nothing to (re-)process this call and the real result - # comes from the cache — not a genuine "no audio" case. - raise ValueError( - "Video doesn't have audio track with `audio_in_video=True`" + hf_processor_mm_kwargs = dict(hf_processor_mm_kwargs) + hf_processor_mm_kwargs["audio_kwargs"] = dict( + hf_processor_mm_kwargs.get("audio_kwargs") or {} + ) + hf_processor_mm_kwargs["text_kwargs"] = dict( + hf_processor_mm_kwargs.get("text_kwargs") or {} ) - hf_inputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, + processed_data = super()._apply_hf_processor_main( + mm_items, + hf_processor_mm_kwargs, ) if ( - "audio_feature_lengths" in hf_inputs - and "feature_attention_mask" in hf_inputs - and (audios := mm_data.get("audio", [])) + "audio_feature_lengths" in processed_data + and "feature_attention_mask" in processed_data ): - audio_num_frames = [] - for _, audio in enumerate(audios): - audio_length = len(audio[0]) if isinstance(audio, tuple) else len(audio) - num_frame = ( - (audio_length // hop_length) - if audio_length % hop_length == 0 - else (audio_length // hop_length - 1) + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + processor_data, _ = self._get_hf_mm_data(valid_mm_items) + audios = processor_data.get("audios", []) + if audios: + feature_extractor = self.info.get_feature_extractor( + **hf_processor_mm_kwargs ) - if mm_kwargs.get("truncation", False): - num_frame = min( - num_frame, feature_extractor.n_samples // hop_length + hop_length = feature_extractor.hop_length + + audio_num_frames = [] + for audio in audios: + audio_length = ( + len(audio[0]) if isinstance(audio, tuple) else len(audio) ) - audio_num_frames.append(num_frame) - hf_inputs["feature_attention_mask"] = [ - torch.ones(num_frame) for num_frame in audio_num_frames - ] - hf_inputs["audio_feature_lengths"] = torch.tensor(audio_num_frames) - return hf_inputs + num_frame = ( + (audio_length // hop_length) + if audio_length % hop_length == 0 + else (audio_length // hop_length - 1) + ) + if hf_processor_mm_kwargs.get("truncation", False): + num_frame = min( + num_frame, feature_extractor.n_samples // hop_length + ) + audio_num_frames.append(num_frame) + + processed_data["feature_attention_mask"] = [ + torch.ones(num_frame) for num_frame in audio_num_frames + ] + processed_data["audio_feature_lengths"] = torch.tensor(audio_num_frames) + + return processed_data def _maybe_apply_prompt_updates( self, diff --git a/vllm/model_executor/models/qwen3_vl.py b/vllm/model_executor/models/qwen3_vl.py index bc31091cfe37..9390b7bacbcd 100644 --- a/vllm/model_executor/models/qwen3_vl.py +++ b/vllm/model_executor/models/qwen3_vl.py @@ -1280,12 +1280,21 @@ def _expands_only_video_token(hf_processor: ProcessorMixin) -> bool: proc_impl = getattr(type(hf_processor), "replace_video_token", None) return proc_impl is not None and proc_impl is not mixin_impl - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + mm_data = dict(mm_data) # Separate video processing from image processing. Because the videos @@ -1314,8 +1323,8 @@ def _call_hf_processor( # NOTE: a copy of is created to update do_sample_frames, # otherwise mm_hash for the object will be incorrect. - video_mm_kwargs = dict(**mm_kwargs) - merged = self.info.ctx.get_merged_mm_kwargs(mm_kwargs) + video_mm_kwargs = dict(**hf_processor_mm_kwargs) + merged = self.info.ctx.get_merged_mm_kwargs(hf_processor_mm_kwargs) if merged.keys() & {"size", "min_pixels", "max_pixels"}: video_size = dict(self.info.get_video_processor().size) size_override = merged.get("size") @@ -1365,10 +1374,13 @@ def _call_hf_processor( if "num_frames" in video_mm_kwargs and "fps" not in video_mm_kwargs: video_mm_kwargs["fps"] = None - video_outputs = super()._call_hf_processor( - prompt="<|vision_start|><|video_pad|><|vision_end|>", - mm_data=video_mm_data, - mm_kwargs=video_mm_kwargs, + video_outputs = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**video_mm_kwargs), + dict( + text="<|vision_start|><|video_pad|><|vision_end|>", + **video_mm_data, + ), + video_mm_kwargs, ) # Discard HF output input_ids — we use get_video_repl below @@ -1424,12 +1436,14 @@ def _call_hf_processor( # fps/num_frames are video-only kwargs already consumed by the loop; # exclude them so the text/image processor call below never gets a list. non_video_mm_kwargs = { - k: v for k, v in mm_kwargs.items() if k not in ("fps", "num_frames") + k: v + for k, v in hf_processor_mm_kwargs.items() + if k not in ("fps", "num_frames") } - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=non_video_mm_kwargs, + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**non_video_mm_kwargs), + dict(text=prompt_text, **mm_data), + non_video_mm_kwargs, ) # Replace each placeholder with pre-computed video tokens. @@ -1443,20 +1457,19 @@ def _call_hf_processor( hf_config.video_token_id, hf_config.vision_end_token_id, ] - input_ids = processed_outputs.pop("input_ids") + input_ids = processed_data.pop("input_ids") if not isinstance(input_ids, list): input_ids = input_ids.tolist() (prompt_ids,) = input_ids expanded_ids = _replace_video_token_placeholders( prompt_ids, video_target, video_input_ids_lst ) - processed_outputs["input_ids"] = [expanded_ids] + processed_data["input_ids"] = [expanded_ids] - combined_outputs = dict( - processed_outputs, - **video_outputs, - ) - return BatchFeature(combined_outputs) + processed_data.update(video_outputs) + processed_data.update(passthrough_data) + + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/transformers/multimodal.py b/vllm/model_executor/models/transformers/multimodal.py index f373a2cccb55..255c67cf1a65 100644 --- a/vllm/model_executor/models/transformers/multimodal.py +++ b/vllm/model_executor/models/transformers/multimodal.py @@ -345,28 +345,6 @@ def _get_mm_fields_config( return mm_fields - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> "BatchFeature": - """Run the HF processor and unpad inputs.""" - try: - hf_inputs = super()._call_hf_processor(prompt, mm_data, mm_kwargs) - except ValueError: - if any(mm_data.values()): - raise - # Some processors reject a prompt holding placeholders with - # no data to go with them, so tokenize it without them - prompt_ids = self.info.get_tokenizer().encode(prompt) - hf_inputs = transformers.BatchFeature( - dict(input_ids=[prompt_ids]), tensor_type="pt" - ) - self._unpad_images(hf_inputs) - self._unpad_audios(hf_inputs, mm_data, mm_kwargs) - return hf_inputs - def _unpad_audios( self, hf_inputs: "BatchFeature", @@ -450,19 +428,6 @@ def _get_prompt_updates( deriving them from updates.""" return [] - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> "BatchFeature": - """Ask which modality owns each token of the expanded prompt, which is the - only marking of the tokens an expansion adds around an item.""" - if any(mm_data.values()): - mm_data = {**mm_data, "return_mm_token_type_ids": True} - - return super()._call_hf_processor(prompt, mm_data, mm_kwargs) - def _get_mm_token_ids(self, modality: str) -> list[int]: """Token ids marking where `modality` sits in the prompt, which for some processors differ from the placeholder written into it. @@ -612,24 +577,55 @@ def apply( with timing_ctx.record("apply_hf_processor"): hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) - prompt = hf_processor.decode(prompt) + prompt_text = hf_processor.decode(prompt) # Bypass cached processor and always apply to the full set of mm inputs # NOTE: we can't just set caching=False because base class method # transforms outputs to `MultiModalKwargs` which is not going to # work for Transformers. The vision path has logic tied to # `mm_tokens_per_modality` in _apply_vision() - prompt_ids, processed_data = self._apply_hf_processor_text_mm( - prompt_text=prompt, - mm_items=mm_items, - hf_processor_mm_kwargs={ - **hf_processor_mm_kwargs, - # HF processors only accept text, and the decoded string already - # contains any special tokens, so don't let them be added again - # (`_call_hf_processor` disables `add_special_tokens`). - "add_special_tokens": False, - }, + hf_processor_mm_kwargs = { + **hf_processor_mm_kwargs, + # HF processors only accept text, and the decoded string already + # contains any special tokens, so don't let them be added again + # (the HF processor call disables `add_special_tokens`). + "add_special_tokens": False, + } + + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} ) + processor_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + # Ask which modality owns each token of the expanded prompt, which + # is the only marking of the tokens an expansion adds around an item + if any(processor_data.values()): + processor_data = {**processor_data, "return_mm_token_type_ids": True} + + try: + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(text=prompt_text, **processor_data), + hf_processor_mm_kwargs, + ) + except ValueError: + if any(processor_data.values()): + raise + # Some processors reject a prompt holding placeholders with + # no data to go with them, so tokenize it without them + prompt_ids = self.info.get_tokenizer().encode(prompt_text) + processed_data = transformers.BatchFeature( + dict(input_ids=[prompt_ids]), tensor_type="pt" + ) + self._unpad_images(processed_data) + self._unpad_audios(processed_data, processor_data, hf_processor_mm_kwargs) + processed_data.update(passthrough_data) + + input_ids = processed_data.pop("input_ids") + if not isinstance(input_ids, list): + input_ids = input_ids.tolist() + + (prompt_ids,) = input_ids # Use overrides if provided; fallback to data-dependent hashing. with timing_ctx.record("get_mm_hashes"): @@ -773,17 +769,41 @@ def _get_num_patches_per_image( except (AttributeError, KeyError, TypeError): return None - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> "BatchFeature": - """Ask for the replacement each placeholder expands to, and record it as - per-item fields: its token ids, and the tokens or patches behind them.""" + mm_counts = mm_items.get_all_counts() + + valid_mm_items = mm_items.select({k for k, c in mm_counts.items() if c > 0}) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_counts) + + # Ask for the replacement each placeholder expands to, and record it as + # per-item fields: its token ids, and the tokens or patches behind them if has_mm_data := any(mm_data.values()): mm_data = {**mm_data, "return_text_replacement_offsets": True} - hf_inputs = super()._call_hf_processor(prompt, mm_data, mm_kwargs) + + try: + hf_inputs = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), + dict(text=prompt_text, **mm_data), + hf_processor_mm_kwargs, + ) + except ValueError: + if any(mm_data.values()): + raise + # Some processors reject a prompt holding placeholders with + # no data to go with them, so tokenize it without them + tokenizer = self.info.get_tokenizer() + hf_inputs = transformers.BatchFeature( + dict(input_ids=[tokenizer.encode(prompt_text)]), tensor_type="pt" + ) + self._unpad_images(hf_inputs) + self._unpad_audios(hf_inputs, mm_data, hf_processor_mm_kwargs) + # Drop the inputs the model would reject hf_inputs.pop("mm_token_type_ids", None) hf_inputs.pop("token_type_ids", None) @@ -802,6 +822,8 @@ def _call_hf_processor( "so it can be fixed, and install `transformers<5.15.0` in the " "meantime to locate placeholders in the expanded prompt instead." ) + hf_inputs.update(passthrough_data) + hf_inputs.pop("input_ids", None) return hf_inputs tokenizer = self.info.get_tokenizer() @@ -829,6 +851,9 @@ def _call_hf_processor( counts.append(int(ids.eq(_get_embed_token_id(ids)).sum())) hf_inputs["num_audio_tokens"] = torch.tensor(counts) + hf_inputs.update(passthrough_data) + hf_inputs.pop("input_ids") + return hf_inputs diff --git a/vllm/model_executor/models/ultravox.py b/vllm/model_executor/models/ultravox.py index 81e97ce4ce37..60dc1fb133fb 100644 --- a/vllm/model_executor/models/ultravox.py +++ b/vllm/model_executor/models/ultravox.py @@ -190,41 +190,84 @@ def get_dummy_mm_data( class UltravoxMultiModalProcessor(BaseMultiModalProcessor[UltravoxProcessingInfo]): - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _preprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - # Text-only input not supported in composite processor - if not mm_data.get("audios", []): - prompt_ids = self.info.get_tokenizer().encode( - prompt, add_special_tokens=False - ) - prompt_ids = self._apply_hf_processor_tokens_only(prompt_ids) - return BatchFeature(dict(input_ids=[prompt_ids]), tensor_type="pt") - - mm_data = dict(mm_data) - audios = mm_data.pop("audios", []) + hf_processor_mm_kwargs: Mapping[str, object], + ) -> tuple[Mapping[str, object], Mapping[str, object]]: + audios = mm_data.get("audios", []) assert isinstance(audios, list) - feature_extractor = self.info.get_feature_extractor(**mm_kwargs) - mm_kwargs = dict( - **mm_kwargs, + feature_extractor = self.info.get_feature_extractor(**hf_processor_mm_kwargs) + hf_processor_mm_kwargs = dict( + **hf_processor_mm_kwargs, sampling_rate=feature_extractor.sampling_rate, include_audio_num_chunks=True, ) - item_processor_data = dict(**mm_data, audios=audios) + return mm_data, hf_processor_mm_kwargs - output = super()._call_hf_processor( - prompt=prompt, - mm_data=item_processor_data, - mm_kwargs=mm_kwargs, + def _apply_hf_processor_main( + self, + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], + ) -> BatchFeature: + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} ) - output["audio_features"] = output.pop("audio_values") + processor_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if processor_data: + processor_data, hf_processor_mm_kwargs = self._preprocess_hf_mm_data( + processor_data, + hf_processor_mm_kwargs, + ) + + prompt_text = self._get_hf_processor_text(mm_items.get_all_counts()) + if prompt_text is not None: + processor_data = dict(text=prompt_text, **processor_data) + + hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) + + def patched_call(**kwargs) -> BatchFeature: + # The remote UltravoxProcessor does not support the + # `truncation` kwarg injected by `call_hf_processor`: it + # passes `truncation` to its audio processor by itself while + # also forwarding `**kwargs` there, causing a duplicate + # keyword argument error. Its hardcoded `truncation=False` + # already matches vLLM's intended behavior. + kwargs.pop("truncation", None) + + return hf_processor(**kwargs) + + processed_data = self.info.ctx.call_hf_processor( + patched_call, + processor_data, + hf_processor_mm_kwargs, + ) + processed_data.update(passthrough_data) + else: + processed_data = BatchFeature(dict(passthrough_data)) + + return self._postprocess_hf_mm_data( + processor_data, + hf_processor_mm_kwargs, + processed_data, + ) + + def _postprocess_hf_mm_data( + self, + mm_data: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, + ) -> BatchFeature: + if "audio_values" in processed_data: + processed_data["audio_features"] = processed_data.pop("audio_values") - return output + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/unlimited_ocr.py b/vllm/model_executor/models/unlimited_ocr.py index 0f29337314e5..c96e5f293d91 100644 --- a/vllm/model_executor/models/unlimited_ocr.py +++ b/vllm/model_executor/models/unlimited_ocr.py @@ -225,7 +225,7 @@ def _cached_apply_hf_processor( self, inputs: ProcessorInputs, timing_ctx: TimingContext, - ) -> tuple[list[int], MultiModalProcessingInfo, bool]: + ) -> MultiModalProcessingInfo: # The processor logic differs for single-image (crop) vs multi-image # (no crop) requests. The processing cache assumes per-item output is # invariant of how many images are passed per prompt, so we only cache diff --git a/vllm/model_executor/models/voxtral.py b/vllm/model_executor/models/voxtral.py index 2013d48ce576..07a7f9c32550 100644 --- a/vllm/model_executor/models/voxtral.py +++ b/vllm/model_executor/models/voxtral.py @@ -47,12 +47,10 @@ from vllm.multimodal.processing.processor import ( BaseMultiModalProcessor, BaseProcessingInfo, - MultiModalProcessingInfo, PlaceholderFeaturesInfo, ProcessorInputs, PromptReplacement, PromptUpdate, - TimingContext, ) from vllm.sequence import IntermediateTensors from vllm.tokenizers import cached_tokenizer_from_config @@ -204,6 +202,9 @@ def get_dummy_processor_inputs( class VoxtralMultiModalProcessor(BaseMultiModalProcessor[VoxtralProcessingInfo]): + # The tokens are already inserted by the chat template + hf_processor_applies_updates = True + def _get_mm_fields_config( self, hf_inputs: Mapping[str, NestedTensors], @@ -220,12 +221,14 @@ def _validate_mm_placeholders( # skip validation here pass - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _preprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: + hf_processor_mm_kwargs: Mapping[str, object], + ) -> tuple[Mapping[str, object], Mapping[str, object]]: mm_data = dict(mm_data) audios = mm_data.pop("audios", []) @@ -233,18 +236,23 @@ def _call_hf_processor( # MistralCommonVoxtralProcessor accepts "audio" mm_data["audio"] = audios - outputs = self.info.ctx.call_hf_processor( - self.info.get_hf_processor(**mm_kwargs), - dict(text=prompt, **mm_data), - # Avoid padding issue - dict(**mm_kwargs, return_tensors=None), - ) + # Avoid padding issue + hf_processor_mm_kwargs = dict(**hf_processor_mm_kwargs, return_tensors=None) + + return mm_data, hf_processor_mm_kwargs + def _postprocess_hf_mm_data( + self, + mm_data: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, + ) -> BatchFeature: # Missing batch dimension - if is_list_of(outputs["input_ids"], int): - outputs["input_ids"] = [outputs["input_ids"]] + input_ids = processed_data.get("input_ids") + if is_list_of(input_ids, int): + processed_data["input_ids"] = [input_ids] - return outputs + return processed_data def _get_prompt_updates( self, @@ -287,16 +295,6 @@ def get_replacement(item_idx: int): ), ] - def _cached_apply_hf_processor( - self, - inputs: ProcessorInputs, - timing_ctx: TimingContext, - ) -> tuple[list[int], MultiModalProcessingInfo, bool]: - prompt_ids, mm_info, _ = super()._cached_apply_hf_processor(inputs, timing_ctx) - - # NOTE: The tokens are already inserted by the chat template - return prompt_ids, mm_info, True - @MULTIMODAL_REGISTRY.register_processor( VoxtralMultiModalProcessor, diff --git a/vllm/model_executor/models/whisper.py b/vllm/model_executor/models/whisper.py index 0486933c9185..d21903bd656c 100644 --- a/vllm/model_executor/models/whisper.py +++ b/vllm/model_executor/models/whisper.py @@ -756,27 +756,36 @@ def create_encoder_prompt( # for encoder profiling. return [0] - def _call_hf_processor( + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) + + def _preprocess_hf_mm_data( self, - prompt: str, mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - if mm_data: - feature_extractor = self.info.get_feature_extractor(**mm_kwargs) - mm_data = dict(audio=mm_data.pop("audios")) - mm_kwargs = dict( - **mm_kwargs, - sampling_rate=feature_extractor.sampling_rate, - ) - processed_outputs = super()._call_hf_processor( - prompt=prompt, - mm_data=mm_data, - mm_kwargs=mm_kwargs, + hf_processor_mm_kwargs: Mapping[str, object], + ) -> tuple[Mapping[str, object], Mapping[str, object]]: + feature_extractor = self.info.get_feature_extractor(**hf_processor_mm_kwargs) + + mm_data = dict(mm_data) + mm_data["audio"] = mm_data.pop("audios") + + hf_processor_mm_kwargs = dict( + **hf_processor_mm_kwargs, + sampling_rate=feature_extractor.sampling_rate, ) - if "labels" in processed_outputs: - processed_outputs["input_ids"] = processed_outputs.pop("labels") - return processed_outputs + + return mm_data, hf_processor_mm_kwargs + + def _postprocess_hf_mm_data( + self, + mm_data: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, + ) -> BatchFeature: + if "labels" in processed_data: + processed_data["input_ids"] = processed_data.pop("labels") + + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/models/dots3_note/common/processor.py b/vllm/models/dots3_note/common/processor.py index 25c8c5de5820..724d9b5ac17c 100644 --- a/vllm/models/dots3_note/common/processor.py +++ b/vllm/models/dots3_note/common/processor.py @@ -649,13 +649,8 @@ def _get_hf_mm_data( processor_data["videos"] = raw_videos return processor_data, passthrough_data - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - return super()._call_hf_processor(prompt, mm_data, mm_kwargs) + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: + return self.dummy_inputs.get_dummy_text(mm_counts) def _get_mm_fields_config( self, diff --git a/vllm/models/inkling/common/mm_preprocess.py b/vllm/models/inkling/common/mm_preprocess.py index df1464e319f9..249d7edeb43f 100644 --- a/vllm/models/inkling/common/mm_preprocess.py +++ b/vllm/models/inkling/common/mm_preprocess.py @@ -22,7 +22,7 @@ MultiModalFieldConfig, MultiModalKwargsItems, ) -from vllm.multimodal.parse import MultiModalDataParser +from vllm.multimodal.parse import MultiModalDataItems, MultiModalDataParser from vllm.multimodal.processing import ( BaseDummyInputsBuilder, BaseMultiModalProcessor, @@ -135,8 +135,8 @@ def get_dummy_text(self, mm_counts: Mapping[str, int]) -> str: num_images = mm_counts.get("image", 0) num_audios = mm_counts.get("audio", 0) # Use spellings the renderer would emit; tokenization is bypassed in - # _call_hf_processor (we build input_ids directly), so the exact text - # only needs to be a stable per-item marker. + # _apply_hf_processor_main (we build input_ids directly), so the exact + # text only needs to be a stable per-item marker. return ("<|content_image|>" * num_images) + ( "<|content_audio_input|>" * num_audios ) @@ -179,17 +179,19 @@ def get_dummy_mm_data( class InklingMultiModalProcessor(BaseMultiModalProcessor[InklingProcessingInfo]): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: - # Inkling is not a standard HF processor (no fused text+mm call), so we run - # the vendored extractors ourselves and tokenize the text separately. - # The MM placeholders in `prompt` are expanded later by the prompt - # updates, so here we emit ONE placeholder id per media item. - processor = self.info.get_hf_processor(**mm_kwargs) + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + + processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) tokenizer = self.info.get_tokenizer() images = mm_data.get("images") or [] @@ -200,7 +202,7 @@ def _call_hf_processor( audios = list(cast(Iterable[Any], audios)) prompt_ids = self._tokenize_with_placeholders( - prompt, tokenizer, len(images), len(audios) + prompt_text, tokenizer, len(images), len(audios) ) data: dict[str, Any] = {"input_ids": [prompt_ids]} @@ -232,7 +234,9 @@ def _call_hf_processor( data["input_audio_features"] = input_audio_features data["num_audio_tokens"] = torch.tensor(num_audio_tokens, dtype=torch.int64) - return BatchFeature(data=data, tensor_type=None) + processed_data = BatchFeature(data=data, tensor_type=None) + processed_data.update(passthrough_data) + return processed_data def _tokenize_with_placeholders( self, diff --git a/vllm/models/kimi_k3/common/mm_preprocess.py b/vllm/models/kimi_k3/common/mm_preprocess.py index bc8386b4a6b7..70c017fd7258 100644 --- a/vllm/models/kimi_k3/common/mm_preprocess.py +++ b/vllm/models/kimi_k3/common/mm_preprocess.py @@ -231,19 +231,6 @@ def get_dummy_mm_data( class KimiK3MultiModalProcessor(BaseMultiModalProcessor[KimiK3ProcessingInfo]): """Image-only multi-modal processor for Kimi-K3.""" - def _call_hf_processor( - self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: - # Override so the base always routes through the text+mm path - # (`KimiK3Processor.__call__`). Otherwise the mm-only fast path calls - # the checkpoint image processor directly with bare PIL images, but it - # requires `{"type": "image", "image": PIL}` media dicts that only our - # wrapper builds. - return super()._call_hf_processor(prompt, mm_data, mm_kwargs) - def _get_mm_fields_config( self, hf_inputs: BatchFeature, diff --git a/vllm/models/minimax_m3/common/mm_preprocess.py b/vllm/models/minimax_m3/common/mm_preprocess.py index 1481da4dc86f..86b702585fd8 100644 --- a/vllm/models/minimax_m3/common/mm_preprocess.py +++ b/vllm/models/minimax_m3/common/mm_preprocess.py @@ -306,12 +306,21 @@ def get_dummy_mm_data( class MiniMaxM3VLMultiModalProcessor( BaseMultiModalProcessor[MiniMaxM3VLProcessingInfo] ): - def _call_hf_processor( + def _apply_hf_processor_main( self, - prompt: str, - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], + mm_items: MultiModalDataItems, + hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: + valid_mm_items = mm_items.select( + {k for k, c in mm_items.get_all_counts().items() if c > 0} + ) + mm_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) + + if not mm_data: + return BatchFeature(dict(passthrough_data)) + + prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + mm_data = dict(mm_data) # With ``video_needs_metadata=True`` each video arrives as a # ``(frames, metadata)`` tuple. Split the frames back out and forward the @@ -340,15 +349,17 @@ def _call_hf_processor( # Override the video processor's default do_resize=False (set for a # pre-resized pipeline) to True for vLLM's raw-frame inputs. - merged = dict(do_resize=True, **mm_kwargs) - data = dict(text=prompt, **mm_data) + merged = dict(do_resize=True, **hf_processor_mm_kwargs) + data = dict(text=prompt_text, **mm_data) if video_metadata is not None: data["video_metadata"] = video_metadata - return self.info.ctx.call_hf_processor( - self.info.get_hf_processor(**mm_kwargs), + processed_data = self.info.ctx.call_hf_processor( + self.info.get_hf_processor(**hf_processor_mm_kwargs), data, merged, ) + processed_data.update(passthrough_data) + return processed_data def _get_mm_fields_config( self, diff --git a/vllm/multimodal/processing/context.py b/vllm/multimodal/processing/context.py index 45acdf8f4d7b..f13ca6ab88f6 100644 --- a/vllm/multimodal/processing/context.py +++ b/vllm/multimodal/processing/context.py @@ -262,9 +262,6 @@ def call_hf_processor( hf_processor: Callable[..., BatchFeature] | ProcessorMixin, data: Mapping[str, object], kwargs: Mapping[str, object] = {}, - *, - num_tries: int = 1, - max_tries: int = 5, ) -> BatchFeature: """ Call `hf_processor` on the prompt `data` @@ -274,6 +271,12 @@ def call_hf_processor( merged_kwargs = self.get_merged_mm_kwargs(kwargs) + # vLLM needs the full untruncated sequence to keep multi-modal + # placeholder tokens aligned; note that the text inputs in + # call_hf_processor are just dummy text, not the original prompt. + # The original prompt is already tokenized by the renderer. + merged_kwargs.setdefault("truncation", False) + allowed_kwargs = get_allowed_kwarg_only_overrides( hf_processor, merged_kwargs, diff --git a/vllm/multimodal/processing/dummy_inputs.py b/vllm/multimodal/processing/dummy_inputs.py index 09df2e05328f..ac3fe547217f 100644 --- a/vllm/multimodal/processing/dummy_inputs.py +++ b/vllm/multimodal/processing/dummy_inputs.py @@ -91,10 +91,7 @@ def get_dummy_processor_inputs( # prompt tokens. dummy_prompt = [] else: - dummy_prompt = tokenizer.encode( - dummy_text, - **self.info.default_tok_params.get_encode_kwargs(), - ) + dummy_prompt = tokenizer.encode(dummy_text, truncation=False) return ProcessorInputs( prompt=dummy_prompt, diff --git a/vllm/multimodal/processing/processor.py b/vllm/multimodal/processing/processor.py index 599db7406a36..17bcc2f6dd56 100644 --- a/vllm/multimodal/processing/processor.py +++ b/vllm/multimodal/processing/processor.py @@ -5,9 +5,10 @@ from collections.abc import Callable, Generator, ItemsView, Iterable, Mapping, Sequence from dataclasses import dataclass, field, replace from enum import Enum -from functools import lru_cache, partial +from functools import lru_cache from typing import ( TYPE_CHECKING, + ClassVar, Generic, NamedTuple, Protocol, @@ -28,7 +29,6 @@ ) from vllm.logger import init_logger from vllm.tokenizers import TokenizerLike -from vllm.transformers_utils.processor import call_hf_processor_mm_only from vllm.utils.collection_utils import flatten_2d_lists, full_groupby from ..inputs import ( @@ -1177,6 +1177,12 @@ class BaseMultiModalProcessor(ABC, Generic[_I]): Not to be confused with `transformers.ProcessorMixin`. """ + hf_processor_applies_updates: ClassVar[bool] = False + """ + Only set to True if the tokenizer or chat template + (which are run before MM processing) inserts placeholder tokens. + """ + def __init__( self, info: _I, @@ -1302,131 +1308,100 @@ def _get_hf_mm_data( return processor_data, passthrough_data - def _call_hf_processor( - self, - prompt: str, - # Not to be confused with `mm_data` in `self.apply`. - # This refers to the data to be passed to HF processor. - mm_data: Mapping[str, object], - mm_kwargs: Mapping[str, object], - ) -> BatchFeature: + def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str | None: """ - Call the HF processor on the prompt text and - associated multi-modal data. + Get the text to pass to the HF processor alongside the multi-modal + data. + + By default, no text is passed. If the HF processor requires that + text and multi-modal items correspond to each other, you should + override this method to return dummy text generated by + [`DummyInputsBuilder`][vllm.multimodal.processing.BaseDummyInputsBuilder]. """ - return self.info.ctx.call_hf_processor( - self.info.get_hf_processor(**mm_kwargs), - dict(text=prompt, **mm_data), - mm_kwargs, - ) + return None - def _apply_hf_processor_text_mm( + def _preprocess_hf_mm_data( self, - prompt_text: str, - mm_items: MultiModalDataItems, + mm_data: Mapping[str, object], hf_processor_mm_kwargs: Mapping[str, object], - ) -> tuple[list[int], BatchFeature]: + ) -> tuple[Mapping[str, object], Mapping[str, object]]: """ - Apply the HF processor on the prompt text and multi-modal data - together. + Pre-process the multi-modal data and HF processor keyword arguments + before they are passed to the HF processor. + + By default, both are returned as-is. If the HF processor expects the + multi-modal data under different keys than those provided by the + multi-modal items (e.g. `audio` instead of `audios`), or requires + additional keyword arguments (e.g. `sampling_rate`), you should + override this method. """ - valid_mm_items = mm_items.select( - {k for k, c in mm_items.get_all_counts().items() if c > 0} - ) - processor_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) - - processed_data = self._call_hf_processor( - prompt=prompt_text, - mm_data=processor_data, - mm_kwargs=hf_processor_mm_kwargs, - ) - processed_data.update(passthrough_data) - - input_ids = processed_data.pop("input_ids") - if not isinstance(input_ids, list): - input_ids = input_ids.tolist() - - (prompt_ids,) = input_ids + return mm_data, hf_processor_mm_kwargs - return prompt_ids, processed_data - - def _apply_hf_processor_tokens_only( + def _postprocess_hf_mm_data( self, - prompt_tokens: list[int], - ) -> list[int]: + mm_data: Mapping[str, object], + hf_processor_mm_kwargs: Mapping[str, object], + processed_data: BatchFeature, + ) -> BatchFeature: """ - Apply the HF processor on the prompt tokens only. + Post-process the output of the HF processor. - Most HF processors accept prompt text but not prompt tokens. - If the HF processor adds or removes tokens that are not related to - multi-modal data, you should override this method so it is consistent - with the tokenization of the corresponding text. + By default, the output is returned as-is. If you need to modify the + output of the HF processor before it is converted into multi-modal + keyword arguments, you should override this method. """ - return prompt_tokens + return processed_data - def _apply_hf_processor_mm_only( + def _apply_hf_processor_main( self, mm_items: MultiModalDataItems, hf_processor_mm_kwargs: Mapping[str, object], ) -> BatchFeature: """ - Apply the HF processor on the multi-modal data only. - - Since HF processor requires that text and multi-modal items - correspond to each other, we generate dummy text using - [`DummyInputsBuilder`][vllm.multimodal.processing.BaseDummyInputsBuilder] - to go along with the multi-modal data. + Apply the HF processor on the multi-modal data. """ - # Custom logic based on text inputs - if type(self)._call_hf_processor != BaseMultiModalProcessor._call_hf_processor: - mm_counts = mm_items.get_all_counts() - - _, mm_processed_data = self._apply_hf_processor_text_mm( - prompt_text=self.dummy_inputs.get_dummy_text(mm_counts), - mm_items=mm_items, - hf_processor_mm_kwargs=hf_processor_mm_kwargs, - ) - - return mm_processed_data - valid_mm_items = mm_items.select( {k for k, c in mm_items.get_all_counts().items() if c > 0} ) processor_data, passthrough_data = self._get_hf_mm_data(valid_mm_items) - processed_data = self.info.ctx.call_hf_processor( - partial( - call_hf_processor_mm_only, + if processor_data: + processor_data, hf_processor_mm_kwargs = self._preprocess_hf_mm_data( + processor_data, hf_processor_mm_kwargs + ) + + prompt_text = self._get_hf_processor_text(mm_items.get_all_counts()) + if prompt_text is not None: + processor_data = dict(text=prompt_text, **processor_data) + + processed_data = self.info.ctx.call_hf_processor( self.info.get_hf_processor(**hf_processor_mm_kwargs), - ), + processor_data, + hf_processor_mm_kwargs, + ) + processed_data.update(passthrough_data) + else: + from transformers.feature_extraction_utils import BatchFeature + + processed_data = BatchFeature(dict(passthrough_data)) + + return self._postprocess_hf_mm_data( processor_data, hf_processor_mm_kwargs, + processed_data, ) - processed_data.update(passthrough_data) - - return processed_data - def _apply_hf_processor_main( - self, - prompt: list[int], - mm_items: MultiModalDataItems, - hf_processor_mm_kwargs: Mapping[str, object], - ) -> tuple[list[int], BatchFeature, bool]: + def _postprocess_prompt(self, prompt: list[int]) -> list[int]: """ - Apply the HF processor on the prompt tokens and multi-modal data. + Post-process the prompt token IDs before locating or applying + multi-modal prompt updates. - In addition, return whether prompt updates have been applied - (`False` here, since for token prompts the updates are normally - applied separately via `_maybe_apply_prompt_updates`). + By default, the prompt is returned as-is. If the HF processor (or + chat template) applies additional transformations to the prompt + that are not reflected in its multi-modal outputs, you should + override this method to replicate them. """ - prompt_ids = self._apply_hf_processor_tokens_only(prompt) - - mm_processed_data = self._apply_hf_processor_mm_only( - mm_items=mm_items, - hf_processor_mm_kwargs=hf_processor_mm_kwargs, - ) - - return prompt_ids, mm_processed_data, False + return prompt def _get_cache_missing_items( self, @@ -1531,14 +1506,9 @@ def _apply_hf_processor( self, inputs: ProcessorInputs, timing_ctx: TimingContext, - ) -> tuple[list[int], MultiModalProcessingInfo, bool]: + ) -> MultiModalProcessingInfo: with timing_ctx.record("apply_hf_processor"): - ( - prompt_ids, - mm_processed_data, - is_update_applied, - ) = self._apply_hf_processor_main( - prompt=inputs.prompt, + mm_processed_data = self._apply_hf_processor_main( mm_items=inputs.mm_data_items, hf_processor_mm_kwargs=inputs.hf_processor_mm_kwargs, ) @@ -1569,13 +1539,13 @@ def _apply_hf_processor( prompt_updates=mm_prompt_updates, ) - return prompt_ids, mm_info, is_update_applied + return mm_info def _cached_apply_hf_processor( self, inputs: ProcessorInputs, timing_ctx: TimingContext, - ) -> tuple[list[int], MultiModalProcessingInfo, bool]: + ) -> MultiModalProcessingInfo: """ Apply the HF processor on the full prompt text, caching the results and reusing cached results. @@ -1599,16 +1569,11 @@ def _cached_apply_hf_processor( mm_hashes=mm_hashes, ) - # NOTE: `prompt` does not correspond to `mm_missing_data_items`, + # NOTE: The prompt does not correspond to `mm_missing_data_items`, # so we can't apply prompt updates until the new multimodal # items are combined with the cached multimodal items with timing_ctx.record("apply_hf_processor"): - ( - prompt_ids, - mm_missing_processed_data, - is_update_applied, - ) = self._apply_hf_processor_main( - prompt=inputs.prompt, + mm_missing_processed_data = self._apply_hf_processor_main( mm_items=mm_missing_data_items, hf_processor_mm_kwargs=inputs.hf_processor_mm_kwargs, ) @@ -1641,7 +1606,7 @@ def _cached_apply_hf_processor( prompt_updates=mm_prompt_updates, ) - return prompt_ids, mm_info, is_update_applied + return mm_info def _apply_token_matches( self, @@ -1770,7 +1735,7 @@ def _validate_mm_kwargs( "There is likely a problem with your " "implementation of merged multi-modal processor for this " "model (usually arising from an inconsistency between " - "`_call_hf_processor` and `_get_mm_fields_config`)." + "`_apply_hf_processor_main` and `_get_mm_fields_config`)." ) def _validate_mm_updates( @@ -1805,8 +1770,8 @@ def _validate_mm_placeholders( f"Expected there to be {item_count} prompt placeholders " f"corresponding to {item_count} {modality} items, but " f"instead found {len(placeholders)} prompt placeholders! " - "Make sure the implementation of `_call_hf_processor` and " - "`_get_mm_fields_config` are consistent with each other." + "Make sure the implementation of `_apply_hf_processor_main` " + "and `_get_mm_fields_config` are consistent with each other." ) def _maybe_apply_prompt_updates( @@ -1854,11 +1819,8 @@ def apply( 3. Extract information about the placeholder tokens from the processed token IDs. """ - ( - prompt_ids, - mm_info, - is_update_applied, - ) = self._cached_apply_hf_processor(inputs, timing_ctx) + prompt_ids = self._postprocess_prompt(inputs.prompt) + mm_info = self._cached_apply_hf_processor(inputs, timing_ctx) with timing_ctx.record("apply_prompt_updates"): prompt_ids, mm_placeholders = self._maybe_apply_prompt_updates( @@ -1866,7 +1828,7 @@ def apply( prompt_ids=prompt_ids, mm_kwargs=mm_info.kwargs, mm_prompt_updates=mm_info.prompt_updates, - is_update_applied=is_update_applied, + is_update_applied=self.hf_processor_applies_updates, ) mm_placeholder_ranges = { diff --git a/vllm/transformers_utils/processor.py b/vllm/transformers_utils/processor.py index ee73393f6b9f..64a94950d4ff 100644 --- a/vllm/transformers_utils/processor.py +++ b/vllm/transformers_utils/processor.py @@ -11,17 +11,13 @@ AutoImageProcessor, AutoProcessor, AutoVideoProcessor, - BatchFeature, processing_utils, ) -from transformers.audio_utils import AudioInput from transformers.feature_extraction_utils import FeatureExtractionMixin from transformers.image_processing_utils import BaseImageProcessor -from transformers.image_utils import ImageInput from transformers.models.auto.video_processing_auto import VIDEO_PROCESSOR_MAPPING_NAMES from transformers.processing_utils import ProcessorMixin from transformers.video_processing_utils import BaseVideoProcessor -from transformers.video_utils import VideoInput from typing_extensions import TypeVar from vllm.logger import init_logger @@ -554,43 +550,3 @@ def cached_video_processor_from_config( processor_cls_overrides=processor_cls, # type: ignore[arg-type] **_merge_mm_kwargs(model_config, AutoVideoProcessor, **kwargs), ) - - -def call_hf_processor_mm_only( - processor: ProcessorMixin, - images: ImageInput | None = None, - videos: VideoInput | None = None, - audio: AudioInput | None = None, - **kwargs, -) -> BatchFeature: - output_kwargs = processor._merge_kwargs( - get_processor_kwargs_type(processor), - **kwargs, - ) - - if audio is not None and ( - feature_extractor := getattr(processor, "feature_extractor", None) - ): - audio_inputs = feature_extractor(audio, **output_kwargs["audio_kwargs"]) - audio_inputs["feature_attention_mask"] = audio_inputs.pop("attention_mask") - else: - audio_inputs = {} - - if images is not None and ( - image_processor := getattr(processor, "image_processor", None) - ): - images_inputs = image_processor(images=images, **output_kwargs["images_kwargs"]) - else: - images_inputs = {} - - if videos is not None and ( - video_processor := getattr(processor, "video_processor", None) - ): - videos_inputs = video_processor(videos=videos, **output_kwargs["videos_kwargs"]) - else: - videos_inputs = {} - - return BatchFeature( - data={**audio_inputs, **images_inputs, **videos_inputs}, - tensor_type=kwargs.get("return_tensors"), - ) From 040700aaa6cae3ef54dad8ba75db1abfb79db449 Mon Sep 17 00:00:00 2001 From: Cyrus Leung Date: Sat, 22 Aug 2026 18:27:59 +0800 Subject: [PATCH 006/235] [MM] Address comments on #53275 (#53364) Signed-off-by: DarkLight1337 --- vllm/model_executor/models/nano_nemotron_vl.py | 2 +- vllm/model_executor/models/pixtral.py | 2 +- vllm/model_executor/models/voxtral.py | 2 +- vllm/multimodal/processing/processor.py | 15 ++++++++------- 4 files changed, 11 insertions(+), 10 deletions(-) diff --git a/vllm/model_executor/models/nano_nemotron_vl.py b/vllm/model_executor/models/nano_nemotron_vl.py index 0e9eda5c82a4..00538fd7b328 100644 --- a/vllm/model_executor/models/nano_nemotron_vl.py +++ b/vllm/model_executor/models/nano_nemotron_vl.py @@ -734,7 +734,7 @@ def apply( prompt_ids=prompt_ids, mm_kwargs=mm_info.kwargs, mm_prompt_updates=mm_info.prompt_updates, - is_update_applied=self.hf_processor_applies_updates, + is_update_applied=self.renderer_applies_updates, ) mm_placeholder_ranges = { diff --git a/vllm/model_executor/models/pixtral.py b/vllm/model_executor/models/pixtral.py index a24d2484ab29..6a210a811b0f 100644 --- a/vllm/model_executor/models/pixtral.py +++ b/vllm/model_executor/models/pixtral.py @@ -212,7 +212,7 @@ def get_dummy_processor_inputs( class PixtralMultiModalProcessor(BaseMultiModalProcessor[PixtralProcessingInfo]): # The tokens are already inserted by the chat template - hf_processor_applies_updates = True + renderer_applies_updates = True def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/voxtral.py b/vllm/model_executor/models/voxtral.py index 07a7f9c32550..20d166ca2f54 100644 --- a/vllm/model_executor/models/voxtral.py +++ b/vllm/model_executor/models/voxtral.py @@ -203,7 +203,7 @@ def get_dummy_processor_inputs( class VoxtralMultiModalProcessor(BaseMultiModalProcessor[VoxtralProcessingInfo]): # The tokens are already inserted by the chat template - hf_processor_applies_updates = True + renderer_applies_updates = True def _get_mm_fields_config( self, diff --git a/vllm/multimodal/processing/processor.py b/vllm/multimodal/processing/processor.py index 17bcc2f6dd56..1eda431a8ae4 100644 --- a/vllm/multimodal/processing/processor.py +++ b/vllm/multimodal/processing/processor.py @@ -1177,9 +1177,9 @@ class BaseMultiModalProcessor(ABC, Generic[_I]): Not to be confused with `transformers.ProcessorMixin`. """ - hf_processor_applies_updates: ClassVar[bool] = False + renderer_applies_updates: ClassVar[bool] = False """ - Only set to True if the tokenizer or chat template + Only set to True if the tokenizer or chat template in the Renderer (which are run before MM processing) inserts placeholder tokens. """ @@ -1811,12 +1811,13 @@ def apply( The main steps are: - 1. Apply HF Processor on prompt text and multi-modal data together, - outputting token IDs and processed tensors. - 2. Find and update sequences in the token IDs with placeholder tokens. + 1. Perform MM-specific post-processing to the token inputs. + 2. Apply HF Processor on multi-modal data together with optional + dummy text, outputting processed multi-modal data tensors. + 3. Find and update sequences in the token IDs with placeholder tokens. The number of placeholder tokens equals the feature size of the multi-modal data outputted by the multi-modal encoder. - 3. Extract information about the placeholder tokens from the + 4. Extract information about the placeholder tokens from the processed token IDs. """ prompt_ids = self._postprocess_prompt(inputs.prompt) @@ -1828,7 +1829,7 @@ def apply( prompt_ids=prompt_ids, mm_kwargs=mm_info.kwargs, mm_prompt_updates=mm_info.prompt_updates, - is_update_applied=self.hf_processor_applies_updates, + is_update_applied=self.renderer_applies_updates, ) mm_placeholder_ranges = { From 10704541aaf72567fe9d6229b3e3d84d37f2ddba Mon Sep 17 00:00:00 2001 From: "Kevin H. Luu" Date: Sat, 22 Aug 2026 03:50:33 -0700 Subject: [PATCH 007/235] [CI][ARM64] Restore known-good CUDA 13 builder image (#53374) Signed-off-by: Kevin Luu <51931015+khluu@users.noreply.github.com> Co-authored-by: OpenAI Codex --- .buildkite/image_build/image_build_arm64.sh | 2 +- .buildkite/release-pipeline.yaml | 6 +++--- .buildkite/scripts/hardware_ci/run-gh200-test.sh | 2 +- 3 files changed, 5 insertions(+), 5 deletions(-) diff --git a/.buildkite/image_build/image_build_arm64.sh b/.buildkite/image_build/image_build_arm64.sh index eb56f02b89e4..4db8a2f5465c 100755 --- a/.buildkite/image_build/image_build_arm64.sh +++ b/.buildkite/image_build/image_build_arm64.sh @@ -26,7 +26,7 @@ else --platform linux/arm64 \ --build-arg max_jobs=16 \ --build-arg nvcc_threads=4 \ - --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda13.0-78e737ad29420ffc4800e677c51e2a852caf8359 \ + --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda13.0-b8b5f17a7d9ccfc25bbc5cf17b3fcea12964a042 \ --build-arg torch_cuda_arch_list="9.0 10.0 11.0 12.0" \ --build-arg USE_SCCACHE=1 \ --build-arg buildkite_commit="$BUILDKITE_COMMIT" \ diff --git a/.buildkite/release-pipeline.yaml b/.buildkite/release-pipeline.yaml index 1d421d664a50..17b6592ad592 100644 --- a/.buildkite/release-pipeline.yaml +++ b/.buildkite/release-pipeline.yaml @@ -33,7 +33,7 @@ steps: agents: queue: arm64_cpu_queue_release commands: - - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda13.0-78e737ad29420ffc4800e677c51e2a852caf8359 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ." + - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda13.0-b8b5f17a7d9ccfc25bbc5cf17b3fcea12964a042 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ." - "mkdir artifacts" - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'" - "bash .buildkite/scripts/upload-nightly-wheels.sh" @@ -260,7 +260,7 @@ steps: --build-arg CUDA_VERSION=13.0.2 \ --build-arg torch_cuda_arch_list="${CUDA_ARCH_AARCH64}" \ --build-arg INSTALL_KV_CONNECTORS=true \ - --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda13.0-78e737ad29420ffc4800e677c51e2a852caf8359 \ + --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda13.0-b8b5f17a7d9ccfc25bbc5cf17b3fcea12964a042 \ --target vllm-openai \ --progress plain \ -f docker/Dockerfile . @@ -361,7 +361,7 @@ steps: --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 \ --build-arg torch_cuda_arch_list="${CUDA_ARCH_AARCH64}" \ --build-arg INSTALL_KV_CONNECTORS=true \ - --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda13.0-78e737ad29420ffc4800e677c51e2a852caf8359 \ + --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda13.0-b8b5f17a7d9ccfc25bbc5cf17b3fcea12964a042 \ --target vllm-openai \ --progress plain \ -f docker/Dockerfile . diff --git a/.buildkite/scripts/hardware_ci/run-gh200-test.sh b/.buildkite/scripts/hardware_ci/run-gh200-test.sh index dca25b177a93..c95d74d3b4e5 100644 --- a/.buildkite/scripts/hardware_ci/run-gh200-test.sh +++ b/.buildkite/scripts/hardware_ci/run-gh200-test.sh @@ -15,7 +15,7 @@ DOCKER_BUILDKIT=1 docker build . \ -t gh200-test \ --build-arg max_jobs=66 \ --build-arg nvcc_threads=2 \ - --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda13.0-78e737ad29420ffc4800e677c51e2a852caf8359 \ + --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda13.0-b8b5f17a7d9ccfc25bbc5cf17b3fcea12964a042 \ --build-arg torch_cuda_arch_list="9.0+PTX" # Setup cleanup From 704f12aa1ed66db6307659f2996fd905a73570d4 Mon Sep 17 00:00:00 2001 From: Taneem Ibrahim Date: Sat, 22 Aug 2026 08:01:40 -0500 Subject: [PATCH 008/235] [Pooling] Report input throughput for batched requests (#53213) Signed-off-by: Taneem Ibrahim Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> --- rust/src/bench/src/backends/mod.rs | 2 ++ rust/src/bench/src/backends/pooling.rs | 15 +++++++++++++-- rust/src/bench/src/compare.rs | 5 +++++ rust/src/bench/src/metrics/calculator.rs | 6 ++++++ rust/src/bench/src/metrics/mod.rs | 2 ++ rust/src/bench/src/multi_run.rs | 20 +++++++++++++++++--- rust/src/bench/src/output/console.rs | 6 ++++++ rust/src/bench/src/output/json.rs | 10 ++++++++++ rust/src/bench/src/sweep.rs | 3 +++ vllm/benchmarks/datasets/datasets.py | 2 +- vllm/benchmarks/lib/endpoint_request_func.py | 16 ++++++++++++++-- vllm/benchmarks/serve.py | 14 ++++++++++++++ 12 files changed, 93 insertions(+), 8 deletions(-) diff --git a/rust/src/bench/src/backends/mod.rs b/rust/src/bench/src/backends/mod.rs index 698a20e49a29..c01018529ff5 100644 --- a/rust/src/bench/src/backends/mod.rs +++ b/rust/src/bench/src/backends/mod.rs @@ -101,6 +101,7 @@ pub struct RequestFuncOutput { pub prompt_len: usize, pub error: String, pub start_time: f64, + pub num_input_sequences: usize, } impl Default for RequestFuncOutput { @@ -116,6 +117,7 @@ impl Default for RequestFuncOutput { prompt_len: 0, error: String::new(), start_time: 0.0, + num_input_sequences: 1, } } } diff --git a/rust/src/bench/src/backends/pooling.rs b/rust/src/bench/src/backends/pooling.rs index 8f9d4deab996..ad43f480e21d 100644 --- a/rust/src/bench/src/backends/pooling.rs +++ b/rust/src/bench/src/backends/pooling.rs @@ -39,9 +39,22 @@ impl PoolingBackend { input: &RequestFuncInput, client: &reqwest::Client, ) -> Result { + let payload = self.build_payload(input); + // Preserve client-side prompt_len as fallback if server doesn't report usage. let mut output = RequestFuncOutput { prompt_len: input.prompt_len, + num_input_sequences: match self.kind { + BackendKind::OpenaiEmbeddings | BackendKind::VllmPooling => { + input.prompt_list.as_ref().map_or(1, |list| list.len()) + } + BackendKind::OpenaiEmbeddingsChat => 1, + BackendKind::VllmRerank => payload + .get("documents") + .and_then(|documents| documents.as_array()) + .map_or(0, |documents| documents.len()), + _ => unreachable!("PoolingBackend with non-pooling kind"), + }, ..Default::default() }; @@ -51,8 +64,6 @@ impl PoolingBackend { &input.request_id, ); - let payload = self.build_payload(input); - let mut request = client.post(&input.api_url); for (k, v) in &headers { request = request.header(k, v); diff --git a/rust/src/bench/src/compare.rs b/rust/src/bench/src/compare.rs index 6d41b72b85cb..0b0dece1c322 100644 --- a/rust/src/bench/src/compare.rs +++ b/rust/src/bench/src/compare.rs @@ -16,6 +16,11 @@ const METRICS: &[MetricDef] = &[ key: "request_throughput", lower_is_better: false, }, + MetricDef { + label: "Input throughput (inputs/s)", + key: "input_sequence_throughput", + lower_is_better: false, + }, MetricDef { label: "Output throughput (tok/s)", key: "output_throughput", diff --git a/rust/src/bench/src/metrics/calculator.rs b/rust/src/bench/src/metrics/calculator.rs index 3bce8dd1a7f4..11b68a3bfeab 100644 --- a/rust/src/bench/src/metrics/calculator.rs +++ b/rust/src/bench/src/metrics/calculator.rs @@ -184,8 +184,10 @@ pub fn calculate_metrics( completed, failed, total_input, + total_input_sequences: completed, total_output, request_throughput: completed as f64 / dur_s, + input_sequence_throughput: completed as f64 / dur_s, request_goodput, input_throughput: total_input as f64 / dur_s, output_throughput: total_output as f64 / dur_s, @@ -291,6 +293,7 @@ pub fn calculate_embedding_metrics( selected_percentiles: &[f64], ) -> BenchmarkMetrics { let mut total_input: usize = 0; + let mut total_input_sequences: usize = 0; let mut completed: usize = 0; let mut e2els: Vec = Vec::new(); @@ -299,6 +302,7 @@ pub fn calculate_embedding_metrics( e2els.push(output.latency); completed += 1; total_input += output.prompt_len; + total_input_sequences += output.num_input_sequences; } } @@ -344,8 +348,10 @@ pub fn calculate_embedding_metrics( completed, failed, total_input, + total_input_sequences, total_output: 0, request_throughput: completed as f64 / dur_s, + input_sequence_throughput: total_input_sequences as f64 / dur_s, request_goodput: 0.0, input_throughput: total_input as f64 / dur_s, output_throughput: 0.0, diff --git a/rust/src/bench/src/metrics/mod.rs b/rust/src/bench/src/metrics/mod.rs index 829291acc84e..2acce568a86e 100644 --- a/rust/src/bench/src/metrics/mod.rs +++ b/rust/src/bench/src/metrics/mod.rs @@ -25,8 +25,10 @@ pub struct BenchmarkMetrics { pub completed: usize, pub failed: usize, pub total_input: usize, + pub total_input_sequences: usize, pub total_output: usize, pub request_throughput: f64, + pub input_sequence_throughput: f64, pub request_goodput: f64, pub input_throughput: f64, pub output_throughput: f64, diff --git a/rust/src/bench/src/multi_run.rs b/rust/src/bench/src/multi_run.rs index 18ff7753b949..3d4518d7ce92 100644 --- a/rust/src/bench/src/multi_run.rs +++ b/rust/src/bench/src/multi_run.rs @@ -7,6 +7,7 @@ use crate::error::Result; /// Key metrics extracted from a single run's JSON result. struct RunMetrics { request_throughput: f64, + input_sequence_throughput: Option, output_throughput: f64, total_token_throughput: f64, mean_ttft_ms: f64, @@ -38,6 +39,7 @@ impl RunMetrics { fn from_json(json: &serde_json::Value) -> Self { Self { request_throughput: get(json, "request_throughput"), + input_sequence_throughput: get_opt(json, "input_sequence_throughput"), output_throughput: get(json, "output_throughput"), total_token_throughput: get(json, "total_token_throughput"), mean_ttft_ms: get(json, "mean_ttft_ms"), @@ -71,6 +73,10 @@ fn get(json: &serde_json::Value, key: &str) -> f64 { json.get(key).and_then(|v| v.as_f64()).unwrap_or(0.0) } +fn get_opt(json: &serde_json::Value, key: &str) -> Option { + json.get(key).and_then(|v| v.as_f64()) +} + fn get_ss_opt(json: &serde_json::Value, key: &str) -> Option { json.get("steady_state") .and_then(|ss| if ss.is_null() { None } else { Some(ss) }) @@ -119,8 +125,15 @@ fn print_multi_run_summary(runs: &[RunMetrics]) { let n = runs.len(); // Collect each metric into a series, compute stats - let stats = vec![ - compute_stats("Request throughput (req/s)", runs, |r| r.request_throughput), + let mut stats = vec![compute_stats("Request throughput (req/s)", runs, |r| { + r.request_throughput + })]; + if runs.iter().all(|r| r.input_sequence_throughput.is_some()) { + stats.push(compute_stats("Input throughput (inputs/s)", runs, |r| { + r.input_sequence_throughput.unwrap_or_default() + })); + } + stats.extend([ compute_stats("Output throughput (tok/s)", runs, |r| r.output_throughput), compute_stats("Total token throughput (tok/s)", runs, |r| { r.total_token_throughput @@ -138,7 +151,7 @@ fn print_multi_run_summary(runs: &[RunMetrics]) { compute_stats("Completed requests", runs, |r| r.completed), compute_stats("Failed requests", runs, |r| r.failed), compute_stats("Duration (s)", runs, |r| r.duration), - ]; + ]); println!("{:=^80}", format!(" Multi-Run Summary ({n} runs) ")); println!( @@ -307,6 +320,7 @@ mod tests { fn mk_run(ss: Option) -> RunMetrics { RunMetrics { request_throughput: 0.0, + input_sequence_throughput: None, output_throughput: 0.0, total_token_throughput: 0.0, mean_ttft_ms: 0.0, diff --git a/rust/src/bench/src/output/console.rs b/rust/src/bench/src/output/console.rs index ae0184e61541..62d7455d7c1c 100644 --- a/rust/src/bench/src/output/console.rs +++ b/rust/src/bench/src/output/console.rs @@ -51,6 +51,12 @@ pub fn print_results( "{:<40} {:<10.2}", "Request throughput (req/s):", metrics.request_throughput ); + if is_pooling { + println!( + "{:<40} {:<10.2}", + "Input throughput (inputs/s):", metrics.input_sequence_throughput + ); + } if metrics.request_goodput > 0.0 { println!( "{:<40} {:<10.2}", diff --git a/rust/src/bench/src/output/json.rs b/rust/src/bench/src/output/json.rs index 6e92cf8a20f1..d2420529e873 100644 --- a/rust/src/bench/src/output/json.rs +++ b/rust/src/bench/src/output/json.rs @@ -98,6 +98,16 @@ pub fn build_result_json( "request_throughput".into(), serde_json::json!(metrics.request_throughput), ); + if is_pooling { + result.insert( + "total_input_sequences".into(), + serde_json::json!(metrics.total_input_sequences), + ); + result.insert( + "input_sequence_throughput".into(), + serde_json::json!(metrics.input_sequence_throughput), + ); + } if config.goodput.is_empty() { result.insert("request_goodput".into(), Value::Null); } else { diff --git a/rust/src/bench/src/sweep.rs b/rust/src/bench/src/sweep.rs index 2d695356dc15..367e7adfc34e 100644 --- a/rust/src/bench/src/sweep.rs +++ b/rust/src/bench/src/sweep.rs @@ -295,6 +295,7 @@ fn print_sweep_summary(param_name: &str, points: &[SweepPoint], summary_percenti fn build_summary_columns(summary_percentiles: &[f64]) -> Vec { let mut columns = vec![ SummaryColumn::new("Req/s", "request_throughput", 10), + SummaryColumn::new("Inputs/s", "input_sequence_throughput", 10), SummaryColumn::new("Tok/s", "output_throughput", 10), SummaryColumn::new("Total tok/s", "total_token_throughput", 12), SummaryColumn::new("SS req/s", SS_REQUEST_THROUGHPUT_KEY, 10), @@ -452,6 +453,7 @@ mod tests { headers, vec![ "Req/s", + "Inputs/s", "Tok/s", "Total tok/s", "SS req/s", @@ -495,6 +497,7 @@ mod tests { headers, vec![ "Req/s", + "Inputs/s", "Tok/s", "Total tok/s", "SS req/s", diff --git a/vllm/benchmarks/datasets/datasets.py b/vllm/benchmarks/datasets/datasets.py index d928123dde87..3ca9fd670806 100644 --- a/vllm/benchmarks/datasets/datasets.py +++ b/vllm/benchmarks/datasets/datasets.py @@ -79,7 +79,7 @@ class SampleRequest: Represents a single inference request for benchmarking. """ - prompt: str | list[str] | list[int] | list[dict] + prompt: str | list[str] | list[int] | list[list[int]] | list[dict] prompt_len: int expected_output_len: int = 0 multi_modal_data: MultiModalDataDict | dict | list[dict] | None = None diff --git a/vllm/benchmarks/lib/endpoint_request_func.py b/vllm/benchmarks/lib/endpoint_request_func.py index 1023700da4b6..0680d666da39 100644 --- a/vllm/benchmarks/lib/endpoint_request_func.py +++ b/vllm/benchmarks/lib/endpoint_request_func.py @@ -66,7 +66,7 @@ def add_chunk(self, chunk_bytes: bytes) -> list[str]: class RequestFuncInput: """The input for the request function.""" - prompt: str | list[str] | list[int] | list[dict[str, Any]] + prompt: str | list[str] | list[int] | list[list[int]] | list[dict[str, Any]] api_url: str prompt_len: int output_len: int @@ -100,6 +100,7 @@ class RequestFuncOutput: error: str = "" start_time: float = 0.0 input_audio_duration: float = 0.0 # in seconds + num_input_sequences: int = 1 class RequestFunc(Protocol): @@ -586,8 +587,9 @@ async def _run_pooling_request( payload: dict[str, Any], headers: dict[str, Any], pbar: tqdm | None = None, + num_input_sequences: int = 1, ) -> RequestFuncOutput: - output = RequestFuncOutput() + output = RequestFuncOutput(num_input_sequences=num_input_sequences) st = time.perf_counter() output.start_time = st try: @@ -617,6 +619,12 @@ async def _run_pooling_request( return output +def _get_num_input_sequences(prompt: Any) -> int: + if prompt and isinstance(prompt, list) and isinstance(prompt[0], (str, list)): + return len(prompt) + return 1 + + async def async_request_openai_embeddings( request_func_input: RequestFuncInput, session: aiohttp.ClientSession, @@ -645,6 +653,7 @@ async def async_request_openai_embeddings( payload=payload, headers=headers, pbar=pbar, + num_input_sequences=_get_num_input_sequences(request_func_input.prompt), ) @@ -681,6 +690,7 @@ async def async_request_vllm_rerank( payload=payload, headers=headers, pbar=pbar, + num_input_sequences=len(request_func_input.prompt) - 1, ) @@ -818,6 +828,7 @@ async def async_request_infinity_embeddings( payload=payload, headers=headers, pbar=pbar, + num_input_sequences=_get_num_input_sequences(request_func_input.prompt), ) @@ -866,6 +877,7 @@ async def async_request_vllm_pooling( payload=payload, headers=headers, pbar=pbar, + num_input_sequences=_get_num_input_sequences(request_func_input.prompt), ) diff --git a/vllm/benchmarks/serve.py b/vllm/benchmarks/serve.py index 9b993689235f..960668b93af5 100644 --- a/vllm/benchmarks/serve.py +++ b/vllm/benchmarks/serve.py @@ -357,7 +357,9 @@ class EmbedBenchmarkMetrics: completed: int failed: int total_input: int + total_input_sequences: int request_throughput: float + input_sequence_throughput: float total_token_throughput: float mean_e2el_ms: float std_e2el_ms: float @@ -520,6 +522,7 @@ def calculate_metrics_for_embeddings( The calculated benchmark metrics. """ total_input = 0 + total_input_sequences = 0 completed = 0 failed = 0 e2els: list[float] = [] @@ -528,6 +531,7 @@ def calculate_metrics_for_embeddings( e2els.append(outputs[i].latency) completed += 1 total_input += outputs[i].prompt_len + total_input_sequences += outputs[i].num_input_sequences else: failed += 1 @@ -541,7 +545,9 @@ def calculate_metrics_for_embeddings( completed=completed, failed=failed, total_input=total_input, + total_input_sequences=total_input_sequences, request_throughput=completed / dur_s, + input_sequence_throughput=total_input_sequences / dur_s, total_token_throughput=total_input / dur_s, mean_e2el_ms=np.mean(e2els or 0) * 1000, std_e2el_ms=np.std(e2els or 0) * 1000, @@ -1181,6 +1187,12 @@ async def probe_loop(): "Request throughput (req/s):", metrics.request_throughput ) ) + if isinstance(metrics, EmbedBenchmarkMetrics): + print( + "{:<40} {:<10.2f}".format( + "Input throughput (inputs/s):", metrics.input_sequence_throughput + ) + ) if goodput_config_dict and isinstance(metrics, BenchmarkMetrics): print( "{:<40} {:<10.2f}".format( @@ -1284,7 +1296,9 @@ async def probe_loop(): "duration": benchmark_duration, "completed": metrics.completed, "total_input_tokens": metrics.total_input, + "total_input_sequences": metrics.total_input_sequences, "request_throughput": metrics.request_throughput, + "input_sequence_throughput": metrics.input_sequence_throughput, "total_token_throughput": metrics.total_token_throughput, "input_lens": [output.prompt_len for output in outputs], "errors": [output.error for output in outputs], From da329cc303a5233e17fa3d553ce0a3d6ceea87a8 Mon Sep 17 00:00:00 2001 From: Jens Date: Sat, 22 Aug 2026 15:34:22 +0200 Subject: [PATCH 009/235] [Bugfix] Fix speculative decoding for short_conv (LFM2) models (#50272) Signed-off-by: jens.luecke Co-authored-by: Tonoken3 Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> --- .../layers/mamba/mamba_utils.py | 3 +- .../model_executor/layers/mamba/short_conv.py | 42 +++++++++++++++---- .../worker/gpu/model_states/mamba_hybrid.py | 7 +++- vllm/v1/worker/gpu_model_runner.py | 2 + 4 files changed, 43 insertions(+), 11 deletions(-) diff --git a/vllm/model_executor/layers/mamba/mamba_utils.py b/vllm/model_executor/layers/mamba/mamba_utils.py index 18a73c6554e6..71872de47c0b 100644 --- a/vllm/model_executor/layers/mamba/mamba_utils.py +++ b/vllm/model_executor/layers/mamba/mamba_utils.py @@ -235,9 +235,10 @@ def short_conv_state_shape( tp_world_size: int, intermediate_size: int, conv_kernel: int, + num_spec: int = 0, ) -> tuple[tuple[int, int]]: conv_dim = divide(intermediate_size, tp_world_size) - conv_state_shape = cls._orient_conv_shape(conv_dim, conv_kernel - 1) + conv_state_shape = cls._orient_conv_shape(conv_dim, conv_kernel - 1 + num_spec) return (conv_state_shape,) @classmethod diff --git a/vllm/model_executor/layers/mamba/short_conv.py b/vllm/model_executor/layers/mamba/short_conv.py index f5548cc7de58..de2efea50fc8 100644 --- a/vllm/model_executor/layers/mamba/short_conv.py +++ b/vllm/model_executor/layers/mamba/short_conv.py @@ -80,7 +80,9 @@ def __init__( prefix=f"{prefix}.out_proj", ) - compilation_config = get_current_vllm_config().compilation_config + vllm_config = get_current_vllm_config() + self.num_spec = vllm_config.num_speculative_tokens + compilation_config = vllm_config.compilation_config if prefix in compilation_config.static_forward_context: raise ValueError(f"Duplicate layer name: {prefix}") compilation_config.static_forward_context[prefix] = self @@ -234,6 +236,8 @@ def forward_cuda( state_indices_tensor_d = attn_metadata.state_indices_tensor_d has_initial_states_p = attn_metadata.has_initial_states_p query_start_loc_p = attn_metadata.query_start_loc_p + num_accepted_tokens = attn_metadata.num_accepted_tokens + query_start_loc_d = attn_metadata.query_start_loc_d BCx, _ = self.in_proj(hidden_states) @@ -296,14 +300,33 @@ def forward_cuda( if has_decode: Bx_d = (B_d * x_d).contiguous() - Bx = causal_conv1d_update( - Bx_d, - conv_state, - conv_weights, - self.conv.bias, - activation=None, - conv_state_indices=state_indices_tensor_d, - ) + if num_accepted_tokens is not None: + # Speculative decode: the verify step feeds >1 query token per + # decode request, so use the spec-aware conv update path + # (mirrors mamba_mixer2). state_indices_tensor_d is + # (num_decodes, 1 + num_spec_tokens) here. + assert state_indices_tensor_d is not None + Bx = causal_conv1d_update( + Bx_d, + conv_state, + conv_weights, + self.conv.bias, + activation=None, + conv_state_indices=state_indices_tensor_d, + num_accepted_tokens=num_accepted_tokens, + query_start_loc=query_start_loc_d, + max_query_len=state_indices_tensor_d.size(-1), + ) + else: + # Non-spec decode + Bx = causal_conv1d_update( + Bx_d, + conv_state, + conv_weights, + self.conv.bias, + activation=None, + conv_state_indices=state_indices_tensor_d, + ) y = C_d * Bx conv_output_list.insert(0, y) @@ -326,6 +349,7 @@ def get_state_shape(self) -> tuple[tuple[int, ...]]: tp_world_size=get_tensor_model_parallel_world_size(), intermediate_size=self.conv_dim, conv_kernel=self.L_cache, + num_spec=self.num_spec, ) @property diff --git a/vllm/v1/worker/gpu/model_states/mamba_hybrid.py b/vllm/v1/worker/gpu/model_states/mamba_hybrid.py index c61a56f9aabf..420bf5d6dd2e 100644 --- a/vllm/v1/worker/gpu/model_states/mamba_hybrid.py +++ b/vllm/v1/worker/gpu/model_states/mamba_hybrid.py @@ -12,6 +12,7 @@ from vllm.triton_utils import tl, triton from vllm.v1.attention.backends.gdn_attn import GDNAttentionMetadataBuilder from vllm.v1.attention.backends.mamba2_attn import Mamba2AttentionMetadataBuilder +from vllm.v1.attention.backends.short_conv_attn import ShortConvAttentionMetadataBuilder from vllm.v1.core.sched.output import NewRequestData from vllm.v1.kv_cache_interface import KVCacheConfig, MambaSpec from vllm.v1.utils import CpuGpuBuffer @@ -48,7 +49,11 @@ def get_extra_attn_kwargs( ) -> dict[str, Any]: if not isinstance( attn_metadata_builder, - (Mamba2AttentionMetadataBuilder, GDNAttentionMetadataBuilder), + ( + Mamba2AttentionMetadataBuilder, + GDNAttentionMetadataBuilder, + ShortConvAttentionMetadataBuilder, + ), ): return {} return { diff --git a/vllm/v1/worker/gpu_model_runner.py b/vllm/v1/worker/gpu_model_runner.py index 4ef7f5860a5f..8a0e27d3c839 100644 --- a/vllm/v1/worker/gpu_model_runner.py +++ b/vllm/v1/worker/gpu_model_runner.py @@ -148,6 +148,7 @@ BailingLinearAttentionMetadataBuilder, ) from vllm.v1.attention.backends.mamba2_attn import Mamba2AttentionMetadataBuilder +from vllm.v1.attention.backends.short_conv_attn import ShortConvAttentionMetadataBuilder from vllm.v1.attention.backends.utils import ( NULL_BLOCK_ID, create_fast_prefill_custom_backend, @@ -2572,6 +2573,7 @@ def _build_attn_group_metadata( Mamba2AttentionMetadataBuilder, GDNAttentionMetadataBuilder, BailingLinearAttentionMetadataBuilder, + ShortConvAttentionMetadataBuilder, ), ): assert ubid is None, ( From 610bfc58e93c652fc8a075bc814ba1617f5fcf2b Mon Sep 17 00:00:00 2001 From: Rohan Potdar Date: Sat, 22 Aug 2026 12:16:25 -0500 Subject: [PATCH 010/235] [ROCm] Ship rocprofiler-sdk 1.3.2 in Dockerfile.rocm_base to fix torch.profiler traces (#53182) Signed-off-by: Rohan138 --- docker/Dockerfile.rocm_base | 41 +++++++++++++++++++++++++++++++++++++ 1 file changed, 41 insertions(+) diff --git a/docker/Dockerfile.rocm_base b/docker/Dockerfile.rocm_base index 639dadbb74bf..ebf1a4313487 100644 --- a/docker/Dockerfile.rocm_base +++ b/docker/Dockerfile.rocm_base @@ -13,6 +13,8 @@ ARG AITER_BRANCH="v0.1.19" ARG AITER_REPO="https://github.com/ROCm/aiter.git" ARG MORI_BRANCH="v1.1.0" ARG MORI_REPO="https://github.com/ROCm/mori.git" +ARG ROCPROFILER_SDK_REPO="https://github.com/ROCm/rocm-systems.git" +ARG ROCPROFILER_SDK_COMMIT="2b22ab0195cc1461cd9abf3b969e9dd7c10af350" # Sccache configuration (only used in release pipeline) ARG USE_SCCACHE @@ -313,10 +315,45 @@ RUN --mount=type=bind,from=build_aiter,src=/app/install/,target=/install \ RUN --mount=type=bind,from=build_mori,src=/app/install/,target=/install \ cp /install/*.whl /app/debs +# rocprofiler-sdk 1.3.2 fixes torch 2.12 profiler on ROCm 7.2.3 +FROM base AS build_rocprofiler_sdk +ARG ROCPROFILER_SDK_REPO +ARG ROCPROFILER_SDK_COMMIT +ARG USE_SCCACHE +RUN apt-get update \ + && apt-get install -y --no-install-recommends libdw-dev libelf-dev libsqlite3-dev libzstd-dev \ + && rm -rf /var/lib/apt/lists/* +RUN git init -q /src && cd /src \ + && git remote add origin ${ROCPROFILER_SDK_REPO} \ + && git config core.sparseCheckout true \ + && git config remote.origin.promisor true \ + && git config remote.origin.partialclonefilter blob:none \ + && git sparse-checkout init --cone \ + && git sparse-checkout set projects/rocprofiler-sdk shared cmake \ + && git fetch --filter=blob:none --depth 1 origin ${ROCPROFILER_SDK_COMMIT} \ + && git checkout -q FETCH_HEAD +RUN cd /src/projects/rocprofiler-sdk \ + && if [ "$USE_SCCACHE" = "1" ]; then \ + export CMAKE_C_COMPILER_LAUNCHER=sccache \ + && export CMAKE_CXX_COMPILER_LAUNCHER=sccache; \ + fi \ + && cmake -B build -G Ninja \ + -DCMAKE_BUILD_TYPE=Release \ + -DCMAKE_PREFIX_PATH=/opt/rocm \ + -DCMAKE_INSTALL_PREFIX=/opt/rocm \ + -DROCPROFILER_BUILD_TESTS=OFF \ + -DROCPROFILER_BUILD_SAMPLES=OFF \ + -DROCPROFILER_BUILD_DOCS=OFF \ + && cmake --build build --parallel "$(nproc)" \ + && cmake --install build --prefix /staging --strip + FROM base AS final RUN --mount=type=bind,from=debs,src=/app/debs,target=/install \ pip install /install/*.whl +COPY --from=build_rocprofiler_sdk /staging/ /opt/rocm/ +RUN ldconfig + ARG BASE_IMAGE ARG TRITON_BRANCH ARG TRITON_REPO @@ -332,7 +369,11 @@ ARG AITER_BRANCH ARG AITER_REPO ARG MORI_BRANCH ARG MORI_REPO +ARG ROCPROFILER_SDK_REPO +ARG ROCPROFILER_SDK_COMMIT RUN echo "BASE_IMAGE: ${BASE_IMAGE}" > /app/versions.txt \ + && echo "ROCPROFILER_SDK_REPO: ${ROCPROFILER_SDK_REPO}" >> /app/versions.txt \ + && echo "ROCPROFILER_SDK_COMMIT: ${ROCPROFILER_SDK_COMMIT}" >> /app/versions.txt \ && echo "TRITON_BRANCH: ${TRITON_BRANCH}" >> /app/versions.txt \ && echo "TRITON_REPO: ${TRITON_REPO}" >> /app/versions.txt \ && echo "PYTORCH_BRANCH: ${PYTORCH_BRANCH}" >> /app/versions.txt \ From 020829f66baa231a3ed549c7f6b4a3171ef1c759 Mon Sep 17 00:00:00 2001 From: Cyrus Leung Date: Sun, 23 Aug 2026 01:25:55 +0800 Subject: [PATCH 011/235] [MM] Remove renderer_applies_updates flag (#53385) Signed-off-by: DarkLight1337 --- .../model_executor/models/nano_nemotron_vl.py | 1 - vllm/model_executor/models/pixtral.py | 25 ++++++++++- .../models/qwen2_5_omni_thinker.py | 44 +++++++------------ .../models/qwen3_asr_realtime.py | 1 - .../models/qwen3_omni_moe_thinker.py | 43 ++++++------------ vllm/model_executor/models/voxtral.py | 24 +++++++++- .../model_executor/models/voxtral_realtime.py | 1 - vllm/multimodal/processing/processor.py | 27 +++--------- 8 files changed, 80 insertions(+), 86 deletions(-) diff --git a/vllm/model_executor/models/nano_nemotron_vl.py b/vllm/model_executor/models/nano_nemotron_vl.py index 00538fd7b328..f08cf5e80682 100644 --- a/vllm/model_executor/models/nano_nemotron_vl.py +++ b/vllm/model_executor/models/nano_nemotron_vl.py @@ -734,7 +734,6 @@ def apply( prompt_ids=prompt_ids, mm_kwargs=mm_info.kwargs, mm_prompt_updates=mm_info.prompt_updates, - is_update_applied=self.renderer_applies_updates, ) mm_placeholder_ranges = { diff --git a/vllm/model_executor/models/pixtral.py b/vllm/model_executor/models/pixtral.py index 6a210a811b0f..869b7001c8b7 100644 --- a/vllm/model_executor/models/pixtral.py +++ b/vllm/model_executor/models/pixtral.py @@ -40,6 +40,7 @@ from vllm.multimodal import MULTIMODAL_REGISTRY, MultiModalKwargsItems from vllm.multimodal.inputs import ( MultiModalFieldConfig, + MultiModalKwargsOptionalItems, NestedTensors, ) from vllm.multimodal.parse import ( @@ -51,6 +52,8 @@ from vllm.multimodal.processing.processor import ( BaseMultiModalProcessor, BaseProcessingInfo, + MultiModalPromptUpdates, + PlaceholderFeaturesInfo, ProcessorInputs, PromptReplacement, PromptUpdate, @@ -211,8 +214,26 @@ def get_dummy_processor_inputs( class PixtralMultiModalProcessor(BaseMultiModalProcessor[PixtralProcessingInfo]): - # The tokens are already inserted by the chat template - renderer_applies_updates = True + # The tokens are already inserted by the chat template, + # so we just double check that they exist + def _maybe_apply_prompt_updates( + self, + mm_items: MultiModalDataItems, + prompt_ids: list[int], + mm_kwargs: MultiModalKwargsOptionalItems, + mm_prompt_updates: MultiModalPromptUpdates, + ) -> tuple[list[int], Mapping[str, list[PlaceholderFeaturesInfo]]]: + mm_item_counts = mm_items.get_all_counts() + self._validate_mm_kwargs(mm_kwargs, mm_item_counts) + self._validate_mm_updates(mm_prompt_updates, mm_item_counts) + + mm_placeholders = self._find_mm_placeholders( + prompt_ids, + mm_prompt_updates, + ) + self._validate_mm_placeholders(mm_placeholders, mm_item_counts) + + return prompt_ids, mm_placeholders def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/qwen2_5_omni_thinker.py b/vllm/model_executor/models/qwen2_5_omni_thinker.py index 1f335cfc2420..58df408a5e22 100644 --- a/vllm/model_executor/models/qwen2_5_omni_thinker.py +++ b/vllm/model_executor/models/qwen2_5_omni_thinker.py @@ -575,7 +575,6 @@ def _maybe_apply_prompt_updates( prompt_ids: list[int], mm_kwargs: MultiModalKwargsItems, mm_prompt_updates: MultiModalPromptUpdates, - is_update_applied: bool, ) -> tuple[list[int], Mapping[str, list[PlaceholderFeaturesInfo]]]: """ Qwen2.5-Omni reimplements this function to handle `use_audio_in_video`. @@ -605,40 +604,27 @@ def _maybe_apply_prompt_updates( if video_audio_item_num != video_updates_num + audio_updates_num: use_audio_in_video = True - if is_update_applied: - mm_placeholders = self._find_mm_placeholders( + if use_audio_in_video and "audio" in mm_prompt_updates: + # Filter out audio updates - they are embedded in video + filtered_updates = { + k: v for k, v in mm_prompt_updates.items() if k != "audio" + } + prompt_ids, mm_placeholders = self._apply_prompt_updates( prompt_ids, - mm_prompt_updates, + filtered_updates, ) - self._validate_mm_placeholders( - mm_placeholders, - mm_item_counts, + # Derive audio placeholders from video placeholders + mm_placeholders = self._derive_audio_from_video_placeholders( + mm_placeholders, mm_prompt_updates ) else: - if use_audio_in_video and "audio" in mm_prompt_updates: - # Filter out audio updates - they are embedded in video - filtered_updates = { - k: v for k, v in mm_prompt_updates.items() if k != "audio" - } - prompt_ids, mm_placeholders = self._apply_prompt_updates( - prompt_ids, - filtered_updates, - ) - # Derive audio placeholders from video placeholders - mm_placeholders = self._derive_audio_from_video_placeholders( - mm_placeholders, mm_prompt_updates - ) - else: - prompt_ids, mm_placeholders = self._apply_prompt_updates( - prompt_ids, - mm_prompt_updates, - ) - - self._validate_mm_placeholders( - mm_placeholders, - mm_item_counts, + prompt_ids, mm_placeholders = self._apply_prompt_updates( + prompt_ids, + mm_prompt_updates, ) + self._validate_mm_placeholders(mm_placeholders, mm_item_counts) + return prompt_ids, mm_placeholders @classmethod diff --git a/vllm/model_executor/models/qwen3_asr_realtime.py b/vllm/model_executor/models/qwen3_asr_realtime.py index 6405c101b4a2..917bfce13446 100644 --- a/vllm/model_executor/models/qwen3_asr_realtime.py +++ b/vllm/model_executor/models/qwen3_asr_realtime.py @@ -116,7 +116,6 @@ def _maybe_apply_prompt_updates( prompt_ids: list[int], mm_kwargs: MultiModalKwargsOptionalItems, mm_prompt_updates: MultiModalPromptUpdates, - is_update_applied: bool, ) -> tuple[list[int], Mapping[str, list[PlaceholderFeaturesInfo]]]: audios = mm_kwargs.get("audio", []) assert len(audios) == 1, ( diff --git a/vllm/model_executor/models/qwen3_omni_moe_thinker.py b/vllm/model_executor/models/qwen3_omni_moe_thinker.py index a3691742ffc5..0e2aab634cbc 100755 --- a/vllm/model_executor/models/qwen3_omni_moe_thinker.py +++ b/vllm/model_executor/models/qwen3_omni_moe_thinker.py @@ -1291,7 +1291,6 @@ def _maybe_apply_prompt_updates( prompt_ids: list[int], mm_kwargs: MultiModalKwargsItems, mm_prompt_updates: MultiModalPromptUpdates, - is_update_applied: bool, ) -> tuple[list[int], Mapping[str, list[PlaceholderFeaturesInfo]]]: """ Qwen3-Omni reimplements this function to handle `use_audio_in_video`. @@ -1318,40 +1317,26 @@ def _maybe_apply_prompt_updates( if video_audio_item_num != video_updates_num + audio_updates_num: use_audio_in_video = True - # normal case with `use_audio_in_video=False` - if is_update_applied: - mm_placeholders = self._find_mm_placeholders( + if use_audio_in_video and "audio" in mm_prompt_updates: + filtered_updates = { + k: v for k, v in mm_prompt_updates.items() if k != "audio" + } + prompt_ids, mm_placeholders = self._apply_prompt_updates( prompt_ids, - mm_prompt_updates, + filtered_updates, ) - self._validate_mm_placeholders( - mm_placeholders, - mm_item_counts, + # Derive audio placeholders from video placeholders + mm_placeholders = self._derive_audio_from_video_placeholders( + mm_placeholders, mm_prompt_updates ) else: - if use_audio_in_video and "audio" in mm_prompt_updates: - filtered_updates = { - k: v for k, v in mm_prompt_updates.items() if k != "audio" - } - prompt_ids, mm_placeholders = self._apply_prompt_updates( - prompt_ids, - filtered_updates, - ) - # Derive audio placeholders from video placeholders - mm_placeholders = self._derive_audio_from_video_placeholders( - mm_placeholders, mm_prompt_updates - ) - else: - prompt_ids, mm_placeholders = self._apply_prompt_updates( - prompt_ids, - mm_prompt_updates, - ) - - self._validate_mm_placeholders( - mm_placeholders, - mm_item_counts, + prompt_ids, mm_placeholders = self._apply_prompt_updates( + prompt_ids, + mm_prompt_updates, ) + self._validate_mm_placeholders(mm_placeholders, mm_item_counts) + return prompt_ids, mm_placeholders def get_updates_use_audio_in_video( diff --git a/vllm/model_executor/models/voxtral.py b/vllm/model_executor/models/voxtral.py index 20d166ca2f54..e078867af36c 100644 --- a/vllm/model_executor/models/voxtral.py +++ b/vllm/model_executor/models/voxtral.py @@ -36,6 +36,7 @@ from vllm.multimodal.inputs import ( MultiModalFieldConfig, MultiModalKwargsItems, + MultiModalKwargsOptionalItems, NestedTensors, ) from vllm.multimodal.parse import ( @@ -47,6 +48,7 @@ from vllm.multimodal.processing.processor import ( BaseMultiModalProcessor, BaseProcessingInfo, + MultiModalPromptUpdates, PlaceholderFeaturesInfo, ProcessorInputs, PromptReplacement, @@ -202,8 +204,26 @@ def get_dummy_processor_inputs( class VoxtralMultiModalProcessor(BaseMultiModalProcessor[VoxtralProcessingInfo]): - # The tokens are already inserted by the chat template - renderer_applies_updates = True + # The tokens are already inserted by the chat template, + # so we just double check that they exist + def _maybe_apply_prompt_updates( + self, + mm_items: MultiModalDataItems, + prompt_ids: list[int], + mm_kwargs: MultiModalKwargsOptionalItems, + mm_prompt_updates: MultiModalPromptUpdates, + ) -> tuple[list[int], Mapping[str, list[PlaceholderFeaturesInfo]]]: + mm_item_counts = mm_items.get_all_counts() + self._validate_mm_kwargs(mm_kwargs, mm_item_counts) + self._validate_mm_updates(mm_prompt_updates, mm_item_counts) + + mm_placeholders = self._find_mm_placeholders( + prompt_ids, + mm_prompt_updates, + ) + self._validate_mm_placeholders(mm_placeholders, mm_item_counts) + + return prompt_ids, mm_placeholders def _get_mm_fields_config( self, diff --git a/vllm/model_executor/models/voxtral_realtime.py b/vllm/model_executor/models/voxtral_realtime.py index b54a4dccba4c..738185a7aecf 100644 --- a/vllm/model_executor/models/voxtral_realtime.py +++ b/vllm/model_executor/models/voxtral_realtime.py @@ -63,7 +63,6 @@ def _maybe_apply_prompt_updates( prompt_ids: list[int], mm_kwargs: MultiModalKwargsOptionalItems, mm_prompt_updates: MultiModalPromptUpdates, - is_update_applied: bool, ) -> tuple[list[int], Mapping[str, list[PlaceholderFeaturesInfo]]]: # there are no placeholder audio tokens for streaming # so we need to build the place placeholder positions manually diff --git a/vllm/multimodal/processing/processor.py b/vllm/multimodal/processing/processor.py index 1eda431a8ae4..dd8a088afa25 100644 --- a/vllm/multimodal/processing/processor.py +++ b/vllm/multimodal/processing/processor.py @@ -8,7 +8,6 @@ from functools import lru_cache from typing import ( TYPE_CHECKING, - ClassVar, Generic, NamedTuple, Protocol, @@ -1177,12 +1176,6 @@ class BaseMultiModalProcessor(ABC, Generic[_I]): Not to be confused with `transformers.ProcessorMixin`. """ - renderer_applies_updates: ClassVar[bool] = False - """ - Only set to True if the tokenizer or chat template in the Renderer - (which are run before MM processing) inserts placeholder tokens. - """ - def __init__( self, info: _I, @@ -1780,24 +1773,17 @@ def _maybe_apply_prompt_updates( prompt_ids: list[int], mm_kwargs: MultiModalKwargsOptionalItems, mm_prompt_updates: MultiModalPromptUpdates, - is_update_applied: bool, ) -> tuple[list[int], Mapping[str, list[PlaceholderFeaturesInfo]]]: mm_item_counts = mm_items.get_all_counts() self._validate_mm_kwargs(mm_kwargs, mm_item_counts) self._validate_mm_updates(mm_prompt_updates, mm_item_counts) - if is_update_applied: - mm_placeholders = self._find_mm_placeholders( - prompt_ids, - mm_prompt_updates, - ) - self._validate_mm_placeholders(mm_placeholders, mm_item_counts) - else: - prompt_ids, mm_placeholders = self._apply_prompt_updates( - prompt_ids, - mm_prompt_updates, - ) - self._validate_mm_placeholders(mm_placeholders, mm_item_counts) + prompt_ids, mm_placeholders = self._apply_prompt_updates( + prompt_ids, + mm_prompt_updates, + ) + + self._validate_mm_placeholders(mm_placeholders, mm_item_counts) return prompt_ids, mm_placeholders @@ -1829,7 +1815,6 @@ def apply( prompt_ids=prompt_ids, mm_kwargs=mm_info.kwargs, mm_prompt_updates=mm_info.prompt_updates, - is_update_applied=self.renderer_applies_updates, ) mm_placeholder_ranges = { From 9eb9d9d3953959695108600c8ed33d36bc6a1e5f Mon Sep 17 00:00:00 2001 From: Wentao Ye <44945378+yewentao256@users.noreply.github.com> Date: Sat, 22 Aug 2026 13:32:40 -0400 Subject: [PATCH 012/235] [Perf] Support internal prefill checkpoints for Mamba prefix caching, 9%~25% TTFT improvement (#52789) Signed-off-by: yewentao256 Signed-off-by: Ziming Huang Signed-off-by: Wentao Ye <44945378+yewentao256@users.noreply.github.com> Co-authored-by: Ziming Huang Co-authored-by: OpenAI Codex --- cmake/external_projects/flashkda.cmake | 2 +- csrc/flashkda_registration.cpp | 3 +- tests/models/kimi_k3/test_kda.py | 81 ++++++++ tests/models/kimi_k3/test_kda_metadata.py | 50 +++++ .../test_partial_prefix_cache_hits.py | 85 +++++++++ tests/v1/core/test_mamba_align_chunk_split.py | 55 +++++- tests/v1/core/test_prefix_caching.py | 1 + vllm/models/kimi_k3/nvidia/kda.py | 180 ++++++++++++++++-- vllm/models/kimi_k3/nvidia/kda_metadata.py | 67 +++++++ vllm/v1/core/sched/scheduler.py | 19 +- vllm/v1/core/single_type_kv_cache_manager.py | 61 ++++-- vllm/v1/kv_cache_interface.py | 6 +- 12 files changed, 562 insertions(+), 48 deletions(-) diff --git a/cmake/external_projects/flashkda.cmake b/cmake/external_projects/flashkda.cmake index 988a27d7c755..b70042f2f4e5 100644 --- a/cmake/external_projects/flashkda.cmake +++ b/cmake/external_projects/flashkda.cmake @@ -13,7 +13,7 @@ else() FetchContent_Declare( flashkda GIT_REPOSITORY https://github.com/vllm-project/FlashKDA.git - GIT_TAG 053de1b716ef3255873e02d2d28f4adf09951978 + GIT_TAG ee0be888cd0e972f9409bf53756f8c38c6652173 GIT_PROGRESS TRUE GIT_SUBMODULES cutlass ) diff --git a/csrc/flashkda_registration.cpp b/csrc/flashkda_registration.cpp index f3da713a09f1..6a9f56b8d2dd 100644 --- a/csrc/flashkda_registration.cpp +++ b/csrc/flashkda_registration.cpp @@ -10,7 +10,8 @@ STABLE_TORCH_LIBRARY(_flashkda_C, m) { "Tensor(a!) out, Tensor(c!) workspace, Tensor A_log, Tensor dt_bias, " "float lower_bound, " "Tensor? initial_state=None, Tensor(b!)? final_state=None, " - "Tensor? cu_seqlens=None) -> ()"); + "Tensor? cu_seqlens=None, Tensor(d!)? checkpoint_state=None, " + "Tensor? checkpoint_offsets=None) -> ()"); } STABLE_TORCH_LIBRARY_IMPL(_flashkda_C, CompositeExplicitAutograd, m) { diff --git a/tests/models/kimi_k3/test_kda.py b/tests/models/kimi_k3/test_kda.py index 7416cca02a32..c8454e9e6f90 100644 --- a/tests/models/kimi_k3/test_kda.py +++ b/tests/models/kimi_k3/test_kda.py @@ -21,6 +21,8 @@ fused_recurrent_kda_packed_decode as fused_recurrent_kda_packed_decode_amd, ) from vllm.models.kimi_k3.nvidia.kda import ( + _flashkda_prefill, + _store_cache_checkpoints_kernel, is_flashkda_supported, is_fused_kda_decode_supported, ) @@ -40,6 +42,7 @@ ) from vllm.platforms import current_platform from vllm.third_party.flash_linear_attention.ops.l2norm import l2norm_fwd +from vllm.v1.attention.backends.utils import NULL_BLOCK_ID DEVICE = current_platform.device_type @@ -1088,6 +1091,17 @@ def test_flashkda_correctness(): expected_states.append(final_state) expected_out = torch.cat(expected_outputs, dim=1) expected_state = torch.cat(expected_states).transpose(-1, -2).contiguous() + _, expected_checkpoint = naive_recurrent_kda( + q_norm[:, :16], + k_norm[:, :16], + v[:, :16], + gate[:, :16], + beta[:, :16], + initial_state=initial_state[0:1].transpose(-1, -2), + output_final_state=True, + ) + assert expected_checkpoint is not None + expected_checkpoint = expected_checkpoint.transpose(-1, -2).contiguous() actual_out = torch.empty_like(v) actual_state = torch.empty_like(initial_state) @@ -1115,3 +1129,70 @@ def test_flashkda_correctness(): assert_close("o", expected_out, actual_out, 0.01) assert_close("ht", expected_state, actual_state, 0.01) + + checkpoint_out = torch.empty_like(v) + checkpoint_final_state = torch.empty_like(initial_state) + checkpoint_state = torch.empty_like(initial_state) + checkpoint_offsets = torch.tensor([16, 31], dtype=torch.int32, device=DEVICE) + _flashkda_prefill( + q=q, + k=k, + v=v, + g=raw_g, + beta=beta_logits, + A_log=A_log, + dt_bias=dt_bias, + lower_bound=lower_bound, + initial_state=initial_state, + cu_seqlens=cu_seqlens, + out=checkpoint_out, + final_state=checkpoint_final_state, + workspace=workspace, + checkpoint_state=checkpoint_state, + checkpoint_offsets=checkpoint_offsets, + ) + + assert_close("checkpoint_o", expected_out, checkpoint_out, 0.01) + assert_close("checkpoint_ht", expected_state, checkpoint_final_state, 0.01) + assert_close("checkpoint", expected_checkpoint, checkpoint_state[:1], 0.01) + + conv_state = torch.zeros(2, H * D, 3, dtype=q.dtype, device=DEVICE) + recurrent_storage = torch.zeros(2, H * D * D + 8, device=DEVICE) + recurrent_state = recurrent_storage[:, : H * D * D].view(2, H, D, D) + conv_input = q[0].flatten(1) + checkpoint_state_indices = torch.tensor( + [1, NULL_BLOCK_ID], dtype=torch.int32, device=DEVICE + ) + state_len = conv_state.shape[-1] + width = H * D + recurrent_row_size = checkpoint_state[0].numel() + block_size = 256 + _store_cache_checkpoints_kernel[ + ( + checkpoint_state_indices.numel(), + (max(width * state_len, recurrent_row_size) + block_size - 1) // block_size, + ) + ]( + conv_input, + conv_state, + checkpoint_state, + recurrent_state, + cu_seqlens, + checkpoint_offsets, + checkpoint_state_indices, + conv_input.stride(0), + conv_input.stride(1), + conv_state.stride(0), + conv_state.stride(1), + conv_state.stride(2), + checkpoint_state.stride(0), + recurrent_state.stride(0), + checkpoint_offsets.stride(0), + state_len, + width, + recurrent_row_size, + NULL_BLOCK_ID, + block_size, + ) + torch.testing.assert_close(conv_state[1], q[0, 13:16].flatten(1).transpose(0, 1)) + torch.testing.assert_close(recurrent_state[1], checkpoint_state[0]) diff --git a/tests/models/kimi_k3/test_kda_metadata.py b/tests/models/kimi_k3/test_kda_metadata.py index 069c9a4d7a8e..ec0046774682 100644 --- a/tests/models/kimi_k3/test_kda_metadata.py +++ b/tests/models/kimi_k3/test_kda_metadata.py @@ -87,6 +87,7 @@ def _make_builder( device: torch.device = DEVICE, mamba_cache_mode: str = "none", use_recoverssm: bool = False, + num_prefill_checkpoint_blocks: int = 0, ) -> AttentionMetadataBuilder: vllm_config = create_vllm_config( model_name="Qwen/Qwen3.5-0.8B", @@ -110,6 +111,7 @@ def _make_builder( dtypes=(torch.float16,), mamba_cache_mode=mamba_cache_mode, num_speculative_blocks=(0 if use_recoverssm else num_speculative_tokens), + num_prefill_checkpoint_blocks=num_prefill_checkpoint_blocks, ), layer_names=["layer.0"], vllm_config=vllm_config, @@ -121,6 +123,54 @@ def _make_builder( return builder +@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA") +def test_internal_checkpoint_metadata_targets_last_aligned_boundary(): + device = torch.device("cuda") + batch = BatchSpec(seq_lens=[50, 32], query_lens=[50, 16]) + common_attn_metadata = create_common_attn_metadata( + batch, BLOCK_SIZE, device, arange_block_indices=True + ).replace(is_prefilling=torch.tensor([True, True])) + actual = _make_builder( + KimiK3KDAMetadataBuilder, + num_speculative_tokens=0, + full_cuda_graph=False, + mamba_cache_mode="align", + num_prefill_checkpoint_blocks=1, + device=device, + ).build(0, common_attn_metadata) + + assert actual.checkpoint is not None + torch.testing.assert_close( + actual.checkpoint.state_indices, + torch.tensor([2, NULL_BLOCK_ID], dtype=torch.int32, device=device), + ) + torch.testing.assert_close( + actual.checkpoint.checkpoint_offsets, + torch.tensor([48, 0], dtype=torch.int32, device=device), + ) + + +@pytest.mark.skipif(not torch.cuda.is_available(), reason="requires CUDA") +def test_internal_checkpoint_metadata_skips_unaligned_offset(): + device = torch.device("cuda") + # The checkpoint block boundary is 48, but this query starts at token 1, + # making the real checkpoint offset 47, which is not FlashKDA-aligned. + batch = BatchSpec(seq_lens=[50], query_lens=[49]) + common_attn_metadata = create_common_attn_metadata( + batch, BLOCK_SIZE, device, arange_block_indices=True + ).replace(is_prefilling=torch.tensor([True])) + actual = _make_builder( + KimiK3KDAMetadataBuilder, + num_speculative_tokens=0, + full_cuda_graph=False, + mamba_cache_mode="align", + num_prefill_checkpoint_blocks=1, + device=device, + ).build(0, common_attn_metadata) + + assert actual.checkpoint is None + + @pytest.mark.parametrize( ( "batch", diff --git a/tests/v1/core/prefix_cache/test_partial_prefix_cache_hits.py b/tests/v1/core/prefix_cache/test_partial_prefix_cache_hits.py index da5bd5b5a1f3..eff889b98740 100644 --- a/tests/v1/core/prefix_cache/test_partial_prefix_cache_hits.py +++ b/tests/v1/core/prefix_cache/test_partial_prefix_cache_hits.py @@ -77,6 +77,7 @@ def make_full_mamba_manager( mamba_block_size: int = 4, num_blocks: int = 32, use_eagle: bool = False, + num_prefill_checkpoint_blocks: int = 0, ): kv_cache_config = KVCacheConfig( num_blocks=num_blocks, @@ -98,6 +99,7 @@ def make_full_mamba_manager( shapes=(1, 1), dtypes=(torch.float32,), mamba_cache_mode="align", + num_prefill_checkpoint_blocks=num_prefill_checkpoint_blocks, ), ), ], @@ -135,6 +137,7 @@ def test_mamba_align_split_partial_tail_schedule(dcp_world_size: int): dcp_world_size=dcp_world_size, scheduler_block_size=scheduler_block_size, mamba_partial_cache_hit=True, + mamba_has_prefill_checkpoint_blocks=False, ) split = Scheduler._mamba_block_aligned_split @@ -180,6 +183,7 @@ def test_mamba_align_split_when_block_exceeds_scheduling_budget(): use_eagle=False, hash_block_size=32, mamba_partial_cache_hit=False, + mamba_has_prefill_checkpoint_blocks=False, ) req = make_request("0", [0] * prompt_length, 32, sha256) split = Scheduler._mamba_block_aligned_split @@ -218,6 +222,7 @@ def test_mamba_align_split_when_block_exceeds_long_prefill_threshold(): use_eagle=False, hash_block_size=32, mamba_partial_cache_hit=False, + mamba_has_prefill_checkpoint_blocks=False, ) req = make_request("0", [0] * prompt_length, 32, sha256) split = Scheduler._mamba_block_aligned_split @@ -462,6 +467,85 @@ def test_hybrid_mamba_partial_tail_owner_uses_cow_on_continue(): assert moved[0].block_hash_num_tokens == 6 +def test_partial_hit_then_internal_checkpoint_uses_distinct_mamba_blocks(): + hash_block_size = 2 + mamba_block_size = 4 + manager = make_full_mamba_manager( + dcp_world_size=1, + hash_block_size=hash_block_size, + full_block_size=hash_block_size, + mamba_block_size=mamba_block_size, + num_prefill_checkpoint_blocks=1, + ) + + owner = make_request("owner", [0, 0, 1, 1, 2, 2], hash_block_size, sha256) + computed_blocks, num_computed, _ = manager.get_computed_blocks(owner) + assert manager.allocate_slots(owner, 6, num_computed, computed_blocks) is not None + manager.free(owner) + manager.new_step_starts() + + partial_hash = owner.block_hashes[2] + partial_block = manager.block_pool.get_cached_block(partial_hash, [1]) + assert partial_block is not None + partial_block_id = partial_block[0].block_id + + replay = make_request( + "replay", + [0, 0, 1, 1, 2, 2, 3, 3, 4, 4, 5, 5, 6, 6], + hash_block_size, + sha256, + ) + computed_blocks, num_computed, _ = manager.get_computed_blocks(replay) + assert num_computed == 6 + + realigned_blocks = manager.allocate_slots(replay, 2, num_computed, computed_blocks) + assert realigned_blocks is not None + mamba_cow_block_id = realigned_blocks.get_block_ids()[1][0] + copies, retained = manager.take_kv_cache_block_copies() + assert KVCacheBlockCopy(partial_block_id, mamba_cow_block_id) in copies + manager.block_pool.free_blocks(retained) + + replay.num_computed_tokens = 8 + manager.new_step_starts() + final_blocks = manager.allocate_slots(replay, 6) + assert final_blocks is not None + + checkpoint_block_id, running_block_id = final_blocks.get_block_ids()[1] + assert len({mamba_cow_block_id, checkpoint_block_id, running_block_id}) == 3 + mamba_blocks = manager.get_blocks(replay.request_id).blocks[1] + assert mamba_blocks[2].block_id == checkpoint_block_id + assert mamba_blocks[3].block_id == running_block_id + + +def test_internal_checkpoint_requires_block_aligned_start(): + hash_block_size = 2 + mamba_block_size = 16 + manager = make_full_mamba_manager( + dcp_world_size=1, + hash_block_size=hash_block_size, + full_block_size=hash_block_size, + mamba_block_size=mamba_block_size, + num_prefill_checkpoint_blocks=1, + ) + request = make_request("producer", list(range(50)), hash_block_size, sha256) + + # Compute token 0 first, so the next query starts at token 1, which is not + # aligned to the Mamba block size. + assert manager.allocate_slots(request, 1) is not None + request.num_computed_tokens = 1 + manager.new_step_starts() + + new_blocks = manager.allocate_slots(request, 49) + + assert new_blocks is not None + mamba_blocks = manager.get_blocks(request.request_id).blocks[1] + checkpoint_block_idx = 48 // mamba_block_size - 1 + assert mamba_blocks[checkpoint_block_idx].is_null + assert not mamba_blocks[-1].is_null + checkpoint_hash = request.block_hashes[48 // hash_block_size - 1] + assert manager.block_pool.get_cached_block(checkpoint_hash, [1]) is None + + def test_external_mamba_hit_same_block_uses_running_cow_on_continue(): """An external mid-block hit must become a running request even when its first continuation does not need another Mamba block.""" @@ -560,6 +644,7 @@ def test_take_partial_tail_offloads_returns_cow_target(): shapes=(1, 1), dtypes=(torch.float32,), mamba_cache_mode="align", + num_prefill_checkpoint_blocks=1, ), ), ], diff --git a/tests/v1/core/test_mamba_align_chunk_split.py b/tests/v1/core/test_mamba_align_chunk_split.py index 31671caf766b..bc809fc3b530 100644 --- a/tests/v1/core/test_mamba_align_chunk_split.py +++ b/tests/v1/core/test_mamba_align_chunk_split.py @@ -37,7 +37,9 @@ MAMBA_GROUP_ID = 1 -def _make_hybrid_kv_cache_manager() -> KVCacheManager: +def _make_hybrid_kv_cache_manager( + num_prefill_checkpoint_blocks: int = 0, +) -> KVCacheManager: config = KVCacheConfig( num_blocks=10000, kv_cache_tensors=[], @@ -59,6 +61,7 @@ def _make_hybrid_kv_cache_manager() -> KVCacheManager: dtypes=(torch.float32,), mamba_cache_mode="align", num_speculative_blocks=NUM_SPEC, + num_prefill_checkpoint_blocks=num_prefill_checkpoint_blocks, ), ), ], @@ -78,6 +81,7 @@ def _split( num_new_tokens: int, use_eagle: bool = True, partial_hit: bool = False, + num_prefill_checkpoint_blocks: int = 0, ) -> int: """Call the real `Scheduler._mamba_block_aligned_split` on a stub self.""" stub = SimpleNamespace( @@ -88,10 +92,42 @@ def _split( # `prefix_match_unit` finer than the block size (#46384). mamba_partial_cache_hit=partial_hit, hash_block_size=ATTN_BLOCK_SIZE, + mamba_has_prefill_checkpoint_blocks=( + num_prefill_checkpoint_blocks > 0 and not use_eagle + ), ) return Scheduler._mamba_block_aligned_split(stub, request, num_new_tokens) +@pytest.mark.parametrize( + ("prompt_len", "num_new_tokens", "use_eagle", "expected"), + [ + (2002, 2002, False, 2002), + (3602, 2000, False, 2000), + (3602, 3602, True, MAMBA_BLOCK_SIZE), + ], +) +def test_internal_checkpoint_split( + prompt_len: int, num_new_tokens: int, use_eagle: bool, expected: int +) -> None: + (request,) = create_requests(1, num_tokens=prompt_len, block_size=ATTN_BLOCK_SIZE) + assert ( + _split( + request, + num_new_tokens, + use_eagle=use_eagle, + num_prefill_checkpoint_blocks=1, + ) + == expected + ) + if not use_eagle: + manager = _make_hybrid_kv_cache_manager(num_prefill_checkpoint_blocks=1) + assert manager.allocate_slots(request, expected) is not None + mamba_manager = manager.coordinator.single_type_managers[MAMBA_GROUP_ID] + blocks = mamba_manager.req_to_blocks[request.request_id] + assert all(not block.is_null for block in blocks) # checkpoint + running state + + def _run_chunked_prefill( manager: KVCacheManager, request: Request, budgets: list[int] ) -> dict[int, int]: @@ -211,8 +247,15 @@ def test_poisoning_is_block_size_independent( @pytest.mark.parametrize("partial_hit", [False, True]) @pytest.mark.parametrize("resume_at", [331, 1599, 1601, 2531, 3011]) +@pytest.mark.parametrize( + ("num_prefill_checkpoint_blocks", "use_eagle"), + [(0, True), (1, False)], +) def test_unaligned_resume_never_runs_past_its_block( - partial_hit: bool, resume_at: int + partial_hit: bool, + resume_at: int, + num_prefill_checkpoint_blocks: int, + use_eagle: bool, ) -> None: """A prefill resuming mid-block must re-align before crossing a boundary. @@ -227,7 +270,13 @@ def test_unaligned_resume_never_runs_past_its_block( pos, ends = resume_at, [] while pos < prompt_len: request.num_computed_tokens = pos - num_new = _split(request, prompt_len - pos, partial_hit=partial_hit) + num_new = _split( + request, + prompt_len - pos, + use_eagle=use_eagle, + partial_hit=partial_hit, + num_prefill_checkpoint_blocks=num_prefill_checkpoint_blocks, + ) assert num_new > 0, f"no progress at {pos}" if pos % MAMBA_BLOCK_SIZE != 0: block_end = (pos // MAMBA_BLOCK_SIZE + 1) * MAMBA_BLOCK_SIZE diff --git a/tests/v1/core/test_prefix_caching.py b/tests/v1/core/test_prefix_caching.py index 038d45be30a2..6fc2131b6e87 100644 --- a/tests/v1/core/test_prefix_caching.py +++ b/tests/v1/core/test_prefix_caching.py @@ -1089,6 +1089,7 @@ def test_hybrid_cache_mamba_align_shared_prefix_detection(): use_eagle=False, hash_block_size=block_size, mamba_partial_cache_hit=False, + mamba_has_prefill_checkpoint_blocks=False, ) req_2.shared_prefix_boundary = shared_prefix_boundary num_new_tokens_adjusted = Scheduler._mamba_block_aligned_split( diff --git a/vllm/models/kimi_k3/nvidia/kda.py b/vllm/models/kimi_k3/nvidia/kda.py index ec6b514f880d..d2ca6f42f525 100644 --- a/vllm/models/kimi_k3/nvidia/kda.py +++ b/vllm/models/kimi_k3/nvidia/kda.py @@ -2,6 +2,7 @@ # SPDX-FileCopyrightText: Copyright contributors to the vLLM project from collections.abc import Callable +from dataclasses import replace import torch from einops import rearrange @@ -45,7 +46,10 @@ from vllm.platforms import current_platform from vllm.third_party.flash_linear_attention.ops.kda import FusedRMSNormGated from vllm.transformers_utils.configs.kimi_linear import KimiLinearConfig +from vllm.triton_utils import tl, triton from vllm.v1.attention.backend import AttentionBackend +from vllm.v1.attention.backends.utils import NULL_BLOCK_ID +from vllm.v1.kv_cache_interface import MambaSpec from vllm.v1.worker.workspace import current_workspace_manager logger = init_logger(__name__) @@ -199,6 +203,8 @@ def _flashkda_prefill( out: torch.Tensor, final_state: torch.Tensor, workspace: torch.Tensor, + checkpoint_state: torch.Tensor | None = None, + checkpoint_offsets: torch.Tensor | None = None, ) -> tuple[torch.Tensor, torch.Tensor]: import vllm._flashkda_C # noqa: F401 @@ -221,10 +227,75 @@ def _flashkda_prefill( initial_state.contiguous(), final_state, cu_seqlens.contiguous(), + checkpoint_state, + checkpoint_offsets.contiguous() if checkpoint_offsets is not None else None, ) return out, final_state +@triton.jit +def _store_cache_checkpoints_kernel( + x_ptr, + conv_state_ptr, + recurrent_checkpoint_ptr, + recurrent_state_ptr, + query_start_loc_ptr, + checkpoint_offsets_ptr, + checkpoint_state_indices_ptr, + x_stride_0: tl.constexpr, + x_stride_1: tl.constexpr, + state_stride_0: tl.constexpr, + state_stride_1: tl.constexpr, + state_stride_2: tl.constexpr, + checkpoint_stride_0: tl.constexpr, + recurrent_state_stride_0: tl.constexpr, + checkpoint_offset_stride: tl.constexpr, + STATE_LEN: tl.constexpr, + WIDTH: tl.constexpr, + RECURRENT_ROW_SIZE: tl.constexpr, + NULL_STATE_IDX: tl.constexpr, + BLOCK_SIZE: tl.constexpr, +): + # store checkpoints to cache + seq_idx = tl.program_id(0) + cols = tl.program_id(1) * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) + state_idx = tl.load(checkpoint_state_indices_ptr + seq_idx) + checkpoint_offset = tl.load( + checkpoint_offsets_ptr + seq_idx * checkpoint_offset_stride + ) + valid_checkpoint = (state_idx != NULL_STATE_IDX) & (checkpoint_offset > 0) + valid_conv = ( + (cols < WIDTH * STATE_LEN) & valid_checkpoint & (checkpoint_offset >= STATE_LEN) + ) + width_idx = cols // STATE_LEN + history_idx = cols % STATE_LEN + checkpoint_end = tl.load(query_start_loc_ptr + seq_idx) + checkpoint_offset + token_idx = checkpoint_end - STATE_LEN + history_idx + values = tl.load( + x_ptr + token_idx * x_stride_0 + width_idx * x_stride_1, + mask=valid_conv, + ) + tl.store( + conv_state_ptr + + state_idx * state_stride_0 + + width_idx * state_stride_1 + + history_idx * state_stride_2, + values, + mask=valid_conv, + ) + + valid_recurrent = (cols < RECURRENT_ROW_SIZE) & valid_checkpoint + recurrent = tl.load( + recurrent_checkpoint_ptr + seq_idx * checkpoint_stride_0 + cols, + mask=valid_recurrent, + ) + tl.store( + recurrent_state_ptr + state_idx * recurrent_state_stride_0 + cols, + recurrent, + mask=valid_recurrent, + ) + + def resolve_kda_prefill_backend( backend: str, head_dim: int, @@ -471,6 +542,7 @@ def __init__( self._flashkda_buffer_specs = ( ((1, T, H, D), self.model_config.dtype), ((N, H, D, D), self.get_state_dtype()[1]), + ((N, H, D, D), self.get_state_dtype()[1]), ((workspace_size,), torch.uint8), ) @@ -518,6 +590,14 @@ def __init__( raise ValueError(f"Duplicate layer name: {prefix}") compilation_config.static_forward_context[prefix] = self + def get_kv_cache_spec(self, vllm_config: VllmConfig) -> MambaSpec: + spec = super().get_kv_cache_spec(vllm_config) + assert isinstance(spec, MambaSpec) + return replace( + spec, + num_prefill_checkpoint_blocks=int(self.kda_prefill_backend == "flashkda"), + ) + def forward( self, hidden_states: torch.Tensor, @@ -590,6 +670,7 @@ def _forward( spec_query_start_loc = m.spec_query_start_loc num_accepted_tokens = m.num_accepted_tokens num_actual_tokens = m.num_actual_tokens + checkpoint = m.checkpoint has_spec_decode = m.num_spec_decodes > 0 mixed_qkv = mixed_qkv[:num_actual_tokens] g1 = g1[:, :num_actual_tokens] @@ -780,7 +861,7 @@ def _prefill_conv( if self.kda_prefill_backend == "flashkda": assert self.gate_lower_bound is not None assert self._flashkda_buffer_specs is not None - workspace_out, final_state, workspace = ( + workspace_out, final_state, checkpoint_state, workspace = ( current_workspace_manager().get_simultaneous( *self._flashkda_buffer_specs ) @@ -788,24 +869,85 @@ def _prefill_conv( flashkda_out = ( workspace_out if has_spec_decode else core_attn_out )[:, : q_ns.shape[1]] - ( - core_attn_out_non_spec, - last_recurrent_state, - ) = _flashkda_prefill( - q=q_ns, - k=k_ns, - v=v_ns, - g=g1_ns, - beta=beta_ns, - A_log=self.A_log, - dt_bias=self.dt_bias, - lower_bound=self.gate_lower_bound, - initial_state=initial_state, - cu_seqlens=non_spec_query_start_loc, - out=flashkda_out, - final_state=final_state[: initial_state.shape[0]], - workspace=workspace, - ) + if checkpoint is not None: + assert non_spec_query_start_loc is not None + num_sequences = initial_state.shape[0] + assert checkpoint.checkpoint_offsets.shape == (num_sequences,) + final_state = final_state[:num_sequences] + checkpoint_state = checkpoint_state[:num_sequences] + checkpoint_offsets = checkpoint.checkpoint_offsets + _flashkda_prefill( + q=q_ns, + k=k_ns, + v=v_ns, + g=g1_ns, + beta=beta_ns, + A_log=self.A_log, + dt_bias=self.dt_bias, + lower_bound=self.gate_lower_bound, + initial_state=initial_state, + cu_seqlens=non_spec_query_start_loc, + out=flashkda_out, + final_state=final_state, + workspace=workspace, + checkpoint_state=checkpoint_state, + checkpoint_offsets=checkpoint_offsets, + ) + core_attn_out_non_spec = flashkda_out + last_recurrent_state = final_state + state_len = conv_state.shape[-1] + width = mixed_qkv_ns.shape[-1] + recurrent_row_size = checkpoint_state[0].numel() + block_size = 256 + _store_cache_checkpoints_kernel[ + ( + checkpoint_offsets.numel(), + triton.cdiv( + max(width * state_len, recurrent_row_size), + block_size, + ), + ) + ]( + mixed_qkv_ns, + conv_state, + checkpoint_state, + recurrent_state, + non_spec_query_start_loc, + checkpoint_offsets, + checkpoint.state_indices, + mixed_qkv_ns.stride(0), + mixed_qkv_ns.stride(1), + conv_state.stride(0), + conv_state.stride(1), + conv_state.stride(2), + checkpoint_state.stride(0), + recurrent_state.stride(0), + checkpoint_offsets.stride(0), + state_len, + width, + recurrent_row_size, + NULL_BLOCK_ID, + block_size, + ) + else: + ( + core_attn_out_non_spec, + last_recurrent_state, + ) = _flashkda_prefill( + q=q_ns, + k=k_ns, + v=v_ns, + g=g1_ns, + beta=beta_ns, + A_log=self.A_log, + dt_bias=self.dt_bias, + lower_bound=self.gate_lower_bound, + initial_state=initial_state, + cu_seqlens=non_spec_query_start_loc, + out=flashkda_out, + final_state=final_state[: initial_state.shape[0]], + workspace=workspace, + ) else: ( core_attn_out_non_spec, diff --git a/vllm/models/kimi_k3/nvidia/kda_metadata.py b/vllm/models/kimi_k3/nvidia/kda_metadata.py index 2154eaa3c51f..b903beac249a 100644 --- a/vllm/models/kimi_k3/nvidia/kda_metadata.py +++ b/vllm/models/kimi_k3/nvidia/kda_metadata.py @@ -44,6 +44,9 @@ ) +FLASHKDA_CHUNK_SIZE = 16 + + @cache def _metadata_launch_pdl() -> bool: return current_platform.is_arch_support_pdl() @@ -257,12 +260,19 @@ class KDARecoverSSMCommitMetadata: align: KDARecoverSSMAlignMetadata | None +@dataclass +class KDACheckpointMetadata: + checkpoint_offsets: torch.Tensor + state_indices: torch.Tensor + + @dataclass class KimiK3KDAMetadata(GDNAttentionMetadata, RecoverSSMMetadata): recoverssm_commit: KDARecoverSSMCommitMetadata | None = None recoverssm_context: "KDARecoverSSMCommitContext | None" = field( default=None, repr=False, compare=False ) + checkpoint: KDACheckpointMetadata | None = None def commit_recoverssm_state( self, num_accepted_tokens: torch.Tensor @@ -557,6 +567,62 @@ def build( # type: ignore[override] else: has_initial_state = None + checkpoint = None + if ( + num_prefills > 0 + and self.kv_cache_spec.num_prefill_checkpoint_blocks > 0 + and self.vllm_config.cache_config.mamba_cache_mode == "align" + ): + # prepare checkpoint metadata + assert m.seq_lens_cpu_upper_bound is not None + request_rows = list(range(m.num_reqs)) + if spec_sequence_masks_cpu is not None: + # get non spec request rows + request_rows = active_non_spec_mask_cpu.nonzero().flatten().tolist() + all_query_lens = query_start_loc_cpu.diff().tolist() + query_lens = [all_query_lens[row] for row in request_rows] + seq_lens = m.seq_lens_cpu_upper_bound.tolist() + block_size = self.kv_cache_spec.block_size + checkpoint_splits = [] + checkpoint_cols = [] + for row, query_len in zip(request_rows, query_lens): + seq_len = seq_lens[row] + offset = seq_len // block_size * block_size - (seq_len - query_len) + # offset should be less than query_len + valid = ( + seq_len % block_size != 0 + and 0 < offset < query_len + and offset % FLASHKDA_CHUNK_SIZE == 0 + ) + offset = offset if valid else 0 + first_len = offset or query_len + checkpoint_splits.append((first_len, query_len - first_len)) + checkpoint_cols.append(seq_len // block_size - 1 if valid else -1) + if any(tail for _, tail in checkpoint_splits): + checkpoint_offsets_tensor = async_tensor_h2d( + [first if tail else 0 for first, tail in checkpoint_splits], + dtype=torch.int32, + device=query_start_loc.device, + ) + request_rows_tensor = async_tensor_h2d( + request_rows, dtype=torch.int64, device=query_start_loc.device + ) + checkpoint_cols_tensor = async_tensor_h2d( + checkpoint_cols, dtype=torch.int64, device=query_start_loc.device + ) + checkpoint_state_indices = m.block_table_tensor[ + request_rows_tensor, checkpoint_cols_tensor + ] + checkpoint_state_indices = torch.where( + checkpoint_cols_tensor >= 0, + checkpoint_state_indices, + NULL_BLOCK_ID, + ) + checkpoint = KDACheckpointMetadata( + checkpoint_offsets_tensor, + checkpoint_state_indices, + ) + # Prepare per-request tensors for cudagraph replay. num_actual_tokens # may be token-padded, while state/query/acceptance metadata is indexed # by request. @@ -647,6 +713,7 @@ def build( # type: ignore[override] nums_dict=nums_dict, batch_ptr=batch_ptr, token_chunk_offset_ptr=token_chunk_offset_ptr, + checkpoint=checkpoint, ) diff --git a/vllm/v1/core/sched/scheduler.py b/vllm/v1/core/sched/scheduler.py index b41727f273cc..c06760b80687 100644 --- a/vllm/v1/core/sched/scheduler.py +++ b/vllm/v1/core/sched/scheduler.py @@ -53,7 +53,7 @@ ) from vllm.v1.core.sched.utils import check_stop, remove_all from vllm.v1.engine import EngineCoreEventType, EngineCoreOutput, EngineCoreOutputs -from vllm.v1.kv_cache_interface import KVCacheConfig +from vllm.v1.kv_cache_interface import KVCacheConfig, MambaSpec from vllm.v1.metrics.perf import ModelMetrics, PerfStats from vllm.v1.metrics.stats import ( PrefixCacheStats, @@ -324,6 +324,16 @@ def __init__( self.need_mamba_block_aligned_split = ( self.has_mamba_layers and self.cache_config.mamba_cache_mode == "align" ) + self.mamba_has_prefill_checkpoint_blocks = ( + self.has_mamba_layers + # TODO: support spec decoding + and not self.use_eagle + and all( + not isinstance(group.kv_cache_spec, MambaSpec) + or group.kv_cache_spec.num_prefill_checkpoint_blocks > 0 + for group in kv_cache_config.kv_cache_groups + ) + ) # A finer prefix_match_unit is configured: a mamba partial tail entry # can only be registered by a step ending exactly at the prompt's last # hash boundary, so the split adds that stop. @@ -406,12 +416,17 @@ def _mamba_block_aligned_split( last_cache_position = max(last_cache_position - block_size, 0) end = start + num_new_tokens + use_internal_checkpoint = ( + self.mamba_has_prefill_checkpoint_blocks and start % block_size == 0 + ) + if use_internal_checkpoint: + last_cache_position = 0 # Invariant: slot p holds the state after exactly (p + 1) * block_size # tokens. State is written at chunk ends, so chunk ends must be block # aligned. Exempt: the prompt's last chunk, whose slot decode advances # to the boundary. A block too wide for one chunk advances sub-block # and re-aligns at the next boundary. - if end < prefill_end: + if end < prefill_end and not use_internal_checkpoint: max_prefill_tokens = self.max_num_scheduled_tokens long_prefill_threshold = self.scheduler_config.long_prefill_token_threshold if long_prefill_threshold > 0: diff --git a/vllm/v1/core/single_type_kv_cache_manager.py b/vllm/v1/core/single_type_kv_cache_manager.py index 75acb7f4973e..ceb56a59875f 100644 --- a/vllm/v1/core/single_type_kv_cache_manager.py +++ b/vllm/v1/core/single_type_kv_cache_manager.py @@ -1285,6 +1285,9 @@ def __init__( self.last_state_block_idx: dict[str, int] = {} # The set of the requests that have been allocated blocks self._allocated_block_reqs: set[str] = set() + # Number of internal checkpoint blocks required by each request's + # current allocation. + self._num_checkpoint_blocks: dict[str, int] = {} # Requests that registered their own last-prompt-boundary partial # tail (producers). On the next step's CoW the boundary state moves # into a private cow_block; we record that block for connector @@ -1464,6 +1467,23 @@ def get_num_common_prefix_blocks(self, running_request_id: str) -> int: """ return 0 + def _needs_internal_checkpoint( + self, + request_id: str, + num_tokens: int, + num_computed_tokens: int, + ) -> bool: + assert isinstance(self.kv_cache_spec, MambaSpec) + checkpoint_idx = cdiv(num_tokens, self.block_size) - 2 + blocks = self.req_to_blocks[request_id] + return ( + self.kv_cache_spec.num_prefill_checkpoint_blocks > 0 + and num_tokens % self.block_size != 0 + and num_computed_tokens % self.block_size == 0 + and checkpoint_idx >= 0 + and (checkpoint_idx >= len(blocks) or blocks[checkpoint_idx].is_null) + ) + def get_num_blocks_to_allocate( self, request_id: str, @@ -1526,18 +1546,17 @@ def get_num_blocks_to_allocate( ) if has_partial_hit: num_new_blocks = max(num_new_blocks, 0) + 1 + checkpoint_block = int( + self._needs_internal_checkpoint( + request_id, num_tokens, total_computed_tokens + ) + ) + if not apply_admission_cap: + self._num_checkpoint_blocks[request_id] = checkpoint_block if num_new_blocks > 0: - if request_id in self._allocated_block_reqs: - # Old request. Needs at most 1 more blocks as we can reuse the - # speculative blocks in previous step. - num_new_blocks = 1 + int(has_partial_hit) - else: - # First prefill. Allocate 1 block for running state, the - # speculative blocks, and one extra block if a partial cache - # hit must be copy-on-written before the new tokens run. - num_new_blocks = ( - 1 + self.num_speculative_blocks + int(has_partial_hit) - ) + num_new_blocks = 1 + int(has_partial_hit) + checkpoint_block + if request_id not in self._allocated_block_reqs: + num_new_blocks += self.num_speculative_blocks num_evictable_computed_blocks = self._get_num_evictable_blocks( new_computed_blocks @@ -1569,6 +1588,7 @@ def allocate_new_blocks( num_required_blocks = ( cdiv(num_tokens, self.block_size) + self.num_speculative_blocks ) + checkpoint_block = self._num_checkpoint_blocks.get(request_id, 0) partial_hit = self._partial_hit_reqs.get(request_id) has_partial_hit = partial_hit is not None # `num_required_blocks` might be less than `len(req_blocks)` if blocks are @@ -1596,11 +1616,11 @@ def allocate_new_blocks( ) # null blocks if prev_block_len < num_skipped_blocks: + # minus the internal checkpoint block + # so we don't set null for that block + null_end = num_skipped_blocks - checkpoint_block req_blocks.extend( - [ - self._null_block - for _ in range(prev_block_len, num_skipped_blocks) - ] + [self._null_block for _ in range(prev_block_len, null_end)] ) if blocks_allocated: @@ -1616,12 +1636,10 @@ def allocate_new_blocks( num_new_blocks = num_required_blocks - len(req_blocks) if has_partial_hit: num_new_blocks = max(num_new_blocks, 0) + 1 - if blocks_allocated: - assert num_new_blocks <= 1 + int(has_partial_hit) - else: - assert num_new_blocks <= self.num_speculative_blocks + 1 + int( - has_partial_hit - ) + max_new_blocks = 1 + int(has_partial_hit) + checkpoint_block + if not blocks_allocated: + max_new_blocks += self.num_speculative_blocks + assert num_new_blocks <= max_new_blocks new_blocks = self.block_pool.get_new_blocks(num_new_blocks) returned_blocks = req_blocks[prev_block_len:] if partial_hit is not None: @@ -1670,6 +1688,7 @@ def pop_blocks_for_free(self, request_id: str) -> list[KVCacheBlock]: if self.mamba_cache_mode == "align": self._allocated_block_reqs.discard(request_id) self.last_state_block_idx.pop(request_id, None) + self._num_checkpoint_blocks.pop(request_id, None) self._producer_partial_tail_reqs.pop(request_id, None) # A hand-off whose request died in this same scheduling pass must # not reach the connector: its unpin hook (free) has already run. diff --git a/vllm/v1/kv_cache_interface.py b/vllm/v1/kv_cache_interface.py index b660f4931709..55c9e806a5e7 100644 --- a/vllm/v1/kv_cache_interface.py +++ b/vllm/v1/kv_cache_interface.py @@ -815,6 +815,7 @@ class MambaSpec(KVCacheSpec): mamba_type: MambaAttentionBackendEnum = MambaAttentionBackendEnum.MAMBA2 mamba_cache_mode: str = "none" num_speculative_blocks: int = 0 + num_prefill_checkpoint_blocks: int = 0 num_heads: int = 1 tokens_per_state: int = -1 @@ -843,7 +844,9 @@ def max_memory_usage_bytes(self, vllm_config: VllmConfig) -> int: cdiv(max_model_len, self.block_size) + self.num_speculative_blocks ) * self.page_size_bytes elif vllm_config.cache_config.mamba_cache_mode == "align": - return self.page_size_bytes * (2 + self.num_speculative_blocks) + return self.page_size_bytes * ( + 2 + self.num_speculative_blocks + self.num_prefill_checkpoint_blocks + ) else: return self.page_size_bytes * (1 + self.num_speculative_blocks) @@ -865,6 +868,7 @@ def is_uniform_with_collection( return all( isinstance(spec, MambaSpec) and spec.num_speculative_blocks == self.num_speculative_blocks + and spec.num_prefill_checkpoint_blocks == self.num_prefill_checkpoint_blocks for spec in kv_cache_specs.values() ) From 236f78cc5c39362f80c0c7df1ffe810466525f33 Mon Sep 17 00:00:00 2001 From: Shuolei Wang <46160365+ShuoleiWang@users.noreply.github.com> Date: Sun, 23 Aug 2026 01:39:58 +0800 Subject: [PATCH 013/235] [Core][RL] Support sparse checkpoint updates through native weight loaders (#50723) Signed-off-by: Shuolei Wang Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> --- .../test_checkpoint_weight_patch.py | 110 +++++++ .../model_loader/checkpoint_weight_patch.py | 273 ++++++++++++++++++ 2 files changed, 383 insertions(+) create mode 100644 tests/model_executor/model_loader/test_checkpoint_weight_patch.py create mode 100644 vllm/model_executor/model_loader/checkpoint_weight_patch.py diff --git a/tests/model_executor/model_loader/test_checkpoint_weight_patch.py b/tests/model_executor/model_loader/test_checkpoint_weight_patch.py new file mode 100644 index 000000000000..8a3579fea3b0 --- /dev/null +++ b/tests/model_executor/model_loader/test_checkpoint_weight_patch.py @@ -0,0 +1,110 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project + +import pytest +import torch + +from vllm.model_executor.model_loader.checkpoint_weight_patch import ( + CheckpointWeightPatch, + load_checkpoint_weight_patches, +) + +pytestmark = pytest.mark.cpu_test + + +class _PackedTPModel(torch.nn.Module): + """Model that packs the TP rank 1 shards of Q and K into one parameter.""" + + def __init__(self): + super().__init__() + self.packed_weight = torch.nn.Parameter(torch.full((4,), -1.0)) + self.load_calls: list[list[str]] = [] + + def load_weights(self, weights): + weights = list(weights) + self.load_calls.append([name for name, _ in weights]) + loaded_names = set() + for name, loaded_weight in weights: + if name == "q_proj.weight": + destination = self.packed_weight.data[:2] + elif name == "k_proj.weight": + destination = self.packed_weight.data[2:] + else: + raise AssertionError(name) + # Slice each full checkpoint tensor and load its second half for TP rank 1. + destination.copy_(loaded_weight.narrow(0, 2, 2)) + loaded_names.add(name) + return loaded_names + + +def _make_patch( + name: str, + *, + values: list[float], + indices: list[int] | None = None, +) -> CheckpointWeightPatch: + return CheckpointWeightPatch( + name=name, + shape=(4,), + dtype=torch.float32, + values=torch.tensor(values), + indices=None if indices is None else torch.tensor(indices, dtype=torch.int32), + ) + + +def test_dense_and_sparse_patches_follow_packed_tp_loader(): + model = _PackedTPModel() + + # Seed the TP-local packed weights through the dense path. + dense_loaded = load_checkpoint_weight_patches( + model, + [ + _make_patch("q_proj.weight", values=[0.0, 1.0, 2.0, 3.0]), + _make_patch("k_proj.weight", values=[10.0, 11.0, 12.0, 13.0]), + ], + ) + assert dense_loaded == {"q_proj.weight", "k_proj.weight"} + assert torch.equal(model.packed_weight, torch.tensor([2.0, 3.0, 12.0, 13.0])) + + # Sparse indices address the full checkpoint tensor; index 0 is outside this shard. + model.load_calls.clear() + original_copy = torch.Tensor.copy_ + sparse_loaded = load_checkpoint_weight_patches( + model, + [ + _make_patch( + "q_proj.weight", + indices=[0, 3], + values=[100.0, 30.0], + ), + _make_patch("k_proj.weight", indices=[2], values=[20.0]), + _make_patch("q_proj.weight", indices=[2], values=[22.0]), + ], + ) + + assert sparse_loaded == {"q_proj.weight", "k_proj.weight"} + assert model.load_calls == [ + ["q_proj.weight", "k_proj.weight"], + ["q_proj.weight"], + ] + assert torch.equal(model.packed_weight, torch.tensor([22.0, 30.0, 20.0, 13.0])) + assert torch.Tensor.copy_ is original_copy + + # The empty result tells the helper that the loader intentionally made no write. + class NoLocalWeightModel(torch.nn.Module): + """Model that deliberately owns none of the supplied weights.""" + + def __init__(self): + super().__init__() + self.weight = torch.nn.Parameter(torch.tensor([1.0])) + + def load_weights(self, _weights): + return set() + + no_local_weight_model = NoLocalWeightModel() + no_local_weight_loaded = load_checkpoint_weight_patches( + no_local_weight_model, + [_make_patch("q_proj.weight", indices=[3], values=[30.0])], + ) + assert no_local_weight_loaded == set() + assert torch.equal(no_local_weight_model.weight, torch.tensor([1.0])) diff --git a/vllm/model_executor/model_loader/checkpoint_weight_patch.py b/vllm/model_executor/model_loader/checkpoint_weight_patch.py new file mode 100644 index 000000000000..0425aabc7dd1 --- /dev/null +++ b/vllm/model_executor/model_loader/checkpoint_weight_patch.py @@ -0,0 +1,273 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Apply dense or sparse weight updates in checkpoint coordinates. + +The model's ``load_weights`` method still handles checkpoint-name mapping, +TP slicing, and packed parameters. Sparse patches use NaN to mark unchanged +checkpoint elements. The supported sparse-loader contract permits only final +same-shaped floating-point ``Tensor.copy_`` writes. Composed, multi-stage, and +custom-write loaders are unsupported. +""" + +from __future__ import annotations + +import math +from collections.abc import Iterable +from typing import NamedTuple + +import torch + +_DEFAULT_PATCH_CHUNK_BYTES = 512 << 20 + +__all__ = [ + "CheckpointWeightPatch", + "load_checkpoint_weight_patches", +] + + +class CheckpointWeightPatch(NamedTuple): + """Describe one dense or sparse checkpoint-coordinate update. + + Attributes: + name: Checkpoint weight name passed to the model loader. + shape: Full checkpoint tensor shape. + dtype: Checkpoint tensor dtype. + values: The flattened full tensor for a dense patch, or the values at + ``indices`` for a sparse patch. + indices: Flat indices into the full checkpoint tensor described by + ``shape``. ``None`` makes ``values`` a dense replacement. + """ + + name: str + shape: tuple[int, ...] + dtype: torch.dtype + values: torch.Tensor + indices: torch.Tensor | None = None + + +def _load_nan_masked_weights( + model: torch.nn.Module, + weights: list[tuple[str, torch.Tensor]], +) -> Iterable[str] | None: + """Load NaN-masked checkpoint tensors while preserving runtime values at NaNs.""" + # Save the original copy_ so it can be restored after loading. + original_copy = torch.Tensor.copy_ + + # Enforce the sparse-loader contract and preserve destination values where + # the corresponding source is NaN. + def copy_non_nan_( + destination: torch.Tensor, + source: torch.Tensor, + non_blocking: bool = False, + ) -> torch.Tensor: + if not ( + destination.is_floating_point() + and source.is_floating_point() + and destination.shape == source.shape + ): + raise NotImplementedError( + "Sparse checkpoint patches require a same-shaped " + "floating-point final copy" + ) + source = source.to(dtype=destination.dtype, device=destination.device) + return torch.where( + torch.isnan(source), + destination, + source, + out=destination, + ) + + torch.Tensor.copy_ = copy_non_nan_ + try: + return model.load_weights(weights) + finally: + torch.Tensor.copy_ = original_copy + + +def _validate_patch_structure(patch: CheckpointWeightPatch) -> int: + if not patch.name: + raise ValueError("Checkpoint weight name must be non-empty") + if not isinstance(patch.dtype, torch.dtype): + raise TypeError(f"{patch.name}: dtype must be a torch.dtype") + if any(dim < 0 for dim in patch.shape): + raise ValueError(f"{patch.name}: shape dimensions must be non-negative") + if patch.values.ndim != 1: + raise ValueError(f"{patch.name}: patch values must be a 1D tensor") + + numel = math.prod(patch.shape) + if patch.indices is None: + if patch.values.numel() != numel: + raise ValueError( + f"{patch.name}: dense patch has {patch.values.numel()} values " + f"for a {numel}-element tensor" + ) + return numel + + if not patch.dtype.is_floating_point: + raise TypeError( + f"{patch.name}: sparse checkpoint patches require a floating dtype " + "because NaN is the unchanged-value sentinel" + ) + if patch.indices.ndim != 1: + raise ValueError(f"{patch.name}: sparse patch indices must be a 1D tensor") + if patch.indices.dtype not in (torch.int32, torch.int64): + raise TypeError(f"{patch.name}: sparse patch indices must be int32 or int64") + if patch.indices.numel() != patch.values.numel(): + raise ValueError( + f"{patch.name}: sparse indices and values must have matching lengths" + ) + + return numel + + +def _load_weight_chunk( + model: torch.nn.Module, + weights: list[tuple[str, torch.Tensor]], + *, + sparse: bool, +) -> set[str]: + if not weights: + return set() + if sparse: + loaded_names = _load_nan_masked_weights(model, weights) + else: + loaded_names = model.load_weights(weights) + return set() if loaded_names is None else set(loaded_names) + + +@torch.no_grad() +def load_checkpoint_weight_patches( + model: torch.nn.Module, + patches: Iterable[CheckpointWeightPatch], + *, + max_chunk_bytes: int = _DEFAULT_PATCH_CHUNK_BYTES, + validate_unique_indices: bool = True, +) -> set[str]: + """Load ordered patches through the model's checkpoint loader. + + A call may contain either dense or sparse patches, but not both. A repeated + name starts a new loader call, so patches for that weight are applied in input + order. A later patch may update positions changed by an earlier patch. Each + sparse patch creates a full checkpoint-shaped tensor whose NaNs mark unchanged + elements. Each locally applied sparse patch must use one final same-shaped + floating-point ``Tensor.copy_``. Intermediate or unrelated ``copy_`` calls, + composed loaders, and custom write paths are unsupported. ``max_chunk_bytes`` + is a batching target; one tensor may exceed it. + + Online checkpoint-format dense updates require the caller to manage vLLM's + layerwise reload lifecycle. Sparse online updates modify initialized model + tensors in place and must not use that lifecycle. Patch shapes, dtypes, value + lengths, and sparse indices are validated before loading begins. + + This function does not roll back a partial update. If a loader fails after a + write, some destinations may already be changed. Do not serve from the affected + worker until a known baseline has been restored or the worker has been + restarted. + + Args: + model: Model whose native ``load_weights`` method applies the patches. + patches: Dense replacements or sparse updates in checkpoint coordinates. + max_chunk_bytes: Target checkpoint tensor bytes per + ``model.load_weights`` call. A single tensor may exceed this target. + validate_unique_indices: Whether to reject duplicate sparse indices within + each patch. Repeated patches may update the same positions. Disable + only for a trusted producer that already guarantees unique positions + within each patch. + + Returns: + Union of the weight names reported by all ``model.load_weights`` calls. + """ + + if max_chunk_bytes <= 0: + raise ValueError("max_chunk_bytes must be positive") + + patch_list = list(patches) + patch_numels = [_validate_patch_structure(patch) for patch in patch_list] + # All patches in one call must use the same dense or sparse representation. + sparse_flags = {patch.indices is not None for patch in patch_list} + if len(sparse_flags) > 1: + raise ValueError("Dense and sparse checkpoint patches cannot be mixed") + sparse = sparse_flags == {True} + + # Keep validation results on their devices and read one combined result per + # device, instead of synchronizing the CPU once for every patch. + checks_by_device: dict[torch.device, list[torch.Tensor]] = {} + if sparse: + for patch, numel in zip(patch_list, patch_numels, strict=True): + assert patch.indices is not None + + # NaN is reserved to mean "leave this checkpoint position unchanged," + # so it cannot also be supplied as a new sparse value. + patch_checks = [torch.isnan(patch.values).any()] + indices = patch.indices.to(device=patch.values.device) + if indices.numel(): + # Sparse indices address the flattened full checkpoint tensor. + patch_checks.append( + torch.logical_or(indices < 0, indices >= numel).any() + ) + if validate_unique_indices: + # Equal neighboring indices after sorting indicate duplicates. + sorted_indices = torch.sort(indices).values + patch_checks.append( + (sorted_indices[1:] == sorted_indices[:-1]).any() + ) + + checks_by_device.setdefault(patch.values.device, []).extend(patch_checks) + + for device_checks in checks_by_device.values(): + if torch.stack(device_checks).any().item(): + invalid_reasons = ( + "NaN values, out-of-range indices, or duplicate indices" + if validate_unique_indices + else "NaN values or out-of-range indices" + ) + raise ValueError(f"Sparse checkpoint patches contain {invalid_reasons}") + + loaded_names: set[str] = set() + # Collect checkpoint tensors for one model.load_weights() call. + weight_chunk: list[tuple[str, torch.Tensor]] = [] + chunk_names: set[str] = set() + chunk_bytes = 0 + + for patch, numel in zip(patch_list, patch_numels, strict=True): + # Count the full checkpoint tensor, including sparse NaN staging. + checkpoint_tensor_bytes = ( + numel * torch.empty((), dtype=patch.dtype).element_size() + ) + # Flush before exceeding the batching target. Repeated names also start a + # new loader call so their patches are applied in input order. + if weight_chunk and ( + chunk_bytes + checkpoint_tensor_bytes > max_chunk_bytes + or patch.name in chunk_names + ): + loaded_names.update(_load_weight_chunk(model, weight_chunk, sparse=sparse)) + weight_chunk = [] + chunk_names = set() + chunk_bytes = 0 + + if patch.indices is None: + checkpoint_tensor = patch.values.to(dtype=patch.dtype).view(patch.shape) + else: + # NaNs mark unchanged checkpoint positions; index_copy_ fills only + # the positions supplied by this sparse patch. + flat = torch.full( + (numel,), + float("nan"), + dtype=patch.dtype, + device=patch.values.device, + ) + if patch.values.numel(): + flat.index_copy_( + 0, + patch.indices.to(device=flat.device, dtype=torch.long), + patch.values.to(dtype=patch.dtype, device=flat.device), + ) + checkpoint_tensor = flat.view(patch.shape) + + weight_chunk.append((patch.name, checkpoint_tensor)) + chunk_names.add(patch.name) + chunk_bytes += checkpoint_tensor_bytes + + # Load the final partial chunk; the helper also accepts an empty chunk. + loaded_names.update(_load_weight_chunk(model, weight_chunk, sparse=sparse)) + return loaded_names From 811d12ee26fa119cf8faae44d80330becf1e7fe3 Mon Sep 17 00:00:00 2001 From: Wentao Ye <44945378+yewentao256@users.noreply.github.com> Date: Sat, 22 Aug 2026 15:41:27 -0400 Subject: [PATCH 014/235] [Mypy Fix] Mypy fix for "vllm/model_executor/models/[eE][fF]" (#53381) Signed-off-by: yewentao256 --- tools/pre_commit/mypy.py | 2 - vllm/compilation/decorators.py | 6 ++- vllm/model_executor/models/eagle2_5_vl.py | 3 +- vllm/model_executor/models/ernie45_moe.py | 3 +- vllm/model_executor/models/ernie45_vl.py | 47 +++++++++++-------- vllm/model_executor/models/ernie45_vl_moe.py | 6 ++- vllm/model_executor/models/exaone4_5.py | 1 + vllm/model_executor/models/exaone4_5_mtp.py | 9 +--- vllm/model_executor/models/exaone_moe.py | 23 ++------- vllm/model_executor/models/exaone_moe_mtp.py | 9 +--- .../models/extract_hidden_states.py | 6 ++- vllm/model_executor/models/falcon.py | 1 + vllm/model_executor/models/fireredasr2.py | 3 +- vllm/model_executor/models/fireredlid.py | 19 ++++---- vllm/model_executor/models/funasr.py | 13 +++-- vllm/model_executor/models/funaudiochat.py | 4 +- 16 files changed, 71 insertions(+), 84 deletions(-) diff --git a/tools/pre_commit/mypy.py b/tools/pre_commit/mypy.py index b5cbc9e136f4..ab220da8a382 100755 --- a/tools/pre_commit/mypy.py +++ b/tools/pre_commit/mypy.py @@ -112,8 +112,6 @@ ] EXCLUDE = [ - r"vllm/model_executor/models/[eE]", - r"vllm/model_executor/models/[fF]", r"vllm/model_executor/models/[gG]", r"vllm/model_executor/models/[hH]", r"vllm/model_executor/models/[iI]", diff --git a/vllm/compilation/decorators.py b/vllm/compilation/decorators.py index ad710a6fe3fc..ceb3ad31705c 100644 --- a/vllm/compilation/decorators.py +++ b/vllm/compilation/decorators.py @@ -106,8 +106,10 @@ def support_torch_compile( @overload def support_torch_compile( *, - dynamic_arg_dims: dict[str, int | list[int] | dict[int, str]] | None, - mark_unbacked_dims: dict[str, int | list[int]] | None, + dynamic_arg_dims: dict[str, int | list[int] | dict[int, str]] | None = None, + mark_unbacked_dims: dict[str, int | list[int]] | None = None, + enable_if: Callable[[VllmConfig], bool] | None = None, + is_encoder: bool = False, ) -> Callable[[type[_T]], type[_T]]: ... diff --git a/vllm/model_executor/models/eagle2_5_vl.py b/vllm/model_executor/models/eagle2_5_vl.py index 30b8173f19cf..8fdd7098f224 100644 --- a/vllm/model_executor/models/eagle2_5_vl.py +++ b/vllm/model_executor/models/eagle2_5_vl.py @@ -150,6 +150,7 @@ def __init__(self, *, vllm_config: VllmConfig, prefix: str = "") -> None: config = vllm_config.model_config.hf_config quant_config = vllm_config.quant_config multimodal_config = vllm_config.model_config.multimodal_config + assert multimodal_config is not None self.config = config self.multimodal_config = multimodal_config @@ -187,7 +188,7 @@ def __init__(self, *, vllm_config: VllmConfig, prefix: str = "") -> None: prefix=maybe_prefix(prefix, "language_model"), ) - self.img_context_token_id = None + self.img_context_token_id: int | None = None self.make_empty_intermediate_tensors = ( self.language_model.make_empty_intermediate_tensors diff --git a/vllm/model_executor/models/ernie45_moe.py b/vllm/model_executor/models/ernie45_moe.py index 746d12c596f1..2683e94f51b9 100644 --- a/vllm/model_executor/models/ernie45_moe.py +++ b/vllm/model_executor/models/ernie45_moe.py @@ -127,7 +127,7 @@ def __init__( self.layer_idx = layer_idx self.tp_size = get_tensor_model_parallel_world_size() - self.moe_num_shared_experts = getattr(config, "moe_num_shared_experts", None) + self.moe_num_shared_experts = getattr(config, "moe_num_shared_experts", 0) self.ep_group = get_ep_group().device_group self.ep_size = self.ep_group.size() self.n_routed_experts: int = config.moe_num_experts @@ -163,6 +163,7 @@ def __init__( torch.empty(config.moe_num_experts, dtype=torch.float32) ) + self.shared_experts: Ernie4_5_MoeMLP | None if self.has_shared_experts: intermediate_size = ( config.moe_intermediate_size * config.moe_num_shared_experts diff --git a/vllm/model_executor/models/ernie45_vl.py b/vllm/model_executor/models/ernie45_vl.py index a1de2ae54e50..96466bad09a5 100644 --- a/vllm/model_executor/models/ernie45_vl.py +++ b/vllm/model_executor/models/ernie45_vl.py @@ -1076,7 +1076,7 @@ class Ernie4_5VLMultiModalProcessor(BaseMultiModalProcessor[Ernie4_5_VLProcessin def _pixel_values_norm( self, pixel_values: torch.Tensor, - mm_kwargs: object, + mm_kwargs: Mapping[str, object], ) -> torch.Tensor: hf_config = self.info.get_hf_config() vision_config = hf_config.vision_config @@ -1125,6 +1125,7 @@ def _apply_hf_processor_main( prompt_text = self.dummy_inputs.get_dummy_text(mm_items.get_all_counts()) + mm_data = dict(mm_data) if "images" not in mm_data: mm_data["images"] = [] if "videos" not in mm_data: @@ -1136,15 +1137,15 @@ def _apply_hf_processor_main( hf_processor, "supports_video_metadata", False ) - if mm_data["videos"] and not supports_video_metadata: + videos = mm_data["videos"] + assert isinstance(videos, Sequence) + if videos and not supports_video_metadata: # Old HF processor, unwrap tuple to pure frames logger.warning_once( "HF processor doesn't support video metadata. " "Timestamps will NOT be rendered. Please upgrade the model." ) - mm_data["videos"] = [ - v[0] if isinstance(v, tuple) else v for v in mm_data["videos"] - ] + mm_data["videos"] = [v[0] if isinstance(v, tuple) else v for v in videos] processor_output = self.info.ctx.call_hf_processor( hf_processor, @@ -1286,6 +1287,7 @@ def get_dummy_mm_data( image_overrides = mm_options.get("image") video_overrides = mm_options.get("video") + assert video_overrides is None or isinstance(video_overrides, VideoDummyOptions) return { "image": self._get_dummy_images( @@ -1541,10 +1543,14 @@ def iter_mm_grid_thw( offset = mm_feature.mm_position.offset if mm_feature.modality == "image": - t, h, w = mm_feature.data["image_grid_thw"].data.tolist() + grid_thw = mm_feature.data["image_grid_thw"].data + assert isinstance(grid_thw, torch.Tensor) + t, h, w = grid_thw.tolist() yield offset, t, h // spatial_conv_size, w // spatial_conv_size elif mm_feature.modality == "video": - t, h, w = mm_feature.data["video_grid_thw"].data.tolist() + grid_thw = mm_feature.data["video_grid_thw"].data + assert isinstance(grid_thw, torch.Tensor) + t, h, w = grid_thw.tolist() yield ( offset, t // temporal_conv_size, @@ -1563,12 +1569,11 @@ def _parse_and_validate_image_input( if pixel_values is None: return None - if pixel_values is not None: - return Ernie4_5_VLImagePixelInputs( - type="pixel_values", - pixel_values=pixel_values, - image_grid_thw=image_grid_thw, - ) + return Ernie4_5_VLImagePixelInputs( + type="pixel_values", + pixel_values=pixel_values, + image_grid_thw=image_grid_thw, + ) def _parse_and_validate_video_input( self, **kwargs: object @@ -1579,12 +1584,11 @@ def _parse_and_validate_video_input( if pixel_values_videos is None: return None - if pixel_values_videos is not None: - return Ernie4_5_VLVideoPixelInputs( - type="pixel_values_videos", - pixel_values_videos=pixel_values_videos, - video_grid_thw=video_grid_thw, - ) + return Ernie4_5_VLVideoPixelInputs( + type="pixel_values_videos", + pixel_values_videos=pixel_values_videos, + video_grid_thw=video_grid_thw, + ) # -- SupportsEncoderCudaGraph protocol methods -- # @@ -1748,6 +1752,7 @@ def postprocess_encoder_output( # the actual batch grid_thw, then scatter the post-merge embeddings. # Ernie only uses the single "default" encoder path. output = outputs["default"] + assert batch_mm_kwargs is not None grid_thw_cpu = batch_mm_kwargs["image_grid_thw"] grid_thw = grid_thw_cpu.to(output.device, non_blocking=True) # The valid token count slices the graph output for the eager @@ -1799,7 +1804,9 @@ def _process_video_input( return video_embeds.split(sizes.tolist()) def _parse_and_validate_multimodal_inputs(self, **kwargs: object) -> dict: - modalities = {} + modalities: dict[ + str, Ernie4_5_VLImageInputs | Ernie4_5_VLVideoInputs | None + ] = {} # Preserve the order of modalities if there are multiple of them # from the order of kwargs. diff --git a/vllm/model_executor/models/ernie45_vl_moe.py b/vllm/model_executor/models/ernie45_vl_moe.py index e18a3fec508e..ccc4b9b868b6 100644 --- a/vllm/model_executor/models/ernie45_vl_moe.py +++ b/vllm/model_executor/models/ernie45_vl_moe.py @@ -232,6 +232,7 @@ def __init__( assert text_moe_layer_start_index <= text_moe_layer_end_index + self.shared_experts: Ernie4_5_VLMoeMLP | None if self.has_shared_experts: intermediate_size = ( config.moe_intermediate_size[0] * config.moe_num_shared_experts @@ -771,9 +772,10 @@ def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]: if is_pp_missing_parameter(name, self): continue # Remapping the name of FP8 kv-scale. - name = maybe_remap_kv_scale_name(name, params_dict) - if name is None: + remapped_name = maybe_remap_kv_scale_name(name, params_dict) + if remapped_name is None: continue + name = remapped_name param = params_dict[name] diff --git a/vllm/model_executor/models/exaone4_5.py b/vllm/model_executor/models/exaone4_5.py index 676faa9d4f3a..b325c6a3043e 100644 --- a/vllm/model_executor/models/exaone4_5.py +++ b/vllm/model_executor/models/exaone4_5.py @@ -332,6 +332,7 @@ def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""): config: Exaone4_5_Config = vllm_config.model_config.hf_config self.vllm_config = vllm_config multimodal_config = vllm_config.model_config.multimodal_config + assert multimodal_config is not None self.use_data_parallel = multimodal_config.mm_encoder_tp_mode == "data" self.config = config diff --git a/vllm/model_executor/models/exaone4_5_mtp.py b/vllm/model_executor/models/exaone4_5_mtp.py index 80e1903f9dd8..d9c7960010a5 100644 --- a/vllm/model_executor/models/exaone4_5_mtp.py +++ b/vllm/model_executor/models/exaone4_5_mtp.py @@ -48,17 +48,11 @@ def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""): model_config = vllm_config.model_config quant_config = vllm_config.quant_config - lora_config = vllm_config.lora_config config = model_config.hf_config text_config = config.text_config self.config = config - lora_vocab = ( - (lora_config.lora_extra_vocab_size * (lora_config.max_loras or 1)) - if lora_config - else 0 - ) - self.vocab_size = config.vocab_size + lora_vocab + self.vocab_size = config.vocab_size self.org_vocab_size = config.vocab_size self.mtp_start_layer_idx = text_config.num_hidden_layers @@ -67,7 +61,6 @@ def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""): self.embed_tokens = VocabParallelEmbedding( self.vocab_size, text_config.hidden_size, - org_num_embeddings=config.vocab_size, ) self.fc = ColumnParallelLinear( diff --git a/vllm/model_executor/models/exaone_moe.py b/vllm/model_executor/models/exaone_moe.py index 6fbeb03af8db..4c19f1ed0fcf 100644 --- a/vllm/model_executor/models/exaone_moe.py +++ b/vllm/model_executor/models/exaone_moe.py @@ -41,7 +41,6 @@ from vllm.model_executor.layers.quantization import QuantizationConfig from vllm.model_executor.layers.rotary_embedding import get_rope from vllm.model_executor.layers.vocab_parallel_embedding import ( - DEFAULT_VOCAB_PADDING_SIZE, ParallelLMHead, VocabParallelEmbedding, ) @@ -112,6 +111,7 @@ def __init__( self.n_physical_experts = self.n_logical_experts + self.n_redundant_experts self.n_local_physical_experts = self.n_physical_experts // self.ep_size + self.shared_experts: ExaoneMoeGatedMLP | None if getattr(config, "num_shared_experts", 0) > 0: intermediate_size = config.moe_intermediate_size * config.num_shared_experts self.shared_experts = ExaoneMoeGatedMLP( @@ -292,7 +292,7 @@ def __init__( config: PretrainedConfig, cache_config: CacheConfig | None = None, quant_config: QuantizationConfig | None = None, - is_mtp: bool = None, + is_mtp: bool = False, prefix: str = "", ) -> None: super().__init__() @@ -381,26 +381,19 @@ def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""): config = vllm_config.model_config.hf_config cache_config = vllm_config.cache_config quant_config = vllm_config.quant_config - lora_config = vllm_config.lora_config self.num_redundant_experts = ( vllm_config.parallel_config.eplb_config.num_redundant_experts ) self.config = config self.quant_config = quant_config - lora_vocab = ( - (lora_config.lora_extra_vocab_size * (lora_config.max_loras or 1)) - if lora_config - else 0 - ) - self.vocab_size = config.vocab_size + lora_vocab + self.vocab_size = config.vocab_size if get_pp_group().is_first_rank or ( config.tie_word_embeddings and get_pp_group().is_last_rank ): self.embed_tokens = VocabParallelEmbedding( self.vocab_size, config.hidden_size, - org_num_embeddings=config.vocab_size, quant_config=quant_config, ) else: @@ -497,10 +490,8 @@ def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""): super().__init__() config = vllm_config.model_config.hf_config.get_text_config() quant_config = vllm_config.quant_config - lora_config = vllm_config.lora_config self.config = config - self.lora_config = lora_config self.quant_config = quant_config self.model = ExaoneMoeModel( @@ -509,17 +500,9 @@ def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""): ) if get_pp_group().is_last_rank: self.unpadded_vocab_size = config.vocab_size - if lora_config: - self.unpadded_vocab_size += lora_config.lora_extra_vocab_size self.lm_head = ParallelLMHead( self.unpadded_vocab_size, config.hidden_size, - org_num_embeddings=config.vocab_size, - padding_size=DEFAULT_VOCAB_PADDING_SIZE - # We need bigger padding if using lora for kernel - # compatibility - if not lora_config - else lora_config.lora_vocab_padding_size, quant_config=quant_config, ) if config.tie_word_embeddings: diff --git a/vllm/model_executor/models/exaone_moe_mtp.py b/vllm/model_executor/models/exaone_moe_mtp.py index ebe1f8cbaeb6..108fa37df8d1 100644 --- a/vllm/model_executor/models/exaone_moe_mtp.py +++ b/vllm/model_executor/models/exaone_moe_mtp.py @@ -46,16 +46,10 @@ def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""): model_config = vllm_config.model_config quant_config = vllm_config.quant_config - lora_config = vllm_config.lora_config config = model_config.hf_config self.config = config - lora_vocab = ( - (lora_config.lora_extra_vocab_size * (lora_config.max_loras or 1)) - if lora_config - else 0 - ) - self.vocab_size = config.vocab_size + lora_vocab + self.vocab_size = config.vocab_size self.org_vocab_size = config.vocab_size self.mtp_start_layer_idx = config.num_hidden_layers @@ -64,7 +58,6 @@ def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""): self.embed_tokens = VocabParallelEmbedding( self.vocab_size, config.hidden_size, - org_num_embeddings=config.vocab_size, ) self.fc = ColumnParallelLinear( diff --git a/vllm/model_executor/models/extract_hidden_states.py b/vllm/model_executor/models/extract_hidden_states.py index d94db2c692da..c11c64fcba5a 100644 --- a/vllm/model_executor/models/extract_hidden_states.py +++ b/vllm/model_executor/models/extract_hidden_states.py @@ -231,7 +231,7 @@ def __init__( head_size: int, cache_config: CacheConfig | None = None, prefix: str = "", - attn_type: str = AttentionType.DECODER, + attn_type: AttentionType = AttentionType.DECODER, ): super().__init__() @@ -329,7 +329,9 @@ def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""): super().__init__() self.vllm_config = vllm_config - self.hf_config = vllm_config.speculative_config.draft_model_config.hf_config + speculative_config = vllm_config.speculative_config + assert speculative_config is not None + self.hf_config = speculative_config.draft_model_config.hf_config self.hidden_size = vllm_config.model_config.get_hidden_size() self.target_num_hidden_layers = ( vllm_config.model_config.get_total_num_hidden_layers() diff --git a/vllm/model_executor/models/falcon.py b/vllm/model_executor/models/falcon.py index 97e133a03c08..f2eb776b14e6 100644 --- a/vllm/model_executor/models/falcon.py +++ b/vllm/model_executor/models/falcon.py @@ -413,6 +413,7 @@ def forward( else: hidden_states = self.embed_input_ids(input_ids) else: + assert intermediate_tensors is not None hidden_states = intermediate_tensors["hidden_states"] for layer in islice(self.h, self.start_layer, self.end_layer): hidden_states = layer(positions, hidden_states) diff --git a/vllm/model_executor/models/fireredasr2.py b/vllm/model_executor/models/fireredasr2.py index f652eb7518db..b9d99474597a 100644 --- a/vllm/model_executor/models/fireredasr2.py +++ b/vllm/model_executor/models/fireredasr2.py @@ -12,7 +12,7 @@ ) from vllm.config import ModelConfig, SpeechToTextConfig, VllmConfig -from vllm.config.multimodal import BaseDummyOptions +from vllm.config.multimodal import AudioDummyOptions, BaseDummyOptions from vllm.config.speech_to_text import SpeechToTextParams from vllm.inputs import MultiModalDataDict, PromptType from vllm.logger import init_logger @@ -216,6 +216,7 @@ def get_dummy_mm_data( num_audios = mm_counts.get("audio", 0) audio_overrides = mm_options.get("audio") + assert audio_overrides is None or isinstance(audio_overrides, AudioDummyOptions) ret = { "audio": self._get_dummy_audios( diff --git a/vllm/model_executor/models/fireredlid.py b/vllm/model_executor/models/fireredlid.py index 2225d70f9e18..7376556b3fa6 100644 --- a/vllm/model_executor/models/fireredlid.py +++ b/vllm/model_executor/models/fireredlid.py @@ -18,14 +18,13 @@ from collections.abc import Iterable, Mapping, Sequence from typing import Annotated, Literal -import numpy as np import torch from torch import nn from transformers import BatchFeature from vllm.config import ModelConfig, VllmConfig -from vllm.config.multimodal import BaseDummyOptions -from vllm.config.speech_to_text import SpeechToTextConfig +from vllm.config.multimodal import AudioDummyOptions, BaseDummyOptions +from vllm.config.speech_to_text import SpeechToTextConfig, SpeechToTextParams from vllm.distributed import get_tensor_model_parallel_world_size from vllm.inputs import MultiModalDataDict, PromptType from vllm.logger import init_logger @@ -431,6 +430,7 @@ def get_dummy_mm_data( audio_len = feature_extractor.chunk_length * sampling_rate num_audios = mm_counts.get("audio", 0) audio_overrides = mm_options.get("audio") + assert audio_overrides is None or isinstance(audio_overrides, AudioDummyOptions) return { "audio": self._get_dummy_audios( length=audio_len, @@ -747,13 +747,7 @@ def validate_language(cls, language: str | None) -> str | None: @classmethod def get_generation_prompt( cls, - audio: np.ndarray, - stt_config: SpeechToTextConfig, - model_config: ModelConfig, - language: str | None, - task_type: Literal["transcribe", "translate"], - request_prompt: str, - to_language: str | None, + stt_params: SpeechToTextParams, ) -> PromptType: """Build the prompt for the FireRedLID encoder-decoder model. @@ -764,7 +758,10 @@ def get_generation_prompt( "encoder_prompt": { "prompt": "", "multi_modal_data": { - "audio": (audio, int(stt_config.sample_rate)), + "audio": ( + stt_params.audio, + int(stt_params.stt_config.sample_rate), + ), }, }, "decoder_prompt": { diff --git a/vllm/model_executor/models/funasr.py b/vllm/model_executor/models/funasr.py index e15cde2d7519..b7981f42d364 100644 --- a/vllm/model_executor/models/funasr.py +++ b/vllm/model_executor/models/funasr.py @@ -14,7 +14,7 @@ ) from vllm.config import ModelConfig, SpeechToTextConfig, VllmConfig -from vllm.config.multimodal import BaseDummyOptions +from vllm.config.multimodal import AudioDummyOptions, BaseDummyOptions from vllm.config.speech_to_text import SpeechToTextParams from vllm.distributed import get_tensor_model_parallel_world_size from vllm.inputs import MultiModalDataDict, PromptType @@ -261,8 +261,8 @@ def __init__(self, d_model=80): def encode( self, - positions: torch.Tensor = None, - depth: int = None, + positions: torch.Tensor, + depth: int, dtype: torch.dtype = torch.float32, ): batch_size = positions.size(0) @@ -573,8 +573,8 @@ def __init__( ] ) - def forward(self, hidden_states: torch.Tensor, ilens: int = 0): - max_len = max(ilens) + def forward(self, hidden_states: torch.Tensor, ilens: torch.Tensor): + max_len = ilens.max() hidden_states = hidden_states[:, :max_len, :] batch_size, seq_len, dim = hidden_states.size() chunk_num = (seq_len - 1) // self.k + 1 @@ -682,6 +682,8 @@ def get_encoder_outputs( self.feat_permute = False if self.feat_permute: + if not isinstance(speech, torch.Tensor): + raise NotImplementedError("List audio inputs cannot be permuted") encoder_out, encoder_out_lens = self.encoder.audio_encoder( speech.permute(0, 2, 1), speech_lengths ) @@ -739,6 +741,7 @@ def get_dummy_mm_data( num_audios = mm_counts.get("audio", 0) audio_overrides = mm_options.get("audio") + assert audio_overrides is None or isinstance(audio_overrides, AudioDummyOptions) return { "audio": self._get_dummy_audios( diff --git a/vllm/model_executor/models/funaudiochat.py b/vllm/model_executor/models/funaudiochat.py index 8a0cf1364c6a..4b8ba82d70e6 100644 --- a/vllm/model_executor/models/funaudiochat.py +++ b/vllm/model_executor/models/funaudiochat.py @@ -26,7 +26,7 @@ from transformers.modeling_outputs import BaseModelOutput from vllm.config import VllmConfig -from vllm.config.multimodal import BaseDummyOptions +from vllm.config.multimodal import AudioDummyOptions, BaseDummyOptions from vllm.inputs import MultiModalDataDict from vllm.model_executor.layers.attention.mm_encoder_attention import MMEncoderAttention from vllm.model_executor.layers.linear import QKVParallelLinear, RowParallelLinear @@ -600,6 +600,7 @@ def get_dummy_mm_data( num_audios = int(mm_counts.get("audio", 0)) audio_overrides = mm_options.get("audio") + assert audio_overrides is None or isinstance(audio_overrides, AudioDummyOptions) return { "audio": self._get_dummy_audios( length=audio_len, @@ -635,6 +636,7 @@ def _apply_hf_processor_main( processed_data = BatchFeature({"input_ids": input_ids}) processed_data.update(passthrough_data) return processed_data + assert isinstance(audios, Sequence) feature_extractor = self.info.get_feature_extractor(**hf_processor_mm_kwargs) sr = int(feature_extractor.sampling_rate) From a014e35f38c80fb0652387740193ad2147fed6a3 Mon Sep 17 00:00:00 2001 From: Taneem Ibrahim Date: Sat, 22 Aug 2026 14:45:09 -0500 Subject: [PATCH 015/235] [Pooling] Fix Rust pooling endpoint for benchmark (#53352) Signed-off-by: Taneem Ibrahim --- rust/src/bench/README.md | 2 +- rust/src/bench/src/backends/pooling.rs | 2 +- rust/src/bench/src/cli.rs | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/rust/src/bench/README.md b/rust/src/bench/README.md index ff5b6ebfd55f..d4e0c69b80b0 100644 --- a/rust/src/bench/README.md +++ b/rust/src/bench/README.md @@ -419,7 +419,7 @@ vllm-bench \ | --------- | ------------- | ------------- | | `openai-embeddings` | `/v1/embeddings` | Text embedding (accepts text or token IDs) | | `openai-embeddings-chat` | `/v1/embeddings` | Chat-format embedding (supports multimodal content) | -| `vllm-pooling` | `/v1/pooling` | vLLM native pooling endpoint | +| `vllm-pooling` | `/pooling` | vLLM native pooling endpoint | | `vllm-rerank` | `/v1/rerank` | vLLM reranking (query from prompt, documents via `--extra-body`) | Pooling backends are non-streaming and report E2EL (end-to-end latency) only. Use `--dataset-name sharegpt`, `sonnet`, or `hf` for text-based embedding/rerank benchmarks, or `random` for token-ID-based embedding benchmarks. diff --git a/rust/src/bench/src/backends/pooling.rs b/rust/src/bench/src/backends/pooling.rs index ad43f480e21d..09522b4d67fb 100644 --- a/rust/src/bench/src/backends/pooling.rs +++ b/rust/src/bench/src/backends/pooling.rs @@ -8,7 +8,7 @@ //! - `openai-embeddings`: Standard OpenAI `/v1/embeddings` with text input //! - `openai-embeddings-chat`: OpenAI `/v1/embeddings` with chat message format (supports //! multimodal) -//! - `vllm-pooling`: vLLM `/v1/pooling` endpoint +//! - `vllm-pooling`: vLLM `/pooling` endpoint //! - `vllm-rerank`: vLLM `/v1/rerank` endpoint (query + documents) use std::time::Instant; diff --git a/rust/src/bench/src/cli.rs b/rust/src/bench/src/cli.rs index 24a7b3e2512a..fc0f5a29e483 100644 --- a/rust/src/bench/src/cli.rs +++ b/rust/src/bench/src/cli.rs @@ -682,7 +682,7 @@ impl BenchServeArgs { BackendKind::OpenaiEmbeddings | BackendKind::OpenaiEmbeddingsChat => { "/v1/embeddings".to_string() } - BackendKind::VllmPooling => "/v1/pooling".to_string(), + BackendKind::VllmPooling => "/pooling".to_string(), BackendKind::VllmRerank => "/v1/rerank".to_string(), } } From 2f55ef254c70e110d637beeedf48238977ebb683 Mon Sep 17 00:00:00 2001 From: Zhou248 <62542847+Zhou248@users.noreply.github.com> Date: Sun, 23 Aug 2026 06:35:22 +0800 Subject: [PATCH 016/235] [Model] Add Qwen3-Omni DSpark support (#52560) Signed-off-by: Zhou248 <630563665@qq.com> --- tests/model_executor/test_qwen3_omni.py | 172 +++++++++- tests/models/registry.py | 7 + tests/test_config.py | 121 +++++++ .../test_speculators_dspark_config.py | 58 ++++ vllm/config/speculative.py | 304 +++++++++++++++++- vllm/model_executor/models/qwen3_dflash.py | 16 +- vllm/model_executor/models/qwen3_dspark.py | 22 +- .../models/qwen3_omni_moe_thinker.py | 12 +- vllm/model_executor/models/registry.py | 1 + .../configs/speculators/algos.py | 21 +- .../v1/worker/gpu/spec_decode/dspark/utils.py | 19 +- 11 files changed, 731 insertions(+), 22 deletions(-) create mode 100644 tests/transformers_utils/test_speculators_dspark_config.py diff --git a/tests/model_executor/test_qwen3_omni.py b/tests/model_executor/test_qwen3_omni.py index 1ceaba04aa02..54d2f2672f7b 100644 --- a/tests/model_executor/test_qwen3_omni.py +++ b/tests/model_executor/test_qwen3_omni.py @@ -1,9 +1,13 @@ # SPDX-License-Identifier: Apache-2.0 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project -from unittest.mock import Mock +from contextlib import nullcontext +from types import SimpleNamespace +from unittest.mock import Mock, patch import pytest +import torch +import torch.nn as nn from transformers import PretrainedConfig from vllm.multimodal.processing import InputProcessingContext @@ -218,5 +222,171 @@ def test_qwen3_omni_get_updates_use_audio_in_video( ) +@pytest.mark.skip_global_cleanup +def test_qwen3_omni_exposes_eagle3_to_its_text_backbone(): + from vllm.model_executor.models.interfaces import EagleModelMixin, supports_eagle3 + from vllm.model_executor.models.qwen3_omni_moe_thinker import ( + Qwen3OmniMoeThinkerForConditionalGeneration, + ) + + class DummyBackbone(nn.Module, EagleModelMixin): + def __init__(self): + super().__init__() + self.layers = nn.ModuleList([nn.Identity(), nn.Identity()]) + + class DummyLanguageModel(nn.Module): + def __init__(self): + super().__init__() + self.model = DummyBackbone() + + def embed_input_ids(self, input_ids): + return input_ids + + model = Qwen3OmniMoeThinkerForConditionalGeneration.__new__( + Qwen3OmniMoeThinkerForConditionalGeneration + ) + nn.Module.__init__(model) + model.language_model = DummyLanguageModel() + + assert supports_eagle3(model) + model.set_aux_hidden_state_layers((1, 2)) + assert model.language_model.model.aux_hidden_state_layers == (1, 2) + + +@pytest.mark.skip_global_cleanup +def test_qwen3_omni_text_model_collects_post_deepstack_aux_hidden_states(): + from vllm.model_executor.models.qwen3_omni_moe_thinker import Qwen3MoeLLMModel + + class DummyLayer(nn.Module): + def forward(self, positions, hidden_states, residual): + return hidden_states + 1, torch.full_like(hidden_states, 10) + + class DummyNorm(nn.Module): + def forward(self, hidden_states, residual): + return hidden_states + residual, None + + model = Qwen3MoeLLMModel.__new__(Qwen3MoeLLMModel) + nn.Module.__init__(model) + model.start_layer = 0 + model.end_layer = 1 + model.layers = nn.ModuleList([DummyLayer()]) + model.norm = DummyNorm() + model.aux_hidden_state_layers = (1,) + + pp_group = Mock(is_first_rank=True, is_last_rank=True) + inputs_embeds = torch.tensor([[1.0]]) + deepstack_inputs = {"deepstack_input_embeds_0": torch.tensor([[3.0]])} + with patch( + "vllm.model_executor.models.qwen3_omni_moe_thinker.get_pp_group", + return_value=pp_group, + ): + output, aux_hidden_states = model.forward( + input_ids=None, + positions=torch.tensor([0]), + inputs_embeds=inputs_embeds, + deepstack_input_embeds=deepstack_inputs, + ) + + torch.testing.assert_close(output, torch.tensor([[15.0]])) + assert len(aux_hidden_states) == 1 + torch.testing.assert_close(aux_hidden_states[0], torch.tensor([[15.0]])) + + +@pytest.mark.skip_global_cleanup +@pytest.mark.parametrize( + ("input_vocab_size", "draft_vocab_size", "weights", "error"), + [ + (101, 100, [], "must include embed_tokens weights"), + (99, 99, [], "must include lm_head weights"), + (100, 40, [], "must include lm_head weights"), + ( + 100, + 40, + [("lm_head.weight", torch.empty(40, 8))], + "must include a d2t mapping", + ), + ], +) +def test_qwen3_dspark_rejects_incomplete_vocab_weights( + input_vocab_size, draft_vocab_size, weights, error +): + from vllm.model_executor.models.qwen3_dspark import Qwen3DSparkForCausalLM + + model = Qwen3DSparkForCausalLM.__new__(Qwen3DSparkForCausalLM) + nn.Module.__init__(model) + object.__setattr__( + model, + "config", + SimpleNamespace( + vocab_size=input_vocab_size, + draft_vocab_size=draft_vocab_size, + ), + ) + object.__setattr__(model, "target_vocab_size", 100) + + with pytest.raises(ValueError, match=error): + model.load_weights(weights) + + +@pytest.mark.skip_global_cleanup +def test_dspark_shares_target_embedding_with_smaller_draft_vocabulary(): + from vllm.v1.worker.gpu.spec_decode.dspark import utils as dspark_utils + + target_embedding = nn.Embedding(100, 8) + draft_embedding = nn.Embedding(99, 8) + target_model = SimpleNamespace(model=SimpleNamespace(embed_tokens=target_embedding)) + draft_model = SimpleNamespace( + model=SimpleNamespace(embed_tokens=draft_embedding), + has_own_embed_tokens=False, + ) + draft_model_config = SimpleNamespace( + hf_config=SimpleNamespace(model_type="qwen3"), + get_vocab_size=Mock(return_value=99), + ) + vllm_config = SimpleNamespace( + speculative_config=SimpleNamespace( + draft_model_config=draft_model_config, + attention_backend=None, + kv_cache_dtype=None, + ), + attention_config=SimpleNamespace(backend=None), + cache_config=SimpleNamespace(), + model_config=SimpleNamespace(get_vocab_size=Mock(return_value=100)), + ) + + def fake_replace(config, **changes): + values = vars(config).copy() + values.update(changes) + return SimpleNamespace(**values) + + with ( + patch.object(dspark_utils, "replace", side_effect=fake_replace), + patch.object( + dspark_utils, + "get_pp_group", + return_value=SimpleNamespace(world_size=1), + ), + patch( + "vllm.compilation.backends.set_model_tag", + return_value=nullcontext(), + ), + patch( + "vllm.model_executor.model_loader.get_model", + return_value=draft_model, + ), + patch( + "vllm.model_executor.models.qwen3_dflash.dflash_has_any_non_causal", + return_value=False, + ), + patch( + "vllm.model_executor.models.utils.get_draft_quant_config", + return_value=None, + ), + ): + loaded_model = dspark_utils.load_dspark_model(target_model, vllm_config) + + assert loaded_model.model.embed_tokens is target_embedding + + if __name__ == "__main__": pytest.main([__file__, "-v"]) diff --git a/tests/models/registry.py b/tests/models/registry.py index 24a48693998f..712f160e3db2 100644 --- a/tests/models/registry.py +++ b/tests/models/registry.py @@ -1557,6 +1557,13 @@ def check_available_online( is_available_online=False, use_original_num_layers=True, # DSpark backbone requires all layers ), + "Qwen3OmniDSparkModel": _HfExamplesInfo( + "Qwen/Qwen3-Omni-30B-A3B-Instruct", + speculative_model="Qwen/Qwen3-Omni-30B-A3B-DSpark", + is_available_online=False, + use_original_num_layers=True, + min_transformers_version="4.57", + ), "Gemma4DSparkModel": _HfExamplesInfo( "google/gemma-4-12B-it", speculative_model="deepseek-ai/dspark_gemma4_12b_block7", diff --git a/tests/test_config.py b/tests/test_config.py index eac6a951957a..40ab2fa426ee 100644 --- a/tests/test_config.py +++ b/tests/test_config.py @@ -32,6 +32,7 @@ from vllm.config.kernel import IrOpPriorityConfig from vllm.config.load import LoadConfig from vllm.config.mamba import MambaBackendEnum +from vllm.config.speculative import _validate_qwen3_omni_dspark from vllm.config.utils import get_field from vllm.config.vllm import OPTIMIZATION_LEVEL_TO_CONFIG, OptimizationLevel from vllm.platforms import current_platform @@ -2100,6 +2101,126 @@ def test_mtp_draft_uses_model_weights_not_local_cache(mock_model_config_cls): assert actual_model == s3_url +def _make_qwen3_omni_dspark_configs(): + text_config = SimpleNamespace( + hidden_size=2048, + num_hidden_layers=48, + num_attention_heads=32, + num_key_value_heads=4, + head_dim=128, + vocab_size=152064, + ) + target_model_config = SimpleNamespace( + hf_config=SimpleNamespace( + model_type="qwen3_omni_moe", + architectures=["Qwen3OmniMoeForConditionalGeneration"], + thinker_config=SimpleNamespace(text_config=text_config), + ), + hf_text_config=text_config, + architectures=["Qwen3OmniMoeForConditionalGeneration"], + get_hidden_size=lambda: 2048, + get_total_num_hidden_layers=lambda: 48, + get_vocab_size=lambda: 152064, + ) + draft_hf_config = SimpleNamespace( + model_type="qwen3", + architectures=["Qwen3OmniDSparkModel"], + hidden_size=2048, + num_attention_heads=32, + num_key_value_heads=4, + head_dim=128, + block_size=7, + target_hidden_size=2048, + target_layer_ids=[1, 9, 17, 25, 33], + use_aux_hidden_state=True, + markov_rank=64, + markov_head_type="vanilla", + sample_from_anchor=True, + dspark_bonus_anchor=False, + vocab_size=152064, + draft_vocab_size=32000, + mask_token_id=151669, + rope_parameters={"rope_type": "default", "rope_theta": 1000000.0}, + ) + draft_model_config = SimpleNamespace( + hf_config=draft_hf_config, + architectures=["Qwen3OmniDSparkModel"], + ) + return target_model_config, draft_model_config + + +@pytest.mark.skip_global_cleanup +def test_qwen3_omni_dspark_checkpoint_contract_is_accepted(): + target_config, draft_config = _make_qwen3_omni_dspark_configs() + _validate_qwen3_omni_dspark(target_config, draft_config, 7) + + +@pytest.mark.skip_global_cleanup +def test_qwen3_omni_dspark_rejects_generic_qwen3_architecture(): + target_config, draft_config = _make_qwen3_omni_dspark_configs() + draft_config.hf_config.architectures = ["Qwen3DSparkModel"] + draft_config.architectures = ["Qwen3DSparkModel"] + + with pytest.raises(ValueError, match="must be converted first"): + _validate_qwen3_omni_dspark(target_config, draft_config, 7) + + +@pytest.mark.parametrize( + ("field", "value", "error"), + [ + ("block_size", 5, "trained block_size"), + ("target_hidden_size", 4096, "target_hidden_size"), + ("hidden_size", 4096, "draft hidden_size"), + ("num_attention_heads", 16, "num_attention_heads"), + ("num_key_value_heads", 8, "num_key_value_heads"), + ("head_dim", 64, "head_dim"), + ("target_layer_ids", [7, 48], "zero-based text-layer"), + ("target_layer_ids", [23, 7], "strictly increasing"), + ("use_aux_hidden_state", False, "use_aux_hidden_state=true"), + ("markov_rank", 0, "markov_rank"), + ("markov_head_type", "gated", "markov_head_type='vanilla'"), + ("sample_from_anchor", False, "sample_from_anchor=true"), + ("dspark_bonus_anchor", True, "dspark_bonus_anchor=false"), + ("vocab_size", 0, "input vocab_size must be a positive integer"), + ], +) +@pytest.mark.skip_global_cleanup +def test_qwen3_omni_dspark_rejects_incompatible_checkpoint_fields(field, value, error): + target_config, draft_config = _make_qwen3_omni_dspark_configs() + setattr(draft_config.hf_config, field, value) + + with pytest.raises(ValueError, match=error): + _validate_qwen3_omni_dspark(target_config, draft_config, 7) + + +@pytest.mark.skip_global_cleanup +def test_qwen3_omni_dspark_rejects_mrope_draft_positions(): + target_config, draft_config = _make_qwen3_omni_dspark_configs() + draft_config.hf_config.rope_parameters["mrope_section"] = [24, 20, 20] + + with pytest.raises(ValueError, match="logical 1-D RoPE"): + _validate_qwen3_omni_dspark(target_config, draft_config, 7) + + +@pytest.mark.skip_global_cleanup +def test_qwen3_omni_dspark_allows_draft_only_noise_token_row(): + target_config, draft_config = _make_qwen3_omni_dspark_configs() + draft_config.hf_config.vocab_size = 152065 + draft_config.hf_config.draft_vocab_size = 152064 + draft_config.hf_config.mask_token_id = 152064 + + _validate_qwen3_omni_dspark(target_config, draft_config, 7) + + +@pytest.mark.skip_global_cleanup +def test_qwen3_omni_dspark_allows_smaller_input_vocabulary(): + target_config, draft_config = _make_qwen3_omni_dspark_configs() + draft_config.hf_config.vocab_size = 151936 + draft_config.hf_config.mask_token_id = 151669 + + _validate_qwen3_omni_dspark(target_config, draft_config, 7) + + def test_ir_op_priority_default(): """Test that IR op priority defaults are set correctly.""" from vllm.config.kernel import IrOpPriorityConfig diff --git a/tests/transformers_utils/test_speculators_dspark_config.py b/tests/transformers_utils/test_speculators_dspark_config.py new file mode 100644 index 000000000000..ce32b6e5dfad --- /dev/null +++ b/tests/transformers_utils/test_speculators_dspark_config.py @@ -0,0 +1,58 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project + +import pytest + +from vllm.transformers_utils.configs.speculators.base import SpeculatorsConfig + +pytestmark = pytest.mark.skip_global_cleanup + + +def _make_dspark_config(architecture: str) -> dict: + return { + "speculators_model_type": "dspark", + "architectures": [architecture], + "sample_from_anchor": True, + "use_aux_hidden_state": True, + "draft_vocab_size": 32, + "target_hidden_size": 16, + "mask_token_id": 31, + "markov_rank": 4, + "markov_head_type": "vanilla", + "block_size": 3, + "aux_hidden_state_layer_ids": [2, 5], + "transformer_layer_config": { + "model_type": "qwen3", + "hidden_size": 16, + }, + } + + +def test_dspark_updater_preserves_qwen3_omni_contract() -> None: + config = SpeculatorsConfig.extract_transformers_pre_trained_config( + _make_dspark_config("Qwen3OmniDSparkModel") + ) + + assert config["architectures"] == ["Qwen3OmniDSparkModel"] + assert config["sample_from_anchor"] is True + assert config["dspark_bonus_anchor"] is False + assert config["use_aux_hidden_state"] is True + assert config["target_layer_ids"] == [1, 4] + + +def test_dspark_updater_keeps_legacy_qwen3_checkpoint_loadable() -> None: + config = SpeculatorsConfig.extract_transformers_pre_trained_config( + _make_dspark_config("DSparkSpeculator") + ) + + assert config["architectures"] == ["Qwen3DSparkModel"] + + +def test_dspark_updater_maps_bonus_anchor_semantics() -> None: + outer_config = _make_dspark_config("Qwen3DSparkModel") + outer_config["sample_from_anchor"] = False + + config = SpeculatorsConfig.extract_transformers_pre_trained_config(outer_config) + + assert config["sample_from_anchor"] is False + assert config["dspark_bonus_anchor"] is True diff --git a/vllm/config/speculative.py b/vllm/config/speculative.py index f0eee88429b6..d3b4af1cf060 100644 --- a/vllm/config/speculative.py +++ b/vllm/config/speculative.py @@ -3,7 +3,7 @@ import copy import functools -from collections.abc import Callable +from collections.abc import Callable, Mapping from typing import TYPE_CHECKING, Any, Literal, get_args from pydantic import Field, SkipValidation, field_validator, model_validator @@ -80,6 +80,293 @@ RejectionSampleMethod = Literal["standard", "synthetic", "block"] DraftSampleMethod = Literal["greedy", "probabilistic"] +_QWEN3_OMNI_TARGET_ARCHITECTURES = frozenset( + { + "Qwen3OmniMoeForConditionalGeneration", + "Qwen3OmniMoeThinkerForConditionalGeneration", + } +) +_QWEN3_OMNI_DSPARK_ARCHITECTURE = "Qwen3OmniDSparkModel" + + +def _is_qwen3_omni_target(model_config: ModelConfig) -> bool: + hf_config = model_config.hf_config + architectures = set(getattr(model_config, "architectures", ()) or ()) + architectures.update(getattr(hf_config, "architectures", ()) or ()) + return getattr(hf_config, "model_type", None) == "qwen3_omni_moe" or bool( + architectures & _QWEN3_OMNI_TARGET_ARCHITECTURES + ) + + +def _get_qwen3_omni_text_config(model_config: ModelConfig) -> Any: + thinker_config = getattr(model_config.hf_config, "thinker_config", None) + text_config = getattr(thinker_config, "text_config", None) + if text_config is None: + text_config = getattr(model_config, "hf_text_config", None) + return text_config + + +def _get_attention_head_dim(config: Any) -> int | None: + head_dim = getattr(config, "head_dim", None) + if head_dim is not None: + return head_dim + hidden_size = getattr(config, "hidden_size", None) + num_attention_heads = getattr(config, "num_attention_heads", None) + if ( + isinstance(hidden_size, int) + and isinstance(num_attention_heads, int) + and num_attention_heads > 0 + and hidden_size % num_attention_heads == 0 + ): + return hidden_size // num_attention_heads + return None + + +def _get_nested_config_value(config: Any, section: str, name: str) -> Any: + nested_config = getattr(config, section, None) + if isinstance(nested_config, Mapping): + return nested_config.get(name) + return getattr(nested_config, name, None) + + +def _get_qwen3_dspark_value(config: Any, name: str) -> Any: + # Match DFlashQwen3Model's precedence: dflash_config overrides + # eagle_config, and top-level values are the modern fallback. + value = _get_nested_config_value(config, "dflash_config", name) + if value is None: + value = _get_nested_config_value(config, "eagle_config", name) + if value is None: + value = getattr(config, name, None) + return value + + +def _validate_qwen3_omni_dspark( + target_model_config: ModelConfig, + draft_model_config: ModelConfig, + num_speculative_tokens: int, +) -> None: + """Validate the standalone Qwen3-Omni DSpark checkpoint contract.""" + if not _is_qwen3_omni_target(target_model_config): + return + + draft_hf_config = draft_model_config.hf_config + draft_architectures = set(getattr(draft_model_config, "architectures", ()) or ()) + draft_architectures.update(getattr(draft_hf_config, "architectures", ()) or ()) + if _QWEN3_OMNI_DSPARK_ARCHITECTURE not in draft_architectures: + raise ValueError( + "Qwen3-Omni DSpark requires a standalone draft checkpoint with " + f"architectures=['{_QWEN3_OMNI_DSPARK_ARCHITECTURE}']; generic " + "Qwen3DSparkModel checkpoints must be converted first." + ) + + block_size = _get_qwen3_dspark_value(draft_hf_config, "block_size") + if block_size is None: + block_size = _get_qwen3_dspark_value(draft_hf_config, "dspark_block_size") + if ( + not isinstance(block_size, int) + or isinstance(block_size, bool) + or block_size <= 0 + ): + raise ValueError( + "Qwen3-Omni DSpark requires a positive integer block_size in the " + "draft config." + ) + if num_speculative_tokens != block_size: + raise ValueError( + "Qwen3-Omni DSpark requires num_speculative_tokens to match the " + f"trained block_size ({block_size}); got {num_speculative_tokens}." + ) + + target_text_config = _get_qwen3_omni_text_config(target_model_config) + if target_text_config is None: + raise ValueError( + "Qwen3-Omni DSpark could not resolve the target thinker text_config." + ) + target_hidden_size = getattr( + target_text_config, + "hidden_size", + target_model_config.get_hidden_size(), + ) + draft_target_hidden_size = getattr(draft_hf_config, "target_hidden_size", None) + if draft_target_hidden_size != target_hidden_size: + raise ValueError( + "Qwen3-Omni DSpark draft target_hidden_size must match the target " + f"text hidden size ({target_hidden_size}); got {draft_target_hidden_size}." + ) + draft_hidden_size = getattr(draft_hf_config, "hidden_size", None) + if draft_hidden_size != target_hidden_size: + raise ValueError( + "Qwen3-Omni DSpark draft hidden_size must match the target thinker " + f"hidden size ({target_hidden_size}); got {draft_hidden_size}." + ) + + target_attention = { + "num_attention_heads": getattr(target_text_config, "num_attention_heads", None), + "num_key_value_heads": getattr( + target_text_config, + "num_key_value_heads", + getattr(target_text_config, "num_attention_heads", None), + ), + "head_dim": _get_attention_head_dim(target_text_config), + } + draft_attention = { + "num_attention_heads": getattr(draft_hf_config, "num_attention_heads", None), + "num_key_value_heads": getattr( + draft_hf_config, + "num_key_value_heads", + getattr(draft_hf_config, "num_attention_heads", None), + ), + "head_dim": _get_attention_head_dim(draft_hf_config), + } + for name, target_value in target_attention.items(): + if ( + not isinstance(target_value, int) + or isinstance(target_value, bool) + or target_value <= 0 + ): + raise ValueError( + "Qwen3-Omni DSpark requires a positive integer " + f"{name} in the target thinker text_config; got {target_value}." + ) + draft_value = draft_attention[name] + if draft_value != target_value: + raise ValueError( + f"Qwen3-Omni DSpark draft {name} must match the target thinker " + f"value ({target_value}); got {draft_value}." + ) + + target_layer_ids = _get_nested_config_value( + draft_hf_config, "dflash_config", "target_layer_ids" + ) + if target_layer_ids is None: + target_layer_ids = getattr(draft_hf_config, "dspark_target_layer_ids", None) + if target_layer_ids is None: + target_layer_ids = getattr(draft_hf_config, "target_layer_ids", None) + if not isinstance(target_layer_ids, (list, tuple)) or not target_layer_ids: + raise ValueError( + "Qwen3-Omni DSpark requires a non-empty target_layer_ids list in " + "the draft config." + ) + if any( + not isinstance(layer_id, int) or isinstance(layer_id, bool) + for layer_id in target_layer_ids + ): + raise ValueError( + "Qwen3-Omni DSpark target_layer_ids must contain only integers." + ) + if list(target_layer_ids) != sorted(set(target_layer_ids)): + raise ValueError( + "Qwen3-Omni DSpark target_layer_ids must be unique and strictly increasing." + ) + target_num_layers = getattr( + target_text_config, + "num_hidden_layers", + target_model_config.get_total_num_hidden_layers(), + ) + if target_layer_ids[0] < 0 or target_layer_ids[-1] >= target_num_layers: + raise ValueError( + "Qwen3-Omni DSpark target_layer_ids must be zero-based text-layer " + f"indices in [0, {target_num_layers - 1}]; got {target_layer_ids}." + ) + + if _get_qwen3_dspark_value(draft_hf_config, "use_aux_hidden_state") is not True: + raise ValueError( + "Qwen3-Omni DSpark requires use_aux_hidden_state=true because Omni " + "conditioning is supplied through thinker hidden states." + ) + + markov_rank = getattr(draft_hf_config, "markov_rank", None) + if ( + not isinstance(markov_rank, int) + or isinstance(markov_rank, bool) + or markov_rank <= 0 + ): + raise ValueError( + "Qwen3-Omni DSpark requires a positive integer markov_rank in the " + "draft config." + ) + markov_head_type = getattr(draft_hf_config, "markov_head_type", None) + if markov_head_type != "vanilla": + raise ValueError( + "Qwen3-Omni DSpark currently requires markov_head_type='vanilla'; " + f"got {markov_head_type!r}." + ) + + sample_from_anchor = getattr(draft_hf_config, "sample_from_anchor", True) + bonus_anchor = getattr(draft_hf_config, "dspark_bonus_anchor", False) + if sample_from_anchor is not True or bonus_anchor is not False: + raise ValueError( + "Qwen3-Omni DSpark requires sample_from_anchor=true and " + "dspark_bonus_anchor=false." + ) + + target_vocab_size = getattr( + target_text_config, + "vocab_size", + target_model_config.get_vocab_size(), + ) + draft_input_vocab_size = getattr(draft_hf_config, "vocab_size", None) + if ( + not isinstance(draft_input_vocab_size, int) + or isinstance(draft_input_vocab_size, bool) + or draft_input_vocab_size <= 0 + ): + raise ValueError( + "Qwen3-Omni DSpark input vocab_size must be a positive integer; " + f"got {draft_input_vocab_size}." + ) + draft_output_vocab_size = getattr(draft_hf_config, "draft_vocab_size", None) + if draft_output_vocab_size is None: + draft_output_vocab_size = draft_input_vocab_size + if ( + not isinstance(draft_output_vocab_size, int) + or isinstance(draft_output_vocab_size, bool) + or not 0 < draft_output_vocab_size <= target_vocab_size + ): + raise ValueError( + "Qwen3-Omni DSpark draft_vocab_size must be a positive integer no " + f"larger than target vocab_size ({target_vocab_size}); got " + f"{draft_output_vocab_size}." + ) + + noise_token_id = _get_nested_config_value( + draft_hf_config, "dflash_config", "mask_token_id" + ) + for name in ( + "mask_token_id", + "dspark_noise_token_id", + "pard_token", + "ptd_token_id", + ): + if noise_token_id is None: + noise_token_id = getattr(draft_hf_config, name, None) + if noise_token_id is not None: + break + if ( + not isinstance(noise_token_id, int) + or isinstance(noise_token_id, bool) + or not 0 <= noise_token_id < draft_input_vocab_size + ): + raise ValueError( + "Qwen3-Omni DSpark requires a valid mask/noise token id within the " + f"draft input vocabulary [0, {draft_input_vocab_size - 1}]." + ) + + rope_configs = ( + getattr(draft_hf_config, "rope_parameters", None), + getattr(draft_hf_config, "rope_scaling", None), + ) + has_mrope = getattr(draft_hf_config, "mrope_section", None) is not None + has_mrope = has_mrope or any( + isinstance(rope_config, Mapping) and "mrope_section" in rope_config + for rope_config in rope_configs + ) + if has_mrope: + raise ValueError( + "Qwen3-Omni DSpark draft checkpoints must use logical 1-D RoPE and " + "must not define mrope_section." + ) + @config class SpeculativeConfig: @@ -964,6 +1251,8 @@ def __post_init__(self): elif ( "dspark" in self.draft_model_config.model.lower() or "Qwen3DSparkModel" in self.draft_model_config.architectures + or _QWEN3_OMNI_DSPARK_ARCHITECTURE + in self.draft_model_config.architectures or "Gemma4DSparkModel" in self.draft_model_config.architectures or ( "DSparkDraftModel" in self.draft_model_config.architectures @@ -1037,6 +1326,8 @@ def __post_init__(self): self.update_arch_() elif self.method == "dspark" and ( "Qwen3DSparkModel" not in self.draft_model_config.architectures + and _QWEN3_OMNI_DSPARK_ARCHITECTURE + not in self.draft_model_config.architectures and "Gemma4DSparkModel" not in self.draft_model_config.architectures and "K3DSparkModel" not in self.draft_model_config.architectures ): @@ -1125,13 +1416,22 @@ def __post_init__(self): if ( "Qwen3DSparkModel" not in self.draft_model_config.architectures + and _QWEN3_OMNI_DSPARK_ARCHITECTURE + not in self.draft_model_config.architectures ): raise ValueError( "dspark_draft_topk is only supported by " - "Qwen3DSparkModel" + "Qwen3DSparkModel and Qwen3OmniDSparkModel" ) hf_config.dspark_draft_topk = dspark_draft_topk + assert self.target_model_config is not None + _validate_qwen3_omni_dspark( + self.target_model_config, + self.draft_model_config, + self.num_speculative_tokens, + ) + self.draft_tensor_parallel_size = ( SpeculativeConfig._verify_and_get_draft_tp( self.target_parallel_config, diff --git a/vllm/model_executor/models/qwen3_dflash.py b/vllm/model_executor/models/qwen3_dflash.py index 45f0bf53447d..b7b5b29d41d3 100644 --- a/vllm/model_executor/models/qwen3_dflash.py +++ b/vllm/model_executor/models/qwen3_dflash.py @@ -377,8 +377,6 @@ def forward( @support_torch_compile class DFlashQwen3Model(nn.Module): - decoder_layer_cls = DFlashQwen3DecoderLayer - hf_to_vllm_mapper = WeightsMapper( orig_to_new_substr={ "midlayer.": "layers.0.", @@ -412,10 +410,10 @@ def __init__( drafter_config = getattr(self.config, "eagle_config", {}) drafter_config.update(getattr(self.config, "dflash_config", {})) - if drafter_config is not None and "use_aux_hidden_state" in drafter_config: - self.use_aux_hidden_state = drafter_config["use_aux_hidden_state"] - else: - self.use_aux_hidden_state = True + self.use_aux_hidden_state = drafter_config.get( + "use_aux_hidden_state", + getattr(self.config, "use_aux_hidden_state", True), + ) current_vllm_config = get_current_vllm_config() @@ -430,7 +428,9 @@ def __init__( # at that slot id. Some checkpoints (XiaomiMiMo/MiMo-V2.5-Pro-FP4-DFlash) ship # with a separate mask embedding tensor to use instead. When present, we load it # and substitute it for embed_tokens[mask_token_id] when computing embeddings. - self.mask_token_id = drafter_config.get("mask_token_id") + self.mask_token_id = drafter_config.get( + "mask_token_id", getattr(self.config, "mask_token_id", None) + ) self.mask_embedding = nn.Parameter( torch.zeros(self.config.hidden_size, dtype=vllm_config.model_config.dtype), requires_grad=False, @@ -439,7 +439,7 @@ def __init__( self.layers = nn.ModuleList( [ - self.decoder_layer_cls( + DFlashQwen3DecoderLayer( current_vllm_config, config=self.config, layer_idx=layer_idx, diff --git a/vllm/model_executor/models/qwen3_dspark.py b/vllm/model_executor/models/qwen3_dspark.py index 250d54d9ca88..2cee016124c7 100644 --- a/vllm/model_executor/models/qwen3_dspark.py +++ b/vllm/model_executor/models/qwen3_dspark.py @@ -200,8 +200,8 @@ def __init__(self, *, vllm_config: VllmConfig, prefix: str = "") -> None: self.logits_processor = LogitsProcessor( self.config.draft_vocab_size, scale=logit_scale ) - target_vocab_size = vllm_config.model_config.get_vocab_size() - if self.config.draft_vocab_size != target_vocab_size: + self.target_vocab_size = vllm_config.model_config.get_vocab_size() + if self.config.draft_vocab_size != self.target_vocab_size: self.draft_id_to_target_id = nn.Parameter( torch.zeros(self.config.draft_vocab_size, dtype=torch.long), requires_grad=False, @@ -280,6 +280,24 @@ def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]): # mask_embedding is an unused placeholder param; DSpark masks via the vocab row. # embed_tokens / lm_head are optional; when omitted they are shared from # the target by load_dspark_model, so skip the unloaded params here. + uses_expanded_input_vocab = self.config.vocab_size > self.target_vocab_size + uses_reduced_vocab = self.config.draft_vocab_size != self.target_vocab_size + if uses_expanded_input_vocab and not includes_embed_tokens: + raise ValueError( + "Qwen3 DSpark checkpoints whose input vocab_size is larger than " + "the target vocabulary must include embed_tokens weights." + ) + if uses_reduced_vocab and not includes_lm_head: + raise ValueError( + "Reduced-vocabulary Qwen3 DSpark checkpoints must include " + "lm_head weights; the full target lm_head cannot be shared." + ) + if uses_reduced_vocab and not includes_draft_id_mapping: + raise ValueError( + "Reduced-vocabulary Qwen3 DSpark checkpoints must include a " + "d2t mapping so sampled draft ids can be converted to target ids." + ) + orig_to_new_substr = {"mask_embedding": None} if not includes_embed_tokens: orig_to_new_substr["embed_tokens"] = None diff --git a/vllm/model_executor/models/qwen3_omni_moe_thinker.py b/vllm/model_executor/models/qwen3_omni_moe_thinker.py index 0e2aab634cbc..2bcc07a27655 100755 --- a/vllm/model_executor/models/qwen3_omni_moe_thinker.py +++ b/vllm/model_executor/models/qwen3_omni_moe_thinker.py @@ -83,6 +83,7 @@ from .interfaces import ( MultiModalEmbeddings, + SupportsEagle3, SupportsMRoPE, SupportsMultiModal, SupportsPP, @@ -1067,7 +1068,7 @@ def forward( intermediate_tensors: IntermediateTensors | None = None, inputs_embeds: torch.Tensor | None = None, deepstack_input_embeds: IntermediateTensors | None = None, - ) -> torch.Tensor | IntermediateTensors: + ) -> torch.Tensor | IntermediateTensors | tuple[torch.Tensor, list[torch.Tensor]]: if get_pp_group().is_first_rank: if inputs_embeds is not None: hidden_states = inputs_embeds @@ -1078,6 +1079,9 @@ def forward( assert intermediate_tensors is not None hidden_states = intermediate_tensors["hidden_states"] residual = intermediate_tensors["residual"] + aux_hidden_states = self._maybe_add_hidden_state( + [], self.start_layer, hidden_states, residual + ) for layer_idx, layer in enumerate( self.layers[self.start_layer : self.end_layer] ): @@ -1096,12 +1100,17 @@ def forward( hidden_states + deepstack_input_embeds[f"deepstack_input_embeds_{layer_idx}"] ) + self._maybe_add_hidden_state( + aux_hidden_states, layer_idx + 1, hidden_states, residual + ) if not get_pp_group().is_last_rank: return IntermediateTensors( {"hidden_states": hidden_states, "residual": residual} ) hidden_states, _ = self.norm(hidden_states, residual) + if aux_hidden_states: + return hidden_states, aux_hidden_states return hidden_states @@ -1577,6 +1586,7 @@ class Qwen3OmniMoeThinkerForConditionalGeneration( SupportsMultiModal, SupportsPP, SupportsMRoPE, + SupportsEagle3, Qwen3OmniMoeConditionalGenerationMixin, SupportsTranscription, ): diff --git a/vllm/model_executor/models/registry.py b/vllm/model_executor/models/registry.py index 24fc66a08999..c6043716a809 100644 --- a/vllm/model_executor/models/registry.py +++ b/vllm/model_executor/models/registry.py @@ -643,6 +643,7 @@ "DFlashMuseGlimmerAssistantModel": ("qwen3_dflash", "DFlashQwen3ForCausalLM"), "DSparkDraftModel": ("vllm.models.deepseek_v4", "DSparkDeepseekV4ForCausalLM"), "Qwen3DSparkModel": ("qwen3_dspark", "Qwen3DSparkForCausalLM"), + "Qwen3OmniDSparkModel": ("qwen3_dspark", "Qwen3DSparkForCausalLM"), "K3DSparkModel": ( "vllm.models.kimi_k3.nvidia.dspark_mla", "K3DSparkForCausalLM", diff --git a/vllm/transformers_utils/configs/speculators/algos.py b/vllm/transformers_utils/configs/speculators/algos.py index fbd2c4518196..8ec978d73591 100644 --- a/vllm/transformers_utils/configs/speculators/algos.py +++ b/vllm/transformers_utils/configs/speculators/algos.py @@ -154,10 +154,22 @@ def update_dspark(config_dict: dict, pre_trained_config: dict) -> None: False (anchor is a bonus token, only mask tokens predict, yielding block_size - 1 speculative tokens). """ - pre_trained_config["architectures"] = ["Qwen3DSparkModel"] - pre_trained_config["sample_from_anchor"] = config_dict.get( - "sample_from_anchor", False - ) + architectures = config_dict.get("architectures") or [] + supported_architectures = { + "Qwen3DSparkModel", + "Qwen3OmniDSparkModel", + } + selected_architectures = [ + architecture + for architecture in architectures + if architecture in supported_architectures + ] + # Legacy msModelSpec checkpoints use a training-only architecture name. + pre_trained_config["architectures"] = selected_architectures or ["Qwen3DSparkModel"] + + sample_from_anchor = config_dict.get("sample_from_anchor", False) + pre_trained_config["sample_from_anchor"] = sample_from_anchor + pre_trained_config["dspark_bonus_anchor"] = not sample_from_anchor aux_layer_ids = config_dict["aux_hidden_state_layer_ids"] pre_trained_config["eagle_aux_hidden_state_layer_ids"] = aux_layer_ids @@ -173,6 +185,7 @@ def update_dspark(config_dict: dict, pre_trained_config: dict) -> None: "block_size", "enable_confidence_head", "confidence_head_with_markov", + "use_aux_hidden_state", ): if config_dict.get(key) is not None: pre_trained_config[key] = config_dict[key] diff --git a/vllm/v1/worker/gpu/spec_decode/dspark/utils.py b/vllm/v1/worker/gpu/spec_decode/dspark/utils.py index a33fa259c3e7..177060747ab8 100644 --- a/vllm/v1/worker/gpu/spec_decode/dspark/utils.py +++ b/vllm/v1/worker/gpu/spec_decode/dspark/utils.py @@ -86,11 +86,16 @@ def load_dspark_model(target_model: nn.Module, vllm_config: VllmConfig) -> nn.Mo ) target_inner = target_language_model.model draft_inner = draft_model.model + target_vocab_size = vllm_config.model_config.get_vocab_size() target_embed = getattr(target_inner, "embed_tokens", None) draft_embed = getattr(draft_inner, "embed_tokens", None) - if target_embed is not None and _should_share( - draft_model, "has_own_embed_tokens", draft_embed, target_embed + if ( + target_embed is not None + and draft_model_config.get_vocab_size() <= target_vocab_size + and _should_share( + draft_model, "has_own_embed_tokens", draft_embed, target_embed + ) ): if draft_embed is not None: del draft_inner.embed_tokens @@ -98,8 +103,14 @@ def load_dspark_model(target_model: nn.Module, vllm_config: VllmConfig) -> nn.Mo target_lm_head = get_target_lm_head(target_model, target_language_model) draft_lm_head = getattr(draft_model, "lm_head", None) - if target_lm_head is not None and _should_share( - draft_model, "has_own_lm_head", draft_lm_head, target_lm_head + draft_output_vocab_size = ( + getattr(draft_model_config.hf_config, "draft_vocab_size", None) + or draft_model_config.get_vocab_size() + ) + if ( + target_lm_head is not None + and draft_output_vocab_size == target_vocab_size + and _should_share(draft_model, "has_own_lm_head", draft_lm_head, target_lm_head) ): if draft_lm_head is not None: del draft_model.lm_head From bbe8b23e1a2b32a96240b27f63255170d09ef144 Mon Sep 17 00:00:00 2001 From: Brian Woods Date: Sat, 22 Aug 2026 15:38:10 -0700 Subject: [PATCH 017/235] [Deprecation] Remove dead use_prefill_decode_attention flag (#52557) Signed-off-by: Brian Woods Co-authored-by: Claude Opus 5 (1M context) Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com> --- tests/engine/test_arg_utils.py | 5 ----- tests/v1/attention/utils.py | 1 - vllm/config/attention.py | 4 ---- 3 files changed, 10 deletions(-) diff --git a/tests/engine/test_arg_utils.py b/tests/engine/test_arg_utils.py index 6a799ee912d4..db64c5bb6f03 100644 --- a/tests/engine/test_arg_utils.py +++ b/tests/engine/test_arg_utils.py @@ -383,8 +383,6 @@ def test_attention_config(): "FLASH_ATTN", "--attention-config.flash_attn_version", "3", - "--attention-config.use_prefill_decode_attention", - "true", "--attention-config.flash_attn_max_num_splits_for_cuda_graph", "16", "--attention-config.use_trtllm_attention", @@ -398,7 +396,6 @@ def test_attention_config(): assert engine_args.attention_config.backend is not None assert engine_args.attention_config.backend.name == "FLASH_ATTN" assert engine_args.attention_config.flash_attn_version == 3 - assert engine_args.attention_config.use_prefill_decode_attention is True assert engine_args.attention_config.flash_attn_max_num_splits_for_cuda_graph == 16 assert engine_args.attention_config.use_trtllm_attention is True assert engine_args.attention_config.disable_flashinfer_q_quantization is True @@ -408,7 +405,6 @@ def test_attention_config(): [ "--attention-config=" '{"backend": "FLASHINFER", "flash_attn_version": 2, ' - '"use_prefill_decode_attention": false, ' '"flash_attn_max_num_splits_for_cuda_graph": 8, ' '"use_trtllm_attention": false, ' '"disable_flashinfer_q_quantization": false}', @@ -419,7 +415,6 @@ def test_attention_config(): assert engine_args.attention_config.backend is not None assert engine_args.attention_config.backend.name == "FLASHINFER" assert engine_args.attention_config.flash_attn_version == 2 - assert engine_args.attention_config.use_prefill_decode_attention is False assert engine_args.attention_config.flash_attn_max_num_splits_for_cuda_graph == 8 assert engine_args.attention_config.use_trtllm_attention is False assert engine_args.attention_config.disable_flashinfer_q_quantization is False diff --git a/tests/v1/attention/utils.py b/tests/v1/attention/utils.py index 6aeeb12b4bc8..ee702bdbdfa7 100644 --- a/tests/v1/attention/utils.py +++ b/tests/v1/attention/utils.py @@ -387,7 +387,6 @@ class BackendConfig: name="RocmAttn", attention_config={ "backend": "ROCM_ATTN", - "use_prefill_decode_attention": True, }, comp_config={ "cudagraph_mode": "FULL", diff --git a/vllm/config/attention.py b/vllm/config/attention.py index 907917d091cb..ea92be32cdcd 100644 --- a/vllm/config/attention.py +++ b/vllm/config/attention.py @@ -42,10 +42,6 @@ class AttentionConfig: """Force vllm to use a specific flash-attention version (2, 3, or 4). Only valid when using the flash-attention backend.""" - use_prefill_decode_attention: bool = False - """Use separate prefill and decode kernels for attention instead of - the unified triton kernel.""" - flash_attn_max_num_splits_for_cuda_graph: int = 32 """Flash Attention max number splits for cuda graph decode.""" From 30b34171b113887e0e08d7f6d06e2e5a5c33b9d2 Mon Sep 17 00:00:00 2001 From: Andreas Karatzas Date: Sat, 22 Aug 2026 20:00:25 -0500 Subject: [PATCH 018/235] [ROCm][CI] Restore attention coverage after KV-cache layout refactor (#53351) Signed-off-by: Andreas Karatzas --- tests/kernels/attention/test_minimax_m3.py | 15 +++++++++ tests/v1/attention/test_attention_backends.py | 29 +++++++++++------ tests/v1/attention/test_mla_backends.py | 31 +++++++++---------- 3 files changed, 50 insertions(+), 25 deletions(-) diff --git a/tests/kernels/attention/test_minimax_m3.py b/tests/kernels/attention/test_minimax_m3.py index 42883d98994c..745d0fd587f1 100644 --- a/tests/kernels/attention/test_minimax_m3.py +++ b/tests/kernels/attention/test_minimax_m3.py @@ -971,6 +971,21 @@ def test_aiter_sparse_pa_layout_contract(monkeypatch): assert value_cache.is_contiguous() +def test_aiter_sparse_pa_rejects_multiple_kv_heads(monkeypatch): + """Do not pair AITER's separated cache layout with the Triton fallback.""" + import vllm.models.minimax_m3.common.sparse_attention as sparse_attn_mod + + monkeypatch.setattr(sparse_attn_mod.rocm_aiter_ops, "is_enabled", lambda: True) + monkeypatch.setattr( + sparse_attn_mod.rocm_aiter_ops, + "is_shuffle_kv_cache_enabled", + lambda: True, + ) + + with pytest.raises(ValueError, match="num_kv_heads == 1"): + sparse_attn_mod.minimax_m3_use_aiter_sparse_pa(2) + + def test_indexer_cache_squeezes_to_contiguous_3d(): """The indexer side cache is standardized 4D with H=1: under both layouts the allocator's logical view stays contiguous and squeezes (as diff --git a/tests/v1/attention/test_attention_backends.py b/tests/v1/attention/test_attention_backends.py index cd78911ce812..cc16242753b2 100644 --- a/tests/v1/attention/test_attention_backends.py +++ b/tests/v1/attention/test_attention_backends.py @@ -42,9 +42,11 @@ ] -def _actual_backend(backend) -> AttentionBackendEnum: +def _actual_backend(backend: AttentionBackendEnum | str) -> AttentionBackendEnum: """Resolve pseudo-backends (FLEX_ATTENTION_SLOW) to their real enum.""" - if backend == "FLEX_ATTENTION_SLOW": + if isinstance(backend, str): + if backend != "FLEX_ATTENTION_SLOW": + raise ValueError(f"Unknown pseudo-backend: {backend}") return AttentionBackendEnum.FLEX_ATTENTION return backend @@ -266,7 +268,7 @@ def _clone_kv_cache_in_layout( def run_attention_backend( - backend: AttentionBackendEnum, + backend: AttentionBackendEnum | str, kv_cache_spec: FullAttentionSpec, layer_names: list[str], vllm_config, @@ -283,9 +285,7 @@ def run_attention_backend( ) -> torch.Tensor: """Run attention computation using the specified backend's AttentionImpl.""" - use_direct_block_mask = not current_platform.is_rocm() and is_torch_equal_or_newer( - "2.9.0.dev0" - ) + use_direct_block_mask = is_torch_equal_or_newer("2.9.0.dev0") if backend == "FLEX_ATTENTION_SLOW": use_direct_block_mask = False backend = _actual_backend(backend) @@ -377,7 +377,7 @@ def mock_get_per_layer_parameters(vllm_config, layer_names, impl_cls): def _test_backend_correctness( batch_spec: BatchSpec, model: str, - backend_to_test: list[AttentionBackendEnum], + backend_to_test: list[AttentionBackendEnum | str], mask_mod, *, causal: bool = True, @@ -655,7 +655,7 @@ def _test_backend_correctness( ) # Check numerical similarity - def error_msg(msg: str, backend_name: str): + def error_msg(msg: str, backend_name: AttentionBackendEnum | str): return f"[{backend_name}] output differs from SDPA baseline. {msg}" torch.testing.assert_close( @@ -1046,14 +1046,24 @@ def causal_mask_mod( SLIDING_WINDOW_BACKENDS_TO_TEST = [ AttentionBackendEnum.FLEX_ATTENTION, AttentionBackendEnum.TRITON_ATTN, + "FLEX_ATTENTION_SLOW", ] else: SLIDING_WINDOW_BACKENDS_TO_TEST = [ AttentionBackendEnum.FLASH_ATTN, AttentionBackendEnum.FLEX_ATTENTION, AttentionBackendEnum.TRITON_ATTN, + "FLEX_ATTENTION_SLOW", ] +# Encoder-only FlexAttention always uses the slow builder, so the pseudo-backend +# would run an identical implementation twice. +SLIDING_WINDOW_ENCODER_BACKENDS_TO_TEST = [ + backend + for backend in SLIDING_WINDOW_BACKENDS_TO_TEST + if backend != "FLEX_ATTENTION_SLOW" +] + @pytest.mark.parametrize( "batch_spec_name", @@ -1162,7 +1172,7 @@ def bidi_sliding_window_mask_mod( _test_backend_correctness( batch_spec, model, - SLIDING_WINDOW_BACKENDS_TO_TEST, + SLIDING_WINDOW_ENCODER_BACKENDS_TO_TEST, sliding_window_mask_mod_fn, causal=False, attn_type=AttentionType.ENCODER_ONLY, @@ -1173,6 +1183,7 @@ def bidi_sliding_window_mask_mod( NON_CAUSAL_BACKENDS_TO_TEST = [ AttentionBackendEnum.FLASH_ATTN, AttentionBackendEnum.FLEX_ATTENTION, + "FLEX_ATTENTION_SLOW", ] if current_platform.is_rocm(): diff --git a/tests/v1/attention/test_mla_backends.py b/tests/v1/attention/test_mla_backends.py index 5c621cf6ad47..cbccbfc9a2ac 100644 --- a/tests/v1/attention/test_mla_backends.py +++ b/tests/v1/attention/test_mla_backends.py @@ -62,6 +62,9 @@ AttentionBackendEnum.TOKENSPEED_MLA, ] +if current_platform.is_rocm(): + BACKENDS_TO_TEST.append(AttentionBackendEnum.ROCM_AITER_MLA) + DEVICE_TYPE = current_platform.device_type @@ -201,13 +204,18 @@ def test_mla_post_load_preserves_runtime_weight_addresses(monkeypatch): # Validate parameter combinations during collection, before GPU fixtures run. -PREFILL_BACKENDS_TO_TEST = [ - MLAPrefillBackendEnum.ROCM_AITER_FA, - MLAPrefillBackendEnum.FLASH_ATTN, - MLAPrefillBackendEnum.FLASHINFER, - MLAPrefillBackendEnum.TRTLLM_RAGGED, - MLAPrefillBackendEnum.TOKENSPEED_MLA, -] +PREFILL_BACKENDS_TO_TEST: list[MLAPrefillBackendEnum] = [] +if current_platform.is_cuda(): + PREFILL_BACKENDS_TO_TEST.extend( + [ + MLAPrefillBackendEnum.FLASH_ATTN, + MLAPrefillBackendEnum.FLASHINFER, + MLAPrefillBackendEnum.TRTLLM_RAGGED, + MLAPrefillBackendEnum.TOKENSPEED_MLA, + ] + ) +elif current_platform.is_rocm(): + PREFILL_BACKENDS_TO_TEST.append(MLAPrefillBackendEnum.ROCM_AITER_FA) MLA_DIMENSIONS_TO_TEST = [ ("deepseek", 128, 128), @@ -250,15 +258,6 @@ def _prefill_backend_dimension_params(): ) ) ) - elif ( - current_platform.is_rocm() - and prefill_backend == MLAPrefillBackendEnum.FLASH_ATTN - ): - # Pre-existing on main: flash-attn MLA prefill returns wrong - # outputs and faults on ROCm, so it fails on the base commit too. - marks.append( - pytest.mark.skip(reason="FLASH_ATTN MLA prefill broken on ROCm") - ) params.append( pytest.param( prefill_backend, From 9dba2c9b374af3a47c1cef952195ee114bab94c7 Mon Sep 17 00:00:00 2001 From: Biswa Panda Date: Sat, 22 Aug 2026 19:31:23 -0700 Subject: [PATCH 019/235] [Rust Frontend][RL]: report engine world size over gRPC (#53204) Signed-off-by: Biswa Panda --- rust/proto/control.proto | 1 + rust/src/engine-core-client/src/protocol/handshake.rs | 3 ++- rust/src/server/src/grpc/control.rs | 1 + rust/src/server/src/grpc/tests.rs | 10 +++++++++- 4 files changed, 13 insertions(+), 2 deletions(-) diff --git a/rust/proto/control.proto b/rust/proto/control.proto index 1903405e2749..45eac602eccc 100644 --- a/rust/proto/control.proto +++ b/rust/proto/control.proto @@ -55,6 +55,7 @@ message ParallelismInfo { uint32 data_parallel_size = 3; uint32 data_parallel_rank = 4; uint32 decode_context_parallel_size = 5; + uint64 world_size = 6; } message GetModelInfoRequest {} diff --git a/rust/src/engine-core-client/src/protocol/handshake.rs b/rust/src/engine-core-client/src/protocol/handshake.rs index c7146abdbb6a..c14c19a50373 100644 --- a/rust/src/engine-core-client/src/protocol/handshake.rs +++ b/rust/src/engine-core-client/src/protocol/handshake.rs @@ -61,7 +61,8 @@ pub struct EngineCoreReadyResponse { pub dtype: ModelDtype, /// Python vLLM version reported by the engine process. pub vllm_version: String, - /// World size (TP * PP) from the parallel config. + /// Per-DP-engine world size (TP * PP * PCP) from the parallel config. + /// Deployment-wide world size also multiplies by `data_parallel_size`. pub world_size: u64, /// Data-parallel size from this EngineCore's effective parallel config. /// Dense independent-DP ranks are reconfigured to report `1`; the client diff --git a/rust/src/server/src/grpc/control.rs b/rust/src/server/src/grpc/control.rs index 9637dcbf8efa..ba4ee7b99774 100644 --- a/rust/src/server/src/grpc/control.rs +++ b/rust/src/server/src/grpc/control.rs @@ -46,6 +46,7 @@ impl ControlServiceImpl { data_parallel_size: self.client().data_parallel_size().min(u32::MAX as usize) as u32, data_parallel_rank: ready.data_parallel_rank, decode_context_parallel_size: ready.decode_context_parallel_size, + world_size: ready.world_size, } } diff --git a/rust/src/server/src/grpc/tests.rs b/rust/src/server/src/grpc/tests.rs index 9c1ea33b6b0e..6a1074b5240c 100644 --- a/rust/src/server/src/grpc/tests.rs +++ b/rust/src/server/src/grpc/tests.rs @@ -1626,6 +1626,9 @@ async fn control_aggregates_multi_engine_capacity() { ready_0.max_model_len = 8_192; ready_0.num_gpu_blocks = 10; ready_0.effective_data_parallel_size = 2; + ready_0.tensor_parallel_size = 2; + ready_0.pipeline_parallel_size = 3; + ready_0.world_size = 12; ready_0.weight_transfer_backend = Some("nccl".to_string()); ready_0.enable_sleep_mode = true; ready_0.supports_draft_weight_updates = true; @@ -1634,6 +1637,9 @@ async fn control_aggregates_multi_engine_capacity() { ready_1.max_model_len = 4_096; ready_1.num_gpu_blocks = 20; ready_1.effective_data_parallel_size = 2; + ready_1.tensor_parallel_size = 2; + ready_1.pipeline_parallel_size = 3; + ready_1.world_size = 12; ready_1.data_parallel_rank = 1; let engine_tasks = [ready_0, ready_1].map(|ready| { @@ -1684,7 +1690,9 @@ async fn control_aggregates_multi_engine_capacity() { assert!(rl.weight_transfer_backend.is_empty()); assert!(!rl.sleep_mode_enabled); assert!(!rl.draft_weight_updates_enabled); - assert_eq!(server.parallelism.unwrap().data_parallel_size, 2); + let parallelism = server.parallelism.unwrap(); + assert_eq!(parallelism.data_parallel_size, 2); + assert_eq!(parallelism.world_size, 12); drop(engine_tasks); } From d569e71bd031693ac945c590731a36f0e70789cc Mon Sep 17 00:00:00 2001 From: Andreas Karatzas Date: Sat, 22 Aug 2026 22:37:55 -0500 Subject: [PATCH 020/235] [CI][AMD] Honor single-node Docker workload timeout (#52547) Signed-off-by: Andreas Karatzas --- .../scripts/hardware_ci/run-amd-test.sh | 71 +++++++++++++++++-- 1 file changed, 66 insertions(+), 5 deletions(-) diff --git a/.buildkite/scripts/hardware_ci/run-amd-test.sh b/.buildkite/scripts/hardware_ci/run-amd-test.sh index b9674c108b13..4b96da67c768 100755 --- a/.buildkite/scripts/hardware_ci/run-amd-test.sh +++ b/.buildkite/scripts/hardware_ci/run-amd-test.sh @@ -83,6 +83,8 @@ export PYTHONFAULTHANDLER # depend on their current working directory. export PYTHONPATH="${PYTHONPATH:-..}" +ci_started_at=$SECONDS + ############################################################################### # Helper Functions ############################################################################### @@ -130,6 +132,66 @@ cleanup_network() { fi } +amd_ci_teardown_log() { + local event=$1 + shift + + printf '[amd-ci-teardown] event=%s' "${event}" >&2 + if (($#)); then + printf ' %s' "$@" >&2 + fi + printf '\n' >&2 +} + +run_docker_with_ci_timeout() { + local raw_timeout=${CONTAINER_TIMEOUT_S:-0} + local configured_timeout=0 + local elapsed=0 + local remaining_timeout=0 + local status=0 + + if [[ ! "${raw_timeout}" =~ ^(0|[1-9][0-9]{0,5})$ ]]; then + amd_ci_teardown_log configuration_error \ + "reason=invalid_timeout" "expected=integer_0_to_604800" + return 2 + fi + configured_timeout=$((10#${raw_timeout})) + if ((configured_timeout > 604800)); then + amd_ci_teardown_log configuration_error \ + "reason=invalid_timeout" "expected=integer_0_to_604800" + return 2 + fi + + remaining_timeout=${configured_timeout} + if ((configured_timeout > 0)); then + if ! command -v timeout >/dev/null 2>&1; then + amd_ci_teardown_log configuration_error \ + "reason=timeout_command_not_found" + return 127 + fi + elapsed=$((SECONDS - ci_started_at)) + if ((elapsed >= configured_timeout)); then + amd_ci_teardown_log deadline_exhausted \ + "mode=docker" "configured_s=${configured_timeout}" "elapsed_s=${elapsed}" + return 124 + fi + remaining_timeout=$((configured_timeout - elapsed)) + fi + + amd_ci_teardown_log workload_started \ + "mode=docker" "timeout_s=${remaining_timeout}" + if ((remaining_timeout == 0)); then + "$@" || status=$? + else + # Docker runs interactively, so keep it in the foreground process group. + # --verbose records both TERM and any KILL escalation caused by timeout. + timeout --verbose --foreground --signal=TERM --kill-after=10s \ + "${remaining_timeout}s" "$@" || status=$? + fi + amd_ci_teardown_log workload_finished "mode=docker" "status=${status}" + return "${status}" +} + prepare_artifact_image() { if [[ "${VLLM_CI_USE_ARTIFACTS:-0}" != "1" ]]; then return 1 @@ -1599,8 +1661,9 @@ else echo "ROCm debug agent not enabled, coredumps are disabled in the test container." fi + exit_code=0 # shellcheck disable=SC2086 # word splitting is intentional: both hold multiple docker flags - docker run \ + run_docker_with_ci_timeout docker run \ "${docker_run_terminal_args[@]}" \ --device /dev/kfd $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES \ $RDMA_FLAGS \ @@ -1636,8 +1699,6 @@ else "${standalone_merge_base_env[@]}" \ --name "${container_name}" \ "${image_name}" \ - /bin/bash -c "${CONTAINER_PREFLIGHT} && ${commands}" - - exit_code=$? - handle_pytest_exit "$exit_code" + /bin/bash -c "${CONTAINER_PREFLIGHT} && ${commands}" || exit_code=$? + handle_pytest_exit "${exit_code}" fi From a3561ef8e49d3545c4078df43444beb4c98ae124 Mon Sep 17 00:00:00 2001 From: anmolgupt <14880251+anmolgupt@users.noreply.github.com> Date: Sat, 22 Aug 2026 21:34:38 -0700 Subject: [PATCH 021/235] [Bugfix][Model] Mistral3: fix image placeholder grid for processor size overrides (#52874) Signed-off-by: Anmol Gupta <14880251+anmolgupt@users.noreply.github.com> --- .../multimodal/processing/test_mistral3.py | 143 ++++++++++++++++++ vllm/model_executor/models/lightonocr.py | 15 +- vllm/model_executor/models/mistral3.py | 59 +++++++- 3 files changed, 210 insertions(+), 7 deletions(-) create mode 100644 tests/models/multimodal/processing/test_mistral3.py diff --git a/tests/models/multimodal/processing/test_mistral3.py b/tests/models/multimodal/processing/test_mistral3.py new file mode 100644 index 000000000000..fcb783f22f6e --- /dev/null +++ b/tests/models/multimodal/processing/test_mistral3.py @@ -0,0 +1,143 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +"""Tests for Mistral3's multimodal preprocessing kwargs.""" + +import pytest +import torch +from PIL import Image +from transformers import BatchFeature + +from vllm.model_executor.models.lightonocr import LightOnOCRProcessingInfo +from vllm.model_executor.models.mistral3 import Mistral3HFEncoderInfo +from vllm.model_executor.models.pixtral import PixtralHFEncoderInfo +from vllm.multimodal import MULTIMODAL_REGISTRY +from vllm.multimodal.inputs import MultiModalKwargsItems + +from ...utils import build_model_context + +# This repo ships both params.json (Mistral) and config.json (HF). Auto config +# selects PixtralForConditionalGeneration; force HF to exercise Mistral3. +_MODEL_CONFIG_KWARGS = {"config_format": "hf"} +_MODEL_ID = "mistralai/Mistral-Small-3.1-24B-Instruct-2503" +_LIGHTON_MODEL_ID = "lightonai/LightOnOCR-1B-1025" + + +def _processed_pixel_values( + hf_processor, + images: list[Image.Image], + mm_processor_kwargs: dict[str, object], +) -> list[torch.Tensor]: + """Resize via the HF image processor and un-pad to per-image H×W.""" + image_processor = hf_processor.image_processor + hf_out = image_processor(images=images, return_tensors="pt", **mm_processor_kwargs) + pixel_values = hf_out["pixel_values"] + image_sizes = hf_out["image_sizes"] + return [p[:, :h, :w] for p, (h, w) in zip(pixel_values, image_sizes)] + + +def _placeholder_count_from_prompt_updates( + processor, + images: list[Image.Image], + pixel_values: list[torch.Tensor], + mm_processor_kwargs: dict[str, object], +) -> int: + hf_inputs = BatchFeature({"pixel_values": pixel_values}) + fields_config = processor._get_mm_fields_config(hf_inputs, mm_processor_kwargs) + out_mm_kwargs = MultiModalKwargsItems.from_hf_inputs(hf_inputs, fields_config) + # Prompt updates use raw PIL sizes and must predict the processed grid. + mm_items = processor.info.parse_mm_data({"image": images}) + updates = processor._get_prompt_updates( + mm_items, mm_processor_kwargs, out_mm_kwargs + ) + image_token_id = processor.info.get_hf_config().image_token_index + + total = 0 + for item_idx in range(len(images)): + details = updates[0].resolve(item_idx).content + total += details.full.count(image_token_id) + return total + + +def _expected_placeholder_tokens_per_image( + hf_processor, + pixel_values: torch.Tensor, +) -> int: + """Count projected tokens from the actual HF-processed H×W.""" + image_h, image_w = pixel_values.shape[-2:] + patch_size = hf_processor.image_processor.patch_size + if isinstance(patch_size, dict): + patch_h = patch_size["height"] + patch_w = patch_size["width"] + else: + patch_h = patch_w = int(patch_size) + + spatial_merge_size = getattr(hf_processor, "spatial_merge_size", 1) + assert image_h % patch_h == 0 + assert image_w % patch_w == 0 + + return (image_h // (patch_h * spatial_merge_size)) * ( + image_w // (patch_w * spatial_merge_size) + ) + + +@pytest.mark.parametrize("model_id", [_MODEL_ID]) +@pytest.mark.parametrize( + ("mm_processor_kwargs", "image_size", "expected_toks_per_img"), + [ + ({}, (448, 448), 256), + ({"size": {"longest_edge": 1008}}, (1540, 1540), 1296), + ({"size": {"longest_edge": 1288}}, (1536, 1187), 1656), + ({"size": {"longest_edge": 1008}}, (29, 29), 1), + ({"size": {"longest_edge": 1000}}, (1540, 1700), 1152), + ], +) +@pytest.mark.parametrize("num_imgs", [1, 2]) +@pytest.mark.parametrize("kwargs_on_init", [True, False]) +def test_processor_size_override( + model_id: str, + mm_processor_kwargs: dict[str, object], + image_size: tuple[int, int], + expected_toks_per_img: int, + num_imgs: int, + kwargs_on_init: bool, +): + ctx = build_model_context( + model_id, + mm_processor_kwargs=mm_processor_kwargs if kwargs_on_init else None, + limit_mm_per_prompt={"image": num_imgs}, + model_config_kwargs=_MODEL_CONFIG_KWARGS, + ) + processor = MULTIMODAL_REGISTRY.create_processor(ctx.model_config) + hf_processor_mm_kwargs = {} if kwargs_on_init else mm_processor_kwargs + hf_processor = processor.info.get_hf_processor(**hf_processor_mm_kwargs) + + dummy_image = Image.new("RGB", image_size, color=(127, 127, 127)) + images = [dummy_image] * num_imgs + merged_mm_kwargs = processor.info.ctx.get_merged_mm_kwargs(hf_processor_mm_kwargs) + pixel_values = _processed_pixel_values(hf_processor, images, merged_mm_kwargs) + + image_token_count = _placeholder_count_from_prompt_updates( + processor, images, pixel_values, hf_processor_mm_kwargs + ) + expected_from_pixel_values = _expected_placeholder_tokens_per_image( + hf_processor, pixel_values[0] + ) + assert expected_from_pixel_values == expected_toks_per_img + assert image_token_count == expected_from_pixel_values * num_imgs + + +def test_lightonocr_keeps_vision_config_image_size(): + ctx = build_model_context( + _LIGHTON_MODEL_ID, + mm_processor_kwargs={"size": {"longest_edge": 1008}}, + model_config_kwargs=_MODEL_CONFIG_KWARGS, + ) + processor = MULTIMODAL_REGISTRY.create_processor(ctx.model_config) + + assert isinstance(processor.info, LightOnOCRProcessingInfo) + encoder_info = processor.info.get_vision_encoder_info() + assert isinstance(encoder_info, PixtralHFEncoderInfo) + assert not isinstance(encoder_info, Mistral3HFEncoderInfo) + assert encoder_info.get_image_size() == ( + processor.info.get_hf_config().vision_config.image_size + ) diff --git a/vllm/model_executor/models/lightonocr.py b/vllm/model_executor/models/lightonocr.py index ccb33177fbca..7f108ba556a6 100644 --- a/vllm/model_executor/models/lightonocr.py +++ b/vllm/model_executor/models/lightonocr.py @@ -38,7 +38,18 @@ _I = TypeVar("_I", bound=Mistral3ProcessingInfo) -class LightOnOCRMultiModalProcessor(BaseMultiModalProcessor[Mistral3ProcessingInfo]): +class LightOnOCRProcessingInfo(Mistral3ProcessingInfo): + def get_vision_encoder_info( + self, + mm_processor_kwargs: Mapping[str, object] | None = None, + ) -> PixtralHFEncoderInfo: + # LightOnOCR prompt updates use vision_config.image_size directly; + # preserve that behavior rather than inheriting Mistral3's processor + # longest_edge override. + return PixtralHFEncoderInfo(self.get_hf_config()) + + +class LightOnOCRMultiModalProcessor(BaseMultiModalProcessor[LightOnOCRProcessingInfo]): def _get_hf_processor_text(self, mm_counts: Mapping[str, int]) -> str: return self.dummy_inputs.get_dummy_text(mm_counts) @@ -124,7 +135,7 @@ def replace(item_idx: int): @MULTIMODAL_REGISTRY.register_processor( LightOnOCRMultiModalProcessor, - info=Mistral3ProcessingInfo, + info=LightOnOCRProcessingInfo, dummy_inputs=Mistral3DummyInputsBuilder, ) class LightOnOCRForConditionalGeneration(Mistral3ForConditionalGeneration): diff --git a/vllm/model_executor/models/mistral3.py b/vllm/model_executor/models/mistral3.py index 11f214c65b1a..23c632361904 100644 --- a/vllm/model_executor/models/mistral3.py +++ b/vllm/model_executor/models/mistral3.py @@ -1,6 +1,7 @@ # SPDX-License-Identifier: Apache-2.0 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project +import math from collections.abc import Iterable, Mapping, Sequence from typing import Annotated, Literal @@ -22,7 +23,11 @@ MultiModalFieldConfig, MultiModalKwargsItems, ) -from vllm.multimodal.parse import ImageProcessorItems, ImageSize, MultiModalDataItems +from vllm.multimodal.parse import ( + ImageProcessorItems, + ImageSize, + MultiModalDataItems, +) from vllm.multimodal.processing import ( BaseDummyInputsBuilder, BaseMultiModalProcessor, @@ -50,7 +55,6 @@ init_vllm_registered_model, maybe_prefix, ) -from .vision import get_vision_encoder_info class Mistral3ImagePixelInputs(TensorSchema): @@ -170,12 +174,57 @@ def forward( return hidden_states +class Mistral3HFEncoderInfo(PixtralHFEncoderInfo): + def __init__(self, hf_config: Mistral3Config, image_size: int) -> None: + super().__init__(hf_config) + self.image_size = image_size + + def get_image_size(self) -> int: + return self.image_size + + def get_patch_grid_size( + self, + *, + image_width: int, + image_height: int, + ) -> tuple[int, int]: + max_size = self.get_image_size() + ratio = max(image_width / max_size, image_height / max_size) + + if ratio > 1: + image_width = math.floor(image_width / ratio) + image_height = math.floor(image_height / ratio) + + patch_size = self.vision_config.patch_size + spatial_merge_size = self.hf_config.spatial_merge_size + + # The HF processor rounds each dimension up to the vision patch size + # before the projector drops incomplete spatial-merge groups. This is + # not equivalent to rounding directly to the merged patch size. + num_width_patches = (image_width - 1) // patch_size + 1 + num_height_patches = (image_height - 1) // patch_size + 1 + + ncols = num_width_patches // spatial_merge_size + nrows = num_height_patches // spatial_merge_size + return ncols, nrows + + class Mistral3ProcessingInfo(BaseProcessingInfo): def get_hf_config(self) -> Mistral3Config: return self.ctx.get_hf_config(Mistral3Config) - def get_vision_encoder_info(self): - return get_vision_encoder_info(self.get_hf_config()) + def get_vision_encoder_info( + self, + mm_processor_kwargs: Mapping[str, object] | None = None, + ) -> Mistral3HFEncoderInfo: + processor = self.get_hf_processor() + size = processor.image_processor.size + merged_kwargs = self.ctx.get_merged_mm_kwargs(mm_processor_kwargs or {}) + if override_size := merged_kwargs.get("size"): + size = size | override_size + + image_size = size["longest_edge"] + return Mistral3HFEncoderInfo(self.get_hf_config(), image_size) def get_hf_processor(self, **kwargs: object): return self.ctx.get_hf_processor(PixtralProcessor, **kwargs) @@ -284,7 +333,7 @@ def _get_prompt_updates( image_end_id = vocab[processor.image_end_token] assert isinstance(hf_config.vision_config, PixtralVisionConfig) - encoder_info = PixtralHFEncoderInfo(hf_config) + encoder_info = self.info.get_vision_encoder_info(hf_processor_mm_kwargs) def get_replacement(item_idx: int): images = mm_items.get_items("image", ImageProcessorItems) From ebcd60646741be239026db75358715ba707eb50f Mon Sep 17 00:00:00 2001 From: Cyrus Leung Date: Sun, 23 Aug 2026 14:50:32 +0800 Subject: [PATCH 022/235] [MM] Simplify prompt updates: replace `PromptSeq` with `list[int]` (#53372) Signed-off-by: DarkLight1337 Co-authored-by: Qwen Code --- docs/contributing/model/multimodal.md | 6 + .../multimodal/processing/test_openvla.py | 4 +- tests/multimodal/test_cache.py | 4 +- tests/multimodal/test_processing.py | 495 +++++----------- vllm/model_executor/models/audioflamingo3.py | 2 +- vllm/model_executor/models/cohere2_vision.py | 8 +- vllm/model_executor/models/ernie45_vl.py | 11 +- vllm/model_executor/models/funaudiochat.py | 2 +- vllm/model_executor/models/gemma3_mm.py | 11 +- vllm/model_executor/models/gemma3n_mm.py | 27 +- vllm/model_executor/models/gemma4_mm.py | 18 +- vllm/model_executor/models/glm4_1v.py | 13 +- vllm/model_executor/models/glmasr.py | 2 +- vllm/model_executor/models/h2ovl.py | 5 +- vllm/model_executor/models/idefics3.py | 15 +- vllm/model_executor/models/interns1.py | 31 +- vllm/model_executor/models/internvl.py | 9 +- vllm/model_executor/models/isaac.py | 11 +- vllm/model_executor/models/kanana_v.py | 5 +- vllm/model_executor/models/lfm2_vl.py | 15 +- vllm/model_executor/models/midashenglm.py | 2 +- vllm/model_executor/models/mimo_v2_omni.py | 2 +- vllm/model_executor/models/minicpmo.py | 19 +- vllm/model_executor/models/minicpmv.py | 69 ++- vllm/model_executor/models/minicpmv4_6.py | 77 ++- vllm/model_executor/models/mllama4.py | 12 +- vllm/model_executor/models/molmo.py | 5 +- vllm/model_executor/models/moss_audio.py | 4 +- .../models/moss_transcribe_diarize.py | 7 +- vllm/model_executor/models/muse_glimmer.py | 4 +- .../model_executor/models/nano_nemotron_vl.py | 14 +- vllm/model_executor/models/nvlm_d.py | 13 +- vllm/model_executor/models/openvla.py | 2 +- vllm/model_executor/models/ovis.py | 5 +- vllm/model_executor/models/phi3v.py | 16 +- vllm/model_executor/models/phi4mm.py | 31 +- vllm/model_executor/models/phi4siglip.py | 7 +- vllm/model_executor/models/pixtral.py | 2 +- .../models/qwen2_5_omni_thinker.py | 6 +- vllm/model_executor/models/qwen3_asr.py | 2 +- .../models/qwen3_omni_moe_thinker.py | 6 +- vllm/model_executor/models/qwen3_vl.py | 15 +- .../models/transformers/multimodal.py | 13 +- vllm/model_executor/models/ultravox.py | 12 +- vllm/model_executor/models/voxtral.py | 2 +- vllm/models/dots3_note/common/processor.py | 12 +- vllm/models/kimi_k3/common/mm_preprocess.py | 13 +- vllm/multimodal/processing/__init__.py | 2 + vllm/multimodal/processing/processor.py | 532 ++++++++---------- vllm/renderers/hf.py | 8 +- .../transformers_utils/processors/internvl.py | 31 +- .../processors/nano_nemotron_vl.py | 37 +- vllm/transformers_utils/processors/nvlm_d.py | 8 +- 53 files changed, 830 insertions(+), 854 deletions(-) diff --git a/docs/contributing/model/multimodal.md b/docs/contributing/model/multimodal.md index 36830dd8ea3b..c59c1a80bf06 100644 --- a/docs/contributing/model/multimodal.md +++ b/docs/contributing/model/multimodal.md @@ -532,6 +532,12 @@ return a list of [PromptUpdate][vllm.multimodal.processing.PromptUpdate] instanc Each [PromptUpdate][vllm.multimodal.processing.PromptUpdate] instance specifies an update operation (e.g.: insertion, replacement) performed by the HF processor. +!!! note + The target and content of each update are token sequences. When converting + text to token sequences, remember to encode it with + `cached_encode(..., add_special_tokens=False)`. Otherwise, the updated prompt + may contain duplicated special tokens or fail to match the target. + === "Basic example: LLaVA" Looking at HF's `LlavaProcessor`: diff --git a/tests/models/multimodal/processing/test_openvla.py b/tests/models/multimodal/processing/test_openvla.py index 992608f893cc..36775d53e40f 100644 --- a/tests/models/multimodal/processing/test_openvla.py +++ b/tests/models/multimodal/processing/test_openvla.py @@ -192,11 +192,11 @@ def test_openvla_prompt_update_inserts_image_tokens_after_bos() -> None: assert resolved.modality == "image" assert [ (match.start_idx, match.end_idx) - for match in resolved.iter_matches([1, 10, 11], _FakeTokenizer()) + for match in resolved.iter_token_matches([1, 10, 11]) ] == [(1, 1)] assert content.full == [32000] * 256 - is_embed = content.is_embed(None, content.full) + is_embed = content.is_embed(content.full) assert is_embed.dtype == torch.bool assert is_embed.tolist() == [True] * 256 diff --git a/tests/multimodal/test_cache.py b/tests/multimodal/test_cache.py index c20e554e5794..7c93808f282a 100644 --- a/tests/multimodal/test_cache.py +++ b/tests/multimodal/test_cache.py @@ -93,7 +93,7 @@ def test_cache_item_size(item, expected_size): cache[""] = item assert cache.currsize == expected_size - prompt_update = PromptInsertion("dummy", "target", "insertion").resolve(0) + prompt_update = PromptInsertion("dummy", [0], [1]).resolve(0) cache[""] = MultiModalProcessorCacheItem(item, [prompt_update]) assert cache.currsize == expected_size @@ -151,7 +151,7 @@ def _compare_caches( for item in all_items ] - prompt_update = PromptInsertion("dummy", "target", "insertion").resolve(0) + prompt_update = PromptInsertion("dummy", [0], [1]).resolve(0) for it in range(n_iter): num_items_to_select = rng.randint(0, max_items_per_iter) diff --git a/tests/multimodal/test_processing.py b/tests/multimodal/test_processing.py index 601fe9b8a819..ec6a8f93df22 100644 --- a/tests/multimodal/test_processing.py +++ b/tests/multimodal/test_processing.py @@ -12,18 +12,19 @@ from vllm.multimodal import MULTIMODAL_REGISTRY from vllm.multimodal.processing.context import InputProcessingContext from vllm.multimodal.processing.processor import ( + BaseMultiModalProcessor, PlaceholderFeaturesInfo, PromptIndexTargets, PromptInsertion, PromptReplacement, _apply_matches, _apply_token_matches_with_placeholders, - apply_text_matches, apply_token_matches, find_mm_placeholders, iter_token_matches, replace_token_matches, ) +from vllm.utils.collection_utils import flatten_2d_lists from .utils import random_image @@ -242,7 +243,7 @@ def test_find_token_matches( for key, target in target_by_key.items() } result = { - key: list(update.iter_token_matches(prompt, tokenizer=None)) + key: list(update.iter_token_matches(prompt)) for key, update in prompt_updates.items() } @@ -259,315 +260,6 @@ def test_find_token_matches( } == expected_by_key -@pytest.mark.parametrize( - ("prompt", "target_by_key", "expected_by_key"), - [ - # Detokenized test cases of `test_find_token_matches` - # using the vocab of llava-hf/llava-v1.6-mistral-7b-hf - ( - "", - { - "pattern_1": "", - "pattern_2": "", - "pattern_3": PromptIndexTargets.start(), - "pattern_4": PromptIndexTargets.prefix(""), - "pattern_5": PromptIndexTargets.end(), - }, - { - "pattern_1": [{"start_idx": 0, "end_idx": 0}], - "pattern_2": [], - "pattern_3": [ - {"start_idx": 0, "end_idx": 0}, - ], - "pattern_4": [], - "pattern_5": [ - {"start_idx": 0, "end_idx": 0}, - ], - }, - ), - ( - "", - { - "pattern_1": "", - "pattern_2": "", - "pattern_3": "", - "pattern_4": PromptIndexTargets.start(), - "pattern_5": PromptIndexTargets.prefix(""), - "pattern_6": PromptIndexTargets.end(), - }, - { - "pattern_1": [ - {"start_idx": 0, "end_idx": 7}, - {"start_idx": 7, "end_idx": 14}, - {"start_idx": 14, "end_idx": 21}, - {"start_idx": 21, "end_idx": 28}, - ], - "pattern_2": [ - {"start_idx": 0, "end_idx": 14}, - {"start_idx": 14, "end_idx": 28}, - ], - "pattern_3": [ - {"start_idx": 0, "end_idx": 21}, - ], - "pattern_4": [ - {"start_idx": 0, "end_idx": 0}, - ], - "pattern_5": [ - {"start_idx": 7, "end_idx": 7}, - ], - "pattern_6": [ - {"start_idx": 28, "end_idx": 28}, - ], - }, - ), - ( - "Image:Image:!", - { - "pattern_1": "Image:", - "pattern_2": "Image:", - "pattern_3": "Image:", - "pattern_4": PromptIndexTargets.start(), - "pattern_5": PromptIndexTargets.prefix("Image:"), - "pattern_6": PromptIndexTargets.end(), - }, - { - "pattern_1": [ - {"start_idx": 0, "end_idx": 13}, - {"start_idx": 27, "end_idx": 40}, - ], - "pattern_2": [ - {"start_idx": 0, "end_idx": 27}, - ], - "pattern_3": [], - "pattern_4": [ - {"start_idx": 0, "end_idx": 0}, - ], - "pattern_5": [ - {"start_idx": 13, "end_idx": 13}, - ], - "pattern_6": [ - {"start_idx": 48, "end_idx": 48}, - ], - }, - ), - # Test regex escape - ( - "<|image|><|image|>", - { - "pattern_1": "<|image|>", - "pattern_2": "<|image|>", - "pattern_3": "<|image|><|image|>", - }, - { - "pattern_1": [ - {"start_idx": 0, "end_idx": 9}, - {"start_idx": 16, "end_idx": 25}, - ], - "pattern_2": [ - {"start_idx": 0, "end_idx": 16}, - {"start_idx": 16, "end_idx": 32}, - ], - "pattern_3": [ - {"start_idx": 0, "end_idx": 25}, - ], - }, - ), - ], -) -@pytest.mark.parametrize("update_type", [PromptInsertion, PromptReplacement]) -def test_find_text_matches( - prompt, - target_by_key, - expected_by_key, - update_type, -): - prompt_updates = { - key: update_type(key, target, []).resolve(0) - for key, target in target_by_key.items() - } - result = { - key: list(update.iter_text_matches(prompt, tokenizer=None)) - for key, update in prompt_updates.items() - } - - # Only displayed on error - print("result:", result) - - # Manually constructed results - assert { - key: [ - dict(start_idx=item.start_idx, end_idx=item.end_idx) - for item in result.get(key, []) - ] - for key in expected_by_key - } == expected_by_key - - -@pytest.mark.parametrize( - ("prompt", "target_by_key", "repl_by_key", "expected_by_update_type_mm_count"), # noqa: E501 - [ - ( - "Image:Image:!", - { - # We use `` before `Image:` to test matches that - # occur out of order - "pattern_1": "", - "pattern_2": "Image:", - "pattern_3": "!", - }, - { - # Test whether target is confused with replacement - "pattern_1": "", - # Test empty replacement - "pattern_2": "", - # Test dynamic replacement (beyond the form of `unit * count`) - "pattern_3": "?!?", - }, - { - PromptInsertion: { - 0: "Image:Image:!", - 1: "Image:Image:!?!?", - 2: "Image:Image:!?!??!?", # noqa: E501 - }, - PromptReplacement: { - 0: "Image:Image:!", - 1: "Image:?!?", - 2: "?!?", - }, - }, - ), - # Test index targets - ( - "", - { - "pattern_1": PromptIndexTargets.start(), - "pattern_2": PromptIndexTargets.prefix(""), - "pattern_3": PromptIndexTargets.end(), - }, - { - "pattern_1": "1", - "pattern_2": "2", - "pattern_3": "3", - }, - { - PromptInsertion: { - 0: "", - 1: "13", - 2: "1133", - }, - PromptReplacement: { - 0: "", - 1: "13", - 2: "1133", - }, - }, - ), - ( - "", - { - "pattern_1": PromptIndexTargets.start(), - "pattern_2": PromptIndexTargets.prefix(""), - "pattern_3": PromptIndexTargets.end(), - }, - { - "pattern_1": "1", - "pattern_2": "2", - "pattern_3": "3", - }, - { - PromptInsertion: { - 0: "", - 1: "123", - 2: "112233", - }, - PromptReplacement: { - 0: "", - 1: "123", - 2: "112233", - }, - }, - ), - # Test different replacement per item - ( - "", - { - "pattern_1": "", - }, - { - "pattern_1": lambda idx: str(idx + 1), - }, - { - PromptInsertion: { - 0: "", - 1: "1", - 2: "12", - }, - PromptReplacement: { - 0: "", - 1: "1", - 2: "12", - }, - }, - ), - ( - "", - { - "pattern_1": PromptIndexTargets.prefix(""), - }, - { - "pattern_1": lambda idx: str(idx + 1), - }, - { - PromptInsertion: { - 0: "", - 1: "1", - 2: "12", - }, - PromptReplacement: { - 0: "", - 1: "1", - 2: "12", - }, - }, - ), - ], -) -def test_find_update_text( - prompt, - target_by_key, - repl_by_key, - expected_by_update_type_mm_count, -): - for ( - update_type, - expected_by_mm_count, - ) in expected_by_update_type_mm_count.items(): - for mm_count, expected in expected_by_mm_count.items(): - mm_prompt_updates = { - key: [ - [update_type(key, target, repl_by_key[key]).resolve(i)] - for i in range(mm_count) - ] - for key, target in target_by_key.items() - } - - new_prompt, result = apply_text_matches( - prompt, - mm_prompt_updates, - tokenizer=None, - ) - - # Only displayed on error - print("update_type:", update_type) - print("mm_count:", mm_count) - print("mm_prompt_updates:", mm_prompt_updates) - print("new_prompt:", new_prompt) - print("result:", result) - - # Manually constructed results - assert new_prompt == expected - - FIND_UPDATE_TOKENS_TEST_CASES = [ # Tokenized test cases of `test_find_update_text` # using the vocab of llava-hf/llava-v1.6-mistral-7b-hf @@ -891,11 +583,7 @@ def test_find_update_tokens( for key, target in target_by_key.items() } - new_prompt, result = apply_token_matches( - prompt, - mm_prompt_updates, - tokenizer=None, - ) + new_prompt, result = apply_token_matches(prompt, mm_prompt_updates) # Only displayed on error print("update_type:", update_type) @@ -945,7 +633,6 @@ def test_apply_token_matches_with_placeholders( new_prompt, result, placeholders = _apply_token_matches_with_placeholders( prompt, mm_prompt_updates, - tokenizer=None, ) if any( @@ -1097,7 +784,7 @@ def test_find_mm_placeholders( for key, repl in repl_by_key.items() } - result = find_mm_placeholders(prompt, mm_prompt_updates, tokenizer=None) + result = find_mm_placeholders(prompt, mm_prompt_updates) # Only displayed on error print("result:", result) @@ -1285,24 +972,18 @@ def test_apply_matches_no_match_exits_quickly(): With the fix, it should exit immediately when no match is found. """ # Create a long prompt with no placeholder - long_prompt = "x" * 10000 + long_prompt = [1] * 10000 # Create update looking for a placeholder that doesn't exist - mm_prompt_updates = { - "image": [[PromptReplacement("image", "", "REPLACED").resolve(0)]] - } + mm_prompt_updates = {"image": [[PromptReplacement("image", [0], [-1]).resolve(0)]]} start = time.perf_counter() - result, _ = _apply_matches( - long_prompt, - mm_prompt_updates, - tokenizer=None, - ) + result, _ = _apply_matches(long_prompt, mm_prompt_updates) elapsed = time.perf_counter() - start # Should complete in < 100ms (was taking seconds before the fix) assert elapsed < 0.1, f"_apply_matches took {elapsed:.2f}s, expected < 0.1s" - assert "".join(result) == long_prompt + assert flatten_2d_lists(result) == long_prompt def test_apply_matches_many_shared_targets_scales_linearly(): @@ -1317,11 +998,7 @@ def measure(item_count: int) -> float: prompt = [0] * item_count start = time.perf_counter() - result, match_result = apply_token_matches( - prompt, - mm_prompt_updates, - tokenizer=None, - ) + result, match_result = apply_token_matches(prompt, mm_prompt_updates) elapsed = time.perf_counter() - start assert len(result) == item_count * len(replacement) @@ -1360,7 +1037,7 @@ def test_find_mm_placeholders_avoids_quadratic_false_prefixes(): } start = time.perf_counter() - result = find_mm_placeholders(prompt, mm_prompt_updates, tokenizer=None) + result = find_mm_placeholders(prompt, mm_prompt_updates) elapsed = time.perf_counter() - start assert result == {} @@ -1377,22 +1054,158 @@ def test_find_mm_placeholders_avoids_quadratic_false_prefixes(): [1, 2, 3, 4, 5], ], ) -def test_find_mm_placeholders_resolves_content_lazily(prompt): +def test_find_mm_placeholders_stops_at_missing_item(prompt): """ - Test that content of items the scan never reaches is not resolved. - - With `tokenizer=None`, resolving string content raises; the scan must - return no placeholders instead of raising on the second item. + Test that the scan returns no placeholders once it fails to find + an item's placeholder, leaving later items unresolved. """ result = find_mm_placeholders( prompt, { "image": [ [PromptReplacement("image", [0], [999]).resolve(0)], - [PromptReplacement("image", [0], "never reached").resolve(1)], + [PromptReplacement("image", [0], [998]).resolve(1)], ] }, - tokenizer=None, ) assert result == {} + + +class _FakeTokenizer: + """ + Character-level tokenizer where "foo" merges into one token differently + depending on whether it is followed by "d", like BPE merging "foo" in + "food" across the search-text boundary. + """ + + _MERGES = {"food": (1000,), "foo": (101, 111, 111)} + _INVERSE = {ids: text for text, ids in _MERGES.items()} + + def encode(self, text: str, add_special_tokens: bool = True) -> list[int]: + token_ids = list[int]() + pos = 0 + while pos < len(text): + for length in (4, 3): + word = text[pos : pos + length] + if word in self._MERGES: + token_ids.extend(self._MERGES[word]) + pos += length + break + else: + token_ids.append(ord(text[pos])) + pos += 1 + return token_ids + + def decode( + self, + token_ids: list[int], + skip_special_tokens: bool = False, + ) -> str: + chars = list[str]() + pos = 0 + while pos < len(token_ids): + for length in (3, 1): + key = tuple(token_ids[pos : pos + length]) + if key in self._INVERSE: + chars.append(self._INVERSE[key]) + pos += length + break + else: + chars.append(chr(token_ids[pos])) + pos += 1 + return "".join(chars) + + +class _FakeProcessingInfo: + def __init__(self, tokenizer) -> None: + self._tokenizer = tokenizer + + def get_tokenizer(self): + return self._tokenizer + + +class _TextFallbackProcessor(BaseMultiModalProcessor): + """Only `self.info.get_tokenizer()` is needed by the text fallback.""" + + def __init__(self, tokenizer: _FakeTokenizer) -> None: + self.info = _FakeProcessingInfo(tokenizer) + + def _get_mm_fields_config(self, hf_inputs, hf_processor_mm_kwargs): + raise NotImplementedError + + def _get_prompt_updates(self, mm_items, hf_processor_mm_kwargs, out_mm_kwargs): + raise NotImplementedError + + +def _text_fallback_processor() -> BaseMultiModalProcessor: + return _TextFallbackProcessor(_FakeTokenizer()) + + +def test_apply_prompt_updates_falls_back_to_text_matching(): + """ + Test that the fallback in `_apply_prompt_updates` finds targets that + tokenize differently inside the prompt ("foo" in "food"). + """ + processor = _text_fallback_processor() + + new_token_ids, placeholders = processor._apply_prompt_updates( + [1000], # "food" + { + "image": [ + [PromptReplacement("image", [101, 111, 111], [200, 201]).resolve(0)] + ] + }, + ) + + assert new_token_ids == [200, 201, ord("d")] + assert [p.to_range().offset for p in placeholders["image"]] == [0] + assert [p.tokens for p in placeholders["image"]] == [[200, 201]] + + +def test_apply_prompt_updates_falls_back_with_prefix_target(): + """ + Test that `PromptIndexTargets.prefix` targets are resolved against the + decoded text in the fallback path of `_apply_prompt_updates`. + """ + processor = _text_fallback_processor() + + new_token_ids, placeholders = processor._apply_prompt_updates( + [1000], # "food" + { + "image": [ + [ + PromptInsertion( + "image", + PromptIndexTargets.prefix([101, 111, 111]), + [9], + ).resolve(0) + ] + ] + }, + ) + + assert new_token_ids == [101, 111, 111, 9, ord("d")] + assert [p.tokens for p in placeholders["image"]] == [[9]] + + +def test_apply_prompt_updates_falls_back_with_index_targets(): + """ + Test that the text resolvers of `PromptIndexTargets.start`/`end` + match against the decoded text when another item forces the + fallback in `_apply_prompt_updates`. + """ + processor = _text_fallback_processor() + + new_token_ids, placeholders = processor._apply_prompt_updates( + [1000], # "food" + { + "image": [ + [PromptReplacement("image", [101, 111, 111], [200, 201]).resolve(0)], + [PromptInsertion("image", PromptIndexTargets.end(), [9]).resolve(1)], + ] + }, + ) + + assert new_token_ids == [200, 201, ord("d"), 9] + assert [p.tokens for p in placeholders["image"]] == [[200, 201], [9]] diff --git a/vllm/model_executor/models/audioflamingo3.py b/vllm/model_executor/models/audioflamingo3.py index 43e6212c7aee..7d6e51f4d336 100644 --- a/vllm/model_executor/models/audioflamingo3.py +++ b/vllm/model_executor/models/audioflamingo3.py @@ -483,7 +483,7 @@ def get_replacement_audioflamingo3(item_idx: int): return [ PromptReplacement( modality="audio", - target=audio_token, + target=[audio_token_id], replacement=get_replacement_audioflamingo3, ) ] diff --git a/vllm/model_executor/models/cohere2_vision.py b/vllm/model_executor/models/cohere2_vision.py index 6c13203b2739..42252afd67a2 100644 --- a/vllm/model_executor/models/cohere2_vision.py +++ b/vllm/model_executor/models/cohere2_vision.py @@ -40,6 +40,7 @@ PromptReplacement, PromptUpdate, PromptUpdateDetails, + cached_encode, ) from vllm.sequence import IntermediateTensors from vllm.utils.tensor_schema import TensorSchema, TensorShape @@ -277,10 +278,12 @@ def _get_prompt_updates( ) -> Sequence[PromptUpdate]: hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) image_token = hf_processor.image_token + image_token_id = hf_processor.image_token_id img_tokens_per_tile = int(hf_processor.patch_size**2) img_line_break_token = hf_processor.img_line_break_token boi_token = hf_processor.boi_token eoi_token = hf_processor.eoi_token + tokenizer = self.info.get_tokenizer() def get_replacement(item_idx: int): images = mm_items.get_items("image", ImageProcessorItems) @@ -295,12 +298,13 @@ def get_replacement(item_idx: int): patch_tokens = image_token * img_tokens_per_tile + img_line_break_token repl = f"{boi_token}{patch_tokens * num_patches}{eoi_token}" - return PromptUpdateDetails.select_text(repl, image_token) + repl_ids = cached_encode(tokenizer, repl, add_special_tokens=False) + return PromptUpdateDetails.select_token_id(repl_ids, image_token_id) return [ PromptReplacement( modality="image", - target=image_token, + target=[image_token_id], replacement=get_replacement, ) ] diff --git a/vllm/model_executor/models/ernie45_vl.py b/vllm/model_executor/models/ernie45_vl.py index 96466bad09a5..060392df338f 100644 --- a/vllm/model_executor/models/ernie45_vl.py +++ b/vllm/model_executor/models/ernie45_vl.py @@ -69,6 +69,7 @@ BaseProcessingInfo, PromptReplacement, PromptUpdate, + cached_encode, ) from vllm.sequence import IntermediateTensors from vllm.utils.gpu_sync_debug import gpu_sync_allowed @@ -1197,6 +1198,7 @@ def _get_prompt_updates( out_mm_kwargs: MultiModalKwargsItems, ) -> Sequence[PromptUpdate]: hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) + tokenizer = self.info.get_tokenizer() before_placeholder = { "image": "<|image@placeholder|>", @@ -1223,12 +1225,17 @@ def get_replacement_ernie45vl(item_idx: int, modality: str): ) else: num_tokens = int(grid_thw.prod()) // merge_length - return after_placeholder[modality] * num_tokens + placeholder_ids = cached_encode( + tokenizer, after_placeholder[modality], add_special_tokens=False + ) + return placeholder_ids * num_tokens return [ PromptReplacement( modality=modality, - target=before_placeholder[modality], + target=cached_encode( + tokenizer, before_placeholder[modality], add_special_tokens=False + ), replacement=partial(get_replacement_ernie45vl, modality=modality), ) for modality in ("image", "video") diff --git a/vllm/model_executor/models/funaudiochat.py b/vllm/model_executor/models/funaudiochat.py index 4b8ba82d70e6..607eb4ea5d05 100644 --- a/vllm/model_executor/models/funaudiochat.py +++ b/vllm/model_executor/models/funaudiochat.py @@ -753,7 +753,7 @@ def get_replacement_funaudiochat(item_idx: int): return [ PromptReplacement( modality="audio", - target=audio_token, + target=[audio_token_id], replacement=get_replacement_funaudiochat, ) ] diff --git a/vllm/model_executor/models/gemma3_mm.py b/vllm/model_executor/models/gemma3_mm.py index d698483cb806..080842280126 100644 --- a/vllm/model_executor/models/gemma3_mm.py +++ b/vllm/model_executor/models/gemma3_mm.py @@ -32,6 +32,7 @@ PromptReplacement, PromptUpdate, PromptUpdateDetails, + cached_encode, replace_token_matches, ) from vllm.sequence import IntermediateTensors @@ -165,7 +166,7 @@ def get_image_repl( image_height: int, processor: Gemma3Processor, mm_kwargs: Mapping[str, object], - ) -> PromptUpdateDetails[str]: + ) -> PromptUpdateDetails: boi_token = processor.boi_token num_crops = self.get_num_crops( @@ -189,8 +190,9 @@ def get_image_repl( tokenizer = processor.tokenizer vocab = tokenizer.get_vocab() image_token_id = vocab[tokenizer.image_token] + repl_full_ids = cached_encode(tokenizer, repl_full, add_special_tokens=False) - return PromptUpdateDetails.select_token_id(repl_full, image_token_id) + return PromptUpdateDetails.select_token_id(repl_full_ids, image_token_id) def get_num_image_tokens( self, @@ -313,7 +315,8 @@ def _get_prompt_updates( out_mm_kwargs: MultiModalKwargsItems, ) -> Sequence[PromptUpdate]: hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) - image_token = hf_processor.boi_token + tokenizer = self.info.get_tokenizer() + boi_token_id = tokenizer.get_vocab()[hf_processor.boi_token] def get_replacement_gemma3(item_idx: int): images = mm_items.get_items("image", ImageProcessorItems) @@ -329,7 +332,7 @@ def get_replacement_gemma3(item_idx: int): return [ PromptReplacement( modality="image", - target=image_token, + target=[boi_token_id], replacement=get_replacement_gemma3, ) ] diff --git a/vllm/model_executor/models/gemma3n_mm.py b/vllm/model_executor/models/gemma3n_mm.py index a473d9f437fc..4cbaed8d9be3 100644 --- a/vllm/model_executor/models/gemma3n_mm.py +++ b/vllm/model_executor/models/gemma3n_mm.py @@ -51,6 +51,7 @@ PromptReplacement, PromptUpdate, PromptUpdateDetails, + cached_encode, replace_token_matches, ) from vllm.sequence import IntermediateTensors @@ -183,31 +184,37 @@ def get_image_repl( image_width: int, image_height: int, processor: Gemma3nProcessor, - ) -> str: + ) -> PromptUpdateDetails: """ - Get the replacement text for image tokens. + Get the replacement metadata for image tokens. For Gemma3n, this should return the full_image_sequence which includes BOI token, repeated image tokens, and EOI token. """ + full_token_ids = cached_encode( + processor.tokenizer, processor.full_image_sequence, add_special_tokens=False + ) return PromptUpdateDetails.select_token_id( - processor.full_image_sequence, processor.image_token_id + full_token_ids, processor.image_token_id ) def get_audio_repl( self, *, processor: Gemma3nProcessor, - ) -> str: + ) -> PromptUpdateDetails: """ - Get the replacement text for audio tokens. + Get the replacement metadata for audio tokens. For Gemma3n, this should return the full_audio_sequence which includes BOA token, repeated audio tokens, and EOA token. """ # Return the full audio sequence as defined by the processor + full_token_ids = cached_encode( + processor.tokenizer, processor.full_audio_sequence, add_special_tokens=False + ) return PromptUpdateDetails.select_token_id( - processor.full_audio_sequence, processor.audio_token_id + full_token_ids, processor.audio_token_id ) @@ -345,7 +352,7 @@ def _get_prompt_updates( # Handle image tokens if "image" in mm_items: - image_token = hf_processor.image_token + image_token_id = hf_processor.image_token_id def get_replacement_image(item_idx: int): images = mm_items.get_items("image", ImageProcessorItems) @@ -359,14 +366,14 @@ def get_replacement_image(item_idx: int): prompt_updates.append( PromptReplacement( modality="image", - target=image_token, + target=[image_token_id], replacement=get_replacement_image, ) ) # Handle audio tokens if "audio" in mm_items: - audio_token = hf_processor.audio_token + audio_token_id = hf_processor.audio_token_id def get_replacement_audio(item_idx: int): return self.info.get_audio_repl( @@ -376,7 +383,7 @@ def get_replacement_audio(item_idx: int): prompt_updates.append( PromptReplacement( modality="audio", - target=audio_token, + target=[audio_token_id], replacement=get_replacement_audio, ) ) diff --git a/vllm/model_executor/models/gemma4_mm.py b/vllm/model_executor/models/gemma4_mm.py index 29b9a08aefb8..c636897aaeea 100644 --- a/vllm/model_executor/models/gemma4_mm.py +++ b/vllm/model_executor/models/gemma4_mm.py @@ -332,7 +332,7 @@ def get_image_repl( image_height: int, processor: Gemma4Processor | None, max_soft_tokens: int | None = None, - ) -> PromptUpdateDetails[list[int]]: + ) -> PromptUpdateDetails: """Return the dynamic image token sequence for this image. Computes the exact number of soft tokens the vision tower will @@ -386,7 +386,7 @@ def get_audio_repl( *, audio_len: int, processor: Gemma4Processor | None, - ) -> PromptUpdateDetails[list[int]]: + ) -> PromptUpdateDetails: """Return the dynamic audio token sequence for this audio. Computes the number of soft tokens from the audio waveform @@ -414,7 +414,7 @@ def get_video_repl( timestamps: list[float], num_soft_tokens_per_frame: list[int], processor: Gemma4Processor, - ) -> PromptUpdateDetails[list[int]]: + ) -> PromptUpdateDetails: """Build the full token replacement for one video. Produces the same interleaved sequence as the HF Gemma4Processor: @@ -818,7 +818,7 @@ def _get_prompt_updates( # one image_token exists per image in the token stream. # The replacement expands it to the full image sequence # (boi + N×image_token + eoi, where N = max_soft_tokens). - image_token = hf_processor.image_token + image_token_id = hf_processor.image_token_id def get_replacement_image(item_idx: int): images = mm_items.get_items("image", ImageProcessorItems) @@ -848,13 +848,13 @@ def get_replacement_image(item_idx: int): prompt_updates.append( PromptReplacement( modality="image", - target=image_token, + target=[image_token_id], replacement=get_replacement_image, ) ) if "video" in mm_items: - video_token = hf_processor.video_token + video_token_id = hf_processor.video_token_id def get_replacement_video(item_idx: int): out_item = out_mm_kwargs["video"][item_idx] @@ -869,13 +869,13 @@ def get_replacement_video(item_idx: int): prompt_updates.append( PromptReplacement( modality="video", - target=video_token, + target=[video_token_id], replacement=get_replacement_video, ) ) if "audio" in mm_items: - audio_token = hf_processor.audio_token + audio_token_id = hf_processor.audio_token_id def get_replacement_audio(item_idx: int): audios = mm_items.get_items("audio", AudioProcessorItems) @@ -888,7 +888,7 @@ def get_replacement_audio(item_idx: int): prompt_updates.append( PromptReplacement( modality="audio", - target=audio_token, + target=[audio_token_id], replacement=get_replacement_audio, ) ) diff --git a/vllm/model_executor/models/glm4_1v.py b/vllm/model_executor/models/glm4_1v.py index a4690e9a47d7..8dd0beb66cb4 100644 --- a/vllm/model_executor/models/glm4_1v.py +++ b/vllm/model_executor/models/glm4_1v.py @@ -94,6 +94,7 @@ PromptReplacement, PromptUpdate, PromptUpdateDetails, + cached_encode, ) from vllm.sequence import IntermediateTensors from vllm.transformers_utils.processor import get_processor_cls_name_from_config @@ -1694,6 +1695,8 @@ def _get_prompt_updates( ) -> Sequence[PromptUpdate]: hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) image_processor = self.info.get_image_processor(**hf_processor_mm_kwargs) + tokenizer = self.info.get_tokenizer() + vocab = tokenizer.get_vocab() merge_length = image_processor.merge_size**2 @@ -1722,12 +1725,18 @@ def get_video_replacement(item_idx: int): return [ PromptReplacement( modality="image", - target=hf_processor.image_token, + target=cached_encode( + tokenizer, hf_processor.image_token, add_special_tokens=False + ), replacement=get_image_replacement, ), PromptReplacement( modality="video", - target="<|begin_of_video|><|video|><|end_of_video|>", + target=[ + vocab["<|begin_of_video|>"], + vocab["<|video|>"], + vocab["<|end_of_video|>"], + ], replacement=get_video_replacement, ), ] diff --git a/vllm/model_executor/models/glmasr.py b/vllm/model_executor/models/glmasr.py index fcf17c550ca5..c9d0638e0c91 100644 --- a/vllm/model_executor/models/glmasr.py +++ b/vllm/model_executor/models/glmasr.py @@ -896,7 +896,7 @@ def get_replacement_glmasr(item_idx: int): return [ PromptReplacement( modality="audio", - target=audio_token, + target=[audio_token_id], replacement=get_replacement_glmasr, ) ] diff --git a/vllm/model_executor/models/h2ovl.py b/vllm/model_executor/models/h2ovl.py index dc63276d8269..02e31e719c10 100644 --- a/vllm/model_executor/models/h2ovl.py +++ b/vllm/model_executor/models/h2ovl.py @@ -25,6 +25,7 @@ ProcessorInputs, PromptReplacement, TimingContext, + cached_encode, ) from vllm.transformers_utils.processors.h2ovl import H2OVLImageProcessor, H2OVLProcessor @@ -90,6 +91,8 @@ def _get_prompt_repl_image( hf_processor: H2OVLProcessor, out_mm_data: BatchedTensorInputs, ): + tokenizer = self.info.get_tokenizer() + if "image_num_patches" in out_mm_data: image_num_patches = out_mm_data["image_num_patches"] assert isinstance(image_num_patches, torch.Tensor) @@ -127,7 +130,7 @@ def get_replacement_internvl(item_idx: int): return PromptReplacement( modality="image", - target="", + target=cached_encode(tokenizer, "", add_special_tokens=False), replacement=get_replacement_internvl, ) diff --git a/vllm/model_executor/models/idefics3.py b/vllm/model_executor/models/idefics3.py index fbf0e7a7c3f4..223b881196cd 100644 --- a/vllm/model_executor/models/idefics3.py +++ b/vllm/model_executor/models/idefics3.py @@ -49,6 +49,7 @@ PromptReplacement, PromptUpdate, PromptUpdateDetails, + cached_encode, ) from vllm.sequence import IntermediateTensors from vllm.utils.gpu_sync_debug import gpu_sync_allowed @@ -315,7 +316,8 @@ def _get_prompt_updates( out_mm_kwargs: MultiModalKwargsItems, ) -> Sequence[PromptUpdate]: hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) - image_token, _, _ = self.info._get_image_token(hf_processor) + image_token_id = hf_processor.image_token_id + tokenizer = self.info.get_tokenizer() def get_replacement_idefics3(item_idx: int) -> PromptUpdateDetails: images = mm_items.get_items("image", ImageProcessorItems) @@ -329,15 +331,18 @@ def get_replacement_idefics3(item_idx: int) -> PromptUpdateDetails: mm_kwargs=hf_processor_mm_kwargs, ) - return PromptUpdateDetails.select_text( - image_repl, - embed_text=image_token, + image_repl_ids = cached_encode( + tokenizer, image_repl, add_special_tokens=False + ) + return PromptUpdateDetails.select_token_id( + image_repl_ids, + image_token_id, ) return [ PromptReplacement( modality="image", - target=image_token, + target=[image_token_id], replacement=get_replacement_idefics3, ) ] diff --git a/vllm/model_executor/models/interns1.py b/vllm/model_executor/models/interns1.py index 2248911b2010..f8062cb751dd 100644 --- a/vllm/model_executor/models/interns1.py +++ b/vllm/model_executor/models/interns1.py @@ -45,6 +45,7 @@ PromptReplacement, PromptUpdate, PromptUpdateDetails, + cached_encode, ) from vllm.sequence import IntermediateTensors from vllm.transformers_utils.processor import cached_video_processor_from_config @@ -355,11 +356,9 @@ def _apply_hf_processor_main( hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) tokenizer = hf_processor.tokenizer - video_token_id = tokenizer.encode( - hf_processor.video_token, add_special_tokens=False - ) - assert len(video_token_id) == 1 - video_token_id = video_token_id[0] + vocab = tokenizer.get_vocab() + + video_token_id = vocab[hf_processor.video_token] prompt_text = re.sub( hf_processor.image_token, "", prompt_text @@ -469,10 +468,16 @@ def _get_prompt_updates( ) -> Sequence[PromptUpdate]: hf_processor = self.info.get_hf_processor(**hf_processor_mm_kwargs) img_context_token = hf_processor.image_token + img_context_token_id = hf_processor.image_token_id start_image_token = hf_processor.start_image_token end_image_token = hf_processor.end_image_token video_token = hf_processor.video_token + tokenizer = self.info.get_tokenizer() + video_token_id = cached_encode(tokenizer, video_token, add_special_tokens=False) + assert len(video_token_id) == 1 + video_token_id = video_token_id[0] + out_mm_data = out_mm_kwargs.get_data() if "video_num_patches" in out_mm_data: video_num_patches = out_mm_data["video_num_patches"] @@ -501,7 +506,12 @@ def get_replacement_interns1_image(item_idx: int): repl_features = img_context_token * feature_size repl_full = start_image_token + repl_features + end_image_token - return PromptUpdateDetails.select_text(repl_full, img_context_token) + repl_full_ids = cached_encode( + tokenizer, repl_full, add_special_tokens=False + ) + return PromptUpdateDetails.select_token_id( + repl_full_ids, img_context_token_id + ) def get_replacement_interns1_video(item_idx: int): num_patches = video_num_patches[item_idx] @@ -512,17 +522,20 @@ def get_replacement_interns1_video(item_idx: int): [f"Frame{i + 1}: {repl_features_with_sep}" for i in range(num_patches)] ) - return PromptUpdateDetails.select_text(repl_full, video_token) + repl_full_ids = cached_encode( + tokenizer, repl_full, add_special_tokens=False + ) + return PromptUpdateDetails.select_token_id(repl_full_ids, video_token_id) return [ PromptReplacement( modality="image", - target=img_context_token, + target=[img_context_token_id], replacement=get_replacement_interns1_image, ), PromptReplacement( modality="video", - target=video_token, + target=[video_token_id], replacement=get_replacement_interns1_video, ), ] diff --git a/vllm/model_executor/models/internvl.py b/vllm/model_executor/models/internvl.py index ae7b05a553c5..54fa28cf5afc 100644 --- a/vllm/model_executor/models/internvl.py +++ b/vllm/model_executor/models/internvl.py @@ -43,6 +43,7 @@ BaseProcessingInfo, PromptReplacement, PromptUpdate, + cached_encode, ) from vllm.sequence import IntermediateTensors from vllm.transformers_utils.processors.internvl import ( @@ -264,6 +265,8 @@ def _get_prompt_repl_image( hf_processor: InternVLProcessor, out_mm_data: BatchedTensorInputs, ): + tokenizer = self.info.get_tokenizer() + if "image_num_patches" in out_mm_data: image_num_patches = out_mm_data["image_num_patches"] assert isinstance(image_num_patches, torch.Tensor) @@ -298,7 +301,7 @@ def get_replacement_internvl(item_idx: int): return PromptReplacement( modality="image", - target="", + target=cached_encode(tokenizer, "", add_special_tokens=False), replacement=get_replacement_internvl, ) @@ -498,6 +501,8 @@ def _get_prompt_repl_video( hf_processor: InternVLProcessor, out_mm_data: BatchedTensorInputs, ): + tokenizer = self.info.get_tokenizer() + if "video_num_patches" in out_mm_data: video_num_patches = out_mm_data["video_num_patches"] assert isinstance(video_num_patches, torch.Tensor) @@ -514,7 +519,7 @@ def get_video_replacement_internvl(item_idx: int): return PromptReplacement( modality="video", - target="