remove multiple dead codes

Signed-off-by: yewentao256 <zhyanwentao@126.com>

.buildkite/performance-benchmarks/tests/serving-tests.json

@@ -28,11 +28,6 @@
         "dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json"
       }
     },
-    {
-        "dataset_name": "sharegpt",
-        "dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json"
-      }
-    },
     {
       "test_name": "serving_llama8B_tp1_random_128_128",
       "server_parameters": {

.buildkite/scripts/hardware_ci/run-intel-test.sh

@@ -369,7 +369,7 @@ export HF_TOKEN ZE_AFFINITY_MASK
     -e CMDS \
     --name "${container_name}" \
     "${IMAGE}" \
-    bash -c 'set -e; source /opt/intel/oneapi/setvars.sh --force; source /opt/intel/oneapi/ccl/2021.15/env/vars.sh --force; echo "ZE_AFFINITY_MASK is ${ZE_AFFINITY_MASK:-}"; eval "$CMDS"' \
+    bash -c 'set -e; echo "ZE_AFFINITY_MASK is ${ZE_AFFINITY_MASK:-}"; eval "$CMDS"' \
     >/dev/null
 } 9>/tmp/docker-pull.lock
 

docker/Dockerfile.xpu

@@ -63,21 +63,6 @@ RUN apt-get update -y && \
     python3-pip && \
     rm -rf /var/lib/apt/lists/*
 
-# Add oneAPI repo, pin oneAPI to 2025.3, then install pinned packages in one layer.
-RUN wget -O- https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS.PUB | gpg --dearmor | tee /usr/share/keyrings/oneapi-archive-keyring.gpg > /dev/null && \
-    echo "deb [signed-by=/usr/share/keyrings/oneapi-archive-keyring.gpg] https://apt.repos.intel.com/oneapi all main" | tee /etc/apt/sources.list.d/oneAPI.list && \
-    printf '%s\n' \
-    'Package: intel-oneapi-* intel-deep-learning-essentials* intel-pti*' \
-    'Pin: version 2025.3*' \
-    'Pin-Priority: 1001' \
-    > /etc/apt/preferences.d/oneapi-2025.3.pref && \
-    apt-get update -y && \
-    apt-get install -y --no-install-recommends \
-    intel-oneapi-compiler-dpcpp-cpp-2025.3 \
-    intel-oneapi-mkl-devel-2025.3 \
-    intel-oneapi-dnnl-devel-2025.3 && \
-    rm -rf /var/lib/apt/lists/*
-
 # Install UMD
 RUN mkdir neo && \
     cd neo && \
@@ -100,22 +85,6 @@ RUN curl -LsSf https://astral.sh/uv/install.sh | sh \
     && uv venv --python ${PYTHON_VERSION} --seed ${VIRTUAL_ENV}
 ENV PATH="$VIRTUAL_ENV/bin:$PATH"
 
-# This oneccl contains the BMG support which is not the case for default version of oneapi 2025.3.
-ARG ONECCL_INSTALLER="intel-oneccl-2021.15.9.14_offline.sh"
-RUN wget "https://github.com/uxlfoundation/oneCCL/releases/download/2021.15.9/${ONECCL_INSTALLER}" && \
-    bash "${ONECCL_INSTALLER}" -a --silent --eula accept && \
-    rm "${ONECCL_INSTALLER}" && \
-    echo "source /opt/intel/oneapi/setvars.sh --force" >> /root/.bashrc && \
-    echo "source /opt/intel/oneapi/ccl/2021.15/env/vars.sh --force" >> /root/.bashrc && \
-    rm -f /opt/intel/oneapi/ccl/latest && \
-    ln -s /opt/intel/oneapi/ccl/2021.15 /opt/intel/oneapi/ccl/latest && \
-    printf '%s\n' \
-    '/opt/intel/oneapi/ccl/2021.15/lib' \
-    '/opt/intel/oneapi/mpi/2021.15/lib' \
-    '/opt/intel/oneapi/compiler/2025.3/lib' \
-    > /etc/ld.so.conf.d/oneapi-ccl.conf && \
-    ldconfig
-
 SHELL ["bash", "-c"]
 CMD ["bash", "-c", "source /root/.bashrc && exec bash"]
 
@@ -135,8 +104,7 @@ RUN --mount=type=cache,target=/root/.cache/uv \
     uv pip install --upgrade pip
 
 
-
-ENV LD_LIBRARY_PATH=/opt/intel/oneapi/ccl/2021.15/lib:/opt/intel/oneapi/mpi/2021.15/lib:/opt/intel/oneapi/compiler/2025.3/lib:/usr/local/lib
+ENV LD_LIBRARY_PATH=/opt/venv/lib:/usr/local/lib
 CMD ["/bin/bash"]
 
 ######################### UCX + NIXL BUILD STAGE #########################
@@ -216,8 +184,7 @@ RUN --mount=type=cache,target=/root/.cache/uv \
     uv pip install -r /workspace/vllm/requirements/xpu.txt && \
     uv pip install --no-build-isolation -r /workspace/vllm/requirements/test/xpu.txt && \
     uv pip uninstall triton triton-xpu && \
-    uv pip install triton-xpu==3.7.1 && \
-    uv pip uninstall oneccl oneccl-devel
+    uv pip install triton-xpu==3.7.2
 
 # Keep source-dependent layers near the end so frequent code-only changes
 # don't invalidate heavy dependency and UCX/NIXL layers.

docs/getting_started/installation/gpu.xpu.inc.md

@@ -42,12 +42,12 @@ pip install -v -r requirements/xpu.txt
 
     ```bash
     pip uninstall -y triton triton-xpu
-    pip install triton-xpu==3.7.1 --extra-index-url https://download.pytorch.org/whl/xpu
+    pip install triton-xpu==3.7.2 --extra-index-url https://download.pytorch.org/whl/xpu
     ```
 
     !!! note
         - `triton` (without suffix) is for NVIDIA GPUs only. On XPU, using it instead of `triton-xpu` can cause correctness or runtime issues.
-        - For torch 2.12 (the version used in `requirements/xpu.txt`), the matching package is `triton-xpu==3.7.1`. If you use a different version of torch, check the corresponding `triton-xpu` version in [docker/Dockerfile.xpu](https://github.com/vllm-project/vllm/blob/main/docker/Dockerfile.xpu).
+        - For torch 2.13 (the version used in `requirements/xpu.txt`), the matching package is `triton-xpu==3.7.2`. If you use a different version of torch, check the corresponding `triton-xpu` version in [docker/Dockerfile.xpu](https://github.com/vllm-project/vllm/blob/main/docker/Dockerfile.xpu).
 
 - Finally, build and install vLLM XPU backend:
 

requirements/test/xpu.txt

@@ -140,7 +140,7 @@ docopt==0.6.2
     # via num2words
 docstring-parser==0.18.0
     # via anthropic
-dpcpp-cpp-rt==2025.3.2
+dpcpp-cpp-rt==2026.0.0
     # via
     #   onemkl-sycl-blas
     #   onemkl-sycl-dft
@@ -253,27 +253,27 @@ ijson==3.5.0
     # via -r requirements/test/../common.txt
 imageio==2.37.3
     # via scikit-image
-impi-rt==2021.17.2
+impi-rt==2021.18.0
     # via
     #   oneccl
     #   torch
 iniconfig==2.3.0
     # via pytest
-intel-cmplr-lib-rt==2025.3.2
+intel-cmplr-lib-rt==2026.0.0
     # via
     #   intel-sycl-rt
     #   torch
-intel-cmplr-lib-ur==2025.3.2
+intel-cmplr-lib-ur==2026.0.0
     # via
     #   intel-openmp
     #   intel-sycl-rt
     #   torch
-intel-cmplr-lic-rt==2025.3.2
+intel-cmplr-lic-rt==2026.0.0
     # via
     #   intel-opencl-rt
     #   intel-sycl-rt
     #   torch
-intel-opencl-rt==2025.3.2
+intel-opencl-rt==2026.0.0
     # via
     #   dpcpp-cpp-rt
     #   onemkl-sycl-blas
@@ -282,14 +282,14 @@ intel-opencl-rt==2025.3.2
     #   onemkl-sycl-rng
     #   onemkl-sycl-sparse
     #   torch
-intel-openmp==2025.3.2
+intel-openmp==2026.0.0
     # via
     #   dpcpp-cpp-rt
     #   mkl
     #   torch
-intel-pti==0.16.0
+intel-pti==0.17.0
     # via torch
-intel-sycl-rt==2025.3.2
+intel-sycl-rt==2026.0.0
     # via
     #   dpcpp-cpp-rt
     #   oneccl
@@ -378,7 +378,7 @@ mistral-common==1.11.5
     #   -c requirements/common.txt
     #   -r requirements/test/../common.txt
     #   -r requirements/test/xpu.in
-mkl==2025.3.1
+mkl==2026.0.0
     # via
     #   onemkl-sycl-blas
     #   onemkl-sycl-dft
@@ -453,28 +453,28 @@ numpy==2.2.6
     #   torchvision
     #   transformers
     #   xgrammar
-oneccl==2021.17.2
+oneccl==2022.0.0
     # via
     #   oneccl-devel
     #   torch
-oneccl-devel==2021.17.2
+oneccl-devel==2022.0.0
     # via torch
-onemkl-license==2025.3.1
+onemkl-license==2026.0.0
     # via
     #   mkl
     #   torch
-onemkl-sycl-blas==2025.3.1
+onemkl-sycl-blas==2026.0.0
     # via
     #   onemkl-sycl-lapack
     #   onemkl-sycl-sparse
     #   torch
-onemkl-sycl-dft==2025.3.1
+onemkl-sycl-dft==2026.0.0
     # via torch
-onemkl-sycl-lapack==2025.3.1
+onemkl-sycl-lapack==2026.0.0
     # via torch
-onemkl-sycl-rng==2025.3.1
+onemkl-sycl-rng==2026.0.0
     # via torch
-onemkl-sycl-sparse==2025.3.1
+onemkl-sycl-sparse==2026.0.0
     # via torch
 openai==2.44.0
     # via
@@ -719,6 +719,8 @@ pyyaml==6.0.3
     #   timm
     #   transformers
     #   uvicorn
+pyzes==0.1.1
+    # via torch
 pyzmq==27.1.0
     # via
     #   -c requirements/common.txt
@@ -871,14 +873,14 @@ tabledata==1.3.4
     # via pytablewriter
 tabulate==0.10.0
     # via sacrebleu
-tbb==2022.3.1
+tbb==2023.0.0
     # via
     #   intel-opencl-rt
     #   mkl
     #   torch
 tblib==3.1.0
     # via -r requirements/test/xpu.in
-tcmlib==1.4.1
+tcmlib==1.5.0
     # via
     #   tbb
     #   torch
@@ -910,7 +912,7 @@ tokenizers==0.22.2
     #   -c requirements/common.txt
     #   -r requirements/test/../common.txt
     #   transformers
-torch==2.12.0+xpu
+torch==2.13.0+xpu
     # via
     #   -c requirements/xpu.txt
     #   accelerate
@@ -920,7 +922,7 @@ torch==2.12.0+xpu
     #   timm
     #   torchvision
     #   xgrammar
-torchvision==0.27.0+xpu
+torchvision==0.28.0+xpu
     # via timm
 tqdm==4.67.3
     # via
@@ -946,7 +948,7 @@ transformers==5.14.1
     #   xgrammar
 triton==3.7.1
     # via xgrammar
-triton-xpu==3.7.1
+triton-xpu==3.7.2
     # via torch
 typepy==1.3.4
     # via
@@ -1001,7 +1003,7 @@ typing-inspection==0.4.2
     #   mcp
     #   pydantic
     #   pydantic-settings
-umf==1.0.3
+umf==1.1.0
     # via
     #   intel-cmplr-lib-ur
     #   torch

requirements/xpu.txt

@@ -12,10 +12,10 @@ jinja2>=3.1.6
 datasets # for benchmark scripts
 numba == 0.65.0 # Required for N-gram speculative decoding
 --extra-index-url=https://download.pytorch.org/whl/xpu
-torch==2.12.0
+torch==2.13.0
 torchaudio
 torchvision
 torchcodec >= 0.14 # Required for the torchcodec video decoding backend
 
-auto_round_lib==0.14.1
-vllm_xpu_kernels @ https://github.com/vllm-project/vllm-xpu-kernels/releases/download/v0.1.11.1/vllm_xpu_kernels-0.1.11.1-cp38-abi3-manylinux_2_28_x86_64.whl
+auto_round_lib==0.14.2
+vllm_xpu_kernels @ https://github.com/vllm-project/vllm-xpu-kernels/releases/download/v0.1.12/vllm_xpu_kernels-0.1.12-cp38-abi3-manylinux_2_28_x86_64.whl

rust/src/chat/src/backend/hf.rs

@@ -57,6 +57,7 @@ impl HfChatBackend {
                     processor_config: files.processor_config_path.as_deref(),
                 },
                 tokenizer.clone(),
+                options.limit_mm_per_prompt.clone(),
             )?
         };
         let multimodal_render_info = resolve_multimodal_render_info(multimodal_model_info.as_ref());
@@ -231,6 +232,7 @@ mod tests {
                 chat_template_content_format: Default::default(),
                 chat_template: None,
                 default_chat_template_kwargs: HashMap::new(),
+                limit_mm_per_prompt: HashMap::new(),
             },
             test_tokenizer(),
         )

rust/src/chat/src/backend/mod.rs

@@ -8,7 +8,7 @@ use serde_json::Value;
 use vllm_text::{DynTextBackend, TextBackend};
 
 use crate::error::Result;
-use crate::multimodal::MultimodalModelInfo;
+use crate::multimodal::{MmLimitPerPrompt, MultimodalModelInfo};
 use crate::output::DynChatOutputProcessor;
 use crate::renderer::DynChatRenderer;
 use crate::request::ChatRequest;
@@ -74,6 +74,9 @@ pub struct LoadModelBackendsOptions {
     /// Optional server-default keyword arguments merged into every
     /// chat-template render before request-level `chat_template_kwargs`.
     pub default_chat_template_kwargs: HashMap<String, Value>,
+    /// Maximum number of input items allowed per prompt for each modality.
+    /// Unspecified modalities are unlimited.
+    pub limit_mm_per_prompt: MmLimitPerPrompt,
 }
 
 /// Shared backends loaded from a model id.

rust/src/chat/src/error.rs

@@ -23,6 +23,8 @@ pub enum Error {
     UnsupportedMultimodalContent(&'static str),
     #[error("`{modality}` input is not supported by this model")]
     UnsupportedModality { modality: String },
+    #[error("At most {limit} {modality}(s) may be provided in one prompt.")]
+    MmLimitExceeded { modality: String, limit: usize },
     #[error("multimodal preprocessing error: {0}")]
     Multimodal(#[message] String),
     #[error("{kind} parsing is not available for model `{model_id}`")]
@@ -87,7 +89,8 @@ impl Error {
             Self::Text(error) => error.is_request_validation_error(),
             Self::UnsupportedMultimodalRenderer
             | Self::UnsupportedMultimodalContent(_)
-            | Self::UnsupportedModality { .. } => true,
+            | Self::UnsupportedModality { .. }
+            | Self::MmLimitExceeded { .. } => true,
 
             _ => false,
         }

rust/src/chat/src/multimodal.rs

@@ -11,7 +11,7 @@
 //! Raw media stays above `vllm-text`; this module lowers it into token IDs and
 //! opaque tensor payloads before the request is handed to text generation.
 
-use std::collections::HashSet;
+use std::collections::{BTreeMap, HashMap, HashSet};
 use std::fs;
 use std::path::Path;
 use std::sync::{Arc, LazyLock};
@@ -24,6 +24,7 @@ use llm_multimodal::{
     PromptReplacement, Tokenizer as TokenResolver, TrackedMedia, VideoClip, VisionPreProcessor,
     VisionProcessorRegistry,
 };
+use serde::{Deserialize, Serialize};
 use thiserror_ext::AsReport as _;
 use tracing::warn;
 use vllm_engine_core_client::protocol::dtype::ModelDtype;
@@ -52,6 +53,71 @@ pub struct MultimodalModelInfo {
     video: Option<ModalitySupport>,
     audio: Option<AudioModalitySupport>,
     media_connector: Arc<MediaConnector>,
+    /// Maximum number of input items allowed per prompt for each modality.
+    limit_mm_per_prompt: MmLimitPerPrompt,
+}
+
+/// Per-modality item-count limits configured by `--limit-mm-per-prompt`.
+///
+/// Modalities absent from the map are unlimited.
+pub type MmLimitPerPrompt = HashMap<MmLimitModality, MmLimitSpec>;
+
+/// Modalities that `--limit-mm-per-prompt` can be keyed by.
+///
+/// Closed on purpose: these are exactly the keys Python accepts, per
+/// `MultiModalDummyOptionsBuiltins` in `vllm/config/multimodal.py`.
+#[derive(Debug, Clone, Copy, PartialEq, Eq, Hash, Serialize, Deserialize)]
+#[serde(rename_all = "snake_case")]
+pub enum MmLimitModality {
+    Image,
+    Audio,
+    Video,
+}
+
+impl MmLimitModality {
+    /// The wire name, matching Python's modality strings.
+    pub fn as_str(self) -> &'static str {
+        match self {
+            Self::Image => "image",
+            Self::Audio => "audio",
+            Self::Video => "video",
+        }
+    }
+}
+
+/// One modality's limit, in either of the two shapes Python accepts.
+#[derive(Debug, Clone, PartialEq, Eq, Serialize, Deserialize)]
+#[serde(
+    untagged,
+    expecting = "an item count, or an object with an optional `count` field"
+)]
+pub enum MmLimitSpec {
+    /// Legacy form: `"image": 16`
+    Count(usize),
+
+    /// Configurable form:
+    /// `"video": {"count": 1, "num_frames": 32}`
+    Options {
+        /// Absent means unlimited, matching an absent modality.
+        #[serde(default, skip_serializing_if = "Option::is_none")]
+        count: Option<usize>,
+
+        /// Preserve Python-owned options for forwarding. Never interpreted
+        /// here: they size the engine's dummy-profiling encoder cache, which
+        /// has no Rust counterpart.
+        #[serde(flatten)]
+        extra: BTreeMap<String, serde_json::Value>,
+    },
+}
+
+impl MmLimitSpec {
+    /// The configured item count, or `None` when this modality is unlimited.
+    pub fn count(&self) -> Option<usize> {
+        match self {
+            Self::Count(count) => Some(*count),
+            Self::Options { count, .. } => *count,
+        }
+    }
 }
 
 /// Model metadata and tokenizer access shared by all multimodal specs.
@@ -287,6 +353,7 @@ impl MultimodalModelInfo {
         model_type: Option<String>,
         files: MultimodalConfigFiles<'_>,
         tokenizer: DynTokenizer,
+        limit_mm_per_prompt: MmLimitPerPrompt,
     ) -> Result<Option<Self>> {
         let config = match files.config {
             Some(path) => {
@@ -319,7 +386,12 @@ impl MultimodalModelInfo {
             tokenizer: TokenizerResolver(tokenizer),
         };
 
-        Self::from_loaded(context, preprocessor_config, video_preprocessor_config)
+        Self::from_loaded(
+            context,
+            preprocessor_config,
+            video_preprocessor_config,
+            limit_mm_per_prompt,
+        )
     }
 
     /// Resolve multimodal support from an assembled context and parsed
@@ -328,6 +400,7 @@ impl MultimodalModelInfo {
         context: MultimodalModelContext,
         preprocessor_config: PreProcessorConfig,
         video_preprocessor_config: PreProcessorConfig,
+        limit_mm_per_prompt: MmLimitPerPrompt,
     ) -> Result<Option<Self>> {
         let (image, video) = Self::resolve_vision_lanes(
             &context,
@@ -356,6 +429,7 @@ impl MultimodalModelInfo {
             video,
             audio,
             media_connector,
+            limit_mm_per_prompt,
         }))
     }
 
@@ -567,7 +641,55 @@ fn input_audio_data_url(data: &str, format: Option<&str>) -> Result<String> {
     Ok(format!("data:{mime_type};base64,{data}"))
 }
 
+/// The modality a content part counts against, or `None` for plain text.
+///
+/// Embedding inputs share their base modality's budget rather than getting one
+/// of their own, matching Python's `modality.replace("_embeds", "")` in
+/// `vllm/entrypoints/chat_utils.py`.
+fn media_part_limit_modality(part: &MediaContentPart) -> Option<MmLimitModality> {
+    match part {
+        MediaContentPart::Text { .. } => None,
+        MediaContentPart::ImageUrl { .. }
+        | MediaContentPart::ImageData { .. }
+        | MediaContentPart::ImageEmbeds { .. } => Some(MmLimitModality::Image),
+        MediaContentPart::AudioUrl { .. } | MediaContentPart::AudioData { .. } => {
+            Some(MmLimitModality::Audio)
+        }
+        MediaContentPart::VideoUrl { .. } | MediaContentPart::VideoData { .. } => {
+            Some(MmLimitModality::Video)
+        }
+    }
+}
+
 impl MultimodalModelInfo {
+    /// Reject requests exceeding `--limit-mm-per-prompt`'s configured
+    /// per-modality item count, before any fetch/decode work is spent on them.
+    ///
+    /// Modalities without a configured count are unlimited.
+    fn validate_mm_limits(&self, media_parts: &[MediaContentPart]) -> Result<()> {
+        let mut counts: HashMap<MmLimitModality, usize> = HashMap::new();
+        for part in media_parts {
+            if let Some(modality) = media_part_limit_modality(part) {
+                *counts.entry(modality).or_default() += 1;
+            }
+        }
+
+        for (modality, count) in counts {
+            let Some(limit) = self.limit_mm_per_prompt.get(&modality).and_then(MmLimitSpec::count)
+            else {
+                continue;
+            };
+            if count > limit {
+                return Err(Error::MmLimitExceeded {
+                    modality: modality.as_str().to_string(),
+                    limit,
+                });
+            }
+        }
+
+        Ok(())
+    }
+
     /// Run media fetch, per-modality preprocessing, prompt expansion, and
     /// feature build.
     ///
@@ -584,8 +706,7 @@ impl MultimodalModelInfo {
         if media_parts_len == 0 {
             return Ok(Vec::new());
         }
-        // TODO: enforce per-modality item-count limits, aligned with the
-        // engine's `--limit-mm-per-prompt` semantics.
+        self.validate_mm_limits(&media_parts)?;
         let fetched = self.fetch_media(media_parts).await?;
 
         let mut prepared = Vec::new();
@@ -753,10 +874,11 @@ mod tests {
             .with_regular_token("<|video_pad|>", QWEN3_VIDEO_PAD_ID)
     }
 
-    fn test_info(
+    fn test_info_with_limits(
         model_type: &str,
         config: serde_json::Value,
         tokenizer: TestTokenizer,
+        limit_mm_per_prompt: MmLimitPerPrompt,
     ) -> MultimodalModelInfo {
         let context = MultimodalModelContext {
             model_id: format!("{model_type}-test"),
@@ -769,11 +891,20 @@ mod tests {
             context,
             PreProcessorConfig::default(),
             PreProcessorConfig::default(),
+            limit_mm_per_prompt,
         )
         .unwrap()
         .unwrap_or_else(|| panic!("{model_type} multimodal support should resolve"))
     }
 
+    fn test_info(
+        model_type: &str,
+        config: serde_json::Value,
+        tokenizer: TestTokenizer,
+    ) -> MultimodalModelInfo {
+        test_info_with_limits(model_type, config, tokenizer, HashMap::new())
+    }
+
     fn llama4_info() -> MultimodalModelInfo {
         let config = serde_json::json!({
             "model_type": "llama4",
@@ -783,16 +914,19 @@ mod tests {
         test_info("llama4", config, llama4_tokenizer())
     }
 
-    pub(super) fn qwen3_vl_info() -> MultimodalModelInfo {
-        let config = serde_json::json!({
+    fn qwen3_vl_config() -> serde_json::Value {
+        serde_json::json!({
             "model_type": "qwen3_vl",
             "image_token_id": QWEN3_IMAGE_PAD_ID,
             "video_token_id": QWEN3_VIDEO_PAD_ID,
             "vision_start_token_id": 151652,
             "vision_end_token_id": 151653,
             "vision_config": {"patch_size": 16}
-        });
-        test_info("qwen3_vl", config, qwen3_vl_tokenizer())
+        })
+    }
+
+    pub(super) fn qwen3_vl_info() -> MultimodalModelInfo {
+        test_info("qwen3_vl", qwen3_vl_config(), qwen3_vl_tokenizer())
     }
 
     #[test]
@@ -853,4 +987,123 @@ mod tests {
         );
         assert!(input_audio_data_url("AAAA", Some("flac")).is_err());
     }
+
+    fn image_url_part() -> MediaContentPart {
+        MediaContentPart::ImageUrl {
+            url: "https://example.com/image.png".to_string(),
+            detail: None,
+            uuid: None,
+        }
+    }
+
+    fn qwen3_vl_info_with_limits(limit_mm_per_prompt: MmLimitPerPrompt) -> MultimodalModelInfo {
+        test_info_with_limits(
+            "qwen3_vl",
+            qwen3_vl_config(),
+            qwen3_vl_tokenizer(),
+            limit_mm_per_prompt,
+        )
+    }
+
+    #[test]
+    fn validate_mm_limits_ignores_text_parts() {
+        let info = qwen3_vl_info();
+        let parts = vec![
+            MediaContentPart::Text {
+                text: "hello".to_string(),
+            },
+            MediaContentPart::Text {
+                text: "world".to_string(),
+            },
+        ];
+        assert!(info.validate_mm_limits(&parts).is_ok());
+    }
+
+    #[test]
+    fn validate_mm_limits_leaves_unconfigured_modalities_unlimited() {
+        let info = qwen3_vl_info();
+        let parts: Vec<_> = std::iter::repeat_with(image_url_part).take(1_000).collect();
+
+        assert!(info.validate_mm_limits(&parts).is_ok());
+    }
+
+    #[test]
+    fn validate_mm_limits_enforces_configured_limit_at_the_boundary() {
+        let info = qwen3_vl_info_with_limits(HashMap::from([(
+            MmLimitModality::Image,
+            MmLimitSpec::Count(1),
+        )]));
+        assert!(info.validate_mm_limits(&[image_url_part()]).is_ok());
+
+        let error = info.validate_mm_limits(&[image_url_part(), image_url_part()]).unwrap_err();
+        assert_eq!(
+            error.to_report_string(),
+            "At most 1 image(s) may be provided in one prompt."
+        );
+        // Confirms the HTTP-mapping bug found during implementation stays fixed:
+        // this must map to 400, not 500.
+        assert!(error.is_request_validation_error());
+    }
+
+    #[test]
+    fn validate_mm_limits_counts_image_embeds_against_the_image_limit() {
+        let info = qwen3_vl_info_with_limits(HashMap::from([(
+            MmLimitModality::Image,
+            MmLimitSpec::Count(1),
+        )]));
+        let image_embeds_part = MediaContentPart::ImageEmbeds {
+            payload: serde_json::Value::String("AAAA".to_string()),
+            uuid: None,
+        };
+
+        let error = info.validate_mm_limits(&[image_url_part(), image_embeds_part]).unwrap_err();
+        assert_eq!(
+            error.to_report_string(),
+            "At most 1 image(s) may be provided in one prompt."
+        );
+    }
+
+    /// An options object without a `count` carries only profiling keys, which
+    /// say nothing about how many items are allowed.
+    #[test]
+    fn validate_mm_limits_treats_a_count_less_options_object_as_unlimited() {
+        let info = qwen3_vl_info_with_limits(HashMap::from([(
+            MmLimitModality::Image,
+            MmLimitSpec::Options {
+                count: None,
+                extra: BTreeMap::from([("width".to_string(), serde_json::json!(512))]),
+            },
+        )]));
+
+        assert!(info.validate_mm_limits(&[image_url_part(), image_url_part()]).is_ok());
+    }
+
+    fn parse_limits(json: &str) -> MmLimitPerPrompt {
+        serde_json::from_str(json).expect("limit map should parse")
+    }
+
+    #[test]
+    fn limit_map_parses_both_shapes_python_accepts() {
+        let limits = parse_limits(r#"{"image": 16, "video": {"count": 1, "num_frames": 32}}"#);
+
+        assert_eq!(limits[&MmLimitModality::Image].count(), Some(16));
+        assert_eq!(limits[&MmLimitModality::Video].count(), Some(1));
+        assert_eq!(limits.get(&MmLimitModality::Audio), None);
+    }
+
+    #[test]
+    fn limit_map_rejects_keys_python_does_not_accept() {
+        assert!(serde_json::from_str::<MmLimitPerPrompt>(r#"{"image_embeds": 1}"#).is_err());
+    }
+
+    /// Managed mode forwards this map back to Python as JSON, where
+    /// `BaseDummyOptions.count` is a non-optional `int` under `extra="forbid"`.
+    /// Emitting `"count": null` would make the engine subprocess fail to start.
+    #[test]
+    fn limit_map_round_trips_without_emitting_a_null_count() {
+        let source = r#"{"video":{"num_frames":32}}"#;
+        let encoded = serde_json::to_string(&parse_limits(source)).expect("map should serialize");
+
+        assert_eq!(encoded, source);
+    }
 }

rust/src/chat/src/multimodal/audio.rs

@@ -107,6 +107,7 @@ mod tests {
             context,
             PreProcessorConfig::default(),
             PreProcessorConfig::default(),
+            HashMap::new(),
         )
         .unwrap()
         .expect("Inkling multimodal support")
@@ -126,6 +127,7 @@ mod tests {
             context,
             PreProcessorConfig::default(),
             PreProcessorConfig::default(),
+            HashMap::new(),
         )
         .unwrap()
         .expect("Qwen3-ASR multimodal support")

rust/src/chat/src/multimodal/expand.rs

@@ -222,7 +222,9 @@ mod tests {
         assert_eq!(tensor.shape, vec![expected.len()]);
         assert_eq!(
             tensor.data,
-            WireArrayData::RawView(expected.iter().map(|value| u8::from(*value)).collect())
+            WireArrayData::RawView(
+                expected.iter().map(|value| u8::from(*value)).collect::<Vec<_>>().into(),
+            )
         );
     }
 

rust/src/chat/src/multimodal/tensor.rs

@@ -2,28 +2,58 @@
 // SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
 use std::collections::HashMap;
+use std::mem::size_of;
 
 use half::{bf16, f16};
 use llm_multimodal::{ModelSpecificValue, PreprocessedEncoderInputs};
 use vllm_engine_core_client::protocol::dtype::ModelDtype;
 use vllm_engine_core_client::protocol::multimodal::MmKwargValue as ProtocolKwargValue;
-use vllm_engine_core_client::protocol::tensor::{ShapeExt as _, WireTensor};
+use vllm_engine_core_client::protocol::tensor::{ShapeExt as _, WireArrayData, WireTensor};
 
 use crate::error::{Error, Result, bail_multimodal, multimodal};
 
+/// Element type retained alongside an encoded tensor during multimodal lowering.
+#[derive(Debug, Clone, Copy)]
+pub(super) enum TensorKind {
+    /// 32-bit floating point.
+    F32,
+    /// IEEE 16-bit floating point.
+    F16,
+    /// Brain floating point.
+    Bf16,
+    /// Signed 64-bit integer.
+    I64,
+    /// Unsigned 32-bit integer.
+    U32,
+}
+
+impl TensorKind {
+    const fn element_size(self) -> usize {
+        match self {
+            Self::F32 => size_of::<f32>(),
+            Self::F16 => size_of::<f16>(),
+            Self::Bf16 => size_of::<bf16>(),
+            Self::I64 => size_of::<i64>(),
+            Self::U32 => size_of::<u32>(),
+        }
+    }
+
+    const fn wire_dtype(self) -> &'static str {
+        match self {
+            Self::F32 => "float32",
+            Self::F16 => "float16",
+            Self::Bf16 => "bfloat16",
+            Self::I64 => "int64",
+            Self::U32 => "uint32",
+        }
+    }
+}
+
 /// Representation for multimodal kwarg values for transformation.
 #[derive(Debug)]
 pub(super) enum KwargValue {
-    /// Float tensor with row-major flat data and shape.
-    F32Tensor { data: Vec<f32>, shape: Vec<usize> },
-    /// Float16 tensor with row-major flat data and shape.
-    F16Tensor { data: Vec<f16>, shape: Vec<usize> },
-    /// BFloat16 tensor with row-major flat data and shape.
-    Bf16Tensor { data: Vec<bf16>, shape: Vec<usize> },
-    /// Signed integer tensor with row-major flat data and shape.
-    I64Tensor { data: Vec<i64>, shape: Vec<usize> },
-    /// Unsigned integer tensor with row-major flat data and shape.
-    U32Tensor { data: Vec<u32>, shape: Vec<usize> },
+    /// Tensor with row-major flat data and shape.
+    Tensor { kind: TensorKind, wire: WireTensor },
     /// Non-tensor kwarg value that is shared or copied as-is.
     Passthrough(ProtocolKwargValue),
 }
@@ -67,8 +97,14 @@ impl KwargValue {
             ModelSpecificValue::Tensor { data, shape } => {
                 Self::from_f32_tensor(data, shape, float_dtype)?
             }
-            ModelSpecificValue::IntTensor { data, shape } => Self::I64Tensor { data, shape },
-            ModelSpecificValue::UintTensor { data, shape } => Self::U32Tensor { data, shape },
+            ModelSpecificValue::IntTensor { data, shape } => {
+                let wire = WireTensor::from_i64(shape, data).map_err(Error::Multimodal)?;
+                Self::tensor(TensorKind::I64, wire)
+            }
+            ModelSpecificValue::UintTensor { data, shape } => {
+                let wire = WireTensor::from_u32(shape, data).map_err(Error::Multimodal)?;
+                Self::tensor(TensorKind::U32, wire)
+            }
             ModelSpecificValue::Int(value) => Self::Passthrough(Int(value)),
             ModelSpecificValue::Float(value) => Self::Passthrough(Float(value)),
             ModelSpecificValue::IntVec(values) => {
@@ -93,17 +129,23 @@ impl KwargValue {
     /// Convert a float tensor to the target float dtype if needed, keeping the
     /// same shape.
     fn from_f32_tensor(data: Vec<f32>, shape: Vec<usize>, float_dtype: ModelDtype) -> Result<Self> {
-        match float_dtype {
-            ModelDtype::Float16 => Ok(Self::F16Tensor {
-                data: data.into_iter().map(f16::from_f32).collect(),
-                shape,
-            }),
-            ModelDtype::BFloat16 => Ok(Self::Bf16Tensor {
-                data: data.into_iter().map(bf16::from_f32).collect(),
-                shape,
-            }),
-            ModelDtype::Float32 => Ok(Self::F32Tensor { data, shape }),
-        }
+        let (kind, wire) = match float_dtype {
+            ModelDtype::Float16 => (
+                TensorKind::F16,
+                WireTensor::from_f16(shape, data.into_iter().map(f16::from_f32).collect()),
+            ),
+            ModelDtype::BFloat16 => (
+                TensorKind::Bf16,
+                WireTensor::from_bf16(shape, data.into_iter().map(bf16::from_f32).collect()),
+            ),
+            ModelDtype::Float32 => (TensorKind::F32, WireTensor::from_f32(shape, data)),
+        };
+        wire.map(|wire| Self::tensor(kind, wire)).map_err(Error::Multimodal)
+    }
+
+    fn tensor(kind: TensorKind, wire: WireTensor) -> Self {
+        debug_assert_eq!(wire.dtype, kind.wire_dtype());
+        Self::Tensor { kind, wire }
     }
 }
 
@@ -111,15 +153,11 @@ impl TryFrom<&KwargValue> for ProtocolKwargValue {
     type Error = Error;
 
     fn try_from(value: &KwargValue) -> Result<Self> {
-        let tensor = match value {
-            KwargValue::F32Tensor { data, shape } => WireTensor::from_f32(shape.clone(), data),
-            KwargValue::F16Tensor { data, shape } => WireTensor::from_f16(shape.clone(), data),
-            KwargValue::Bf16Tensor { data, shape } => WireTensor::from_bf16(shape.clone(), data),
-            KwargValue::I64Tensor { data, shape } => WireTensor::from_i64(shape.clone(), data),
-            KwargValue::U32Tensor { data, shape } => WireTensor::from_u32(shape.clone(), data),
+        let wire = match value {
+            KwargValue::Tensor { wire, .. } => wire.clone(),
             KwargValue::Passthrough(value) => return Ok(value.clone()),
         };
-        tensor.map(ProtocolKwargValue::Tensor).map_err(Error::Multimodal)
+        Ok(ProtocolKwargValue::Tensor(wire))
     }
 }
 
@@ -127,11 +165,7 @@ impl KwargValue {
     /// First-axis length for tensor values; `None` for passthrough kwargs.
     pub(super) fn first_dim(&self) -> Option<usize> {
         match self {
-            Self::F32Tensor { shape, .. }
-            | Self::F16Tensor { shape, .. }
-            | Self::Bf16Tensor { shape, .. }
-            | Self::I64Tensor { shape, .. }
-            | Self::U32Tensor { shape, .. } => shape.first().copied(),
+            Self::Tensor { wire, .. } => wire.shape.first().copied(),
             Self::Passthrough(_) => None,
         }
     }
@@ -161,30 +195,13 @@ impl KwargValue {
         end: usize,
         drop_axis: bool,
     ) -> Result<ProtocolKwargValue> {
-        let tensor = match self {
-            Self::F32Tensor { data, shape } => {
-                let (shape, data) = slice_first_axis_range(shape, data, start, end, drop_axis)?;
-                WireTensor::from_f32(shape, data)
-            }
-            Self::F16Tensor { data, shape } => {
-                let (shape, data) = slice_first_axis_range(shape, data, start, end, drop_axis)?;
-                WireTensor::from_f16(shape, data)
-            }
-            Self::Bf16Tensor { data, shape } => {
-                let (shape, data) = slice_first_axis_range(shape, data, start, end, drop_axis)?;
-                WireTensor::from_bf16(shape, data)
-            }
-            Self::I64Tensor { data, shape } => {
-                let (shape, data) = slice_first_axis_range(shape, data, start, end, drop_axis)?;
-                WireTensor::from_i64(shape, data)
-            }
-            Self::U32Tensor { data, shape } => {
-                let (shape, data) = slice_first_axis_range(shape, data, start, end, drop_axis)?;
-                WireTensor::from_u32(shape, data)
+        let wire = match self {
+            Self::Tensor { kind, wire } => {
+                slice_first_axis_range(wire, kind.element_size(), start, end, drop_axis)
             }
             Self::Passthrough(value) => return Ok(value.clone()),
         };
-        tensor.map(ProtocolKwargValue::Tensor).map_err(Error::Multimodal)
+        wire.map(ProtocolKwargValue::Tensor)
     }
 }
 
@@ -208,44 +225,53 @@ pub(super) fn flat_range_for_index(
 /// Read a tensor value as per-image sizes for flat slicing.
 fn tensor_as_usize_vec(tensor: &KwargValue) -> Result<Vec<usize>> {
     match tensor {
-        KwargValue::I64Tensor { data, .. } => data
-            .iter()
+        KwargValue::Tensor {
+            kind: TensorKind::I64,
+            wire,
+        } => raw_tensor_bytes(wire, size_of::<i64>())?
+            .chunks_exact(size_of::<i64>())
+            .map(|bytes| i64::from_ne_bytes(bytes.try_into().expect("exact int64 chunk")))
             .map(|value| {
-                usize::try_from(*value)
+                usize::try_from(value)
                     .map_err(|_| multimodal!("negative flat tensor size `{value}`"))
             })
             .collect(),
-        KwargValue::U32Tensor { data, .. } => {
-            Ok(data.iter().map(|value| *value as usize).collect())
-        }
+        KwargValue::Tensor {
+            kind: TensorKind::U32,
+            wire,
+        } => Ok(raw_tensor_bytes(wire, size_of::<u32>())?
+            .chunks_exact(size_of::<u32>())
+            .map(|bytes| u32::from_ne_bytes(bytes.try_into().expect("exact uint32 chunk")) as usize)
+            .collect()),
         _ => Err(multimodal!("flat tensor sizes must be int64 or uint32")),
     }
 }
 
 /// Slice a flat row-major tensor along its first axis.
-fn slice_first_axis_range<'a, T>(
-    shape: &[usize],
-    data: &'a [T],
+fn slice_first_axis_range(
+    tensor: &WireTensor,
+    element_size: usize,
     start: usize,
     end: usize,
     drop_axis: bool,
-) -> Result<(Vec<usize>, &'a [T])> {
+) -> Result<WireTensor> {
+    let shape = tensor.shape.as_slice();
+    raw_tensor_bytes(tensor, element_size)?;
     let first_dim = *shape.first().ok_or_else(|| multimodal!("tensor has no first dimension"))?;
     if start > end || end > first_dim {
         bail_multimodal!("invalid tensor slice {start}..{end} for first dimension {first_dim}");
     }
-    let expected_len = shape
-        .checked_numel()
-        .ok_or_else(|| multimodal!("tensor shape {shape:?} has too many elements"))?;
-    if expected_len != data.len() {
-        bail_multimodal!(
-            "tensor shape {shape:?} expects {expected_len} elements, got {}",
-            data.len()
-        );
-    }
-    let stride = shape[1..].iter().product::<usize>();
-    let data_start = start * stride;
-    let data_end = end * stride;
+    let stride = shape[1..]
+        .iter()
+        .try_fold(1usize, |acc, dim| acc.checked_mul(*dim))
+        .and_then(|stride| stride.checked_mul(element_size))
+        .ok_or_else(|| multimodal!("tensor shape {shape:?} byte stride overflowed usize"))?;
+    let data_start = start
+        .checked_mul(stride)
+        .ok_or_else(|| multimodal!("tensor slice start byte offset overflowed usize"))?;
+    let data_end = end
+        .checked_mul(stride)
+        .ok_or_else(|| multimodal!("tensor slice end byte offset overflowed usize"))?;
     let out_shape = if drop_axis {
         shape[1..].to_vec()
     } else {
@@ -253,7 +279,38 @@ fn slice_first_axis_range<'a, T>(
         shape[0] = end - start;
         shape
     };
-    Ok((out_shape, &data[data_start..data_end]))
+    let WireArrayData::RawView(data) = &tensor.data else {
+        return Err(multimodal!("cannot slice an aux tensor buffer"));
+    };
+    Ok(WireTensor::from_raw_bytes(
+        tensor.dtype.clone(),
+        out_shape,
+        data.slice(data_start..data_end),
+    ))
+}
+
+fn raw_tensor_bytes(tensor: &WireTensor, element_size: usize) -> Result<&[u8]> {
+    let WireArrayData::RawView(data) = &tensor.data else {
+        return Err(multimodal!("expected an inline tensor buffer"));
+    };
+    let expected_bytes = tensor
+        .shape
+        .checked_numel()
+        .and_then(|numel| numel.checked_mul(element_size))
+        .ok_or_else(|| {
+            multimodal!(
+                "tensor shape {:?} byte length overflowed usize",
+                tensor.shape
+            )
+        })?;
+    if expected_bytes != data.len() {
+        bail_multimodal!(
+            "tensor shape {:?} expects {expected_bytes} bytes, got {}",
+            tensor.shape,
+            data.len()
+        );
+    }
+    Ok(data)
 }
 
 #[cfg(test)]
@@ -262,28 +319,32 @@ mod tests {
 
     #[test]
     fn batched_wire_value_at_drops_first_axis() {
-        let value = KwargValue::F32Tensor {
-            data: vec![1.0, 2.0, 3.0, 4.0],
-            shape: vec![2, 2],
-        };
+        let data = vec![1.0_f32, 2.0, 3.0, 4.0];
+        let expected_ptr = data.as_ptr().cast::<u8>().wrapping_add(2 * size_of::<f32>());
+        let value = KwargValue::tensor(
+            TensorKind::F32,
+            WireTensor::from_f32(vec![2, 2], data).unwrap(),
+        );
 
         let ProtocolKwargValue::Tensor(tensor) = value.batched_wire_value_at(1).unwrap() else {
             panic!("expected tensor");
         };
 
         assert_eq!(tensor.shape, vec![2]);
+        let raw_view = tensor.data.into_raw_view().unwrap();
+        assert_eq!(raw_view.as_ptr(), expected_ptr);
         assert_eq!(
-            tensor.data.into_raw_view().unwrap(),
+            raw_view,
             [3.0_f32, 4.0].into_iter().flat_map(f32::to_ne_bytes).collect::<Vec<_>>()
         );
     }
 
     #[test]
     fn flat_wire_value_range_keeps_first_axis() {
-        let value = KwargValue::U32Tensor {
-            data: (0..10).collect(),
-            shape: vec![5, 2],
-        };
+        let value = KwargValue::tensor(
+            TensorKind::U32,
+            WireTensor::from_u32(vec![5, 2], (0..10_u32).collect()).unwrap(),
+        );
 
         let ProtocolKwargValue::Tensor(tensor) = value.flat_wire_value_range(1, 3).unwrap() else {
             panic!("expected tensor");
@@ -298,10 +359,10 @@ mod tests {
 
     #[test]
     fn flat_range_for_index_uses_size_tensor() {
-        let sizes = KwargValue::I64Tensor {
-            data: vec![2, 3, 4],
-            shape: vec![3],
-        };
+        let sizes = KwargValue::tensor(
+            TensorKind::I64,
+            WireTensor::from_i64(vec![3], vec![2_i64, 3, 4]).unwrap(),
+        );
 
         assert_eq!(
             flat_range_for_index(&sizes, "image_grid_thw", 1).unwrap(),
@@ -311,10 +372,11 @@ mod tests {
 
     #[test]
     fn slice_first_axis_range_errors_on_shape_data_mismatch() {
-        let error = slice_first_axis_range(&[2, 2], &[1.0_f32, 2.0, 3.0], 0, 1, true).unwrap_err();
+        let tensor = WireTensor::from_raw("float32", vec![2, 2], vec![0; 3 * size_of::<f32>()]);
+        let error = slice_first_axis_range(&tensor, size_of::<f32>(), 0, 1, true).unwrap_err();
 
         assert!(
-            matches!(error, Error::Multimodal(message) if message.contains("expects 4 elements"))
+            matches!(error, Error::Multimodal(message) if message.contains("expects 16 bytes"))
         );
     }
 

rust/src/chat/src/multimodal/video.rs

@@ -207,6 +207,7 @@ mod tests {
                 Some("qwen3_vl".to_string()),
                 files,
                 Arc::new(qwen3_vl_tokenizer()),
+                std::collections::HashMap::new(),
             )
         };
 

rust/src/cmd/src/cli.rs

@@ -23,6 +23,7 @@ use serde_with::{DefaultOnNull, OneOrMany, serde_as};
 use thiserror_ext::AsReport as _;
 use uuid::Uuid;
 use vllm_chat::ReasoningParserFactory;
+use vllm_chat::multimodal::MmLimitPerPrompt;
 use vllm_engine_core_client::TransportMode;
 use vllm_managed_engine::ManagedEngineConfig;
 use vllm_managed_engine::cli::{ManagedEngineArgs, repartition_managed_engine_args};
@@ -186,6 +187,17 @@ pub struct SharedRuntimeArgs {
     #[serde(default)]
     pub default_chat_template_kwargs: Option<HashMap<String, Value>>,
 
+    /// The maximum number of input items allowed per prompt for each
+    /// modality, as a JSON object (e.g. `{"image": 16, "video": 2}`).
+    ///
+    /// Also accepts the engine's configurable form
+    /// (e.g. `{"video": {"count": 1, "num_frames": 32}}`); the extra
+    /// profiling options are forwarded to the engine untouched.
+    /// Unspecified modalities are unlimited.
+    #[arg(long, value_parser = parse_json::<MmLimitPerPrompt>, value_name = "JSON", default_value = "{}")]
+    #[serde(default)]
+    pub limit_mm_per_prompt: MmLimitPerPrompt,
+
     /// The format to render message content within a chat template.
     ///
     /// * "auto" detects the format from the template
@@ -348,6 +360,18 @@ impl SharedRuntimeArgs {
             .expect("profiler config serialization should not fail")
     }
 
+    /// Return the per-modality limits as JSON for managed Python engine
+    /// forwarding, or `None` when nothing is configured.
+    ///
+    /// Round-tripping the parsed map rather than the raw argument keeps the
+    /// engine's own profiling options (`num_frames`, `width`, ...) intact.
+    pub fn limit_mm_per_prompt_json(&self) -> Option<String> {
+        (!self.limit_mm_per_prompt.is_empty()).then(|| {
+            serde_json::to_string(&self.limit_mm_per_prompt)
+                .expect("limit-mm-per-prompt serialization should not fail")
+        })
+    }
+
     /// Apply fallback logic for API key configuration from env variables.
     fn apply_env_api_key_fallback(&mut self) {
         if self.api_key.is_empty()
@@ -399,6 +423,7 @@ impl SharedRuntimeArgs {
             language_model_only: self.language_model_only,
             chat_template: self.chat_template,
             default_chat_template_kwargs: self.default_chat_template_kwargs,
+            limit_mm_per_prompt: self.limit_mm_per_prompt,
             chat_template_content_format: self.chat_template_content_format,
             max_logprobs: self.max_logprobs,
             api_server_options,
@@ -451,6 +476,7 @@ impl SharedRuntimeArgs {
             language_model_only: self.language_model_only,
             chat_template: self.chat_template,
             default_chat_template_kwargs: self.default_chat_template_kwargs,
+            limit_mm_per_prompt: self.limit_mm_per_prompt,
             chat_template_content_format: self.chat_template_content_format,
             max_logprobs: self.max_logprobs,
             api_server_options,
@@ -666,6 +692,7 @@ impl ServeArgs {
             self.runtime.disable_log_stats,
             self.runtime.shutdown_timeout,
             handshake_port,
+            self.runtime.limit_mm_per_prompt_json(),
         )
     }
 }

rust/src/cmd/src/cli/tests.rs

@@ -64,6 +64,7 @@ fn serve_args_forward_python_flags_with_separator() {
                         http_timeout_keep_alive: None,
                         chat_template: None,
                         default_chat_template_kwargs: None,
+                        limit_mm_per_prompt: {},
                         chat_template_content_format: Auto,
                         enable_log_requests: false,
                         enable_prompt_tokens_details: false,
@@ -762,6 +763,7 @@ fn frontend_args_accept_json() {
                         http_timeout_keep_alive: None,
                         chat_template: None,
                         default_chat_template_kwargs: None,
+                        limit_mm_per_prompt: {},
                         chat_template_content_format: Auto,
                         enable_log_requests: false,
                         enable_prompt_tokens_details: false,
@@ -1117,6 +1119,24 @@ fn frontend_args_json_rejects_malformed_json() {
     "#]].assert_eq(&error.to_string());
 }
 
+#[test]
+fn serve_args_reject_unsupported_modality_in_limit_mm_per_prompt() {
+    let error = Cli::try_parse_from([
+        "vllm-rs",
+        "serve",
+        "Qwen/Qwen3-0.6B",
+        "--limit-mm-per-prompt",
+        r#"{"unsupported_modality": 1}"#,
+    ])
+    .unwrap_err();
+
+    expect![[r#"
+        error: invalid value '{"unsupported_modality": 1}' for '--limit-mm-per-prompt <JSON>': invalid JSON object: unknown variant `unsupported_modality`, expected one of `image`, `audio`, `video` at line 1 column 23
+
+        For more information, try '--help'.
+    "#]].assert_eq(&error.to_string());
+}
+
 #[test]
 fn serve_args_reject_flags_before_model() {
     let error = Cli::try_parse_from(["vllm-rs", "serve", "--python", "python3", "Qwen/Qwen3-0.6B"])
@@ -1331,6 +1351,7 @@ fn serve_args_accept_handshake_aliases() {
                         http_timeout_keep_alive: None,
                         chat_template: None,
                         default_chat_template_kwargs: None,
+                        limit_mm_per_prompt: {},
                         chat_template_content_format: Auto,
                         enable_log_requests: false,
                         enable_prompt_tokens_details: false,
@@ -1474,6 +1495,7 @@ fn serve_frontend_config_uses_dp_address_as_advertised_host() {
             language_model_only: false,
             chat_template: None,
             default_chat_template_kwargs: None,
+            limit_mm_per_prompt: {},
             chat_template_content_format: Auto,
             max_logprobs: None,
             api_server_options: ApiServerOptions {
@@ -1558,6 +1580,7 @@ fn serve_frontend_config_keeps_tcp_transport_for_non_local_only_topology() {
             language_model_only: false,
             chat_template: None,
             default_chat_template_kwargs: None,
+            limit_mm_per_prompt: {},
             chat_template_content_format: Auto,
             max_logprobs: None,
             api_server_options: ApiServerOptions {
@@ -1660,6 +1683,7 @@ fn frontend_config_uses_external_coordinator_when_coordinator_address_is_present
             language_model_only: false,
             chat_template: None,
             default_chat_template_kwargs: None,
+            limit_mm_per_prompt: {},
             chat_template_content_format: Auto,
             max_logprobs: None,
             api_server_options: ApiServerOptions {

rust/src/cmd/src/cli/unsupported.rs

@@ -299,11 +299,6 @@ pub struct EngineUnsupportedArgs {
     )]
     pub kv_sharing_fast_prefill: Option<Unsupported>,
 
-    /// The maximum number of input items and options allowed per
-    /// prompt for each modality.
-    #[arg(long)]
-    pub limit_mm_per_prompt: Option<Unsupported>,
-
     /// Additional args passed to process media inputs, keyed by modalities.
     #[arg(long)]
     pub media_io_kwargs: Option<Unsupported>,

rust/src/engine-core-client/src/client.rs

@@ -519,7 +519,7 @@ impl EngineCoreClient {
                 "registered request to engine"
             );
 
-            self.inner.send_to_engine(&engine_id, EngineCoreRequestType::Add, &req).await?;
+            self.inner.send_request_to_engine(&engine_id, req).await?;
             Ok(())
         }
         .await;

rust/src/engine-core-client/src/client/imp.rs

@@ -6,6 +6,7 @@ use std::sync::Arc;
 use std::time::{SystemTime, UNIX_EPOCH};
 
 use arc_swap::ArcSwapOption;
+use bytes::Bytes;
 use parking_lot::Mutex;
 use thiserror_ext::AsReport as _;
 use tokio::runtime::Handle;
@@ -21,17 +22,30 @@ use crate::error::{client_closed, dispatcher_closed, unexpected_dispatcher_outpu
 use crate::metrics::{LoraInfoExporter, SchedulerStatsRecorder};
 use crate::protocol::encode_msgpack;
 use crate::protocol::output::{EngineCoreOutput, EngineCoreOutputs};
-use crate::protocol::request::EngineCoreRequestType;
+use crate::protocol::request::{EngineCoreRequest, EngineCoreRequestType};
 use crate::protocol::stats::SchedulerStats;
 use crate::protocol::utility::UtilityOutput;
 use crate::transport::{ConnectedEngine, EngineId};
 use crate::{Error, Result, transport};
 
+const MSGPACK_ZERO_COPY_THRESHOLD_ENV: &str = "VLLM_MSGPACK_ZERO_COPY_THRESHOLD";
+const DEFAULT_MSGPACK_ZERO_COPY_THRESHOLD: usize = 256;
+
+fn msgpack_zero_copy_threshold() -> usize {
+    std::env::var(MSGPACK_ZERO_COPY_THRESHOLD_ENV)
+        .ok()
+        .and_then(|value| value.parse().ok())
+        .unwrap_or(DEFAULT_MSGPACK_ZERO_COPY_THRESHOLD)
+}
+
 pub(crate) struct ClientInner {
     input_send: RouterSendHalf,
     /// The runtime handle used for sending messages to the engine.
     handle: Handle,
     model_name: String,
+    /// Per-tensor byte threshold loaded from env variable
+    /// `VLLM_MSGPACK_ZERO_COPY_THRESHOLD` when this inner client is created.
+    msgpack_zero_copy_threshold: usize,
     scheduler_stats_recorder: SchedulerStatsRecorder,
     request_reg: Mutex<RequestRegistry>,
     utility_reg: Mutex<UtilityRegistry>,
@@ -54,6 +68,7 @@ impl ClientInner {
             input_send,
             handle,
             model_name,
+            msgpack_zero_copy_threshold: msgpack_zero_copy_threshold(),
             scheduler_stats_recorder,
             request_reg: Mutex::new(RequestRegistry::new(engines)),
             utility_reg: Mutex::new(UtilityRegistry::default()),
@@ -229,9 +244,29 @@ impl ClientInner {
     where
         T: serde::Serialize + std::fmt::Debug,
     {
-        // TODO: for `EngineCoreRequest`, split outbound tensor raw views into aux
-        // frames instead of always producing a single msgpack frame.
-        let payload = encode_msgpack(payload)?;
+        let payload = Bytes::from(encode_msgpack(payload)?);
+        self.send_encoded_to_engine(engine_id, request_type, payload, Vec::new()).await
+    }
+
+    /// Send an add request, moving large tensor buffers into auxiliary frames.
+    pub async fn send_request_to_engine(
+        &self,
+        engine_id: &EngineId,
+        mut payload: EngineCoreRequest,
+    ) -> Result<()> {
+        let aux_frames = payload.extract_aux_frames(self.msgpack_zero_copy_threshold);
+        let payload = Bytes::from(encode_msgpack(&payload)?);
+        self.send_encoded_to_engine(engine_id, EngineCoreRequestType::Add, payload, aux_frames)
+            .await
+    }
+
+    async fn send_encoded_to_engine(
+        &self,
+        engine_id: &EngineId,
+        request_type: EngineCoreRequestType,
+        payload: Bytes,
+        aux_frames: Vec<Bytes>,
+    ) -> Result<()> {
         let mut input_send = self.input_send.clone();
         let engine_id = engine_id.clone();
 
@@ -242,6 +277,7 @@ impl ClientInner {
                     &engine_id,
                     request_type.to_frame(),
                     payload,
+                    aux_frames,
                 )
                 .await
             })

rust/src/engine-core-client/src/protocol/logprobs.rs

@@ -209,17 +209,17 @@ impl WireLogprobs {
             logprob_token_ids: WireNdArray {
                 dtype: "<i8".to_string(),
                 shape: vec![rows, cols],
-                data: WireArrayData::RawView(token_ids),
+                data: WireArrayData::RawView(token_ids.into()),
             },
             logprobs: WireNdArray {
                 dtype: "<f4".to_string(),
                 shape: vec![rows, cols],
-                data: WireArrayData::RawView(logprobs),
+                data: WireArrayData::RawView(logprobs.into()),
             },
             token_ranks: WireNdArray {
                 dtype: "<i8".to_string(),
                 shape: vec![rows],
-                data: WireArrayData::RawView(token_ranks),
+                data: WireArrayData::RawView(token_ranks.into()),
             },
             cu_num_generated_tokens: None,
         })

rust/src/engine-core-client/src/protocol/logprobs/array.rs

@@ -4,6 +4,7 @@
 use std::io::Cursor;
 
 use byteorder::{BigEndian, LittleEndian, NativeEndian, ReadBytesExt};
+use bytes::Bytes;
 use itertools::Itertools as _;
 
 use crate::error::{Error, Result, ext_value_decode};
@@ -126,7 +127,7 @@ pub(super) fn decode_array_metadata<Frame>(
     field: &str,
     frames: &[Frame],
     expected_scalars: &[ScalarType],
-) -> Result<(Vec<usize>, Vec<u8>, ScalarType, Endianness)>
+) -> Result<(Vec<usize>, Bytes, ScalarType, Endianness)>
 where
     Frame: AsRef<[u8]>,
 {
@@ -171,7 +172,7 @@ pub(super) fn resolve_array_bytes<Frame>(
     value: WireArrayData,
     field: &str,
     frames: &[Frame],
-) -> Result<Vec<u8>>
+) -> Result<Bytes>
 where
     Frame: AsRef<[u8]>,
 {
@@ -187,7 +188,7 @@ where
                     ),
                 )
             })?;
-            Ok(frame.as_ref().to_vec())
+            Ok(Bytes::copy_from_slice(frame.as_ref()))
         }
     }
 }

rust/src/engine-core-client/src/protocol/multimodal.rs

@@ -3,6 +3,7 @@
 
 use std::collections::BTreeMap;
 
+use bytes::Bytes;
 use serde::{Deserialize, Serialize};
 use serde_tuple::{Deserialize_tuple, Serialize_tuple};
 
@@ -105,6 +106,44 @@ pub enum MmKwargValue {
     List(Vec<MmKwargValue>),
 }
 
+impl MmFeatureSpec {
+    /// Extract large tensor buffers from this feature in serialized field order.
+    pub(crate) fn extract_aux_frames(&mut self, aux_frames: &mut Vec<Bytes>, threshold: usize) {
+        if let Some(data) = &mut self.data {
+            for elem in data.values_mut() {
+                elem.extract_aux_frames(aux_frames, threshold);
+            }
+        }
+        if let Some(is_embed) = &mut self.mm_position.is_embed {
+            is_embed.extract_aux_frame(aux_frames, threshold);
+        }
+    }
+}
+
+impl MmFieldElem {
+    /// Extract large tensor buffers from this field element.
+    fn extract_aux_frames(&mut self, aux_frames: &mut Vec<Bytes>, threshold: usize) {
+        if let Some(data) = &mut self.data {
+            data.extract_aux_frames(aux_frames, threshold);
+        }
+    }
+}
+
+impl MmKwargValue {
+    /// Recursively extract large tensor buffers from this nested value.
+    fn extract_aux_frames(&mut self, aux_frames: &mut Vec<Bytes>, threshold: usize) {
+        match self {
+            Self::Tensor(tensor) => tensor.extract_aux_frame(aux_frames, threshold),
+            Self::List(values) => {
+                for value in values {
+                    value.extract_aux_frames(aux_frames, threshold);
+                }
+            }
+            Self::Int(_) | Self::Float(_) => {}
+        }
+    }
+}
+
 /// Defines how to interpret tensor data belonging to a keyword argument for
 /// `MultiModalKwargsItems`, and vice versa.
 ///

rust/src/engine-core-client/src/protocol/request.rs

@@ -137,6 +137,17 @@ impl EngineCoreRequest {
         }
         Ok(())
     }
+
+    /// Extract large request tensors into ordered auxiliary frames.
+    pub(crate) fn extract_aux_frames(&mut self, threshold: usize) -> Vec<Bytes> {
+        let mut aux_frames = Vec::new();
+        if let Some(features) = &mut self.mm_features {
+            for feature in features {
+                feature.extract_aux_frames(&mut aux_frames, threshold);
+            }
+        }
+        aux_frames
+    }
 }
 
 #[cfg(test)]
@@ -144,9 +155,15 @@ mod tests {
     use rmpv::Value;
 
     use super::*;
+    use crate::protocol::multimodal::{
+        MmBatchedField, MmFeatureSpec, MmField, MmFieldElem, MmKwargValue, PlaceholderRange,
+    };
     use crate::protocol::sampling::EngineCoreSamplingParams;
+    use crate::protocol::tensor::{WireArrayData, WireTensor};
     use crate::protocol::{decode_value, encode_msgpack};
 
+    const AUX_FRAME_THRESHOLD: usize = 256;
+
     #[test]
     fn engine_core_request_serializes_as_full_array() {
         let request = EngineCoreRequest {
@@ -175,4 +192,84 @@ mod tests {
         assert_eq!(array[10], Value::Nil);
         assert_eq!(array[11], Value::from(7));
     }
+
+    #[test]
+    fn engine_core_request_extracts_large_nested_tensors_in_wire_order() {
+        let inline = vec![1_u8; AUX_FRAME_THRESHOLD - 1];
+        let first_aux = vec![2_u8; AUX_FRAME_THRESHOLD];
+        let second_aux = vec![3_u8; AUX_FRAME_THRESHOLD + 1];
+        let first_aux_ptr = first_aux.as_ptr();
+        let second_aux_ptr = second_aux.as_ptr();
+        let mut request = EngineCoreRequest {
+            mm_features: Some(vec![MmFeatureSpec {
+                data: Some(BTreeMap::from([
+                    (
+                        "inline".to_string(),
+                        MmFieldElem {
+                            data: Some(MmKwargValue::Tensor(WireTensor::from_raw(
+                                "uint8",
+                                vec![inline.len()],
+                                inline,
+                            ))),
+                            field: MmField::Batched(MmBatchedField { keep_on_cpu: false }),
+                        },
+                    ),
+                    (
+                        "nested".to_string(),
+                        MmFieldElem {
+                            data: Some(MmKwargValue::List(vec![
+                                MmKwargValue::Int(7),
+                                MmKwargValue::Tensor(WireTensor::from_raw(
+                                    "uint8",
+                                    vec![first_aux.len()],
+                                    first_aux,
+                                )),
+                            ])),
+                            field: MmField::Batched(MmBatchedField { keep_on_cpu: false }),
+                        },
+                    ),
+                ])),
+                modality: "image".to_string(),
+                identifier: "id".to_string(),
+                mm_position: PlaceholderRange {
+                    offset: 0,
+                    length: second_aux.len(),
+                    is_embed: Some(WireTensor::from_raw(
+                        "bool",
+                        vec![second_aux.len()],
+                        second_aux,
+                    )),
+                },
+                mm_hash: None,
+            }]),
+            ..EngineCoreRequest::default()
+        };
+
+        let aux_frames = request.extract_aux_frames(AUX_FRAME_THRESHOLD);
+
+        assert_eq!(aux_frames.len(), 2);
+        assert_eq!(aux_frames[0].as_ptr(), first_aux_ptr);
+        assert_eq!(aux_frames[1].as_ptr(), second_aux_ptr);
+        let feature = &request.mm_features.as_ref().unwrap()[0];
+        let MmKwargValue::Tensor(inline) =
+            feature.data.as_ref().unwrap()["inline"].data.as_ref().unwrap()
+        else {
+            panic!("expected inline tensor");
+        };
+        assert!(matches!(inline.data, WireArrayData::RawView(_)));
+        let MmKwargValue::List(nested) =
+            feature.data.as_ref().unwrap()["nested"].data.as_ref().unwrap()
+        else {
+            panic!("expected nested tensor list");
+        };
+        let MmKwargValue::Tensor(nested_tensor) = &nested[1] else {
+            panic!("expected nested tensor");
+        };
+        assert_eq!(nested_tensor.data, WireArrayData::AuxIndex(1));
+        assert_eq!(
+            feature.mm_position.is_embed.as_ref().unwrap().data,
+            WireArrayData::AuxIndex(2)
+        );
+        assert!(request.extract_aux_frames(AUX_FRAME_THRESHOLD).is_empty());
+    }
 }

rust/src/engine-core-client/src/protocol/tensor.rs

@@ -1,7 +1,8 @@
 // SPDX-License-Identifier: Apache-2.0
 // SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
-use bytemuck::allocation::pod_collect_to_vec;
+use bytemuck::{Pod, cast_slice};
+use bytes::Bytes;
 use enum_as_inner::EnumAsInner;
 use half::{bf16, f16};
 use rmpv::Value;
@@ -20,6 +21,21 @@ struct MsgpackExtRef<'a>((i8, ByteSlice<'a>));
 
 struct ByteSlice<'a>(&'a [u8]);
 
+struct PodVec<T: Pod>(Vec<T>);
+
+impl<T: Pod> AsRef<[u8]> for PodVec<T> {
+    fn as_ref(&self) -> &[u8] {
+        cast_slice(&self.0)
+    }
+}
+
+fn bytes_from_pod_vec<T>(data: Vec<T>) -> Bytes
+where
+    T: Pod + Send + 'static,
+{
+    Bytes::from_owner(PodVec(data))
+}
+
 impl Serialize for ByteSlice<'_> {
     fn serialize<S>(&self, serializer: S) -> std::result::Result<S::Ok, S::Error>
     where
@@ -55,58 +71,63 @@ pub struct WireNdArray {
 
 impl WireNdArray {
     /// Build a float32 tensor/ndarray backed by native-endian raw-view bytes.
-    pub fn from_f32(shape: Vec<usize>, data: impl AsRef<[f32]>) -> Result<Self, String> {
-        let data = data.as_ref();
+    ///
+    /// Takes ownership of the backing buffer without copying its data.
+    pub fn from_f32(shape: Vec<usize>, data: Vec<f32>) -> Result<Self, String> {
         validate_element_count(&shape, data.len())?;
-        Ok(Self {
-            dtype: "float32".to_string(),
+        Ok(Self::from_raw_bytes(
+            "float32",
             shape,
-            data: WireArrayData::RawView(pod_collect_to_vec::<f32, u8>(data)),
-        })
+            bytes_from_pod_vec(data),
+        ))
     }
 
     /// Build a float16 tensor/ndarray backed by native-endian raw-view bytes.
-    pub fn from_f16(shape: Vec<usize>, data: impl AsRef<[f16]>) -> Result<Self, String> {
-        let data = data.as_ref();
+    ///
+    /// Takes ownership of the backing buffer without copying its data.
+    pub fn from_f16(shape: Vec<usize>, data: Vec<f16>) -> Result<Self, String> {
         validate_element_count(&shape, data.len())?;
-        Ok(Self {
-            dtype: "float16".to_string(),
+        Ok(Self::from_raw_bytes(
+            "float16",
             shape,
-            data: WireArrayData::RawView(pod_collect_to_vec::<f16, u8>(data)),
-        })
+            bytes_from_pod_vec(data),
+        ))
     }
 
     /// Build a bfloat16 tensor/ndarray backed by native-endian raw-view bytes.
-    pub fn from_bf16(shape: Vec<usize>, data: impl AsRef<[bf16]>) -> Result<Self, String> {
-        let data = data.as_ref();
+    ///
+    /// Takes ownership of the backing buffer without copying its data.
+    pub fn from_bf16(shape: Vec<usize>, data: Vec<bf16>) -> Result<Self, String> {
         validate_element_count(&shape, data.len())?;
-        Ok(Self {
-            dtype: "bfloat16".to_string(),
+        Ok(Self::from_raw_bytes(
+            "bfloat16",
             shape,
-            data: WireArrayData::RawView(pod_collect_to_vec::<bf16, u8>(data)),
-        })
+            bytes_from_pod_vec(data),
+        ))
     }
 
     /// Build an int64 tensor/ndarray backed by native-endian raw-view bytes.
-    pub fn from_i64(shape: Vec<usize>, data: impl AsRef<[i64]>) -> Result<Self, String> {
-        let data = data.as_ref();
+    ///
+    /// Takes ownership of the backing buffer without copying its data.
+    pub fn from_i64(shape: Vec<usize>, data: Vec<i64>) -> Result<Self, String> {
         validate_element_count(&shape, data.len())?;
-        Ok(Self {
-            dtype: "int64".to_string(),
+        Ok(Self::from_raw_bytes(
+            "int64",
             shape,
-            data: WireArrayData::RawView(pod_collect_to_vec::<i64, u8>(data)),
-        })
+            bytes_from_pod_vec(data),
+        ))
     }
 
     /// Build a uint32 tensor/ndarray backed by native-endian raw-view bytes.
-    pub fn from_u32(shape: Vec<usize>, data: impl AsRef<[u32]>) -> Result<Self, String> {
-        let data = data.as_ref();
+    ///
+    /// Takes ownership of the backing buffer without copying its data.
+    pub fn from_u32(shape: Vec<usize>, data: Vec<u32>) -> Result<Self, String> {
         validate_element_count(&shape, data.len())?;
-        Ok(Self {
-            dtype: "uint32".to_string(),
+        Ok(Self::from_raw_bytes(
+            "uint32",
             shape,
-            data: WireArrayData::RawView(pod_collect_to_vec::<u32, u8>(data)),
-        })
+            bytes_from_pod_vec(data),
+        ))
     }
 
     /// Build a bool tensor/ndarray backed by raw-view bytes.
@@ -118,7 +139,9 @@ impl WireNdArray {
         Ok(Self {
             dtype: "bool".to_string(),
             shape,
-            data: WireArrayData::RawView(data.iter().map(|value| u8::from(*value)).collect()),
+            data: WireArrayData::RawView(Bytes::from(
+                data.into_iter().map(u8::from).collect::<Vec<_>>(),
+            )),
         })
     }
 
@@ -127,12 +150,22 @@ impl WireNdArray {
     /// Use this as an escape hatch when the caller already owns bytes that
     /// match the requested `dtype` and `shape`.
     pub fn from_raw(dtype: impl Into<String>, shape: Vec<usize>, data: Vec<u8>) -> Self {
+        Self::from_raw_bytes(dtype, shape, Bytes::from(data))
+    }
+
+    /// Build a tensor/ndarray from an owned immutable raw-view buffer.
+    pub fn from_raw_bytes(dtype: impl Into<String>, shape: Vec<usize>, data: Bytes) -> Self {
         Self {
             dtype: dtype.into(),
             shape,
             data: WireArrayData::RawView(data),
         }
     }
+
+    /// Move a sufficiently large inline buffer into the ordered auxiliary-frame list.
+    pub(crate) fn extract_aux_frame(&mut self, aux_frames: &mut Vec<Bytes>, threshold: usize) {
+        self.data.extract_aux_frame(aux_frames, threshold);
+    }
 }
 
 /// Validate that the number of elements implied by the shape matches the length
@@ -170,7 +203,25 @@ pub enum WireArrayData {
     /// stored.
     AuxIndex(usize),
     /// The raw bytes of this array/tensor.
-    RawView(Vec<u8>),
+    RawView(Bytes),
+}
+
+impl WireArrayData {
+    /// Replace a sufficiently large raw view with its one-based auxiliary-frame index.
+    fn extract_aux_frame(&mut self, aux_frames: &mut Vec<Bytes>, threshold: usize) {
+        let Self::RawView(bytes) = self else {
+            return;
+        };
+        if bytes.len() < threshold {
+            return;
+        }
+
+        let index = aux_frames.len() + 1;
+        let bytes = std::mem::replace(self, Self::AuxIndex(index))
+            .into_raw_view()
+            .expect("raw view was matched above");
+        aux_frames.push(bytes);
+    }
 }
 
 impl<'de> Deserialize<'de> for WireArrayData {
@@ -180,7 +231,9 @@ impl<'de> Deserialize<'de> for WireArrayData {
     {
         let value = Value::deserialize(deserializer)?;
         match value {
-            Value::Ext(tag, bytes) if tag == CUSTOM_TYPE_RAW_VIEW => Ok(Self::RawView(bytes)),
+            Value::Ext(tag, bytes) if tag == CUSTOM_TYPE_RAW_VIEW => {
+                Ok(Self::RawView(Bytes::from(bytes)))
+            }
             Value::Ext(tag, _) => Err(serde::de::Error::custom(format!(
                 "unsupported extension type code {tag}"
             ))),
@@ -201,9 +254,6 @@ impl Serialize for WireArrayData {
     where
         S: Serializer,
     {
-        // TODO: outbound request serialization currently only supports inline
-        // raw-view bytes. Emitting aux frames needs transport-level plumbing;
-        // serializing `AuxIndex` here only preserves an already-built reference.
         match self {
             Self::AuxIndex(index) => serializer.serialize_u64(*index as u64),
             Self::RawView(bytes) => {
@@ -221,7 +271,7 @@ mod tests {
     fn raw_view_serializes_as_msgpack_ext() {
         let bytes = vec![1, 2, 3, 4];
         let encoded =
-            rmp_serde::to_vec_named(&WireArrayData::RawView(bytes.clone())).expect("encode");
+            rmp_serde::to_vec_named(&WireArrayData::RawView(bytes.clone().into())).expect("encode");
         let expected = rmp_serde::to_vec_named(&Value::Ext(CUSTOM_TYPE_RAW_VIEW, bytes.clone()))
             .expect("encode expected");
 
@@ -234,11 +284,15 @@ mod tests {
 
     #[test]
     fn constructors_build_raw_view_tensors() {
-        let f32_tensor = WireNdArray::from_f32(vec![2], vec![1.0, 2.5]).unwrap();
+        let f32_data = vec![1.0, 2.5];
+        let f32_data_ptr = f32_data.as_ptr().cast::<u8>();
+        let f32_tensor = WireNdArray::from_f32(vec![2], f32_data).unwrap();
         assert_eq!(f32_tensor.dtype, "float32");
         assert_eq!(f32_tensor.shape, vec![2]);
+        let f32_raw_view = f32_tensor.data.into_raw_view().expect("raw view");
+        assert_eq!(f32_raw_view.as_ptr(), f32_data_ptr);
         assert_eq!(
-            f32_tensor.data.into_raw_view().expect("raw view"),
+            f32_raw_view,
             [1.0_f32, 2.5].into_iter().flat_map(f32::to_ne_bytes).collect::<Vec<_>>()
         );
 
@@ -258,15 +312,15 @@ mod tests {
         let i64_tensor = WireNdArray::from_i64(vec![1], vec![-7]).unwrap();
         assert_eq!(i64_tensor.dtype, "int64");
         assert_eq!(
-            i64_tensor.data.into_raw_view().expect("raw view"),
-            (-7_i64).to_ne_bytes()
+            i64_tensor.data.into_raw_view().expect("raw view").as_ref(),
+            (-7_i64).to_ne_bytes().as_ref()
         );
 
         let u32_tensor = WireNdArray::from_u32(vec![1], vec![42]).unwrap();
         assert_eq!(u32_tensor.dtype, "uint32");
         assert_eq!(
-            u32_tensor.data.into_raw_view().expect("raw view"),
-            42_u32.to_ne_bytes()
+            u32_tensor.data.into_raw_view().expect("raw view").as_ref(),
+            42_u32.to_ne_bytes().as_ref()
         );
 
         let bool_tensor = WireNdArray::from_bool(vec![2], vec![false, true]).unwrap();

rust/src/engine-core-client/src/tests/client.rs

@@ -32,7 +32,7 @@ use crate::protocol::output::{
 use crate::protocol::request::{EngineCoreRequest, EngineCoreRequestType};
 use crate::protocol::sampling::EngineCoreSamplingParams;
 use crate::protocol::stats::SchedulerStats;
-use crate::protocol::tensor::WireTensor;
+use crate::protocol::tensor::{WireArrayData, WireTensor};
 use crate::protocol::utility::{UtilityOutput, UtilityResultEnvelope};
 use crate::test_utils::{
     IpcNamespace, setup_bootstrapped_mock_engine, setup_mock_engine_sockets,
@@ -1728,6 +1728,86 @@ async fn client_decodes_multipart_logprob_outputs() {
     client.shutdown().await.unwrap();
 }
 
+#[tokio::test(flavor = "multi_thread", worker_threads = 2)]
+async fn client_sends_large_multimodal_tensor_as_aux_frame() {
+    init_tracing();
+    let ipc = IpcNamespace::new().unwrap();
+    let handshake_address = ipc.handshake_endpoint();
+    let engine_id = b"engine-multimodal-aux".to_vec();
+    let tensor_data = (0..64).map(|value| value as f32).collect::<Vec<_>>();
+    let expected_bytes =
+        tensor_data.iter().flat_map(|value| value.to_ne_bytes()).collect::<Vec<_>>();
+    let mut request = sample_multimodal_request();
+    request.mm_features.as_mut().unwrap()[0]
+        .data
+        .as_mut()
+        .unwrap()
+        .get_mut("pixel_values")
+        .unwrap()
+        .data = Some(MmKwargValue::Tensor(
+        WireTensor::from_f32(vec![64], tensor_data).unwrap(),
+    ));
+
+    let (shutdown_tx, engine_task) = spawn_mock_engine_task(
+        handshake_address.clone(),
+        engine_id.clone(),
+        move |dealer, push| {
+            Box::pin(async move {
+                let add = recv_engine_message(dealer).await;
+                assert_eq!(add.len(), 3);
+                assert_eq!(add[0].as_ref(), &[0x00]);
+                let request: EngineCoreRequest = rmp_serde::from_slice(&add[1]).unwrap();
+                let MmKwargValue::Tensor(tensor) =
+                    request.mm_features.as_ref().unwrap()[0].data.as_ref().unwrap()["pixel_values"]
+                        .data
+                        .as_ref()
+                        .unwrap()
+                else {
+                    panic!("expected tensor");
+                };
+                assert_eq!(tensor.data, WireArrayData::AuxIndex(1));
+                assert_eq!(add[2].as_ref(), expected_bytes);
+
+                send_outputs(
+                    push,
+                    RequestBatchOutputs {
+                        outputs: vec![request_output(
+                            "req-mm",
+                            vec![],
+                            Some(EngineCoreFinishReason::Length),
+                        )],
+                        finished_requests: Some(BTreeSet::from(["req-mm".to_string()])),
+                        ..Default::default()
+                    }
+                    .into(),
+                )
+                .await;
+            })
+        },
+    );
+
+    let client = connect_client_with_ipc(
+        handshake_test_config(
+            handshake_address,
+            1,
+            "test-model",
+            Duration::from_secs(2),
+            0,
+            None,
+        ),
+        &ipc,
+    )
+    .await;
+
+    let outputs = client.call(request).await.unwrap().collect::<Vec<_>>().await;
+    assert_eq!(outputs.len(), 1);
+    assert!(outputs[0].is_ok());
+
+    let _ = shutdown_tx.send(());
+    engine_task.await.unwrap();
+    client.shutdown().await.unwrap();
+}
+
 #[tokio::test(flavor = "multi_thread", worker_threads = 2)]
 async fn multi_engine_client_shares_transport_and_routes_by_inflight_count() {
     init_tracing();

rust/src/engine-core-client/src/transport.rs

@@ -512,14 +512,14 @@ pub async fn send_message(
     input_send: &mut RouterSendHalf,
     engine_id: &EngineId,
     request_type: Bytes,
-    payload: Vec<u8>,
+    payload: Bytes,
+    aux_frames: Vec<Bytes>,
 ) -> Result<()> {
-    let message = ZmqMessage::try_from(vec![
-        engine_id.to_frame(),
-        request_type,
-        Bytes::from(payload),
-    ])
-    .expect("router messages must contain identity and payload");
+    let mut frames = Vec::with_capacity(3 + aux_frames.len());
+    frames.extend([engine_id.to_frame(), request_type, payload]);
+    frames.extend(aux_frames);
+    let message =
+        ZmqMessage::try_from(frames).expect("router messages must contain identity and payload");
 
     trace!(
         ?engine_id,

rust/src/managed-engine/src/cli.rs

@@ -91,6 +91,7 @@ impl ManagedEngineArgs {
         disable_log_stats: bool,
         shutdown_timeout: u64,
         handshake_port: u16,
+        limit_mm_per_prompt: Option<String>,
     ) -> ManagedEngineConfig {
         let mut python_args = self.python_args;
         // Manually forward some args to the Python engine.
@@ -126,6 +127,10 @@ impl ManagedEngineArgs {
             python_args.push("--data-parallel-size-local".to_string());
             python_args.push(data_parallel_size_local.to_string());
         }
+        if let Some(limit_mm_per_prompt) = limit_mm_per_prompt {
+            python_args.push("--limit-mm-per-prompt".to_string());
+            python_args.push(limit_mm_per_prompt);
+        }
 
         ManagedEngineConfig {
             python: self.python,

rust/src/server/examples/external_engine_openai_qwen.rs

@@ -1,6 +1,7 @@
 // SPDX-License-Identifier: Apache-2.0
 // SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+use std::collections::HashMap;
 use std::time::Duration;
 
 use anyhow::{Context, Result, bail};
@@ -70,6 +71,7 @@ async fn main() -> Result<()> {
         language_model_only: false,
         chat_template: None,
         default_chat_template_kwargs: None,
+        limit_mm_per_prompt: HashMap::new(),
         chat_template_content_format: ChatTemplateContentFormatOption::Auto,
         max_logprobs: None,
         api_server_options: ApiServerOptions::default(),

rust/src/server/src/config.rs

@@ -10,6 +10,7 @@ use axum::http::{HeaderName, HeaderValue, Method};
 use educe::Educe;
 use serde::Serialize;
 use serde_json::Value;
+use vllm_chat::multimodal::MmLimitPerPrompt;
 use vllm_chat::{ChatTemplateContentFormatOption, ParserSelection, RendererSelection};
 use vllm_engine_core_client::{CoordinatorMode as EngineCoreCoordinatorMode, TransportMode};
 
@@ -184,6 +185,9 @@ pub struct Config {
     pub chat_template: Option<String>,
     /// Server-default keyword arguments merged into every chat-template render.
     pub default_chat_template_kwargs: Option<HashMap<String, Value>>,
+    /// Maximum number of input items allowed per prompt for each modality.
+    /// Unspecified modalities are unlimited.
+    pub limit_mm_per_prompt: MmLimitPerPrompt,
     /// How to serialize `message.content` for chat-template rendering.
     pub chat_template_content_format: ChatTemplateContentFormatOption,
     /// Optional maximum number of top log probabilities accepted by the

rust/src/server/src/lib.rs

@@ -101,6 +101,7 @@ async fn build_state(config: &Config) -> Result<Arc<AppState>> {
                 .default_chat_template_kwargs
                 .clone()
                 .unwrap_or_default(),
+            limit_mm_per_prompt: config.limit_mm_per_prompt.clone(),
         },
     )
     .await

rust/src/server/src/routes/tests.rs

@@ -577,6 +577,12 @@ fn render_fake_content(content: &ChatContent, placeholder: &str) -> vllm_chat::R
 }
 
 fn qwen_multimodal_model_info() -> vllm_chat::multimodal::MultimodalModelInfo {
+    qwen_multimodal_model_info_with_limits(std::collections::HashMap::new())
+}
+
+fn qwen_multimodal_model_info_with_limits(
+    limit_mm_per_prompt: vllm_chat::multimodal::MmLimitPerPrompt,
+) -> vllm_chat::multimodal::MultimodalModelInfo {
     let config_path = std::env::temp_dir().join(format!(
         "vllm-server-qwen-config-{}.json",
         uuid::Uuid::new_v4()
@@ -594,6 +600,7 @@ fn qwen_multimodal_model_info() -> vllm_chat::multimodal::MultimodalModelInfo {
             ..Default::default()
         },
         Arc::new(fake_chat_tokenizer()),
+        limit_mm_per_prompt,
     )
     .expect("load multimodal info")
     .expect("qwen multimodal info is registered");
@@ -2307,6 +2314,74 @@ async fn non_stream_chat_image_url_reaches_engine_mm_features() {
     assert_eq!(json["choices"][0]["message"]["content"], "hi");
 }
 
+#[tokio::test(flavor = "multi_thread", worker_threads = 2)]
+#[serial]
+async fn non_stream_chat_rejects_when_image_count_exceeds_limit_mm_per_prompt() {
+    // The request is rejected by `--limit-mm-per-prompt` validation before
+    // ever reaching the engine, so the mock engine task is never awaited.
+    let (chat, _engine_task) = test_models_with_engine_outputs_and_backend(
+        b"engine-openai-mm-limit",
+        default_stream_output_specs(),
+        Arc::new(FakeChatBackend::with_multimodal_model_info(
+            qwen_multimodal_model_info_with_limits(std::collections::HashMap::from([(
+                vllm_chat::multimodal::MmLimitModality::Image,
+                vllm_chat::multimodal::MmLimitSpec::Count(1),
+            )])),
+        )),
+    )
+    .await;
+    let app = build_router(Arc::new(AppState::new(
+        vec!["Qwen/Qwen1.5-0.5B-Chat".to_string()],
+        chat,
+    )));
+
+    let response = app
+        .clone()
+        .call(
+            Request::builder()
+                .method("POST")
+                .uri("/v1/chat/completions")
+                .header("content-type", "application/json")
+                .body(Body::from(
+                    json!({
+                        "model": "Qwen/Qwen1.5-0.5B-Chat",
+                        "stream": false,
+                        "messages": [{
+                            "role": "user",
+                            "content": [
+                                {"type": "text", "text": "describe "},
+                                {
+                                    "type": "image_url",
+                                    "image_url": {
+                                        "url": "data:image/png;base64,iVBORw0KGgoAAAANSUhEUgAAAAEAAAABCAQAAAC1HAwCAAAAC0lEQVR42mNk+A8AAQUBAScY42YAAAAASUVORK5CYII="
+                                    }
+                                },
+                                {
+                                    "type": "image_url",
+                                    "image_url": {
+                                        "url": "data:image/png;base64,iVBORw0KGgoAAAANSUhEUgAAAAEAAAABCAQAAAC1HAwCAAAAC0lEQVR42mNk+A8AAQUBAScY42YAAAAASUVORK5CYII="
+                                    }
+                                }
+                            ]
+                        }]
+                    })
+                    .to_string(),
+                ))
+                .expect("build request"),
+        )
+        .await
+        .expect("call app");
+
+    assert_eq!(response.status(), StatusCode::BAD_REQUEST);
+    let body = to_bytes(response.into_body(), usize::MAX).await.expect("read body");
+    let json: serde_json::Value = serde_json::from_slice(&body).expect("decode json");
+    assert_eq!(json["error"]["type"], "invalid_request_error");
+    assert_eq!(
+        json["error"]["message"],
+        "At most 1 image(s) may be provided in one prompt."
+    );
+}
+
 #[tokio::test(flavor = "multi_thread", worker_threads = 2)]
 #[serial]
 async fn non_stream_chat_includes_logprobs_and_prompt_logprobs() {

setup.py

@@ -1122,7 +1122,7 @@ if _is_cuda() or _is_hip():
     # copying the relevant .py files from the source repository.
     ext_modules.append(CMakeExtension(name="vllm.triton_kernels", optional=True))
 
-if sys.version_info >= (3, 11):
+if not _is_xpu() and sys.version_info >= (3, 11):
     ext_modules.append(CMakeExtension(name="vllm.spinloop"))
     ext_modules.append(CMakeExtension(name="vllm.fs_io_C"))
 

tests/config/test_speculative_draft_max_position_embeddings.py

@@ -0,0 +1,115 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Tests for the EAGLE draft ``max_position_embeddings`` override (#48894).
+
+EAGLE drafts share the target's positional space, but some draft
+checkpoints (e.g. ``yuhuili/EAGLE3-LLaMA3.1-Instruct-8B``) ship a
+``max_position_embeddings`` (2048) far smaller than the target's context.
+That value sizes the draft's rotary ``cos_sin_cache`` while the proposer
+feeds positions up to the target's ``max_model_len``, so the cache gather
+goes out of bounds — a device-side assert under torch.compile and silent
+garbage reads in eager mode. ``SpeculativeConfig`` must raise the draft's
+value to the target's ``max_model_len``, with a log, for the eagle/eagle3
+methods only.
+"""
+
+import logging
+
+import pytest
+from transformers import PretrainedConfig
+
+from vllm.config.model import ModelConfig
+from vllm.config.parallel import ParallelConfig
+from vllm.config.speculative import SpeculativeConfig
+
+# All repos are public; only config/tokenizer-config files are fetched.
+EAGLE3_DRAFT = "yuhuili/EAGLE3-LLaMA3.1-Instruct-8B"  # max_position_embeddings=2048
+LLAMA3_TARGET = "unsloth/Meta-Llama-3.1-8B-Instruct"  # max_position_embeddings=131072
+AR_MODEL = "JackFram/llama-68m"  # max_position_embeddings=2048
+
+_LOGGER = "vllm.config.speculative"
+_OVERRIDE_MSG = "Overriding draft model max_position_embeddings"
+
+
+@pytest.fixture
+def vllm_caplog(caplog: pytest.LogCaptureFixture, monkeypatch: pytest.MonkeyPatch):
+    """Make caplog see vLLM logger records (vLLM sets propagate=False)."""
+    monkeypatch.setattr(logging.getLogger("vllm"), "propagate", True)
+    with caplog.at_level(logging.INFO, logger=_LOGGER):
+        yield caplog
+
+
+def _override_logged(caplog: pytest.LogCaptureFixture) -> bool:
+    return any(_OVERRIDE_MSG in record.getMessage() for record in caplog.records)
+
+
+@pytest.mark.cpu_test
+def test_override_raises_smaller_value(vllm_caplog: pytest.LogCaptureFixture):
+    hf_config = PretrainedConfig(max_position_embeddings=2048)
+    SpeculativeConfig._maybe_override_draft_max_position_embeddings(
+        hf_config, target_max_model_len=8192
+    )
+    assert hf_config.max_position_embeddings == 8192
+    assert _override_logged(vllm_caplog)
+
+
+@pytest.mark.cpu_test
+def test_override_keeps_sufficient_value(vllm_caplog: pytest.LogCaptureFixture):
+    hf_config = PretrainedConfig(max_position_embeddings=8192)
+    SpeculativeConfig._maybe_override_draft_max_position_embeddings(
+        hf_config, target_max_model_len=8192
+    )
+    assert hf_config.max_position_embeddings == 8192
+    assert not _override_logged(vllm_caplog)
+
+
+@pytest.mark.cpu_test
+def test_override_ignores_missing_attribute(vllm_caplog: pytest.LogCaptureFixture):
+    hf_config = PretrainedConfig()
+    hf_config.__dict__.pop("max_position_embeddings", None)
+    SpeculativeConfig._maybe_override_draft_max_position_embeddings(
+        hf_config, target_max_model_len=8192
+    )
+    assert not hasattr(hf_config, "max_position_embeddings")
+    assert not _override_logged(vllm_caplog)
+
+
+@pytest.mark.cpu_test
+@pytest.mark.parametrize("method", ["eagle", "eagle3"])
+def test_eagle_draft_inherits_target_max_model_len(
+    method: str, vllm_caplog: pytest.LogCaptureFixture
+):
+    target_model_config = ModelConfig(LLAMA3_TARGET)
+    assert target_model_config.max_model_len > 2048
+    speculative_config = SpeculativeConfig(
+        target_model_config=target_model_config,
+        target_parallel_config=ParallelConfig(),
+        model=EAGLE3_DRAFT,
+        method=method,
+        num_speculative_tokens=3,
+    )
+    draft_hf_config = speculative_config.draft_model_config.hf_config
+    assert draft_hf_config.max_position_embeddings == target_model_config.max_model_len
+    assert _override_logged(vllm_caplog)
+
+
+@pytest.mark.cpu_test
+def test_independent_draft_model_keeps_its_own_limit(
+    vllm_caplog: pytest.LogCaptureFixture,
+):
+    """An independent AR draft may genuinely have a smaller context than the
+    target; its max_position_embeddings must not be resized."""
+    target_model_config = ModelConfig(
+        AR_MODEL, hf_overrides={"max_position_embeddings": 8192}
+    )
+    assert target_model_config.max_model_len == 8192
+    speculative_config = SpeculativeConfig(
+        target_model_config=target_model_config,
+        target_parallel_config=ParallelConfig(),
+        model=AR_MODEL,
+        method="draft_model",
+        num_speculative_tokens=3,
+    )
+    draft_hf_config = speculative_config.draft_model_config.hf_config
+    assert draft_hf_config.max_position_embeddings == 2048
+    assert not _override_logged(vllm_caplog)

tests/kernels/mamba/test_precopy_mamba_align.py

@@ -2,9 +2,9 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 """Equivalence test for ``precopy_mamba_align_fused_kernel``.
 
-The V2 "align" pre-copy must migrate mamba state across block boundaries with
-byte-identical semantics to the V1 copy specs (``get_conv_copy_spec`` /
-``get_temporal_copy_spec``):
+The fused "align" pre-copy must migrate mamba state across block boundaries
+with byte-identical semantics to the scalar V1 copy specs
+(``get_conv_copy_spec`` / ``get_temporal_copy_spec``):
 
 * conv state (SD layout, conv_width > 0): shift the sliding window by
   ``token_bias`` tokens -- ``state[bt[src_col], token_bias:]`` ->
@@ -14,20 +14,27 @@ byte-identical semantics to the V1 copy specs (``get_conv_copy_spec`` /
   ``state[bt[dst_col]]``.
 
 The kernel must also no-op when ``src_col < 0`` (fresh request) or
-``src_col == dst_col`` (no boundary crossed).
+``src_col == dst_col`` (no boundary crossed). V2 callers pass an explicit
+``idx_mapping``; V1 align preprocessing launches in batch order with
+``idx_mapping=None``.
 """
 
 from __future__ import annotations
 
+from types import SimpleNamespace
+
+import numpy as np
 import torch
 
+from vllm.model_executor.layers.mamba import mamba_utils as layer_mamba_utils
 from vllm.platforms import current_platform
+from vllm.v1.worker import mamba_utils as worker_mamba_utils
 from vllm.v1.worker.mamba_utils import precopy_mamba_align_fused_kernel
 
 try:
     import pytest
 
-    pytestmark = pytest.mark.skipif(
+    _cuda_required = pytest.mark.skipif(
         not current_platform.is_cuda(),
         reason="precopy_mamba_align_fused_kernel needs CUDA/Triton",
     )
@@ -35,6 +42,9 @@ try:
 except ModuleNotFoundError:  # allow running directly as ``python <thisfile>``
     pytest = None
 
+    def _cuda_required(fn):
+        return fn
+
     def _parametrize(_name, _values):
         def _deco(fn):
             return fn
@@ -49,13 +59,20 @@ SSM_SHAPE = (4, 16, 16)
 MAX_COLS = 8
 
 
-def _build_state(num_blocks, device):
+def _build_state(num_blocks, device, conv_state_dim_first):
     """Per-layer (conv SD [nb, width, dim] bf16, ssm [nb, *shape] fp32) pools."""
     convs, ssms = [], []
     for _ in range(NUM_LAYERS):
+        conv_shape = (
+            (num_blocks, CONV_DIM, CONV_WIDTH)
+            if conv_state_dim_first
+            else (num_blocks, CONV_WIDTH, CONV_DIM)
+        )
         convs.append(
             torch.randn(
-                num_blocks, CONV_WIDTH, CONV_DIM, dtype=torch.bfloat16, device=device
+                *conv_shape,
+                dtype=torch.bfloat16,
+                device=device,
             )
         )
         ssms.append(
@@ -64,7 +81,7 @@ def _build_state(num_blocks, device):
     return convs, ssms
 
 
-def _build_meta(convs, ssms, device):
+def _build_meta(convs, ssms, device, conv_state_dim_first):
     """Flattened per-(layer, state-type) metadata, ordered conv, ssm per layer."""
     n = NUM_LAYERS * 2
     base = torch.zeros(n, dtype=torch.int64, device=device)
@@ -82,8 +99,14 @@ def _build_meta(convs, ssms, device):
         base[i] = conv.data_ptr()
         blk_stride[i] = conv.stride(0) * conv.element_size()
         elem[i] = conv.element_size()
-        width[i] = conv.size(1)
-        inner[i] = conv.stride(1)
+        if conv_state_dim_first:
+            width[i] = conv.size(2)
+            inner[i] = 1
+            drc[i] = conv.size(1)
+            drs[i] = conv.stride(1) * conv.element_size()
+        else:
+            width[i] = conv.size(1)
+            inner[i] = conv.stride(1)
         i += 1
         # ssm (temporal): width = 0, inner = elems per block
         base[i] = ssm.data_ptr()
@@ -95,7 +118,7 @@ def _build_meta(convs, ssms, device):
     return base, blk_stride, elem, inner, width, group, drc, drs
 
 
-def _reference(convs, ssms, bt, src_col, dst_col, bias, num_reqs):
+def _reference(convs, ssms, bt, src_col, dst_col, bias, num_reqs, conv_dim_first):
     """Apply the V1 copy semantics on clones, reading from the pre-copy state."""
     conv_pre = [c.clone() for c in convs]
     ssm_pre = [s.clone() for s in ssms]
@@ -108,14 +131,24 @@ def _reference(convs, ssms, bt, src_col, dst_col, bias, num_reqs):
         sblk, dblk = int(bt[r, sc]), int(bt[r, dc])
         tblk = int(bt[r, sc + tb])  # temporal src column shifted by bias
         for layer in range(NUM_LAYERS):
-            conv_ref[layer][dblk, : CONV_WIDTH - tb] = conv_pre[layer][sblk, tb:]
+            if conv_dim_first:
+                conv_ref[layer][dblk, :, : CONV_WIDTH - tb] = conv_pre[layer][
+                    sblk, :, tb:
+                ]
+            else:
+                conv_ref[layer][dblk, : CONV_WIDTH - tb] = conv_pre[layer][sblk, tb:]
             ssm_ref[layer][dblk] = ssm_pre[layer][tblk]
     return conv_ref, ssm_ref
 
 
+@_parametrize("conv_state_dim_first", [False, True])
 @_parametrize("num_reqs", [1, 4, 16])
 @_parametrize("token_bias", [0, 1, 2])
-def test_precopy_matches_v1_copy_specs(num_reqs, token_bias):
+@_parametrize("has_idx_mapping", [True, False])
+@_cuda_required
+def test_precopy_matches_v1_copy_specs(
+    num_reqs, token_bias, has_idx_mapping, conv_state_dim_first
+):
     device = torch.device("cuda")
     torch.manual_seed(0)
     # Distinct physical block per (req, col) so copies never alias.
@@ -136,13 +169,20 @@ def test_precopy_matches_v1_copy_specs(num_reqs, token_bias):
     if num_reqs >= 2:
         dst_col[1] = 1  # src_col == dst_col -> no copy
 
-    convs, ssms = _build_state(num_blocks, device)
+    convs, ssms = _build_state(num_blocks, device, conv_state_dim_first)
     conv_ref, ssm_ref = _reference(
-        convs, ssms, bt.cpu(), src_col.cpu(), dst_col.cpu(), bias.cpu(), num_reqs
+        convs,
+        ssms,
+        bt.cpu(),
+        src_col.cpu(),
+        dst_col.cpu(),
+        bias.cpu(),
+        num_reqs,
+        conv_state_dim_first,
     )
 
     base, blk_stride, elem, inner, width, group, drc, drs = _build_meta(
-        convs, ssms, device
+        convs, ssms, device, conv_state_dim_first
     )
     bt_ptrs = torch.tensor([bt.data_ptr()], dtype=torch.int64, device=device)
     idx_mapping = torch.arange(num_reqs, dtype=torch.int32, device=device)
@@ -161,10 +201,11 @@ def test_precopy_matches_v1_copy_specs(num_reqs, token_bias):
         group,
         drc,
         drs,
-        idx_mapping,
+        idx_mapping if has_idx_mapping else None,
         num_reqs,
         COPY_BLOCK_SIZE=1024,
-        CONV_STATE_DIM_FIRST=False,
+        CONV_STATE_DIM_FIRST=conv_state_dim_first,
+        HAS_IDX_MAPPING=has_idx_mapping,
     )
     torch.accelerator.synchronize()
 
@@ -173,8 +214,204 @@ def test_precopy_matches_v1_copy_specs(num_reqs, token_bias):
         torch.testing.assert_close(ssms[layer], ssm_ref[layer], rtol=0, atol=0)
 
 
+def test_ds_conv_copy_spec_reproduces_multi_accept_assert(monkeypatch):
+    monkeypatch.setattr(
+        layer_mamba_utils,
+        "is_conv_state_dim_first",
+        lambda: True,
+    )
+    state = torch.empty((2, CONV_DIM, CONV_WIDTH), dtype=torch.bfloat16)
+
+    with pytest.raises(AssertionError, match="num_accepted_tokens > 1"):
+        layer_mamba_utils.get_conv_copy_spec(
+            state=state,
+            block_ids=[0, 1],
+            cur_block_idx=0,
+            num_accepted_tokens=3,
+        )
+
+
+class _FakeCpuGpuBuffer:
+    def __init__(self, n):
+        self.np = np.zeros(n, dtype=np.int32)
+        self.gpu = object()
+        self.copy_sizes = []
+
+    def copy_to_gpu(self, n=None):
+        self.copy_sizes.append(n)
+        return self.gpu
+
+
+class _FakePrecopyContext:
+    def __init__(self, n):
+        self.is_initialized = True
+        self.mamba_group_ids = [0]
+        self.mamba_state_idx_buf = _FakeCpuGpuBuffer(n)
+        self.precopy_src_col_buf = _FakeCpuGpuBuffer(n)
+        self.precopy_token_bias_buf = _FakeCpuGpuBuffer(n)
+        self.calls = []
+
+    def initialize_from_forward_context(self, *args, **kwargs):
+        raise AssertionError("test context is pre-initialized")
+
+    def run_fused_precopy(
+        self,
+        *,
+        num_reqs,
+        state_idx_gpu,
+        src_col_gpu,
+        token_bias_gpu,
+        idx_mapping,
+    ):
+        self.calls.append(
+            {
+                "num_reqs": num_reqs,
+                "state_idx": self.mamba_state_idx_buf.np[:num_reqs].copy(),
+                "src_col": self.precopy_src_col_buf.np[:num_reqs].copy(),
+                "token_bias": self.precopy_token_bias_buf.np[:num_reqs].copy(),
+                "idx_mapping": idx_mapping,
+            }
+        )
+
+
+def _make_preprocess_case(token_bias):
+    req_ids = ["fresh", "same", "cross_a", "cross_b"]
+    scheduler_output = SimpleNamespace(
+        finished_req_ids=set(),
+        preempted_req_ids=set(),
+        scheduled_cached_reqs=SimpleNamespace(resumed_req_ids=set()),
+        num_scheduled_tokens={
+            "fresh": 1,
+            "same": 1,
+            "cross_a": 1,
+            "cross_b": 2,
+        },
+    )
+    input_batch = SimpleNamespace(
+        req_ids=req_ids,
+        num_accepted_tokens_cpu=np.array(
+            [token_bias + 1, token_bias + 1, token_bias + 1, 2],
+            dtype=np.int32,
+        ),
+    )
+    requests = {
+        "fresh": SimpleNamespace(req_id="fresh", num_computed_tokens=0),
+        "same": SimpleNamespace(req_id="same", num_computed_tokens=5),
+        "cross_a": SimpleNamespace(req_id="cross_a", num_computed_tokens=8),
+        "cross_b": SimpleNamespace(req_id="cross_b", num_computed_tokens=7),
+    }
+    mamba_state_idx = {"same": 1, "cross_a": 0, "cross_b": 1}
+    return scheduler_output, input_batch, requests, mamba_state_idx
+
+
+@_parametrize("token_bias", [1, 2])
+def test_preprocess_fused_align_matches_scalar_bookkeeping(monkeypatch, token_bias):
+    block_size = 4
+    mamba_spec = SimpleNamespace(block_size=block_size, num_speculative_blocks=1)
+    cache_config = SimpleNamespace(enable_prefix_caching=True)
+    kv_cache_config = SimpleNamespace()
+    scalar_copy_calls = []
+
+    def fake_collect(
+        copy_bufs,
+        kv_cache_config,
+        mamba_state_copy_funcs,
+        mamba_group_ids,
+        src_block_idx,
+        dest_block_idx,
+        accept_token_bias,
+        req_state,
+        forward_context,
+    ):
+        scalar_copy_calls.append(
+            (
+                req_state.req_id,
+                src_block_idx,
+                dest_block_idx,
+                accept_token_bias,
+            )
+        )
+
+    monkeypatch.setattr(worker_mamba_utils, "collect_mamba_copy_meta", fake_collect)
+    monkeypatch.setattr(
+        worker_mamba_utils, "do_mamba_copy_block", lambda copy_bufs: None
+    )
+
+    scalar_case = _make_preprocess_case(token_bias)
+    fused_case = _make_preprocess_case(token_bias)
+
+    scalar_copy_bufs = SimpleNamespace(
+        mamba_group_ids=[0],
+        mamba_spec=mamba_spec,
+        offset=0,
+    )
+    worker_mamba_utils.preprocess_mamba(
+        scheduler_output=scalar_case[0],
+        kv_cache_config=kv_cache_config,
+        cache_config=cache_config,
+        mamba_state_idx=scalar_case[3],
+        input_batch=scalar_case[1],
+        requests=scalar_case[2],
+        forward_context={},
+        mamba_state_copy_funcs=(),
+        copy_bufs=scalar_copy_bufs,
+    )
+
+    ctx = _FakePrecopyContext(len(fused_case[1].req_ids))
+    fused_copy_bufs = SimpleNamespace(
+        mamba_group_ids=[0],
+        mamba_spec=mamba_spec,
+        offset=0,
+    )
+    worker_mamba_utils.preprocess_mamba(
+        scheduler_output=fused_case[0],
+        kv_cache_config=kv_cache_config,
+        cache_config=cache_config,
+        mamba_state_idx=fused_case[3],
+        input_batch=fused_case[1],
+        requests=fused_case[2],
+        forward_context={},
+        mamba_state_copy_funcs=(),
+        copy_bufs=fused_copy_bufs,
+        align_ctx=ctx,
+    )
+
+    assert fused_case[3] == scalar_case[3]
+    np.testing.assert_array_equal(
+        fused_case[1].num_accepted_tokens_cpu,
+        scalar_case[1].num_accepted_tokens_cpu,
+    )
+    assert scalar_copy_calls == [
+        ("cross_a", 0, 2, token_bias),
+        ("cross_b", 1, 2, 1),
+    ]
+    assert len(ctx.calls) == 1
+    call = ctx.calls[0]
+    assert call["num_reqs"] == len(fused_case[1].req_ids)
+    assert call["idx_mapping"] is None
+    np.testing.assert_array_equal(call["state_idx"], np.array([0, 1, 2, 2]))
+    np.testing.assert_array_equal(call["src_col"], np.array([-1, -1, 0, 1]))
+    np.testing.assert_array_equal(call["token_bias"], np.array([0, 0, token_bias, 1]))
+    fused_copy_calls = [
+        (req_id, int(src), int(dst), int(bias))
+        for req_id, src, dst, bias in zip(
+            fused_case[1].req_ids,
+            call["src_col"],
+            call["state_idx"],
+            call["token_bias"],
+        )
+        if int(src) != -1 and int(src) != int(dst)
+    ]
+    assert fused_copy_calls == scalar_copy_calls
+
+
 if __name__ == "__main__":
     for nr in (1, 4, 16):
         for tb in (0, 1, 2):
-            test_precopy_matches_v1_copy_specs(nr, tb)
-            print(f"OK num_reqs={nr} token_bias={tb}")
+            for mapping in (True, False):
+                for dim_first in (False, True):
+                    test_precopy_matches_v1_copy_specs(nr, tb, mapping, dim_first)
+                    print(
+                        f"OK num_reqs={nr} token_bias={tb} "
+                        f"has_idx_mapping={mapping} conv_dim_first={dim_first}"
+                    )

tests/kernels/moe/test_mxfp8_aiter_backend_selection.py

@@ -27,8 +27,8 @@ from vllm.model_executor.layers.fused_moe.experts.aiter_mxfp8_moe import (  # no
     _AITER_SWIGLU_BETA,
     AiterMxfp8Experts,
 )
-from vllm.model_executor.layers.fused_moe.experts.mxfp8_native_moe import (  # noqa: E402
-    Mxfp8NativeTritonExperts,
+from vllm.model_executor.layers.fused_moe.experts.mxfp8_emulation_moe import (  # noqa: E402
+    Mxfp8EmulationTritonExperts,
 )
 from vllm.model_executor.layers.fused_moe.modular_kernel import (  # noqa: E402
     FusedMoEActivationFormat,
@@ -150,16 +150,20 @@ def test_explicit_moe_backend_aiter():
 
 def test_gfx950_picks_aiter():
     """Auto-select on real ROCm hardware with flydsl usable -> FlyDSL wins."""
-    with _flydsl_installed(True):
+    # NOTE: Fp8MoeBackend.AITER_MXFP8 does not require VLLM_ROCM_USE_AITER=1
+    with (
+        patch(f"{_AITER_MOD}.current_platform.supports_mx", return_value=True),
+        _flydsl_installed(True),
+    ):
         backend, experts_cls = select_mxfp8_moe_backend(_config())
     assert backend is Fp8MoeBackend.AITER_MXFP8
     assert experts_cls is AiterMxfp8Experts
 
 
-def test_gfx942_picks_triton():
+def test_gfx942_picks_emulation():
     """flydsl unusable (e.g. gfx942, no FlyDSL support) -> native Triton
     dot_scaled backend wins instead."""
-    with _flydsl_installed(False):
+    with patch(f"{_AITER_MOD}.current_platform.supports_mx", return_value=False):
         backend, experts_cls = select_mxfp8_moe_backend(_config())
-    assert backend is Fp8MoeBackend.TRITON_MXFP8
-    assert experts_cls is Mxfp8NativeTritonExperts
+    assert backend is Fp8MoeBackend.EMULATION
+    assert experts_cls is Mxfp8EmulationTritonExperts

tests/models/language/pooling/test_all_pooling_plus_chunked_prefill.py

@@ -15,7 +15,9 @@ from vllm.config import PoolerConfig
     ["Qwen/Qwen3-Embedding-0.6B"],
 )
 @torch.inference_mode
-def test_embed_models(hf_runner, vllm_runner, model: str):
+def test_embed_models(hf_runner, vllm_runner, monkeypatch, model: str):
+    # Keep token_embed on MRV1 when sequence pooling becomes MRV2 by default.
+    monkeypatch.setenv("VLLM_USE_V2_MODEL_RUNNER", "0")
     chunk_size = 10
     n_prompt_tokens = [55, 56, 57]
     token_prompts = [[1024 + i for i in range(n)] for n in n_prompt_tokens]
@@ -104,3 +106,44 @@ def test_last_pool_score_chunked_prefill_matches_unchunked(vllm_runner, model: s
     assert chunked == pytest.approx(unchunked, abs=5e-2), (
         f"chunked score {chunked} diverged from unchunked {unchunked}"
     )
+
+
+@torch.inference_mode
+def test_sequence_embed_model_runner_v2(hf_runner, vllm_runner, monkeypatch) -> None:
+    model = "Qwen/Qwen3-Embedding-0.6B"
+    chunk_size = 10
+    token_prompts = [[1024 + i for i in range(n)] for n in (25, 27)]
+    prompts = [TokensPrompt(prompt_token_ids=t) for t in token_prompts]
+
+    with hf_runner(model, auto_cls=AutoModel) as hf_model:
+        hf_outputs = []
+        for token_prompt in token_prompts:
+            inputs = hf_model.wrap_device({"input_ids": torch.tensor([token_prompt])})
+            output = hf_model.model(inputs["input_ids"])
+            embedding = torch.nn.functional.normalize(
+                output.last_hidden_state.float()[0, -1], dim=0
+            )
+            hf_outputs.append(embedding.cpu().tolist())
+
+    monkeypatch.setenv("VLLM_USE_V2_MODEL_RUNNER", "1")
+    with vllm_runner(
+        model,
+        runner="pooling",
+        pooler_config=PoolerConfig(task="embed"),
+        max_model_len=64,
+        max_num_batched_tokens=chunk_size,
+        max_num_seqs=2,
+        gpu_memory_utilization=0.25,
+        enforce_eager=True,
+        enable_chunked_prefill=True,
+    ) as vllm_model:
+        assert vllm_model.llm.llm_engine.vllm_config.use_v2_model_runner
+        vllm_outputs = vllm_model.embed(prompts)
+
+    check_embeddings_close(
+        embeddings_0_lst=hf_outputs,
+        embeddings_1_lst=vllm_outputs,
+        name_0="hf",
+        name_1="vllm",
+        tol=1e-2,
+    )

tests/models/language/pooling/test_classification.py

@@ -49,3 +49,60 @@ def test_models(
             vllm_output,
             rtol=2e-3 if dtype == "float" else 1e-2,
         )
+
+
+@pytest.mark.core_model
+def test_bert_model_runner_v2(hf_runner, vllm_runner, monkeypatch) -> None:
+    model = "cross-encoder/ms-marco-TinyBERT-L-2-v2"
+    score_inputs = (
+        "What is the capital of France?",
+        [
+            "Paris.",
+            "Paris is the capital and largest city of France.",
+            "William Shakespeare wrote Hamlet in the early seventeenth century.",
+        ],
+    )
+    prompt_batches = [
+        ["short input"],
+        [
+            "short input",
+            "a longer input that exercises mixed sequence lengths",
+        ],
+    ]
+
+    with hf_runner(
+        model, dtype="half", auto_cls=AutoModelForSequenceClassification
+    ) as hf_model:
+        # HfRunner uses problem_type to preserve the model's
+        # sbert_ce_default_activation_function=Identity raw logits.
+        hf_model.config.problem_type = "regression"
+        hf_outputs = [hf_model.classify(prompts) for prompts in prompt_batches]
+
+    text_1, text_2 = score_inputs
+    text_pairs = [[text_1, document] for document in text_2]
+    with hf_runner(model, dtype="half", is_cross_encoder=True) as hf_model:
+        hf_scores = hf_model.predict(text_pairs).tolist()
+
+    monkeypatch.setenv("VLLM_USE_V2_MODEL_RUNNER", "1")
+    with vllm_runner(
+        model,
+        runner="pooling",
+        dtype="half",
+        max_model_len=64,
+    ) as vllm_model:
+        assert vllm_model.llm.llm_engine.vllm_config.use_v2_model_runner
+        vllm_outputs = [vllm_model.classify(prompts) for prompts in prompt_batches]
+        vllm_scores = vllm_model.score(*score_inputs)
+
+    for hf_batch, vllm_batch in zip(hf_outputs, vllm_outputs):
+        hf_tensor = torch.tensor(hf_batch)
+        vllm_tensor = torch.tensor(vllm_batch)
+        assert vllm_tensor.shape == hf_tensor.shape
+        assert torch.allclose(vllm_tensor, hf_tensor, rtol=1e-2, atol=1e-4)
+
+    assert torch.allclose(
+        torch.tensor(vllm_scores),
+        torch.tensor(hf_scores),
+        rtol=1e-2,
+        atol=1e-4,
+    )

tests/models/language/pooling/test_embedding.py

@@ -5,6 +5,7 @@ import pytest
 import torch
 from transformers import AutoModel
 
+from vllm import PoolingParams
 from vllm.config import PoolerConfig
 
 from ...utils import check_embeddings_close
@@ -132,6 +133,7 @@ def test_encoder_only_model_runner_v2_attention(
             task="embed", seq_pooling_type="LAST", use_activation=True
         ),
     ) as vllm_model:
+        assert vllm_model.llm.llm_engine.vllm_config.use_v2_model_runner
         vllm_outputs = vllm_model.embed(prompts)
 
     check_embeddings_close(
@@ -141,3 +143,77 @@ def test_encoder_only_model_runner_v2_attention(
         name_1="vllm",
         tol=1e-2,
     )
+
+
+@pytest.mark.core_model
+def test_encoder_model_runner_v2(hf_runner, vllm_runner, monkeypatch) -> None:
+    model = "sentence-transformers/all-MiniLM-L6-v2"
+    prompt_batches = [
+        ["short input"],
+        [
+            "short input",
+            "a longer input that exercises mixed sequence lengths",
+        ],
+    ]
+
+    with hf_runner(model, is_sentence_transformer=True) as hf_model:
+        hf_outputs = [hf_model.encode(prompts) for prompts in prompt_batches]
+
+    monkeypatch.setenv("VLLM_USE_V2_MODEL_RUNNER", "1")
+    with vllm_runner(
+        model,
+        runner="pooling",
+        max_model_len=64,
+    ) as vllm_model:
+        assert vllm_model.llm.llm_engine.vllm_config.use_v2_model_runner
+        vllm_outputs = [vllm_model.embed(prompts) for prompts in prompt_batches]
+
+    for hf_batch, vllm_batch in zip(hf_outputs, vllm_outputs):
+        check_embeddings_close(
+            embeddings_0_lst=hf_batch,
+            embeddings_1_lst=vllm_batch,
+            name_0="hf",
+            name_1="vllm",
+            tol=1e-2,
+        )
+
+
+@pytest.mark.core_model
+def test_matryoshka_dimensions_model_runner_v2(
+    hf_runner, vllm_runner, monkeypatch
+) -> None:
+    model = "Snowflake/snowflake-arctic-embed-m-v1.5"
+    prompts = ["short input", "a longer input for a different output width"]
+    dimensions = [None, 256]
+
+    with hf_runner(model, is_sentence_transformer=True) as hf_model:
+        hf_outputs = hf_model.encode(prompts)
+
+    monkeypatch.setenv("VLLM_USE_V2_MODEL_RUNNER", "1")
+    with vllm_runner(
+        model,
+        runner="pooling",
+        max_model_len=64,
+        gpu_memory_utilization=0.25,
+    ) as vllm_model:
+        assert vllm_model.llm.llm_engine.vllm_config.use_v2_model_runner
+        vllm_outputs = vllm_model.embed(
+            prompts,
+            pooling_params=[PoolingParams(dimensions=d) for d in dimensions],
+        )
+
+    expected_outputs = []
+    for output, dimension in zip(hf_outputs, dimensions):
+        output = torch.as_tensor(output)
+        if dimension is not None:
+            output = torch.nn.functional.normalize(output[:dimension], dim=0)
+        expected_outputs.append(output.tolist())
+
+    assert [len(output) for output in vllm_outputs] == [768, 256]
+    check_embeddings_close(
+        embeddings_0_lst=expected_outputs,
+        embeddings_1_lst=vllm_outputs,
+        name_0="hf",
+        name_1="vllm",
+        tol=1e-2,
+    )

tests/models/language/pooling/test_splade_sparse_pooler.py

@@ -1,6 +1,9 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+from unittest.mock import MagicMock, patch
+
+import numpy as np
 import pytest
 import torch
 import torch.nn as nn
@@ -9,8 +12,13 @@ from vllm.model_executor.models.bert import (
     BertMLMHead,
     SPLADESparsePooler,
 )
+from vllm.platforms import current_platform
 from vllm.pooling_params import PoolingParams
+from vllm.utils.torch_utils import PIN_MEMORY
 from vllm.v1.pool.metadata import PoolingMetadata, PoolingStates
+from vllm.v1.worker.gpu.input_batch import InputBatch
+from vllm.v1.worker.gpu.pool.pooling_runner import PoolingRunner
+from vllm.v1.worker.gpu.states import RequestState
 
 # ---------------------------------------------------------------------
 # Functional test: SPLADE formula correctness (no HF download needed)
@@ -91,3 +99,70 @@ def test_splade_pooler_matches_reference_formula(B, T, H, V):
         rtol=1e-4,
         atol=1e-4,
     )
+
+
+def test_pooling_runner_gathers_required_token_ids() -> None:
+    runner = PoolingRunner.__new__(PoolingRunner)
+    pooling_params = PoolingParams(task="embed", requires_token_ids=True)
+    runner.pooling_params = {1: pooling_params, 3: pooling_params}
+    runner.pooling_states = {1: PoolingStates(), 3: PoolingStates()}
+    runner.prompt_token_ids = {
+        1: torch.tensor([101, 102]),
+        3: torch.tensor([101, 11, 102]),
+    }
+
+    input_batch = MagicMock(spec=InputBatch)
+    input_batch.idx_mapping_np = np.array([3, 1], dtype=np.int32)
+    input_batch.num_reqs = 2
+    req_states = MagicMock(spec=RequestState)
+    req_states.prompt_len = MagicMock(np=np.array([0, 2, 0, 3], dtype=np.int32))
+    metadata = runner._get_pooling_metadata(
+        input_batch, req_states, torch.device(current_platform.device_type)
+    )
+
+    expected = torch.tensor([[101, 11, 102], [101, 102, 0]])
+    assert metadata.prompt_token_ids_cpu is not None
+    assert metadata.prompt_token_ids is not None
+    assert metadata.prompt_token_ids_cpu.is_pinned() == PIN_MEMORY
+    torch.testing.assert_close(
+        metadata.prompt_lens, torch.tensor([3, 2], dtype=torch.int32)
+    )
+    torch.testing.assert_close(metadata.prompt_token_ids_cpu, expected)
+    torch.testing.assert_close(metadata.prompt_token_ids.cpu(), expected)
+
+
+def test_pooling_runner_stores_only_required_token_ids() -> None:
+    runner = PoolingRunner.__new__(PoolingRunner)
+    runner.model = MagicMock()
+    runner.supported_tasks = frozenset({"embed"})
+    runner.pooling_params = {}
+    runner.pooling_states = {}
+    runner.prompt_token_ids = {}
+
+    runner.add_request(1, PoolingParams(task="embed"), [101, 102])
+    runner.add_request(
+        2,
+        PoolingParams(task="embed", requires_token_ids=True),
+        [101, 11, 102],
+    )
+
+    assert 1 not in runner.prompt_token_ids
+    torch.testing.assert_close(runner.prompt_token_ids[2], torch.tensor([101, 11, 102]))
+
+
+def test_pooling_runner_rejects_unsupported_selected_task() -> None:
+    model = MagicMock()
+    model.pooler.get_supported_tasks.return_value = {
+        "embed",
+        "embed&token_classify",
+        "token_classify",
+    }
+    vllm_config = MagicMock()
+    vllm_config.scheduler_config.max_num_seqs = 2
+    vllm_config.model_config.get_pooling_task.return_value = "embed&token_classify"
+
+    with (
+        patch.object(PoolingRunner, "get_supported_tasks", return_value=["embed"]),
+        pytest.raises(ValueError, match="selects 'embed&token_classify'"),
+    ):
+        PoolingRunner(model, vllm_config)

tests/models/language/pooling/test_truncation_control.py

@@ -2,6 +2,8 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 import pytest
 
+from vllm.exceptions import VLLMValidationError
+
 MODEL_NAME = "sentence-transformers/all-MiniLM-L12-v2"
 max_model_len = 128
 
@@ -60,7 +62,7 @@ def test_bigger_truncation_size(
     truncate_prompt_tokens = max_model_len + 1
 
     with (
-        pytest.raises(ValueError),
+        pytest.raises(VLLMValidationError),
         vllm_runner(
             model_name, runner="pooling", max_model_len=max_model_len
         ) as vllm_model,

tests/models/registry.py

@@ -497,6 +497,8 @@ _TEXT_GENERATION_EXAMPLE_MODELS = {
     "Qwen2MoeForCausalLM": _HfExamplesInfo("Qwen/Qwen1.5-MoE-A2.7B-Chat"),
     "Qwen3ForCausalLM": _HfExamplesInfo("Qwen/Qwen3-8B"),
     "Qwen3MoeForCausalLM": _HfExamplesInfo("Qwen/Qwen3-30B-A3B"),
+    "Qwen3_5ForCausalLM": _HfExamplesInfo("codecho/Qwen3.5-0.8B-text-only"),
+    "Qwen3_5MoeForCausalLM": _HfExamplesInfo("codecho/Qwen3.5-35B-A3B-text-only"),
     "MellumForCausalLM": _HfExamplesInfo("JetBrains/Mellum2-12B-A2.5B-Base"),
     "Qwen3NextForCausalLM": _HfExamplesInfo(
         "Qwen/Qwen3-Next-80B-A3B-Instruct",

tests/multimodal/test_hasher.py

@@ -1,6 +1,7 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 import uuid
+from io import BytesIO
 from pathlib import Path
 
 import numpy as np
@@ -10,6 +11,7 @@ from PIL import Image, ImageDraw
 
 from vllm.multimodal.hasher import MultiModalHasher
 from vllm.multimodal.media.base import MediaWithBytes
+from vllm.multimodal.media.image import ImageMediaIO
 from vllm.multimodal.parse import MultiModalDataParser
 
 pytestmark = pytest.mark.cpu_test
@@ -134,3 +136,37 @@ def test_hash_image_exif_id():
     assert hasher.hash_kwargs(image=image1) == hasher.hash_kwargs(image=id.bytes)
     # second image has non-UUID in ImageID, so it should hash to the image data
     assert hasher.hash_kwargs(image=image2) == hasher.hash_kwargs(image=image2a)
+
+
+def _rgba_png_bytes() -> bytes:
+    image = Image.new("RGBA", (8, 8), (255, 0, 0, 128))
+    buf = BytesIO()
+    image.save(buf, format="PNG")
+    return buf.getvalue()
+
+
+def test_hash_collision_media_io_config():
+    data = _rgba_png_bytes()
+    white = ImageMediaIO(rgba_background_color=(255, 255, 255)).load_bytes(data)
+    black = ImageMediaIO(rgba_background_color=(0, 0, 0)).load_bytes(data)
+    white2 = ImageMediaIO(rgba_background_color=(255, 255, 255)).load_bytes(data)
+    keep = ImageMediaIO(image_mode=None).load_bytes(data)
+
+    hasher = MultiModalHasher
+    assert hasher.hash_kwargs(image=white) != hasher.hash_kwargs(image=black)
+    assert hasher.hash_kwargs(image=white) != hasher.hash_kwargs(image=keep)
+    assert hasher.hash_kwargs(image=white) == hasher.hash_kwargs(image=white2)
+
+
+def test_hash_media_io_noop_config_preserves_hash():
+    image = Image.new("RGB", (8, 8), (0, 128, 255))
+    buf = BytesIO()
+    image.save(buf, format="PNG")
+    data = buf.getvalue()
+
+    loaded = ImageMediaIO().load_bytes(data)
+    assert loaded.io_config is None
+
+    plain = MediaWithBytes(loaded.media, data)
+    hasher = MultiModalHasher
+    assert hasher.hash_kwargs(image=loaded) == hasher.hash_kwargs(image=plain)

tests/multimodal/test_video.py

@@ -14,6 +14,7 @@ from transformers import AutoVideoProcessor
 from transformers.video_utils import VideoMetadata
 
 from vllm.assets.base import get_vllm_public_assets
+from vllm.models.minimax_m3.common.mm_preprocess import MiniMaxM3VideoBackend
 from vllm.multimodal.video import (
     PYNVVIDEOCODEC_DECODER_CACHE_SIZE,
     PYNVVIDEOCODEC_VIDEO_BACKEND,
@@ -388,6 +389,12 @@ def test_cosmos3_edge_uses_qwen3_vl_video_backend():
             {"fps": 2},
             id="qwen2_5_vl",
         ),
+        pytest.param(
+            "MiniMaxAI/MiniMax-M3",
+            MiniMaxM3VideoBackend,
+            None,
+            id="minimax_m3_vl",
+        ),
     ],
 )
 def test_video_processor_from_model_repo(

tests/v1/core/test_async_scheduler.py

@@ -1,6 +1,7 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-from collections import deque
+from collections import defaultdict, deque
+from collections.abc import Callable
 from unittest.mock import Mock
 
 import pytest
@@ -329,46 +330,311 @@ def test_abort_request_when_structured_output_fsm_cannot_advance():
     assert not scheduler.running
 
 
-def test_no_placeholder_underflow_on_discarded_spec_frame():
-    num_spec = 5
+class PipelinedEngine:
+    """Drive a real AsyncScheduler like EngineCore.step_with_batch_queue:
+    schedule until the batch queue is full, then process the oldest step's
+    output. Async PP runs pp_size+1 concurrent batches, so up to pp_size
+    steps are in flight at each schedule() call -- the window in which
+    preemption must handle output that has not yet returned. (Single-GPU e2e
+    tests can never create this window: at PP=1, exactly one step is in
+    flight and it is processed before a preempted request can resume.)
+
+    The model runner is emulated with the V2 runner's own bookkeeping, from
+    only what the scheduler serializes to it: slots flushed on
+    preempted_req_ids, resumed requests re-added from the NewRequestData
+    snapshot, sampling when a step reaches the end of the runner's own view
+    of the sequence. This makes preemption races observable: a stale token
+    delivered after a resume is scheduled extends the scheduler's sequence
+    but not the runner's.
+
+    Every sample emits a globally unique token tagged with its sampled
+    position, so tests can assert exact delivery.
+    """
+
+    def __init__(
+        self,
+        scheduler: AsyncScheduler,
+        queue_size: int,
+        accept_drafts: Callable[[int, str, int], int] | None = None,
+    ):
+        self.scheduler = scheduler
+        self.queue_size = queue_size
+        self.accept_drafts = accept_drafts
+        # In-flight steps: (scheduler_output, new_reqs snapshot) in FIFO order.
+        self.queue: deque[tuple[SchedulerOutput, list[tuple[str, int, int]]]] = deque()
+        # Runner-side request state: req_id -> [seq_len, num_computed] as the
+        # runner sees them (its own sampled tokens, not the scheduler's).
+        self.runner_view: dict[str, list[int]] = {}
+        # All tokens the fake runner ever sampled, per request, in order.
+        self.emitted: dict[str, list[int]] = defaultdict(list)
+        # Sequence position each (globally unique) token was sampled for.
+        self.emitted_position: dict[int, int] = {}
+        self.step_idx = 0
+        self._next_token = 1000
+
+    def _schedule(self) -> bool:
+        scheduler_output = self.scheduler.schedule()
+        self.step_idx += 1
+        # Snapshot what NewRequestData serializes at schedule time (both new
+        # and resumed requests for the V2 runner).
+        new_reqs = [
+            (r.req_id, len(r.prefill_token_ids), r.num_computed_tokens)
+            for r in scheduler_output.scheduled_new_reqs
+        ]
+        # Enqueue empty steps too (the engine executes them), so the runner
+        # still observes their preempted/finished request ids in step order.
+        self.queue.appendleft((scheduler_output, new_reqs))
+        return True
+
+    def _process_oldest_step(self) -> None:
+        scheduler_output, new_reqs = self.queue.pop()
+        # Worker-side state updates, in step order: flush preempted/finished
+        # slots, then (re-)add new/resumed requests.
+        for req_id in scheduler_output.preempted_req_ids or ():
+            self.runner_view.pop(req_id, None)
+        for req_id in scheduler_output.finished_req_ids or ():
+            self.runner_view.pop(req_id, None)
+        for req_id, seq_len, num_computed in new_reqs:
+            self.runner_view[req_id] = [seq_len, num_computed]
+
+        req_ids = list(scheduler_output.num_scheduled_tokens.keys())
+        sampled_token_ids: list[list[int]] = []
+        for req_id in req_ids:
+            num_scheduled = scheduler_output.num_scheduled_tokens[req_id]
+            view = self.runner_view.get(req_id)
+            if view is None:
+                # Slot already flushed (request finished/aborted mid-flight).
+                sampled_token_ids.append([])
+                continue
+            seq_len, num_computed = view
+            end = num_computed + num_scheduled
+            if end < seq_len:
+                # Partial prefill by the runner's own bookkeeping: no sample.
+                view[1] = end
+                sampled_token_ids.append([])
+                continue
+            drafts = scheduler_output.scheduled_spec_decode_tokens.get(req_id, ())
+            num_accepted = (
+                min(self.accept_drafts(self.step_idx, req_id, len(drafts)), len(drafts))
+                if drafts and self.accept_drafts
+                else 0
+            )
+            num_rejected = len(drafts) - num_accepted
+            tokens = list(range(self._next_token, self._next_token + 1 + num_accepted))
+            self._next_token += 1 + num_accepted
+            self.emitted[req_id].extend(tokens)
+            sampled_token_ids.append(tokens)
+            # Rejected drafts roll back computed; the sampled tokens extend
+            # the runner's sequence.
+            view[1] = end - num_rejected
+            view[0] = view[1] + 1
+            for offset, token in enumerate(tokens):
+                self.emitted_position[token] = view[0] - len(tokens) + offset
+        model_runner_output = ModelRunnerOutput(
+            req_ids=req_ids,
+            req_id_to_index={req_id: i for i, req_id in enumerate(req_ids)},
+            sampled_token_ids=sampled_token_ids,
+            logprobs=None,
+            prompt_logprobs_dict={},
+            pooler_output=[],
+        )
+        self.scheduler.update_from_output(scheduler_output, model_runner_output)
+
+    def run(
+        self,
+        max_steps: int = 2000,
+        before_step: Callable[[int, "PipelinedEngine"], None] | None = None,
+    ) -> None:
+        for i in range(max_steps):
+            if not self.scheduler.has_requests() and not self.queue:
+                return
+            if before_step is not None:
+                before_step(i, self)
+            scheduled = (
+                self.scheduler.has_requests()
+                and len(self.queue) < self.queue_size
+                and self._schedule()
+            )
+            if scheduled and len(self.queue) < self.queue_size:
+                # Queue not yet full: the engine returns without blocking.
+                continue
+            if self.queue:
+                self._process_oldest_step()
+        raise AssertionError("engine loop did not converge")
+
+
+def _create_async_pp_scheduler(
+    num_spec: int, pp_size: int = 3, num_blocks: int = 5
+) -> AsyncScheduler:
     scheduler = create_scheduler(
         async_scheduling=True,
-        num_speculative_tokens=num_spec,
-        speculative_method="ngram_gpu",
+        num_speculative_tokens=num_spec or None,
+        speculative_method="ngram_gpu" if num_spec else None,
+        use_v2_model_runner=True,
+        num_blocks=num_blocks,
+        block_size=16,
+        max_num_batched_tokens=512,
     )
-    req = create_requests(num_requests=1, max_tokens=20)[0]
-    req.num_computed_tokens = req.num_tokens
-    scheduler.requests[req.request_id] = req
-    scheduler.running.append(req)
-    req.status = RequestStatus.RUNNING
+    # Emulate PP at the scheduler level; constructing with
+    # pipeline_parallel_size>1 requires that many visible GPUs. Drive with
+    # queue_size=pp_size+1 (V2 async PP runs pp_size+1 concurrent batches).
+    scheduler.pp_size = pp_size
+    scheduler.use_pp = pp_size > 1
+    return scheduler
 
-    req.num_output_placeholders = 1
-    req.async_tokens_to_discard = num_spec
-    computed_before = req.num_computed_tokens
 
-    scheduler_output = SchedulerOutput(
-        scheduled_new_reqs=[],
-        scheduled_cached_reqs=CachedRequestData.make_empty(),
-        num_scheduled_tokens={req.request_id: num_spec + 1},
-        total_num_scheduled_tokens=num_spec + 1,
-        scheduled_encoder_inputs={},
-        scheduled_spec_decode_tokens={req.request_id: [10] * num_spec},
-        num_common_prefix_blocks=[],
-        finished_req_ids=set(),
-        free_encoder_mm_hashes=[],
+def _assert_ordered_subset(delivered: list[int], emitted: list[int]) -> None:
+    """Delivered tokens must be an order-preserving subset of the emitted
+    tokens with no duplicates (tokens are globally unique)."""
+    it = iter(emitted)
+    for token in delivered:
+        assert token in it, f"token {token} delivered out of order or twice"
+
+
+def _assert_positions_consistent(req, engine: PipelinedEngine) -> None:
+    """The i-th delivered output token must be one the runner sampled for
+    exactly sequence position prompt_len + i: catches a preempted request's
+    stale output landing on a position the resumed request resampled (or
+    vice versa), which token-stream equality alone cannot see."""
+    for i, token in enumerate(req.output_token_ids):
+        expected = req.num_prompt_tokens + i
+        actual = engine.emitted_position[token]
+        assert actual == expected, (
+            f"output {i} of {req.request_id}: token sampled for position "
+            f"{actual}, delivered as position {expected}"
+        )
+
+
+@pytest.mark.parametrize("num_spec", [0, 3])
+def test_kv_pressure_preemption_with_inflight_output(num_spec: int):
+    """KV-pressure preemption of requests with in-flight async output.
+
+    PP=3 + async scheduling (batch queue of 4), a block pool small enough
+    that decodes contend and preempt mid-flight, and staggered arrivals so
+    the batch queue actually pipelines. A preempted request's in-flight steps
+    still return: their tokens must be delivered exactly once, their stale
+    spec-rejection counts must not corrupt the rolled-back counters, and the
+    resume must not resample a position that output later delivers.
+
+    Regression for the num_output_placeholders underflow EngineCore crash:
+    with the fix reverted, the num_spec=3 variant fails with exactly
+    ``assert request.num_output_placeholders >= 0`` when a stale spec output
+    returns after the preempted request was resumed and sampled.
+    """
+    max_tokens = 24
+    scheduler = _create_async_pp_scheduler(num_spec)
+    requests = create_requests(
+        num_requests=8, num_tokens=8, max_tokens=max_tokens, ignore_eos=True
     )
-    model_runner_output = ModelRunnerOutput(
-        req_ids=[req.request_id],
-        req_id_to_index={req.request_id: 0},
-        sampled_token_ids=[[999]],
-        logprobs=None,
-        prompt_logprobs_dict={},
-        pooler_output=[],
+    pending = list(requests)
+    for _ in range(2):
+        scheduler.add_request(pending.pop(0))
+
+    # Observe that the scenario under test actually occurs.
+    preempts_with_inflight_output = 0
+    orig_preempt = scheduler._preempt_request
+
+    def counting_preempt(request, timestamp, **kwargs):
+        nonlocal preempts_with_inflight_output
+        if request.num_in_flight_tokens > 0:
+            preempts_with_inflight_output += 1
+        return orig_preempt(request, timestamp, **kwargs)
+
+    scheduler._preempt_request = counting_preempt
+
+    def add_requests(step: int, engine: PipelinedEngine):
+        if pending:
+            scheduler.add_request(pending.pop(0))
+
+    engine = PipelinedEngine(
+        scheduler,
+        queue_size=4,
+        # Deterministically vary spec acceptance so stale outputs carry
+        # nonzero rejection counts.
+        accept_drafts=lambda step, req_id, n: (step + int(req_id)) % (n + 1),
     )
+    engine.run(before_step=add_requests)
 
-    scheduler.update_from_output(scheduler_output, model_runner_output)
+    assert preempts_with_inflight_output > 0, (
+        "test did not exercise preemption with in-flight output"
+    )
+    for req in requests:
+        assert req.is_finished()
+        assert req.num_output_tokens == max_tokens
+        # Lossless: delivered tokens are exactly the sampled tokens, in order
+        # (the excluded tail was emitted after the request finished).
+        emitted = engine.emitted[req.request_id]
+        assert list(req.output_token_ids) == emitted[:max_tokens]
+        _assert_positions_consistent(req, engine)
 
-    assert req.num_output_placeholders == 1
-    assert req.num_computed_tokens == computed_before
-    assert req.async_tokens_to_discard == num_spec - 1
-    assert req.status == RequestStatus.RUNNING
+
+@pytest.mark.parametrize("pp_size", [1, 3])
+def test_reset_prefix_cache_with_inflight_output_under_kv_pressure(pp_size: int):
+    """reset_prefix_cache(reset_running_requests=True) resumes requests in
+    the same step it preempts them, so in-flight output must be dropped (the
+    resume resamples those positions).
+
+    pp_size=1: regression for the frame-based discard this fix replaces,
+    which with spec decode drained one *token* count per output frame and
+    over-discarded, corrupting the fresh frames after the resume.
+    pp_size=3: back-to-back resets, so the second re-preempts requests whose
+    dropped stale share is still in flight -- it must be recorded once (not
+    accumulated) and stay dropped.
+    """
+    max_tokens = 24
+    scheduler = _create_async_pp_scheduler(num_spec=3, pp_size=pp_size)
+    requests = create_requests(
+        num_requests=8, num_tokens=8, max_tokens=max_tokens, ignore_eos=True
+    )
+    pending = list(requests)
+    for _ in range(2):
+        scheduler.add_request(pending.pop(0))
+
+    # Observe re-preemptions with an undrained stale share (the
+    # double-count hazard).
+    repreempts_with_stale = 0
+    orig_preempt = scheduler._preempt_request
+
+    def counting_preempt(request, timestamp, **kwargs):
+        nonlocal repreempts_with_stale
+        if getattr(request, "num_stale_output_tokens", 0) > 0:
+            repreempts_with_stale += 1
+        return orig_preempt(request, timestamp, **kwargs)
+
+    scheduler._preempt_request = counting_preempt
+
+    resets = 0
+    reset_steps = {6, 14} if pp_size == 1 else {6, 7, 18, 19}
+
+    def before_step(step: int, engine: PipelinedEngine):
+        nonlocal resets
+        if pending:
+            scheduler.add_request(pending.pop(0))
+        if step in reset_steps and (engine.queue or scheduler.running):
+            scheduler.reset_prefix_cache(reset_running_requests=True)
+            resets += 1
+
+    engine = PipelinedEngine(
+        scheduler,
+        queue_size=pp_size + 1,
+        accept_drafts=lambda step, req_id, n: (step + int(req_id)) % (n + 1),
+    )
+    engine.run(before_step=before_step)
+
+    assert resets > 0, "test did not exercise reset_prefix_cache"
+    if pp_size > 1:
+        # The re-preempt-while-stale-pending window needs pipeline depth.
+        assert repreempts_with_stale > 0, (
+            "test did not exercise re-preemption with an undrained stale share"
+        )
+    for req in requests:
+        assert req.is_finished()
+        assert req.num_output_tokens == max_tokens
+        # Dropped tokens are never delivered; order must be preserved with
+        # no duplicates.
+        _assert_ordered_subset(
+            list(req.output_token_ids), engine.emitted[req.request_id]
+        )
+        _assert_positions_consistent(req, engine)
+        # All stale shares fully drained by the end.
+        assert getattr(req, "num_stale_output_tokens", 0) == 0

tests/v1/core/test_scheduler.py

@@ -1107,7 +1107,7 @@ def test_preemption_re_records_prefix_cache_query():
     request = create_requests(num_requests=1)[0]
     scheduler.add_request(request)
 
-    scheduler.schedule()
+    scheduler_output = scheduler.schedule()
     stats = scheduler.kv_cache_manager.prefix_cache_stats
     assert stats is not None
     assert (stats.requests, stats.preempted_requests) == (1, 0)
@@ -1116,6 +1116,21 @@ def test_preemption_re_records_prefix_cache_query():
     scheduler._preempt_request(request, 0.0)
     assert request.status == RequestStatus.PREEMPTED
 
+    scheduler.update_from_output(
+        scheduler_output,
+        ModelRunnerOutput(
+            req_ids=[request.request_id],
+            req_id_to_index={request.request_id: 0},
+            sampled_token_ids=[[1000]],
+            logprobs=None,
+            prompt_logprobs_dict={},
+            pooler_output=[],
+        ),
+    )
+    assert request.num_stale_output_tokens == 0
+    stats = scheduler.kv_cache_manager.prefix_cache_stats
+    assert stats is not None
+
     scheduler.schedule()
     assert request.status == RequestStatus.RUNNING
     assert stats.preempted_requests == 1

tests/v1/e2e/general/test_mamba_prefix_cache.py

@@ -338,6 +338,30 @@ def get_fake_process_mamba_fn(
             assert copy_info[0][-1] == expected_temporal_src
             assert copy_info[1][-1] == expected_temporal_dest
 
+    def check_fused_copy_info(
+        action: tuple[int, int],
+        align_ctx: mamba_utils.MambaSpecDecodeGPUContext,
+    ):
+        # Align + spec-decode on a hybrid model routes the pre-copy through the
+        # fused kernel (preprocess_mamba -> run_fused_precopy) instead of
+        # do_mamba_copy_block, so copy_info is never populated. Verify from the
+        # fused buffers (req-0 scope, mirroring check_copy_info):
+        #   - the copy DECISION: src_col is -1 iff no pre-copy is scheduled;
+        #   - the DESTINATION column: state_idx == action[1] (curr_state_idx;
+        #     maps directly through block_ids, exactly as check_copy_info's dst).
+        # The source column is NOT asserted here: on the scalar path the source
+        # address is produced by the per-state copy func with an accept-token
+        # bias offset (collect_mamba_copy_meta), so prev_state_idx does not map
+        # to action[0] by plain equality. Source block-level exactness (incl.
+        # the accept-bias) is covered by test_precopy_mamba_align.py.
+        src_col = int(align_ctx.precopy_src_col_buf.np[0])
+        state_idx = int(align_ctx.mamba_state_idx_buf.np[0])
+        if action == (-1, -1):
+            assert src_col == -1
+        else:
+            assert src_col != -1
+            assert state_idx == action[1]
+
     def fake_preprocess_mamba_fn(
         scheduler_output: SchedulerOutput,
         kv_cache_config: KVCacheConfig,
@@ -348,6 +372,7 @@ def get_fake_process_mamba_fn(
         forward_context: dict[str, Any],
         mamba_state_copy_funcs: tuple[MambaStateCopyFunc, ...],
         copy_bufs: mamba_utils.MambaCopyBuffers,
+        align_ctx: mamba_utils.MambaSpecDecodeGPUContext | None = None,
     ):
         nonlocal copy_info
         copy_info = None
@@ -361,14 +386,21 @@ def get_fake_process_mamba_fn(
             forward_context,
             mamba_state_copy_funcs,
             copy_bufs,
+            align_ctx,
         )
         if cur_step_action is not None:
-            check_copy_info(
-                cur_step_action.preprocess_copy_idx,
-                kv_cache_config,
-                forward_context,
-                input_batch,
-            )
+            if align_ctx is not None:
+                check_fused_copy_info(
+                    cur_step_action.preprocess_copy_idx,
+                    align_ctx,
+                )
+            else:
+                check_copy_info(
+                    cur_step_action.preprocess_copy_idx,
+                    kv_cache_config,
+                    forward_context,
+                    input_batch,
+                )
         return ret
 
     def fake_copy_fn(copy_bufs: mamba_utils.MambaCopyBuffers):

tests/v1/kv_offload/test_factory.py

@@ -136,8 +136,12 @@ def test_create_cpu_offloading_spec():
 
 
 def test_cpu_spec_sizes_normalized_worker_layout():
+    # The CPU spec now rounds the offloaded row up to the mmap page size
+    # (matching the shared region), so kv_bytes_per_chunk picks up padding
+    # while cpu_page_size_per_worker stays the un-padded per-worker slot.
+    alignment = SharedOffloadRegion.BLOCK_SIZE_ALIGNMENT
     spec = _create_spec(
-        cpu_bytes_to_use=1920,
+        cpu_bytes_to_use=alignment * 3,
         worker_kv_bytes_per_block=16,
         blocks_per_chunk=2,
         world_size=6,
@@ -147,8 +151,8 @@ def test_cpu_spec_sizes_normalized_worker_layout():
 
     assert isinstance(spec, CPUOffloadingSpec)
     assert spec.cpu_page_size_per_worker == 32
-    assert spec.kv_bytes_per_chunk == 192
-    assert spec.num_blocks == 10
+    assert spec.kv_bytes_per_chunk == alignment
+    assert spec.num_blocks == 3
 
 
 def test_cpu_spec_zero_worker_bytes_produces_empty_cache():
@@ -267,7 +271,9 @@ def test_tiering_spec_create_worker_folds_device_index_for_sharded_layout(monkey
 
 @pytest.mark.parametrize("world_size", [2, 4, 8])
 def test_cpu_spec_replicated_config_preserves_per_rank_sizing(world_size: int):
-    worker_kv_bytes_per_block = 4096
+    # Use a page-multiple per-worker block so the row size is unaffected by the
+    # page-alignment rounding regardless of the host mmap page size.
+    worker_kv_bytes_per_block = SharedOffloadRegion.BLOCK_SIZE_ALIGNMENT
     spec = _create_spec(
         cpu_bytes_to_use=worker_kv_bytes_per_block * world_size * 2,
         worker_kv_bytes_per_block=worker_kv_bytes_per_block,
@@ -282,6 +288,107 @@ def test_cpu_spec_replicated_config_preserves_per_rank_sizing(world_size: int):
     assert spec.num_blocks == 2
 
 
+def test_cpu_spec_create_worker_uses_mmap_on_cuda_alike(monkeypatch):
+    import vllm.v1.kv_offload.cpu.spec as cpu_spec_module
+
+    worker_kv_bytes_per_block = SharedOffloadRegion.BLOCK_SIZE_ALIGNMENT
+    spec = _create_spec(
+        cpu_bytes_to_use=worker_kv_bytes_per_block * 8,
+        worker_kv_bytes_per_block=worker_kv_bytes_per_block,
+        world_size=4,
+    )
+    assert isinstance(spec, CPUOffloadingSpec)
+
+    region = MagicMock()
+    region_calls: list[dict[str, Any]] = []
+    worker_calls: list[dict[str, Any]] = []
+
+    def fake_region_ctor(**kwargs):
+        region_calls.append(kwargs)
+        return region
+
+    def fake_worker_ctor(**kwargs):
+        worker_calls.append(kwargs)
+        return MagicMock()
+
+    monkeypatch.setattr(cpu_spec_module.current_platform, "is_cuda_alike", lambda: True)
+    monkeypatch.setattr(cpu_spec_module, "SharedOffloadRegion", fake_region_ctor)
+    monkeypatch.setattr(cpu_spec_module, "CPUOffloadingWorker", fake_worker_ctor)
+    monkeypatch.setattr(
+        cpu_spec_module.torch.accelerator, "current_device_index", lambda: 5
+    )
+
+    kv_caches = MagicMock()
+    spec.create_worker(kv_caches)
+
+    # rank folds the physical device index into [0, world_size): 5 % 4 == 1.
+    assert region_calls[0]["rank"] == 1
+    assert region_calls[0]["engine_id"] == "test-engine"
+    assert region_calls[0]["kv_bytes_per_block"] == worker_kv_bytes_per_block * 4
+    assert worker_calls[0]["kv_caches"] is kv_caches
+    assert worker_calls[0]["mmap_region"] is region
+
+
+def test_cpu_spec_create_worker_uses_tensor_path_off_cuda_alike(monkeypatch):
+    import vllm.v1.kv_offload.cpu.spec as cpu_spec_module
+
+    spec = _create_spec(worker_kv_bytes_per_block=4096, world_size=4)
+    assert isinstance(spec, CPUOffloadingSpec)
+
+    region_calls: list[dict[str, Any]] = []
+    worker_calls: list[dict[str, Any]] = []
+
+    def fake_region_ctor(**kwargs):
+        region_calls.append(kwargs)
+        return MagicMock()
+
+    def fake_worker_ctor(**kwargs):
+        worker_calls.append(kwargs)
+        return MagicMock()
+
+    monkeypatch.setattr(
+        cpu_spec_module.current_platform, "is_cuda_alike", lambda: False
+    )
+    monkeypatch.setattr(cpu_spec_module, "SharedOffloadRegion", fake_region_ctor)
+    monkeypatch.setattr(cpu_spec_module, "CPUOffloadingWorker", fake_worker_ctor)
+
+    spec.create_worker(MagicMock())
+
+    # Non-CUDA-alike platforms keep the per-rank pinned-tensor path.
+    assert region_calls == []
+    assert worker_calls[0]["mmap_region"] is None
+
+
+def test_cpu_spec_create_worker_skips_mmap_for_empty_cache(monkeypatch):
+    import vllm.v1.kv_offload.cpu.spec as cpu_spec_module
+
+    # worker_kv_bytes_per_block=0 yields num_blocks=0; a zero-byte region cannot
+    # be mmap'd, so even on CUDA-alike this must fall back to the tensor path.
+    spec = _create_spec(worker_kv_bytes_per_block=0, world_size=4)
+    assert isinstance(spec, CPUOffloadingSpec)
+    assert spec.num_blocks == 0
+
+    region_calls: list[dict[str, Any]] = []
+    worker_calls: list[dict[str, Any]] = []
+
+    monkeypatch.setattr(cpu_spec_module.current_platform, "is_cuda_alike", lambda: True)
+    monkeypatch.setattr(
+        cpu_spec_module,
+        "SharedOffloadRegion",
+        lambda **kwargs: region_calls.append(kwargs),
+    )
+    monkeypatch.setattr(
+        cpu_spec_module,
+        "CPUOffloadingWorker",
+        lambda **kwargs: worker_calls.append(kwargs),
+    )
+
+    spec.create_worker(MagicMock())
+
+    assert region_calls == []
+    assert worker_calls[0]["mmap_region"] is None
+
+
 def test_offloading_spec_has_replicated_layout_default():
     spec = SingleArgExternalOffloadingSpec(_make_offloading_config())
     assert spec.replicated_layout is False

tests/v1/spec_decode/test_llm_base_proposer.py

@@ -0,0 +1,112 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Tests for SpecDecodeBaseProposer.initialize_attn_backend.
+
+Block tables are stored at kernel-block granularity, so the proposer's
+``block_size`` (used for slot-mapping math) must be the kernel block size,
+not the KV cache manager's block size — the two differ when manager blocks
+are split for the attention kernel. The value must also be deterministic:
+``_draft_attn_layer_names`` is a set, whose iteration order varies across
+processes, so anything derived from iteration order must not leak into
+``block_size``.
+"""
+
+from types import SimpleNamespace
+
+import pytest
+
+import vllm.v1.spec_decode.llm_base_proposer as llm_base_proposer
+from vllm.v1.spec_decode.eagle import EagleProposer
+
+SCHEDULER_BLOCK_SIZE = 256
+KERNEL_BLOCK_SIZE = 64
+
+
+class _FakeAttentionGroup:
+    def __init__(self, backend, layer_names, kv_cache_spec, kv_cache_group_id):
+        self.backend = backend
+        self.layer_names = list(layer_names)
+        self.kv_cache_spec = kv_cache_spec
+        self.kv_cache_group_id = kv_cache_group_id
+        self.kernel_block_size = None
+
+    def create_metadata_builders(self, vllm_config, device, kernel_block_size=None):
+        self.kernel_block_size = kernel_block_size
+
+    def get_metadata_builder(self):
+        return SimpleNamespace(kv_cache_spec=self.kv_cache_spec)
+
+
+def _make_proposer(
+    monkeypatch: pytest.MonkeyPatch, layer_names: set[str]
+) -> EagleProposer:
+    fake_layers = {}
+    for name in layer_names:
+        backend = SimpleNamespace(full_cls_name=lambda: "FakeBackend")
+        fake_layers[name] = SimpleNamespace(
+            get_attn_backend=lambda backend=backend: backend
+        )
+    monkeypatch.setattr(
+        llm_base_proposer, "get_layers_from_vllm_config", lambda *a, **k: fake_layers
+    )
+    monkeypatch.setattr(llm_base_proposer, "AttentionGroup", _FakeAttentionGroup)
+
+    proposer = EagleProposer.__new__(EagleProposer)
+    proposer.vllm_config = None
+    proposer.device = None
+    proposer._draft_attn_layer_names = set(layer_names)
+    proposer.kv_cache_gid = -1
+    proposer.draft_attn_groups = []
+    proposer.block_size = -1
+    return proposer
+
+
+def _make_kv_cache_config(layer_names: set[str]) -> SimpleNamespace:
+    spec = SimpleNamespace(block_size=SCHEDULER_BLOCK_SIZE)
+    group = SimpleNamespace(layer_names=list(layer_names), kv_cache_spec=spec)
+    return SimpleNamespace(kv_cache_groups=[group])
+
+
+def test_block_size_uses_kernel_block_size(monkeypatch: pytest.MonkeyPatch):
+    """The proposer's slot-mapping math runs against the kernel-granularity
+    block table, so block_size must come from kernel_block_sizes."""
+    layer_names = {"draft.0.self_attn.attn"}
+    proposer = _make_proposer(monkeypatch, layer_names)
+
+    proposer.initialize_attn_backend(
+        _make_kv_cache_config(layer_names),
+        kernel_block_sizes=[KERNEL_BLOCK_SIZE],
+    )
+
+    assert proposer.block_size == KERNEL_BLOCK_SIZE
+    assert proposer.block_size != SCHEDULER_BLOCK_SIZE
+    # The metadata builder keeps receiving the kernel block size as well.
+    assert proposer.draft_attn_groups[0].kernel_block_size == KERNEL_BLOCK_SIZE
+
+
+def test_block_size_falls_back_to_kv_cache_spec(monkeypatch: pytest.MonkeyPatch):
+    layer_names = {"draft.0.self_attn.attn"}
+    proposer = _make_proposer(monkeypatch, layer_names)
+
+    proposer.initialize_attn_backend(
+        _make_kv_cache_config(layer_names), kernel_block_sizes=None
+    )
+
+    assert proposer.block_size == SCHEDULER_BLOCK_SIZE
+
+
+def test_draft_layer_iteration_is_deterministic(monkeypatch: pytest.MonkeyPatch):
+    """_draft_attn_layer_names is a set; the attention groups built from it
+    must not depend on its (process-random) iteration order."""
+    layer_names = {"draft.c.attn", "draft.a.attn", "draft.b.attn"}
+    expected_order = sorted(layer_names)
+
+    for insertion_order in (expected_order, expected_order[::-1]):
+        proposer = _make_proposer(monkeypatch, set(insertion_order))
+        proposer.initialize_attn_backend(
+            _make_kv_cache_config(set(insertion_order)),
+            kernel_block_sizes=[KERNEL_BLOCK_SIZE],
+        )
+        assert len(proposer.draft_attn_groups) == 1
+        assert proposer.draft_attn_groups[0].layer_names == expected_order
+        assert proposer.block_size == KERNEL_BLOCK_SIZE

tests/v1/streaming_input/test_gpu_model_runner_v2_streaming.py

@@ -41,6 +41,7 @@ def mock_model_runner_with_req_states():
     runner.sampler = None
     runner.prompt_logprobs_worker = None
     runner.is_last_pp_rank = False
+    runner.pooling_runner = None
 
     # Mock staged writes — they use Triton kernels that require GPU
     runner.req_states.apply_staged_writes = Mock()

vllm/_custom_ops.py

@@ -2391,17 +2391,6 @@ def topk_softmax(
     e_score_correction_bias: torch.Tensor | None = None,
     is_padding: torch.Tensor | None = None,
 ) -> None:
-    if current_platform.is_xpu():
-        # TODO: Remove after vllm-xpu-kernels supports is_padding.
-        torch.ops._moe_C.topk_softmax(
-            topk_weights,
-            topk_ids,
-            token_expert_indices,
-            gating_output,
-            renormalize,
-            e_score_correction_bias,
-        )
-        return
     torch.ops._moe_C.topk_softmax(
         topk_weights,
         topk_ids,
@@ -2447,21 +2436,6 @@ def topk_hash_softplus_sqrt(
     hash_indices_table: torch.Tensor | None = None,
     is_padding: torch.Tensor | None = None,
 ) -> None:
-    if current_platform.is_xpu():
-        # TODO: Remove after vllm-xpu-kernels supports is_padding.
-        torch.ops._moe_C.topk_softplus_sqrt(
-            topk_weights,
-            topk_indices,
-            token_expert_indices,
-            gating_output,
-            renormalize,
-            routed_scaling_factor,
-            e_score_correction_bias,
-            input_tokens,
-            hash_indices_table,
-        )
-        return
-
     torch.ops._moe_C.topk_softplus_sqrt(
         topk_weights,
         topk_indices,

vllm/compilation/passes/fx_utils.py

@@ -14,10 +14,6 @@ def is_func(node: fx.Node, target: Target) -> bool:
     return bool(node.op == "call_function" and node.target == target)
 
 
-def is_auto_func(node: fx.Node, op: OpOverload) -> bool:
-    return is_func(node, auto_functionalized) and node.args[0] == op
-
-
 # Returns the first auto_functionalized node with the given op (if it exists)
 def find_auto_fn_maybe(nodes: Iterable[fx.Node], op: OpOverload) -> fx.Node | None:
     for node in nodes:
@@ -42,13 +38,6 @@ def find_getitem_maybe(node: fx.Node, idx: int) -> fx.Node | None:
     return None
 
 
-# Returns the getitem node that extracts the idx-th element from node
-def find_getitem(node: fx.Node, idx: int) -> fx.Node:
-    ret = find_getitem_maybe(node, idx)
-    assert ret is not None, f"Could not find getitem {idx} in node {node}"
-    return ret
-
-
 # An auto-functionalization-aware utility for finding nodes with a specific op
 # Also handles op overload packets and finds all overloads
 def find_op_nodes(
@@ -67,11 +56,3 @@ def find_op_nodes(
     for n in graph.find_nodes(op="call_function", target=auto_functionalized):
         if n.args[0] == op:
             yield n
-
-
-# Asserts that the node only has one user and returns it
-# Even if a node has only 1 user, it might share storage with another node,
-# which might need to be taken into account.
-def get_only_user(node: fx.Node) -> fx.Node:
-    assert len(node.users) == 1
-    return next(iter(node.users))

vllm/compilation/passes/vllm_inductor_pass.py

@@ -333,12 +333,3 @@ class VllmFusionPatternMatcherPass(VllmPatternMatcherPass):
     def __call__(self, graph: torch.fx.Graph) -> None:
         self.matched_count = self.pm_pass.apply(graph)
         VllmPatternMatcherPass.match_table[self.pass_name] += self.matched_count
-
-
-class PrinterInductorPass(VllmInductorPass):
-    def __init__(self, name: str, config: VllmConfig) -> None:
-        super().__init__(config)
-        self.name = name
-
-    def __call__(self, graph: torch.fx.Graph) -> None:
-        self.dump_graph(graph, self.name)

vllm/config/speculative.py

@@ -910,6 +910,15 @@ class SpeculativeConfig:
                         f"Unsupported speculative method: '{self.method}'"
                     )
 
+                if self.method in ("eagle", "eagle3"):
+                    # EAGLE drafts share the target's positional space; a
+                    # draft checkpoint with a smaller max_position_embeddings
+                    # than the target under-sizes its rotary cache (#48894).
+                    SpeculativeConfig._maybe_override_draft_max_position_embeddings(
+                        self.draft_model_config.hf_config,
+                        self.target_model_config.max_model_len,
+                    )
+
                 # Replace hf_config for EAGLE draft_model
                 if self.method in ("eagle", "eagle3", "dflash"):
                     from vllm.transformers_utils.configs.eagle import EAGLEConfig
@@ -1130,6 +1139,40 @@ class SpeculativeConfig:
             )
         return result
 
+    @staticmethod
+    def _maybe_override_draft_max_position_embeddings(
+        draft_hf_config: PretrainedConfig,
+        target_max_model_len: int,
+    ) -> None:
+        """Raise an EAGLE draft's max_position_embeddings up to the target's.
+
+        The proposer feeds the draft positions up to the target's
+        max_model_len, while max_position_embeddings sizes the draft's
+        rotary cos_sin_cache. A smaller checkpoint value (e.g. 2048 for
+        yuhuili/EAGLE3-LLaMA3.1-Instruct-8B) makes that cache gather go
+        out of bounds (#48894).
+
+        Args:
+            draft_hf_config: The draft model's HF config, mutated in place.
+            target_max_model_len: The target model's max_model_len.
+        """
+        draft_max_position_embeddings = getattr(
+            draft_hf_config, "max_position_embeddings", None
+        )
+        if (
+            draft_max_position_embeddings is None
+            or draft_max_position_embeddings >= target_max_model_len
+        ):
+            return
+        logger.info(
+            "Overriding draft model max_position_embeddings from %d to the "
+            "target model's max_model_len (%d); EAGLE drafts share the "
+            "target's positional space.",
+            draft_max_position_embeddings,
+            target_max_model_len,
+        )
+        draft_hf_config.max_position_embeddings = target_max_model_len
+
     @staticmethod
     def _verify_and_get_draft_tp(
         target_parallel_config: ParallelConfig,

vllm/config/utils.py

@@ -415,31 +415,6 @@ class Range:
         return self.__str__()
 
 
-def handle_deprecated(
-    config: ConfigT,
-    old_name: str,
-    new_name_or_names: str | list[str],
-    removal_version: str,
-) -> None:
-    old_val = getattr(config, old_name)
-    if old_val is None:
-        return
-
-    if isinstance(new_name_or_names, str):
-        new_names = [new_name_or_names]
-    else:
-        new_names = new_name_or_names
-
-    msg = (
-        f"{old_name} is deprecated and will be removed in {removal_version}. "
-        f"Use {', '.join(new_names)} instead."
-    )
-    logger.warning(msg)
-
-    for new_name in new_names:
-        setattr(config, new_name, old_val)
-
-
 def get_from_deprecated_env_if_set(
     env_name: str,
     removal_version: str,

vllm/distributed/ec_transfer/ec_connector/base.py

@@ -275,3 +275,15 @@ class ECConnectorBase(ABC):
             get_finished().
         """
         return False, None
+
+    def has_pending_push_work(self) -> bool:
+        """Return True if the connector has push-mode work that requires
+        the engine main loop to keep stepping (e.g. for EPD,
+        Producer has push work when Xfer is in progress - Consumer
+        is reading it).
+        This mirrors exactly the KV Connector's has_pending_push_work().
+
+        Connectors that don't implement push-based EC transfer should
+        leave this as False.
+        """
+        return False

vllm/distributed/parallel_state.py

@@ -2081,9 +2081,6 @@ def model_parallel_is_initialized():
     return _TP is not None and _PP is not None
 
 
-_TP_STATE_PATCHED = False
-
-
 def get_tensor_model_parallel_world_size() -> int:
     """Return world size for the tensor model parallel group."""
     return get_tp_group().world_size

vllm/entrypoints/chat_utils.py

@@ -2015,15 +2015,6 @@ async def parse_chat_messages_async(
     return conversation, mm_data, mm_uuids
 
 
-def get_history_tool_calls_cnt(conversation: list[ConversationMessage]):
-    idx = 0
-    for msg in conversation:
-        if msg["role"] == "assistant":
-            tool_calls = msg.get("tool_calls")
-            idx += len(list(tool_calls)) if tool_calls is not None else 0  # noqa
-    return idx
-
-
 _KIMI_MODEL_TYPES = ("kimi_k2", "kimi_k25", "kimi_k3")
 
 

vllm/entrypoints/speech_to_text/base/serving.py

@@ -644,7 +644,7 @@ class SpeechToTextBaseServing(GenerateBaseServing):
                         TranscriptionResponseVerbose(
                             text=text,
                             language=request.language,
-                            duration=str(duration_s),
+                            duration=duration_s,
                             segments=total_segments,
                         ),
                     )
@@ -658,7 +658,7 @@ class SpeechToTextBaseServing(GenerateBaseServing):
                         TranslationResponseVerbose(
                             text=text,
                             language=request.language,
-                            duration=str(duration_s),
+                            duration=duration_s,
                             segments=total_segments,
                         ),
                     )

vllm/entrypoints/speech_to_text/transcription/protocol.py

@@ -384,7 +384,7 @@ class TranscriptionSegment(OpenAIBaseModel):
 
 
 class TranscriptionResponseVerbose(OpenAIBaseModel):
-    duration: str
+    duration: float
     """The duration of the input audio."""
 
     language: str

vllm/entrypoints/speech_to_text/translation/protocol.py

@@ -357,7 +357,7 @@ class TranslationSegment(OpenAIBaseModel):
 
 
 class TranslationResponseVerbose(OpenAIBaseModel):
-    duration: str
+    duration: float
     """The duration of the input audio."""
 
     language: str

vllm/kernels/xpu_ops.py

@@ -30,9 +30,6 @@ def rms_norm(
     x: Tensor, weight: Tensor | None, epsilon: float, variance_size: int | None = None
 ) -> Tensor:
     assert variance_size is None
-    if weight is None:
-        # Kernel requires weight tensor, pass ones
-        weight = torch.ones(x.shape[-1], device=x.device, dtype=x.dtype)
     output = torch.empty(x.shape, device=x.device, dtype=x.dtype)
     torch.ops._C.rms_norm(output, x, weight, epsilon)
     return output
@@ -58,8 +55,5 @@ def fused_add_rms_norm(
     variance_size: int | None = None,
 ) -> tuple[Tensor, Tensor]:
     assert variance_size is None
-    if weight is None:
-        # Kernel requires weight tensor, pass ones
-        weight = torch.ones(x.shape[-1], device=x.device, dtype=x.dtype)
     torch.ops._C.fused_add_rms_norm(x, x_residual, weight, epsilon)
     return x, x_residual

vllm/model_executor/layers/activation.py

@@ -8,14 +8,8 @@ import torch
 import torch.nn as nn
 import torch.nn.functional as F
 
-from vllm.distributed import (
-    divide,
-    get_tensor_model_parallel_rank,
-    get_tensor_model_parallel_world_size,
-)
 from vllm.logger import init_logger
 from vllm.model_executor.custom_op import CustomOp
-from vllm.model_executor.utils import set_weight_attrs
 from vllm.platforms import CpuArchEnum, current_platform
 from vllm.triton_utils import tl, triton
 from vllm.utils.collection_utils import LazyDict
@@ -788,48 +782,6 @@ class XIELU(CustomOp):
         return self.forward_native(input)
 
 
-class ScaledActivation(nn.Module):
-    """An activation function with post-scale parameters.
-
-    This is used for some quantization methods like AWQ.
-    """
-
-    def __init__(
-        self,
-        act_module: nn.Module,
-        intermediate_size: int,
-        input_is_parallel: bool = True,
-        params_dtype: torch.dtype | None = None,
-    ):
-        super().__init__()
-        self.act = act_module
-        self.input_is_parallel = input_is_parallel
-        if input_is_parallel:
-            tp_size = get_tensor_model_parallel_world_size()
-            intermediate_size_per_partition = divide(intermediate_size, tp_size)
-        else:
-            intermediate_size_per_partition = intermediate_size
-        if params_dtype is None:
-            params_dtype = torch.get_default_dtype()
-        self.scales = nn.Parameter(
-            torch.empty(intermediate_size_per_partition, dtype=params_dtype)
-        )
-        set_weight_attrs(self.scales, {"weight_loader": self.weight_loader})
-
-    def forward(self, x: torch.Tensor) -> torch.Tensor:
-        return self.act(x) / self.scales
-
-    def weight_loader(self, param: nn.Parameter, loaded_weight: torch.Tensor):
-        param_data = param.data
-        if self.input_is_parallel:
-            tp_rank = get_tensor_model_parallel_rank()
-            shard_size = param_data.shape[0]
-            start_idx = tp_rank * shard_size
-            loaded_weight = loaded_weight.narrow(0, start_idx, shard_size)
-        assert param_data.shape == loaded_weight.shape
-        param_data.copy_(loaded_weight)
-
-
 _ACTIVATION_REGISTRY = LazyDict(
     {
         "gelu": lambda: GELU(),

vllm/model_executor/layers/conv.py

@@ -155,58 +155,6 @@ class Conv2dLayer(ConvLayerBase):
         return self._forward_conv(x)
 
 
-class CausalConv2dLayer(Conv2dLayer):
-    """
-    A causal version of nn.Conv2d where each location in the 2D matrix would
-    have no access to locations on its right or down
-    All arguments are the same as nn.Conv2d except padding which should be
-    set as None
-    """
-
-    def __init__(
-        self,
-        in_channels: int,
-        out_channels: int,
-        kernel_size: int,
-        stride: int,
-        padding: int = 0,
-        dilation: int = 1,
-        groups: int = 1,
-        bias: bool = True,
-        padding_mode: str = "zeros",
-        *,
-        params_dtype: torch.dtype | None = None,
-    ) -> None:
-        if padding is not None:
-            raise ValueError(
-                "Argument padding should be set to None for CausalConv2dLayer."
-            )
-        self._left_padding: int = kernel_size - 1
-        self._right_padding: int = stride - 1
-        padding = 0
-
-        super().__init__(
-            in_channels,
-            out_channels,
-            kernel_size,
-            stride,
-            padding,
-            dilation,
-            groups,
-            bias,
-            padding_mode,
-            params_dtype=params_dtype,
-        )
-
-    def forward(
-        self,
-        x: torch.Tensor,
-    ) -> torch.Tensor:
-        x = F.pad(x, pad=(self._left_padding, self._right_padding, 0, 0))
-        x = super().forward(x)
-        return x
-
-
 # --8<-- [start:conv3d]
 @CustomOp.register("conv3d")
 class Conv3dLayer(ConvLayerBase):

vllm/model_executor/layers/quantization/fp8.py

@@ -4,7 +4,6 @@
 from typing import TYPE_CHECKING, Any
 
 import torch
-from torch.utils._python_dispatch import TorchDispatchMode
 
 import vllm.envs as envs
 import vllm.model_executor.layers.fused_moe.modular_kernel as mk
@@ -234,26 +233,6 @@ class Fp8Config(QuantizationConfig):
         return cache_scale_mapper | QuantizationConfig.get_cache_scale_mapper()
 
 
-class CopyNumelCounter(TorchDispatchMode):
-    """
-    Tracks total number of elements modified with `copy_`. Useful for keeping
-    track of weight loading where underlying weights can be arbitrarily
-    transformed (such as with `narrow`) before calling copy.
-    """
-
-    def __init__(self):
-        super().__init__()
-        self.copied_numel = 0
-
-    def __torch_dispatch__(self, func, types, args=(), kwargs=None):
-        if kwargs is None:
-            kwargs = {}
-        out = func(*args, **kwargs)
-        if func == torch.ops.aten.copy_.default:
-            self.copied_numel += args[0].numel()
-        return out
-
-
 class Fp8LinearMethod(LinearMethodBase):
     """Linear method for FP8.
     Supports loading FP8 checkpoints with static weight scale and

vllm/model_executor/layers/utils.py

@@ -25,12 +25,6 @@ MOE_LAYER_ROUTER_GATE_SUFFIXES = {
 }
 
 
-def is_layer_moe_router_gate(prefix: str) -> bool:
-    if not prefix:
-        return False
-    return prefix.rsplit(".", 1)[-1] in MOE_LAYER_ROUTER_GATE_SUFFIXES
-
-
 def get_token_bin_counts_and_mask(
     tokens: torch.Tensor,
     vocab_size: int,

vllm/model_executor/models/config.py

@@ -768,6 +768,20 @@ class Qwen3_5ForConditionalGenerationConfig(VerifyAndUpdateConfig):
             )
 
 
+class Qwen3_5ForCausalLMConfig(Qwen3_5ForConditionalGenerationConfig):
+    @staticmethod
+    def verify_and_update_config(vllm_config: "VllmConfig") -> None:
+        Qwen3_5ForConditionalGenerationConfig.verify_and_update_config(vllm_config)
+
+        # Text-only Qwen3.5 models use one-dimensional positions. Remove the
+        # M-RoPE fields inherited from the multimodal configuration.
+        hf_text_config = vllm_config.model_config.hf_text_config
+        rope_parameters = getattr(hf_text_config, "rope_parameters", None)
+        if rope_parameters is not None:
+            rope_parameters.pop("mrope_section", None)
+            rope_parameters.pop("mrope_interleaved", None)
+
+
 class ColQwen3_5Config(Qwen3_5ForConditionalGenerationConfig):
     """Apply the attention contract declared by a ColQwen3.5 checkpoint."""
 
@@ -884,7 +898,9 @@ MODELS_CONFIG_MAP: dict[str, type[VerifyAndUpdateConfig]] = {
     "Qwen2ForRewardModel": Qwen2ForRewardModelConfig,
     "Qwen3ForSequenceClassification": Qwen3ForSequenceClassificationConfig,
     "Qwen3VLForSequenceClassification": Qwen3VLForSequenceClassificationConfig,
+    "Qwen3_5ForCausalLM": Qwen3_5ForCausalLMConfig,
     "Qwen3_5ForConditionalGeneration": Qwen3_5ForConditionalGenerationConfig,
+    "Qwen3_5MoeForCausalLM": Qwen3_5ForCausalLMConfig,
     "Qwen3_5MoeForConditionalGeneration": Qwen3_5ForConditionalGenerationConfig,
     "UnlimitedOCRForCausalLM": UnlimitedOCRForCausalLMConfig,
     "VoyageQwen3BidirectionalEmbedModel": VoyageQwen3BidirectionalEmbedModelConfig,

vllm/model_executor/models/glm_ocr.py

@@ -242,10 +242,6 @@ class GlmOcrVisionBlock(Glm4vVisionBlock):
         )
 
 
-class GlmOcrVisionPatchEmbed(Glm4vVisionPatchEmbed):
-    pass
-
-
 class GlmOcrPatchMerger(Glm4vPatchMerger):
     pass
 

vllm/model_executor/models/idefics3.py

@@ -103,41 +103,6 @@ class Idefics3ProcessingInfo(BaseProcessingInfo):
     def get_supported_mm_limits(self) -> Mapping[str, int | None]:
         return {"image": None}
 
-    def _resize_output_size(
-        self,
-        *,
-        height: int,
-        width: int,
-        max_len: int | None = None,
-        min_len: int = 1,
-        max_size: int | None = None,
-    ) -> tuple[int, int]:
-        # Set default value for max_len if not provided
-        max_len = max(height, width) if max_len is None else max_len
-        aspect_ratio = width / height
-
-        # Handle the maximum size constraint
-        if max_size is not None:
-            max_len = min(max_len, max_size)
-
-        # Adjust dimensions according to the aspect ratio
-        if width >= height:
-            width = max_len
-            height = int(width / aspect_ratio)
-        else:
-            height = max_len
-            width = int(height * aspect_ratio)
-
-        # Ensure both width and height are even (if needed)
-        height += height % 2
-        width += width % 2
-
-        # Ensure dimensions are not smaller than the minimum length
-        height = max(height, min_len)
-        width = max(width, min_len)
-
-        return height, width
-
     def _get_image_feature_grid_size(
         self,
         *,

vllm/model_executor/models/llava_onevision2.py

@@ -1531,13 +1531,6 @@ class LlavaOnevision2MultiModalDataParser(MultiModalDataParser):
 class LlavaOnevision2MultiModalProcessor(
     BaseMultiModalProcessor[LlavaOnevision2ProcessingInfo]
 ):
-    def _get_data_parser(self) -> MultiModalDataParser:
-        # Retained for symmetry; vLLM actually fetches the parser via
-        # info.get_data_parser() (see ProcessingInfo override above).
-        return LlavaOnevision2MultiModalDataParser(
-            self.info.get_hf_config().vision_config.spatial_merge_size
-        )
-
     def _call_hf_processor(
         self,
         prompt: str,

vllm/model_executor/models/nemotron_parse.py

@@ -77,24 +77,6 @@ class BartScaledWordEmbedding(VocabParallelEmbedding):
         return super().forward(input_ids) * self.embed_scale
 
 
-class BartParallelLMHead(ParallelLMHead):
-    """
-    This module overrides ParallelLMHead's
-    forward by dividing by embeddings scale,
-    yielding effectively the inverse of
-    BartScaledWordEmbedding
-    """
-
-    def __init__(
-        self, num_embeddings: int, embedding_dim: int, embed_scale: float = 1.0
-    ):
-        super().__init__(num_embeddings, embedding_dim)
-        self.embed_scale = embed_scale
-
-    def forward(self, input_ids: torch.Tensor) -> torch.Tensor:
-        return super().forward(input_ids) / self.embed_scale
-
-
 class BartDecoderLayer(nn.Module):
     def __init__(
         self,

vllm/model_executor/models/phi4mm.py

@@ -521,32 +521,6 @@ class Phi4MMAudioEmbeddingInputs(TensorSchema):
 Phi4MMAudioInputs: TypeAlias = Phi4MMAudioFeatureInputs | Phi4MMAudioEmbeddingInputs
 
 
-def cat_with_pad(tensors, dim, padding_value=0):
-    """
-    cat along dim, while pad to max for all other dims
-    """
-    ndim = tensors[0].dim()
-    assert all(t.dim() == ndim for t in tensors[1:]), (
-        "All tensors must have the same number of dimensions"
-    )
-
-    out_size = [max(t.shape[i] for t in tensors) for i in range(ndim)]
-    out_size[dim] = sum(t.shape[dim] for t in tensors)
-    output = tensors[0].new_full(out_size, padding_value)
-
-    index = 0
-    for t in tensors:
-        # Create a slice list where every dimension except dim is full slice
-        slices = [slice(0, t.shape[d]) for d in range(ndim)]
-        # Update only the concat dimension slice
-        slices[dim] = slice(index, index + t.shape[dim])
-
-        output[slices] = t
-        index += t.shape[dim]
-
-    return output
-
-
 def stack_with_pad(
     tensors: torch.Tensor | list[torch.Tensor],
     padding_value: int | float = 0,

vllm/model_executor/models/phi4mm_utils.py

@@ -1595,14 +1595,6 @@ class AttModule(nn.Module):
         return x, memory, pos_emb, att_mask
 
 
-class AttBlock(BlockBase, AttModule):
-    """Attention Block module to support both Attention and Block module."""
-
-    def memory_dims(self, max_len: bool = False) -> tuple[int, int]:
-        """memory dimensions"""
-        return (1, self.input_size)
-
-
 def masked_softmax(
     scores: Tensor,
     mask: Tensor | None,

vllm/model_executor/models/qwen3_5.py

@@ -282,6 +282,7 @@ class Qwen3_5Model(Qwen3NextModel):
 class Qwen3_5ForCausalLMBase(
     nn.Module,
     HasInnerState,
+    IsHybrid,
     SupportsEagle3,
     SupportsLoRA,
     SupportsPP,
@@ -361,6 +362,45 @@ class Qwen3_5ForCausalLMBase(
 
         return hidden_states
 
+    @classmethod
+    def get_mamba_state_dtype_from_config(
+        cls,
+        vllm_config: "VllmConfig",
+    ) -> tuple[torch.dtype, torch.dtype]:
+        return MambaStateDtypeCalculator.gated_delta_net_state_dtype(
+            vllm_config.model_config.dtype,
+            vllm_config.cache_config.mamba_cache_dtype,
+            vllm_config.cache_config.mamba_ssm_cache_dtype,
+        )
+
+    @classmethod
+    def get_mamba_state_shape_from_config(
+        cls, vllm_config: "VllmConfig"
+    ) -> tuple[tuple[int, int], tuple[int, int]]:
+        parallel_config = vllm_config.parallel_config
+        hf_config = vllm_config.model_config.hf_text_config
+        tp_size = parallel_config.tensor_parallel_size
+        num_spec = (
+            vllm_config.speculative_config.num_speculative_tokens
+            if vllm_config.speculative_config
+            else 0
+        )
+        return MambaStateShapeCalculator.gated_delta_net_state_shape(
+            tp_size,
+            hf_config.linear_num_key_heads,
+            hf_config.linear_num_value_heads,
+            hf_config.linear_key_head_dim,
+            hf_config.linear_value_head_dim,
+            hf_config.linear_conv_kernel_dim,
+            num_spec,
+        )
+
+    @classmethod
+    def get_mamba_state_copy_func(
+        cls,
+    ) -> tuple[MambaStateCopyFunc, MambaStateCopyFunc]:
+        return MambaStateCopyFuncCalculator.gated_delta_net_state_copy_func()
+
     def compute_logits(
         self,
         hidden_states: torch.Tensor,

vllm/model_executor/models/registry.py

@@ -195,6 +195,8 @@ _TEXT_GENERATION_MODELS = {
     "Qwen2MoeForCausalLM": ("qwen2_moe", "Qwen2MoeForCausalLM"),
     "Qwen3ForCausalLM": ("qwen3", "Qwen3ForCausalLM"),
     "Qwen3MoeForCausalLM": ("qwen3_moe", "Qwen3MoeForCausalLM"),
+    "Qwen3_5ForCausalLM": ("qwen3_5", "Qwen3_5ForCausalLM"),
+    "Qwen3_5MoeForCausalLM": ("qwen3_5", "Qwen3_5MoeForCausalLM"),
     "RWForCausalLM": ("falcon", "FalconForCausalLM"),
     "SarvamMoEForCausalLM": ("sarvam", "SarvamMoEForCausalLM"),
     "SarvamMLAForCausalLM": ("sarvam", "SarvamMLAForCausalLM"),

vllm/model_executor/models/skyworkr1v.py

@@ -7,7 +7,7 @@
 # Copyright (c) 2025 Skywork
 # Licensed under The MIT License [see LICENSE for details]
 # --------------------------------------------------------
-from collections.abc import Iterable, Mapping
+from collections.abc import Iterable
 from typing import Annotated, Literal, TypeAlias
 
 import torch
@@ -15,8 +15,6 @@ import torch.nn as nn
 from transformers import PretrainedConfig
 
 from vllm.config import VllmConfig
-from vllm.config.multimodal import BaseDummyOptions
-from vllm.inputs import MultiModalDataDict
 from vllm.model_executor.layers.linear import ReplicatedLinear
 from vllm.model_executor.layers.quantization import QuantizationConfig
 from vllm.model_executor.layers.quantization.auto_awq import AutoAWQConfig
@@ -24,7 +22,6 @@ from vllm.model_executor.models.intern_vit import (
     InternVisionModel,
 )
 from vllm.multimodal import MULTIMODAL_REGISTRY
-from vllm.multimodal.processing import BaseDummyInputsBuilder
 from vllm.sequence import IntermediateTensors
 from vllm.transformers_utils.processors.internvl import (
     InternVLImageProcessor,
@@ -117,33 +114,6 @@ class SkyworkR1VProcessingInfo(BaseInternVLProcessingInfo):
         )
 
 
-class SkyworkR1VDummyInputsBuilder(BaseDummyInputsBuilder[SkyworkR1VProcessingInfo]):
-    def get_dummy_text(self, mm_counts: Mapping[str, int]) -> str:
-        num_images = mm_counts.get("image", 0)
-
-        return "<image>" * num_images
-
-    def get_dummy_mm_data(
-        self,
-        seq_len: int,
-        mm_counts: Mapping[str, int],
-        mm_options: Mapping[str, BaseDummyOptions],
-    ) -> MultiModalDataDict:
-        target_width, target_height = self.info.get_image_size_with_most_features()
-        num_images = mm_counts.get("image", 0)
-
-        image_overrides = mm_options.get("image")
-
-        return {
-            "image": self._get_dummy_images(
-                width=target_width,
-                height=target_height,
-                num_images=num_images,
-                overrides=image_overrides,
-            )
-        }
-
-
 @MULTIMODAL_REGISTRY.register_processor(
     BaseInternVLMultiModalProcessor,
     info=SkyworkR1VProcessingInfo,

vllm/model_executor/models/utils.py

@@ -27,7 +27,6 @@ from vllm.multimodal import NestedTensors
 from vllm.sequence import IntermediateTensors
 from vllm.utils.math_utils import cdiv
 from vllm.utils.torch_utils import (
-    async_tensor_h2d,
     direct_register_custom_op,
 )
 
@@ -673,17 +672,6 @@ def _merge_multimodal_embeddings(
     return inputs_embeds
 
 
-def isin_list(
-    elements: torch.Tensor,
-    test_elements_list: list[int],
-) -> torch.Tensor:
-    test_elements = async_tensor_h2d(
-        test_elements_list, dtype=torch.int64, device=elements.device
-    )
-
-    return torch.isin(elements, test_elements)
-
-
 class StageMissingLayer(nn.Module):
     def __init__(self, stage_name: str, module: nn.Module | None = None) -> None:
         super().__init__()

vllm/model_executor/models/vision.py

@@ -580,40 +580,3 @@ def run_dp_sharded_mrope_vision_model(
         "Found unassigned embeddings"
     )
     return out_embeddings
-
-
-def get_llm_pos_ids_for_vision(
-    start_idx: int,
-    vision_idx: int,
-    spatial_merge_size: int,
-    t_index: list[int],
-    grid_hs: torch.Tensor,
-    grid_ws: torch.Tensor,
-) -> torch.Tensor:
-    llm_pos_ids_list = []
-    llm_grid_h = grid_hs[vision_idx] // spatial_merge_size
-    llm_grid_w = grid_ws[vision_idx] // spatial_merge_size
-    h_index = (
-        torch.arange(llm_grid_h)
-        .view(1, -1, 1)
-        .expand(len(t_index), -1, llm_grid_w)
-        .flatten()
-    )
-    w_index = (
-        torch.arange(llm_grid_w)
-        .view(1, 1, -1)
-        .expand(len(t_index), llm_grid_h, -1)
-        .flatten()
-    )
-    t_index_tensor = (
-        torch.Tensor(t_index)
-        .to(llm_grid_h.device)
-        .view(-1, 1)
-        .expand(-1, llm_grid_h * llm_grid_w)
-        .long()
-        .flatten()
-    )
-    _llm_pos_ids = torch.stack([t_index_tensor, h_index, w_index])
-    llm_pos_ids_list.append(_llm_pos_ids + start_idx)
-    llm_pos_ids = torch.cat(llm_pos_ids_list, dim=1)
-    return llm_pos_ids

vllm/models/minimax_m3/common/mm_preprocess.py

@@ -3,8 +3,9 @@
 
 import math
 from collections.abc import Mapping, Sequence
-from typing import cast
+from typing import Any, Literal, cast
 
+import numpy.typing as npt
 import torch
 from transformers import BatchFeature
 from transformers.video_utils import VideoMetadata
@@ -469,10 +470,39 @@ class MiniMaxM3VLMultiModalProcessor(
         ]
 
 
-# TODO(Isotr0py): Tie with MinimaxVideoProcessor
-# after https://github.com/vllm-project/vllm/pull/44126
-@VIDEO_LOADER_REGISTRY.register("minimax_m3_vl")
+@VIDEO_LOADER_REGISTRY.register(
+    name="minimax_m3_vl",
+    video_processor="MiniMaxM3VLVideoProcessor",
+)
 class MiniMaxM3VideoBackend(VideoBackend):
+    @classmethod
+    def load_bytes(
+        cls,
+        data: bytes,
+        num_frames: int = -1,
+        fps: int = 1,
+        max_duration: int = 300,
+        frame_recovery: bool = False,
+        *,
+        backend: Literal[
+            "opencv",
+            "pyav",
+            "torchcodec",
+            "pynvvideocodec",
+            "deepstream",
+        ] = "opencv",
+        **kwargs,
+    ) -> tuple[npt.NDArray, dict[str, Any]]:
+        return super().load_bytes(
+            data,
+            num_frames=num_frames,
+            fps=fps,
+            max_duration=max_duration,
+            frame_recovery=frame_recovery,
+            backend=backend,
+            **kwargs,
+        )
+
     @classmethod
     def compute_frames_index_to_sample(
         cls,
@@ -483,7 +513,6 @@ class MiniMaxM3VideoBackend(VideoBackend):
         total_frames = source.total_frames_num
         video_fps = source.original_fps
         fps = target.fps
-
         if total_frames <= 0 or video_fps <= 0 or fps <= 0:
             return [0] if total_frames > 0 else []
 
@@ -503,8 +532,9 @@ class MiniMaxM3VideoBackend(VideoBackend):
                 break
             indices.append(target_frame)
             prev_kept_ts = target_frame / video_fps
-
-        last_frame_idx = total_frames - 1
+        # Because HF sample_frames includes the last frame,
+        # we will use HF as the standard.
+        last_frame_idx = total_frames
         last_ts = last_frame_idx / video_fps
         if indices and indices[-1] != last_frame_idx and last_ts - prev_kept_ts > eps:
             indices.append(last_frame_idx)

vllm/multimodal/hasher.py

@@ -81,6 +81,11 @@ class MultiModalHasher:
             ):
                 return (exif[Image.ExifTags.Base.ImageID].bytes,)
 
+            if obj.io_config:
+                return cls.iter_item_to_bytes(
+                    "image",
+                    {"io_config": obj.io_config, "data": obj.original_bytes},
+                )
             return cls.iter_item_to_bytes("image", obj.original_bytes)
 
         if isinstance(obj, MediaWithBytes) and isinstance(obj.media, np.ndarray):

vllm/multimodal/media/base.py

@@ -29,6 +29,9 @@ class MediaWithBytes(Generic[_T]):
 
     media: _T
     original_bytes: bytes = field(repr=False)
+    io_config: dict[str, Any] | None = None
+    """Decode settings that altered the media relative to `original_bytes`
+    (e.g. `image_mode` conversion), so they participate in cache hashing."""
 
     def __array__(self, *args, **kwargs) -> np.ndarray:
         """Allow np.array(obj) to return np.array(obj.media)."""

vllm/multimodal/media/image.py

@@ -86,10 +86,17 @@ class ImageMediaIO(MediaIO[Image.Image]):
                 )
             image = normalize_image(image)
             image.load()
-            image = self._convert_image_mode(image)
+            converted = self._convert_image_mode(image)
         except (OSError, Image.UnidentifiedImageError) as e:
             raise ValueError(f"Failed to load image: {e}") from e
-        return MediaWithBytes(image, data)
+
+        io_config = None
+        if converted is not image:
+            io_config = {
+                "image_mode": self.image_mode,
+                "rgba_background_color": self.rgba_background_color,
+            }
+        return MediaWithBytes(converted, data, io_config)
 
     def load_base64(self, media_type: str, data: str) -> MediaWithBytes[Image.Image]:
         return self.load_bytes(pybase64.b64decode(data, validate=True))

vllm/multimodal/processing/inputs.py

@@ -65,7 +65,7 @@ class ProcessorInputs:
                         **{modality: item},
                         **hf_processor_mm_kwargs,
                     )
-                    for item in data_items
+                    for item in data_items.get_all_items_for_hash()
                 ]
 
         return mm_hashes

vllm/profiler/utils.py

@@ -12,15 +12,6 @@ from torch._C._profiler import _EventType, _ProfilerEvent, _TensorMetadata
 #
 
 
-def trim_string_front(string: str, width: int) -> str:
-    if len(string) > width:
-        offset = len(string) - width + 3
-        string = string[offset:]
-        if len(string) > 3:
-            string = "..." + string[3:]
-    return string
-
-
 def trim_string_back(string: str, width: int) -> str:
     if len(string) > width:
         offset = len(string) - width + 3

vllm/transformers_utils/config.py

@@ -124,7 +124,9 @@ _CONFIG_REGISTRY: dict[str, type[PretrainedConfig]] = LazyConfigDict(
     qwen3_asr="Qwen3ASRConfig",
     qwen3_next="Qwen3NextConfig",
     qwen3_5="Qwen3_5Config",
+    qwen3_5_text="Qwen3_5TextConfig",
     qwen3_5_moe="Qwen3_5MoeConfig",
+    qwen3_5_moe_text="Qwen3_5MoeTextConfig",
     laguna="LagunaConfig",
     lfm2_moe="Lfm2MoeConfig",
     **{"unlimited-ocr": "UnlimitedOCRConfig"},

vllm/utils/cache.py

@@ -182,10 +182,6 @@ class LRUCache(cachetools.LRUCache[_K, _V]):
 
         self.popitem(remove_pinned=remove_pinned)
 
-    def _remove_old_if_needed(self) -> None:
-        while self.currsize > self.capacity:
-            self.remove_oldest()
-
     def popitem(self, remove_pinned: bool = False):
         """Remove and return the `(key, value)` pair least recently used."""
         if not remove_pinned:

vllm/utils/multi_stream_utils.py

@@ -8,10 +8,6 @@ from typing import Any
 import torch
 
 
-class AuxStreamType(Enum):
-    Attention = 1
-
-
 class EventType(Enum):
     Main = 0
     Attention = 1

vllm/v1/core/kv_cache_utils.py

@@ -427,26 +427,6 @@ class FreeKVCacheBlockQueue:
             curr_block = curr_block.next_free_block
 
 
-def need_extra_keys(request: Request) -> bool:
-    """Check whether the blocks allocated to this request need extra hash keys.
-
-    Args:
-        request (Request): The request.
-
-    Returns:
-        bool: Whether blocks allocated to this request need extra hash keys.
-    """
-
-    # Multimodal requests need to include the MM hash.
-    # LoRA requests need to include the LoRA name.
-    # Request with provided cache salt need to include the salt.
-    return (
-        bool(request.mm_features)
-        or (request.lora_request is not None)
-        or (request.cache_salt is not None)
-    )
-
-
 def _gen_mm_extra_hash_keys(
     request: Request, start_token_idx: int, end_token_idx: int, start_mm_idx: int
 ) -> tuple[list[Any], int]:
@@ -1053,20 +1033,6 @@ def _get_kv_cache_groups_uniform_type(
     return [KVCacheGroupSpec(list(spec.kv_cache_specs.keys()), spec)]
 
 
-def is_kv_cache_page_size_uniform(kv_cache_spec: dict[str, KVCacheSpec]) -> bool:
-    """
-    Whether all layers in the given KVCacheSpec have the same page size.
-    Args:
-        kv_cache_spec: The KVCacheSpec of each attention layer in the model
-
-    Returns:
-        True if all layers have the same page size, False otherwise.
-    """
-
-    page_sizes = {layer.page_size_bytes for layer in kv_cache_spec.values()}
-    return len(page_sizes) == 1
-
-
 def unify_kv_cache_spec_page_size(
     kv_cache_spec: dict[str, KVCacheSpec],
 ) -> dict[str, KVCacheSpec]:

vllm/v1/core/sched/async_scheduler.py

@@ -49,23 +49,18 @@ class AsyncScheduler(Scheduler):
                 request.next_decode_eligible_step = self.current_step + self.pp_size
 
     def _update_request_with_output(
-        self, request: Request, new_token_ids: list[int]
+        self, request: Request, new_token_ids: list[int], is_stale: bool = False
     ) -> tuple[list[int], bool]:
-        if request.async_tokens_to_discard > 0:
-            # The request was force-preempted in reset_prefix_cache; drop one
-            # stale in-flight async output frame per call until the counter
-            # is drained.
-            request.async_tokens_to_discard -= 1
-            return [], False
-
         status_before_update = request.status
         new_token_ids, stopped = super()._update_request_with_output(
             request, new_token_ids
         )
 
-        # Update the number of output placeholders.
-        request.num_output_placeholders -= len(new_token_ids)
-        assert request.num_output_placeholders >= 0
+        # Placeholders were zeroed at preemption; a stale delivery must not
+        # decrement them (it would underflow).
+        if not is_stale:
+            request.num_output_placeholders -= len(new_token_ids)
+            assert request.num_output_placeholders >= 0
 
         # Cache the new tokens. Preempted requests should be skipped.
         if status_before_update == RequestStatus.RUNNING:

vllm/v1/core/sched/scheduler.py

@@ -694,6 +694,17 @@ class Scheduler(SchedulerInterface):
                     step_skipped_waiting.prepend_request(request)
                     continue
 
+                if (
+                    request.num_stale_output_tokens > 0
+                    and not request.drop_stale_output
+                ):
+                    # Deliverable stale output still in flight: resuming now
+                    # could resample a position that output later delivers.
+                    # It drains within the pipeline depth.
+                    request_queue.pop_request()
+                    step_skipped_waiting.prepend_request(request)
+                    continue
+
                 # Check that adding the request still respects the max_loras
                 # constraint.
                 if (
@@ -1244,11 +1255,17 @@ class Scheduler(SchedulerInterface):
 
         return new_block_ids_to_zero or None
 
-    def _preempt_request(self, request: Request, timestamp: float) -> None:
+    def _preempt_request(
+        self, request: Request, timestamp: float, drop_stale_output: bool = False
+    ) -> None:
         """Preempt a request and put it back to the waiting queue.
 
         NOTE: The request should be popped from the running queue outside of this
         method.
+
+        drop_stale_output: drop (rather than deliver) any in-flight output; used
+        by reset_prefix_cache, whose same-step resume would otherwise deliver
+        tokens out of order.
         """
         assert request.status == RequestStatus.RUNNING, (
             "Only running requests can be preempted"
@@ -1260,6 +1277,18 @@ class Scheduler(SchedulerInterface):
         request.num_computed_tokens = 0
         if request.spec_token_ids:
             request.spec_token_ids = []
+        # Async scheduling: mark all in-flight output as stale. Its tokens are
+        # still delivered on return (dropping them would perturb spec-decode
+        # acceptance) but must not mutate the reset counters; each step drains
+        # its share in update_from_output. num_in_flight_tokens already
+        # includes any undrained stale share, so assign rather than accumulate.
+        # An undrained drop-mode share stays dropped: its positions have
+        # already been resampled.
+        request.drop_stale_output = drop_stale_output or (
+            request.drop_stale_output and request.num_stale_output_tokens > 0
+        )
+        request.num_stale_output_tokens = request.num_in_flight_tokens
+        request.num_output_placeholders = 0
         request.num_preemptions += 1
         if self.log_stats:
             request.record_event(EngineCoreEventType.PREEMPTED, timestamp)
@@ -1687,8 +1716,14 @@ class Scheduler(SchedulerInterface):
         for req_id, num_tokens_scheduled in num_scheduled_tokens.items():
             assert num_tokens_scheduled > 0
             request = self.requests.get(req_id)
+            output_is_stale = False
             if request is not None:
                 request.num_in_flight_tokens -= num_tokens_scheduled
+                # Drain any stale share (see _preempt_request) in lockstep.
+                if request.num_stale_output_tokens > 0:
+                    output_is_stale = True
+                    request.num_stale_output_tokens -= num_tokens_scheduled
+                    assert request.num_stale_output_tokens >= 0
             if failed_kv_load_req_ids and req_id in failed_kv_load_req_ids:
                 # skip failed or rescheduled requests from KV load failure
                 continue
@@ -1702,6 +1737,10 @@ class Scheduler(SchedulerInterface):
                 # In this case, we use is_finished() to check.
                 continue
 
+            # Drop-mode stale output (same-step resume) is discarded entirely.
+            if output_is_stale and request.drop_stale_output:
+                continue
+
             req_index = model_runner_output.req_id_to_index[req_id]
             generated_token_ids = (
                 sampled_token_ids[req_index] if sampled_token_ids else []
@@ -1710,28 +1749,22 @@ class Scheduler(SchedulerInterface):
             scheduled_spec_token_ids = (
                 scheduler_output.scheduled_spec_decode_tokens.get(req_id)
             )
-            # Skip a stale frame still pending discard (async_tokens_to_discard
-            # > 0): its pre-reset rejection count would underflow the counters.
-            if (
-                scheduled_spec_token_ids
-                and (generated_token_ids or self.num_sampled_tokens_per_step == 0)
-                and request.async_tokens_to_discard == 0
+            if scheduled_spec_token_ids and (
+                generated_token_ids or self.num_sampled_tokens_per_step == 0
             ):
                 num_draft_tokens = len(scheduled_spec_token_ids)
                 num_sampled = self.num_sampled_tokens_per_step
                 num_accepted = max(len(generated_token_ids) - num_sampled, 0)
                 num_rejected = num_draft_tokens - num_accepted
-                # num_computed_tokens represents the number of tokens
-                # processed in the current step, considering scheduled
-                # tokens and rejections. If some tokens are rejected,
-                # num_computed_tokens is decreased by the number of rejected
-                # tokens.
-                if request.num_computed_tokens > 0:
-                    request.num_computed_tokens -= num_rejected
-                # If async scheduling, num_output_placeholders also includes
-                # the scheduled spec tokens count and so is similarly adjusted.
-                if request.num_output_placeholders > 0:
-                    request.num_output_placeholders -= num_rejected
+                # Rejections roll back num_computed_tokens (and, under async
+                # scheduling, num_output_placeholders, which covers the spec
+                # tokens). A stale rejection count predates the preemption
+                # rollback and must not apply.
+                if not output_is_stale:
+                    if request.num_computed_tokens > 0:
+                        request.num_computed_tokens -= num_rejected
+                    if request.num_output_placeholders > 0:
+                        request.num_output_placeholders -= num_rejected
                 spec_decoding_stats = self.make_spec_decoding_stats(
                     spec_decoding_stats,
                     num_draft_tokens=num_draft_tokens,
@@ -1757,7 +1790,7 @@ class Scheduler(SchedulerInterface):
             # Check for stop and update request status.
             if new_token_ids:
                 new_token_ids, stopped = self._update_request_with_output(
-                    request, new_token_ids
+                    request, new_token_ids, is_stale=output_is_stale
                 )
             elif request.pooling_params and pooler_output is not None:
                 # Pooling stops as soon as there is output.
@@ -1899,6 +1932,7 @@ class Scheduler(SchedulerInterface):
         if stopped_preempted_reqs:
             # This is a rare case and unlikely to impact performance.
             self.waiting.remove_requests(stopped_preempted_reqs)
+            self.skipped_waiting.remove_requests(stopped_preempted_reqs)
 
         error_req_ids = set(self.grammar_compile_error_reqs)
         self.grammar_compile_error_reqs.clear()
@@ -2041,8 +2075,9 @@ class Scheduler(SchedulerInterface):
         return False
 
     def _update_request_with_output(
-        self, request: Request, new_token_ids: list[int]
+        self, request: Request, new_token_ids: list[int], is_stale: bool = False
     ) -> tuple[list[int], bool]:
+        # is_stale is only used by the AsyncScheduler override.
         # Append generated tokens and check for stop. Note that if
         # a request is still being prefilled, we expect the model runner
         # to return empty token ids for the request.
@@ -2362,6 +2397,10 @@ class Scheduler(SchedulerInterface):
             self.has_unfinished_requests()
             or self.has_finished_requests()
             or (self.connector is not None and self.connector.has_pending_push_work())
+            or (
+                self.ec_connector is not None
+                and self.ec_connector.has_pending_push_work()
+            )
         )
 
     def reset_prefix_cache(
@@ -2384,15 +2423,7 @@ class Scheduler(SchedulerInterface):
             # running queue in FIFO order.
             while self.running:
                 request = self.running.pop()
-                self._preempt_request(request, timestamp)
-                # For async scheduling, any output frames already in flight at
-                # preemption time are now stale and must be discarded when they
-                # return. num_output_placeholders is exactly that count: 0 if
-                # the engine has drained (e.g. pause_generation(keep) waited
-                # for idle), 1 for vanilla async mid-step, or 1 + spec/PP frames
-                # otherwise.
-                request.async_tokens_to_discard = request.num_output_placeholders
-                request.num_output_placeholders = 0
+                self._preempt_request(request, timestamp, drop_stale_output=True)
 
             # Clear scheduled request ids cache. Since we are forcing preemption
             # + resumption in the same step, we must act as if these requests were

vllm/v1/kv_offload/cpu/spec.py

@@ -2,6 +2,7 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 from typing import Any
 
+import torch
 from typing_extensions import override
 
 from vllm.platforms import current_platform
@@ -20,10 +21,11 @@ from vllm.v1.kv_offload.config import OffloadingConfig
 from vllm.v1.kv_offload.cpu.common import CPUOffloadingMetrics
 from vllm.v1.kv_offload.cpu.gpu_worker import CPUOffloadingWorker
 from vllm.v1.kv_offload.cpu.manager import CPUOffloadingManager
+from vllm.v1.kv_offload.cpu.shared_offload_region import SharedOffloadRegion
 
 
 class CPUOffloadingSpec(OffloadingSpec):
-    BLOCK_SIZE_ALIGNMENT = 1
+    BLOCK_SIZE_ALIGNMENT = SharedOffloadRegion.BLOCK_SIZE_ALIGNMENT
     SUPPORTS_REPLICATED_LAYOUT = False
 
     @classmethod
@@ -143,10 +145,28 @@ class CPUOffloadingSpec(OffloadingSpec):
         return self._manager
 
     def create_worker(self, kv_caches: CanonicalKVCaches) -> CPUOffloadingWorker:
+        mmap_region: SharedOffloadRegion | None = None
+        # num_blocks == 0 would size the region to zero bytes, which cannot be
+        # mmap'd; fall back to the tensor path (empty tensors) as before.
+        if current_platform.is_cuda_alike() and self.num_blocks > 0:
+            # Back each worker's CPU buffer with a private slot in a single
+            # shared mmap region instead of a per-rank pinned tensor. Fold the
+            # global physical device index into this replica's
+            # [0, world_size) slot range.
+            world_size = self.config.parallel.world_size
+            rank = torch.accelerator.current_device_index() % world_size
+            mmap_region = SharedOffloadRegion(
+                engine_id=self.config.engine_id,
+                num_blocks=self.num_blocks,
+                rank=rank,
+                kv_bytes_per_block=self.kv_bytes_per_chunk,
+                cpu_page_size=self.cpu_page_size_per_worker,
+            )
         return CPUOffloadingWorker(
             kv_caches=kv_caches,
             blocks_per_chunk=self.blocks_per_chunk,
             num_cpu_blocks=self.num_blocks,
+            mmap_region=mmap_region,
         )
 
     @override

vllm/v1/request.py

@@ -149,7 +149,12 @@ class Request:
 
         # Used in async scheduling.
         self.num_output_placeholders = 0
-        self.async_tokens_to_discard = 0
+        # Tokens of output in flight when the request was preempted: delivered
+        # on return, but must not mutate the reset counters.
+        self.num_stale_output_tokens = 0
+        # Drop the stale output instead, for same-step preempt + resume
+        # (reset_prefix_cache).
+        self.drop_stale_output = False
 
         # Tokens of steps whose output is not yet processed (async scheduling
         # and PP run ahead of the GPU); `num_computed_tokens` counts them

vllm/v1/spec_decode/llm_base_proposer.py

@@ -1740,7 +1740,10 @@ class SpecDecodeBaseProposer:
 
         attention_groups: dict[tuple[str, str], AttentionGroup] = {}
         if kv_cache_spec is not None:
-            for layer_name in self._draft_attn_layer_names:
+            # _draft_attn_layer_names is a set; iterate in sorted order so
+            # that attention_groups (and anything derived from its first
+            # element) is deterministic across processes.
+            for layer_name in sorted(self._draft_attn_layer_names):
                 attn_backend = all_attn_layers[layer_name].get_attn_backend()
                 backend_key = attn_backend.full_cls_name()
                 if backend_key not in attention_groups:
@@ -1772,9 +1775,20 @@ class SpecDecodeBaseProposer:
                     attention_groups[backend_key].layer_names.append(layer_name)
 
         self.draft_attn_groups = list(attention_groups.values())
-        self.block_size = (
-            self.draft_attn_groups[0].get_metadata_builder().kv_cache_spec.block_size
-        )
+        if kernel_block_sizes is not None and 0 <= self.kv_cache_gid < len(
+            kernel_block_sizes
+        ):
+            # Slot mappings are computed against the block table, which is
+            # stored at kernel-block granularity. Use the kernel block size
+            # rather than the KV cache manager's block size; the two differ
+            # when manager blocks are split for the attention kernel.
+            self.block_size = kernel_block_sizes[self.kv_cache_gid]
+        else:
+            self.block_size = (
+                self.draft_attn_groups[0]
+                .get_metadata_builder()
+                .kv_cache_spec.block_size
+            )
         logger.debug("Using block size %d for drafting layers", self.block_size)
 
     def _determine_batch_execution_and_padding(

vllm/v1/worker/gpu/async_utils.py

@@ -6,7 +6,12 @@ import numpy as np
 import torch
 
 from vllm.model_executor.layers.fused_moe.all2all_utils import get_ep_all2all_manager
-from vllm.v1.outputs import AsyncModelRunnerOutput, LogprobsTensors, ModelRunnerOutput
+from vllm.v1.outputs import (
+    AsyncModelRunnerOutput,
+    LogprobsTensors,
+    ModelRunnerOutput,
+    PoolerOutput,
+)
 from vllm.v1.worker.gpu.sample.output import SamplerOutput
 
 
@@ -89,34 +94,42 @@ class AsyncPoolingOutput(AsyncModelRunnerOutput):
     def __init__(
         self,
         model_runner_output: ModelRunnerOutput,
-        pooler_output: torch.Tensor,
-        is_valid: torch.Tensor | None,
+        pooler_output: PoolerOutput,
+        finished_mask: list[bool],
         main_stream: torch.cuda.Stream,
         copy_stream: torch.cuda.Stream,
     ):
         self.model_runner_output = model_runner_output
         self.pooler_output = pooler_output
-        self.is_valid = is_valid
         # Blocking (sleep) event to avoid busy-polling the CUDA driver lock.
         self.copy_event = torch.cuda.Event(blocking=True)
 
         with stream(copy_stream, main_stream):
             copy_stream.wait_stream(main_stream)
-            self.pooler_output_cpu = self.pooler_output.to("cpu", non_blocking=True)
-            if self.is_valid is not None:
-                self.is_valid_cpu = self.is_valid.to("cpu", non_blocking=True)
+            if isinstance(self.pooler_output, torch.Tensor) and all(finished_mask):
+                self.pooler_output_cpu: PoolerOutput = self.pooler_output.to(
+                    "cpu", non_blocking=True
+                )
             else:
-                self.is_valid_cpu = None
+                outputs = (
+                    self.pooler_output.unbind()
+                    if isinstance(self.pooler_output, torch.Tensor)
+                    else self.pooler_output
+                )
+                self.pooler_output_cpu = [
+                    None
+                    if output is None or not is_finished
+                    else output.to("cpu", non_blocking=True)
+                    for output, is_finished in zip(outputs, finished_mask, strict=True)
+                ]
             self.copy_event.record(copy_stream)
 
     def get_output(self) -> ModelRunnerOutput:
-        pooler_output = list(self.pooler_output_cpu.unbind(dim=0))
+        if isinstance(self.pooler_output_cpu, torch.Tensor):
+            pooler_output = list(self.pooler_output_cpu.unbind(dim=0))
+        else:
+            pooler_output = self.pooler_output_cpu
         self.copy_event.synchronize()
-        if self.is_valid_cpu is not None:
-            is_valid_cpu = self.is_valid_cpu.tolist()
-            for i, is_valid in enumerate(is_valid_cpu):
-                if not is_valid:
-                    pooler_output[i] = None
         self.model_runner_output.pooler_output = pooler_output
         return self.model_runner_output
 

vllm/v1/worker/gpu/model_runner.py

@@ -366,7 +366,7 @@ class GPUModelRunner(LoRAModelRunnerMixin):
             )
 
         if self.is_pooling_model and self.is_last_pp_rank:
-            self.pooling_runner = PoolingRunner(self.model)
+            self.pooling_runner = PoolingRunner(self.model, self.vllm_config)
         eplb_models_added |= self.eplb.maybe_register_model(
             self.model,
             self.model_config,
@@ -787,6 +787,8 @@ class GPUModelRunner(LoRAModelRunnerMixin):
         req_idx = self.req_states.remove_request(req_id)
         if req_idx is None:
             return False
+        if self.pooling_runner is not None:
+            self.pooling_runner.remove_request(req_idx)
         if self.pp_handler is not None:
             self.pp_handler.on_req_idx_freed(req_idx)
         if self.encoder_cache is not None:
@@ -839,6 +841,14 @@ class GPUModelRunner(LoRAModelRunnerMixin):
             )
             req_index = self.req_states.req_id_to_index[req_id]
 
+            if self.pooling_runner is not None:
+                assert new_req_data.pooling_params is not None
+                self.pooling_runner.add_request(
+                    req_index,
+                    new_req_data.pooling_params,
+                    new_req_data.prompt_token_ids,
+                )
+
             if self.encoder_cache is not None:
                 self.encoder_cache.add_request(req_id, new_req_data.mm_features)
 
@@ -1591,7 +1601,7 @@ class GPUModelRunner(LoRAModelRunnerMixin):
             return ModelRunnerOutput.with_kv_conn_output_only(kv_connector_output)
 
         assert self.pooling_runner is not None
-        pooler_output, is_valid = self.pooling_runner.pool(
+        pooler_output, finished_mask = self.pooling_runner.pool(
             hidden_states, input_batch, self.req_states
         )
 
@@ -1604,7 +1614,7 @@ class GPUModelRunner(LoRAModelRunnerMixin):
         async_output = AsyncPoolingOutput(
             model_runner_output=model_runner_output,
             pooler_output=pooler_output,
-            is_valid=is_valid,
+            finished_mask=finished_mask,
             main_stream=self.main_stream,
             copy_stream=self.output_copy_stream,
         )

vllm/v1/worker/gpu/model_states/encoder_only.py

@@ -8,12 +8,13 @@ import torch.nn as nn
 from vllm.config import VllmConfig, get_layers_from_vllm_config
 from vllm.config.compilation import CUDAGraphMode
 from vllm.model_executor.layers.attention import Attention
-from vllm.utils.torch_utils import STR_DTYPE_TO_TORCH_DTYPE
+from vllm.utils.torch_utils import PIN_MEMORY, STR_DTYPE_TO_TORCH_DTYPE
 from vllm.v1.attention.backend import (
     AttentionCGSupport,
     AttentionType,
     CommonAttentionMetadata,
 )
+from vllm.v1.core.sched.output import NewRequestData
 from vllm.v1.kv_cache_interface import (
     AttentionSpec,
     EncoderOnlyAttentionSpec,
@@ -22,6 +23,7 @@ from vllm.v1.kv_cache_interface import (
 from vllm.v1.worker.gpu.input_batch import InputBatch
 from vllm.v1.worker.gpu.mm.encoder_cache import EncoderCache
 from vllm.v1.worker.gpu.model_states.default import DefaultModelState
+from vllm.v1.worker.gpu.states import RequestState
 from vllm.v1.worker.utils import AttentionGroup
 
 
@@ -95,6 +97,52 @@ class EncoderOnlyModelState(DefaultModelState):
         self._dummy_slot_mapping = torch.zeros(
             self.max_num_tokens, dtype=torch.int64, device=device
         )
+        self.token_type_ids: dict[str, torch.Tensor] = {}
+
+    def add_request(self, req_index: int, new_req_data: NewRequestData) -> None:
+        super().add_request(req_index, new_req_data)
+        pooling_params = new_req_data.pooling_params
+        if pooling_params is None or pooling_params.extra_kwargs is None:
+            return
+
+        token_type_start = pooling_params.extra_kwargs.get("compressed_token_type_ids")
+        if token_type_start is not None:
+            assert new_req_data.prompt_token_ids is not None
+            self.token_type_ids[new_req_data.req_id] = (
+                torch.arange(len(new_req_data.prompt_token_ids), dtype=torch.int32)
+                >= token_type_start
+            ).to(torch.int32)
+
+    def remove_request(self, req_id: str) -> None:
+        super().remove_request(req_id)
+        self.token_type_ids.pop(req_id, None)
+
+    def prepare_inputs(
+        self, input_batch: InputBatch, req_states: RequestState
+    ) -> dict[str, torch.Tensor | None]:
+        model_inputs = super().prepare_inputs(input_batch, req_states)
+        if not self.token_type_ids:
+            return model_inputs
+
+        token_type_ids_cpu = torch.zeros(
+            input_batch.num_tokens_after_padding,
+            dtype=torch.int32,
+            pin_memory=PIN_MEMORY,
+        )
+        offset = 0
+        for i, req_id in enumerate(input_batch.req_ids):
+            num_tokens = int(input_batch.num_scheduled_tokens[i])
+            request_token_type_ids = self.token_type_ids.get(req_id)
+            if request_token_type_ids is not None:
+                start = int(input_batch.num_computed_tokens_np[i])
+                token_type_ids_cpu[offset : offset + num_tokens].copy_(
+                    request_token_type_ids[start : start + num_tokens]
+                )
+            offset += num_tokens
+        model_inputs["token_type_ids"] = token_type_ids_cpu.to(
+            self.device, non_blocking=True
+        )
+        return model_inputs
 
     def get_additional_cg_support(self) -> tuple[AttentionCGSupport, str | None]:
         # Encoder groups are built here rather than in init_attn_backend, so

vllm/v1/worker/gpu/model_states/mamba_hybrid.py

@@ -9,10 +9,6 @@ import torch.nn as nn
 
 from vllm.config import VllmConfig
 from vllm.config.compilation import CUDAGraphMode
-from vllm.model_executor.layers.mamba.mamba_utils import (
-    get_conv_copy_spec,
-    is_conv_state_dim_first,
-)
 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
@@ -133,14 +129,11 @@ class MambaHybridModelState(DefaultModelState):
     ) -> MambaSpecDecodeGPUContext:
         if self._mamba_ctx is None:
             copy_funcs = self.model.get_mamba_state_copy_func()
-            # The fused copy kernels shift conv windows assuming the SD layout;
-            # the DS layout cannot express a >0 spec-decode shift as a single
-            # contiguous copy (mirrors get_conv_copy_spec's NotImplementedError).
-            if get_conv_copy_spec in copy_funcs and is_conv_state_dim_first():
-                assert self.vllm_config.speculative_config is None, (
-                    "DS conv state layout does not support mamba align state "
-                    "copies with speculative decoding"
-                )
+            # Both SD and DS conv layouts support a >0 spec-decode shift: the
+            # fused pre-copy kernel (``_copy_mamba_state_block``) applies the
+            # ``token_bias = num_accepted - 1`` window shift per conv layout
+            # (SD: contiguous slice; DS: per-dim-row strided slice), matching
+            # the V1 ``get_conv_copy_spec`` semantics.
             self._mamba_ctx = MambaSpecDecodeGPUContext.create(
                 max_num_reqs=self.max_num_reqs,
                 kv_cache_config=kv_cache_config,

vllm/v1/worker/gpu/pool/pooling_runner.py

@@ -2,45 +2,190 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 from typing import cast
 
+import numpy as np
 import torch
 import torch.nn as nn
-import torch.nn.functional as F
 
+from vllm.config import VllmConfig
 from vllm.model_executor.models import VllmModelForPooling, is_pooling_model
+from vllm.pooling_params import PoolingParams
 from vllm.tasks import PoolingTask
+from vllm.utils.torch_utils import PIN_MEMORY
+from vllm.v1.outputs import PoolerOutput
+from vllm.v1.pool.metadata import PoolingMetadata, PoolingStates
 from vllm.v1.worker.gpu.input_batch import InputBatch
 from vllm.v1.worker.gpu.states import RequestState
 
+_SUPPORTED_TASKS: frozenset[PoolingTask] = frozenset({"embed", "classify"})
+
 
-# NOTE(woosuk): Currently, this class only supports the "LAST" pooling task
-# on decoder-only models. How to support other pooling tasks and models
-# is to be determined.
 class PoolingRunner:
-    def __init__(self, model: nn.Module):
+    def __init__(self, model: nn.Module, vllm_config: VllmConfig):
         self.model = cast(VllmModelForPooling, model)
+        self.model_config = vllm_config.model_config
+        self.max_num_reqs = vllm_config.scheduler_config.max_num_seqs
+        model_tasks = tuple(sorted(self.model.pooler.get_supported_tasks()))
+        selected_task = self.model_config.get_pooling_task(model_tasks)
+        if selected_task not in _SUPPORTED_TASKS:
+            raise ValueError(
+                "Model Runner V2 supports only sequence-level pooling tasks "
+                f"{sorted(_SUPPORTED_TASKS)}, but this model selects "
+                f"{selected_task!r} from {list(model_tasks)}. Set an explicitly "
+                "supported task or VLLM_USE_V2_MODEL_RUNNER=0."
+            )
+        self.supported_tasks = frozenset(self.get_supported_tasks(model))
+        if not self.supported_tasks:
+            raise ValueError(
+                "Model Runner V2 supports only sequence-level pooling tasks "
+                f"{sorted(_SUPPORTED_TASKS)}, but this model supports "
+                f"{list(model_tasks)}. "
+                "Set VLLM_USE_V2_MODEL_RUNNER=0 to use this model."
+            )
+        self.pooling_params: dict[int, PoolingParams] = {}
+        self.pooling_states: dict[int, PoolingStates] = {}
+        self.prompt_token_ids: dict[int, torch.Tensor] = {}
 
     @staticmethod
     def get_supported_tasks(model: nn.Module) -> list[PoolingTask]:
         if not is_pooling_model(model):
             return []
-        assert "embed" in model.pooler.get_supported_tasks()
-        return ["embed"]
+        return sorted(model.pooler.get_supported_tasks() & _SUPPORTED_TASKS)
+
+    def add_request(
+        self,
+        req_index: int,
+        pooling_params: PoolingParams,
+        prompt_token_ids: list[int],
+    ) -> None:
+        task = pooling_params.task
+        if task not in self.supported_tasks:
+            raise ValueError(
+                f"Unsupported task: {task!r}. "
+                f"Supported tasks: {sorted(self.supported_tasks)}"
+            )
+        self.model.pooler.get_pooling_updates(task).apply(pooling_params)
+        self.pooling_params[req_index] = pooling_params
+        self.pooling_states[req_index] = PoolingStates()
+        if pooling_params.requires_token_ids:
+            self.prompt_token_ids[req_index] = torch.tensor(
+                prompt_token_ids, dtype=torch.int64
+            )
+
+    def remove_request(self, req_index: int) -> None:
+        self.pooling_params.pop(req_index, None)
+        if state := self.pooling_states.pop(req_index, None):
+            state.clean()
+        self.prompt_token_ids.pop(req_index, None)
+
+    def _get_pooling_metadata(
+        self,
+        input_batch: InputBatch,
+        req_states: RequestState,
+        device: torch.device,
+    ) -> PoolingMetadata:
+        req_indices = input_batch.idx_mapping_np.tolist()
+        pooling_params = [self.pooling_params[i] for i in req_indices]
+        pooling_states = [self.pooling_states[i] for i in req_indices]
+        prompt_lens = torch.from_numpy(
+            req_states.prompt_len.np[input_batch.idx_mapping_np].copy()
+        )
+
+        prompt_token_ids_cpu = None
+        prompt_token_ids = None
+        if any(params.requires_token_ids for params in pooling_params):
+            max_prompt_len = int(prompt_lens.max())
+            prompt_token_ids_cpu = torch.zeros(
+                (input_batch.num_reqs, max_prompt_len),
+                dtype=torch.int64,
+                pin_memory=PIN_MEMORY,
+            )
+            for i, (req_index, params) in enumerate(zip(req_indices, pooling_params)):
+                if not params.requires_token_ids:
+                    continue
+                token_ids = self.prompt_token_ids[req_index]
+                prompt_token_ids_cpu[i, : token_ids.numel()] = token_ids
+            prompt_token_ids = prompt_token_ids_cpu.to(device, non_blocking=True)
+
+        return PoolingMetadata(
+            prompt_lens=prompt_lens,
+            prompt_token_ids=prompt_token_ids,
+            prompt_token_ids_cpu=prompt_token_ids_cpu,
+            pooling_params=pooling_params,
+            pooling_states=pooling_states,
+        )
 
     def pool(
         self,
         hidden_states: torch.Tensor,
         input_batch: InputBatch,
         req_states: RequestState,
-    ) -> tuple[torch.Tensor, torch.Tensor | None]:
-        # TODO(woosuk): Support different types of pooling tasks.
-        last_hidden_states = hidden_states[input_batch.logits_indices]
-        # TODO(woosuk): Make normalization optional.
-        last_hidden_states = F.normalize(last_hidden_states, p=2, dim=-1)
+    ) -> tuple[PoolerOutput, list[bool]]:
+        hidden_states = hidden_states[: input_batch.num_tokens]
+        pooling_metadata = self._get_pooling_metadata(
+            input_batch, req_states, hidden_states.device
+        )
+        num_reqs = input_batch.num_reqs
+        # Pooling has no speculative tokens, so this CPU upper bound is exact.
+        seq_lens_cpu = input_batch.seq_lens_cpu_upper_bound[:num_reqs]
+        pooling_metadata.build_pooling_cursor(
+            input_batch.num_scheduled_tokens,
+            seq_lens_cpu,
+            device=hidden_states.device,
+            query_start_loc_gpu=input_batch.query_start_loc[: num_reqs + 1],
+        )
+        pooler_output = self.model.pooler(hidden_states, pooling_metadata)
 
-        prompt_len = req_states.prompt_len.gpu[input_batch.idx_mapping]
-        is_valid = input_batch.seq_lens == prompt_len
-        return last_hidden_states, is_valid
+        finished_mask = pooling_metadata.get_pooling_cursor().is_finished().tolist()
+        return pooler_output, finished_mask
+
+    def _dummy_pooler_run_task(
+        self, hidden_states: torch.Tensor, task: PoolingTask
+    ) -> PoolerOutput:
+        num_tokens = hidden_states.shape[0]
+        num_reqs = min(num_tokens, self.max_num_reqs)
+        base_tokens = num_tokens // num_reqs
+        num_extra = num_tokens % num_reqs
+        num_scheduled_tokens = np.full(num_reqs, base_tokens, dtype=np.int32)
+        if num_extra > 0:
+            num_scheduled_tokens[-num_extra:] += 1
+        prompt_lens = torch.from_numpy(num_scheduled_tokens)
+
+        pooling_params = PoolingParams(task=task)
+        pooling_params.verify(self.model_config)
+        self.model.pooler.get_pooling_updates(task).apply(pooling_params)
+        prompt_token_ids = None
+        if pooling_params.requires_token_ids:
+            prompt_token_ids = torch.zeros(
+                (num_reqs, int(prompt_lens.max())),
+                dtype=torch.int64,
+                device=hidden_states.device,
+            )
+        pooling_metadata = PoolingMetadata(
+            prompt_lens=prompt_lens,
+            prompt_token_ids=prompt_token_ids,
+            prompt_token_ids_cpu=None
+            if prompt_token_ids is None
+            else prompt_token_ids.cpu(),
+            pooling_params=[pooling_params] * num_reqs,
+            pooling_states=[PoolingStates() for _ in range(num_reqs)],
+        )
+        pooling_metadata.build_pooling_cursor(
+            num_scheduled_tokens,
+            seq_lens_cpu=prompt_lens,
+            device=hidden_states.device,
+        )
+        try:
+            return self.model.pooler(hidden_states, pooling_metadata)
+        except RuntimeError as e:
+            if "out of memory" not in str(e):
+                raise
+            raise RuntimeError(
+                "CUDA out of memory occurred when warming up pooler "
+                f"({task=}) with {num_reqs} dummy requests. Please try "
+                "lowering `max_num_seqs` or `gpu_memory_utilization` when "
+                "initializing the engine."
+            ) from e
 
     def dummy_pooler_run(self, hidden_states: torch.Tensor) -> None:
-        F.normalize(hidden_states, p=2, dim=-1)
-        return
+        for task in sorted(self.supported_tasks):
+            self._dummy_pooler_run_task(hidden_states, task)

vllm/v1/worker/gpu/warmup.py

@@ -230,7 +230,11 @@ def warmup_kernels(
     # SamplingParams exercising all sampling features.
     if model_runner.is_pooling_model:
         sampling_params = None
-        pooling_params = PoolingParams()
+        pooling_task = model_runner.model_config.get_pooling_task(
+            model_runner.get_supported_tasks()
+        )
+        pooling_params = PoolingParams(task=pooling_task)
+        pooling_params.verify(model_runner.model_config)
     else:
         sampling_params = SamplingParams.for_sampler_warmup()
         pooling_params = None

vllm/v1/worker/gpu_model_runner.py

@@ -4335,6 +4335,7 @@ class GPUModelRunner(
                     self.compilation_config.static_forward_context,
                     self.model.get_mamba_state_copy_func(),
                     mamba_bufs.preprocess,
+                    align_ctx=mamba_bufs.postprocess_align,
                 )
                 # preprocess_mamba resets num_accepted_tokens_cpu to 1
                 # for requests whose state was copied to a new block.

vllm/v1/worker/mamba_utils.py

@@ -3,7 +3,7 @@
 import dataclasses
 import itertools
 from collections.abc import Callable
-from typing import Any
+from typing import Any, NamedTuple
 
 import torch
 
@@ -348,8 +348,9 @@ def precopy_mamba_align_fused_kernel(
     num_reqs,
     COPY_BLOCK_SIZE: tl.constexpr,
     CONV_STATE_DIM_FIRST: tl.constexpr,
+    HAS_IDX_MAPPING: tl.constexpr = True,
 ):
-    """Pre-copy mamba "align" state across block boundaries on the V2 runner.
+    """Pre-copy mamba "align" state across block boundaries.
 
     Before the forward pass, copy each request's last SSM/conv state from its
     previous block column into the new window block column, so the kernels read
@@ -359,16 +360,20 @@ def precopy_mamba_align_fused_kernel(
     copy specs), but driven by the GPU-resident src columns so it needs no
     CPU-GPU sync (async-scheduling safe).
 
-    Grid: (num_reqs, num_layers * num_state_types); block tables are indexed by
-    batch row, per-request state by req_idx via idx_mapping (V2 layout).
+    Grid: (num_reqs, num_layers * num_state_types). V2 passes a batch-to-state
+    idx_mapping; V1 already stores the staged arrays in batch order and uses
+    HAS_IDX_MAPPING=False.
     """
     batch_idx = tl.program_id(0)
     state_idx = tl.program_id(1)
     if batch_idx >= num_reqs:
         return
-    req_idx = tl.load(idx_mapping_ptr + batch_idx)
-    if req_idx < 0:
-        return
+    if HAS_IDX_MAPPING:
+        req_idx = tl.load(idx_mapping_ptr + batch_idx)
+        if req_idx < 0:
+            return
+    else:
+        req_idx = batch_idx
 
     src_col = tl.load(src_col_ptr + req_idx)
     dst_col = tl.load(mamba_state_idx_ptr + req_idx)
@@ -527,6 +532,8 @@ class MambaSpecDecodeGPUContext:
     num_scheduled_tokens_buf: CpuGpuBuffer | None = None
     num_computed_tokens_buf: CpuGpuBuffer | None = None
     num_draft_tokens_buf: CpuGpuBuffer | None = None
+    precopy_src_col_buf: CpuGpuBuffer | None = None
+    precopy_token_bias_buf: CpuGpuBuffer | None = None
 
     # Flag to track if metadata has been populated
     is_initialized: bool = False
@@ -590,6 +597,8 @@ class MambaSpecDecodeGPUContext:
             num_scheduled_tokens_buf=make_buffer(max_num_reqs, dtype=torch.int32),
             num_computed_tokens_buf=make_buffer(max_num_reqs, dtype=torch.int32),
             num_draft_tokens_buf=make_buffer(max_num_reqs, dtype=torch.int32),
+            precopy_src_col_buf=make_buffer(max_num_reqs, dtype=torch.int32),
+            precopy_token_bias_buf=make_buffer(max_num_reqs, dtype=torch.int32),
             is_initialized=False,
         )
 
@@ -797,17 +806,18 @@ class MambaSpecDecodeGPUContext:
         state_idx_gpu: torch.Tensor,
         src_col_gpu: torch.Tensor,
         token_bias_gpu: torch.Tensor,
-        idx_mapping: torch.Tensor,
+        idx_mapping: torch.Tensor | None,
     ) -> None:
         """Pre-copy each request's previous running block into its new window
-        block before the forward pass (V2 align boundary migration).
+        block before the forward pass (align boundary migration).
 
         Args:
             num_reqs: Number of active requests (batch order).
             state_idx_gpu: [max_reqs] post-advance dst block column per req slot.
             src_col_gpu: [max_reqs] pre-advance src block column (-1 = fresh).
             token_bias_gpu: [max_reqs] accepted-token bias (num_accepted - 1).
-            idx_mapping: [num_reqs] batch_idx -> req_state_idx (-1 to skip).
+            idx_mapping: optional [num_reqs] batch_idx -> req_state_idx.
+                None means V1 batch order already equals request state order.
         """
         if num_reqs == 0 or not self.is_initialized:
             return
@@ -831,6 +841,7 @@ class MambaSpecDecodeGPUContext:
             num_reqs,
             COPY_BLOCK_SIZE=1024,
             CONV_STATE_DIM_FIRST=is_conv_state_dim_first(),
+            HAS_IDX_MAPPING=idx_mapping is not None,
         )
 
     def run_fused_postprocess_align(
@@ -989,6 +1000,36 @@ def cleanup_mamba_state_idx(
         mamba_state_idx.pop(req_id, None)
 
 
+class _FusedPrecopy(NamedTuple):
+    """Resolved fused align pre-copy resources (all non-None once resolved)."""
+
+    ctx: "MambaSpecDecodeGPUContext"
+    state_idx: CpuGpuBuffer
+    src_col: CpuGpuBuffer
+    token_bias: CpuGpuBuffer
+
+
+def _resolve_fused_precopy(
+    align_ctx: "MambaSpecDecodeGPUContext | None",
+) -> _FusedPrecopy | None:
+    """Bundle the fused-path buffers, or None for the scalar path.
+
+    Returning one non-None bundle lets callers narrow all four members with a
+    single ``is not None`` check instead of re-asserting each buffer per use.
+    """
+    if align_ctx is None:
+        return None
+    assert align_ctx.mamba_state_idx_buf is not None
+    assert align_ctx.precopy_src_col_buf is not None
+    assert align_ctx.precopy_token_bias_buf is not None
+    return _FusedPrecopy(
+        align_ctx,
+        align_ctx.mamba_state_idx_buf,
+        align_ctx.precopy_src_col_buf,
+        align_ctx.precopy_token_bias_buf,
+    )
+
+
 def preprocess_mamba(
     scheduler_output: SchedulerOutput,
     kv_cache_config: KVCacheConfig,
@@ -999,11 +1040,13 @@ def preprocess_mamba(
     forward_context: dict[str, Any],
     mamba_state_copy_funcs: tuple[MambaStateCopyFunc, ...],
     copy_bufs: MambaCopyBuffers,
+    align_ctx: MambaSpecDecodeGPUContext | None = None,
 ):
     """
     Copy the mamba state of previous step to the last
     (1 + num_speculative_blocks) block.
     """
+    fused = _resolve_fused_precopy(align_ctx)
     mamba_group_ids = copy_bufs.mamba_group_ids
     mamba_spec = copy_bufs.mamba_spec
     num_speculative_blocks = mamba_spec.num_speculative_blocks
@@ -1013,20 +1056,37 @@ def preprocess_mamba(
     cleanup_mamba_state_idx(scheduler_output, mamba_state_idx)
 
     copy_bufs.offset = 0
+    num_reqs = len(input_batch.req_ids)
+
+    if fused is not None:
+        if num_reqs == 0:
+            return
+        if not fused.ctx.is_initialized:
+            fused.ctx.initialize_from_forward_context(
+                kv_cache_config,
+                forward_context,
+                mamba_state_copy_funcs,
+                [
+                    input_batch.block_table[gid].get_device_tensor(num_reqs)
+                    for gid in fused.ctx.mamba_group_ids
+                ],
+            )
+
+        fused.src_col.np[:num_reqs] = -1
+        fused.token_bias.np[:num_reqs] = 0
+
     for i, req_id in enumerate(input_batch.req_ids):
         req_state = requests[req_id]
         prev_state_idx = mamba_state_idx.get(req_id)
         if prev_state_idx is None:
-            # new / resumed request, no previous state
-            # if num_computed_tokens is 0, prev_state_idx will be -1
+            # New / resumed request; num_computed_tokens == 0 gives -1.
             prev_state_idx = (req_state.num_computed_tokens - 1) // block_size
 
         num_scheduled_tokens = scheduler_output.num_scheduled_tokens[req_id]
-        num_blocks: int = (
+        num_blocks = (
             cdiv(req_state.num_computed_tokens + num_scheduled_tokens, block_size)
             + num_speculative_blocks
         )
-
         # We always save the current running state at the last
         # (1 + num_speculative_blocks) block.
         # A corner case worth mention here: assume we have block_size = 4 and
@@ -1039,20 +1099,42 @@ def preprocess_mamba(
         # And use block 1 to save the running state.
         curr_state_idx = num_blocks - 1 - num_speculative_blocks
         mamba_state_idx[req_id] = curr_state_idx
+        if fused is not None:
+            fused.state_idx.np[i] = curr_state_idx
+
         if prev_state_idx != -1 and prev_state_idx != curr_state_idx:
-            collect_mamba_copy_meta(
-                copy_bufs,
-                kv_cache_config,
-                mamba_state_copy_funcs,
-                mamba_group_ids,
-                prev_state_idx,
-                curr_state_idx,
-                input_batch.num_accepted_tokens_cpu[i] - 1,
-                req_state,
-                forward_context,
-            )
+            accept_token_bias = int(input_batch.num_accepted_tokens_cpu[i]) - 1
+            if fused is not None:
+                assert accept_token_bias >= 0
+                fused.src_col.np[i] = prev_state_idx
+                fused.token_bias.np[i] = accept_token_bias
+            else:
+                collect_mamba_copy_meta(
+                    copy_bufs,
+                    kv_cache_config,
+                    mamba_state_copy_funcs,
+                    mamba_group_ids,
+                    prev_state_idx,
+                    curr_state_idx,
+                    accept_token_bias,
+                    req_state,
+                    forward_context,
+                )
             input_batch.num_accepted_tokens_cpu[i] = 1
-    do_mamba_copy_block(copy_bufs)
+
+    if fused is not None:
+        fused.state_idx.copy_to_gpu(num_reqs)
+        fused.src_col.copy_to_gpu(num_reqs)
+        fused.token_bias.copy_to_gpu(num_reqs)
+        fused.ctx.run_fused_precopy(
+            num_reqs=num_reqs,
+            state_idx_gpu=fused.state_idx.gpu,
+            src_col_gpu=fused.src_col.gpu,
+            token_bias_gpu=fused.token_bias.gpu,
+            idx_mapping=None,
+        )
+    else:
+        do_mamba_copy_block(copy_bufs)
 
 
 def postprocess_mamba_all(