tests: de-duplicate online scoring API format checks

Refactor repeated score request format tests into a shared assertion helper and parametrized cases in bi-encoder and cross-encoder online suites, reducing bloat while preserving coverage.

Co-authored-by: Claude <noreply@anthropic.com>

Signed-off-by: Cursor Agent <cursoragent@cursor.com>

.buildkite/ci_config_intel.yaml

@@ -1,23 +0,0 @@
-name: vllm_intel_ci
-job_dirs:
-  - ".buildkite/intel_jobs"
-run_all_patterns:
-  - "docker/Dockerfile"
-  - "CMakeLists.txt"
-  - "requirements/common.txt"
-  - "requirements/xpu.txt"
-  - "requirements/build.txt"
-  - "requirements/test.txt"
-  - "setup.py"
-  - "csrc/"
-  - "cmake/"
-run_all_exclude_patterns:
-  - "docker/Dockerfile."
-  - "csrc/cpu/"
-  - "csrc/rocm/"
-  - "cmake/hipify.py"
-  - "cmake/cpu_extension.cmake"
-registries: public.ecr.aws/q9t5s3a7
-repositories:
-  main: "vllm-ci-test-repo"
-  premerge: "vllm-ci-test-repo"

.buildkite/image_build/image_build_xpu.sh

@@ -1,34 +0,0 @@
-#!/bin/bash
-set -e
-
-if [[ $# -lt 3 ]]; then
-  echo "Usage: $0 <registry> <repo> <commit>"
-  exit 1
-fi
-
-REGISTRY=$1
-REPO=$2
-BUILDKITE_COMMIT=$3
-
-# authenticate with AWS ECR
-aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY"
-aws ecr get-login-password --region us-east-1 | docker login --username AWS --password-stdin 936637512419.dkr.ecr.us-east-1.amazonaws.com
-
-# skip build if image already exists
-if ! docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu &> /dev/null; then
-  echo "Image not found, proceeding with build..."
-else
-  echo "Image found"
-  exit 0
-fi
-
-# build
-docker build \
-  --file docker/Dockerfile.xpu \
-  --build-arg max_jobs=16 \
-  --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
-  --tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu \
-  --progress plain .
-
-# push
-docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu

.buildkite/intel_jobs/test-intel.yaml

@@ -1,64 +0,0 @@
-group: Intel
-steps:
-  - label: ":docker: Build XPU image"
-    soft_fail: true
-    depends_on: []
-    key: image-build-xpu
-    commands:
-      - bash -lc '.buildkite/image_build/image_build_xpu.sh "public.ecr.aws/q9t5s3a7" "vllm-ci-test-repo" "$BUILDKITE_COMMIT"'
-    env:
-      DOCKER_BUILDKIT: "1"
-    retry:
-      automatic:
-        - exit_status: -1  # Agent was lost
-          limit: 2
-        - exit_status: -10  # Agent was lost
-          limit: 2
-  - label: "XPU example Test"
-    depends_on:
-      - image-build-xpu
-    timeout_in_minutes: 30
-    device: intel_gpu
-    no_plugin: true
-    env:
-      REGISTRY: "public.ecr.aws/q9t5s3a7"
-      REPO: "vllm-ci-test-repo"
-    source_file_dependencies:
-      - vllm/
-      - .buildkite/intel_jobs/test-intel.yaml 
-    commands:
-      - >-
-        bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-        'pip install tblib==3.1.0 &&
-        python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager &&
-        python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 -O3 -cc.cudagraph_mode=NONE &&
-        python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager -tp 2 --distributed-executor-backend mp &&
-        python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --attention-backend=TRITON_ATTN &&
-        python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --quantization fp8 &&
-        python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --block-size 64 --enforce-eager --max-model-len 8192 &&
-        python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 &&
-        python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 --enable-expert-parallel'
-  - label: "XPU V1 test"
-    depends_on:
-      - image-build-xpu
-    timeout_in_minutes: 30
-    device: intel_gpu
-    no_plugin: true
-    env:
-      REGISTRY: "public.ecr.aws/q9t5s3a7"
-      REPO: "vllm-ci-test-repo"
-    source_file_dependencies:
-      - vllm/
-      - .buildkite/intel_jobs/test-intel.yaml 
-    commands:
-      - >-
-        bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-        'cd tests &&
-        pytest -v -s v1/core --ignore=v1/core/test_reset_prefix_cache_e2e.py --ignore=v1/core/test_scheduler_e2e.py &&
-        pytest -v -s v1/engine --ignore=v1/engine/test_output_processor.py &&
-        pytest -v -s v1/sample --ignore=v1/sample/test_logprobs.py --ignore=v1/sample/test_logprobs_e2e.py &&
-        pytest -v -s v1/worker --ignore=v1/worker/test_gpu_model_runner.py --ignore=v1/worker/test_worker_memory_snapshot.py &&
-        pytest -v -s v1/structured_output &&
-        pytest -v -s v1/test_serial_utils.py &&
-        pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_tree_attention.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py &&
-        pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_nixl_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --ignore=v1/kv_connector/unit/test_lmcache_integration.py'

.buildkite/release-pipeline.yaml

@@ -90,14 +90,6 @@ steps:
         env:
           DOCKER_BUILDKIT: "1"
 
-  - label: "Generate and upload wheel indices"
-    depends_on: "build-wheels"
-    allow_dependency_failure: true
-    agents:
-      queue: cpu_queue_release
-    commands:
-      - "bash .buildkite/scripts/generate-and-upload-nightly-index.sh"
-
   - group: "Build release Docker images"
     key: "build-release-images"
     steps:
@@ -611,7 +603,7 @@ steps:
       - "bash tools/vllm-rocm/generate-rocm-wheels-root-index.sh"
     env:
       S3_BUCKET: "vllm-wheels"
-      VARIANT: "rocm721"
+      VARIANT: "rocm700"
 
   # ROCm Job 6: Build ROCm Release Docker Image
   - label: ":docker: Build release image - x86_64 - ROCm"
@@ -681,7 +673,6 @@ steps:
   - label: "Publish nightly ROCm image to DockerHub"
     depends_on:
       - build-rocm-release-image
-    if: build.env("NIGHTLY") == "1"
     agents:
       queue: small_cpu_queue_release
     commands:

.buildkite/scripts/generate-and-upload-nightly-index.sh

@@ -1,84 +0,0 @@
-#!/usr/bin/env bash
-
-set -ex
-
-# Generate and upload wheel indices for all wheels in the commit directory.
-# This script should run once after all wheels have been built and uploaded.
-
-# ======== setup ========
-
-BUCKET="vllm-wheels"
-INDICES_OUTPUT_DIR="indices"
-DEFAULT_VARIANT_ALIAS="cu129" # align with vLLM_MAIN_CUDA_VERSION in vllm/envs.py
-PYTHON="${PYTHON_PROG:-python3}" # try to read from env var, otherwise use python3
-SUBPATH=$BUILDKITE_COMMIT
-S3_COMMIT_PREFIX="s3://$BUCKET/$SUBPATH/"
-
-# detect if python3.12+ is available
-has_new_python=$($PYTHON -c "print(1 if __import__('sys').version_info >= (3,12) else 0)")
-if [[ "$has_new_python" -eq 0 ]]; then
-    # use new python from docker
-    docker pull python:3-slim
-    PYTHON="docker run --rm -v $(pwd):/app -w /app python:3-slim python3"
-fi
-
-echo "Using python interpreter: $PYTHON"
-echo "Python version: $($PYTHON --version)"
-
-# ======== generate and upload indices ========
-
-# list all wheels in the commit directory
-echo "Existing wheels on S3:"
-aws s3 ls "$S3_COMMIT_PREFIX"
-obj_json="objects.json"
-aws s3api list-objects-v2 --bucket "$BUCKET" --prefix "$SUBPATH/" --delimiter / --output json > "$obj_json"
-mkdir -p "$INDICES_OUTPUT_DIR"
-
-# call script to generate indices for all existing wheels
-# these indices have relative paths that work as long as they are next to the wheel directory in s3
-# i.e., the wheels are always in s3://vllm-wheels/<commit>/
-# and indices can be placed in /<commit>/, or /nightly/, or /<version>/
-alias_args=()
-if [[ -n "$DEFAULT_VARIANT_ALIAS" ]]; then
-    alias_args=(--alias-to-default "$DEFAULT_VARIANT_ALIAS")
-fi
-
-# HACK: we do not need regex module here, but it is required by pre-commit hook
-# To avoid any external dependency, we simply replace it back to the stdlib re module
-sed -i 's/import regex as re/import re/g' .buildkite/scripts/generate-nightly-index.py
-$PYTHON .buildkite/scripts/generate-nightly-index.py --version "$SUBPATH" --current-objects "$obj_json" --output-dir "$INDICES_OUTPUT_DIR" --comment "commit $BUILDKITE_COMMIT" "${alias_args[@]}"
-
-# copy indices to /<commit>/ unconditionally
-echo "Uploading indices to $S3_COMMIT_PREFIX"
-aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "$S3_COMMIT_PREFIX"
-
-# copy to /nightly/ only if it is on the main branch and not a PR
-if [[ "$BUILDKITE_BRANCH" == "main" && "$BUILDKITE_PULL_REQUEST" == "false" ]]; then
-    echo "Uploading indices to overwrite /nightly/"
-    aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "s3://$BUCKET/nightly/"
-fi
-
-# detect version from any wheel in the commit directory
-# download the first wheel we find to extract version metadata
-first_wheel_key=$($PYTHON -c "import json; obj=json.load(open('$obj_json')); print(next((c['Key'] for c in obj.get('Contents', []) if c['Key'].endswith('.whl')), ''))")
-if [[ -z "$first_wheel_key" ]]; then
-    echo "Error: No wheels found in $S3_COMMIT_PREFIX"
-    exit 1
-fi
-first_wheel=$(basename "$first_wheel_key")
-aws s3 cp "s3://$BUCKET/${first_wheel_key}" "/tmp/${first_wheel}"
-version=$(unzip -p "/tmp/${first_wheel}" '**/METADATA' | grep '^Version: ' | cut -d' ' -f2)
-rm -f "/tmp/${first_wheel}"
-echo "Version in wheel: $version"
-pure_version="${version%%+*}"
-echo "Pure version (without variant): $pure_version"
-
-# re-generate and copy to /<pure_version>/ only if it does not have "dev" in the version
-if [[ "$version" != *"dev"* ]]; then
-    echo "Re-generating indices for /$pure_version/"
-    rm -rf "${INDICES_OUTPUT_DIR:?}"
-    mkdir -p "$INDICES_OUTPUT_DIR"
-    # wheel-dir is overridden to be the commit directory, so that the indices point to the correct wheel path
-    $PYTHON .buildkite/scripts/generate-nightly-index.py --version "$pure_version" --wheel-dir "$SUBPATH" --current-objects "$obj_json" --output-dir "$INDICES_OUTPUT_DIR" --comment "version $pure_version" "${alias_args[@]}"
-    aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "s3://$BUCKET/$pure_version/"
-fi

.buildkite/scripts/hardware_ci/run-cpu-distributed-smoke-test.sh

@@ -1,10 +1,9 @@
 #!/bin/bash
 set -euox pipefail
 export VLLM_CPU_CI_ENV=0
-export VLLM_CPU_KVCACHE_SPACE=1 # avoid OOM
 
 echo "--- PP+TP"
-vllm serve meta-llama/Llama-3.2-3B-Instruct -tp=2 -pp=2 --max-model-len=4096 &
+vllm serve meta-llama/Llama-3.2-3B-Instruct -tp=2 -pp=2 &
 server_pid=$!
 timeout 600 bash -c "until curl localhost:8000/v1/models > /dev/null 2>&1; do sleep 1; done" || exit 1
 vllm bench serve \
@@ -24,7 +23,7 @@ if [ "$failed_req" -ne 0 ]; then
 fi
 
 echo "--- DP+TP"
-vllm serve meta-llama/Llama-3.2-3B-Instruct -tp=2 -dp=2 --max-model-len=4096 &
+vllm serve meta-llama/Llama-3.2-3B-Instruct -tp=2 -dp=2 &
 server_pid=$!
 timeout 600 bash -c "until curl localhost:8000/v1/models > /dev/null 2>&1; do sleep 1; done" || exit 1
 vllm bench serve \

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

@@ -16,5 +16,5 @@ echo "--- :docker: Building Docker image"
 docker build --progress plain --tag "$IMAGE_NAME" --target vllm-test -f docker/Dockerfile.cpu .
 
 # Run the image, setting --shm-size=4g for tensor parallel.
-docker run --rm --cpuset-cpus="$CORE_RANGE" --cpuset-mems="$NUMA_NODE" -v ~/.cache/huggingface:/root/.cache/huggingface --privileged=true -e HF_TOKEN -e VLLM_CPU_KVCACHE_SPACE=16 -e VLLM_CPU_CI_ENV=1 -e VLLM_CPU_SIM_MULTI_NUMA=1 -e VLLM_CPU_ATTN_SPLIT_KV=0 --shm-size=4g "$IMAGE_NAME" \
+docker run --rm --cpuset-cpus="$CORE_RANGE" --cpuset-mems="$NUMA_NODE" -v ~/.cache/huggingface:/root/.cache/huggingface --privileged=true -e HF_TOKEN -e VLLM_CPU_KVCACHE_SPACE=16 -e VLLM_CPU_CI_ENV=1 -e VLLM_CPU_SIM_MULTI_NUMA=1 --shm-size=4g "$IMAGE_NAME" \
         timeout "$TIMEOUT_VAL" bash -c "set -euox pipefail; echo \"--- Print packages\"; pip list; echo \"--- Running tests\"; ${TEST_COMMAND}"

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

@@ -1,276 +0,0 @@
-#!/bin/bash
-
-# This script runs tests inside the Intel XPU docker container.
-# It mirrors the structure of run-amd-test.sh while keeping Intel-specific
-# container setup and allowing commands to be sourced from YAML or env.
-#
-# Command sources (in priority order):
-#   1) VLLM_TEST_COMMANDS env var (preferred, preserves quoting)
-#   2) Positional args (legacy)
-#   3) One or more YAML files with a commands list (test-area style)
-###############################################################################
-set -o pipefail
-
-DRY_RUN=${DRY_RUN:-0}
-if [[ "${1:-}" == "--dry-run" ]]; then
-  DRY_RUN=1
-  shift
-fi
-
-# Export Python path
-export PYTHONPATH=".."
-
-###############################################################################
-# Helper Functions
-###############################################################################
-
-cleanup_docker() {
-  docker_root=$(docker info -f '{{.DockerRootDir}}')
-  if [ -z "$docker_root" ]; then
-    echo "Failed to determine Docker root directory." >&2
-    exit 1
-  fi
-  echo "Docker root directory: $docker_root"
-
-  disk_usage=$(df "$docker_root" | tail -1 | awk '{print $5}' | sed 's/%//')
-  threshold=70
-  if [ "$disk_usage" -gt "$threshold" ]; then
-    echo "Disk usage is above $threshold%. Cleaning up Docker images and volumes..."
-    docker image prune -f
-    docker volume prune -f && docker system prune --force --filter "until=72h" --all
-    echo "Docker images and volumes cleanup completed."
-  else
-    echo "Disk usage is below $threshold%. No cleanup needed."
-  fi
-}
-
-re_quote_pytest_markers() {
-  local input="$1"
-  local output=""
-  local collecting=false
-  local marker_buf=""
-
-  local flat="${input//$'\n'/ }"
-  local restore_glob
-  restore_glob="$(shopt -p -o noglob 2>/dev/null || true)"
-  set -o noglob
-  local -a words
-  read -ra words <<< "$flat"
-  eval "$restore_glob"
-
-  for word in "${words[@]}"; do
-    if $collecting; then
-      if [[ "$word" == *"'"* ]]; then
-        if [[ -n "$marker_buf" ]]; then
-          output+="${marker_buf} "
-          marker_buf=""
-        fi
-        output+="${word} "
-        collecting=false
-        continue
-      fi
-
-      local is_boundary=false
-      case "$word" in
-        "&&"|"||"|";"|"|")
-          is_boundary=true ;;
-        --*)
-          is_boundary=true ;;
-        -[a-zA-Z])
-          is_boundary=true ;;
-        */*)
-          is_boundary=true ;;
-        *.py|*.py::*)
-          is_boundary=true ;;
-        *=*)
-          if [[ "$word" =~ ^[A-Z_][A-Z0-9_]*= ]]; then
-            is_boundary=true
-          fi
-          ;;
-      esac
-
-      if $is_boundary; then
-        if [[ "$marker_buf" == *" "* || "$marker_buf" == *"("* ]]; then
-          output+="'${marker_buf}' "
-        else
-          output+="${marker_buf} "
-        fi
-        collecting=false
-        marker_buf=""
-        if [[ "$word" == "-m" || "$word" == "-k" ]]; then
-          output+="${word} "
-          collecting=true
-        else
-          output+="${word} "
-        fi
-      else
-        if [[ -n "$marker_buf" ]]; then
-          marker_buf+=" ${word}"
-        else
-          marker_buf="${word}"
-        fi
-      fi
-    elif [[ "$word" == "-m" || "$word" == "-k" ]]; then
-      output+="${word} "
-      collecting=true
-      marker_buf=""
-    else
-      output+="${word} "
-    fi
-  done
-
-  if $collecting && [[ -n "$marker_buf" ]]; then
-    if [[ "$marker_buf" == *" "* || "$marker_buf" == *"("* ]]; then
-      output+="'${marker_buf}'"
-    else
-      output+="${marker_buf}"
-    fi
-  fi
-
-  echo "${output% }"
-}
-
-apply_intel_test_overrides() {
-  local cmds="$1"
-  # Placeholder for Intel-specific exclusions/overrides.
-  echo "$cmds"
-}
-
-is_yaml_file() {
-  local p="$1"
-  [[ -f "$p" && "$p" == *.yaml ]]
-}
-
-extract_yaml_commands() {
-  local yaml_path="$1"
-  awk '
-    $1 == "commands:" { in_cmds=1; next }
-    in_cmds && $0 ~ /^[[:space:]]*-[[:space:]]/ {
-      sub(/^[[:space:]]*-[[:space:]]/, "");
-      print;
-      next
-    }
-    in_cmds && $0 ~ /^[^[:space:]]/ { exit }
-  ' "$yaml_path"
-}
-
-###############################################################################
-# Main
-###############################################################################
-
-default_image_name="${REGISTRY}/${REPO}:${BUILDKITE_COMMIT}-xpu"
-#default_image_name="public.ecr.aws/q9t5s3a7/vllm-ci-test-repo:${BUILDKITE_COMMIT}-xpu"
-image_name="${IMAGE_TAG_XPU:-${default_image_name}}"
-container_name="xpu_${BUILDKITE_COMMIT}_$(tr -dc A-Za-z0-9 < /dev/urandom | head -c 10; echo)"
-
-# ---- Command source selection ----
-commands=""
-if [[ -n "${VLLM_TEST_COMMANDS:-}" ]]; then
-  commands="${VLLM_TEST_COMMANDS}"
-  echo "Commands sourced from VLLM_TEST_COMMANDS (quoting preserved)"
-elif [[ $# -gt 0 ]]; then
-  all_yaml=true
-  for arg in "$@"; do
-    if ! is_yaml_file "$arg"; then
-      all_yaml=false
-      break
-    fi
-  done
-
-  if $all_yaml; then
-    for yaml in "$@"; do
-      mapfile -t COMMANDS < <(extract_yaml_commands "$yaml")
-      if [[ ${#COMMANDS[@]} -eq 0 ]]; then
-        echo "Error: No commands found in ${yaml}" >&2
-        exit 1
-      fi
-      for cmd in "${COMMANDS[@]}"; do
-        if [[ -z "$commands" ]]; then
-          commands="${cmd}"
-        else
-          commands+=" && ${cmd}"
-        fi
-      done
-    done
-    echo "Commands sourced from YAML files: $*"
-  else
-    commands="$*"
-    echo "Commands sourced from positional args (legacy mode)"
-  fi
-else
-  SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
-  DEFAULT_YAML="${SCRIPT_DIR}/intel-test.yaml"
-  if [[ ! -f "${DEFAULT_YAML}" ]]; then
-    echo "Error: YAML file not found: ${DEFAULT_YAML}" >&2
-    exit 1
-  fi
-  mapfile -t COMMANDS < <(extract_yaml_commands "${DEFAULT_YAML}")
-  if [[ ${#COMMANDS[@]} -eq 0 ]]; then
-    echo "Error: No commands found in ${DEFAULT_YAML}" >&2
-    exit 1
-  fi
-  for cmd in "${COMMANDS[@]}"; do
-    if [[ -z "$commands" ]]; then
-      commands="${cmd}"
-    else
-      commands+=" && ${cmd}"
-    fi
-  done
-  echo "Commands sourced from default YAML: ${DEFAULT_YAML}"
-fi
-
-if [[ -z "$commands" ]]; then
-  echo "Error: No test commands provided." >&2
-  exit 1
-fi
-
-echo "Raw commands: $commands"
-commands=$(re_quote_pytest_markers "$commands")
-echo "After re-quoting: $commands"
-commands=$(apply_intel_test_overrides "$commands")
-echo "Final commands: $commands"
-
-# Dry-run mode prints final commands and exits before Docker.
-if [[ "$DRY_RUN" == "1" ]]; then
-  echo "DRY_RUN=1 set, skipping Docker execution."
-  exit 0
-fi
-
-# --- Docker housekeeping ---
-cleanup_docker
-
-# --- Build or pull test image ---
-if [[ -n "${IMAGE_TAG_XPU:-}" ]]; then
-  echo "Using prebuilt XPU image: ${IMAGE_TAG_XPU}"
-  docker pull "${IMAGE_TAG_XPU}"
-else
-  echo "Using prebuilt XPU image: ${image_name}"
-  docker pull "${image_name}"
-fi
-
-remove_docker_container() {
-  docker rm -f "${container_name}" || true
-  docker image rm -f "${image_name}" || true
-  docker system prune -f || true
-}
-trap remove_docker_container EXIT
-
-# --- Single-node job ---
-
-if [[ -z "${ZE_AFFINITY_MASK:-}" ]]; then
-  echo "Warning: ZE_AFFINITY_MASK is not set. Proceeding without device affinity." >&2
-fi
-
-docker run \
-    --device /dev/dri:/dev/dri \
-    --net=host \
-    --ipc=host \
-    --privileged \
-    -v /dev/dri/by-path:/dev/dri/by-path \
-    --entrypoint="" \
-    -e "HF_TOKEN=${HF_TOKEN:-}" \
-    -e "ZE_AFFINITY_MASK=${ZE_AFFINITY_MASK:-}" \
-    -e "CMDS=${commands}" \
-    --name "${container_name}" \
-    "${image_name}" \
-    bash -c 'set -e; echo "ZE_AFFINITY_MASK is ${ZE_AFFINITY_MASK:-}"; eval "$CMDS"'

.buildkite/scripts/upload-nightly-wheels.sh

@@ -2,14 +2,27 @@
 
 set -ex
 
-# Upload a single wheel to S3 (rename linux -> manylinux).
-# Index generation is handled separately by generate-and-upload-nightly-index.sh.
+# ======== part 0: setup ========
 
 BUCKET="vllm-wheels"
+INDICES_OUTPUT_DIR="indices"
+DEFAULT_VARIANT_ALIAS="cu129" # align with vLLM_MAIN_CUDA_VERSION in vllm/envs.py
+PYTHON=${PYTHON_PROG:=python3} # try to read from env var, otherwise use python3
 SUBPATH=$BUILDKITE_COMMIT
 S3_COMMIT_PREFIX="s3://$BUCKET/$SUBPATH/"
 
-# ========= collect, rename & upload the wheel ==========
+# detect if python3.10+ is available
+has_new_python=$($PYTHON -c "print(1 if __import__('sys').version_info >= (3,12) else 0)")
+if [[ "$has_new_python" -eq 0 ]]; then
+    # use new python from docker
+    docker pull python:3-slim
+    PYTHON="docker run --rm -v $(pwd):/app -w /app python:3-slim python3"
+fi
+
+echo "Using python interpreter: $PYTHON"
+echo "Python version: $($PYTHON --version)"
+
+# ========= part 1: collect, rename & upload the wheel ==========
 
 # Assume wheels are in artifacts/dist/*.whl
 wheel_files=(artifacts/dist/*.whl)
@@ -39,8 +52,56 @@ echo "Renamed wheel to: $wheel"
 # Extract the version from the wheel
 version=$(unzip -p "$wheel" '**/METADATA' | grep '^Version: ' | cut -d' ' -f2)
 echo "Version in wheel: $version"
+pure_version="${version%%+*}"
+echo "Pure version (without variant): $pure_version"
 
 # copy wheel to its own bucket
 aws s3 cp "$wheel" "$S3_COMMIT_PREFIX"
 
-echo "Wheel uploaded. Index generation is handled by a separate step."
+# ========= part 2: generate and upload indices ==========
+# generate indices for all existing wheels in the commit directory
+# this script might be run multiple times if there are multiple variants being built
+# so we need to guarantee there is little chance for "TOCTOU" issues
+# i.e., one process is generating indices while another is uploading a new wheel
+# so we need to ensure no time-consuming operations happen below
+
+# list all wheels in the commit directory
+echo "Existing wheels on S3:"
+aws s3 ls "$S3_COMMIT_PREFIX"
+obj_json="objects.json"
+aws s3api list-objects-v2 --bucket "$BUCKET" --prefix "$SUBPATH/" --delimiter / --output json > "$obj_json"
+mkdir -p "$INDICES_OUTPUT_DIR"
+
+# call script to generate indices for all existing wheels
+# this indices have relative paths that could work as long as it is next to the wheel directory in s3
+# i.e., the wheels are always in s3://vllm-wheels/<commit>/
+# and indices can be placed in /<commit>/, or /nightly/, or /<version>/
+alias_args=()
+if [[ -n "$DEFAULT_VARIANT_ALIAS" ]]; then
+    alias_args=(--alias-to-default "$DEFAULT_VARIANT_ALIAS")
+fi
+
+# HACK: we do not need regex module here, but it is required by pre-commit hook
+# To avoid any external dependency, we simply replace it back to the stdlib re module
+sed -i 's/import regex as re/import re/g' .buildkite/scripts/generate-nightly-index.py
+$PYTHON .buildkite/scripts/generate-nightly-index.py --version "$SUBPATH" --current-objects "$obj_json" --output-dir "$INDICES_OUTPUT_DIR" --comment "commit $BUILDKITE_COMMIT" "${alias_args[@]}"
+
+# copy indices to /<commit>/ unconditionally
+echo "Uploading indices to $S3_COMMIT_PREFIX"
+aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "$S3_COMMIT_PREFIX"
+
+# copy to /nightly/ only if it is on the main branch and not a PR 
+if [[ "$BUILDKITE_BRANCH" == "main" && "$BUILDKITE_PULL_REQUEST" == "false" ]]; then
+    echo "Uploading indices to overwrite /nightly/"
+    aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "s3://$BUCKET/nightly/"
+fi
+
+# re-generate and copy to /<pure_version>/ only if it does not have "dev" in the version
+if [[ "$version" != *"dev"* ]]; then
+    echo "Re-generating indices for /$pure_version/"
+    rm -rf "${INDICES_OUTPUT_DIR:?}/*"
+    mkdir -p "$INDICES_OUTPUT_DIR"
+    # wheel-dir is overridden to be the commit directory, so that the indices point to the correct wheel path
+    $PYTHON .buildkite/scripts/generate-nightly-index.py --version "$pure_version" --wheel-dir "$SUBPATH" --current-objects "$obj_json" --output-dir "$INDICES_OUTPUT_DIR" --comment "version $pure_version" "${alias_args[@]}"
+    aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "s3://$BUCKET/$pure_version/"
+fi

.buildkite/test-amd.yaml

@@ -812,7 +812,7 @@ steps:
   commands:
   - apt-get update && apt-get install -y curl libsodium23
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s model_executor -m '(not slow_test)'
+  - pytest -v -s model_executor
   - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py
 
 
@@ -1242,7 +1242,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - TARGET_TEST_SUITE=L4 pytest basic_correctness/ -v -s -m 'distributed(num_gpus=2)'
-  - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s model_executor/model_loader/test_sharded_state_loader.py -m '(not slow_test)'
+  - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s model_executor/model_loader/test_sharded_state_loader.py
   - pytest models/test_transformers.py -v -s -m 'distributed(num_gpus=2)'
   - pytest models/language -v -s -m 'distributed(num_gpus=2)'
   - pytest models/multimodal -v -s -m 'distributed(num_gpus=2)' --ignore models/multimodal/generation/test_whisper.py
@@ -1801,19 +1801,6 @@ steps:
   - tests/v1/e2e
   commands:
     - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "eagle_correctness_heavy"
-  
-
-- label: V1 e2e (4xH100-4xMI325) # TBD
-  timeout_in_minutes: 180
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
-  agent_pool: mi325_4
-  optional: true
-  source_file_dependencies:
-    - vllm/v1/attention/backends/utils.py
-    - vllm/v1/worker/gpu_model_runner.py
-    - tests/v1/e2e/test_hybrid_chunked_prefill.py
-  commands:
-    - pytest -v -s v1/e2e/test_hybrid_chunked_prefill.py
 
 
 - label: V1 Spec Decode # TBD
@@ -2514,7 +2501,7 @@ steps:
   - tests/models/
   commands:
   - TARGET_TEST_SUITE=L4 pytest basic_correctness/ -v -s -m 'distributed(num_gpus=2)'
-  - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s model_executor/model_loader/test_sharded_state_loader.py -m '(not slow_test)'
+  - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s model_executor/model_loader/test_sharded_state_loader.py
   - pytest models/test_transformers.py -v -s -m 'distributed(num_gpus=2)'
   - pytest models/language -v -s -m 'distributed(num_gpus=2)'
   - pytest models/multimodal -v -s -m 'distributed(num_gpus=2)' --ignore models/multimodal/generation/test_whisper.py

.buildkite/test_areas/expert_parallelism.yaml

@@ -8,10 +8,8 @@ steps:
   source_file_dependencies:
   - vllm/distributed/eplb
   - tests/distributed/test_eplb_algo.py
-  - tests/distributed/test_eplb_utils.py
   commands:
   - pytest -v -s distributed/test_eplb_algo.py
-  - pytest -v -s distributed/test_eplb_utils.py
 
 - label: EPLB Execution
   timeout_in_minutes: 20

.buildkite/test_areas/kernels.yaml

@@ -10,20 +10,7 @@ steps:
   - tests/kernels/test_top_k_per_row.py
   - tests/kernels/test_concat_mla_q.py
   commands:
-    - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py kernels/test_top_k_per_row.py kernels/test_concat_mla_q.py
-
-- label: Kernels MiniMax Reduce RMS Test (2 GPUs)
-  timeout_in_minutes: 15
-  num_devices: 2
-  device: h100
-  source_file_dependencies:
-  - csrc/minimax_reduce_rms_kernel.cu
-  - csrc/minimax_reduce_rms_kernel.h
-  - vllm/model_executor/layers/mamba/linear_attn.py
-  - vllm/model_executor/layers/mamba/lamport_workspace.py
-  - tests/kernels/core/test_minimax_reduce_rms.py
-  commands:
-    - pytest -v -s kernels/core/test_minimax_reduce_rms.py
+    - pytest -v -s kernels/core kernels/test_top_k_per_row.py kernels/test_concat_mla_q.py
 
 - label: Kernels Attention Test %N
   timeout_in_minutes: 35

.buildkite/test_areas/model_executor.yaml

@@ -13,5 +13,5 @@ steps:
   commands:
     - apt-get update && apt-get install -y curl libsodium23
     - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-    - pytest -v -s model_executor -m '(not slow_test)'
+    - pytest -v -s model_executor
     - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py

.buildkite/test_areas/models_basic.yaml

@@ -69,18 +69,3 @@ steps:
     - python3 examples/offline_inference/vision_language.py --model-type qwen2_5_vl
     # Whisper needs spawn method to avoid deadlock
     - VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/offline_inference/audio_language.py --model-type whisper
-
-- label: Transformers Backward Compatibility Models Test
-  working_dir: "/vllm-workspace/"
-  optional: true
-  soft_fail: true
-  commands:
-    - pip install transformers==4.57.5
-    - pytest -v -s tests/models/test_initialization.py
-    - pytest -v -s tests/models/test_transformers.py
-    - pytest -v -s tests/models/multimodal/processing/
-    - pytest -v -s tests/models/multimodal/test_mapping.py
-    - python3 examples/offline_inference/basic/chat.py
-    - python3 examples/offline_inference/vision_language.py --model-type qwen2_5_vl
-    # Whisper needs spawn method to avoid deadlock
-    - VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/offline_inference/audio_language.py --model-type whisper

.buildkite/test_areas/models_distributed.yaml

@@ -14,7 +14,7 @@ steps:
   - tests/models/
   commands:
   - TARGET_TEST_SUITE=L4 pytest basic_correctness/ -v -s -m 'distributed(num_gpus=2)'
-  - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s model_executor/model_loader/test_sharded_state_loader.py -m '(not slow_test)'
+  - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s model_executor/model_loader/test_sharded_state_loader.py
   # Avoid importing model tests that cause CUDA reinitialization error
   - pytest models/test_transformers.py -v -s -m 'distributed(num_gpus=2)'
   - pytest models/language -v -s -m 'distributed(num_gpus=2)'

.github/CODEOWNERS

@@ -9,7 +9,6 @@
 /vllm/model_executor/layers/fused_moe @mgoin @pavanimajety
 /vllm/model_executor/layers/quantization @mgoin @robertgshaw2-redhat @tlrmchlsmth @yewentao256 @pavanimajety
 /vllm/model_executor/layers/mamba @tdoublep
-/vllm/model_executor/layers/mamba/gdn_linear_attn.py @tdoublep @ZJY0516
 /vllm/model_executor/model_loader @22quinn
 /vllm/model_executor/layers/batch_invariant.py @yewentao256 
 /vllm/multimodal @DarkLight1337 @ywang96 @NickLucche @tjtanaa
@@ -49,7 +48,6 @@ CMakeLists.txt @tlrmchlsmth @LucasWilkinson
 /vllm/v1/attention/backends/mla @pavanimajety
 /vllm/v1/attention/backends/flashinfer.py @mgoin @pavanimajety
 /vllm/v1/attention/backends/triton_attn.py @tdoublep
-/vllm/v1/attention/backends/gdn_attn.py @ZJY0516
 /vllm/v1/core @WoosukKwon @robertgshaw2-redhat @njhill @ywang96 @alexm-redhat @heheda12345 @ApostaC @orozery
 /vllm/v1/sample @22quinn @houseroad @njhill
 /vllm/v1/spec_decode @benchislett @luccafong @MatthewBonanni
@@ -144,7 +142,6 @@ mkdocs.yaml @hmellor
 # Kernels
 /vllm/v1/attention/ops/chunked_prefill_paged_decode.py @tdoublep
 /vllm/v1/attention/ops/triton_unified_attention.py @tdoublep
-/vllm/model_executor/layers/fla @ZJY0516
 
 # ROCm related: specify owner with write access to notify AMD folks for careful code review
 /vllm/**/*rocm* @tjtanaa

.github/mergify.yml

@@ -234,36 +234,6 @@ pull_request_rules:
       add:
         - rocm
 
