Use PDL

Signed-off-by: yewentao256 <zhyanwentao@126.com>
add assert
2026-07-10 15:17:05 +00:00 · 2026-07-10 15:12:30 +00:00 · 2026-07-10 15:10:05 +00:00 · 2026-07-10 14:54:33 +00:00 · 2026-07-10 12:27:58 +00:00 · 2026-07-10 04:42:53 -07:00
83 changed files with 3367 additions and 1350 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}"
@@ -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.")
@@ -6,6 +6,7 @@
 #include "libtorch_stable/quantization/w8a8/per_token_group_quant_8bit.h"

 #include <cmath>
+#include <cstdint>

 #ifdef USE_ROCM
  #include <hip/hip_fp8.h>
@@ -18,6 +19,14 @@
 #include "libtorch_stable/dispatch_utils.h"
 #include "libtorch_stable/torch_utils.h"

+namespace {
+
+inline bool Is16ByteAligned(const void* ptr) {
+  return (reinterpret_cast<std::uintptr_t>(ptr) & 0xFu) == 0;
+}
+
+}  // namespace
+
 __device__ __forceinline__ float GroupReduceMax(float val) {
 #ifdef USE_ROCM
  // 16-thread logical groups may pack up to four per 64-lane wavefront; use a
@@ -95,6 +104,80 @@ __device__ __forceinline__ void QuantizeGroup(
      scalar_op_quant);   // scalar handler
 }

