[Core][Model] Unify XD-RoPE into M-RoPE and derive the channel count (#56078)
Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
This commit is contained in:
co-authored by
Claude Opus 5
parent
c55e15a44e
commit
719284fe15
@@ -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",
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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."""
|
||||
|
||||
@@ -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:
|
||||
|
||||
@@ -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"]
|
||||
|
||||
@@ -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
|
||||
@@ -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
|
||||
|
||||
@@ -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,
|
||||
)
|
||||
|
||||
|
||||
@@ -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}
|
||||
)
|
||||
|
||||
@@ -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:
|
||||
|
||||
@@ -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,
|
||||
|
||||
@@ -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
|
||||
]
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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
|
||||
|
||||
@@ -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()
|
||||
|
||||
@@ -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.
|
||||
|
||||
@@ -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]
|
||||
|
||||
|
||||
Reference in New Issue
Block a user