-- name: label-xpu
-  description: Automatically apply intel-gpu label
-  conditions:
-    - label != stale
-    - or:
-      - files~=^docker/Dockerfile.xpu
-      - files~=^\\.buildkite/intel_jobs/
-      - files=\.buildkite/ci_config_intel.yaml
-      - files=vllm/model_executor/layers/fused_moe/xpu_fused_moe.py
-      - files=vllm/model_executor/kernels/linear/mixed_precision/xpu.py
-      - files=vllm/model_executor/kernels/linear/scaled_mm/xpu.py
-      - files=vllm/distributed/device_communicators/xpu_communicator.py
-      - files=vllm/v1/attention/backends/mla/xpu_mla_sparse.py
-      - files=vllm/v1/attention/ops/xpu_mla_sparse.py
-      - files=vllm/v1/worker/xpu_worker.py
-      - files=vllm/v1/worker/xpu_model_runner.py
-      - files=vllm/_xpu_ops.py
-      - files~=^vllm/lora/ops/xpu_ops
-      - files=vllm/lora/punica_wrapper/punica_xpu.py
-      - files=vllm/platforms/xpu.py
-      - title~=(?i)Intel gpu
-      - title~=(?i)XPU
-      - title~=(?i)Intel
-      - title~=(?i)BMG
-      - title~=(?i)Arc
-  actions:
-    label:
-      add:
-        - intel-gpu
-
 - name: label-cpu
   description: Automatically apply cpu label
   conditions:

CMakeLists.txt

@@ -306,8 +306,6 @@ set(VLLM_EXT_SRC
   "csrc/torch_bindings.cpp")
 
 if(VLLM_GPU_LANG STREQUAL "CUDA")
-  list(APPEND VLLM_EXT_SRC "csrc/minimax_reduce_rms_kernel.cu")
-
   SET(CUTLASS_ENABLE_HEADERS_ONLY ON CACHE BOOL "Enable only the header library")
 
   # Set CUTLASS_REVISION. Used for FetchContent. Also fixes some bogus messages when building.
@@ -365,7 +363,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
   # - sm80 doesn't support fp8 computation
   # - sm90 and sm100 don't support QMMA.16832.F32.E4M3.E4M3 SAAS instruction
   # so we only enable fp8 computation for SM89 (e.g. RTX 40x0)  and 12.0 (e.g. RTX 50x0)
-  cuda_archs_loose_intersection(MARLIN_FP8_ARCHS "8.9;12.0;12.1" "${CUDA_ARCHS}")
+  cuda_archs_loose_intersection(MARLIN_FP8_ARCHS "8.9;12.0" "${CUDA_ARCHS}")
   # marlin arches for other files
   cuda_archs_loose_intersection(MARLIN_OTHER_ARCHS "7.5;8.0+PTX" "${CUDA_ARCHS}")
 
@@ -525,12 +523,12 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
   endif()
 
 
-  # The cutlass_scaled_mm kernels for Blackwell SM12x (c3x, i.e. CUTLASS 3.x) require
+  # The cutlass_scaled_mm kernels for Geforce Blackwell SM120 (c3x, i.e. CUTLASS 3.x) require
   # CUDA 12.8 or later
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
     cuda_archs_loose_intersection(SCALED_MM_ARCHS "12.0f" "${CUDA_ARCHS}")
   else()
-    cuda_archs_loose_intersection(SCALED_MM_ARCHS "12.0a;12.1a" "${CUDA_ARCHS}")
+    cuda_archs_loose_intersection(SCALED_MM_ARCHS "12.0a" "${CUDA_ARCHS}")
   endif()
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND SCALED_MM_ARCHS)
     set(SRCS
@@ -618,12 +616,12 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
     endif()
   endif()
 
-  # The nvfp4_scaled_mm_sm120 kernels for Blackwell SM12x require
+  # The nvfp4_scaled_mm_sm120 kernels for Geforce Blackwell SM120 require
   # CUDA 12.8 or later
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
     cuda_archs_loose_intersection(FP4_ARCHS "12.0f" "${CUDA_ARCHS}")
   else()
-    cuda_archs_loose_intersection(FP4_ARCHS "12.0a;12.1a" "${CUDA_ARCHS}")
+    cuda_archs_loose_intersection(FP4_ARCHS "12.0a" "${CUDA_ARCHS}")
   endif()
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND FP4_ARCHS)
     set(SRCS
@@ -1052,7 +1050,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
   # - sm80 doesn't support fp8 computation
   # - sm90 and sm100 don't support QMMA.16832.F32.E4M3.E4M3 SAAS instruction
   # so we only enable fp8 computation for SM89 (e.g. RTX 40x0)  and 12.0 (e.g. RTX 50x0)
-  cuda_archs_loose_intersection(MARLIN_MOE_FP8_ARCHS "8.9;12.0;12.1" "${CUDA_ARCHS}")
+  cuda_archs_loose_intersection(MARLIN_MOE_FP8_ARCHS "8.9;12.0" "${CUDA_ARCHS}")
   # moe marlin arches for other files
   cuda_archs_loose_intersection(MARLIN_MOE_OTHER_ARCHS "7.5;8.0+PTX" "${CUDA_ARCHS}")
   if (MARLIN_MOE_OTHER_ARCHS)

benchmarks/attention_benchmarks/benchmark.py

@@ -546,7 +546,10 @@ def main():
         args.prefill_backends = yaml_config.get("prefill_backends", None)
 
         # Check for special modes
-        args.mode = yaml_config.get("mode", None)
+        if "mode" in yaml_config:
+            args.mode = yaml_config["mode"]
+        else:
+            args.mode = None
 
         # Batch specs and sizes
         # Support both explicit batch_specs and generated batch_spec_ranges
@@ -569,7 +572,10 @@ def main():
             elif "batch_specs" in yaml_config:
                 args.batch_specs = yaml_config["batch_specs"]
 
-        args.batch_sizes = yaml_config.get("batch_sizes", None)
+        if "batch_sizes" in yaml_config:
+            args.batch_sizes = yaml_config["batch_sizes"]
+        else:
+            args.batch_sizes = None
 
         # Model config
         if "model" in yaml_config:

benchmarks/kernels/benchmark_moe.py

@@ -627,8 +627,9 @@ class BenchmarkWorker:
                 need_device_guard = True
 
         with (
-            # Ray restricts each worker to one GPU; use local index 0
-            torch.accelerator.device_index(0) if need_device_guard else nullcontext()
+            torch.accelerator.device_index(self.device_id)
+            if need_device_guard
+            else nullcontext()
         ):
             for idx, config in enumerate(tqdm(search_space)):
                 try:

cmake/external_projects/qutlass.cmake

@@ -32,16 +32,16 @@ endif()
 message(STATUS "[QUTLASS] QuTLASS is available at ${qutlass_SOURCE_DIR}")
 
 if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-  cuda_archs_loose_intersection(QUTLASS_ARCHS "10.0f;12.0f" "${CUDA_ARCHS}")
+  cuda_archs_loose_intersection(QUTLASS_ARCHS "12.0a;10.0f" "${CUDA_ARCHS}")
 else()
-  cuda_archs_loose_intersection(QUTLASS_ARCHS "12.0a;12.1a;10.0a;10.3a" "${CUDA_ARCHS}")
+  cuda_archs_loose_intersection(QUTLASS_ARCHS "12.0a;10.0a;10.3a" "${CUDA_ARCHS}")
 endif()
 
 if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND QUTLASS_ARCHS)
 
   if(QUTLASS_ARCHS MATCHES "10\\.(0a|3a|0f)")
     set(QUTLASS_TARGET_CC 100)
-  elseif(QUTLASS_ARCHS MATCHES "12\\.[01][af]?")
+  elseif(QUTLASS_ARCHS MATCHES "12\\.0a")
     set(QUTLASS_TARGET_CC 120)
   else()
     message(FATAL_ERROR "[QUTLASS] internal error parsing CUDA_ARCHS='${QUTLASS_ARCHS}'.")
@@ -96,7 +96,7 @@ else()
       "[QUTLASS] Skipping build: CUDA 12.8 or newer is required (found ${CMAKE_CUDA_COMPILER_VERSION}).")
   else()
     message(STATUS
-      "[QUTLASS] Skipping build: no supported arch (12.0f / 10.0f) found in "
+      "[QUTLASS] Skipping build: no supported arch (12.0a / 10.0a) found in "
       "CUDA_ARCHS='${CUDA_ARCHS}'.")
   endif()
 endif()

cmake/utils.cmake

@@ -355,11 +355,8 @@ function(cuda_archs_loose_intersection OUT_CUDA_ARCHS SRC_CUDA_ARCHS TGT_CUDA_AR
   list(REMOVE_DUPLICATES _PTX_ARCHS)
   list(REMOVE_DUPLICATES _SRC_CUDA_ARCHS)
 
-  # Handle architecture-specific suffixes (a/f) for SRC entries.
-  # First try exact base match (x.y), then cross-suffix match (x.ya / x.yf).
-  # For 'f' (family) suffix: if no exact/cross match, fall back to major-version
-  # match — e.g. SRC="12.0f" matches TGT="12.1a" since SM121 is in the SM12x
-  # family. The output uses TGT's value to preserve the user's compilation flags.
+  # If x.0a or x.0f is in SRC_CUDA_ARCHS and x.0 is in CUDA_ARCHS then we should
+  # remove x.0a or x.0f from SRC_CUDA_ARCHS and add x.0a or x.0f to _CUDA_ARCHS
   set(_CUDA_ARCHS)
   foreach(_arch ${_SRC_CUDA_ARCHS})
     if(_arch MATCHES "[af]$")
@@ -368,38 +365,6 @@ function(cuda_archs_loose_intersection OUT_CUDA_ARCHS SRC_CUDA_ARCHS TGT_CUDA_AR
       if ("${_base}" IN_LIST TGT_CUDA_ARCHS)
         list(REMOVE_ITEM _TGT_CUDA_ARCHS "${_base}")
         list(APPEND _CUDA_ARCHS "${_arch}")
-      elseif("${_base}a" IN_LIST _TGT_CUDA_ARCHS)
-        list(REMOVE_ITEM _TGT_CUDA_ARCHS "${_base}a")
-        list(APPEND _CUDA_ARCHS "${_base}a")
-      elseif("${_base}f" IN_LIST _TGT_CUDA_ARCHS)
-        list(REMOVE_ITEM _TGT_CUDA_ARCHS "${_base}f")
-        list(APPEND _CUDA_ARCHS "${_base}f")
-      elseif(_arch MATCHES "f$")
-        # Family suffix: match any TGT entry in the same major version family.
-        string(REGEX REPLACE "^([0-9]+)\\..*$" "\\1" _src_major "${_base}")
-        foreach(_tgt ${_TGT_CUDA_ARCHS})
-          string(REGEX REPLACE "[af]$" "" _tgt_base "${_tgt}")
-          string(REGEX REPLACE "^([0-9]+)\\..*$" "\\1" _tgt_major "${_tgt_base}")
-          if(_tgt_major STREQUAL _src_major)
-            list(REMOVE_ITEM _TGT_CUDA_ARCHS "${_tgt}")
-            list(APPEND _CUDA_ARCHS "${_tgt}")
-            break()
-          endif()
-        endforeach()
-      endif()
-    endif()
-  endforeach()
-
-  # Symmetric handling: if TGT has x.ya/f and SRC has x.y (without suffix),
-  # preserve TGT's suffix in the output.
-  set(_tgt_copy ${_TGT_CUDA_ARCHS})
-  foreach(_arch ${_tgt_copy})
-    if(_arch MATCHES "[af]$")
-      string(REGEX REPLACE "[af]$" "" _base "${_arch}")
-      if ("${_base}" IN_LIST _SRC_CUDA_ARCHS)
-        list(REMOVE_ITEM _TGT_CUDA_ARCHS "${_arch}")
-        list(REMOVE_ITEM _SRC_CUDA_ARCHS "${_base}")
-        list(APPEND _CUDA_ARCHS "${_arch}")
       endif()
     endif()
   endforeach()

csrc/minimax_reduce_rms_kernel.cu

@@ -1,879 +0,0 @@
-
-/*
- * Copyright (c) 2026, NVIDIA CORPORATION.  All rights reserved.
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#include <cooperative_groups.h>
-#include <cuda_runtime.h>
-
-#include <torch/cuda.h>
-#include <ATen/cuda/CUDAContext.h>
-#include <c10/cuda/CUDAGuard.h>
-
-#include "cuda_compat.h"
-#include "cuda_utils.h"
-#include "core/registration.h"
-#include "minimax_reduce_rms_kernel.h"
-
-#include <algorithm>
-
-#define FINAL_MASK 0xffffffff
-#define MINIMAX_REDUCE_RMS_WARP_SIZE 32
-
-namespace vllm {
-namespace tensorrt_llm {
-
-template <int NRanks>
-struct LamportComm {
-  __device__ __forceinline__ LamportComm(void** workspace, int rank) {
-    counter_ptr = &reinterpret_cast<int*>(workspace[NRanks * 3])[0];
-    flag_ptr = &reinterpret_cast<int*>(workspace[NRanks * 3])[2];
-    clear_ptr = &reinterpret_cast<int64_t*>(workspace[NRanks * 3 + 1])[0];
-    flag_value = *flag_ptr;
-    auto comm_size = reinterpret_cast<int64_t*>(workspace[NRanks * 3 + 1])[1];
-    clear_size = *clear_ptr;
-    int data_offset = flag_value % 3;
-    int clear_offset = (flag_value + 2) % 3;
-    for (int r = 0; r < NRanks; ++r) {
-      data_bufs[r] = reinterpret_cast<uint8_t*>(workspace[2 * NRanks + r]) +
-                     data_offset * comm_size;
-    }
-    clear_buf = reinterpret_cast<uint8_t*>(workspace[2 * NRanks + rank]) +
-                clear_offset * comm_size;
-    __syncthreads();
-    if (threadIdx.x == 0) {
-      atomicAdd(counter_ptr, 1);
-    }
-  }
-
-  __device__ __forceinline__ void update(int64_t new_clear_size) {
-    if (blockIdx.x == 0 && threadIdx.x == 0) {
-      while (*reinterpret_cast<int volatile*>(counter_ptr) != gridDim.x) {
-      }
-      *flag_ptr = (flag_value + 1) % 3;
-      *clear_ptr = new_clear_size;
-      *counter_ptr = 0;
-    }
-  }
-
-  int* counter_ptr;
-  int* flag_ptr;
-  int64_t* clear_ptr;
-  uint8_t* data_bufs[NRanks];
-  uint8_t* clear_buf;
-  int64_t clear_size;
-  int flag_value;
-};
-
-__device__ __forceinline__ bool is_neg_zero(float v) {
-  return *reinterpret_cast<uint32_t*>(&v) == 0x80000000;
-}
-
-__device__ __forceinline__ bool is_neg_zero(float4 v) {
-  return is_neg_zero(v.x) || is_neg_zero(v.y) || is_neg_zero(v.z) ||
-         is_neg_zero(v.w);
-}
-
-__device__ __forceinline__ float4 get_neg_zero() {
-  float4 vec;
-#pragma unroll
-  for (int i = 0; i < 4; ++i) {
-    reinterpret_cast<uint32_t*>(&vec)[i] = 0x80000000;
-  }
-  return vec;
-}
-
-template <int Dim>
-__device__ __forceinline__ float rms_rsqrt(float& v, float eps) {
-  constexpr float kInvDim = 1.0F / static_cast<float>(Dim);
-  v = rsqrtf((v * kInvDim) + eps);
-  return v;
-}
-
-template <int Dim>
-__device__ __forceinline__ float4 rms_rsqrt(float4& v, float eps) {
-  constexpr float kInvDim = 1.0F / static_cast<float>(Dim);
-  v.x = rsqrtf((v.x * kInvDim) + eps);
-  v.y = rsqrtf((v.y * kInvDim) + eps);
-  v.z = rsqrtf((v.z * kInvDim) + eps);
-  v.w = rsqrtf((v.w * kInvDim) + eps);
-  return v;
-}
-__device__ __forceinline__ float4 ld_global_volatile(float4* addr) {
-  float4 val;
-  asm volatile("ld.volatile.global.v4.f32 {%0, %1, %2, %3}, [%4];"
-               : "=f"(val.x), "=f"(val.y), "=f"(val.z), "=f"(val.w)
-               : "l"(addr));
-  return val;
-}
-
-__device__ __forceinline__ float ld_global_volatile(float* addr) {
-  float val;
-  asm volatile("ld.volatile.global.f32 %0, [%1];" : "=f"(val) : "l"(addr));
-  return val;
-}
-
-// Used by the scalar (non-float4) kernel only
-template <typename T, int NUM>
-__inline__ __device__ T warpReduceSumV2(T* val) {
-#pragma unroll
-  for (int i = 0; i < NUM; i++) {
-#pragma unroll
-    for (int mask = 16; mask > 0; mask >>= 1)
-      val[i] += __shfl_xor_sync(FINAL_MASK, val[i], mask, 32);
-  }
-  return (T)(0.0f);
-}
-
-template <typename T, int NUM>
-__inline__ __device__ T blockReduceSumV2(T* val) {
-  static __shared__ T shared[NUM][33];
-  int lane = threadIdx.x & 0x1f;
-  int wid = threadIdx.x >> 5;
-
-  warpReduceSumV2<T, NUM>(val);
-
-  if (lane == 0) {
-#pragma unroll
-    for (int i = 0; i < NUM; i++) {
-      shared[i][wid] = val[i];
-    }
-  }
-
-  __syncthreads();
-
-  bool is_mask = threadIdx.x < (blockDim.x / 32.f);
-#pragma unroll
-  for (int i = 0; i < NUM; i++) {
-    val[i] = is_mask ? shared[i][lane] : (T)(0.0f);
-  }
-  warpReduceSumV2<T, NUM>(val);
-  return (T)0.0f;
-}
-
-// for float4 version
-template <uint32_t kNumThreads, typename T, int ArraySize = 4>
-__device__ __forceinline__ void local_warp_reduce_sum_array(
-    T* value_ptr, uint32_t active_mask = 0xffffffffu) {
-  static_assert(kNumThreads >= 1 &&
-                kNumThreads <= MINIMAX_REDUCE_RMS_WARP_SIZE);
-#pragma unroll
-  for (int i = 0; i < ArraySize; ++i) {
-#pragma unroll
-    for (int mask = kNumThreads / 2; mask > 0; mask >>= 1) {
-      value_ptr[i] += __shfl_xor_sync(active_mask, value_ptr[i], mask,
-                                      MINIMAX_REDUCE_RMS_WARP_SIZE);
-    }
-  }
-}
-
-constexpr int next_pow2(int val) {
-  int result = 1;
-  while (result < val) {
-    result <<= 1;
-  }
-  return result;
-}
-
-// ---------------------------------------------------------------------------
-
-template <typename DType>
-class IndexHelper {
- public:
-  __device__ __forceinline__ IndexHelper(MiniMaxReduceRMSParams const& params) {
-#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
-    namespace cg = cooperative_groups;
-    cg::cluster_group cluster = cg::this_cluster();
-    cg::grid_group grid = cg::this_grid();
-    token_id = grid.cluster_rank();
-    access_id_in_token = cluster.thread_rank();
-    token_stride = grid.num_clusters();
-#else
-    token_id = blockIdx.x;
-    access_id_in_token = threadIdx.x;
-    token_stride = gridDim.x;
-#endif
-    access_id = token_id * params.hidden_dim / kElemsPerAccess<DType> +
-                access_id_in_token;
-    access_stride = token_stride * params.hidden_dim / kElemsPerAccess<DType>;
-    tot_access = params.size_q / kElemsPerAccess<DType>;
-  }
-
-  int token_id;
-  int access_id_in_token;
-  int token_stride;
-  int access_id;
-  int access_stride;
-  int tot_access;
-};
-
-/**
-* this kernel is used to for minimax attention module
-* input tensor [total_tokens, hidden_dim / tp_size], fp32
-* rms weight [hidden_dim / tp_size], bf16
-step 1: reduce from single rank to get the variance sum (reduce(input^2,
-dim=-1)) step 2: reduce from all ranks to get the variance sum
-(all_reduce(variance_sum)) step 3: calculate the rms norm (input *
-rsqrt(variance + eps)) in this case, max hidden_dim is 6144 (float data), for
-each token, we only need 6144 / 4 / tp_size = (1536 / tp_size) threads so we can
-assume cluster size is 1 (tp_size >= 2)
- */
-template <typename DType, int NRanks>
-__global__ void __launch_bounds__(1024)
-    minimax_reduce_rms_kernel_lamport(MiniMaxReduceRMSParams params) {
-  IndexHelper<DType> index_helper(params);
-  int token_id = index_helper.token_id;
-  int access_id_in_token = index_helper.access_id_in_token;
-  int token_stride = index_helper.token_stride;
-  int access_id = index_helper.access_id;
-  int access_stride = index_helper.access_stride;
-  int tot_access = index_helper.tot_access;
-  int tot_tokens = params.size_q / params.hidden_dim;
-  float4 clear_vec = get_neg_zero();
-
-  LamportComm<NRanks> comm(params.workspace, params.rank);
-  int clear_access = comm.clear_size / kElemsPerAccess<DType>;
-#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
-  asm volatile("griddepcontrol.wait;");
-#endif
-  for (int idx = access_id; idx < tot_access;
-       idx += access_stride, token_id += token_stride) {
-    alignas(16) DType vals[kElemsPerAccess<DType>];
-    float sum_variance = 0.F;
-    *reinterpret_cast<float4*>(vals) =
-        reinterpret_cast<float4*>(params.allreduce_in)[idx];
-#pragma unroll
-    for (int i = 0; i < kElemsPerAccess<DType>; ++i) {
-      sum_variance += static_cast<float>(vals[i]) * static_cast<float>(vals[i]);
-    }
-    blockReduceSumV2<float, 1>(&sum_variance);
-    if (is_neg_zero(sum_variance)) {
-      sum_variance = 0.F;
-    }
-    if (threadIdx.x == 0) {
-      for (int r = 0; r < NRanks; ++r) {
-        reinterpret_cast<float*>(
-            comm.data_bufs[r])[(params.rank * tot_tokens) + token_id] =
-            (sum_variance);
-      }
-    }
-
-    bool done = false;
-    float vars_all_ranks[NRanks];
-    while (!done) {
-      done = true;
-#pragma unroll
-      for (int r = 0; r < NRanks; ++r) {
-        vars_all_ranks[r] = ld_global_volatile(&reinterpret_cast<float*>(
-            comm.data_bufs[params.rank])[(r * tot_tokens) + token_id]);
-        done &= !is_neg_zero(vars_all_ranks[r]);
-      }
-    }
-    sum_variance = 0.F;
-#pragma unroll
-    for (int r = 0; r < NRanks; ++r) {
-      sum_variance += vars_all_ranks[r];
-    }
-
-    DType norm_weight[kElemsPerAccess<DType>];
-    *reinterpret_cast<typename ElemsPerAccess<DType>::vec_type*>(norm_weight) =
-        reinterpret_cast<typename ElemsPerAccess<DType>::vec_type*>(
-            params.rms_gamma)[access_id_in_token];
-
-#pragma unroll
-    for (int i = 0; i < kElemsPerAccess<DType>; ++i) {
-      vals[i] = static_cast<DType>(
-          static_cast<float>(vals[i]) *
-          rsqrtf(
-              (sum_variance / static_cast<float>(params.hidden_dim) / NRanks) +
-              params.rms_eps) *
-          static_cast<float>(norm_weight[i]));
-    }
-
-    reinterpret_cast<float4*>(params.rms_norm_out)[idx] =
-        *reinterpret_cast<float4*>(vals);
-  }
-  for (int idx = access_id; idx < clear_access; idx += access_stride) {
-    reinterpret_cast<float4*>(comm.clear_buf)[idx] = clear_vec;
-  }
-  comm.update(params.size_q * NRanks);
-#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
-  asm volatile("griddepcontrol.launch_dependents;");
-#endif
-}
-
-/**
- * Float4 variant: process 4 rows at once, allreduce variance sums as float4 for
- * better memory coalescing. sum_variance is always float; applies to all DTypes
- * (half, bf16, float). When tot_tokens % 4 != 0, the last group pads rows with
- * zeros; padded rows are not written to rms_norm_out. IsQK: when true, process
- * Q+K in one loop with doubled comm buffer; when false, single-matrix (Q only).
- */
-template <typename DType, int NRanks, int OriginQDim, int OriginKDim>
-__global__ void __launch_bounds__(1024)
-    minimax_reduce_qk_rms_kernel_lamport_float4(MiniMaxReduceRMSParams params) {
-  // Compile-time per-rank dimensions
-  constexpr int RankQDim = OriginQDim / NRanks;
-  constexpr int RankKDim = OriginKDim / NRanks;
-  // Threads needed to cover one row of Q / K with float4 accesses
-  constexpr int ThreadsPerRowQ = RankQDim / kElemsPerAccess<DType>;
-  constexpr int ThreadsPerRowK = RankKDim / kElemsPerAccess<DType>;
-  // Number of warps dedicated to Q / K
-  constexpr int NumWarpQ = (ThreadsPerRowQ + MINIMAX_REDUCE_RMS_WARP_SIZE - 1) /
-                           MINIMAX_REDUCE_RMS_WARP_SIZE;
-  constexpr int NumWarpK = (ThreadsPerRowK + MINIMAX_REDUCE_RMS_WARP_SIZE - 1) /
-                           MINIMAX_REDUCE_RMS_WARP_SIZE;
-
-  int tot_tokens = params.size_q / RankQDim;
-  int tot_groups = (tot_tokens + 3) / 4;  // ceiling; last group may be partial
-
-  // Memory strides for strided qkv tensors (elements -> float4-access units)
-  int access_stride_q = (params.stride_q > 0 ? params.stride_q : RankQDim) /
-                        kElemsPerAccess<DType>;
-  int access_stride_k = (params.stride_k > 0 ? params.stride_k : RankKDim) /
-                        kElemsPerAccess<DType>;
-  // Output strides: default to contiguous (hidden_dim / hidden_dim_k)
-  int access_stride_q_out =
-      (params.stride_q_out > 0 ? params.stride_q_out : params.hidden_dim) /
-      kElemsPerAccess<DType>;
-  int access_stride_k_out =
-      (params.stride_k_out > 0 ? params.stride_k_out : params.hidden_dim_k) /
-      kElemsPerAccess<DType>;
-
-#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
-  namespace cg = cooperative_groups;
-  cg::cluster_group cluster = cg::this_cluster();
-  cg::grid_group grid = cg::this_grid();
-  int group_id = grid.cluster_rank();
-  int access_id_in_token = cluster.thread_rank();
-  int group_stride = grid.num_clusters();
-#else
-  int group_id = blockIdx.x;
-  int access_id_in_token = threadIdx.x;
-  int group_stride = gridDim.x;
-#endif
-
-  bool is_q = (access_id_in_token < NumWarpQ * MINIMAX_REDUCE_RMS_WARP_SIZE);
-  int k_thread_idx =
-      access_id_in_token - (NumWarpQ * MINIMAX_REDUCE_RMS_WARP_SIZE);
-  bool is_valid_q = (access_id_in_token < ThreadsPerRowQ);
-  bool is_valid_k = (k_thread_idx >= 0 && k_thread_idx < ThreadsPerRowK);
-  float4 clear_vec = get_neg_zero();
-
-  // Shared memory for two-level block reduction and scale broadcast
-  __shared__ float block_reduce_sum[4][MINIMAX_REDUCE_RMS_WARP_SIZE + 1];
-  __shared__ float global_scale_q[4];
-  __shared__ float global_scale_k[4];
-
-  LamportComm<NRanks> comm(params.workspace, params.rank);
-
-  DType norm_weight[kElemsPerAccess<DType>]{};
-#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
-  asm volatile("griddepcontrol.wait;");
-#endif
-  if (is_q) {
-    if (is_valid_q) {
-      *reinterpret_cast<typename ElemsPerAccess<DType>::vec_type*>(
-          norm_weight) =
-          reinterpret_cast<typename ElemsPerAccess<DType>::vec_type const*>(
-              params.rms_gamma)[access_id_in_token];
-    }
-  } else {
-    if (is_valid_k) {
-      *reinterpret_cast<typename ElemsPerAccess<DType>::vec_type*>(
-          norm_weight) =
-          reinterpret_cast<typename ElemsPerAccess<DType>::vec_type const*>(
-              params.rms_gamma_k)[k_thread_idx];
-    }
-  }
-
-  // Main loop: process one group of 4 tokens per iteration.
-  for (int g = group_id; g < tot_groups; g += group_stride) {
-    alignas(16) DType vals[4][kElemsPerAccess<DType>]{};
-    float warp_sum_variance[4]{0.F, 0.F, 0.F, 0.F};
-
-    if (is_q) {
-#pragma unroll
-      for (int row = 0; row < 4; ++row) {
-        int token_r = g * 4 + row;
-        if (token_r >= tot_tokens || !is_valid_q) {
-          continue;
-        }
-        int idx_r = token_r * access_stride_q + access_id_in_token;
-        *reinterpret_cast<float4*>(&vals[row][0]) =
-            reinterpret_cast<float4 const*>(params.allreduce_in)[idx_r];
-#pragma unroll
-        for (int i = 0; i < kElemsPerAccess<DType>; ++i) {
-          float x = static_cast<float>(vals[row][i]);
-          warp_sum_variance[row] += x * x;
-        }
-      }
-    } else {
-#pragma unroll
-      for (int row = 0; row < 4; ++row) {
-        int token_r = g * 4 + row;
-        if (token_r >= tot_tokens || !is_valid_k) {
-          continue;
-        }
-        int idx_r = token_r * access_stride_k + k_thread_idx;
-        *reinterpret_cast<float4*>(&vals[row][0]) =
-            reinterpret_cast<float4 const*>(params.allreduce_in_k)[idx_r];
-#pragma unroll
-        for (int i = 0; i < kElemsPerAccess<DType>; ++i) {
-          float x = static_cast<float>(vals[row][i]);
-          warp_sum_variance[row] += x * x;
-        }
-      }
-    }
-
-    local_warp_reduce_sum_array<MINIMAX_REDUCE_RMS_WARP_SIZE, float, 4>(
-        warp_sum_variance);
-    // Warp lane 0 writes its warp's partial sum to shared memory
-    int lane = threadIdx.x & (MINIMAX_REDUCE_RMS_WARP_SIZE - 1);
-    if (lane == 0) {
-#pragma unroll
-      for (int t = 0; t < 4; ++t) {
-        block_reduce_sum[t][threadIdx.x / MINIMAX_REDUCE_RMS_WARP_SIZE] =
-            warp_sum_variance[t];
-      }
-    }
-    __syncthreads();
-
-    int tid = threadIdx.x;
-
-    if (tid < MINIMAX_REDUCE_RMS_WARP_SIZE) {
-      constexpr int kNumWarpQPow2 =
-          (next_pow2(NumWarpQ) > NRanks) ? next_pow2(NumWarpQ) : NRanks;
-      float local_sum[4];
-#pragma unroll
-      for (int t = 0; t < 4; ++t) {
-        local_sum[t] = (tid < NumWarpQ) ? block_reduce_sum[t][tid] : 0.F;
-      }
-      // After this, all kNumWarpQPow2 lanes (including tid 0..NRanks-1) have
-      // the total Q sum-of-squares for all 4 tokens.
-      local_warp_reduce_sum_array<kNumWarpQPow2, float, 4>(local_sum);
-
-      if (tid < NRanks) {
-#pragma unroll
-        for (int t = 0; t < 4; ++t) {
-          if (is_neg_zero(local_sum[t])) {
-            local_sum[t] = 0.F;
-          }
-        }
-        // Parallel push: thread tid writes this rank's Q sum to rank tid's buf
-        reinterpret_cast<float4*>(
-            comm.data_bufs[tid])[(params.rank * tot_groups * 2) + (2 * g)] =
-            *reinterpret_cast<float4*>(local_sum);
-
-        // Parallel pull: thread tid reads rank tid's contribution from
-        // this rank's (params.rank's) buffer
-        bool done = false;
-        float4 var_all_ranks;
-        while (!done) {
-          done = true;
-          var_all_ranks = ld_global_volatile(&reinterpret_cast<float4*>(
-              comm.data_bufs[params.rank])[(tid * tot_groups * 2) + (2 * g)]);
-          done &= !is_neg_zero(var_all_ranks);
-        }
-
-        // Warp-level allreduce: each of the NRanks threads holds one rank's
-        // partial sum; after this all NRanks threads have the global total.
-        constexpr uint32_t kQActiveMask = (1u << NRanks) - 1u;
-        local_warp_reduce_sum_array<NRanks, float, 4>(
-            reinterpret_cast<float*>(&var_all_ranks), kQActiveMask);
-
-        // Thread 0 computes rsqrt with compile-time Dim and writes to smem
-        if (tid == 0) {
-          *reinterpret_cast<float4*>(global_scale_q) =
-              rms_rsqrt<OriginQDim>(var_all_ranks, params.rms_eps);
-        }
-      }
-    } else if (tid >= MINIMAX_REDUCE_RMS_WARP_SIZE * NumWarpQ &&
-               tid < MINIMAX_REDUCE_RMS_WARP_SIZE * (NumWarpQ + 1)) {
-      // --- K leader warp ---
-      constexpr int kNumWarpKPow2 =
-          (next_pow2(NumWarpK) > NRanks) ? next_pow2(NumWarpK) : NRanks;
-      float local_sum[4];
-#pragma unroll
-      for (int t = 0; t < 4; ++t) {
-        local_sum[t] = (k_thread_idx < NumWarpK)
-                           ? block_reduce_sum[t][NumWarpQ + k_thread_idx]
-                           : 0.F;
-      }
-      local_warp_reduce_sum_array<kNumWarpKPow2, float, 4>(local_sum);
-
-      if (k_thread_idx < NRanks) {
-#pragma unroll
-        for (int t = 0; t < 4; ++t) {
-          if (is_neg_zero(local_sum[t])) {
-            local_sum[t] = 0.F;
-          }
-        }
-        reinterpret_cast<float4*>(
-            comm.data_bufs[k_thread_idx])[(params.rank * tot_groups * 2) +
-                                          (2 * g + 1)] =
-            *reinterpret_cast<float4*>(local_sum);
-
-        bool done = false;
-        float4 var_all_ranks;
-        while (!done) {
-          done = true;
-          var_all_ranks = ld_global_volatile(&reinterpret_cast<float4*>(
-              comm.data_bufs[params.rank])[(k_thread_idx * tot_groups * 2) +
-                                           (2 * g + 1)]);
-          done &= !is_neg_zero(var_all_ranks);
-        }
-
-        constexpr uint32_t kKActiveMask = (1u << NRanks) - 1u;
-        local_warp_reduce_sum_array<NRanks, float, 4>(
-            reinterpret_cast<float*>(&var_all_ranks), kKActiveMask);
-
-        if (k_thread_idx == 0) {
-          *reinterpret_cast<float4*>(global_scale_k) =
-              rms_rsqrt<OriginKDim>(var_all_ranks, params.rms_eps);
-        }
-      }
-    }
-    __syncthreads();
-
-    if (is_q) {
-#pragma unroll
-      for (int t = 0; t < 4; ++t) {
-        warp_sum_variance[t] = global_scale_q[t];
-      }
-#pragma unroll
-      for (int r = 0; r < 4; ++r) {
-#pragma unroll
-        for (int i = 0; i < kElemsPerAccess<DType>; ++i) {
-          vals[r][i] = static_cast<DType>(static_cast<float>(vals[r][i]) *
-                                          warp_sum_variance[r] *
-                                          static_cast<float>(norm_weight[i]));
-        }
-        int token_r = g * 4 + r;
-        if (token_r >= tot_tokens || !is_valid_q) {
-          continue;
-        }
-        int idx_out = token_r * access_stride_q_out + access_id_in_token;
-        reinterpret_cast<float4*>(params.rms_norm_out)[idx_out] =
-            *reinterpret_cast<float4*>(&vals[r][0]);
-      }
-    } else {
-#pragma unroll
-      for (int t = 0; t < 4; ++t) {
-        warp_sum_variance[t] = global_scale_k[t];
-      }
-#pragma unroll
-      for (int r = 0; r < 4; ++r) {
-#pragma unroll
-        for (int i = 0; i < kElemsPerAccess<DType>; ++i) {
-          vals[r][i] = static_cast<DType>(static_cast<float>(vals[r][i]) *
-                                          warp_sum_variance[r] *
-                                          static_cast<float>(norm_weight[i]));
-        }
-        int token_r = g * 4 + r;
-        if (token_r >= tot_tokens || !is_valid_k) {
-          continue;
-        }
-        int idx_out = token_r * access_stride_k_out + k_thread_idx;
-        reinterpret_cast<float4*>(params.rms_norm_out_k)[idx_out] =
-            *reinterpret_cast<float4*>(&vals[r][0]);
-      }
-    }
-  }  // end group loop
-#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900))
-  asm volatile("griddepcontrol.launch_dependents;");
-#endif
-
-  int clear_access = static_cast<int>(comm.clear_size / kElemsPerAccess<DType>);
-  int clear_stride = group_stride * blockDim.x;
-  for (int idx = group_id * blockDim.x + threadIdx.x; idx < clear_access;
-       idx += clear_stride) {
-    reinterpret_cast<float4*>(comm.clear_buf)[idx] = clear_vec;
-  }
-
-  comm.update(static_cast<int64_t>(2) * tot_groups * kElemsPerAccess<DType> *
-              NRanks);
-}
-
-int get_sm_count() {
-  static int sm_count = 0;
-  if (sm_count == 0) {
-    int device_id;
-    CUDA_CHECK(cudaGetDevice(&device_id));
-    cudaDeviceProp device_prop;
-    cudaGetDeviceProperties(&device_prop, device_id);
-    sm_count = device_prop.multiProcessorCount;
-  }
-  return sm_count;
-}
-
-inline int getSMVersion(bool queryRealSmArch = false) {
-  int device{-1};
-  CUDA_CHECK(cudaGetDevice(&device));
-  int sm_major = 0;
-  int sm_minor = 0;
-  CUDA_CHECK(cudaDeviceGetAttribute(&sm_major,
-                                    cudaDevAttrComputeCapabilityMajor, device));
-  CUDA_CHECK(cudaDeviceGetAttribute(&sm_minor,
-                                    cudaDevAttrComputeCapabilityMinor, device));
-  int sm = sm_major * 10 + sm_minor;
-  if (sm == 121 && !queryRealSmArch) {
-    return 120;
-  }
-  return sm;
-}
-
-template <typename KernelFunc>
-int get_max_active_blocks(KernelFunc kernel, int block_size,
-                          int dynamic_smem = 0) {
-  int max_active = 0;
-  CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
-      &max_active, kernel, block_size, dynamic_smem));
-  return std::max(max_active, 1);
-}
-
-template <typename DType, int NRanks>
-void minimax_reduce_rms_kernel_launcher(MiniMaxReduceRMSParams const& params) {
-  static int SM = getSMVersion();
-  int token_num = params.size_q / params.hidden_dim;
-  int sm_count = get_sm_count();
-  int cluster_size = 1;
-  int cluster_num = token_num;
-  int threads_per_token = params.hidden_dim / kElemsPerAccess<DType>;
-  int block_size = threads_per_token;
-
-  int max_blocks_per_sm = get_max_active_blocks(
-      minimax_reduce_rms_kernel_lamport<DType, NRanks>, block_size);
-  int max_grid = max_blocks_per_sm * sm_count;
-
-  int grid_size =
-      (std::min(max_grid, cluster_num * cluster_size) / cluster_size) *
-      cluster_size;
-
-  cudaLaunchConfig_t cfg;
-  cfg.gridDim = grid_size;
-  cfg.blockDim = block_size;
-  cfg.dynamicSmemBytes = 0;
-  cfg.stream = params.stream;
-
-  cudaLaunchAttribute attribute[2];
-  attribute[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
-  attribute[0].val.programmaticStreamSerializationAllowed = 1;
-  attribute[1].id = cudaLaunchAttributeClusterDimension;
-  attribute[1].val.clusterDim.x = cluster_size;
-  attribute[1].val.clusterDim.y = 1;
-  attribute[1].val.clusterDim.z = 1;
-  cfg.attrs = attribute;
-  cfg.numAttrs = SM >= 90 ? 2 : 0;
-
-  CUDA_CHECK(cudaLaunchKernelEx(
-      &cfg, minimax_reduce_rms_kernel_lamport<DType, NRanks>, params));
-}
-
-template <typename DType, int NRanks, int OriginQDim, int OriginKDim>
-void minimax_reduce_rms_kernel_launcher_float4(
-    MiniMaxReduceRMSParams const& params) {
-  TORCH_CHECK(params.size_q % params.hidden_dim == 0);
-  TORCH_CHECK(params.hidden_dim % kElemsPerAccess<DType> == 0);
-  if (params.stride_q > 0) {
-    TORCH_CHECK(params.stride_q % kElemsPerAccess<DType> == 0);
-  }
-  TORCH_CHECK(params.allreduce_in_k != nullptr,
-              "float4 QK kernel requires K input");
-  TORCH_CHECK(params.hidden_dim >= params.hidden_dim_k);
-  TORCH_CHECK(params.size_k % params.hidden_dim_k == 0);
-  TORCH_CHECK(params.hidden_dim_k % kElemsPerAccess<DType> == 0);
-  TORCH_CHECK(params.size_q / params.hidden_dim ==
-              params.size_k / params.hidden_dim_k);
-  if (params.stride_k > 0) {
-    TORCH_CHECK(params.stride_k % kElemsPerAccess<DType> == 0);
-  }
-
-  int token_num = params.size_q / params.hidden_dim;
-  int tot_groups = (token_num + 3) / 4;
-  if (tot_groups == 0) {
-    return;
-  }
-
-  static int SM = getSMVersion();
-  int sm_count = get_sm_count();
-  int cluster_size = 1;
-  int cluster_num = tot_groups;
-
-  int access_per_row_q = params.hidden_dim / kElemsPerAccess<DType>;
-  int access_per_row_k = params.hidden_dim_k / kElemsPerAccess<DType>;
-
-  // Round each section up to a warp boundary
-  auto divUp = [](int a, int b) { return (a + b - 1) / b * b; };
-  int block_size = divUp(access_per_row_q, MINIMAX_REDUCE_RMS_WARP_SIZE) +
-                   divUp(access_per_row_k, MINIMAX_REDUCE_RMS_WARP_SIZE);
-
-  auto kfn =
-      minimax_reduce_qk_rms_kernel_lamport_float4<DType, NRanks, OriginQDim,
-                                                  OriginKDim>;
-
-  int max_blocks_per_sm = get_max_active_blocks(kfn, block_size);
-  int max_grid = max_blocks_per_sm * sm_count;
-  int grid_size =
-      (std::min(max_grid, cluster_num * cluster_size) / cluster_size) *
-      cluster_size;
-
-  cudaLaunchConfig_t cfg;
-  cfg.gridDim = grid_size;
-  cfg.blockDim = block_size;
-  cfg.dynamicSmemBytes = 0;
-  cfg.stream = params.stream;
-
-  cudaLaunchAttribute attribute[2];
-  attribute[0].id = cudaLaunchAttributeProgrammaticStreamSerialization;
-  attribute[0].val.programmaticStreamSerializationAllowed = 1;
-  attribute[1].id = cudaLaunchAttributeClusterDimension;
-  attribute[1].val.clusterDim.x = cluster_size;
-  attribute[1].val.clusterDim.y = 1;
-  attribute[1].val.clusterDim.z = 1;
-  cfg.attrs = attribute;
-  cfg.numAttrs = SM >= 90 ? 2 : 0;
-
-  CUDA_CHECK(cudaLaunchKernelEx(&cfg, kfn, params));
-}
-
-template <int NRanks>
-void dispatch_dtype(MiniMaxReduceRMSParams const& params) {
-  // Use the optimized QK float4 kernel when:
-  //  - K input is present, AND
-  //  - the full (NRanks * per-rank) dimensions match the MiniMax M2 shape.
-  // Otherwise fall back to the scalar kernel.
-  bool use_float4 = (params.allreduce_in_k != nullptr) &&
-                    (params.hidden_dim * params.nranks == 6144) &&
-                    (params.hidden_dim_k * params.nranks == 1024);
-
-  if (params.dtype == at::ScalarType::Half) {
-    if (use_float4) {
-      minimax_reduce_rms_kernel_launcher_float4<half, NRanks, 6144, 1024>(
-          params);
-    } else {
-      minimax_reduce_rms_kernel_launcher<half, NRanks>(params);
-    }
-  } else if (params.dtype == at::ScalarType::BFloat16) {
-    if (use_float4) {
-      minimax_reduce_rms_kernel_launcher_float4<__nv_bfloat16, NRanks, 6144,
-                                                1024>(params);
-    } else {
-      minimax_reduce_rms_kernel_launcher<__nv_bfloat16, NRanks>(params);
-    }
-  } else if (params.dtype == at::ScalarType::Float) {
-    if (use_float4) {
-      minimax_reduce_rms_kernel_launcher_float4<float, NRanks, 6144, 1024>(
-          params);
-    } else {
-      minimax_reduce_rms_kernel_launcher<float, NRanks>(params);
-    }
-  } else {
-    TORCH_CHECK(false, "Unsupported data type for minimax_reduce_rms_op");
-  }
-}
-
-void minimax_reduce_rms_op(MiniMaxReduceRMSParams const& params) {
-  if (params.nranks == 2) {
-    dispatch_dtype<2>(params);
-  } else if (params.nranks == 4) {
-    dispatch_dtype<4>(params);
-  } else if (params.nranks == 8) {
-    dispatch_dtype<8>(params);
-  } else if (params.nranks == 16) {
-    dispatch_dtype<16>(params);
-  } else {
-    TORCH_CHECK(false, "minimax_reduce_rms_op: unsupported ranks number!");
-  }
-}
-}  // namespace tensorrt_llm
-}  // namespace vllm
-
-torch::Tensor minimax_allreduce_rms(torch::Tensor const& input,
-                                    torch::Tensor const& norm_weight,
-                                    torch::Tensor workspace, int64_t const rank,
-                                    int64_t const nranks, double const eps) {
-  auto allreduce_params = vllm::tensorrt_llm::MiniMaxReduceRMSParams();
-
-  allreduce_params.nranks = static_cast<int>(nranks);
-  allreduce_params.rank = static_cast<int>(rank);
-  allreduce_params.dtype = input.scalar_type();
-  allreduce_params.size_q = static_cast<int>(input.numel());
-  allreduce_params.hidden_dim = static_cast<int>(input.size(-1));
-  allreduce_params.stride_q = allreduce_params.hidden_dim;
-  allreduce_params.workspace =
-      reinterpret_cast<void**>(workspace.mutable_data_ptr());
-  allreduce_params.allreduce_in = input.data_ptr();
-  allreduce_params.rms_gamma = norm_weight.data_ptr();
-  allreduce_params.rms_eps = static_cast<float>(eps);
-  allreduce_params.stream = at::cuda::getCurrentCUDAStream(input.get_device());
-
-  torch::Tensor rms_norm_out = torch::empty_like(input);
-  allreduce_params.rms_norm_out = rms_norm_out.mutable_data_ptr();
-
-  vllm::tensorrt_llm::minimax_reduce_rms_op(allreduce_params);
-
-  return rms_norm_out;
-}
-
-std::tuple<torch::Tensor, torch::Tensor> minimax_allreduce_rms_qk(
-    torch::Tensor qkv, torch::Tensor const& norm_weight_q,
-    torch::Tensor const& norm_weight_k, torch::Tensor workspace,
-    int64_t const q_size, int64_t const kv_size, int64_t const rank,
-    int64_t const nranks, double const eps) {
-  TORCH_CHECK(qkv.dim() == 2, "minimax_allreduce_rms_qk: qkv must be 2D");
-  TORCH_CHECK(qkv.is_contiguous(),
-              "minimax_allreduce_rms_qk: qkv must be contiguous");
-  int64_t qkv_dim = qkv.size(-1);
-  TORCH_CHECK(qkv_dim == q_size + 2 * kv_size,
-              "minimax_allreduce_rms_qk: qkv last dim must equal "
-              "q_size + 2 * kv_size");
-  TORCH_CHECK(rank < nranks,
-              "minimax_allreduce_rms_qk: rank must be less than nranks");
-
-  int64_t num_tokens = qkv.size(0);
-  int elem_bytes = qkv.element_size();
-
-  torch::Tensor q_out = torch::empty({num_tokens, q_size}, qkv.options());
-  torch::Tensor k_out = torch::empty({num_tokens, kv_size}, qkv.options());
-
-  auto params = vllm::tensorrt_llm::MiniMaxReduceRMSParams();
-  params.nranks = static_cast<int>(nranks);
-  params.rank = static_cast<int>(rank);
-  params.dtype = qkv.scalar_type();
-  params.size_q = static_cast<int>(num_tokens * q_size);
-  params.hidden_dim = static_cast<int>(q_size);
-  params.size_k = static_cast<int>(num_tokens * kv_size);
-  params.hidden_dim_k = static_cast<int>(kv_size);
-  params.stride_q = static_cast<int>(qkv_dim);
-  params.stride_k = static_cast<int>(qkv_dim);
-  params.stride_q_out = 0;  // q_out is contiguous; kernel uses hidden_dim
-  params.stride_k_out = 0;  // k_out is contiguous; kernel uses hidden_dim_k
-  params.workspace = reinterpret_cast<void**>(workspace.mutable_data_ptr());
-
-  uint8_t* base = static_cast<uint8_t*>(qkv.data_ptr());
-  params.allreduce_in = base;
-  params.allreduce_in_k = base + q_size * elem_bytes;
-  params.rms_gamma = norm_weight_q.data_ptr();
-  params.rms_gamma_k = norm_weight_k.data_ptr();
-  params.rms_eps = static_cast<float>(eps);
-  params.stream = at::cuda::getCurrentCUDAStream(qkv.get_device());
-
-  params.rms_norm_out = q_out.mutable_data_ptr();
-  params.rms_norm_out_k = k_out.mutable_data_ptr();
-
-  vllm::tensorrt_llm::minimax_reduce_rms_op(params);
-  return {q_out, k_out};
-}