+template <typename T, int VEC_SIZE>
+__device__ __forceinline__ float LoadRegisterGroupAndComputeAbsmax(
+    const T* __restrict__ group_input, const bool is_valid_group,
+    T (&regs)[VEC_SIZE], const float eps) {
+  static_assert(VEC_SIZE * sizeof(T) == 32,
+                "LoadRegisterGroupAndComputeAbsmax assumes two uint4 loads.");
+  float local_absmax = eps;
+  if (!is_valid_group) {
+    return local_absmax;
+  }
+
+  uint4* dst = reinterpret_cast<uint4*>(&regs[0]);
+  const uint4* src = reinterpret_cast<const uint4*>(group_input);
+  dst[0] = src[0];
+  dst[1] = src[1];
+#pragma unroll
+  for (int i = 0; i < VEC_SIZE; ++i) {
+    float v = fabsf(static_cast<float>(regs[i]));
+    local_absmax = fmaxf(local_absmax, v);
+  }
+
+  return local_absmax;
+}
+
+__device__ __forceinline__ float GroupReduceMax8(float val) {
+#ifdef USE_ROCM
+  const int lane_in_wave = threadIdx.x % warpSize;
+  const unsigned long long mask = 0xFFull << (lane_in_wave & ~7);
+  val = fmaxf(val, __shfl_xor_sync(mask, val, 4, 8));
+  val = fmaxf(val, __shfl_xor_sync(mask, val, 2, 8));
+  val = fmaxf(val, __shfl_xor_sync(mask, val, 1, 8));
+#else
+  unsigned mask = 0xffu << (threadIdx.x & 24u);
+  val = fmaxf(val, __shfl_xor_sync(mask, val, 4));
+  val = fmaxf(val, __shfl_xor_sync(mask, val, 2));
+  val = fmaxf(val, __shfl_xor_sync(mask, val, 1));
+#endif
+  return val;
+}
+
+template <typename T, typename DST_DTYPE, int VEC_SIZE, bool SCALE_IS_INVERTED>
+__device__ __forceinline__ uint4
+QuantizeRegisterGroup(const T (&regs)[VEC_SIZE], const float scale,
+                      const float min_8bit, const float max_8bit) {
+  uint32_t packed_lo = 0;
+  uint32_t packed_lo_hi = 0;
+  uint32_t packed_hi_lo = 0;
+  uint32_t packed_hi = 0;
+#pragma unroll
+  for (int i = 0; i < VEC_SIZE; ++i) {
+    float q = static_cast<float>(regs[i]);
+    if constexpr (SCALE_IS_INVERTED) {
+      q *= scale;
+    } else {
+      q /= scale;
+    }
+    q = fminf(fmaxf(q, min_8bit), max_8bit);
+    DST_DTYPE qb = DST_DTYPE(q);
+    uint8_t byte = *reinterpret_cast<uint8_t*>(&qb);
+    const int shift = (i & 3) * 8;
+    if (i < 4) {
+      packed_lo |= static_cast<uint32_t>(byte) << shift;
+    } else if (i < 8) {
+      packed_lo_hi |= static_cast<uint32_t>(byte) << shift;
+    } else if (i < 12) {
+      packed_hi_lo |= static_cast<uint32_t>(byte) << shift;
+    } else {
+      packed_hi |= static_cast<uint32_t>(byte) << shift;
+    }
+  }
+
+  return make_uint4(packed_lo, packed_lo_hi, packed_hi_lo, packed_hi);
+}
+
 template <typename T, typename DST_DTYPE, bool IS_COLUMN_MAJOR = false,
          bool SCALE_UE8M0 = false, typename scale_packed_t = float>
 __global__ void per_token_group_quant_8bit_kernel(
@@ -190,6 +273,96 @@ inline int GetGroupsPerBlockX(int64_t padded_groups_per_row) {
  return 4;
 }

+template <typename T, typename DST_DTYPE, bool IS_COLUMN_MAJOR,
+          bool SCALE_UE8M0, int GROUP_SIZE, int kGroupsPerBlockX,
+          int kRowsPerBlock>
+__global__ void per_token_group_quant_8bit_register_kernel(
+    const T* __restrict__ input, void* __restrict__ output_q,
+    float* __restrict__ output_s, const int groups_per_row, const int mn,
+    const float eps, const float min_8bit, const float max_8bit,
+    const int scale_stride) {
+  static_assert(GROUP_SIZE == 128, "fast path supports GROUP_SIZE==128 only");
+  constexpr int THREADS_PER_GROUP = 8;
+  constexpr int VEC_SIZE = 32 / sizeof(T);  // 16 for bf16/fp16
+  static_assert(GROUP_SIZE == THREADS_PER_GROUP * VEC_SIZE,
+                "GROUP_SIZE must equal THREADS_PER_GROUP * VEC_SIZE");
+  static_assert(32 % THREADS_PER_GROUP == 0,
+                "THREADS_PER_GROUP must divide warp size for the shuffle "
+                "mask to be valid");
+  static_assert(
+      kGroupsPerBlockX > 0 && (kGroupsPerBlockX & (kGroupsPerBlockX - 1)) == 0,
+      "kGroupsPerBlockX must be a positive power of 2");
+  static_assert(kRowsPerBlock > 0, "kRowsPerBlock must be positive");
+
+  const int local_group_id = threadIdx.x / THREADS_PER_GROUP;
+  const int lane_id = threadIdx.x % THREADS_PER_GROUP;
+
+  const int sf_k_local = local_group_id % kGroupsPerBlockX;
+  const int row_local = local_group_id / kGroupsPerBlockX;
+  const int sf_k_idx = blockIdx.y * kGroupsPerBlockX + sf_k_local;
+  const int mn_idx = blockIdx.x * kRowsPerBlock + row_local;
+
+#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
+  cudaGridDependencySynchronize();
+#endif
+
+  if (mn_idx >= mn) {
+#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
+    cudaTriggerProgrammaticLaunchCompletion();
+#endif
+    return;
+  }
+
+  const bool is_valid_group = sf_k_idx < groups_per_row;
+
+  // register-only fast path for group_size == 128.
+  alignas(16) T regs[VEC_SIZE];
+  const T* group_input =
+      input + static_cast<int64_t>(mn_idx) * groups_per_row * GROUP_SIZE +
+      sf_k_idx * GROUP_SIZE + lane_id * VEC_SIZE;
+  // each 8-lane subgroup processes one logical group.
+  float local_absmax = LoadRegisterGroupAndComputeAbsmax<T, VEC_SIZE>(
+      group_input, is_valid_group, regs, eps);
+  local_absmax = GroupReduceMax8(local_absmax);
+
+  float y_s = local_absmax / max_8bit;
+  if constexpr (SCALE_UE8M0) {
+    y_s = exp2f(ceilf(log2f(fmaxf(fabsf(y_s), 1e-10f))));
+  }
+
+  if (lane_id == 0 && is_valid_group) {
+    float* scale_output;
+    if constexpr (IS_COLUMN_MAJOR) {
+      scale_output =
+          output_s + static_cast<int64_t>(sf_k_idx) * scale_stride + mn_idx;
+    } else {
+      scale_output =
+          output_s + static_cast<int64_t>(mn_idx) * groups_per_row + sf_k_idx;
+    }
+    *scale_output = y_s;
+  }
+
+  if (!is_valid_group) {
+#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
+    cudaTriggerProgrammaticLaunchCompletion();
+#endif
+    return;
+  }
+
+  // pack into 16 fp8/int8 bytes (= uint4)
+  uint4 packed_out = QuantizeRegisterGroup<T, DST_DTYPE, VEC_SIZE, false>(
+      regs, y_s, min_8bit, max_8bit);
+  DST_DTYPE* group_output =
+      static_cast<DST_DTYPE*>(output_q) +
+      static_cast<int64_t>(mn_idx) * groups_per_row * GROUP_SIZE +
+      sf_k_idx * GROUP_SIZE + lane_id * VEC_SIZE;
+  *reinterpret_cast<uint4*>(group_output) = packed_out;
+
+#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
+  cudaTriggerProgrammaticLaunchCompletion();
+#endif
+}
+
 void per_token_group_quant_8bit(const torch::stable::Tensor& input,
                                torch::stable::Tensor& output_q,
                                torch::stable::Tensor& output_s,
@@ -218,6 +391,12 @@ void per_token_group_quant_8bit(const torch::stable::Tensor& input,
  const bool is_column_major = output_s.stride(0) < output_s.stride(1);
  const int scale_num_rows = output_s.size(1);
  const int scale_stride = output_s.stride(1);
+  const bool register_fast_path_is_aligned =
+      Is16ByteAligned(input.data_ptr()) && Is16ByteAligned(output_q.data_ptr());
+  const bool use_register_fast_path =
+      register_fast_path_is_aligned && group_size == 128 &&
+      (input.scalar_type() == torch::headeronly::ScalarType::Half ||
+       input.scalar_type() == torch::headeronly::ScalarType::BFloat16);

 #ifndef USE_ROCM
  #define LAUNCH_KERNEL_INST(T, DST_DTYPE, COL_MAJOR, UE8M0, SMEM_BYTES)     \
@@ -271,17 +450,121 @@ void per_token_group_quant_8bit(const torch::stable::Tensor& input,
    }                                                                   \
  } while (0)

-  VLLM_STABLE_DISPATCH_FLOATING_TYPES(
-      input.scalar_type(), "per_token_group_quant_8bit", ([&] {
-        if (dst_type == torch::headeronly::ScalarType::Char) {
-          LAUNCH_KERNEL(scalar_t, int8_t);
-        } else {
-          VLLM_STABLE_DISPATCH_FP8_TYPES(
-              dst_type, "per_token_group_quant_8bit_fp8",
-              ([&] { LAUNCH_KERNEL(scalar_t, fp8_t); }));
-        }
-      }));
+#ifndef USE_ROCM
+  #define LAUNCH_REG_KERNEL_INST(T, DST_DTYPE, IS_COLUMN_MAJOR, SCALE_UE8M0, \
+                                 KX, RY)                                     \
+    do {                                                                     \
+      cudaLaunchConfig_t config = {};                                        \
+      config.gridDim = dim3(static_cast<unsigned int>(blocks_x),             \
+                            static_cast<unsigned int>(blocks_y));            \
+      config.blockDim = dim3(num_threads_fast);                              \
+      config.dynamicSmemBytes = 0;                                           \
+      config.stream = stream;                                                \
+      cudaLaunchAttribute attrs[1];                                          \
+      attrs[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;      \
+      attrs[0].val.programmaticStreamSerializationAllowed = 1;               \
+      config.numAttrs = 1;                                                   \
+      config.attrs = attrs;                                                  \
+      cudaLaunchKernelEx(                                                    \
+          &config,                                                           \
+          per_token_group_quant_8bit_register_kernel<                        \
+              T, DST_DTYPE, IS_COLUMN_MAJOR, SCALE_UE8M0, 128, KX, RY>,      \
+          static_cast<const T*>(input.data_ptr()), output_q.data_ptr(),      \
+          static_cast<float*>(output_s.data_ptr()),                          \
+          static_cast<int>(groups_per_row), static_cast<int>(mn),            \
+          static_cast<float>(eps), static_cast<float>(min_8bit),             \
+          static_cast<float>(max_8bit), scale_stride);                       \
+    } while (0)
+#else
+  #define LAUNCH_REG_KERNEL_INST(T, DST_DTYPE, IS_COLUMN_MAJOR, SCALE_UE8M0, \
+                                 KX, RY)                                     \
+    do {                                                                     \
+      dim3 grid(static_cast<unsigned int>(blocks_x),                         \
+                static_cast<unsigned int>(blocks_y));                        \
+      dim3 block(num_threads_fast);                                          \
+      per_token_group_quant_8bit_register_kernel<                            \
+          T, DST_DTYPE, IS_COLUMN_MAJOR, SCALE_UE8M0, 128, KX, RY>           \
+          <<<grid, block, 0, stream>>>(                                      \
+              static_cast<const T*>(input.data_ptr()), output_q.data_ptr(),  \
+              static_cast<float*>(output_s.data_ptr()),                      \
+              static_cast<int>(groups_per_row), static_cast<int>(mn),        \
+              static_cast<float>(eps), static_cast<float>(min_8bit),         \
+              static_cast<float>(max_8bit), scale_stride);                   \
+    } while (0)
+#endif

+#define LAUNCH_REG_KERNEL(T, DST_DTYPE, IS_COLUMN_MAJOR, SCALE_UE8M0)        \
+  do {                                                                       \
+    if (kx == 16) {                                                          \
+      LAUNCH_REG_KERNEL_INST(T, DST_DTYPE, IS_COLUMN_MAJOR, SCALE_UE8M0, 16, \
+                             1);                                             \
+    } else if (kx == 8) {                                                    \
+      LAUNCH_REG_KERNEL_INST(T, DST_DTYPE, IS_COLUMN_MAJOR, SCALE_UE8M0, 8,  \
+                             2);                                             \
+    } else if (kx == 4) {                                                    \
+      LAUNCH_REG_KERNEL_INST(T, DST_DTYPE, IS_COLUMN_MAJOR, SCALE_UE8M0, 4,  \
+                             4);                                             \
+    } else {                                                                 \
+      STD_TORCH_CHECK(false, "Unsupported kx value ", kx);                   \
+    }                                                                        \
+  } while (0)
+
+#define LAUNCH_REG_KERNEL_FLAGS(T, DST_DTYPE)         \
+  do {                                                \
+    if (is_column_major) {                            \
+      if (scale_ue8m0) {                              \
+        LAUNCH_REG_KERNEL(T, DST_DTYPE, true, true);  \
+      } else {                                        \
+        LAUNCH_REG_KERNEL(T, DST_DTYPE, true, false); \
+      }                                               \
+    } else if (scale_ue8m0) {                         \
+      LAUNCH_REG_KERNEL(T, DST_DTYPE, false, true);   \
+    } else {                                          \
+      LAUNCH_REG_KERNEL(T, DST_DTYPE, false, false);  \
+    }                                                 \
+  } while (0)
+
+  if (use_register_fast_path) {
+    const int64_t k = input.size(-1);
+    const int64_t mn = input.numel() / k;
+    const int64_t groups_per_row = k / group_size;
+    const int64_t padded_groups_per_row = ((groups_per_row + 3) / 4) * 4;
+    const int kx = GetGroupsPerBlockX(padded_groups_per_row);
+    const int ry = 16 / kx;
+    const int64_t blocks_x = (mn + ry - 1) / ry;
+    const int64_t blocks_y = padded_groups_per_row / kx;
+    const int num_threads_fast = (kx * ry) * 8;
+    STD_TORCH_CHECK(
+        blocks_x <= static_cast<int64_t>(INT32_MAX) && blocks_y <= 65535,
+        "per_token_group_quant_8bit fast-path grid too large: (", blocks_x,
+        ", ", blocks_y, ").");
+
+    VLLM_STABLE_DISPATCH_HALF_TYPES(
+        input.scalar_type(), "per_token_group_quant_8bit_register", ([&] {
+          if (dst_type == torch::headeronly::ScalarType::Char) {
+            LAUNCH_REG_KERNEL_FLAGS(scalar_t, int8_t);
+          } else {
+            VLLM_STABLE_DISPATCH_FP8_TYPES(
+                dst_type, "per_token_group_quant_8bit_register_fp8",
+                ([&] { LAUNCH_REG_KERNEL_FLAGS(scalar_t, fp8_t); }));
+          }
+        }));
+  } else {
+    VLLM_STABLE_DISPATCH_FLOATING_TYPES(
+        input.scalar_type(), "per_token_group_quant_8bit", ([&] {
+          if (dst_type == torch::headeronly::ScalarType::Char) {
+            LAUNCH_KERNEL(scalar_t, int8_t);
+          } else {
+            VLLM_STABLE_DISPATCH_FP8_TYPES(
+                dst_type, "per_token_group_quant_8bit_fp8",
+                ([&] { LAUNCH_KERNEL(scalar_t, fp8_t); }));
+          }
+        }));
+  }
+
+#undef LAUNCH_REG_KERNEL
+#undef LAUNCH_REG_KERNEL_INST
+#undef LAUNCH_REG_KERNEL_FLAGS
 #undef LAUNCH_KERNEL
 #undef LAUNCH_KERNEL_INST
 }
