[Core] Add explicit layer parallel plans

Co-authored-by: OpenAI Codex <codex@openai.com>
Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>

.buildkite/hardware_tests/cpu.yaml

@@ -141,9 +141,22 @@ steps:
   commands:
     - |
       bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 45m "
-      pytest -x -v -s tests/models/multimodal/generation --ignore=tests/models/multimodal/generation/test_pixtral.py -m cpu_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB"
+      pytest -x -v -s tests/models/multimodal/generation --ignore=tests/models/multimodal/generation/test_pixtral.py --ignore=tests/models/multimodal/generation/test_qwen2_5_vl.py -m cpu_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB"
   parallelism: 4
 
+- label: CPU-Qwen2.5-VL Multimodal Tests
+  depends_on: []
+  device: intel_cpu
+  no_plugin: true
+  source_file_dependencies:
+  # - vllm/
+  - vllm/model_executor/layers/rotary_embedding
+  - tests/models/multimodal/generation/
+  commands:
+    - |
+      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 40m "
+      VLLM_CI_ENV=0 pytest -x -v -s tests/models/multimodal/generation/test_qwen2_5_vl.py"
+
 - label: "Arm CPU Test"
   depends_on: []
   soft_fail: false

.buildkite/image_build/image_build.sh

@@ -79,12 +79,18 @@ setup_buildx_builder() {
     docker buildx ls | grep -E '^\*|^NAME' || docker buildx ls
 }
 
+annotate_image_tags() {
+    .buildkite/scripts/annotate-image-build.sh \
+        "${IMAGE_TAG:-}" "${IMAGE_TAG_LATEST:-}"
+}
+
 check_and_skip_if_image_exists() {
     if [[ -n "${IMAGE_TAG:-}" ]]; then
         echo "--- :mag: Checking if image exists"
         if docker manifest inspect "${IMAGE_TAG}" >/dev/null 2>&1; then
             echo "Image already exists: ${IMAGE_TAG}"
             echo "Skipping build"
+            annotate_image_tags
             exit 0
         fi
         echo "Image not found, proceeding with build"
@@ -171,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}"
@@ -254,3 +272,5 @@ echo "--- :docker: Building ${TARGET}"
 docker --debug buildx bake -f "${VLLM_BAKE_FILE_PATH}" -f "${CI_HCL_PATH}" --progress plain "${TARGET}"
 
 echo "--- :white_check_mark: Build complete"
+
+annotate_image_tags

.buildkite/image_build/image_build_arm64.sh

@@ -9,30 +9,31 @@ fi
 REGISTRY=$1
 REPO=$2
 BUILDKITE_COMMIT=$3
+IMAGE="$REGISTRY/$REPO:$BUILDKITE_COMMIT-arm64"
 
 # authenticate with AWS ECR
 aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
 
 # skip build if image already exists
-if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64) ]]; then
-  echo "Image not found, proceeding with build..."
-else
+if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
   echo "Image found"
-  exit 0
+else
+  echo "Image not found, proceeding with build..."
+  # build for arm64 GPU targets: Grace/GH200 (sm_90),
+  # Blackwell/Thor (sm_100/sm_103/sm_110), and DGX Spark/GB10
+  # (sm_121, family-covered by 12.0 under CUDA 13)
+  docker build --file docker/Dockerfile \
+    --platform linux/arm64 \
+    --build-arg max_jobs=16 \
+    --build-arg nvcc_threads=4 \
+    --build-arg torch_cuda_arch_list="9.0 10.0 11.0 12.0" \
+    --build-arg USE_SCCACHE=1 \
+    --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
+    --tag "$IMAGE" \
+    --target test \
+    --progress plain .
+  # push
+  docker push "$IMAGE"
 fi
 
-# build for arm64 GPU targets: Grace/GH200 (sm_90) and DGX Spark/GB10
-# (sm_121, family-covered by 12.0 under CUDA 13)
-docker build --file docker/Dockerfile \
-  --platform linux/arm64 \
-  --build-arg max_jobs=16 \
-  --build-arg nvcc_threads=4 \
-  --build-arg torch_cuda_arch_list="9.0 12.0" \
-  --build-arg USE_SCCACHE=1 \
-  --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
-  --tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64 \
-  --target test \
-  --progress plain .
-
-# push
-docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64
+.buildkite/scripts/annotate-image-build.sh "$IMAGE"

.buildkite/image_build/image_build_cpu.sh

@@ -9,26 +9,26 @@ fi
 REGISTRY=$1
 REPO=$2
 BUILDKITE_COMMIT=$3
+IMAGE="$REGISTRY/$REPO:$BUILDKITE_COMMIT-cpu"
 
 # authenticate with AWS ECR
 aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
 
 # skip build if image already exists
-if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-cpu) ]]; then
-  echo "Image not found, proceeding with build..."
-else
+if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
   echo "Image found"
-  exit 0
+else
+  echo "Image not found, proceeding with build..."
+  # build
+  docker build --file docker/Dockerfile.cpu \
+    --build-arg max_jobs=16 \
+    --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
+    --build-arg VLLM_CPU_X86=true \
+    --tag "$IMAGE" \
+    --target vllm-test \
+    --progress plain .
+  # push
+  docker push "$IMAGE"
 fi
 
-# build
-docker build --file docker/Dockerfile.cpu \
-  --build-arg max_jobs=16 \
-  --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
-  --build-arg VLLM_CPU_X86=true \
-  --tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-cpu \
-  --target vllm-test \
-  --progress plain .
-
-# push
-docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-cpu
+.buildkite/scripts/annotate-image-build.sh "$IMAGE"

.buildkite/image_build/image_build_cpu_arm64.sh

@@ -9,25 +9,25 @@ fi
 REGISTRY=$1
 REPO=$2
 BUILDKITE_COMMIT=$3
+IMAGE="$REGISTRY/$REPO:$BUILDKITE_COMMIT-arm64-cpu"
 
 # authenticate with AWS ECR
 aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
 
 # skip build if image already exists
-if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64-cpu) ]]; then
-  echo "Image not found, proceeding with build..."
-else
+if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
   echo "Image found"
-  exit 0
+else
+  echo "Image not found, proceeding with build..."
+  # build
+  docker build --file docker/Dockerfile.cpu \
+    --build-arg max_jobs=16 \
+    --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
+    --tag "$IMAGE" \
+    --target vllm-test \
+    --progress plain .
+  # push
+  docker push "$IMAGE"
 fi
 
-# build
-docker build --file docker/Dockerfile.cpu \
-  --build-arg max_jobs=16 \
-  --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
-  --tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64-cpu \
-  --target vllm-test \
-  --progress plain .
-
-# push
-docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64-cpu
+.buildkite/scripts/annotate-image-build.sh "$IMAGE"

.buildkite/image_build/image_build_hpu.sh

@@ -9,26 +9,26 @@ fi
 REGISTRY=$1
 REPO=$2
 BUILDKITE_COMMIT=$3
+IMAGE="$REGISTRY/$REPO:$BUILDKITE_COMMIT-hpu"
 
 # authenticate with AWS ECR
 aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
 
 # skip build if image already exists
-if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-hpu) ]]; then
-  echo "Image not found, proceeding with build..."
-else
+if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
   echo "Image found"
-  exit 0
+else
+  echo "Image not found, proceeding with build..."
+  # build
+  docker build \
+    --file tests/pytorch_ci_hud_benchmark/Dockerfile.hpu \
+    --build-arg max_jobs=16 \
+    --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
+    --tag "$IMAGE" \
+    --progress plain \
+    https://github.com/vllm-project/vllm-gaudi.git
+  # push
+  docker push "$IMAGE"
 fi
 
-# build
-docker build \
-  --file tests/pytorch_ci_hud_benchmark/Dockerfile.hpu \
-  --build-arg max_jobs=16 \
-  --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
-  --tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-hpu \
-  --progress plain \
-  https://github.com/vllm-project/vllm-gaudi.git
-
-# push
-docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-hpu
+.buildkite/scripts/annotate-image-build.sh "$IMAGE"

.buildkite/image_build/image_build_torch_nightly.sh

@@ -40,6 +40,7 @@ docker buildx ls
 echo "--- :mag: Checking if image already exists"
 if docker manifest inspect "$IMAGE_TAG" >/dev/null 2>&1; then
   echo "Image found: $IMAGE_TAG — skipping build"
+  .buildkite/scripts/annotate-image-build.sh "$IMAGE_TAG"
   exit 0
 fi
 echo "Image not found, proceeding with build..."
@@ -66,3 +67,5 @@ docker buildx build --file docker/Dockerfile \
   --progress plain .
 
 echo "--- :white_check_mark: Torch nightly image build complete: $IMAGE_TAG"
+
+.buildkite/scripts/annotate-image-build.sh "$IMAGE_TAG"

.buildkite/image_build/image_build_xpu.sh

@@ -9,26 +9,26 @@ fi
 REGISTRY=$1
 REPO=$2
 BUILDKITE_COMMIT=$3
+IMAGE="$REGISTRY/$REPO:$BUILDKITE_COMMIT-xpu"
 
 # authenticate with AWS ECR
 aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
 aws ecr get-login-password --region us-east-1 | docker login --username AWS --password-stdin 936637512419.dkr.ecr.us-east-1.amazonaws.com || true
 
 # skip build if image already exists
-if ! docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu &> /dev/null; then
-  echo "Image not found, proceeding with build..."
-else
+if docker manifest inspect "$IMAGE" &> /dev/null; then
   echo "Image found"
-  exit 0
+else
+  echo "Image not found, proceeding with build..."
+  # build
+  docker build \
+    --file docker/Dockerfile.xpu \
+    --build-arg max_jobs=16 \
+    --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
+    --tag "$IMAGE" \
+    --progress plain .
+  # push
+  docker push "$IMAGE"
 fi
 
-# build
-docker build \
-  --file docker/Dockerfile.xpu \
-  --build-arg max_jobs=16 \
-  --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
-  --tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu \
-  --progress plain .
-
-# push
-docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu
+.buildkite/scripts/annotate-image-build.sh "$IMAGE"

.buildkite/intel_jobs/misc_intel.yaml

@@ -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

.buildkite/release-pipeline.yaml

@@ -350,6 +350,25 @@ steps:
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"
           - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"'
 
+      - label: ":docker: Build release image - x86_64 - XPU"
+        depends_on: ~
+        id: build-xpu-release-image
+        agents:
+          queue: cpu_queue_release
+        commands:
+          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
+          - |
+            DOCKER_BUILDKIT=1 docker build \
+              $(bash .buildkite/scripts/docker-build-metadata-args.sh xpu) \
+              --build-arg GIT_REPO_CHECK=1 \
+              --target vllm-openai \
+              --progress plain \
+              -f docker/Dockerfile.xpu .
+          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-xpu"
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-xpu"'
+        env:
+          DOCKER_BUILDKIT: "1"
+
       - block: "Build release image for x86_64 CPU"
         key: block-cpu-release-image-build
         depends_on: ~
@@ -445,6 +464,16 @@ steps:
           - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"
           - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 12.9 Ubuntu 24.04" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"'
 
+      - label: "Create manifest - XPU"
+        depends_on:
+          - build-xpu-release-image
+        id: create-manifest-xpu
+        agents:
+          queue: small_cpu_queue_release
+        commands:
+          - "bash .buildkite/scripts/xpu/create-xpu-ecr-manifest.sh"
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: XPU" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-xpu"'
+
       - label: "Publish nightly multi-arch image to DockerHub"
         depends_on:
           - create-multi-arch-manifest
@@ -819,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
@@ -849,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

.buildkite/scripts/annotate-image-build.sh

@@ -0,0 +1,36 @@
+#!/bin/bash
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+#
+# Append the Docker image tag(s) an image-build step pushed to a Buildkite
+# annotation, so the built image tags show up on the build page instead of
+# being buried in the job logs.
+#
+# Usage: annotate-image-build.sh <image_tag> [<image_tag> ...]
+set -euo pipefail
+
+# buildkite-agent only exists on Buildkite agents; no-op elsewhere so the
+# image build scripts stay runnable locally.
+if ! command -v buildkite-agent >/dev/null 2>&1; then
+    echo "buildkite-agent not found; skipping image tag annotation"
+    exit 0
+fi
+
+label="${BUILDKITE_LABEL:-Image build}"
+content=""
+for image in "$@"; do
+    [[ -n "$image" ]] || continue
+    content+="- **${label}**: \`${image}\`"$'\n'
+done
+
+if [[ -z "$content" ]]; then
+    echo "No image tags provided; nothing to annotate"
+    exit 0
+fi
+
+# Best-effort: a flaky annotation must never fail an otherwise successful
+# (and expensive) image build.
+if ! printf '%s' "$content" | \
+    buildkite-agent annotate --append --style 'info' --context 'docker-images'; then
+    echo "warning: failed to annotate build with image tags"
+fi

.buildkite/scripts/check-ray-compatibility.sh

@@ -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}"
 

.buildkite/scripts/publish-release-images.sh

@@ -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

.buildkite/scripts/run-rust-frontend-cargo-ci.sh

@@ -21,16 +21,20 @@ export CARGO_HOME="${CARGO_HOME:-$HOME/.cargo}"
 export RUSTUP_HOME="${RUSTUP_HOME:-$HOME/.rustup}"
 export PATH="$CARGO_HOME/bin:$PATH"
 
+PROTOC_VERSION="${PROTOC_VERSION:-31.1}"
+CARGO_BINSTALL_VERSION="${CARGO_BINSTALL_VERSION:-1.20.1}"
+UV_VERSION="${UV_VERSION:-0.11.28}"
+PYO3_PYTHON_VERSION="${PYO3_PYTHON_VERSION:-3.12}"
+
+CARGO_SORT_VERSION_REQ="${CARGO_SORT_VERSION_REQ:-2}"
+CARGO_DENY_VERSION_REQ="${CARGO_DENY_VERSION_REQ:-0.20}"
+CARGO_NEXTEST_VERSION_REQ="${CARGO_NEXTEST_VERSION_REQ:-0.9}"
+
 log_section() {
   echo "--- $*"
 }
 
 install_protoc() {
-  if command -v protoc >/dev/null 2>&1; then
-    return
-  fi
-
-  local version="${PROTOC_VERSION:-31.1}"
   local arch
   case "$(uname -m)" in
     x86_64)