csrc/minimax_reduce_rms_kernel.h

@@ -1,79 +0,0 @@
-/*
- * Copyright (c) 2026, NVIDIA CORPORATION.  All rights reserved.
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- *     http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#pragma once
-
-#include <cuda_bf16.h>
-#include <cuda_fp16.h>
-
-#include <torch/types.h>
-
-namespace vllm {
-namespace tensorrt_llm {
-
-template <typename DType>
-struct ElemsPerAccess;
-
-template <>
-struct ElemsPerAccess<half> {
-  static constexpr int value = 8;
-  using vec_type = float4;
-};
-
-template <>
-struct ElemsPerAccess<nv_bfloat16> {
-  static constexpr int value = 8;
-  using vec_type = float4;
-};
-
-template <>
-struct ElemsPerAccess<float> {
-  static constexpr int value = 4;
-  using vec_type = float4;
-};
-
-template <typename DType>
-static constexpr int kElemsPerAccess = ElemsPerAccess<DType>::value;
-
-struct MiniMaxReduceRMSParams {
-  int nranks{};
-  int rank{};
-  at::ScalarType dtype{at::ScalarType::Undefined};
-  int size_q{};
-  int hidden_dim{};
-  int size_k{};
-  int hidden_dim_k{};
-  int stride_q{};  // row stride for q input (elements); when > hidden_dim,
-                   // q is part of a wider qkv tensor
-  int stride_k{};  // row stride for k input (elements); when > hidden_dim_k,
-                   // k is part of a wider qkv tensor
-  int stride_q_out{};  // row stride for q output (elements); 0 = contiguous
-  int stride_k_out{};  // row stride for k output (elements); 0 = contiguous
-  void** workspace{};
-  void* allreduce_in{};
-  void* rms_norm_out{};
-  void* rms_gamma{};
-  void* allreduce_in_k{};
-  void* rms_norm_out_k{};
-  void* rms_gamma_k{};
-  float rms_eps{};
-  cudaStream_t stream{};
-};
-
-void minimax_reduce_rms_op(MiniMaxReduceRMSParams const& params);
-
-}  // namespace tensorrt_llm
-}  // namespace vllm

csrc/moe/marlin_moe_wna16/kernel.h

@@ -13,7 +13,7 @@
       const int4 *__restrict__ b_bias_ptr,                            \
       const float *__restrict__ a_scales_ptr,                         \
       const int4 *__restrict__ scales_ptr,                            \
-      const float *__restrict__ global_scale_ptr,                     \
+      const uint16_t *__restrict__ global_scale_ptr,                  \
       const int4 *__restrict__ zp_ptr, const int *__restrict__ g_idx, \
       const int32_t *__restrict__ sorted_token_ids_ptr,               \
       const int32_t *__restrict__ expert_ids_ptr,                     \

csrc/moe/marlin_moe_wna16/marlin_template.h

@@ -260,7 +260,7 @@ __global__ void Marlin(
     // fp16 quantization scales. shape (k/groupsize, n)
     const int4* __restrict__ scales_ptr,
     // fp16 global scale (for nvfp4// only)
-    const float* __restrict__ global_scale_ptr,
+    const uint16_t* __restrict__ global_scale_ptr,
     // 4bit packed zero-points of shape
     // (k/groupsize, n/pack_factor)
     const int4* __restrict__ zp_ptr,
@@ -308,14 +308,7 @@ __global__ void Marlin(
   constexpr int moe_block_size = m_block_size_8 ? 8 : (16 * thread_m_blocks);
 
   #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == 750
-  static constexpr auto num_bits =
-      vllm::ScalarType::from_id(b_type_id).size_bits();
-  // Disable use_fp16_accum for NVFP4 and cases when group_size == -1 &&
-  // num_bits == 4
-  constexpr bool use_fp16_accum =
-      a_type_id == vllm::kFloat16.id() &&
-      (!(b_type_id == vllm::kFE2M1f.id() && s_type_id == vllm::kFE4M3fn.id()) &&
-       !(group_blocks == -1 && num_bits == 4));
+  constexpr bool use_fp16_accum = a_type_id == vllm::kFloat16.id();
   #else
   constexpr bool use_fp16_accum = false;
   #endif
@@ -364,7 +357,7 @@ __global__ void Marlin(
       has_zp && !is_zp_float && !std::is_same<scalar_t, nv_bfloat16>::value ||
       has_zp && !is_zp_float && !(b_type == vllm::kU8);
 
-  float global_scale_f32 = 1.0f;
+  c_scalar_t2 global_scale;
 
   constexpr bool has_act_order = group_blocks == 0;
 
@@ -514,12 +507,11 @@ __global__ void Marlin(
 
       if (mul_topk_weights) {
         idx = idx < prob_m_top_k ? idx : 0;
-        float topk_weight_tmp = topk_weights_ptr[idx];
-        if constexpr (b_type == vllm::kFE2M1f && s_type == vllm::kFE4M3fn) {
-          topk_weight_tmp *= global_scale_f32;
-        }
         c_scalar_t2 topk_weight_val =
-            Cdtype::num2num2(Cdtype::float2num(topk_weight_tmp));
+            Cdtype::num2num2(Cdtype::float2num(topk_weights_ptr[idx]));
+        if constexpr (b_type == vllm::kFE2M1f && s_type == vllm::kFE4M3fn) {
+          topk_weight_val = __hmul2(topk_weight_val, global_scale);
+        }
         sh_block_topk_weights[threadIdx.x] = topk_weight_val;
       }
     }
@@ -540,7 +532,8 @@ __global__ void Marlin(
     expert_id = expert_ids_ptr[block_id];
 
     if constexpr (b_type == vllm::kFE2M1f && s_type == vllm::kFE4M3fn) {
-      global_scale_f32 = global_scale_ptr[expert_id];
+      uint16_t val = global_scale_ptr[expert_id];
+      global_scale = Cdtype::num2num2(*reinterpret_cast<c_scalar_t*>(&val));
     }
 
     B_expert_off = expert_id * prob_n * prob_k / (pack_factor * 4);
@@ -1791,13 +1784,6 @@ __global__ void Marlin(
     // We first reorder in shared memory to guarantee the most efficient final
     // global write patterns
     auto write = [&](int idx, float c0, float c1, FragS& s, FragS& b_bias) {
-      if constexpr (b_type == vllm::kFE2M1f && s_type == vllm::kFE4M3fn) {
-        if (!mul_topk_weights) {
-          c0 *= global_scale_f32;
-          c1 *= global_scale_f32;
-        }
-      }
-
       c_scalar_t2 res =
           Cdtype::nums2num2(Cdtype::float2num(c0), Cdtype::float2num(c1));
 
@@ -1814,6 +1800,11 @@ __global__ void Marlin(
         res = __hmul2(res, tmp_scale);
       }
 
+      if constexpr (b_type == vllm::kFE2M1f && s_type == vllm::kFE4M3fn) {
+        if (!mul_topk_weights) {
+          res = __hmul2(res, global_scale);
+        }
+      }
       if (has_bias && last) {
         c_scalar_t2 tmp_bias = b_bias[0];
         if constexpr (m_block_size_8) {

csrc/moe/marlin_moe_wna16/ops.cu

@@ -382,7 +382,7 @@ void marlin_mm(const void* A, const void* B, void* C, void* C_tmp, void* b_bias,
   const int4* bias_ptr = (const int4*)b_bias;
   const float* a_s_ptr = (const float*)a_s;
   const int4* b_s_ptr = (const int4*)b_s;
-  const float* g_s_ptr = (const float*)g_s;
+  const uint16_t* g_s_ptr = (const uint16_t*)g_s;
   const int4* zp_ptr = (const int4*)zp;
   const int* g_idx_ptr = (const int*)g_idx;
   const int* perm_ptr = (const int*)perm;
@@ -759,7 +759,7 @@ torch::Tensor moe_wna16_marlin_gemm(
     TORCH_CHECK(b_type == vllm::kFE2M1f && s_type == vllm::kFE4M3fn,
                 "global_scale can only be used for nvfp4 format.");
   } else {
-    global_scale = torch::empty({0}, options_fp32);
+    global_scale = torch::empty({0}, options);
     TORCH_CHECK(!(b_type == vllm::kFE2M1f && s_type == vllm::kFE4M3fn),
                 "the global_scale parameter must be passed for nvfp4 format.");
   }
@@ -842,8 +842,8 @@ torch::Tensor moe_wna16_marlin_gemm(
 
   TORCH_CHECK(a_scales.scalar_type() == at::ScalarType::Float,
               "scalar type of a_scales must be float");
-  TORCH_CHECK(global_scale.scalar_type() == at::ScalarType::Float,
-              "scalar type of global_scale must be float");
+  TORCH_CHECK(global_scale.scalar_type() == c.scalar_type(),
+              "scalar type of global_scale must be the same with c");
   if (a_type.size_bits() == 16) {
     TORCH_CHECK(
         a.scalar_type() == c.scalar_type(),

csrc/ops.h

@@ -391,16 +391,4 @@ int64_t qr_max_size();
 #ifndef USE_ROCM
 void dsv3_fused_a_gemm(torch::Tensor& output, torch::Tensor const& mat_a,
                        torch::Tensor const& mat_b);
-#endif
-
-#ifndef USE_ROCM
-torch::Tensor minimax_allreduce_rms(torch::Tensor const& input,
-                                    torch::Tensor const& norm_weight,
-                                    torch::Tensor workspace, int64_t const rank,
-                                    int64_t const nranks, double const eps);
-std::tuple<torch::Tensor, torch::Tensor> minimax_allreduce_rms_qk(
-    torch::Tensor qkv, torch::Tensor const& norm_weight_q,
-    torch::Tensor const& norm_weight_k, torch::Tensor workspace,
-    int64_t const q_size, int64_t const kv_size, int64_t const rank,
-    int64_t const nranks, double const eps);
 #endif

csrc/quantization/activation_kernels.cu

@@ -189,7 +189,10 @@ __device__ __forceinline__ void cp_async_wait<0>() {
 }
 
 __device__ __forceinline__ float clip(float v, float mmin, float mmax) {
+#if __CUDACC_VER_MAJOR__ >= 11 && __CUDA_ARCH__ >= 800
   return fminf(mmax, fmaxf(v, mmin));
+#else
+#endif
 }
 
 __device__ __forceinline__ __nv_bfloat16 clip(__nv_bfloat16 v,

csrc/quantization/marlin/kernel.h

@@ -13,7 +13,7 @@
       const int4 *__restrict__ b_bias_ptr,                                     \
       const float *__restrict__ a_scales_ptr,                                  \
       const int4 *__restrict__ scales_ptr,                                     \
-      const float *__restrict__ global_scale_ptr,                              \
+      const uint16_t *__restrict__ global_scale_ptr,                           \
       const int4 *__restrict__ zp_ptr, const int *__restrict__ g_idx,          \
       int num_groups, int prob_m, int prob_n, int prob_k, int lda, int *locks, \
       bool has_bias, bool use_atomic_add, bool use_fp32_reduce,                \

csrc/quantization/marlin/marlin.cu

@@ -57,7 +57,7 @@ torch::Tensor marlin_gemm(
     int64_t size_k, bool is_k_full, bool use_atomic_add, bool use_fp32_reduce,
     bool is_zp_float) {
   TORCH_CHECK_NOT_IMPLEMENTED(false,
-                              "marlin_gemm(..) requires CUDA_ARCH >= 7.5");
+                              "marlin_gemm(..) requires CUDA_ARCH >= 8.0");
   return torch::empty({1, 1});
 }
 
@@ -356,7 +356,7 @@ void marlin_mm(const void* A, const void* B, void* C, void* C_tmp, void* b_bias,
   const int4* bias_ptr = (const int4*)b_bias;
   const float* a_s_ptr = (const float*)a_s;
   const int4* b_s_ptr = (const int4*)b_s;
-  const float* g_s_ptr = (const float*)g_s;
+  const uint16_t* g_s_ptr = (const uint16_t*)g_s;
 
   const int4* zp_ptr = (const int4*)zp;
   const int* g_idx_ptr = (const int*)g_idx;
@@ -751,7 +751,7 @@ torch::Tensor marlin_gemm(
     TORCH_CHECK(b_type == vllm::kFE2M1f && s_type == vllm::kFE4M3fn,
                 "global_scale can only be used for nvfp4 format.");
   } else {
-    global_scale = torch::empty({0}, options_fp32);
+    global_scale = torch::empty({0}, options);
     TORCH_CHECK(!(b_type == vllm::kFE2M1f && s_type == vllm::kFE4M3fn),
                 "the global_scale parameter must be passed for nvfp4 format.");
   }
@@ -832,8 +832,8 @@ torch::Tensor marlin_gemm(
 
   TORCH_CHECK(a_scales.scalar_type() == at::ScalarType::Float,
               "scalar type of a_scales must be float");
-  TORCH_CHECK(global_scale.scalar_type() == at::ScalarType::Float,
-              "scalar type of global_scale must be float");
+  TORCH_CHECK(global_scale.scalar_type() == c.scalar_type(),
+              "scalar type of global_scale must be the same with c");
   if (a_type.size_bits() == 16) {
     TORCH_CHECK(
         a.scalar_type() == c.scalar_type(),

csrc/quantization/marlin/marlin_template.h

@@ -251,8 +251,8 @@ __global__ void Marlin(
     const float* __restrict__ a_scales_ptr,
     // fp16 quantization scales. shape (k/groupsize, n)
     const int4* __restrict__ scales_ptr,
-    // float global scale (for nvfp4// only)
-    const float* __restrict__ global_scale_ptr,
+    // fp16 global scale (for nvfp4// only)
+    const uint16_t* __restrict__ global_scale_ptr,
     // 4bit packed zero-points of shape
     // (k/groupsize, n/pack_factor)
     const int4* __restrict__ zp_ptr,
@@ -292,13 +292,7 @@ __global__ void Marlin(
   #endif
 
   #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == 750
-  constexpr auto num_bits = vllm::ScalarType::from_id(b_type_id).size_bits();
-  // Disable use_fp16_accum for NVFP4 and cases when group_size == -1 &&
-  // num_bits == 4
-  constexpr bool use_fp16_accum =
-      a_type_id == vllm::kFloat16.id() &&
-      (!(b_type_id == vllm::kFE2M1f.id() && s_type_id == vllm::kFE4M3fn.id()) &&
-       !(group_blocks == -1 && num_bits == 4));
+  constexpr bool use_fp16_accum = a_type_id == vllm::kFloat16.id();
   #else
   constexpr bool use_fp16_accum = false;
   #endif
@@ -348,10 +342,11 @@ __global__ void Marlin(
       has_zp && !is_zp_float && !std::is_same<scalar_t, nv_bfloat16>::value ||
       has_zp && !is_zp_float && !(b_type == vllm::kU8);
 
-  float global_scale_f32 = 1.0f;
+  c_scalar_t2 global_scale;
 
   if constexpr (b_type == vllm::kFE2M1f && s_type == vllm::kFE4M3fn) {
-    global_scale_f32 = global_scale_ptr[0];
+    uint16_t val = global_scale_ptr[0];
+    global_scale = Cdtype::num2num2(*reinterpret_cast<c_scalar_t*>(&val));
   }
 
   constexpr bool has_act_order = group_blocks == 0;
@@ -1649,10 +1644,6 @@ __global__ void Marlin(
     // We first reorder in shared memory to guarantee the most efficient final
     // global write patterns
     auto write = [&](int idx, float c0, float c1, FragS& s, FragS& b_bias) {
-      if constexpr (b_type == vllm::kFE2M1f && s_type == vllm::kFE4M3fn) {
-        c0 *= global_scale_f32;
-        c1 *= global_scale_f32;
-      }
       c_scalar_t2 res =
           Cdtype::nums2num2(Cdtype::float2num(c0), Cdtype::float2num(c1));
 
@@ -1668,6 +1659,10 @@ __global__ void Marlin(
         }
         res = __hmul2(res, tmp_scale);
       }
+
+      if constexpr (b_type == vllm::kFE2M1f && s_type == vllm::kFE4M3fn) {
+        res = __hmul2(res, global_scale);
+      }
       if (has_bias && last) {
         c_scalar_t2 tmp_bias = b_bias[0];
         if constexpr (m_block_size_8) {

csrc/torch_bindings.cpp

@@ -668,29 +668,6 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
       "Tensor? b_qzeros, "
       "SymInt n, SymInt group_size, SymInt sm_count, SymInt sm_version, SymInt "
       "CUBLAS_M_THRESHOLD, bool has_zp, bool n32k16_reorder) -> Tensor");
-
-  ops.def(
-      "minimax_allreduce_rms("
-      "Tensor input,"
-      "Tensor norm_weight,"
-      "Tensor workspace,"
-      "int rank,"
-      "int nranks,"
-      "float eps) -> Tensor");
-  ops.impl("minimax_allreduce_rms", torch::kCUDA, &minimax_allreduce_rms);
-  ops.def(
-      "minimax_allreduce_rms_qk("
-      "Tensor qkv,"
-      "Tensor norm_weight_q,"
-      "Tensor norm_weight_k,"
-      "Tensor workspace,"
-      "int q_size,"
-      "int kv_size,"
-      "int rank,"
-      "int nranks,"
-      "float eps) -> (Tensor, Tensor)");
-  ops.impl("minimax_allreduce_rms_qk", torch::kCUDA, &minimax_allreduce_rms_qk);
-
   //  conditionally compiled so impl in source file
 #endif
 }

docker/Dockerfile

@@ -596,25 +596,6 @@ RUN --mount=type=cache,target=/root/.cache/uv \
         --extra-index-url https://flashinfer.ai/whl/cu$(echo $CUDA_VERSION | cut -d. -f1,2 | tr -d '.') \
     && flashinfer show-config
 
-# Pre-download FlashInfer TRTLLM BMM headers for air-gapped environments.
-# At runtime, MoE JIT compilation downloads these from edge.urm.nvidia.com
-# which fails without internet. This step caches them at build time.
-RUN python3 <<'PYEOF'
-from flashinfer.jit import env as jit_env
-from flashinfer.jit.cubin_loader import download_trtllm_headers, get_cubin
-from flashinfer.artifacts import ArtifactPath, CheckSumHash
-
-download_trtllm_headers(
-    'bmm',
-    jit_env.FLASHINFER_CUBIN_DIR / 'flashinfer' / 'trtllm' / 'batched_gemm' / 'trtllmGen_bmm_export',
-    f'{ArtifactPath.TRTLLM_GEN_BMM}/include/trtllmGen_bmm_export',
-    ArtifactPath.TRTLLM_GEN_BMM,
-    get_cubin(f'{ArtifactPath.TRTLLM_GEN_BMM}/checksums.txt', CheckSumHash.TRTLLM_GEN_BMM),
-)
-
-print('FlashInfer TRTLLM BMM headers downloaded successfully')
-PYEOF
-
 # ============================================================
 # OPENAI API SERVER DEPENDENCIES
 # Pre-install these to avoid reinstalling on every vLLM wheel rebuild
@@ -649,7 +630,7 @@ RUN --mount=type=cache,target=/root/.cache/uv \
     else \
         BITSANDBYTES_VERSION="${BITSANDBYTES_VERSION_X86}"; \
     fi; \
-    uv pip install --system accelerate modelscope \
+    uv pip install --system accelerate hf_transfer modelscope \
         "bitsandbytes>=${BITSANDBYTES_VERSION}" "timm${TIMM_VERSION}" "runai-model-streamer[s3,gcs,azure]${RUNAI_MODEL_STREAMER_VERSION}"
 
 # ============================================================
@@ -772,10 +753,9 @@ RUN --mount=type=cache,target=/root/.cache/uv \
     uv pip install --system -e tests/vllm_test_utils
 
 # enable fast downloads from hf (for testing)
-ENV HF_XET_HIGH_PERFORMANCE 1
-
-# increase timeout for hf downloads (for testing)
-ENV HF_HUB_DOWNLOAD_TIMEOUT 60
+RUN --mount=type=cache,target=/root/.cache/uv \
+    uv pip install --system hf_transfer
+ENV HF_HUB_ENABLE_HF_TRANSFER 1
 
 # Copy in the v1 package for testing (it isn't distributed yet)
 COPY vllm/v1 /usr/local/lib/python${PYTHON_VERSION}/dist-packages/vllm/v1

docker/Dockerfile.cpu

@@ -140,11 +140,9 @@ RUN \
     esac; \
     }; \
     remove_packages_not_supported_on_aarch64 && \
-    sed -i 's/^torch==.*/torch==2.11.0/g' requirements/cpu-test.in && \
+    sed -i 's/^torch==.*/torch==2.10.0/g' requirements/cpu-test.in && \
     sed -i 's/torchaudio.*/torchaudio/g' requirements/cpu-test.in && \
     sed -i 's/torchvision.*/torchvision/g' requirements/cpu-test.in && \
-    # Related issue: https://github.com/vllm-project/vllm/pull/38800#issuecomment-4228314305
-    sed -i 's/^sentence-transformers.*/sentence-transformers==5.3.0/g' requirements/cpu-test.in && \
     uv pip compile requirements/cpu-test.in -o requirements/cpu-test.txt --index-strategy unsafe-best-match --torch-backend cpu
 
 RUN --mount=type=cache,target=/root/.cache/uv \