@@ -347,36 +630,15 @@ __global__ void per_token_group_quant_8bit_packed_register_kernel(
  // alignas(16) is required so the uint4* reinterpret_cast below is
  // well-defined for T == bf16/fp16 (default alignof is 2).
  alignas(16) T regs[VEC_SIZE];
-  float local_absmax = eps;
-  if (is_valid_group) {
-    const T* group_input =
-        input + static_cast<int64_t>(mn_idx) * groups_per_row * GROUP_SIZE +
-        sf_k_idx * GROUP_SIZE + lane_id * VEC_SIZE;
-    uint4* dst = reinterpret_cast<uint4*>(&regs[0]);
-    const uint4* src = reinterpret_cast<const uint4*>(group_input);
-    dst[0] = src[0];
-    dst[1] = src[1];
-#pragma unroll
-    for (int i = 0; i < VEC_SIZE; ++i) {
-      float v = fabsf(static_cast<float>(regs[i]));
-      local_absmax = fmaxf(local_absmax, v);
-    }
-  }
+  const T* group_input =
+      input + static_cast<int64_t>(mn_idx) * groups_per_row * GROUP_SIZE +
+      sf_k_idx * GROUP_SIZE + lane_id * VEC_SIZE;
+  float local_absmax = LoadRegisterGroupAndComputeAbsmax<T, VEC_SIZE>(
+      group_input, is_valid_group, regs, eps);

  // 8-lane subgroup shuffle reduce (octet of the warp). The mask selects the
  // 8 lanes within the warp that share a group.
-#ifdef USE_ROCM
-  const int lane_in_wave = threadIdx.x % warpSize;
-  const unsigned long long mask = 0xFFull << (lane_in_wave & ~7);
-  local_absmax = fmaxf(local_absmax, __shfl_xor_sync(mask, local_absmax, 4, 8));
-  local_absmax = fmaxf(local_absmax, __shfl_xor_sync(mask, local_absmax, 2, 8));
-  local_absmax = fmaxf(local_absmax, __shfl_xor_sync(mask, local_absmax, 1, 8));
-#else
-  unsigned mask = 0xffu << (threadIdx.x & 24u);
-  local_absmax = fmaxf(local_absmax, __shfl_xor_sync(mask, local_absmax, 4));
-  local_absmax = fmaxf(local_absmax, __shfl_xor_sync(mask, local_absmax, 2));
-  local_absmax = fmaxf(local_absmax, __shfl_xor_sync(mask, local_absmax, 1));
-#endif
+  local_absmax = GroupReduceMax8(local_absmax);

  float y_s = local_absmax / max_8bit;
  y_s = fmaxf(y_s, 1e-10f);
@@ -419,30 +681,8 @@ __global__ void per_token_group_quant_8bit_packed_register_kernel(

  // Quantize and pack into 16 fp8/int8 bytes (= uint4). VEC_SIZE==16 so we
  // fill four 32-bit words, four bytes each.
-  uint32_t packed_lo = 0;
-  uint32_t packed_lo_hi = 0;
-  uint32_t packed_hi_lo = 0;
-  uint32_t packed_hi = 0;
-#pragma unroll
-  for (int i = 0; i < VEC_SIZE; ++i) {
-    float q =
-        fminf(fmaxf(static_cast<float>(regs[i]) * inv_y, min_8bit), max_8bit);
-    DST_DTYPE qb = DST_DTYPE(q);
-    uint8_t byte = *reinterpret_cast<uint8_t*>(&qb);
-    const int shift = (i & 3) * 8;
-    if (i < 4) {
-      packed_lo |= static_cast<uint32_t>(byte) << shift;
-    } else if (i < 8) {
-      packed_lo_hi |= static_cast<uint32_t>(byte) << shift;
-    } else if (i < 12) {
-      packed_hi_lo |= static_cast<uint32_t>(byte) << shift;
-    } else {
-      packed_hi |= static_cast<uint32_t>(byte) << shift;
-    }
-  }
-
-  uint4 packed_out =
-      make_uint4(packed_lo, packed_lo_hi, packed_hi_lo, packed_hi);
+  uint4 packed_out = QuantizeRegisterGroup<T, DST_DTYPE, VEC_SIZE, true>(
+      regs, inv_y, min_8bit, max_8bit);
  DST_DTYPE* group_output =
      static_cast<DST_DTYPE*>(output_q) +
      static_cast<int64_t>(mn_idx) * groups_per_row * GROUP_SIZE +
@@ -473,6 +713,14 @@ void per_token_group_quant_8bit_packed(const torch::stable::Tensor& input,

  STD_TORCH_CHECK(input.is_contiguous());
  STD_TORCH_CHECK(output_q.is_contiguous());
+  STD_TORCH_CHECK(
+      Is16ByteAligned(input.data_ptr()),
+      "per_token_group_quant_8bit_packed requires 16-byte aligned input "
+      "data_ptr for uint4 vectorized loads.");
+  STD_TORCH_CHECK(
+      Is16ByteAligned(output_q.data_ptr()),
+      "per_token_group_quant_8bit_packed requires 16-byte aligned output_q "
+      "data_ptr for uint4 vectorized stores.");

  const int64_t k = input.size(-1);
  STD_TORCH_CHECK(k % group_size == 0, "input last dim k=", k,
@@ -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`).
@@ -512,7 +512,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:
@@ -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"
@@ -11,7 +11,17 @@ from vllm.platforms import current_platform


@pytest.mark.parametrize(
-    "shape", [(31, 128), (32, 128), (63, 256), (64, 256), (16, 512)]
+    "shape",
+    [
+        (31, 128),
+        (32, 128),
+        (63, 256),
+        (64, 256),
+        (16, 512),
+        (31, 384),
+        (33, 640),
+        (65, 768),
+    ],
 )
@pytest.mark.parametrize("column_major", [False, True])
@pytest.mark.parametrize("tma_aligned", [False, True])
@@ -402,7 +412,10 @@ def test_per_token_group_quant_fp8_packed_large_mn():
    assert torch.equal(out_s_packed.cpu(), expected.cpu()), "Packed scale mismatch"


-@pytest.mark.parametrize("shape", [(32, 128), (64, 256), (16, 512)])
+@pytest.mark.parametrize(
+    "shape",
+    [(32, 128), (64, 256), (16, 512), (31, 384), (33, 640), (65, 768)],
+)
@pytest.mark.parametrize("group_size", [64, 128])
@pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA not available")
 def test_per_token_group_quant_int8(shape, group_size: int):
@@ -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']}")
@@ -799,9 +799,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 +1403,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
@@ -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
@@ -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,
@@ -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.
@@ -2,62 +2,91 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 from collections.abc import Iterable

-import regex as re
+import regex
 import torch

 from vllm.model_executor.models.deepseek_v2 import DeepseekV3ForCausalLM
+from vllm.model_executor.models.utils import AutoWeightsLoader, WeightsMapper


 class MistralLarge3ForCausalLM(DeepseekV3ForCausalLM):
-    # fmt: off
-    remapping = {
-        r"layers\.(\d+)\.attention_norm\.weight": r"model.layers.\1.input_layernorm.weight",  # noqa: E501
-        r"layers\.(\d+)\.attention\.wq_a\.(\w+)": r"model.layers.\1.self_attn.q_a_proj.\2",  # noqa: E501
-        r"layers\.(\d+)\.attention\.q_a_norm\.weight": r"model.layers.\1.self_attn.q_a_layernorm.weight",  # noqa: E501
-        r"layers\.(\d+)\.attention\.wq_b\.(\w+)": r"model.layers.\1.self_attn.q_b_proj.\2",  # noqa: E501
-        r"layers\.(\d+)\.attention\.wkv_a_with_mqa\.(\w+)": r"model.layers.\1.self_attn.kv_a_proj_with_mqa.\2",  # noqa: E501
-        r"layers\.(\d+)\.attention\.kv_a_norm\.weight": r"model.layers.\1.self_attn.kv_a_layernorm.weight",  # noqa: E501
-        r"layers\.(\d+)\.attention\.wkv_b\.(\w+)": r"model.layers.\1.self_attn.kv_b_proj.\2",  # noqa: E501
-        r"layers\.(\d+)\.attention\.wo\.(\w+)": r"model.layers.\1.self_attn.o_proj.\2",  # noqa: E501
-        r"layers\.(\d+)\.ffn_norm\.weight": r"model.layers.\1.post_attention_layernorm.weight",  # noqa: E501
-        r"layers\.(\d+)\.feed_forward\.w1\.(\w+)": r"model.layers.\1.mlp.gate_proj.\2",  # noqa: E501
-        r"layers\.(\d+)\.feed_forward\.w2\.(\w+)": r"model.layers.\1.mlp.down_proj.\2",  # noqa: E501
-        r"layers\.(\d+)\.feed_forward\.w3\.(\w+)": r"model.layers.\1.mlp.up_proj.\2",  # noqa: E501
-        r"layers\.(\d+)\.gate\.weight": r"model.layers.\1.mlp.gate.weight",  # noqa: E501
-        r"layers\.(\d+)\.shared_experts\.w1\.(\w+)": r"model.layers.\1.mlp.shared_experts.gate_proj.\2",  # noqa: E501
-        r"layers\.(\d+)\.shared_experts\.w2\.(\w+)": r"model.layers.\1.mlp.shared_experts.down_proj.\2",  # noqa: E501
-        r"layers\.(\d+)\.shared_experts\.w3\.(\w+)": r"model.layers.\1.mlp.shared_experts.up_proj.\2",  # noqa: E501
-        r"layers\.(\d+)\.experts\.(\d+)\.w1\.(\w+)": r"model.layers.\1.mlp.experts.\2.gate_proj.\3",  # noqa: E501
-        r"layers\.(\d+)\.experts\.(\d+)\.w2\.(\w+)": r"model.layers.\1.mlp.experts.\2.down_proj.\3",  # noqa: E501
-        r"layers\.(\d+)\.experts\.(\d+)\.w3\.(\w+)": r"model.layers.\1.mlp.experts.\2.up_proj.\3",  # noqa: E501
-        r"norm\.weight": "model.norm.weight",  # noqa: E501
-        r"tok_embeddings\.weight": "model.embed_tokens.weight",  # noqa: E501
-        r"output\.weight": "lm_head.weight",  # noqa: E501
-    }
-    # fmt: on
+    # WeightsMapper applies all matching patterns sequentially (no break on first
+    # match). This is safe here because every pattern is anchored at both ends
+    # (\A...\Z) and after substitution the resulting key always starts with
+    # "model." or "lm_head.", so no later pattern can accidentally match again.
+    hf_to_vllm_mapper = WeightsMapper(
+        orig_to_new_regex={  # noqa: B950
+            regex.compile(
+                r"\Alayers\.(\d+)\.attention_norm\.weight\Z"
+            ): r"model.layers.\1.input_layernorm.weight",
+            regex.compile(
+                r"\Alayers\.(\d+)\.attention\.wq_a\.(\w+)\Z"
+            ): r"model.layers.\1.self_attn.q_a_proj.\2",
+            regex.compile(
+                r"\Alayers\.(\d+)\.attention\.q_a_norm\.weight\Z"
+            ): r"model.layers.\1.self_attn.q_a_layernorm.weight",
+            regex.compile(
+                r"\Alayers\.(\d+)\.attention\.wq_b\.(\w+)\Z"
+            ): r"model.layers.\1.self_attn.q_b_proj.\2",
+            regex.compile(
+                r"\Alayers\.(\d+)\.attention\.wkv_a_with_mqa\.(\w+)\Z"
+            ): r"model.layers.\1.self_attn.kv_a_proj_with_mqa.\2",
+            regex.compile(
+                r"\Alayers\.(\d+)\.attention\.kv_a_norm\.weight\Z"
+            ): r"model.layers.\1.self_attn.kv_a_layernorm.weight",
+            regex.compile(
+                r"\Alayers\.(\d+)\.attention\.wkv_b\.(\w+)\Z"
+            ): r"model.layers.\1.self_attn.kv_b_proj.\2",
+            regex.compile(
+                r"\Alayers\.(\d+)\.attention\.wo\.(\w+)\Z"
+            ): r"model.layers.\1.self_attn.o_proj.\2",
+            regex.compile(
+                r"\Alayers\.(\d+)\.ffn_norm\.weight\Z"
+            ): r"model.layers.\1.post_attention_layernorm.weight",
+            regex.compile(
+                r"\Alayers\.(\d+)\.feed_forward\.w1\.(\w+)\Z"
+            ): r"model.layers.\1.mlp.gate_proj.\2",
+            regex.compile(
+                r"\Alayers\.(\d+)\.feed_forward\.w2\.(\w+)\Z"
+            ): r"model.layers.\1.mlp.down_proj.\2",
+            regex.compile(
+                r"\Alayers\.(\d+)\.feed_forward\.w3\.(\w+)\Z"
+            ): r"model.layers.\1.mlp.up_proj.\2",
+            regex.compile(
+                r"\Alayers\.(\d+)\.gate\.weight\Z"
+            ): r"model.layers.\1.mlp.gate.weight",
+            regex.compile(
+                r"\Alayers\.(\d+)\.shared_experts\.w1\.(\w+)\Z"
+            ): r"model.layers.\1.mlp.shared_experts.gate_proj.\2",
+            regex.compile(
+                r"\Alayers\.(\d+)\.shared_experts\.w2\.(\w+)\Z"
+            ): r"model.layers.\1.mlp.shared_experts.down_proj.\2",
+            regex.compile(
+                r"\Alayers\.(\d+)\.shared_experts\.w3\.(\w+)\Z"
+            ): r"model.layers.\1.mlp.shared_experts.up_proj.\2",
+            regex.compile(
+                r"\Alayers\.(\d+)\.experts\.(\d+)\.w1\.(\w+)\Z"
+            ): r"model.layers.\1.mlp.experts.\2.gate_proj.\3",
+            regex.compile(
+                r"\Alayers\.(\d+)\.experts\.(\d+)\.w2\.(\w+)\Z"
+            ): r"model.layers.\1.mlp.experts.\2.down_proj.\3",
+            regex.compile(
+                r"\Alayers\.(\d+)\.experts\.(\d+)\.w3\.(\w+)\Z"
+            ): r"model.layers.\1.mlp.experts.\2.up_proj.\3",
+            regex.compile(r"\Anorm\.weight\Z"): "model.norm.weight",
+            regex.compile(r"\Atok_embeddings\.weight\Z"): "model.embed_tokens.weight",
+            regex.compile(r"\Aoutput\.weight\Z"): "lm_head.weight",
+        },
+        orig_to_new_suffix={
+            ".qscale_act": ".input_scale",
+            ".qscale_weight": ".weight_scale",
+        },
+    )

+    # Bypass super().load_weights() and construct AutoWeightsLoader(self)
+    # directly (same pattern as Qwen2ForCausalLM). Any logic in the parent
+    # class's load_weights is a thin wrapper around AutoWeightsLoader, and
+    # we must apply hf_to_vllm_mapper before the loader walks the tree.
    def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]:
-        return super().load_weights(map(self._remap_mistral_to_ds, weights))
-
-    def _remap_mistral_to_ds(
-        self, weight: tuple[str, torch.Tensor]
-    ) -> tuple[str, torch.Tensor]:
-        """Remap Mistral parameters to DeepseekV2 parameters."""
-        name, loaded_weight = weight
-
-        for k, v in self.remapping.items():
-            match = re.fullmatch(k, name)
-            if match:
-                name = re.sub(k, v, name)
-                break
-        else:
-            raise ValueError(f"Cannot remap {name}")
-
-        # Remapping scale names. We could do this in the regex above but it
-        # would triple the number of lines for most layers.
-        if name.endswith(".qscale_act"):
-            name = re.sub(r"\.qscale_act$", ".input_scale", name)
-        elif name.endswith(".qscale_weight"):
-            name = re.sub(r"\.qscale_weight$", ".weight_scale", name)
-
-        return name, loaded_weight
+        loader = AutoWeightsLoader(self)
+        return loader.load_weights(weights, mapper=self.hf_to_vllm_mapper)
@@ -5,6 +5,7 @@ import copy
 from collections.abc import Iterable
 from functools import partial

+import regex
 import torch
 import torch.nn as nn

@@ -22,7 +23,7 @@ from vllm.model_executor.models.deepseek_v2 import (
 from vllm.model_executor.models.mistral_large_3 import MistralLarge3ForCausalLM

 from .interfaces import SupportsMultiModal
-from .utils import make_empty_intermediate_tensors_factory, maybe_prefix
+from .utils import WeightsMapper, make_empty_intermediate_tensors_factory, maybe_prefix

 logger = init_logger(__name__)

@@ -107,11 +108,13 @@ class EagleMistralLarge3Model(DeepseekV2Model):


 class EagleMistralLarge3ForCausalLM(MistralLarge3ForCausalLM):
-    remapping = MistralLarge3ForCausalLM.remapping | {
-        r"eagle_linear\.weight": r"model.fc.weight",
-        r"eagle_linear\.qscale_act": r"model.fc.input_scale",
-        r"eagle_linear\.qscale_weight": r"model.fc.weight_scale",
-    }
+    hf_to_vllm_mapper = MistralLarge3ForCausalLM.hf_to_vllm_mapper | WeightsMapper(
+        orig_to_new_regex={
+            regex.compile(r"\Aeagle_linear\.weight\Z"): r"model.fc.weight",
+            regex.compile(r"\Aeagle_linear\.qscale_act\Z"): r"model.fc.input_scale",
+            regex.compile(r"\Aeagle_linear\.qscale_weight\Z"): r"model.fc.weight_scale",
+        },
+    )

    def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""):
        target_layer_num = vllm_config.model_config.get_num_layers(
@@ -754,6 +754,14 @@ class DFlashQwen3ForCausalLM(Qwen3ForCausalLM):
        needs_squeeze = hidden_states.dim() == 1
        if needs_squeeze:
            hidden_states = hidden_states.unsqueeze(0)
+        expected = self.model.fc.input_size
+        if hidden_states.shape[-1] != expected:
+            raise ValueError(
+                f"DFlash drafter expects {expected} concatenated aux hidden "
+                f"features but received {hidden_states.shape[-1]}. This usually "
+                "means the draft model's target_layer_ids reference layers that "
+                "do not exist in the target model (incompatible draft/target pair)."
+            )
        result = self.model.fc(hidden_states)
        if needs_squeeze:
            result = result.squeeze(0)
@@ -381,15 +381,15 @@ class Qwen3NextAttention(nn.Module):
    def forward(
        self,
        positions: torch.Tensor,
-        output: torch.Tensor,
        hidden_states: torch.Tensor,
-    ):
+    ) -> torch.Tensor:
        qkv, _ = self.qkv_proj(hidden_states)
        q, k, v, gate = self._project_qkv_gate(qkv, positions)
        attn_output = self.attn(q, k, v)
        if gate is not None:
            attn_output = attn_output * torch.sigmoid(gate)
-        output[:], _ = self.o_proj(attn_output)
+        output, _ = self.o_proj(attn_output)
+        return output


 class Qwen3NextDecoderLayer(nn.Module):
@@ -484,21 +484,15 @@ class Qwen3NextDecoderLayer(nn.Module):
        else:
            hidden_states, residual = self.input_layernorm(hidden_states, residual)

-        self_attention_output = torch.empty_like(hidden_states)
        if self.layer_type == "linear_attention":
-            self.linear_attn(
-                hidden_states=hidden_states,
-                output=self_attention_output,
-            )
+            hidden_states = self.linear_attn(hidden_states=hidden_states)
        elif self.layer_type == "full_attention":
-            self.self_attn(
+            hidden_states = self.self_attn(
                hidden_states=hidden_states,
-                output=self_attention_output,
                positions=positions,
            )
        else:
            raise ValueError("Invalid layer_type")
-        hidden_states = self_attention_output

        if self.layer_scale:
            if len(hidden_states.shape) == 2:
@@ -826,6 +826,114 @@ class PyNvVideoCodecVideoBackendMixin:
        return frames, source, frame_idx, valid_frame_indices


+class DeepStreamVideoBackendMixin:
+    """NVIDIA DeepStream (NVDEC) GPU-decode codec utilities.
+
+    Decoding runs on a shared pool of daemon threads inside one CUDA
+    context (see the ``nvidia-deepstream-videodecode-cu13`` package). The
+    container bytes are pushed into an ``appsrc`` GStreamer pipeline, so no
+    local file path is required — HTTP and base64 sources decode identically
+    to local files.
+
+    Like the OpenCV/PyAV mixins, this provides only the codec layer.
+    Frame *selection* lives in the loader's
+    ``compute_frames_index_to_sample`` and arrives here as an explicit
+    list of frame indices.
+    """
+
+    # Process-wide lazy decode pool, shared across all DeepStream backends.
+    _pool: ClassVar[Any] = None
+    _pool_lock: ClassVar[Any] = None
+
+    @classmethod
+    def _get_pool(cls, pool_size: int | None = None):
+        """Lazy-initialize the shared decode pool on first use.
+
+        ``pool_size`` (number of decode worker threads) comes from
+        ``--media-io-kwargs`` (``{"video": {"pool_size": N}}``); when unset it
+        defaults to the existing ``VLLM_MEDIA_LOADING_THREAD_COUNT`` so no
+        DeepStream-specific env var is needed. The pool is a process-wide
+        singleton, so the first decode's value wins.
+        """
+        if cls._pool is not None:
+            return cls._pool
+        if cls._pool_lock is None:
+            cls._pool_lock = threading.Lock()
+        with cls._pool_lock:
+            if cls._pool is not None:
+                return cls._pool
+            import os
+
+            from nvidia.deepstream_videodecode import DecodePool
+
+            if pool_size is None:
+                pool_size = int(os.environ.get("VLLM_MEDIA_LOADING_THREAD_COUNT", 8))
+            pool_size = max(1, min(int(pool_size), 16))
+            logger.info(
+                "[DeepStream] initializing decode pool with %d workers",
+                pool_size,
+            )
+            cls._pool = DecodePool(num_workers=pool_size)
+            return cls._pool
+
+    @classmethod
+    def decode_indices(
+        cls,
+        data: bytes,
+        frame_indices: list[int],
+        source: VideoSourceMetadata,
+        codec: str = "",
+        pool_size: int | None = None,
+        timeout_sec: float = 120.0,
+    ) -> tuple[npt.NDArray, list[int]]:
+        """Decode the requested frame indices from raw container bytes.
+
+        The whole stream is decoded; the pool keeps exactly the frames whose
+        decode-order index is in ``frame_indices`` (1:1, frame-exact) and
+        sends EOS once the last one is matched.
+
+        ``codec`` (e.g. ``"h264"``/``"hevc"``) lets the pool keep its NVDEC
+        session warm across same-codec streams and rebuild only on a codec
+        change. Frames are returned as a CPU NHWC uint8 array so the
+        upstream multimodal parser sees the same shape as the other
+        backends.
+        """
+        if not frame_indices:
+            raise ValueError("DeepStream backend received no frame indices")
+
+        result = cls._get_pool(pool_size).decode(
+            data,
+            target_indices=frame_indices,
+            codec=codec,
+            max_frames=len(frame_indices),
+            timeout_sec=timeout_sec,
+        )
+        if result.error:
+            raise ValueError(f"DeepStream decode failed: {result.error}")
+        if result.frames is None or result.n_kept == 0:
+            raise ValueError("DeepStream decode produced no frames")
+
+        valid = frame_indices[: result.n_kept]
+        # GPU -> CPU NHWC uint8 at the codec boundary (one PCIe copy); keeps
+        # the array shape identical to the OpenCV/PyAV backends. Copy into
+        # PINNED host memory (reused across calls by PyTorch's pinned caching
+        # allocator) so the D2H runs at full PCIe bandwidth (~13 GB/s) rather
+        # than the ~1 GB/s pageable path that plain ``.cpu()`` takes — ~12x
+        # faster for a 1080p x8 frame batch (~46ms -> ~4ms). ``numpy()`` keeps
+        # the pinned tensor alive via the array's base.
+        import torch
+
+        gpu = result.frames
+        if gpu.is_cuda:
+            host = torch.empty(gpu.shape, dtype=gpu.dtype, pin_memory=True)
+            host.copy_(gpu, non_blocking=True)
+            torch.cuda.current_stream().synchronize()
+            arr = host.numpy()
+        else:
+            arr = gpu.numpy()
+        return arr, valid
+
+
@VIDEO_LOADER_REGISTRY.register("opencv")
 class VideoBackend(
    VideoLoader,
@@ -833,14 +941,15 @@ class VideoBackend(
    PyAVVideoBackendMixin,
    TorchCodecVideoBackendMixin,
    PyNvVideoCodecVideoBackendMixin,
+    DeepStreamVideoBackendMixin,
 ):
    """Uniform-sampling video backend.

    Samples ``num_frames`` uniformly across the video (or one frame every
    ``1/fps`` seconds, whichever produces fewer frames). The decoding codec
    is selected via the ``backend`` kwarg (``"opencv"``, ``"pyav"``,
-    ``"torchcodec"`` or ``"pynvvideocodec"``), which can be passed through
-    ``--media-io-kwargs``. Defaults to ``"opencv"``.
+    ``"torchcodec"``, ``"pynvvideocodec"``, or ``"deepstream"``),
+    which can be passed through ``--media-io-kwargs``. Defaults to ``"opencv"``.
    """

    _sampling_suffix: ClassVar[str] = ""
@@ -885,7 +994,9 @@ class VideoBackend(
        max_duration: int = 300,
        frame_recovery: bool = False,
        *,
-        backend: Literal["opencv", "pyav", "torchcodec", "pynvvideocodec"] = "opencv",
+        backend: Literal[
+            "opencv", "pyav", "torchcodec", "pynvvideocodec", "deepstream"
+        ] = "opencv",
        num_ffmpeg_threads: int = 0,
        seek_mode: Literal["exact", "approximate"] = "exact",
        **kwargs,
@@ -901,7 +1012,7 @@ class VideoBackend(
            frame_recovery: Enable forward-scan recovery for failed frames.
                Only honored by the OpenCV codec.
            backend: Decoding codec — ``"opencv"``, ``"pyav"``,
-                ``"torchcodec"`` or ``"pynvvideocodec"``.
+                ``"torchcodec"``, ``"pynvvideocodec"`` or ``"deepstream"``.
            num_ffmpeg_threads: Number of FFmpeg decoding threads, only used by
                TorchCodec: ``0`` (default) relies on the FFmpeg default value
                which is ``min(cpu_count + 1, 16)``.
@@ -982,11 +1093,37 @@ class VideoBackend(
                target,
                **kwargs,
            )
+        elif backend == "deepstream":
+            assert not frame_recovery, (
+                "frame_recovery is only available for `opencv` backend"
+            )
+            # Decode-pool size comes from media-io-kwargs (no env var); the
+            # pool is a process-wide singleton so the first decode's value
+            # wins. Pop it so it isn't forwarded to the frame sampler.
+            pool_size = kwargs.pop("pool_size", None)
+            # Probe container metadata from the bytes via GStreamer (in
+            # the deepstream video-decode wheel) — no PyAV/pymediainfo, no path.
+            from nvidia.deepstream_videodecode import probe_metadata
+
+            total_frames, original_fps, duration, _w, _h, codec = probe_metadata(data)
+            source = cls._prepare_source(
+                VideoSourceMetadata(
+                    total_frames_num=total_frames,
+                    original_fps=original_fps,
+                    duration=duration,
+                )
+            )
+            frame_idx = cls.compute_frames_index_to_sample(
+                source=source, target=target, **kwargs
+            )
+            frames, valid = cls.decode_indices(
+                data, frame_idx, source, codec=codec, pool_size=pool_size
+            )
        else:
            raise ValueError(
                f"Unknown video codec backend {backend!r}; "
                "valid options: 'opencv', 'pyav', 'torchcodec', "
-                "'pynvvideocodec'."
+                "'pynvvideocodec' and 'deepstream'."
            )

        if len(valid) < len(frame_idx):
@@ -1073,7 +1210,9 @@ class Qwen3VLVideoBackend(VideoBackend):
        max_duration: int = 300,
        frame_recovery: bool = False,
        *,
-        backend: Literal["opencv", "pyav", "torchcodec", "pynvvideocodec"] = "opencv",
+        backend: Literal[
+            "opencv", "pyav", "torchcodec", "pynvvideocodec", "deepstream"
+        ] = "opencv",
        **kwargs,
    ) -> tuple[npt.NDArray, dict[str, Any]]:
        return super().load_bytes(
@@ -1152,7 +1291,9 @@ class Qwen2VLVideoBackend(VideoBackend):
        max_duration: int = 300,
        frame_recovery: bool = False,
        *,
-        backend: Literal["opencv", "pyav", "torchcodec", "pynvvideocodec"] = "opencv",
+        backend: Literal[
+            "opencv", "pyav", "torchcodec", "pynvvideocodec", "deepstream"
+        ] = "opencv",
        **kwargs,
    ) -> tuple[npt.NDArray, dict[str, Any]]:
        return super().load_bytes(
@@ -1244,7 +1385,9 @@ class DynamicVideoBackend(VideoBackend):
        max_duration: int = 300,
        frame_recovery: bool = False,
        *,
-        backend: Literal["opencv", "pyav", "torchcodec", "pynvvideocodec"] = "opencv",
+        backend: Literal[
+            "opencv", "pyav", "torchcodec", "pynvvideocodec", "deepstream"
+        ] = "opencv",
        **kwargs,
    ) -> tuple[npt.NDArray, dict[str, Any]]:
        return super().load_bytes(
@@ -1369,7 +1512,9 @@ class GLM46VVideoBackend(VideoBackend):
        max_duration: int = 300,
        frame_recovery: bool = False,
        *,
-        backend: Literal["opencv", "pyav", "torchcodec", "pynvvideocodec"] = "opencv",
+        backend: Literal[
+            "opencv", "pyav", "torchcodec", "pynvvideocodec", "deepstream"
+        ] = "opencv",
        **kwargs,
    ) -> tuple[npt.NDArray, dict[str, Any]]:
        return super().load_bytes(
@@ -1467,7 +1612,9 @@ class GLMGAVideoBackend(VideoBackend):
        max_duration: int = 300,
        frame_recovery: bool = False,
        *,
-        backend: Literal["opencv", "pyav", "torchcodec", "pynvvideocodec"] = "opencv",
+        backend: Literal[
+            "opencv", "pyav", "torchcodec", "pynvvideocodec", "deepstream"
+        ] = "opencv",
        **kwargs,
    ) -> tuple[npt.NDArray, dict[str, Any]]:
        frames, metadata = super().load_bytes(
@@ -1790,7 +1937,9 @@ class NemotronVLVideoBackend(VideoBackend):
        max_duration: int = 300,
        frame_recovery: bool = False,
        *,
-        backend: Literal["opencv", "pyav", "torchcodec", "pynvvideocodec"] = "opencv",
+        backend: Literal[
+            "opencv", "pyav", "torchcodec", "pynvvideocodec", "deepstream"
+        ] = "opencv",
        **kwargs,
    ) -> tuple[npt.NDArray, dict[str, Any]]:
        frames, metadata = super().load_bytes(
@@ -447,7 +447,7 @@ class DelegatingParser(Parser):

        tool_calls = list[FunctionCall]()
        if is_named_tool_choice and supports_required_and_named:
-            if content is None:
+            if content is None or (isinstance(content, str) and not content.strip()):
                return [], None
            function_name = self._get_function_name(request)
            tool_calls.append(
@@ -496,6 +496,14 @@ class DelegatingParser(Parser):
                    content = None
            else:
                # No tool calls.
+                # For required/named tool choice (when falling back to auto
+                # parsing), if content is empty or whitespace-only, return
+                # empty list with None content.
+                if (is_required_tool_choice or is_named_tool_choice) and (
+                    content is None
+                    or (isinstance(content, str) and not content.strip())
+                ):
+                    return [], None
                return None, content

        return tool_calls, content
@@ -917,6 +925,15 @@ class DelegatingParser(Parser):
            delta_message = self.finalize_generation(delta_message, request, state)
            delta_message = self._flush_engine_parsers(delta_message)

+        # Suppress reasoning deltas if not requested
+        if delta_message and not request.include_reasoning:
+            delta_message.reasoning = None
+
+            # If only reasoning was in the message (no content, no tool_calls)
+            # skip emitting entirely
+            if not delta_message.content and not delta_message.tool_calls:
+                delta_message = None
+
        return delta_message

    def _flush_engine_parsers(
@@ -443,7 +443,15 @@ class ParserEngine(Parser):
        if finished:
            events.extend(self._engine.finish())
        result = self._events_to_delta(events, finished=finished)
-        return self._strip_trailing_reasoning(result)
+        result = self._strip_trailing_reasoning(result)
+
+        # Suppress reasoning deltas if not requested
+        if result and not request.include_reasoning:
+            result.reasoning = None
+            if not result.content and not result.tool_calls:
+                result = None
+
+        return result

    def _strip_trailing_reasoning(
        self,
@@ -293,6 +293,16 @@ class HarmonyParser(DelegatingParser):
            delta_message.reasoning = combined_reasoning
        if tool_messages:
            delta_message.tool_calls = tool_messages
+
+        # Suppress reasoning deltas if not requested
+        if delta_message and not request.include_reasoning:
+            delta_message.reasoning = None
+
+            # If only reasoning was in the message (no content, no tool_calls)
+            # skip emitting entirely
+            if not delta_message.content and not delta_message.tool_calls:
+                return None
+
        return delta_message

    def process_chunk(self, token_ids: Sequence[int]) -> ChunkResult:
Author	SHA1	Message	Date
yewentao256	3ecf840f0f	Use PDL Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-07-10 15:17:05 +00:00
yewentao256	00aa9038db	add assert Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-07-10 15:12:30 +00:00
yewentao256	6d9242a17c	make sure aligned Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-07-10 15:10:05 +00:00
yewentao256	d2688ce612	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-07-10 14:54:33 +00: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
Wentao YeandGitHub	5dfbdc4d43	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-07-08 11:17:26 -04:00
Wentao YeandGitHub	3cdb41453a	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-07-04 10:21:50 -04:00
Wentao YeandGitHub	5723461449	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-07-02 16:34:46 -04:00
yewentao256	87e62fc11f	fix cuda error Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-07-02 20:26:10 +00:00
Wentao YeandGitHub	7ec53deb8b	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-07-02 10:44:36 -04:00
Wentao YeandGitHub	ec5cec657e	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-06-30 14:33:40 -04:00
Wentao YeandGitHub	b41f7d8ffd	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-06-29 17:37:43 -04:00
yewentao256	69655da6d9	fix bug Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-06-25 20:23:15 +00:00
yewentao256	61eabe9f5e	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-06-25 19:54:03 +00:00
mergify[bot]andGitHub	f3e37e3c50	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-06-24 14:27:03 +00:00
mergify[bot]andGitHub	9307364c39	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-06-23 15:15:06 +00:00
mergify[bot]andGitHub	c0de337f95	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-06-22 15:42:46 +00:00
yewentao256	4bbe1b7a21	solve conflicts Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-06-20 15:16:38 +00:00
Wentao YeandGitHub	b6a4252d66	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-06-17 13:56:11 -04:00
mergify[bot]andGitHub	349af588f0	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-06-12 16:00:35 +00:00
Wentao YeandGitHub	c4bf7dadb9	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-06-11 11:09:10 -04:00
Wentao YeandGitHub	45dfecbe55	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-06-04 15:09:36 -04:00
yewentao256	f185d7d7ac	fix conflicts Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-06-03 17:24:48 +00:00
Wentao YeandGitHub	887eaedb3d	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-06-01 16:20:04 -04:00
yewentao256	f6fb856228	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-05-29 20:53:15 +00:00
Wentao YeandGitHub	a55cb117a7	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-05-29 11:12:12 -04:00
yewentao256	74f4d24f4f	refactor Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-05-28 22:48:37 +00:00
yewentao256	9475c0bba1	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-05-28 19:37:46 +00:00
Wentao YeandGitHub	cfbc2282a7	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-05-26 14:47:47 -04:00
Wentao YeandGitHub	673c7901fc	Merge branch 'main' into wentao-optimize-per-token-group-quant	2026-05-21 14:18:15 -04:00
yewentao256	e20a7be4d0	optimize per token group quant Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-05-19 20:13:19 +00:00