diff --git a/tests/kernels/core/test_apply_rotary_emb.py b/tests/kernels/core/test_apply_rotary_emb.py index 23c722fa5e..5c765bbb4a 100644 --- a/tests/kernels/core/test_apply_rotary_emb.py +++ b/tests/kernels/core/test_apply_rotary_emb.py @@ -46,7 +46,6 @@ def get_test_cases() -> list[RotaryEmbeddingTestCase]: Ernie4_5_VLRotaryEmbedding, ) from vllm.model_executor.layers.rotary_embedding.mrope import MRotaryEmbedding - from vllm.model_executor.layers.rotary_embedding.xdrope import XDRotaryEmbedding common_kwargs = { "head_size": 128, @@ -77,20 +76,6 @@ def get_test_cases() -> list[RotaryEmbeddingTestCase]: expect_forward_native=False, expect_forward=True, ), - # XDRotaryEmbedding tests - RotaryEmbeddingTestCase( - name="XDRotaryEmbedding.forward", - rope_class=XDRotaryEmbedding, - rope_kwargs={ - **common_kwargs, - "scaling_alpha": 1.0, - "xdrope_section": [16, 16, 16, 16], - }, - method_name="forward", - positions_shape=(4, 32), # 4D for P/W/H/T - expect_forward_native=False, - expect_forward=True, - ), # Ernie4_5_VLRotaryEmbedding tests RotaryEmbeddingTestCase( name="Ernie4_5_VLRotaryEmbedding.forward_native", diff --git a/tests/models/transformers/test_backend.py b/tests/models/transformers/test_backend.py index 1eaf21b4cf..fd0d980def 100644 --- a/tests/models/transformers/test_backend.py +++ b/tests/models/transformers/test_backend.py @@ -769,3 +769,22 @@ def test_attention_scale_rejects_unresolvable_expression(): with pytest.raises(ValueError, match="Cannot resolve attention scaling expression"): fuser.scale(Attention()) + + +class _ImageOnlyMRoPEModel: + """`get_rope_index` without `video_grid_thw` or `**kwargs` (e.g. HunYuanVL).""" + + def get_rope_index(self, input_ids, image_grid_thw): + seq_len = input_ids.shape[-1] + positions = torch.arange(seq_len).view(1, 1, -1).expand(4, 1, -1) + return positions, torch.tensor([0]) + + +def test_get_mrope_input_positions_omits_unsupported_grid_kwargs(): + """Optional grids the model can't accept must not be passed when empty.""" + mixin = SimpleNamespace(model=_ImageOnlyMRoPEModel()) + + positions, delta = MultiModalMixin.get_mrope_input_positions(mixin, [1, 2, 3], []) + + assert positions.shape == (4, 3) + assert delta == 0 diff --git a/tests/transformers_utils/test_config.py b/tests/transformers_utils/test_config.py index 727017ef41..ba1f0ef8dd 100644 --- a/tests/transformers_utils/test_config.py +++ b/tests/transformers_utils/test_config.py @@ -18,8 +18,10 @@ from vllm.tokenizers import get_tokenizer from vllm.transformers_utils import config as config_module from vllm.transformers_utils.config import ( get_safetensors_params_metadata, + mrope_num_dims, patch_legacy_rope_type, try_get_generation_config, + uses_mrope, ) from vllm.transformers_utils.configs.glm5_next import ( Glm5NextConfig, @@ -176,3 +178,49 @@ def test_safetensors_metadata_of_repo_without_safetensors(): assert get_safetensors_params_metadata("some/pytorch-only-model") == {} get_safetensors_metadata.assert_called_once() + + +@pytest.mark.parametrize( + ("section_key", "mrope_section", "expected_num_dims"), + [ + ("mrope_section", [16, 24, 24], 3), + ("mrope_section", [16, 16, 16, 16], 4), + # Interleaved M-RoPE takes 2 sections but still consumes 3D positions + ("mrope_section", [32, 32], 3), + # HunYuan-VL checkpoints ship the section under its legacy name + ("xdrope_section", [16, 16, 16, 16], 4), + ], +) +def test_mrope_num_dims(section_key, mrope_section, expected_num_dims): + config = PretrainedConfig() + config.rope_parameters = {"rope_type": "default", section_key: mrope_section} + + assert uses_mrope(config) + assert mrope_num_dims(config) == expected_num_dims + + +@pytest.mark.parametrize("section_name", ["mrope_section", "xdrope_section"]) +def test_mrope_num_dims_from_config_attribute(section_name): + """Some configs expose the section as an attribute rather than under + `rope_parameters`.""" + config = PretrainedConfig() + setattr(config, section_name, [16, 16, 16, 16]) + + assert uses_mrope(config) + assert mrope_num_dims(config) == 4 + + +def test_mrope_num_dims_from_nested_rope_parameters(): + """Sections nested by layer type must be found, not silently defaulted.""" + config = PretrainedConfig() + config.rope_parameters = { + "full_attention": {"mrope_section": [16, 16, 16, 16]}, + "linear_attention": {"rope_type": "default"}, + } + + assert uses_mrope(config) + assert mrope_num_dims(config) == 4 + + +def test_mrope_num_dims_without_mrope(): + assert mrope_num_dims(PretrainedConfig()) == 0 diff --git a/tests/v1/core/test_output.py b/tests/v1/core/test_output.py index 15157e1e07..ee8a4a0e78 100644 --- a/tests/v1/core/test_output.py +++ b/tests/v1/core/test_output.py @@ -105,24 +105,6 @@ def _mm_feature_mixed(offset: int, length: int) -> MultiModalFeatureSpec: ) -def test_strip_covered_mm_data_xdrope() -> None: - """XD-RoPE models (e.g. HunyuanOCR) compute positions the same way, so a - covered item must keep its grid dims: stripping them made the worker index - an empty ``image_grid_thw`` and crash the engine on the second request with - an identical prompt.""" - covered = _mm_feature_mixed(offset=0, length=100) - uncovered = _mm_feature_mixed(offset=300, length=100) - - stripped = strip_covered_mm_data( - [covered, uncovered], num_computed_tokens=250, uses_xdrope=True - ) - - assert stripped[0].data is not None - assert list(stripped[0].data.keys()) == ["image_grid_thw"] - assert stripped[1].data is not None - assert set(stripped[1].data.keys()) == {"pixel_values", "image_grid_thw"} - - def test_strip_covered_mm_data_mrope() -> None: """For M-RoPE models, covered items keep their keep_on_cpu metadata fields (the worker needs them to compute positions); payload fields are dropped.""" diff --git a/vllm/config/model.py b/vllm/config/model.py index 0ebd3dcaba..ad8bb06a08 100644 --- a/vllm/config/model.py +++ b/vllm/config/model.py @@ -38,11 +38,11 @@ from vllm.transformers_utils.config import ( get_sentence_transformer_tokenizer_config, is_encoder_decoder, is_rope_parameters_nested, + mrope_num_dims, try_get_dense_modules, try_get_generation_config, try_get_tokenizer_config, uses_mrope, - uses_xdrope_dim, ) from vllm.transformers_utils.model_arch_config_convertor import ( MODEL_ARCH_CONFIG_CONVERTORS, @@ -1850,12 +1850,8 @@ class ModelConfig: return uses_mrope(self.hf_config) @property - def uses_xdrope_dim(self) -> int: - return uses_xdrope_dim(self.hf_config) - - @property - def uses_xdrope(self) -> bool: - return self.uses_xdrope_dim > 0 + def mrope_num_dims(self) -> int: + return mrope_num_dims(self.hf_config) @property def is_multimodal_model(self) -> bool: diff --git a/vllm/model_executor/layers/rotary_embedding/__init__.py b/vllm/model_executor/layers/rotary_embedding/__init__.py index 0a69b3b3f1..110c9790f0 100644 --- a/vllm/model_executor/layers/rotary_embedding/__init__.py +++ b/vllm/model_executor/layers/rotary_embedding/__init__.py @@ -24,7 +24,6 @@ from .mrope_interleaved import MRotaryEmbeddingInterleaved from .ntk_scaling_rope import NTKScalingRotaryEmbedding from .phi3_long_rope_scaled_rope import Phi3LongRoPEScaledRotaryEmbedding from .telechat3_scaling_rope import TeleChat3RoPEScaledRotaryEmbedding -from .xdrope import XDRotaryEmbedding from .yarn_scaling_rope import YaRNScalingRotaryEmbedding _ROPE_DICT: dict[tuple[Any, ...], RotaryEmbedding] = {} @@ -228,18 +227,6 @@ def get_rope( raise ValueError( "Dynamic rope scaling must contain either 'alpha' or 'factor' field" ) - elif scaling_type == "xdrope": - scaling_alpha = rope_parameters["alpha"] - rotary_emb = XDRotaryEmbedding( - head_size, - rotary_dim, - max_position, - base, - is_neox_style, - scaling_alpha, - dtype, - xdrope_section=rope_parameters["xdrope_section"], - ) elif scaling_type == "yarn": scaling_factor = rope_parameters["factor"] original_max_position = rope_parameters["original_max_position_embeddings"] diff --git a/vllm/model_executor/layers/rotary_embedding/xdrope.py b/vllm/model_executor/layers/rotary_embedding/xdrope.py deleted file mode 100644 index a799b10a19..0000000000 --- a/vllm/model_executor/layers/rotary_embedding/xdrope.py +++ /dev/null @@ -1,172 +0,0 @@ -# SPDX-License-Identifier: Apache-2.0 -# SPDX-FileCopyrightText: Copyright contributors to the vLLM project - -import numpy as np -import torch - -from .dynamic_ntk_alpha_rope import DynamicNTKAlphaRotaryEmbedding - - -class XDRotaryEmbedding(DynamicNTKAlphaRotaryEmbedding): - """DynamicNTKAlphaRotaryEmbedding extended with MultiModal(XD) Sections. - - Based on the original DynamicNTKAlphaRotaryEmbedding implementation. - """ - - def __init__( - self, - head_size: int, - rotary_dim: int, - max_position_embeddings: int, - base: float, - is_neox_style: bool, - scaling_alpha: float, - dtype: torch.dtype, - xdrope_section: list[int], - ) -> None: - self.xdrope_section = xdrope_section - super().__init__( - head_size, - rotary_dim, - max_position_embeddings, - base, - is_neox_style, - scaling_alpha, - dtype, - ) - - def forward_native( - self, - positions: torch.Tensor, - query: torch.Tensor, - key: torch.Tensor | None = None, - offsets: torch.Tensor | None = None, - ) -> tuple[torch.Tensor, torch.Tensor | None]: - """PyTorch-native implementation equivalent to forward(). - - Args: - positions: - [4, num_tokens] (P/W/H/T positions with multimodal inputs) - query: [num_tokens, num_heads * head_size] - key: [num_tokens, num_kv_heads * head_size] - """ - assert positions.ndim == 2 - assert key is not None - - num_tokens = positions.shape[-1] - cos_sin = self.cos_sin_cache[positions] - cos, sin = cos_sin.chunk(2, dim=-1) - cos = torch.cat( - [m[i] for i, m in enumerate(cos.split(self.xdrope_section, dim=-1))], dim=-1 - ) - sin = torch.cat( - [m[i] for i, m in enumerate(sin.split(self.xdrope_section, dim=-1))], dim=-1 - ) - - query_shape = query.shape - query = query.view(num_tokens, -1, self.head_size) - query_rot = query[..., : self.rotary_dim] - query_pass = query[..., self.rotary_dim :] - query_rot = self.apply_rotary_emb.forward_native( - query_rot, - cos, - sin, - ) - query = torch.cat((query_rot, query_pass), dim=-1).reshape(query_shape) - - key_shape = key.shape - key = key.view(num_tokens, -1, self.head_size) - key_rot = key[..., : self.rotary_dim] - key_pass = key[..., self.rotary_dim :] - key_rot = self.apply_rotary_emb.forward_native( - key_rot, - cos, - sin, - ) - key = torch.cat((key_rot, key_pass), dim=-1).reshape(key_shape) - return query, key - - def forward_cuda( - self, - positions: torch.Tensor, - query: torch.Tensor, - key: torch.Tensor | None = None, - offsets: torch.Tensor | None = None, - ) -> tuple[torch.Tensor, torch.Tensor | None]: - """PyTorch-native implementation equivalent to forward(). - - Args: - positions: - [4, num_tokens] (P/W/H/T positions with multimodal inputs) - query: [num_tokens, num_heads * head_size] - key: [num_tokens, num_kv_heads * head_size] - """ - assert positions.ndim == 2 - assert key is not None - - num_tokens = positions.shape[-1] - cos_sin = self.cos_sin_cache[positions] - cos, sin = cos_sin.chunk(2, dim=-1) - cos = torch.cat( - [m[i] for i, m in enumerate(cos.split(self.xdrope_section, dim=-1))], dim=-1 - ) - sin = torch.cat( - [m[i] for i, m in enumerate(sin.split(self.xdrope_section, dim=-1))], dim=-1 - ) - - query_shape = query.shape - query = query.view(num_tokens, -1, self.head_size) - query_rot = query[..., : self.rotary_dim] - query_pass = query[..., self.rotary_dim :] - query_rot = self.apply_rotary_emb( - query_rot, - cos, - sin, - ) - query = torch.cat((query_rot, query_pass), dim=-1).reshape(query_shape) - - key_shape = key.shape - key = key.view(num_tokens, -1, self.head_size) - key_rot = key[..., : self.rotary_dim] - key_pass = key[..., self.rotary_dim :] - key_rot = self.apply_rotary_emb( - key_rot, - cos, - sin, - ) - key = torch.cat((key_rot, key_pass), dim=-1).reshape(key_shape) - return query, key - - def forward_xpu( - self, - positions: torch.Tensor, - query: torch.Tensor, - key: torch.Tensor | None = None, - offsets: torch.Tensor | None = None, - ) -> tuple[torch.Tensor, torch.Tensor | None]: - # No fused XPU kernel for XD-sectioned RoPE; the base - # RotaryEmbedding.forward_xpu calls the generic C++ kernel, which - # rejects the 2-D [4, num_tokens] xdrope positions. Use the native path. - return self.forward_native(positions, query, key, offsets) - - @staticmethod - def get_next_input_positions( - context_len: int, - seq_len: int, - xd_sections: int = 4, - ) -> list[list[int]]: - return [list(range(context_len, seq_len)) for _ in range(xd_sections)] - - @staticmethod - def get_next_input_positions_tensor( - out: np.ndarray, - out_offset: int, - context_len: int, - num_new_tokens: int, - ): - values = np.arange( - context_len, - context_len + num_new_tokens, - dtype=out.dtype, - ) - out[:, out_offset : out_offset + num_new_tokens] = values diff --git a/vllm/model_executor/models/interfaces.py b/vllm/model_executor/models/interfaces.py index a9a9bc2450..9ff3492bfc 100644 --- a/vllm/model_executor/models/interfaces.py +++ b/vllm/model_executor/models/interfaces.py @@ -1783,7 +1783,8 @@ class SupportsMRoPE(Protocol): Returns: Tuple of `(llm_positions, mrope_position_delta)` - - llm_positions: Tensor of shape `[3, num_tokens]` with T/H/W positions + - llm_positions: Tensor of shape `[num_dims, num_tokens]`, one row + per M-RoPE position channel (e.g. T/H/W) - mrope_position_delta: Delta for position calculations """ ... @@ -1803,55 +1804,6 @@ def supports_mrope( return isinstance(model, SupportsMRoPE) -@runtime_checkable -class SupportsXDRoPE(Protocol): - """The interface required for all models that support XD-RoPE.""" - - supports_xdrope: ClassVar[Literal[True]] = True - """ - A flag that indicates this model supports XD-RoPE. - - Note: - There is no need to redefine this flag if this class is in the - XDRope of your model class. - """ - - def get_xdrope_input_positions( - self, - input_tokens: list[int], - mm_features: list["MultiModalFeatureSpec"], - ) -> torch.Tensor: - """ - Get XD-RoPE input positions and delta value for this specific model. - - This method should be implemented by each model that supports XD-RoPE - to provide model-specific logic for computing input positions. - - Args: - input_tokens: List of input token IDs - mm_features: Information about each multi-modal data item - - Returns: - llm_positions: Tensor of shape `[xdrope_dim, num_tokens]` with - 4D(P/W/H/T) or 3D(W/H/T) positions. - """ - ... - - -@overload -def supports_xdrope(model: type[object]) -> TypeIs[type[SupportsXDRoPE]]: ... - - -@overload -def supports_xdrope(model: object) -> TypeIs[SupportsXDRoPE]: ... - - -def supports_xdrope( - model: type[object] | object, -) -> TypeIs[type[SupportsXDRoPE]] | TypeIs[SupportsXDRoPE]: - return isinstance(model, SupportsXDRoPE) - - @runtime_checkable class SupportsEncoderCudaGraph(Protocol): """Interface for models whose vision encoder supports CUDA graph diff --git a/vllm/model_executor/models/transformers/multimodal.py b/vllm/model_executor/models/transformers/multimodal.py index befc2ffffe..1cdbadb34b 100644 --- a/vllm/model_executor/models/transformers/multimodal.py +++ b/vllm/model_executor/models/transformers/multimodal.py @@ -1236,17 +1236,31 @@ class MultiModalMixin(SupportsMultiModal, SupportsMRoPE): video_grid_thw = torch.stack(video_grid_thw) if video_grid_thw else None # `get_rope_index` doesn't always accept arbitrary `kwargs` - kwargs = {} - if not hasattr(self, "_get_rope_index_accepts_mm_token_type_ids"): + if not hasattr(self, "_get_rope_index_kwarg_names"): import inspect - sig = inspect.signature(self.model.get_rope_index) - params = sig.parameters - self._get_rope_index_accepts_mm_token_type_ids = ( - "mm_token_type_ids" in params - or any(p.kind == inspect.Parameter.VAR_KEYWORD for p in params.values()) + params = inspect.signature(self.model.get_rope_index).parameters + self._get_rope_index_kwarg_names = ( + None + if any(p.kind == inspect.Parameter.VAR_KEYWORD for p in params.values()) + else frozenset(params) ) - if self._get_rope_index_accepts_mm_token_type_ids: + + kwarg_names = self._get_rope_index_kwarg_names + + def accepts_kwarg(name: str) -> bool: + return kwarg_names is None or name in kwarg_names + + # Drop a grid the model can't accept only when there is nothing to pass. + kwargs = { + name: value + for name, value in ( + ("image_grid_thw", image_grid_thw), + ("video_grid_thw", video_grid_thw), + ) + if value is not None or accepts_kwarg(name) + } + if accepts_kwarg("mm_token_type_ids"): mm_token_type_ids = torch.zeros(len(input_tokens), dtype=torch.int) for feature in mm_features: position = feature.mm_position @@ -1257,8 +1271,6 @@ class MultiModalMixin(SupportsMultiModal, SupportsMRoPE): mrope_positions, mrope_position_delta = self.model.get_rope_index( input_ids=torch.tensor(input_tokens).unsqueeze(0), - image_grid_thw=image_grid_thw, - video_grid_thw=video_grid_thw, **kwargs, ) diff --git a/vllm/multimodal/utils.py b/vllm/multimodal/utils.py index 35ba5844a4..4de183e0dc 100644 --- a/vllm/multimodal/utils.py +++ b/vllm/multimodal/utils.py @@ -216,12 +216,11 @@ def strip_covered_mm_data( mm_features: list[MultiModalFeatureSpec], num_computed_tokens: int, uses_mrope: bool = False, - uses_xdrope: bool = False, ) -> list[MultiModalFeatureSpec]: """Drop the tensor data of mm items whose placeholder span is fully inside a prefix-cache-covered region: no encoder run can be scheduled for them, - so the workers never consume the payload fields. M-RoPE and XD-RoPE models - keep CPU-side metadata fields used to compute positions. SHM address items + so the workers never consume the payload fields. M-RoPE models keep + CPU-side metadata fields used to compute positions. SHM address items are also kept so workers can balance the sender's reference count. The scheduler-side ``Request`` keeps the full features.""" if not mm_features or num_computed_tokens == 0: @@ -239,7 +238,7 @@ def strip_covered_mm_data( return f data = None - if uses_mrope or uses_xdrope: + if uses_mrope: data = MultiModalKwargsItem( {k: elem for k, elem in f.data.items() if elem.field.keep_on_cpu} ) diff --git a/vllm/transformers_utils/config.py b/vllm/transformers_utils/config.py index 90d1e3c176..d82e0d000a 100644 --- a/vllm/transformers_utils/config.py +++ b/vllm/transformers_utils/config.py @@ -2,7 +2,7 @@ # SPDX-FileCopyrightText: Copyright contributors to the vLLM project import os -from collections.abc import Callable, Iterator +from collections.abc import Callable, Iterator, Sequence from contextlib import contextmanager from dataclasses import asdict from functools import cache, partial, wraps @@ -569,15 +569,47 @@ def patch_rope_parameters(config: PretrainedConfig) -> None: config.validate_rope() -def _uses_mrope(config: PretrainedConfig) -> bool: +def _iter_rope_parameters(config: PretrainedConfig) -> Iterator[dict[str, Any]]: + """Yield a config's rope parameters, one dict per layer type if nested.""" rope_parameters = getattr(config, "rope_parameters", None) - if rope_parameters is None: - return False + if not isinstance(rope_parameters, dict): + return - return "mrope_section" in rope_parameters or any( - isinstance(params, dict) and "mrope_section" in params - for params in rope_parameters.values() - ) + if is_rope_parameters_nested(rope_parameters): + yield from (p for p in rope_parameters.values() if isinstance(p, dict)) + else: + yield rope_parameters + + +def _mrope_section(config: PretrainedConfig) -> Sequence[int] | None: + """Return the M-RoPE section this config declares, if any. + + `xdrope_section` is the legacy name HunYuan-VL checkpoints use for the same + field; upstream Transformers normalises it to `mrope_section`. + """ + from vllm.config.utils import getattr_iter + + names = ("mrope_section", "xdrope_section") + + for params in _iter_rope_parameters(config): + for i, name in enumerate(names): + section = params.get(name) + if isinstance(section, (list, tuple)): + if i > 0: + logger.warning_once( + "rope_parameters contains a deprecated key '%s'. " + "Please use the preferred key '%s' instead.", + name, + names[0], + ) + return section + + section = getattr_iter(config, names, None, warn=True) + return section if isinstance(section, (list, tuple)) else None + + +def _uses_mrope(config: PretrainedConfig) -> bool: + return _mrope_section(config) is not None def uses_mrope(config: PretrainedConfig) -> bool: @@ -602,21 +634,23 @@ def thinker_uses_mrope(config: PretrainedConfig) -> bool: return uses_mrope(thinker_text_config) -def uses_xdrope_dim(config: PretrainedConfig) -> int: - """Detect if the model with this config uses XD-ROPE.""" - xdrope_section = getattr(config, "xdrope_section", None) - if xdrope_section is not None and isinstance(xdrope_section, list): - return len(xdrope_section) - rope_scaling = getattr(config, "rope_scaling", None) - if rope_scaling is None: +def mrope_num_dims(config: PretrainedConfig) -> int: + """Number of M-RoPE position channels the model consumes. + + Each section entry sizes one channel, so the section length is the channel + count. Interleaved M-RoPE also accepts a 2 section variant whose positions + are still 3D, so never return fewer than 3. + """ + if not uses_mrope(config): return 0 - if isinstance(rope_scaling, dict) and "xdrope_section" in rope_scaling: - xdrope_section = rope_scaling["xdrope_section"] - if xdrope_section is not None and isinstance(xdrope_section, list): - return len(xdrope_section) + for candidate in (config.get_text_config(), config): + mrope_section = _mrope_section(candidate) + if mrope_section is not None: + return max(len(mrope_section), 3) - return 0 + # Custom configs may declare M-RoPE without exposing the sections. + return 3 def is_encoder_decoder(config: PretrainedConfig) -> bool: diff --git a/vllm/v1/core/sched/output.py b/vllm/v1/core/sched/output.py index 77904f880a..93fc2d83c1 100644 --- a/vllm/v1/core/sched/output.py +++ b/vllm/v1/core/sched/output.py @@ -56,7 +56,6 @@ class NewRequestData: block_ids: tuple[list[int], ...], prefill_token_ids: list[int] | None = None, uses_mrope: bool = False, - uses_xdrope: bool = False, ) -> "NewRequestData": return cls( req_id=request.request_id, @@ -65,7 +64,6 @@ class NewRequestData: request.mm_features, request.num_computed_tokens, uses_mrope=uses_mrope, - uses_xdrope=uses_xdrope, ), sampling_params=request.sampling_params, pooling_params=request.pooling_params, diff --git a/vllm/v1/core/sched/scheduler.py b/vllm/v1/core/sched/scheduler.py index 838358a364..0142321299 100644 --- a/vllm/v1/core/sched/scheduler.py +++ b/vllm/v1/core/sched/scheduler.py @@ -94,7 +94,6 @@ class Scheduler(SchedulerInterface): self.cache_config = vllm_config.cache_config self.lora_config = vllm_config.lora_config self.model_uses_mrope = vllm_config.model_config.uses_mrope - self.model_uses_xdrope = vllm_config.model_config.uses_xdrope self.kv_cache_config = kv_cache_config self.kv_events_config = vllm_config.kv_events_config self.parallel_config = vllm_config.parallel_config @@ -1324,7 +1323,6 @@ class Scheduler(SchedulerInterface): req_to_new_blocks[req.request_id].get_block_ids(), req._all_token_ids, uses_mrope=self.model_uses_mrope, - uses_xdrope=self.model_uses_xdrope, ) for req in scheduled_new_reqs ] @@ -1334,7 +1332,6 @@ class Scheduler(SchedulerInterface): req, req_to_new_blocks[req.request_id].get_block_ids(), uses_mrope=self.model_uses_mrope, - uses_xdrope=self.model_uses_xdrope, ) for req in scheduled_new_reqs ] diff --git a/vllm/v1/spec_decode/llm_base_proposer.py b/vllm/v1/spec_decode/llm_base_proposer.py index 1556ddb441..1b30118069 100644 --- a/vllm/v1/spec_decode/llm_base_proposer.py +++ b/vllm/v1/spec_decode/llm_base_proposer.py @@ -173,25 +173,17 @@ class SpecDecodeBaseProposer: # Use draft model's M-RoPE setting, not target model's # Draft models may be text-only even if target is multimodal self.uses_mrope = self.draft_model_config.uses_mrope - self.uses_xdrope_dim = self.vllm_config.model_config.uses_xdrope_dim - self.draft_uses_xdrope_dim = self.draft_model_config.uses_xdrope_dim if self.uses_mrope: # NOTE: `mrope_positions` is implemented with one additional dummy # position on purpose to make it non-contiguous so that it can work # with torch compile. # See detailed explanation in https://github.com/vllm-project/vllm/pull/12128#discussion_r1926431923 - # NOTE: When M-RoPE is enabled, position ids are 3D regardless of - # the modality of inputs. For text-only inputs, each dimension has - # identical position IDs, making M-RoPE functionally equivalent to - # 1D-RoPE. + # NOTE: For text-only inputs, each dimension has identical position + # IDs, making M-RoPE functionally equivalent to 1D-RoPE. # See page 5 of https://arxiv.org/abs/2409.12191 self.mrope_positions = torch.zeros( - (3, self.max_positions + 1), dtype=torch.int64, device=device - ) - elif self.uses_xdrope_dim > 0 and self.draft_uses_xdrope_dim > 0: - self.xdrope_positions = torch.zeros( - (self.uses_xdrope_dim, self.max_positions + 1), + (self.draft_model_config.mrope_num_dims, self.max_positions + 1), dtype=torch.int64, device=device, ) @@ -381,15 +373,11 @@ class SpecDecodeBaseProposer: def _get_positions(self, num_tokens: int): if self.uses_mrope: return self.mrope_positions[:, :num_tokens] - if self.uses_xdrope_dim > 0 and self.draft_uses_xdrope_dim > 0: - return self.xdrope_positions[:, :num_tokens] return self.positions[:num_tokens] def _set_positions(self, num_tokens: int, positions: torch.Tensor): if self.uses_mrope: self.mrope_positions[:, :num_tokens] = positions - elif self.uses_xdrope_dim > 0 and self.draft_uses_xdrope_dim > 0: - self.xdrope_positions[:, :num_tokens] = positions else: # Convert M-RoPE positions if target model uses M-RoPE # but draft doesn't, For text inputs, all M-RoPE @@ -787,8 +775,6 @@ class SpecDecodeBaseProposer: positions_1d = positions[0] if self.uses_mrope else positions if self.uses_mrope: out_pos = self.mrope_positions[0, :batch_size] - elif self.uses_xdrope_dim > 0 and self.draft_uses_xdrope_dim > 0: - out_pos = self.xdrope_positions[0, :batch_size] else: out_pos = self.positions[:batch_size] eagle_step_update_slot_mapping_and_metadata( @@ -805,11 +791,6 @@ class SpecDecodeBaseProposer: if self.uses_mrope: self.mrope_positions[1:, :batch_size] = self.mrope_positions[0, :batch_size] positions = self.mrope_positions[:, :batch_size] - elif self.uses_xdrope_dim > 0 and self.draft_uses_xdrope_dim > 0: - self.xdrope_positions[1:, :batch_size] = self.xdrope_positions[ - 0, :batch_size - ] - positions = self.xdrope_positions[0, :batch_size] else: positions = self.positions[:batch_size] common_attn_metadata.max_seq_len = min( @@ -859,8 +840,6 @@ class SpecDecodeBaseProposer: self.input_ids[token_indices_to_sample] = next_token_ids # copy inputs to buffer for cudagraph - if self.uses_xdrope_dim > 0 and self.draft_uses_xdrope_dim == 0: - target_positions = target_positions[0] self._set_positions(num_tokens, target_positions) self.hidden_states[:num_tokens] = target_hidden_states diff --git a/vllm/v1/worker/gpu/mm/rope.py b/vllm/v1/worker/gpu/mm/rope.py index e5de282234..637ad41400 100644 --- a/vllm/v1/worker/gpu/mm/rope.py +++ b/vllm/v1/worker/gpu/mm/rope.py @@ -6,39 +6,36 @@ import torch import torch.nn as nn from vllm.config import ModelConfig -from vllm.model_executor.models.interfaces import SupportsMRoPE, SupportsXDRoPE +from vllm.model_executor.models.interfaces import SupportsMRoPE from vllm.triton_utils import tl, triton from vllm.v1.worker.gpu.buffer_utils import StagedWriteTensor, UvaBackedTensor class RopeState: - """Unified state for multi-dimensional RoPE variants (M-RoPE, XD-RoPE). + """State for multi-dimensional (M-RoPE) positions. - M-RoPE: 3 dims, uses position delta for decode. - XD-RoPE: 3 or 4 dims, delta is 0 (decode uses orig_pos for all dims). + `num_dims` is the number of position channels the model consumes, which + the model config derives from its M-RoPE sections. NOTE: `positions` is implemented with one additional dummy position on purpose to make it non-contiguous so that it can work with torch compile. See detailed explanation in https://github.com/vllm-project/vllm/pull/12128#discussion_r1926431923 - NOTE: When M-RoPE is enabled, position ids are 3D regardless of the - modality of inputs. For text-only inputs, each dimension has identical - position IDs, making M-RoPE functionally equivalent to 1D-RoPE. + NOTE: For text-only inputs, each dimension has identical position IDs, + making M-RoPE functionally equivalent to 1D-RoPE. See page 5 of https://arxiv.org/abs/2409.12191 """ def __init__( self, num_dims: int, - has_delta: bool, max_num_reqs: int, max_num_tokens: int, max_model_len: int, device: torch.device, ): self.num_dims = num_dims - self.has_delta = has_delta self.max_num_reqs = max_num_reqs self.max_num_tokens = max_num_tokens self.max_model_len = max_model_len @@ -56,7 +53,6 @@ class RopeState: (num_dims, max_num_tokens + 1), dtype=torch.int64, device=device ) - # Delta is non-zero for M-RoPE, always 0 for XD-RoPE. self.prefill_delta = UvaBackedTensor(max_num_reqs, dtype=torch.int32) def init_prefill_positions( @@ -66,17 +62,11 @@ class RopeState: prefill_token_ids: list[int], mm_features: list, ) -> None: - if self.has_delta: - mrope_model = cast(SupportsMRoPE, model) - prefill_positions, delta = mrope_model.get_mrope_input_positions( - prefill_token_ids, mm_features - ) - self.prefill_delta.np[req_idx] = delta - else: - xdrope_model = cast(SupportsXDRoPE, model) - prefill_positions = xdrope_model.get_xdrope_input_positions( - prefill_token_ids, mm_features - ) + mrope_model = cast(SupportsMRoPE, model) + prefill_positions, delta = mrope_model.get_mrope_input_positions( + prefill_token_ids, mm_features + ) + self.prefill_delta.np[req_idx] = delta for i in range(self.num_dims): pos = prefill_positions[i].tolist() @@ -84,8 +74,7 @@ class RopeState: def apply_staged_writes(self) -> None: self.prefill_positions.apply_write() - if self.has_delta: - self.prefill_delta.copy_to_uva() + self.prefill_delta.copy_to_uva() def get_positions(self, num_tokens: int) -> torch.Tensor: return self.positions[:, :num_tokens] @@ -104,8 +93,7 @@ class RopeState: self.prefill_positions.gpu[base : base + self.num_dims, :length].copy_( positions ) - if self.has_delta: - self.prefill_delta.np[req_idx] = delta + self.prefill_delta.np[req_idx] = delta def prepare_positions( self, @@ -140,27 +128,17 @@ def get_rope_state( device: torch.device, ) -> RopeState | None: """Create a RopeState if the model uses multi-dimensional RoPE.""" - if model_config.uses_mrope: - assert isinstance(model, SupportsMRoPE) - return RopeState( - num_dims=3, - has_delta=True, - max_num_reqs=max_num_reqs, - max_num_tokens=max_num_tokens, - max_model_len=max_model_len, - device=device, - ) - if model_config.uses_xdrope_dim > 0: - assert isinstance(model, SupportsXDRoPE) - return RopeState( - num_dims=model_config.uses_xdrope_dim, - has_delta=False, - max_num_reqs=max_num_reqs, - max_num_tokens=max_num_tokens, - max_model_len=max_model_len, - device=device, - ) - return None + if not model_config.uses_mrope: + return None + + assert isinstance(model, SupportsMRoPE) + return RopeState( + num_dims=model_config.mrope_num_dims, + max_num_reqs=max_num_reqs, + max_num_tokens=max_num_tokens, + max_model_len=max_model_len, + device=device, + ) @triton.jit diff --git a/vllm/v1/worker/gpu/model_states/mm_pruning.py b/vllm/v1/worker/gpu/model_states/mm_pruning.py index 8bdb85d4a1..e73a7c0ad3 100644 --- a/vllm/v1/worker/gpu/model_states/mm_pruning.py +++ b/vllm/v1/worker/gpu/model_states/mm_pruning.py @@ -132,7 +132,6 @@ def maybe_create_mm_pruner( """Create a MultiModalPruner if the model prunes embeddings and uses M-RoPE.""" if ( rope_state is None - or not rope_state.has_delta or encoder_cache is None or model_config.multimodal_config is None or not model_config.multimodal_config.is_multimodal_pruning_enabled() diff --git a/vllm/v1/worker/gpu_input_batch.py b/vllm/v1/worker/gpu_input_batch.py index 7585f4d892..01b6b5b0cc 100644 --- a/vllm/v1/worker/gpu_input_batch.py +++ b/vllm/v1/worker/gpu_input_batch.py @@ -46,8 +46,6 @@ class CachedRequestState: mrope_positions: torch.Tensor | None = None mrope_position_delta: int | None = None - xdrope_positions: torch.Tensor | None = None - lora_request: LoRARequest | None = None prompt_embeds: torch.Tensor | None = None # To accumulate prompt logprobs tensor chunks across prefill steps. diff --git a/vllm/v1/worker/gpu_model_runner.py b/vllm/v1/worker/gpu_model_runner.py index 6d35dc73ca..d6b8ed025c 100644 --- a/vllm/v1/worker/gpu_model_runner.py +++ b/vllm/v1/worker/gpu_model_runner.py @@ -70,10 +70,7 @@ from vllm.model_executor.layers.mamba.mamba_utils import MambaStateCopyFuncsByTy from vllm.model_executor.layers.mamba.ops.ssu_dispatch import ( initialize_mamba_ssu_backend, ) -from vllm.model_executor.layers.rotary_embedding import ( - MRotaryEmbedding, - XDRotaryEmbedding, -) +from vllm.model_executor.layers.rotary_embedding import MRotaryEmbedding from vllm.model_executor.model_loader import get_model_loader from vllm.model_executor.model_loader.reload import ( finalize_layerwise_reload, @@ -84,14 +81,12 @@ from vllm.model_executor.models.interfaces import ( MultiModalEmbeddings, SupportsMRoPE, SupportsMultiModal, - SupportsXDRoPE, get_mixture_of_experts_model, supports_eagle3, supports_mrope, supports_multimodal_pruning, supports_realtime, supports_transcription, - supports_xdrope, ) from vllm.model_executor.models.interfaces_base import ( VllmModelForPooling, @@ -576,7 +571,7 @@ class GPUModelRunner( # Multi-modal data support self.mm_registry = MULTIMODAL_REGISTRY self.uses_mrope = model_config.uses_mrope - self.uses_xdrope_dim = model_config.uses_xdrope_dim + self.mrope_num_dims = model_config.mrope_num_dims self.supports_mm_inputs = self.mm_registry.supports_multimodal_inputs( model_config ) @@ -872,20 +867,11 @@ class GPUModelRunner( # with torch compile. # See detailed explanation in https://github.com/vllm-project/vllm/pull/12128#discussion_r1926431923 - # NOTE: When M-RoPE is enabled, position ids are 3D regardless of - # the modality of inputs. For text-only inputs, each dimension has - # identical position IDs, making M-RoPE functionally equivalent to - # 1D-RoPE. + # NOTE: For text-only inputs, each dimension has identical position + # IDs, making M-RoPE functionally equivalent to 1D-RoPE. # See page 5 of https://arxiv.org/abs/2409.12191 self.mrope_positions = self._make_buffer( - (3, self.max_num_tokens + 1), dtype=torch.int64 - ) - - # Only relevant for models using XD-RoPE (e.g, HunYuan-VL) - if self.uses_xdrope_dim > 0: - # Similar to mrope but use assigned dimension number for RoPE, 4 as default. - self.xdrope_positions = self._make_buffer( - (self.uses_xdrope_dim, self.max_num_tokens + 1), dtype=torch.int64 + (self.mrope_num_dims, self.max_num_tokens + 1), dtype=torch.int64 ) # None in the first PP rank. The rest are set after load_model. @@ -1036,14 +1022,10 @@ class GPUModelRunner( if isinstance(num_tokens, int): if self.uses_mrope: return self.mrope_positions.gpu[:, :num_tokens] - if self.uses_xdrope_dim > 0: - return self.xdrope_positions.gpu[:, :num_tokens] return self.positions[:num_tokens] else: if self.uses_mrope: return self.mrope_positions.gpu[:, num_tokens] - if self.uses_xdrope_dim > 0: - return self.xdrope_positions.gpu[:, num_tokens] return self.positions[num_tokens] def _make_buffer( @@ -1332,10 +1314,6 @@ class GPUModelRunner( if self.uses_mrope: self._init_mrope_positions(req_state) - # Only relevant for models using XD-RoPE (e.g, HunYuan-VL) - if self.uses_xdrope_dim > 0: - self._init_xdrope_positions(req_state) - reqs_to_add.append(req_state) # Track new requests for ngram_gpu full tensor copy if is_ngram_gpu: @@ -1702,19 +1680,6 @@ class GPUModelRunner( ) ) - def _init_xdrope_positions(self, req_state: CachedRequestState): - model = self.get_model() - xdrope_model = cast(SupportsXDRoPE, model) - assert req_state.prompt_token_ids is not None, ( - "XD-RoPE requires prompt_token_ids to be available." - ) - assert supports_xdrope(model), "XD-RoPE support is not implemented." - - req_state.xdrope_positions = xdrope_model.get_xdrope_input_positions( - req_state.prompt_token_ids, - req_state.mm_features, - ) - def _extract_mm_kwargs( self, scheduler_output: "SchedulerOutput", @@ -2017,11 +1982,6 @@ class GPUModelRunner( if self.uses_mrope: self._calc_mrope_positions(scheduler_output) - # Calculate XD-RoPE positions. - # Only relevant for models using XD-RoPE (e.g, HunYuan-VL) - if self.uses_xdrope_dim > 0: - self._calc_xdrope_positions(scheduler_output) - # Get token indices. # E.g., [0, 1, 0, 1, 2, 3, 4, 0, 1, 2] # -> [0, 1, M, M + 1, M + 2, M + 3, M + 4, 2 * M, 2 * M + 1, 2 * M + 2] @@ -2229,7 +2189,7 @@ class GPUModelRunner( if self.uses_mrope: # Only relevant for models using M-RoPE (e.g, Qwen2-VL) # Copy one row at a time. mrope_positions is allocated as - # [3, max_num_tokens + 1] with a dummy trailing column to keep it + # [num_dims, max_num_tokens + 1] with a dummy trailing column to keep it # non-contiguous for torch.compile, so cpu[:, :N] is a strided view. # copy_() cannot express a strided source as a single # cudaMemcpyAsync, so it first gathers into a contiguous *pageable* @@ -2242,20 +2202,7 @@ class GPUModelRunner( self.mrope_positions.cpu[row, :total_num_scheduled_tokens], non_blocking=True, ) - elif self.uses_xdrope_dim > 0: - # Only relevant for models using XD-RoPE (e.g, HunYuan-VL). - # xdrope_positions is allocated as [uses_xdrope_dim, max_num_tokens - # + 1] with the same trailing-column trick as mrope_positions above, - # so cpu[:, :N] is a strided view of the pinned buffer and the - # single-slice copy_() runs into the same pageable-fallback silent - # sync described in PR #51841. Split into per-row copies for the - # same reason. - for row in range(self.xdrope_positions.gpu.shape[0]): - self.xdrope_positions.gpu[row, :total_num_scheduled_tokens].copy_( - self.xdrope_positions.cpu[row, :total_num_scheduled_tokens], - non_blocking=True, - ) - if self.use_async_spec_decode and (self.uses_mrope or self.uses_xdrope_dim > 0): + if self.use_async_spec_decode and self.uses_mrope: drift = self.num_computed_tokens[req_indices_gpu].to( torch.int64 ) - async_tensor_h2d( @@ -2263,8 +2210,7 @@ class GPUModelRunner( device=self.device, dtype=torch.int64, ) - target = self.mrope_positions if self.uses_mrope else self.xdrope_positions - target.gpu[:, :total_num_scheduled_tokens] += drift + self.mrope_positions.gpu[:, :total_num_scheduled_tokens] += drift use_spec_decode = len(scheduler_output.scheduled_spec_decode_tokens) > 0 if not use_spec_decode: @@ -2862,53 +2808,6 @@ class GPUModelRunner( mrope_pos_ptr += completion_part_len - def _calc_xdrope_positions(self, scheduler_output: "SchedulerOutput"): - xdrope_pos_ptr = 0 - for index, req_id in enumerate(self.input_batch.req_ids): - req = self.requests[req_id] - assert req.xdrope_positions is not None - - num_computed_tokens = self.input_batch.num_computed_tokens_cpu[index] - num_scheduled_tokens = scheduler_output.num_scheduled_tokens[req_id] - num_prompt_tokens = length_from_prompt_token_ids_or_embeds( - req.prompt_token_ids, req.prompt_embeds - ) - - if num_computed_tokens + num_scheduled_tokens > num_prompt_tokens: - prompt_part_len = max(0, num_prompt_tokens - num_computed_tokens) - completion_part_len = max(0, num_scheduled_tokens - prompt_part_len) - else: - prompt_part_len = num_scheduled_tokens - completion_part_len = 0 - - assert num_scheduled_tokens == prompt_part_len + completion_part_len - - if prompt_part_len > 0: - # prompt's xdrope_positions are pre-computed - dst_start = xdrope_pos_ptr - dst_end = xdrope_pos_ptr + prompt_part_len - src_start = num_computed_tokens - src_end = num_computed_tokens + prompt_part_len - - self.xdrope_positions.cpu[:, dst_start:dst_end] = req.xdrope_positions[ - :, src_start:src_end - ] - xdrope_pos_ptr += prompt_part_len - - if completion_part_len > 0: - # compute completion's xdrope_positions on-the-fly - dst_start = xdrope_pos_ptr - dst_end = xdrope_pos_ptr + completion_part_len - - XDRotaryEmbedding.get_next_input_positions_tensor( - out=self.xdrope_positions.np, - out_offset=dst_start, - context_len=num_computed_tokens + prompt_part_len, - num_new_tokens=completion_part_len, - ) - - xdrope_pos_ptr += completion_part_len - def _calc_spec_decode_metadata( self, num_draft_tokens: np.ndarray, @@ -3303,7 +3202,6 @@ class GPUModelRunner( req_start_idx = 0 should_sync_mrope_positions = False - should_sync_xdrope_positions = False for req_id in self.input_batch.req_ids: mm_embeds_req: list[torch.Tensor] = [] @@ -3396,10 +3294,6 @@ class GPUModelRunner( self._calc_mrope_positions(scheduler_output) self.mrope_positions.copy_to_gpu(total_num_scheduled_tokens) - if should_sync_xdrope_positions: - self._calc_xdrope_positions(scheduler_output) - self.xdrope_positions.copy_to_gpu(total_num_scheduled_tokens) - return mm_embeds, is_mm_embed def get_model(self) -> nn.Module: @@ -3709,8 +3603,6 @@ class GPUModelRunner( if self.uses_mrope: positions = self.mrope_positions.gpu[:, :num_input_tokens] - elif self.uses_xdrope_dim > 0: - positions = self.xdrope_positions.gpu[:, :num_input_tokens] else: positions = self.positions[:num_input_tokens] if num_input_tokens > num_scheduled_tokens: @@ -6202,8 +6094,6 @@ class GPUModelRunner( if self.uses_mrope: positions = self.mrope_positions.gpu[:, :num_tokens_padded] - elif self.uses_xdrope_dim > 0: - positions = self.xdrope_positions.gpu[:, :num_tokens_padded] else: positions = self.positions[:num_tokens_padded]