@@ -197,12 +195,6 @@ ADD ./.buildkite/ ./.buildkite/
 RUN --mount=type=cache,target=/root/.cache/uv \
     uv pip install -e tests/vllm_test_utils
 
-# enable fast downloads from hf (for testing)
-ENV HF_XET_HIGH_PERFORMANCE 1
-
-# increase timeout for hf downloads (for testing)
-ENV HF_HUB_DOWNLOAD_TIMEOUT 60
-
 ######################### RELEASE IMAGE #########################
 FROM base AS vllm-openai
 

docker/Dockerfile.nightly_torch

@@ -269,10 +269,9 @@ RUN --mount=type=cache,target=/root/.cache/uv \
     uv pip install --system -e tests/vllm_test_utils
 
 # enable fast downloads from hf (for testing)
-ENV HF_XET_HIGH_PERFORMANCE 1
-
-# increase timeout for hf downloads (for testing)
-ENV HF_HUB_DOWNLOAD_TIMEOUT 60
+RUN --mount=type=cache,target=/root/.cache/uv \
+    uv pip install --system hf_transfer
+ENV HF_HUB_ENABLE_HF_TRANSFER 1
 
 RUN --mount=type=cache,target=/root/.cache/uv \
     uv pip install --system -r requirements/nightly_torch_test.txt

docker/Dockerfile.rocm

@@ -29,11 +29,8 @@ RUN if [ "$USE_SCCACHE" != "1" ]; then \
         rm -f "$(which sccache)" || true; \
     fi
 
-# Install UV — download first, then run, so a curl failure is not masked by the pipe
-RUN curl -LsSf --retry 3 --retry-delay 5 https://astral.sh/uv/install.sh -o /tmp/uv-install.sh \
-    && env UV_INSTALL_DIR="/usr/local/bin" sh /tmp/uv-install.sh \
-    && rm -f /tmp/uv-install.sh \
-    && uv --version
+# Install UV
+RUN curl -LsSf https://astral.sh/uv/install.sh | env UV_INSTALL_DIR="/usr/local/bin" sh
 
 # This timeout (in seconds) is necessary when installing some dependencies via uv since it's likely to time out
 # Reference: https://github.com/astral-sh/uv/pull/1694
@@ -364,10 +361,9 @@ RUN cd /vllm-workspace \
     && python3 -m pip install pytest-shard
 
 # enable fast downloads from hf (for testing)
-ENV HF_XET_HIGH_PERFORMANCE=1
-
-# increase timeout for hf downloads (for testing)
-ENV HF_HUB_DOWNLOAD_TIMEOUT 60
+RUN --mount=type=cache,target=/root/.cache/uv \
+    uv pip install --system hf_transfer
+ENV HF_HUB_ENABLE_HF_TRANSFER=1
 
 # install audio decode package `torchcodec` from source (required due to 
 # ROCm and torch version mismatch) for tests with datasets package
@@ -390,9 +386,6 @@ ENV MIOPEN_DEBUG_CONV_GEMM=0
 # will not be imported by other tests
 RUN mkdir src && mv vllm src/vllm
 
-# This is a workaround to ensure pytest exits with the correct status code in CI tests.
-RUN echo "import os\n\ndef pytest_sessionfinish(session, exitstatus):\n    os._exit(int(exitstatus))" > /vllm-workspace/conftest.py
-
 # -----------------------
 # Final vLLM image
 FROM base AS final

docker/Dockerfile.rocm_base

@@ -1,7 +1,7 @@
-ARG BASE_IMAGE=rocm/dev-ubuntu-22.04:7.2.1-complete
-ARG TRITON_BRANCH="ba5c1517"
+ARG BASE_IMAGE=rocm/dev-ubuntu-22.04:7.0-complete
+ARG TRITON_BRANCH="57c693b6"
 ARG TRITON_REPO="https://github.com/ROCm/triton.git"
-ARG PYTORCH_BRANCH="8514f051" # release/2.10 as of 3/17
+ARG PYTORCH_BRANCH="89075173"
 ARG PYTORCH_REPO="https://github.com/ROCm/pytorch.git"
 ARG PYTORCH_VISION_BRANCH="v0.24.1"
 ARG PYTORCH_VISION_REPO="https://github.com/pytorch/vision.git"
@@ -114,8 +114,6 @@ ARG TRITON_REPO
 RUN git clone ${TRITON_REPO}
 RUN cd triton \
     && git checkout ${TRITON_BRANCH} \
