[Model] Add LongCat-2.0 (sparse attention, MTP, n-gram fixes)

Support meituan-longcat/LongCat-2.0-FP8: the LongCat-Flash-Lite backbone plus a DeepSeek-V3.2-style sparse attention indexer, reusing vLLM's existing DSA infrastructure. - Resolve the checkpoint's null model_type via a registered config class (arch LongcatCausalLM); alias the oe_* n-gram config fields. - Wire the indexer into the dual-attention layer: only the first attention of each layer computes top-k indices (cli_factor), the second reuses them via the shared buffer. Adds a generic skip_topk override to DeepseekV2MLAAttention; the schedule lives in the model. - Streaming-aware indexing: force the first index_init_tokens and last index_local_tokens into the top-k set (sparse_attn_indexer). - MTP speculative decoding (one module iterated up to 3 draft steps, plain token embedding, FP8 indexer weights). - Fix a latent LongCat weight-loading bug: the post-load mla_scale_q_lora/kv_lora fold breaks under incremental load_weights calls; fold at weight-load time instead. - Fix two n-gram embedding bugs (also affect LongCat-Flash-Lite): rebuild the left-context from the accepted-token history so rejected draft tokens don't pollute it, and hash EOS-current positions with full look-back per the HF reference (unit test included). - Add --reasoning-parser longcat (<longcat_think> tags) and docs rows. - Add longcat model types to the DeepGEMM Blackwell blocklist (non-ue8m0 scales; measured GSM8K-neutral on vLLM but matches SGLang's guard). Co-authored-by: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: mgoin <mgoin64@gmail.com>
[ROCm] Enable DeepSeek-V4 DSpark speculative decoding on AMD (MI350X / MI355X, gfx950) (#47419 )
2026-07-11 02:12:49 +00:00 · 2026-07-10 23:22:19 +08:00 · 2026-07-10 07:17:50 -07:00 · 2026-07-10 10:12:50 -04:00 · 2026-07-10 14:08:20 +00:00 · 2026-07-10 06:55:28 -07:00
144 changed files with 5958 additions and 1736 deletions
@@ -177,6 +177,18 @@ BRANCH=$4
 IMAGE_TAG=$5
 IMAGE_TAG_LATEST=${6:-} # only used for main branch, optional

+# When TORCH_NIGHTLY=1, build the base CI image against PyTorch nightly so the
+# entire existing pipeline runs on nightly torch (CUDA/GPU lane only). Delegate
+# to the dedicated nightly build (PYTORCH_NIGHTLY=1, CUDA 13.0) and tag it at the
+# normal IMAGE_TAG that every test step already pulls -- no separate image tag,
+# no duplicate "vLLM Against PyTorch Nightly" pipeline section.
+if [[ "${TORCH_NIGHTLY:-0}" == "1" ]]; then
+    echo "--- :warning: TORCH_NIGHTLY=1 -- building base image on PyTorch nightly"
+    SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+    exec "${SCRIPT_DIR}/image_build_torch_nightly.sh" \
+        "${REGISTRY}" "${REPO}" "${BUILDKITE_COMMIT}" "${BRANCH}" "${IMAGE_TAG}"
+fi
+
 # build config
 TARGET="test-ci"
 VLLM_BAKE_FILE_PATH="${VLLM_BAKE_FILE_PATH:-docker/docker-bake.hcl}"
@@ -19,13 +19,14 @@ if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
  echo "Image found"
 else
  echo "Image not found, proceeding with build..."
-  # build for arm64 GPU targets: Grace/GH200 (sm_90) and DGX Spark/GB10
+  # build for arm64 GPU targets: Grace/GH200 (sm_90),
+  # Blackwell/Thor (sm_100/sm_103/sm_110), and DGX Spark/GB10
  # (sm_121, family-covered by 12.0 under CUDA 13)
  docker build --file docker/Dockerfile \
    --platform linux/arm64 \
    --build-arg max_jobs=16 \
    --build-arg nvcc_threads=4 \
-    --build-arg torch_cuda_arch_list="9.0 12.0" \
+    --build-arg torch_cuda_arch_list="9.0 10.0 11.0 12.0" \
    --build-arg USE_SCCACHE=1 \
    --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
    --tag "$IMAGE" \
@@ -72,9 +72,7 @@ steps:
      pytest -v -s v1/test_oracle.py &&
      pytest -v -s v1/test_request.py &&
      pytest -v -s v1/test_outputs.py &&
-      pytest -v -s v1/sample/test_topk_topp_sampler.py &&
-      pytest -v -s v1/sample/test_logprobs.py &&
-      pytest -v -s v1/sample/test_logprobs_e2e.py'
+      pytest -v -s v1/sample'

 - label: Basic Models Tests (Initialization)
  timeout_in_minutes: 60
@@ -848,6 +848,23 @@ steps:
      DOCKER_BUILDKIT: "1"
      S3_BUCKET: "vllm-wheels"

+  - label: "Publish nightly XPU image to DockerHub"
+    depends_on:
+      - create-manifest-xpu
+    if: build.env("NIGHTLY") == "1"
+    agents:
+      queue: small_cpu_queue_release
+    commands:
+      - "bash .buildkite/scripts/xpu/push-nightly-builds-xpu.sh"
+      - "bash .buildkite/scripts/cleanup-nightly-builds.sh nightly- vllm/vllm-openai-xpu"
+    plugins:
+      - docker-login#v3.0.0:
+          username: vllmbot
+          password-env: DOCKERHUB_TOKEN
+    env:
+      DOCKER_BUILDKIT: "1"
+      DOCKERHUB_USERNAME: "vllmbot"
+
  - label: "Publish nightly ROCm image to DockerHub"
    depends_on:
      - build-rocm-release-image
@@ -878,6 +895,7 @@ steps:
      - create-multi-arch-manifest-cuda-12-9
      - create-multi-arch-manifest-ubuntu2404
      - create-multi-arch-manifest-cuda-12-9-ubuntu2404
+      - create-manifest-xpu
      - build-rocm-release-image
      - input-release-version
      # Wait for CPU builds if their block steps were unblocked, so publish
@@ -29,7 +29,11 @@ if python3 -c "import torch; assert torch.version.hip" 2>/dev/null; then
        TORCH_INDEX_URL=""
    fi
 else
-    TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
+    if [ "${TORCH_NIGHTLY:-0}" = "1" ]; then
+        TORCH_INDEX_URL="https://download.pytorch.org/whl/nightly/cu130"
+    else
+        TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
+    fi
 fi
 echo ">>> Using PyTorch index: ${TORCH_INDEX_URL:-PyPI default}"

@@ -130,6 +130,22 @@ docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm
 docker push vllm/vllm-openai-rocm:latest-base
 docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base

+# ---- XPU ----
+
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu vllm/vllm-openai-xpu:latest-x86_64
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64
+docker push vllm/vllm-openai-xpu:latest-x86_64
+docker push vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64
+
+docker manifest rm vllm/vllm-openai-xpu:latest || true
+docker manifest rm vllm/vllm-openai-xpu:v${RELEASE_VERSION} || true
+docker manifest create vllm/vllm-openai-xpu:latest vllm/vllm-openai-xpu:latest-x86_64 --amend
+docker manifest create vllm/vllm-openai-xpu:v${RELEASE_VERSION} vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64 --amend
+docker manifest push vllm/vllm-openai-xpu:latest
+docker manifest push vllm/vllm-openai-xpu:v${RELEASE_VERSION}
+
 # ---- CPU ----
 # CPU images are behind separate block steps and may not have been built.
 # All-or-nothing: inspect both arches first, then either publish everything
@@ -23,6 +23,7 @@ NC='\033[0m' # No Color
 # Default configuration
 PIPELINE="ci"
 DRY_RUN=true
+TORCH_NIGHTLY=false

 usage() {
    cat <<EOF
@@ -34,12 +35,14 @@ Sets RUN_ALL=1 and NIGHTLY=1 environment variables.
 SAFETY: Dry-run by default. Use --execute to actually trigger a build.

 Options:
-    --execute       Actually trigger the build (default: dry-run)
-    --pipeline      Buildkite pipeline slug (default: ${PIPELINE})
-    --commit        Override commit SHA (default: current HEAD)
-    --branch        Override branch name (default: current branch)
-    --message       Custom build message (default: auto-generated)
-    --help          Show this help message
+    --execute        Actually trigger the build (default: dry-run)
+    --pipeline       Buildkite pipeline slug (default: ${PIPELINE})
+    --commit         Override commit SHA (default: current HEAD)
+    --branch         Override branch name (default: current branch)
+    --message        Custom build message (default: auto-generated)
+    --torch-nightly  Also build and run the full suite against torch nightly
+                     (sets TORCH_NIGHTLY=1)
+    --help           Show this help message

 Prerequisites:
    - bk CLI installed: brew tap buildkite/buildkite && brew install buildkite/buildkite/bk
@@ -49,6 +52,7 @@ Examples:
    $(basename "$0")                        # Dry-run, show what would happen
    $(basename "$0") --execute              # Actually trigger the build
    $(basename "$0") --pipeline ci-shadow   # Dry-run with different pipeline
+    $(basename "$0") --torch-nightly        # Dry-run a full torch-nightly run
 EOF
    exit 1
 }
@@ -96,6 +100,10 @@ while [[ $# -gt 0 ]]; do
            MESSAGE="$2"
            shift 2
            ;;
+        --torch-nightly)
+            TORCH_NIGHTLY=true
+            shift
+            ;;
        --help|-h)
            usage
            ;;
@@ -171,11 +179,17 @@ if [[ $(echo "$REMOTE_BRANCHES" | wc -l) -gt 5 ]]; then
 fi
 echo ""

+# Environment variables passed to the build.
+BUILD_ENV=("RUN_ALL=1" "NIGHTLY=1")
+if [[ "$TORCH_NIGHTLY" == true ]]; then
+    BUILD_ENV+=("TORCH_NIGHTLY=1")
+fi
+
 log_info "Pipeline: ${PIPELINE}"
 log_info "Branch: ${BRANCH}"
 log_info "Commit: ${COMMIT}"
 log_info "Message: ${MESSAGE}"
-log_info "Environment: RUN_ALL=1, NIGHTLY=1"
+log_info "Environment: ${BUILD_ENV[*]}"
 echo ""

 # Build the command
@@ -187,9 +201,10 @@ CMD=(bk build create
    --commit "${COMMIT}"
    --branch "${BRANCH}"
    --message "${MESSAGE}"
-    --env "RUN_ALL=1"
-    --env "NIGHTLY=1"
 )
+for env_var in "${BUILD_ENV[@]}"; do
+    CMD+=(--env "${env_var}")
+done

 if [[ "$DRY_RUN" == true ]]; then
    echo "=========================================="
@@ -210,8 +225,14 @@ if [[ "$DRY_RUN" == true ]]; then
    echo "    --commit '$(escape_for_shell "${COMMIT}")' \\"
    echo "    --branch '$(escape_for_shell "${BRANCH}")' \\"
    echo "    --message '$(escape_for_shell "${MESSAGE}")' \\"
-    echo "    --env 'RUN_ALL=1' \\"
-    echo "    --env 'NIGHTLY=1'"
+    last_idx=$(( ${#BUILD_ENV[@]} - 1 ))
+    for i in "${!BUILD_ENV[@]}"; do
+        if [[ $i -eq $last_idx ]]; then
+            echo "    --env '$(escape_for_shell "${BUILD_ENV[$i]}")'"
+        else
+            echo "    --env '$(escape_for_shell "${BUILD_ENV[$i]}")' \\"
+        fi
+    done
    echo ""
    echo "=========================================="
    echo -e "${YELLOW}To actually trigger this build, run:${NC}"
@@ -0,0 +1,21 @@
+#!/bin/bash
+
+set -ex
+
+ORIG_TAG_NAME="$BUILDKITE_COMMIT"
+REPO="vllm/vllm-openai-xpu"
+
+echo "Pushing original XPU tag ${ORIG_TAG_NAME}-xpu to nightly tags in ${REPO}"
+
+aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:"$ORIG_TAG_NAME"-x86_64-xpu
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:"$ORIG_TAG_NAME"-x86_64-xpu ${REPO}:nightly-x86_64
+docker push ${REPO}:nightly-x86_64
+
+docker manifest rm ${REPO}:nightly || true
+docker manifest rm ${REPO}:nightly-"$BUILDKITE_COMMIT" || true
+docker manifest create ${REPO}:nightly ${REPO}:nightly-x86_64 --amend
+docker manifest create ${REPO}:nightly-"$BUILDKITE_COMMIT" ${REPO}:nightly-x86_64 --amend
+docker manifest push ${REPO}:nightly
+docker manifest push ${REPO}:nightly-"$BUILDKITE_COMMIT"
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: V1 attention (H100-MI300)
  key: v1-attention-h100-mi300
-  timeout_in_minutes: 30
+  timeout_in_minutes: 85
  device: h100
  source_file_dependencies:
    - vllm/config/attention.py
@@ -16,7 +16,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 70
+      timeout_in_minutes: 95
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -30,7 +30,7 @@ steps:

 - label: V1 attention (B200)
  key: v1-attention-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 80
  device: b200-k8s
  source_file_dependencies:
    - vllm/config/attention.py
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Basic Correctness
  key: basic-correctness
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -19,6 +19,6 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 50
+      timeout_in_minutes: 70
      depends_on:
      - image-build-amd
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Benchmarks CLI Test
  key: benchmarks-cli-test
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -23,7 +23,7 @@ steps:
  num_gpus: 2
  optional: true
  working_dir: "/vllm-workspace/"
-  timeout_in_minutes: 10
+  timeout_in_minutes: 20
  source_file_dependencies:
  - benchmarks/attention_benchmarks/
  - vllm/v1/attention/
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Sequence Parallel Correctness Tests (2 GPUs)
  key: sequence-parallel-correctness-tests-2-gpus
-  timeout_in_minutes: 50
+  timeout_in_minutes: 80
  working_dir: "/vllm-workspace/"
  num_devices: 2
  source_file_dependencies:
@@ -19,7 +19,7 @@ steps:

 - label: Sequence Parallel Correctness Tests (2xH100)
  key: sequence-parallel-correctness-tests-2xh100
-  timeout_in_minutes: 50
+  timeout_in_minutes: 75
  working_dir: "/vllm-workspace/"
  device: h100
  optional: true
@@ -30,7 +30,7 @@ steps:

 - label: AsyncTP Correctness Tests (2xH100)
  key: asynctp-correctness-tests-2xh100
-  timeout_in_minutes: 50
+  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/"
  device: h100
  optional: true
@@ -41,7 +41,7 @@ steps:

 - label: AsyncTP Correctness Tests (B200)
  key: asynctp-correctness-tests-b200
-  timeout_in_minutes: 50
+  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/"
  device: b200-k8s
  optional: true
@@ -52,7 +52,7 @@ steps:

 - label: Distributed Compile Unit Tests (2xH100)
  key: distributed-compile-unit-tests-2xh100
-  timeout_in_minutes: 20
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/"
  device: h100
  num_devices: 2
@@ -66,7 +66,7 @@ steps:

 - label: Fusion and Compile Unit Tests (2xB200)
  key: fusion-and-compile-unit-tests-2xb200
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/"
  device: b200-k8s
  source_file_dependencies:
@@ -96,7 +96,7 @@ steps:

 - label: Fusion E2E Quick (H100)
  key: fusion-e2e-quick-h100
-  timeout_in_minutes: 15
+  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/"
  device: h100
  num_devices: 1
@@ -115,7 +115,7 @@ steps:

 - label: Fusion E2E Config Sweep (H100)
  key: fusion-e2e-config-sweep-h100
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/"
  device: h100
  num_devices: 1
@@ -149,7 +149,7 @@ steps:

 - label: Fusion E2E TP2 Quick (H100)
  key: fusion-e2e-tp2-quick-h100
-  timeout_in_minutes: 20
+  timeout_in_minutes: 35
  working_dir: "/vllm-workspace/"
  device: h100
  num_devices: 2
@@ -167,7 +167,7 @@ steps:

 - label: Fusion E2E TP2 AR-RMS Config Sweep (H100)
  key: fusion-e2e-tp2-ar-rms-config-sweep-h100
-  timeout_in_minutes: 40
+  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/"
  device: h100
  num_devices: 2
@@ -207,7 +207,7 @@ steps:

 - label: Fusion E2E TP2 (B200)
  key: fusion-e2e-tp2-b200
-  timeout_in_minutes: 20
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/"
  device: b200-k8s
  num_devices: 2
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Platform Tests
  key: platform-tests
-  timeout_in_minutes: 15
+  timeout_in_minutes: 20
  device: h200_18gb
  source_file_dependencies:
  - vllm/envs.py
@@ -19,7 +19,7 @@ steps:

 - label: Cudagraph
  key: cudagraph
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  source_file_dependencies:
  - tests/v1/cudagraph
  - vllm/v1/cudagraph_dispatcher.py
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed NixlConnector PD accuracy (4 GPUs)
  key: distributed-nixlconnector-pd-accuracy-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 55
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -16,7 +16,7 @@ steps:
  mirror:
    amd:
      device: mi300_4
-      timeout_in_minutes: 110
+      timeout_in_minutes: 85
      depends_on:
        - image-build-amd
      source_file_dependencies:
@@ -29,7 +29,7 @@ steps:

 - label: Distributed FlashInfer NixlConnector PD accuracy (4 GPUs)
  key: distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 55
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -66,7 +66,7 @@ steps:
  mirror:
    amd:
      device: mi300_4
-      timeout_in_minutes: 50
+      timeout_in_minutes: 60
      depends_on:
        - image-build-amd
      source_file_dependencies:
@@ -79,7 +79,7 @@ steps:

 - label: CrossLayer KV layout Distributed NixlConnector PD accuracy tests (4 GPUs)
  key: crosslayer-kv-layout-distributed-nixlconnector-pd-accuracy-tests-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 55
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -91,7 +91,7 @@ steps:
  mirror:
    amd:
      device: mi300_4
-      timeout_in_minutes: 110
+      timeout_in_minutes: 85
      depends_on:
        - image-build-amd
      source_file_dependencies:
@@ -104,7 +104,7 @@ steps:

 - label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs)
  key: hybrid-ssm-nixlconnector-pd-accuracy-tests-4-gpus
-  timeout_in_minutes: 25
+  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -116,7 +116,7 @@ steps:
  mirror:
    amd:
      device: mi300_4
-      timeout_in_minutes: 60
+      timeout_in_minutes: 80
      depends_on:
        - image-build-amd
      source_file_dependencies:
@@ -143,7 +143,7 @@ steps:

 - label: MultiConnector (Nixl+Offloading) PD accuracy (2 GPUs)
  key: multiconnector-nixl-offloading-pd-accuracy-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -158,7 +158,7 @@ steps:

 - label: NixlConnector PD + Spec Decode acceptance (2 GPUs)
  key: nixlconnector-pd-spec-decode-acceptance-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
  device: a100
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
@@ -172,7 +172,7 @@ steps:
  mirror:
    amd:
      device: mi300_2
-      timeout_in_minutes: 60
+      timeout_in_minutes: 70
      depends_on:
        - image-build-amd
      source_file_dependencies:
@@ -186,7 +186,7 @@ steps:

 - label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs)
  key: multiconnector-nixl-offloading-pd-edge-cases-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed Comm Ops
  key: distributed-comm-ops
-  timeout_in_minutes: 20
+  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -18,7 +18,7 @@ steps:

 - label: Distributed DP Tests (2 GPUs)
  key: distributed-dp-tests-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 35
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -55,7 +55,7 @@ steps:

 - label: Distributed Compile + RPC Tests (2 GPUs)
  key: distributed-compile-rpc-tests-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 65
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -78,7 +78,7 @@ steps:

 - label: Distributed Torchrun + Shutdown Tests (2 GPUs)
  key: distributed-torchrun-shutdown-tests-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -133,7 +133,7 @@ steps:

 - label: Distributed DP Tests (4 GPUs)
  key: distributed-dp-tests-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -154,7 +154,7 @@ steps:

 - label: Distributed Compile + Comm (4 GPUs)
  key: distributed-compile-comm-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 70
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -176,7 +176,7 @@ steps:

 - label: Distributed Tests (8xH100)
  key: distributed-tests-8xh100
-  timeout_in_minutes: 10
+  timeout_in_minutes: 20
  device: h100
  num_devices: 8
  working_dir: "/vllm-workspace/tests"
@@ -212,7 +212,7 @@ steps:

 - label: Distributed Tests (2xH100-2xMI300)
  key: distributed-tests-2xh100-2xmi300
-  timeout_in_minutes: 15
+  timeout_in_minutes: 30
  device: h100
  optional: true
  working_dir: "/vllm-workspace/"
@@ -259,7 +259,7 @@ steps:

 - label: Pipeline + Context Parallelism (4 GPUs)
  key: pipeline-context-parallelism-4-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 55
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -274,7 +274,7 @@ steps:

 - label: RayExecutorV2 (4 GPUs)
  key: rayexecutorv2-4-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -3,7 +3,7 @@ depends_on:
  - image-build-cpu
 steps:
 - label: Docker Build Metadata
-  timeout_in_minutes: 10
+  timeout_in_minutes: 20
  device: cpu-small
  source_file_dependencies:
    - .buildkite/release-pipeline.yaml
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: DeepSeek V2-Lite Sync EPLB Accuracy (4xH100)
  key: deepseek-v2-lite-sync-eplb-accuracy-4xh100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
  device: h100
  optional: true
  num_devices: 4
@@ -14,7 +14,7 @@ steps:

 - label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (4xH100)
  key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-4xh100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
  device: h100
  optional: true
  num_devices: 4
@@ -24,7 +24,7 @@ steps:

 - label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (2xB200)
  key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-2xb200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 20
  device: b200-k8s
  optional: true
  num_devices: 2
@@ -34,7 +34,7 @@ steps:

 - label: Qwen3-30B-A3B-FP8 DP4 Async EPLB Accuracy
  key: qwen3-30b-a3b-fp8-dp4-async-eplb-accuracy
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
  device: h100
  optional: true
  num_devices: 4
@@ -44,7 +44,7 @@ steps:

 - label: DeepSeek V2-Lite Prefetch Offload Accuracy (H100)
  key: deepseek-v2-lite-prefetch-offload-accuracy-h100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 20
  device: h100
  optional: true
  num_devices: 1
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Engine
  key: engine
-  timeout_in_minutes: 15
+  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
  - vllm/compilation/
@@ -29,13 +29,13 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 50
      depends_on:
      - image-build-amd

 - label: Engine (1 GPU)
  key: engine-1-gpu
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
  source_file_dependencies:
    - vllm/v1/engine/
    - tests/v1/engine/
@@ -45,13 +45,13 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 40
+      timeout_in_minutes: 55
      depends_on:
      - image-build-amd

 - label: e2e Scheduling (1 GPU)
  key: e2e-scheduling-1-gpu
-  timeout_in_minutes: 30
+  timeout_in_minutes: 35
  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/
@@ -61,14 +61,14 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 70
      depends_on:
      - image-build-amd

 - label: e2e Core (1 GPU)
  device: h200_35gb
  key: e2e-core-1-gpu
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
  source_file_dependencies:
    - vllm/v1/
    - tests/v1/e2e/general/
@@ -77,7 +77,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 35
+      timeout_in_minutes: 60
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -87,7 +87,7 @@ steps:

 - label: V1 e2e (2 GPUs)
  key: v1-e2e-2-gpus
-  timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
+  timeout_in_minutes: 25 # TODO: Fix timeout after we have more confidence in the test stability
  optional: true
  num_devices: 2
  source_file_dependencies:
@@ -120,7 +120,7 @@ steps:

 - label: V1 e2e (4 GPUs)
  key: v1-e2e-4-gpus
-  timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
+  timeout_in_minutes: 20 # TODO: Fix timeout after we have more confidence in the test stability
  optional: true
  num_devices: 4
  source_file_dependencies:
@@ -148,7 +148,7 @@ steps:

 - label: V1 e2e (4xH100)
  key: v1-e2e-4xh100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 35
  device: h100
  num_devices: 4
  optional: true
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Entrypoints Unit Tests
  key: entrypoints-unit-tests
-  timeout_in_minutes: 10
+  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/entrypoints
@@ -16,7 +16,7 @@ steps:

 - label: Entrypoints Integration (LLM)
  key: entrypoints-integration-llm
-  timeout_in_minutes: 40
+  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -37,7 +37,7 @@ steps:
 - label: Entrypoints Integration (API Server)
  key: entrypoints-integration-api-server
  device: h200_35gb
-  timeout_in_minutes: 130
+  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -56,7 +56,7 @@ steps:

 - label: Entrypoints Integration (API Server OpenAI - Part 1)
  key: entrypoints-integration-api-server-openai-part-1
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -68,13 +68,13 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 80
+      timeout_in_minutes: 65
      depends_on:
      - image-build-amd

 - label: Entrypoints Integration (API Server OpenAI - Part 2)
  key: entrypoints-integration-api-server-openai-part-2
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -109,7 +109,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 65
      depends_on:
      - image-build-amd

@@ -126,7 +126,7 @@ steps:
 - label: Entrypoints Integration (Speech to Text)
  device: h200_35gb
  key: entrypoints-integration-speech_to_text
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -138,7 +138,7 @@ steps:
 - label: Entrypoints Integration (Multimodal)
  device: h200_35gb
  key: entrypoints-integration-multimodal
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -160,7 +160,7 @@ steps:

 - label: OpenAI API Correctness
  key: openai-api-correctness
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: h200_18gb
  source_file_dependencies:
  - csrc/
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: EPLB Algorithm
  key: eplb-algorithm
-  timeout_in_minutes: 15
+  timeout_in_minutes: 20
  device: h200_18gb
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -27,7 +27,7 @@ steps:

 - label: EPLB Execution # 17min
  key: eplb-execution
-  timeout_in_minutes: 27
+  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -39,7 +39,7 @@ steps:

 - label: Elastic EP Scaling Test
  key: elastic-ep-scaling-test
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  device: h100
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: vLLM IR Tests
  key: vllm-ir-tests
-  timeout_in_minutes: 10
+  timeout_in_minutes: 35
  device: h200_18gb
  working_dir: "/vllm-workspace/"
  source_file_dependencies:
@@ -16,7 +16,7 @@ steps:

 - label: Kernels Core Operation Test
  key: kernels-core-operation-test
-  timeout_in_minutes: 75
+  timeout_in_minutes: 120
  source_file_dependencies:
  - csrc/
  - tests/kernels/core
@@ -27,7 +27,7 @@ steps:

 - label: Kernels MiniMax Reduce RMS Test (2 GPUs)
  key: kernels-minimax-reduce-rms-test-2-gpus
-  timeout_in_minutes: 15
+  timeout_in_minutes: 20
  num_devices: 2
  device: h100
  source_file_dependencies:
@@ -41,7 +41,7 @@ steps:

 - label: Deepseek V4 Kernel Test (H100)
  key: deepseek-v4-kernel-test-h100
-  timeout_in_minutes: 15
+  timeout_in_minutes: 30
  device: h100
  source_file_dependencies:
  - csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
@@ -54,7 +54,7 @@ steps:

 - label: Deepseek V4 Kernel Test (B200)
  key: deepseek-v4-kernel-test-b200
-  timeout_in_minutes: 15
+  timeout_in_minutes: 20
  device: b200-k8s
  source_file_dependencies:
  - csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
@@ -65,7 +65,7 @@ steps:

 - label: Kernels Attention Test %N
  key: kernels-attention-test
-  timeout_in_minutes: 35
+  timeout_in_minutes: 65
  source_file_dependencies:
  - csrc/attention/
  - vllm/v1/attention
@@ -79,7 +79,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 55
+      timeout_in_minutes: 90
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -106,7 +106,7 @@ steps:

 - label: Kernels Quantization Test %N
  key: kernels-quantization-test
-  timeout_in_minutes: 90
+  timeout_in_minutes: 60
  source_file_dependencies:
  - csrc/quantization/
  - vllm/model_executor/layers/quantization
@@ -131,7 +131,7 @@ steps:

 - label: Kernels MoE Test %N
  key: kernels-moe-test
-  timeout_in_minutes: 25
+  timeout_in_minutes: 50
  source_file_dependencies:
  - csrc/quantization/cutlass_w8a8/moe/
  - csrc/moe/
@@ -147,7 +147,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 50
+      timeout_in_minutes: 65
      source_file_dependencies:
      - csrc/quantization/cutlass_w8a8/moe/
      - csrc/moe/
@@ -163,7 +163,7 @@ steps:

 - label: Kernels Mamba Test
  key: kernels-mamba-test
-  timeout_in_minutes: 45
+  timeout_in_minutes: 40
  source_file_dependencies:
  - csrc/mamba/
  - tests/kernels/mamba
@@ -172,7 +172,7 @@ steps:
    - pytest -v -s kernels/mamba

 - label: Kernels KDA Test
-  timeout_in_minutes: 20
+  timeout_in_minutes: 25
  device: h200_18gb
  source_file_dependencies:
  - vllm/model_executor/layers/fla/ops/kda.py
@@ -184,7 +184,7 @@ steps:

 - label: Kernels DeepGEMM Test (H100)
  key: kernels-deepgemm-test-h100
-  timeout_in_minutes: 45
+  timeout_in_minutes: 35
  device: h100
  num_devices: 1
  source_file_dependencies:
@@ -211,7 +211,7 @@ steps:

 - label: Kernels (B200)
  key: kernels-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 80
  working_dir: "/vllm-workspace/"
  device: b200-k8s
  # optional: true
@@ -264,7 +264,7 @@ steps:

 - label: Kernels Helion Test
  key: kernels-helion-test
-  timeout_in_minutes: 30
+  timeout_in_minutes: 115
  device: h100
  source_file_dependencies:
  - vllm/utils/import_utils.py
@@ -276,7 +276,7 @@ steps:
 
 - label: Kernels FP8 MoE Test (1xH100)
  key: kernels-fp8-moe-test-1xh100
-  timeout_in_minutes: 90
+  timeout_in_minutes: 40
  device: h100
  num_devices: 1
  optional: true
@@ -293,7 +293,7 @@ steps:

 - label: Kernels FP8 MoE Test (2xH100)
  key: kernels-fp8-moe-test-2xh100
-  timeout_in_minutes: 90
+  timeout_in_minutes: 45
  device: h100
  num_devices: 2
  optional: true
@@ -303,7 +303,7 @@ steps:

 - label: Kernels Fp4 MoE Test (B200)
  key: kernels-fp4-moe-test-b200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
  device: b200-k8s
  num_devices: 1
  optional: true
@@ -316,7 +316,7 @@ steps:

 - label: Kernels FusedMoE Layer Test (2 H100s)
  key: kernels-fusedmoe-layer-test-2-h100s
-  timeout_in_minutes: 90
+  timeout_in_minutes: 30
  device: h100
  num_devices: 2
  source_file_dependencies:
@@ -5,7 +5,7 @@ steps:
 - label: LM Eval Small Models
  device: h200_35gb
  key: lm-eval-small-models
-  timeout_in_minutes: 75
+  timeout_in_minutes: 45
  source_file_dependencies:
  - csrc/
  - vllm/model_executor/layers/quantization
@@ -56,7 +56,7 @@ steps:

 - label: LM Eval Small Models (1xB200)
  key: lm-eval-small-models-1xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 50
  device: b200-k8s
  optional: true
  source_file_dependencies:
@@ -80,7 +80,7 @@ steps:

 - label: LM Eval Large Models EP (2xB200)
  key: lm-eval-large-models-ep-2xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 60
  device: b200-k8s
  optional: true
  num_devices: 2
@@ -92,7 +92,7 @@ steps:

 - label: LM Eval Qwen3.5 Models (2xB200)
  key: lm-eval-qwen3-5-models-2xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 45
  device: b200-k8s
  optional: true
  num_devices: 2
@@ -109,7 +109,7 @@ steps:

 - label: LM Eval Large Models (8xH200)
  key: lm-eval-large-models-8xh200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 50
  device: h200
  optional: true
  num_devices: 8
@@ -118,7 +118,7 @@ steps:
  mirror:
    amd:
      device: mi300_8
-      timeout_in_minutes: 180
+      timeout_in_minutes: 60
      depends_on:
      - image-build-amd
      commands:
@@ -152,7 +152,7 @@ steps:

 - label: LM Eval Humming f16 (A100 - TEMPORARY)
  key: lm-eval-humming-f16-a100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 75
  device: a100
  optional: true
  num_devices: 1
@@ -167,7 +167,7 @@ steps:

 - label: LM Eval Humming Act int8 (A100 - TEMPORARY)
  key: lm-eval-humming-act-a100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 45
  device: a100
  optional: true
  num_devices: 1
@@ -182,7 +182,7 @@ steps:

 - label: LM Eval Humming f16 (H100 - TEMPORARY)
  key: lm-eval-humming-f16-h100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 70
  device: h100
  optional: true
  num_devices: 1
@@ -197,7 +197,7 @@ steps:

 - label: LM Eval Humming Act fp8/int8 (H100 - TEMPORARY)
  key: lm-eval-humming-act-h100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 70
  device: h100
  optional: true
  num_devices: 1
@@ -213,7 +213,7 @@ steps:

 - label: LM Eval Humming f16 (B200 - TEMPORARY)
  key: lm-eval-humming-f16-b200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 50
  device: b200-k8s
  optional: true
  num_devices: 1
@@ -228,7 +228,7 @@ steps:

 - label: LM Eval Humming Act fp8/int8 (B200 - TEMPORARY)
  key: lm-eval-humming-act-b200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 50
  device: b200-k8s
  optional: true
  num_devices: 1
@@ -244,7 +244,7 @@ steps:

 - label: LM Eval TurboQuant KV Cache
  key: lm-eval-turboquant-kv-cache
-  timeout_in_minutes: 75
+  timeout_in_minutes: 55
  device: h200_18gb
  source_file_dependencies:
  - vllm/model_executor/layers/quantization/turboquant/
@@ -256,7 +256,7 @@ steps:

 - label: GPQA Eval (GPT-OSS) (2xH100)
  key: gpqa-eval-gpt-oss-2xh100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 35
  device: h100
  optional: true
  num_devices: 2
@@ -270,7 +270,7 @@ steps:

 - label: GPQA Eval (GPT-OSS) (2xB200)
  key: gpqa-eval-gpt-oss-2xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 30
  device: b200-k8s
  optional: true
  num_devices: 2
@@ -284,7 +284,7 @@ steps:

 - label: GPQA Eval (GPT-OSS) (DGX Spark)
  key: gpqa-eval-gpt-oss-spark
-  timeout_in_minutes: 120
+  timeout_in_minutes: 35
  device: dgx-spark
  optional: true
  num_devices: 1
@@ -313,7 +313,7 @@ steps:

 - label: LM Eval KV-Offload (2xH100)
  key: kv-offload-medium
-  timeout_in_minutes: 60
+  timeout_in_minutes: 30
  device: h100
  num_devices: 2
  source_file_dependencies:
@@ -327,7 +327,7 @@ steps:

 - label: LM Eval KV-Offload (4xH100)
  key: kv-offload-large
-  timeout_in_minutes: 60
+  timeout_in_minutes: 40
  device: h100
  num_devices: 4
  source_file_dependencies:
@@ -341,7 +341,7 @@ steps:

 - label: MRCR Eval Small Models
  device: h200_35gb
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
  source_file_dependencies:
  - tests/evals/mrcr/
  commands:
@@ -5,7 +5,7 @@ steps:
 - label: LoRA %N
  device: h200_35gb
  key: lora
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
  source_file_dependencies:
  - vllm/lora
  - tests/lora
@@ -16,7 +16,7 @@ steps:
    amd:
      device: mi325_1
      working_dir: "/vllm-workspace/tests"
-      timeout_in_minutes: 60
+      timeout_in_minutes: 65
      source_file_dependencies:
      - vllm/lora
      - tests/lora
@@ -27,7 +27,7 @@ steps:

 - label: LoRA TP (Distributed)
  key: lora-tp-distributed
-  timeout_in_minutes: 30
+  timeout_in_minutes: 60
  num_devices: 4
  source_file_dependencies:
  - vllm/lora
@@ -5,7 +5,7 @@ steps:
 - label: V1 Spec Decode
  device: h200_35gb
  key: v1-spec-decode
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
  source_file_dependencies:
    - vllm/config/
    - vllm/distributed/
@@ -24,13 +24,13 @@ steps:
  mirror:
    amd:
      device: mi300_1
-      timeout_in_minutes: 65
+      timeout_in_minutes: 75
      depends_on:
      - image-build-amd

 - label: V1 Sample + Logits
  key: v1-sample-logits
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
  device: h200_18gb
  source_file_dependencies:
    - vllm/config/
@@ -64,7 +64,7 @@ steps:

 - label: V1 Core + KV + Metrics
  key: v1-core-kv-metrics
-  timeout_in_minutes: 30
+  timeout_in_minutes: 60
  source_file_dependencies:
    - vllm/config/
    - vllm/distributed/
@@ -108,7 +108,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 75
      depends_on:
      - image-build-amd

@@ -172,7 +172,7 @@ steps:

 - label: Regression
  key: regression
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
  - vllm/config/
@@ -195,7 +195,7 @@ steps:
 - label: Examples
  device: h200_35gb
  key: examples
-  timeout_in_minutes: 45
+  timeout_in_minutes: 40
  working_dir: "/vllm-workspace/examples"
  source_file_dependencies:
  - vllm/entrypoints
@@ -237,7 +237,7 @@ steps:

 - label: Metrics, Tracing (2 GPUs)
  key: metrics-tracing-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 25
  num_devices: 2
  source_file_dependencies:
  - vllm/config/
@@ -281,7 +281,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 20
+      timeout_in_minutes: 45
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -292,7 +292,7 @@ steps:
 - label: Async Engine, Inputs, Utils, Worker
  device: h200_35gb
  key: async-engine-inputs-utils-worker
-  timeout_in_minutes: 50
+  timeout_in_minutes: 25
  source_file_dependencies:
  - vllm/assets/
  - vllm/config/
@@ -319,7 +319,7 @@ steps:
  key: async-engine-inputs-utils-worker-config-cpu
  depends_on:
  - image-build-cpu
-  timeout_in_minutes: 30
+  timeout_in_minutes: 65
  source_file_dependencies:
  - vllm/assets/
  - vllm/config/
@@ -381,7 +381,7 @@ steps:

 - label: Batch Invariance (A100)
  key: batch-invariance-a100
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
  device: a100
  source_file_dependencies:
    - vllm/v1/attention
@@ -395,7 +395,7 @@ steps:

 - label: Batch Invariance (H100)
  key: batch-invariance-h100
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
  device: h100
  source_file_dependencies:
    - vllm/v1/attention
@@ -411,7 +411,7 @@ steps:

 - label: Batch Invariance (B200)
  key: batch-invariance-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 35
  device: b200-k8s
  source_file_dependencies:
    - vllm/v1/attention
@@ -430,7 +430,7 @@ steps:
 - label: Acceptance Length Test (Large Models) # optional
  device: h200_35gb
  key: acceptance-length-test-large-models
-  timeout_in_minutes: 25
+  timeout_in_minutes: 20
  gpu: h100
  optional: true
  num_gpus: 1
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Model Executor
  key: model-executor
-  timeout_in_minutes: 35
+  timeout_in_minutes: 45
  source_file_dependencies:
  - vllm/engine/arg_utils.py
  - vllm/config/model.py
@@ -5,7 +5,7 @@ steps:
 - label: Model Runner V2 Core Tests
  device: h200_35gb
  key: model-runner-v2-core-tests
-  timeout_in_minutes: 45
+  timeout_in_minutes: 35
  source_file_dependencies:
  - vllm/v1/worker/gpu/
  - vllm/v1/worker/gpu_worker.py
@@ -27,7 +27,7 @@ steps:
 - label: Model Runner V2 Examples
  device: h200_35gb
  key: model-runner-v2-examples
-  timeout_in_minutes: 45
+  timeout_in_minutes: 35
  working_dir: "/vllm-workspace/examples"
  source_file_dependencies:
    - vllm/v1/worker/gpu/
@@ -63,7 +63,7 @@ steps:

 - label: Model Runner V2 Distributed (2 GPUs)
  key: model-runner-v2-distributed-2-gpus
-  timeout_in_minutes: 45
+  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -84,7 +84,7 @@ steps:

 - label: Model Runner V2 Pipeline Parallelism (4 GPUs)
  key: model-runner-v2-pipeline-parallelism-4-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Basic Models Tests (Initialization)
  key: basic-models-tests-initialization
-  timeout_in_minutes: 45
+  timeout_in_minutes: 25
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -17,7 +17,7 @@ steps:
 - label: Basic Models Tests (Extra Initialization) %N
  device: h200_35gb
  key: basic-models-tests-extra-initialization
-  timeout_in_minutes: 45
+  timeout_in_minutes: 100
  source_file_dependencies:
  - vllm/model_executor/models/
  - tests/models/test_initialization.py
@@ -32,7 +32,7 @@ steps:
 - label: Basic Models Tests (Other)
  device: h200_35gb
  key: basic-models-tests-other
-  timeout_in_minutes: 45
+  timeout_in_minutes: 35
  source_file_dependencies:
  - vllm/
  - tests/models/test_terratorch.py
@@ -50,7 +50,7 @@ steps:
  key: basic-models-test-other-cpu
  depends_on:
  - image-build-cpu
-  timeout_in_minutes: 10
+  timeout_in_minutes: 20
  source_file_dependencies:
  - vllm/
  - tests/models/test_utils.py
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed Model Tests (2 GPUs)
  key: distributed-model-tests-2-gpus
-  timeout_in_minutes: 50
+  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Language Models Tests (Standard)
  key: language-models-tests-standard
-  timeout_in_minutes: 25
+  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -21,7 +21,7 @@ steps:

 - label: Language Models Tests (Extra Standard) %N
  key: language-models-tests-extra-standard
-  timeout_in_minutes: 45
+  timeout_in_minutes: 40
  source_file_dependencies:
  - vllm/model_executor/models/
  - tests/models/language/pooling/test_embedding.py
@@ -52,7 +52,7 @@ steps:

 - label: Language Models Tests (Hybrid) %N
  key: language-models-tests-hybrid
-  timeout_in_minutes: 75
+  timeout_in_minutes: 65
  source_file_dependencies:
  - vllm/
  - tests/models/language/generation
@@ -67,7 +67,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 90
+      timeout_in_minutes: 70
      depends_on:
      - image-build-amd
      commands:
@@ -78,7 +78,7 @@ steps:
 - label: Language Models Test (Extended Generation) # 80min
  device: h200_35gb
  key: language-models-test-extended-generation
-  timeout_in_minutes: 110
+  timeout_in_minutes: 65
  optional: true
  source_file_dependencies:
  - vllm/
@@ -92,7 +92,7 @@ steps:

 - label: Language Models Test (PPL)
  key: language-models-test-ppl
-  timeout_in_minutes: 110
+  timeout_in_minutes: 30
  device: h200_18gb
  optional: true
  source_file_dependencies:
@@ -104,7 +104,7 @@ steps:
 - label: Language Models Test (Extended Pooling)  # 36min
  device: h200_35gb
  key: language-models-test-extended-pooling
-  timeout_in_minutes: 50
+  timeout_in_minutes: 70
  optional: true
  source_file_dependencies:
  - vllm/
@@ -120,7 +120,7 @@ steps:

 - label: Language Models Test (MTEB)
  key: language-models-test-mteb
-  timeout_in_minutes: 110
+  timeout_in_minutes: 45
  device: h200_18gb
  optional: true
  source_file_dependencies:
@@ -20,7 +20,7 @@ steps:

 - label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
  key: multi-modal-models-standard-2-qwen3-gemma
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -38,7 +38,7 @@ steps:
 - label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
  device: h200_35gb
  key: multi-modal-models-standard-3-llava-qwen2-vl
-  timeout_in_minutes: 45
+  timeout_in_minutes: 40
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal
@@ -54,7 +54,7 @@ steps:
 - label: "Multi-Modal Models (Standard) 4: other + whisper"
  device: h200_35gb
  key: multi-modal-models-standard-4-other-whisper
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal
@@ -73,7 +73,7 @@ steps:
  key: multi-modal-processor-cpu
  depends_on:
  - image-build-cpu
-  timeout_in_minutes: 60
+  timeout_in_minutes: 125
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal
@@ -84,7 +84,7 @@ steps:

 - label: Multi-Modal Processor # 44min
  key: multi-modal-processor
-  timeout_in_minutes: 60
+  timeout_in_minutes: 65
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -96,7 +96,7 @@ steps:
 - label: Multi-Modal Accuracy Eval (Small Models) # 50min
  device: h200_35gb
  key: multi-modal-accuracy-eval-small-models
-  timeout_in_minutes: 70
+  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
  source_file_dependencies:
  - vllm/multimodal/
@@ -164,7 +164,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 75
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Plugin Tests (2 GPUs)
  key: plugin-tests-2-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 35
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -5,7 +5,7 @@ steps:
 - label: PyTorch Compilation Unit Tests
  device: h200_35gb
  key: pytorch-compilation-unit-tests
-  timeout_in_minutes: 10
+  timeout_in_minutes: 90
  source_file_dependencies:
    - vllm/__init__.py
    - vllm/_aiter_ops.py
@@ -78,7 +78,7 @@ steps:

 - label: PyTorch Compilation Passes Unit Tests
  key: pytorch-compilation-passes-unit-tests
-  timeout_in_minutes: 20
+  timeout_in_minutes: 45
  source_file_dependencies:
    - vllm/__init__.py
    - vllm/_aiter_ops.py
@@ -110,13 +110,13 @@ steps:
  mirror:
    amd:
      device: mi300_1
-      timeout_in_minutes: 180
+      timeout_in_minutes: 65
      depends_on:
      - image-build-amd

 - label: PyTorch Fullgraph Smoke Test
  key: pytorch-fullgraph-smoke-test
-  timeout_in_minutes: 35
+  timeout_in_minutes: 60
  source_file_dependencies:
  - vllm/__init__.py
  - vllm/_aiter_ops.py
@@ -152,7 +152,7 @@ steps:

 - label: PyTorch Fullgraph
  key: pytorch-fullgraph
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
  device: h200_18gb
  source_file_dependencies:
  - vllm/__init__.py
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Quantization
  key: quantization
-  timeout_in_minutes: 90
+  timeout_in_minutes: 60
  source_file_dependencies:
  - csrc/
  - vllm/model_executor/layers/quantization
@@ -23,7 +23,7 @@ steps:

 - label: Quantized Fusions
  key: quantized-fusions
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  source_file_dependencies:
  - tests/fusion
  - vllm/model_executor/layers/fusion
@@ -35,7 +35,7 @@ steps:

 - label: Quantized MoE Test (B200)
  key: quantized-moe-test-b200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 120
  working_dir: "/vllm-workspace/"
  device: b200-k8s
  source_file_dependencies:
@@ -53,7 +53,7 @@ steps:

 - label: Quantized Models Test
  key: quantized-models-test
-  timeout_in_minutes: 60
+  timeout_in_minutes: 50
  source_file_dependencies:
  - vllm/model_executor/layers/quantization
  - tests/models/quantization
@@ -3,7 +3,7 @@ depends_on:
  - image-build
 steps:
 - label: Rust Frontend OpenAI Coverage
-  timeout_in_minutes: 90
+  timeout_in_minutes: 30
  device: h200_18gb
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -38,7 +38,7 @@ steps:
  - pytest -v -s v1/sample/test_logprobs_e2e.py -k "test_prompt_logprobs_e2e_server"

 - label: Rust Frontend Serve/Admin Coverage
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
  device: h200_18gb
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -67,7 +67,7 @@ steps:
  - pytest -v -s entrypoints/serve/tokenize/test_tokenization.py -k "not tokenizer_info"

 - label: Rust Frontend Core Correctness
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: h200_18gb
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -81,7 +81,7 @@ steps:
  - pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine

 - label: Rust Frontend Tool Use
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - rust/
@@ -95,7 +95,7 @@ steps:
  - pytest -v -s tool_use --ignore=tool_use/mistral --models llama3.2 -k "not test_response_format_with_tool_choice_required and not test_parallel_tool_calls_false and not test_tool_call_and_choice"

 - label: Rust Frontend Distributed
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
  num_devices: 4
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -4,7 +4,7 @@ steps:
 - label: Rust Frontend Cargo Style + Clippy
  key: rust-frontend-cargo-style-clippy
  depends_on: []
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: cpu-medium
  no_plugin: true
  source_file_dependencies:
@@ -18,7 +18,7 @@ steps:
 - label: Rust Frontend Cargo Tests
  key: rust-frontend-cargo-tests
  depends_on: []
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: cpu-medium
  no_plugin: true
  source_file_dependencies:
@@ -5,7 +5,7 @@ steps:
 - label: Samplers Test
  device: h200_35gb
  key: samplers-test
-  timeout_in_minutes: 75
+  timeout_in_minutes: 40
  source_file_dependencies:
  - vllm/model_executor/layers
  - vllm/sampling_metadata.py
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Spec Decode Eagle
  key: spec-decode-eagle
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/spec_decode/
@@ -15,7 +15,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 45
+      timeout_in_minutes: 60
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -29,7 +29,7 @@ steps:

 - label: Spec Decode Eagle Nightly B200
  key: spec-decode-eagle-nightly-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
  device: b200-k8s
  optional: true
  source_file_dependencies:
@@ -41,7 +41,7 @@ steps:

 - label: Spec Decode Speculators + MTP
  key: spec-decode-speculators-mtp
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/spec_decode/
@@ -82,7 +82,7 @@ steps:
  
 - label: Spec Decode Ngram + Suffix
  key: spec-decode-ngram-suffix
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/spec_decode/
@@ -93,7 +93,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 65
+      timeout_in_minutes: 55
      # TODO(akaratza): Test after Torch >= 2.12 bump
      soft_fail: true
      depends_on:
@@ -109,7 +109,7 @@ steps:

 - label: Spec Decode Draft Model
  key: spec-decode-draft-model
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/spec_decode/
@@ -120,7 +120,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 50
+      timeout_in_minutes: 55
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -134,7 +134,7 @@ steps:

 - label: Spec Decode Draft Model Nightly B200
  key: spec-decode-draft-model-nightly-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
  device: b200-k8s
  optional: true
  source_file_dependencies:
@@ -146,7 +146,7 @@ steps:

 - label: Speculators Correctness
  key: speculators-correctness
-  timeout_in_minutes: 60
+  timeout_in_minutes: 30
  device: h100
  optional: true
  num_devices: 1
@@ -159,7 +159,7 @@ steps:
    - pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test

 - label: Spec Decode MTP hybrid (B200)
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: b200-k8s
  optional: true
  source_file_dependencies:
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Weight Loading Multiple GPU  # 33min
  key: weight-loading-multiple-gpu
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  optional: true
@@ -70,6 +70,15 @@ endif()
 #
 set(TORCH_SUPPORTED_VERSION_CUDA "2.11.0")
 set(TORCH_SUPPORTED_VERSION_ROCM "2.11.0")
+# TORCH_NIGHTLY=1 builds run against unpinned nightly wheels, so the supported-
+# version check would always warn. Only treat it as a nightly build when the
+# value is exactly "1" (the bootstrap exports TORCH_NIGHTLY=0 by default, which
+# must NOT suppress the warning for normal builds).
+if (DEFINED ENV{TORCH_NIGHTLY} AND "$ENV{TORCH_NIGHTLY}" STREQUAL "1")
+  set(TORCH_NIGHTLY_BUILD TRUE)
+else()
+  set(TORCH_NIGHTLY_BUILD FALSE)
+endif()

 #
 # Try to find python package with an executable that exactly matches
@@ -175,7 +184,7 @@ endif()
 if (NOT HIP_FOUND AND NOT PYTORCH_FOUND_HIP AND CUDA_FOUND)
  set(VLLM_GPU_LANG "CUDA")

-  if (NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
+  if (NOT TORCH_NIGHTLY_BUILD AND NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
    message(WARNING "Pytorch version ${TORCH_SUPPORTED_VERSION_CUDA} "
      "expected for CUDA build, saw ${Torch_VERSION} instead.")
  endif()
@@ -188,7 +197,7 @@ elseif(HIP_FOUND OR PYTORCH_FOUND_HIP)
  enable_language(HIP)

  # ROCm 5.X and 6.X
-  if (ROCM_VERSION_DEV_MAJOR GREATER_EQUAL 5 AND
+  if (NOT TORCH_NIGHTLY_BUILD AND ROCM_VERSION_DEV_MAJOR GREATER_EQUAL 5 AND
      Torch_VERSION VERSION_LESS ${TORCH_SUPPORTED_VERSION_ROCM})
    message(WARNING "Pytorch version >= ${TORCH_SUPPORTED_VERSION_ROCM} "
      "expected for ROCm build, saw ${Torch_VERSION} instead.")
@@ -381,6 +390,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
    "csrc/libtorch_stable/cuda_view.cu"
    "csrc/libtorch_stable/cuda_utils_kernels.cu"
    "csrc/libtorch_stable/activation_kernels.cu"
+    "csrc/libtorch_stable/ngram_embedding_kernels.cu"
    "csrc/libtorch_stable/quantization/activation_kernels.cu"
    "csrc/libtorch_stable/quantization/w8a8/int8/scaled_quant.cu"
    "csrc/libtorch_stable/quantization/w8a8/fp8/common.cu"
@@ -0,0 +1,96 @@
+// N-gram embedding index kernel for LongCat-Flash (n-gram embedding variant).
+//
+// Adapted from SGLang:
+// https://github.com/sgl-project/sglang/blob/main/python/sglang/jit_kernel/csrc/ngram_embedding.cuh
+//
+// For each position, computes the hashed n-gram embedding ids that index the
+// concatenated embedder table. Integer tensors are int32 except ``row_indices``
+// (int64); the token table is ``[max_running_reqs, max_context_len]`` int32,
+// where a negative entry marks an ignored token (e.g. an EOS boundary).
+
+#include "torch_utils.h"
+
+#include "ops.h"
+
+#include <cstdint>
+
+namespace vllm::ngram_embedding {
+
+constexpr int kBlockThreads = 256;
+
+__global__ void ComputeNGramIdsKernel(
+    int batch_size, int ne_n, int ne_k,
+    int* ne_weights,                       // [ne_n-1, ne_k, ne_n]
+    int* ne_mods,                          // [ne_n-1, ne_k]
+    int* exclusive_ne_embedder_size_sums,  // [(ne_n-1)*ne_k + 1]
+    int* exclusive_req_len_sums,           // [batch_size + 1]
+    int* ne_token_table,  // [max_running_reqs, max_context_len]
+    int max_context_len,
+    const int64_t* __restrict__ row_indices,  // [batch_size]
+    int* column_starts,                       // [batch_size]
+    int* n_gram_ids                           // [token_num, (ne_n-1)*ne_k]
+) {
+  const int req_id = blockIdx.x % batch_size;
+  const int config_id = (blockIdx.x - req_id) / batch_size;
+  // n and k are offset from their physical meaning: n = real_n - 2, k = real_k
+  // - 1 (they index into ne_weights / ne_mods).
+  const int k = config_id % ne_k;
+  const int n = (config_id - config_id % ne_k) / ne_k;
+  const int ne_weight_base_idx = n * ne_k * ne_n + k * ne_n;
+  const int ne_mod = ne_mods[n * ne_k + k];
+  for (int i = exclusive_req_len_sums[req_id] + threadIdx.x;
+       i < exclusive_req_len_sums[req_id + 1]; i += blockDim.x) {
+    uint64_t n_gram_id = 0;
+    const int64_t current_token_offset = i - exclusive_req_len_sums[req_id];
+    const int64_t req_token_table_index =
+        row_indices[req_id] * static_cast<int64_t>(max_context_len);
+    const int64_t current_token_table_index =
+        req_token_table_index + column_starts[req_id] + current_token_offset;
+    for (int j = 0; j < n + 2; j++) {
+      if (current_token_table_index - j < req_token_table_index) {
+        break;  // outside this request's range
+      }
+      if (ne_token_table[current_token_table_index - j] < 0) {
+        break;  // ignored token
+      }
+      const uint64_t term =
+          (uint64_t)ne_token_table[current_token_table_index - j] *
+          (uint64_t)ne_weights[ne_weight_base_idx + j];
+      n_gram_id += term % ne_mod;
+    }
+    n_gram_id %= ne_mod;
+    n_gram_id += exclusive_ne_embedder_size_sums[n * ne_k + k];
+    n_gram_ids[i * (ne_n - 1) * ne_k + n * ne_k + k] = (int)(n_gram_id);
+  }
+}
+
+}  // namespace vllm::ngram_embedding
+
+void ngram_compute_n_gram_ids(
+    int64_t ne_n, int64_t ne_k, torch::stable::Tensor& ne_weights,
+    torch::stable::Tensor& ne_mods,
+    torch::stable::Tensor& exclusive_ne_embedder_size_sums,
+    torch::stable::Tensor& exclusive_req_len_sums,
+    torch::stable::Tensor& ne_token_table, torch::stable::Tensor& row_indices,
+    torch::stable::Tensor& column_starts, torch::stable::Tensor& n_gram_ids) {
+  const int batch_size = static_cast<int>(exclusive_req_len_sums.size(0) - 1);
+  const int max_context_len = static_cast<int>(ne_token_table.size(1));
+  const int num_configs = (static_cast<int>(ne_n) - 1) * static_cast<int>(ne_k);
+  const int grid_size = num_configs * batch_size;
+  if (grid_size <= 0) return;
+
+  const torch::stable::accelerator::DeviceGuard device_guard(
+      ne_weights.get_device_index());
+  const cudaStream_t stream = get_current_cuda_stream();
+  vllm::ngram_embedding::ComputeNGramIdsKernel<<<
+      grid_size, vllm::ngram_embedding::kBlockThreads, 0, stream>>>(
+      batch_size, static_cast<int>(ne_n), static_cast<int>(ne_k),
+      ne_weights.mutable_data_ptr<int32_t>(),
+      ne_mods.mutable_data_ptr<int32_t>(),
+      exclusive_ne_embedder_size_sums.mutable_data_ptr<int32_t>(),
+      exclusive_req_len_sums.mutable_data_ptr<int32_t>(),
+      ne_token_table.mutable_data_ptr<int32_t>(), max_context_len,
+      row_indices.const_data_ptr<int64_t>(),
+      column_starts.mutable_data_ptr<int32_t>(),
+      n_gram_ids.mutable_data_ptr<int32_t>());
+}
@@ -554,3 +554,12 @@ void cp_gather_indexer_k_quant_cache(
                                                // quant_block_size * 4]
    const torch::stable::Tensor& block_table,   // [batch_size, num_blocks]
    const torch::stable::Tensor& cu_seq_lens);  // [batch_size + 1]
+
+// LongCat n-gram embedding index kernel (see ngram_embedding_kernels.cu).
+void ngram_compute_n_gram_ids(
+    int64_t ne_n, int64_t ne_k, torch::stable::Tensor& ne_weights,
+    torch::stable::Tensor& ne_mods,
+    torch::stable::Tensor& exclusive_ne_embedder_size_sums,
+    torch::stable::Tensor& exclusive_req_len_sums,
+    torch::stable::Tensor& ne_token_table, torch::stable::Tensor& row_indices,
+    torch::stable::Tensor& column_starts, torch::stable::Tensor& n_gram_ids);
@@ -598,9 +598,22 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
      "Tensor? initial_state_idx,"
      "Tensor? cu_chunk_seqlen,"
      "Tensor? last_chunk_indices) -> ()");
+
+  // LongCat n-gram embedding index kernel. All tensor args are marked mutable
+  // to match the (non-const) stable-Tensor& C++ signature; only ne_token_table
+  // and n_gram_ids are actually written in place.
+  ops.def(
+      "ngram_compute_n_gram_ids(int ne_n, int ne_k, Tensor(a!) ne_weights, "
+      "Tensor(b!) ne_mods, Tensor(c!) exclusive_ne_embedder_size_sums, "
+      "Tensor(d!) exclusive_req_len_sums, Tensor(e!) ne_token_table, "
+      "Tensor(f!) row_indices, Tensor(g!) column_starts, "
+      "Tensor(h!) n_gram_ids) -> ()");
 }

 STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
+  // LongCat n-gram embedding index kernel.
+  ops.impl("ngram_compute_n_gram_ids", TORCH_BOX(&ngram_compute_n_gram_ids));
+
  // Per-token group quantization
  ops.impl("per_token_group_fp8_quant", TORCH_BOX(&per_token_group_quant_fp8));
  ops.impl("per_token_group_fp8_quant_packed",
@@ -93,6 +93,28 @@ To address this, manually trigger a build on Buildkite to accomplish two objecti
 <img width="60%" alt="Buildkite new build popup" src="https://github.com/user-attachments/assets/3b07f71b-bb18-4ca3-aeaf-da0fe79d315f" />
 </p>

+You can also trigger this build from the command line with
+[`.buildkite/scripts/trigger-ci-build.sh`](../../../.buildkite/scripts/trigger-ci-build.sh)
+(dry-run by default; pass `--execute` to actually trigger it).
+
+## Test against PyTorch nightly
+
+The steps above test against a specific PyTorch RC/stable wheel pinned in the
+requirements files. To instead build and run the CI suite against the latest
+PyTorch **nightly** wheels, set the `TORCH_NIGHTLY=1` environment variable on
+the build (or apply the `ready-torch-nightly` label to the PR).
+
+When `TORCH_NIGHTLY=1`, the base CI image is built against PyTorch nightly
+(`image_build_torch_nightly.sh`, `PYTORCH_NIGHTLY=1`, CUDA 13.0) and tagged at
+the normal image tag, so the entire existing pipeline runs on nightly torch --
+there is no separate pipeline section to trigger. Combine it with `RUN_ALL=1`
+to run the full suite (the `ready-torch-nightly` label and
+`trigger-ci-build.sh --torch-nightly` both set this for you). This is the
+configuration to use for a scheduled "vLLM vs PyTorch nightly" run.
+
+Use `.buildkite/scripts/trigger-ci-build.sh --torch-nightly` to trigger it from
+the command line.
+
 ## Update all the different vLLM platforms

 Rather than attempting to update all vLLM platforms in a single pull request, it's more manageable
@@ -879,6 +879,55 @@ vllm serve Qwen/Qwen3-VL-30B-A3B-Instruct \

 Works with common video formats like MP4 when using OpenCV backends.

+#### GPU Video Decoding with DeepStream (NVDEC)
+
+By default vLLM decodes video on the CPU. On NVIDIA GPUs you can instead decode
+directly on the hardware video engine (NVDEC) with the DeepStream backend, which
+keeps decoding off the CPU and can significantly increase video throughput.
+
+Install the backend (Linux x86-64 only):
+
+```bash
+pip install vllm[deepstream]
+```
+
+The pip wheel bundles the DeepStream libraries but still relies on a few system
+packages that pip cannot install. On Ubuntu:
+
+```bash
+apt-get install -y \
+  gstreamer1.0-tools gstreamer1.0-plugins-base gstreamer1.0-plugins-good \
+  gstreamer1.0-plugins-bad gstreamer1.0-libav \
+  python3-gi python3-gst-1.0 libv4l-0 cuda-libraries-13-0
+```
+
+Select the backend either with an environment variable:
+
+```bash
+export VLLM_VIDEO_LOADER_BACKEND=deepstream
+vllm serve Qwen/Qwen3-VL-30B-A3B-Instruct
+```
+
+or per request via `--media-io-kwargs`:
+
+```bash
+vllm serve Qwen/Qwen3-VL-30B-A3B-Instruct \
+  --media-io-kwargs '{"video": {"backend": "deepstream"}}'
+```
+
+**Parameters:**
+
+- `pool_size`: Number of GPU decode workers in the process-wide decode pool
+  (clamped to `[1, 16]`). When unset it defaults to
+  `VLLM_MEDIA_LOADING_THREAD_COUNT` (default `8`). The pool is a singleton, so
+  the first request's value wins.
+
+```bash
+# Example: 12 decode workers
+vllm serve Qwen/Qwen3-VL-30B-A3B-Instruct \
+  --media-io-kwargs '{"video": {"backend": "deepstream", "pool_size": 12}}'
+```
+
 #### Pre-extracted Frame Sequences with `media_io_kwargs`

 When you extract video frames on the client side and send them as `video/jpeg` (base64-concatenated JPEG frames), you can preserve the original video metadata by using `media_io_kwargs` in your request. This enables more accurate video understanding by preserving temporal information that would otherwise be lost during client-side frame extraction.
@@ -351,6 +351,69 @@ print(response.choices[0].message.reasoning)
 print(response.choices[0].message.content)
 ```

+## Suppressing Reasoning Output
+
+You can suppress reasoning content from API responses using the `include_reasoning` parameter. When set to `false`, reasoning tokens are still generated (so model quality is unaffected) but excluded from the response. This reduces network traffic without changing inference behavior.
+
+The parameter is supported in both the Chat Completions API and the Responses API, for streaming and non-streaming requests.
+
+When `include_reasoning=false`, vLLM also suppresses per-token metadata (logprobs and token IDs) to prevent leaking reasoning content through decoded token text in logprob entries or raw token IDs.
+
+### Chat Completions API
+
+```python
+from openai import OpenAI
+
+client = OpenAI(base_url="http://localhost:8000/v1", api_key="EMPTY")
+model = client.models.list().data[0].id
+
+# Reasoning is included by default (include_reasoning=True)
+response = client.chat.completions.create(
+    model=model,
+    messages=[{"role": "user", "content": "What is 15 * 37?"}],
+    extra_body={"include_reasoning": False},
+)
+
+msg = response.choices[0].message
+assert msg.content  # Content is still present
+assert not getattr(msg, "reasoning", None)  # Reasoning is suppressed
+```
+
+Streaming works the same way, reasoning deltas are omitted from chunks:
+
+```python
+stream = client.chat.completions.create(
+    model=model,
+    messages=[{"role": "user", "content": "What is 15 * 37?"}],
+    stream=True,
+    extra_body={"include_reasoning": False},
+)
+
+for chunk in stream:
+    delta = chunk.choices[0].delta
+    # delta.reasoning will always be None
+    if delta.content:
+        print(delta.content, end="", flush=True)
+```
+
+### Responses API
+
+```python
+from openai import OpenAI
+
+client = OpenAI(base_url="http://localhost:8000/v1", api_key="EMPTY")
+
+response = client.responses.create(
+    model=client.models.list().data[0].id,
+    input="What is 15 * 37?",
+    include_reasoning=False,
+)
+
+# No "reasoning" items in output
+types = [item.type for item in response.output]
+assert "reasoning" not in types
+```
+
 ## Limitations

 - The reasoning content is only available for online serving's chat completion endpoint (`/v1/chat/completions`), Anthropic Messages API (`/v1/messages`) and the Responses API (`/v1/responses`).
@@ -405,6 +405,8 @@ th {
 | `Lfm2MoeForCausalLM` | LFM2MoE | `LiquidAI/LFM2-8B-A1B-preview`, etc. | ✅︎ | ✅︎ |
 | `LlamaForCausalLM` | Llama 3.1, Llama 3, Llama 2, LLaMA, Yi | `meta-llama/Meta-Llama-3.1-405B-Instruct`, `meta-llama/Meta-Llama-3.1-70B`, `meta-llama/Meta-Llama-3-70B-Instruct`, `meta-llama/Llama-2-70b-hf`, `01-ai/Yi-34B`, etc. | ✅︎ | ✅︎ |
 | `LongcatFlashForCausalLM` | LongCat-Flash | `meituan-longcat/LongCat-Flash-Chat`, `meituan-longcat/LongCat-Flash-Chat-FP8` | ✅︎ | ✅︎ |
+| `LongcatFlashNgramForCausalLM` | LongCat-Flash-Lite | `meituan-longcat/LongCat-Flash-Lite` | ✅︎ | ✅︎ |
+| `LongcatCausalLM` | LongCat-2.0 | `meituan-longcat/LongCat-2.0-FP8` | ✅︎ | ✅︎ |
 | `MambaForCausalLM` | Mamba | `state-spaces/mamba-130m-hf`, `state-spaces/mamba-790m-hf`, `state-spaces/mamba-2.8b-hf`, etc. | | ✅︎ |
 | `Mamba2ForCausalLM` | Mamba2 | `mistralai/Mamba-Codestral-7B-v0.1`, etc. | | ✅︎ |
 | `MellumForCausalLM` | Mellum 2 | `JetBrains/Mellum2-12B-A2.5B-Base`, etc. | | ✅︎ |
@@ -512,7 +514,7 @@ These models primarily accept the [`LLM.generate`](./generative_models.md#llmgen
 | `ChameleonForConditionalGeneration` | Chameleon | T + I | `facebook/chameleon-7b`, etc. | | ✅︎ |
 | `CheersForConditionalGeneration` | Cheers | T + I | `ai9stars/Cheers` | | ✅︎ |
 | `Cohere2VisionForConditionalGeneration` | Command A Vision, Command-A+ | T + I<sup>+</sup> | `CohereLabs/command-a-vision-07-2025`, `CohereLabs/command-a-plus-05-2026`, etc. | | ✅︎ |
-| `Cosmos3ForConditionalGeneration` | Cosmos3 (understanding tower) | T + I<sup>E+</sup> + V<sup>E+</sup> | `nvidia/Cosmos3-Nano` | | ✅︎ |
+| `Cosmos3ForConditionalGeneration` | Cosmos3 (understanding tower) | T + I<sup>E+</sup> + V<sup>E+</sup> | `nvidia/Cosmos3-Nano`, `nvidia/Cosmos3-Super` | | ✅︎ |
 | `DeepseekVLV2ForCausalLM` | DeepSeek-VL2 | T + I<sup>+</sup> | `deepseek-ai/deepseek-vl2-tiny`, `deepseek-ai/deepseek-vl2-small`, `deepseek-ai/deepseek-vl2`, etc. | | ✅︎ |
 | `DeepseekOCRForCausalLM` | DeepSeek-OCR | T + I<sup>+</sup> | `deepseek-ai/DeepSeek-OCR`, etc. | ✅︎ | ✅︎ |
 | `DeepseekOCR2ForCausalLM` | DeepSeek-OCR-2 | T + I<sup>+</sup> | `deepseek-ai/DeepSeek-OCR-2`, etc. | ✅︎ | ✅︎ |
@@ -2167,7 +2167,7 @@ checksum = "11d3d7f243d5c5a8b9bb5d6dd2b1602c0cb0b9db1621bafc7ed66e35ff9fe092"
 [[package]]
 name = "llm-multimodal"
 version = "1.7.1"
-source = "git+https://github.com/smg-project/llm-multimodal?rev=7d74582aeaf0e4086a44964382655d22f1af0686#7d74582aeaf0e4086a44964382655d22f1af0686"
+source = "git+https://github.com/smg-project/llm-multimodal?rev=c8a29dcc755139fdc26185f400ea48c6d6d48273#c8a29dcc755139fdc26185f400ea48c6d6d48273"
 dependencies = [
 "anyhow",
 "base64 0.22.1",
@@ -2473,6 +2473,7 @@ dependencies = [
 "portable-atomic",
 "portable-atomic-util",
 "rawpointer",
+ "serde",
 ]

 [[package]]
@@ -5093,6 +5094,7 @@ dependencies = [
 "llm-multimodal",
 "minijinja",
 "minijinja-contrib",
+ "ndarray 0.17.2",
 "oss-harmony",
 "paste",
 "reqwest",
@@ -53,12 +53,12 @@ hyper-util = { version = "0.1.20", features = [
 indexmap = "2.13.0"
 itertools = "0.14.0"
 libc = "0.2.177"
-llm-multimodal = { git = "https://github.com/smg-project/llm-multimodal", rev = "7d74582aeaf0e4086a44964382655d22f1af0686" }
+llm-multimodal = { git = "https://github.com/smg-project/llm-multimodal", rev = "c8a29dcc755139fdc26185f400ea48c6d6d48273" }
 mimalloc = "0.1.52"
 minijinja = { version = "2.0", features = ["unstable_machinery", "json", "builtins", "loader", "loop_controls", "preserve_order"] }
 minijinja-contrib = { version = "2.0", features = ["pycompat"] }
 native-tls-vendored = { package = "native-tls", version = "0.2.18", features = ["vendored"] }
-ndarray = { version = "0.16.1", features = ["serde"] }
+ndarray = { version = "0.17", features = ["serde"] }
 openai-harmony = { package = "oss-harmony", git = "https://github.com/oss-harmony/harmony", tag = "v0.0.11", default-features = false }
 openai-protocol = "1.6.0"
 openssl = "0.10"
@@ -42,6 +42,7 @@ anyhow.workspace = true
 bytes.workspace = true
 clap.workspace = true
 expect-test.workspace = true
+ndarray.workspace = true
 paste.workspace = true
 rmp-serde.workspace = true
 serial_test.workspace = true
@@ -10,7 +10,7 @@ use crate::backend::{
    NewChatOutputProcessorOptions,
 };
 use crate::error::Result;
-use crate::multimodal::MultimodalModelInfo;
+use crate::multimodal::{MultimodalConfigFiles, MultimodalModelInfo};
 use crate::output::{
    DefaultChatOutputProcessor, HarmonyChatOutputProcessor, validate_harmony_parser_overrides,
 };
@@ -46,8 +46,12 @@ impl HfChatBackend {
            MultimodalModelInfo::from_paths(
                model_id.clone(),
                (!model_type.is_empty()).then_some(model_type.to_string()),
-                files.config_path.as_deref(),
-                files.preprocessor_config_path.as_deref(),
+                MultimodalConfigFiles {
+                    config: files.config_path.as_deref(),
+                    preprocessor_config: files.preprocessor_config_path.as_deref(),
+                    video_preprocessor_config: files.video_preprocessor_config_path.as_deref(),
+                    processor_config: files.processor_config_path.as_deref(),
+                },
                tokenizer.clone(),
            )?
        };
@@ -139,8 +143,11 @@ pub(super) async fn load_model_backends(
 fn resolve_multimodal_render_info(
    info: Option<&MultimodalModelInfo>,
 ) -> Option<MultimodalRenderInfo> {
+    use llm_multimodal::Modality;
+
    info.map(|info| MultimodalRenderInfo {
-        placeholder_token: info.placeholder_token().to_string(),
+        image_token: info.placeholder_token(Modality::Image).map(str::to_string),
+        video_token: info.placeholder_token(Modality::Video).map(str::to_string),
    })
 }

@@ -192,6 +199,8 @@ mod tests {
            tokenizer_config_path: Some(tokenizer_config_path),
            generation_config_path: None,
            preprocessor_config_path: None,
+            video_preprocessor_config_path: None,
+            processor_config_path: None,
            chat_template_path: None,
            config_path: Some(config_path),
        }
@@ -1,5 +1,5 @@
 use thiserror::Error;
-use thiserror_ext::Macro;
+use thiserror_ext::{AsReport as _, Macro};

 type BoxedError = Box<dyn std::error::Error + Send + Sync>;

@@ -18,6 +18,8 @@ pub enum Error {
    UnsupportedMultimodalRenderer,
    #[error("unsupported multimodal content: {0}")]
    UnsupportedMultimodalContent(&'static str),
+    #[error("`{modality}` input is not supported by this model")]
+    UnsupportedModality { modality: String },
    #[error("multimodal preprocessing error: {0}")]
    Multimodal(#[message] String),
    #[error("{kind} parsing is not available for model `{model_id}`")]
@@ -80,11 +82,39 @@ impl Error {
        match self {
            Self::PromptTooLong { .. } => true,
            Self::Text(error) => error.is_request_validation_error(),
+            Self::UnsupportedMultimodalRenderer
+            | Self::UnsupportedMultimodalContent(_)
+            | Self::UnsupportedModality { .. } => true,
+
            _ => false,
        }
    }
 }

+impl From<llm_multimodal::MediaConnectorError> for Error {
+    fn from(error: llm_multimodal::MediaConnectorError) -> Self {
+        Self::Multimodal(error.to_report_string())
+    }
+}
+
+impl From<llm_multimodal::MultiModalError> for Error {
+    fn from(error: llm_multimodal::MultiModalError) -> Self {
+        Self::Multimodal(error.to_report_string())
+    }
+}
+
+impl From<llm_multimodal::TransformError> for Error {
+    fn from(error: llm_multimodal::TransformError) -> Self {
+        Self::Multimodal(error.to_report_string())
+    }
+}
+
+impl From<llm_multimodal::registry::ModelRegistryError> for Error {
+    fn from(error: llm_multimodal::registry::ModelRegistryError) -> Self {
+        Self::Multimodal(error.to_report_string())
+    }
+}
+
 /// Format the available-parser suffix used in user-facing error messages.
 fn available_parser_hint(available_names: &[String]) -> String {
    if available_names.is_empty() {
@@ -0,0 +1,446 @@
+//! Prompt placeholder expansion shared across modalities.
+
+use std::collections::{HashMap, VecDeque};
+
+use llm_multimodal::{Modality, PromptReplacement};
+use vllm_engine_core_client::protocol::multimodal::PlaceholderRange;
+use vllm_engine_core_client::protocol::tensor::WireTensor;
+
+use super::PreparedMedia;
+use crate::error::{Error, Result, bail_multimodal};
+
+/// One modality's queue of pending placeholder replacements for prompt
+/// expansion.
+struct ExpansionLane<'a> {
+    modality: Modality,
+    marker_token_id: u32,
+    embed_token_id: u32,
+    placeholder_token: String,
+    replacements: VecDeque<&'a PromptReplacement>,
+}
+
+impl<'a> ExpansionLane<'a> {
+    fn from_prepared(media: &'a PreparedMedia) -> Option<Self> {
+        if media.replacements.is_empty() {
+            return None;
+        }
+
+        Some(Self {
+            modality: media.modality,
+            marker_token_id: media.placeholder.marker_token_id,
+            embed_token_id: media.placeholder.embed_token_id,
+            placeholder_token: media.placeholder.token.clone(),
+            replacements: media.replacements.iter().collect(),
+        })
+    }
+}
+
+/// Replace rendered placeholder markers with model-specific replacement
+/// tokens across all modalities in one left-to-right pass.
+///
+/// Each prepared modality consumes its own marker occurrences in order,
+/// matching the original media-part order within that modality; markers of
+/// different modalities may interleave freely.
+///
+/// The returned ranges point into the already-expanded prompt, grouped per
+/// modality in item order.
+pub(super) fn expand_prompt_token_ids(
+    prompt_token_ids: &mut Vec<u32>,
+    prepared: &[PreparedMedia],
+) -> Result<HashMap<Modality, Vec<PlaceholderRange>>> {
+    let mut lanes = prepared.iter().filter_map(ExpansionLane::from_prepared).collect::<Vec<_>>();
+    if lanes.is_empty() {
+        return Ok(HashMap::new());
+    }
+
+    let replacement_growth = lanes
+        .iter()
+        .flat_map(|lane| lane.replacements.iter())
+        .fold(0usize, |total, replacement| {
+            total.saturating_add(replacement.tokens.len().saturating_sub(1))
+        });
+    let expanded_len = prompt_token_ids.len().saturating_add(replacement_growth);
+
+    let mut expanded = Vec::with_capacity(expanded_len);
+    let mut ranges = HashMap::<Modality, Vec<PlaceholderRange>>::new();
+
+    for &token in prompt_token_ids.iter() {
+        let lane = lanes
+            .iter_mut()
+            .find(|lane| lane.marker_token_id == token && !lane.replacements.is_empty());
+        let Some(lane) = lane else {
+            expanded.push(token);
+            continue;
+        };
+
+        let replacement = lane.replacements.pop_front().expect("lane queue is non-empty");
+        debug_assert_eq!(replacement.modality, lane.modality);
+        if replacement.tokens.is_empty() {
+            bail_multimodal!(
+                "placeholder token `{}` expanded to no tokens",
+                lane.placeholder_token
+            );
+        }
+
+        let replacement_len = replacement.tokens.len();
+        let is_embed = {
+            let mask = replacement
+                .tokens
+                .iter()
+                .map(|&token| token as u32 == lane.embed_token_id)
+                .collect::<Vec<_>>();
+            WireTensor::from_bool(vec![replacement_len], mask).map_err(Error::Multimodal)?
+        };
+
+        let expanded_offset = expanded.len();
+        expanded.extend(replacement.tokens.iter().map(|&token| token as u32));
+        ranges.entry(lane.modality).or_default().push(PlaceholderRange {
+            offset: expanded_offset,
+            length: replacement_len,
+            is_embed: Some(is_embed),
+        });
+    }
+
+    for lane in &lanes {
+        if !lane.replacements.is_empty() {
+            bail_multimodal!(
+                "placeholder token `{}` was not found in tokenized prompt for {} remaining `{}` item(s)",
+                lane.placeholder_token,
+                lane.replacements.len(),
+                lane.modality
+            );
+        }
+    }
+
+    *prompt_token_ids = expanded;
+
+    Ok(ranges)
+}
+
+#[cfg(test)]
+mod tests {
+    use llm_multimodal::TokenId;
+    use vllm_engine_core_client::protocol::tensor::WireArrayData;
+
+    use super::super::tests::{
+        LLAMA4_IMAGE_END_ID, LLAMA4_IMAGE_ID, LLAMA4_IMAGE_START_ID, LLAMA4_PATCH_ID,
+        LLAMA4_TILE_X_SEPARATOR_ID, LLAMA4_TILE_Y_SEPARATOR_ID, QWEN3_IMAGE_PAD_ID,
+        QWEN3_VIDEO_PAD_ID,
+    };
+    use super::super::{PreparedMedia, ResolvedPlaceholder};
+    use super::*;
+
+    /// Build prepared media directly from placeholder token IDs.
+    fn prepared_media(
+        modality: Modality,
+        placeholder_token: &str,
+        marker_token_id: u32,
+        embed_token_id: u32,
+        replacements: Vec<PromptReplacement>,
+    ) -> PreparedMedia {
+        PreparedMedia {
+            modality,
+            placeholder: ResolvedPlaceholder {
+                token: placeholder_token.to_string(),
+                marker_token_id,
+                embed_token_id,
+            },
+            replacements,
+            items: Vec::new(),
+        }
+    }
+
+    /// Llama4 image prepared media: the `<|image|>` marker expands to
+    /// sequences whose embed positions are the `<|patch|>` tokens.
+    fn llama4_prepared(replacements: Vec<PromptReplacement>) -> PreparedMedia {
+        prepared_media(
+            Modality::Image,
+            "<|image|>",
+            LLAMA4_IMAGE_ID,
+            LLAMA4_PATCH_ID,
+            replacements,
+        )
+    }
+
+    fn qwen3_image_prepared(replacements: Vec<PromptReplacement>) -> PreparedMedia {
+        prepared_media(
+            Modality::Image,
+            "<|image_pad|>",
+            QWEN3_IMAGE_PAD_ID,
+            QWEN3_IMAGE_PAD_ID,
+            replacements,
+        )
+    }
+
+    fn qwen3_video_prepared(replacements: Vec<PromptReplacement>) -> PreparedMedia {
+        prepared_media(
+            Modality::Video,
+            "<|video_pad|>",
+            QWEN3_VIDEO_PAD_ID,
+            QWEN3_VIDEO_PAD_ID,
+            replacements,
+        )
+    }
+
+    fn llama4_single_tile_replacement() -> PromptReplacement {
+        PromptReplacement::sequence(
+            Modality::Image,
+            "<|image|>",
+            vec![
+                LLAMA4_IMAGE_START_ID as TokenId,
+                LLAMA4_IMAGE_ID as TokenId,
+                LLAMA4_PATCH_ID as TokenId,
+                LLAMA4_PATCH_ID as TokenId,
+                LLAMA4_IMAGE_END_ID as TokenId,
+            ],
+        )
+    }
+
+    fn llama4_multi_tile_replacement() -> PromptReplacement {
+        PromptReplacement::sequence(
+            Modality::Image,
+            "<|image|>",
+            vec![
+                LLAMA4_IMAGE_START_ID as TokenId,
+                LLAMA4_PATCH_ID as TokenId,
+                LLAMA4_TILE_X_SEPARATOR_ID as TokenId,
+                LLAMA4_PATCH_ID as TokenId,
+                LLAMA4_TILE_Y_SEPARATOR_ID as TokenId,
+                LLAMA4_IMAGE_ID as TokenId,
+                LLAMA4_PATCH_ID as TokenId,
+                LLAMA4_IMAGE_END_ID as TokenId,
+            ],
+        )
+    }
+
+    fn assert_bool_mask(range: &PlaceholderRange, expected: &[bool]) {
+        let tensor = range.is_embed.as_ref().expect("is_embed mask");
+        assert_eq!(tensor.dtype, "bool");
+        assert_eq!(tensor.shape, vec![expected.len()]);
+        assert_eq!(
+            tensor.data,
+            WireArrayData::RawView(expected.iter().map(|value| u8::from(*value)).collect())
+        );
+    }
+
+    #[test]
+    fn expand_prompt_tokens_marks_only_llama4_patch_tokens_as_embed() {
+        let mut prompt_token_ids = vec![1, LLAMA4_IMAGE_ID, 2];
+        let prepared = vec![llama4_prepared(vec![llama4_multi_tile_replacement()])];
+
+        let ranges = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap();
+        let ranges = &ranges[&Modality::Image];
+
+        assert_eq!(
+            prompt_token_ids,
+            vec![
+                1,
+                LLAMA4_IMAGE_START_ID,
+                LLAMA4_PATCH_ID,
+                LLAMA4_TILE_X_SEPARATOR_ID,
+                LLAMA4_PATCH_ID,
+                LLAMA4_TILE_Y_SEPARATOR_ID,
+                LLAMA4_IMAGE_ID,
+                LLAMA4_PATCH_ID,
+                LLAMA4_IMAGE_END_ID,
+                2,
+            ]
+        );
+        assert_eq!(ranges[0].offset, 1);
+        assert_eq!(ranges[0].length, 8);
+        assert_bool_mask(
+            &ranges[0],
+            &[false, true, false, true, false, false, true, false],
+        );
+    }
+
+    #[test]
+    fn expand_prompt_tokens_errors_when_placeholder_missing() {
+        let mut prompt_token_ids = vec![1, 2, 3];
+        let prepared = vec![llama4_prepared(vec![llama4_single_tile_replacement()])];
+
+        let error = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap_err();
+
+        assert!(matches!(error, Error::Multimodal(message) if message.contains("not found")));
+    }
+
+    #[test]
+    fn expand_prompt_tokens_ignores_empty_replacements() {
+        let mut prompt_token_ids = vec![1, LLAMA4_IMAGE_ID, 2];
+        let original_prompt_token_ids = prompt_token_ids.clone();
+        let prepared = vec![llama4_prepared(Vec::new())];
+
+        let ranges = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap();
+
+        assert!(ranges.is_empty());
+        assert_eq!(prompt_token_ids, original_prompt_token_ids);
+    }
+
+    #[test]
+    fn expand_prompt_tokens_leaves_prompt_unchanged_when_later_placeholder_missing() {
+        let mut prompt_token_ids = vec![1, LLAMA4_IMAGE_ID, 2];
+        let original_prompt_token_ids = prompt_token_ids.clone();
+        let prepared = vec![llama4_prepared(vec![
+            llama4_single_tile_replacement(),
+            llama4_single_tile_replacement(),
+        ])];
+
+        let error = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap_err();
+
+        assert!(matches!(error, Error::Multimodal(message) if message.contains("not found")));
+        assert_eq!(prompt_token_ids, original_prompt_token_ids);
+    }
+
+    #[test]
+    fn expand_prompt_tokens_errors_when_replacement_is_empty() {
+        let mut prompt_token_ids = vec![1, LLAMA4_IMAGE_ID, 2];
+        let original_prompt_token_ids = prompt_token_ids.clone();
+        let prepared = vec![llama4_prepared(vec![PromptReplacement::sequence(
+            Modality::Image,
+            "<|image|>",
+            Vec::new(),
+        )])];
+
+        let error = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap_err();
+
+        assert!(
+            matches!(error, Error::Multimodal(message) if message.contains("expanded to no tokens"))
+        );
+        assert_eq!(prompt_token_ids, original_prompt_token_ids);
+    }
+
+    #[test]
+    fn expand_prompt_tokens_skips_llama4_image_marker_inside_replacement() {
+        let mut prompt_token_ids = vec![1, LLAMA4_IMAGE_ID, 2, LLAMA4_IMAGE_ID, 3];
+        let prepared = vec![llama4_prepared(vec![
+            llama4_single_tile_replacement(),
+            llama4_single_tile_replacement(),
+        ])];
+
+        let ranges = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap();
+        let ranges = &ranges[&Modality::Image];
+
+        assert_eq!(
+            prompt_token_ids,
+            vec![
+                1,
+                LLAMA4_IMAGE_START_ID,
+                LLAMA4_IMAGE_ID,
+                LLAMA4_PATCH_ID,
+                LLAMA4_PATCH_ID,
+                LLAMA4_IMAGE_END_ID,
+                2,
+                LLAMA4_IMAGE_START_ID,
+                LLAMA4_IMAGE_ID,
+                LLAMA4_PATCH_ID,
+                LLAMA4_PATCH_ID,
+                LLAMA4_IMAGE_END_ID,
+                3,
+            ]
+        );
+        assert_eq!(ranges[0].offset, 1);
+        assert_eq!(ranges[0].length, 5);
+        assert_bool_mask(&ranges[0], &[false, false, true, true, false]);
+        assert_eq!(ranges[1].offset, 7);
+        assert_eq!(ranges[1].length, 5);
+        assert_bool_mask(&ranges[1], &[false, false, true, true, false]);
+    }
+
+    #[test]
+    fn expand_prompt_tokens_interleaves_image_and_video_prepared_media() {
+        let mut prompt_token_ids = vec![
+            1,
+            QWEN3_IMAGE_PAD_ID,
+            2,
+            QWEN3_VIDEO_PAD_ID,
+            3,
+            QWEN3_IMAGE_PAD_ID,
+            4,
+        ];
+        let prepared = vec![
+            qwen3_image_prepared(vec![
+                PromptReplacement::repeated(
+                    Modality::Image,
+                    "<|image_pad|>",
+                    QWEN3_IMAGE_PAD_ID as TokenId,
+                    2,
+                ),
+                PromptReplacement::repeated(
+                    Modality::Image,
+                    "<|image_pad|>",
+                    QWEN3_IMAGE_PAD_ID as TokenId,
+                    3,
+                ),
+            ]),
+            qwen3_video_prepared(vec![PromptReplacement::repeated(
+                Modality::Video,
+                "<|video_pad|>",
+                QWEN3_VIDEO_PAD_ID as TokenId,
+                4,
+            )]),
+        ];
+
+        let ranges = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap();
+
+        assert_eq!(
+            prompt_token_ids,
+            vec![
+                1,
+                QWEN3_IMAGE_PAD_ID,
+                QWEN3_IMAGE_PAD_ID,
+                2,
+                QWEN3_VIDEO_PAD_ID,
+                QWEN3_VIDEO_PAD_ID,
+                QWEN3_VIDEO_PAD_ID,
+                QWEN3_VIDEO_PAD_ID,
+                3,
+                QWEN3_IMAGE_PAD_ID,
+                QWEN3_IMAGE_PAD_ID,
+                QWEN3_IMAGE_PAD_ID,
+                4,
+            ]
+        );
+
+        let image_ranges = &ranges[&Modality::Image];
+        assert_eq!(image_ranges[0].offset, 1);
+        assert_eq!(image_ranges[0].length, 2);
+        assert_bool_mask(&image_ranges[0], &[true, true]);
+        assert_eq!(image_ranges[1].offset, 9);
+        assert_eq!(image_ranges[1].length, 3);
+        assert_bool_mask(&image_ranges[1], &[true, true, true]);
+
+        let video_ranges = &ranges[&Modality::Video];
+        assert_eq!(video_ranges[0].offset, 4);
+        assert_eq!(video_ranges[0].length, 4);
+        assert_bool_mask(&video_ranges[0], &[true, true, true, true]);
+    }
+
+    #[test]
+    fn expand_prompt_tokens_error_names_modality_with_leftover_replacements() {
+        let mut prompt_token_ids = vec![1, QWEN3_IMAGE_PAD_ID, 2];
+        let original_prompt_token_ids = prompt_token_ids.clone();
+        let prepared = vec![
+            qwen3_image_prepared(vec![PromptReplacement::repeated(
+                Modality::Image,
+                "<|image_pad|>",
+                QWEN3_IMAGE_PAD_ID as TokenId,
+                2,
+            )]),
+            qwen3_video_prepared(vec![PromptReplacement::repeated(
+                Modality::Video,
+                "<|video_pad|>",
+                QWEN3_VIDEO_PAD_ID as TokenId,
+                4,
+            )]),
+        ];
+
+        let error = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap_err();
+
+        assert!(matches!(
+            error,
+            Error::Multimodal(message)
+                if message.contains("<|video_pad|>") && message.contains("`video`")
+        ));
+        assert_eq!(prompt_token_ids, original_prompt_token_ids);
+    }
+}
@@ -0,0 +1,141 @@
+//! Image-modality preparation: batch preprocessing and per-item feature
+//! build.
+
+use std::sync::Arc;
+
+use itertools::izip;
+use llm_multimodal::{FieldLayout, ImageFrame, Modality, PreprocessedEncoderInputs};
+use vllm_engine_core_client::protocol::dtype::ModelDtype;
+use vllm_engine_core_client::protocol::multimodal::{
+    MmBatchedField, MmField, MmFieldElem, MmFlatField, MmKwargsItem, MmSharedField, MmSlice,
+    SliceSpec,
+};
+
+use super::{ModalitySupport, MultimodalModelInfo, PreparedItem, PreparedMedia, tensor};
+use crate::error::{Error, Result, bail_multimodal, multimodal};
+
+impl MultimodalModelInfo {
+    /// Preprocess all fetched image frames as one batch and build per-item
+    /// features.
+    pub(super) async fn prepare_images(
+        &self,
+        frames: Vec<Arc<ImageFrame>>,
+        uuids: Vec<Option<String>>,
+        model_dtype: ModelDtype,
+    ) -> Result<PreparedMedia> {
+        let support = self.image.as_ref().ok_or_else(|| Error::UnsupportedModality {
+            modality: Modality::Image.to_string(),
+        })?;
+        let preprocessed = self.preprocess_images(support, &frames).await?;
+        let replacements =
+            self.spec
+                .prompt_replacements_for(&self.context, &preprocessed, Modality::Image)?;
+        if replacements.len() != frames.len() {
+            bail_multimodal!(
+                "number of image prompt replacements {} does not match number of images {}",
+                replacements.len(),
+                frames.len()
+            );
+        }
+        let items = self.build_image_items(preprocessed, &frames, uuids, model_dtype)?;
+
+        Ok(PreparedMedia {
+            modality: Modality::Image,
+            placeholder: support.placeholder.clone(),
+            replacements,
+            items,
+        })
+    }
+
+    /// Preprocess fetched image frames with the model's resolved vision
+    /// processor.
+    ///
+    /// The processor work is CPU-heavy relative to request wiring, so it runs
+    /// in a blocking task and returns owned tensors ready for wire
+    /// conversion.
+    async fn preprocess_images(
+        &self,
+        support: &ModalitySupport,
+        image_frames: &[Arc<ImageFrame>],
+    ) -> Result<PreprocessedEncoderInputs> {
+        let config = support.config.clone();
+        let processor = support.processor;
+        let images = image_frames.iter().map(|frame| frame.data().clone()).collect::<Vec<_>>();
+
+        // TODO: is it still necessary given that we've already in a dedicated runtime?
+        tokio::task::spawn_blocking(move || Ok(processor.preprocess(&images, &config)?))
+            .await
+            .map_err(|error| multimodal!("image preprocessing task failed: {error}"))?
+    }
+
+    /// Convert one batch of preprocessed image tensors into per-item engine
+    /// kwargs.
+    ///
+    /// Tensor fields are sliced per item according to the model spec's field
+    /// layout declarations.
+    fn build_image_items(
+        &self,
+        preprocessed: PreprocessedEncoderInputs,
+        frames: &[Arc<ImageFrame>],
+        uuids: Vec<Option<String>>,
+        model_dtype: ModelDtype,
+    ) -> Result<Vec<PreparedItem>> {
+        let len = frames.len();
+        let tensors = tensor::collect_tensors(preprocessed, "pixel_values", model_dtype)?;
+
+        let mut items = Vec::with_capacity(len);
+        for (index, (frame, uuid)) in izip!(frames, uuids).enumerate() {
+            let mut data = MmKwargsItem::new();
+            for (key, tensor) in &tensors {
+                let keep_on_cpu = self.spec.keep_on_cpu_keys.contains(key);
+                let (value, field) = match self.spec.field_layouts.get(key) {
+                    Some(FieldLayout::Batched) => (
+                        tensor.batched_value_at(index)?,
+                        MmField::Batched(MmBatchedField { keep_on_cpu }),
+                    ),
+                    Some(FieldLayout::Flat { sizes_key }) => {
+                        let sizes = tensors.get(sizes_key).ok_or_else(|| {
+                            multimodal!("flat tensor sizes key `{sizes_key}` is missing")
+                        })?;
+                        let (start, end) = tensor::flat_range_for_index(sizes, sizes_key, index)?;
+                        (
+                            tensor.flat_value_range(start, end)?,
+                            MmField::Flat(MmFlatField {
+                                slices: vec![MmSlice::Slice(SliceSpec {
+                                    start: Some(0),
+                                    stop: Some((end - start) as isize),
+                                    step: None,
+                                })],
+                                dim: 0,
+                                keep_on_cpu,
+                            }),
+                        )
+                    }
+                    None => (
+                        tensor.clone(),
+                        MmField::Shared(MmSharedField {
+                            batch_size: len,
+                            keep_on_cpu,
+                        }),
+                    ),
+                };
+
+                data.insert(
+                    key.clone(),
+                    MmFieldElem {
+                        data: Some(value.try_into()?),
+                        field,
+                    },
+                );
+            }
+
+            items.push(PreparedItem {
+                data,
+                hash: frame.hash.clone(),
+                uuid,
+            });
+        }
+
+        Ok(items)
+    }
+}
@@ -1,7 +1,7 @@
 use std::collections::HashMap;

 use half::{bf16, f16};
-use llm_multimodal::{ModelSpecificValue, PreprocessedEncoderInputs as PreprocessedImages};
+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};
@@ -25,25 +25,31 @@ pub(super) enum KwargValue {
    Passthrough(ProtocolKwargValue),
 }

-/// Collect `pixel_values` and model-specific outputs into one tensor map.
+/// Collect the primary encoder input and model-specific outputs into one
+/// tensor map.
+///
+/// `primary_key` names the encoder-input tensor as the model's forward kwargs
+/// expect it (e.g. `pixel_values` for images, `pixel_values_videos` for
+/// videos).
 pub(super) fn collect_tensors(
-    preprocessed: PreprocessedImages,
+    preprocessed: PreprocessedEncoderInputs,
+    primary_key: &str,
    float_dtype: ModelDtype,
 ) -> Result<HashMap<String, KwargValue>> {
-    let PreprocessedImages {
+    let PreprocessedEncoderInputs {
        encoder_input,
        model_specific,
        ..
    } = preprocessed;

-    let pixel_values = {
+    let primary_value = {
        let shape = encoder_input.shape().to_vec();
        let data = encoder_input.into_iter().collect();
        KwargValue::from_f32_tensor(data, shape, float_dtype)?
    };

    let mut tensors = HashMap::new();
-    tensors.insert("pixel_values".to_string(), pixel_values);
+    tensors.insert(primary_key.to_string(), primary_value);
    for (key, value) in model_specific {
        tensors.insert(key, KwargValue::from_model_specific(value, float_dtype)?);
    }
@@ -124,10 +130,22 @@ impl TryFrom<KwargValue> for ProtocolKwargValue {
 }

 impl KwargValue {
-    /// Extract one image from a batched tensor field.
+    /// 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::Passthrough(_) => None,
+        }
+    }
+
+    /// Extract one media item from a batched tensor field.
    ///
-    /// Batched fields use their first axis as image index and drop that axis in
-    /// the per-feature value, matching vLLM's batched-field semantics.
+    /// Batched fields use their first axis as media-item index and drop that
+    /// axis in the per-feature value, matching vLLM's batched-field semantics.
    pub(super) fn batched_value_at(&self, index: usize) -> Result<Self> {
        match self {
            Self::F32Tensor { data, shape } => {
@@ -154,9 +172,9 @@ impl KwargValue {
        }
    }

-    /// Extract one image's variable-length range from a flat tensor field.
+    /// Extract one media item's variable-length range from a flat tensor field.
    ///
-    /// Flat fields keep the first axis as the sliced length for this image.
+    /// Flat fields keep the first axis as the sliced length for this item.
    pub(super) fn flat_value_range(&self, start: usize, end: usize) -> Result<Self> {
        match self {
            Self::F32Tensor { data, shape } => {
@@ -184,10 +202,10 @@ impl KwargValue {
    }
 }

-/// Compute the first-axis range for one image in a flat tensor.
+/// Compute the first-axis range for one media item in a flat tensor.
 ///
 /// `sizes_key` names a companion tensor whose entries are cumulative slice
-/// sizes per image.
+/// sizes per media item.
 pub(super) fn flat_range_for_index(
    sizes: &KwargValue,
    sizes_key: &str,
@@ -195,7 +213,7 @@ pub(super) fn flat_range_for_index(
 ) -> Result<(usize, usize)> {
    let sizes = tensor_as_usize_vec(sizes)?;
    let size = *sizes.get(index).ok_or_else(|| {
-        multimodal!("flat tensor sizes key `{sizes_key}` has no entry for image {index}")
+        multimodal!("flat tensor sizes key `{sizes_key}` has no entry for media item {index}")
    })?;
    let start = sizes[..index].iter().sum::<usize>();
    Ok((start, start + size))
@@ -0,0 +1,316 @@
+//! Video-modality preparation: per-clip preprocessing, config resolution,
+//! and per-item feature build.
+
+use std::sync::Arc;
+
+use itertools::izip;
+use llm_multimodal::{FieldLayout, Modality, PreprocessedEncoderInputs, VideoClip};
+use thiserror_ext::AsReport as _;
+use tracing::warn;
+use vllm_engine_core_client::protocol::dtype::ModelDtype;
+use vllm_engine_core_client::protocol::multimodal::{
+    MmBatchedField, MmField, MmFieldElem, MmFlatField, MmKwargsItem, MmSharedField, MmSlice,
+    SliceSpec,
+};
+
+use super::{ModalitySupport, MultimodalModelInfo, PreparedItem, PreparedMedia, tensor};
+use crate::error::{Error, Result, bail_multimodal, multimodal};
+
+/// Forward-kwargs name of the primary video encoder input.
+///
+/// Video-capable vLLM models read `pixel_values_videos` alongside
+/// `video_grid_thw`, mirroring the HF processor output naming.
+const VIDEO_PRIMARY_KEY: &str = "pixel_values_videos";
+
+impl MultimodalModelInfo {
+    /// Preprocess fetched video clips one at a time and build per-item
+    /// features.
+    ///
+    /// Unlike images, each clip runs through the preprocessor independently
+    /// (a batch of one), so its tensors are complete per item and need no
+    /// cross-item slicing.
+    pub(super) async fn prepare_videos(
+        &self,
+        clips: Vec<Arc<VideoClip>>,
+        uuids: Vec<Option<String>>,
+        model_dtype: ModelDtype,
+    ) -> Result<PreparedMedia> {
+        let support = self.video.as_ref().ok_or_else(|| Error::UnsupportedModality {
+            modality: Modality::Video.to_string(),
+        })?;
+        let mut replacements = Vec::with_capacity(clips.len());
+        let mut items = Vec::with_capacity(clips.len());
+
+        for (clip, uuid) in izip!(&clips, uuids) {
+            let preprocessed = self.preprocess_video_clip(support, Arc::clone(clip)).await?;
+            let mut clip_replacements =
+                self.spec
+                    .prompt_replacements_for(&self.context, &preprocessed, Modality::Video)?;
+            if clip_replacements.len() != 1 {
+                bail_multimodal!(
+                    "expected exactly one prompt replacement per video clip, got {}",
+                    clip_replacements.len()
+                );
+            }
+            replacements.push(clip_replacements.pop().unwrap());
+            items.push(self.build_video_item(
+                preprocessed,
+                clip.hash.clone(),
+                uuid,
+                model_dtype,
+            )?);
+        }
+
+        Ok(PreparedMedia {
+            modality: Modality::Video,
+            placeholder: support.placeholder.clone(),
+            replacements,
+            items,
+        })
+    }
+
+    /// Preprocess one decoded video clip with the model's resolved vision
+    /// processor.
+    async fn preprocess_video_clip(
+        &self,
+        support: &ModalitySupport,
+        clip: Arc<VideoClip>,
+    ) -> Result<PreprocessedEncoderInputs> {
+        let config = support.config.clone();
+        let processor = support.processor;
+
+        tokio::task::spawn_blocking(move || {
+            // Prefer the borrowed-RGB fast path, which avoids materializing a
+            // `DynamicImage` per sampled frame after media decode.
+            if let Some(rgb_video) = clip.rgb_video() {
+                match rgb_video.frame_refs() {
+                    Ok(frame_refs) => match processor.preprocess_video_rgb(&frame_refs, &config) {
+                        Ok(preprocessed) => return Ok(preprocessed),
+                        Err(error) => warn!(
+                            error = %error.as_report(),
+                            "RGB video preprocessing fast path failed; falling back to materialized frames"
+                        ),
+                    },
+                    Err(error) => warn!(
+                        error,
+                        "RGB video frame refs are invalid; falling back to materialized frames"
+                    ),
+                }
+            }
+
+            let frames = clip.materialized_frames().map_err(|error| multimodal!("{error}"))?;
+            Ok(processor.preprocess_video(&frames, &config)?)
+        })
+        .await
+        .map_err(|error| multimodal!("video preprocessing task failed: {error}"))?
+    }
+
+    /// Convert one preprocessed video clip into engine kwargs.
+    ///
+    /// The clip is a batch of one, so no per-item slicing is required: the
+    /// primary tensor ships as a full-range flat field (the engine re-batches
+    /// flat fields by concatenating along the declared dim, matching vLLM's
+    /// `flat_from_sizes` treatment of video patches), and batched metadata
+    /// tensors drop their singleton batch axis.
+    fn build_video_item(
+        &self,
+        preprocessed: PreprocessedEncoderInputs,
+        hash: String,
+        uuid: Option<String>,
+        model_dtype: ModelDtype,
+    ) -> Result<PreparedItem> {
+        let tensors = tensor::collect_tensors(preprocessed, VIDEO_PRIMARY_KEY, model_dtype)?;
+
+        let mut data = MmKwargsItem::new();
+        for (key, tensor) in tensors {
+            let keep_on_cpu = self.spec.keep_on_cpu_keys.contains(&key);
+            let (value, field) = if key == VIDEO_PRIMARY_KEY {
+                let len = tensor
+                    .first_dim()
+                    .ok_or_else(|| multimodal!("video encoder input `{key}` is not a tensor"))?;
+                (
+                    tensor,
+                    MmField::Flat(MmFlatField {
+                        slices: vec![MmSlice::Slice(SliceSpec {
+                            start: Some(0),
+                            stop: Some(len as isize),
+                            step: None,
+                        })],
+                        dim: 0,
+                        keep_on_cpu,
+                    }),
+                )
+            } else if matches!(
+                self.spec.field_layouts.get(&key),
+                Some(FieldLayout::Batched)
+            ) {
+                (
+                    tensor.batched_value_at(0)?,
+                    MmField::Batched(MmBatchedField { keep_on_cpu }),
+                )
+            } else {
+                (
+                    tensor,
+                    MmField::Shared(MmSharedField {
+                        batch_size: 1,
+                        keep_on_cpu,
+                    }),
+                )
+            };
+
+            data.insert(
+                key,
+                MmFieldElem {
+                    data: Some(value.try_into()?),
+                    field,
+                },
+            );
+        }
+
+        Ok(PreparedItem { data, hash, uuid })
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use std::collections::HashMap;
+    use std::sync::Arc;
+
+    use llm_multimodal::ModelSpecificValue;
+    use ndarray::ArrayD;
+    use vllm_engine_core_client::protocol::multimodal::MmKwargValue;
+
+    use super::super::tests::{
+        QWEN3_IMAGE_PAD_ID, QWEN3_VIDEO_PAD_ID, qwen3_vl_info, qwen3_vl_tokenizer,
+    };
+    use super::super::{MultimodalConfigFiles, MultimodalModelInfo};
+    use super::*;
+
+    #[test]
+    fn from_paths_resolves_video_config_from_dedicated_file_or_processor_config() {
+        let dir = tempfile::tempdir().unwrap();
+        let config_path = dir.path().join("config.json");
+        std::fs::write(
+            &config_path,
+            serde_json::json!({
+                "model_type": "qwen3_vl",
+                "image_token_id": QWEN3_IMAGE_PAD_ID,
+                "video_token_id": QWEN3_VIDEO_PAD_ID,
+            })
+            .to_string(),
+        )
+        .unwrap();
+
+        let info_for = |files: MultimodalConfigFiles<'_>| {
+            MultimodalModelInfo::from_paths(
+                "qwen3-vl-test".to_string(),
+                Some("qwen3_vl".to_string()),
+                files,
+                Arc::new(qwen3_vl_tokenizer()),
+            )
+        };
+
+        // Dedicated video preprocessor config file.
+        let video_config_path = dir.path().join("video_preprocessor_config.json");
+        std::fs::write(&video_config_path, r#"{"size":{"shortest_edge":128}}"#).unwrap();
+        let info = info_for(MultimodalConfigFiles {
+            config: Some(&config_path),
+            video_preprocessor_config: Some(&video_config_path),
+            ..Default::default()
+        })
+        .unwrap()
+        .unwrap();
+        assert!(info.video.is_some());
+
+        // `video_processor` section of the combined processor config.
+        let processor_config_path = dir.path().join("processor_config.json");
+        std::fs::write(
+            &processor_config_path,
+            r#"{"video_processor":{"size":{"shortest_edge":128}}}"#,
+        )
+        .unwrap();
+        let info = info_for(MultimodalConfigFiles {
+            config: Some(&config_path),
+            processor_config: Some(&processor_config_path),
+            ..Default::default()
+        })
+        .unwrap()
+        .unwrap();
+        assert!(info.video.is_some());
+
+        // Neither source: video support still resolves on the image config.
+        let info = info_for(MultimodalConfigFiles {
+            config: Some(&config_path),
+            ..Default::default()
+        })
+        .unwrap()
+        .unwrap();
+        assert!(info.video.is_some());
+
+        // Malformed dedicated file is a real error, not a silent fallback.
+        std::fs::write(&video_config_path, r#"{"size""#).unwrap();
+        let error = match info_for(MultimodalConfigFiles {
+            config: Some(&config_path),
+            video_preprocessor_config: Some(&video_config_path),
+            ..Default::default()
+        }) {
+            Err(error) => error,
+            Ok(_) => panic!("malformed video preprocessor config should fail"),
+        };
+        assert!(matches!(
+            error,
+            Error::Multimodal(message)
+                if message.contains("failed to parse video_preprocessor_config.json")
+        ));
+    }
+
+    #[test]
+    fn build_video_item_names_primary_tensor_and_layouts() {
+        let info = qwen3_vl_info();
+        // One clip flattened to 6 patches with 4 features each.
+        let preprocessed = PreprocessedEncoderInputs {
+            encoder_input: ArrayD::zeros(vec![6, 4]),
+            feature_token_counts: vec![6],
+            item_sizes: vec![(32, 32)],
+            model_specific: HashMap::from([
+                (
+                    "video_grid_thw".to_string(),
+                    ModelSpecificValue::int_2d(vec![1, 2, 3], 1, 3),
+                ),
+                (
+                    "patches_per_video".to_string(),
+                    ModelSpecificValue::int_1d(vec![6]),
+                ),
+            ]),
+        };
+
+        let item = info
+            .build_video_item(
+                preprocessed,
+                "<hash>".to_string(),
+                None,
+                ModelDtype::Float32,
+            )
+            .unwrap();
+
+        let primary = &item.data[VIDEO_PRIMARY_KEY];
+        assert!(matches!(
+            &primary.field,
+            MmField::Flat(MmFlatField { slices, dim: 0, .. })
+                if matches!(
+                    slices.as_slice(),
+                    [MmSlice::Slice(SliceSpec { start: Some(0), stop: Some(6), step: None })]
+                )
+        ));
+
+        // Batched metadata drops its singleton batch axis per item.
+        let grid = &item.data["video_grid_thw"];
+        assert!(matches!(&grid.field, MmField::Batched(_)));
+        let MmKwargValue::Tensor(grid_tensor) = grid.data.as_ref().unwrap() else {
+            panic!("expected tensor value for video_grid_thw");
+        };
+        assert_eq!(grid_tensor.shape, vec![3]);
+
+        assert_eq!(item.hash, "<hash>");
+    }
+}
@@ -31,9 +31,14 @@ pub use template::{load_chat_template, resolve_chat_template};

 pub use self::format::ChatTemplateContentFormatOption;

-#[derive(Debug, Clone)]
+/// Template-visible placeholder tokens per supported modality.
+///
+/// A `None` token means the loaded model does not support that modality, and
+/// content parts of that modality are rejected during rendering.
+#[derive(Debug, Clone, Default)]
 pub struct MultimodalRenderInfo {
-    pub placeholder_token: String,
+    pub image_token: Option<String>,
+    pub video_token: Option<String>,
 }

 /// Hugging Face chat-template renderer backed by the local Jinja chat-template
@@ -254,6 +259,7 @@ enum TemplateContent {
 enum TemplateContentPart {
    Text { text: String },
    Image,
+    Video,
 }

 #[derive(Debug, Serialize)]
@@ -417,9 +423,17 @@ fn to_template_openai_content(
                }
                // All multimodal contents are normalized to `{ "type": <modality> }`.
                ChatContentPart::ImageUrl { .. } => {
-                    multimodal.ok_or(Error::UnsupportedMultimodalContent("image_url"))?;
+                    multimodal
+                        .and_then(|multimodal| multimodal.image_token.as_ref())
+                        .ok_or(Error::UnsupportedMultimodalContent("image_url"))?;
                    Ok(TemplateContentPart::Image)
                }
+                ChatContentPart::VideoUrl { .. } => {
+                    multimodal
+                        .and_then(|multimodal| multimodal.video_token.as_ref())
+                        .ok_or(Error::UnsupportedMultimodalContent("video_url"))?;
+                    Ok(TemplateContentPart::Video)
+                }
            })
            .collect(),
    }
@@ -437,9 +451,16 @@ fn to_template_string_content(
                match part {
                    ChatContentPart::Text { text } => out.push_str(text),
                    ChatContentPart::ImageUrl { .. } => {
-                        let multimodal =
-                            multimodal.ok_or(Error::UnsupportedMultimodalContent("image_url"))?;
-                        out.push_str(&multimodal.placeholder_token);
+                        let image_token = multimodal
+                            .and_then(|multimodal| multimodal.image_token.as_ref())
+                            .ok_or(Error::UnsupportedMultimodalContent("image_url"))?;
+                        out.push_str(image_token);
+                    }
+                    ChatContentPart::VideoUrl { .. } => {
+                        let video_token = multimodal
+                            .and_then(|multimodal| multimodal.video_token.as_ref())
+                            .ok_or(Error::UnsupportedMultimodalContent("video_url"))?;
+                        out.push_str(video_token);
                    }
                }
            }
@@ -468,7 +489,7 @@ fn append_continue_final_message_tag(message: &mut TemplateMessage) -> Result<St
        // Pick the last text part in the message.
        TemplateContent::OpenAi(parts) => parts.iter_mut().rev().find_map(|part| match part {
            TemplateContentPart::Text { text } => Some(text),
-            TemplateContentPart::Image => None,
+            TemplateContentPart::Image | TemplateContentPart::Video => None,
        }),
    };
    let text = text.ok_or_else(|| {
@@ -577,7 +598,8 @@ mod tests {
    ) -> Result<crate::RenderedPrompt> {
        HfChatRenderer::new(Some(template.to_string()), HashMap::new(), content_format)?
            .with_multimodal(Some(MultimodalRenderInfo {
-                placeholder_token: "<image>".to_string(),
+                image_token: Some("<image>".to_string()),
+                video_token: Some("<video>".to_string()),
            }))
            .render(request)
    }
@@ -590,6 +612,14 @@ mod tests {
        ])])
    }

+    fn video_request() -> ChatRequest {
+        sample_request(vec![ChatMessage::user(vec![
+            ChatContentPart::text("a"),
+            ChatContentPart::video_url("https://example.com/demo.mp4"),
+            ChatContentPart::text("b"),
+        ])])
+    }
+
    #[test]
    fn string_content_format_replaces_image_with_placeholder_text() {
        let rendered = render_mm(
@@ -602,6 +632,18 @@ mod tests {
        assert_eq!(rendered.prompt, Prompt::Text("a<image>b".to_string()));
    }

+    #[test]
+    fn string_content_format_replaces_video_with_placeholder_text() {
+        let rendered = render_mm(
+            "{{ messages[0].content }}",
+            &video_request(),
+            ChatTemplateContentFormatOption::String,
+        )
+        .unwrap();
+
+        assert_eq!(rendered.prompt, Prompt::Text("a<video>b".to_string()));
+    }
+
    #[test]
    fn openai_content_format_normalizes_image_url_for_template() {
        let rendered = render_mm(
@@ -614,6 +656,39 @@ mod tests {
        assert_eq!(rendered.prompt, Prompt::Text("a<|image_pad|>b".to_string()));
    }

+    #[test]
+    fn openai_content_format_normalizes_video_url_for_template() {
+        let rendered = render_mm(
+            "{% for item in messages[0].content %}{% if item.type == 'video' %}<|video_pad|>{% else %}{{ item.text }}{% endif %}{% endfor %}",
+            &video_request(),
+            ChatTemplateContentFormatOption::OpenAi,
+        )
+        .unwrap();
+
+        assert_eq!(rendered.prompt, Prompt::Text("a<|video_pad|>b".to_string()));
+    }
+
+    #[test]
+    fn video_parts_are_rejected_when_model_lacks_video_support() {
+        let error = HfChatRenderer::new(
+            Some("{{ messages[0].content }}".to_string()),
+            HashMap::new(),
+            ChatTemplateContentFormatOption::String,
+        )
+        .unwrap()
+        .with_multimodal(Some(MultimodalRenderInfo {
+            image_token: Some("<image>".to_string()),
+            video_token: None,
+        }))
+        .render(&video_request())
+        .unwrap_err();
+
+        assert!(matches!(
+            error,
+            Error::UnsupportedMultimodalContent("video_url")
+        ));
+    }
+
    #[test]
    fn chat_template_supports_pycompat_templates() {
        let request = sample_request(vec![ChatMessage::text(ChatRole::User, "<think>hello")]);
@@ -36,7 +36,13 @@ pub enum ChatContentPart {
        detail: Option<ImageDetail>,
        uuid: Option<String>,
    },
+    /// One video URL/data URL content block.
+    VideoUrl {
+        video_url: String,
+        uuid: Option<String>,
+    },
    // ImageData...
+    // VideoData...
    // ImageEmbeds...
 }

@@ -55,12 +61,21 @@ impl ChatContentPart {
        }
    }

+    /// Construct one video URL content part with the given URL string.
+    pub fn video_url(video_url: impl Into<String>) -> Self {
+        Self::VideoUrl {
+            video_url: video_url.into(),
+            uuid: None,
+        }
+    }
+
    /// Return the text content of this part when it's a text block, or an
    /// "unsupported multimodal content" error otherwise.
    pub(crate) fn as_text(&self) -> Result<&str> {
        match self {
            Self::Text { text } => Ok(text),
            Self::ImageUrl { .. } => Err(Error::UnsupportedMultimodalContent("image_url")),
+            Self::VideoUrl { .. } => Err(Error::UnsupportedMultimodalContent("video_url")),
        }
    }

@@ -73,7 +88,7 @@ impl ChatContentPart {
    pub(crate) fn is_multimodal(&self) -> bool {
        match self {
            Self::Text { .. } => false,
-            Self::ImageUrl { .. } => true,
+            Self::ImageUrl { .. } | Self::VideoUrl { .. } => true,
        }
    }
 }
@@ -555,6 +570,26 @@ mod tests {
        assert_eq!(content, ChatContent::Text("hello".to_string()));
    }

+    #[test]
+    fn chat_content_video_url_part_round_trips_through_serde() {
+        let content = ChatContent::Parts(vec![ChatContentPart::VideoUrl {
+            video_url: "https://example.com/demo.mp4".to_string(),
+            uuid: Some("video-1".to_string()),
+        }]);
+
+        let value = to_value(&content).unwrap();
+        assert_eq!(
+            value,
+            json!([{
+                "type": "video_url",
+                "video_url": "https://example.com/demo.mp4",
+                "uuid": "video-1",
+            }])
+        );
+        let decoded: ChatContent = serde_json::from_value(value).unwrap();
+        assert_eq!(decoded, content);
+    }
+
    #[test]
    fn chat_content_deserializes_from_openai_text_blocks() {
        let content: ChatContent =
@@ -117,6 +117,19 @@ impl RoundtripCase {
        }
    }

+    /// MiniMax M3 invoke format with `<mm:think>` reasoning tags.
+    fn minimax_m3() -> Self {
+        Self {
+            model_id: "MiniMaxAI/MiniMax-M3",
+            assistant_stop_suffix: "[e~[\n",
+            tool_call_parser: ParserSelection::Auto,
+            reasoning_parser: ParserSelection::Auto,
+            thinking_behavior: ThinkingBehavior::Always { value: true },
+            json_fmt: compact_json_fmt(),
+            sort_json_keys: false,
+        }
+    }
+
    /// DeepSeek V4 DSML tool-call format.
    fn deepseek_v4() -> Self {
        Self {
@@ -143,6 +156,19 @@ impl RoundtripCase {
        }
    }

+    /// GLM-4.5 XML-like argument format with `<think>` reasoning tags.
+    fn glm45() -> Self {
+        Self {
+            model_id: "zai-org/GLM-4.5",
+            assistant_stop_suffix: "",
+            tool_call_parser: ParserSelection::Auto,
+            reasoning_parser: ParserSelection::Auto,
+            thinking_behavior: ThinkingBehavior::Toggleable { default: true },
+            json_fmt: compact_json_fmt(),
+            sort_json_keys: false,
+        }
+    }
+
    /// GLM-4.7 XML-like argument format with `<think>` reasoning tags.
    fn glm47() -> Self {
        Self {
@@ -209,6 +235,19 @@ impl RoundtripCase {
        }
    }

+    /// Nemotron V3 with `<think>` / `</think>` reasoning tags.
+    fn nemotron_v3() -> Self {
+        Self {
+            model_id: "nvidia/NVIDIA-Nemotron-3-Nano-30B-A3B-BF16",
+            assistant_stop_suffix: "<|im_end|>\n",
+            tool_call_parser: ParserSelection::Auto,
+            reasoning_parser: ParserSelection::Auto,
+            thinking_behavior: ThinkingBehavior::Always { value: true },
+            json_fmt: compact_json_fmt(),
+            sort_json_keys: false,
+        }
+    }
+
    /// GPT-OSS Harmony token-id renderer and native Harmony output processor.
    fn gpt_oss() -> Self {
        Self {
@@ -247,11 +286,14 @@ roundtrip_tests! {
    qwen3 => [reasoning_and_content, tool_call_mix],
    qwen35 => [reasoning_and_content, tool_call_mix],
    minimax_m25 => [reasoning_and_content, tool_call_mix],
+    minimax_m3 => [reasoning_and_content, tool_call_mix],
    deepseek_v4 => [reasoning_and_content, tool_call_mix],
    deepseek_v32 => [tool_call_mix],
+    glm45 => [reasoning_and_content, tool_call_mix],
    glm47 => [reasoning_and_content, tool_call_mix],
    seed_oss => [reasoning_and_content],
    step3p5 => [reasoning_and_content],
+    nemotron_v3 => [reasoning_and_content],
    gemma4 => [tool_call_mix], // Gemma4 strips reasoning in history if there's no tool call
    kimi_k25 => [tool_call_mix], // Kimi K2.5 strips reasoning in history
    gpt_oss => [tool_call_mix], // Harmony strips reasoning in history if there's no tool call
@@ -290,7 +290,10 @@ fn convert_content(content: MessageContent) -> Result<ChatContent, ApiError> {
                    detail: image_url.detail,
                    uuid,
                }),
-                _ => bail_invalid_request!("Only text and image_url content parts are supported."),
+                ContentPart::VideoUrl { video_url, uuid } => Ok(ChatContentPart::VideoUrl {
+                    video_url: video_url.url,
+                    uuid,
+                }),
            })
            .try_collect()
            .map(ChatContent::Parts),
@@ -772,28 +775,34 @@ mod tests {
    }

    #[test]
-    fn prepare_chat_request_rejects_video_content_parts() {
+    fn prepare_chat_request_accepts_video_content_parts() {
        let request = ChatCompletionRequest {
            messages: vec![ChatMessage::User {
                content: MessageContent::Parts(vec![ContentPart::VideoUrl {
                    video_url: VideoUrl {
                        url: "https://example.com/video.mp4".to_string(),
                    },
+                    uuid: Some("video-uuid".to_string()),
                }]),
                name: None,
            }],
            ..base_request()
        };

-        let error = prepare_chat_request(
+        let prepared = prepare_chat_request(
            request,
            &served(&["Qwen/Qwen1.5-0.5B-Chat"]),
            ResolvedRequestContext::default(),
        )
-        .unwrap_err();
+        .expect("request is valid");

-        expect!["Only text and image_url content parts are supported."]
-            .assert_eq(&error.to_error_response().error.message);
+        assert_eq!(
+            prepared.chat_request.messages,
+            vec![VllmChatMessage::user(vec![ChatContentPart::VideoUrl {
+                video_url: "https://example.com/video.mp4".to_string(),
+                uuid: Some("video-uuid".to_string()),
+            }])]
+        );
    }

    #[test]
@@ -91,7 +91,11 @@ pub enum ContentPart {
        uuid: Option<String>,
    },
    #[serde(rename = "video_url")]
-    VideoUrl { video_url: VideoUrl },
+    VideoUrl {
+        video_url: VideoUrl,
+        #[serde(skip_serializing_if = "Option::is_none")]
+        uuid: Option<String>,
+    },
 }

 #[serde_with::skip_serializing_none]
@@ -504,7 +504,7 @@ impl ChatRenderer for FakeChatBackend {
        let placeholder = self
            .multimodal_model_info
            .as_ref()
-            .map(|info| info.placeholder_token())
+            .and_then(|info| info.placeholder_token(llm_multimodal::Modality::Image))
            .unwrap_or("<image>");
        let mut prompt = String::new();
        for message in &request.messages {
@@ -544,7 +544,9 @@ fn render_fake_content(content: &ChatContent, placeholder: &str) -> vllm_chat::R
            for part in parts {
                match part {
                    ChatContentPart::Text { text } => out.push_str(text),
-                    ChatContentPart::ImageUrl { .. } => out.push_str(placeholder),
+                    ChatContentPart::ImageUrl { .. } | ChatContentPart::VideoUrl { .. } => {
+                        out.push_str(placeholder)
+                    }
                }
            }
            out
@@ -565,8 +567,10 @@ fn qwen_multimodal_model_info() -> vllm_chat::multimodal::MultimodalModelInfo {
    let info = vllm_chat::multimodal::MultimodalModelInfo::from_paths(
        "qwen2-vl-test".to_string(),
        Some("qwen2_vl".to_string()),
-        Some(&config_path),
-        None,
+        vllm_chat::multimodal::MultimodalConfigFiles {
+            config: Some(&config_path),
+            ..Default::default()
+        },
        Arc::new(fake_chat_tokenizer()),
    )
    .expect("load multimodal info")
@@ -41,6 +41,10 @@ pub struct ResolvedModelFiles {
    pub tokenizer_config_path: Option<PathBuf>,
    pub generation_config_path: Option<PathBuf>,
    pub preprocessor_config_path: Option<PathBuf>,
+    /// Video-specific preprocessor config, when provided by the model repo.
+    pub video_preprocessor_config_path: Option<PathBuf>,
+    /// Combined processor config, which may embed a `video_processor` section.
+    pub processor_config_path: Option<PathBuf>,
    pub chat_template_path: Option<PathBuf>,
    pub config_path: Option<PathBuf>,
 }
@@ -70,6 +74,11 @@ fn resolve_local_model_files(model_dir: &Path) -> Result<ResolvedModelFiles> {
        tokenizer_config_path,
        generation_config_path: local_file_if_exists(model_dir, "generation_config.json"),
        preprocessor_config_path: local_file_if_exists(model_dir, "preprocessor_config.json"),
+        video_preprocessor_config_path: local_file_if_exists(
+            model_dir,
+            "video_preprocessor_config.json",
+        ),
+        processor_config_path: local_file_if_exists(model_dir, "processor_config.json"),
        chat_template_path: discover_chat_template_in_dir(model_dir),
        config_path: local_file_if_exists(model_dir, "config.json"),
    })
@@ -107,6 +116,10 @@ async fn resolve_remote_model_files(model_id: &str) -> Result<ResolvedModelFiles
        download_if_present(&repo, model_id, &siblings, "generation_config.json").await?;
    let preprocessor_config_path =
        download_if_present(&repo, model_id, &siblings, "preprocessor_config.json").await?;
+    let video_preprocessor_config_path =
+        download_if_present(&repo, model_id, &siblings, "video_preprocessor_config.json").await?;
+    let processor_config_path =
+        download_if_present(&repo, model_id, &siblings, "processor_config.json").await?;
    let chat_template_name = siblings
        .contains("chat_template.json")
        .then_some("chat_template.json")
@@ -123,6 +136,8 @@ async fn resolve_remote_model_files(model_id: &str) -> Result<ResolvedModelFiles
        tokenizer_config_path,
        generation_config_path,
        preprocessor_config_path,
+        video_preprocessor_config_path,
+        processor_config_path,
        chat_template_path,
        config_path,
    })
@@ -143,6 +158,8 @@ fn resolve_cached_model_files(model_id: &str) -> Result<Option<ResolvedModelFile
    })?;
    let generation_config_path = cache_repo.get("generation_config.json");
    let preprocessor_config_path = cache_repo.get("preprocessor_config.json");
+    let video_preprocessor_config_path = cache_repo.get("video_preprocessor_config.json");
+    let processor_config_path = cache_repo.get("processor_config.json");
    let chat_template_path = discover_chat_template_in_dir(model_dir);
    let config_path = cache_repo.get("config.json");

@@ -151,6 +168,8 @@ fn resolve_cached_model_files(model_id: &str) -> Result<Option<ResolvedModelFile
        tokenizer_config_path,
        generation_config_path,
        preprocessor_config_path,
+        video_preprocessor_config_path,
+        processor_config_path,
        chat_template_path,
        config_path,
    }))
@@ -1257,6 +1257,9 @@ setup(
            "mistral_common[audio]",
        ],  # Required for audio processing
        "video": [],  # Kept for backwards compatibility
+        # NVIDIA DeepStream (NVDEC) GPU video-decode backend. Linux x86-64
+        # only; also needs system GStreamer + libv4l (see docs).
+        "deepstream": ["nvidia-deepstream-videodecode-cu13>=9.0.2"],
        "flashinfer": [],  # Kept for backwards compatibility
        # Optional deps for Helion kernel development
        # NOTE: When updating helion version, also update CI files:
@@ -33,6 +33,7 @@ CP_TEST_MODELS = [
    # [LANGUAGE GENERATION]
    "deepseek-ai/DeepSeek-V2-Lite-Chat",
    "Qwen/Qwen2.5-1.5B-Instruct",
+    "Qwen/Qwen3.5-0.8B",  # hybrid attention model
 ]

 # GSM8K eval configuration
@@ -46,6 +47,7 @@ MIN_ACCURACY = {
    "deepseek-ai/DeepSeek-V2-Lite-Chat": 0.64,
    # .buildkite/lm-eval-harness/configs/Qwen2.5-1.5B-Instruct.yaml
    "Qwen/Qwen2.5-1.5B-Instruct": 0.52,
+    "Qwen/Qwen3.5-0.8B": 0.33,
 }


@@ -151,6 +153,12 @@ else:
                cp_kv_cache_interleave_size=16, attn_backend="FLASHINFER"
            ),
        ],
+        "Qwen/Qwen3.5-0.8B": [
+            CPTestSettings.detailed(
+                cp_kv_cache_interleave_size=16,
+                attn_backend="FLASH_ATTN",
+            ),
+        ],
    }


@@ -16,6 +16,7 @@ from vllm.distributed.device_communicators.pynccl import PyNcclCommunicator
 from vllm.distributed.device_communicators.pynccl_wrapper import NCCLLibrary
 from vllm.distributed.parallel_state import (
    ensure_model_parallel_initialized,
+    get_tp_group,
    get_world_group,
    graph_capture,
    init_distributed_environment,
@@ -199,6 +200,52 @@ def test_pynccl_all_gather():
    distributed_run(all_gather_worker_fn, 2)


+@worker_fn_wrapper
+def cuda_communicator_all_gather_dim_worker_fn():
+    with ensure_current_vllm_config():
+        ensure_model_parallel_initialized(2, 1)
+
+    tp_group = get_tp_group()
+    comm = tp_group.device_communicator
+    assert comm is not None
+
+    rank = tp_group.rank_in_group
+    world_size = tp_group.world_size
+    device = tp_group.device
+
+    shape = (2, 3, 4)
+    num_elems = 1
+    for size in shape:
+        num_elems *= size
+
+    for dim in (1, -1):
+        tensor = (
+            torch.arange(num_elems, dtype=torch.float32, device=device).reshape(shape)
+            + rank * num_elems
+        )
+        expected = torch.cat(
+            [
+                torch.arange(num_elems, dtype=torch.float32, device=device).reshape(
+                    shape
+                )
+                + r * num_elems
+                for r in range(world_size)
+            ],
+            dim=dim,
+        )
+
+        result = comm.all_gather(tensor, dim=dim)
+        torch.accelerator.synchronize()
+        torch.testing.assert_close(result, expected, rtol=1e-5, atol=1e-8)
+
+
+@pytest.mark.skipif(
+    torch.accelerator.device_count() < 2, reason="Need at least 2 GPUs to run the test."
+)
+def test_cuda_communicator_all_gather_dim_not_zero():
+    distributed_run(cuda_communicator_all_gather_dim_worker_fn, 2)
+
+
@worker_fn_wrapper
 def all_gatherv_worker_fn():
    pynccl_comm = PyNcclCommunicator(
@@ -0,0 +1,158 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+"""E2E tests for ``include_reasoning`` with non-Harmony reasoning models.
+
+Verifies that reasoning content is included by default and suppressed
+when ``include_reasoning=False``, for both streaming and non-streaming
+Chat Completions.
+"""
+
+import openai
+import pytest
+import pytest_asyncio
+
+from tests.utils import RemoteOpenAIServer
+
+MODEL_NAME = "Qwen/Qwen3-0.6B"
+MESSAGES = [{"role": "user", "content": "What is 1+1? Be concise."}]
+
+
+@pytest.fixture(scope="module")
+def server():
+    args = [
+        "--reasoning-parser",
+        "qwen3",
+        "--max-model-len",
+        "2048",
+        "--enforce-eager",
+        "--gpu-memory-utilization",
+        "0.4",
+    ]
+    with RemoteOpenAIServer(MODEL_NAME, args) as remote_server:
+        yield remote_server
+
+
+@pytest_asyncio.fixture
+async def client(server):
+    async with server.get_async_client() as async_client:
+        yield async_client
+
+
+@pytest.mark.asyncio
+async def test_include_reasoning_true_non_streaming(client: openai.AsyncOpenAI):
+    """Default: reasoning content appears in non-streaming response."""
+    response = await client.chat.completions.create(
+        model=MODEL_NAME,
+        messages=MESSAGES,
+        max_tokens=200,
+        extra_body={"include_reasoning": True},
+    )
+
+    msg = response.choices[0].message
+    reasoning = getattr(msg, "reasoning", None) or getattr(
+        msg, "reasoning_content", None
+    )
+    assert reasoning, "Expected reasoning content when include_reasoning=True"
+    assert msg.content, "Expected content in response"
+
+
+@pytest.mark.asyncio
+async def test_include_reasoning_false_non_streaming(client: openai.AsyncOpenAI):
+    """Reasoning content is suppressed when include_reasoning=False."""
+    response = await client.chat.completions.create(
+        model=MODEL_NAME,
+        messages=MESSAGES,
+        max_tokens=200,
+        extra_body={"include_reasoning": False},
+    )
+
+    msg = response.choices[0].message
+    reasoning = getattr(msg, "reasoning", None) or getattr(
+        msg, "reasoning_content", None
+    )
+    assert not reasoning, (
+        f"Expected no reasoning when include_reasoning=False, got: {reasoning}"
+    )
+    assert msg.content, "Expected content in response even without reasoning"
+
+
+@pytest.mark.asyncio
+async def test_include_reasoning_true_streaming(client: openai.AsyncOpenAI):
+    """Default: reasoning deltas appear in streaming response."""
+    stream = await client.chat.completions.create(
+        model=MODEL_NAME,
+        messages=MESSAGES,
+        max_tokens=200,
+        stream=True,
+        extra_body={"include_reasoning": True},
+    )
+
+    reasoning_parts = []
+    content_parts = []
+    async for chunk in stream:
+        delta = chunk.choices[0].delta if chunk.choices else None
+        if delta:
+            r = getattr(delta, "reasoning", None) or getattr(
+                delta, "reasoning_content", None
+            )
+            if r:
+                reasoning_parts.append(r)
+            if delta.content:
+                content_parts.append(delta.content)
+
+    reasoning_text = "".join(reasoning_parts)
+    content_text = "".join(content_parts)
+
+    assert reasoning_text, "Expected reasoning deltas when include_reasoning=True"
+    assert content_text, "Expected content deltas in streaming response"
+
+
+@pytest.mark.asyncio
+async def test_include_reasoning_false_streaming(client: openai.AsyncOpenAI):
+    """Reasoning deltas are suppressed in streaming when include_reasoning=False."""
+    stream = await client.chat.completions.create(
+        model=MODEL_NAME,
+        messages=MESSAGES,
+        max_tokens=200,
+        stream=True,
+        extra_body={"include_reasoning": False},
+    )
+
+    reasoning_parts = []
+    content_parts = []
+    async for chunk in stream:
+        delta = chunk.choices[0].delta if chunk.choices else None
+        if delta:
+            r = getattr(delta, "reasoning", None) or getattr(
+                delta, "reasoning_content", None
+            )
+            if r:
+                reasoning_parts.append(r)
+            if delta.content:
+                content_parts.append(delta.content)
+
+    reasoning_text = "".join(reasoning_parts)
+    content_text = "".join(content_parts)
+
+    assert not reasoning_text, (
+        f"Expected no reasoning deltas when include_reasoning=False, "
+        f"got: {reasoning_text[:100]}"
+    )
+    assert content_text, "Expected content deltas even without reasoning"
+
+
+@pytest.mark.asyncio
+async def test_default_includes_reasoning(client: openai.AsyncOpenAI):
+    """Without specifying include_reasoning, reasoning appears (default=True)."""
+    response = await client.chat.completions.create(
+        model=MODEL_NAME,
+        messages=MESSAGES,
+        max_tokens=200,
+    )
+
+    msg = response.choices[0].message
+    reasoning = getattr(msg, "reasoning", None) or getattr(
+        msg, "reasoning_content", None
+    )
+    assert reasoning, "Expected reasoning content by default"
@@ -197,46 +197,6 @@ def _ragged_from_rows(
    )


-def _ref_combine_topk_swa_ragged(
-    device: torch.device,
-) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
-    expected_ragged = torch.tensor(
-        [
-            100,
-            101,
-            7,
-            8,
-            9,
-            110,
-            111,
-            8,
-            9,
-            10,
-            120,
-            121,
-            122,
-            9,
-            10,
-            11,
-            150,
-            27,
-            28,
-            29,
-            160,
-            161,
-            28,
-            29,
-            30,
-        ],
-        dtype=torch.int32,
-        device=device,
-    )
-    expected_lens = torch.tensor([5, 5, 6, 4, 5], dtype=torch.int32, device=device)
-    expected_indptr = torch.zeros(6, dtype=torch.int32, device=device)
-    torch.cumsum(expected_lens, dim=0, out=expected_indptr[1:])
-    return expected_ragged, expected_indptr, expected_lens
-
-
@torch.inference_mode()
 def test_compute_global_topk_ragged_indices_and_indptr() -> None:
    from vllm.models.deepseek_v4.amd.rocm import (
@@ -369,55 +329,6 @@ def test_sparse_attn_decode_ragged_kernel() -> None:
    torch.testing.assert_close(actual, expected, atol=2e-2, rtol=2e-2)


-@torch.inference_mode()
-def test_combine_topk_swa_indices_ragged() -> None:
-    from vllm.models.deepseek_v4.amd.rocm import (
-        combine_topk_swa_indices_ragged,
-    )
-
-    device = torch.device("cuda")
-    topk_indices = torch.tensor(
-        [
-            [100, 101, 102, 103],
-            [110, 111, 112, 113],
-            [120, 121, 122, 123],
-            [130, 131, 132, 133],
-            [140, 141, 142, 143],
-        ],
-        dtype=torch.int32,
-        device=device,
-    )
-    query_start_loc = torch.tensor([0, 3, 5], dtype=torch.int32, device=device)
-    seq_lens = torch.tensor([6, 4], dtype=torch.int32, device=device)
-    gather_lens = torch.tensor([4, 3], dtype=torch.int32, device=device)
-    window_size = 3
-    compress_ratio = 2
-    topk = 4
-    M = 20
-    N = 8
-
-    actual_ragged, actual_indptr, actual_lens = combine_topk_swa_indices_ragged(
-        topk_indices,
-        query_start_loc,
-        seq_lens,
-        gather_lens,
-        window_size,
-        compress_ratio,
-        topk,
-        M,
-        N,
-    )
-    expected_ragged, expected_indptr, expected_lens = _ref_combine_topk_swa_ragged(
-        device
-    )
-
-    torch.testing.assert_close(
-        actual_ragged[: expected_ragged.numel()], expected_ragged
-    )
-    torch.testing.assert_close(actual_indptr, expected_indptr)
-    torch.testing.assert_close(actual_lens, expected_lens)
-
-
@requires_gfx950
@torch.inference_mode()
 def test_decode_num_splits_heuristic(monkeypatch) -> None:
@@ -0,0 +1,117 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""LongCat n-gram embedding id computation vs a pure-Python reference.
+
+Guards the hash-id semantics of ``ngram_compute_n_gram_ids`` plus the
+EOS-position fixup (``compute_eos_position_ngram_ids``): an EOS *current*
+token hashes with its full look-back, while later positions' look-back stops
+at the EOS boundary (LongCat reference behavior).
+"""
+
+from types import SimpleNamespace
+
+import pytest
+import torch
+
+from vllm import _custom_ops as ops
+from vllm.model_executor.models.longcat_flash_ngram import (
+    compute_eos_position_ngram_ids,
+)
+from vllm.platforms import current_platform
+
+VOCAB = 163840
+N, K = 5, 4  # oe_neighbor_num, oe_split_num (LongCat-2.0)
+M = int(100.567 * VOCAB)
+EOS = 2
+NUM_EMB = K * (N - 1)
+
+SIZES = [M + c * 2 + 1 for c in range(NUM_EMB)]
+OFFSETS = [0]
+for s in SIZES:
+    OFFSETS.append(OFFSETS[-1] + s)
+
+
+def _ref_ids(tokens: list[int]) -> list[list[int]]:
+    """Reference n-gram ids on raw tokens (HF LongCat semantics)."""
+    out = []
+    for pos in range(len(tokens)):
+        row = []
+        for i in range(N - 1):
+            n_real = i + 2
+            for j in range(K):
+                cfg = i * K + j
+                mod = SIZES[cfg]
+                h = 0
+                for delta in range(n_real):
+                    p = pos - delta
+                    if p < 0:
+                        break
+                    t = tokens[p]
+                    if delta > 0 and t == EOS:
+                        break  # look-back stops at an EOS boundary
+                    h += (t * pow(VOCAB, delta, mod)) % mod
+                row.append(h % mod + OFFSETS[cfg])
+        out.append(row)
+    return out
+
+
+@pytest.mark.skipif(not current_platform.is_cuda(), reason="CUDA kernel")
+def test_ngram_ids_match_reference_with_eos():
+    device = "cuda"
+    torch.manual_seed(0)
+    tokens = torch.randint(3, VOCAB, (14,), dtype=torch.int32).tolist()
+    tokens[5] = EOS
+    tokens[6] = EOS  # double EOS: second one's look-back stops at the first
+    tokens[13] = EOS  # trailing EOS (chat-template turn boundary)
+
+    ctx_len = N - 1
+    toks_neg = [-t if t == EOS else t for t in tokens]
+    width = ctx_len + len(tokens)
+    table = torch.full((1, width), -1, dtype=torch.int32, device=device)
+    table[0, ctx_len:] = torch.tensor(toks_neg, dtype=torch.int32, device=device)
+
+    ne_weights = torch.zeros(N - 1, K, N, dtype=torch.int32)
+    ne_mods = torch.zeros(N - 1, K, dtype=torch.int32)
+    for i in range(N - 1):
+        for j in range(K):
+            mod = SIZES[i * K + j]
+            ne_mods[i, j] = mod
+            for delta in range(N):
+                ne_weights[i, j, delta] = pow(VOCAB, delta, mod)
+    ngram = SimpleNamespace(
+        n=N,
+        k=K,
+        num_embedders=NUM_EMB,
+        ne_weights=ne_weights.to(device),
+        ne_mods=ne_mods.to(device),
+        exclusive_sizes=torch.tensor(OFFSETS, dtype=torch.int32, device=device),
+    )
+
+    T = len(tokens)
+    qsl = torch.tensor([0, T], dtype=torch.int32, device=device)
+    row_indices = torch.zeros(1, dtype=torch.int64, device=device)
+    column_starts = torch.full((1,), ctx_len, dtype=torch.int32, device=device)
+    got = torch.empty(T, NUM_EMB, dtype=torch.int32, device=device)
+    ops.ngram_compute_n_gram_ids(
+        N,
+        K,
+        ngram.ne_weights,
+        ngram.ne_mods,
+        ngram.exclusive_sizes,
+        qsl,
+        table,
+        row_indices,
+        column_starts,
+        got,
+    )
+
+    cur = torch.tensor(tokens, dtype=torch.int32, device=device)
+    tok_req = torch.zeros(T, dtype=torch.int64, device=device)
+    col = ctx_len + torch.arange(T, device=device)
+    eos_tok = (cur == EOS).nonzero(as_tuple=True)[0]
+    got[eos_tok] = compute_eos_position_ngram_ids(
+        ngram, EOS, table, tok_req, col, eos_tok
+    )
+
+    want = torch.tensor(_ref_ids(tokens), dtype=torch.int32)
+    torch.testing.assert_close(got.cpu(), want, rtol=0, atol=0)
@@ -0,0 +1,79 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import json
+import os
+import tempfile
+
+import pytest
+
+from vllm.model_executor.model_loader.weight_utils import (
+    filter_duplicate_safetensors_files,
+)
+
+
+def test_filter_duplicate_safetensors_files_missing_weight():
+    with tempfile.TemporaryDirectory() as tmpdir:
+        existing_file = os.path.join(tmpdir, "model-00001-of-00002.safetensors")
+        with open(existing_file, "wb") as f:
+            f.write(b"")
+
+        existing_file2 = os.path.join(tmpdir, "model-00002-of-00002.safetensors")
+        with open(existing_file2, "wb") as f:
+            f.write(b"")
+
+        index_file = os.path.join(tmpdir, "model.safetensors.index.json")
+        index_content = {
+            "weight_map": {
+                "layer.0.weight": "model-00001-of-00002.safetensors",
+                "layer.1.weight": "model-00002-of-00002.safetensors",
+                "layer.2.weight": "model-00003-of-00002.safetensors",
+            }
+        }
+        with open(index_file, "w") as f:
+            json.dump(index_content, f)
+
+        hf_weights_files = [
+            os.path.join(tmpdir, "model-00001-of-00002.safetensors"),
+            os.path.join(tmpdir, "model-00002-of-00002.safetensors"),
+        ]
+
+        with pytest.raises(FileNotFoundError) as exc_info:
+            filter_duplicate_safetensors_files(
+                hf_weights_files=hf_weights_files,
+                hf_folder=tmpdir,
+                index_file="model.safetensors.index.json",
+            )
+
+        assert "model-00003-of-00002.safetensors" in str(exc_info.value)
+
+
+def test_filter_duplicate_safetensors_files_all_exist():
+    with tempfile.TemporaryDirectory() as tmpdir:
+        existing_files = []
+        for i in range(1, 3):
+            file_path = os.path.join(tmpdir, f"model-0000{i}-of-00002.safetensors")
+            with open(file_path, "wb") as f:
+                f.write(b"")
+            existing_files.append(file_path)
+
+        index_file = os.path.join(tmpdir, "model.safetensors.index.json")
+        index_content = {
+            "weight_map": {
+                "layer.0.weight": "model-00001-of-00002.safetensors",
+                "layer.1.weight": "model-00002-of-00002.safetensors",
+            }
+        }
+        with open(index_file, "w") as f:
+            json.dump(index_content, f)
+
+        filter_duplicate_safetensors_files(
+            hf_weights_files=existing_files,
+            hf_folder=tmpdir,
+            index_file="model.safetensors.index.json",
+        )
+
+
+if __name__ == "__main__":
+    test_filter_duplicate_safetensors_files_missing_weight()
+    test_filter_duplicate_safetensors_files_all_exist()
@@ -43,6 +43,11 @@ HYBRID_MODELS = [
    "tiny-random/qwen3-next-moe",
 ]

+HYBRID_MODELS_REQUIRING_CHUNKED_PREFILL = {
+    "LiquidAI/LFM2-1.2B",
+    "tiny-random/qwen3-next-moe",
+}
+
 FULL_CUDA_GRAPH_MODELS = [
    "ai21labs/Jamba-tiny-dev",
    "pfnet/plamo-2-1b",
@@ -92,8 +97,15 @@ def test_models(
            example_prompts, max_tokens, num_logprobs
        )

+    extra_kwargs = {}
+    if model in HYBRID_MODELS_REQUIRING_CHUNKED_PREFILL:
+        extra_kwargs["enable_chunked_prefill"] = True
+
    with vllm_runner(
-        model, max_num_seqs=MAX_NUM_SEQS, attention_backend=ATTN_BACKEND
+        model,
+        max_num_seqs=MAX_NUM_SEQS,
+        attention_backend=ATTN_BACKEND,
+        **extra_kwargs,
    ) as vllm_model:
        vllm_outputs = vllm_model.generate_greedy_logprobs(
            example_prompts, max_tokens, num_logprobs
@@ -369,6 +369,12 @@ VLM_TEST_SETTINGS = {
        vllm_output_post_proc=model_utils.qwen2_vllm_to_hf_output,
        patch_hf_runner=model_utils.qwen3_vl_patch_hf_runner,
        image_size_factors=[(0.25,), (0.25, 0.25, 0.25), (0.25, 0.2, 0.15)],
+        marks=[
+            pytest.mark.skip(
+                reason="Transformers has no cosmos3_omni mapping, so the HF "
+                "reference runner cannot load the checkpoint."
+            )
+        ],
    ),
    "deepseek_vl_v2": VLMTestInfo(
        models=["Isotr0py/deepseek-vl2-tiny"],  # model repo using dynamic module
@@ -1,438 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""
-Consolidated test for ViT attention backend functionality across multiple models.
-
-This test validates that each multimodal model can successfully generate outputs
-using different ViT attention backends. Tests are parametrized by model and backend.
-"""
-
-from typing import Any
-
-import pytest
-from transformers import AutoProcessor
-
-from vllm import LLM, SamplingParams
-from vllm.multimodal.utils import encode_image_url
-from vllm.multimodal.video import sample_frames_from_video
-from vllm.platforms import current_platform
-from vllm.v1.attention.backends.registry import AttentionBackendEnum
-
-from ....utils import create_new_process_for_each_test
-from ...utils import dummy_hf_overrides
-
-# Dots.OCR prompt from official repository
-# https://github.com/rednote-hilab/dots.ocr/blob/d72d1d8c5bdd0362eb264f714cdbd1e5daa7cdff/dots_ocr/utils/prompts.py#L3
-# ruff: noqa: E501
-DOTS_OCR_PROMPT = """Please output the layout information from the PDF image, including each layout element's bbox, its category, and the corresponding text content within the bbox.
-
-1. Bbox format: [x1, y1, x2, y2]
-
-2. Layout Categories: The possible categories are ['Caption', 'Footnote', 'Formula', 'List-item', 'Page-footer', 'Page-header', 'Picture', 'Section-header', 'Table', 'Text', 'Title'].
-
-3. Text Extraction & Formatting Rules:
-    - Picture: For the 'Picture' category, the text field should be omitted.
-    - Formula: Format its text as LaTeX.
-    - Table: Format its text as HTML.
-    - All Others (Text, Title, etc.): Format their text as Markdown.
-
-4. Constraints:
-    - The output text must be the original text from the image, with no translation.
-    - All layout elements must be sorted according to human reading order.
-
-5. Final Output: The entire output must be a single JSON object.
-"""
-
-VIDEO_PLACEHOLDER = "<|vision_start|><|video_pad|><|vision_end|>"
-
-
-# Model configurations
-MODEL_CONFIGS: dict[str, dict[str, Any]] = {
-    "dots_ocr": {
-        "model_name": "rednote-hilab/dots.ocr",
-        "interface": "llm_chat",
-        "max_model_len": 32768,
-        "max_num_seqs": 1,
-        "limit_mm_per_prompt": {"image": 1},
-        "sampling_params": {
-            "temperature": 0.1,
-            "max_tokens": 16384,
-            "top_p": 0.9,
-            "stop_token_ids": None,
-        },
-        "use_specific_image": "stop_sign",
-        "prompt_builder": "build_dots_ocr_prompt",
-        "output_validator": lambda x: len(x) > 10 and "stop" in x.lower(),
-    },
-    "ernie45_vl": {
-        "model_name": "baidu/ERNIE-4.5-VL-28B-A3B-PT",
-        "interface": "llm_generate",
-        "max_model_len": 16384,
-        "max_num_seqs": 2,
-        "sampling_params": {
-            "temperature": 0.0,
-            "max_tokens": 256,
-            "stop_token_ids": None,
-        },
-        "use_processor": True,
-        "question": "What is the content of each image?",
-    },
-    "glm4_1v": {
-        "model_name": "zai-org/GLM-4.1V-9B-Thinking",
-        "interface": "llm_generate",
-        "max_model_len": 32768,
-        "max_num_seqs": 2,
-        "sampling_params": {
-            "temperature": 0.0,
-            "max_tokens": 256,
-            "stop_token_ids": None,
-        },
-        "use_processor": True,
-        "question": "What is the content of each image?",
-    },
-    "glm_ocr": {
-        "model_name": "zai-org/GLM-OCR",
-        "interface": "llm_generate",
-        "max_model_len": 131072,
-        "max_num_seqs": 2,
-        "sampling_params": {
-            "temperature": 0.0,
-            "max_tokens": 256,
-            "stop_token_ids": None,
-        },
-        "use_processor": True,
-        "question": "Text Recognition:",
-    },
-    "keye_vl": {
-        "model_name": "Kwai-Keye/Keye-VL-8B-Preview",
-        "interface": "llm_generate",
-        "max_model_len": 8192,
-        "max_num_seqs": 5,
-        "sampling_params": {
-            "temperature": 0.0,
-            "max_tokens": 256,
-            "stop_token_ids": None,
-        },
-        "supported_backends": {
-            AttentionBackendEnum.FLASH_ATTN,
-            AttentionBackendEnum.ROCM_AITER_FA,
-        },
-        "use_processor": True,
-        "question": "What is the content of each image?",
-    },
-    "ovis2_5": {
-        "model_name": "AIDC-AI/Ovis2.5-2B",
-        "interface": "llm_generate",
-        "max_model_len": 8192,
-        "max_num_seqs": 2,
-        "sampling_params": {
-            "temperature": 0.0,
-            "max_tokens": 256,
-            "stop_token_ids": None,
-        },
-        "prompt_builder": "build_ovis_prompt",
-        "question": "What is the content of each image?",
-    },
-    "qwen2_5_vl": {
-        "model_name": "Qwen/Qwen2.5-VL-3B-Instruct",
-        "interface": "vllm_runner",
-        "media_type": "video",
-        "max_model_len": 4000,
-        "max_num_seqs": 1,
-        "limit_mm_per_prompt": {"video": 1},
-        "sampling_params": {
-            "max_tokens": 128,
-        },
-        "runner_kwargs": {
-            "runner": "generate",
-            "dtype": "bfloat16",
-        },
-        "video_params": {
-            "num_frames": 16,
-            "pruning_rates": [0.0, 0.75],
-        },
-    },
-    "qwen2_5_omni": {
-        "model_name": "Qwen/Qwen2.5-Omni-3B",
-        "interface": "llm_generate",
-        "max_model_len": 32768,
-        "max_num_seqs": 2,
-        "limit_mm_per_prompt": {"image": 3, "video": 3, "audio": 3},
-        "sampling_params": {
-            "temperature": 0.6,
-            "top_p": 0.95,
-            "top_k": 20,
-            "max_tokens": 16384,
-        },
-        "use_processor": True,
-        "question": "What is the content of each image?",
-    },
-    "qwen3_omni": {
-        "model_name": "Qwen/Qwen3-Omni-30B-A3B-Instruct",
-        "interface": "llm_generate",
-        "max_model_len": 32768,
-        "max_num_seqs": 2,
-        "limit_mm_per_prompt": {"image": 3, "video": 3, "audio": 3},
-        "sampling_params": {
-            "temperature": 0.6,
-            "top_p": 0.95,
-            "top_k": 20,
-            "max_tokens": 16384,
-        },
-        "use_processor": True,
-        "question": "What is the content of each image?",
-    },
-}
-
-
-# Prompt builder functions
-def build_dots_ocr_prompt(images, config):
-    """Build Dots.OCR specific prompt with OCR instructions."""
-    # Use only stop_sign image for Dots.OCR
-    image = images[0]  # Already filtered to stop_sign
-    image_url = encode_image_url(image)
-
-    placeholders = [{"type": "image_url", "image_url": {"url": image_url}}]
-    messages = [
-        {
-            "role": "user",
-            "content": [
-                *placeholders,
-                {
-                    "type": "text",
-                    "text": f"<|img|><|imgpad|><|endofimg|>{DOTS_OCR_PROMPT}",
-                },
-            ],
-        },
-    ]
-
-    return messages
-
-
-def build_processor_prompt(images, config):
-    """Build prompt using AutoProcessor.apply_chat_template()."""
-    processor = AutoProcessor.from_pretrained(
-        config["model_name"], trust_remote_code=True
-    )
-
-    image_urls = [encode_image_url(img) for img in images]
-    placeholders = [{"type": "image", "image": url} for url in image_urls]
-    messages = [
-        {
-            "role": "user",
-            "content": [
-                *placeholders,
-                {"type": "text", "text": config["question"]},
-            ],
-        },
-    ]
-
-    return processor.apply_chat_template(
-        messages, tokenize=False, add_generation_prompt=True
-    )
-
-
-def build_ovis_prompt(images, config):
-    """Build Ovis2.5 specific prompt with custom format."""
-    image_urls = [encode_image_url(img) for img in images]
-
-    placeholders = "\n".join(
-        f"Image-{i}: <image>\n" for i, _ in enumerate(image_urls, start=1)
-    )
-
-    return (
-        f"<|im_start|>user\n\n{placeholders}\n{config['question']}<|im_end|>\n"
-        "<|im_start|>assistant\n"
-    )
-
-
-def build_qwen2_5_video_prompt():
-    """Build Qwen2.5-VL video prompt with EVS placeholder."""
-    return (
-        f"<|im_start|>system\nYou are a helpful assistant.<|im_end|>\n"
-        f"<|im_start|>user\n{VIDEO_PLACEHOLDER}"
-        "Describe this video with a short sentence (no more than 20 words)"
-        "<|im_end|><|im_start|>assistant\n"
-    )
-
-
-# Handler functions
-def run_llm_generate_test(config, mm_encoder_attn_backend, image_assets):
-    """Standard LLM.generate() interface handler."""
-    images = [asset.pil_image for asset in image_assets]
-
-    # Build prompt
-    if config.get("use_processor"):
-        prompt = build_processor_prompt(images, config)
-    else:
-        prompt_builder_name = config.get("prompt_builder", "build_ovis_prompt")
-        prompt_builder = globals()[prompt_builder_name]
-        prompt = prompt_builder(images, config)
-
-    # Determine limit_mm_per_prompt
-    limit_mm_per_prompt = config.get("limit_mm_per_prompt", {"image": len(images)})
-
-    # Create engine
-    llm = LLM(
-        model=config["model_name"],
-        trust_remote_code=True,
-        max_model_len=config["max_model_len"],
-        max_num_seqs=config["max_num_seqs"],
-        limit_mm_per_prompt=limit_mm_per_prompt,
-        mm_encoder_attn_backend=mm_encoder_attn_backend,
-        hf_overrides=dummy_hf_overrides,
-        load_format="dummy",
-        seed=42,
-    )
-
-    # Generate
-    sampling_params = SamplingParams(**config["sampling_params"])
-    outputs = llm.generate(
-        {
-            "prompt": prompt,
-            "multi_modal_data": {"image": images},
-        },
-        sampling_params=sampling_params,
-    )
-
-    # Validate
-    for o in outputs:
-        generated_text = o.outputs[0].text
-        validator = config.get("output_validator", lambda x: len(x) > 10)
-        assert validator(generated_text), (
-            f"Validation failed for {config['model_name']}: {generated_text}"
-        )
-
-
-def run_llm_chat_test(config, mm_encoder_attn_backend, image_assets):
-    """LLM.chat() interface handler for Dots.OCR."""
-    # Filter to stop_sign image only
-    stop_sign_image = [
-        asset.pil_image for asset in image_assets if asset.name == "stop_sign"
-    ][0]
-
-    # Build messages
-    messages = build_dots_ocr_prompt([stop_sign_image], config)
-
-    # Create engine
-    llm = LLM(
-        model=config["model_name"],
-        trust_remote_code=True,
-        max_model_len=config["max_model_len"],
-        max_num_seqs=config["max_num_seqs"],
-        limit_mm_per_prompt=config["limit_mm_per_prompt"],
-        mm_encoder_attn_backend=mm_encoder_attn_backend,
-        hf_overrides=dummy_hf_overrides,
-        load_format="dummy",
-        seed=42,
-    )
-
-    # Generate using chat
-    sampling_params = SamplingParams(**config["sampling_params"])
-    outputs = llm.chat(messages=messages, sampling_params=sampling_params)
-
-    # Validate
-    for o in outputs:
-        generated_text = o.outputs[0].text
-        validator = config.get("output_validator", lambda x: len(x) > 10)
-        assert validator(generated_text), (
-            f"Validation failed for {config['model_name']}: {generated_text}"
-        )
-
-
-def run_video_test(config, mm_encoder_attn_backend, video_assets, vllm_runner):
-    """Video test with EVS (Efficient Video Sampling) handler."""
-    for pruning_rate in config["video_params"]["pruning_rates"]:
-        num_frames = config["video_params"]["num_frames"]
-
-        # Sample frames from video
-        sampled_vids = [
-            sample_frames_from_video(asset.np_ndarrays, num_frames)
-            for asset in video_assets
-        ]
-
-        # Build prompt and prepare video
-        prompt = build_qwen2_5_video_prompt()
-        prompts = [prompt]
-        videos = [sampled_vids[0]]
-
-        # Run with vllm_runner context manager
-        with vllm_runner(
-            config["model_name"],
-            max_model_len=config["max_model_len"],
-            max_num_seqs=config["max_num_seqs"],
-            limit_mm_per_prompt=config["limit_mm_per_prompt"],
-            tensor_parallel_size=1,
-            video_pruning_rate=pruning_rate,
-            mm_encoder_attn_backend=mm_encoder_attn_backend,
-            hf_overrides=dummy_hf_overrides,
-            load_format="dummy",
-            **config["runner_kwargs"],
-        ) as vllm_model:
-            outputs = vllm_model.generate_greedy(
-                prompts,
-                config["sampling_params"]["max_tokens"],
-                videos=videos,
-            )
-
-            # Validate output
-            assert len(outputs) == 1, f"Expected 1 output, got {len(outputs)}"
-            output_ids, output_text = outputs[0]
-            assert len(output_ids) > 0, "Generated no output IDs"
-            assert len(output_text) > 0, "Generated empty text"
-            assert isinstance(output_text, str), (
-                f"Output is not string: {type(output_text)}"
-            )
-
-
-# Main test function
-@pytest.mark.parametrize("model_key", list(MODEL_CONFIGS.keys()))
-@pytest.mark.parametrize(
-    "mm_encoder_attn_backend",
-    [None] + current_platform.get_supported_vit_attn_backends(),
-)
-@pytest.mark.skip(reason="Broken test due to memory segmentation fault")
-@create_new_process_for_each_test()
-def test_vit_backend_functionality(
-    model_key: str,
-    mm_encoder_attn_backend: AttentionBackendEnum | None,
-    image_assets,
-    video_assets,
-    vllm_runner,
-    request,
-):
-    """Test ViT attention backend functionality for multimodal models.
-
-    This test validates that each model can successfully generate outputs
-    using different ViT attention backends. The test:
-    1. Filters unsupported backends per model
-    2. Applies appropriate GPU marks
-    3. Routes to the correct test handler based on interface
-    4. Validates output meets minimum requirements
-    """
-    config = MODEL_CONFIGS[model_key]
-
-    # Step 1: Backend filtering
-    if (
-        "supported_backends" in config
-        and mm_encoder_attn_backend is not None
-        and mm_encoder_attn_backend not in config["supported_backends"]
-    ):
-        pytest.skip(
-            f"{model_key} does not support {mm_encoder_attn_backend} backend now."
-        )
-
-    # Step 2: Apply GPU marks dynamically
-    if "gpu_marks" in config:
-        for mark in config["gpu_marks"]:
-            request.applymarker(mark)
-
-    # Step 3: Route to appropriate handler
-    if config.get("media_type") == "video":
-        run_video_test(config, mm_encoder_attn_backend, video_assets, vllm_runner)
-    elif config["interface"] == "llm_chat":
-        run_llm_chat_test(config, mm_encoder_attn_backend, image_assets)
-    elif config["interface"] == "llm_generate":
-        run_llm_generate_test(config, mm_encoder_attn_backend, image_assets)
-    else:
-        raise ValueError(f"Unknown interface: {config['interface']}")
@@ -171,6 +171,7 @@ MODEL_CONFIGS: dict[str, VitCudagraphTestConfig] = {
            "Describe this video in one sentence."
        ),
        needs_video_metadata=True,
+        vllm_runner_kwargs={"enable_chunked_prefill": True},
        marks=[pytest.mark.core_model],
    ),
    "internvl": VitCudagraphTestConfig(
@@ -382,6 +382,19 @@ _TEXT_GENERATION_EXAMPLE_MODELS = {
    "LongcatFlashForCausalLM": _HfExamplesInfo(
        "meituan-longcat/LongCat-Flash-Chat", trust_remote_code=True
    ),
+    "LongcatFlashNgramForCausalLM": _HfExamplesInfo(
+        "meituan-longcat/LongCat-Flash-Lite",
+        trust_remote_code=True,
+        # Shrink the ~62GB n-gram tables (ngram_vocab_size_ratio * vocab_size)
+        # so the dummy-weight init test fits in CI memory.
+        hf_overrides={"ngram_vocab_size_ratio": 1},
+    ),
+    "LongcatCausalLM": _HfExamplesInfo(
+        "meituan-longcat/LongCat-2.0-FP8",
+        # Shrink the huge n-gram tables (~264M rows at the checkpoint's
+        # oe_vocab_size_ratio=100.567) so dummy-weight init fits in CI memory.
+        hf_overrides={"ngram_vocab_size_ratio": 1},
+    ),
    "MambaForCausalLM": _HfExamplesInfo("state-spaces/mamba-130m-hf"),
    "Mamba2ForCausalLM": _HfExamplesInfo(
        "mistralai/Mamba-Codestral-7B-v0.1",
@@ -799,9 +812,9 @@ _MULTIMODAL_EXAMPLE_MODELS = {
    ),
    "Cosmos3ForConditionalGeneration": _HfExamplesInfo(
        "nvidia/Cosmos3-Nano",
+        extras={"super": "nvidia/Cosmos3-Super"},
        max_model_len=4096,
        min_transformers_version="4.57",
-        is_available_online=False,
    ),
    "DeepseekVLV2ForCausalLM": _HfExamplesInfo(
        "deepseek-ai/deepseek-vl2-tiny",
@@ -1403,7 +1416,7 @@ _SPECULATIVE_DECODING_EXAMPLE_MODELS = {
    # ),
    # [DFlash]
    "DFlashDraftModel": _HfExamplesInfo(
-        "Qwen/Qwen3.5-4B",
+        "Qwen/Qwen3-4B",
        speculative_model="z-lab/Qwen3-4B-DFlash-b16",
        use_original_num_layers=True,  # Need all layers since DFlash has >1 layer,
        max_model_len=8192,  # Reduce max len to ensure test runs in low-VRAM CI env
@@ -48,9 +48,11 @@ def test_registry_imports(model_arch):
                "(see #41376)"
            )

-    # DSpark draft model is NVIDIA-only; class is stubbed to None on ROCm/XPU.
-    if model_arch == "DSparkDraftModel" and not current_platform.is_cuda():
-        pytest.skip("DSparkDraftModel is only supported on CUDA")
+    # DSpark draft model is supported on CUDA and ROCm; stubbed to None on XPU.
+    if model_arch == "DSparkDraftModel" and not (
+        current_platform.is_cuda() or current_platform.is_rocm()
+    ):
+        pytest.skip("DSparkDraftModel is only supported on CUDA and ROCm")

    # Ensure all model classes can be imported successfully
    model_cls = ModelRegistry._try_load_model_cls(model_arch)
@@ -529,8 +529,13 @@ def dummy_hf_overrides(
            }
        )

-    # Update num_hidden_layers for non-Longcat architectures
-    if model_arch != "LongcatFlashForCausalLM" and model_arch != "LongCatFlashMTPModel":
+    # Update num_hidden_layers for non-Longcat architectures (Longcat derives it
+    # from num_layers for its dual-attention layers).
+    if model_arch not in (
+        "LongcatFlashForCausalLM",
+        "LongCatFlashMTPModel",
+        "LongcatFlashNgramForCausalLM",
+    ):
        update_dict["num_hidden_layers"] = num_hidden_layers

    text_config.update(update_dict)
@@ -49,4 +49,5 @@ def mock_request():
    req = MagicMock(spec=ChatCompletionRequest)
    req.tools = []
    req.tool_choice = "auto"
+    req.include_reasoning = True
    return req
@@ -43,6 +43,7 @@ def _make_request(**chat_template_kwargs):
    request = MagicMock(spec=ChatCompletionRequest)
    request.tools = []
    request.tool_choice = "auto"
+    request.include_reasoning = True
    request.chat_template_kwargs = chat_template_kwargs or None
    return request

@@ -896,6 +896,7 @@ def _make_delegating_request():
    req = MagicMock(spec=ChatCompletionRequest)
    req.tools = []
    req.tool_choice = "auto"
+    req.include_reasoning = True
    return req


@@ -0,0 +1,456 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""Tests for include_reasoning suppression in the unified Parser interface.
+
+Covers non-streaming (parser.parse() + build_response_output_items),
+streaming (parse_delta), and ParsableContext.append_output() paths.
+"""
+
+import json
+import os
+
+import pytest
+
+_STRICT_TOOL_CALLING_ENV = "VLLM_ENFORCE_STRICT_TOOL_CALLING"
+_STRICT_TOOL_CALLING_ENV_VALUE = os.environ.get(_STRICT_TOOL_CALLING_ENV)
+os.environ[_STRICT_TOOL_CALLING_ENV] = "0"
+
+from vllm.entrypoints.openai.chat_completion.protocol import (  # noqa: E402
+    ChatCompletionRequest,
+)
+from vllm.entrypoints.openai.engine.protocol import DeltaMessage  # noqa: E402
+from vllm.entrypoints.openai.responses.protocol import ResponsesRequest  # noqa: E402
+from vllm.parser.abstract_parser import DelegatingParser  # noqa: E402
+from vllm.reasoning.basic_parsers import BaseThinkingReasoningParser  # noqa: E402
+from vllm.tool_parsers.hermes_tool_parser import Hermes2ProToolParser  # noqa: E402
+
+
+@pytest.fixture(scope="module", autouse=True)
+def restore_strict_tool_calling_env():
+    yield
+    if _STRICT_TOOL_CALLING_ENV_VALUE is None:
+        os.environ.pop(_STRICT_TOOL_CALLING_ENV, None)
+    else:
+        os.environ[_STRICT_TOOL_CALLING_ENV] = _STRICT_TOOL_CALLING_ENV_VALUE
+
+
+class ThinkReasoningParser(BaseThinkingReasoningParser):
+    @property
+    def start_token(self) -> str:
+        return "<think>"
+
+    @property
+    def end_token(self) -> str:
+        return "</think>"
+
+
+MODEL_OUTPUT_REASONING_AND_CONTENT = (
+    "<think>let me think about this</think>The answer is 42."
+)
+
+MODEL_OUTPUT_REASONING_AND_TOOL = (
+    "<think>I need to call a tool</think>"
+    '<tool_call>\n{"name": "get_weather", '
+    '"arguments": {"city": "Dallas"}}\n</tool_call>'
+)
+
+MODEL_OUTPUT_CONTENT_ONLY = "The answer is 42."
+
+
+@pytest.fixture(scope="module")
+def tokenizer():
+    from vllm.tokenizers import get_tokenizer
+
+    return get_tokenizer("Qwen/Qwen3-32B")
+
+
+def make_responses_request(**kwargs) -> ResponsesRequest:
+    defaults = dict(model="test-model", input="test input")
+    defaults.update(kwargs)
+    return ResponsesRequest(**defaults)
+
+
+def make_chat_request(**kwargs) -> ChatCompletionRequest:
+    defaults = dict(
+        model="test-model",
+        messages=[{"role": "user", "content": "hi"}],
+    )
+    defaults.update(kwargs)
+    return ChatCompletionRequest(**defaults)
+
+
+def make_parser(tokenizer, reasoning=False, tool=False):
+    class TestParser(DelegatingParser):
+        reasoning_parser_cls = ThinkReasoningParser if reasoning else None
+        tool_parser_cls = Hermes2ProToolParser if tool else None
+
+    return TestParser(tokenizer)
+
+
+# ── Non-streaming: parser.parse() + build_response_output_items ──────
+
+
+def parse_and_build(parser, request, model_output, enable_auto_tools=False):
+    """Mirrors the non-streaming path in _make_response_output_items /
+    ParsableContext.append_output(): parse → suppress reasoning → build items.
+    """
+    from vllm.entrypoints.openai.responses.utils import (
+        build_response_output_items,
+    )
+
+    reasoning, content, tool_calls = parser.parse(
+        model_output, request, enable_auto_tools=enable_auto_tools
+    )
+    if not request.include_reasoning:
+        reasoning = None
+    return build_response_output_items(
+        reasoning=reasoning,
+        content=content,
+        tool_calls=tool_calls,
+    )
+
+
+class TestNonStreamingIncludeReasoning:
+    def test_include_reasoning_true_has_reasoning_item(self, tokenizer):
+        """Default: reasoning items appear in output."""
+        parser = make_parser(tokenizer, reasoning=True)
+        request = make_responses_request(include_reasoning=True)
+
+        outputs = parse_and_build(parser, request, MODEL_OUTPUT_REASONING_AND_CONTENT)
+
+        types = [o.type for o in outputs]
+        assert "reasoning" in types
+        assert "message" in types
+
+    def test_include_reasoning_false_no_reasoning_item(self, tokenizer):
+        """Reasoning item is suppressed when include_reasoning=False."""
+        parser = make_parser(tokenizer, reasoning=True)
+        request = make_responses_request(include_reasoning=False)
+
+        outputs = parse_and_build(parser, request, MODEL_OUTPUT_REASONING_AND_CONTENT)
+
+        types = [o.type for o in outputs]
+        assert "reasoning" not in types
+        assert "message" in types
+        assert outputs[0].content[0].text == "The answer is 42."
+
+    def test_include_reasoning_false_content_preserved(self, tokenizer):
+        """Content is extracted correctly even when reasoning is suppressed."""
+        parser = make_parser(tokenizer, reasoning=True)
+        request = make_responses_request(include_reasoning=False)
+
+        outputs = parse_and_build(parser, request, MODEL_OUTPUT_REASONING_AND_CONTENT)
+
+        message = next(o for o in outputs if o.type == "message")
+        assert message.content[0].text == "The answer is 42."
+
+    def test_include_reasoning_false_tool_calls_preserved(self, tokenizer):
+        """Tool calls still work when reasoning is suppressed."""
+        parser = make_parser(tokenizer, reasoning=True, tool=True)
+        request = make_responses_request(
+            include_reasoning=False,
+            tools=[
+                {
+                    "type": "function",
+                    "name": "get_weather",
+                    "parameters": {
+                        "type": "object",
+                        "properties": {"city": {"type": "string"}},
+                    },
+                }
+            ],
+        )
+
+        outputs = parse_and_build(
+            parser,
+            request,
+            MODEL_OUTPUT_REASONING_AND_TOOL,
+            enable_auto_tools=True,
+        )
+
+        types = [o.type for o in outputs]
+        assert "reasoning" not in types
+        assert "function_call" in types
+        fc = next(o for o in outputs if o.type == "function_call")
+        assert fc.name == "get_weather"
+        assert json.loads(fc.arguments) == {"city": "Dallas"}
+
+    def test_no_reasoning_parser_include_false_is_noop(self, tokenizer):
+        """include_reasoning=False is harmless when no reasoning parser."""
+        parser = make_parser(tokenizer, reasoning=False)
+        request = make_responses_request(include_reasoning=False)
+
+        outputs = parse_and_build(parser, request, MODEL_OUTPUT_CONTENT_ONLY)
+
+        assert len(outputs) == 1
+        assert outputs[0].type == "message"
+        assert outputs[0].content[0].text == MODEL_OUTPUT_CONTENT_ONLY
+
+    def test_default_include_reasoning_is_true(self, tokenizer):
+        """ResponsesRequest defaults to include_reasoning=True."""
+        request = make_responses_request()
+        assert request.include_reasoning is True
+
+    def test_include_reasoning_false_suppresses_all_reasoning(self, tokenizer):
+        """Reasoning is suppressed regardless of request type."""
+        parser = make_parser(tokenizer, reasoning=True)
+        request = make_responses_request(include_reasoning=False)
+
+        outputs = parse_and_build(parser, request, MODEL_OUTPUT_REASONING_AND_CONTENT)
+
+        assert all(o.type != "reasoning" for o in outputs)
+
+
+# ── Streaming: parse_delta ───────────────────────────────────────────
+
+
+def stream_text(parser, tokenizer, text, request, prompt_token_ids=None):
+    token_ids = tokenizer.encode(text, add_special_tokens=False)
+    results: list[DeltaMessage | None] = []
+    for i, tid in enumerate(token_ids):
+        delta_text = tokenizer.decode([tid])
+        is_last = i == len(token_ids) - 1
+        result = parser.parse_delta(
+            delta_text,
+            [tid],
+            request,
+            prompt_token_ids=prompt_token_ids,
+            finished=is_last,
+        )
+        prompt_token_ids = None
+        results.append(result)
+    return results
+
+
+def collect_fields(results):
+    all_reasoning = "".join(r.reasoning for r in results if r and r.reasoning)
+    all_content = "".join(r.content for r in results if r and r.content)
+    all_tool_calls = [tc for r in results if r and r.tool_calls for tc in r.tool_calls]
+    return all_reasoning, all_content, all_tool_calls
+
+
+class TestParseDeltaIncludeReasoning:
+    def test_streaming_include_true_emits_reasoning(self, tokenizer):
+        """With include_reasoning=True, reasoning deltas are emitted."""
+        parser = make_parser(tokenizer, reasoning=True)
+        request = make_responses_request(include_reasoning=True)
+
+        results = stream_text(
+            parser,
+            tokenizer,
+            MODEL_OUTPUT_REASONING_AND_CONTENT,
+            request,
+            prompt_token_ids=[],
+        )
+        reasoning, content, _ = collect_fields(results)
+
+        assert "let me think about this" in reasoning
+        assert "42" in content
+
+    def test_streaming_include_false_suppresses_reasoning(self, tokenizer):
+        """With include_reasoning=False, no reasoning deltas are emitted."""
+        parser = make_parser(tokenizer, reasoning=True)
+        request = make_responses_request(include_reasoning=False)
+
+        results = stream_text(
+            parser,
+            tokenizer,
+            MODEL_OUTPUT_REASONING_AND_CONTENT,
+            request,
+            prompt_token_ids=[],
+        )
+        reasoning, content, _ = collect_fields(results)
+
+        assert reasoning == ""
+        assert "42" in content
+
+    def test_streaming_include_false_content_still_works(self, tokenizer):
+        """Content is correctly extracted in streaming even with suppression."""
+        parser = make_parser(tokenizer, reasoning=True)
+        request = make_responses_request(include_reasoning=False)
+
+        results = stream_text(
+            parser,
+            tokenizer,
+            MODEL_OUTPUT_REASONING_AND_CONTENT,
+            request,
+            prompt_token_ids=[],
+        )
+        _, content, _ = collect_fields(results)
+
+        assert "The answer is 42" in content
+
+    def test_streaming_include_false_tool_calls_preserved(self, tokenizer):
+        """Tool calls stream correctly when reasoning is suppressed."""
+        parser = make_parser(tokenizer, reasoning=True, tool=True)
+        request = make_responses_request(
+            include_reasoning=False,
+            tools=[
+                {
+                    "type": "function",
+                    "name": "get_weather",
+                    "parameters": {
+                        "type": "object",
+                        "properties": {"city": {"type": "string"}},
+                    },
+                }
+            ],
+        )
+
+        results = stream_text(
+            parser,
+            tokenizer,
+            MODEL_OUTPUT_REASONING_AND_TOOL,
+            request,
+            prompt_token_ids=[],
+        )
+        reasoning, content, tool_calls = collect_fields(results)
+
+        assert reasoning == ""
+        assert len(tool_calls) > 0
+        assert tool_calls[0].function.name == "get_weather"
+        tool_args = "".join(
+            tc.function.arguments for tc in tool_calls if tc.function.arguments
+        )
+        assert json.loads(tool_args) == {"city": "Dallas"}
+
+    def test_streaming_no_reasoning_parser_include_false(self, tokenizer):
+        """No crash when reasoning parser absent and include_reasoning=False."""
+        parser = make_parser(tokenizer, reasoning=False)
+        request = make_responses_request(include_reasoning=False)
+
+        results = stream_text(
+            parser,
+            tokenizer,
+            MODEL_OUTPUT_CONTENT_ONLY,
+            request,
+            prompt_token_ids=[],
+        )
+        reasoning, content, _ = collect_fields(results)
+
+        assert reasoning == ""
+        assert "42" in content
+
+    def test_streaming_chat_completion_include_false(self, tokenizer):
+        """parse_delta also respects ChatCompletionRequest.include_reasoning."""
+        parser = make_parser(tokenizer, reasoning=True)
+        request = make_chat_request(include_reasoning=False)
+
+        results = stream_text(
+            parser,
+            tokenizer,
+            MODEL_OUTPUT_REASONING_AND_CONTENT,
+            request,
+            prompt_token_ids=[],
+        )
+        reasoning, content, _ = collect_fields(results)
+
+        assert reasoning == ""
+        assert "42" in content
+
+    def test_streaming_reasoning_only_deltas_become_none(self, tokenizer):
+        """Deltas that carry only reasoning become None (not empty)."""
+        parser = make_parser(tokenizer, reasoning=True)
+        request = make_responses_request(include_reasoning=False)
+
+        results = stream_text(
+            parser,
+            tokenizer,
+            MODEL_OUTPUT_REASONING_AND_CONTENT,
+            request,
+            prompt_token_ids=[],
+        )
+
+        for r in results:
+            if r is not None:
+                assert r.reasoning is None
+
+
+# ── ParsableContext.append_output() ───────────────────────────────────
+
+
+class TestParsableContextIncludeReasoning:
+    def _make_context(self, tokenizer, request):
+        from vllm.entrypoints.openai.responses.context import ParsableContext
+
+        class TestParser(DelegatingParser):
+            reasoning_parser_cls = ThinkReasoningParser
+            tool_parser_cls = None
+
+        return ParsableContext(
+            tokenizer=tokenizer,
+            parser_cls=TestParser,
+            response_messages=[],
+            request=request,
+            available_tools=None,
+            chat_template=None,
+            chat_template_content_format="auto",
+        )
+
+    def test_process_include_false_suppresses_reasoning(self, tokenizer):
+        """ParsableContext.process() suppresses reasoning items."""
+        from vllm.outputs import CompletionOutput, RequestOutput
+
+        request = make_responses_request(include_reasoning=False)
+        ctx = self._make_context(tokenizer, request)
+
+        output = RequestOutput(
+            request_id="test",
+            prompt=None,
+            prompt_token_ids=[],
+            prompt_logprobs=None,
+            outputs=[
+                CompletionOutput(
+                    index=0,
+                    text=MODEL_OUTPUT_REASONING_AND_CONTENT,
+                    token_ids=tokenizer.encode(
+                        MODEL_OUTPUT_REASONING_AND_CONTENT,
+                        add_special_tokens=False,
+                    ),
+                    cumulative_logprob=None,
+                    logprobs=None,
+                    finish_reason="stop",
+                )
+            ],
+            finished=True,
+        )
+
+        ctx.append_output(output)
+
+        types = [getattr(m, "type", None) for m in ctx.response_messages]
+        assert "reasoning" not in types
+        assert "message" in types
+
+    def test_process_include_true_has_reasoning(self, tokenizer):
+        """ParsableContext.process() includes reasoning by default."""
+        from vllm.outputs import CompletionOutput, RequestOutput
+
+        request = make_responses_request(include_reasoning=True)
+        ctx = self._make_context(tokenizer, request)
+
+        output = RequestOutput(
+            request_id="test",
+            prompt=None,
+            prompt_token_ids=[],
+            prompt_logprobs=None,
+            outputs=[
+                CompletionOutput(
+                    index=0,
+                    text=MODEL_OUTPUT_REASONING_AND_CONTENT,
+                    token_ids=tokenizer.encode(
+                        MODEL_OUTPUT_REASONING_AND_CONTENT,
+                        add_special_tokens=False,
+                    ),
+                    cumulative_logprob=None,
+                    logprobs=None,
+                    finish_reason="stop",
+                )
+            ],
+            finished=True,
+        )
+
+        ctx.append_output(output)
+
+        types = [getattr(m, "type", None) for m in ctx.response_messages]
+        assert "reasoning" in types
+        assert "message" in types
@@ -1121,14 +1121,14 @@ def test_is_chunked_prefill_supported(
        (
            "Qwen/Qwen3-Next-80B-A3B-Instruct",
            "hybrid",
-            False,
-            "Hybrid models do not support prefix caching since the feature is still experimental.",  # noqa: E501
+            True,
+            "Generative hybrid models support prefix caching.",  # noqa: E501
        ),
        (
            "ibm-granite/granite-4.0-h-small",
            "hybrid",
-            False,
-            "Hybrid models do not support prefix caching since the feature is still experimental.",  # noqa: E501
+            True,
+            "Generative hybrid models support prefix caching.",  # noqa: E501
        ),
        (
            "state-spaces/mamba-130m-hf",
@@ -12,7 +12,12 @@ import torch

 import vllm.v1.core.kv_cache_manager as kv_cache_manager
 import vllm.v1.core.kv_cache_utils as kv_cache_utils
-from vllm.distributed.kv_events import AllBlocksCleared, BlockRemoved, BlockStored
+from vllm.distributed.kv_events import (
+    MEDIUM_GPU,
+    AllBlocksCleared,
+    BlockRemoved,
+    BlockStored,
+)
 from vllm.lora.request import LoRARequest
 from vllm.multimodal.inputs import (
    MultiModalFeatureSpec,
@@ -2454,6 +2459,118 @@ def test_block_removed_event_group_idx(group_id: int):
        assert event.group_idx == group_id


+def test_emit_cached_block_events():
+    """emit_cached_block_events emits one BlockStored for already-cached
+    (reused) prefix blocks, carrying the correct group_idx /
+    parent_block_hash / token_ids, and without mutating block state."""
+    block_size = 4
+    num_cached_blocks = 3
+    kv_cache_group_id = 1
+    num_tokens = block_size * 4  # 4 full blocks; reuse the first 3
+
+    pool = BlockPool(
+        num_gpu_blocks=8,
+        enable_caching=True,
+        hash_block_size=block_size,
+        enable_kv_cache_events=True,
+    )
+
+    req = make_request(
+        "req_emit_cached",
+        prompt_token_ids=list(range(num_tokens)),
+        block_size=block_size,
+        hash_fn=sha256,
+    )
+    assert len(req.block_hashes) >= num_cached_blocks
+
+    # Snapshot block state to prove emit_cached_block_events does not mutate it.
+    free_before = pool.get_num_free_blocks()
+    assert len(pool.cached_block_hash_to_block) == 0
+
+    pool.emit_cached_block_events(
+        request=req,
+        num_cached_blocks=num_cached_blocks,
+        block_size=block_size,
+        kv_cache_group_id=kv_cache_group_id,
+    )
+
+    # No block-state mutation: nothing allocated, nothing inserted into the
+    # prefix-cache map.
+    assert pool.get_num_free_blocks() == free_before
+    assert len(pool.cached_block_hash_to_block) == 0
+
+    events = pool.take_events()
+    assert len(events) == 1
+    event = events[0]
+    assert isinstance(event, BlockStored)
+
+    expected_hashes = [
+        kv_cache_utils.maybe_convert_block_hash(req.block_hashes[i])
+        for i in range(num_cached_blocks)
+    ]
+    assert event.block_hashes == expected_hashes
+    # Reused blocks start from block 0, so there is no parent block hash.
+    assert event.parent_block_hash is None
+    assert event.token_ids == list(req.all_token_ids[: num_cached_blocks * block_size])
+    assert event.group_idx == kv_cache_group_id
+    assert event.block_size == block_size
+    assert event.medium == MEDIUM_GPU
+    assert event.lora_id is None
+    assert event.lora_name is None
+
+
+def test_emit_cached_block_events_disabled():
+    """No events are emitted when enable_kv_cache_events is False."""
+    block_size = 4
+    pool = BlockPool(
+        num_gpu_blocks=8,
+        enable_caching=True,
+        hash_block_size=block_size,
+        enable_kv_cache_events=False,
+    )
+    req = make_request(
+        "req_emit_disabled",
+        prompt_token_ids=list(range(block_size * 4)),
+        block_size=block_size,
+        hash_fn=sha256,
+    )
+
+    pool.emit_cached_block_events(
+        request=req,
+        num_cached_blocks=3,
+        block_size=block_size,
+        kv_cache_group_id=0,
+    )
+
+    assert pool.take_events() == []
+
+
+def test_emit_cached_block_events_zero_cached():
+    """No events are emitted when num_cached_blocks == 0."""
+    block_size = 4
+    pool = BlockPool(
+        num_gpu_blocks=8,
+        enable_caching=True,
+        hash_block_size=block_size,
+        enable_kv_cache_events=True,
+    )
+    req = make_request(
+        "req_emit_zero",
+        prompt_token_ids=list(range(block_size * 4)),
+        block_size=block_size,
+        hash_fn=sha256,
+    )
+
+    pool.emit_cached_block_events(
+        request=req,
+        num_cached_blocks=0,
+        block_size=block_size,
+        kv_cache_group_id=0,
+    )
+
+    assert pool.take_events() == []
+
+
 def test_eagle_enabled_removes_last_block():
    """Verify Eagle does NOT remove blocks when request
    length is divisible by block size."""
@@ -38,6 +38,7 @@ def mock_model_runner_with_input_batch():
        vocab_size=32000,
        block_sizes=[16],
        kernel_block_sizes=[16],
+        max_num_blocks_per_req=[64],
        logitsprocs=None,
        is_pooling_model=False,
    )
@@ -0,0 +1,45 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+import pytest
+import torch
+
+from vllm.v1.attention.backends.utils import get_dcp_local_seq_lens
+from vllm.v1.worker.cp_utils import should_skip_dcp_context_attention
+
+
+def test_skip_gate_only_for_zero_context():
+    assert should_skip_dcp_context_attention(torch.zeros(3, dtype=torch.int32))
+    assert not should_skip_dcp_context_attention(
+        torch.tensor([0, 5, 0], dtype=torch.int32)
+    )
+
+
+@pytest.mark.parametrize(
+    "dcp_world_size,interleave_size,context_len",
+    [(2, 16, 10), (4, 16, 10), (8, 16, 10), (4, 1, 2)],
+)
+def test_skip_gate_rank_invariant_with_divergent_local_context(
+    dcp_world_size: int, interleave_size: int, context_len: int
+):
+    """Contexts shorter than a full interleave round land entirely on a
+    subset of DCP ranks, so the per-rank local context lengths diverge:
+    some ranks hold zero local context while others hold all of it. Ranks
+    with zero local context must still take the collective (non-skip) path,
+    otherwise the query all-gather in _forward_with_dcp deadlocks across
+    ranks. The skip gate must therefore depend only on the rank-invariant
+    global context lengths, never on get_dcp_local_seq_lens output.
+    """
+    context_kv_lens = torch.tensor([context_len], dtype=torch.int32)
+    local_maxes = [
+        int(
+            get_dcp_local_seq_lens(
+                context_kv_lens, dcp_world_size, rank, interleave_size
+            ).max()
+        )
+        for rank in range(dcp_world_size)
+    ]
+    # Precondition: the local view diverges across ranks.
+    assert 0 in local_maxes
+    assert max(local_maxes) > 0
+    # The batch still has context globally, so no rank may skip.
+    assert not should_skip_dcp_context_attention(context_kv_lens)
@@ -238,6 +238,7 @@ def test_sampling_metadata_in_input_batch(device: str, batch_size: int):
        vocab_size=1024,
        block_sizes=[1],
        kernel_block_sizes=[1],
+        max_num_blocks_per_req=[1024],
    )
    reqs: list[CachedRequestState] = []
    req_id_reqs = {}
@@ -332,6 +333,7 @@ def test_swap_states_in_input_batch(device: str, batch_size: int, swap_list: lis
        vocab_size=1024,
        block_sizes=[1],
        kernel_block_sizes=[1],
+        max_num_blocks_per_req=[1024],
    )
    ref_input_batch: InputBatch = InputBatch(
        max_num_reqs=batch_size,
@@ -341,6 +343,7 @@ def test_swap_states_in_input_batch(device: str, batch_size: int, swap_list: lis
        vocab_size=1024,
        block_sizes=[1],
        kernel_block_sizes=[1],
+        max_num_blocks_per_req=[1024],
    )

    reqs: list[CachedRequestState] = []
@@ -409,6 +412,7 @@ def test_pooling_prompt_lens_not_aliased(device: str):
        vocab_size=VOCAB_SIZE,
        block_sizes=[16],
        kernel_block_sizes=[16],
+        max_num_blocks_per_req=[64],
        is_pooling_model=True,
    )

@@ -444,6 +448,7 @@ def test_placeholder_spec_token_ids_written_verbatim():
        vocab_size=VOCAB_SIZE,
        block_sizes=[16],
        kernel_block_sizes=[16],
+        max_num_blocks_per_req=[1],
    )
    req = CachedRequestState(
        req_id="req",
@@ -491,6 +496,7 @@ def test_pooling_metadata_token_id_buffers(
        vocab_size=VOCAB_SIZE,
        block_sizes=[16],
        kernel_block_sizes=[16],
+        max_num_blocks_per_req=[64],
        is_pooling_model=True,
    )
    req = _construct_pooling_request(0, PoolingParams(**pooling_params))
@@ -89,6 +89,7 @@ def initialize_kv_cache(runner: GPUModelRunner):
        kernel_block_sizes=[
            kv_cache_config.kv_cache_groups[0].kv_cache_spec.block_size
        ],
+        max_num_blocks_per_req=[NUM_BLOCKS],
    )
    runner.initialize_attn_backend(kv_cache_config)

@@ -1397,6 +1398,7 @@ def test_input_batch_with_kernel_block_sizes():
        vocab_size=vocab_size,
        block_sizes=block_sizes,
        kernel_block_sizes=kernel_block_sizes,
+        max_num_blocks_per_req=[16, 8],
    )

    # Verify that block tables were created with kernel block sizes
@@ -1457,6 +1459,7 @@ def test_hybrid_cache_integration(default_vllm_config, dist_init):
        vocab_size=runner.model_config.get_vocab_size(),
        block_sizes=[kv_cache_config.kv_cache_groups[0].kv_cache_spec.block_size],
        kernel_block_sizes=[16],
+        max_num_blocks_per_req=[NUM_BLOCKS],
    )  # Use kernel block size

    runner.initialize_attn_backend(kv_cache_config)
@@ -235,6 +235,37 @@ def rms_norm(
    torch.ops._C.rms_norm(out, input, weight, epsilon)


+# LongCat n-gram embedding index kernel (see csrc/.../ngram_embedding_kernels.cu).
+def ngram_compute_n_gram_ids(
+    ne_n: int,
+    ne_k: int,
+    ne_weights: torch.Tensor,
+    ne_mods: torch.Tensor,
+    exclusive_ne_embedder_size_sums: torch.Tensor,
+    exclusive_req_len_sums: torch.Tensor,
+    ne_token_table: torch.Tensor,
+    row_indices: torch.Tensor,
+    column_starts: torch.Tensor,
+    n_gram_ids: torch.Tensor,
+) -> None:
+    """Compute concatenated (offset) n-gram ids for a ragged prefill batch.
+
+    Writes ``n_gram_ids`` of shape ``[token_num, (ne_n-1)*ne_k]``.
+    """
+    torch.ops._C.ngram_compute_n_gram_ids(
+        ne_n,
+        ne_k,
+        ne_weights,
+        ne_mods,
+        exclusive_ne_embedder_size_sums,
+        exclusive_req_len_sums,
+        ne_token_table,
+        row_indices,
+        column_starts,
+        n_gram_ids,
+    )
+
+
 def fused_add_rms_norm(
    input: torch.Tensor,
    residual: torch.Tensor,
@@ -1881,11 +1881,8 @@ class ModelConfig:
        else:
            # for generative models
            if attn_type == "hybrid":
-                logger.debug(
-                    "Hybrid models do not support prefix caching since the feature "
-                    "is still experimental."
-                )
-                return False
+                logger.debug("Generative hybrid models support prefix caching.")
+                return True
            elif attn_type == "attention_free":
                logger.debug(
                    "Attention free models do not support prefix caching since the "
@@ -518,9 +518,15 @@ class SpeculativeConfig:
                    "architectures": ["Qwen3_5MoeMTP" if is_moe else "Qwen3_5MTP"],
                }
            )
-        if hf_config.model_type == "longcat_flash":
+        if hf_config.model_type in ("longcat_flash", "longcat_flash_ngram"):
            hf_config.model_type = "longcat_flash_mtp"
-            n_predict = getattr(hf_config, "num_nextn_predict_layers", 1)
+            # LongCat-2.0 ships one MTP module applied for up to
+            # mtp_num_layers draft steps (mtp_replicate_modules).
+            n_predict = (
+                getattr(hf_config, "num_nextn_predict_layers", None)
+                or getattr(hf_config, "mtp_num_layers", None)
+                or 1
+            )
            hf_config.update(
                {"n_predict": n_predict, "architectures": ["LongCatFlashMTPModel"]}
            )
@@ -942,6 +948,31 @@ class SpeculativeConfig:
                        "`num_speculative_tokens` was not provided"
                    )

+                if self.method == "dspark":
+                    # DSpark is a semi-autoregressive *block* drafter. A
+                    # speculative length smaller than the checkpoint's block
+                    # feeds the block / Markov-head machinery an unsupported
+                    # layout and yields incorrect (garbled) output rather than
+                    # merely lower acceptance. Require num_speculative_tokens to
+                    # be at least the block size (e.g. 5 or 7 for DeepSeek-V4).
+                    dspark_block_size = getattr(
+                        self.draft_model_config.hf_config,
+                        "dspark_block_size",
+                        None,
+                    )
+                    if (
+                        dspark_block_size is not None
+                        and self.num_speculative_tokens < dspark_block_size
+                    ):
+                        raise ValueError(
+                            "DSpark requires num_speculative_tokens >= "
+                            f"dspark_block_size ({dspark_block_size}); got "
+                            f"{self.num_speculative_tokens}. Smaller values "
+                            "produce incorrect output. Use "
+                            f"num_speculative_tokens={dspark_block_size} or "
+                            "larger (e.g. 7)."
+                        )
+
                self.draft_tensor_parallel_size = (
                    SpeculativeConfig._verify_and_get_draft_tp(
                        self.target_parallel_config,
@@ -70,6 +70,8 @@ DEFAULT_V2_MODEL_RUNNER_ARCHITECTURES = frozenset(
        "DeepseekV2ForCausalLM",
        "Qwen2MoeForCausalLM",
        "GraniteMoeForCausalLM",
+        "LongcatFlashNgramForCausalLM",
+        "LongcatCausalLM",
    }
 )

@@ -341,13 +341,30 @@ class CudaCommunicator(DeviceCommunicatorBase):
    def all_gather(self, input_: torch.Tensor, dim: int = -1) -> torch.Tensor:
        # Route uniform dim-0 all-gathers through NVLS symmetric memory when
        # enabled (mirrors reduce_scatter); otherwise fall back to the
-        # base-class ring all-gather. Sequence parallelism's gather-before-GEMM
-        # uses dim=0 with tp-aligned (uniform) shards.
+        # PyNccl/base-class all-gather. Sequence parallelism's
+        # gather-before-GEMM uses dim=0 with tp-aligned (uniform) shards.
        if dim < 0:
            dim += input_.dim()
        if dim == 0 and should_nccl_symm_mem_ag_rs():
            return self._all_gather_symm_mem(input_.contiguous())
-        return super().all_gather(input_, dim)
+
+        pynccl_comm = self.pynccl_comm
+        if pynccl_comm is None or pynccl_comm.disabled:
+            return super().all_gather(input_, dim)
+
+        input_size = input_.size()
+        output_size = (input_size[0] * self.world_size,) + input_size[1:]
+        output_tensor = torch.empty(
+            output_size, dtype=input_.dtype, device=input_.device
+        )
+        pynccl_comm.all_gather(output_tensor, input_.contiguous())
+        output_tensor = output_tensor.reshape((self.world_size,) + input_size)
+        output_tensor = output_tensor.movedim(0, dim)
+        return output_tensor.reshape(
+            input_size[:dim]
+            + (self.world_size * input_size[dim],)
+            + input_size[dim + 1 :]
+        )

    def reduce_scatter(self, input_: torch.Tensor, dim: int = -1):
        world_size = self.world_size
@@ -2493,7 +2493,11 @@ class EngineArgs:
        self, model_config: ModelConfig
    ) -> None:
        default_chunked_prefill = model_config.is_chunked_prefill_supported
-        default_prefix_caching = model_config.is_prefix_caching_supported
+        # Hybrid models support prefix caching but keep it opt-in for now
+        # while the feature matures.
+        default_prefix_caching = (
+            model_config.is_prefix_caching_supported and not model_config.is_hybrid
+        )

        if self.enable_chunked_prefill is None:
            self.enable_chunked_prefill = default_chunked_prefill
@@ -600,19 +600,8 @@ class OpenAIServingChat(GenerateBaseServing):
                            prompt_token_ids=res.prompt_token_ids,
                            finished=output.finish_reason is not None,
                        )
-                        if delta_message is not None:
-                            if delta_message.tool_calls:
-                                tools_streamed[i] = True
-
-                            if (
-                                delta_message.reasoning
-                                and not request.include_reasoning
-                            ):
-                                delta_message.reasoning = None
-                                if not (
-                                    delta_message.content or delta_message.tool_calls
-                                ):
-                                    delta_message = None
+                        if delta_message is not None and delta_message.tool_calls:
+                            tools_streamed[i] = True

                    # handle streaming just a content delta (no parsers)
                    else:
@@ -627,13 +616,22 @@ class OpenAIServingChat(GenerateBaseServing):
                    # "control token" for tool calls or the parser otherwise
                    # wasn't ready to send a token, then
                    #   get the next token without streaming a chunk
+                    # When reasoning is hidden, suppress per-token
+                    # metadata (logprobs, token_ids) on every chunk to
+                    # prevent leaking reasoning tokens through decoded
+                    # token text in logprob entries or raw token IDs.
+                    hide_stream_metadata = (
+                        not request.include_reasoning and parser is not None
+                    )
+                    if hide_stream_metadata:
+                        logprobs = None
+
                    if delta_message is None:
                        # NOTE: If return_token_ids is enabled, we still need to
                        # send a chunk with token_ids even if delta_message is None
                        # to ensure all tokens are included in the response
-                        if (
-                            output.finish_reason is None
-                            and not request.return_token_ids
+                        if output.finish_reason is None and (
+                            not request.return_token_ids or hide_stream_metadata
                        ):
                            continue
                        delta_message = DeltaMessage()
@@ -666,6 +664,10 @@ class OpenAIServingChat(GenerateBaseServing):
                                delta=True,
                            )

+                    include_token_ids = (
+                        request.return_token_ids and not hide_stream_metadata
+                    )
+
                    if output.finish_reason is None:
                        # Send token-by-token response for each request.n
                        choice_data = ChatCompletionResponseStreamChoice(
@@ -674,9 +676,7 @@ class OpenAIServingChat(GenerateBaseServing):
                            logprobs=logprobs,
                            finish_reason=None,
                            token_ids=(
-                                as_list(output.token_ids)
-                                if request.return_token_ids
-                                else None
+                                as_list(output.token_ids) if include_token_ids else None
                            ),
                        )

@@ -704,9 +704,7 @@ class OpenAIServingChat(GenerateBaseServing):
                            finish_reason=finish_reason_,
                            stop_reason=output.stop_reason,
                            token_ids=(
-                                as_list(output.token_ids)
-                                if request.return_token_ids
-                                else None
+                                as_list(output.token_ids) if include_token_ids else None
                            ),
                        )

@@ -883,12 +881,16 @@ class OpenAIServingChat(GenerateBaseServing):
                    enable_auto_tools=self.enable_auto_tools,
                    model_output_token_ids=token_ids,
                )
+                suppress_metadata = not request.include_reasoning and parser is not None
                if not request.include_reasoning:
                    reasoning = None
+                if suppress_metadata:
+                    logprobs = None
            else:
                reasoning = None
                content = output.text
                tool_calls = []
+                suppress_metadata = False

            auto_tools_called = False
            is_named_tool_choice = (
@@ -982,7 +984,9 @@ class OpenAIServingChat(GenerateBaseServing):
                else "stop",
                stop_reason=output.stop_reason,
                token_ids=(
-                    as_list(output.token_ids) if request.return_token_ids else None
+                    as_list(output.token_ids)
+                    if request.return_token_ids and not suppress_metadata
+                    else None
                ),
                routed_experts=routed_experts_b64,
            )
@@ -310,7 +310,9 @@ class ParsableContext(ConversationContext):
        self.finish_reason: str | None = None
        self.enable_auto_tools = enable_auto_tools

-        self.response_parser = response_parser
+        self.response_parser = response_parser or (
+            parser_cls(tokenizer, request.tools) if parser_cls is not None else None
+        )
        self.parser_cls = parser_cls
        self.request = request

@@ -344,6 +346,8 @@ class ParsableContext(ConversationContext):
                enable_auto_tools=self.enable_auto_tools,
                model_output_token_ids=completion.token_ids,
            )
+            if not self.request.include_reasoning:
+                reasoning = None
            self.response_messages.extend(
                build_response_output_items(
                    reasoning=reasoning,
@@ -161,6 +161,15 @@ class ResponsesRequest(OpenAIBaseModel):
    previous_response_id: str | None = None
    prompt: ResponsePrompt | None = None
    reasoning: Reasoning | None = None
+    include_reasoning: bool = Field(
+        default=True,
+        description=(
+            "Whether to include reasoning content in the response. "
+            "When false, reasoning tokens are still generated but "
+            "excluded from the output. This reduces network traffic "
+            "without affecting model inference."
+        ),
+    )
    service_tier: Literal["auto", "default", "flex", "scale", "priority"] = "auto"
    store: bool | None = True
    stream: bool | None = False
@@ -1068,6 +1068,9 @@ class OpenAIServingResponses(GenerateBaseServing):
                enable_auto_tools=self.enable_auto_tools,
                model_output_token_ids=final_output.token_ids,
            )
+            if not request.include_reasoning:
+                reasoning = None
+                logprobs = None
            return build_response_output_items(
                reasoning=reasoning,
                content=content,
@@ -1342,11 +1345,15 @@ class OpenAIServingResponses(GenerateBaseServing):
    ) -> AsyncGenerator[StreamingResponsesResponse, None]:
        processor = SimpleStreamingEventProcessor(tools=request.tools)

+        hide_stream_metadata = not request.include_reasoning and self.parser is not None
+
        def _get_logprobs(
            output: CompletionOutput,
        ) -> list[response_text_delta_event.Logprob]:
            if not request.is_include_output_logprobs():
                return []
+            if hide_stream_metadata:
+                return []
            return self._create_stream_response_logprobs(
                token_ids=output.token_ids,
                logprobs=output.logprobs,
@@ -19,6 +19,7 @@ class AttentionLayerBase(ABC):
    """

    impl: "AttentionImpl"
+    supports_dcp: bool = True

    @abstractmethod
    def get_attn_backend(self) -> type[AttentionBackend]:
@@ -22,6 +22,7 @@ class MambaBase(AttentionLayerBase):
    # Contains the KV cache (mamba state) for the layer
    # in the shape specified by `self.get_state_shape`.
    kv_cache: tuple[torch.Tensor, ...]
+    supports_dcp: bool = False

    @abstractmethod
    def get_state_shape(self) -> Iterable[tuple[int, ...]]:
@@ -844,17 +844,14 @@ class QwenGatedDeltaNetAttention(GatedDeltaNetAttention):
    def forward(
        self,
        hidden_states: torch.Tensor,
-        output: torch.Tensor,
-    ):
-        self._forward_method(hidden_states, output)
+    ) -> torch.Tensor:
+        return self._forward_method(hidden_states)

    def _output_projection(
        self,
        core_attn_out: torch.Tensor,
        z: torch.Tensor,
-        output: torch.Tensor,
-        num_tokens: int,
-    ):
+    ) -> torch.Tensor:
        """Part 3: RMSNormGated + output linear projection.

        The RMSNormGated + quant sequence is eligible for fusion
@@ -866,13 +863,13 @@ class QwenGatedDeltaNetAttention(GatedDeltaNetAttention):
        core_attn_out = self.norm(core_attn_out, z)
        core_attn_out = core_attn_out.reshape(z_shape_og)
        core_attn_out = core_attn_out.flatten(-2)  # ... h d -> ... (h d)
-        output[:num_tokens], _ = self.out_proj(core_attn_out)
+        output, _ = self.out_proj(core_attn_out)
+        return output

    def forward_hip(
        self,
        hidden_states: torch.Tensor,
-        output: torch.Tensor,
-    ):
+    ) -> torch.Tensor:
        """ROCm forward using AITER Triton fused projection+attention when
        available, otherwise falling back to the generic CUDA path."""
        if GDN_AITER_TRITON_AVAILABLE:
@@ -901,15 +898,14 @@ class QwenGatedDeltaNetAttention(GatedDeltaNetAttention):
                use_aiter=True,
            )

-            self._output_projection(core_attn_out, z, output, num_tokens)
+            return self._output_projection(core_attn_out, z)
        else:
-            self.forward_cuda(hidden_states, output)
+            return self.forward_cuda(hidden_states)

    def forward_cuda(
        self,
        hidden_states: torch.Tensor,
-        output: torch.Tensor,
-    ):
+    ) -> torch.Tensor:
        """
        Forward pass with three parts:
        1. Input projection
@@ -964,13 +960,12 @@ class QwenGatedDeltaNetAttention(GatedDeltaNetAttention):
        # ============================================================
        # Part 3: Output Projection
        # ============================================================
-        self._output_projection(core_attn_out, z, output, num_tokens)
+        return self._output_projection(core_attn_out, z)

    def forward_xpu(
        self,
        hidden_states: torch.Tensor,
-        output: torch.Tensor,
-    ):
+    ) -> torch.Tensor:
        """
        Forward pass with three parts:
        1. Input projection
@@ -1013,13 +1008,13 @@ class QwenGatedDeltaNetAttention(GatedDeltaNetAttention):
        core_attn_out = self.norm(core_attn_out, z)
        core_attn_out = core_attn_out.reshape(z_shape_og)
        core_attn_out = core_attn_out.flatten(-2)  # ... h d -> ... (h d)
-        output[:num_tokens], _ = self.out_proj(core_attn_out)
+        out, _ = self.out_proj(core_attn_out)
+        return out

    def forward_cpu(
        self,
        hidden_states: torch.Tensor,
-        output: torch.Tensor,
-    ):
+    ) -> torch.Tensor:
        assert not hasattr(self, "in_proj_qkv"), "lora isn't supported on CPU."

        mixed_qkvz, _ = self.in_proj_qkvz(hidden_states)
@@ -1063,7 +1058,8 @@ class QwenGatedDeltaNetAttention(GatedDeltaNetAttention):
        core_attn_out = self.norm(core_attn_out, z)
        core_attn_out = core_attn_out.reshape(z_shape_og)
        core_attn_out = core_attn_out.flatten(-2)  # ... h d -> ... (h d)
-        output[:num_tokens], _ = self.out_proj(core_attn_out)
+        out, _ = self.out_proj(core_attn_out)
+        return out

    def _warmup_prefill_kernels(self, qkv_or_qkvz: torch.Tensor, v_dim: int) -> None:
        """Warm up GDN prefill kernels during V1 profiling.
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
mgoinandClaude Fable 5	62a0750892	[Model] Add LongCat-2.0 (sparse attention, MTP, n-gram fixes) Support meituan-longcat/LongCat-2.0-FP8: the LongCat-Flash-Lite backbone plus a DeepSeek-V3.2-style sparse attention indexer, reusing vLLM's existing DSA infrastructure. - Resolve the checkpoint's null model_type via a registered config class (arch LongcatCausalLM); alias the oe_* n-gram config fields. - Wire the indexer into the dual-attention layer: only the first attention of each layer computes top-k indices (cli_factor), the second reuses them via the shared buffer. Adds a generic skip_topk override to DeepseekV2MLAAttention; the schedule lives in the model. - Streaming-aware indexing: force the first index_init_tokens and last index_local_tokens into the top-k set (sparse_attn_indexer). - MTP speculative decoding (one module iterated up to 3 draft steps, plain token embedding, FP8 indexer weights). - Fix a latent LongCat weight-loading bug: the post-load mla_scale_q_lora/kv_lora fold breaks under incremental load_weights calls; fold at weight-load time instead. - Fix two n-gram embedding bugs (also affect LongCat-Flash-Lite): rebuild the left-context from the accepted-token history so rejected draft tokens don't pollute it, and hash EOS-current positions with full look-back per the HF reference (unit test included). - Add --reasoning-parser longcat (<longcat_think> tags) and docs rows. - Add longcat model types to the DeepGEMM Blackwell blocklist (non-ue8m0 scales; measured GSM8K-neutral on vLLM but matches SGLang's guard). Co-authored-by: Claude Fable 5 <noreply@anthropic.com> Signed-off-by: mgoin <mgoin64@gmail.com>	2026-07-11 02:12:49 +00:00
larryli2-amd GitHub Andreas Karatzas	c227aaa3f8	[ROCm] Enable DeepSeek-V4 DSpark speculative decoding on AMD (MI350X / MI355X, gfx950) (#47419 ) Signed-off-by: larryli2-amd <larryli2@amd.com> Signed-off-by: larryli2-amd <Larry.Li@amd.com> Co-authored-by: Andreas Karatzas <akaratza@amd.com>	2026-07-10 23:22:19 +08:00
Michael GoinandGitHub	08dfd68610	[Model] Add LongCat-Flash-Lite (n-gram embedding) (#47857 ) Signed-off-by: mgoin <mgoin64@gmail.com>	2026-07-10 07:17:50 -07:00
Tyler Michael SmithandGitHub	978a6dfa3f	[Build/CI] Build arm64 PR and postmerge image builds for Blackwell SM10x and SM110 (#48041 ) Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>	2026-07-10 10:12:50 -04:00
HuYiPeng GitHub Yipeng Hu Isotr0py	85c09e9885	fix: correct load_weights track logic and enable weight integrity for… (#41811 ) Signed-off-by: Yipeng Hu <i26268@metax-tech.com> Signed-off-by: HuYiPeng <144002351+MynameFelix@users.noreply.github.com> Signed-off-by: Isotr0py <Isotr0py@outlook.com> Co-authored-by: Yipeng Hu <i26268@metax-tech.com> Co-authored-by: Isotr0py <Isotr0py@outlook.com>	2026-07-10 14:08:20 +00:00
XuZhou GitHub Xu Zhou Hoseung Kim	b12cca6a23	[Bugfix] Fix turboquant FP8 cast failure for BF16 models on Ampere GPUs (#39988 ) Signed-off-by: Xu Zhou <xuzhou9417@163.com> Co-authored-by: Xu Zhou <xuzhou9417@163.com> Co-authored-by: Hoseung Kim <ghyutjik123@gmail.com>	2026-07-10 06:55:28 -07:00
Wentao YeandGitHub	e257faf87d	[Refactor] Remove unused rocm kernel `combine_topk_swa_indices_ragged` (#48158 ) Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-07-10 09:33:07 -04:00
FAN YUCHENandGitHub	fabec87f63	[Model] Migrate MistralLarge3ForCausalLM to AutoWeightsLoader (#48153 ) Signed-off-by: Yuchen Fan <functionhx@gmail.com>	2026-07-10 12:27:58 +00:00
Nick Hill GitHub Claude Opus 4.8	7614b88ebd	[Bugfix][Spec Decode] Fix DFlash draft/target layer-count mismatch (#48113 ) Signed-off-by: Nick Hill <nickhill123@gmail.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>	2026-07-10 04:42:53 -07:00
Isotr0pyandGitHub	68ea76e780	[Misc] Remove dead code in ViT functionality test (#48220 ) Signed-off-by: Isotr0py <Isotr0py@outlook.com>	2026-07-10 11:17:42 +00:00
Reid GitHub Bugen Zhao	c241c7a2b0	[Rust Frontend] Add roundtrip fixtures for more chat parsers (#47883 ) Co-authored-by: Bugen Zhao <i@bugenzhao.com> Signed-off-by: reidliu41 <reid201711@gmail.com> Signed-off-by: Bugen Zhao <i@bugenzhao.com>	2026-07-10 10:03:58 +00:00
ViranjanPagar GitHub Claude Opus 4.8 Isotr0py Isotr0py Roger Wang	e23b19309b	Deepstream video backend (#42424 ) Signed-off-by: Viranjan Pagar <vpagar@nvidia.com> Signed-off-by: Isotr0py <Isotr0py@outlook.com> Signed-off-by: Isotr0py <mozf@mail2.sysu.edu.cn> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Co-authored-by: Isotr0py <mozf@mail2.sysu.edu.cn> Co-authored-by: Isotr0py <Isotr0py@outlook.com> Co-authored-by: Roger Wang <hey@rogerw.io>	2026-07-10 02:23:30 -07:00
Andrey Talman GitHub Andrey Talman Kevin H. Luu	f36284a8d2	[CI] Add TORCH_NIGHTLY=1 build mode (run full suite on torch nightly) (#47180 ) Co-authored-by: Andrey Talman <atalman@users.noreply.github.com> Co-authored-by: Kevin H. Luu <khluu000@gmail.com>	2026-07-10 01:38:35 -07:00
Mingfei GuoandGitHub	424df4f65d	[Model][CI/Build] Cosmos3: enable registry tests and register Cosmos3-Super (#48211 ) Signed-off-by: Mingfei Guo <1800012773@pku.edu.cn>	2026-07-10 16:35:15 +08:00
Bugen ZhaoandGitHub	074bdd0d99	[Rust Frontend] Integrate MM video support (#47959 ) Signed-off-by: Bugen Zhao <i@bugenzhao.com>	2026-07-10 08:15:33 +00:00
wenjun liu GitHubjun,du <jun.du@intel.com>Kunshang Ji	216ee58780	Add XPU nightly and release image publishing to DockerHub (#48126 ) Signed-off-by: wenjun.liu <wenjun.liu@intel.com> Signed-off-by: jun,du <jun.du@intel.com> Co-authored-by: jun,du <jun.du@intel.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>	2026-07-10 00:56:00 -07:00
Kevin H. Luu GitHub Claude Opus 4.8	433f291195	[CI] Right-size test-area timeouts from nightly durations (#48186 ) Signed-off-by: khluu <khluu000@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>	2026-07-10 00:53:16 -07:00
Chaojun ZhangandGitHub	28eaf05d56	[XPU] Enable v1/sample tests on XPU CI (#44472 ) Signed-off-by: Chaojun Zhang <chaojun.zhang@intel.com>	2026-07-10 15:40:51 +08:00
Jiangyun ZhuandGitHub	300e33797f	[Perf] fuse more rmsnorm and all-reduce in qwen3.5 (#46998 ) Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>	2026-07-10 15:37:51 +08:00
alberto GitHub Chauncey	5715fde12c	[Feature][Parser] Support include_reasoning param for non-Harmony models (#44301 ) Signed-off-by: Alberto Perdomo <aperdomo@redhat.com> Co-authored-by: Chauncey <chaunceyjiang@gmail.com>	2026-07-10 15:34:02 +08:00
GongLei-HW GitHub Lei Gong Claude	e5588e49bc	[Core][KV events] Report prefix-cache-reused blocks in full report mode (#45261 ) Signed-off-by: Lei Gong <gonglei25@huawei.com> Co-authored-by: Lei Gong <gonglei25@huawei.com> Co-authored-by: Claude <noreply@anthropic.com>	2026-07-09 22:46:54 -07:00
Yan Xu GitHub Jingyi Yang	95ed0feaa5	DCP supports hybrid attention (#40996 ) Signed-off-by: YanXu <yancey.yx@alibaba-inc.com> Signed-off-by: Jingyi Yang <girasoleyang@gmail.com> Co-authored-by: Jingyi Yang <girasoleyang@gmail.com>	2026-07-09 21:34:45 -07:00