[Core] Add explicit layer parallel plans

Co-authored-by: OpenAI Codex <codex@openai.com> Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
[Model] Optimize Qwen3.5 on H20 (#48350 )
2026-07-13 19:12:30 +00:00 · 2026-07-13 03:30:48 +00:00 · 2026-07-12 19:27:29 -07:00 · 2026-07-12 19:18:03 -07:00 · 2026-07-12 20:00:04 +03:00 · 2026-07-12 16:34:26 +01:00
293 changed files with 11211 additions and 3099 deletions
@@ -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
@@ -177,6 +177,18 @@ BRANCH=$4
 IMAGE_TAG=$5
 IMAGE_TAG_LATEST=${6:-} # only used for main branch, optional
 # When TORCH_NIGHTLY=1, build the base CI image against PyTorch nightly so the
 # entire existing pipeline runs on nightly torch (CUDA/GPU lane only). Delegate
 # to the dedicated nightly build (PYTORCH_NIGHTLY=1, CUDA 13.0) and tag it at the
 # normal IMAGE_TAG that every test step already pulls -- no separate image tag,
 # no duplicate "vLLM Against PyTorch Nightly" pipeline section.
 if [[ "${TORCH_NIGHTLY:-0}" == "1" ]]; then
    echo "--- :warning: TORCH_NIGHTLY=1 -- building base image on PyTorch nightly"
    SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
    exec "${SCRIPT_DIR}/image_build_torch_nightly.sh" \
        "${REGISTRY}" "${REPO}" "${BUILDKITE_COMMIT}" "${BRANCH}" "${IMAGE_TAG}"
 fi
 # build config
 TARGET="test-ci"
 VLLM_BAKE_FILE_PATH="${VLLM_BAKE_FILE_PATH:-docker/docker-bake.hcl}"
@@ -19,13 +19,14 @@ if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
  echo "Image found"
 else
  echo "Image not found, proceeding with build..."
-  # build for arm64 GPU targets: Grace/GH200 (sm_90) and DGX Spark/GB10
+  # build for arm64 GPU targets: Grace/GH200 (sm_90),
  # Blackwell/Thor (sm_100/sm_103/sm_110), and DGX Spark/GB10
  # (sm_121, family-covered by 12.0 under CUDA 13)
  docker build --file docker/Dockerfile \
    --platform linux/arm64 \
    --build-arg max_jobs=16 \
    --build-arg nvcc_threads=4 \
-    --build-arg torch_cuda_arch_list="9.0 12.0" \
+    --build-arg torch_cuda_arch_list="9.0 10.0 11.0 12.0" \
    --build-arg USE_SCCACHE=1 \
    --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
    --tag "$IMAGE" \
@@ -72,9 +72,7 @@ steps:
      pytest -v -s v1/test_oracle.py &&
      pytest -v -s v1/test_request.py &&
      pytest -v -s v1/test_outputs.py &&
-      pytest -v -s v1/sample/test_topk_topp_sampler.py &&
+      pytest -v -s v1/sample'
      pytest -v -s v1/sample/test_logprobs.py &&
      pytest -v -s v1/sample/test_logprobs_e2e.py'
 - label: Basic Models Tests (Initialization)
  timeout_in_minutes: 60
@@ -848,6 +848,23 @@ steps:
      DOCKER_BUILDKIT: "1"
      S3_BUCKET: "vllm-wheels"
  - label: "Publish nightly XPU image to DockerHub"
    depends_on:
      - create-manifest-xpu
    if: build.env("NIGHTLY") == "1"
    agents:
      queue: small_cpu_queue_release
    commands:
      - "bash .buildkite/scripts/xpu/push-nightly-builds-xpu.sh"
      - "bash .buildkite/scripts/cleanup-nightly-builds.sh nightly- vllm/vllm-openai-xpu"
    plugins:
      - docker-login#v3.0.0:
          username: vllmbot
          password-env: DOCKERHUB_TOKEN
    env:
      DOCKER_BUILDKIT: "1"
      DOCKERHUB_USERNAME: "vllmbot"
  - label: "Publish nightly ROCm image to DockerHub"
    depends_on:
      - build-rocm-release-image
@@ -878,6 +895,7 @@ steps:
      - create-multi-arch-manifest-cuda-12-9
      - create-multi-arch-manifest-ubuntu2404
      - create-multi-arch-manifest-cuda-12-9-ubuntu2404
      - create-manifest-xpu
      - build-rocm-release-image
      - input-release-version
      # Wait for CPU builds if their block steps were unblocked, so publish
@@ -29,7 +29,11 @@ if python3 -c "import torch; assert torch.version.hip" 2>/dev/null; then
        TORCH_INDEX_URL=""
    fi
 else
-    TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
+    if [ "${TORCH_NIGHTLY:-0}" = "1" ]; then
        TORCH_INDEX_URL="https://download.pytorch.org/whl/nightly/cu130"
    else
        TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
    fi
 fi
 echo ">>> Using PyTorch index: ${TORCH_INDEX_URL:-PyPI default}"
@@ -130,6 +130,22 @@ docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm
 docker push vllm/vllm-openai-rocm:latest-base
 docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
 # ---- XPU ----
 docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu
 docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu vllm/vllm-openai-xpu:latest-x86_64
 docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64
 docker push vllm/vllm-openai-xpu:latest-x86_64
 docker push vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64
 docker manifest rm vllm/vllm-openai-xpu:latest || true
 docker manifest rm vllm/vllm-openai-xpu:v${RELEASE_VERSION} || true
 docker manifest create vllm/vllm-openai-xpu:latest vllm/vllm-openai-xpu:latest-x86_64 --amend
 docker manifest create vllm/vllm-openai-xpu:v${RELEASE_VERSION} vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64 --amend
 docker manifest push vllm/vllm-openai-xpu:latest
 docker manifest push vllm/vllm-openai-xpu:v${RELEASE_VERSION}
 # ---- CPU ----
 # CPU images are behind separate block steps and may not have been built.
 # All-or-nothing: inspect both arches first, then either publish everything
@@ -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
+  log_section "Installing cargo-binstall ${CARGO_BINSTALL_VERSION}"
    return
  fi
  log_section "Installing cargo-binstall"
  curl -L --proto '=https' --tlsv1.2 -sSf \
-    https://raw.githubusercontent.com/cargo-bins/cargo-binstall/main/install-from-binstall-release.sh \
+    "https://raw.githubusercontent.com/cargo-bins/cargo-binstall/v${CARGO_BINSTALL_VERSION}/install-from-binstall-release.sh" \
-    | bash
+    | env BINSTALL_VERSION="$CARGO_BINSTALL_VERSION" bash
  cargo-binstall -V
 }
 install_cargo_sort() {
-  if command -v cargo-sort >/dev/null 2>&1; then
+  log_section "Installing cargo-sort ${CARGO_SORT_VERSION_REQ}"
-    return
+  cargo binstall --no-confirm --force "cargo-sort@${CARGO_SORT_VERSION_REQ}"
  fi
  log_section "Installing cargo-sort"
  install_cargo_binstall
  cargo binstall --no-confirm cargo-sort
 }
 install_cargo_deny() {
-  if command -v cargo-deny >/dev/null 2>&1; then
+  log_section "Installing cargo-deny ${CARGO_DENY_VERSION_REQ}"
-    return
+  cargo binstall --no-confirm --force "cargo-deny@${CARGO_DENY_VERSION_REQ}"
  fi
  log_section "Installing cargo-deny"
  install_cargo_binstall
  cargo binstall --no-confirm cargo-deny
 }
 install_cargo_nextest() {
-  if command -v cargo-nextest >/dev/null 2>&1; then
+  log_section "Installing cargo-nextest ${CARGO_NEXTEST_VERSION_REQ}"
-    return
+  cargo binstall \
-  fi
+    --no-confirm \
-
+    --force \
-  log_section "Installing cargo-nextest"
+    --secure \
-  install_cargo_binstall
+    "cargo-nextest@${CARGO_NEXTEST_VERSION_REQ}"
  cargo binstall --no-confirm --secure cargo-nextest
 }
 install_uv() {
-  if command -v uv >/dev/null 2>&1; then
+  log_section "Installing uv ${UV_VERSION}"
-    return
+  curl -L --proto '=https' --tlsv1.2 -sSf \
-  fi
+    "https://github.com/astral-sh/uv/releases/download/${UV_VERSION}/uv-installer.sh" \
  log_section "Installing uv"
  curl -LsSf --proto '=https' --tlsv1.2 https://astral.sh/uv/install.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 ${PYO3_PYTHON_VERSION} for PyO3 tests"
-
+  uv python install "$PYO3_PYTHON_VERSION"
  log_section "Installing Python ${python_version} for PyO3 tests"
  uv python install "$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
@@ -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"
@@ -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)
+    --execute        Actually trigger the build (default: dry-run)
-    --pipeline      Buildkite pipeline slug (default: ${PIPELINE})
+    --pipeline       Buildkite pipeline slug (default: ${PIPELINE})
-    --commit        Override commit SHA (default: current HEAD)
+    --commit         Override commit SHA (default: current HEAD)
-    --branch        Override branch name (default: current branch)
+    --branch         Override branch name (default: current branch)
-    --message       Custom build message (default: auto-generated)
+    --message        Custom build message (default: auto-generated)
-    --help          Show this help message
+    --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' \\"
+    last_idx=$(( ${#BUILD_ENV[@]} - 1 ))
-    echo "    --env 'NIGHTLY=1'"
+    for i in "${!BUILD_ENV[@]}"; do
        if [[ $i -eq $last_idx ]]; then
            echo "    --env '$(escape_for_shell "${BUILD_ENV[$i]}")'"
        else
            echo "    --env '$(escape_for_shell "${BUILD_ENV[$i]}")' \\"
        fi
    done
    echo ""
    echo "=========================================="
    echo -e "${YELLOW}To actually trigger this build, run:${NC}"
@@ -0,0 +1,21 @@
 #!/bin/bash
 set -ex
 ORIG_TAG_NAME="$BUILDKITE_COMMIT"
 REPO="vllm/vllm-openai-xpu"
 echo "Pushing original XPU tag ${ORIG_TAG_NAME}-xpu to nightly tags in ${REPO}"
 aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
 docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:"$ORIG_TAG_NAME"-x86_64-xpu
 docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:"$ORIG_TAG_NAME"-x86_64-xpu ${REPO}:nightly-x86_64
 docker push ${REPO}:nightly-x86_64
 docker manifest rm ${REPO}:nightly || true
 docker manifest rm ${REPO}:nightly-"$BUILDKITE_COMMIT" || true
 docker manifest create ${REPO}:nightly ${REPO}:nightly-x86_64 --amend
 docker manifest create ${REPO}:nightly-"$BUILDKITE_COMMIT" ${REPO}:nightly-x86_64 --amend
 docker manifest push ${REPO}:nightly
 docker manifest push ${REPO}:nightly-"$BUILDKITE_COMMIT"
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: V1 attention (H100-MI300)
  key: v1-attention-h100-mi300
-  timeout_in_minutes: 30
+  timeout_in_minutes: 85
  device: h100
  source_file_dependencies:
    - vllm/config/attention.py
@@ -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
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Basic Correctness
  key: basic-correctness
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -19,6 +19,6 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 50
+      timeout_in_minutes: 70
      depends_on:
      - image-build-amd
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Benchmarks CLI Test
  key: benchmarks-cli-test
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -23,7 +23,7 @@ steps:
  num_gpus: 2
  optional: true
  working_dir: "/vllm-workspace/"
-  timeout_in_minutes: 10
+  timeout_in_minutes: 20
  source_file_dependencies:
  - benchmarks/attention_benchmarks/
  - vllm/v1/attention/
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Sequence Parallel Correctness Tests (2 GPUs)
  key: sequence-parallel-correctness-tests-2-gpus
-  timeout_in_minutes: 50
+  timeout_in_minutes: 80
  working_dir: "/vllm-workspace/"
  num_devices: 2
  source_file_dependencies:
@@ -19,7 +19,7 @@ steps:
 - label: Sequence Parallel Correctness Tests (2xH100)
  key: sequence-parallel-correctness-tests-2xh100
-  timeout_in_minutes: 50
+  timeout_in_minutes: 75
  working_dir: "/vllm-workspace/"
  device: h100
  optional: true
@@ -30,7 +30,7 @@ steps:
 - label: AsyncTP Correctness Tests (2xH100)
  key: asynctp-correctness-tests-2xh100
-  timeout_in_minutes: 50
+  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/"
  device: h100
  optional: true
@@ -41,7 +41,7 @@ steps:
 - label: AsyncTP Correctness Tests (B200)
  key: asynctp-correctness-tests-b200
-  timeout_in_minutes: 50
+  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/"
  device: b200-k8s
  optional: true
@@ -52,7 +52,7 @@ steps:
 - label: Distributed Compile Unit Tests (2xH100)
  key: distributed-compile-unit-tests-2xh100
-  timeout_in_minutes: 20
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/"
  device: h100
  num_devices: 2
@@ -66,7 +66,7 @@ steps:
 - label: Fusion and Compile Unit Tests (2xB200)
  key: fusion-and-compile-unit-tests-2xb200
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/"
  device: b200-k8s
  source_file_dependencies:
@@ -96,7 +96,7 @@ steps:
 - label: Fusion E2E Quick (H100)
  key: fusion-e2e-quick-h100
-  timeout_in_minutes: 15
+  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/"
  device: h100
  num_devices: 1
@@ -115,7 +115,7 @@ steps:
 - label: Fusion E2E Config Sweep (H100)
  key: fusion-e2e-config-sweep-h100
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/"
  device: h100
  num_devices: 1
@@ -149,7 +149,7 @@ steps:
 - label: Fusion E2E TP2 Quick (H100)
  key: fusion-e2e-tp2-quick-h100
-  timeout_in_minutes: 20
+  timeout_in_minutes: 35
  working_dir: "/vllm-workspace/"
  device: h100
  num_devices: 2
@@ -167,7 +167,7 @@ steps:
 - label: Fusion E2E TP2 AR-RMS Config Sweep (H100)
  key: fusion-e2e-tp2-ar-rms-config-sweep-h100
-  timeout_in_minutes: 40
+  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/"
  device: h100
  num_devices: 2
@@ -207,7 +207,7 @@ steps:
 - label: Fusion E2E TP2 (B200)
  key: fusion-e2e-tp2-b200
-  timeout_in_minutes: 20
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/"
  device: b200-k8s
  num_devices: 2
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Platform Tests
  key: platform-tests
-  timeout_in_minutes: 15
+  timeout_in_minutes: 20
  device: h200_18gb
  source_file_dependencies:
  - vllm/envs.py
@@ -19,7 +19,7 @@ steps:
 - label: Cudagraph
  key: cudagraph
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  source_file_dependencies:
  - tests/v1/cudagraph
  - vllm/v1/cudagraph_dispatcher.py
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed NixlConnector PD accuracy (4 GPUs)
  key: distributed-nixlconnector-pd-accuracy-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 55
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -16,7 +16,7 @@ steps:
  mirror:
    amd:
      device: mi300_4
-      timeout_in_minutes: 110
+      timeout_in_minutes: 85
      depends_on:
        - image-build-amd
      source_file_dependencies:
@@ -29,7 +29,7 @@ steps:
 - label: Distributed FlashInfer NixlConnector PD accuracy (4 GPUs)
  key: distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 55
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -66,7 +66,7 @@ steps:
  mirror:
    amd:
      device: mi300_4
-      timeout_in_minutes: 50
+      timeout_in_minutes: 60
      depends_on:
        - image-build-amd
      source_file_dependencies:
@@ -79,7 +79,7 @@ steps:
 - label: CrossLayer KV layout Distributed NixlConnector PD accuracy tests (4 GPUs)
  key: crosslayer-kv-layout-distributed-nixlconnector-pd-accuracy-tests-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 55
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -91,7 +91,7 @@ steps:
  mirror:
    amd:
      device: mi300_4
-      timeout_in_minutes: 110
+      timeout_in_minutes: 85
      depends_on:
        - image-build-amd
      source_file_dependencies:
@@ -104,7 +104,7 @@ steps:
 - label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs)
  key: hybrid-ssm-nixlconnector-pd-accuracy-tests-4-gpus
-  timeout_in_minutes: 25
+  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -116,7 +116,7 @@ steps:
  mirror:
    amd:
      device: mi300_4
-      timeout_in_minutes: 60
+      timeout_in_minutes: 80
      depends_on:
        - image-build-amd
      source_file_dependencies:
@@ -143,7 +143,7 @@ steps:
 - label: MultiConnector (Nixl+Offloading) PD accuracy (2 GPUs)
  key: multiconnector-nixl-offloading-pd-accuracy-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -158,7 +158,7 @@ steps:
 - label: NixlConnector PD + Spec Decode acceptance (2 GPUs)
  key: nixlconnector-pd-spec-decode-acceptance-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
  device: a100
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
@@ -172,7 +172,7 @@ steps:
  mirror:
    amd:
      device: mi300_2
-      timeout_in_minutes: 60
+      timeout_in_minutes: 70
      depends_on:
        - image-build-amd
      source_file_dependencies:
@@ -186,7 +186,7 @@ steps:
 - label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs)
  key: multiconnector-nixl-offloading-pd-edge-cases-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed Comm Ops
  key: distributed-comm-ops
-  timeout_in_minutes: 20
+  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -18,7 +18,7 @@ steps:
 - label: Distributed DP Tests (2 GPUs)
  key: distributed-dp-tests-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 35
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -55,7 +55,7 @@ steps:
 - label: Distributed Compile + RPC Tests (2 GPUs)
  key: distributed-compile-rpc-tests-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 65
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -78,7 +78,7 @@ steps:
 - label: Distributed Torchrun + Shutdown Tests (2 GPUs)
  key: distributed-torchrun-shutdown-tests-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -133,7 +133,7 @@ steps:
 - label: Distributed DP Tests (4 GPUs)
  key: distributed-dp-tests-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -154,7 +154,7 @@ steps:
 - label: Distributed Compile + Comm (4 GPUs)
  key: distributed-compile-comm-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 70
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -176,7 +176,7 @@ steps:
 - label: Distributed Tests (8xH100)
  key: distributed-tests-8xh100
-  timeout_in_minutes: 10
+  timeout_in_minutes: 20
  device: h100
  num_devices: 8
  working_dir: "/vllm-workspace/tests"
@@ -212,7 +212,7 @@ steps:
 - label: Distributed Tests (2xH100-2xMI300)
  key: distributed-tests-2xh100-2xmi300
-  timeout_in_minutes: 15
+  timeout_in_minutes: 30
  device: h100
  optional: true
  working_dir: "/vllm-workspace/"
@@ -259,7 +259,7 @@ steps:
 - label: Pipeline + Context Parallelism (4 GPUs)
  key: pipeline-context-parallelism-4-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 55
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -274,7 +274,7 @@ steps:
 - label: RayExecutorV2 (4 GPUs)
  key: rayexecutorv2-4-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -3,7 +3,7 @@ depends_on:
  - image-build-cpu
 steps:
 - label: Docker Build Metadata
-  timeout_in_minutes: 10
+  timeout_in_minutes: 20
  device: cpu-small
  source_file_dependencies:
    - .buildkite/release-pipeline.yaml
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: DeepSeek V2-Lite Sync EPLB Accuracy (4xH100)
  key: deepseek-v2-lite-sync-eplb-accuracy-4xh100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
  device: h100
  optional: true
  num_devices: 4
@@ -14,7 +14,7 @@ steps:
 - label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (4xH100)
  key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-4xh100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
  device: h100
  optional: true
  num_devices: 4
@@ -24,7 +24,7 @@ steps:
 - label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (2xB200)
  key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-2xb200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 20
  device: b200-k8s
  optional: true
  num_devices: 2
@@ -34,7 +34,7 @@ steps:
 - label: Qwen3-30B-A3B-FP8 DP4 Async EPLB Accuracy
  key: qwen3-30b-a3b-fp8-dp4-async-eplb-accuracy
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
  device: h100
  optional: true
  num_devices: 4
@@ -44,7 +44,7 @@ steps:
 - label: DeepSeek V2-Lite Prefetch Offload Accuracy (H100)
  key: deepseek-v2-lite-prefetch-offload-accuracy-h100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 20
  device: h100
  optional: true
  num_devices: 1
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Engine
  key: engine
-  timeout_in_minutes: 15
+  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
  - vllm/compilation/
@@ -29,13 +29,13 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 50
      depends_on:
      - image-build-amd
 - label: Engine (1 GPU)
  key: engine-1-gpu
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
  source_file_dependencies:
    - vllm/v1/engine/
    - tests/v1/engine/
@@ -45,13 +45,13 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 40
+      timeout_in_minutes: 55
      depends_on:
      - image-build-amd
 - label: e2e Scheduling (1 GPU)
  key: e2e-scheduling-1-gpu
-  timeout_in_minutes: 30
+  timeout_in_minutes: 35
  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/
@@ -61,14 +61,14 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 70
      depends_on:
      - image-build-amd
 - label: e2e Core (1 GPU)
  device: h200_35gb
  key: e2e-core-1-gpu
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
  source_file_dependencies:
    - vllm/v1/
    - tests/v1/e2e/general/
@@ -77,7 +77,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 35
+      timeout_in_minutes: 60
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -87,7 +87,7 @@ steps:
 - label: V1 e2e (2 GPUs)
  key: v1-e2e-2-gpus
-  timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
+  timeout_in_minutes: 25 # TODO: Fix timeout after we have more confidence in the test stability
  optional: true
  num_devices: 2
  source_file_dependencies:
@@ -120,7 +120,7 @@ steps:
 - label: V1 e2e (4 GPUs)
  key: v1-e2e-4-gpus
-  timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
+  timeout_in_minutes: 20 # TODO: Fix timeout after we have more confidence in the test stability
  optional: true
  num_devices: 4
  source_file_dependencies:
@@ -148,7 +148,7 @@ steps:
 - label: V1 e2e (4xH100)
  key: v1-e2e-4xh100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 35
  device: h100
  num_devices: 4
  optional: true
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Entrypoints Unit Tests
  key: entrypoints-unit-tests
-  timeout_in_minutes: 10
+  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/entrypoints
@@ -16,7 +16,7 @@ steps:
 - label: Entrypoints Integration (LLM)
  key: entrypoints-integration-llm
-  timeout_in_minutes: 40
+  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -37,7 +37,7 @@ steps:
 - label: Entrypoints Integration (API Server)
  key: entrypoints-integration-api-server
  device: h200_35gb
-  timeout_in_minutes: 130
+  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -56,7 +56,7 @@ steps:
 - label: Entrypoints Integration (API Server OpenAI - Part 1)
  key: entrypoints-integration-api-server-openai-part-1
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -68,13 +68,13 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 80
+      timeout_in_minutes: 65
      depends_on:
      - image-build-amd
 - label: Entrypoints Integration (API Server OpenAI - Part 2)
  key: entrypoints-integration-api-server-openai-part-2
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -109,7 +109,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 65
      depends_on:
      - image-build-amd
@@ -126,7 +126,7 @@ steps:
 - label: Entrypoints Integration (Speech to Text)
  device: h200_35gb
  key: entrypoints-integration-speech_to_text
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -138,7 +138,7 @@ steps:
 - label: Entrypoints Integration (Multimodal)
  device: h200_35gb
  key: entrypoints-integration-multimodal
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -160,7 +160,7 @@ steps:
 - label: OpenAI API Correctness
  key: openai-api-correctness
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: h200_18gb
  source_file_dependencies:
  - csrc/
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: EPLB Algorithm
  key: eplb-algorithm
-  timeout_in_minutes: 15
+  timeout_in_minutes: 20
  device: h200_18gb
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -27,7 +27,7 @@ steps:
 - label: EPLB Execution # 17min
  key: eplb-execution
-  timeout_in_minutes: 27
+  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -39,7 +39,7 @@ steps:
 - label: Elastic EP Scaling Test
  key: elastic-ep-scaling-test
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  device: h100
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: vLLM IR Tests
  key: vllm-ir-tests
-  timeout_in_minutes: 10
+  timeout_in_minutes: 35
  device: h200_18gb
  working_dir: "/vllm-workspace/"
  source_file_dependencies:
@@ -16,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:
@@ -5,7 +5,7 @@ steps:
 - label: LM Eval Small Models
  device: h200_35gb
  key: lm-eval-small-models
-  timeout_in_minutes: 75
+  timeout_in_minutes: 45
  source_file_dependencies:
  - csrc/
  - vllm/model_executor/layers/quantization
@@ -56,7 +56,7 @@ steps:
 - label: LM Eval Small Models (1xB200)
  key: lm-eval-small-models-1xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 50
  device: b200-k8s
  optional: true
  source_file_dependencies:
@@ -80,7 +80,7 @@ steps:
 - label: LM Eval Large Models EP (2xB200)
  key: lm-eval-large-models-ep-2xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 60
  device: b200-k8s
  optional: true
  num_devices: 2
@@ -92,7 +92,7 @@ steps:
 - label: LM Eval Qwen3.5 Models (2xB200)
  key: lm-eval-qwen3-5-models-2xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 45
  device: b200-k8s
  optional: true
  num_devices: 2
@@ -109,7 +109,7 @@ steps:
 - label: LM Eval Large Models (8xH200)
  key: lm-eval-large-models-8xh200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 50
  device: h200
  optional: true
  num_devices: 8
@@ -118,7 +118,7 @@ steps:
  mirror:
    amd:
      device: mi300_8
-      timeout_in_minutes: 180
+      timeout_in_minutes: 60
      depends_on:
      - image-build-amd
      commands:
@@ -152,7 +152,7 @@ steps:
 - label: LM Eval Humming f16 (A100 - TEMPORARY)
  key: lm-eval-humming-f16-a100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 75
  device: a100
  optional: true
  num_devices: 1
@@ -167,7 +167,7 @@ steps:
 - label: LM Eval Humming Act int8 (A100 - TEMPORARY)
  key: lm-eval-humming-act-a100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 45
  device: a100
  optional: true
  num_devices: 1
@@ -182,7 +182,7 @@ steps:
 - label: LM Eval Humming f16 (H100 - TEMPORARY)
  key: lm-eval-humming-f16-h100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 70
  device: h100
  optional: true
  num_devices: 1
@@ -197,7 +197,7 @@ steps:
 - label: LM Eval Humming Act fp8/int8 (H100 - TEMPORARY)
  key: lm-eval-humming-act-h100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 70
  device: h100
  optional: true
  num_devices: 1
@@ -213,7 +213,7 @@ steps:
 - label: LM Eval Humming f16 (B200 - TEMPORARY)
  key: lm-eval-humming-f16-b200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 50
  device: b200-k8s
  optional: true
  num_devices: 1
@@ -228,7 +228,7 @@ steps:
 - label: LM Eval Humming Act fp8/int8 (B200 - TEMPORARY)
  key: lm-eval-humming-act-b200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 50
  device: b200-k8s
  optional: true
  num_devices: 1
@@ -244,7 +244,7 @@ steps:
 - label: LM Eval TurboQuant KV Cache
  key: lm-eval-turboquant-kv-cache
-  timeout_in_minutes: 75
+  timeout_in_minutes: 55
  device: h200_18gb
  source_file_dependencies:
  - vllm/model_executor/layers/quantization/turboquant/
@@ -256,7 +256,7 @@ steps:
 - label: GPQA Eval (GPT-OSS) (2xH100)
  key: gpqa-eval-gpt-oss-2xh100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 35
  device: h100
  optional: true
  num_devices: 2
@@ -270,7 +270,7 @@ steps:
 - label: GPQA Eval (GPT-OSS) (2xB200)
  key: gpqa-eval-gpt-oss-2xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 30
  device: b200-k8s
  optional: true
  num_devices: 2
@@ -284,7 +284,7 @@ steps:
 - label: GPQA Eval (GPT-OSS) (DGX Spark)
  key: gpqa-eval-gpt-oss-spark
-  timeout_in_minutes: 120
+  timeout_in_minutes: 35
  device: dgx-spark
  optional: true
  num_devices: 1
@@ -313,7 +313,7 @@ steps:
 - label: LM Eval KV-Offload (2xH100)
  key: kv-offload-medium
-  timeout_in_minutes: 60
+  timeout_in_minutes: 30
  device: h100
  num_devices: 2
  source_file_dependencies:
@@ -327,7 +327,7 @@ steps:
 - label: LM Eval KV-Offload (4xH100)
  key: kv-offload-large
-  timeout_in_minutes: 60
+  timeout_in_minutes: 40
  device: h100
  num_devices: 4
  source_file_dependencies:
@@ -341,7 +341,7 @@ steps:
 - label: MRCR Eval Small Models
  device: h200_35gb
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
  source_file_dependencies:
  - tests/evals/mrcr/
  commands:
@@ -5,7 +5,7 @@ steps:
 - label: LoRA %N
  device: h200_35gb
  key: lora
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
  source_file_dependencies:
  - vllm/lora
  - tests/lora
@@ -16,7 +16,7 @@ steps:
    amd:
      device: mi325_1
      working_dir: "/vllm-workspace/tests"
-      timeout_in_minutes: 60
+      timeout_in_minutes: 65
      source_file_dependencies:
      - vllm/lora
      - tests/lora
@@ -27,7 +27,7 @@ steps:
 - label: LoRA TP (Distributed)
  key: lora-tp-distributed
-  timeout_in_minutes: 30
+  timeout_in_minutes: 60
  num_devices: 4
  source_file_dependencies:
  - vllm/lora
@@ -5,7 +5,7 @@ steps:
 - label: V1 Spec Decode
  device: h200_35gb
  key: v1-spec-decode
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
  source_file_dependencies:
    - vllm/config/
    - vllm/distributed/
@@ -24,13 +24,13 @@ steps:
  mirror:
    amd:
      device: mi300_1
-      timeout_in_minutes: 65
+      timeout_in_minutes: 75
      depends_on:
      - image-build-amd
 - label: V1 Sample + Logits
  key: v1-sample-logits
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
  device: h200_18gb
  source_file_dependencies:
    - vllm/config/
@@ -64,7 +64,7 @@ steps:
 - label: V1 Core + KV + Metrics
  key: v1-core-kv-metrics
-  timeout_in_minutes: 30
+  timeout_in_minutes: 60
  source_file_dependencies:
    - vllm/config/
    - vllm/distributed/
@@ -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
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Model Executor
  key: model-executor
-  timeout_in_minutes: 35
+  timeout_in_minutes: 45
  source_file_dependencies:
  - vllm/engine/arg_utils.py
  - vllm/config/model.py
@@ -5,7 +5,7 @@ steps:
 - label: Model Runner V2 Core Tests
  device: h200_35gb
  key: model-runner-v2-core-tests
-  timeout_in_minutes: 45
+  timeout_in_minutes: 35
  source_file_dependencies:
  - vllm/v1/worker/gpu/
  - vllm/v1/worker/gpu_worker.py
@@ -27,7 +27,7 @@ steps:
 - label: Model Runner V2 Examples
  device: h200_35gb
  key: model-runner-v2-examples
-  timeout_in_minutes: 45
+  timeout_in_minutes: 35
  working_dir: "/vllm-workspace/examples"
  source_file_dependencies:
    - vllm/v1/worker/gpu/
@@ -63,7 +63,7 @@ steps:
 - label: Model Runner V2 Distributed (2 GPUs)
  key: model-runner-v2-distributed-2-gpus
-  timeout_in_minutes: 45
+  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -84,7 +84,7 @@ steps:
 - label: Model Runner V2 Pipeline Parallelism (4 GPUs)
  key: model-runner-v2-pipeline-parallelism-4-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Basic Models Tests (Initialization)
  key: basic-models-tests-initialization
-  timeout_in_minutes: 45
+  timeout_in_minutes: 25
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -17,7 +17,7 @@ steps:
 - label: Basic Models Tests (Extra Initialization) %N
  device: h200_35gb
  key: basic-models-tests-extra-initialization
-  timeout_in_minutes: 45
+  timeout_in_minutes: 100
  source_file_dependencies:
  - vllm/model_executor/models/
  - tests/models/test_initialization.py
@@ -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
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed Model Tests (2 GPUs)
  key: distributed-model-tests-2-gpus
-  timeout_in_minutes: 50
+  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Language Models Tests (Standard)
  key: language-models-tests-standard
-  timeout_in_minutes: 25
+  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -21,7 +21,7 @@ steps:
 - label: Language Models Tests (Extra Standard) %N
  key: language-models-tests-extra-standard
-  timeout_in_minutes: 45
+  timeout_in_minutes: 40
  source_file_dependencies:
  - vllm/model_executor/models/
  - tests/models/language/pooling/test_embedding.py
@@ -52,7 +52,7 @@ steps:
 - label: Language Models Tests (Hybrid) %N
  key: language-models-tests-hybrid
-  timeout_in_minutes: 75
+  timeout_in_minutes: 65
  source_file_dependencies:
  - vllm/
  - tests/models/language/generation
@@ -67,7 +67,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 90
+      timeout_in_minutes: 70
      depends_on:
      - image-build-amd
      commands:
@@ -78,7 +78,7 @@ steps:
 - label: Language Models Test (Extended Generation) # 80min
  device: h200_35gb
  key: language-models-test-extended-generation
-  timeout_in_minutes: 110
+  timeout_in_minutes: 65
  optional: true
  source_file_dependencies:
  - vllm/
@@ -92,7 +92,7 @@ steps:
 - label: Language Models Test (PPL)
  key: language-models-test-ppl
-  timeout_in_minutes: 110
+  timeout_in_minutes: 30
  device: h200_18gb
  optional: true
  source_file_dependencies:
@@ -104,7 +104,7 @@ steps:
 - label: Language Models Test (Extended Pooling)  # 36min
  device: h200_35gb
  key: language-models-test-extended-pooling
-  timeout_in_minutes: 50
+  timeout_in_minutes: 70
  optional: true
  source_file_dependencies:
  - vllm/
@@ -120,7 +120,7 @@ steps:
 - label: Language Models Test (MTEB)
  key: language-models-test-mteb
-  timeout_in_minutes: 110
+  timeout_in_minutes: 45
  device: h200_18gb
  optional: true
  source_file_dependencies:
@@ -20,7 +20,7 @@ steps:
 - label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
  key: multi-modal-models-standard-2-qwen3-gemma
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -38,7 +38,7 @@ steps:
 - label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
  device: h200_35gb
  key: multi-modal-models-standard-3-llava-qwen2-vl
-  timeout_in_minutes: 45
+  timeout_in_minutes: 40
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal
@@ -54,7 +54,7 @@ steps:
 - label: "Multi-Modal Models (Standard) 4: other + whisper"
  device: h200_35gb
  key: multi-modal-models-standard-4-other-whisper
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal
@@ -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:
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Plugin Tests (2 GPUs)
  key: plugin-tests-2-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 35
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -5,7 +5,7 @@ steps:
 - label: PyTorch Compilation Unit Tests
  device: h200_35gb
  key: pytorch-compilation-unit-tests
-  timeout_in_minutes: 10
+  timeout_in_minutes: 90
  source_file_dependencies:
    - vllm/__init__.py
    - vllm/_aiter_ops.py
@@ -78,7 +78,7 @@ steps:
 - label: PyTorch Compilation Passes Unit Tests
  key: pytorch-compilation-passes-unit-tests
-  timeout_in_minutes: 20
+  timeout_in_minutes: 45
  source_file_dependencies:
    - vllm/__init__.py
    - vllm/_aiter_ops.py
@@ -110,13 +110,13 @@ steps:
  mirror:
    amd:
      device: mi300_1
-      timeout_in_minutes: 180
+      timeout_in_minutes: 65
      depends_on:
      - image-build-amd
 - label: PyTorch Fullgraph Smoke Test
  key: pytorch-fullgraph-smoke-test
-  timeout_in_minutes: 35
+  timeout_in_minutes: 60
  source_file_dependencies:
  - vllm/__init__.py
  - vllm/_aiter_ops.py
@@ -152,7 +152,7 @@ steps:
 - label: PyTorch Fullgraph
  key: pytorch-fullgraph
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
  device: h200_18gb
  source_file_dependencies:
  - vllm/__init__.py
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Quantization
  key: quantization
-  timeout_in_minutes: 90
+  timeout_in_minutes: 60
  source_file_dependencies:
  - csrc/
  - vllm/model_executor/layers/quantization
@@ -23,7 +23,7 @@ steps:
 - label: Quantized Fusions
  key: quantized-fusions
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  source_file_dependencies:
  - tests/fusion
  - vllm/model_executor/layers/fusion
@@ -35,7 +35,7 @@ steps:
 - label: Quantized MoE Test (B200)
  key: quantized-moe-test-b200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 120
  working_dir: "/vllm-workspace/"
  device: b200-k8s
  source_file_dependencies:
@@ -53,7 +53,7 @@ steps:
 - label: Quantized Models Test
  key: quantized-models-test
-  timeout_in_minutes: 60
+  timeout_in_minutes: 50
  source_file_dependencies:
  - vllm/model_executor/layers/quantization
  - tests/models/quantization
@@ -3,7 +3,7 @@ depends_on:
  - image-build
 steps:
 - label: Rust Frontend OpenAI Coverage
-  timeout_in_minutes: 90
+  timeout_in_minutes: 30
  device: h200_18gb
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -38,7 +38,7 @@ steps:
  - pytest -v -s v1/sample/test_logprobs_e2e.py -k "test_prompt_logprobs_e2e_server"
 - label: Rust Frontend Serve/Admin Coverage
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
  device: h200_18gb
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -67,7 +67,7 @@ steps:
  - pytest -v -s entrypoints/serve/tokenize/test_tokenization.py -k "not tokenizer_info"
 - label: Rust Frontend Core Correctness
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: h200_18gb
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -81,7 +81,7 @@ steps:
  - pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
 - label: Rust Frontend Tool Use
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - rust/
@@ -95,7 +95,7 @@ steps:
  - pytest -v -s tool_use --ignore=tool_use/mistral --models llama3.2 -k "not test_response_format_with_tool_choice_required and not test_parallel_tool_calls_false and not test_tool_call_and_choice"
 - label: Rust Frontend Distributed
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
  num_devices: 4
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
@@ -4,7 +4,7 @@ steps:
 - label: Rust Frontend Cargo Style + Clippy
  key: rust-frontend-cargo-style-clippy
  depends_on: []
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: cpu-medium
  no_plugin: true
  source_file_dependencies:
@@ -18,7 +18,7 @@ steps:
 - label: Rust Frontend Cargo Tests
  key: rust-frontend-cargo-tests
  depends_on: []
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: cpu-medium
  no_plugin: true
  source_file_dependencies:
@@ -5,7 +5,7 @@ steps:
 - label: Samplers Test
  device: h200_35gb
  key: samplers-test
-  timeout_in_minutes: 75
+  timeout_in_minutes: 40
  source_file_dependencies:
  - vllm/model_executor/layers
  - vllm/sampling_metadata.py
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Spec Decode Eagle
  key: spec-decode-eagle
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/spec_decode/
@@ -15,7 +15,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 45
+      timeout_in_minutes: 60
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -29,7 +29,7 @@ steps:
 - label: Spec Decode Eagle Nightly B200
  key: spec-decode-eagle-nightly-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
  device: b200-k8s
  optional: true
  source_file_dependencies:
@@ -41,7 +41,7 @@ steps:
 - label: Spec Decode Speculators + MTP
  key: spec-decode-speculators-mtp
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/spec_decode/
@@ -82,7 +82,7 @@ steps:
 - label: Spec Decode Ngram + Suffix
  key: spec-decode-ngram-suffix
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/spec_decode/
@@ -93,7 +93,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 65
+      timeout_in_minutes: 55
      # TODO(akaratza): Test after Torch >= 2.12 bump
      soft_fail: true
      depends_on:
@@ -109,7 +109,7 @@ steps:
 - label: Spec Decode Draft Model
  key: spec-decode-draft-model
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/spec_decode/
@@ -120,7 +120,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 50
+      timeout_in_minutes: 55
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -134,7 +134,7 @@ steps:
 - label: Spec Decode Draft Model Nightly B200
  key: spec-decode-draft-model-nightly-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
  device: b200-k8s
  optional: true
  source_file_dependencies:
@@ -146,7 +146,7 @@ steps:
 - label: Speculators Correctness
  key: speculators-correctness
-  timeout_in_minutes: 60
+  timeout_in_minutes: 30
  device: h100
  optional: true
  num_devices: 1
@@ -159,7 +159,7 @@ steps:
    - pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test
 - label: Spec Decode MTP hybrid (B200)
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: b200-k8s
  optional: true
  source_file_dependencies:
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Weight Loading Multiple GPU  # 33min
  key: weight-loading-multiple-gpu
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  optional: true
@@ -70,6 +70,15 @@ endif()
 #
 set(TORCH_SUPPORTED_VERSION_CUDA "2.11.0")
 set(TORCH_SUPPORTED_VERSION_ROCM "2.11.0")
 # TORCH_NIGHTLY=1 builds run against unpinned nightly wheels, so the supported-
 # version check would always warn. Only treat it as a nightly build when the
 # value is exactly "1" (the bootstrap exports TORCH_NIGHTLY=0 by default, which
 # must NOT suppress the warning for normal builds).
 if (DEFINED ENV{TORCH_NIGHTLY} AND "$ENV{TORCH_NIGHTLY}" STREQUAL "1")
  set(TORCH_NIGHTLY_BUILD TRUE)
 else()
  set(TORCH_NIGHTLY_BUILD FALSE)
 endif()
 #
 # Try to find python package with an executable that exactly matches
@@ -175,7 +184,7 @@ endif()
 if (NOT HIP_FOUND AND NOT PYTORCH_FOUND_HIP AND CUDA_FOUND)
  set(VLLM_GPU_LANG "CUDA")
-  if (NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
+  if (NOT TORCH_NIGHTLY_BUILD AND NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
    message(WARNING "Pytorch version ${TORCH_SUPPORTED_VERSION_CUDA} "
      "expected for CUDA build, saw ${Torch_VERSION} instead.")
  endif()
@@ -188,7 +197,7 @@ elseif(HIP_FOUND OR PYTORCH_FOUND_HIP)
  enable_language(HIP)
  # ROCm 5.X and 6.X
-  if (ROCM_VERSION_DEV_MAJOR GREATER_EQUAL 5 AND
+  if (NOT TORCH_NIGHTLY_BUILD AND ROCM_VERSION_DEV_MAJOR GREATER_EQUAL 5 AND
      Torch_VERSION VERSION_LESS ${TORCH_SUPPORTED_VERSION_ROCM})
    message(WARNING "Pytorch version >= ${TORCH_SUPPORTED_VERSION_ROCM} "
      "expected for ROCm build, saw ${Torch_VERSION} instead.")
@@ -381,6 +390,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
    "csrc/libtorch_stable/cuda_view.cu"
    "csrc/libtorch_stable/cuda_utils_kernels.cu"
    "csrc/libtorch_stable/activation_kernels.cu"
    "csrc/libtorch_stable/ngram_embedding_kernels.cu"
    "csrc/libtorch_stable/quantization/activation_kernels.cu"
    "csrc/libtorch_stable/quantization/w8a8/int8/scaled_quant.cu"
    "csrc/libtorch_stable/quantization/w8a8/fp8/common.cu"
@@ -0,0 +1,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)
@@ -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 ""
@@ -39,7 +39,7 @@ else()
  FetchContent_Declare(
          vllm-flash-attn
          GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
-          GIT_TAG b3964b1d8b95d8e8447435668ab169a2700bab65
+          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
@@ -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);
 }
@@ -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,
+// Each block computes kEPB expert columns; wider blocks / kEPB > 1 are
-          int kHiddenDim>
+// selected per (shape, M) in invokeFp32RouterGemm (B300-tuned, see below).
-__global__ __launch_bounds__(128, 1) void fp32_router_gemm_kernel(
+// kTGroups > 1 splits the tokens across groups of kBlockSize threads within
-    float* out, InputT const* mat_a, float const* mat_b) {
+// 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 e_base = blockIdx.x * kEPB;
-  int const tid = threadIdx.x;
+  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] = {};
+  float acc[kMG][kEPB] = {};
-  __shared__ float sm_reduction[kNumTokens][kNumWarps];
+  __shared__ float sm_reduction[kNumTokens][kEPB][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;
  }
 #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];
+    float b_float[kEPB][VPT];
-    load_weight<VPT>(b_col + k_base, b_float);
+#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,
+      load_activation<InputT, VPT>(
-                                   a_float);
+          mat_a + (size_t)(m0 + m_idx) * kHiddenDim + k_base, a_float);
 #pragma unroll
-      for (int k = 0; k < VPT; k++) {
+      for (int e = 0; e < kEPB; e++) {
-        acc[m_idx] += a_float[k] * b_float[k];
+#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++) {
+  for (int m = 0; m < kMG; m++) {
-    float sum = acc[m];
+#pragma unroll
-    sum += __shfl_xor_sync(0xffffffff, sum, 16);
+    for (int e = 0; e < kEPB; e++) {
-    sum += __shfl_xor_sync(0xffffffff, sum, 8);
+      float sum = acc[m][e];
-    sum += __shfl_xor_sync(0xffffffff, sum, 4);
+      sum += __shfl_xor_sync(0xffffffff, sum, 16);
-    sum += __shfl_xor_sync(0xffffffff, sum, 2);
+      sum += __shfl_xor_sync(0xffffffff, sum, 8);
-    sum += __shfl_xor_sync(0xffffffff, sum, 1);
+      sum += __shfl_xor_sync(0xffffffff, sum, 4);
-    if (laneId == 0) sm_reduction[m][warpId] = sum;
+      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++) {
+    for (int w = 0; w < kNumWarps; w++) final_sum += sm_reduction[m][e][w];
-      float final_sum = 0.0f;
+    out[m * kNumExperts + e_base + e] = final_sum;
 #pragma unroll
      for (int w = 0; w < kNumWarps; w++) final_sum += sm_reduction[m][w];
      out[m * kNumExperts + n_idx] = final_sum;
    }
  }
 #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
  cudaTriggerProgrammaticLaunchCompletion();
 #endif
 }
 // ---------------------------------------------------------------------------
 // Launcher
 // ---------------------------------------------------------------------------
-template <typename InputT, int kNumTokens, int kNumExperts, int kHiddenDim>
+template <typename InputT, int kBlockSize, int kEPB, int kNumTokens,
-void invokeFp32RouterGemm(float* output, InputT const* mat_a,
+          int kNumExperts, int kHiddenDim, int kTGroups = 1>
-                          float const* mat_b, cudaStream_t stream) {
+static void launchFp32RouterGemm(float* output, InputT const* mat_a,
-  constexpr int kBlockSize = 128;
+                                 float const* mat_b, cudaStream_t stream) {
  static_assert(kNumExperts % kEPB == 0);
  cudaLaunchConfig_t config;
-  config.gridDim = kNumExperts;
+  config.gridDim = kNumExperts / kEPB;
-  config.blockDim = kBlockSize;
+  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,
+  cudaLaunchKernelEx(
-                     fp32_router_gemm_kernel<InputT, kBlockSize, kNumTokens,
+      &config,
-                                             kNumExperts, kHiddenDim>,
+      fp32_router_gemm_kernel<InputT, kBlockSize, kNumTokens, kEPB, kNumExperts,
-                     output, mat_a, mat_b);
+                              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
@@ -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(
@@ -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``
+ * Dense vs sparse row layout and index-branch processing are separate template
- * template bools (3 instantiations: dense <false,false>, sparse-profiling
+ * choices. Skip-index-topk reuse layers still have sparse rows and insert main
- * <true,false>, sparse-serving <true,true>), so the index slots, the V slots
+ * K/V cache entries, but compile away index_q/index_k work and index-cache
- * and the cache inserts fold away entirely on paths that don't use them. The
+ * writes.
 * dense layer passes no caches/index: norm+RoPE happens in place and the
 * generic ``Attention`` layer owns the cache write.
 *
 * 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
+// `kHasIndex`, `kProcessIndex`, and `kInsertKV` are compile-time template
-// branch decisions that distinguish the dense layer from the sparse layer
+// bools, so branch decisions that distinguish the dense layer from the sparse
-// (index slots, KV/index inserts, V slots) fold away per instantiation.
+// layer (index slots, KV/index inserts, V slots) fold away per instantiation.
-// Three instantiations are built: dense <false,false>, sparse-profiling
+// Slots per token:
 // <true,false> and sparse-serving <true,true>.  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)
+//     IQ: niq  only if kProcessIndex    (norm+RoPE)
-//     IK: 1    only if kIsSparse        (norm+RoPE; +index-cache insert)
+//     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].
+    // kv_cache strides (in elements) for logical shape [nb, nkv, bs, 2*128].
-    // The head_dim (last) dim is always innermost-contiguous (stride 1), so the
+    // The content (last) dim is always innermost-contiguous (stride 1), so the
-    // NHD/HND layout choice is fully captured by these four strides: NHD keeps
+    // NHD/HND layout choice is captured by the head/token strides.
-    // s_token < s_head, HND swaps them. dim_base addresses head_dim directly.
+    int64_t const kv_s_block, int64_t const kv_s_head, int64_t const kv_s_token,
-    int64_t const kv_s_block, int64_t const kv_s_kv, int64_t const kv_s_token,
+    int64_t const kv_s_dim) {
    int64_t const kv_s_head) {
 #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)
+      int64_t sm = -1;
-                             ? slot_mapping[tokenIdx]
+      if (isK || isV) {
-                             : (isIK ? index_slot_mapping[tokenIdx] : -1);
+        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 =
+          int64_t const off = b * kv_s_block + head * kv_s_head +
-              b * kv_s_block + kv * kv_s_kv + t * kv_s_token + head * kv_s_head;
+                              t * kv_s_token +
-          storeCacheElems<scalar_t, cache_t, kv_dt>(kv_cache + off + dim_base,
+                              (kv * kHeadDim + dim_base) * kv_s_dim;
-                                                    elems);
+          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_block, int64_t const kv_s_head, int64_t const kv_s_token,
-    int64_t const kv_s_head, bool const has_index, bool const insert_kv,
+    int64_t const kv_s_dim, bool const has_index, bool const insert_kv,
-    bool const fp8_idx, cudaStream_t stream) {
+    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)                              \
+  #define LAUNCH(HAS_INDEX, INSERT, PROCESS_INDEX, FP8, OUT_T)              \
-    fusedMiniMaxM3QNormRopeKVInsertKernel<scalar_t, cache_t, kv_dt, OUT_T,   \
+    fusedMiniMaxM3QNormRopeKVInsertKernel<                                  \
-                                          IS_SPARSE, INSERT, FP8>            \
+        scalar_t, cache_t, kv_dt, OUT_T, HAS_INDEX, INSERT, PROCESS_INDEX,  \
-        <<<grid, kBlockSize, 0, stream>>>(                                   \
+        FP8><<<grid, kBlockSize, 0, stream>>>(                              \
-            qkv, q_out, reinterpret_cast<OUT_T*>(index_q_out), q_norm_w,     \
+        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,        \
+        k_norm_w, iq_norm_w, ik_norm_w, cos_sin_cache, positions,           \
-            slot_mapping, index_slot_mapping, kv_cache,                      \
+        slot_mapping, index_slot_mapping, kv_cache,                         \
-            reinterpret_cast<OUT_T*>(index_cache), eps, rotary_dim,          \
+        reinterpret_cast<OUT_T*>(index_cache), eps, rotary_dim, num_tokens, \
-            num_tokens, nq, nkv, niq, block_size, kv_s_block, kv_s_kv,       \
+        nq, nkv, niq, block_size, kv_s_block, kv_s_head, kv_s_token,         \
-            kv_s_token, kv_s_head)
+        kv_s_dim)
  // clang-format on
 #endif
  if (has_index) {
-    if (insert_kv) {
+    if (!process_index) {
-      if (fp8_idx) {
+      if (insert_kv) {
-        LAUNCH(true, true, true, uint8_t);  // sparse serving, fp8 index outputs
+        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()) \
+      process_index                                                            \
-                : nullptr,                                                     \
+          ? reinterpret_cast<st const*>(index_q_norm_weight->data_ptr())       \
-      has_index ? reinterpret_cast<st const*>(index_k_norm_weight->data_ptr()) \
+          : nullptr,                                                           \
-                : 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*>(                            \
+      (insert_kv && process_index)                                             \
-                      effective_index_slot_mapping->data_ptr())                \
+          ? reinterpret_cast<int64_t const*>(                                  \
-                : nullptr,                                                     \
+                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, \
+      nkv, niq, static_cast<int>(block_size), kv_s_block, kv_s_head,           \
-      kv_s_head, has_index, insert_kv, fp8_idx, stream)
+      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(
+    if (process_index) {
-        !index_slot_mapping.has_value() ||
+      STD_TORCH_CHECK(
-            (index_slot_mapping->is_cuda() &&
+          !index_slot_mapping.has_value() ||
-             index_slot_mapping->scalar_type() ==
+              (index_slot_mapping->is_cuda() &&
-                 torch::headeronly::ScalarType::Long &&
+               index_slot_mapping->scalar_type() ==
-             index_slot_mapping->numel() == slot_mapping->numel()),
+                   torch::headeronly::ScalarType::Long &&
-        "index_slot_mapping must be int64 CUDA with slot_mapping length");
+               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(
+    if (process_index) {
-        index_cache.has_value() &&
+      STD_TORCH_CHECK(
-            (index_cache->scalar_type() == qkv.scalar_type() ||
+          index_cache.has_value() &&
-             index_cache->scalar_type() ==
+              (index_cache->scalar_type() == qkv.scalar_type() ||
-                 torch::headeronly::ScalarType::Float8_e4m3fn),
+               index_cache->scalar_type() ==
-        "insert mode requires index_cache matching qkv dtype or fp8 e4m3");
+                   torch::headeronly::ScalarType::Float8_e4m3fn),
-    STD_TORCH_CHECK(kv_cache->dim() == 5 && kv_cache->stride(4) == 1,
+          "insert mode requires index_cache matching qkv dtype or fp8 e4m3");
-                    "kv_cache must be [nb,2,bs,nkv,head_dim] with contiguous "
+    }
-                    "head_dim (stride(4)==1)");
+    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);
+    kv_s_dim = kv_cache->stride(3);
-    effective_index_slot_mapping = index_slot_mapping.has_value()
+    if (process_index) {
-                                       ? &index_slot_mapping.value()
+      effective_index_slot_mapping = index_slot_mapping.has_value()
-                                       : &slot_mapping.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(
+    STD_TORCH_CHECK(process_index,
-        has_index,
+                    "index_q_out requires index branch processing");
        "index_q_out requires the index branch (num_index_heads > 0)");
    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) ||
+      process_index &&
-      (index_q_out.has_value() && index_q_out->scalar_type() == kFp8);
+      ((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,
@@ -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 (int expert_id = get_local_expert_id(i, topk_ids, expert_map,
-    if (expert_id >= num_experts) {
+                                            num_experts, has_expert_map);
-      continue;
+        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 (int expert_id = get_local_expert_id(i, topk_ids, expert_map,
-    if (has_expert_map) {
+                                            num_experts, has_expert_map);
-      expert_id = expert_map[expert_id];
+        expert_id != -1) {
-      // filter invalid expert
+      int mask = token_mask == nullptr ? 1 : token_mask[i / topk_num];
-      if (expert_id == -1) continue;
+      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 (int expert_id = get_local_expert_id(i, topk_ids, expert_map,
-    if (has_expert_map) {
+                                            num_experts, has_expert_map);
-      expert_id = expert_map[expert_id];
+        expert_id != -1) {
-      // filter invalid expert
+      int32_t rank_post_pad =
-      if (expert_id == -1) continue;
+          tokens_cnts[tid * num_experts + expert_id] + cumsum[expert_id];
    }
    int32_t rank_post_pad =
        tokens_cnts[tid * num_experts + expert_id] + cumsum[expert_id];
-    if (token_mask == nullptr || token_mask[i / topk_num]) {
+      if (token_mask == nullptr || token_mask[i / topk_num]) {
-      sorted_token_ids[sorted_token_ids_offset + rank_post_pad] = i;
+        sorted_token_ids[sorted_token_ids_offset + rank_post_pad] = i;
-      ++tokens_cnts[tid * num_experts + expert_id];
+        ++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 (int expert_id = get_local_expert_id(i, topk_ids, expert_map,
-    if (expert_id >= num_experts) {
+                                            num_experts, has_expert_map);
-      continue;
+        expert_id != -1) {
-    }
+      if (token_mask == nullptr || token_mask[i / topk_num]) {
-
+        int32_t rank_post_pad = atomicAdd(
-    if (has_expert_map) {
+            &cumsum_buffer[(model_offset * (num_experts + 1)) + expert_id], 1);
-      expert_id = expert_map[expert_id];
+        sorted_token_ids[max_num_tokens_padded * model_offset + rank_post_pad] =
-      // filter invalid experts
+            i;
-      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;
    }
  }
 }
@@ -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>());
 }
@@ -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);
@@ -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
@@ -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);
@@ -93,6 +93,28 @@ To address this, manually trigger a build on Buildkite to accomplish two objecti
 <img width="60%" alt="Buildkite new build popup" src="https://github.com/user-attachments/assets/3b07f71b-bb18-4ca3-aeaf-da0fe79d315f" />
 </p>
 You can also trigger this build from the command line with
 [`.buildkite/scripts/trigger-ci-build.sh`](../../../.buildkite/scripts/trigger-ci-build.sh)
 (dry-run by default; pass `--execute` to actually trigger it).
 ## Test against PyTorch nightly
 The steps above test against a specific PyTorch RC/stable wheel pinned in the
 requirements files. To instead build and run the CI suite against the latest
 PyTorch **nightly** wheels, set the `TORCH_NIGHTLY=1` environment variable on
 the build (or apply the `ready-torch-nightly` label to the PR).
 When `TORCH_NIGHTLY=1`, the base CI image is built against PyTorch nightly
 (`image_build_torch_nightly.sh`, `PYTORCH_NIGHTLY=1`, CUDA 13.0) and tagged at
 the normal image tag, so the entire existing pipeline runs on nightly torch --
 there is no separate pipeline section to trigger. Combine it with `RUN_ALL=1`
 to run the full suite (the `ready-torch-nightly` label and
 `trigger-ci-build.sh --torch-nightly` both set this for you). This is the
 configuration to use for a scheduled "vLLM vs PyTorch nightly" run.
 Use `.buildkite/scripts/trigger-ci-build.sh --torch-nightly` to trigger it from
 the command line.
 ## Update all the different vLLM platforms
 Rather than attempting to update all vLLM platforms in a single pull request, it's more manageable
@@ -879,6 +879,55 @@ vllm serve Qwen/Qwen3-VL-30B-A3B-Instruct \
 Works with common video formats like MP4 when using OpenCV backends.
 #### GPU Video Decoding with DeepStream (NVDEC)
 By default vLLM decodes video on the CPU. On NVIDIA GPUs you can instead decode
 directly on the hardware video engine (NVDEC) with the DeepStream backend, which
 keeps decoding off the CPU and can significantly increase video throughput.
 Install the backend (Linux x86-64 only):
 ```bash
 pip install vllm[deepstream]
 ```
 The pip wheel bundles the DeepStream libraries but still relies on a few system
 packages that pip cannot install. On Ubuntu:
 ```bash
 apt-get install -y \
  gstreamer1.0-tools gstreamer1.0-plugins-base gstreamer1.0-plugins-good \
  gstreamer1.0-plugins-bad gstreamer1.0-libav \
  python3-gi python3-gst-1.0 libv4l-0 cuda-libraries-13-0
 ```
 Select the backend either with an environment variable:
 ```bash
 export VLLM_VIDEO_LOADER_BACKEND=deepstream
 vllm serve Qwen/Qwen3-VL-30B-A3B-Instruct
 ```
 or per request via `--media-io-kwargs`:
 ```bash
 vllm serve Qwen/Qwen3-VL-30B-A3B-Instruct \
  --media-io-kwargs '{"video": {"backend": "deepstream"}}'
 ```
 **Parameters:**
 - `pool_size`: Number of GPU decode workers in the process-wide decode pool
  (clamped to `[1, 16]`). When unset it defaults to
  `VLLM_MEDIA_LOADING_THREAD_COUNT` (default `8`). The pool is a singleton, so
  the first request's value wins.
 ```bash
 # Example: 12 decode workers
 vllm serve Qwen/Qwen3-VL-30B-A3B-Instruct \
  --media-io-kwargs '{"video": {"backend": "deepstream", "pool_size": 12}}'
 ```
 #### Pre-extracted Frame Sequences with `media_io_kwargs`
 When you extract video frames on the client side and send them as `video/jpeg` (base64-concatenated JPEG frames), you can preserve the original video metadata by using `media_io_kwargs` in your request. This enables more accurate video understanding by preserving temporal information that would otherwise be lost during client-side frame extraction.
@@ -351,6 +351,69 @@ print(response.choices[0].message.reasoning)
 print(response.choices[0].message.content)
 ```
 ## Suppressing Reasoning Output
 You can suppress reasoning content from API responses using the `include_reasoning` parameter. When set to `false`, reasoning tokens are still generated (so model quality is unaffected) but excluded from the response. This reduces network traffic without changing inference behavior.
 The parameter is supported in both the Chat Completions API and the Responses API, for streaming and non-streaming requests.
 When `include_reasoning=false`, vLLM also suppresses per-token metadata (logprobs and token IDs) to prevent leaking reasoning content through decoded token text in logprob entries or raw token IDs.
 ### Chat Completions API
 ```python
 from openai import OpenAI
 client = OpenAI(base_url="http://localhost:8000/v1", api_key="EMPTY")
 model = client.models.list().data[0].id
 # Reasoning is included by default (include_reasoning=True)
 response = client.chat.completions.create(
    model=model,
    messages=[{"role": "user", "content": "What is 15 * 37?"}],
    extra_body={"include_reasoning": False},
 )
 msg = response.choices[0].message
 assert msg.content  # Content is still present
 assert not getattr(msg, "reasoning", None)  # Reasoning is suppressed
 ```
 Streaming works the same way, reasoning deltas are omitted from chunks:
 ```python
 stream = client.chat.completions.create(
    model=model,
    messages=[{"role": "user", "content": "What is 15 * 37?"}],
    stream=True,
    extra_body={"include_reasoning": False},
 )
 for chunk in stream:
    delta = chunk.choices[0].delta
    # delta.reasoning will always be None
    if delta.content:
        print(delta.content, end="", flush=True)
 ```
 ### Responses API
 ```python
 from openai import OpenAI
 client = OpenAI(base_url="http://localhost:8000/v1", api_key="EMPTY")
 response = client.responses.create(
    model=client.models.list().data[0].id,
    input="What is 15 * 37?",
    include_reasoning=False,
 )
 # No "reasoning" items in output
 types = [item.type for item in response.output]
 assert "reasoning" not in types
 ```
 ## Limitations
 - The reasoning content is only available for online serving's chat completion endpoint (`/v1/chat/completions`), Anthropic Messages API (`/v1/messages`) and the Responses API (`/v1/responses`).
@@ -512,7 +512,7 @@ These models primarily accept the [`LLM.generate`](./generative_models.md#llmgen
 | `ChameleonForConditionalGeneration` | Chameleon | T + I | `facebook/chameleon-7b`, etc. | | ✅︎ |
 | `CheersForConditionalGeneration` | Cheers | T + I | `ai9stars/Cheers` | | ✅︎ |
 | `Cohere2VisionForConditionalGeneration` | Command A Vision, Command-A+ | T + I<sup>+</sup> | `CohereLabs/command-a-vision-07-2025`, `CohereLabs/command-a-plus-05-2026`, etc. | | ✅︎ |
-| `Cosmos3ForConditionalGeneration` | Cosmos3 (understanding tower) | T + I<sup>E+</sup> + V<sup>E+</sup> | `nvidia/Cosmos3-Nano` | | ✅︎ |
+| `Cosmos3ForConditionalGeneration` | Cosmos3 (understanding tower) | T + I<sup>E+</sup> + V<sup>E+</sup> | `nvidia/Cosmos3-Nano`, `nvidia/Cosmos3-Super` | | ✅︎ |
 | `DeepseekVLV2ForCausalLM` | DeepSeek-VL2 | T + I<sup>+</sup> | `deepseek-ai/deepseek-vl2-tiny`, `deepseek-ai/deepseek-vl2-small`, `deepseek-ai/deepseek-vl2`, etc. | | ✅︎ |
 | `DeepseekOCRForCausalLM` | DeepSeek-OCR | T + I<sup>+</sup> | `deepseek-ai/DeepSeek-OCR`, etc. | ✅︎ | ✅︎ |
 | `DeepseekOCR2ForCausalLM` | DeepSeek-OCR-2 | T + I<sup>+</sup> | `deepseek-ai/DeepSeek-OCR-2`, etc. | ✅︎ | ✅︎ |
@@ -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)
@@ -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.
@@ -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:
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -2167,7 +2167,7 @@ checksum = "11d3d7f243d5c5a8b9bb5d6dd2b1602c0cb0b9db1621bafc7ed66e35ff9fe092"
 [[package]]
 name = "llm-multimodal"
 version = "1.7.1"
-source = "git+https://github.com/smg-project/llm-multimodal?rev=7d74582aeaf0e4086a44964382655d22f1af0686#7d74582aeaf0e4086a44964382655d22f1af0686"
+source = "git+https://github.com/smg-project/llm-multimodal?rev=c8a29dcc755139fdc26185f400ea48c6d6d48273#c8a29dcc755139fdc26185f400ea48c6d6d48273"
 dependencies = [
 "anyhow",
 "base64 0.22.1",
@@ -2473,6 +2473,7 @@ dependencies = [
 "portable-atomic",
 "portable-atomic-util",
 "rawpointer",
 "serde",
 ]
 [[package]]
@@ -5093,6 +5094,7 @@ dependencies = [
 "llm-multimodal",
 "minijinja",
 "minijinja-contrib",
 "ndarray 0.17.2",
 "oss-harmony",
 "paste",
 "reqwest",
@@ -53,12 +53,12 @@ hyper-util = { version = "0.1.20", features = [
 indexmap = "2.13.0"
 itertools = "0.14.0"
 libc = "0.2.177"
-llm-multimodal = { git = "https://github.com/smg-project/llm-multimodal", rev = "7d74582aeaf0e4086a44964382655d22f1af0686" }
+llm-multimodal = { git = "https://github.com/smg-project/llm-multimodal", rev = "c8a29dcc755139fdc26185f400ea48c6d6d48273" }
 mimalloc = "0.1.52"
 minijinja = { version = "2.0", features = ["unstable_machinery", "json", "builtins", "loader", "loop_controls", "preserve_order"] }
 minijinja-contrib = { version = "2.0", features = ["pycompat"] }
 native-tls-vendored = { package = "native-tls", version = "0.2.18", features = ["vendored"] }
-ndarray = { version = "0.16.1", features = ["serde"] }
+ndarray = { version = "0.17", features = ["serde"] }
 openai-harmony = { package = "oss-harmony", git = "https://github.com/oss-harmony/harmony", tag = "v0.0.11", default-features = false }
 openai-protocol = "1.6.0"
 openssl = "0.10"
@@ -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
@@ -42,6 +42,7 @@ anyhow.workspace = true
 bytes.workspace = true
 clap.workspace = true
 expect-test.workspace = true
 ndarray.workspace = true
 paste.workspace = true
 rmp-serde.workspace = true
 serial_test.workspace = true
@@ -10,7 +10,7 @@ use crate::backend::{
    NewChatOutputProcessorOptions,
 };
 use crate::error::Result;
-use crate::multimodal::MultimodalModelInfo;
+use crate::multimodal::{MultimodalConfigFiles, MultimodalModelInfo};
 use crate::output::{
    DefaultChatOutputProcessor, HarmonyChatOutputProcessor, validate_harmony_parser_overrides,
 };
@@ -46,8 +46,12 @@ impl HfChatBackend {
            MultimodalModelInfo::from_paths(
                model_id.clone(),
                (!model_type.is_empty()).then_some(model_type.to_string()),
-                files.config_path.as_deref(),
+                MultimodalConfigFiles {
-                files.preprocessor_config_path.as_deref(),
+                    config: files.config_path.as_deref(),
                    preprocessor_config: files.preprocessor_config_path.as_deref(),
                    video_preprocessor_config: files.video_preprocessor_config_path.as_deref(),
                    processor_config: files.processor_config_path.as_deref(),
                },
                tokenizer.clone(),
            )?
        };
@@ -139,8 +143,11 @@ pub(super) async fn load_model_backends(
 fn resolve_multimodal_render_info(
    info: Option<&MultimodalModelInfo>,
 ) -> Option<MultimodalRenderInfo> {
    use llm_multimodal::Modality;
    info.map(|info| MultimodalRenderInfo {
-        placeholder_token: info.placeholder_token().to_string(),
+        image_token: info.placeholder_token(Modality::Image).map(str::to_string),
        video_token: info.placeholder_token(Modality::Video).map(str::to_string),
    })
 }
@@ -192,6 +199,8 @@ mod tests {
            tokenizer_config_path: Some(tokenizer_config_path),
            generation_config_path: None,
            preprocessor_config_path: None,
            video_preprocessor_config_path: None,
            processor_config_path: None,
            chat_template_path: None,
            config_path: Some(config_path),
        }
@@ -1,5 +1,5 @@
 use thiserror::Error;
-use thiserror_ext::Macro;
+use thiserror_ext::{AsReport as _, Macro};
 type BoxedError = Box<dyn std::error::Error + Send + Sync>;
@@ -18,6 +18,8 @@ pub enum Error {
    UnsupportedMultimodalRenderer,
    #[error("unsupported multimodal content: {0}")]
    UnsupportedMultimodalContent(&'static str),
    #[error("`{modality}` input is not supported by this model")]
    UnsupportedModality { modality: String },
    #[error("multimodal preprocessing error: {0}")]
    Multimodal(#[message] String),
    #[error("{kind} parsing is not available for model `{model_id}`")]
@@ -80,11 +82,39 @@ impl Error {
        match self {
            Self::PromptTooLong { .. } => true,
            Self::Text(error) => error.is_request_validation_error(),
            Self::UnsupportedMultimodalRenderer
            | Self::UnsupportedMultimodalContent(_)
            | Self::UnsupportedModality { .. } => true,
            _ => false,
        }
    }
 }
 impl From<llm_multimodal::MediaConnectorError> for Error {
    fn from(error: llm_multimodal::MediaConnectorError) -> Self {
        Self::Multimodal(error.to_report_string())
    }
 }
 impl From<llm_multimodal::MultiModalError> for Error {
    fn from(error: llm_multimodal::MultiModalError) -> Self {
        Self::Multimodal(error.to_report_string())
    }
 }
 impl From<llm_multimodal::TransformError> for Error {
    fn from(error: llm_multimodal::TransformError) -> Self {
        Self::Multimodal(error.to_report_string())
    }
 }
 impl From<llm_multimodal::registry::ModelRegistryError> for Error {
    fn from(error: llm_multimodal::registry::ModelRegistryError) -> Self {
        Self::Multimodal(error.to_report_string())
    }
 }
 /// Format the available-parser suffix used in user-facing error messages.
 fn available_parser_hint(available_names: &[String]) -> String {
    if available_names.is_empty() {
@@ -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>,
    },
 }
@@ -0,0 +1,446 @@
 //! Prompt placeholder expansion shared across modalities.
 use std::collections::{HashMap, VecDeque};
 use llm_multimodal::{Modality, PromptReplacement};
 use vllm_engine_core_client::protocol::multimodal::PlaceholderRange;
 use vllm_engine_core_client::protocol::tensor::WireTensor;
 use super::PreparedMedia;
 use crate::error::{Error, Result, bail_multimodal};
 /// One modality's queue of pending placeholder replacements for prompt
 /// expansion.
 struct ExpansionLane<'a> {
    modality: Modality,
    marker_token_id: u32,
    embed_token_id: u32,
    placeholder_token: String,
    replacements: VecDeque<&'a PromptReplacement>,
 }
 impl<'a> ExpansionLane<'a> {
    fn from_prepared(media: &'a PreparedMedia) -> Option<Self> {
        if media.replacements.is_empty() {
            return None;
        }
        Some(Self {
            modality: media.modality,
            marker_token_id: media.placeholder.marker_token_id,
            embed_token_id: media.placeholder.embed_token_id,
            placeholder_token: media.placeholder.token.clone(),
            replacements: media.replacements.iter().collect(),
        })
    }
 }
 /// Replace rendered placeholder markers with model-specific replacement
 /// tokens across all modalities in one left-to-right pass.
 ///
 /// Each prepared modality consumes its own marker occurrences in order,
 /// matching the original media-part order within that modality; markers of
 /// different modalities may interleave freely.
 ///
 /// The returned ranges point into the already-expanded prompt, grouped per
 /// modality in item order.
 pub(super) fn expand_prompt_token_ids(
    prompt_token_ids: &mut Vec<u32>,
    prepared: &[PreparedMedia],
 ) -> Result<HashMap<Modality, Vec<PlaceholderRange>>> {
    let mut lanes = prepared.iter().filter_map(ExpansionLane::from_prepared).collect::<Vec<_>>();
    if lanes.is_empty() {
        return Ok(HashMap::new());
    }
    let replacement_growth = lanes
        .iter()
        .flat_map(|lane| lane.replacements.iter())
        .fold(0usize, |total, replacement| {
            total.saturating_add(replacement.tokens.len().saturating_sub(1))
        });
    let expanded_len = prompt_token_ids.len().saturating_add(replacement_growth);
    let mut expanded = Vec::with_capacity(expanded_len);
    let mut ranges = HashMap::<Modality, Vec<PlaceholderRange>>::new();
    for &token in prompt_token_ids.iter() {
        let lane = lanes
            .iter_mut()
            .find(|lane| lane.marker_token_id == token && !lane.replacements.is_empty());
        let Some(lane) = lane else {
            expanded.push(token);
            continue;
        };
        let replacement = lane.replacements.pop_front().expect("lane queue is non-empty");
        debug_assert_eq!(replacement.modality, lane.modality);
        if replacement.tokens.is_empty() {
            bail_multimodal!(
                "placeholder token `{}` expanded to no tokens",
                lane.placeholder_token
            );
        }
        let replacement_len = replacement.tokens.len();
        let is_embed = {
            let mask = replacement
                .tokens
                .iter()
                .map(|&token| token as u32 == lane.embed_token_id)
                .collect::<Vec<_>>();
            WireTensor::from_bool(vec![replacement_len], mask).map_err(Error::Multimodal)?
        };
        let expanded_offset = expanded.len();
        expanded.extend(replacement.tokens.iter().map(|&token| token as u32));
        ranges.entry(lane.modality).or_default().push(PlaceholderRange {
            offset: expanded_offset,
            length: replacement_len,
            is_embed: Some(is_embed),
        });
    }
    for lane in &lanes {
        if !lane.replacements.is_empty() {
            bail_multimodal!(
                "placeholder token `{}` was not found in tokenized prompt for {} remaining `{}` item(s)",
                lane.placeholder_token,
                lane.replacements.len(),
                lane.modality
            );
        }
    }
    *prompt_token_ids = expanded;
    Ok(ranges)
 }
 #[cfg(test)]
 mod tests {
    use llm_multimodal::TokenId;
    use vllm_engine_core_client::protocol::tensor::WireArrayData;
    use super::super::tests::{
        LLAMA4_IMAGE_END_ID, LLAMA4_IMAGE_ID, LLAMA4_IMAGE_START_ID, LLAMA4_PATCH_ID,
        LLAMA4_TILE_X_SEPARATOR_ID, LLAMA4_TILE_Y_SEPARATOR_ID, QWEN3_IMAGE_PAD_ID,
        QWEN3_VIDEO_PAD_ID,
    };
    use super::super::{PreparedMedia, ResolvedPlaceholder};
    use super::*;
    /// Build prepared media directly from placeholder token IDs.
    fn prepared_media(
        modality: Modality,
        placeholder_token: &str,
        marker_token_id: u32,
        embed_token_id: u32,
        replacements: Vec<PromptReplacement>,
    ) -> PreparedMedia {
        PreparedMedia {
            modality,
            placeholder: ResolvedPlaceholder {
                token: placeholder_token.to_string(),
                marker_token_id,
                embed_token_id,
            },
            replacements,
            items: Vec::new(),
        }
    }
    /// Llama4 image prepared media: the `<|image|>` marker expands to
    /// sequences whose embed positions are the `<|patch|>` tokens.
    fn llama4_prepared(replacements: Vec<PromptReplacement>) -> PreparedMedia {
        prepared_media(
            Modality::Image,
            "<|image|>",
            LLAMA4_IMAGE_ID,
            LLAMA4_PATCH_ID,
            replacements,
        )
    }
    fn qwen3_image_prepared(replacements: Vec<PromptReplacement>) -> PreparedMedia {
        prepared_media(
            Modality::Image,
            "<|image_pad|>",
            QWEN3_IMAGE_PAD_ID,
            QWEN3_IMAGE_PAD_ID,
            replacements,
        )
    }
    fn qwen3_video_prepared(replacements: Vec<PromptReplacement>) -> PreparedMedia {
        prepared_media(
            Modality::Video,
            "<|video_pad|>",
            QWEN3_VIDEO_PAD_ID,
            QWEN3_VIDEO_PAD_ID,
            replacements,
        )
    }
    fn llama4_single_tile_replacement() -> PromptReplacement {
        PromptReplacement::sequence(
            Modality::Image,
            "<|image|>",
            vec![
                LLAMA4_IMAGE_START_ID as TokenId,
                LLAMA4_IMAGE_ID as TokenId,
                LLAMA4_PATCH_ID as TokenId,
                LLAMA4_PATCH_ID as TokenId,
                LLAMA4_IMAGE_END_ID as TokenId,
            ],
        )
    }
    fn llama4_multi_tile_replacement() -> PromptReplacement {
        PromptReplacement::sequence(
            Modality::Image,
            "<|image|>",
            vec![
                LLAMA4_IMAGE_START_ID as TokenId,
                LLAMA4_PATCH_ID as TokenId,
                LLAMA4_TILE_X_SEPARATOR_ID as TokenId,
                LLAMA4_PATCH_ID as TokenId,
                LLAMA4_TILE_Y_SEPARATOR_ID as TokenId,
                LLAMA4_IMAGE_ID as TokenId,
                LLAMA4_PATCH_ID as TokenId,
                LLAMA4_IMAGE_END_ID as TokenId,
            ],
        )
    }
    fn assert_bool_mask(range: &PlaceholderRange, expected: &[bool]) {
        let tensor = range.is_embed.as_ref().expect("is_embed mask");
        assert_eq!(tensor.dtype, "bool");
        assert_eq!(tensor.shape, vec![expected.len()]);
        assert_eq!(
            tensor.data,
            WireArrayData::RawView(expected.iter().map(|value| u8::from(*value)).collect())
        );
    }
    #[test]
    fn expand_prompt_tokens_marks_only_llama4_patch_tokens_as_embed() {
        let mut prompt_token_ids = vec![1, LLAMA4_IMAGE_ID, 2];
        let prepared = vec![llama4_prepared(vec![llama4_multi_tile_replacement()])];
        let ranges = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap();
        let ranges = &ranges[&Modality::Image];
        assert_eq!(
            prompt_token_ids,
            vec![
                1,
                LLAMA4_IMAGE_START_ID,
                LLAMA4_PATCH_ID,
                LLAMA4_TILE_X_SEPARATOR_ID,
                LLAMA4_PATCH_ID,
                LLAMA4_TILE_Y_SEPARATOR_ID,
                LLAMA4_IMAGE_ID,
                LLAMA4_PATCH_ID,
                LLAMA4_IMAGE_END_ID,
                2,
            ]
        );
        assert_eq!(ranges[0].offset, 1);
        assert_eq!(ranges[0].length, 8);
        assert_bool_mask(
            &ranges[0],
            &[false, true, false, true, false, false, true, false],
        );
    }
    #[test]
    fn expand_prompt_tokens_errors_when_placeholder_missing() {
        let mut prompt_token_ids = vec![1, 2, 3];
        let prepared = vec![llama4_prepared(vec![llama4_single_tile_replacement()])];
        let error = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap_err();
        assert!(matches!(error, Error::Multimodal(message) if message.contains("not found")));
    }
    #[test]
    fn expand_prompt_tokens_ignores_empty_replacements() {
        let mut prompt_token_ids = vec![1, LLAMA4_IMAGE_ID, 2];
        let original_prompt_token_ids = prompt_token_ids.clone();
        let prepared = vec![llama4_prepared(Vec::new())];
        let ranges = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap();
        assert!(ranges.is_empty());
        assert_eq!(prompt_token_ids, original_prompt_token_ids);
    }
    #[test]
    fn expand_prompt_tokens_leaves_prompt_unchanged_when_later_placeholder_missing() {
        let mut prompt_token_ids = vec![1, LLAMA4_IMAGE_ID, 2];
        let original_prompt_token_ids = prompt_token_ids.clone();
        let prepared = vec![llama4_prepared(vec![
            llama4_single_tile_replacement(),
            llama4_single_tile_replacement(),
        ])];
        let error = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap_err();
        assert!(matches!(error, Error::Multimodal(message) if message.contains("not found")));
        assert_eq!(prompt_token_ids, original_prompt_token_ids);
    }
    #[test]
    fn expand_prompt_tokens_errors_when_replacement_is_empty() {
        let mut prompt_token_ids = vec![1, LLAMA4_IMAGE_ID, 2];
        let original_prompt_token_ids = prompt_token_ids.clone();
        let prepared = vec![llama4_prepared(vec![PromptReplacement::sequence(
            Modality::Image,
            "<|image|>",
            Vec::new(),
        )])];
        let error = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap_err();
        assert!(
            matches!(error, Error::Multimodal(message) if message.contains("expanded to no tokens"))
        );
        assert_eq!(prompt_token_ids, original_prompt_token_ids);
    }
    #[test]
    fn expand_prompt_tokens_skips_llama4_image_marker_inside_replacement() {
        let mut prompt_token_ids = vec![1, LLAMA4_IMAGE_ID, 2, LLAMA4_IMAGE_ID, 3];
        let prepared = vec![llama4_prepared(vec![
            llama4_single_tile_replacement(),
            llama4_single_tile_replacement(),
        ])];
        let ranges = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap();
        let ranges = &ranges[&Modality::Image];
        assert_eq!(
            prompt_token_ids,
            vec![
                1,
                LLAMA4_IMAGE_START_ID,
                LLAMA4_IMAGE_ID,
                LLAMA4_PATCH_ID,
                LLAMA4_PATCH_ID,
                LLAMA4_IMAGE_END_ID,
                2,
                LLAMA4_IMAGE_START_ID,
                LLAMA4_IMAGE_ID,
                LLAMA4_PATCH_ID,
                LLAMA4_PATCH_ID,
                LLAMA4_IMAGE_END_ID,
                3,
            ]
        );
        assert_eq!(ranges[0].offset, 1);
        assert_eq!(ranges[0].length, 5);
        assert_bool_mask(&ranges[0], &[false, false, true, true, false]);
        assert_eq!(ranges[1].offset, 7);
        assert_eq!(ranges[1].length, 5);
        assert_bool_mask(&ranges[1], &[false, false, true, true, false]);
    }
    #[test]
    fn expand_prompt_tokens_interleaves_image_and_video_prepared_media() {
        let mut prompt_token_ids = vec![
            1,
            QWEN3_IMAGE_PAD_ID,
            2,
            QWEN3_VIDEO_PAD_ID,
            3,
            QWEN3_IMAGE_PAD_ID,
            4,
        ];
        let prepared = vec![
            qwen3_image_prepared(vec![
                PromptReplacement::repeated(
                    Modality::Image,
                    "<|image_pad|>",
                    QWEN3_IMAGE_PAD_ID as TokenId,
                    2,
                ),
                PromptReplacement::repeated(
                    Modality::Image,
                    "<|image_pad|>",
                    QWEN3_IMAGE_PAD_ID as TokenId,
                    3,
                ),
            ]),
            qwen3_video_prepared(vec![PromptReplacement::repeated(
                Modality::Video,
                "<|video_pad|>",
                QWEN3_VIDEO_PAD_ID as TokenId,
                4,
            )]),
        ];
        let ranges = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap();
        assert_eq!(
            prompt_token_ids,
            vec![
                1,
                QWEN3_IMAGE_PAD_ID,
                QWEN3_IMAGE_PAD_ID,
                2,
                QWEN3_VIDEO_PAD_ID,
                QWEN3_VIDEO_PAD_ID,
                QWEN3_VIDEO_PAD_ID,
                QWEN3_VIDEO_PAD_ID,
                3,
                QWEN3_IMAGE_PAD_ID,
                QWEN3_IMAGE_PAD_ID,
                QWEN3_IMAGE_PAD_ID,
                4,
            ]
        );
        let image_ranges = &ranges[&Modality::Image];
        assert_eq!(image_ranges[0].offset, 1);
        assert_eq!(image_ranges[0].length, 2);
        assert_bool_mask(&image_ranges[0], &[true, true]);
        assert_eq!(image_ranges[1].offset, 9);
        assert_eq!(image_ranges[1].length, 3);
        assert_bool_mask(&image_ranges[1], &[true, true, true]);
        let video_ranges = &ranges[&Modality::Video];
        assert_eq!(video_ranges[0].offset, 4);
        assert_eq!(video_ranges[0].length, 4);
        assert_bool_mask(&video_ranges[0], &[true, true, true, true]);
    }
    #[test]
    fn expand_prompt_tokens_error_names_modality_with_leftover_replacements() {
        let mut prompt_token_ids = vec![1, QWEN3_IMAGE_PAD_ID, 2];
        let original_prompt_token_ids = prompt_token_ids.clone();
        let prepared = vec![
            qwen3_image_prepared(vec![PromptReplacement::repeated(
                Modality::Image,
                "<|image_pad|>",
                QWEN3_IMAGE_PAD_ID as TokenId,
                2,
            )]),
            qwen3_video_prepared(vec![PromptReplacement::repeated(
                Modality::Video,
                "<|video_pad|>",
                QWEN3_VIDEO_PAD_ID as TokenId,
                4,
            )]),
        ];
        let error = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap_err();
        assert!(matches!(
            error,
            Error::Multimodal(message)
                if message.contains("<|video_pad|>") && message.contains("`video`")
        ));
        assert_eq!(prompt_token_ids, original_prompt_token_ids);
    }
 }
@@ -0,0 +1,141 @@
 //! Image-modality preparation: batch preprocessing and per-item feature
 //! build.
 use std::sync::Arc;
 use itertools::izip;
 use llm_multimodal::{FieldLayout, ImageFrame, Modality, PreprocessedEncoderInputs};
 use vllm_engine_core_client::protocol::dtype::ModelDtype;
 use vllm_engine_core_client::protocol::multimodal::{
    MmBatchedField, MmField, MmFieldElem, MmFlatField, MmKwargsItem, MmSharedField, MmSlice,
    SliceSpec,
 };
 use super::{ModalitySupport, MultimodalModelInfo, PreparedItem, PreparedMedia, tensor};
 use crate::error::{Error, Result, bail_multimodal, multimodal};
 impl MultimodalModelInfo {
    /// Preprocess all fetched image frames as one batch and build per-item
    /// features.
    pub(super) async fn prepare_images(
        &self,
        frames: Vec<Arc<ImageFrame>>,
        uuids: Vec<Option<String>>,
        model_dtype: ModelDtype,
    ) -> Result<PreparedMedia> {
        let support = self.image.as_ref().ok_or_else(|| Error::UnsupportedModality {
            modality: Modality::Image.to_string(),
        })?;
        let preprocessed = self.preprocess_images(support, &frames).await?;
        let replacements =
            self.spec
                .prompt_replacements_for(&self.context, &preprocessed, Modality::Image)?;
        if replacements.len() != frames.len() {
            bail_multimodal!(
                "number of image prompt replacements {} does not match number of images {}",
                replacements.len(),
                frames.len()
            );
        }
        let items = self.build_image_items(preprocessed, &frames, uuids, model_dtype)?;
        Ok(PreparedMedia {
            modality: Modality::Image,
            placeholder: support.placeholder.clone(),
            replacements,
            items,
        })
    }
    /// Preprocess fetched image frames with the model's resolved vision
    /// processor.
    ///
    /// The processor work is CPU-heavy relative to request wiring, so it runs
    /// in a blocking task and returns owned tensors ready for wire
    /// conversion.
    async fn preprocess_images(
        &self,
        support: &ModalitySupport,
        image_frames: &[Arc<ImageFrame>],
    ) -> Result<PreprocessedEncoderInputs> {
        let config = support.config.clone();
        let processor = support.processor;
        let images = image_frames.iter().map(|frame| frame.data().clone()).collect::<Vec<_>>();
        // TODO: is it still necessary given that we've already in a dedicated runtime?
        tokio::task::spawn_blocking(move || Ok(processor.preprocess(&images, &config)?))
            .await
            .map_err(|error| multimodal!("image preprocessing task failed: {error}"))?
    }
    /// Convert one batch of preprocessed image tensors into per-item engine
    /// kwargs.
    ///
    /// Tensor fields are sliced per item according to the model spec's field
    /// layout declarations.
    fn build_image_items(
        &self,
        preprocessed: PreprocessedEncoderInputs,
        frames: &[Arc<ImageFrame>],
        uuids: Vec<Option<String>>,
        model_dtype: ModelDtype,
    ) -> Result<Vec<PreparedItem>> {
        let len = frames.len();
        let tensors = tensor::collect_tensors(preprocessed, "pixel_values", model_dtype)?;
        let mut items = Vec::with_capacity(len);
        for (index, (frame, uuid)) in izip!(frames, uuids).enumerate() {
            let mut data = MmKwargsItem::new();
            for (key, tensor) in &tensors {
                let keep_on_cpu = self.spec.keep_on_cpu_keys.contains(key);
                let (value, field) = match self.spec.field_layouts.get(key) {
                    Some(FieldLayout::Batched) => (
                        tensor.batched_value_at(index)?,
                        MmField::Batched(MmBatchedField { keep_on_cpu }),
                    ),
                    Some(FieldLayout::Flat { sizes_key }) => {
                        let sizes = tensors.get(sizes_key).ok_or_else(|| {
                            multimodal!("flat tensor sizes key `{sizes_key}` is missing")
                        })?;
                        let (start, end) = tensor::flat_range_for_index(sizes, sizes_key, index)?;
                        (
                            tensor.flat_value_range(start, end)?,
                            MmField::Flat(MmFlatField {
                                slices: vec![MmSlice::Slice(SliceSpec {
                                    start: Some(0),
                                    stop: Some((end - start) as isize),
                                    step: None,
                                })],
                                dim: 0,
                                keep_on_cpu,
                            }),
                        )
                    }
                    None => (
                        tensor.clone(),
                        MmField::Shared(MmSharedField {
                            batch_size: len,
                            keep_on_cpu,
                        }),
                    ),
                };
                data.insert(
                    key.clone(),
                    MmFieldElem {
                        data: Some(value.try_into()?),
                        field,
                    },
                );
            }
            items.push(PreparedItem {
                data,
                hash: frame.hash.clone(),
                uuid,
            });
        }
        Ok(items)
    }
 }
@@ -1,7 +1,7 @@
 use std::collections::HashMap;
 use half::{bf16, f16};
-use llm_multimodal::{ModelSpecificValue, PreprocessedEncoderInputs as PreprocessedImages};
+use llm_multimodal::{ModelSpecificValue, PreprocessedEncoderInputs};
 use vllm_engine_core_client::protocol::dtype::ModelDtype;
 use vllm_engine_core_client::protocol::multimodal::MmKwargValue as ProtocolKwargValue;
 use vllm_engine_core_client::protocol::tensor::{ShapeExt as _, WireTensor};
@@ -25,25 +25,31 @@ pub(super) enum KwargValue {
    Passthrough(ProtocolKwargValue),
 }
-/// Collect `pixel_values` and model-specific outputs into one tensor map.
+/// Collect the primary encoder input and model-specific outputs into one
 /// tensor map.
 ///
 /// `primary_key` names the encoder-input tensor as the model's forward kwargs
 /// expect it (e.g. `pixel_values` for images, `pixel_values_videos` for
 /// videos).
 pub(super) fn collect_tensors(
-    preprocessed: PreprocessedImages,
+    preprocessed: PreprocessedEncoderInputs,
    primary_key: &str,
    float_dtype: ModelDtype,
 ) -> Result<HashMap<String, KwargValue>> {
-    let PreprocessedImages {
+    let PreprocessedEncoderInputs {
        encoder_input,
        model_specific,
        ..
    } = preprocessed;
-    let pixel_values = {
+    let primary_value = {
        let shape = encoder_input.shape().to_vec();
        let data = encoder_input.into_iter().collect();
        KwargValue::from_f32_tensor(data, shape, float_dtype)?
    };
    let mut tensors = HashMap::new();
-    tensors.insert("pixel_values".to_string(), pixel_values);
+    tensors.insert(primary_key.to_string(), primary_value);
    for (key, value) in model_specific {
        tensors.insert(key, KwargValue::from_model_specific(value, float_dtype)?);
    }
@@ -124,10 +130,22 @@ impl TryFrom<KwargValue> for ProtocolKwargValue {
 }
 impl KwargValue {
-    /// Extract one image from a batched tensor field.
+    /// First-axis length for tensor values; `None` for passthrough kwargs.
    pub(super) fn first_dim(&self) -> Option<usize> {
        match self {
            Self::F32Tensor { shape, .. }
            | Self::F16Tensor { shape, .. }
            | Self::Bf16Tensor { shape, .. }
            | Self::I64Tensor { shape, .. }
            | Self::U32Tensor { shape, .. } => shape.first().copied(),
            Self::Passthrough(_) => None,
        }
    }
    /// Extract one media item from a batched tensor field.
    ///
-    /// Batched fields use their first axis as image index and drop that axis in
+    /// Batched fields use their first axis as media-item index and drop that
-    /// the per-feature value, matching vLLM's batched-field semantics.
+    /// axis in the per-feature value, matching vLLM's batched-field semantics.
    pub(super) fn batched_value_at(&self, index: usize) -> Result<Self> {
        match self {
            Self::F32Tensor { data, shape } => {
@@ -154,9 +172,9 @@ impl KwargValue {
        }
    }
-    /// Extract one image's variable-length range from a flat tensor field.
+    /// Extract one media item's variable-length range from a flat tensor field.
    ///
-    /// Flat fields keep the first axis as the sliced length for this image.
+    /// Flat fields keep the first axis as the sliced length for this item.
    pub(super) fn flat_value_range(&self, start: usize, end: usize) -> Result<Self> {
        match self {
            Self::F32Tensor { data, shape } => {
@@ -184,10 +202,10 @@ impl KwargValue {
    }
 }
-/// Compute the first-axis range for one image in a flat tensor.
+/// Compute the first-axis range for one media item in a flat tensor.
 ///
 /// `sizes_key` names a companion tensor whose entries are cumulative slice
-/// sizes per image.
+/// sizes per media item.
 pub(super) fn flat_range_for_index(
    sizes: &KwargValue,
    sizes_key: &str,
@@ -195,7 +213,7 @@ pub(super) fn flat_range_for_index(
 ) -> Result<(usize, usize)> {
    let sizes = tensor_as_usize_vec(sizes)?;
    let size = *sizes.get(index).ok_or_else(|| {
-        multimodal!("flat tensor sizes key `{sizes_key}` has no entry for image {index}")
+        multimodal!("flat tensor sizes key `{sizes_key}` has no entry for media item {index}")
    })?;
    let start = sizes[..index].iter().sum::<usize>();
    Ok((start, start + size))
@@ -0,0 +1,316 @@
 //! Video-modality preparation: per-clip preprocessing, config resolution,
 //! and per-item feature build.
 use std::sync::Arc;
 use itertools::izip;
 use llm_multimodal::{FieldLayout, Modality, PreprocessedEncoderInputs, VideoClip};
 use thiserror_ext::AsReport as _;
 use tracing::warn;
 use vllm_engine_core_client::protocol::dtype::ModelDtype;
 use vllm_engine_core_client::protocol::multimodal::{
    MmBatchedField, MmField, MmFieldElem, MmFlatField, MmKwargsItem, MmSharedField, MmSlice,
    SliceSpec,
 };
 use super::{ModalitySupport, MultimodalModelInfo, PreparedItem, PreparedMedia, tensor};
 use crate::error::{Error, Result, bail_multimodal, multimodal};
 /// Forward-kwargs name of the primary video encoder input.
 ///
 /// Video-capable vLLM models read `pixel_values_videos` alongside
 /// `video_grid_thw`, mirroring the HF processor output naming.
 const VIDEO_PRIMARY_KEY: &str = "pixel_values_videos";
 impl MultimodalModelInfo {
    /// Preprocess fetched video clips one at a time and build per-item
    /// features.
    ///
    /// Unlike images, each clip runs through the preprocessor independently
    /// (a batch of one), so its tensors are complete per item and need no
    /// cross-item slicing.
    pub(super) async fn prepare_videos(
        &self,
        clips: Vec<Arc<VideoClip>>,
        uuids: Vec<Option<String>>,
        model_dtype: ModelDtype,
    ) -> Result<PreparedMedia> {
        let support = self.video.as_ref().ok_or_else(|| Error::UnsupportedModality {
            modality: Modality::Video.to_string(),
        })?;
        let mut replacements = Vec::with_capacity(clips.len());
        let mut items = Vec::with_capacity(clips.len());
        for (clip, uuid) in izip!(&clips, uuids) {
            let preprocessed = self.preprocess_video_clip(support, Arc::clone(clip)).await?;
            let mut clip_replacements =
                self.spec
                    .prompt_replacements_for(&self.context, &preprocessed, Modality::Video)?;
            if clip_replacements.len() != 1 {
                bail_multimodal!(
                    "expected exactly one prompt replacement per video clip, got {}",
                    clip_replacements.len()
                );
            }
            replacements.push(clip_replacements.pop().unwrap());
            items.push(self.build_video_item(
                preprocessed,
                clip.hash.clone(),
                uuid,
                model_dtype,
            )?);
        }
        Ok(PreparedMedia {
            modality: Modality::Video,
            placeholder: support.placeholder.clone(),
            replacements,
            items,
        })
    }
    /// Preprocess one decoded video clip with the model's resolved vision
    /// processor.
    async fn preprocess_video_clip(
        &self,
        support: &ModalitySupport,
        clip: Arc<VideoClip>,
    ) -> Result<PreprocessedEncoderInputs> {
        let config = support.config.clone();
        let processor = support.processor;
        tokio::task::spawn_blocking(move || {
            // Prefer the borrowed-RGB fast path, which avoids materializing a
            // `DynamicImage` per sampled frame after media decode.
            if let Some(rgb_video) = clip.rgb_video() {
                match rgb_video.frame_refs() {
                    Ok(frame_refs) => match processor.preprocess_video_rgb(&frame_refs, &config) {
                        Ok(preprocessed) => return Ok(preprocessed),
                        Err(error) => warn!(
                            error = %error.as_report(),
                            "RGB video preprocessing fast path failed; falling back to materialized frames"
                        ),
                    },
                    Err(error) => warn!(
                        error,
                        "RGB video frame refs are invalid; falling back to materialized frames"
                    ),
                }
            }
            let frames = clip.materialized_frames().map_err(|error| multimodal!("{error}"))?;
            Ok(processor.preprocess_video(&frames, &config)?)
        })
        .await
        .map_err(|error| multimodal!("video preprocessing task failed: {error}"))?
    }
    /// Convert one preprocessed video clip into engine kwargs.
    ///
    /// The clip is a batch of one, so no per-item slicing is required: the
    /// primary tensor ships as a full-range flat field (the engine re-batches
    /// flat fields by concatenating along the declared dim, matching vLLM's
    /// `flat_from_sizes` treatment of video patches), and batched metadata
    /// tensors drop their singleton batch axis.
    fn build_video_item(
        &self,
        preprocessed: PreprocessedEncoderInputs,
        hash: String,
        uuid: Option<String>,
        model_dtype: ModelDtype,
    ) -> Result<PreparedItem> {
        let tensors = tensor::collect_tensors(preprocessed, VIDEO_PRIMARY_KEY, model_dtype)?;
        let mut data = MmKwargsItem::new();
        for (key, tensor) in tensors {
            let keep_on_cpu = self.spec.keep_on_cpu_keys.contains(&key);
            let (value, field) = if key == VIDEO_PRIMARY_KEY {
                let len = tensor
                    .first_dim()
                    .ok_or_else(|| multimodal!("video encoder input `{key}` is not a tensor"))?;
                (
                    tensor,
                    MmField::Flat(MmFlatField {
                        slices: vec![MmSlice::Slice(SliceSpec {
                            start: Some(0),
                            stop: Some(len as isize),
                            step: None,
                        })],
                        dim: 0,
                        keep_on_cpu,
                    }),
                )
            } else if matches!(
                self.spec.field_layouts.get(&key),
                Some(FieldLayout::Batched)
            ) {
                (
                    tensor.batched_value_at(0)?,
                    MmField::Batched(MmBatchedField { keep_on_cpu }),
                )
            } else {
                (
                    tensor,
                    MmField::Shared(MmSharedField {
                        batch_size: 1,
                        keep_on_cpu,
                    }),
                )
            };
            data.insert(
                key,
                MmFieldElem {
                    data: Some(value.try_into()?),
                    field,
                },
            );
        }
        Ok(PreparedItem { data, hash, uuid })
    }
 }
 #[cfg(test)]
 mod tests {
    use std::collections::HashMap;
    use std::sync::Arc;
    use llm_multimodal::ModelSpecificValue;
    use ndarray::ArrayD;
    use vllm_engine_core_client::protocol::multimodal::MmKwargValue;
    use super::super::tests::{
        QWEN3_IMAGE_PAD_ID, QWEN3_VIDEO_PAD_ID, qwen3_vl_info, qwen3_vl_tokenizer,
    };
    use super::super::{MultimodalConfigFiles, MultimodalModelInfo};
    use super::*;
    #[test]
    fn from_paths_resolves_video_config_from_dedicated_file_or_processor_config() {
        let dir = tempfile::tempdir().unwrap();
        let config_path = dir.path().join("config.json");
        std::fs::write(
            &config_path,
            serde_json::json!({
                "model_type": "qwen3_vl",
                "image_token_id": QWEN3_IMAGE_PAD_ID,
                "video_token_id": QWEN3_VIDEO_PAD_ID,
            })
            .to_string(),
        )
        .unwrap();
        let info_for = |files: MultimodalConfigFiles<'_>| {
            MultimodalModelInfo::from_paths(
                "qwen3-vl-test".to_string(),
                Some("qwen3_vl".to_string()),
                files,
                Arc::new(qwen3_vl_tokenizer()),
            )
        };
        // Dedicated video preprocessor config file.
        let video_config_path = dir.path().join("video_preprocessor_config.json");
        std::fs::write(&video_config_path, r#"{"size":{"shortest_edge":128}}"#).unwrap();
        let info = info_for(MultimodalConfigFiles {
            config: Some(&config_path),
            video_preprocessor_config: Some(&video_config_path),
            ..Default::default()
        })
        .unwrap()
        .unwrap();
        assert!(info.video.is_some());
        // `video_processor` section of the combined processor config.
        let processor_config_path = dir.path().join("processor_config.json");
        std::fs::write(
            &processor_config_path,
            r#"{"video_processor":{"size":{"shortest_edge":128}}}"#,
        )
        .unwrap();
        let info = info_for(MultimodalConfigFiles {
            config: Some(&config_path),
            processor_config: Some(&processor_config_path),
            ..Default::default()
        })
        .unwrap()
        .unwrap();
        assert!(info.video.is_some());
        // Neither source: video support still resolves on the image config.
        let info = info_for(MultimodalConfigFiles {
            config: Some(&config_path),
            ..Default::default()
        })
        .unwrap()
        .unwrap();
        assert!(info.video.is_some());
        // Malformed dedicated file is a real error, not a silent fallback.
        std::fs::write(&video_config_path, r#"{"size""#).unwrap();
        let error = match info_for(MultimodalConfigFiles {
            config: Some(&config_path),
            video_preprocessor_config: Some(&video_config_path),
            ..Default::default()
        }) {
            Err(error) => error,
            Ok(_) => panic!("malformed video preprocessor config should fail"),
        };
        assert!(matches!(
            error,
            Error::Multimodal(message)
                if message.contains("failed to parse video_preprocessor_config.json")
        ));
    }
    #[test]
    fn build_video_item_names_primary_tensor_and_layouts() {
        let info = qwen3_vl_info();
        // One clip flattened to 6 patches with 4 features each.
        let preprocessed = PreprocessedEncoderInputs {
            encoder_input: ArrayD::zeros(vec![6, 4]),
            feature_token_counts: vec![6],
            item_sizes: vec![(32, 32)],
            model_specific: HashMap::from([
                (
                    "video_grid_thw".to_string(),
                    ModelSpecificValue::int_2d(vec![1, 2, 3], 1, 3),
                ),
                (
                    "patches_per_video".to_string(),
                    ModelSpecificValue::int_1d(vec![6]),
                ),
            ]),
        };
        let item = info
            .build_video_item(
                preprocessed,
                "<hash>".to_string(),
                None,
                ModelDtype::Float32,
            )
            .unwrap();
        let primary = &item.data[VIDEO_PRIMARY_KEY];
        assert!(matches!(
            &primary.field,
            MmField::Flat(MmFlatField { slices, dim: 0, .. })
                if matches!(
                    slices.as_slice(),
                    [MmSlice::Slice(SliceSpec { start: Some(0), stop: Some(6), step: None })]
                )
        ));
        // Batched metadata drops its singleton batch axis per item.
        let grid = &item.data["video_grid_thw"];
        assert!(matches!(&grid.field, MmField::Batched(_)));
        let MmKwargValue::Tensor(grid_tensor) = grid.data.as_ref().unwrap() else {
            panic!("expected tensor value for video_grid_thw");
        };
        assert_eq!(grid_tensor.shape, vec![3]);
        assert_eq!(item.hash, "<hash>");
    }
 }
@@ -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,
                },
            ]
        );
@@ -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;
                }
@@ -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,
        })
    );
 }
@@ -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>,
    },
 }
@@ -146,6 +146,7 @@ impl StructuredEventState {
        usage: vllm_llm::TokenUsage,
        finish_reason: FinishReason,
        kv_transfer_params: Option<serde_json::Value>,
        ec_transfer_params: Option<serde_json::Value>,
    ) -> Result<Vec<ChatEvent>> {
        let mut events = Vec::new();
        self.close_open_text_block(&mut events);
@@ -155,6 +156,7 @@ impl StructuredEventState {
            usage,
            finish_reason,
            kv_transfer_params,
            ec_transfer_params,
        });
        Ok(events)
    }
@@ -296,8 +298,11 @@ pub(crate) async fn structured_chat_event_stream(
                usage,
                finish_reason,
                kv_transfer_params,
                ec_transfer_params,
            } => {
-                for next in state.finish(usage, finish_reason, kv_transfer_params)? {
+                for next in
                    state.finish(usage, finish_reason, kv_transfer_params, ec_transfer_params)?
                {
                    y.yield_ok(next).await;
                }
            }
@@ -334,6 +339,7 @@ mod tests {
                },
                finish_reason: FinishReason::stop_eos(),
                kv_transfer_params: None,
                ec_transfer_params: None,
            }),
        ]);
@@ -388,6 +394,7 @@ mod tests {
                },
                finish_reason: FinishReason::stop_eos(),
                kv_transfer_params: None,
                ec_transfer_params: None,
            }),
        ]);
@@ -439,6 +446,7 @@ mod tests {
                },
                finish_reason: FinishReason::stop_eos(),
                kv_transfer_params: None,
                ec_transfer_params: None,
            }),
        ]);
@@ -490,6 +498,7 @@ mod tests {
                },
                finish_reason: FinishReason::stop_eos(),
                kv_transfer_params: None,
                ec_transfer_params: None,
            }),
        ]);
@@ -557,6 +566,7 @@ mod tests {
                },
                finish_reason: FinishReason::stop_eos(),
                kv_transfer_params: None,
                ec_transfer_params: None,
            }),
        ]);
@@ -31,9 +31,14 @@ pub use template::{load_chat_template, resolve_chat_template};
 pub use self::format::ChatTemplateContentFormatOption;
-#[derive(Debug, Clone)]
+/// Template-visible placeholder tokens per supported modality.
 ///
 /// A `None` token means the loaded model does not support that modality, and
 /// content parts of that modality are rejected during rendering.
 #[derive(Debug, Clone, Default)]
 pub struct MultimodalRenderInfo {
-    pub placeholder_token: String,
+    pub image_token: Option<String>,
    pub video_token: Option<String>,
 }
 /// Hugging Face chat-template renderer backed by the local Jinja chat-template
@@ -254,6 +259,7 @@ enum TemplateContent {
 enum TemplateContentPart {
    Text { text: String },
    Image,
    Video,
 }
 #[derive(Debug, Serialize)]
@@ -417,9 +423,17 @@ fn to_template_openai_content(
                }
                // All multimodal contents are normalized to `{ "type": <modality> }`.
                ChatContentPart::ImageUrl { .. } => {
-                    multimodal.ok_or(Error::UnsupportedMultimodalContent("image_url"))?;
+                    multimodal
                        .and_then(|multimodal| multimodal.image_token.as_ref())
                        .ok_or(Error::UnsupportedMultimodalContent("image_url"))?;
                    Ok(TemplateContentPart::Image)
                }
                ChatContentPart::VideoUrl { .. } => {
                    multimodal
                        .and_then(|multimodal| multimodal.video_token.as_ref())
                        .ok_or(Error::UnsupportedMultimodalContent("video_url"))?;
                    Ok(TemplateContentPart::Video)
                }
            })
            .collect(),
    }
@@ -437,9 +451,16 @@ fn to_template_string_content(
                match part {
                    ChatContentPart::Text { text } => out.push_str(text),
                    ChatContentPart::ImageUrl { .. } => {
-                        let multimodal =
+                        let image_token = multimodal
-                            multimodal.ok_or(Error::UnsupportedMultimodalContent("image_url"))?;
+                            .and_then(|multimodal| multimodal.image_token.as_ref())
-                        out.push_str(&multimodal.placeholder_token);
+                            .ok_or(Error::UnsupportedMultimodalContent("image_url"))?;
                        out.push_str(image_token);
                    }
                    ChatContentPart::VideoUrl { .. } => {
                        let video_token = multimodal
                            .and_then(|multimodal| multimodal.video_token.as_ref())
                            .ok_or(Error::UnsupportedMultimodalContent("video_url"))?;
                        out.push_str(video_token);
                    }
                }
            }
@@ -468,7 +489,7 @@ fn append_continue_final_message_tag(message: &mut TemplateMessage) -> Result<St
        // Pick the last text part in the message.
        TemplateContent::OpenAi(parts) => parts.iter_mut().rev().find_map(|part| match part {
            TemplateContentPart::Text { text } => Some(text),
-            TemplateContentPart::Image => None,
+            TemplateContentPart::Image | TemplateContentPart::Video => None,
        }),
    };
    let text = text.ok_or_else(|| {
@@ -577,7 +598,8 @@ mod tests {
    ) -> Result<crate::RenderedPrompt> {
        HfChatRenderer::new(Some(template.to_string()), HashMap::new(), content_format)?
            .with_multimodal(Some(MultimodalRenderInfo {
-                placeholder_token: "<image>".to_string(),
+                image_token: Some("<image>".to_string()),
                video_token: Some("<video>".to_string()),
            }))
            .render(request)
    }
@@ -590,6 +612,14 @@ mod tests {
        ])])
    }
    fn video_request() -> ChatRequest {
        sample_request(vec![ChatMessage::user(vec![
            ChatContentPart::text("a"),
            ChatContentPart::video_url("https://example.com/demo.mp4"),
            ChatContentPart::text("b"),
        ])])
    }
    #[test]
    fn string_content_format_replaces_image_with_placeholder_text() {
        let rendered = render_mm(
@@ -602,6 +632,18 @@ mod tests {
        assert_eq!(rendered.prompt, Prompt::Text("a<image>b".to_string()));
    }
    #[test]
    fn string_content_format_replaces_video_with_placeholder_text() {
        let rendered = render_mm(
            "{{ messages[0].content }}",
            &video_request(),
            ChatTemplateContentFormatOption::String,
        )
        .unwrap();
        assert_eq!(rendered.prompt, Prompt::Text("a<video>b".to_string()));
    }
    #[test]
    fn openai_content_format_normalizes_image_url_for_template() {
        let rendered = render_mm(
@@ -614,6 +656,39 @@ mod tests {
        assert_eq!(rendered.prompt, Prompt::Text("a<|image_pad|>b".to_string()));
    }
    #[test]
    fn openai_content_format_normalizes_video_url_for_template() {
        let rendered = render_mm(
            "{% for item in messages[0].content %}{% if item.type == 'video' %}<|video_pad|>{% else %}{{ item.text }}{% endif %}{% endfor %}",
            &video_request(),
            ChatTemplateContentFormatOption::OpenAi,
        )
        .unwrap();
        assert_eq!(rendered.prompt, Prompt::Text("a<|video_pad|>b".to_string()));
    }
    #[test]
    fn video_parts_are_rejected_when_model_lacks_video_support() {
        let error = HfChatRenderer::new(
            Some("{{ messages[0].content }}".to_string()),
            HashMap::new(),
            ChatTemplateContentFormatOption::String,
        )
        .unwrap()
        .with_multimodal(Some(MultimodalRenderInfo {
            image_token: Some("<image>".to_string()),
            video_token: None,
        }))
        .render(&video_request())
        .unwrap_err();
        assert!(matches!(
            error,
            Error::UnsupportedMultimodalContent("video_url")
        ));
    }
    #[test]
    fn chat_template_supports_pycompat_templates() {
        let request = sample_request(vec![ChatMessage::text(ChatRole::User, "<think>hello")]);
@@ -36,7 +36,13 @@ pub enum ChatContentPart {
        detail: Option<ImageDetail>,
        uuid: Option<String>,
    },
    /// One video URL/data URL content block.
    VideoUrl {
        video_url: String,
        uuid: Option<String>,
    },
    // ImageData...
    // VideoData...
    // ImageEmbeds...
 }
@@ -55,12 +61,21 @@ impl ChatContentPart {
        }
    }
    /// Construct one video URL content part with the given URL string.
    pub fn video_url(video_url: impl Into<String>) -> Self {
        Self::VideoUrl {
            video_url: video_url.into(),
            uuid: None,
        }
    }
    /// Return the text content of this part when it's a text block, or an
    /// "unsupported multimodal content" error otherwise.
    pub(crate) fn as_text(&self) -> Result<&str> {
        match self {
            Self::Text { text } => Ok(text),
            Self::ImageUrl { .. } => Err(Error::UnsupportedMultimodalContent("image_url")),
            Self::VideoUrl { .. } => Err(Error::UnsupportedMultimodalContent("video_url")),
        }
    }
@@ -73,7 +88,7 @@ impl ChatContentPart {
    pub(crate) fn is_multimodal(&self) -> bool {
        match self {
            Self::Text { .. } => false,
-            Self::ImageUrl { .. } => true,
+            Self::ImageUrl { .. } | Self::VideoUrl { .. } => true,
        }
    }
 }
@@ -555,6 +570,26 @@ mod tests {
        assert_eq!(content, ChatContent::Text("hello".to_string()));
    }
    #[test]
    fn chat_content_video_url_part_round_trips_through_serde() {
        let content = ChatContent::Parts(vec![ChatContentPart::VideoUrl {
            video_url: "https://example.com/demo.mp4".to_string(),
            uuid: Some("video-1".to_string()),
        }]);
        let value = to_value(&content).unwrap();
        assert_eq!(
            value,
            json!([{
                "type": "video_url",
                "video_url": "https://example.com/demo.mp4",
                "uuid": "video-1",
            }])
        );
        let decoded: ChatContent = serde_json::from_value(value).unwrap();
        assert_eq!(decoded, content);
    }
    #[test]
    fn chat_content_deserializes_from_openai_text_blocks() {
        let content: ChatContent =
@@ -22,6 +22,9 @@ pub struct CollectedAssistantMessage {
    pub finish_reason: FinishReason,
    /// Connector-specific KV transfer parameters for disaggregated serving.
    pub kv_transfer_params: Option<serde_json::Value>,
    /// Connector-specific encoder cache transfer parameters for disaggregated
    /// serving.
    pub ec_transfer_params: Option<serde_json::Value>,
 }
 /// Per-request stream of chat events.
@@ -77,6 +80,7 @@ impl ChatEventStream {
                    usage,
                    finish_reason,
                    kv_transfer_params,
                    ec_transfer_params,
                } => {
                    return Ok(CollectedAssistantMessage {
                        message: done,
@@ -89,6 +93,7 @@ impl ChatEventStream {
                        usage,
                        finish_reason,
                        kv_transfer_params,
                        ec_transfer_params,
                    });
                }
                ChatEvent::ToolCallEnd { call, .. } => {
@@ -194,6 +199,7 @@ mod tests {
                    },
                    finish_reason: FinishReason::stop_eos(),
                    kv_transfer_params: None,
                    ec_transfer_params: None,
                }),
            ]),
        );
@@ -234,6 +240,7 @@ mod tests {
                },
                finish_reason: FinishReason::stop_eos(),
                kv_transfer_params: None,
                ec_transfer_params: None,
            }
        );
    }
@@ -50,6 +50,7 @@ fn request_output(
        stop_reason,
        events: None,
        kv_transfer_params: None,
        ec_transfer_params: None,
        trace_headers: None,
        prefill_stats: None,
        routed_experts: None,
@@ -75,6 +76,7 @@ fn request_output_with_logprobs(
        stop_reason,
        events: None,
        kv_transfer_params: None,
        ec_transfer_params: None,
        trace_headers: None,
        prefill_stats: None,
        routed_experts: None,
@@ -117,6 +117,19 @@ impl RoundtripCase {
        }
    }
    /// MiniMax M3 invoke format with `<mm:think>` reasoning tags.
    fn minimax_m3() -> Self {
        Self {
            model_id: "MiniMaxAI/MiniMax-M3",
            assistant_stop_suffix: "[e~[\n",
            tool_call_parser: ParserSelection::Auto,
            reasoning_parser: ParserSelection::Auto,
            thinking_behavior: ThinkingBehavior::Always { value: true },
            json_fmt: compact_json_fmt(),
            sort_json_keys: false,
        }
    }
    /// DeepSeek V4 DSML tool-call format.
    fn deepseek_v4() -> Self {
        Self {
@@ -143,6 +156,19 @@ impl RoundtripCase {
        }
    }
    /// GLM-4.5 XML-like argument format with `<think>` reasoning tags.
    fn glm45() -> Self {
        Self {
            model_id: "zai-org/GLM-4.5",
            assistant_stop_suffix: "",
            tool_call_parser: ParserSelection::Auto,
            reasoning_parser: ParserSelection::Auto,
            thinking_behavior: ThinkingBehavior::Toggleable { default: true },
            json_fmt: compact_json_fmt(),
            sort_json_keys: false,
        }
    }
    /// GLM-4.7 XML-like argument format with `<think>` reasoning tags.
    fn glm47() -> Self {
        Self {
@@ -209,6 +235,19 @@ impl RoundtripCase {
        }
    }
    /// Nemotron V3 with `<think>` / `</think>` reasoning tags.
    fn nemotron_v3() -> Self {
        Self {
            model_id: "nvidia/NVIDIA-Nemotron-3-Nano-30B-A3B-BF16",
            assistant_stop_suffix: "<|im_end|>\n",
            tool_call_parser: ParserSelection::Auto,
            reasoning_parser: ParserSelection::Auto,
            thinking_behavior: ThinkingBehavior::Always { value: true },
            json_fmt: compact_json_fmt(),
            sort_json_keys: false,
        }
    }
    /// GPT-OSS Harmony token-id renderer and native Harmony output processor.
    fn gpt_oss() -> Self {
        Self {
@@ -247,11 +286,14 @@ roundtrip_tests! {
    qwen3 => [reasoning_and_content, tool_call_mix],
    qwen35 => [reasoning_and_content, tool_call_mix],
    minimax_m25 => [reasoning_and_content, tool_call_mix],
    minimax_m3 => [reasoning_and_content, tool_call_mix],
    deepseek_v4 => [reasoning_and_content, tool_call_mix],
    deepseek_v32 => [tool_call_mix],
    glm45 => [reasoning_and_content, tool_call_mix],
    glm47 => [reasoning_and_content, tool_call_mix],
    seed_oss => [reasoning_and_content],
    step3p5 => [reasoning_and_content],
    nemotron_v3 => [reasoning_and_content],
    gemma4 => [tool_call_mix], // Gemma4 strips reasoning in history if there's no tool call
    kimi_k25 => [tool_call_mix], // Kimi K2.5 strips reasoning in history
    gpt_oss => [tool_call_mix], // Harmony strips reasoning in history if there's no tool call
@@ -634,6 +676,7 @@ fn decoded_completion_stream(
                    usage: Default::default(),
                    finish_reason: FinishReason::stop_eos(),
                    kv_transfer_params: None,
                    ec_transfer_params: None,
                }),
            }
        });
@@ -644,6 +687,7 @@ fn decoded_completion_stream(
                usage: Default::default(),
                finish_reason: FinishReason::stop_eos(),
                kv_transfer_params: None,
                ec_transfer_params: None,
            });
            events.push(DecodedTextEvent::TextDelta {
                delta: chunk.delta,
@@ -97,6 +97,8 @@ pub struct EngineCoreOutput {
    #[serde(default)]
    pub kv_transfer_params: Option<serde_json::Value>,
    #[serde(default)]
    pub ec_transfer_params: Option<serde_json::Value>,
    #[serde(default)]
    pub trace_headers: Option<OpaqueValue>,
    /// Breakdown of the scheduled prefill computation, set on the first output
    /// of a newly scheduled prefill and elided for subsequent decode outputs.
@@ -374,6 +376,7 @@ mod tests {
                stop_reason: Some(StopReason::Text("stop".to_string())),
                events: None,
                kv_transfer_params: None,
                ec_transfer_params: None,
                trace_headers: None,
                prefill_stats: None,
                routed_experts: None,
@@ -426,6 +429,7 @@ mod tests {
                            stop_reason: None,
                            events: None,
                            kv_transfer_params: None,
                            ec_transfer_params: None,
                            trace_headers: None,
                            prefill_stats: None,
                            routed_experts: None,
@@ -226,6 +226,7 @@ fn request_output(
        stop_reason: None,
        events: None,
        kv_transfer_params: None,
        ec_transfer_params: None,
        trace_headers: None,
        prefill_stats: None,
        routed_experts: None,
@@ -2517,6 +2518,7 @@ fn python_msgpack_fixtures_match_rust_encoding() {
                        stop_reason: None,
                        events: None,
                        kv_transfer_params: None,
                        ec_transfer_params: None,
                        trace_headers: None,
                        prefill_stats: None,
                        routed_experts: None,
@@ -91,6 +91,7 @@ class EngineCoreOutput(
    stop_reason: int | str | None = None
    events: object | None = None
    kv_transfer_params: object | None = None
    ec_transfer_params: object | None = None
    trace_headers: object | None = None
    prefill_stats: object | None = None
    routed_experts: object | None = None
@@ -64,6 +64,13 @@ impl InflightRequests {
            .collect()
    }
    /// Collect the internal engine ids of every in-flight request. Used to
    /// abort all outstanding requests when no external ids are given.
    pub(crate) fn all_internal_ids(&self) -> Vec<String> {
        let map = self.map.lock();
        map.values().flat_map(|internal_ids| internal_ids.keys()).cloned().collect()
    }
    #[cfg(test)]
    fn is_empty(&self) -> bool {
        self.map.lock().is_empty()
@@ -122,7 +122,12 @@ impl Llm {
    /// tracking entries themselves are removed when the corresponding output
    /// streams are dropped, not here.
    pub async fn abort(&self, external_ids: &[String]) -> Result<()> {
-        let internal_ids = self.inflight.resolve(external_ids);
+        // Empty `external_ids` means abort every in-flight request.
        let internal_ids = if external_ids.is_empty() {
            self.inflight.all_internal_ids()
        } else {
            self.inflight.resolve(external_ids)
        };
        if internal_ids.is_empty() {
            return Ok(());
        }
@@ -38,6 +38,9 @@ pub struct CollectedGenerateOutput {
    pub usage: TokenUsage,
    /// Connector-specific KV transfer parameters for disaggregated serving.
    pub kv_transfer_params: Option<serde_json::Value>,
    /// Connector-specific encoder cache transfer parameters for disaggregated
    /// serving.
    pub ec_transfer_params: Option<serde_json::Value>,
 }
 /// Prompt-scoped metadata emitted only once on the first [`GenerateOutput`] for
@@ -146,6 +149,9 @@ pub struct GenerateOutput {
    pub cached_token_count: usize,
    /// Connector-specific KV transfer parameters for disaggregated serving.
    pub kv_transfer_params: Option<serde_json::Value>,
    /// Connector-specific encoder cache transfer parameters for disaggregated
    /// serving.
    pub ec_transfer_params: Option<serde_json::Value>,
 }
 impl GenerateOutput {
@@ -192,6 +198,7 @@ impl GenerateOutput {
            finish_reason,
            cached_token_count: 0,
            kv_transfer_params: None,
            ec_transfer_params: None,
        }
    }
 }
@@ -280,6 +287,7 @@ impl Stream for GenerateOutputStream {
            finish_reason,
            cached_token_count,
            kv_transfer_params: raw.kv_transfer_params,
            ec_transfer_params: raw.ec_transfer_params,
        };
        Poll::Ready(Some(Ok(output)))
@@ -362,6 +370,7 @@ impl<T: Stream<Item = Result<GenerateOutput>> + Send> T {
                            cached_token_count,
                        },
                        kv_transfer_params: None,
                        ec_transfer_params: None,
                    });
                }
@@ -374,6 +383,7 @@ impl<T: Stream<Item = Result<GenerateOutput>> + Send> T {
                        cached_token_count,
                    };
                    collected.kv_transfer_params = output.kv_transfer_params;
                    collected.ec_transfer_params = output.ec_transfer_params;
                    return Ok(collected);
                }
            }
@@ -51,6 +51,7 @@ fn request_output_with_events(
        stop_reason: None,
        events,
        kv_transfer_params: None,
        ec_transfer_params: None,
        trace_headers: None,
        prefill_stats: None,
        routed_experts: None,
@@ -75,6 +76,7 @@ fn request_output_with_logprobs(
        stop_reason: None,
        events: None,
        kv_transfer_params: None,
        ec_transfer_params: None,
        trace_headers: None,
        prefill_stats: None,
        routed_experts: None,
@@ -89,6 +91,7 @@ fn request_output_with_logprobs_and_kv(
    new_logprobs: Option<Logprobs>,
    prompt_logprobs: Option<Logprobs>,
    kv_transfer_params: Option<serde_json::Value>,
    ec_transfer_params: Option<serde_json::Value>,
 ) -> EngineCoreOutput {
    EngineCoreOutput {
        request_id: request_id.to_string(),
@@ -100,6 +103,7 @@ fn request_output_with_logprobs_and_kv(
        stop_reason: None,
        events: None,
        kv_transfer_params,
        ec_transfer_params,
        trace_headers: None,
        prefill_stats: None,
        routed_experts: None,
@@ -356,6 +360,7 @@ async fn collect_output_aggregates_raw_tokens_logprobs_and_terminal_metadata() {
                                Some(logprobs_for_position(44, -0.3, 1, 88, -0.4)),
                                None,
                                Some(serde_json::json!({"connector": "x"})),
                                None,
                            ),
                        ],
                        ..Default::default()
@@ -69,6 +69,11 @@ pub fn to_text_request(
            let map = sampling_params.vllm_xargs.get_or_insert_with(Default::default);
            map.insert("kv_transfer_params".to_string(), kv_json);
        }
        if let Some(ec_struct) = kv.ec_transfer_params.as_ref() {
            let ec_json = proto_struct_to_json(ec_struct);
            let map = sampling_params.vllm_xargs.get_or_insert_with(Default::default);
            map.insert("ec_transfer_params".to_string(), ec_json);
        }
        if kv.bypass_prefix_cache {
            sampling_params.skip_reading_prefix_cache = Some(true);
        }
@@ -343,6 +348,7 @@ fn to_finish_info(finished: &Finished, token_ids: &[u32]) -> pb::FinishInfo {
        finish_reason,
        stop_reason,
        kv_transfer_params: finished.kv_transfer_params.as_ref().and_then(json_to_proto_struct),
        ec_transfer_params: finished.ec_transfer_params.as_ref().and_then(json_to_proto_struct),
    }
 }
@@ -586,6 +592,7 @@ mod tests {
            },
            finish_reason: reason,
            kv_transfer_params: None,
            ec_transfer_params: None,
        }
    }
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
Lucas WilkinsonandOpenAI Codex	6bf03e0d95	[Core] Add explicit layer parallel plans Co-authored-by: OpenAI Codex <codex@openai.com> Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>	2026-07-13 19:12:30 +00:00
zztandGitHub	2595d5cebc	[Model] Optimize Qwen3.5 on H20 (#48350 ) Signed-off-by: zzt <zengzetang.zzt@antgroup.com>	2026-07-13 03:30:48 +00:00
Tan Pin Siang GitHub vllmellm Hongxia Yang Jun Kang Chow tjtanaa	ee5a89f4d7	[ROCm][MiniMax-M3] Add AITER sparse paged attention (#47287 ) Signed-off-by: Tan Pin Siang <tanpinsiang@gmail.com> Signed-off-by: tjtanaa <tunjian.tan@embeddedllm.com> Co-authored-by: vllmellm <vllm.ellm@embeddedllm.com> Co-authored-by: Hongxia Yang <hongxia.yang@amd.com> Co-authored-by: Jun Kang Chow <junkangchow@gmail.com> Co-authored-by: tjtanaa <tunjian.tan@embeddedllm.com>	2026-07-12 19:27:29 -07:00
Tanish Malekar GitHub Claude Sonnet 4.6	e26264f3ef	[Kernel] Implement CUDA kernel for ReLUSquaredActivation (relu^2) (#39058 ) Signed-off-by: Tanish Malekar <tanishmalekar32@gmail.com> Co-authored-by: Claude Sonnet 4.6 <noreply@anthropic.com>	2026-07-12 19:18:03 -07:00
AlexHuangandGitHub	4c81772e8b	[Bugfix][KV Offloading] Fix stale transfer_jobs after reset_cache + harden job completion (#48102 ) Signed-off-by: Alex <alex.tech.lab@outlook.com>	2026-07-12 20:00:04 +03:00
Bugen ZhaoandGitHub	27c3e579f0	[CI][Rust Frontend] Pin cargo tool versions (#48222 ) Signed-off-by: Bugen Zhao <i@bugenzhao.com>	2026-07-12 16:34:26 +01:00
omerpaz95 GitHub Or Ozeri	8df14cfc8c	[EC Connector] Add EC Transfer Params (#42433 ) Signed-off-by: omerpaz95 <omerpaz95@gmail.com> Co-authored-by: Or Ozeri <oro@il.ibm.com>	2026-07-12 14:35:33 +03:00
Jiangyun ZhuandGitHub	370b678a02	[CI][2/N] reduce CI time (#48394 ) Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>	2026-07-12 04:16:55 -07:00
liranschour GitHub Or Ozeri	5c0c987c03	Make tiering offload region DP-replica aware (#47987 ) Signed-off-by: Liran Schour <lirans@il.ibm.com> Co-authored-by: Or Ozeri <oro@il.ibm.com>	2026-07-12 13:10:21 +03:00
Hugo CentenoandGitHub	5f8e73cb8b	[Bugfix] Guard mixed-dtype allreduce RMSNorm quant fusions (#48330 ) Signed-off-by: hcenteno <hugo.centeno@estudiantat.upc.edu>	2026-07-12 09:39:27 +00:00
aoshen02 GitHub Claude Opus 4.8mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	83762b77b0	[Frontend] Add /abort_requests to the RLHF dev API router (#47173 ) Signed-off-by: aoshen02 <aoshen@inferact.ai> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-07-12 14:21:02 +08:00
Canlin Guo GitHub yewentao256	a02984ed47	[Perf][Qwen] Replace MOE all-reduce with reduce-scatter (#47006 ) Signed-off-by: gcanlin <canlinguosdu@gmail.com> Signed-off-by: yewentao256 <zhyanwentao@126.com> Co-authored-by: yewentao256 <zhyanwentao@126.com>	2026-07-12 06:14:49 +00:00
vx120 GitHub crp0128 aoshen02 Claude Opus 4.8	fc1c548093	Runtime Draft Weight Update for Speculative Decoding (#46725 ) Signed-off-by: vx120 <893600387@qq.com> Signed-off-by: vx120 <57470515+vx120@users.noreply.github.com> Signed-off-by: aoshen02 <aoshen@inferact.ai> Co-authored-by: crp0128 <191679376@qq.com> Co-authored-by: aoshen02 <aoshen@inferact.ai> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>	2026-07-11 22:51:53 -07:00
Jiangyun Zhu GitHub Yifan Qiao	481e481be7	[2/N][Core] support partial prefix cache hit for hybrid model (#46384 ) Signed-off-by: zjy0516 <riverclouds.zhu@qq.com> Signed-off-by: Yifan Qiao <yifanqiao@inferact.ai> Co-authored-by: Yifan Qiao <yifanqiao@inferact.ai>	2026-07-12 05:37:51 +00:00
zhao, zhenhuiandGitHub	8e981630c9	[CI][CPU] Add Qwen2-VL multimodal tests for CPU backend and fix incompatibilities (#48072 ) Signed-off-by: Zhenhui Zhao <zhenhui.zhao@intel.com>	2026-07-12 12:30:34 +08:00
Alejandro Paredes La TorreandGitHub	9a48eef89a	[Bugfix][LoRA] Support ark_linear base layer in _get_lora_device (#47690 ) Signed-off-by: AlejandroParedesLT <alejandroparedeslatorre@gmail.com>	2026-07-12 00:13:50 +00:00
Jiangyun ZhuandGitHub	1ef1c7ebba	[CI] split tests to reduce CI time (#48219 ) Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>	2026-07-11 13:00:14 -07:00
Ievgen Bondarenko GitHub Claude Opus 4.7 Isotr0py	54503ecec0	fix(processor): route MiMo-V2-Omni media fetch through MediaConnector (#43117 ) Signed-off-by: Ievgen Bondarenko <ibondarenko@student.sierracollege.edu> Signed-off-by: Ievgen (Jack) Bondarenko <ibondarenko@student.sierracollege.edu> Signed-off-by: Isotr0py <Isotr0py@outlook.com> Co-authored-by: Claude Opus 4.7 (1M context) <noreply@anthropic.com> Co-authored-by: Isotr0py <Isotr0py@outlook.com>	2026-07-11 15:52:52 +00:00
ErenAta16andGitHub	0067311536	fix(entrypoints): stop resolve_items leaking in-flight media fetch tasks on partial failure (#48333 ) Signed-off-by: ErenAta16 <erena6466@gmail.com>	2026-07-11 15:42:08 +00:00
Lucas Wilkinson GitHub Claude OpenAI Codex Matthew Bonanni Nicolò Lucchesi	51878e5b6e	[2/N][KV-Cache Layout Refactor] Pack K/V into the content dim across attention backends (#44455 ) Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com> Signed-off-by: Matthew Bonanni <mbonanni@redhat.com> Signed-off-by: Lucas Wilkinson <LucasWilkinson@users.noreply.github.com> Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: OpenAI Codex <codex@openai.com> Co-authored-by: Matthew Bonanni <mbonanni@redhat.com> Co-authored-by: Nicolò Lucchesi <nlucches@redhat.com>	2026-07-11 11:11:16 -04:00
Yejing LaiandGitHub	76fedaa2a5	[XPU][UT]Fix InternS1ProForConditionalGeneration AssertionError (#48232 ) Signed-off-by: Lai, Yejing <yejing.lai@intel.com>	2026-07-11 13:56:55 +00:00
Jee Jee Li GitHub peiyuanz Claude zhouzhou	19069bcbd5	FP32 router GEMV optimization (#48335 ) Signed-off-by: peiyuanz <peiyuanz@inferact.ai> Signed-off-by: Jee Jee Li <jeejeelee@inferact.ai> Co-authored-by: peiyuanz <peiyuanz@inferact.ai> Co-authored-by: Claude <noreply@anthropic.com> Co-authored-by: zhouzhou <zhouzhou@zhouzhoudeMacBook-Pro.local>	2026-07-11 13:07:48 +00:00
Harry MellorandGitHub	1bd8f80a64	[CI] Point CI at Transformers release rather than release branch (#48328 ) Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>	2026-07-11 02:31:14 -07:00
Qiming Zhang GitHub Kunshang Ji	0b6636cbcb	[XPU]remove is_xxx from moe class and bump up kernels (#48079 ) Signed-off-by: mayuyuace <qiming1.zhang@intel.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>	2026-07-11 09:27:13 +00:00
Harry MellorandGitHub	4a6440acef	Bump Transformers version to 5.13.0 (#47867 ) Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com>	2026-07-11 00:56:14 -07:00
Lucas WilkinsonandGitHub	bec0a4ede6	[Revert] [Build] Update vllm ...builds FA3 with torch stable API (#48269 ) Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>	2026-07-11 05:20:25 +00:00
Robert Shaw GitHub Robert Shaw	3d99b0499a	[Logs] DP Supervisor Log Improvement (#48278 ) Signed-off-by: Robert Shaw <robertgshaw2-redhat@ip-172-31-18-125.us-east-2.compute.internal> Co-authored-by: Robert Shaw <robertgshaw2-redhat@ip-172-31-18-125.us-east-2.compute.internal>	2026-07-11 12:07:00 +08:00
Lucas Wilkinson GitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	04d553f390	[Misc] Use meta tensor for KV cache stride calculation (#47316 ) Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-07-10 23:24:59 -04:00
Lucas Wilkinson GitHubmergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	9c18e90f6c	[BugFix] Fix packed HND KV cache reshape for FlashAttention (#47314 ) Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>	2026-07-10 23:22:39 -04:00
Jimmy LeeandGitHub	092387963c	[BugFix] weights processing peak memory reduction for nvfp4 MoE layers (#46276 ) Signed-off-by: Jimmy Lee <hirejimmylee@gmail.com>	2026-07-11 02:05:35 +00:00
Joe Rowell GitHub mgoin	1bf3997eae	[Quantization] Bound peak memory when repacking FP4 MoE weights for Marlin (#47851 ) Signed-off-by: Joe Rowell <joerowell4@gmail.com> Signed-off-by: mgoin <mgoin64@gmail.com> Co-authored-by: mgoin <mgoin64@gmail.com>	2026-07-10 19:46:13 -06:00
Michael Goin GitHub Claude Fable 5	29fd688892	Add VLLM_FLASHINFER_AUTOTUNE_SKIP_OPS and skip CuTeDSL fp4_gemm autotuning by default (#48268 ) Signed-off-by: mgoin <mgoin64@gmail.com> Signed-off-by: Michael Goin <mgoin64@gmail.com> Co-authored-by: Claude Fable 5 <noreply@anthropic.com>	2026-07-10 18:13:08 -07:00
Ashwin GiridharanandGitHub	ed908cf0a0	[Bugfix] Fix thinking_token_budget not enforced after natural </think> re-entry (#45984 ) Signed-off-by: Ashwin Giridharan <girida@amazon.com>	2026-07-10 22:47:51 +00:00
Nick Hill GitHub Claude Opus 4.8	26ff616bbf	[Bugfix][Test] Register Qwen/Qwen3.5-4B example model (#48276 ) Signed-off-by: Nick Hill <nickhill123@gmail.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>	2026-07-10 17:01:23 -04:00
gnovackandGitHub	f378f79b7c	handle topk_ids padding in align sum kernel (#47785 ) Signed-off-by: gnovack <novackgm@gmail.com>	2026-07-10 13:33:28 -07:00
Matthew Bonanni GitHub Claude	735def4fcf	[Bugfix] Fix FlashMLA dense fp8 metadata crash (num_sm_parts clamp) (#48045 ) Signed-off-by: Matthew Bonanni <mbonanni@redhat.com> Co-authored-by: Claude <noreply@anthropic.com>	2026-07-10 12:24:52 -07:00
larryli2-amd GitHub Andreas Karatzas	c227aaa3f8	[ROCm] Enable DeepSeek-V4 DSpark speculative decoding on AMD (MI350X / MI355X, gfx950) (#47419 ) Signed-off-by: larryli2-amd <larryli2@amd.com> Signed-off-by: larryli2-amd <Larry.Li@amd.com> Co-authored-by: Andreas Karatzas <akaratza@amd.com>	2026-07-10 23:22:19 +08:00
Michael GoinandGitHub	08dfd68610	[Model] Add LongCat-Flash-Lite (n-gram embedding) (#47857 ) Signed-off-by: mgoin <mgoin64@gmail.com>	2026-07-10 07:17:50 -07:00
Tyler Michael SmithandGitHub	978a6dfa3f	[Build/CI] Build arm64 PR and postmerge image builds for Blackwell SM10x and SM110 (#48041 ) Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>	2026-07-10 10:12:50 -04:00
HuYiPeng GitHub Yipeng Hu Isotr0py	85c09e9885	fix: correct load_weights track logic and enable weight integrity for… (#41811 ) Signed-off-by: Yipeng Hu <i26268@metax-tech.com> Signed-off-by: HuYiPeng <144002351+MynameFelix@users.noreply.github.com> Signed-off-by: Isotr0py <Isotr0py@outlook.com> Co-authored-by: Yipeng Hu <i26268@metax-tech.com> Co-authored-by: Isotr0py <Isotr0py@outlook.com>	2026-07-10 14:08:20 +00:00
XuZhou GitHub Xu Zhou Hoseung Kim	b12cca6a23	[Bugfix] Fix turboquant FP8 cast failure for BF16 models on Ampere GPUs (#39988 ) Signed-off-by: Xu Zhou <xuzhou9417@163.com> Co-authored-by: Xu Zhou <xuzhou9417@163.com> Co-authored-by: Hoseung Kim <ghyutjik123@gmail.com>	2026-07-10 06:55:28 -07:00
Wentao YeandGitHub	e257faf87d	[Refactor] Remove unused rocm kernel `combine_topk_swa_indices_ragged` (#48158 ) Signed-off-by: yewentao256 <zhyanwentao@126.com>	2026-07-10 09:33:07 -04:00
FAN YUCHENandGitHub	fabec87f63	[Model] Migrate MistralLarge3ForCausalLM to AutoWeightsLoader (#48153 ) Signed-off-by: Yuchen Fan <functionhx@gmail.com>	2026-07-10 12:27:58 +00:00
Nick Hill GitHub Claude Opus 4.8	7614b88ebd	[Bugfix][Spec Decode] Fix DFlash draft/target layer-count mismatch (#48113 ) Signed-off-by: Nick Hill <nickhill123@gmail.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>	2026-07-10 04:42:53 -07:00
Isotr0pyandGitHub	68ea76e780	[Misc] Remove dead code in ViT functionality test (#48220 ) Signed-off-by: Isotr0py <Isotr0py@outlook.com>	2026-07-10 11:17:42 +00:00
Reid GitHub Bugen Zhao	c241c7a2b0	[Rust Frontend] Add roundtrip fixtures for more chat parsers (#47883 ) Co-authored-by: Bugen Zhao <i@bugenzhao.com> Signed-off-by: reidliu41 <reid201711@gmail.com> Signed-off-by: Bugen Zhao <i@bugenzhao.com>	2026-07-10 10:03:58 +00:00
ViranjanPagar GitHub Claude Opus 4.8 Isotr0py Isotr0py Roger Wang	e23b19309b	Deepstream video backend (#42424 ) Signed-off-by: Viranjan Pagar <vpagar@nvidia.com> Signed-off-by: Isotr0py <Isotr0py@outlook.com> Signed-off-by: Isotr0py <mozf@mail2.sysu.edu.cn> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com> Co-authored-by: Isotr0py <mozf@mail2.sysu.edu.cn> Co-authored-by: Isotr0py <Isotr0py@outlook.com> Co-authored-by: Roger Wang <hey@rogerw.io>	2026-07-10 02:23:30 -07:00
Andrey Talman GitHub Andrey Talman Kevin H. Luu	f36284a8d2	[CI] Add TORCH_NIGHTLY=1 build mode (run full suite on torch nightly) (#47180 ) Co-authored-by: Andrey Talman <atalman@users.noreply.github.com> Co-authored-by: Kevin H. Luu <khluu000@gmail.com>	2026-07-10 01:38:35 -07:00
Mingfei GuoandGitHub	424df4f65d	[Model][CI/Build] Cosmos3: enable registry tests and register Cosmos3-Super (#48211 ) Signed-off-by: Mingfei Guo <1800012773@pku.edu.cn>	2026-07-10 16:35:15 +08:00
Bugen ZhaoandGitHub	074bdd0d99	[Rust Frontend] Integrate MM video support (#47959 ) Signed-off-by: Bugen Zhao <i@bugenzhao.com>	2026-07-10 08:15:33 +00:00
wenjun liu GitHubjun,du <jun.du@intel.com>Kunshang Ji	216ee58780	Add XPU nightly and release image publishing to DockerHub (#48126 ) Signed-off-by: wenjun.liu <wenjun.liu@intel.com> Signed-off-by: jun,du <jun.du@intel.com> Co-authored-by: jun,du <jun.du@intel.com> Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>	2026-07-10 00:56:00 -07:00
Kevin H. Luu GitHub Claude Opus 4.8	433f291195	[CI] Right-size test-area timeouts from nightly durations (#48186 ) Signed-off-by: khluu <khluu000@gmail.com> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>	2026-07-10 00:53:16 -07:00
Chaojun ZhangandGitHub	28eaf05d56	[XPU] Enable v1/sample tests on XPU CI (#44472 ) Signed-off-by: Chaojun Zhang <chaojun.zhang@intel.com>	2026-07-10 15:40:51 +08:00
Jiangyun ZhuandGitHub	300e33797f	[Perf] fuse more rmsnorm and all-reduce in qwen3.5 (#46998 ) Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>	2026-07-10 15:37:51 +08:00
alberto GitHub Chauncey	5715fde12c	[Feature][Parser] Support include_reasoning param for non-Harmony models (#44301 ) Signed-off-by: Alberto Perdomo <aperdomo@redhat.com> Co-authored-by: Chauncey <chaunceyjiang@gmail.com>	2026-07-10 15:34:02 +08:00