-    && git config --global user.email "you@example.com" && git config --global user.name "Your Name" \
-    && git cherry-pick 555d04f \
     && if [ ! -f setup.py ]; then cd python; fi \
     && python3 setup.py bdist_wheel --dist-dir=dist \
     && mkdir -p /app/install && cp dist/*.whl /app/install
@@ -144,14 +142,10 @@ ARG PYTORCH_VISION_REPO
 ARG PYTORCH_AUDIO_REPO
 ARG USE_SCCACHE
 
-RUN apt-get update && apt-get install -y pkg-config liblzma-dev
 RUN git clone ${PYTORCH_REPO} pytorch
-RUN cd pytorch && git checkout ${PYTORCH_BRANCH}
-RUN cd pytorch \
-    && pip install -r requirements.txt && git submodule update --init --recursive
-RUN cd pytorch/third_party/kineto \
-    && git remote add rocm https://github.com/ROCm/kineto && git fetch rocm && git checkout 2d73be3 
-RUN cd pytorch && python3 tools/amd_build/build_amd.py \
+RUN cd pytorch && git checkout ${PYTORCH_BRANCH} \
+    && pip install -r requirements.txt && git submodule update --init --recursive \
+    && python3 tools/amd_build/build_amd.py \
     && if [ "$USE_SCCACHE" = "1" ]; then \
            export HIP_CLANG_PATH=/opt/sccache-wrappers \
            && export CMAKE_C_COMPILER_LAUNCHER=sccache \
@@ -245,7 +239,7 @@ RUN pip install pyyaml && cd aiter \
            export HIP_CLANG_PATH=/opt/sccache-wrappers \
            && sccache --show-stats; \
        fi \
-    && PREBUILD_KERNELS=1 GPU_ARCHS=${AITER_ROCM_ARCH} python3 setup.py bdist_wheel --dist-dir=dist \
+    && GPU_ARCHS=${AITER_ROCM_ARCH} python3 setup.py bdist_wheel --dist-dir=dist \
     && if [ "$USE_SCCACHE" = "1" ]; then sccache --show-stats; fi \
     && ls /app/aiter/dist/*.whl
 RUN mkdir -p /app/install && cp /app/aiter/dist/*.whl /app/install

docs/deployment/frameworks/helm.md

@@ -17,8 +17,6 @@ Before you begin, ensure that you have the following:
 
 ## Installing the chart
 
-This guide uses the Helm chart at [examples/online_serving/chart-helm](../../../examples/online_serving/chart-helm).
-
 To install the chart with the release name `test-vllm`:
 
 ```bash

docs/design/attention_backends.md

@@ -173,9 +173,9 @@ Priority is **1 = highest** (tried first).
 | `FLASH_ATTN` | FA4* | fp16, bf16 | `auto`, `float16`, `bfloat16` | %16 | Any | ❌ | ❌ | ✅ | All | ≥10.0 |
 | `FLASH_ATTN_DIFFKV` | | fp16, bf16 | `auto` | Any | Any | ❌ | ❌ | ✅ | Decoder | Any |
 | `FLEX_ATTENTION` | | fp16, bf16, fp32 | `auto`, `float16`, `bfloat16` | Any | Any | ❌ | ✅ | ❌ | Decoder, Encoder Only | Any |
-| `ROCM_AITER_FA` | | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | 16, 32 | 64, 128, 256 | ❌ | ❌ | ❌ | Decoder | N/A |
+| `ROCM_AITER_FA` | | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | 16, 32 | 64, 128, 256 | ❌ | ❌ | ❌ | Decoder, Enc-Dec | N/A |
 | `ROCM_AITER_UNIFIED_ATTN` | | fp16, bf16 | `auto` | %16 | Any | ✅ | ✅ | ❌ | All | N/A |
-| `ROCM_ATTN` | | fp16, bf16, fp32 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | %16 | 32, 64, 80, 96, 128, 160, 192, 224, 256 | ❌ | ✅ | ❌ | Decoder, Encoder, Encoder Only | N/A |
+| `ROCM_ATTN` | | fp16, bf16, fp32 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | %16 | 32, 64, 80, 96, 128, 160, 192, 224, 256 | ❌ | ✅ | ❌ | All | N/A |
 | `TREE_ATTN` | | fp16, bf16 | `auto`, `float16`, `bfloat16` | %16 | 32, 64, 96, 128, 160, 192, 224, 256 | ❌ | ❌ | ❌ | Decoder | Any |
 | `TRITON_ATTN` | | fp16, bf16, fp32 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | %16 | Any | ✅ | ✅ | ❌ | All | Any |
 

docs/features/reasoning_outputs.md

@@ -244,12 +244,12 @@ response = client.chat.completions.create(
 
 Some models, such as [Qwen3](https://qwen.readthedocs.io/en/latest/getting_started/quickstart.html#thinking-budget), [DeepSeek](https://www.alibabacloud.com/help/en/model-studio/deep-thinking), and [Nemotron3](https://huggingface.co/nvidia/NVIDIA-Nemotron-3-Nano-30B-A3B-BF16), support a thinking budget that limits the maximum number of tokens used for reasoning.
 
-Token counting starts from `reasoning_start_str`. Once the reasoning token count reaches the configured `thinking_token_budget`, vLLM forces the model to produce `reasoning_end_str`, effectively terminating the reasoning block.
+Token counting starts from `think_start_str`. Once the reasoning token count reaches the configured `thinking_token_budget`, vLLM forces the model to produce `think_end_str`, effectively terminating the reasoning block.
 
 To use this feature:
 
 - `--reasoning-parser` enables reasoning extraction.
-- `--reasoning-config` defines the reasoning boundary tokens (e.g., `reasoning_start_str`, `reasoning_end_str`).
+- `--reasoning-config` defines the reasoning boundary tokens (e.g., `think_start_str`, `think_end_str`).
 - `thinking_token_budget` (a sampling parameter) sets the per-request reasoning token limit.
 
 If `thinking_token_budget` is not specified, no explicit reasoning limit is applied beyond normal generation constraints such as `max_tokens`.
@@ -257,20 +257,20 @@ If `thinking_token_budget` is not specified, no explicit reasoning limit is appl
 `--reasoning-config` accepts a JSON object corresponding to  
 [ReasoningConfig][vllm.config.ReasoningConfig] with the following fields:
 
-| Field                 | Type           | Description                                      |
-|-----------------------|----------------|--------------------------------------------------|
-| `reasoning_start_str` | `str \| null`  | String that marks the start of reasoning content |
-| `reasoning_end_str`   | `str \| null`  | String that marks the end of reasoning content   |
+| Field             | Type           | Description                                      |
+|-------------------|----------------|--------------------------------------------------|
+| `think_start_str` | `str \| null`  | String that marks the start of reasoning content |
+| `think_end_str`   | `str \| null`  | String that marks the end of reasoning content   |
 
 !!! note
-    `reasoning_end_str` can include a transition phrase before the reasoning end token. For example, setting `reasoning_end_str` to `"I have to give the solution based on the reasoning directly now.</think>"` instructs the model to emit that phrase when the budget is exhausted, making the reasoning termination more natural.
+    `think_end_str` can include a transition phrase before the think end token. For example, setting `think_end_str` to `"I have to give the solution based on the thinking directly now.</think>"` instructs the model to emit that phrase when the budget is exhausted, making the reasoning termination more natural.
 
 ### Online Serving
 
 ```bash
 vllm serve Qwen/Qwen3-0.6B \
     --reasoning-parser qwen3 \
-    --reasoning-config '{"reasoning_start_str": "<think>", "reasoning_end_str": "I have to give the solution based on the reasoning directly now.</think>"}'
+    --reasoning-config '{"think_start_str": "<think>", "think_end_str": "I have to give the solution based on the thinking directly now.</think>"}'
 ```
 
 Then make a request with `thinking_token_budget` to limit the reasoning tokens:
@@ -298,8 +298,8 @@ from vllm.config import ReasoningConfig
 llm = LLM(
     model="Qwen/Qwen3-0.6B",
     reasoning_config=ReasoningConfig(
-        reasoning_start_str="<think>",
-        reasoning_end_str="I have to give the solution based on the thinking directly now.</think>",
+        think_start_str="<think>",
+        think_end_str="I have to give the solution based on the thinking directly now.</think>",
     ),
 )
 

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

@@ -147,7 +147,7 @@ uv pip install vllm --extra-index-url https://wheels.vllm.ai/rocm/0.15.0/rocm700
         # Install dependencies
         pip install --upgrade numba \
             scipy \
-            huggingface-hub[cli] \
+            huggingface-hub[cli,hf_transfer] \
             setuptools_scm
         pip install -r requirements/rocm.txt
 
@@ -172,11 +172,8 @@ uv pip install vllm --extra-index-url https://wheels.vllm.ai/rocm/0.15.0/rocm700
 --8<-- [end:build-wheel-from-source]
 --8<-- [start:pre-built-images]
 
-vLLM offers official Docker images for deployment.
-The images can be used to run OpenAI compatible server and are available on Docker Hub as [vllm/vllm-openai-rocm](https://hub.docker.com/r/vllm/vllm-openai-rocm/tags).
-
-- `vllm/vllm-openai-rocm:latest` — stable release
-- `vllm/vllm-openai-rocm:nightly` — preview build from the latest development branch, use this if you want the latest features and fixes
+vLLM offers an official Docker image for deployment.
+The image can be used to run OpenAI compatible server and is available on Docker Hub as [vllm/vllm-openai-rocm](https://hub.docker.com/r/vllm/vllm-openai-rocm/tags).
 
 ```bash
 docker run --rm \
@@ -189,18 +186,30 @@ docker run --rm \
     --env "HF_TOKEN=$HF_TOKEN" \
     -p 8000:8000 \
     --ipc=host \
-    vllm/vllm-openai-rocm:<tag> \
+    vllm/vllm-openai-rocm:latest \
     --model Qwen/Qwen3-0.6B
 ```
 
-#### Use AMD's Docker Images (Deprecated)
+#### Use AMD's Docker Images
 
-!!! warning "Deprecated"
-    AMD's Docker images (`rocm/vllm` and `rocm/vllm-dev`) are deprecated in favor of the official vLLM Docker images above (`vllm/vllm-openai-rocm`). Please migrate to the official images.
-
-Prior to January 20th, 2026 when the official docker images became available on [upstream vLLM docker hub](https://hub.docker.com/v2/repositories/vllm/vllm-openai-rocm/tags/), the [AMD Infinity hub for vLLM](https://hub.docker.com/r/rocm/vllm/tags) offered a prebuilt, optimized
+Prior to January 20th, 2026 when the official docker images are available on [upstream vLLM docker hub](https://hub.docker.com/v2/repositories/vllm/vllm-openai-rocm/tags/), the [AMD Infinity hub for vLLM](https://hub.docker.com/r/rocm/vllm/tags) offers a prebuilt, optimized
 docker image designed for validating inference performance on the AMD Instinct MI300X™ accelerator.
-AMD also offered nightly prebuilt docker image from [Docker Hub](https://hub.docker.com/r/rocm/vllm-dev), which has vLLM and all its dependencies installed. The entrypoint of this docker image is `/bin/bash` (different from the vLLM's Official Docker Image).
+AMD also offers nightly prebuilt docker image from [Docker Hub](https://hub.docker.com/r/rocm/vllm-dev), which has vLLM and all its dependencies installed. The entrypoint of this docker image is `/bin/bash` (different from the vLLM's Official Docker Image).
+
+```bash
+docker pull rocm/vllm-dev:nightly # to get the latest image
+docker run -it --rm \
+--network=host \
+--group-add=video \
+--ipc=host \
+--cap-add=SYS_PTRACE \
+--security-opt seccomp=unconfined \
+--device /dev/kfd \
+--device /dev/dri \
+-v <path/to/your/models>:/app/models \
+-e HF_HOME="/app/models" \
+rocm/vllm-dev:nightly
+```
 
 !!! tip
     Please check [LLM inference performance validation on AMD Instinct MI300X](https://rocm.docs.amd.com/en/latest/how-to/performance-validation/mi300x/vllm-benchmark.html)

docs/getting_started/quickstart.md

@@ -56,12 +56,9 @@ This guide will help you quickly get started with vLLM to perform:
     !!! note
         It currently supports Python 3.12, ROCm 7.0 and `glibc >= 2.35`.
 
-    !!! note
+    !!! note    
         Note that, previously, docker images were published using AMD's docker release pipeline and were located `rocm/vllm-dev`. This is being deprecated by using vLLM's docker release pipeline.
 
-    !!! tip
-        A nightly Docker image is also available as [vllm/vllm-openai-rocm:nightly](https://hub.docker.com/r/vllm/vllm-openai-rocm/tags) for testing the latest development builds.
-
 === "Google TPU"
 
     To run vLLM on Google TPUs, you need to install the `vllm-tpu` package.

docs/mkdocs/hooks/generate_argparse.py

@@ -153,7 +153,7 @@ class MarkdownFormatter(HelpFormatter):
             heading_md = f"{self._argument_heading_prefix} {option_strings}\n\n"
             self._markdown_output.append(heading_md)
 
-            if action.choices or isinstance(action.metavar, list | tuple):
+            if action.choices or isinstance(action.metavar, (list, tuple)):
                 choices_iterable = action.choices or action.metavar
                 choices = f"`{'`, `'.join(str(c) for c in choices_iterable)}`"
                 self._markdown_output.append(f":   Possible choices: {choices}\n\n")

docs/models/pooling_models/token_classify.md

@@ -15,7 +15,7 @@ Many classification models support both (sequence) classification and token clas
 
 !!! note
 
-    Pooling multitask support is deprecated and will be removed in v0.20. When the default pooling task (classify) is not
+    Pooling multitask support is deprecated and will be removed in v0.20. When the default pooling task (classify) is not 
     what you want, you need to manually specify it via `PoolerConfig(task="token_classify")` offline or
     `--pooler-config.task token_classify` online.
 
@@ -29,12 +29,6 @@ Offline: [examples/pooling/token_classify/ner_offline.py](../../../examples/pool
 
 Online: [examples/pooling/token_classify/ner_online.py](../../../examples/pooling/token_classify/ner_online.py)
 
-### Forced Alignment
-
-Forced alignment takes audio and reference text as input and produces word-level timestamps.
-
-Offline: [examples/pooling/token_classify/forced_alignment_offline.py](../../../examples/pooling/token_classify/forced_alignment_offline.py)
-
 ### Sparse retrieval (lexical matching)
 
 The BAAI/bge-m3 model leverages token classification for sparse retrieval. For more information, see [this page](specific_models.md#baaibge-m3).
@@ -49,25 +43,12 @@ The BAAI/bge-m3 model leverages token classification for sparse retrieval. For m
 | `Qwen3ForTokenClassification`<sup>C</sup> | Qwen3-based | `bd2lcco/Qwen3-0.6B-finetuned` | | |
 | `*Model`<sup>C</sup>, `*ForCausalLM`<sup>C</sup>, etc. | Generative models | N/A | \* | \* |
 
-<sup>C</sup> Automatically converted into a classification model via `--convert classify`. ([details](./README.md#model-conversion))
+<sup>C</sup> Automatically converted into a classification model via `--convert classify`. ([details](./README.md#model-conversion))  
 \* Feature support is the same as that of the original model.
 
 If your model is not in the above list, we will try to automatically convert the model using
 [as_seq_cls_model][vllm.model_executor.models.adapters.as_seq_cls_model]. By default, the class probabilities are extracted from the softmaxed hidden state corresponding to the last token.
 
-### Multimodal Models
-
-!!! note
-    For more information about multimodal models inputs, see [this page](../supported_models.md#list-of-multimodal-language-models).
-
-| Architecture                                  | Models              | Inputs            | Example HF Models                          | [LoRA](../../features/lora.md) | [PP](../../serving/parallelism_scaling.md) |
-| --------------------------------------------- | ------------------- | ----------------- | ------------------------------------------ | ------------------------------ | ------------------------------------------ |
-| `Qwen3ASRForcedAlignerForTokenClassification` | Qwen3-ForcedAligner | T + A<sup>+</sup> | `Qwen/Qwen3-ForcedAligner-0.6B` (see note) |                                | ✅︎                                         |
-
-!!! note
-    Forced alignment usage requires `--hf-overrides '{"architectures": ["Qwen3ASRForcedAlignerForTokenClassification"]}'`.
-    Please refer to [examples/pooling/token_classify/forced_alignment_offline.py](../../../examples/pooling/token_classify/forced_alignment_offline.py).
-
 ### As Reward Models
 
 Using token classification models as reward models. For details on reward models, see [Reward Models](reward.md).

docs/usage/security.md

@@ -231,18 +231,6 @@ The most effective approach is to deploy vLLM behind a reverse proxy (such as ng
 - Blocks all other endpoints, including the unauthenticated inference and operational control endpoints
 - Implements additional authentication, rate limiting, and logging at the proxy layer
 
-## Request Parameter Resource Limits
-
-Certain API request parameters can have a large impact on resource consumption and may be abused to exhaust server resources. The `n` parameter in the `/v1/completions` and `/v1/chat/completions` endpoints controls how many independent output sequences are generated per request. A very large value causes the engine to allocate memory, CPU, and GPU time proportional to `n`, which can lead to out-of-memory conditions on the host and block the server from processing other requests.
-
-To mitigate this, vLLM enforces a configurable upper bound on the `n` parameter via the `VLLM_MAX_N_SEQUENCES` environment variable (default: **16384**). Requests exceeding this limit are rejected before reaching the engine.
-
-### Recommendations
-
-- **Public-facing deployments:** Consider setting `VLLM_MAX_N_SEQUENCES` to a value appropriate for your workload (e.g., `64` or `128`) to limit the blast radius of a single request.
-- **Reverse proxy layer:** In addition to vLLM's built-in limit, consider enforcing request body validation and rate limiting at your reverse proxy to further constrain abusive payloads.
-- **Monitoring:** Monitor per-request resource consumption to detect anomalous patterns that may indicate abuse.
-
 ## Tool Server and MCP Security
 
 vLLM supports connecting to external tool servers via the `--tool-server` argument. This enables models to call tools through the Responses API (`/v1/responses`). Tool server support works with all models — it is not limited to specific model architectures.

examples/pooling/token_classify/forced_alignment_offline.py

@@ -1,90 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-# Adapted from Qwen3-ForcedAligner inference:
-# https://github.com/QwenLM/Qwen3-ASR
-
-"""
-Offline forced alignment example using Qwen3-ForcedAligner-0.6B.
-
-Forced alignment takes audio and reference text as input and produces
-word-level timestamps. The model predicts a time bin at each <timestamp>
-token position; multiplying by ``timestamp_segment_time`` gives milliseconds.
-
-Usage::
-
-    python forced_alignment_offline.py \
-        --model Qwen/Qwen3-ForcedAligner-0.6B
-"""
-
-from argparse import Namespace
-
-import numpy as np
-
-from vllm import LLM, EngineArgs
-from vllm.utils.argparse_utils import FlexibleArgumentParser
-
-
-def parse_args():
-    parser = FlexibleArgumentParser()
-    parser = EngineArgs.add_cli_args(parser)
-    parser.set_defaults(
-        model="Qwen/Qwen3-ForcedAligner-0.6B",
-        runner="pooling",
-        enforce_eager=True,
-        hf_overrides={"architectures": ["Qwen3ASRForcedAlignerForTokenClassification"]},
-    )
-    return parser.parse_args()
-
-
-def build_prompt(words: list[str]) -> str:
-    """Build the forced alignment prompt from a word list.
-
-    Format: <|audio_start|><|audio_pad|><|audio_end|>
-            word1<timestamp><timestamp>word2<timestamp><timestamp>...
-    """
-    body = "<timestamp><timestamp>".join(words) + "<timestamp><timestamp>"
-    return f"<|audio_start|><|audio_pad|><|audio_end|>{body}"
-
-
-def main(args: Namespace):
-    llm = LLM(**vars(args))
-
-    config = llm.llm_engine.vllm_config.model_config.hf_config
-    timestamp_token_id = config.timestamp_token_id
-    timestamp_segment_time = config.timestamp_segment_time
-
-    # Example: align these words against a 5-second audio clip
-    words = ["Hello", "world"]
-    prompt = build_prompt(words)
-
-    # Use a 5-second silent audio as placeholder (replace with real audio)
-    sample_rate = 16000
-    audio = np.zeros(sample_rate * 5, dtype=np.float32)
-
-    outputs = llm.encode(
-        [{"prompt": prompt, "multi_modal_data": {"audio": audio}}],
-        pooling_task="token_classify",
-    )
-
-    for output in outputs:
-        logits = output.outputs.data  # [num_tokens, classify_num]
-        predictions = logits.argmax(dim=-1)
-        token_ids = output.prompt_token_ids
-
-        # Extract timestamps at <timestamp> positions
-        ts_predictions = [
-            pred.item() * timestamp_segment_time
-            for tid, pred in zip(token_ids, predictions)
-            if tid == timestamp_token_id
-        ]
-
-        # Pair up start/end times per word
-        for i, word in enumerate(words):
-            start_ms = ts_predictions[i * 2]
-            end_ms = ts_predictions[i * 2 + 1]
-            print(f"{word:15s} {start_ms / 1000:.3f}s - {end_ms / 1000:.3f}s")
-
-
-if __name__ == "__main__":
-    args = parse_args()
-    main(args)

examples/tool_chat_template_gemma4.jinja

@@ -1,331 +0,0 @@
-{%- macro format_parameters(properties, required) -%}
-    {%- set standard_keys = ['description', 'type', 'properties', 'required', 'nullable'] -%}
-    {%- set ns = namespace(found_first=false) -%}
-    {%- for key, value in properties | dictsort -%}
-        {%- set add_comma = false -%}
-        {%- if key not in standard_keys -%}
-            {%- if ns.found_first %},{% endif -%}
-            {%- set ns.found_first = true -%}
-            {{ key }}:{
-            {%- if value['description'] -%}
-                description:<|"|>{{ value['description'] }}<|"|>
-                {%- set add_comma = true -%}
-            {%- endif -%}
-            {%- if value['nullable'] %}
-                {%- if add_comma %},{%- else -%} {%- set add_comma = true -%} {% endif -%}
-                nullable:true
-            {%- endif -%}
-            {%- if value['type'] | upper == 'STRING' -%}
-                {%- if value['enum'] -%}
-                    {%- if add_comma %},{%- else -%} {%- set add_comma = true -%} {% endif -%}
-                    enum:{{ format_argument(value['enum']) }}
-                {%- endif -%}
-            {%- elif value['type'] | upper == 'OBJECT' -%}
-                ,properties:{
-                {%- if value['properties'] is defined and value['properties'] is mapping -%}
-                    {{- format_parameters(value['properties'], value['required'] | default([])) -}}
-                {%- elif value is mapping -%}
-                    {{- format_parameters(value, value['required'] | default([])) -}}
-                {%- endif -%}
-                }
-                {%- if value['required'] -%}
-                    ,required:[
-                    {%- for item in value['required'] | default([]) -%}
-                        <|"|>{{- item -}}<|"|>
-                        {%- if not loop.last %},{% endif -%}
-                    {%- endfor -%}
-                    ]
-                {%- endif -%}
-            {%- elif value['type'] | upper == 'ARRAY' -%}
-                {%- if value['items'] is mapping and value['items'] -%}
-                    ,items:{
-                    {%- set ns_items = namespace(found_first=false) -%}
-                    {%- for item_key, item_value in value['items'] | dictsort -%}
-                        {%- if item_value is not none -%}
-                            {%- if ns_items.found_first %},{% endif -%}
-                            {%- set ns_items.found_first = true -%}
-                            {%- if item_key == 'properties' -%}
-                                properties:{
-                                {%- if item_value is mapping -%}
-                                    {{- format_parameters(item_value, value['items']['required'] | default([])) -}}
-                                {%- endif -%}
-                                }
-                            {%- elif item_key == 'required' -%}
-                                required:[
-                                {%- for req_item in item_value -%}
-                                    <|"|>{{- req_item -}}<|"|>
-                                    {%- if not loop.last %},{% endif -%}
-                                {%- endfor -%}
-                                ]
-                            {%- elif item_key == 'type' -%}
-                                {%- if item_value is string -%}
-                                    type:{{ format_argument(item_value | upper) }}
-                                {%- else -%}
-                                    type:{{ format_argument(item_value | map('upper') | list) }}
-                                {%- endif -%}
-                            {%- else -%}
-                                {{ item_key }}:{{ format_argument(item_value) }}
-                            {%- endif -%}
-                        {%- endif -%}
-                    {%- endfor -%}
-                    }
-                {%- endif -%}
-            {%- endif -%}
-            {%- if add_comma %},{%- else -%} {%- set add_comma = true -%} {% endif -%}
-            type:<|"|>{{ value['type'] | upper }}<|"|>}
-        {%- endif -%}
-    {%- endfor -%}
-{%- endmacro -%}
-{%- macro format_function_declaration(tool_data) -%}
-    declaration:{{- tool_data['function']['name'] -}}{description:<|"|>{{- tool_data['function']['description'] -}}<|"|>
-    {%- set params = tool_data['function']['parameters'] -%}
-    {%- if params -%}
-        ,parameters:{
-        {%- if params['properties'] -%}
-            properties:{ {{- format_parameters(params['properties'], params['required']) -}} },
-        {%- endif -%}
-        {%- if params['required'] -%}
-            required:[
-            {%- for item in params['required'] -%}
-                <|"|>{{- item -}}<|"|>
-                {{- ',' if not loop.last -}}
-            {%- endfor -%}
-            ],
-        {%- endif -%}
-        {%- if params['type'] -%}
-            type:<|"|>{{- params['type'] | upper -}}<|"|>}
-        {%- endif -%}
-    {%- endif -%}
-    {%- if 'response' in tool_data['function'] -%}
-        {%- set response_declaration = tool_data['function']['response'] -%}
-        ,response:{
-        {%- if response_declaration['description'] -%}
-            description:<|"|>{{- response_declaration['description'] -}}<|"|>,
-        {%- endif -%}
-        {%- if response_declaration['type'] | upper == 'OBJECT' -%}
-            type:<|"|>{{- response_declaration['type'] | upper -}}<|"|>}
-        {%- endif -%}
-    {%- endif -%}
-    }
-{%- endmacro -%}
-{%- macro format_argument(argument, escape_keys=True) -%}
-    {%- if argument is string -%}
-        {{- '<|"|>' + argument + '<|"|>' -}}
-    {%- elif argument is boolean -%}
-        {{- 'true' if argument else 'false' -}}
-    {%- elif argument is mapping -%}
-        {{- '{' -}}
-        {%- set ns = namespace(found_first=false) -%}
-        {%- for key, value in argument | dictsort -%}
-            {%- if ns.found_first %},{% endif -%}
-            {%- set ns.found_first = true -%}
-            {%- if escape_keys -%}
-                {{- '<|"|>' + key + '<|"|>' -}}
-            {%- else -%}
-                {{- key -}}
-            {%- endif -%}
-            :{{- format_argument(value, escape_keys=escape_keys) -}}
-        {%- endfor -%}
-        {{- '}' -}}
-    {%- elif argument is sequence -%}
-        {{- '[' -}}
-        {%- for item in argument -%}
-            {{- format_argument(item, escape_keys=escape_keys) -}}
-            {%- if not loop.last %},{% endif -%}
-        {%- endfor -%}
-        {{- ']' -}}
-    {%- else -%}
-        {{- argument -}}
-    {%- endif -%}
-{%- endmacro -%}
-{%- macro strip_thinking(text) -%}
-    {%- set ns = namespace(result='') -%}
-    {%- for part in text.split('<channel|>') -%}
-        {%- if '<|channel>' in part -%}
-            {%- set ns.result = ns.result + part.split('<|channel>')[0] -%}
-        {%- else -%}
-            {%- set ns.result = ns.result + part -%}
-        {%- endif -%}
-    {%- endfor -%}
-    {{- ns.result | trim -}}
-{%- endmacro -%}
-
-{%- macro format_tool_response_block(tool_name, response) -%}
-    {{- '<|tool_response>' -}}
-    {%- if response is mapping -%}
-        {{- 'response:' + tool_name + '{' -}}
-        {%- for key, value in response | dictsort -%}
-            {{- key -}}:{{- format_argument(value, escape_keys=False) -}}
-            {%- if not loop.last %},{% endif -%}
-        {%- endfor -%}
-        {{- '}' -}}
-    {%- else -%}
-        {{- 'response:' + tool_name + '{value:' + format_argument(response, escape_keys=False) + '}' -}}
-    {%- endif -%}
-    {{- '<tool_response|>' -}}
-{%- endmacro -%}
-
-{%- set ns = namespace(prev_message_type=None) -%}
-{%- set loop_messages = messages -%}
-{{ bos_token }}
-{%- if (enable_thinking is defined and enable_thinking) or tools or messages[0]['role'] in ['system', 'developer'] -%}
-    {{- '<|turn>system\n' -}}
-
-    {%- if enable_thinking is defined and enable_thinking -%}
-        {{- '<|think|>' -}}
-        {%- set ns.prev_message_type = 'think' -%}
-    {%- endif -%}
-
-    {%- if messages[0]['role'] in ['system', 'developer'] -%}
-        {{- messages[0]['content'] | trim -}}
-        {%- set loop_messages = messages[1:] -%}
-    {%- endif -%}
-
-    {%- if tools -%}
-        {%- for tool in tools %}
-            {{- '<|tool>' -}}
-            {{- format_function_declaration(tool) | trim -}}
-            {{- '<tool|>' -}}
-        {%- endfor %}
-        {%- set ns.prev_message_type = 'tool' -%}
-    {%- endif -%}
-
-    {{- '<turn|>\n' -}}
-{%- endif %}
-
-{%- set ns_turn = namespace(last_user_idx=-1) -%}
-{%- for i in range(loop_messages | length) -%}
-    {%- if loop_messages[i]['role'] == 'user' -%}
-        {%- set ns_turn.last_user_idx = i -%}
-    {%- endif -%}
-{%- endfor -%}
-
-{%- for message in loop_messages -%}
-    {%- if message['role'] != 'tool' -%}
-    {%- set ns.prev_message_type = None -%}
-    {%- set role = 'model' if message['role'] == 'assistant' else message['role'] -%}
-    {#- OpenAI may emit multiple assistant messages in one tool loop (user → asst → tool → asst → tool).
-        Only the first of those should open <|turn>model; later ones continue the same model turn. -#}
-    {%- set prev_nt = namespace(role=None, found=false) -%}
-    {%- if loop.index0 > 0 -%}
-        {%- for j in range(loop.index0 - 1, -1, -1) -%}
-            {%- if not prev_nt.found -%}
-                {%- if loop_messages[j]['role'] != 'tool' -%}
-                    {%- set prev_nt.role = loop_messages[j]['role'] -%}
-                    {%- set prev_nt.found = true -%}
-                {%- endif -%}
-            {%- endif -%}
-        {%- endfor -%}
-    {%- endif -%}
-    {%- set continue_same_model_turn = (role == 'model' and prev_nt.role == 'assistant') -%}
-    {%- if not continue_same_model_turn -%}
-        {{- '<|turn>' + role + '\n' }}
-    {%- endif -%}
-
-    {%- if message.get('reasoning') and loop.index0 > ns_turn.last_user_idx and message.get('tool_calls') -%}
-        {{- '<|channel>thought\n' + message['reasoning'] + '\n<channel|>'}}
-    {%- endif -%}
-
-            {%- if message['tool_calls'] -%}
-                {%- for tool_call in message['tool_calls'] -%}
-                    {%- set function = tool_call['function'] -%}
-                    {{- '<|tool_call>call:' + function['name'] + '{' -}}
-                    {%- if function['arguments'] is mapping -%}
-                        {%- set ns_args = namespace(found_first=false) -%}
-                        {%- for key, value in function['arguments'] | dictsort -%}
-                            {%- if ns_args.found_first %},{% endif -%}
-                            {%- set ns_args.found_first = true -%}
-                            {{- key -}}:{{- format_argument(value, escape_keys=False) -}}
-                        {%- endfor -%}
-                    {%- elif function['arguments'] is string -%}
-                        {{- function['arguments'] -}}
-                    {%- endif -%}
-                    {{- '}<tool_call|>' -}}
-                {%- endfor -%}
-                {%- set ns.prev_message_type = 'tool_call' -%}
-            {%- endif -%}
-
-            {%- set ns_tr_out = namespace(flag=false) -%}
-            {%- if message.get('tool_responses') -%}
-                {#- Legacy: tool_responses embedded on the assistant message -#}
-                {%- for tool_response in message['tool_responses'] -%}
-                    {{- format_tool_response_block(tool_response['name'] | default('unknown'), tool_response['response']) -}}
-                    {%- set ns_tr_out.flag = true -%}
-                    {%- set ns.prev_message_type = 'tool_response' -%}
-                {%- endfor -%}
-            {%- elif message.get('tool_calls') -%}
-                {#- OpenAI Chat Completions: consecutive following messages with role "tool" (no break/continue; range scan) -#}
-                {%- set ns_tool_scan = namespace(stopped=false) -%}
-                {%- for k in range(loop.index0 + 1, loop_messages | length) -%}
-                    {%- if ns_tool_scan.stopped -%}
-                    {%- elif loop_messages[k]['role'] != 'tool' -%}
-                        {%- set ns_tool_scan.stopped = true -%}
-                    {%- else -%}
-                        {%- set follow = loop_messages[k] -%}
-                        {%- set ns_tname = namespace(name=follow.get('name') | default('unknown')) -%}
-                        {%- for tc in message['tool_calls'] -%}
-                            {%- if tc.get('id') == follow.get('tool_call_id') -%}
-                                {%- set ns_tname.name = tc['function']['name'] -%}
-                            {%- endif -%}
-                        {%- endfor -%}
-                        {%- set tool_body = follow.get('content') -%}
-                        {%- if tool_body is string -%}
-                            {{- format_tool_response_block(ns_tname.name, tool_body) -}}
-                        {%- elif tool_body is sequence and tool_body is not string -%}
-                            {%- set ns_txt = namespace(s='') -%}
-                            {%- for part in tool_body -%}
-                                {%- if part.get('type') == 'text' -%}
-                                    {%- set ns_txt.s = ns_txt.s + (part.get('text') | default('')) -%}
-                                {%- endif -%}
-                            {%- endfor -%}
-                            {{- format_tool_response_block(ns_tname.name, ns_txt.s) -}}
-                        {%- else -%}
-                            {{- format_tool_response_block(ns_tname.name, tool_body) -}}
-                        {%- endif -%}
-                        {%- set ns_tr_out.flag = true -%}
-                        {%- set ns.prev_message_type = 'tool_response' -%}
-                    {%- endif -%}
-                {%- endfor -%}
-            {%- endif -%}
-
-            {%- if message['content'] is string -%}
-                {%- if role == 'model' -%}
-                    {{- strip_thinking(message['content']) -}}
-                {%- else -%}
-                    {{- message['content'] | trim -}}
-                {%- endif -%}
-            {%- elif message['content'] is sequence -%}
-                {%- for item in message['content'] -%}
-                    {%- if item['type'] == 'text' -%}
-                        {%- if role == 'model' -%}
-                            {{- strip_thinking(item['text']) -}}
-                        {%- else -%}
-                            {{- item['text'] | trim -}}
-                        {%- endif -%}
-                    {%- elif item['type'] == 'image' -%}
-                        {{- '\n\n<|image|>\n\n' -}}
-                        {%- set ns.prev_message_type = 'image' -%}
-                    {%- elif item['type'] == 'audio' -%}
-                        {{- '<|audio|>' -}}
-                        {%- set ns.prev_message_type = 'audio' -%}
-                    {%- elif item['type'] == 'video' -%}
-                        {{- '\n\n<|video|>\n\n' -}}
-                        {%- set ns.prev_message_type = 'video' -%}
-                    {%- endif -%}
-                {%- endfor -%}
-            {%- endif -%}
-
-        {%- if not (ns_tr_out.flag and not message.get('content')) -%}
-            {{- '<turn|>\n' -}}
-        {%- endif -%}
-    {%- endif -%}
-{%- endfor -%}
-
-{%- if add_generation_prompt -%}
-    {%- if ns.prev_message_type != 'tool_response' -%}
-        {{- '<|turn>model\n' -}}
-    {%- endif -%}
-    {%- if not enable_thinking | default(false) -%}
-        {{- '<|channel>thought\n<channel|>' -}}
-    {%- endif -%}
-{%- endif -%}

requirements/common.txt

@@ -7,7 +7,7 @@ requests >= 2.26.0
 tqdm
 blake3
 py-cpuinfo
-transformers >= 4.56.0, != 5.0.*, != 5.1.*, != 5.2.*, != 5.3.*, != 5.4.*, != 5.5.0
+transformers >= 4.56.0, < 5
 tokenizers >= 0.21.1  # Required for fast incremental detokenization.
 protobuf >= 5.29.6, !=6.30.*, !=6.31.*, !=6.32.*, !=6.33.0.*, !=6.33.1.*, !=6.33.2.*, !=6.33.3.*, !=6.33.4.* # Required by LlamaTokenizer, gRPC. CVE-2026-0994
 fastapi[standard] >= 0.115.0 # Required by FastAPI's form models in the OpenAI API server's audio transcriptions endpoint.
@@ -37,7 +37,7 @@ pyyaml
 six>=1.16.0; python_version > '3.11' # transitive dependency of pandas that needs to be the latest version for python 3.12
 setuptools>=77.0.3,<81.0.0; python_version > '3.11' # Setuptools is used by triton, we need to ensure a modern version is installed for 3.12+ so that it does not try to import distutils, which was removed in 3.12
 einops # Required for Qwen2-VL.
-compressed-tensors == 0.15.0.1 # required for compressed-tensors
+compressed-tensors == 0.14.0.1 # required for compressed-tensors
 depyf==0.20.0 # required for profiling and debugging with compilation config
 cloudpickle # allows pickling lambda functions in model_executor/models/registry.py
 watchfiles # required for http server to monitor the updates of TLS files

requirements/kv_connectors.txt

@@ -1,5 +1,3 @@
 lmcache >= 0.3.9
 nixl >= 0.7.1, < 0.10.0 # Required for disaggregated prefill
-nixl-cu12 >= 0.7.1, < 0.10.0
-nixl-cu13 >= 0.7.1, < 0.10.0
 mooncake-transfer-engine >= 0.3.8

requirements/nightly_torch_test.txt

@@ -29,8 +29,8 @@ opencv-python-headless >= 4.13.0 # required for video test
 datamodel_code_generator # required for minicpm3 test
 lm-eval[api]>=0.4.11 # required for model evaluation test
 mteb[bm25s]>=2, <3 # required for mteb test
-transformers==5.5.3
-tokenizers==0.22.2
+transformers==4.57.5
+tokenizers==0.22.0
 schemathesis>=3.39.15 # Required for openai schema test.
 # quantization
 bitsandbytes>=0.49.2

requirements/rocm-test.in

@@ -36,8 +36,8 @@ opencv-python-headless>=4.13.0 # required for video test
 datamodel_code_generator # required for minicpm3 test
 lm-eval[api]>=0.4.11 # required for model evaluation test
 mteb[bm25s]>=2, <3 # required for mteb test
-transformers==5.5.3
-tokenizers==0.22.2
+transformers==4.57.5
+tokenizers==0.22.0
 schemathesis>=3.39.15 # Required for openai schema test
 # quantization
 bitsandbytes==0.49.2
@@ -80,3 +80,4 @@ plotly # required for perf comparison html report
 rapidfuzz
 torchgeo==0.7.0
 multiprocess==0.70.16
+huggingface-hub==0.36.2

requirements/rocm-test.txt

@@ -232,6 +232,7 @@ filelock==3.25.2
     #   python-discovery
     #   ray
     #   torch
+    #   transformers
     #   virtualenv
 fiona==1.10.1
     # via torchgeo
@@ -317,7 +318,7 @@ h5py==3.16.0
     # via terratorch
 harfile==0.4.0
     # via schemathesis
-hf-xet==1.4.3
+hf-xet==1.4.2
     # via huggingface-hub
 hiredis==3.3.1
     # via tensorizer
@@ -331,11 +332,11 @@ httpx==0.27.2
     # via
     #   -r requirements/rocm-test.in
     #   diffusers
-    #   huggingface-hub
     #   perceptron
     #   schemathesis
-huggingface-hub==1.10.2
+huggingface-hub==0.36.2
     # via
+    #   -r requirements/rocm-test.in
     #   accelerate
     #   datasets
     #   diffusers
@@ -969,6 +970,7 @@ requests==2.32.5
     #   google-api-core
     #   google-cloud-storage
     #   gpt-oss
+    #   huggingface-hub
     #   lightly
     #   lm-eval
     #   mistral-common
@@ -981,6 +983,7 @@ requests==2.32.5
     #   starlette-testclient
     #   tacoreader
     #   tiktoken
+    #   transformers
     #   wandb
 resampy==0.4.3
     # via -r requirements/rocm-test.in
@@ -1188,7 +1191,7 @@ timm==1.0.17
     #   segmentation-models-pytorch
     #   terratorch
     #   torchgeo
-tokenizers==0.22.2
+tokenizers==0.22.0
     # via
     #   -c requirements/common.txt
     #   -r requirements/rocm-test.in
@@ -1227,7 +1230,7 @@ tqdm==4.67.3
     #   tacoreader
     #   terratorch
     #   transformers
-transformers==5.5.3
+transformers==4.57.5
     # via
     #   -c requirements/common.txt
     #   -r requirements/rocm-test.in
@@ -1249,9 +1252,7 @@ typepy==1.3.4
 typer==0.24.1
     # via
     #   fastsafetensors
-    #   huggingface-hub
     #   perceptron
-    #   transformers
 typeshed-client==2.9.0
     # via jsonargparse
 typing-extensions==4.15.0

requirements/test.in

@@ -18,7 +18,7 @@ httpx
 librosa # required for audio tests
 vector_quantize_pytorch # required for minicpmo_26 test
 vocos # required for minicpmo_26 test
-peft>=0.18.1 # required for phi-4-mm test
+peft>=0.15.0 # required for phi-4-mm test
 pqdm
 ray[cgraph,default]>=2.48.0 # Ray Compiled Graph, required by pipeline parallelism tests
 resampy # required for audio tests
@@ -39,8 +39,8 @@ opencv-python-headless >= 4.13.0 # required for video test
 datamodel_code_generator # required for minicpm3 test
 lm-eval[api]>=0.4.11 # required for model evaluation test
 mteb[bm25s]>=2, <3 # required for mteb test
-transformers==5.5.3
-tokenizers==0.22.2
+transformers==4.57.5
+tokenizers==0.22.0
 schemathesis>=3.39.15 # Required for openai schema test.
 # quantization
 bitsandbytes==0.49.2

requirements/test.txt

@@ -4,7 +4,7 @@ absl-py==2.1.0
     # via
     #   rouge-score
     #   tensorboard
-accelerate==1.13.0
+accelerate==1.0.1
     # via peft
 aenum==3.1.16
     # via lightly
@@ -240,6 +240,7 @@ filelock==3.16.1
     #   huggingface-hub
     #   ray
     #   torch
+    #   transformers
     #   virtualenv
 fiona==1.10.1
     # via torchgeo
@@ -322,7 +323,7 @@ h5py==3.13.0
     # via terratorch
 harfile==0.3.0
     # via schemathesis
-hf-xet==1.4.3
+hf-xet==1.1.7
     # via huggingface-hub
 hiredis==3.0.0
     # via tensorizer
@@ -336,10 +337,9 @@ httpx==0.27.2
     # via
     #   -r requirements/test.in
     #   diffusers
-    #   huggingface-hub
     #   perceptron
     #   schemathesis
-huggingface-hub==1.10.2
+huggingface-hub==0.36.2
     # via
     #   accelerate
     #   datasets
@@ -740,7 +740,7 @@ pathvalidate==3.2.1
     # via pytablewriter
 patsy==1.0.1
     # via statsmodels
-peft==0.18.1
+peft==0.16.0
     # via -r requirements/test.in
 perceptron==0.1.4
     # via -r requirements/test.in
@@ -963,7 +963,7 @@ referencing==0.35.1
     # via
     #   jsonschema
     #   jsonschema-specifications
-regex==2026.2.28
+regex==2024.9.11
     # via
     #   diffusers
     #   nltk
@@ -982,6 +982,7 @@ requests==2.32.3
     #   google-api-core
     #   google-cloud-storage
     #   gpt-oss
+    #   huggingface-hub
     #   lightly
     #   lm-eval
     #   mistral-common
@@ -994,6 +995,7 @@ requests==2.32.3
     #   starlette-testclient
     #   tacoreader
     #   tiktoken
+    #   transformers
     #   wandb
 resampy==0.4.3
     # via -r requirements/test.in
@@ -1191,7 +1193,7 @@ timm==1.0.17
     #   segmentation-models-pytorch
     #   terratorch
     #   torchgeo
-tokenizers==0.22.2
+tokenizers==0.22.0
     # via
     #   -r requirements/test.in
     #   transformers
@@ -1267,7 +1269,7 @@ tqdm==4.67.3
     #   tacoreader
     #   terratorch
     #   transformers
-transformers==5.5.3
+transformers==4.57.5
     # via
     #   -r requirements/test.in
     #   genai-perf
@@ -1288,9 +1290,7 @@ typepy==1.3.2
 typer==0.15.2
     # via
     #   fastsafetensors
-    #   huggingface-hub
     #   perceptron
-    #   transformers
 types-python-dateutil==2.9.0.20241206
     # via arrow
 typeshed-client==2.8.2

requirements/test/xpu.txt

@@ -1,736 +0,0 @@
-# This file was autogenerated by uv via the following command:
-#    uv pip compile requirements/test/xpu.in -c requirements/xpu.txt -o requirements/test/xpu.txt --index-strategy unsafe-best-match --torch-backend xpu --python-platform x86_64-manylinux_2_39 --python-version 3.12
-absl-py==2.4.0
-    # via
-    #   -r requirements/test/xpu.in
-    #   rouge-score
-accelerate==1.13.0
-    # via -r requirements/test/xpu.in
-aiohappyeyeballs==2.6.1
-    # via aiohttp
-aiohttp==3.13.4
-    # via
-    #   -c requirements/common.txt
-    #   fsspec
-    #   gpt-oss
-    #   lm-eval
-aiosignal==1.4.0
-    # via aiohttp
-albumentations==1.4.6
-    # via -r requirements/test/xpu.in
-annotated-doc==0.0.4
-    # via
-    #   fastapi
-    #   typer
-annotated-types==0.7.0
-    # via pydantic
-anyio==4.13.0
-    # via
-    #   httpx
-    #   starlette
-arctic-inference==0.1.1
-    # via -r requirements/test/xpu.in
-attrs==26.1.0
-    # via
-    #   aiohttp
-    #   jsonlines
-    #   jsonschema
-    #   referencing
-audioread==3.0.1
-    # via
-    #   -r requirements/test/xpu.in
-    #   librosa
-blobfile==3.0.0
-    # via -r requirements/test/xpu.in
-bm25s==0.2.13
-    # via
-    #   -r requirements/test/xpu.in
-    #   mteb
-bounded-pool-executor==0.0.3
-    # via pqdm
-certifi==2026.2.25
-    # via
-    #   httpcore
-    #   httpx
-    #   requests
-cffi==2.0.0
-    # via soundfile
-chardet==5.2.0
-    # via mbstrdecoder
-charset-normalizer==3.4.6
-    # via requests
-chz==0.4.0
-    # via gpt-oss
-click==8.3.1
-    # via
-    #   jiwer
-    #   nltk
-    #   schemathesis
-    #   typer
-    #   uvicorn
-colorama==0.4.6
-    # via sacrebleu
-coverage==7.13.5
-    # via pytest-cov
-dataproperty==1.1.0
-    # via
-    #   pytablewriter
-    #   tabledata
-datasets==4.8.4
-    # via
-    #   evaluate
-    #   lm-eval
-    #   mteb
-decorator==5.2.1
-    # via librosa
-dill==0.4.1
-    # via
-    #   datasets
-    #   evaluate
-    #   lm-eval
-    #   multiprocess
-docker==7.1.0
-    # via gpt-oss
-docopt==0.6.2
-    # via num2words
-dpcpp-cpp-rt==2025.3.1
-    # via
-    #   onemkl-sycl-blas
-    #   onemkl-sycl-dft
-    #   onemkl-sycl-lapack
-    #   onemkl-sycl-rng
-    #   onemkl-sycl-sparse
-    #   torch
-evaluate==0.4.6
-    # via lm-eval
-fastapi==0.135.2
-    # via
-    #   -c requirements/common.txt
-    #   gpt-oss
-filelock==3.25.2
-    # via
-    #   -c requirements/common.txt
-    #   blobfile
-    #   datasets
-    #   huggingface-hub
-    #   modelscope
-    #   torch
-frozenlist==1.8.0
-    # via
-    #   aiohttp
-    #   aiosignal
-fsspec==2026.2.0
-    # via
-    #   datasets
-    #   evaluate
-    #   huggingface-hub
-    #   torch
-gpt-oss==0.0.8
-    # via -r requirements/test/xpu.in
-graphql-core==3.2.8
-    # via hypothesis-graphql
-h11==0.16.0
-    # via
-    #   httpcore
-    #   uvicorn
-harfile==0.4.0
-    # via schemathesis
-hf-xet==1.4.3
-    # via huggingface-hub
-html2text==2025.4.15
-    # via gpt-oss
-httpcore==1.0.9
-    # via httpx
-httpx==0.28.1
-    # via
-    #   datasets
-    #   huggingface-hub
-    #   schemathesis
-huggingface-hub==1.10.2
-    # via
-    #   accelerate
-    #   datasets
-    #   evaluate
-    #   sentence-transformers
-    #   timm
-    #   tokenizers
-    #   transformers
-hypothesis==6.151.10
-    # via
-    #   hypothesis-graphql
-    #   hypothesis-jsonschema
-    #   schemathesis
-hypothesis-graphql==0.12.0
-    # via schemathesis
-hypothesis-jsonschema==0.23.1
-    # via schemathesis
-idna==3.11
-    # via
-    #   anyio
-    #   httpx
-    #   requests
-    #   yarl
-imageio==2.37.3
-    # via scikit-image
-impi-rt==2021.17.0
-    # via
-    #   oneccl
-    #   torch
-iniconfig==2.3.0
-    # via pytest
-intel-cmplr-lib-rt==2025.3.1
-    # via
-    #   intel-sycl-rt
-    #   torch
-intel-cmplr-lib-ur==2025.3.1
-    # via
-    #   intel-openmp
-    #   intel-sycl-rt
-    #   torch
-intel-cmplr-lic-rt==2025.3.1
-    # via
-    #   intel-opencl-rt
-    #   intel-sycl-rt
-    #   torch
-intel-opencl-rt==2025.3.1
-    # via
-    #   dpcpp-cpp-rt
-    #   onemkl-sycl-blas
-    #   onemkl-sycl-dft
-    #   onemkl-sycl-lapack
-    #   onemkl-sycl-rng
-    #   onemkl-sycl-sparse
-    #   torch
-intel-openmp==2025.3.1
-    # via
-    #   dpcpp-cpp-rt
-    #   mkl
-    #   torch
-intel-pti==0.15.0
-    # via torch
-intel-sycl-rt==2025.3.1
-    # via
-    #   dpcpp-cpp-rt
-    #   oneccl
-    #   torch
-jinja2==3.1.6
-    # via
-    #   -c requirements/xpu.txt
-    #   lm-eval
-    #   torch
-jiwer==4.0.0
-    # via -r requirements/test/xpu.in
-joblib==1.5.3
-    # via
-    #   librosa
-    #   nltk
-    #   scikit-learn
-jsonlines==4.0.0
-    # via lm-eval
-jsonschema==4.26.0
-    # via
-    #   hypothesis-jsonschema
-    #   mistral-common
-    #   schemathesis
-jsonschema-rs==0.45.0
-    # via schemathesis
-jsonschema-specifications==2025.9.1
-    # via jsonschema
-junit-xml==1.9
-    # via schemathesis
-lazy-loader==0.5
-    # via
-    #   librosa
-    #   scikit-image
-librosa==0.10.2.post1
-    # via -r requirements/test/xpu.in
-llvmlite==0.44.0
-    # via numba
-lm-eval==0.4.11
-    # via -r requirements/test/xpu.in
-lxml==6.0.2
-    # via
-    #   blobfile
-    #   gpt-oss
-    #   sacrebleu
-markdown-it-py==4.0.0
-    # via rich
-markupsafe==3.0.3
-    # via
-    #   jinja2
-    #   werkzeug
-mbstrdecoder==1.1.4
-    # via
-    #   dataproperty
-    #   pytablewriter
-    #   typepy
-mdurl==0.1.2
-    # via markdown-it-py
-mistral-common==1.11.0
-    # via
-    #   -c requirements/common.txt
-    #   -r requirements/test/xpu.in
-mkl==2025.3.0
-    # via
-    #   onemkl-sycl-blas
-    #   onemkl-sycl-dft
-    #   onemkl-sycl-lapack
-    #   onemkl-sycl-rng
-    #   onemkl-sycl-sparse
-    #   torch
-modelscope==1.35.3
-    # via -r requirements/test/xpu.in
-more-itertools==10.8.0
-    # via lm-eval
-mpmath==1.3.0
-    # via sympy
-msgpack==1.1.2
-    # via librosa
-mteb==2.12.7
-    # via -r requirements/test/xpu.in
-multidict==6.7.1
-    # via
-    #   aiohttp
-    #   yarl
-multiprocess==0.70.19
-    # via
-    #   datasets
-    #   evaluate
-networkx==3.6.1
-    # via
-    #   scikit-image
-    #   torch
-nltk==3.9.4
-    # via rouge-score
-num2words==0.5.14
-    # via -r requirements/test/xpu.in
-numba==0.61.2
-    # via
-    #   -c requirements/xpu.txt
-    #   librosa
-numpy==2.2.6
-    # via
-    #   accelerate
-    #   albumentations
-    #   bm25s
-    #   datasets
-    #   evaluate
-    #   imageio
-    #   librosa
-    #   lm-eval
-    #   mistral-common
-    #   mteb
-    #   numba
-    #   opencv-python-headless
-    #   pandas
-    #   pytrec-eval-terrier
-    #   rouge-score
-    #   sacrebleu
-    #   scikit-image
-    #   scikit-learn
-    #   scipy
-    #   sentence-transformers
-    #   soundfile
-    #   soxr
-    #   tifffile
-    #   torchvision
-    #   transformers
-oneccl==2021.17.1
-    # via
-    #   oneccl-devel
-    #   torch
-oneccl-devel==2021.17.1
-    # via torch
-onemkl-license==2025.3.0
-    # via
-    #   mkl
-    #   torch
-onemkl-sycl-blas==2025.3.0
-    # via
-    #   onemkl-sycl-lapack
-    #   onemkl-sycl-sparse
-    #   torch
-onemkl-sycl-dft==2025.3.0
-    # via torch
-onemkl-sycl-lapack==2025.3.0
-    # via torch
-onemkl-sycl-rng==2025.3.0
-    # via torch
-onemkl-sycl-sparse==2025.3.0
-    # via torch
-openai-harmony==0.0.8
-    # via
-    #   -c requirements/common.txt
-    #   gpt-oss
-opencv-python-headless==4.13.0.92
-    # via
-    #   -c requirements/common.txt
-    #   albumentations
-    #   mistral-common
-packaging==26.0
-    # via
-    #   -c requirements/xpu.txt
-    #   accelerate
-    #   datasets
-    #   evaluate
-    #   huggingface-hub
-    #   lazy-loader
-    #   modelscope
-    #   pooch
-    #   pytest
-    #   pytest-rerunfailures
-    #   scikit-image
-    #   transformers
-    #   typepy
-pandas==3.0.1
-    # via
-    #   datasets
-    #   evaluate
-pathvalidate==3.3.1
-    # via pytablewriter
-pillow==12.1.1
-    # via
-    #   imageio
-    #   mistral-common
-    #   scikit-image
-    #   torchvision
-platformdirs==4.9.4
-    # via pooch
-pluggy==1.6.0
-    # via
-    #   pytest
-    #   pytest-cov
-polars==1.39.3
-    # via mteb
-polars-runtime-32==1.39.3
-    # via polars
-pooch==1.8.2
-    # via
-    #   -r requirements/test/xpu.in
-    #   librosa
-portalocker==3.2.0
-    # via sacrebleu
-pqdm==0.2.0
-    # via -r requirements/test/xpu.in
-propcache==0.4.1
-    # via
-    #   aiohttp
-    #   yarl
-psutil==7.2.2
-    # via accelerate
-py==1.11.0
-    # via pytest-forked
-pyarrow==23.0.1
-    # via datasets
-pycountry==26.2.16
-    # via pydantic-extra-types
-pycparser==3.0
-    # via cffi
-pycryptodomex==3.23.0
-    # via blobfile
-pydantic==2.12.5
-    # via
-    #   -c requirements/common.txt
-    #   albumentations
-    #   fastapi
-    #   gpt-oss
-    #   mistral-common
-    #   mteb
-    #   openai-harmony
-    #   pydantic-extra-types
-pydantic-core==2.41.5
-    # via pydantic
-pydantic-extra-types==2.11.1
-    # via mistral-common
-pyelftools==0.32
-    # via triton-xpu
-pygments==2.20.0
-    # via
-    #   pytest
-    #   rich
-pyrate-limiter==4.1.0
-    # via schemathesis
-pystemmer==3.0.0
-    # via
-    #   -r requirements/test/xpu.in
-    #   mteb
-pytablewriter==1.2.1
-    # via lm-eval
-pytest==9.0.2
-    # via
-    #   -r requirements/test/xpu.in
-    #   pytest-asyncio
-    #   pytest-cov
-    #   pytest-forked
-    #   pytest-rerunfailures
-    #   pytest-shard
-    #   pytest-timeout
-    #   schemathesis
-pytest-asyncio==1.3.0
-    # via -r requirements/test/xpu.in
-pytest-cov==6.3.0
-    # via -r requirements/test/xpu.in
-pytest-forked==1.6.0
-    # via -r requirements/test/xpu.in
-pytest-rerunfailures==14.0
-    # via -r requirements/test/xpu.in
-pytest-shard==0.1.2
-    # via -r requirements/test/xpu.in
-pytest-timeout==2.3.1
-    # via -r requirements/test/xpu.in
-python-dateutil==2.9.0.post0
-    # via
-    #   pandas
-    #   typepy
-pytrec-eval-terrier==0.5.10
-    # via mteb
-pytz==2026.1.post1
-    # via typepy
-pyyaml==6.0.3
-    # via
-    #   accelerate
-    #   albumentations
-    #   datasets
-    #   huggingface-hub
-    #   schemathesis
-    #   timm
-    #   transformers
-rapidfuzz==3.12.1
-    # via
-    #   -r requirements/test/xpu.in
-    #   jiwer
-referencing==0.37.0
-    # via
-    #   jsonschema
-    #   jsonschema-specifications
-regex==2026.3.32
-    # via
-    #   nltk
-    #   sacrebleu
-    #   tiktoken
-    #   transformers
-requests==2.33.1
-    # via
-    #   -c requirements/common.txt
-    #   datasets
-    #   docker
-    #   evaluate
-    #   gpt-oss
-    #   lm-eval
-    #   mistral-common
-    #   modelscope
-    #   mteb
-    #   pooch
-    #   schemathesis
-    #   starlette-testclient
-    #   tiktoken
-rich==14.3.3
-    # via
-    #   mteb
-    #   schemathesis
-    #   typer
-rouge-score==0.1.2
-    # via lm-eval
-rpds-py==0.30.0
-    # via
-    #   jsonschema
-    #   referencing
-sacrebleu==2.6.0
-    # via lm-eval
-safetensors==0.7.0
-    # via
-    #   accelerate
-    #   timm
-    #   transformers
-schemathesis==4.14.2
-    # via -r requirements/test/xpu.in
-scikit-image==0.26.0
-    # via albumentations
-scikit-learn==1.8.0
-    # via
-    #   albumentations
-    #   librosa
-    #   lm-eval
-    #   mteb
-    #   sentence-transformers
-scipy==1.17.1
-    # via
-    #   albumentations
-    #   bm25s
-    #   librosa
-    #   mteb
-    #   pytrec-eval-terrier
-    #   scikit-image
-    #   scikit-learn
-    #   sentence-transformers
-sentence-transformers==5.3.0
-    # via mteb
-setuptools==80.10.2
-    # via
-    #   -c requirements/common.txt
-    #   -c requirements/xpu.txt
-    #   modelscope
-    #   pytablewriter
-    #   torch
-shellingham==1.5.4
-    # via typer
-six==1.17.0
-    # via
-    #   -c requirements/common.txt
-    #   junit-xml
-    #   python-dateutil
-    #   rouge-score
-sortedcontainers==2.4.0
-    # via hypothesis
-soundfile==0.13.1
-    # via
-    #   -r requirements/test/xpu.in
-    #   librosa
-    #   mistral-common
-soxr==0.5.0.post1
-    # via
-    #   -r requirements/test/xpu.in
-    #   librosa
-    #   mistral-common
-sqlitedict==2.1.0
-    # via lm-eval
-starlette==1.0.0
-    # via
-    #   fastapi
-    #   starlette-testclient
-starlette-testclient==0.4.1
-    # via schemathesis
-structlog==25.5.0
-    # via gpt-oss
-sympy==1.14.0
-    # via torch
-tabledata==1.3.4
-    # via pytablewriter
-tabulate==0.10.0
-    # via sacrebleu
-tbb==2022.3.0
-    # via
-    #   intel-opencl-rt
-    #   mkl
-    #   torch
-tblib==3.1.0
-    # via -r requirements/test/xpu.in
-tcmlib==1.4.1
-    # via
-    #   tbb
-    #   torch
-    #   umf
-tcolorpy==0.1.7
-    # via pytablewriter
-tenacity==9.1.4
-    # via
-    #   gpt-oss
-    #   lm-eval
-    #   schemathesis
-termcolor==3.3.0
-    # via gpt-oss
-threadpoolctl==3.6.0
-    # via scikit-learn
-tifffile==2026.3.3
-    # via scikit-image
-tiktoken==0.12.0
-    # via
-    #   -c requirements/common.txt
-    #   gpt-oss
-    #   lm-eval
-    #   mistral-common
-timm==1.0.17
-    # via -r requirements/test/xpu.in
-tokenizers==0.22.2
-    # via
-    #   -c requirements/common.txt
-    #   transformers
-torch==2.10.0+xpu
-    # via
-    #   -c requirements/xpu.txt
-    #   accelerate
-    #   mteb
-    #   sentence-transformers
-    #   timm
-    #   torchvision
-torchvision==0.25.0+xpu
-    # via timm
-tqdm==4.67.3
-    # via
-    #   datasets
-    #   evaluate
-    #   huggingface-hub
-    #   lm-eval
-    #   modelscope
-    #   mteb
-    #   nltk
-    #   pqdm
-    #   sentence-transformers
-    #   transformers
-transformers==5.5.3
-    # via
-    #   -c requirements/common.txt
-    #   sentence-transformers
-triton-xpu==3.6.0
-    # via torch
-typepy==1.3.4
-    # via
-    #   dataproperty
-    #   pytablewriter
-    #   tabledata
-typer==0.24.1
-    # via
-    #   huggingface-hub
-    #   transformers
-typing-extensions==4.15.0
-    # via
-    #   -c requirements/common.txt
-    #   aiosignal
-    #   albumentations
-    #   anyio
-    #   chz
-    #   fastapi
-    #   huggingface-hub
-    #   librosa
-    #   lm-eval
-    #   mistral-common
-    #   mteb
-    #   pqdm
-    #   pydantic
-    #   pydantic-core
-    #   pydantic-extra-types
-    #   pytest-asyncio
-    #   referencing
-    #   schemathesis
-    #   sentence-transformers
-    #   starlette
-    #   torch
-    #   typing-inspection
-typing-inspection==0.4.2
-    # via
-    #   fastapi
-    #   pydantic
-umf==1.0.2
-    # via
-    #   intel-cmplr-lib-ur
-    #   torch
-urllib3==2.6.3
-    # via
-    #   blobfile
-    #   docker
-    #   modelscope
-    #   requests
-uvicorn==0.42.0
-    # via gpt-oss
-werkzeug==3.1.7
-    # via schemathesis
-word2number==1.1
-    # via lm-eval
-xxhash==3.6.0
-    # via
-    #   datasets
-    #   evaluate
-yarl==1.23.0
-    # via aiohttp
-zstandard==0.25.0
-    # via lm-eval

requirements/xpu-test.in

@@ -9,8 +9,6 @@ pytest-shard
 # --- Core Tools & Bindings ---
 absl-py
 arctic-inference
-lm_eval[api]
-modelscope
 
 # --- Audio Processing ---
 librosa

tests/conftest.py

@@ -409,15 +409,6 @@ class HfRunner:
             model_name,
             trust_remote_code=trust_remote_code,
         )
-        # HF runner should use the HF config so that it's consistent with the HF model
-        if self.config.__module__.startswith("vllm.transformers_utils.configs"):
-            from transformers.models.auto.configuration_auto import CONFIG_MAPPING
-
-            del CONFIG_MAPPING._extra_content[self.config.model_type]
-            self.config = AutoConfig.from_pretrained(
-                model_name,
-                trust_remote_code=trust_remote_code,
-            )
         self.device = self.get_default_device()
         self.dtype = dtype = _get_and_verify_dtype(
             self.model_name,

tests/distributed/test_eplb_utils.py

@@ -1,154 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-from unittest.mock import MagicMock
-
-import pytest
-import torch
-
-from vllm.distributed.eplb.eplb_state import (
-    _commit_eplb_maps,
-    _commit_eplb_maps_for_layer,
-)
-
-
-def _make_model_state(
-    phy2log: torch.Tensor,
-    log2phy: torch.Tensor,
-    logcnt: torch.Tensor,
-) -> MagicMock:
-    """Build a minimal EplbModelState mock with only the three map tensors."""
-    state = MagicMock()
-    state.physical_to_logical_map = phy2log
-    state.logical_to_physical_map = log2phy
-    state.logical_replica_count = logcnt
-    return state
-
-
-def test_commit_eplb_maps_shape_change():
-    """
-    The normal path copies the physical_to_logical map in-place. When the number of
-    physical experts changes, the old map should be replaced entirely.
-    """
-    num_layers, num_logical, num_physical = 2, 4, 6
-    max_replicas = 3
-
-    # Build current state tensors
-    model_state = _make_model_state(
-        phy2log=torch.zeros(num_layers, num_physical, dtype=torch.long),
-        log2phy=torch.full(
-            (num_layers, num_logical, max_replicas), -1, dtype=torch.long
-        ),
-        logcnt=torch.zeros(num_layers, num_logical, dtype=torch.long),
-    )
-
-    # The new map has two more physical experts. These new physical experts will
-    # automatically map to the first two logical experts
-    new_phy2log_larger = (
-        (torch.arange(num_physical + 2, dtype=torch.long) % num_logical)
-        .unsqueeze(0)
-        .expand(num_layers, -1)
-    )
-    _commit_eplb_maps(model_state, new_phy2log_larger)
-
-    # Check that the number of physical experts has been updated and that the values
-    # match
-    assert model_state.physical_to_logical_map.shape[1] == num_physical + 2
-    assert torch.equal(model_state.physical_to_logical_map, new_phy2log_larger)
-
-
-def test_commit_eplb_maps_for_layer_logical_padding():
-    """
-    Test that logical_to_physical_map is padded with -1 to fill the
-    pre-allocated slots when the new map has fewer replicas than the max.
-    """
-    num_layers, num_logical, num_physical = 2, 4, 6
-    max_replicas = 3
-
-    model_state = _make_model_state(
-        phy2log=torch.zeros(num_layers, num_physical, dtype=torch.long),
-        log2phy=torch.full(
-            (num_layers, num_logical, max_replicas), -1, dtype=torch.long
-        ),
-        logcnt=torch.zeros(num_layers, num_logical, dtype=torch.long),
-    )
-
-    new_phy2log = (
-        (torch.arange(num_physical, dtype=torch.long) % num_logical)
-        .unsqueeze(0)
-        .expand(num_layers, -1)
-        .contiguous()
-    )
-    layer = 0
-    _commit_eplb_maps_for_layer(model_state, new_phy2log, layer)
-
-    assert torch.all(model_state.logical_to_physical_map[layer, :, 2] == -1)
-
-
-def test_commit_eplb_maps_for_layer_shape_assert():
-    """Test that a mismatched number of physical experts triggers an assertion error."""
-    num_layers, num_logical, num_physical = 2, 4, 6
-
-    model_state = _make_model_state(
-        phy2log=torch.zeros(num_layers, num_physical, dtype=torch.long),
-        log2phy=torch.full((num_layers, num_logical, 2), -1, dtype=torch.long),
-        logcnt=torch.zeros(num_layers, num_logical, dtype=torch.long),
-    )
-    bad_phy2log = torch.zeros(num_layers, num_physical + 1, dtype=torch.long)
-    with pytest.raises(AssertionError):
-        _commit_eplb_maps_for_layer(model_state, bad_phy2log, layer=0)
-
-
-def test_commit_eplb_maps():
-    """Test that all values are copied correctly into model_state."""
-    num_layers, num_logical, num_physical, max_replicas = 2, 3, 4, 2
-
-    model_state = _make_model_state(
-        phy2log=torch.zeros(num_layers, num_physical, dtype=torch.long),
-        log2phy=torch.full(
-            (num_layers, num_logical, max_replicas), -1, dtype=torch.long
-        ),
-        logcnt=torch.zeros(num_layers, num_logical, dtype=torch.long),
-    )
-
-    new_phy2log = torch.tensor([[0, 1, 2, 0], [1, 2, 0, 1]], dtype=torch.long)
-    new_log2phy = torch.tensor(
-        [[[0, 3], [1, -1], [2, -1]], [[2, -1], [0, 3], [1, -1]]], dtype=torch.long
-    )
-    new_logcnt = torch.tensor([[2, 1, 1], [1, 2, 1]], dtype=torch.long)
-
-    _commit_eplb_maps(model_state, new_phy2log)
-
-    assert torch.equal(model_state.physical_to_logical_map, new_phy2log)
-    assert torch.equal(model_state.logical_to_physical_map, new_log2phy)
-    assert torch.equal(model_state.logical_replica_count, new_logcnt)
-
-
-def test_commit_eplb_maps_for_layer():
-    """Test that only the target layer is updated"""
-    num_layers, num_logical, max_replicas = 2, 3, 2
-
-    original_phy2log = torch.tensor([[9, 9, 9, 9], [8, 8, 8, 8]], dtype=torch.long)
-    model_state = _make_model_state(
-        phy2log=original_phy2log.clone(),
-        log2phy=torch.full(
-            (num_layers, num_logical, max_replicas), -1, dtype=torch.long
-        ),
-        logcnt=torch.zeros(num_layers, num_logical, dtype=torch.long),
-    )
-
-    new_phy2log = torch.tensor([[0, 1, 2, 0], [1, 2, 0, 1]], dtype=torch.long)
-    new_log2phy = torch.tensor(
-        [[[0, 3], [1, -1], [2, -1]], [[2, -1], [0, 3], [1, -1]]], dtype=torch.long
-    )
-    new_logcnt = torch.tensor([[2, 1, 1], [1, 2, 1]], dtype=torch.long)
-
-    _commit_eplb_maps_for_layer(model_state, new_phy2log, layer=0)
-
-    # Layer 0 updated
-    assert torch.equal(model_state.physical_to_logical_map[0], new_phy2log[0])
-    assert torch.equal(model_state.logical_to_physical_map[0], new_log2phy[0])
-    assert torch.equal(model_state.logical_replica_count[0], new_logcnt[0])
-
-    # Layer 1 untouched
-    assert torch.equal(model_state.physical_to_logical_map[1], original_phy2log[1])

tests/entrypoints/openai/chat_completion/test_audio_in_video.py

@@ -64,12 +64,11 @@ async def test_online_audio_in_video(
     ]
 
     # multi-turn to test mm processor cache as well
-    for turn in range(2):
+    for _ in range(2):
         chat_completion = await client.chat.completions.create(
             model=MODEL_NAME,
             messages=messages,
-            max_tokens=8,
-            temperature=0.0,
+            max_tokens=16,
             extra_body={
                 "mm_processor_kwargs": {
                     "use_audio_in_video": True,
@@ -79,12 +78,6 @@ async def test_online_audio_in_video(
 
         assert len(chat_completion.choices) == 1
         choice = chat_completion.choices[0]
-        print(
-            f"[DEBUG][single-video] turn={turn} "
-            f"finish_reason={choice.finish_reason!r} "
-            f"content={choice.message.content!r} "
-            f"usage={chat_completion.usage}"
-        )
         assert choice.finish_reason == "length"
 
 
@@ -118,12 +111,11 @@ async def test_online_audio_in_video_multi_videos(
     ]
 
     # multi-turn to test mm processor cache as well
-    for turn in range(2):
+    for _ in range(2):
         chat_completion = await client.chat.completions.create(
             model=MODEL_NAME,
             messages=messages,
-            max_tokens=8,
-            temperature=0.0,
+            max_tokens=16,
             extra_body={
                 "mm_processor_kwargs": {
                     "use_audio_in_video": True,
@@ -133,12 +125,6 @@ async def test_online_audio_in_video_multi_videos(
 
         assert len(chat_completion.choices) == 1
         choice = chat_completion.choices[0]
-        print(
-            f"[DEBUG][multi-video] turn={turn} "
-            f"finish_reason={choice.finish_reason!r} "
-            f"content={choice.message.content!r} "
-            f"usage={chat_completion.usage}"
-        )
         assert choice.finish_reason == "length"
 
 

tests/entrypoints/openai/chat_completion/test_chat.py

@@ -1020,114 +1020,3 @@ def test_chat_completion_request_n_parameter_various_values():
         assert sampling_params.n == n_value, (
             f"Expected n={n_value}, got n={sampling_params.n}"
         )
-
-
-def test_chat_completion_request_n_parameter_exceeds_default_limit(
-    monkeypatch: pytest.MonkeyPatch,
-):
-    """Test that n values exceeding the default limit are rejected."""
-    import vllm.envs as envs
-
-    monkeypatch.delenv("VLLM_MAX_N_SEQUENCES", raising=False)
-    if hasattr(envs.__getattr__, "cache_clear"):
-        envs.__getattr__.cache_clear()
-
-    max_n = envs.VLLM_MAX_N_SEQUENCES
-    request = ChatCompletionRequest(
-        model="test-model",
-        messages=[{"role": "user", "content": "Test"}],
-        n=max_n + 1,
-        max_tokens=10,
-    )
-
-    with pytest.raises(ValueError, match="n must be at most"):
-        request.to_sampling_params(
-            max_tokens=10,
-            default_sampling_params={},
-        )
-
-
-def test_chat_completion_request_n_parameter_at_limit(
-    monkeypatch: pytest.MonkeyPatch,
-):
-    """Test that n at exactly the limit is accepted."""
-    import vllm.envs as envs
-
-    monkeypatch.delenv("VLLM_MAX_N_SEQUENCES", raising=False)
-    if hasattr(envs.__getattr__, "cache_clear"):
-        envs.__getattr__.cache_clear()
-
-    max_n = envs.VLLM_MAX_N_SEQUENCES
-    request = ChatCompletionRequest(
-        model="test-model",
-        messages=[{"role": "user", "content": "Test"}],
-        n=max_n,
-        max_tokens=10,
-    )
-
-    sampling_params = request.to_sampling_params(
-        max_tokens=10,
-        default_sampling_params={},
-    )
-    assert sampling_params.n == max_n
-
-
-def test_chat_completion_request_n_parameter_custom_limit(
-    monkeypatch: pytest.MonkeyPatch,
-):
-    """Test that VLLM_MAX_N_SEQUENCES env var overrides the default limit."""
-    import vllm.envs as envs
-
-    monkeypatch.setenv("VLLM_MAX_N_SEQUENCES", "128")
-    if hasattr(envs.__getattr__, "cache_clear"):
-        envs.__getattr__.cache_clear()
-
-    request = ChatCompletionRequest(
-        model="test-model",
-        messages=[{"role": "user", "content": "Test"}],
-        n=128,
-        max_tokens=10,
-    )
-
-    sampling_params = request.to_sampling_params(
-        max_tokens=10,
-        default_sampling_params={},
-    )
-    assert sampling_params.n == 128
-
-    request_over = ChatCompletionRequest(
-        model="test-model",
-        messages=[{"role": "user", "content": "Test"}],
-        n=129,
-        max_tokens=10,
-    )
-
-    with pytest.raises(ValueError, match="n must be at most 128"):
-        request_over.to_sampling_params(
-            max_tokens=10,
-            default_sampling_params={},
-        )
-
-
-def test_chat_completion_request_n_parameter_massive_value(
-    monkeypatch: pytest.MonkeyPatch,
-):
-    """Test that astronomically large n values are rejected (CVE fix)."""
-    import vllm.envs as envs
-
-    monkeypatch.delenv("VLLM_MAX_N_SEQUENCES", raising=False)
-    if hasattr(envs.__getattr__, "cache_clear"):
-        envs.__getattr__.cache_clear()
-
-    request = ChatCompletionRequest(
-        model="test-model",
-        messages=[{"role": "user", "content": "Test"}],
-        n=100_000_000,
-        max_tokens=1,
-    )
-
-    with pytest.raises(ValueError, match="n must be at most"):
-        request.to_sampling_params(
-            max_tokens=1,
-            default_sampling_params={},
-        )

tests/entrypoints/openai/chat_completion/test_chat_error.py

@@ -55,7 +55,6 @@ class MockModelConfig:
     skip_tokenizer_init = False
     is_encoder_decoder: bool = False
     is_multimodal_model: bool = False
-    renderer_num_workers: int = 1
 
     def get_diff_sampling_param(self):
         return self.diff_sampling_param or {}

tests/entrypoints/openai/chat_completion/test_serving_chat.py

@@ -536,7 +536,6 @@ class MockModelConfig:
     skip_tokenizer_init: bool = False
     is_encoder_decoder: bool = False
     is_multimodal_model: bool = False
-    renderer_num_workers: int = 1
 
     def get_diff_sampling_param(self):
         return self.diff_sampling_param or {}

tests/entrypoints/openai/completion/test_completion_error.py

@@ -54,7 +54,6 @@ class MockModelConfig:
     skip_tokenizer_init = False
     is_encoder_decoder: bool = False
     is_multimodal_model: bool = False
-    renderer_num_workers: int = 1
 
     def get_diff_sampling_param(self):
         return self.diff_sampling_param or {}

tests/entrypoints/openai/completion/test_lora_resolvers.py

@@ -54,7 +54,6 @@ class MockModelConfig:
     skip_tokenizer_init: bool = False
     is_encoder_decoder: bool = False
     is_multimodal_model: bool = False
-    renderer_num_workers: int = 1
 
     def get_diff_sampling_param(self):
         return self.diff_sampling_param or {}

tests/entrypoints/openai/speech_to_text/test_transcription_validation_whisper.py

@@ -14,62 +14,13 @@ import pytest_asyncio
 import soundfile as sf
 
 from tests.utils import RemoteOpenAIServer
-from vllm.platforms import current_platform
 
 MODEL_NAME = "openai/whisper-large-v3-turbo"
 
-# Disable prefix caching on ROCm to reduce non-determinism in
-# streaming-vs-non-streaming comparisons.
-_ROCM_ARGS = ["--no-enable-prefix-caching"] if current_platform.is_rocm() else []
 
-
-def _get_attention_backend_params() -> list[str | None]:
-    """Return attention backends to parametrize the server fixture with.
-
-    On ROCm, we test multiple backends explicitly:
-      - None: default auto-selection (ROCM_ATTN for decoder self-attention,
-              falls back to ROCM_AITER_UNIFIED_ATTN or TRITON_ATTN for
-              cross-attention since ROCM_ATTN doesn't support ENCODER_DECODER)
-      - TRITON_ATTN: always available on ROCm
-      - ROCM_AITER_UNIFIED_ATTN: only on gfx942/gfx950
-
-    On non-ROCm platforms, we just run with the default backend.
-    """
-    try:
-        from vllm.platforms import current_platform
-
-        if current_platform.is_rocm():
-            backends: list[str | None] = [None, "TRITON_ATTN"]
-            from vllm.platforms.rocm import _ON_MI3XX
-
-            if _ON_MI3XX:
-                backends.append("ROCM_AITER_UNIFIED_ATTN")
-            return backends
-    except Exception:
-        pass
-    return [None]
-
-
-# Aiter backends need VLLM_ROCM_USE_AITER=1 (and MHA=1 for ROCM_AITER_FA)
-# to be enabled in the server subprocess.
-_AITER_ENV = {
-    "VLLM_ROCM_USE_AITER": "1",
-    "VLLM_ROCM_USE_AITER_MHA": "1",
-}
-
-_ATTN_BACKENDS = _get_attention_backend_params()
-_ATTN_IDS = [b or "default" for b in _ATTN_BACKENDS]
-
-
-@pytest.fixture(scope="module", params=_ATTN_BACKENDS, ids=_ATTN_IDS)
-def server(request):
-    args = [*_ROCM_ARGS]
-    env_dict = None
-    if request.param is not None:
-        args += ["--attention-backend", request.param]
-        if "AITER" in request.param:
-            env_dict = _AITER_ENV
-    with RemoteOpenAIServer(MODEL_NAME, args, env_dict=env_dict) as remote_server:
+@pytest.fixture(scope="module")
+def server():
+    with RemoteOpenAIServer(MODEL_NAME, []) as remote_server:
         yield remote_server
 
 

tests/entrypoints/pooling/embed/test_cohere_openai_parity.py

@@ -57,25 +57,16 @@ def _openai_embed(
     return [item["embedding"] for item in resp.json()["data"]]
 
 
-def _cosine_sim(a: list[float], b: list[float]) -> float:
-    va, vb = np.array(a), np.array(b)
-    return float(np.dot(va, vb) / (np.linalg.norm(va) * np.linalg.norm(vb)))
-
-
 def test_single_text_parity(server: RemoteOpenAIServer):
-    """A single text should produce equivalent embeddings via both APIs."""
+    """A single text should produce identical embeddings via both APIs."""
     texts = ["the quick brown fox jumps over the lazy dog"]
     v2 = _cohere_embed(server, texts)
     v1 = _openai_embed(server, texts)
-    # Full-suite BF16 runs can introduce tiny numerical drift even when both
-    # endpoints are functionally equivalent, so compare semantic equivalence
-    # instead of exact elementwise equality.
-    cos = _cosine_sim(v2[0], v1[0])
-    assert cos > 0.9999, f"single-text parity failed, cosine={cos}"
+    np.testing.assert_allclose(v2[0], v1[0], rtol=1e-5)
 
 
 def test_batch_parity(server: RemoteOpenAIServer):
-    """A batch of texts should produce equivalent embeddings via both APIs,
+    """A batch of texts should produce identical embeddings via both APIs,
     in the same order."""
     texts = [
         "machine learning",
@@ -85,18 +76,8 @@ def test_batch_parity(server: RemoteOpenAIServer):
     v2 = _cohere_embed(server, texts)
     v1 = _openai_embed(server, texts)
     assert len(v2) == len(v1) == 3
-
-    similarities = np.array(
-        [[_cosine_sim(v2_emb, v1_emb) for v1_emb in v1] for v2_emb in v2]
-    )
     for i in range(3):
-        assert int(np.argmax(similarities[i])) == i, (
-            f"batch parity order mismatch at index {i}: "
-            f"similarities={similarities[i].tolist()}"
-        )
-        assert similarities[i, i] > 0.9999, (
-            f"batch parity failed at index {i}, cosine={similarities[i, i]}"
-        )
+        np.testing.assert_allclose(v2[i], v1[i], rtol=1e-5, err_msg=f"index {i}")
 
 
 def test_token_count_parity(server: RemoteOpenAIServer):

tests/entrypoints/pooling/embed/test_io_processor.py

@@ -6,11 +6,8 @@ import pytest
 
 from vllm.entrypoints.pooling.embed.io_processor import EmbedIOProcessor
 from vllm.entrypoints.pooling.embed.protocol import (
-    CohereEmbedContent,
-    CohereEmbedInput,
     CohereEmbedRequest,
 )
-from vllm.entrypoints.pooling.typing import PoolingServeContext
 
 
 class TestResolveTruncation:
@@ -209,116 +206,3 @@ class TestValidateInputType:
         handler = self._make_handler({"a": "", "b": ""})
         with pytest.raises(ValueError, match="Supported values: a, b"):
             handler._validate_input_type("z")
-
-
-class TestPreProcessCohereOnline:
-    """Unit tests for EmbedIOProcessor._pre_process_cohere_online."""
-
-    @staticmethod
-    def _make_context(**request_kwargs) -> PoolingServeContext[CohereEmbedRequest]:
-        return PoolingServeContext(
-            request=CohereEmbedRequest(model="test", **request_kwargs),
-            model_name="test",
-            request_id="embd-test",
-        )
-
-    @staticmethod
-    def _make_handler():
-        handler = object.__new__(EmbedIOProcessor)
-        handler._validate_input_type = lambda _input_type: None
-        return handler
-
-    def test_text_only_without_task_prefix_uses_completion_path(self):
-        handler = self._make_handler()
-        ctx = self._make_context(texts=["hello"])
-        calls: list[tuple[str, object]] = []
-
-        def preprocess_completion(request, prompt_input, prompt_embeds):
-            calls.append(("completion", prompt_input))
-            return ["completion"]
-
-        handler._get_task_instruction_prefix = lambda _input_type: None
-        handler._has_chat_template = lambda: False
-        handler._preprocess_completion_online = preprocess_completion
-        handler._batch_render_chat = lambda *_args, **_kwargs: (
-            pytest.fail("text-only request should not require chat rendering")
-        )
-
-        handler._pre_process_cohere_online(ctx)
-
-        assert ctx.engine_inputs == ["completion"]
-        assert calls == [("completion", ["hello"])]
-
-    def test_text_only_falls_back_to_prefixed_completion_without_template(self):
-        handler = self._make_handler()
-        ctx = self._make_context(texts=["hello"], input_type="query")
-        calls: list[tuple[str, object]] = []
-
-        def preprocess_completion(request, prompt_input, prompt_embeds):
-            calls.append(("completion", prompt_input))
-            return ["fallback"]
-
-        handler._get_task_instruction_prefix = lambda _input_type: "query: "
-        handler._has_chat_template = lambda: False
-        handler._batch_render_chat = lambda *_args, **_kwargs: (
-            pytest.fail("chat rendering should be skipped without a template")
-        )
-        handler._preprocess_completion_online = preprocess_completion
-
-        handler._pre_process_cohere_online(ctx)
-
-        assert ctx.engine_inputs == ["fallback"]
-        assert calls == [("completion", ["query: hello"])]
-
-    def test_text_only_with_template_uses_chat_path(self):
-        handler = self._make_handler()
-        ctx = self._make_context(texts=["hello"], input_type="query")
-        calls: list[tuple[str, object]] = []
-
-        def batch_render_chat(
-            request,
-            all_messages,
-            truncate_prompt_tokens,
-            truncation_side,
-        ):
-            calls.append(
-                (
-                    "chat",
-                    {
-                        "request": request,
-                        "all_messages": all_messages,
-                        "truncate_prompt_tokens": truncate_prompt_tokens,
-                        "truncation_side": truncation_side,
-                    },
-                )
-            )
-            return ["chat"]
-
-        handler._get_task_instruction_prefix = lambda _input_type: "query: "
-        handler._has_chat_template = lambda: True
-        handler._batch_render_chat = batch_render_chat
-        handler._preprocess_completion_online = lambda *_args, **_kwargs: (
-            pytest.fail("completion path should be skipped when a template exists")
-        )
-
-        handler._pre_process_cohere_online(ctx)
-
-        assert ctx.engine_inputs == ["chat"]
-        assert calls == [
-            (
-                "chat",
-                {
-                    "request": ctx.request,
-                    "all_messages": [
-                        handler._mixed_input_to_messages(
-                            CohereEmbedInput(
-                                content=[CohereEmbedContent(type="text", text="hello")]
-                            ),
-                            task_prefix="query: ",
-                        )
-                    ],
-                    "truncate_prompt_tokens": -1,
-                    "truncation_side": None,
-                },
-            )
-        ]

tests/entrypoints/pooling/scoring/test_bi_encoder_online.py

@@ -26,6 +26,25 @@ TEXTS_2 = [
 ]
 
 
+def _assert_score_output_matches_hf(hf_model, server: RemoteOpenAIServer,
+                                    payload: dict, text_pairs: list[list[str]]):
+    score_response = requests.post(
+        server.url_for("score"),
+        json={"model": MODEL_NAME, **payload},
+    )
+    score_response.raise_for_status()
+    score = ScoreResponse.model_validate(score_response.json())
+
+    assert score.id is not None
+    assert score.data is not None
+
+    vllm_outputs = [d.score for d in score.data]
+    hf_outputs = hf_model.predict(text_pairs).tolist()
+    assert len(vllm_outputs) == len(hf_outputs)
+    for hf_output, vllm_output in zip(hf_outputs, vllm_outputs):
+        assert hf_output == pytest.approx(vllm_output, rel=0.01)
+
+
 @pytest.fixture(scope="module")
 def server():
     args = ["--enforce-eager", "--max-model-len", "100", "--dtype", DTYPE]
@@ -44,207 +63,78 @@ def hf_model():
 
 
 @pytest.mark.asyncio
-async def test_score_api_queries_str_1_documents_str_1(
-    hf_model, server: RemoteOpenAIServer
-):
-    score_response = requests.post(
-        server.url_for("score"),
-        json={
-            "model": MODEL_NAME,
-            "queries": TEXTS_1[0],
-            "documents": TEXTS_2[0],
-        },
-    )
-    score_response.raise_for_status()
-    score = ScoreResponse.model_validate(score_response.json())
-
-    assert score.id is not None
-    assert score.data is not None
-    assert len(score.data) == 1
-
-    vllm_outputs = [d.score for d in score.data]
-    hf_outputs = hf_model.predict([[TEXTS_1[0], TEXTS_2[0]]]).tolist()
-
-    for i in range(len(vllm_outputs)):
-        assert hf_outputs[i] == pytest.approx(vllm_outputs[i], rel=0.01)
-
-
-@pytest.mark.asyncio
-async def test_score_api_queries_str_1_documents_str_n(
-    hf_model, server: RemoteOpenAIServer
-):
-    text_pairs = [
-        [TEXTS_1[0], TEXTS_2[0]],
-        [TEXTS_1[0], TEXTS_2[1]],
-    ]
-
-    score_response = requests.post(
-        server.url_for("score"),
-        json={
-            "model": MODEL_NAME,
-            "queries": TEXTS_1[0],
-            "documents": TEXTS_2,
-        },
-    )
-    score_response.raise_for_status()
-    score = ScoreResponse.model_validate(score_response.json())
-
-    assert score.id is not None
-    assert score.data is not None
-    assert len(score.data) == 2
-
-    vllm_outputs = [d.score for d in score.data]
-    hf_outputs = hf_model.predict(text_pairs).tolist()
-
-    for i in range(len(vllm_outputs)):
-        assert hf_outputs[i] == pytest.approx(vllm_outputs[i], rel=0.01)
-
-
-@pytest.mark.asyncio
-async def test_score_api_queries_str_n_documents_str_n(
-    hf_model, server: RemoteOpenAIServer
-):
-    text_pairs = [
-        [TEXTS_1[0], TEXTS_2[0]],
-        [TEXTS_1[1], TEXTS_2[1]],
-    ]
-
-    score_response = requests.post(
-        server.url_for("score"),
-        json={
-            "model": MODEL_NAME,
-            "queries": TEXTS_1,
-            "documents": TEXTS_2,
-        },
-    )
-    score_response.raise_for_status()
-    score = ScoreResponse.model_validate(score_response.json())
-
-    assert score.id is not None
-    assert score.data is not None
-    assert len(score.data) == 2
-
-    vllm_outputs = [d.score for d in score.data]
-    hf_outputs = hf_model.predict(text_pairs).tolist()
-
-    for i in range(len(vllm_outputs)):
-        assert hf_outputs[i] == pytest.approx(vllm_outputs[i], rel=0.01)
-
-
-@pytest.mark.asyncio
-async def test_score_api_queries_vs_documents(hf_model, server: RemoteOpenAIServer):
-    text_pairs = [
-        [TEXTS_1[0], TEXTS_2[0]],
-        [TEXTS_1[1], TEXTS_2[1]],
-    ]
-
-    score_response = requests.post(
-        server.url_for("score"),
-        json={
-            "model": MODEL_NAME,
-            "queries": TEXTS_1,
-            "documents": TEXTS_2,
-        },
-    )
-    score_response.raise_for_status()
-    score = ScoreResponse.model_validate(score_response.json())
-
-    assert score.id is not None
-    assert score.data is not None
-    assert len(score.data) == 2
-
-    vllm_outputs = [d.score for d in score.data]
-    hf_outputs = hf_model.predict(text_pairs).tolist()
-
-    for i in range(len(vllm_outputs)):
-        assert hf_outputs[i] == pytest.approx(vllm_outputs[i], rel=0.01)
-
-
-@pytest.mark.asyncio
-async def test_score_api_queries_vs_items(hf_model, server: RemoteOpenAIServer):
-    text_pairs = [
-        [TEXTS_1[0], TEXTS_2[0]],
-        [TEXTS_1[1], TEXTS_2[1]],
-    ]
-
-    score_response = requests.post(
-        server.url_for("score"),
-        json={
-            "model": MODEL_NAME,
-            "queries": TEXTS_1,
-            "items": TEXTS_2,
-        },
-    )
-    score_response.raise_for_status()
-    score = ScoreResponse.model_validate(score_response.json())
-
-    assert score.id is not None
-    assert score.data is not None
-    assert len(score.data) == 2
-
-    vllm_outputs = [d.score for d in score.data]
-    hf_outputs = hf_model.predict(text_pairs).tolist()
-
-    for i in range(len(vllm_outputs)):
-        assert hf_outputs[i] == pytest.approx(vllm_outputs[i], rel=0.01)
-
-
-@pytest.mark.asyncio
-async def test_score_api_text_1_vs_text_2(hf_model, server: RemoteOpenAIServer):
-    text_pairs = [
-        [TEXTS_1[0], TEXTS_2[0]],
-        [TEXTS_1[1], TEXTS_2[1]],
-    ]
-
-    score_response = requests.post(
-        server.url_for("score"),
-        json={
-            "model": MODEL_NAME,
-            "text_1": TEXTS_1,
-            "text_2": TEXTS_2,
-        },
-    )
-    score_response.raise_for_status()
-    score = ScoreResponse.model_validate(score_response.json())
-
-    assert score.id is not None
-    assert score.data is not None
-    assert len(score.data) == 2
-
-    vllm_outputs = [d.score for d in score.data]
-    hf_outputs = hf_model.predict(text_pairs).tolist()
-
-    for i in range(len(vllm_outputs)):
-        assert hf_outputs[i] == pytest.approx(vllm_outputs[i], rel=0.01)
-
-
-@pytest.mark.asyncio
-async def test_score_api_data_1_vs_data_2(hf_model, server: RemoteOpenAIServer):
-    text_pairs = [
-        [TEXTS_1[0], TEXTS_2[0]],
-        [TEXTS_1[1], TEXTS_2[1]],
-    ]
-
-    score_response = requests.post(
-        server.url_for("score"),
-        json={
-            "model": MODEL_NAME,
-            "data_1": TEXTS_1,
-            "data_2": TEXTS_2,
-        },
-    )
-    score_response.raise_for_status()
-    score = ScoreResponse.model_validate(score_response.json())
-
-    assert score.id is not None
-    assert score.data is not None
-    assert len(score.data) == 2
-
-    vllm_outputs = [d.score for d in score.data]
-    hf_outputs = hf_model.predict(text_pairs).tolist()
-
-    for i in range(len(vllm_outputs)):
-        assert hf_outputs[i] == pytest.approx(vllm_outputs[i], rel=0.01)
+@pytest.mark.parametrize(
+    ("payload", "text_pairs"),
+    [
+        pytest.param(
+            {
+                "queries": TEXTS_1[0],
+                "documents": TEXTS_2[0],
+            },
+            [[TEXTS_1[0], TEXTS_2[0]]],
+            id="queries-str-documents-str",
+        ),
+        pytest.param(
+            {
+                "queries": TEXTS_1[0],
+                "documents": TEXTS_2,
+            },
+            [
+                [TEXTS_1[0], TEXTS_2[0]],
+                [TEXTS_1[0], TEXTS_2[1]],
+            ],
+            id="queries-str-documents-list",
+        ),
+        pytest.param(
+            {
+                "queries": TEXTS_1,
+                "documents": TEXTS_2,
+            },
+            [
+                [TEXTS_1[0], TEXTS_2[0]],
+                [TEXTS_1[1], TEXTS_2[1]],
+            ],
+            id="queries-list-documents-list",
+        ),
+        pytest.param(
+            {
+                "queries": TEXTS_1,
+                "items": TEXTS_2,
+            },
+            [
+                [TEXTS_1[0], TEXTS_2[0]],
+                [TEXTS_1[1], TEXTS_2[1]],
+            ],
+            id="queries-list-items-list",
+        ),
+        pytest.param(
+            {
+                "text_1": TEXTS_1,
+                "text_2": TEXTS_2,
+            },
+            [
+                [TEXTS_1[0], TEXTS_2[0]],
+                [TEXTS_1[1], TEXTS_2[1]],
+            ],
+            id="text-1-vs-text-2",
+        ),
+        pytest.param(
+            {
+                "data_1": TEXTS_1,
+                "data_2": TEXTS_2,
+            },
+            [
+                [TEXTS_1[0], TEXTS_2[0]],
+                [TEXTS_1[1], TEXTS_2[1]],
+            ],
+            id="data-1-vs-data-2",
+        ),
+    ],
+)
+async def test_score_api_request_formats(hf_model, server: RemoteOpenAIServer,
+                                         payload: dict,
+                                         text_pairs: list[list[str]]):
+    _assert_score_output_matches_hf(hf_model, server, payload, text_pairs)
 
 
 @pytest.mark.asyncio

tests/entrypoints/pooling/scoring/test_cross_encoder_online.py

@@ -28,6 +28,25 @@ TEXTS_2 = [
 ]
 
 
+def _assert_score_output_matches_hf(hf_model, server: RemoteOpenAIServer,
+                                    payload: dict, text_pairs: list[list[str]]):
+    score_response = requests.post(
+        server.url_for("score"),
+        json={"model": MODEL_NAME, **payload},
+    )
+    score_response.raise_for_status()
+    score = ScoreResponse.model_validate(score_response.json())
+
+    assert score.id is not None
+    assert score.data is not None
+
+    vllm_outputs = [d.score for d in score.data]
+    hf_outputs = hf_model.predict(text_pairs).tolist()
+    assert len(vllm_outputs) == len(hf_outputs)
+    for hf_output, vllm_output in zip(hf_outputs, vllm_outputs):
+        assert hf_output == pytest.approx(vllm_output, rel=0.01)
+
+
 @pytest.fixture(scope="module")
 def server():
     args = ["--enforce-eager", "--max-model-len", "100", "--dtype", DTYPE]
@@ -61,207 +80,78 @@ async def test_basic(server: RemoteOpenAIServer):
 
 
 @pytest.mark.asyncio
-async def test_score_api_queries_str_1_documents_str_1(
-    hf_model, server: RemoteOpenAIServer
-):
-    score_response = requests.post(
-        server.url_for("score"),
-        json={
-            "model": MODEL_NAME,
-            "queries": TEXTS_1[0],
-            "documents": TEXTS_2[0],
-        },
-    )
-    score_response.raise_for_status()
-    score = ScoreResponse.model_validate(score_response.json())
-
-    assert score.id is not None
-    assert score.data is not None
-    assert len(score.data) == 1
-
-    vllm_outputs = [d.score for d in score.data]
-    hf_outputs = hf_model.predict([[TEXTS_1[0], TEXTS_2[0]]]).tolist()
-
-    for i in range(len(vllm_outputs)):
-        assert hf_outputs[i] == pytest.approx(vllm_outputs[i], rel=0.01)
-
-
-@pytest.mark.asyncio
-async def test_score_api_queries_str_1_documents_str_n(
-    hf_model, server: RemoteOpenAIServer
-):
-    text_pairs = [
-        [TEXTS_1[0], TEXTS_2[0]],
-        [TEXTS_1[0], TEXTS_2[1]],
-    ]
-
-    score_response = requests.post(
-        server.url_for("score"),
-        json={
-            "model": MODEL_NAME,
-            "queries": TEXTS_1[0],
-            "documents": TEXTS_2,
-        },
-    )
-    score_response.raise_for_status()
-    score = ScoreResponse.model_validate(score_response.json())
-
-    assert score.id is not None
-    assert score.data is not None
-    assert len(score.data) == 2
-
-    vllm_outputs = [d.score for d in score.data]
-    hf_outputs = hf_model.predict(text_pairs).tolist()
-
-    for i in range(len(vllm_outputs)):
-        assert hf_outputs[i] == pytest.approx(vllm_outputs[i], rel=0.01)
-
-
-@pytest.mark.asyncio
-async def test_score_api_queries_str_n_documents_str_n(
-    hf_model, server: RemoteOpenAIServer
-):
-    text_pairs = [
-        [TEXTS_1[0], TEXTS_2[0]],
-        [TEXTS_1[1], TEXTS_2[1]],
-    ]
-
-    score_response = requests.post(
-        server.url_for("score"),
-        json={
-            "model": MODEL_NAME,
-            "queries": TEXTS_1,
-            "documents": TEXTS_2,
-        },
-    )
-    score_response.raise_for_status()
-    score = ScoreResponse.model_validate(score_response.json())
-
-    assert score.id is not None
-    assert score.data is not None
-    assert len(score.data) == 2
-
-    vllm_outputs = [d.score for d in score.data]
-    hf_outputs = hf_model.predict(text_pairs).tolist()
-
-    for i in range(len(vllm_outputs)):
-        assert hf_outputs[i] == pytest.approx(vllm_outputs[i], rel=0.01)
-
-
-@pytest.mark.asyncio
-async def test_score_api_queries_vs_documents(hf_model, server: RemoteOpenAIServer):
-    text_pairs = [
-        [TEXTS_1[0], TEXTS_2[0]],
-        [TEXTS_1[1], TEXTS_2[1]],
-    ]
-
-    score_response = requests.post(
-        server.url_for("score"),
-        json={
-            "model": MODEL_NAME,
-            "queries": TEXTS_1,
-            "documents": TEXTS_2,
-        },
-    )
-    score_response.raise_for_status()
-    score = ScoreResponse.model_validate(score_response.json())
-
-    assert score.id is not None
-    assert score.data is not None
-    assert len(score.data) == 2
-
-    vllm_outputs = [d.score for d in score.data]
-    hf_outputs = hf_model.predict(text_pairs).tolist()
-
-    for i in range(len(vllm_outputs)):
-        assert hf_outputs[i] == pytest.approx(vllm_outputs[i], rel=0.01)
-
-
-@pytest.mark.asyncio
-async def test_score_api_queries_vs_items(hf_model, server: RemoteOpenAIServer):
-    text_pairs = [
-        [TEXTS_1[0], TEXTS_2[0]],
-        [TEXTS_1[1], TEXTS_2[1]],
-    ]
-
-    score_response = requests.post(
-        server.url_for("score"),
-        json={
-            "model": MODEL_NAME,
-            "queries": TEXTS_1,
-            "items": TEXTS_2,
-        },
-    )
-    score_response.raise_for_status()
-    score = ScoreResponse.model_validate(score_response.json())
-
-    assert score.id is not None
-    assert score.data is not None
-    assert len(score.data) == 2
-
-    vllm_outputs = [d.score for d in score.data]
-    hf_outputs = hf_model.predict(text_pairs).tolist()
-
-    for i in range(len(vllm_outputs)):
-        assert hf_outputs[i] == pytest.approx(vllm_outputs[i], rel=0.01)
-
-
-@pytest.mark.asyncio
-async def test_score_api_text_1_vs_text_2(hf_model, server: RemoteOpenAIServer):
-    text_pairs = [
-        [TEXTS_1[0], TEXTS_2[0]],
-        [TEXTS_1[1], TEXTS_2[1]],
-    ]
-
-    score_response = requests.post(
-        server.url_for("score"),
-        json={
-            "model": MODEL_NAME,
-            "text_1": TEXTS_1,
-            "text_2": TEXTS_2,
-        },
-    )
-    score_response.raise_for_status()
-    score = ScoreResponse.model_validate(score_response.json())
-
-    assert score.id is not None
-    assert score.data is not None
-    assert len(score.data) == 2
-
-    vllm_outputs = [d.score for d in score.data]
-    hf_outputs = hf_model.predict(text_pairs).tolist()
-
-    for i in range(len(vllm_outputs)):
-        assert hf_outputs[i] == pytest.approx(vllm_outputs[i], rel=0.01)
-
-
-@pytest.mark.asyncio
-async def test_score_api_data_1_vs_data_2(hf_model, server: RemoteOpenAIServer):
-    text_pairs = [
-        [TEXTS_1[0], TEXTS_2[0]],
-        [TEXTS_1[1], TEXTS_2[1]],
-    ]
-
-    score_response = requests.post(
-        server.url_for("score"),
-        json={
-            "model": MODEL_NAME,
-            "data_1": TEXTS_1,
-            "data_2": TEXTS_2,
-        },
-    )
-    score_response.raise_for_status()
-    score = ScoreResponse.model_validate(score_response.json())
-
-    assert score.id is not None
-    assert score.data is not None
-    assert len(score.data) == 2
-
-    vllm_outputs = [d.score for d in score.data]
-    hf_outputs = hf_model.predict(text_pairs).tolist()
-
-    for i in range(len(vllm_outputs)):
-        assert hf_outputs[i] == pytest.approx(vllm_outputs[i], rel=0.01)
+@pytest.mark.parametrize(
+    ("payload", "text_pairs"),
+    [
+        pytest.param(
+            {
+                "queries": TEXTS_1[0],
+                "documents": TEXTS_2[0],
+            },
+            [[TEXTS_1[0], TEXTS_2[0]]],
+            id="queries-str-documents-str",
+        ),
+        pytest.param(
+            {
+                "queries": TEXTS_1[0],
+                "documents": TEXTS_2,
+            },
+            [
+                [TEXTS_1[0], TEXTS_2[0]],
+                [TEXTS_1[0], TEXTS_2[1]],
+            ],
+            id="queries-str-documents-list",
+        ),
+        pytest.param(
+            {
+                "queries": TEXTS_1,
+                "documents": TEXTS_2,
+            },
+            [
+                [TEXTS_1[0], TEXTS_2[0]],
+                [TEXTS_1[1], TEXTS_2[1]],
+            ],
+            id="queries-list-documents-list",
+        ),
+        pytest.param(
+            {
+                "queries": TEXTS_1,
+                "items": TEXTS_2,
+            },
+            [
+                [TEXTS_1[0], TEXTS_2[0]],
+                [TEXTS_1[1], TEXTS_2[1]],
+            ],
+            id="queries-list-items-list",
+        ),
+        pytest.param(
+            {
+                "text_1": TEXTS_1,
+                "text_2": TEXTS_2,
+            },
+            [
+                [TEXTS_1[0], TEXTS_2[0]],
+                [TEXTS_1[1], TEXTS_2[1]],
+            ],
+            id="text-1-vs-text-2",
+        ),
+        pytest.param(
+            {
+                "data_1": TEXTS_1,
+                "data_2": TEXTS_2,
+            },
+            [
+                [TEXTS_1[0], TEXTS_2[0]],
+                [TEXTS_1[1], TEXTS_2[1]],
+            ],
+            id="data-1-vs-data-2",
+        ),
+    ],
+)
+async def test_score_api_request_formats(hf_model, server: RemoteOpenAIServer,
+                                         payload: dict,
+                                         text_pairs: list[list[str]]):
+    _assert_score_output_matches_hf(hf_model, server, payload, text_pairs)
 
 
 @pytest.mark.asyncio

tests/evals/gsm8k/configs/Qwen3.5-397B-A17B-NVFP4-DEP2.yaml

@@ -7,4 +7,3 @@ server_args: >-
   --max-model-len 4096
   --data-parallel-size 2
   --enable-expert-parallel
-  --max-num-seqs 512

tests/kernels/core/test_minimax_reduce_rms.py

@@ -1,152 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""Tests for MiniMax QK RMS-norm: NCCL reference vs Lamport fused kernel."""
-
-import pytest
-import torch
-import torch.nn as nn
-from torch.multiprocessing import spawn
-
-from tests.kernels.utils import opcheck
-from tests.utils import ensure_current_vllm_config, init_test_distributed_environment
-from vllm.distributed import cleanup_dist_env_and_memory
-from vllm.model_executor.layers.mamba.linear_attn import MiniMaxText01RMSNormTP
-from vllm.platforms import current_platform
-from vllm.utils.network_utils import get_open_port
-from vllm.utils.torch_utils import set_random_seed
-
-
-@ensure_current_vllm_config()
-def _worker_forward_qk(
-    local_rank,
-    world_size,
-    port,
-    num_tokens,
-    hidden_q_full,
-    hidden_k_full,
-    dtype,
-    seed,
-    eps,
-):
-    """Per-rank worker: compare NCCL allreduce path vs Lamport fused kernel."""
-
-    if not hasattr(torch.ops._C, "minimax_allreduce_rms_qk"):
-        cleanup_dist_env_and_memory()
-        return
-    device = torch.device(f"cuda:{local_rank}")
-    torch.accelerator.set_device_index(device)
-    init_test_distributed_environment(
-        world_size, 1, local_rank, port, local_rank=local_rank
-    )
-
-    hq = hidden_q_full // world_size
-    hk = hidden_k_full // world_size
-
-    q_norm = MiniMaxText01RMSNormTP(hidden_q_full, eps=eps).cuda()
-    k_norm = MiniMaxText01RMSNormTP(hidden_k_full, eps=eps).cuda()
-
-    set_random_seed(seed)
-    qw = torch.randn(hidden_q_full, dtype=dtype, device="cuda")
-    kw = torch.randn(hidden_k_full, dtype=dtype, device="cuda")
-    q_norm.weight = nn.Parameter(qw[local_rank * hq : (local_rank + 1) * hq])
-    k_norm.weight = nn.Parameter(kw[local_rank * hk : (local_rank + 1) * hk])
-
-    torch.manual_seed(seed + 1000 + local_rank)
-    qkv = torch.randn(num_tokens, hq + hk + hk, dtype=dtype, device="cuda")
-
-    q_ref, k_ref, v_ref = qkv.clone().split([hq, hk, hk], dim=-1)
-    ref_q, ref_k = MiniMaxText01RMSNormTP.forward_qk(q_norm, k_norm, q_ref, k_ref)
-
-    # Set up Lamport workspace.
-    from vllm.distributed.parallel_state import get_tp_group
-    from vllm.model_executor.layers.mamba.lamport_workspace import (
-        get_allreduce_workspace,
-    )
-
-    workspace = get_allreduce_workspace(
-        rank=local_rank,
-        world_size=world_size,
-        max_tokens=num_tokens,
-        process_group=get_tp_group().cpu_group,
-    )
-
-    opcheck(
-        torch.ops._C.minimax_allreduce_rms_qk,
-        (
-            qkv.clone(),
-            q_norm.weight,
-            k_norm.weight,
-            workspace,
-            hq,
-            hk,
-            local_rank,
-            world_size,
-            eps,
-        ),
-    )
-    fused_q, fused_k = torch.ops._C.minimax_allreduce_rms_qk(
-        qkv.clone(),
-        q_norm.weight,
-        k_norm.weight,
-        workspace,
-        hq,
-        hk,
-        local_rank,
-        world_size,
-        eps,
-    )
-    _, _, fused_v = qkv.split([hq, hk, hk], dim=-1)
-    torch.accelerator.synchronize()
-
-    torch.testing.assert_close(
-        fused_q,
-        ref_q,
-        atol=3e-2,
-        rtol=3e-2,
-    )
-    torch.testing.assert_close(fused_k, ref_k, atol=3e-2, rtol=3e-2)
-
-    cleanup_dist_env_and_memory()
-
-
-@pytest.mark.skipif(
-    not current_platform.is_cuda(),
-    reason="CUDA required",
-)
-@pytest.mark.parametrize("world_size", [2, 4, 8])
-@pytest.mark.parametrize("num_tokens", [1, 128, 333])
-@pytest.mark.parametrize(
-    "hidden_dims",
-    [(6144, 1024)],
-)
-@pytest.mark.parametrize("dtype", [torch.bfloat16, torch.float16])
-@pytest.mark.parametrize("eps", [1e-6])
-@pytest.mark.parametrize("seed", [42])
-def test_minimax_reduce_rms_qk(
-    world_size,
-    num_tokens,
-    hidden_dims,
-    dtype,
-    eps,
-    seed,
-):
-    num_gpus = current_platform.device_count()
-    if num_gpus < world_size:
-        pytest.skip(f"Need >= {world_size} GPUs, have {num_gpus}")
-    hidden_q_full, hidden_k_full = hidden_dims
-    port = str(get_open_port())
-    spawn(
-        _worker_forward_qk,
-        args=(
-            world_size,
-            port,
-            num_tokens,
-            hidden_q_full,
-            hidden_k_full,
-            dtype,
-            seed,
-            eps,
-        ),
-        nprocs=world_size,
-        join=True,
-    )

tests/lora/conftest.py

@@ -3,7 +3,6 @@
 
 import tempfile
 from collections import OrderedDict
-from importlib import reload
 from unittest.mock import MagicMock
 
 import pytest
@@ -44,18 +43,6 @@ def cleanup_fixture(should_do_global_cleanup_after_test: bool):
         cleanup_dist_env_and_memory(shutdown_ray=True)
 
 
-@pytest.fixture
-def maybe_enable_lora_dual_stream(monkeypatch: pytest.MonkeyPatch):
-    if current_platform.is_cuda():
-        monkeypatch.setenv("VLLM_LORA_ENABLE_DUAL_STREAM", "1")
-        import vllm.lora.layers.base_linear
-
-        if not hasattr(vllm.lora.layers.base_linear, "lora_linear_async"):
-            # Reload the module to ensure the environment variable takes effect.
-            reload(vllm.lora.layers.base_linear)
-    yield
-
-
 @pytest.fixture
 def dist_init():
     from tests.utils import ensure_current_vllm_config

tests/lora/test_lora_checkpoints.py

@@ -5,9 +5,7 @@ import pytest
 
 from vllm.lora.lora_model import LoRAModel
 from vllm.lora.peft_helper import PEFTHelper
-from vllm.lora.utils import parse_fine_tuned_lora_name
 from vllm.model_executor.models.baichuan import BaiChuanBaseForCausalLM
-from vllm.model_executor.models.gemma4 import Gemma4ForCausalLM
 from vllm.model_executor.models.utils import WeightsMapper
 
 lora_lst = ["baichuan7B", "baichuan7B-zero", "baichuan7B-zero-regex", "chatglm3-6b"]
@@ -130,24 +128,3 @@ def test_lora_weights_mapping(baichuan_lora_files):
     for name in lora_model.loras:
         assert name.startswith(hf_to_vllm_mapper.orig_to_new_prefix["model."])
         assert ".baichuan_layers." in name
-
-
-def test_gemma4_lora_weights_mapping():
-    mapper = Gemma4ForCausalLM.hf_to_vllm_mapper
-    name = "base_model.model.model.language_model.layers.9.mlp.down_proj.lora_A.weight"
-    assert parse_fine_tuned_lora_name(name, mapper) == (
-        "model.layers.9.mlp.down_proj",
-        True,
-    )
-
-
-def test_gemma4_moe_lora_weights_mapping():
-    mapper = Gemma4ForCausalLM.hf_to_vllm_mapper
-    name = (
-        "base_model.model.model.language_model.layers.9.moe.experts."
-        "gate_up_proj.lora_B.weight"
-    )
-    assert parse_fine_tuned_lora_name(name, mapper) == (
-        "model.layers.9.moe.gate_up_proj",
-        False,
-    )

tests/lora/test_minicpmv_tp.py

@@ -1,10 +1,7 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
-from importlib.metadata import version
-
 import pytest
-from packaging.version import Version
 
 import vllm
 from vllm.assets.image import ImageAsset
@@ -13,14 +10,6 @@ from vllm.platforms import current_platform
 
 from ..utils import multi_gpu_test
 
-pytestmark = pytest.mark.skipif(
-    Version("5.0") <= Version(version("transformers")),
-    reason=(
-        "MiniCPMV custom processor uses tokenizer.im_start_id which is not "
-        "available on TokenizersBackend in transformers v5.0+"
-    ),
-)
-
 MODEL_PATH = "openbmb/MiniCPM-Llama3-V-2_5"
 
 PROMPT_TEMPLATE = (

tests/model_executor/model_loader/test_reload.py

@@ -38,10 +38,7 @@ def test_move_metatensors():
 
 def test_reload_lifecycle():
     layer = torch.nn.Linear(2, 3)
-    info = LayerReloadingInfo(
-        restore_metadata=capture_layer_to_meta(layer),
-        restore_device=torch.device("cpu"),
-    )
+    info = LayerReloadingInfo(restore_metadata=capture_layer_to_meta(layer))
 
     restore_layer_on_meta(layer, info)
     for name, tensor in get_layer_tensors(layer).items():
@@ -51,7 +48,7 @@ def test_reload_lifecycle():
         assert tensor.__class__ == meta_tensor.__class__
         assert tensor.__dict__ == meta_tensor.__dict__
 
-    materialize_layer(layer, info)
+    materialize_layer(layer)
     for name, tensor in get_layer_tensors(layer).items():
         materialized_tensor = getattr(layer, name)
         assert tensor.dtype == materialized_tensor.dtype
@@ -63,10 +60,7 @@ def test_reload_lifecycle():
 def test_model_cleanup(dist_init, default_vllm_config):
     layer = QKVParallelLinear(2, 3, 4)
     assert layer.weight.weight_loader.__self__ is layer
-    info = LayerReloadingInfo(
-        restore_metadata=capture_layer_to_meta(layer),
-        restore_device=torch.device("cpu"),
-    )
+    info = LayerReloadingInfo(restore_metadata=capture_layer_to_meta(layer))
 
     mock_info_dict: WeakKeyDictionary[torch.nn.Module, LayerReloadingInfo] = (
         WeakKeyDictionary()
@@ -96,46 +90,39 @@ def test_get_numel_loaded():
     assert ret == "value"
 
 
-@pytest.mark.parametrize(
-    "tp_size", [pytest.param(1), pytest.param(2, marks=[pytest.mark.slow_test])]
-)
+@pytest.mark.parametrize("tp_size", [2])
 @pytest.mark.parametrize(
     "base_model,mul_model,add_model",
     [
-        pytest.param(
+        (
             "Qwen/Qwen3-0.6B",
             "inference-optimization/Qwen3-0.6B-debug-multiply",
             "inference-optimization/Qwen3-0.6B-debug-add",
-            marks=[pytest.mark.slow_test],
         ),
-        pytest.param(
+        (
             "inference-optimization/Qwen3-0.6B-FP8_BLOCK",
             "inference-optimization/Qwen3-0.6B-debug-multiply-FP8_BLOCK",
             "inference-optimization/Qwen3-0.6B-debug-add-FP8_BLOCK",
-            marks=[pytest.mark.slow_test],
         ),
-        pytest.param(
+        (
             "inference-optimization/Qwen3-0.6B-W4A16-G128",
             "inference-optimization/Qwen3-0.6B-debug-multiply-W4A16-G128",
             "inference-optimization/Qwen3-0.6B-debug-add-W4A16-G128",
-            marks=[pytest.mark.slow_test],
         ),
-        pytest.param(
+        (
             "inference-optimization/DeepSeek-V3-debug-empty",
             "inference-optimization/DeepSeek-V3-debug-multiply",
             "inference-optimization/DeepSeek-V3-debug-add",
-            marks=[pytest.mark.slow_test],
         ),
-        pytest.param(
+        (
             "inference-optimization/DeepSeek-V3-debug-empty-FP8_DYNAMIC",
             "inference-optimization/DeepSeek-V3-debug-multiply-FP8_DYNAMIC",
             "inference-optimization/DeepSeek-V3-debug-add-FP8_DYNAMIC",
         ),
-        pytest.param(
+        (
             "inference-optimization/DeepSeek-V3-debug-empty-NVFP4A16",
             "inference-optimization/DeepSeek-V3-debug-multiply-NVFP4A16",
             "inference-optimization/DeepSeek-V3-debug-add-NVFP4A16",
-            marks=[pytest.mark.slow_test],
         ),
     ],
 )
@@ -151,75 +138,6 @@ def test_reload_weights(base_model, mul_model, add_model, tp_size, vllm_runner):
         tensor_parallel_size=tp_size,
         enable_expert_parallel=(tp_size > 1 and "DeepSeek" in base_model),
         enable_prefix_caching=False,
-        max_model_len=16,
-        max_num_seqs=1,
-    ) as llm:
-        llm.collective_rpc("reload_weights", kwargs={"weights_path": mul_model})
-        mul_perp = llm.generate_prompt_perplexity(["3 4 = 12"], mask=["3 4 ="])[0]
-        add_perp = llm.generate_prompt_perplexity(["3 4 = 7"], mask=["3 4 ="])[0]
-        assert mul_perp < add_perp
-
-        llm.collective_rpc("reload_weights", kwargs={"weights_path": add_model})
-        mul_perp = llm.generate_prompt_perplexity(["3 4 = 12"], mask=["3 4 ="])[0]
-        add_perp = llm.generate_prompt_perplexity(["3 4 = 7"], mask=["3 4 ="])[0]
-        assert add_perp < mul_perp
-
-
-@pytest.mark.parametrize(
-    "tp_size", [pytest.param(1), pytest.param(2, marks=[pytest.mark.slow_test])]
-)
-@pytest.mark.parametrize(
-    "base_model,mul_model,add_model,quantization",
-    [
-        pytest.param(
-            "Qwen/Qwen3-0.6B",
-            "inference-optimization/Qwen3-0.6B-debug-multiply",
-            "inference-optimization/Qwen3-0.6B-debug-add",
-            "fp8",
-        ),
-        pytest.param(
-            "inference-optimization/DeepSeek-V3-debug-empty",
-            "inference-optimization/DeepSeek-V3-debug-multiply",
-            "inference-optimization/DeepSeek-V3-debug-add",
-            "fp8",
-            marks=[pytest.mark.slow_test],
-        ),
-        pytest.param(
-            "Qwen/Qwen3-0.6B",
-            "inference-optimization/Qwen3-0.6B-debug-multiply",
-            "inference-optimization/Qwen3-0.6B-debug-add",
-            "mxfp8",
-            marks=[pytest.mark.slow_test],
-        ),
-        pytest.param(
-            "inference-optimization/DeepSeek-V3-debug-empty",
-            "inference-optimization/DeepSeek-V3-debug-multiply",
-            "inference-optimization/DeepSeek-V3-debug-add",
-            "mxfp8",
-            marks=[
-                pytest.mark.slow_test,
-                pytest.mark.xfail(reason="mxfp4 & mla is not supported yet"),
-            ],
-        ),
-    ],
-)
-def test_online_quantize_reload(
-    base_model, mul_model, add_model, quantization, tp_size, vllm_runner
-):
-    if cuda_device_count_stateless() < tp_size:
-        pytest.skip(reason="Not enough CUDA devices")
-
-    if quantization == "fp8" and not current_platform.supports_fp8():
-        pytest.skip(reason="Requires FP8 support")
-
-    with vllm_runner(
-        model_name=base_model,
-        quantization=quantization,
-        tensor_parallel_size=tp_size,
-        enable_expert_parallel=(tp_size > 1 and "DeepSeek" in base_model),
-        enable_prefix_caching=False,
-        max_model_len=16,
-        max_num_seqs=1,
     ) as llm:
         llm.collective_rpc("reload_weights", kwargs={"weights_path": mul_model})
         mul_perp = llm.generate_prompt_perplexity(["3 4 = 12"], mask=["3 4 ="])[0]

tests/model_executor/test_weight_utils.py

@@ -1,6 +1,7 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+import os
 import tempfile
 
 import huggingface_hub.constants
@@ -9,10 +10,26 @@ from huggingface_hub.utils import LocalEntryNotFoundError
 
 from vllm.model_executor.model_loader.weight_utils import (
     download_weights_from_hf,
+    enable_hf_transfer,
     maybe_remap_kv_scale_name,
 )
 
 
+def test_hf_transfer_auto_activation():
+    if "HF_HUB_ENABLE_HF_TRANSFER" in os.environ:
+        # in case it is already set, we can't test the auto activation
+        pytest.skip("HF_HUB_ENABLE_HF_TRANSFER is set, can't test auto activation")
+    enable_hf_transfer()
+    try:
+        # enable hf hub transfer if available
+        import hf_transfer  # type: ignore # noqa
+
+        HF_TRANSFER_ACTIVE = True
+    except ImportError:
+        HF_TRANSFER_ACTIVE = False
+    assert huggingface_hub.constants.HF_HUB_ENABLE_HF_TRANSFER == HF_TRANSFER_ACTIVE
+
+
 def test_download_weights_from_hf():
     with tempfile.TemporaryDirectory() as tmpdir:
         # assert LocalEntryNotFoundError error is thrown
@@ -161,4 +178,5 @@ class TestMaybeRemapKvScaleName:
 
 
 if __name__ == "__main__":
+    test_hf_transfer_auto_activation()
     test_download_weights_from_hf()

tests/models/language/generation/test_common.py

@@ -143,11 +143,6 @@ def test_models(
         # in parts of the operators
         pytest.skip(f"Skipping '{model}' model test with AITER kernel.")
 
-    if current_platform.is_cpu() and model == "TitanML/tiny-mixtral":
-        # This untrained model is sensitive to the rounding error
-        # Fuse ops to reduce bfloat16 rounding
-        monkeypatch.setenv("VLLM_CPU_CI_ENV", "0")
-
     with hf_runner(model) as hf_model:
         hf_outputs = hf_model.generate_greedy_logprobs_limit(
             example_prompts, max_tokens, num_logprobs

tests/models/language/pooling/test_colbert.py

@@ -109,14 +109,6 @@ def _load_hf_model(model_name: str, hf_spec: dict, device: torch.device):
         **extra,
     ).to(device)
     model.eval()
-
-    # Transformers 5.0 weight materialization can clear non-persistent
-    # buffers (e.g. rotary inv_freq) that were registered with
-    # persistent=False.  Re-compute them so the model produces valid output.
-    for mod in model.modules():
-        if hasattr(mod, "_compute_inv_freq") and hasattr(mod, "inv_freq"):
-            mod.inv_freq = mod._compute_inv_freq(device=device)
-
     return model
 
 

tests/models/language/pooling/test_nomic_max_model_len.py

@@ -8,13 +8,7 @@ import pytest
 from ...utils import EmbedModelInfo
 
 MODELS = [
-    EmbedModelInfo(
-        "nomic-ai/nomic-embed-text-v1",
-        # Fixme:
-        #  Update nomic-embed code to support the latest
-        #  HF version and remove revision set.
-        revision="720244025c1a7e15661a174c63cce63c8218e52b",
-    ),
+    EmbedModelInfo("nomic-ai/nomic-embed-text-v1"),
     # EmbedModelInfo("nomic-ai/nomic-embed-text-v1.5"),
     # EmbedModelInfo("nomic-ai/CodeRankEmbed"),
     EmbedModelInfo("nomic-ai/nomic-embed-text-v2-moe"),
@@ -30,10 +24,7 @@ max_model_len = int(original_max_position_embeddings * factor)
 @pytest.mark.parametrize("model_info", MODELS)
 def test_default(model_info, vllm_runner):
     with vllm_runner(
-        model_info.name,
-        revision=model_info.revision,
-        runner="pooling",
-        max_model_len=None,
+        model_info.name, runner="pooling", max_model_len=None
     ) as vllm_model:
         model_config = vllm_model.llm.llm_engine.model_config
         if model_info.name == "nomic-ai/nomic-embed-text-v2-moe":
@@ -48,10 +39,7 @@ def test_default(model_info, vllm_runner):
 def test_set_max_model_len_legal(model_info, vllm_runner):
     # set max_model_len <= 512
     with vllm_runner(
-        model_info.name,
-        revision=model_info.revision,
-        runner="pooling",
-        max_model_len=256,
+        model_info.name, runner="pooling", max_model_len=256
     ) as vllm_model:
         model_config = vllm_model.llm.llm_engine.model_config
         assert model_config.max_model_len == 256
@@ -61,19 +49,11 @@ def test_set_max_model_len_legal(model_info, vllm_runner):
         # For nomic-embed-text-v2-moe the length is set to 512
         # by sentence_bert_config.json.
         with pytest.raises(ValueError):
-            with vllm_runner(
-                model_info.name,
-                revision=model_info.revision,
-                runner="pooling",
-                max_model_len=1024,
-            ):
+            with vllm_runner(model_info.name, runner="pooling", max_model_len=1024):
                 pass
     else:
         with vllm_runner(
-            model_info.name,
-            revision=model_info.revision,
-            runner="pooling",
-            max_model_len=1024,
+            model_info.name, runner="pooling", max_model_len=1024
         ) as vllm_model:
             model_config = vllm_model.llm.llm_engine.model_config
             assert model_config.max_model_len == 1024
@@ -83,12 +63,7 @@ def test_set_max_model_len_legal(model_info, vllm_runner):
 def test_set_max_model_len_illegal(model_info, vllm_runner):
     # set max_model_len > 2048
     with pytest.raises(ValueError):
-        with vllm_runner(
-            model_info.name,
-            revision=model_info.revision,
-            runner="pooling",
-            max_model_len=4096,
-        ):
+        with vllm_runner(model_info.name, runner="pooling", max_model_len=4096):
             pass
 
     # set max_model_len > 2048 by hf_overrides
@@ -96,7 +71,6 @@ def test_set_max_model_len_illegal(model_info, vllm_runner):
     with pytest.raises(ValueError):
         with vllm_runner(
             model_info.name,
-            revision=model_info.revision,
             runner="pooling",
             max_model_len=None,
             hf_overrides=hf_overrides,
@@ -117,11 +91,7 @@ def test_use_rope_scaling_legal(model_info, vllm_runner):
     }
 
     with vllm_runner(
-        model_info.name,
-        revision=model_info.revision,
-        runner="pooling",
-        max_model_len=None,
-        hf_overrides=hf_overrides,
+        model_info.name, runner="pooling", max_model_len=None, hf_overrides=hf_overrides
     ):
         pass
 
@@ -140,7 +110,6 @@ def test_use_rope_scaling_illegal(model_info, vllm_runner):
     with pytest.raises(ValueError):
         with vllm_runner(
             model_info.name,
-            revision=model_info.revision,
             runner="pooling",
             max_model_len=max_model_len + 1,
             hf_overrides=hf_overrides,
@@ -160,7 +129,6 @@ def test_use_rope_scaling_illegal(model_info, vllm_runner):
     with pytest.raises(ValueError):
         with vllm_runner(
             model_info.name,
-            revision=model_info.revision,
             runner="pooling",
             max_model_len=None,
             hf_overrides=hf_overrides,

tests/models/language/pooling/test_splade_sparse_pooler.py

@@ -1,6 +1,8 @@
 # SPDX-License-Identifier: Apache-2.0
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 
+import types
+
 import pytest
 import torch
 import torch.nn as nn
@@ -9,8 +11,6 @@ from vllm.model_executor.models.bert import (
     BertMLMHead,
     SPLADESparsePooler,
 )
-from vllm.pooling_params import PoolingParams
-from vllm.v1.pool.metadata import PoolingMetadata, PoolingStates
 
 # ---------------------------------------------------------------------
 # Functional test: SPLADE formula correctness (no HF download needed)
@@ -38,12 +38,8 @@ def test_splade_pooler_matches_reference_formula(B, T, H, V):
         ],
         dtype=torch.long,
     )
-    meta = PoolingMetadata(
-        prompt_lens=prompt_lens_tenser,
-        prompt_token_ids=token_ids,
-        prompt_token_ids_cpu=token_ids,
-        pooling_params=[PoolingParams(task="embed")] * B,
-        pooling_states=[PoolingStates() for _ in range(B)],
+    meta = types.SimpleNamespace(
+        prompt_lens=prompt_lens_tenser, prompt_token_ids=token_ids
     )
 
     # MLM head (prefer BertMLMHead, fallback to Linear if unavailable)

tests/models/language/pooling_mteb_test/mteb_embed_utils.py

@@ -151,7 +151,6 @@ def mteb_test_embed_models(
 
     with vllm_runner(
         model_info.name,
-        revision=model_info.revision,
         runner="pooling",
         max_model_len=model_info.max_model_len,
         **vllm_extra_kwargs,
@@ -202,7 +201,6 @@ def mteb_test_embed_models(
     if model_info.mteb_score is None:
         with hf_runner(
             model_info.name,
-            revision=model_info.revision,
             is_sentence_transformer=True,
             dtype=ci_envs.VLLM_CI_HF_DTYPE or model_info.hf_dtype,
         ) as hf_model:

tests/models/language/pooling_mteb_test/mteb_score_utils.py

@@ -241,7 +241,6 @@ def mteb_test_rerank_models(
 
     with vllm_runner(
         model_info.name,
-        revision=model_info.revision,
         runner="pooling",
         max_model_len=None,
         max_num_seqs=8,
@@ -287,9 +286,7 @@ def mteb_test_rerank_models(
     # Accelerate mteb test by setting
     # SentenceTransformers mteb score to a constant
     if model_info.mteb_score is None:
-        with hf_runner(
-            model_info.name, revision=model_info.revision, dtype=model_info.hf_dtype
-        ) as hf_model:
+        with hf_runner(model_info.name, dtype=model_info.hf_dtype) as hf_model:
             hf_model.chat_template = chat_template
             st_main_score = run_mteb_rerank(
                 hf_model,

tests/models/language/pooling_mteb_test/test_baai.py

@@ -69,10 +69,7 @@ MODELS = [
         attn_type="decoder",
         is_prefix_caching_supported=True,
         is_chunked_prefill_supported=True,
-        # Skip: model's custom tokenizer on HF hub is incompatible with
-        # transformers v5 (sets attrs before super().__init__, triggering
-        # AttributeError on 'verbose' in __getattr__).
-        enable_test=False,
+        enable_test=True,
     ),
 ]
 

tests/models/language/pooling_mteb_test/test_gte.py

@@ -72,8 +72,7 @@ MODELS = [
         attn_type="encoder_only",
         is_prefix_caching_supported=False,
         is_chunked_prefill_supported=False,
-        # Skip: numerical regression with transformers v5.
-        enable_test=False,
+        enable_test=True,
     ),
     ########## ModernBertModel
     EmbedModelInfo(

tests/models/language/pooling_mteb_test/test_jina.py

@@ -75,10 +75,6 @@ def test_rerank_models_mteb(vllm_runner, model_info: RerankModelInfo) -> None:
     mteb_test_rerank_models(vllm_runner, model_info)
 
 
-@pytest.mark.skip(
-    reason="jinaai/jina-embeddings-v3 custom XLMRobertaLoRA model on HF hub "
-    "is incompatible with transformers v5 (missing all_tied_weights_keys)"
-)
 @pytest.mark.parametrize("model_info", EMBEDDING_MODELS)
 @pytest.mark.parametrize("dtype", ["half"])
 @pytest.mark.parametrize("dimensions", [16, 32])

tests/models/language/pooling_mteb_test/test_nomic.py

@@ -12,10 +12,6 @@ MODELS = [
     EmbedModelInfo(
         "nomic-ai/nomic-embed-text-v1",
         architecture="NomicBertModel",
-        # Fixme:
-        #  Update nomic-embed code to support the latest
-        #  HF version and remove revision set.
-        revision="720244025c1a7e15661a174c63cce63c8218e52b",
         mteb_score=0.737568559,
         enable_test=True,
         seq_pooling_type="MEAN",

tests/models/multimodal/generation/test_common.py

@@ -186,14 +186,7 @@ VLM_TEST_SETTINGS = {
         max_num_seqs=2,
         auto_cls=AutoModel,
         hf_output_post_proc=model_utils.ultravox_trunc_hf_output,
-        marks=[
-            pytest.mark.core_model,
-            pytest.mark.cpu_model,
-            # TODO: Remove skip once model has been upstreamed to Transformers
-            pytest.mark.skip(
-                reason="Custom model code is not compatible with Transformers v5"
-            ),
-        ],
+        marks=[pytest.mark.core_model, pytest.mark.cpu_model],
     ),
     #### Transformers fallback to test
     ## To reduce test burden, we only test batching arbitrary image size
@@ -401,22 +394,6 @@ VLM_TEST_SETTINGS = {
         vllm_runner_kwargs={"mm_processor_kwargs": {"do_pan_and_scan": True}},
         patch_hf_runner=model_utils.gemma3_patch_hf_runner,
     ),
-    "gemma4": VLMTestInfo(
-        models=["google/gemma-4-E2B-it"],
-        test_type=(VLMTestType.IMAGE, VLMTestType.MULTI_IMAGE),
-        prompt_formatter=lambda img_prompt: f"<bos><|turn>user\n{img_prompt}<turn|>\n<|turn>model\n",  # noqa: E501
-        single_image_prompts=IMAGE_ASSETS.prompts(
-            {
-                "stop_sign": "<|image|>What's the content in the center of the image?",  # noqa: E501
-                "cherry_blossom": "<|image|>What is the season?",
-            }
-        ),
-        multi_image_prompt="<|image|><|image|>Describe the two images in detail.",  # noqa: E501
-        max_model_len=4096,
-        max_num_seqs=2,
-        auto_cls=AutoModelForImageTextToText,
-        vllm_runner_kwargs={"limit_mm_per_prompt": {"image": 4}},
-    ),
     "granite_vision": VLMTestInfo(
         models=["ibm-granite/granite-vision-3.3-2b"],
         test_type=(VLMTestType.IMAGE),
@@ -540,12 +517,6 @@ VLM_TEST_SETTINGS = {
         max_model_len=4096,
         use_tokenizer_eos=True,
         patch_hf_runner=model_utils.internvl_patch_hf_runner,
-        # TODO: Remove skip once model has been upstreamed to Transformers
-        marks=[
-            pytest.mark.skip(
-                reason="Custom model code tries to access data from meta-tensor"
-            )
-        ],
     ),
     "intern_vl-video": VLMTestInfo(
         models=[
@@ -558,12 +529,6 @@ VLM_TEST_SETTINGS = {
         use_tokenizer_eos=True,
         patch_hf_runner=model_utils.internvl_patch_hf_runner,
         num_logprobs=10 if current_platform.is_rocm() else 5,
-        # TODO: Remove skip once model has been upstreamed to Transformers
-        marks=[
-            pytest.mark.skip(
-                reason="Custom model code tries to access data from meta-tensor"
-            )
-        ],
     ),
     "intern_vl-hf": VLMTestInfo(
         models=["OpenGVLab/InternVL3-1B-hf"],
@@ -610,8 +575,6 @@ VLM_TEST_SETTINGS = {
         hf_model_kwargs={"device_map": "auto"},
         patch_hf_runner=model_utils.isaac_patch_hf_runner,
         image_size_factors=[(0.25,), (0.25, 0.25, 0.25), (0.25, 0.2, 0.15)],
-        # TODO: Remove skip once model has been upstreamed to Transformers
-        marks=[pytest.mark.skip(reason="Custom model imports deleted object")],  # noqa: E501
     ),
     "kimi_vl": VLMTestInfo(
         models=["moonshotai/Kimi-VL-A3B-Instruct"],
@@ -827,12 +790,7 @@ VLM_TEST_SETTINGS = {
             pytest.mark.skipif(
                 Version(TRANSFORMERS_VERSION) == Version("4.57.3"),
                 reason="This model is broken in Transformers v4.57.3",
-            ),
-            pytest.mark.skipif(
-                Version(TRANSFORMERS_VERSION) >= Version("5.0.0"),
-                reason="Model's custom code uses ROPE_INIT_FUNCTIONS"
-                "['default'] which was removed in transformers v5",
-            ),
+            )
         ],
     ),
     "phi3v": VLMTestInfo(
@@ -986,12 +944,6 @@ VLM_TEST_SETTINGS = {
             )
             for inp in custom_inputs.different_patch_input_cases_internvl()
         ],
-        # TODO: Remove skip once model has been upstreamed to Transformers
-        marks=[
-            pytest.mark.skip(
-                reason="Custom model code tries to access data from meta-tensor"
-            )
-        ],
     ),
     "llava_onevision-multiple-images": VLMTestInfo(
         models=["llava-hf/llava-onevision-qwen2-0.5b-ov-hf"],

tests/models/multimodal/generation/test_nemotron_parse.py

@@ -103,10 +103,6 @@ def run_test(
         )
 
 
-@pytest.mark.skip(
-    reason="Model's custom MBart decoder has head count mismatch with "
-    "transformers v5's GQA-aware cross-attention (8 vs 16 heads)"
-)
 @pytest.mark.parametrize("model", ["nvidia/NVIDIA-Nemotron-Parse-v1.1"])
 @pytest.mark.parametrize("dtype", ["bfloat16"])
 @pytest.mark.parametrize("num_logprobs", [5])

tests/models/multimodal/generation/test_phi4siglip.py

@@ -1,198 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-from collections.abc import Sequence
-from importlib.metadata import version
-
-import pytest
-import regex as re
-from packaging.version import Version
-from transformers import AutoModelForCausalLM, AutoTokenizer
-
-from vllm.logprobs import SampleLogprobs
-from vllm.multimodal.image import rescale_image_size
-
-from ....conftest import (
-    IMAGE_ASSETS,
-    HfRunner,
-    PromptImageInput,
-    VllmRunner,
-)
-from ....utils import multi_gpu_test
-from ...utils import check_logprobs_close
-
-pytestmark = pytest.mark.skipif(
-    Version("5.0") <= Version(version("transformers")),
-    reason=(
-        "vllm upgraded transformers above v5.4 where HF model custom code uses siglip2 "
-        "internals (filter_out_non_signature_kwargs) removed by "
-        "huggingface/transformers#43514"
-    ),
-)
-
-MODEL_ID = "microsoft/Phi-4-reasoning-vision-15B"
-
-HF_IMAGE_PROMPTS = IMAGE_ASSETS.prompts(
-    {
-        "stop_sign": "<|user|>\n<image>\nWhat's the content of the image?<|end|>\n<|assistant|>\n",  # noqa: E501
-        "cherry_blossom": "<|user|>\n<image>\nPlease infer the season with reason in details.<|end|>\n<|assistant|>\n",  # noqa: E501
-    }
-)
-HF_MULTIIMAGE_IMAGE_PROMPT = (
-    "<|user|>\n<image>\n<image>\nDescribe these images.<|end|>\n<|assistant|>\n"  # noqa: E501
-)
-
-DTYPE = "half"
-MAX_TOKENS = 128
-NUM_LOGPROBS = 10
-
-
-def vllm_to_hf_output(
-    vllm_output: tuple[list[int], str, SampleLogprobs | None], model: str
-):
-    """Sanitize vllm output to be comparable with hf output."""
-    _, output_str, out_logprobs = vllm_output
-
-    output_str_without_image = re.sub(r"(<image>)+", "", output_str)
-    if output_str_without_image and output_str_without_image[0] == " ":
-        output_str_without_image = output_str_without_image[1:]
-
-    hf_output_str = output_str_without_image + "<|end|><|endoftext|>"
-
-    tokenizer = AutoTokenizer.from_pretrained(model, trust_remote_code=True)
-    hf_output_ids = tokenizer.encode(output_str_without_image)
-    if hf_output_ids and hf_output_ids[0] == tokenizer.bos_token_id:
-        hf_output_ids = hf_output_ids[1:]
-
-    return hf_output_ids, hf_output_str, out_logprobs
-
-
-def _build_single_image_inputs(
-    image_assets,
-) -> list[tuple[list[str], PromptImageInput]]:
-    """Build single-image inputs for all size_factors at once."""
-    images = [asset.pil_image for asset in image_assets]
-    all_inputs: list[tuple[list[str], PromptImageInput]] = []
-    for size_factors in [[1.0], [0.25, 0.5, 1.0]]:
-        for image, prompt in zip(images, HF_IMAGE_PROMPTS):
-            all_inputs.append(
-                (
-                    [prompt for _ in size_factors],
-                    [rescale_image_size(image, f) for f in size_factors],
-                )
-            )
-    return all_inputs
-
-
-def _build_multi_image_inputs(
-    image_assets,
-) -> list[tuple[list[str], PromptImageInput]]:
-    """Build multi-image inputs for all size_factors at once."""
-    images = [asset.pil_image for asset in image_assets]
-    all_inputs: list[tuple[list[str], PromptImageInput]] = []
-    for size_factors in [[0.5], [0.15, 0.30]]:
-        all_inputs.append(
-            (
-                [HF_MULTIIMAGE_IMAGE_PROMPT for _ in size_factors],
-                [
-                    [rescale_image_size(image, factor) for image in images]
-                    for factor in size_factors
-                ],
-            )
-        )
-    return all_inputs
-
-
-def _run_and_compare(
-    hf_runner: type[HfRunner],
-    vllm_runner: type[VllmRunner],
-    all_inputs: Sequence[tuple[list[str], PromptImageInput]],
-    model: str,
-    max_model_len: int,
-    max_num_seqs: int,
-    mm_limit: int,
-    gpu_memory_utilization: float,
-):
-    """Load each runner once, run all inputs, then compare."""
-    # NOTE: run vLLM first, then HF.  vLLM needs a fresh process without
-    # cuda initialization; running HF first would break the multiprocessing
-    # backend with fork method.
-    with vllm_runner(
-        model,
-        runner="generate",
-        max_model_len=max_model_len,
-        max_num_seqs=max_num_seqs,
-        gpu_memory_utilization=gpu_memory_utilization,
-        dtype=DTYPE,
-        limit_mm_per_prompt={"image": mm_limit},
-        tensor_parallel_size=2,
-        trust_remote_code=True,
-        enforce_eager=True,
-    ) as vllm_model:
-        vllm_outputs_per_case = [
-            vllm_model.generate_greedy_logprobs(
-                prompts,
-                MAX_TOKENS,
-                num_logprobs=NUM_LOGPROBS,
-                images=images,
-            )
-            for prompts, images in all_inputs
-        ]
-
-    hf_model_kwargs = {"_attn_implementation": "sdpa", "device_map": "auto"}
-    with hf_runner(
-        model,
-        dtype=DTYPE,
-        model_kwargs=hf_model_kwargs,
-        auto_cls=AutoModelForCausalLM,
-        trust_remote_code=True,
-    ) as hf_model:
-        hf_outputs_per_case = [
-            hf_model.generate_greedy_logprobs_limit(
-                prompts,
-                MAX_TOKENS,
-                num_logprobs=NUM_LOGPROBS,
-                images=images,
-            )
-            for prompts, images in all_inputs
-        ]
-
-    for hf_outputs, vllm_outputs in zip(hf_outputs_per_case, vllm_outputs_per_case):
-        check_logprobs_close(
-            outputs_0_lst=hf_outputs,
-            outputs_1_lst=vllm_outputs,
-            name_0="hf",
-            name_1="vllm",
-        )
-
-
-@multi_gpu_test(num_gpus=2)
-@pytest.mark.parametrize("model", [MODEL_ID])
-def test_models(hf_runner, vllm_runner, image_assets, model) -> None:
-    all_inputs = _build_single_image_inputs(image_assets)
-    _run_and_compare(
-        hf_runner,
-        vllm_runner,
-        all_inputs,
-        model,
-        max_model_len=8192,
-        max_num_seqs=2,
-        mm_limit=1,
-        gpu_memory_utilization=0.80,
-    )
-
-
-@multi_gpu_test(num_gpus=2)
-@pytest.mark.parametrize("model", [MODEL_ID])
-def test_multi_images_models(hf_runner, vllm_runner, image_assets, model) -> None:
-    all_inputs = _build_multi_image_inputs(image_assets)
-    _run_and_compare(
-        hf_runner,
-        vllm_runner,
-        all_inputs,
-        model,
-        max_model_len=8192,
-        max_num_seqs=2,
-        mm_limit=2,
-        gpu_memory_utilization=0.80,
-    )

tests/models/multimodal/generation/test_voxtral.py

@@ -149,10 +149,6 @@ def test_online_serving(vllm_runner, audio_assets: AudioTestAssets):
     )
 
 
-@pytest.mark.skip(
-    reason="VoxtralProcessor.apply_chat_template() in transformers v5 "
-    "doesn't resolve chat_template=None to the default template"
-)
 def test_hf_reference(hf_runner, vllm_runner, audio_assets: AudioTestAssets):
     """Compare vLLM Mistral-format output against HF Transformers reference.
 

tests/models/multimodal/generation/vlm_utils/core.py

@@ -80,11 +80,6 @@ def run_test(
     if vllm_runner_kwargs:
         vllm_runner_kwargs_.update(vllm_runner_kwargs)
 
-    # Avoid passing limit_mm_per_prompt twice when vllm_runner_kwargs
-    # already contains it (e.g. gemma4 sets it via vllm_runner_kwargs).
-    if "limit_mm_per_prompt" in vllm_runner_kwargs_:
-        limit_mm_per_prompt = vllm_runner_kwargs_.pop("limit_mm_per_prompt")
-
     with vllm_runner(
         model,
         max_model_len=max_model_len,

tests/models/multimodal/pooling/test_colmodernvbert.py

@@ -15,10 +15,6 @@ from vllm.entrypoints.pooling.score.utils import compute_maxsim_score
 MODEL_NAME = "ModernVBERT/colmodernvbert-merged"
 COLBERT_DIM = 128
 DTYPE = "half"
-# Fixme:
-#  Update colmodernvbert code to support the latest HF version
-#  and remove revision set.
-REVISION = "4a0a9f3ac7a7992fec410bfa8e3d080ac9a5bcee"
 
 
 # -----------------------------------------------------------------------
@@ -30,7 +26,6 @@ def test_colmodernvbert_text_token_embed(vllm_runner):
     """Text query produces per-token embeddings with shape (seq_len, 128)."""
     with vllm_runner(
         MODEL_NAME,
-        revision=REVISION,
         runner="pooling",
         dtype=DTYPE,
         enforce_eager=True,
@@ -54,7 +49,6 @@ def test_colmodernvbert_text_relevance_ordering(vllm_runner):
 
     with vllm_runner(
         MODEL_NAME,
-        revision=REVISION,
         runner="pooling",
         dtype=DTYPE,
         enforce_eager=True,
@@ -72,7 +66,6 @@ def test_colmodernvbert_text_late_interaction(vllm_runner):
 
     with vllm_runner(
         MODEL_NAME,
-        revision=REVISION,
         runner="pooling",
         dtype=DTYPE,
         enforce_eager=True,
@@ -99,7 +92,6 @@ def test_colmodernvbert_image_token_embed(vllm_runner, image_assets):
     """Image input produces per-token embeddings including vision tokens."""
     with vllm_runner(
         MODEL_NAME,
-        revision=REVISION,
         runner="pooling",
         dtype=DTYPE,
         enforce_eager=True,

tests/models/multimodal/pooling/test_colqwen3.py

@@ -22,11 +22,6 @@ from vllm.entrypoints.pooling.score.utils import ScoreMultiModalParam
 
 from ....conftest import VllmRunner
 
-pytestmark = pytest.mark.skip(
-    reason="ColQwen3 model's weight tying is incompatible with "
-    "transformers v5 (missing all_tied_weights_keys)"
-)
-
 MODELS = [
     "TomoroAI/tomoro-colqwen3-embed-4b",
     "OpenSearch-AI/Ops-Colqwen3-4B",

tests/models/multimodal/pooling/test_intern_vit.py

@@ -11,11 +11,6 @@ from vllm.utils.torch_utils import STR_DTYPE_TO_TORCH_DTYPE
 
 from ....conftest import ImageTestAssets
 
-pytestmark = pytest.mark.skip(
-    reason="InternVisionModel's custom code is incompatible with "
-    "transformers v5 (missing all_tied_weights_keys)"
-)
-
 # we use snapshot_download to prevent conflicts between
 # dynamic_module and trust_remote_code for hf_runner
 DOWNLOAD_PATTERN = ["*.json", "*.py", "*.safetensors", "*.txt", "*.model"]

tests/models/multimodal/pooling/test_jinavl_reranker.py

@@ -15,11 +15,6 @@ from vllm.entrypoints.pooling.score.utils import ScoreMultiModalParam
 
 from ....conftest import HfRunner, VllmRunner
 
-pytestmark = pytest.mark.skip(
-    reason="jinaai/jina-reranker-m0 custom code is incompatible with "
-    "transformers v5 (missing all_tied_weights_keys)"
-)
-
 MODELS = ["jinaai/jina-reranker-m0"]
 
 MM_PROCESSOR_KWARGS = {

tests/models/multimodal/pooling/test_qwen3_asr_forced_aligner.py

@@ -1,67 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-import numpy as np
-import pytest
-import torch
-
-MODEL = "Qwen/Qwen3-ForcedAligner-0.6B"
-CLASSIFY_NUM = 5000
-TIMESTAMP_TOKEN_ID = 151705
-
-
-def build_prompt(words: list[str]) -> str:
-    body = "<timestamp><timestamp>".join(words) + "<timestamp><timestamp>"
-    return f"<|audio_start|><|audio_pad|><|audio_end|>{body}"
-
-
-@pytest.mark.parametrize("model", [MODEL])
-@pytest.mark.parametrize("dtype", ["bfloat16"])
-@torch.inference_mode()
-def test_qwen3_forced_aligner(
-    vllm_runner,
-    model: str,
-    dtype: str,
-) -> None:
-    words = ["Hello", "world"]
-    prompt = build_prompt(words)
-
-    # 5-second silent audio at 16kHz
-    audio = np.zeros(16000 * 5, dtype=np.float32)
-
-    with vllm_runner(
-        model,
-        runner="pooling",
-        dtype=dtype,
-        enforce_eager=True,
-        max_model_len=512,
-        hf_overrides={
-            "architectures": [
-                "Qwen3ASRForcedAlignerForTokenClassification",
-            ],
-        },
-    ) as vllm_model:
-        outputs = vllm_model.llm.encode(
-            [{"prompt": prompt, "multi_modal_data": {"audio": audio}}],
-            pooling_task="token_classify",
-        )
-
-    # Validate output structure
-    assert len(outputs) == 1
-    logits = outputs[0].outputs.data
-    assert logits.dim() == 2
-    assert logits.shape[1] == CLASSIFY_NUM
-
-    # Validate timestamp extraction
-    token_ids = outputs[0].prompt_token_ids
-    predictions = logits.argmax(dim=-1)
-    ts_indices = [i for i, t in enumerate(token_ids) if t == TIMESTAMP_TOKEN_ID]
-
-    # 2 words x 2 timestamps each (start + end) = 4
-    assert len(ts_indices) == 4
-
-    ts_preds = [predictions[i].item() for i in ts_indices]
-    assert all(p >= 0 for p in ts_preds)
-    # end >= start for each word
-    assert ts_preds[1] >= ts_preds[0]  # Hello
-    assert ts_preds[3] >= ts_preds[2]  # world

tests/models/multimodal/processing/test_gemma4.py

@@ -1,44 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-import pytest
-
-from vllm.multimodal import MULTIMODAL_REGISTRY
-
-from ....conftest import ImageTestAssets
-from ...utils import build_model_context
-
-# TODO: to be updated to "google/gemma-4-e2b-it" once the models are available
-GEMMA4_MODEL_ID = "google/gemma-4-E2B-it"
-
-
-@pytest.mark.parametrize("model_id", [GEMMA4_MODEL_ID])
-def test_limit_mm_per_prompt(
-    image_assets: ImageTestAssets,
-    model_id: str,
-):
-    """Test that limit_mm_per_prompt accurately restricts multiple images."""
-    # We only allow 1 image
-    ctx = build_model_context(
-        model_id,
-        mm_processor_kwargs={},
-        limit_mm_per_prompt={"image": 1},
-    )
-    processor = MULTIMODAL_REGISTRY.create_processor(ctx.model_config)
-
-    # Provide 2 images in the prompt
-    prompt = "<image><image>"
-    # image_assets usually has multiple images
-    images = [asset.pil_image for asset in image_assets][:2]
-    if len(images) < 2:
-        images = [images[0], images[0]]
-
-    mm_data = {"image": images}
-
-    # Expect ValueError when exceeding limit
-    with pytest.raises(ValueError, match="At most 1 image"):
-        processor(
-            prompt,
-            mm_items=processor.info.parse_mm_data(mm_data),
-            hf_processor_mm_kwargs={},
-        )

tests/models/multimodal/processing/test_musicflamingo.py

@@ -17,13 +17,11 @@
 # See the License for the specific language governing permissions and
 # limitations under the License.
 
-from importlib.metadata import version
 from unittest.mock import MagicMock
 
 import numpy as np
 import pytest
 import torch
-from packaging.version import Version
 from transformers import PretrainedConfig
 
 from tests.models.registry import HF_EXAMPLE_MODELS
@@ -124,11 +122,6 @@ def test_musicflamingo_dummy_text_uses_plain_audio_tokens(mock_ctx):
     assert builder.get_dummy_text({"audio": 2}) == "<sound><sound>"
 
 
-@pytest.mark.skipif(
-    Version(version("transformers")) >= Version("5.5"),
-    reason="transformers v5.5 added native MusicFlamingoForConditionalGeneration "
-    "with a different get_audio_features signature (requires input_ids)",
-)
 def test_musicflamingo_audio_feature_pipeline_matches_hf_small_config():
     from transformers.models.musicflamingo import (
         modeling_musicflamingo as hf_musicflamingo_modeling,

tests/models/registry.py

@@ -277,10 +277,6 @@ _TEXT_GENERATION_EXAMPLE_MODELS = {
         "google/gemma-2-9b", extras={"tiny": "google/gemma-2-2b-it"}
     ),
     "Gemma3ForCausalLM": _HfExamplesInfo("google/gemma-3-1b-it"),
-    "Gemma4ForCausalLM": _HfExamplesInfo(
-        "google/gemma-4-E2B-it",
-        min_transformers_version="5.0.0",
-    ),
     "Gemma3nForCausalLM": _HfExamplesInfo("google/gemma-3n-E2B-it"),
     "GlmForCausalLM": _HfExamplesInfo("zai-org/glm-4-9b-chat-hf"),
     "Glm4ForCausalLM": _HfExamplesInfo("zai-org/GLM-4-9B-0414"),
@@ -334,15 +330,7 @@ _TEXT_GENERATION_EXAMPLE_MODELS = {
         "internlm/internlm2-chat-7b", trust_remote_code=True
     ),
     "InternLM2VEForCausalLM": _HfExamplesInfo(
-        "OpenGVLab/Mono-InternVL-2B",
-        trust_remote_code=True,
-        max_transformers_version="4.57",
-        transformers_version_reason={
-            "vllm": (
-                "Custom config cannot be loaded with Transformers "
-                "v5 because `vision_config` is not always set"
-            )
-        },
+        "OpenGVLab/Mono-InternVL-2B", trust_remote_code=True
     ),
     "InternLM3ForCausalLM": _HfExamplesInfo(
         "internlm/internlm3-8b-instruct", trust_remote_code=True
@@ -477,13 +465,6 @@ _TEXT_GENERATION_EXAMPLE_MODELS = {
     "Plamo2ForCausalLM": _HfExamplesInfo(
         "pfnet/plamo-2-1b",
         trust_remote_code=True,
-        max_transformers_version="4.57",
-        transformers_version_reason={
-            "hf": (
-                "Custom model code uses `_tied_weight_keys: list[str]` but "
-                "Transformers v5 now expects `_tied_weight_keys: dict[str, str]`"
-            )
-        },
     ),
     "Plamo3ForCausalLM": _HfExamplesInfo(
         "pfnet/plamo-3-nict-2b-base",
@@ -524,13 +505,6 @@ _TEXT_GENERATION_EXAMPLE_MODELS = {
         trust_remote_code=True,
         max_model_len=4096,
         is_available_online=True,
-        max_transformers_version="5.3",
-        transformers_version_reason={
-            "vllm": (
-                "vllm upgraded transformers above v5.4 where "
-                "validate_rope() no longer accepts ignore_keys param"
-            )
-        },
     ),
     "SeedOssForCausalLM": _HfExamplesInfo(
         "ByteDance-Seed/Seed-OSS-36B-Instruct",
@@ -566,11 +540,6 @@ _TEXT_GENERATION_EXAMPLE_MODELS = {
         "xverse/XVERSE-7B-Chat",
         tokenizer="meta-llama/Llama-2-7b",
         trust_remote_code=True,
-        max_transformers_version="4.57",
-        transformers_version_reason={
-            "vllm": "XVERSE tokenizer is incompatible with transformers v5 "
-            "(add_prefix_space / prepend_scheme mismatch).",
-        },
     ),
     "Zamba2ForCausalLM": _HfExamplesInfo("Zyphra/Zamba2-7B-instruct"),
     "MiMoForCausalLM": _HfExamplesInfo("XiaomiMiMo/MiMo-7B-RL", trust_remote_code=True),
@@ -667,7 +636,6 @@ _LATE_INTERACTION_EXAMPLE_MODELS = {
     # [Multimodal]
     "ColModernVBertForRetrieval": _HfExamplesInfo(
         "ModernVBERT/colmodernvbert-merged",
-        revision="4a0a9f3ac7a7992fec410bfa8e3d080ac9a5bcee",
     ),
     "ColPaliForRetrieval": _HfExamplesInfo("vidore/colpali-v1.3-hf"),
     "ColQwen3": _HfExamplesInfo(
@@ -781,18 +749,10 @@ _MULTIMODAL_EXAMPLE_MODELS = {
     # [Decoder-only]
     "AriaForConditionalGeneration": _HfExamplesInfo("rhymes-ai/Aria"),
     "AudioFlamingo3ForConditionalGeneration": _HfExamplesInfo(
-        "nvidia/audio-flamingo-3-hf",
-        min_transformers_version="5.3.0",
-        transformers_version_reason={
-            "vllm": "Needs https://github.com/huggingface/transformers/pull/43538"
-        },
+        "nvidia/audio-flamingo-3-hf", min_transformers_version="5.0.0"
     ),
     "MusicFlamingoForConditionalGeneration": _HfExamplesInfo(
-        "nvidia/music-flamingo-2601-hf",
-        min_transformers_version="5.3.0",
-        transformers_version_reason={
-            "vllm": "Needs https://github.com/huggingface/transformers/pull/43538"
-        },
+        "nvidia/music-flamingo-2601-hf", min_transformers_version="5.3.0"
     ),
     "AyaVisionForConditionalGeneration": _HfExamplesInfo("CohereLabs/aya-vision-8b"),
     "BagelForConditionalGeneration": _HfExamplesInfo("ByteDance-Seed/BAGEL-7B-MoT"),
@@ -831,44 +791,18 @@ _MULTIMODAL_EXAMPLE_MODELS = {
     "Ernie4_5_VLMoeForConditionalGeneration": _HfExamplesInfo(
         "baidu/ERNIE-4.5-VL-28B-A3B-PT",
         trust_remote_code=True,
-        revision="refs/pr/17",
     ),
     "FireRedASR2ForConditionalGeneration": _HfExamplesInfo(
         "allendou/FireRedASR2-LLM-vllm",
-        trust_remote_code=True,
-        max_transformers_version="5.1",
-        transformers_version_reason={
-            "vllm": "Incompatible with transformers v5.2+ "
-            "(dict object has no attribute '__name__').",
-        },
-    ),
-    "FireRedLIDForConditionalGeneration": _HfExamplesInfo(
-        "PatchyTisa/FireRedLID-vllm",
-        trust_remote_code=True,
-        max_transformers_version="5.1",
-        transformers_version_reason={
-            "vllm": "Incompatible with transformers v5.2+ "
-            "(dict object has no attribute '__name__').",
-        },
     ),
     "FunASRForConditionalGeneration": _HfExamplesInfo(
         "allendou/Fun-ASR-Nano-2512-vllm",
-        trust_remote_code=True,
-        max_transformers_version="5.1",
-        transformers_version_reason={
-            "vllm": "Incompatible with transformers v5.2+ "
-            "(dict object has no attribute '__name__').",
-        },
     ),
     "FunAudioChatForConditionalGeneration": _HfExamplesInfo(
         "funaudiochat", is_available_online=False
     ),
     "FuyuForCausalLM": _HfExamplesInfo("adept/fuyu-8b"),
     "Gemma3ForConditionalGeneration": _HfExamplesInfo("google/gemma-3-4b-it"),
-    "Gemma4ForConditionalGeneration": _HfExamplesInfo(
-        "google/gemma-4-E2B-it",
-        min_transformers_version="5.5.0",
-    ),
     "Gemma3nForConditionalGeneration": _HfExamplesInfo("google/gemma-3n-E2B-it"),
     "GlmAsrForConditionalGeneration": _HfExamplesInfo(
         "zai-org/GLM-ASR-Nano-2512",
@@ -900,13 +834,6 @@ _MULTIMODAL_EXAMPLE_MODELS = {
     "HCXVisionForCausalLM": _HfExamplesInfo(
         "naver-hyperclovax/HyperCLOVAX-SEED-Vision-Instruct-3B",
         trust_remote_code=True,
-        max_transformers_version="4.57",
-        transformers_version_reason={
-            "vllm": (
-                "Custom config cannot be loaded with Transformers "
-                "v5 because `text_config` is not always set"
-            )
-        },
     ),
     "HCXVisionV2ForCausalLM": _HfExamplesInfo(
         "naver-hyperclovax/HyperCLOVAX-SEED-Think-32B",
@@ -926,12 +853,7 @@ _MULTIMODAL_EXAMPLE_MODELS = {
         extras={"0.2-2B-Preview": "PerceptronAI/Isaac-0.2-2B-Preview"},
     ),
     "InternS1ForConditionalGeneration": _HfExamplesInfo(
-        "internlm/Intern-S1",
-        trust_remote_code=True,
-        max_transformers_version="4.57",
-        transformers_version_reason={
-            "vllm": "Custom tokenizer code is not compatible with Transformers v5."
-        },
+        "internlm/Intern-S1", trust_remote_code=True
     ),
     "InternS1ProForConditionalGeneration": _HfExamplesInfo(
         "internlm/Intern-S1-Pro",
@@ -1020,14 +942,7 @@ _MULTIMODAL_EXAMPLE_MODELS = {
     "MiDashengLMModel": _HfExamplesInfo(
         "mispeech/midashenglm-7b", trust_remote_code=True
     ),
-    "MiniCPMO": _HfExamplesInfo(
-        "openbmb/MiniCPM-o-2_6",
-        trust_remote_code=True,
-        max_transformers_version="4.57",
-        transformers_version_reason={
-            "hf": "Custom processor code is not compatible with Transformers v5."
-        },
-    ),
+    "MiniCPMO": _HfExamplesInfo("openbmb/MiniCPM-o-2_6", trust_remote_code=True),
     "MiniCPMV": _HfExamplesInfo(
         "openbmb/MiniCPM-Llama3-V-2_5",
         extras={
@@ -1035,13 +950,6 @@ _MULTIMODAL_EXAMPLE_MODELS = {
             "4.0": "openbmb/MiniCPM-V-4",
             "4.5": "openbmb/MiniCPM-V-4_5",
         },
-        max_transformers_version="4.57",
-        transformers_version_reason={
-            "vllm": (
-                "MiniCPMVBatchFeature is incompatible with its base class in "
-                "Transformers v5. See https://huggingface.co/openbmb/MiniCPM-Llama3-V-2_5/discussions/78"
-            )
-        },
         trust_remote_code=True,
     ),
     "MiniMaxVL01ForConditionalGeneration": _HfExamplesInfo(
@@ -1078,25 +986,13 @@ _MULTIMODAL_EXAMPLE_MODELS = {
         "nano_vl_dummy", is_available_online=False, trust_remote_code=True
     ),
     "OpenCUAForConditionalGeneration": _HfExamplesInfo(
-        "xlangai/OpenCUA-7B",
-        trust_remote_code=True,
-        max_transformers_version="4.57",
-        transformers_version_reason={
-            "vllm": "Tokenizer cannot be initialised in Transformers v5."
-        },
+        "xlangai/OpenCUA-7B", trust_remote_code=True
     ),
     "OpenPanguVLForConditionalGeneration": _HfExamplesInfo(
         "FreedomIntelligence/openPangu-VL-7B",
         trust_remote_code=True,
         max_model_len=4096,
         enforce_eager=True,
-        max_transformers_version="4.57",
-        transformers_version_reason={
-            "vllm": (
-                "OpenPanguVLVideoProcessorInitKwargs does not specify total=False, "
-                "making all kwargs required. See https://huggingface.co/FreedomIntelligence/openPangu-VL-7B/discussions/2"
-            )
-        },
     ),
     "Ovis": _HfExamplesInfo(
         "AIDC-AI/Ovis2-1B",
@@ -1108,24 +1004,12 @@ _MULTIMODAL_EXAMPLE_MODELS = {
             "1.6-gemma": "AIDC-AI/Ovis1.6-Gemma2-9B",
         },
     ),
-    "Ovis2_5": _HfExamplesInfo(
-        "AIDC-AI/Ovis2.5-2B",
-        trust_remote_code=True,
-        max_transformers_version="4.57",
-        transformers_version_reason={
-            "vllm": "Custom processor code is not compatible with Transformers v5."
-        },
-    ),
+    "Ovis2_5": _HfExamplesInfo("AIDC-AI/Ovis2.5-2B", trust_remote_code=True),
     "Ovis2_6ForCausalLM": _HfExamplesInfo(
         "AIDC-AI/Ovis2.6-2B", is_available_online=False, trust_remote_code=True
     ),
     "Ovis2_6_MoeForCausalLM": _HfExamplesInfo(
-        "AIDC-AI/Ovis2.6-30B-A3B",
-        trust_remote_code=True,
-        max_transformers_version="4.57",
-        transformers_version_reason={
-            "vllm": "Custom processor code is not compatible with Transformers v5."
-        },
+        "AIDC-AI/Ovis2.6-30B-A3B", trust_remote_code=True
     ),
     "PaddleOCRVLForConditionalGeneration": _HfExamplesInfo(
         "PaddlePaddle/PaddleOCR-VL",
@@ -1144,19 +1028,6 @@ _MULTIMODAL_EXAMPLE_MODELS = {
         },  # noqa: E501
         extras={"phi3.5": "microsoft/Phi-3.5-vision-instruct"},
     ),
-    "Phi4ForCausalLMV": _HfExamplesInfo(
-        "microsoft/Phi-4-reasoning-vision-15B",
-        trust_remote_code=True,
-        max_transformers_version="5.3",
-        transformers_version_reason={
-            "vllm": (
-                "vllm upgraded transformers above v5.4 where HF model "
-                "custom code uses siglip2 internals "
-                "(filter_out_non_signature_kwargs) removed "
-                "by huggingface/transformers#43514"
-            )
-        },
-    ),
     "Phi4MMForCausalLM": _HfExamplesInfo(
         "microsoft/Phi-4-multimodal-instruct", trust_remote_code=True
     ),
@@ -1222,12 +1093,6 @@ _MULTIMODAL_EXAMPLE_MODELS = {
         min_transformers_version="4.57",
         hf_overrides={"architectures": ["Qwen3ASRRealtimeGeneration"]},
     ),
-    "Qwen3ASRForcedAlignerForTokenClassification": _HfExamplesInfo(
-        "Qwen/Qwen3-ForcedAligner-0.6B",
-        max_model_len=4096,
-        min_transformers_version="4.57",
-        hf_overrides={"architectures": ["Qwen3ASRForcedAlignerForTokenClassification"]},
-    ),
     "RForConditionalGeneration": _HfExamplesInfo("YannQi/R-4B", trust_remote_code=True),
     "SkyworkR1VChatModel": _HfExamplesInfo(
         "Skywork/Skywork-R1V-38B", trust_remote_code=True
@@ -1252,14 +1117,6 @@ _MULTIMODAL_EXAMPLE_MODELS = {
             "architectures": ["Tarsier2ForConditionalGeneration"],
             "model_type": "tarsier2",
         },
-        max_transformers_version="5.3",
-        transformers_version_reason={
-            "vllm": (
-                "Qwen2VLConfig was split into Qwen2VLConfig + "
-                "Qwen2VLTextConfig in transformers v5, breaking "
-                "attribute access (num_attention_heads, hidden_size, etc.)"
-            )
-        },
     ),
     "VoxtralForConditionalGeneration": _HfExamplesInfo(
         "mistralai/Voxtral-Mini-3B-2507",