@@ -45,16 +49,17 @@ install_protoc() {
       ;;
   esac
 
-  local url="https://github.com/protocolbuffers/protobuf/releases/download/v${version}/protoc-${version}-linux-${arch}.zip"
+  local url="https://github.com/protocolbuffers/protobuf/releases/download/v${PROTOC_VERSION}/protoc-${PROTOC_VERSION}-linux-${arch}.zip"
   local tmp_dir
   tmp_dir="$(mktemp -d)"
 
-  log_section "Installing protoc ${version}"
+  log_section "Installing protoc ${PROTOC_VERSION}"
   curl -L --proto '=https' --tlsv1.2 -sSf "$url" -o "$tmp_dir/protoc.zip"
   mkdir -p "$CARGO_HOME/bin"
   unzip -q "$tmp_dir/protoc.zip" bin/protoc 'include/*' -d "$CARGO_HOME"
   chmod +x "$CARGO_HOME/bin/protoc"
   rm -rf "$tmp_dir"
+  protoc --version
 }
 
 rust_toolchain() {
@@ -75,66 +80,48 @@ install_rust_toolchain() {
 }
 
 install_cargo_binstall() {
-  if command -v cargo-binstall >/dev/null 2>&1; then
-    return
-  fi
-
-  log_section "Installing cargo-binstall"
+  log_section "Installing cargo-binstall ${CARGO_BINSTALL_VERSION}"
   curl -L --proto '=https' --tlsv1.2 -sSf \
-    https://raw.githubusercontent.com/cargo-bins/cargo-binstall/main/install-from-binstall-release.sh \
-    | bash
+    "https://raw.githubusercontent.com/cargo-bins/cargo-binstall/v${CARGO_BINSTALL_VERSION}/install-from-binstall-release.sh" \
+    | env BINSTALL_VERSION="$CARGO_BINSTALL_VERSION" bash
+  cargo-binstall -V
 }
 
 install_cargo_sort() {
-  if command -v cargo-sort >/dev/null 2>&1; then
-    return
-  fi
-
-  log_section "Installing cargo-sort"
-  install_cargo_binstall
-  cargo binstall --no-confirm cargo-sort
+  log_section "Installing cargo-sort ${CARGO_SORT_VERSION_REQ}"
+  cargo binstall --no-confirm --force "cargo-sort@${CARGO_SORT_VERSION_REQ}"
 }
 
 install_cargo_deny() {
-  if command -v cargo-deny >/dev/null 2>&1; then
-    return
-  fi
-
-  log_section "Installing cargo-deny"
-  install_cargo_binstall
-  cargo binstall --no-confirm cargo-deny
+  log_section "Installing cargo-deny ${CARGO_DENY_VERSION_REQ}"
+  cargo binstall --no-confirm --force "cargo-deny@${CARGO_DENY_VERSION_REQ}"
 }
 
 install_cargo_nextest() {
-  if command -v cargo-nextest >/dev/null 2>&1; then
-    return
-  fi
-
-  log_section "Installing cargo-nextest"
-  install_cargo_binstall
-  cargo binstall --no-confirm --secure cargo-nextest
+  log_section "Installing cargo-nextest ${CARGO_NEXTEST_VERSION_REQ}"
+  cargo binstall \
+    --no-confirm \
+    --force \
+    --secure \
+    "cargo-nextest@${CARGO_NEXTEST_VERSION_REQ}"
 }
 
 install_uv() {
-  if command -v uv >/dev/null 2>&1; then
-    return
-  fi
-
-  log_section "Installing uv"
-  curl -LsSf --proto '=https' --tlsv1.2 https://astral.sh/uv/install.sh \
+  log_section "Installing uv ${UV_VERSION}"
+  curl -L --proto '=https' --tlsv1.2 -sSf \
+    "https://github.com/astral-sh/uv/releases/download/${UV_VERSION}/uv-installer.sh" \
     | env UV_INSTALL_DIR="$CARGO_HOME/bin" sh
+  uv --version
 }
 
 setup_pyo3_python() {
-  local python_version="${PYO3_PYTHON_VERSION:-3.12}"
-
-  log_section "Installing Python ${python_version} for PyO3 tests"
-  uv python install "$python_version"
+  log_section "Installing Python ${PYO3_PYTHON_VERSION} for PyO3 tests"
+  uv python install "$PYO3_PYTHON_VERSION"
   PYO3_PYTHON="$(uv python find \
     --managed-python \
     --no-project \
     --resolve-links \
-    "$python_version")"
+    "$PYO3_PYTHON_VERSION")"
   export PYO3_PYTHON
 
   local python_libdir
@@ -156,6 +143,7 @@ PY
 }
 
 run_style_clippy() {
+  install_cargo_binstall
   install_cargo_sort
   install_cargo_deny
 
@@ -168,8 +156,8 @@ run_style_clippy() {
   log_section "Checking Rust dependency bans"
   cargo deny \
     --manifest-path rust/Cargo.toml \
-    check \
     --config rust/deny.toml \
+    check \
     bans
 
   log_section "Running clippy"
@@ -186,6 +174,7 @@ run_style_clippy() {
 run_tests() {
   install_uv
   setup_pyo3_python
+  install_cargo_binstall
   install_cargo_nextest
 
   log_section "Running cargo nextest"

.buildkite/scripts/trigger-ci-build.sh

@@ -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}"

.buildkite/scripts/xpu/create-xpu-ecr-manifest.sh

@@ -0,0 +1,13 @@
+#!/bin/bash
+
+set -euo pipefail
+
+REGISTRY="public.ecr.aws/q9t5s3a7"
+REPO="vllm-release-repo"
+ARCH_TAG="${BUILDKITE_COMMIT}-$(uname -m)-xpu"
+PLATFORM_TAG="${BUILDKITE_COMMIT}-xpu"
+
+aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin ${REGISTRY}
+docker manifest rm ${REGISTRY}/${REPO}:${PLATFORM_TAG} || true
+docker manifest create ${REGISTRY}/${REPO}:${PLATFORM_TAG} ${REGISTRY}/${REPO}:${ARCH_TAG} --amend
+docker manifest push ${REGISTRY}/${REPO}:${PLATFORM_TAG}

.buildkite/scripts/xpu/push-nightly-builds-xpu.sh

@@ -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"

.buildkite/test-amd.yaml

@@ -275,7 +275,7 @@ steps:
   - pytest -v -s v1/cudagraph/test_cudagraph_mode.py
 
 - label: e2e Core (1 GPU) # TBD
-  timeout_in_minutes: 35
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
   agent_pool: mi250_1
   optional: true
@@ -439,7 +439,7 @@ steps:
 #-----------------------------------------------------------  mi250 · docker  ----------------------------------------------------------#
 
 - label: Docker Build Metadata (ROCm) # TBD
-  timeout_in_minutes: 20
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
   agent_pool: mi250_1
   no_gpu: true
@@ -467,7 +467,7 @@ steps:
 #-----------------------------------------------------  mi300 · basic_correctness  -----------------------------------------------------#
 
 - label: Basic Correctness # TBD
-  timeout_in_minutes: 95
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   fast_check: true
@@ -484,7 +484,7 @@ steps:
   - pytest -v -s basic_correctness/test_cpu_offload.py
 
 - label: Distributed Model Tests (2 GPUs) # TBD
-  timeout_in_minutes: 110
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_2
   num_gpus: 2
@@ -724,7 +724,7 @@ steps:
   - pytest -v -s distributed/test_eplb_spec_decode.py
 
 - label: Distributed Tests (2xH100-2xMI300) # TBD
-  timeout_in_minutes: 75
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_2
   num_gpus: 2
@@ -1297,7 +1297,7 @@ steps:
 #---------------------------------------------------------  mi300 · examples  ----------------------------------------------------------#
 
 - label: Examples # TBD
-  timeout_in_minutes: 90
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   optional: true
@@ -1348,7 +1348,7 @@ steps:
   - pytest -v -s tests/kernels/ir
 
 - label: Kernels Attention Test %N # TBD
-  timeout_in_minutes: 100
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   optional: true
@@ -1397,7 +1397,7 @@ steps:
   - pytest -v -s kernels/test_kda.py
 
 - label: Kernels MoE Test %N # TBD
-  timeout_in_minutes: 95
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   optional: true
@@ -1511,7 +1511,7 @@ steps:
 #----------------------------------------------------  mi300 · model_runner_v2  -------------------------------------------------------#
 
 - label: Model Runner V2 Core Tests # TBD
-  timeout_in_minutes: 60
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   optional: true
@@ -1535,7 +1535,7 @@ steps:
   - pytest -v -s entrypoints/llm/test_struct_output_generate.py -k "xgrammar and not speculative_config6 and not speculative_config7 and not speculative_config8 and not speculative_config0"
 
 - label: Model Runner V2 Examples # TBD
-  timeout_in_minutes: 60
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   optional: true
@@ -1569,7 +1569,7 @@ steps:
   - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
 
 - label: Model Runner V2 Distributed (2 GPUs) # TBD
-  timeout_in_minutes: 60
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_2
   num_gpus: 2
@@ -1661,7 +1661,7 @@ steps:
   - pytest -v -s models/test_initialization.py::test_can_initialize_small_subset
 
 - label: Basic Models Tests (Other) # TBD
-  timeout_in_minutes: 90
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   optional: true
@@ -1806,7 +1806,7 @@ steps:
   - cd .. && VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s tests/models/multimodal/generation/test_whisper.py -m core_model
 
 - label: Multi-Modal Processor # 1h 42m
-  timeout_in_minutes: 138
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   optional: true
@@ -2255,7 +2255,7 @@ steps:
   - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
 
 - label: Spec Decode Eagle # TBD
-  timeout_in_minutes: 90
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   optional: true
@@ -2472,7 +2472,7 @@ steps:
   - DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py
 
 - label: Metrics, Tracing (2 GPUs) # TBD
-  timeout_in_minutes: 65
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_2
   optional: true
@@ -2678,7 +2678,7 @@ steps:
 #------------------------------------------------------  mi300 · weight_loading  -------------------------------------------------------#
 
 - label: Weight Loading Multiple GPU # TBD
-  timeout_in_minutes: 75
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_2
   num_gpus: 2
@@ -2690,7 +2690,7 @@ steps:
   - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-amd.txt
 
 - label: Weight Loading Multiple GPU - Large Models # TBD
-  timeout_in_minutes: 75
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_2
   num_gpus: 2
@@ -3089,7 +3089,7 @@ steps:
     - pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-gfx950.txt
 
 - label: LM Eval Qwen3-5 Models (B200-MI355) # TBD
-  timeout_in_minutes: 120
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_2
   num_gpus: 2
@@ -3171,7 +3171,7 @@ steps:
 #---------------------------------------------------------  mi355 · examples  ----------------------------------------------------------#
 
 - label: Examples # TBD
-  timeout_in_minutes: 100
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
   working_dir: "/vllm-workspace/examples"
@@ -3206,7 +3206,7 @@ steps:
 #----------------------------------------------------------  mi355 · kernels  ----------------------------------------------------------#
 
 - label: Kernels (B200-MI355) # TBD
-  timeout_in_minutes: 15
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
   working_dir: "/vllm-workspace/"
@@ -3231,7 +3231,7 @@ steps:
   - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py
 
 - label: Kernels Attention Test %N # TBD
-  timeout_in_minutes: 100
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
   parallelism: 2
@@ -3248,7 +3248,7 @@ steps:
   - pytest -v -s kernels/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
 
 - label: Kernels MoE Test %N # TBD
-  timeout_in_minutes: 50
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
   parallelism: 5
@@ -3512,7 +3512,7 @@ steps:
   - pytest -v -s v1/attention
 
 - label: V1 Core + KV + Metrics # TBD
-  timeout_in_minutes: 60
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
   optional: true
@@ -3538,7 +3538,7 @@ steps:
   - pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
 
 - label: V1 Sample + Logits # TBD
-  timeout_in_minutes: 60
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
   optional: true
@@ -3558,7 +3558,7 @@ steps:
   - pytest -v -s v1/test_outputs.py
 
 - label: V1 Spec Decode # TBD
-  timeout_in_minutes: 60
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
   working_dir: "/vllm-workspace/tests"
@@ -3571,7 +3571,7 @@ steps:
 #------------------------------------------------------  mi355 · weight_loading  -------------------------------------------------------#
 
 - label: Weight Loading Multiple GPU # TBD
-  timeout_in_minutes: 75
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_2
   num_gpus: 2
@@ -3583,7 +3583,7 @@ steps:
   - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-amd.txt
 
 - label: Weight Loading Multiple GPU - Large Models # TBD
-  timeout_in_minutes: 75
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_2
   working_dir: "/vllm-workspace/tests"

.buildkite/test_areas/attention.yaml

@@ -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
@@ -12,11 +12,12 @@ steps:
     - vllm/v1/attention
     - tests/v1/attention
   commands:
-    - pytest -v -s v1/attention
+    - pytest -v -s v1/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  parallelism: 2
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 70
+      timeout_in_minutes: 95
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -30,7 +31,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
@@ -38,4 +39,5 @@ steps:
     - vllm/v1/attention
     - tests/v1/attention
   commands:
-    - pytest -v -s v1/attention
+    - pytest -v -s v1/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  parallelism: 2

.buildkite/test_areas/basic_correctness.yaml

@@ -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

.buildkite/test_areas/benchmarks.yaml

@@ -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/

.buildkite/test_areas/compile.yaml

@@ -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

.buildkite/test_areas/cuda.yaml

@@ -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

.buildkite/test_areas/disaggregated.yaml

@@ -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:
@@ -39,6 +39,19 @@ steps:
     - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
     - FLASHINFER=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
+- label: Push NixlConnector PP prefill PD accuracy (4 GPUs)
+  key: push-nixlconnector-pp-prefill-pd-accuracy-4-gpus
+  timeout_in_minutes: 30
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 4
+  source_file_dependencies:
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - tests/v1/kv_connector/nixl_integration/
+    - tests/v1/kv_connector/nixl_push_integration/
+  commands:
+    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+    - bash v1/kv_connector/nixl_push_integration/config_sweep_accuracy_test.sh
+
 - label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs)
   key: dp-ep-distributed-nixlconnector-pd-accuracy-tests-4-gpus
   timeout_in_minutes: 30
@@ -53,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:
@@ -66,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:
@@ -78,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:
@@ -91,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:
@@ -103,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:
@@ -130,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:
@@ -145,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
@@ -159,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:
@@ -173,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:

.buildkite/test_areas/distributed.yaml

@@ -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:

.buildkite/test_areas/docker.yaml

@@ -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

.buildkite/test_areas/e2e_integration.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

.buildkite/test_areas/engine.yaml

@@ -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/
@@ -60,15 +60,15 @@ steps:
     - pytest -v -s v1/e2e/general/test_async_scheduling.py
   mirror:
     amd:
-      device: mi250_1
-      timeout_in_minutes: 60
+      device: mi325_1
+      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/
@@ -76,8 +76,8 @@ steps:
     - pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py
   mirror:
     amd:
-      device: mi250_1
-      timeout_in_minutes: 35
+      device: mi325_1
+      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

.buildkite/test_areas/entrypoints.yaml

@@ -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/

.buildkite/test_areas/expert_parallelism.yaml

@@ -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

.buildkite/test_areas/kernels.yaml

@@ -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,18 +16,19 @@ 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
   - tests/kernels/test_concat_mla_q.py
   - tests/kernels/test_fused_qk_norm_rope_gate.py
   commands:
-    - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py  kernels/test_concat_mla_q.py kernels/test_fused_qk_norm_rope_gate.py
+    - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py kernels/test_concat_mla_q.py kernels/test_fused_qk_norm_rope_gate.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  parallelism: 3
 
 - 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 +42,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 +55,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 +66,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 +80,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 +107,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 +132,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 +148,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 +164,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 +173,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 +185,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 +212,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,19 +265,20 @@ 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
   - tests/kernels/helion/
   commands:
     - pip install helion==1.1.0
-    - pytest -v -s kernels/helion/
+    - pytest -v -s kernels/helion/ --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+  parallelism: 2
 
  
 - 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 +295,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 +305,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 +318,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:

.buildkite/test_areas/lm_eval.yaml

@@ -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
@@ -298,9 +298,50 @@ steps:
     - uv pip install --system 'gpt-oss[eval]==0.0.5'
     - pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-spark.txt
 
+- label: LM Eval KV-Offload (1xH200)
+  key: kv-offload-small
+  timeout_in_minutes: 30
+  device: h200_35gb
+  source_file_dependencies:
+  - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
+  - vllm/distributed/kv_transfer/kv_connector/v1/simple_cpu_offload_connector.py
+  - vllm/v1/kv_offload/
+  - vllm/v1/simple_kv_offload/
+  - tests/evals/gsm8k/test_gsm8k_offloading.py
+  commands:
+  - pytest -s -v evals/gsm8k/test_gsm8k_offloading.py -k "nemotron-h-8b or gemma-4-e4b-it"
+
+- label: LM Eval KV-Offload (2xH100)
+  key: kv-offload-medium
+  timeout_in_minutes: 30
+  device: h100
+  num_devices: 2
+  source_file_dependencies:
+  - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
+  - vllm/distributed/kv_transfer/kv_connector/v1/simple_cpu_offload_connector.py
+  - vllm/v1/kv_offload/
+  - vllm/v1/simple_kv_offload/
+  - tests/evals/gsm8k/test_gsm8k_offloading.py
+  commands:
+  - pytest -s -v evals/gsm8k/test_gsm8k_offloading.py -k "qwen3.5-35b"
+
+- label: LM Eval KV-Offload (4xH100)
+  key: kv-offload-large
+  timeout_in_minutes: 40
+  device: h100
+  num_devices: 4
+  source_file_dependencies:
+  - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
+  - vllm/distributed/kv_transfer/kv_connector/v1/simple_cpu_offload_connector.py
+  - vllm/v1/kv_offload/
+  - vllm/v1/simple_kv_offload/
+  - tests/evals/gsm8k/test_gsm8k_offloading.py
+  commands:
+  - pytest -s -v evals/gsm8k/test_gsm8k_offloading.py -k "deepseek-v4-flash"
+
 - label: MRCR Eval Small Models
   device: h200_35gb
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
   source_file_dependencies:
   - tests/evals/mrcr/
   commands:

.buildkite/test_areas/lora.yaml

@@ -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

.buildkite/test_areas/misc.yaml

@@ -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/
@@ -89,6 +89,7 @@ steps:
     - tests/v1/simple_kv_offload
     - tests/v1/worker
     - tests/v1/kv_connector/unit
+    - tests/v1/ec_connector/unit
     - tests/v1/metrics
     - tests/entrypoints/openai/correctness/test_lmeval.py
   commands:
@@ -101,6 +102,7 @@ steps:
     - pytest -v -s v1/simple_kv_offload
     - pytest -v -s v1/worker
     - pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
+    - pytest -v -s -m 'not cpu_test' v1/ec_connector/unit
     - pytest -v -s -m 'not cpu_test' v1/metrics
     # Integration test for streaming correctness (requires special branch).
     - pip install -U git+https://github.com/vllm-project/lm-evaluation-harness.git@streaming-api
@@ -108,7 +110,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 75
       depends_on:
       - image-build-amd
 
@@ -172,7 +174,7 @@ steps:
 
 - label: Regression
   key: regression
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
   - vllm/config/
@@ -195,7 +197,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 +239,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 +283,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 +294,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 +321,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 +383,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 +397,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 +413,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 +432,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

.buildkite/test_areas/model_executor.yaml

@@ -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

.buildkite/test_areas/model_runner_v2.yaml

@@ -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:

.buildkite/test_areas/models_basic.yaml

@@ -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
@@ -27,12 +27,12 @@ steps:
     # subset of supported models (the complement of the small subset in the above
     # test.) Also run if model initialization test file is modified
     - pytest -v -s models/test_initialization.py -k 'not test_can_initialize_small_subset' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
-  parallelism: 2
+  parallelism: 4
 
 - 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

.buildkite/test_areas/models_distributed.yaml

@@ -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:

.buildkite/test_areas/models_language.yaml

@@ -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:

.buildkite/test_areas/models_multimodal.yaml

@@ -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
@@ -69,22 +69,23 @@ steps:
       depends_on:
       - image-build-amd
 
-- label: Multi-Modal Processor (CPU)
+- label: Multi-Modal Processor (CPU) %N
   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
   - tests/models/registry.py
   device: cpu-medium
   commands:
-    - pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py
+    - pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
+  parallelism: 4
 
 - 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 +97,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 +165,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 75
       depends_on:
       - image-build-amd
       source_file_dependencies:

.buildkite/test_areas/plugins.yaml

@@ -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:

.buildkite/test_areas/pytorch.yaml

@@ -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

.buildkite/test_areas/quantization.yaml

@@ -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

.buildkite/test_areas/rust_frontend.yaml

@@ -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:

.buildkite/test_areas/rust_frontend_cargo.yaml

@@ -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:

.buildkite/test_areas/samplers.yaml

@@ -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
@@ -19,7 +19,7 @@ steps:
     - VLLM_USE_FLASHINFER_SAMPLER=1 pytest -v -s samplers
   mirror:
     amd:
-      device: mi250_1
+      device: mi325_1
       depends_on:
       - image-build-amd
       commands:

.buildkite/test_areas/spec_decode.yaml

@@ -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:

.buildkite/test_areas/weight_loading.yaml

@@ -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

CMakeLists.txt

@@ -70,6 +70,15 @@ endif()
 #
 set(TORCH_SUPPORTED_VERSION_CUDA "2.11.0")
 set(TORCH_SUPPORTED_VERSION_ROCM "2.11.0")
+# TORCH_NIGHTLY=1 builds run against unpinned nightly wheels, so the supported-
+# version check would always warn. Only treat it as a nightly build when the
+# value is exactly "1" (the bootstrap exports TORCH_NIGHTLY=0 by default, which
+# must NOT suppress the warning for normal builds).
+if (DEFINED ENV{TORCH_NIGHTLY} AND "$ENV{TORCH_NIGHTLY}" STREQUAL "1")
+  set(TORCH_NIGHTLY_BUILD TRUE)
+else()
+  set(TORCH_NIGHTLY_BUILD FALSE)
+endif()
 
 #
 # Try to find python package with an executable that exactly matches
@@ -175,7 +184,7 @@ endif()
 if (NOT HIP_FOUND AND NOT PYTORCH_FOUND_HIP AND CUDA_FOUND)
   set(VLLM_GPU_LANG "CUDA")
 
-  if (NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
+  if (NOT TORCH_NIGHTLY_BUILD AND NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
     message(WARNING "Pytorch version ${TORCH_SUPPORTED_VERSION_CUDA} "
       "expected for CUDA build, saw ${Torch_VERSION} instead.")
   endif()
@@ -188,7 +197,7 @@ elseif(HIP_FOUND OR PYTORCH_FOUND_HIP)
   enable_language(HIP)
 
   # ROCm 5.X and 6.X
-  if (ROCM_VERSION_DEV_MAJOR GREATER_EQUAL 5 AND
+  if (NOT TORCH_NIGHTLY_BUILD AND ROCM_VERSION_DEV_MAJOR GREATER_EQUAL 5 AND
       Torch_VERSION VERSION_LESS ${TORCH_SUPPORTED_VERSION_ROCM})
     message(WARNING "Pytorch version >= ${TORCH_SUPPORTED_VERSION_ROCM} "
       "expected for ROCm build, saw ${Torch_VERSION} instead.")
@@ -381,6 +390,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     "csrc/libtorch_stable/cuda_view.cu"
     "csrc/libtorch_stable/cuda_utils_kernels.cu"
     "csrc/libtorch_stable/activation_kernels.cu"
+    "csrc/libtorch_stable/ngram_embedding_kernels.cu"
     "csrc/libtorch_stable/quantization/activation_kernels.cu"
     "csrc/libtorch_stable/quantization/w8a8/int8/scaled_quant.cu"
     "csrc/libtorch_stable/quantization/w8a8/fp8/common.cu"

benchmarks/kernels/benchmark_relu_squared.py

@@ -0,0 +1,108 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+# Benchmark ReLUSquaredActivation: custom CUDA kernel vs forward_native, both
+# eager and under torch.compile (Inductor fuses relu+square into one kernel).
+
+import itertools
+
+import torch
+import torch.nn.functional as F
+
+import vllm.model_executor.layers.activation  # noqa: F401
+from vllm.benchmarks.lib.utils import default_vllm_config
+from vllm.triton_utils import triton
+from vllm.utils.argparse_utils import FlexibleArgumentParser
+from vllm.utils.torch_utils import STR_DTYPE_TO_TORCH_DTYPE, set_random_seed
+
+# Capped so the largest tensor stays under 2**31 elements: the shared activation
+# kernel computes the per-token pointer offset (blockIdx.x * d) in 32-bit, which
+# overflows for tensors with >2**32 elements. Realistic token counts are well
+# below this; the kernel-vs-native gap is already clear at these sizes.
+batch_size_range = [1, 16, 128]
+seq_len_range = [1, 16, 64, 1024]
+intermediate_size = [3072, 9728, 12288]
+configs = list(itertools.product(batch_size_range, seq_len_range, intermediate_size))
+
+
+@default_vllm_config()
+def benchmark_relu_squared(
+    batch_size: int,
+    seq_len: int,
+    intermediate_size: int,
+    provider: str,
+    dtype: torch.dtype,
+):
+    device = "cuda"
+    num_tokens = batch_size * seq_len
+    set_random_seed(42)
+    torch.set_default_device(device)
+
+    x = torch.randn(num_tokens, intermediate_size, dtype=dtype, device=device)
+    out = torch.empty_like(x)
+
+    def native(x: torch.Tensor) -> torch.Tensor:
+        return torch.square(F.relu(x))
+
+    # Verify the custom kernel matches the native implementation before timing.
+    ref = native(x)
+    torch.ops._C.relu_squared(out, x)
+    torch.testing.assert_close(out, ref)
+
+    if provider == "custom":
+        # Custom CUDA kernel — single fused kernel.
+        fn = lambda: torch.ops._C.relu_squared(out, x)
+    elif provider == "native":
+        # forward_native, eager — relu and square as separate ops.
+        fn = lambda: native(x)
+    elif provider == "native_compiled":
+        # forward_native under torch.compile — Inductor fuses relu+square.
+        # This is the real production baseline (custom ops are off when
+        # Inductor is enabled), so it is the comparison reviewers care about.
+        compiled = torch.compile(native)
+        compiled(x)  # warm up / trigger compilation before timing
+        fn = lambda: compiled(x)
+
+    ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
+        fn, quantiles=[0.5, 0.2, 0.8]
+    )
+    return ms, max_ms, min_ms
+
+
+if __name__ == "__main__":
+    parser = FlexibleArgumentParser(
+        description="Benchmark ReLUSquaredActivation: custom kernel vs native."
+    )
+    parser.add_argument(
+        "--dtype",
+        type=str,
+        choices=["half", "bfloat16", "float"],
+        default="bfloat16",
+    )
+    args = parser.parse_args()
+
+    dtype = STR_DTYPE_TO_TORCH_DTYPE[args.dtype]
+
+    perf_report = triton.testing.perf_report(
+        triton.testing.Benchmark(
+            x_names=["batch_size", "seq_len", "intermediate_size"],
+            x_vals=configs,
+            line_arg="provider",
+            line_vals=["custom", "native_compiled", "native"],
+            line_names=[
+                "Custom Kernel",
+                "Native (torch.compile)",
+                "Native (eager)",
+            ],
+            styles=[("blue", "-"), ("green", "-"), ("red", "-")],
+            ylabel="ms",
+            plot_name="relu_squared-eager-performance",
+            args={},
+        )
+    )
+
+    perf_report(
+        lambda batch_size, seq_len, intermediate_size, provider: benchmark_relu_squared(
+            batch_size, seq_len, intermediate_size, provider, dtype
+        )
+    ).run(print_data=True)

cmake/external_projects/flashmla.cmake

@@ -19,7 +19,7 @@ else()
   FetchContent_Declare(
         flashmla
         GIT_REPOSITORY https://github.com/vllm-project/FlashMLA
-        GIT_TAG a6ec2ba7bd0a7dff98b3f4d3e6b52b159c48d78b
+        GIT_TAG b70aff3d110a2b1a037e62eac295166b5143643a
         GIT_PROGRESS TRUE
         CONFIGURE_COMMAND ""
         BUILD_COMMAND ""

cmake/external_projects/vllm_flash_attn.cmake

@@ -39,7 +39,7 @@ else()
   FetchContent_Declare(
           vllm-flash-attn
           GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
-          GIT_TAG 2c839c33742309ec41e620bf837495ec9926c56e
+          GIT_TAG bb9a72e7dde0dc614ffc663e052cd6a19ce73a42
           GIT_PROGRESS TRUE
           # Don't share the vllm-flash-attn build between build types
           BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn

csrc/libtorch_stable/activation_kernels.cu

@@ -669,6 +669,14 @@ __device__ __forceinline__ T gelu_quick_kernel(const T& x) {
   return (T)(((float)x) / (1.0f + expf(-1.702f * (float)x)));
 }
 
+template <typename T>
+__device__ __forceinline__ T relu_squared_kernel(const T& x) {
+  // relu(x)^2 — introduced in https://arxiv.org/abs/2109.08668v2
+  const float f = (float)x;
+  const float val = f > 0.0f ? f : 0.0f;
+  return (T)(val * val);
+}
+
 }  // namespace vllm
 
 void gelu_new(torch::stable::Tensor& out,    // [..., d]
@@ -688,3 +696,9 @@ void gelu_quick(torch::stable::Tensor& out,    // [..., d]
 {
   LAUNCH_ACTIVATION_KERNEL(vllm::gelu_quick_kernel);
 }
+
+void relu_squared(torch::stable::Tensor& out,    // [..., d]
+                  torch::stable::Tensor& input)  // [..., d]
+{
+  LAUNCH_ACTIVATION_KERNEL(vllm::relu_squared_kernel);
+}

csrc/libtorch_stable/fp32_router_gemm.cu

@@ -1,13 +1,18 @@
 // SPDX-License-Identifier: Apache-2.0
 // SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 //
-// Router GEMM: activation(T) x weight(fp32) -> fp32, H=3072, E=256, M<=32.
+// Router GEMM: activation(T) x weight(fp32) -> fp32, M<=32, for the
+// supported (E, H) pairs listed at the bottom of this file.
 // Supports bf16 or fp32 activation; weight is always fp32.
 // Adapted from dsv3_router_gemm_float_out.cu.
+// (E=256, H=6144) bf16 uses a B300-tuned wide-block geometry; see
+// invokeFp32RouterGemm.
 
 #include <cuda_bf16.h>
 #include <cuda_runtime.h>
 
+#include <type_traits>
+
 // ---------------------------------------------------------------------------
 // Load helpers
 // ---------------------------------------------------------------------------
@@ -73,94 +78,113 @@ __device__ __forceinline__ void load_activation<__nv_bfloat16, 8>(
 // InputT : type of activation (float or __nv_bfloat16)
 // Weight is always fp32; output is always fp32.
 // VPT = 16 / sizeof(InputT):  4 for fp32, 8 for bf16
-template <typename InputT, int kBlockSize, int kNumTokens, int kNumExperts,
-          int kHiddenDim>
-__global__ __launch_bounds__(128, 1) void fp32_router_gemm_kernel(
-    float* out, InputT const* mat_a, float const* mat_b) {
+// Each block computes kEPB expert columns; wider blocks / kEPB > 1 are
+// selected per (shape, M) in invokeFp32RouterGemm (B300-tuned, see below).
+// kTGroups > 1 splits the tokens across groups of kBlockSize threads within
+// the block: all groups scan the same weight K-slices (group 0 misses to
+// DRAM, later groups hit L1) so weight traffic stays 1x, while per-thread
+// accumulator registers drop by kTGroups (at M=16 the 32 fp32 accumulators
+// push the kernel to 128 regs/thread and 1 block/SM).
+template <typename InputT, int kBlockSize, int kNumTokens, int kEPB,
+          int kNumExperts, int kHiddenDim, int kTGroups = 1>
+__global__ __launch_bounds__(
+    kBlockSize* kTGroups, 1) void fp32_router_gemm_kernel(float* out,
+                                                          InputT const* mat_a,
+                                                          float const* mat_b) {
   constexpr int VPT = 16 / sizeof(InputT);
   constexpr int k_elems_per_k_iteration = VPT * kBlockSize;
   constexpr int k_iterations = kHiddenDim / k_elems_per_k_iteration;
+  static_assert(kHiddenDim % k_elems_per_k_iteration == 0);
+  static_assert(kNumTokens % kTGroups == 0);
   constexpr int kWarpSize = 32;
-  constexpr int kNumWarps = kBlockSize / kWarpSize;
+  constexpr int kNumWarps = kBlockSize / kWarpSize;  // per token group
+  constexpr int kMG = kNumTokens / kTGroups;         // tokens per group
 
-  int const n_idx = blockIdx.x;
-  int const tid = threadIdx.x;
+  int const e_base = blockIdx.x * kEPB;
+  int const tid = threadIdx.x % kBlockSize;
+  int const m0 = (threadIdx.x / kBlockSize) * kMG;
   int const warpId = tid / kWarpSize;
   int const laneId = tid % kWarpSize;
 
-  float acc[kNumTokens] = {};
-  __shared__ float sm_reduction[kNumTokens][kNumWarps];
-
-  float const* b_col = mat_b + n_idx * kHiddenDim;
-
-  int k_bases[k_iterations];
-#pragma unroll
-  for (int ki = 0; ki < k_iterations; ki++) {
-    k_bases[ki] = ki * k_elems_per_k_iteration + tid * VPT;
-  }
+  float acc[kMG][kEPB] = {};
+  __shared__ float sm_reduction[kNumTokens][kEPB][kNumWarps];
 
 #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
   cudaGridDependencySynchronize();
+  // Fire the PDL trigger right after our own wait instead of at kernel end:
+  // a gridsync-ing consumer is unaffected (its wait always targets full grid
+  // completion), while a consumer that reads none of our outputs (e.g. the
+  // NVFP4 activation quant, which reads the same hidden_states) can launch
+  // now and fully overlap this kernel's body.
+  cudaTriggerProgrammaticLaunchCompletion();
 #endif
 
+#pragma unroll
   for (int ki = 0; ki < k_iterations; ki++) {
-    int const k_base = k_bases[ki];
+    int const k_base = ki * k_elems_per_k_iteration + tid * VPT;
 
-    float b_float[VPT];
-    load_weight<VPT>(b_col + k_base, b_float);
+    float b_float[kEPB][VPT];
+#pragma unroll
+    for (int e = 0; e < kEPB; e++) {
+      load_weight<VPT>(mat_b + (e_base + e) * kHiddenDim + k_base, b_float[e]);
+    }
 
 #pragma unroll
-    for (int m_idx = 0; m_idx < kNumTokens; m_idx++) {
+    for (int m_idx = 0; m_idx < kMG; m_idx++) {
       float a_float[VPT];
-      load_activation<InputT, VPT>(mat_a + m_idx * kHiddenDim + k_base,
-                                   a_float);
+      load_activation<InputT, VPT>(
+          mat_a + (size_t)(m0 + m_idx) * kHiddenDim + k_base, a_float);
 #pragma unroll
-      for (int k = 0; k < VPT; k++) {
-        acc[m_idx] += a_float[k] * b_float[k];
+      for (int e = 0; e < kEPB; e++) {
+#pragma unroll
+        for (int k = 0; k < VPT; k++) {
+          acc[m_idx][e] += a_float[k] * b_float[e][k];
+        }
       }
     }
   }
 
   // Warp-level butterfly reduction
 #pragma unroll
-  for (int m = 0; m < kNumTokens; m++) {
-    float sum = acc[m];
-    sum += __shfl_xor_sync(0xffffffff, sum, 16);
-    sum += __shfl_xor_sync(0xffffffff, sum, 8);
-    sum += __shfl_xor_sync(0xffffffff, sum, 4);
-    sum += __shfl_xor_sync(0xffffffff, sum, 2);
-    sum += __shfl_xor_sync(0xffffffff, sum, 1);
-    if (laneId == 0) sm_reduction[m][warpId] = sum;
+  for (int m = 0; m < kMG; m++) {
+#pragma unroll
+    for (int e = 0; e < kEPB; e++) {
+      float sum = acc[m][e];
+      sum += __shfl_xor_sync(0xffffffff, sum, 16);
+      sum += __shfl_xor_sync(0xffffffff, sum, 8);
+      sum += __shfl_xor_sync(0xffffffff, sum, 4);
+      sum += __shfl_xor_sync(0xffffffff, sum, 2);
+      sum += __shfl_xor_sync(0xffffffff, sum, 1);
+      if (laneId == 0) sm_reduction[m0 + m][e][warpId] = sum;
+    }
   }
 
   __syncthreads();
 
-  if (tid == 0) {
+  // Parallel finalize: one thread per (m, e) output.
+  for (int idx = threadIdx.x; idx < kNumTokens * kEPB;
+       idx += kBlockSize * kTGroups) {
+    int const m = idx / kEPB;
+    int const e = idx % kEPB;
+    float final_sum = 0.0f;
 #pragma unroll
-    for (int m = 0; m < kNumTokens; m++) {
-      float final_sum = 0.0f;
-#pragma unroll
-      for (int w = 0; w < kNumWarps; w++) final_sum += sm_reduction[m][w];
-      out[m * kNumExperts + n_idx] = final_sum;
-    }
+    for (int w = 0; w < kNumWarps; w++) final_sum += sm_reduction[m][e][w];
+    out[m * kNumExperts + e_base + e] = final_sum;
   }
-
-#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
-  cudaTriggerProgrammaticLaunchCompletion();
-#endif
 }
 
 // ---------------------------------------------------------------------------
 // Launcher
 // ---------------------------------------------------------------------------
 
-template <typename InputT, int kNumTokens, int kNumExperts, int kHiddenDim>
-void invokeFp32RouterGemm(float* output, InputT const* mat_a,
-                          float const* mat_b, cudaStream_t stream) {
-  constexpr int kBlockSize = 128;
+template <typename InputT, int kBlockSize, int kEPB, int kNumTokens,
+          int kNumExperts, int kHiddenDim, int kTGroups = 1>
+static void launchFp32RouterGemm(float* output, InputT const* mat_a,
+                                 float const* mat_b, cudaStream_t stream) {
+  static_assert(kNumExperts % kEPB == 0);
   cudaLaunchConfig_t config;
-  config.gridDim = kNumExperts;
-  config.blockDim = kBlockSize;
+  config.gridDim = kNumExperts / kEPB;
+  config.blockDim = kBlockSize * kTGroups;
   config.dynamicSmemBytes = 0;
   config.stream = stream;
   cudaLaunchAttribute attrs[1];
@@ -168,15 +192,112 @@ void invokeFp32RouterGemm(float* output, InputT const* mat_a,
   attrs[0].val.programmaticStreamSerializationAllowed = 1;
   config.numAttrs = 1;
   config.attrs = attrs;
-  cudaLaunchKernelEx(&config,
-                     fp32_router_gemm_kernel<InputT, kBlockSize, kNumTokens,
-                                             kNumExperts, kHiddenDim>,
-                     output, mat_a, mat_b);
+  cudaLaunchKernelEx(
+      &config,
+      fp32_router_gemm_kernel<InputT, kBlockSize, kNumTokens, kEPB, kNumExperts,
+                              kHiddenDim, kTGroups>,
+      output, mat_a, mat_b);
+}
+
+static bool isBlackwellFamily() {
+  static int sm = []() {
+    int dev = 0, major = 0, minor = 0;
+    cudaGetDevice(&dev);
+    cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, dev);
+    cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, dev);
+    return major * 10 + minor;
+  }();
+  return sm >= 100;
+}
+
+template <typename InputT, int kNumTokens, int kNumExperts, int kHiddenDim>
+void invokeFp32RouterGemm(float* output, InputT const* mat_a,
+                          float const* mat_b, cudaStream_t stream) {
+  // Geometry tuned on B300 per supported shape, bf16 activation, under a
+  // production-fidelity harness (CUDA-graph replay, per-layer cold weights).
+  // GLM-5.2 (E=256, H=6144):
+  //   M <= 4        : BS=768, EPB=1 (2.7us vs cast+cuBLAS 8.1us at M=1)
+  //   M in [5, 15]
+  //   or odd        : BS=384, EPB=2 (crossover vs BS=768 measured in (4, 8))
+  //   M >= 16, even : BS=192, EPB=2, 2 token groups (M=16 4.79us vs 5.04,
+  //                   M=24 5.71 vs 6.38, M=32 6.79 vs 7.72; M=12 loses at
+  //                   0.97x, so the boundary is 16).
+  // Only enabled on the Blackwell family where it was validated; Hopper and
+  // other shapes / fp32 activation keep the legacy geometry.
+  if constexpr (std::is_same_v<InputT, __nv_bfloat16> && kNumExperts == 256 &&
+                kHiddenDim == 6144) {
+    if (!isBlackwellFamily()) {
+      launchFp32RouterGemm<InputT, 128, 1, kNumTokens, kNumExperts, kHiddenDim>(
+          output, mat_a, mat_b, stream);
+      return;
+    }
+    if constexpr (kNumTokens <= 4) {
+      launchFp32RouterGemm<InputT, 768, 1, kNumTokens, kNumExperts, kHiddenDim>(
+          output, mat_a, mat_b, stream);
+    } else if constexpr (kNumTokens >= 16 && kNumTokens % 2 == 0) {
+      launchFp32RouterGemm<InputT, 192, 2, kNumTokens, kNumExperts, kHiddenDim,
+                           2>(output, mat_a, mat_b, stream);
+    } else {
+      launchFp32RouterGemm<InputT, 384, 2, kNumTokens, kNumExperts, kHiddenDim>(
+          output, mat_a, mat_b, stream);
+    }
+  } else if constexpr (std::is_same_v<InputT, __nv_bfloat16> &&
+                       kNumExperts == 128 && kHiddenDim == 6144) {
+    // MiniMax-M3. Legacy 128/1 only fills 128 blocks and pays the same
+    // accumulator register cliffs; B300 sweep:
+    //   even M in [6, 10] : BS=384, EPB=1, 2 token groups (1.26-1.43x)
+    //   even M >= 12      : BS=192, EPB=1, 2 token groups (1.59-1.66x at
+    //                       M >= 18; re-measured on B300+B200: 192 also wins
+    //                       M=12/14 by 5-11%% on both, ties 384 at 16)
+    //   M <= 5 / odd      : BS=384, EPB=1 (1.03-1.19x)
+    if (!isBlackwellFamily()) {
+      launchFp32RouterGemm<InputT, 128, 1, kNumTokens, kNumExperts, kHiddenDim>(
+          output, mat_a, mat_b, stream);
+      return;
+    }
+    if constexpr (kNumTokens >= 12 && kNumTokens % 2 == 0) {
+      launchFp32RouterGemm<InputT, 192, 1, kNumTokens, kNumExperts, kHiddenDim,
+                           2>(output, mat_a, mat_b, stream);
+    } else if constexpr (kNumTokens >= 6 && kNumTokens % 2 == 0) {
+      launchFp32RouterGemm<InputT, 384, 1, kNumTokens, kNumExperts, kHiddenDim,
+                           2>(output, mat_a, mat_b, stream);
+    } else {
+      launchFp32RouterGemm<InputT, 384, 1, kNumTokens, kNumExperts, kHiddenDim>(
+          output, mat_a, mat_b, stream);
+    }
+  } else if constexpr (std::is_same_v<InputT, __nv_bfloat16> &&
+                       kNumExperts == 256 && kHiddenDim == 3072) {
+    // MiniMax-M2/M2.5. The 3.1MB weight is latency-floor bound at small M
+    // (legacy already optimal); token groups win only at even M >= 8
+    // (1.05-1.17x). EPB crossover measured between 12 and 16.
+    if (!isBlackwellFamily()) {
+      launchFp32RouterGemm<InputT, 128, 1, kNumTokens, kNumExperts, kHiddenDim>(
+          output, mat_a, mat_b, stream);
+      return;
+    }
+    if constexpr (kNumTokens >= 14 && kNumTokens % 2 == 0) {
+      // M=14 originally measured 0.91x and stayed on legacy; two fresh
+      // sweeps (B300 dev1 + B200) both put 192/2/tg2 ahead by 3.5-4%%.
+      launchFp32RouterGemm<InputT, 192, 2, kNumTokens, kNumExperts, kHiddenDim,
+                           2>(output, mat_a, mat_b, stream);
+    } else if constexpr (kNumTokens >= 8 && kNumTokens <= 12 &&
+                         kNumTokens % 2 == 0) {
+      launchFp32RouterGemm<InputT, 192, 1, kNumTokens, kNumExperts, kHiddenDim,
+                           2>(output, mat_a, mat_b, stream);
+    } else {
+      launchFp32RouterGemm<InputT, 128, 1, kNumTokens, kNumExperts, kHiddenDim>(
+          output, mat_a, mat_b, stream);
+    }
+  } else {
+    launchFp32RouterGemm<InputT, 128, 1, kNumTokens, kNumExperts, kHiddenDim>(
+        output, mat_a, mat_b, stream);
+  }
 }
 
 // ---------------------------------------------------------------------------
 // Explicit instantiations: M=1..32, for both input types, for the supported
-// (E, H) pairs:  (256, 3072) [MiniMax-M2/M2.5]  and  (128, 6144) [MiniMax-M3].
+// (E, H) pairs:  (256, 3072) [MiniMax-M2/M2.5],  (128, 6144) [MiniMax-M3]
+// and  (256, 6144) [GLM-5.2].
 // ---------------------------------------------------------------------------
 
 #define INSTANTIATE(T, M, E, H)                                    \
@@ -221,6 +342,8 @@ INSTANTIATE_ALL(float, 256, 3072)
 INSTANTIATE_ALL(__nv_bfloat16, 256, 3072)
 INSTANTIATE_ALL(float, 128, 6144)
 INSTANTIATE_ALL(__nv_bfloat16, 128, 6144)
+INSTANTIATE_ALL(float, 256, 6144)
+INSTANTIATE_ALL(__nv_bfloat16, 256, 6144)
 
 #undef INSTANTIATE_ALL
 #undef INSTANTIATE

csrc/libtorch_stable/fp32_router_gemm_entry.cu

@@ -25,10 +25,12 @@ inline int getSMVersion() {
 static constexpr int FP32_MAX_TOKENS = 32;
 
 // Supported (hidden_dim, num_experts) pairs (must match the instantiations in
-// fp32_router_gemm.cu): (3072, 256) for MiniMax-M2/M2.5, (6144, 128) for M3.
+// fp32_router_gemm.cu): (3072, 256) for MiniMax-M2/M2.5, (6144, 128) for M3,
+// (6144, 256) for GLM-5.2.
 static inline bool fp32_router_gemm_supported(int hidden_dim, int num_experts) {
   return (hidden_dim == 3072 && num_experts == 256) ||
-         (hidden_dim == 6144 && num_experts == 128);
+         (hidden_dim == 6144 && num_experts == 128) ||
+         (hidden_dim == 6144 && num_experts == 256);
 }
 
 // Forward declarations — 4 template params must match fp32_router_gemm.cu
@@ -77,6 +79,9 @@ void dispatchFp32RouterGemm(int num_experts, int hidden_dim, int num_tokens,
   } else if (num_experts == 128 && hidden_dim == 6144) {
     Fp32LoopUnroller<InputT, 128, 6144, 1, FP32_MAX_TOKENS>::unroll(
         num_tokens, output, mat_a, mat_b, stream);
+  } else if (num_experts == 256 && hidden_dim == 6144) {
+    Fp32LoopUnroller<InputT, 256, 6144, 1, FP32_MAX_TOKENS>::unroll(
+        num_tokens, output, mat_a, mat_b, stream);
   } else {
     throw std::invalid_argument(
         "fp32_router_gemm: unsupported (hidden_dim, num_experts) pair");
@@ -111,7 +116,7 @@ void fp32_router_gemm(
   STD_TORCH_CHECK(
       fp32_router_gemm_supported(hidden_dim, num_experts),
       "fp32_router_gemm: supported (hidden_dim, num_experts) pairs are "
-      "(3072, 256) and (6144, 128)");
+      "(3072, 256), (6144, 128) and (6144, 256)");
   STD_TORCH_CHECK(num_tokens <= FP32_MAX_TOKENS,
                   "fp32_router_gemm: num_tokens must be in [0, 32]");
   STD_TORCH_CHECK(

csrc/libtorch_stable/fused_minimax_m3_qknorm_rope_kv_insert_kernel.cu

@@ -39,12 +39,10 @@
  * The IQ/IK warps address the index_q/index_k sub-blocks *inside* qkv at the
  * fixed physical offsets (nq+2*nkv)*128 and (nq+2*nkv+niq)*128.
  *
- * Dense vs sparse is a compile-time choice via the ``kIsSparse``/``kInsertKV``
- * template bools (3 instantiations: dense <false,false>, sparse-profiling
- * <true,false>, sparse-serving <true,true>), so the index slots, the V slots
- * and the cache inserts fold away entirely on paths that don't use them. The
- * dense layer passes no caches/index: norm+RoPE happens in place and the
- * generic ``Attention`` layer owns the cache write.
+ * Dense vs sparse row layout and index-branch processing are separate template
+ * choices. Skip-index-topk reuse layers still have sparse rows and insert main
+ * K/V cache entries, but compile away index_q/index_k work and index-cache
+ * writes.
  *
  * Q/K and (sparse) index_q/index_k are all rewritten in place inside the fused
  * ``qkv`` tensor.  Caches (bf16) are scatter-written by slot.
@@ -223,10 +221,25 @@ __device__ __forceinline__ void storeCacheElems(
     // model dtype directly. FP8 cache dtypes use the conversion path below.
     storeElems<scalar_t>(reinterpret_cast<scalar_t*>(dst), elems);
   } else {
-#pragma unroll
+#ifdef USE_ROCM
+    // Match ROCm's model-dtype materialization before FP8 cache conversion.
+    using Converter = vllm::_typeConvert<scalar_t>;
+    using rounded_t = typename Converter::hip_type;
+    rounded_t rounded[kElemsPerLane];
+  #pragma unroll
+    for (int i = 0; i < kElemsPerLane; i++) {
+      rounded[i] = Converter::convert(elems[i]);
+    }
+  #pragma unroll
+    for (int i = 0; i < kElemsPerLane; i++) {
+      dst[i] = fp8::scaled_convert<cache_t, rounded_t, kv_dt>(rounded[i], 1.0f);
+    }
+#else
+  #pragma unroll
     for (int i = 0; i < kElemsPerLane; i++) {
       dst[i] = fp8::scaled_convert<cache_t, float, kv_dt>(elems[i], 1.0f);
     }
+#endif
   }
 }
 
@@ -262,20 +275,25 @@ __device__ __forceinline__ void storeElemsFp8(
 // Grid: 1D, ceil(num_tokens * slots_per_token / warps_per_block).
 // Each warp = one (token, slot).
 //
-// `kIsSparse` and `kInsertKV` are compile-time template bools, so all the
-// branch decisions that distinguish the dense layer from the sparse layer
-// (index slots, KV/index inserts, V slots) fold away per instantiation.
-// Three instantiations are built: dense <false,false>, sparse-profiling
-// <true,false> and sparse-serving <true,true>.  Slots per token:
+// `kHasIndex`, `kProcessIndex`, and `kInsertKV` are compile-time template
+// bools, so branch decisions that distinguish the dense layer from the sparse
+// layer (index slots, KV/index inserts, V slots) fold away per instantiation.
+// Slots per token:
 //     Q : nq                            (always — norm+RoPE)
 //     K : nkv                           (always — norm+RoPE; +K-cache insert)
 //     V : nkv  only if kInsertKV        (V-cache insert; no warps in dense)
-//     IQ: niq  only if kIsSparse        (norm+RoPE)
-//     IK: 1    only if kIsSparse        (norm+RoPE; +index-cache insert)
+//     IQ: niq  only if kProcessIndex    (norm+RoPE)
+//     IK: 1    only if kProcessIndex    (norm+RoPE; +index-cache insert)
 // cache_t/kv_dt: main attention KV-cache dtype (auto/fp8). out_idx_t/kFp8Idx:
 // indexer index-K cache + index-Q output dtype (scalar_t or e4m3 byte).
+// kHasIndex means the qkv row is laid out as sparse [q|k|v|index_q|index_k].
+// kProcessIndex controls whether this launch actually norms/ropes the index
+// branch and writes index_q/index_k outputs. Skip-index-topk reuse layers keep
+// kHasIndex=true but set kProcessIndex=false.
 template <typename scalar_t, typename cache_t, Fp8KVCacheDataType kv_dt,
-          typename out_idx_t, bool kIsSparse, bool kInsertKV, bool kFp8Idx>
+          typename out_idx_t, bool kHasIndex, bool kInsertKV,
+          bool kProcessIndex,
+          bool kFp8Idx>
 __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
     scalar_t* __restrict__ qkv,  // [N, qkv_row] in/out (packs index if sparse)
     scalar_t* __restrict__ q_out,         // [N, nq*128] contiguous, or nullptr
@@ -288,16 +306,15 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
     int64_t const* __restrict__ positions,       // [N] i64
     int64_t const* __restrict__ slot_mapping,    // main K/V slots or nullptr
     int64_t const* __restrict__ index_slot_mapping,  // index K slots/nullptr
-    cache_t* __restrict__ kv_cache,       // [nb,2,bs,nkv,128] or nullptr
+    cache_t* __restrict__ kv_cache,       // [nb,nkv,bs,2*128] or nullptr
     out_idx_t* __restrict__ index_cache,  // [nb*bs, 128]; scalar_t or e4m3 byte
     float const eps, int const rotary_dim, int const num_tokens, int const nq,
     int const nkv, int const niq, int const block_size,
-    // kv_cache strides (in elements) for logical shape [nb, 2, bs, nkv, 128].
-    // The head_dim (last) dim is always innermost-contiguous (stride 1), so the
-    // NHD/HND layout choice is fully captured by these four strides: NHD keeps
-    // s_token < s_head, HND swaps them. dim_base addresses head_dim directly.
-    int64_t const kv_s_block, int64_t const kv_s_kv, int64_t const kv_s_token,
-    int64_t const kv_s_head) {
+    // kv_cache strides (in elements) for logical shape [nb, nkv, bs, 2*128].
+    // The content (last) dim is always innermost-contiguous (stride 1), so the
+    // NHD/HND layout choice is captured by the head/token strides.
+    int64_t const kv_s_block, int64_t const kv_s_head, int64_t const kv_s_token,
+    int64_t const kv_s_dim) {
 #if (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ < 800) && !defined(USE_ROCM)
   // _typeConvert<BFloat16> is unavailable on pre-Ampere; the M3 kernel only
   // runs with bf16/fp16 inputs in practice.  Discard the bf16 body there.
@@ -309,9 +326,12 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
     int const laneId = threadIdx.x % 32;
     int const globalWarpIdx = blockIdx.x * warpsPerBlock + (threadIdx.x / 32);
 
+    static_assert(!kProcessIndex || kHasIndex,
+                  "index processing requires sparse row layout");
+
     // Slot layout (compile-time gated: dense has neither V nor index slots).
     int const v_slots = kInsertKV ? nkv : 0;
-    int const idx_slots = kIsSparse ? niq + 1 : 0;
+    int const idx_slots = kProcessIndex ? niq + 1 : 0;
     int const slots_per_token = nq + nkv + v_slots + idx_slots;
 
     int const tokenIdx = globalWarpIdx / slots_per_token;
@@ -322,14 +342,14 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
     int const k_begin = nq;
     int const v_begin = nq + nkv;             // valid only when kInsertKV
     int const iq_begin = nq + nkv + v_slots;  // index block start
-    int const ik_slot = iq_begin + niq;       // valid only when kIsSparse
+    int const ik_slot = iq_begin + niq;       // valid only when kProcessIndex
 
     bool const isQ = slot < k_begin;
     bool const isK = slot >= k_begin && slot < v_begin;
     bool isV = false;
     if constexpr (kInsertKV) isV = slot >= v_begin && slot < v_begin + nkv;
     bool isIQ = false, isIK = false;
-    if constexpr (kIsSparse) {
+    if constexpr (kProcessIndex) {
       isIQ = slot >= iq_begin && slot < ik_slot;
       isIK = slot == ik_slot;
     }
@@ -337,7 +357,7 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
     int const dim_base = laneId * kElemsPerLane;
     // Physical row width of qkv: the dense layer packs [q|k|v]; the sparse
     // layer additionally packs [index_q (niq heads) | index_k (1 head)].
-    int const qkv_row = (nq + 2 * nkv + (kIsSparse ? (niq + 1) : 0)) * kHeadDim;
+    int const qkv_row = (nq + 2 * nkv + (kHasIndex ? (niq + 1) : 0)) * kHeadDim;
 
     // ── Resolve source pointer + per-branch parameters. ────────────────────
     scalar_t* row_ptr = nullptr;       // in-place output location
@@ -368,10 +388,13 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
       row_ptr = qkv + static_cast<int64_t>(tokenIdx) * qkv_row +
                 (nq + 2 * nkv + ih) * kHeadDim;
       norm_w = iq_norm_w;
-    } else {  // isIK -- single shared index key at (nq+2*nkv+niq)*128.
+    } else if (isIK) {
+      // Single shared index key at (nq+2*nkv+niq)*128.
       row_ptr = qkv + static_cast<int64_t>(tokenIdx) * qkv_row +
                 (nq + 2 * nkv + niq) * kHeadDim;
       norm_w = ik_norm_w;
+    } else {
+      return;
     }
 
     // Store destination.  Q and index_q are gathered into dedicated contiguous
@@ -427,9 +450,12 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
     // ── Cache inserts (sparse serving only). ───────────────────────────────
     if constexpr (kInsertKV) {
       // Guard (not early-return) so every thread reaches the PDL trigger below.
-      int64_t const sm = (isK || isV)
-                             ? slot_mapping[tokenIdx]
-                             : (isIK ? index_slot_mapping[tokenIdx] : -1);
+      int64_t sm = -1;
+      if (isK || isV) {
+        sm = slot_mapping[tokenIdx];
+      } else if constexpr (kProcessIndex) {
+        if (isIK) sm = index_slot_mapping[tokenIdx];
+      }
       if (sm >= 0) {  // skip padded / unscheduled tokens
         if (isIK) {
           if constexpr (kFp8Idx) {
@@ -438,16 +464,16 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
             storeElems<scalar_t>(index_cache + sm * kHeadDim + dim_base, elems);
           }
         } else if (isK || isV) {
-          // kv_cache logical shape [num_blocks, 2, block_size, nkv, head_dim].
+          // kv_cache logical shape [num_blocks, nkv, block_size, 2*head_dim].
           // Paging is logical (block = sm/block_size, token = sm%block_size);
           // the physical NHD/HND layout is honoured via the passed strides.
           int64_t const b = sm / block_size;
           int64_t const t = sm % block_size;
           int const kv = isK ? 0 : 1;
-          int64_t const off =
-              b * kv_s_block + kv * kv_s_kv + t * kv_s_token + head * kv_s_head;
-          storeCacheElems<scalar_t, cache_t, kv_dt>(kv_cache + off + dim_base,
-                                                    elems);
+          int64_t const off = b * kv_s_block + head * kv_s_head +
+                              t * kv_s_token +
+                              (kv * kHeadDim + dim_base) * kv_s_dim;
+          storeCacheElems<scalar_t, cache_t, kv_dt>(kv_cache + off, elems);
         }
       }
     }
@@ -474,14 +500,14 @@ void launchFusedMiniMaxM3(
     int64_t const* index_slot_mapping, cache_t* kv_cache, void* index_cache,
     float const eps, int const rotary_dim, int const num_tokens, int const nq,
     int const nkv, int const niq, int const block_size,
-    int64_t const kv_s_block, int64_t const kv_s_kv, int64_t const kv_s_token,
-    int64_t const kv_s_head, bool const has_index, bool const insert_kv,
-    bool const fp8_idx, cudaStream_t stream) {
+    int64_t const kv_s_block, int64_t const kv_s_head, int64_t const kv_s_token,
+    int64_t const kv_s_dim, bool const has_index, bool const insert_kv,
+    bool const process_index, bool const fp8_idx, cudaStream_t stream) {
   // Index outputs are scalar_t (bf16) or e4m3 bytes (uint8_t); reinterpret the
   // void* pointers per instantiation in the LAUNCH macro.
   // Slot count must match the kernel's compile-time gating.
   int const v_slots = insert_kv ? nkv : 0;
-  int const idx_slots = has_index ? niq + 1 : 0;
+  int const idx_slots = process_index ? niq + 1 : 0;
   int const slots_per_token = nq + nkv + v_slots + idx_slots;
 
   constexpr int kBlockSize = 256;
@@ -508,50 +534,59 @@ void launchFusedMiniMaxM3(
   config.attrs = attrs;
   config.numAttrs = (sm_version >= 90) ? 1 : 0;
 
-  #define LAUNCH(IS_SPARSE, INSERT, FP8, OUT_T)                                \
+  #define LAUNCH(HAS_INDEX, INSERT, PROCESS_INDEX, FP8, OUT_T)                 \
     cudaLaunchKernelEx(                                                        \
         &config,                                                               \
         fusedMiniMaxM3QNormRopeKVInsertKernel<scalar_t, cache_t, kv_dt, OUT_T, \
-                                              IS_SPARSE, INSERT, FP8>,         \
+                                              HAS_INDEX, INSERT,               \
+                                              PROCESS_INDEX, FP8>,             \
         qkv, q_out, reinterpret_cast<OUT_T*>(index_q_out), q_norm_w, k_norm_w, \
         iq_norm_w, ik_norm_w, cos_sin_cache, positions, slot_mapping,          \
         index_slot_mapping, kv_cache, reinterpret_cast<OUT_T*>(index_cache),   \
         eps, rotary_dim, num_tokens, nq, nkv, niq, block_size, kv_s_block,     \
-        kv_s_kv, kv_s_token, kv_s_head)
+        kv_s_head, kv_s_token, kv_s_dim)
 #else
   // ROCm: standard kernel launch syntax (no PDL/stream serialization).
   // clang-format off
-  #define LAUNCH(IS_SPARSE, INSERT, FP8, OUT_T)                              \
-    fusedMiniMaxM3QNormRopeKVInsertKernel<scalar_t, cache_t, kv_dt, OUT_T,   \
-                                          IS_SPARSE, INSERT, FP8>            \
-        <<<grid, kBlockSize, 0, stream>>>(                                   \
-            qkv, q_out, reinterpret_cast<OUT_T*>(index_q_out), q_norm_w,     \
-            k_norm_w, iq_norm_w, ik_norm_w, cos_sin_cache, positions,        \
-            slot_mapping, index_slot_mapping, kv_cache,                      \
-            reinterpret_cast<OUT_T*>(index_cache), eps, rotary_dim,          \
-            num_tokens, nq, nkv, niq, block_size, kv_s_block, kv_s_kv,       \
-            kv_s_token, kv_s_head)
+  #define LAUNCH(HAS_INDEX, INSERT, PROCESS_INDEX, FP8, OUT_T)              \
+    fusedMiniMaxM3QNormRopeKVInsertKernel<                                  \
+        scalar_t, cache_t, kv_dt, OUT_T, HAS_INDEX, INSERT, PROCESS_INDEX,  \
+        FP8><<<grid, kBlockSize, 0, stream>>>(                              \
+        qkv, q_out, reinterpret_cast<OUT_T*>(index_q_out), q_norm_w,        \
+        k_norm_w, iq_norm_w, ik_norm_w, cos_sin_cache, positions,           \
+        slot_mapping, index_slot_mapping, kv_cache,                         \
+        reinterpret_cast<OUT_T*>(index_cache), eps, rotary_dim, num_tokens, \
+        nq, nkv, niq, block_size, kv_s_block, kv_s_head, kv_s_token,         \
+        kv_s_dim)
   // clang-format on
 #endif
 
   if (has_index) {
-    if (insert_kv) {
-      if (fp8_idx) {
-        LAUNCH(true, true, true, uint8_t);  // sparse serving, fp8 index outputs
+    if (!process_index) {
+      if (insert_kv) {
+        LAUNCH(true, true, false, false, scalar_t);
       } else {
-        LAUNCH(true, true, false, scalar_t);  // sparse serving, bf16
+        LAUNCH(true, false, false, false, scalar_t);
+      }
+    } else if (insert_kv) {
+      if (fp8_idx) {
+        LAUNCH(true, true, true, true,
+               uint8_t);  // sparse serving, fp8 index outputs
+      } else {
+        LAUNCH(true, true, true, false, scalar_t);  // sparse serving, bf16
       }
     } else {
       if (fp8_idx) {
-        LAUNCH(true, false, true, uint8_t);  // sparse profiling, fp8 index_q
+        LAUNCH(true, false, true, true,
+               uint8_t);  // sparse profiling, fp8 index_q
       } else {
-        LAUNCH(true, false, false, scalar_t);  // sparse profiling, bf16
+        LAUNCH(true, false, true, false, scalar_t);  // sparse profiling, bf16
       }
     }
   } else {
     // Dense layer: never has an index branch and never inserts here (the
     // generic Attention layer owns the KV insert).
-    LAUNCH(false, false, false, scalar_t);
+    LAUNCH(false, false, false, false, scalar_t);
   }
 #undef LAUNCH
 }
@@ -559,6 +594,7 @@ void launchFusedMiniMaxM3(
 }  // namespace minimax_m3_fused_ops
 }  // namespace vllm
 
+// clang-format off
 #define CALL_FUSED_MINIMAX_M3(_RAW_T, CACHE_T, KV_DTYPE)                       \
   vllm::minimax_m3_fused_ops::launchFusedMiniMaxM3<st, CACHE_T, KV_DTYPE>(     \
       reinterpret_cast<st*>(qkv.data_ptr()),                                   \
@@ -568,24 +604,29 @@ void launchFusedMiniMaxM3(
           : nullptr,                                                           \
       reinterpret_cast<st const*>(q_norm_weight.data_ptr()),                   \
       reinterpret_cast<st const*>(k_norm_weight.data_ptr()),                   \
-      has_index ? reinterpret_cast<st const*>(index_q_norm_weight->data_ptr()) \
-                : nullptr,                                                     \
-      has_index ? reinterpret_cast<st const*>(index_k_norm_weight->data_ptr()) \
-                : nullptr,                                                     \
+      process_index                                                            \
+          ? reinterpret_cast<st const*>(index_q_norm_weight->data_ptr())       \
+          : nullptr,                                                           \
+      process_index                                                            \
+          ? reinterpret_cast<st const*>(index_k_norm_weight->data_ptr())       \
+          : nullptr,                                                           \
       reinterpret_cast<st const*>(cos_sin_cache.data_ptr()),                   \
       reinterpret_cast<int64_t const*>(positions.data_ptr()),                  \
       insert_kv ? reinterpret_cast<int64_t const*>(slot_mapping->data_ptr())   \
                 : nullptr,                                                     \
-      insert_kv ? reinterpret_cast<int64_t const*>(                            \
-                      effective_index_slot_mapping->data_ptr())                \
-                : nullptr,                                                     \
+      (insert_kv && process_index)                                             \
+          ? reinterpret_cast<int64_t const*>(                                  \
+                effective_index_slot_mapping->data_ptr())                      \
+          : nullptr,                                                           \
       insert_kv ? reinterpret_cast<CACHE_T*>(kv_cache->data_ptr()) : nullptr,  \
-      (insert_kv && has_index)                                                 \
+      (insert_kv && process_index)                                             \
           ? reinterpret_cast<void*>(index_cache->data_ptr())                   \
           : nullptr,                                                           \
       static_cast<float>(eps), static_cast<int>(rotary_dim), num_tokens, nq,   \
-      nkv, niq, static_cast<int>(block_size), kv_s_block, kv_s_kv, kv_s_token, \
-      kv_s_head, has_index, insert_kv, fp8_idx, stream)
+      nkv, niq, static_cast<int>(block_size), kv_s_block, kv_s_head,           \
+      kv_s_token, kv_s_dim, has_index, insert_kv, process_index, fp8_idx,       \
+      stream)
+// clang-format on
 
 // ────────────────────────────────────────────────────────────────────────────
 // Torch op wrapper
@@ -602,13 +643,13 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
     int64_t num_index_heads,                                  // niq; 0 => dense
     std::optional<torch::stable::Tensor> slot_mapping,        // [N] i64
     std::optional<torch::stable::Tensor> index_slot_mapping,  // [N] i64
-    std::optional<torch::stable::Tensor> kv_cache,     // [nb,2,bs,nkv,128]
+    std::optional<torch::stable::Tensor> kv_cache,     // [nb,nkv,bs,2*128]
     std::optional<torch::stable::Tensor> index_cache,  // [nb,bs,128]
     int64_t block_size,
     std::optional<torch::stable::Tensor> q_out,  // [N, nq*128] contiguous
     std::optional<torch::stable::Tensor>
         index_q_out,  // [N, niq*128] contiguous
-    const std::string& kv_cache_dtype) {
+    const std::string& kv_cache_dtype, bool skip_index_branch) {
   STD_TORCH_CHECK(qkv.is_cuda() && qkv.is_contiguous(),
                   "qkv must be contiguous CUDA");
   STD_TORCH_CHECK(
@@ -647,6 +688,7 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
   // (1 head)]) right after [q|k|v] in the same row; the dense layer does not.
   bool const has_index = niq > 0;
   bool const insert_kv = kv_cache.has_value();
+  bool const process_index = has_index && !skip_index_branch;
   vllm::Fp8KVCacheDataType const kv_dt =
       vllm::get_fp8_kv_cache_data_type(kv_cache_dtype);
   int const kHeadDim = vllm::minimax_m3_fused_ops::kHeadDim;
@@ -662,7 +704,9 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
   STD_TORCH_CHECK(
       !insert_kv || has_index,
       "insert mode (kv_cache) requires the index branch (sparse layer)");
-  if (has_index) {
+  STD_TORCH_CHECK(has_index || !skip_index_branch,
+                  "skip_index_branch requires sparse qkv rows");
+  if (process_index) {
     STD_TORCH_CHECK(
         index_q_norm_weight.has_value() && index_k_norm_weight.has_value(),
         "index branch requires both index norm weights");
@@ -673,24 +717,26 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
                         index_k_norm_weight->numel() == kHeadDim,
                     "index norm weights must have 128 elements");
   }
-  // kv_cache strides (logical shape [nb, 2, bs, nkv, head_dim]). Read straight
+  // kv_cache strides (logical shape [nb, nkv, bs, 2*head_dim]). Read straight
   // off the tensor so the kernel honours whatever physical layout the attention
-  // backend allocated (NHD: stride order (0,1,2,3,4); HND: (0,1,3,2,4)). No new
+  // backend allocated (NHD: stride order (0,2,1,3); HND: (0,1,2,3)). No new
   // op argument is needed -- the strides ride along with the tensor itself.
-  int64_t kv_s_block = 0, kv_s_kv = 0, kv_s_token = 0, kv_s_head = 0;
+  int64_t kv_s_block = 0, kv_s_head = 0, kv_s_token = 0, kv_s_dim = 0;
   torch::stable::Tensor const* effective_index_slot_mapping = nullptr;
   if (insert_kv) {
     STD_TORCH_CHECK(
         slot_mapping.has_value() && slot_mapping->is_cuda() &&
             slot_mapping->scalar_type() == torch::headeronly::ScalarType::Long,
         "insert mode requires int64 CUDA slot_mapping");
-    STD_TORCH_CHECK(
-        !index_slot_mapping.has_value() ||
-            (index_slot_mapping->is_cuda() &&
-             index_slot_mapping->scalar_type() ==
-                 torch::headeronly::ScalarType::Long &&
-             index_slot_mapping->numel() == slot_mapping->numel()),
-        "index_slot_mapping must be int64 CUDA with slot_mapping length");
+    if (process_index) {
+      STD_TORCH_CHECK(
+          !index_slot_mapping.has_value() ||
+              (index_slot_mapping->is_cuda() &&
+               index_slot_mapping->scalar_type() ==
+                   torch::headeronly::ScalarType::Long &&
+               index_slot_mapping->numel() == slot_mapping->numel()),
+          "index_slot_mapping must be int64 CUDA with slot_mapping length");
+    }
     // Main attention KV cache: auto matches qkv, fp8 uses uint8 storage.
     if (kv_dt == vllm::Fp8KVCacheDataType::kAuto) {
       STD_TORCH_CHECK(kv_cache->scalar_type() == qkv.scalar_type(),
@@ -701,22 +747,26 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
           "fp8 kv_cache must use uint8 storage");
     }
     // Indexer index-K cache: independent dtype -- qkv dtype or fp8 e4m3.
-    STD_TORCH_CHECK(
-        index_cache.has_value() &&
-            (index_cache->scalar_type() == qkv.scalar_type() ||
-             index_cache->scalar_type() ==
-                 torch::headeronly::ScalarType::Float8_e4m3fn),
-        "insert mode requires index_cache matching qkv dtype or fp8 e4m3");
-    STD_TORCH_CHECK(kv_cache->dim() == 5 && kv_cache->stride(4) == 1,
-                    "kv_cache must be [nb,2,bs,nkv,head_dim] with contiguous "
-                    "head_dim (stride(4)==1)");
+    if (process_index) {
+      STD_TORCH_CHECK(
+          index_cache.has_value() &&
+              (index_cache->scalar_type() == qkv.scalar_type() ||
+               index_cache->scalar_type() ==
+                   torch::headeronly::ScalarType::Float8_e4m3fn),
+          "insert mode requires index_cache matching qkv dtype or fp8 e4m3");
+    }
+    STD_TORCH_CHECK(kv_cache->dim() == 4 && kv_cache->stride(3) == 1,
+                    "kv_cache must be [nb,nkv,bs,2*head_dim] with contiguous "
+                    "content dim (stride(3)==1)");
     kv_s_block = kv_cache->stride(0);
-    kv_s_kv = kv_cache->stride(1);
+    kv_s_head = kv_cache->stride(1);
     kv_s_token = kv_cache->stride(2);
-    kv_s_head = kv_cache->stride(3);
-    effective_index_slot_mapping = index_slot_mapping.has_value()
-                                       ? &index_slot_mapping.value()
-                                       : &slot_mapping.value();
+    kv_s_dim = kv_cache->stride(3);
+    if (process_index) {
+      effective_index_slot_mapping = index_slot_mapping.has_value()
+                                         ? &index_slot_mapping.value()
+                                         : &slot_mapping.value();
+    }
   }
   // Optional contiguous gather targets: when given, the normed/roped q (and
   // index_q) are written here instead of in place, so callers avoid a separate
@@ -731,9 +781,8 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
         "q_out must have num_tokens * num_heads * 128 elements");
   }
   if (index_q_out.has_value()) {
-    STD_TORCH_CHECK(
-        has_index,
-        "index_q_out requires the index branch (num_index_heads > 0)");
+    STD_TORCH_CHECK(process_index,
+                    "index_q_out requires index branch processing");
     STD_TORCH_CHECK(
         index_q_out->is_cuda() && index_q_out->is_contiguous() &&
             (index_q_out->scalar_type() == qkv.scalar_type() ||
@@ -750,8 +799,9 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
   // q/k/v + q_out stay qkv dtype. Both index outputs must agree.
   auto const kFp8 = torch::headeronly::ScalarType::Float8_e4m3fn;
   bool const fp8_idx =
-      (index_cache.has_value() && index_cache->scalar_type() == kFp8) ||
-      (index_q_out.has_value() && index_q_out->scalar_type() == kFp8);
+      process_index &&
+      ((index_cache.has_value() && index_cache->scalar_type() == kFp8) ||
+       (index_q_out.has_value() && index_q_out->scalar_type() == kFp8));
   if (fp8_idx) {
     STD_TORCH_CHECK(
         !index_cache.has_value() || index_cache->scalar_type() == kFp8,

csrc/libtorch_stable/moe/moe_align_sum_kernels.cu

@@ -82,6 +82,21 @@ __global__ void batched_moe_align_block_size_kernel(
 }
 }  // namespace batched_moe_align_block_size
 
+template <typename scalar_t>
+__device__ __forceinline__ int get_local_expert_id(
+    size_t idx, const scalar_t* __restrict__ topk_ids,
+    int32_t* __restrict__ expert_map, int32_t num_experts,
+    bool has_expert_map) {
+  int expert_id = topk_ids[idx];
+  if (expert_id >= num_experts || expert_id < 0) {
+    return -1;
+  }
+  if (has_expert_map) {
+    expert_id = expert_map[expert_id];
+  }
+  return expert_id;
+}
+
 template <typename scalar_t>
 __device__ void _moe_align_block_size(
     const scalar_t* __restrict__ topk_ids,
@@ -126,20 +141,15 @@ __device__ void _moe_align_block_size(
   const size_t stride = blockDim.x;
 
   for (size_t i = tid; i < numel; i += stride) {
-    int expert_id = topk_ids[i];
-    if (expert_id >= num_experts) {
-      continue;
+    if (int expert_id = get_local_expert_id(i, topk_ids, expert_map,
+                                            num_experts, has_expert_map);
+        expert_id != -1) {
+      int warp_idx = expert_id / experts_per_warp;
+      int expert_offset = expert_id % experts_per_warp;
+      int mask = token_mask == nullptr ? 1 : token_mask[i / topk_num];
+      atomicAdd(&shared_counts[warp_idx * experts_per_warp + expert_offset],
+                mask);
     }
-    if (has_expert_map) {
-      expert_id = expert_map[expert_id];
-      // filter invalid experts
-      if (expert_id == -1) continue;
-    }
-    int warp_idx = expert_id / experts_per_warp;
-    int expert_offset = expert_id % experts_per_warp;
-    int mask = token_mask == nullptr ? 1 : token_mask[i / topk_num];
-    atomicAdd(&shared_counts[warp_idx * experts_per_warp + expert_offset],
-              mask);
   }
 
   __syncthreads();
@@ -227,14 +237,12 @@ __device__ void _moe_align_block_size_small_batch_expert(
   }
 
   for (size_t i = tid; i < numel; i += stride) {
-    int32_t expert_id = topk_ids[i];
-    if (has_expert_map) {
-      expert_id = expert_map[expert_id];
-      // filter invalid expert
-      if (expert_id == -1) continue;
+    if (int expert_id = get_local_expert_id(i, topk_ids, expert_map,
+                                            num_experts, has_expert_map);
+        expert_id != -1) {
+      int mask = token_mask == nullptr ? 1 : token_mask[i / topk_num];
+      tokens_cnts[(tid + 1) * num_experts + expert_id] += mask;
     }
-    int mask = token_mask == nullptr ? 1 : token_mask[i / topk_num];
-    tokens_cnts[(tid + 1) * num_experts + expert_id] += mask;
   }
 
   __syncthreads();
@@ -276,18 +284,16 @@ __device__ void _moe_align_block_size_small_batch_expert(
   }
 
   for (size_t i = tid; i < numel; i += stride) {
-    int32_t expert_id = topk_ids[i];
-    if (has_expert_map) {
-      expert_id = expert_map[expert_id];
-      // filter invalid expert
-      if (expert_id == -1) continue;
-    }
-    int32_t rank_post_pad =
-        tokens_cnts[tid * num_experts + expert_id] + cumsum[expert_id];
+    if (int expert_id = get_local_expert_id(i, topk_ids, expert_map,
+                                            num_experts, has_expert_map);
+        expert_id != -1) {
+      int32_t rank_post_pad =
+          tokens_cnts[tid * num_experts + expert_id] + cumsum[expert_id];
 
-    if (token_mask == nullptr || token_mask[i / topk_num]) {
-      sorted_token_ids[sorted_token_ids_offset + rank_post_pad] = i;
-      ++tokens_cnts[tid * num_experts + expert_id];
+      if (token_mask == nullptr || token_mask[i / topk_num]) {
+        sorted_token_ids[sorted_token_ids_offset + rank_post_pad] = i;
+        ++tokens_cnts[tid * num_experts + expert_id];
+      }
     }
   }
 }
@@ -303,22 +309,15 @@ __device__ void _count_and_sort_expert_tokens(
   const size_t stride = blockDim.x * gridDim.y;
 
   for (size_t i = tid; i < numel; i += stride) {
-    int32_t expert_id = topk_ids[i];
-    if (expert_id >= num_experts) {
-      continue;
-    }
-
-    if (has_expert_map) {
-      expert_id = expert_map[expert_id];
-      // filter invalid experts
-      if (expert_id == -1) continue;
-    }
-
-    if (token_mask == nullptr || token_mask[i / topk_num]) {
-      int32_t rank_post_pad = atomicAdd(
-          &cumsum_buffer[(model_offset * (num_experts + 1)) + expert_id], 1);
-      sorted_token_ids[max_num_tokens_padded * model_offset + rank_post_pad] =
-          i;
+    if (int expert_id = get_local_expert_id(i, topk_ids, expert_map,
+                                            num_experts, has_expert_map);
+        expert_id != -1) {
+      if (token_mask == nullptr || token_mask[i / topk_num]) {
+        int32_t rank_post_pad = atomicAdd(
+            &cumsum_buffer[(model_offset * (num_experts + 1)) + expert_id], 1);
+        sorted_token_ids[max_num_tokens_padded * model_offset + rank_post_pad] =
+            i;
+      }
     }
   }
 }

csrc/libtorch_stable/ngram_embedding_kernels.cu

@@ -0,0 +1,96 @@
+// N-gram embedding index kernel for LongCat-Flash (n-gram embedding variant).
+//
+// Adapted from SGLang:
+// https://github.com/sgl-project/sglang/blob/main/python/sglang/jit_kernel/csrc/ngram_embedding.cuh
+//
+// For each position, computes the hashed n-gram embedding ids that index the
+// concatenated embedder table. Integer tensors are int32 except ``row_indices``
+// (int64); the token table is ``[max_running_reqs, max_context_len]`` int32,
+// where a negative entry marks an ignored token (e.g. an EOS boundary).
+
+#include "torch_utils.h"
+
+#include "ops.h"
+
+#include <cstdint>
+
+namespace vllm::ngram_embedding {
+
+constexpr int kBlockThreads = 256;
+
+__global__ void ComputeNGramIdsKernel(
+    int batch_size, int ne_n, int ne_k,
+    int* ne_weights,                       // [ne_n-1, ne_k, ne_n]
+    int* ne_mods,                          // [ne_n-1, ne_k]
+    int* exclusive_ne_embedder_size_sums,  // [(ne_n-1)*ne_k + 1]
+    int* exclusive_req_len_sums,           // [batch_size + 1]
+    int* ne_token_table,  // [max_running_reqs, max_context_len]
+    int max_context_len,
+    const int64_t* __restrict__ row_indices,  // [batch_size]
+    int* column_starts,                       // [batch_size]
+    int* n_gram_ids                           // [token_num, (ne_n-1)*ne_k]
+) {
+  const int req_id = blockIdx.x % batch_size;
+  const int config_id = (blockIdx.x - req_id) / batch_size;
+  // n and k are offset from their physical meaning: n = real_n - 2, k = real_k
+  // - 1 (they index into ne_weights / ne_mods).
+  const int k = config_id % ne_k;
+  const int n = (config_id - config_id % ne_k) / ne_k;
+  const int ne_weight_base_idx = n * ne_k * ne_n + k * ne_n;
+  const int ne_mod = ne_mods[n * ne_k + k];
+  for (int i = exclusive_req_len_sums[req_id] + threadIdx.x;
+       i < exclusive_req_len_sums[req_id + 1]; i += blockDim.x) {
+    uint64_t n_gram_id = 0;
+    const int64_t current_token_offset = i - exclusive_req_len_sums[req_id];
+    const int64_t req_token_table_index =
+        row_indices[req_id] * static_cast<int64_t>(max_context_len);
+    const int64_t current_token_table_index =
+        req_token_table_index + column_starts[req_id] + current_token_offset;
+    for (int j = 0; j < n + 2; j++) {
+      if (current_token_table_index - j < req_token_table_index) {
+        break;  // outside this request's range
+      }
+      if (ne_token_table[current_token_table_index - j] < 0) {
+        break;  // ignored token
+      }
+      const uint64_t term =
+          (uint64_t)ne_token_table[current_token_table_index - j] *
+          (uint64_t)ne_weights[ne_weight_base_idx + j];
+      n_gram_id += term % ne_mod;
+    }
+    n_gram_id %= ne_mod;
+    n_gram_id += exclusive_ne_embedder_size_sums[n * ne_k + k];
+    n_gram_ids[i * (ne_n - 1) * ne_k + n * ne_k + k] = (int)(n_gram_id);
+  }
+}
+
+}  // namespace vllm::ngram_embedding
+
+void ngram_compute_n_gram_ids(
+    int64_t ne_n, int64_t ne_k, torch::stable::Tensor& ne_weights,
+    torch::stable::Tensor& ne_mods,
+    torch::stable::Tensor& exclusive_ne_embedder_size_sums,
+    torch::stable::Tensor& exclusive_req_len_sums,
+    torch::stable::Tensor& ne_token_table, torch::stable::Tensor& row_indices,
+    torch::stable::Tensor& column_starts, torch::stable::Tensor& n_gram_ids) {
+  const int batch_size = static_cast<int>(exclusive_req_len_sums.size(0) - 1);
+  const int max_context_len = static_cast<int>(ne_token_table.size(1));
+  const int num_configs = (static_cast<int>(ne_n) - 1) * static_cast<int>(ne_k);
+  const int grid_size = num_configs * batch_size;
+  if (grid_size <= 0) return;
+
+  const torch::stable::accelerator::DeviceGuard device_guard(
+      ne_weights.get_device_index());
+  const cudaStream_t stream = get_current_cuda_stream();
+  vllm::ngram_embedding::ComputeNGramIdsKernel<<<
+      grid_size, vllm::ngram_embedding::kBlockThreads, 0, stream>>>(
+      batch_size, static_cast<int>(ne_n), static_cast<int>(ne_k),
+      ne_weights.mutable_data_ptr<int32_t>(),
+      ne_mods.mutable_data_ptr<int32_t>(),
+      exclusive_ne_embedder_size_sums.mutable_data_ptr<int32_t>(),
+      exclusive_req_len_sums.mutable_data_ptr<int32_t>(),
+      ne_token_table.mutable_data_ptr<int32_t>(), max_context_len,
+      row_indices.const_data_ptr<int64_t>(),
+      column_starts.mutable_data_ptr<int32_t>(),
+      n_gram_ids.mutable_data_ptr<int32_t>());
+}

csrc/libtorch_stable/ops.h

@@ -313,7 +313,7 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
     std::optional<torch::stable::Tensor> index_cache, int64_t block_size,
     std::optional<torch::stable::Tensor> q_out,
     std::optional<torch::stable::Tensor> index_q_out,
-    const std::string& kv_cache_dtype);
+    const std::string& kv_cache_dtype, bool skip_index_branch);
 
 // Sampler kernels (shared CUDA/ROCm)
 void apply_repetition_penalties_(
@@ -414,6 +414,8 @@ void gelu_new(torch::stable::Tensor& out, torch::stable::Tensor& input);
 void gelu_fast(torch::stable::Tensor& out, torch::stable::Tensor& input);
 void gelu_quick(torch::stable::Tensor& out, torch::stable::Tensor& input);
 
+void relu_squared(torch::stable::Tensor& out, torch::stable::Tensor& input);
+
 // INT8 quantization kernels (shared CUDA/ROCm)
 void static_scaled_int8_quant(torch::stable::Tensor& out,
                               torch::stable::Tensor const& input,
@@ -554,3 +556,12 @@ void cp_gather_indexer_k_quant_cache(
                                                 // quant_block_size * 4]
     const torch::stable::Tensor& block_table,   // [batch_size, num_blocks]
     const torch::stable::Tensor& cu_seq_lens);  // [batch_size + 1]
+
+// LongCat n-gram embedding index kernel (see ngram_embedding_kernels.cu).
+void ngram_compute_n_gram_ids(
+    int64_t ne_n, int64_t ne_k, torch::stable::Tensor& ne_weights,
+    torch::stable::Tensor& ne_mods,
+    torch::stable::Tensor& exclusive_ne_embedder_size_sums,
+    torch::stable::Tensor& exclusive_req_len_sums,
+    torch::stable::Tensor& ne_token_table, torch::stable::Tensor& row_indices,
+    torch::stable::Tensor& column_starts, torch::stable::Tensor& n_gram_ids);

csrc/libtorch_stable/torch_bindings.cpp

@@ -466,7 +466,7 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
       "Tensor? slot_mapping, Tensor? index_slot_mapping, "
       "Tensor!? kv_cache, Tensor!? index_cache, "
       "int block_size, Tensor!? q_out, Tensor!? index_q_out, "
-      "str kv_cache_dtype) -> ()");
+      "str kv_cache_dtype, bool skip_index_branch=False) -> ()");
 
   // Apply repetition penalties to logits in-place.
   ops.def(
@@ -539,6 +539,9 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
   // Quick GELU implementation.
   ops.def("gelu_quick(Tensor! out, Tensor input) -> ()");
 
+  // relu(x)^2 activation from https://arxiv.org/abs/2109.08668v2
+  ops.def("relu_squared(Tensor! out, Tensor input) -> ()");
+
   // Compute int8 quantized tensor for given scaling factor.
   ops.def(
       "static_scaled_int8_quant(Tensor! result, Tensor input, Tensor scale,"
@@ -598,9 +601,22 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
       "Tensor? initial_state_idx,"
       "Tensor? cu_chunk_seqlen,"
       "Tensor? last_chunk_indices) -> ()");
+
+  // LongCat n-gram embedding index kernel. All tensor args are marked mutable
+  // to match the (non-const) stable-Tensor& C++ signature; only ne_token_table
+  // and n_gram_ids are actually written in place.
+  ops.def(
+      "ngram_compute_n_gram_ids(int ne_n, int ne_k, Tensor(a!) ne_weights, "
+      "Tensor(b!) ne_mods, Tensor(c!) exclusive_ne_embedder_size_sums, "
+      "Tensor(d!) exclusive_req_len_sums, Tensor(e!) ne_token_table, "
+      "Tensor(f!) row_indices, Tensor(g!) column_starts, "
+      "Tensor(h!) n_gram_ids) -> ()");
 }
 
 STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
+  // LongCat n-gram embedding index kernel.
+  ops.impl("ngram_compute_n_gram_ids", TORCH_BOX(&ngram_compute_n_gram_ids));
+
   // Per-token group quantization
   ops.impl("per_token_group_fp8_quant", TORCH_BOX(&per_token_group_quant_fp8));
   ops.impl("per_token_group_fp8_quant_packed",
@@ -702,6 +718,7 @@ STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
   ops.impl("gelu_new", TORCH_BOX(&gelu_new));
   ops.impl("gelu_fast", TORCH_BOX(&gelu_fast));
   ops.impl("gelu_quick", TORCH_BOX(&gelu_quick));
+  ops.impl("relu_squared", TORCH_BOX(&relu_squared));
   ops.impl("silu_and_mul_with_clamp", TORCH_BOX(&silu_and_mul_clamp));
 
   // INT8 quantization kernels

csrc/ops.h

@@ -43,6 +43,8 @@ void gelu_fast(torch::Tensor& out, torch::Tensor& input);
 
 void gelu_quick(torch::Tensor& out, torch::Tensor& input);
 
+void relu_squared(torch::Tensor& out, torch::Tensor& input);
+
 void static_scaled_int8_quant(torch::Tensor& out, torch::Tensor const& input,
                               torch::Tensor const& scale,
                               std::optional<torch::Tensor> const& azp);

docs/configuration/conserving_memory.md

@@ -42,7 +42,7 @@ and the maximum batch size (`max_num_seqs` option).
 ```python
 from vllm import LLM
 
-llm = LLM(model="adept/fuyu-8b", max_model_len=2048, max_num_seqs=2)
+llm = LLM(model="Qwen/Qwen2.5-VL-3B-Instruct", max_model_len=2048, max_num_seqs=2)
 ```
 
 ## Reduce CUDA Graphs

docs/contributing/ci/update_pytorch_version.md

@@ -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

docs/contributing/model/multimodal.md

@@ -324,154 +324,44 @@ Assuming that the memory usage increases with the number of tokens, the dummy in
         return image_token * num_images
     ```
 
-=== "No input placeholders: Fuyu"
+=== "No input placeholders: PaliGemma"
 
-    Looking at the code of HF's `FuyuForCausalLM`:
-
-    ??? code
-
-        ```python
-        # https://github.com/huggingface/transformers/blob/v4.48.3/src/transformers/models/fuyu/modeling_fuyu.py#L311-L322
-        if image_patches is not None and past_key_values is None:
-            patch_embeddings = [
-                self.vision_embed_tokens(patch.to(self.vision_embed_tokens.weight.dtype))
-                .squeeze(0)
-                .to(inputs_embeds.device)
-                for patch in image_patches
-            ]
-            inputs_embeds = self.gather_continuous_embeddings(
-                word_embeddings=inputs_embeds,
-                continuous_embeddings=patch_embeddings,
-                image_patch_input_indices=image_patches_indices,
-            )
-        ```
-
-    The number of placeholder feature tokens for the `i`th item in the batch is `patch_embeddings[i].shape[0]`,
-    which is the same as `image_patches[i].shape[0]`, i.e. `num_total_patches`.
-
-    Unlike LLaVA, Fuyu does not define the number of patches inside the modeling file. Where can we get more information?
-    Considering that the model input comes from the output of `FuyuProcessor`, let's **look at the preprocessing files**.
-
-    The image outputs are obtained by calling `FuyuImageProcessor.preprocess` and then
-    `FuyuImageProcessor.preprocess_with_tokenizer_info` inside `FuyuProcessor`.
-
-    In `FuyuImageProcessor.preprocess`, the images are resized and padded to the target `FuyuImageProcessor.size`,
-    returning the dimensions after resizing (but before padding) as metadata.
-
-    ??? code
-
-        ```python
-        # https://github.com/huggingface/transformers/blob/v4.48.3/src/transformers/models/fuyu/processing_fuyu.py#L541-L544
-        image_encoding = self.image_processor.preprocess(images, **output_kwargs["images_kwargs"])
-        batch_images = image_encoding["images"]
-        image_unpadded_heights = image_encoding["image_unpadded_heights"]
-        image_unpadded_widths = image_encoding["image_unpadded_widths"]
-
-        # https://github.com/huggingface/transformers/blob/v4.48.3/src/transformers/models/fuyu/image_processing_fuyu.py#L480-L
-        if do_resize:
-            batch_images = [
-                [self.resize(image, size=size, input_data_format=input_data_format) for image in images]
-                for images in batch_images
-            ]
-
-        image_sizes = [get_image_size(images[0], channel_dim=input_data_format) for images in batch_images]
-        image_unpadded_heights = [[image_size[0]] for image_size in image_sizes]
-        image_unpadded_widths = [[image_size[1]] for image_size in image_sizes]
-
-        if do_pad:
-            batch_images = [
-                [
-                    self.pad_image(
-                        image,
-                        size=size,
-                        mode=padding_mode,
-                        constant_values=padding_value,
-                        input_data_format=input_data_format,
-                    )
-                    for image in images
-                ]
-                for images in batch_images
-            ]
-        ```
-
-    In `FuyuImageProcessor.preprocess_with_tokenizer_info`, the images are split into patches based on this metadata:
-
-    ??? code
-
-        ```python
-        # https://github.com/huggingface/transformers/blob/v4.48.3/src/transformers/models/fuyu/processing_fuyu.py#L417-L425
-        model_image_input = self.image_processor.preprocess_with_tokenizer_info(
-            image_input=tensor_batch_images,
-            image_present=image_present,
-            image_unpadded_h=image_unpadded_heights,
-            image_unpadded_w=image_unpadded_widths,
-            image_placeholder_id=image_placeholder_id,
-            image_newline_id=image_newline_id,
-            variable_sized=True,
-        )
-
-        # https://github.com/huggingface/transformers/blob/v4.48.3/src/transformers/models/fuyu/image_processing_fuyu.py#L638-L658
-        image_height, image_width = image.shape[1], image.shape[2]
-        if variable_sized:  # variable_sized=True
-            new_h = min(
-                image_height,
-                math.ceil(image_unpadded_h[batch_index, subseq_index] / patch_height) * patch_height,
-            )
-            new_w = min(
-                image_width,
-                math.ceil(image_unpadded_w[batch_index, subseq_index] / patch_width) * patch_width,
-            )
-            image = image[:, :new_h, :new_w]
-            image_height, image_width = new_h, new_w
-
-        num_patches = self.get_num_patches(image_height=image_height, image_width=image_width)
-        tensor_of_image_ids = torch.full(
-            [num_patches], image_placeholder_id, dtype=torch.int32, device=image_input.device
-        )
-        patches = self.patchify_image(image=image.unsqueeze(0)).squeeze(0)
-        assert num_patches == patches.shape[0]
-        ```
-
-    The number of patches is in turn defined by `FuyuImageProcessor.get_num_patches`:
-
-    ??? code
-
-        ```python
-        # https://github.com/huggingface/transformers/blob/v4.48.3/src/transformers/models/fuyu/image_processing_fuyu.py#L552-L562
-        patch_size = patch_size if patch_size is not None else self.patch_size
-        patch_height, patch_width = self.patch_size["height"], self.patch_size["width"]
-
-        if image_height % patch_height != 0:
-            raise ValueError(f"{image_height=} must be divisible by {patch_height}")
-        if image_width % patch_width != 0:
-            raise ValueError(f"{image_width=} must be divisible by {patch_width}")
-
-        num_patches_per_dim_h = image_height // patch_height
-        num_patches_per_dim_w = image_width // patch_width
-        num_patches = num_patches_per_dim_h * num_patches_per_dim_w
-        ```
-
-    These image patches correspond to placeholder tokens (`|SPEAKER|`). So, we just need to maximize the number of image patches. Since input images are first resized
-    to fit within `image_processor.size`, we can maximize the number of image patches by inputting an image with size equal to `image_processor.size`.
-
-    ```python
-    def get_image_size_with_most_features(self) -> ImageSize:
-        image_processor = self.get_image_processor()
-        return ImageSize(
-            width=image_processor.size["width"],
-            height=image_processor.size["height"],
-        )
-    ```
-
-    Fuyu does not expect image placeholders in the inputs to HF processor, so
-    the dummy prompt text is empty regardless of the number of images.
+    Unlike LLaVA, PaliGemma's HF processor does not expect image placeholder
+    tokens in the input prompt; the placeholder feature tokens are instead
+    inserted afterwards (see [Prompt updates](#prompt-updates)). So the dummy
+    prompt text is empty regardless of the number of images:
 
     ```python
     def get_dummy_text(self, mm_counts: Mapping[str, int]) -> str:
         return ""
     ```
 
-    For the multimodal image profiling data, the logic is very similar to LLaVA:
+    PaliGemma resizes every image to a square of `vision_config.image_size`, so
+    the number of placeholder feature tokens per image is fixed at
+    `(image_size // patch_size) ** 2`. This is computed by the SigLIP vision
+    encoder that PaliGemma uses:
+
+    ??? code
+
+        ```python
+        # vllm/model_executor/models/siglip.py
+        class SiglipEncoderInfo(VisionEncoderInfo[SiglipVisionConfig]):
+            def get_num_image_tokens(
+                self,
+                *,
+                image_width: int,
+                image_height: int,
+            ) -> int:
+                return self.get_patch_grid_length() ** 2
+
+            def get_patch_grid_length(self) -> int:
+                image_size, patch_size = self.get_image_size(), self.get_patch_size()
+                return image_size // patch_size
+        ```
+
+    Since the number of image tokens doesn't depend on the input image dimensions,
+    we can simply use a dummy image of the model's expected input size for the
+    multimodal profiling data:
 
     ??? code
 
@@ -482,16 +372,18 @@ Assuming that the memory usage increases with the number of tokens, the dummy in
             mm_counts: Mapping[str, int],
             mm_options: Mapping[str, BaseDummyOptions],
         ) -> MultiModalDataDict:
-            target_width, target_height = \
-                self.info.get_image_size_with_most_features()
+            hf_config = self.info.get_hf_config()
+            vision_config = hf_config.vision_config
+            max_image_size = vision_config.image_size
+
             num_images = mm_counts.get("image", 0)
 
             image_overrides = mm_options.get("image")
 
             return {
                 "image": self._get_dummy_images(
-                    width=target_width,
-                    height=target_height,
+                    width=max_image_size,
+                    height=max_image_size,
                     num_images=num_images,
                     overrides=image_overrides,
                 )
@@ -545,28 +437,15 @@ return a schema of the tensors outputted by the HF processor that are related to
         Our [actual code](../../../vllm/model_executor/models/llava.py) additionally supports
         pre-computed image embeddings, which can be passed to be model via the `image_embeds` argument.
 
-=== "With postprocessing: Fuyu"
+=== "With postprocessing: Mistral3"
 
-    The `image_patches` output of `FuyuImageProcessor.preprocess_with_tokenizer_info` concatenates
-    the patches from each image belonging to an item in the batch:
+    The `pixel_values` output of Mistral3's HF processor pads every image in the
+    batch to a common size, so that they can be stacked into a single tensor.
 
-    ```python
-    # https://github.com/huggingface/transformers/blob/v4.48.3/src/transformers/models/fuyu/image_processing_fuyu.py#L673-L679
-            image_input_ids.append(tensor_of_image_ids)
-            image_patches.append(patches)
-        else:
-            image_input_ids.append(torch.tensor([], dtype=torch.int32, device=image_input.device))
-
-    batch_image_input_ids.append(image_input_ids)
-    batch_image_patches.append(image_patches)
-    ```
-
-    The shape of `image_patches` outputted by `FuyuImageProcessor` is therefore
-    `(1, num_images, num_patches, patch_width * patch_height * num_channels)`.
-
-    In order to support the use of
-    [MultiModalFieldConfig.batched][vllm.multimodal.inputs.MultiModalFieldConfig.batched]
-    like in LLaVA, we remove the extra batch dimension by overriding
+    To use [MultiModalFieldConfig.batched][vllm.multimodal.inputs.MultiModalFieldConfig.batched]
+    like in LLaVA, each image's features must be independent of the others (which
+    is also required for prefix caching to work correctly). So, we un-pad each image
+    back to its own size by overriding
     [BaseMultiModalProcessor._call_hf_processor][vllm.multimodal.processing.BaseMultiModalProcessor._call_hf_processor]:
 
     ??? code
@@ -586,33 +465,27 @@ return a schema of the tensors outputted by the HF processor that are related to
                 tok_kwargs=tok_kwargs,
             )
 
-            image_patches = processed_outputs.get("image_patches")
-            if image_patches is not None:
-                images = mm_data["images"]
-                assert isinstance(images, list)
+            pixel_values = processed_outputs.get("pixel_values")
+            if pixel_values is not None:
+                # Avoid padding since we need the output for each image to be
+                # independent of other images for the cache to work correctly
+                image_sizes = processed_outputs["image_sizes"]
+                assert len(pixel_values) == len(image_sizes)
 
-                # Original output: (1, num_images, Pn, Px * Py * C)
-                # New output: (num_images, Pn, Px * Py * C)
-                assert (isinstance(image_patches, list)
-                        and len(image_patches) == 1)
-                assert (isinstance(image_patches[0], torch.Tensor)
-                        and len(image_patches[0]) == len(images))
-
-                processed_outputs["image_patches"] = image_patches[0]
+                processed_outputs["pixel_values"] = [
+                    p[:, :h, :w] for p, (h, w) in zip(pixel_values, image_sizes)
+                ]
 
             return processed_outputs
         ```
 
-    !!! note
-        Our [actual code](../../../vllm/model_executor/models/fuyu.py) has special handling
-        for text-only inputs to prevent unnecessary warnings from HF processor.
-
     !!! note
         The `_call_hf_processor` method specifies both `mm_kwargs` and `tok_kwargs` for
         processing. `mm_kwargs` is used to both initialize and call the huggingface
         processor, whereas `tok_kwargs` is only used to call the huggingface processor.
 
-    This lets us override [_get_mm_fields_config][vllm.multimodal.processing.BaseMultiModalProcessor._get_mm_fields_config] as follows:
+    Since `pixel_values` is now a list with one tensor per image, we can override
+    [_get_mm_fields_config][vllm.multimodal.processing.BaseMultiModalProcessor._get_mm_fields_config] as follows:
 
     ```python
     def _get_mm_fields_config(
@@ -620,9 +493,15 @@ return a schema of the tensors outputted by the HF processor that are related to
         hf_inputs: BatchFeature,
         hf_processor_mm_kwargs: Mapping[str, object],
     ) -> Mapping[str, MultiModalFieldConfig]:
-        return dict(image_patches=MultiModalFieldConfig.batched("image"))
+        return dict(
+            pixel_values=MultiModalFieldConfig.batched("image"),
+            image_embeds=MultiModalFieldConfig.batched("image"),
+        )
     ```
 
+    !!! note
+        See our [actual code](../../../vllm/model_executor/models/mistral3.py) for the full implementation.
+
 ### Prompt updates
 
 Override [_get_prompt_updates][vllm.multimodal.processing.BaseMultiModalProcessor._get_prompt_updates] to
@@ -678,121 +557,53 @@ Each [PromptUpdate][vllm.multimodal.processing.PromptUpdate] instance specifies
             ]
         ```
 
-=== "Handling additional tokens: Fuyu"
+=== "Handling additional tokens: PaliGemma"
 
-    Recall the layout of feature tokens from Step 2:
-
-    ```
-    |SPEAKER||SPEAKER|...|SPEAKER||NEWLINE|
-    |SPEAKER||SPEAKER|...|SPEAKER||NEWLINE|
-    ...
-    |SPEAKER||SPEAKER|...|SPEAKER||NEWLINE|
-    ```
-
-    We define a helper function to return `ncols` and `nrows` directly:
+    PaliGemma's HF processor inserts, after the prompt's leading `<bos>` token, a
+    run of image tokens followed by a second `<bos>` token that marks the start of
+    the text prompt. We start by building the run of image tokens, one per
+    placeholder feature token:
 
     ??? code
 
         ```python
-        def get_image_feature_grid_size(
-            self,
-            *,
-            image_width: int,
-            image_height: int,
-        ) -> tuple[int, int]:
-            image_processor = self.get_image_processor()
-            target_width = image_processor.size["width"]
-            target_height = image_processor.size["height"]
-            patch_width = image_processor.patch_size["width"]
-            patch_height = image_processor.patch_size["height"]
-
-            if not (image_width <= target_width and image_height <= target_height):
-                height_scale_factor = target_height / image_height
-                width_scale_factor = target_width / image_width
-                optimal_scale_factor = min(height_scale_factor, width_scale_factor)
-
-                image_height = int(image_height * optimal_scale_factor)
-                image_width = int(image_width * optimal_scale_factor)
-
-            ncols = math.ceil(image_width / patch_width)
-            nrows = math.ceil(image_height / patch_height)
-            return ncols, nrows
-        ```
-
-    Based on this, we can initially define our replacement tokens as:
-
-    ??? code
-
-        ```python
-        def get_replacement(item_idx: int):
-            images = mm_items.get_items("image", ImageProcessorItems)
-            image_size = images.get_image_size(item_idx)
-
-            ncols, nrows = self.info.get_image_feature_grid_size(
-                image_width=image_size.width,
-                image_height=image_size.height,
+        def get_insertion(item_idx: int):
+            images = mm_items.get_items(
+                "image", (ImageEmbeddingItems, ImageProcessorItems)
             )
 
-            # `_IMAGE_TOKEN_ID` corresponds to `|SPEAKER|`
-            # `_NEWLINE_TOKEN_ID` corresponds to `|NEWLINE|`
-            return ([_IMAGE_TOKEN_ID] * ncols + [_NEWLINE_TOKEN_ID]) * nrows
+            if isinstance(images, ImageEmbeddingItems):
+                num_image_tokens = images.get_feature_size(item_idx)
+            else:
+                image_size = images.get_image_size(item_idx)
+                num_image_tokens = self.info.get_num_image_tokens(
+                    image_width=image_size.width,
+                    image_height=image_size.height,
+                )
+
+            image_tokens = [image_token_id] * num_image_tokens
+            ...
         ```
 
-    However, this is not entirely correct. After `FuyuImageProcessor.preprocess_with_tokenizer_info` is called,
-    a BOS token (`<s>`) is also added to the prompt:
+    The trailing `<bos>` token is an additional token that must **not** receive a
+    vision embedding. To assign the vision embeddings to only the image tokens,
+    instead of returning the token ids directly you can return an instance of
+    [PromptUpdateDetails][vllm.multimodal.processing.PromptUpdateDetails] and mark
+    the embedding tokens with `embed_token_id`:
 
     ??? code
 
         ```python
-        # https://github.com/huggingface/transformers/blob/v4.48.3/src/transformers/models/fuyu/processing_fuyu.py#L417-L435
-        model_image_input = self.image_processor.preprocess_with_tokenizer_info(
-            image_input=tensor_batch_images,
-            image_present=image_present,
-            image_unpadded_h=image_unpadded_heights,
-            image_unpadded_w=image_unpadded_widths,
-            image_placeholder_id=image_placeholder_id,
-            image_newline_id=image_newline_id,
-            variable_sized=True,
-        )
-        prompt_tokens, prompts_length = _tokenize_prompts_with_image_and_batch(
-            tokenizer=self.tokenizer,
-            prompts=prompts,
-            scale_factors=scale_factors,
-            max_tokens_to_generate=self.max_tokens_to_generate,
-            max_position_embeddings=self.max_position_embeddings,
-            add_BOS=True,
-            add_beginning_of_answer_token=True,
+        return PromptUpdateDetails.select_token_id(
+            image_tokens + [bos_token_id],
+            embed_token_id=image_token_id,
         )
         ```
 
-    To assign the vision embeddings to only the image tokens, instead of a string
-    you can return an instance of [PromptUpdateDetails][vllm.multimodal.processing.PromptUpdateDetails]:
-
-    ??? code
-
-        ```python
-        hf_config = self.info.get_hf_config()
-        bos_token_id = hf_config.bos_token_id  # `<s>`
-        assert isinstance(bos_token_id, int)
-
-        def get_replacement_fuyu(item_idx: int):
-            images = mm_items.get_items("image", ImageProcessorItems)
-            image_size = images.get_image_size(item_idx)
-
-            ncols, nrows = self.info.get_image_feature_grid_size(
-                image_width=image_size.width,
-                image_height=image_size.height,
-            )
-            image_tokens = ([_IMAGE_TOKEN_ID] * ncols + [_NEWLINE_TOKEN_ID]) * nrows
-
-            return PromptUpdateDetails.select_token_id(
-                image_tokens + [bos_token_id],
-                embed_token_id=_IMAGE_TOKEN_ID,
-            )
-        ```
-
-    Finally, noticing that the HF processor removes the `|ENDOFTEXT|` token from the tokenized prompt,
-    we can search for it to conduct the replacement at the start of the string:
+    Putting it together, we override [_get_prompt_updates][vllm.multimodal.processing.BaseMultiModalProcessor._get_prompt_updates].
+    Since these tokens are inserted (rather than replacing an existing placeholder)
+    after the prompt's leading `<bos>`, we use [PromptInsertion][vllm.multimodal.processing.PromptInsertion]
+    with a prefix target:
 
     ??? code
 
@@ -804,33 +615,41 @@ Each [PromptUpdate][vllm.multimodal.processing.PromptUpdate] instance specifies
             out_mm_kwargs: MultiModalKwargsItems,
         ) -> Sequence[PromptUpdate]:
             hf_config = self.info.get_hf_config()
-            bos_token_id = hf_config.bos_token_id
-            assert isinstance(bos_token_id, int)
+            image_token_id = hf_config.image_token_index
 
             tokenizer = self.info.get_tokenizer()
-            eot_token_id = tokenizer.bos_token_id
-            assert isinstance(eot_token_id, int)
 
-            def get_replacement_fuyu(item_idx: int):
-                images = mm_items.get_items("image", ImageProcessorItems)
-                image_size = images.get_image_size(item_idx)
+            bos_token_id = tokenizer.bos_token_id
+            assert isinstance(bos_token_id, int)
 
-                ncols, nrows = self.info.get_image_feature_grid_size(
-                    image_width=image_size.width,
-                    image_height=image_size.height,
+            def get_insertion(item_idx: int):
+                images = mm_items.get_items(
+                    "image", (ImageEmbeddingItems, ImageProcessorItems)
                 )
-                image_tokens = ([_IMAGE_TOKEN_ID] * ncols + [_NEWLINE_TOKEN_ID]) * nrows
+
+                if isinstance(images, ImageEmbeddingItems):
+                    num_image_tokens = images.get_feature_size(item_idx)
+                else:
+                    image_size = images.get_image_size(item_idx)
+                    num_image_tokens = self.info.get_num_image_tokens(
+                        image_width=image_size.width,
+                        image_height=image_size.height,
+                    )
+
+                image_tokens = [image_token_id] * num_image_tokens
 
                 return PromptUpdateDetails.select_token_id(
                     image_tokens + [bos_token_id],
-                    embed_token_id=_IMAGE_TOKEN_ID,
+                    embed_token_id=image_token_id,
                 )
 
             return [
-                PromptReplacement(
+                PromptInsertion(
                     modality="image",
-                    target=[eot_token_id],
-                    replacement=get_replacement_fuyu,
+                    target=PromptIndexTargets.prefix(
+                        [bos_token_id] if tokenizer.add_bos_token else []
+                    ),
+                    insertion=get_insertion,
                 )
             ]
         ```
@@ -873,7 +692,7 @@ Examples:
 Examples:
 
 - Chameleon (appends `sep_token`): [vllm/model_executor/models/chameleon.py](../../../vllm/model_executor/models/chameleon.py)
-- Fuyu (appends `boa_token`): [vllm/model_executor/models/fuyu.py](../../../vllm/model_executor/models/fuyu.py)
+- Molmo2 (prepends `bos_token`): [vllm/model_executor/models/molmo2.py](../../../vllm/model_executor/models/molmo2.py)
 - Molmo (applies chat template which is not defined elsewhere): [vllm/model_executor/models/molmo.py](../../../vllm/model_executor/models/molmo.py)
 
 ### Custom HF processor

docs/features/kv_offloading_usage.md

@@ -81,6 +81,8 @@ vllm serve <model> \
 
 Each entry in `secondary_tiers` is a dict with a required `type` field plus tier-specific fields.
 
+The filesystem and object-store tiers can publish hash-only `BlockStored` KV events for blocks they successfully store, tagged with a stable per-tier `medium` (`FS` for the filesystem tier, `OBJ` for the object-store tier). Set `enable_kv_events: true` in the tier's entry to opt in; events are published only when KV cache events are also enabled globally via `--kv-events-config`.
+
 ### Filesystem (FS)
 
 The filesystem tier (`type: "fs"`) writes blocks to a directory on local storage.
@@ -91,6 +93,7 @@ The filesystem tier (`type: "fs"`) writes blocks to a directory on local storage
 | `root_dir` | yes | — | Base directory; vLLM creates subdirectories beneath it (see [On-Disk Layout](#on-disk-layout)). |
 | `n_read_threads` | no | `16` | Read-priority I/O threads (load path). |
 | `n_write_threads` | no | `16` | Write-priority I/O threads (store path). |
+| `enable_kv_events` | no | `false` | Publish `BlockStored` KV events (medium `FS`) for successfully stored blocks. Requires KV cache events to be enabled globally. |
 
 Each thread group prefers its own queue but pulls from the other when its primary queue is empty, so a write-heavy or read-heavy burst won't leave the off-priority queue waiting. Size the totals to your storage's effective concurrency.
 
@@ -120,6 +123,31 @@ To enable KV cache sharing between multiple vLLM instances using the same `root_
 PYTHONHASHSEED=0 vllm serve ...
 ```
 
+### Object Store (OBJ)
+
+The object-store tier (`type: "obj"`) offloads blocks to an S3-compatible object store through the NIXL OBJ backend.
+
+| Key | Required | Default | Notes |
+| --- | --- | --- | --- |
+| `type` | yes | — | Must be `obj`. |
+| `store_config` | yes | — | Object store connection parameters (see below). |
+| `prefix` | no | `""` | Key prefix prepended to all object keys. |
+| `io_threads` | no | `4` | Number of NIXL OBJ backend I/O threads. |
+| `enable_kv_events` | no | `false` | Publish `BlockStored` KV events (medium `OBJ`) for successfully stored blocks. Requires KV cache events to be enabled globally. |
+
+`store_config` fields:
+
+| Key | Required | Default | Notes |
+| --- | --- | --- | --- |
+| `bucket` | yes | — | Bucket name. |
+| `endpoint_override` | yes | — | Object store endpoint host; the URL scheme is set separately via `scheme`. |
+| `scheme` | no | `http` | `http` or `https`. |
+| `access_key`, `secret_key`, `session_token` | no | `""` | Explicit credentials. When left empty, the NIXL OBJ plugin falls back to the AWS SDK default credential provider chain (IAM roles, environment variables, credential files), which enables workload-identity auth on Kubernetes. |
+| `region` | no | `""` | Bucket region, if the endpoint requires one. |
+| `ca_bundle` | no | `""` | CA bundle path for TLS verification. |
+
+Object keys follow the same run-configuration digest scheme as the filesystem tier (see [On-Disk Layout](#on-disk-layout)) and are stored under the optional `prefix`. The [Cross-Process Sharing](#cross-process-sharing) requirement (`PYTHONHASHSEED`) applies to shared buckets as well, so instances sharing a bucket produce identical keys for identical content. At startup the tier probes object store connectivity and fails fast with a configuration error if the bucket is unreachable.
+
 ### P2P (Including P/D)
 
 The P2P tier (`type: "p2p"`) shares completed KV blocks between vLLM instances over RDMA via NIXL. Each instance binds a control socket on `host:port` and exchanges blocks directly with peers — no shared filesystem required.

docs/features/multimodal_inputs.md

@@ -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.

docs/features/reasoning_outputs.md

@@ -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`).

docs/models/supported_models.md

@@ -421,8 +421,6 @@ th {
 | `MPTForCausalLM` | MPT, MPT-Instruct, MPT-Chat, MPT-StoryWriter | `mosaicml/mpt-7b`, `mosaicml/mpt-7b-storywriter`, `mosaicml/mpt-30b`, etc. | | ✅︎ |
 | `NemotronForCausalLM` | Nemotron-3, Nemotron-4, Minitron | `nvidia/Minitron-8B-Base`, `mgoin/Nemotron-4-340B-Base-hf-FP8`, etc. | ✅︎ | ✅︎ |
 | `NemotronHForCausalLM` | Nemotron-H | `nvidia/Nemotron-H-8B-Base-8K`, `nvidia/Nemotron-H-47B-Base-8K`, `nvidia/Nemotron-H-56B-Base-8K`, etc. | ✅︎ | ✅︎ |
-| `OlmoForCausalLM` | OLMo | `allenai/OLMo-1B-hf`, `allenai/OLMo-7B-hf`, etc. | ✅︎ | ✅︎ |
-| `Olmo2ForCausalLM` | OLMo2 | `allenai/OLMo-2-0425-1B`, etc. | ✅︎ | ✅︎ |
 | `Olmo3ForCausalLM` | OLMo3 | `allenai/Olmo-3-7B-Instruct`, `allenai/Olmo-3-32B-Think`, etc. | ✅︎ | ✅︎ |
 | `OlmoHybridForCausalLM` | OLMo Hybrid | `allenai/Olmo-Hybrid-7B` | ✅︎ | ✅︎ |
 | `OlmoeForCausalLM` | OLMoE | `allenai/OLMoE-1B-7B-0924`, `allenai/OLMoE-1B-7B-0924-Instruct`, etc. | | ✅︎ |
@@ -436,7 +434,6 @@ th {
 | `PhiForCausalLM` | Phi | `microsoft/phi-1_5`, `microsoft/phi-2`, etc. | ✅︎ | ✅︎ |
 | `Phi3ForCausalLM` | Phi-4, Phi-3 | `microsoft/Phi-4-mini-instruct`, `microsoft/Phi-4`, `microsoft/Phi-3-mini-4k-instruct`, `microsoft/Phi-3-mini-128k-instruct`, `microsoft/Phi-3-medium-128k-instruct`, etc. | ✅︎ | ✅︎ |
 | `PhiMoEForCausalLM` | Phi-3.5-MoE | `microsoft/Phi-3.5-MoE-instruct`, etc. | ✅︎ | ✅︎ |
-| `PersimmonForCausalLM` | Persimmon | `adept/persimmon-8b-base`, `adept/persimmon-8b-chat`, etc. | | ✅︎ |
 | `Plamo2ForCausalLM` | PLaMo2 | `pfnet/plamo-2-1b`, `pfnet/plamo-2-8b`, etc. | ✅ | ✅︎ |
 | `Plamo3ForCausalLM` | PLaMo3 | `pfnet/plamo-3-nict-2b-base`, `pfnet/plamo-3-nict-8b-base`, etc. | ✅ | ✅︎ |
 | `Qwen2ForCausalLM` | QwQ, Qwen2 | `Qwen/QwQ-32B-Preview`, `Qwen/Qwen2-7B-Instruct`, `Qwen/Qwen2-7B`, etc. | ✅︎ | ✅︎ |
@@ -454,7 +451,6 @@ th {
 | `StableLMEpochForCausalLM` | StableLM Epoch | `stabilityai/stablelm-zephyr-3b`, etc. | | ✅︎ |
 | `Step1ForCausalLM` | Step-Audio | `stepfun-ai/Step-Audio-EditX`, etc. | ✅︎ | ✅︎ |
 | `Step3p5ForCausalLM` | Step-3.5-flash | `stepfun-ai/Step-3.5-Flash`, etc. | | ✅︎ |
-| `TeleChatForCausalLM` | TeleChat | `chuhac/TeleChat2-35B`, etc. | ✅︎ | ✅︎ |
 | `TeleChat2ForCausalLM` | TeleChat2 | `Tele-AI/TeleChat2-3B`, `Tele-AI/TeleChat2-7B`, `Tele-AI/TeleChat2-35B`, etc. | ✅︎ | ✅︎ |
 | `TeleChat3ForCausalLM` | TeleChat3 | `Tele-AI/TeleChat3-36B-Thinking`, `Tele-AI/TeleChat3-Coder-36B-Thinking`, etc. | ✅︎ | ✅︎ |
 | `TeleFLMForCausalLM` | TeleFLM | `CofeAI/FLM-2-52B-Instruct-2407`, `CofeAI/Tele-FLM`, etc. | ✅︎ | ✅︎ |
@@ -465,6 +461,8 @@ Some models are supported only via the [Transformers modeling backend](#transfor
 | Architecture | Models | Example HF Models | [LoRA](../features/lora.md) | [PP](../serving/parallelism_scaling.md) |
 | ------------ | ------ | ----------------- | -------------------- | ------------------------- |
 | `GPTBigCodeForCausalLM` | StarCoder, SantaCoder, WizardCoder | `bigcode/starcoder`, `bigcode/gpt_bigcode-santacoder`, `WizardLM/WizardCoder-15B-V1.0`, etc. | ✅︎ | |
+| `OlmoForCausalLM` | OLMo | `allenai/OLMo-1B-hf`, `allenai/OLMo-7B-hf`, etc. | ✅︎ | ✅︎ |
+| `Olmo2ForCausalLM` | OLMo2 | `allenai/OLMo-2-0425-1B`, etc. | ✅︎ | ✅︎ |
 | `SmolLM3ForCausalLM` | SmolLM3 | `HuggingFaceTB/SmolLM3-3B` | ✅︎ | ✅︎ |
 | `Starcoder2ForCausalLM` | Starcoder2 | `bigcode/starcoder2-3b`, `bigcode/starcoder2-7b`, `bigcode/starcoder2-15b`, etc. | ✅︎ | ✅︎ |
 
@@ -514,14 +512,13 @@ 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. | ✅︎ | ✅︎ |
 | `Eagle2_5_VLForConditionalGeneration` | Eagle2.5-VL | T + I<sup>E+</sup> | `nvidia/Eagle2.5-8B`, etc. | ✅︎ | ✅︎ |
 | `Ernie4_5_VLMoeForConditionalGeneration` | Ernie4.5-VL | T + I<sup>+</sup>/ V<sup>+</sup> | `baidu/ERNIE-4.5-VL-28B-A3B-PT`, `baidu/ERNIE-4.5-VL-424B-A47B-PT` | | ✅︎ |
 | `Exaone4_5_ForConditionalGeneration` | EXAONE-4.5 | T + I<sup>E+</sup> | `LGAI-EXAONE/EXAONE-4.5-33B`, etc. | ✅︎ | ✅︎ |
-| `FuyuForCausalLM` | Fuyu | T + I | `adept/fuyu-8b`, etc. | | ✅︎ |
 | `Gemma3ForConditionalGeneration` | Gemma 3 | T + I<sup>E+</sup> | `google/gemma-3-4b-it`, `google/gemma-3-27b-it`, etc. | ✅︎ | ✅︎ |
 | `Gemma3nForConditionalGeneration` | Gemma 3n | T + I + A | `google/gemma-3n-E2B-it`, `google/gemma-3n-E4B-it`, etc. | | |
 | `Gemma4ForConditionalGeneration` | Gemma 4 | T + I<sup>+</sup> + V + A<sup>*</sup> | `google/gemma-4-E2B-it`, etc. | | ✅︎ |

docs/serving/online_serving/README.md

@@ -170,6 +170,7 @@ For further details on Weight Transfer, please refer to [this page](../../traini
 - `/pause` - Pause generation (causes denial of service)
 - `/resume` - Resume generation
 - `/is_paused` - Check if generation is paused
+- `/abort_requests` - Abort in-flight requests (all in-flight, or the given `request_ids`) without pausing the scheduler
 - `/init_weight_transfer_engine` - Initialize weight transfer engine for RLHF
 - `/start_weight_update` - Prepares the inference engine for a weight update.
 - `/update_weights` - Update model weights (can alter model behavior)

docs/training/async_rl.md

@@ -42,6 +42,7 @@ When using the vLLM HTTP server, the same functionality is available via:
 
 - `POST /pause?mode=keep` - Pause generation
 - `POST /resume` - Resume generation
+- `POST /abort_requests` - Abort in-flight requests without pausing the scheduler (send `{}` to abort all, or `{"request_ids": [...]}`)
 
 !!! note "Data Parallelism"
     When using data parallelism with vLLM's **internal load balancer** (i.e. `data_parallel_backend="ray"`), pause and resume are handled automatically across all DP ranks -- a single call is sufficient. When using an **external load balancer** (i.e. multiple independent vLLM instances behind a proxy), you must send pause and resume requests to **every** engine instance individually before and after the weight update.

docs/training/weight_transfer/base.md

@@ -49,6 +49,10 @@ update_request = WeightTransferUpdateRequest(
 )
 ```
 
+At the LLM/API layer, call `start_draft_weight_update()` instead of
+`start_weight_update()` to target the speculative draft model;
+`update_weights` / `finish_weight_update` are unchanged.
+
 ### WeightTransferUpdateInfo
 
 The base `WeightTransferUpdateInfo` is a marker class for backend-specific update info:

docs/usage/security.md

@@ -191,6 +191,7 @@ The following endpoints **do not require authentication** even when `--api-key`
 - `/pause` - Pause generation (causes denial of service)
 - `/resume` - Resume generation
 - `/is_paused` - Check if generation is paused
+- `/abort_requests` - Abort in-flight requests (causes loss of in-flight work)
 - `/scale_elastic_ep` - Trigger scaling operations
 - `/is_scaling_elastic_ep` - Check if scaling is in progress
 - `/init_weight_transfer_engine` - Initialize weight transfer engine for RLHF

examples/disaggregated/disaggregated_serving/disagg_proxy_pushconnector_demo.py

@@ -9,9 +9,9 @@ the same; the difference is in how P and D coordinate the KV transfer:
 * Pull mode: proxy forwards P's ``kv_transfer_params`` (including
   ``remote_block_ids``) to D, and D pulls KV from P via NIXL READ.
 * Push mode: proxy hands D **only** P's coordinates
-  (``remote_engine_id``, ``remote_host``, ``remote_port``, ``tp_size``)
-  and the shared ``remote_request_id``. D registers its locally allocated
-  blocks with P over a NIXL notification; P then pushes the KV to D via
+  (``remote_engine_id``, ``remote_host``, ``remote_port``, ``tp_size``,
+  ``pp_size``) and the shared ``remote_request_id``. D registers its locally
+  allocated blocks with P over a NIXL notification; P then pushes the KV to D via
   NIXL WRITE.
 
 Launch multiple vLLM instances configured with ``NixlPushConnector`` and
@@ -26,6 +26,7 @@ disagg_proxy_pushconnector_demo.py \
        --prefill-kv-host  10.0.0.1 \
        --prefill-side-channel-port 5600 \
        --prefill-tp-size 1 \
+       --prefill-pp-size 1 \
        --port 8000
 """
 
@@ -89,6 +90,7 @@ class PushProxy:
         prefill_kv_host: str,
         prefill_side_channel_port: int,
         prefill_tp_size: int,
+        prefill_pp_size: int,
         custom_create_completion: Callable[[Request], StreamingResponse] | None = None,
         custom_create_chat_completion: Callable[[Request], StreamingResponse]
         | None = None,
@@ -110,6 +112,7 @@ class PushProxy:
             "remote_host": prefill_kv_host,
             "remote_port": prefill_side_channel_port,
             "tp_size": prefill_tp_size,
+            "pp_size": prefill_pp_size,
         }
 
         self.custom_create_completion = custom_create_completion
@@ -198,6 +201,7 @@ class PushProxy:
             "prefill_kv_host": self.push_metadata["remote_host"],
             "prefill_side_channel_port": self.push_metadata["remote_port"],
             "prefill_tp_size": self.push_metadata["tp_size"],
+            "prefill_pp_size": self.push_metadata["pp_size"],
         }
 
     # ── push-mode request handling ──────────────────────────────────── #
@@ -265,7 +269,7 @@ class PushProxy:
         generator = self.forward_request(
             f"http://{decode_instance}{path}", decode_request, headers
         )
-        return StreamingResponse(generator)
+        return StreamingResponse(generator, media_type="application/json")
 
     async def create_completion(self, raw_request: Request):
         try:
@@ -316,6 +320,7 @@ class PushProxyServer:
             prefill_kv_host=args.prefill_kv_host,
             prefill_side_channel_port=args.prefill_side_channel_port,
             prefill_tp_size=args.prefill_tp_size,
+            prefill_pp_size=args.prefill_pp_size,
             custom_create_completion=create_completion,
             custom_create_chat_completion=create_chat_completion,
         )
@@ -420,6 +425,12 @@ def parse_args():
         default=1,
         help="Tensor parallel size of the prefill vLLM instance",
     )
+    parser.add_argument(
+        "--prefill-pp-size",
+        type=int,
+        default=1,
+        help="Pipeline parallel size of the prefill vLLM instance",
+    )
     return parser.parse_args()
 
 

examples/generate/multimodal/vision_language_offline.py

@@ -448,24 +448,6 @@ def run_exaone4_5(questions: list[str], modality: str) -> ModelRequestData:
     )
 
 
-# Fuyu
-def run_fuyu(questions: list[str], modality: str) -> ModelRequestData:
-    assert modality == "image"
-
-    prompts = [f"{question}\n" for question in questions]
-    engine_args = EngineArgs(
-        model="adept/fuyu-8b",
-        max_model_len=2048,
-        max_num_seqs=2,
-        limit_mm_per_prompt={modality: 1},
-    )
-
-    return ModelRequestData(
-        engine_args=engine_args,
-        prompts=prompts,
-    )
-
-
 # Gemma 3
 def run_gemma3(questions: list[str], modality: str) -> ModelRequestData:
     assert modality == "image"
@@ -2318,7 +2300,6 @@ model_example_map = {
     "eagle2_5": run_eagle2_5,
     "ernie45_vl": run_ernie45_vl,
     "exaone4_5": run_exaone4_5,
-    "fuyu": run_fuyu,
     "gemma3": run_gemma3,
     "gemma3n": run_gemma3n,
     "glm4v": run_glm4v,

requirements/cuda.txt

@@ -9,7 +9,7 @@ torchaudio==2.11.0
 # These must be updated alongside torch
 torchvision==0.26.0 # Required for phi3v processor. See https://github.com/pytorch/vision?tab=readme-ov-file#installation for corresponding version
 torchcodec >= 0.14
-PyNvVideoCodec==2.1.0
+PyNvVideoCodec==2.0.4
 # FlashInfer should be updated together with the Dockerfile
 flashinfer-python==0.6.13
 flashinfer-cubin==0.6.13

requirements/test/cpu.txt

@@ -962,7 +962,7 @@ s3transfer==0.10.3
     # via boto3
 sacrebleu==2.4.3
     # via lm-eval
-safetensors==0.7.0
+safetensors==0.8.0
     # via
     #   -r requirements/test/../common.txt
     #   accelerate
@@ -1159,7 +1159,7 @@ tqdm==4.67.3
     #   segmentation-models-pytorch
     #   sentence-transformers
     #   transformers
-transformers==5.10.4
+transformers==5.13.1
     # via
     #   -r requirements/test/../common.txt
     #   -r requirements/test/cuda.in

requirements/test/cuda.in

@@ -39,7 +39,7 @@ open_clip_torch==2.32.0 # Required for nemotron_vl test, Nemotron Parse in test_
 datamodel_code_generator # required for minicpm3 test
 lm-eval[api]>=0.4.12 # required for model evaluation test
 mteb[bm25s]>=2, <3 # required for mteb test
-transformers==5.10.4
+transformers==5.13.1
 tokenizers==0.22.2
 schemathesis>=4.0.0 # Required for openai schema test.
 # quantization

requirements/test/cuda.txt

@@ -1053,7 +1053,7 @@ s3transfer==0.10.3
     # via boto3
 sacrebleu==2.4.3
     # via lm-eval
-safetensors==0.7.0
+safetensors==0.8.0
     # via
     #   -c requirements/common.txt
     #   -r requirements/test/../common.txt
@@ -1261,7 +1261,7 @@ tqdm==4.67.3
     #   segmentation-models-pytorch
     #   sentence-transformers
     #   transformers
-transformers==5.10.4
+transformers==5.13.1
     # via
     #   -c requirements/common.txt
     #   -r requirements/test/../common.txt

requirements/test/nightly-torch.txt

@@ -29,7 +29,7 @@ opencv-python-headless >= 4.13.0 # required for video test
 datamodel_code_generator # required for minicpm3 test
 lm-eval[api]>=0.4.12 # required for model evaluation test
 mteb[bm25s]>=2, <3 # required for mteb test
-transformers==5.10.4
+transformers==5.13.1
 tokenizers==0.22.2
 schemathesis>=4.0.0 # Required for openai schema test.
 # quantization

requirements/test/rocm.in

@@ -35,7 +35,7 @@ open_clip_torch==2.32.0 # Required for nemotron_vl test, Nemotron Parse in test_
 datamodel_code_generator # required for minicpm3 test
 lm-eval[api]>=0.4.12 # required for model evaluation test
 mteb[bm25s]>=2, <3 # required for mteb test
-transformers==5.10.4
+transformers==5.13.1
 tokenizers==0.22.2
 schemathesis>=4.0.0 # Required for openai schema test
 # quantization

requirements/test/rocm.txt

@@ -1035,7 +1035,7 @@ s3transfer==0.16.0
     # via boto3
 sacrebleu==2.6.0
     # via lm-eval
-safetensors==0.7.0
+safetensors==0.8.0
     # via
     #   -c requirements/common.txt
     #   -r requirements/test/../common.txt
@@ -1218,7 +1218,7 @@ tqdm==4.67.3
     #   sentence-transformers
     #   tilelang
     #   transformers
-transformers==5.10.4
+transformers==5.13.1
     # via
     #   -c requirements/common.txt
     #   -r requirements/test/../common.txt

requirements/test/xpu.in

@@ -17,7 +17,7 @@ accelerate
 arctic-inference
 lm_eval[api]>=0.4.12
 modelscope<1.38
-transformers==5.10.4
+transformers==5.13.1
 
 # --- Audio Processing ---
 librosa

requirements/test/xpu.txt

@@ -779,7 +779,7 @@ rpds-py==0.30.0
     #   referencing
 sacrebleu==2.6.0
     # via lm-eval
-safetensors==0.7.0
+safetensors==0.8.0
     # via
     #   -c requirements/common.txt
     #   -r requirements/test/../common.txt
@@ -940,7 +940,7 @@ tqdm==4.67.3
     #   pqdm
     #   sentence-transformers
     #   transformers
-transformers==5.10.4
+transformers==5.13.1
     # via
     #   -c requirements/common.txt
     #   -r requirements/test/../common.txt

requirements/xpu.txt

@@ -18,4 +18,4 @@ torchvision
 torchcodec >= 0.14 # Required for the torchcodec video decoding backend
 
 auto_round_lib>=0.14.0
-vllm_xpu_kernels @ https://github.com/vllm-project/vllm-xpu-kernels/releases/download/v0.1.10.1/vllm_xpu_kernels-0.1.10.1-cp38-abi3-manylinux_2_28_x86_64.whl
+vllm_xpu_kernels @ https://github.com/vllm-project/vllm-xpu-kernels/releases/download/v0.1.11/vllm_xpu_kernels-0.1.11-cp38-abi3-manylinux_2_28_x86_64.whl

rust/Cargo.lock

@@ -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",

rust/Cargo.toml

@@ -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"

rust/proto/vllm_grpc.proto

@@ -107,6 +107,9 @@ message KVCacheParameters {
 
   // KV Connector transfer parameters
   google.protobuf.Struct kv_transfer_params = 3;
+
+  // Encoder cache connector transfer parameters
+  google.protobuf.Struct ec_transfer_params = 4;
 }
 
 // Controls which extra candidate tokens at each position should be returned
@@ -173,6 +176,7 @@ message FinishInfo {
 
   google.protobuf.Struct kv_transfer_params = 6;
   //uint64 seed = 7;
+  google.protobuf.Struct ec_transfer_params = 8;
 }
 
 // Info for candidate tokens other than the input/sampled

rust/src/chat/Cargo.toml

@@ -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

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

@@ -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),
         }

rust/src/chat/src/error.rs

@@ -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() {

rust/src/chat/src/event.rs

@@ -202,5 +202,8 @@ pub enum ChatEvent {
         finish_reason: FinishReason,
         /// Connector-specific KV transfer parameters for disaggregated serving.
         kv_transfer_params: Option<serde_json::Value>,
+        /// Connector-specific encoder cache transfer parameters for
+        /// disaggregated serving.
+        ec_transfer_params: Option<serde_json::Value>,
     },
 }

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

@@ -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);
+    }
+}

rust/src/chat/src/multimodal/image.rs

@@ -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)
+    }
+}

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

@@ -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))

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

@@ -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>");
+    }
+}

rust/src/chat/src/output/default/unified.rs

@@ -257,6 +257,7 @@ pub(crate) async fn unified_event_stream(
                         usage: finished.usage,
                         finish_reason: finished.finish_reason,
                         kv_transfer_params: finished.kv_transfer_params,
+                        ec_transfer_params: finished.ec_transfer_params,
                     })
                     .await;
                 }
@@ -387,6 +388,7 @@ mod tests {
                 usage: vllm_llm::TokenUsage::default(),
                 finish_reason: crate::FinishReason::Stop(None),
                 kv_transfer_params: None,
+                ec_transfer_params: None,
             }),
         }
     }
@@ -628,6 +630,7 @@ mod tests {
                     usage: vllm_llm::TokenUsage::default(),
                     finish_reason: crate::FinishReason::Stop(None),
                     kv_transfer_params: None,
+                    ec_transfer_params: None,
                 },
             ]
         );
@@ -671,6 +674,7 @@ mod tests {
                     usage: vllm_llm::TokenUsage::default(),
                     finish_reason: crate::FinishReason::Stop(None),
                     kv_transfer_params: None,
+                    ec_transfer_params: None,
                 },
             ]
         );

rust/src/chat/src/output/harmony/mod.rs

@@ -370,6 +370,7 @@ async fn harmony_assistant_event_stream(
                         usage: finished.usage,
                         finish_reason: finished.finish_reason,
                         kv_transfer_params: finished.kv_transfer_params,
+                        ec_transfer_params: finished.ec_transfer_params,
                     })
                     .await;
                 }

rust/src/chat/src/output/harmony/tests.rs

@@ -52,6 +52,7 @@ fn finished() -> Finished {
         },
         finish_reason: FinishReason::stop_eos(),
         kv_transfer_params: None,
+        ec_transfer_params: None,
     }
 }
 
@@ -115,6 +116,7 @@ fn interrupted_final_message_is_preserved() {
             },
             finish_reason: FinishReason::stop_eos(),
             kv_transfer_params: None,
+            ec_transfer_params: None,
         })
     );
 }
@@ -175,6 +177,7 @@ fn interrupted_analysis_message_is_preserved() {
             },
             finish_reason: FinishReason::stop_eos(),
             kv_transfer_params: None,
+            ec_transfer_params: None,
         })
     );
 }

rust/src/chat/src/output/mod.rs

@@ -48,6 +48,9 @@ pub(crate) enum AssistantEvent {
         finish_reason: FinishReason,
         /// Connector-specific KV transfer parameters for disaggregated serving.
         kv_transfer_params: Option<serde_json::Value>,
+        /// Connector-specific encoder cache transfer parameters for
+        /// disaggregated serving.
+        ec_transfer_params: Option<serde_json::Value>,
     },
 }