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

Support meituan-longcat/LongCat-2.0-FP8: the LongCat-Flash-Lite backbone
plus a DeepSeek-V3.2-style sparse attention indexer, reusing vLLM's
existing DSA infrastructure.

- Resolve the checkpoint's null model_type via a registered config class
  (arch LongcatCausalLM); alias the oe_* n-gram config fields.
- Wire the indexer into the dual-attention layer: only the first
  attention of each layer computes top-k indices (cli_factor), the
  second reuses them via the shared buffer. Adds a generic skip_topk
  override to DeepseekV2MLAAttention; the schedule lives in the model.
- Streaming-aware indexing: force the first index_init_tokens and last
  index_local_tokens into the top-k set (sparse_attn_indexer).
- MTP speculative decoding (one module iterated up to 3 draft steps,
  plain token embedding, FP8 indexer weights).
- Fix a latent LongCat weight-loading bug: the post-load
  mla_scale_q_lora/kv_lora fold breaks under incremental load_weights
  calls; fold at weight-load time instead.
- Fix two n-gram embedding bugs (also affect LongCat-Flash-Lite):
  rebuild the left-context from the accepted-token history so rejected
  draft tokens don't pollute it, and hash EOS-current positions with
  full look-back per the HF reference (unit test included).
- Add --reasoning-parser longcat (<longcat_think> tags) and docs rows.
- Add longcat model types to the DeepGEMM Blackwell blocklist (non-ue8m0
  scales; measured GSM8K-neutral on vLLM but matches SGLang's guard).

Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: mgoin <mgoin64@gmail.com>

.buildkite/ci_config_rocm.yaml

@@ -8,6 +8,7 @@ run_all_patterns:
   - "docker/docker-bake-rocm.hcl"
   - ".buildkite/hardware_tests/amd.yaml"
   - ".buildkite/scripts/ci-bake-rocm.sh"
+  - ".buildkite/scripts/rocm/"
   - ".buildkite/scripts/hardware_ci/run-amd-test.py"
   - ".buildkite/scripts/hardware_ci/run-amd-test.sh"
   - "CMakeLists.txt"

.buildkite/hardware_tests/amd.yaml

@@ -1,5 +1,28 @@
 group: Hardware - AMD Build
+
+# ROCm image flow:
+# 1. Refresh the long-lived ROCm base image only when Dockerfile.rocm_base changes.
+# 2. Build ci_base from either the stable base or the freshly refreshed base.
+# 3. Build the per-commit ROCm CI image and smoke-test it before GPU jobs run.
 steps:
+  - label: "AMD: :docker: refresh ROCm base"
+    key: refresh-rocm-base-amd
+    depends_on: []
+    device: amd_cpu
+    no_plugin: true
+    commands:
+      - bash .buildkite/scripts/rocm/refresh-base-image.sh
+    env:
+      DOCKER_BUILDKIT: "1"
+      BUILDKIT_PROGRESS: "tty"
+      TERM: "xterm-256color"
+    retry:
+      automatic:
+        - exit_status: -1  # Agent was lost
+          limit: 1
+        - exit_status: -10  # Agent was lost
+          limit: 1
+
   # Ensure ci_base is up-to-date before building the test image.
   # Compares a content hash of ci_base-affecting files against the remote
   # image label. If hashes match the build is skipped (< 30 s); if they
@@ -7,13 +30,16 @@ steps:
   - label: "AMD: :docker: ensure ci_base"
     key: ensure-ci-base-amd
     soft_fail: false
-    depends_on: []
+    depends_on:
+      - refresh-rocm-base-amd
     device: amd_cpu
     no_plugin: true
     commands:
-      - bash .buildkite/scripts/ci-bake-rocm.sh ci-base-rocm-ci-with-deps
+      - bash .buildkite/scripts/rocm/build-ci-base.sh
     env:
       DOCKER_BUILDKIT: "1"
+      BUILDKIT_PROGRESS: "tty"
+      TERM: "xterm-256color"
       VLLM_BAKE_FILE: "docker/docker-bake-rocm.hcl"
       PYTORCH_ROCM_ARCH: "gfx90a;gfx942;gfx950"
       REMOTE_VLLM: "1"
@@ -33,35 +59,12 @@ steps:
     device: amd_cpu
     no_plugin: true
     commands:
-      - |
-        if [[ "${ROCM_CI_ARTIFACT_ONLY:-0}" == "1" ]]; then
-          echo "ROCM_CI_ARTIFACT_ONLY=1; building ROCm wheel artifact only"
-          IMAGE_TAG="" bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-artifacts
-        else
-          bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-wheel
-        fi
-      - |
-        docker run --rm --network=none --entrypoint /bin/bash "rocm/vllm-ci:${BUILDKITE_COMMIT}" -ec '
-          if [ ! -d /vllm-workspace ]; then echo Missing directory: /vllm-workspace >&2; exit 1; fi
-          if [ ! -d /vllm-workspace/tests ]; then echo Missing directory: /vllm-workspace/tests >&2; exit 1; fi
-          if [ ! -d /vllm-workspace/src/vllm ]; then echo Missing directory: /vllm-workspace/src/vllm >&2; exit 1; fi
-          if [ ! -x /vllm-workspace/src/vllm/vllm-rs ]; then echo Missing executable: /vllm-workspace/src/vllm/vllm-rs >&2; exit 1; fi
-          command -v python3
-          command -v uv
-          command -v pytest
-          if ! command -v amd-smi >/dev/null 2>&1 && ! command -v rocminfo >/dev/null 2>&1; then
-            echo No ROCm CLI found in image >&2
-            exit 1
-          fi
-          python3 - <<PY
-        import torch, vllm
-        print(torch.__version__)
-        print(vllm.__version__)
-        PY
-          echo AMD image smoke OK
-        '
+      - bash .buildkite/scripts/rocm/build-test-image.sh
+      - bash .buildkite/scripts/rocm/smoke-test-image.sh
     env:
       DOCKER_BUILDKIT: "1"
+      BUILDKIT_PROGRESS: "tty"
+      TERM: "xterm-256color"
       VLLM_BAKE_FILE: "docker/docker-bake-rocm.hcl"
       PYTORCH_ROCM_ARCH: "gfx90a;gfx942;gfx950"
       IMAGE_TAG: "rocm/vllm-ci:$BUILDKITE_COMMIT"

.buildkite/hardware_tests/cpu.yaml

@@ -17,12 +17,14 @@ steps:
   - tests/kernels/test_awq_int4_to_int8.py
   - tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
   - tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
+  - tests/kernels/mamba/test_cpu_short_conv.py
   commands:
     - |
       bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
       pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
       pytest -x -v -s tests/kernels/moe/test_cpu_fused_moe.py
       pytest -x -v -s tests/kernels/moe/test_cpu_quant_fused_moe.py
+      pytest -x -v -s tests/kernels/mamba/test_cpu_short_conv.py
       pytest -x -v -s tests/kernels/test_onednn.py
       pytest -x -v -s tests/kernels/test_awq_int4_to_int8.py
       pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py

.buildkite/hardware_tests/intel_xpu_ci/test-intel.yaml

@@ -45,7 +45,7 @@ steps:
     agent_tags:
       label: production
       gpu: 1+
-      mem: 16+
+      mem: 24+
     no_plugin: true
     env:
       REGISTRY: "public.ecr.aws/q9t5s3a7"

.buildkite/image_build/image_build.sh

@@ -79,12 +79,18 @@ setup_buildx_builder() {
     docker buildx ls | grep -E '^\*|^NAME' || docker buildx ls
 }
 
+annotate_image_tags() {
+    .buildkite/scripts/annotate-image-build.sh \
+        "${IMAGE_TAG:-}" "${IMAGE_TAG_LATEST:-}"
+}
+
 check_and_skip_if_image_exists() {
     if [[ -n "${IMAGE_TAG:-}" ]]; then
         echo "--- :mag: Checking if image exists"
         if docker manifest inspect "${IMAGE_TAG}" >/dev/null 2>&1; then
             echo "Image already exists: ${IMAGE_TAG}"
             echo "Skipping build"
+            annotate_image_tags
             exit 0
         fi
         echo "Image not found, proceeding with build"
@@ -171,6 +177,18 @@ BRANCH=$4
 IMAGE_TAG=$5
 IMAGE_TAG_LATEST=${6:-} # only used for main branch, optional
 
+# When TORCH_NIGHTLY=1, build the base CI image against PyTorch nightly so the
+# entire existing pipeline runs on nightly torch (CUDA/GPU lane only). Delegate
+# to the dedicated nightly build (PYTORCH_NIGHTLY=1, CUDA 13.0) and tag it at the
+# normal IMAGE_TAG that every test step already pulls -- no separate image tag,
+# no duplicate "vLLM Against PyTorch Nightly" pipeline section.
+if [[ "${TORCH_NIGHTLY:-0}" == "1" ]]; then
+    echo "--- :warning: TORCH_NIGHTLY=1 -- building base image on PyTorch nightly"
+    SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+    exec "${SCRIPT_DIR}/image_build_torch_nightly.sh" \
+        "${REGISTRY}" "${REPO}" "${BUILDKITE_COMMIT}" "${BRANCH}" "${IMAGE_TAG}"
+fi
+
 # build config
 TARGET="test-ci"
 VLLM_BAKE_FILE_PATH="${VLLM_BAKE_FILE_PATH:-docker/docker-bake.hcl}"
@@ -254,3 +272,5 @@ echo "--- :docker: Building ${TARGET}"
 docker --debug buildx bake -f "${VLLM_BAKE_FILE_PATH}" -f "${CI_HCL_PATH}" --progress plain "${TARGET}"
 
 echo "--- :white_check_mark: Build complete"
+
+annotate_image_tags

.buildkite/image_build/image_build_arm64.sh

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

.buildkite/image_build/image_build_cpu.sh

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

.buildkite/image_build/image_build_cpu_arm64.sh

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

.buildkite/image_build/image_build_hpu.sh

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

.buildkite/image_build/image_build_torch_nightly.sh

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

.buildkite/image_build/image_build_xpu.sh

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

.buildkite/intel_jobs/misc_intel.yaml

@@ -38,7 +38,7 @@ steps:
   agent_tags:
     label: production
     gpu: 1+
-    mem: 16+
+    mem: 24+
   no_plugin: true
   working_dir: "."
   env:
@@ -72,9 +72,31 @@ steps:
       pytest -v -s v1/test_oracle.py &&
       pytest -v -s v1/test_request.py &&
       pytest -v -s v1/test_outputs.py &&
-      pytest -v -s v1/sample/test_topk_topp_sampler.py &&
-      pytest -v -s v1/sample/test_logprobs.py &&
-      pytest -v -s v1/sample/test_logprobs_e2e.py'
+      pytest -v -s v1/sample'
+
+- label: Basic Models Tests (Initialization)
+  timeout_in_minutes: 60
+  device: intel_gpu
+  agent_tags:
+    label: production
+    gpu: 1+
+    mem: 16+
+  no_plugin: true
+  working_dir: "."
+  env:
+    REGISTRY: "public.ecr.aws/q9t5s3a7"
+    REPO: "vllm-ci-test-repo"
+    VLLM_TEST_DEVICE: "xpu"
+  source_file_dependencies:
+    - vllm/
+    - tests/models/test_initialization.py
+    - tests/models/registry.py
+  commands:
+    - >-
+      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+      'export VLLM_XPU_FUSED_MOE_USE_REF=1 &&
+      cd tests &&
+      pytest -v -s models/test_initialization.py::test_can_initialize_large_subset[Eagle3MiniMaxM2ForCausalLM]'
 
 - label: XPU CPU Offload
   timeout_in_minutes: 60
@@ -110,7 +132,7 @@ steps:
   agent_tags:
     label: production
     gpu: 2+
-    mem: 24+
+    mem: 16+
   no_plugin: true
   working_dir: "."
   env:

.buildkite/intel_jobs/models_distributed_intel.yaml

@@ -9,7 +9,7 @@ steps:
   agent_tags:
     label: production
     gpu: 2+
-    mem: 24+
+    mem: 16+
   no_plugin: true
   working_dir: "."
   env:

.buildkite/intel_jobs/models_multimodal_intel.yaml

@@ -9,7 +9,7 @@ steps:
   agent_tags:
     label: production
     gpu: 1+
-    mem: 16+
+    mem: 24+
   no_plugin: true
   working_dir: "."
   env:
@@ -81,7 +81,7 @@ steps:
   agent_tags:
     label: production
     gpu: 1+
-    mem: 16+
+    mem: 24+
   no_plugin: true
   working_dir: "."
   env:
@@ -105,7 +105,7 @@ steps:
   agent_tags:
     label: production
     gpu: 1+
-    mem: 16+
+    mem: 24+
   no_plugin: true
   working_dir: "."
   env:

.buildkite/intel_jobs/test-intel.yaml

@@ -49,6 +49,30 @@ steps:
         VLLM_XPU_FUSED_MOE_USE_REF=1 python3 examples/basic/offline_inference/generate.py --model Qwen/Qwen3-30B-A3B-Instruct-2507-FP8 --enforce-eager -tp 2 --max-model-len 8192 &&
         python3 examples/basic/offline_inference/generate.py --model INCModel/Qwen3-30B-A3B-Instruct-2507-MXFP4-LLMC --enforce-eager -tp 2 --max-model-len 8192
         '
+  - label: "XPU W8A8 FP8 Linear Examples"
+    depends_on:
+      - image-build-xpu
+    timeout_in_minutes: 60
+    device: intel_gpu
+    agent_tags:
+      label: production
+      gpu: 1+
+      mem: 24+
+    no_plugin: true
+    env:
+      REGISTRY: "public.ecr.aws/q9t5s3a7"
+      REPO: "vllm-ci-test-repo"
+      VLLM_TEST_DEVICE: "xpu"
+    source_file_dependencies:
+      - vllm/
+      - .buildkite/intel_jobs/test-intel.yaml
+    commands:
+      - >-
+        bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+        'python3 examples/basic/offline_inference/generate.py --linear-backend xpu --model RedHatAI/Meta-Llama-3.1-8B-Instruct-FP8 --enforce-eager --max-model-len 4096 &&
+        python3 examples/basic/offline_inference/generate.py --linear-backend xpu --model neuralmagic/Llama-3.2-1B-Instruct-FP8-dynamic --enforce-eager --max-model-len 4096 &&
+        python3 examples/basic/offline_inference/generate.py --linear-backend xpu --model meta-llama/Llama-3.2-1B-Instruct --quantization fp8 --enforce-eager --max-model-len 4096
+        '
   - label: "XPU V1 test"
     depends_on:
       - image-build-xpu
@@ -57,7 +81,7 @@ steps:
     agent_tags:
       label: production
       gpu: 1+
-      mem: 16+
+      mem: 24+
     no_plugin: true
     env:
       REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -120,4 +144,27 @@ steps:
       - >-
         bash .buildkite/scripts/hardware_ci/run-intel-test.sh
         'cd tests &&
-        pytest -v -s quantization/test_auto_round.py'
+        pytest -v -s quantization/test_auto_round.py'
+  - label: "XPU compressed tensors FP8 test"
+    depends_on:
+      - image-build-xpu
+    timeout_in_minutes: 60
+    device: intel_gpu
+    agent_tags:
+      label: production
+      gpu: 1+
+      mem: 16+
+    no_plugin: true
+    env:
+      REGISTRY: "public.ecr.aws/q9t5s3a7"
+      REPO: "vllm-ci-test-repo"
+      VLLM_TEST_DEVICE: "xpu"
+    source_file_dependencies:
+      - vllm/
+      - tests/quantization/test_compressed_tensors.py
+      - .buildkite/intel_jobs/test-intel.yaml
+    commands:
+      - >-
+        bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+        'cd tests &&
+        pytest -v -s quantization/test_compressed_tensors.py::test_compressed_tensors_fp8'

.buildkite/release-pipeline.yaml

@@ -350,6 +350,25 @@ steps:
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"
           - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"'
 
+      - label: ":docker: Build release image - x86_64 - XPU"
+        depends_on: ~
+        id: build-xpu-release-image
+        agents:
+          queue: cpu_queue_release
+        commands:
+          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
+          - |
+            DOCKER_BUILDKIT=1 docker build \
+              $(bash .buildkite/scripts/docker-build-metadata-args.sh xpu) \
+              --build-arg GIT_REPO_CHECK=1 \
+              --target vllm-openai \
+              --progress plain \
+              -f docker/Dockerfile.xpu .
+          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-xpu"
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-xpu"'
+        env:
+          DOCKER_BUILDKIT: "1"
+
       - block: "Build release image for x86_64 CPU"
         key: block-cpu-release-image-build
         depends_on: ~
@@ -445,6 +464,16 @@ steps:
           - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"
           - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 12.9 Ubuntu 24.04" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"'
 
+      - label: "Create manifest - XPU"
+        depends_on:
+          - build-xpu-release-image
+        id: create-manifest-xpu
+        agents:
+          queue: small_cpu_queue_release
+        commands:
+          - "bash .buildkite/scripts/xpu/create-xpu-ecr-manifest.sh"
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: XPU" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-xpu"'
+
       - label: "Publish nightly multi-arch image to DockerHub"
         depends_on:
           - create-multi-arch-manifest
@@ -819,6 +848,23 @@ steps:
       DOCKER_BUILDKIT: "1"
       S3_BUCKET: "vllm-wheels"
 
+  - label: "Publish nightly XPU image to DockerHub"
+    depends_on:
+      - create-manifest-xpu
+    if: build.env("NIGHTLY") == "1"
+    agents:
+      queue: small_cpu_queue_release
+    commands:
+      - "bash .buildkite/scripts/xpu/push-nightly-builds-xpu.sh"
+      - "bash .buildkite/scripts/cleanup-nightly-builds.sh nightly- vllm/vllm-openai-xpu"
+    plugins:
+      - docker-login#v3.0.0:
+          username: vllmbot
+          password-env: DOCKERHUB_TOKEN
+    env:
+      DOCKER_BUILDKIT: "1"
+      DOCKERHUB_USERNAME: "vllmbot"
+
   - label: "Publish nightly ROCm image to DockerHub"
     depends_on:
       - build-rocm-release-image
@@ -849,6 +895,7 @@ steps:
       - create-multi-arch-manifest-cuda-12-9
       - create-multi-arch-manifest-ubuntu2404
       - create-multi-arch-manifest-cuda-12-9-ubuntu2404
+      - create-manifest-xpu
       - build-rocm-release-image
       - input-release-version
       # Wait for CPU builds if their block steps were unblocked, so publish

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

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

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

@@ -29,7 +29,11 @@ if python3 -c "import torch; assert torch.version.hip" 2>/dev/null; then
         TORCH_INDEX_URL=""
     fi
 else
-    TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
+    if [ "${TORCH_NIGHTLY:-0}" = "1" ]; then
+        TORCH_INDEX_URL="https://download.pytorch.org/whl/nightly/cu130"
+    else
+        TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
+    fi
 fi
 echo ">>> Using PyTorch index: ${TORCH_INDEX_URL:-PyPI default}"
 

.buildkite/scripts/ci-bake-rocm.sh

@@ -15,7 +15,7 @@ set -euo pipefail
 
 DEFAULT_REPO_SLUG="vllm-project/vllm"
 DEFAULT_CI_HCL_SOURCE="docker/ci-rocm.hcl"
-DEFAULT_CI_BASE_CONTENT_FILES="requirements/common.txt requirements/rocm.txt requirements/test/rocm.txt docker/Dockerfile.rocm_base docker/ci-rocm.hcl docker/docker-bake-rocm.hcl tools/install_torchcodec_rocm.sh tests/vllm_test_utils .buildkite/scripts/ci-bake-rocm.sh"
+DEFAULT_CI_BASE_CONTENT_FILES="requirements/common.txt requirements/rocm.txt requirements/test/rocm.txt docker/Dockerfile.rocm_base docker/ci-rocm.hcl docker/docker-bake-rocm.hcl tools/install_torchcodec_rocm.sh tests/vllm_test_utils .buildkite/scripts/ci-bake-rocm.sh .buildkite/scripts/rocm/build-ci-base.sh"
 DEFAULT_CI_BASE_DOCKERFILE="docker/Dockerfile.rocm"
 DEFAULT_CI_BASE_DOCKERFILE_STAGES="base build_rixl build_rocshmem build_deepep mori_base ci_base"
 DEFAULT_CI_BASE_METADATA_VERSION="1"
@@ -393,6 +393,16 @@ should_upload_wheel_artifacts() {
         || "${TARGET}" == *"artifact"* ]]
 }
 
+set_buildkite_metadata() {
+    local key="$1"
+    local value="$2"
+
+    [[ -n "${value}" ]] || return 0
+    if command -v buildkite-agent >/dev/null 2>&1; then
+        buildkite-agent meta-data set "${key}" "${value}" || true
+    fi
+}
+
 get_remote_image_label() {
     local image_ref="$1"
     local label_key="$2"
@@ -733,6 +743,8 @@ configure_ci_base_image_refs() {
 
     if is_ci_base_target; then
         IMAGE_TAG="${primary_tag}"
+        CI_BASE_IMAGE="${primary_tag}"
+        export CI_BASE_IMAGE
         export IMAGE_TAG
 
         echo "ci_base primary image tag: ${CI_BASE_IMAGE_TAG}"
@@ -750,6 +762,10 @@ configure_ci_base_image_refs() {
             echo "ci_base stable alias will not be pushed for this build"
             echo "Set NIGHTLY=1 on ${CI_BASE_STABLE_BRANCH:-main} to refresh ${stable_tag}"
         fi
+        set_buildkite_metadata "rocm-ci-base-image" "${CI_BASE_IMAGE_TAG}"
+        set_buildkite_metadata "rocm-ci-base-image-content" "${content_tag}"
+        set_buildkite_metadata "rocm-ci-base-image-commit" "${CI_BASE_IMAGE_TAG_COMMIT:-}"
+        set_buildkite_metadata "rocm-ci-base-image-stable" "${CI_BASE_IMAGE_TAG_STABLE:-}"
         return 0
     fi
 
@@ -1779,7 +1795,10 @@ seed_dependency_caches_if_needed() {
 
         echo "--- :docker: Seeding ${target}"
         echo "Expected cache ref: ${cache_ref}"
-        docker buildx bake "${BAKE_FILES[@]}" --progress plain "${target}"
+        docker buildx bake \
+            "${BAKE_FILES[@]}" \
+            --progress "${BUILDKIT_PROGRESS:-plain}" \
+            "${target}"
         verify_dependency_cache_ref "${cache_ref}"
     done
 }
@@ -1807,7 +1826,10 @@ run_bake() {
     local build_rc=0
 
     echo "--- :docker: Building ${TARGET}"
-    docker buildx bake "${BAKE_FILES[@]}" --progress plain "${BAKE_TARGETS[@]}" || build_rc=$?
+    docker buildx bake \
+        "${BAKE_FILES[@]}" \
+        --progress "${BUILDKIT_PROGRESS:-plain}" \
+        "${BAKE_TARGETS[@]}" || build_rc=$?
 
     if [[ ${build_rc} -eq 0 ]]; then
         echo "--- :white_check_mark: Build complete"

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

@@ -28,6 +28,17 @@
 ###############################################################################
 set -o pipefail
 
+: "${BUILDKIT_PROGRESS:=plain}"
+: "${TERM:=xterm-256color}"
+: "${FORCE_COLOR:=1}"
+: "${CLICOLOR_FORCE:=1}"
+: "${PY_COLORS:=1}"
+: "${ROCM_DOCKER_TTY:=1}"
+if [[ " ${PYTEST_ADDOPTS:-} " != *" --color"* ]]; then
+  PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }--color=yes"
+fi
+export BUILDKIT_PROGRESS TERM FORCE_COLOR CLICOLOR_FORCE PY_COLORS PYTEST_ADDOPTS ROCM_DOCKER_TTY
+
 # Export Python path for commands that run directly on the host. Containerized
 # tests set this to /vllm-workspace below so spawned Python processes do not
 # depend on their current working directory.
@@ -149,6 +160,7 @@ EOF
   echo "--- Building local ROCm test image"
   docker build \
     --pull=false \
+    --progress "${BUILDKIT_PROGRESS}" \
     --build-arg "BASE_IMAGE=${base_image}" \
     -t "${artifact_image}" \
     "${context_dir}" || return 1
@@ -448,9 +460,11 @@ checkout="${BUILDKITE_BUILD_CHECKOUT_PATH:-}"
 if [[ -z "${checkout}" || ! -d "${checkout}" ]]; then
   checkout="."
 fi
-if git -C "${checkout}" rev-parse --is-inside-work-tree >/dev/null 2>&1; then
+# Pass safe.directory per-command (-c) because buildkite runs will always fail
+# the next check on git 2.35.2+ due to mixed uses of root and buildkite-agent/uids.
+if git -c "safe.directory=${checkout}" -C "${checkout}" rev-parse --is-inside-work-tree >/dev/null 2>&1; then
   vllm_standalone_merge_base="$(
-    git -C "${checkout}" merge-base HEAD origin/main 2>/dev/null || true
+    git -c "safe.directory=${checkout}" -C "${checkout}" merge-base HEAD origin/main 2>/dev/null || true
   )"
 fi
 if [[ -z "${vllm_standalone_merge_base}" ]]; then
@@ -533,6 +547,13 @@ if is_multi_node "$commands"; then
 else
   echo "--- Single-node job"
   echo "Render devices: $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES"
+  docker_run_terminal_args=(-i)
+  if [[ "${ROCM_DOCKER_TTY}" == "1" ]]; then
+    docker_run_terminal_args+=(-t)
+    echo "Docker interactive stdin: enabled; TTY allocation: enabled"
+  else
+    echo "Docker interactive stdin: enabled; TTY allocation: disabled"
+  fi
 
   ulimit_core_hard=$(ulimit -H -c)
   if [[ "$ulimit_core_hard" == "unlimited" ]]; then
@@ -549,6 +570,7 @@ else
   fi
 
   docker run \
+    "${docker_run_terminal_args[@]}" \
     --device /dev/kfd $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES \
     $RDMA_FLAGS \
     --network=host \
@@ -563,6 +585,11 @@ else
     -e AWS_SECRET_ACCESS_KEY \
     -e BUILDKITE_PARALLEL_JOB \
     -e BUILDKITE_PARALLEL_JOB_COUNT \
+    -e TERM \
+    -e FORCE_COLOR \
+    -e CLICOLOR_FORCE \
+    -e PY_COLORS \
+    -e PYTEST_ADDOPTS \
     -v "${HF_CACHE}:${HF_MOUNT}" \
     -e "HF_HOME=${HF_MOUNT}" \
     -e "PYTHONPATH=${MYPYTHONPATH}" \

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

@@ -38,7 +38,9 @@ function cpu_tests() {
     pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
     pytest -x -v -s tests/kernels/core/test_cpu_activation.py
     pytest -x -v -s tests/kernels/moe/test_cpu_fused_moe.py
-    pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py"
+    pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
+    pytest -x -v -s tests/kernels/moe/test_cpu_int4_moe.py
+    pytest -x -v -s tests/kernels/mamba/test_cpu_short_conv.py"
 
   # skip tests requiring model downloads if HF_TOKEN is not set
   # due to rate-limits
@@ -62,7 +64,6 @@ function cpu_tests() {
     set -e
     pytest -x -v -s tests/quantization/test_compressed_tensors.py::test_compressed_tensors_w8a8_logprobs"
 
-
   # basic online serving
   docker exec cpu-test bash -c '
     set -e

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

@@ -130,6 +130,22 @@ docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm
 docker push vllm/vllm-openai-rocm:latest-base
 docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
 
+# ---- XPU ----
+
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu vllm/vllm-openai-xpu:latest-x86_64
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64
+docker push vllm/vllm-openai-xpu:latest-x86_64
+docker push vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64
+
+docker manifest rm vllm/vllm-openai-xpu:latest || true
+docker manifest rm vllm/vllm-openai-xpu:v${RELEASE_VERSION} || true
+docker manifest create vllm/vllm-openai-xpu:latest vllm/vllm-openai-xpu:latest-x86_64 --amend
+docker manifest create vllm/vllm-openai-xpu:v${RELEASE_VERSION} vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64 --amend
+docker manifest push vllm/vllm-openai-xpu:latest
+docker manifest push vllm/vllm-openai-xpu:v${RELEASE_VERSION}
+
 # ---- CPU ----
 # CPU images are behind separate block steps and may not have been built.
 # All-or-nothing: inspect both arches first, then either publish everything

.buildkite/scripts/rocm/build-ci-base.sh

@@ -0,0 +1,32 @@
+#!/usr/bin/env bash
+# Build the ROCm ci_base image, optionally from a freshly rebuilt ROCm base.
+
+set -euo pipefail
+
+metadata_get() {
+    local key="$1"
+    if command -v buildkite-agent >/dev/null 2>&1; then
+        buildkite-agent meta-data get "${key}" 2>/dev/null || true
+    fi
+}
+
+main() {
+    local base_refreshed=""
+
+    base_refreshed="$(metadata_get rocm-base-refresh)"
+    if [[ "${base_refreshed}" == "1" ]]; then
+        export BASE_IMAGE
+        export CI_BASE_PUSH_STABLE_TAG
+
+        BASE_IMAGE="$(metadata_get rocm-base-image)"
+        CI_BASE_PUSH_STABLE_TAG="$(metadata_get rocm-base-push-stable-tag)"
+        CI_BASE_PUSH_STABLE_TAG="${CI_BASE_PUSH_STABLE_TAG:-0}"
+
+        echo "Using refreshed ROCm base image for ci_base: ${BASE_IMAGE}"
+        echo "Push stable ci_base tag: ${CI_BASE_PUSH_STABLE_TAG}"
+    fi
+
+    bash .buildkite/scripts/ci-bake-rocm.sh ci-base-rocm-ci-with-deps
+}
+
+main "$@"

.buildkite/scripts/rocm/build-test-image.sh

@@ -0,0 +1,57 @@
+#!/usr/bin/env bash
+# Build the ROCm CI test image or wheel artifact.
+#
+# When Dockerfile.rocm_base changes, always build the full image so downstream
+# ROCm tests can validate the freshly rebuilt base -> ci_base -> ci image chain.
+
+set -euo pipefail
+
+metadata_get() {
+    local key="$1"
+    if command -v buildkite-agent >/dev/null 2>&1; then
+        buildkite-agent meta-data get "${key}" 2>/dev/null || true
+    fi
+}
+
+use_refreshed_base_if_present() {
+    local base_refreshed=""
+
+    base_refreshed="$(metadata_get rocm-base-refresh)"
+    if [[ "${base_refreshed}" != "1" ]]; then
+        return 1
+    fi
+
+    export BASE_IMAGE
+    export CI_BASE_IMAGE
+    export IMAGE_TAG_LATEST
+
+    BASE_IMAGE="$(metadata_get rocm-base-image)"
+    CI_BASE_IMAGE="$(metadata_get rocm-ci-base-image)"
+    IMAGE_TAG_LATEST="$(metadata_get rocm-ci-image-descriptive)"
+
+    echo "Using refreshed ROCm base image for test image: ${BASE_IMAGE}"
+    echo "Using refreshed ROCm ci_base image for test image: ${CI_BASE_IMAGE}"
+    if [[ -n "${IMAGE_TAG_LATEST}" ]]; then
+        echo "Also tagging full ROCm CI image as: ${IMAGE_TAG_LATEST}"
+    fi
+
+    return 0
+}
+
+main() {
+    local base_refreshed=0
+
+    if use_refreshed_base_if_present; then
+        base_refreshed=1
+    fi
+
+    if [[ "${ROCM_CI_ARTIFACT_ONLY:-0}" == "1" && "${base_refreshed}" != "1" ]]; then
+        echo "ROCM_CI_ARTIFACT_ONLY=1; building ROCm wheel artifact only"
+        IMAGE_TAG="" bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-artifacts
+        return
+    fi
+
+    bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-wheel
+}
+
+main "$@"

.buildkite/scripts/rocm/refresh-base-image.sh

@@ -0,0 +1,513 @@
+#!/usr/bin/env bash
+# Build and publish a fresh ROCm base image when Dockerfile.rocm_base changes.
+#
+# Normal AMD CI builds should not pay for this path. The script no-ops unless
+# docker/Dockerfile.rocm_base changed relative to the branch base, the previous
+# main commit, or ROCM_BASE_REFRESH_FORCE=1 is set.
+
+set -euo pipefail
+
+DOCKERFILE="${ROCM_BASE_DOCKERFILE:-docker/Dockerfile.rocm_base}"
+BASE_REPO="${ROCM_BASE_IMAGE_REPO:-rocm/vllm-dev}"
+CI_IMAGE_REPO="${ROCM_CI_IMAGE_REPO:-rocm/vllm-ci}"
+BUILDER_NAME="${ROCM_BASE_BUILDER_NAME:-vllm-rocm-base-builder}"
+DEFAULT_ROCM_BASE_METADATA_VERSION="1"
+DEFAULT_ROCM_BASE_CONTENT_FILES="${DOCKERFILE}"
+DEFAULT_ROCM_BASE_CONTENT_ARGS="BASE_IMAGE TRITON_BRANCH TRITON_REPO PYTORCH_BRANCH PYTORCH_REPO PYTORCH_VISION_BRANCH PYTORCH_VISION_REPO PYTORCH_AUDIO_BRANCH PYTORCH_AUDIO_REPO FA_BRANCH FA_REPO AITER_BRANCH AITER_REPO MORI_BRANCH MORI_REPO PYTORCH_ROCM_ARCH PYTHON_VERSION USE_SCCACHE"
+
+metadata_set() {
+    local key="$1"
+    local value="$2"
+
+    [[ -n "${value}" ]] || return 0
+    if command -v buildkite-agent >/dev/null 2>&1; then
+        buildkite-agent meta-data set "${key}" "${value}" || true
+    fi
+}
+
+compute_content_hash() {
+    local path=""
+    local file=""
+
+    for path in "$@"; do
+        if [[ -d "${path}" ]]; then
+            while IFS= read -r -d '' file; do
+                printf 'file:%s\n' "${file}"
+                sha256sum "${file}"
+            done < <(find "${path}" -type f -print0 | sort -z)
+        elif [[ -f "${path}" ]]; then
+            printf 'file:%s\n' "${path}"
+            sha256sum "${path}"
+        else
+            printf 'missing:%s\n' "${path}"
+        fi
+    done | sha256sum | cut -d' ' -f1
+}
+
+clean_docker_tag() {
+    local input="$1"
+    echo "${input}" | sed 's/[^a-zA-Z0-9._-]/_/g' | cut -c1-128
+}
+
+tag_component() {
+    local input="$1"
+    local max_chars="${2:-24}"
+
+    clean_docker_tag "${input:-unknown}" | cut -c1-"${max_chars}"
+}
+
+extract_arg_default() {
+    local arg_name="$1"
+
+    sed -n -E "s/^[[:space:]]*ARG[[:space:]]+${arg_name}=\"?([^\"[:space:]]+)\"?.*/\\1/p" \
+        "${DOCKERFILE}" | head -1
+}
+
+resolve_image_digest() {
+    local image_ref="$1"
+
+    docker buildx imagetools inspect "${image_ref}" 2>/dev/null \
+        | sed -n -E 's/^Digest:[[:space:]]+//p' \
+        | head -1 || true
+}
+
+resolve_rocm_base_arg_value() {
+    local arg_name="$1"
+    local use_sccache="$2"
+
+    case "${arg_name}" in
+        USE_SCCACHE)
+            printf '%s\n' "${use_sccache}"
+            ;;
+        *)
+            extract_arg_default "${arg_name}"
+            ;;
+    esac
+}
+
+hash_rocm_base_arg_values() {
+    local use_sccache="$1"
+    local base_image_digest="$2"
+    local arg_name=""
+    local arg_value=""
+    shift 2 || true
+
+    for arg_name in "$@"; do
+        [[ -n "${arg_name}" ]] || continue
+        arg_value=$(resolve_rocm_base_arg_value "${arg_name}" "${use_sccache}")
+        printf 'arg:%s=%s\n' "${arg_name}" "${arg_value:-<empty>}"
+        if [[ "${arg_name}" == "BASE_IMAGE" && -n "${arg_value}" ]]; then
+            printf 'arg:%s.digest=%s\n' "${arg_name}" "${base_image_digest:-unknown}"
+        fi
+    done
+}
+
+rocm_version_from_base_image() {
+    local base_image="$1"
+    local version=""
+
+    version="$(sed -n -E 's/.*:([0-9]+\.[0-9]+(\.[0-9]+)?)-.*/\1/p' <<<"${base_image}")"
+    tag_component "${version:-${base_image}}" 16
+}
+
+git_diff_changed_base() {
+    local range="$1"
+    [[ -n "$(git diff --name-only "${range}" -- "${DOCKERFILE}" 2>/dev/null)" ]]
+}
+
+short_git_ref() {
+    local ref="$1"
+
+    git rev-parse --short "${ref}" 2>/dev/null || printf '%s\n' "${ref}"
+}
+
+extract_arg_default_from_ref() {
+    local ref="$1"
+    local arg_name="$2"
+    local content=""
+
+    content="$(git show "${ref}:${DOCKERFILE}" 2>/dev/null || true)"
+    sed -n -E "s/^[[:space:]]*ARG[[:space:]]+${arg_name}=\"?([^\"[:space:]]+)\"?.*/\\1/p" \
+        <<<"${content}" | head -1
+}
+
+log_arg_default_changes() {
+    local old_ref="$1"
+    local new_ref="$2"
+    local content_args="${ROCM_BASE_CONTENT_ARGS:-${DEFAULT_ROCM_BASE_CONTENT_ARGS}}"
+    local arg_name=""
+    local old_value=""
+    local new_value=""
+    local changed=0
+
+    echo "Changed ROCm base ARG defaults:"
+    for arg_name in ${content_args}; do
+        old_value="$(extract_arg_default_from_ref "${old_ref}" "${arg_name}")"
+        new_value="$(extract_arg_default_from_ref "${new_ref}" "${arg_name}")"
+        if [[ "${old_value}" != "${new_value}" ]]; then
+            echo "  - ${arg_name}: ${old_value:-<unset>} -> ${new_value:-<unset>}"
+            changed=1
+        fi
+    done
+
+    if [[ "${changed}" == "0" ]]; then
+        echo "  - none detected; Dockerfile instructions changed outside tracked ARG defaults"
+    fi
+}
+
+log_arg_line_diff() {
+    local range="$1"
+    local arg_diff=""
+
+    arg_diff="$(
+        git diff --unified=0 "${range}" -- "${DOCKERFILE}" 2>/dev/null \
+            | awk '/^[+-][[:space:]]*ARG[[:space:]]/ && $0 !~ /^(---|\+\+\+)/ { print "  " $0 }' \
+            || true
+    )"
+
+    if [[ -n "${arg_diff}" ]]; then
+        echo "Changed Dockerfile ARG lines:"
+        printf '%s\n' "${arg_diff}"
+    fi
+}
+
+log_rocm_base_change_check() {
+    local context="$1"
+    local range="$2"
+    local old_ref="$3"
+    local old_short=""
+    local head_short=""
+
+    old_short="$(short_git_ref "${old_ref}")"
+    head_short="$(short_git_ref HEAD)"
+
+    echo "--- :mag: ROCm base refresh check"
+    echo "Context: ${context}"
+    echo "Dockerfile: ${DOCKERFILE}"
+    echo "Base revision: ${old_short}"
+    echo "Head revision: ${head_short}"
+    echo "Git diff range: ${range}"
+}
+
+log_rocm_base_rebuild_reason() {
+    local context="$1"
+    local range="$2"
+    local old_ref="$3"
+    local changed_files=""
+
+    log_rocm_base_change_check "${context}" "${range}" "${old_ref}"
+
+    changed_files="$(git diff --name-only "${range}" -- "${DOCKERFILE}" 2>/dev/null || true)"
+    echo "Changed files:"
+    if [[ -n "${changed_files}" ]]; then
+        sed 's/^/  - /' <<<"${changed_files}"
+    else
+        echo "  - ${DOCKERFILE}"
+    fi
+    log_arg_default_changes "${old_ref}" HEAD
+    log_arg_line_diff "${range}"
+    echo "Decision: rebuilding ROCm base image because ${DOCKERFILE} changed."
+}
+
+rocm_base_changed_in_range() {
+    local context="$1"
+    local range="$2"
+    local old_ref="$3"
+
+    if git_diff_changed_base "${range}"; then
+        log_rocm_base_rebuild_reason "${context}" "${range}" "${old_ref}"
+        return 0
+    fi
+
+    log_rocm_base_change_check "${context}" "${range}" "${old_ref}"
+    echo "Decision: ROCm base refresh not required; ${DOCKERFILE} is unchanged."
+    return 1
+}
+
+rocm_base_changed() {
+    local base_branch="${BUILDKITE_PULL_REQUEST_BASE_BRANCH:-main}"
+    local base_ref="refs/remotes/origin/${base_branch}"
+    local merge_base=""
+
+    if [[ "${ROCM_BASE_REFRESH_SKIP:-0}" == "1" ]]; then
+        echo "ROCM_BASE_REFRESH_SKIP=1 set; skipping ROCm base refresh"
+        return 1
+    fi
+
+    if [[ "${ROCM_BASE_REFRESH_FORCE:-0}" == "1" ]]; then
+        echo "ROCM_BASE_REFRESH_FORCE=1 set; refreshing ROCm base image"
+        return 0
+    fi
+
+    if ! git rev-parse --is-inside-work-tree >/dev/null 2>&1; then
+        echo "Not in a git checkout; skipping ROCm base refresh unless forced"
+        return 1
+    fi
+
+    if [[ "${BUILDKITE_PULL_REQUEST:-false}" != "false" ]]; then
+        git fetch --no-tags --depth=200 origin \
+            "+refs/heads/${base_branch}:${base_ref}" >/dev/null 2>&1 || true
+        merge_base=$(git merge-base HEAD "${base_ref}" 2>/dev/null || true)
+        if [[ -z "${merge_base}" ]]; then
+            echo "Unable to determine merge base with PR base ${base_ref}; skipping ROCm base refresh unless forced"
+            return 1
+        fi
+        if rocm_base_changed_in_range \
+            "pull request build against ${base_ref}" \
+            "${merge_base}...HEAD" \
+            "${merge_base}"; then
+            return 0
+        fi
+    elif [[ "${BUILDKITE_BRANCH:-}" == "${ROCM_BASE_STABLE_BRANCH:-main}" ]] \
+        && git rev-parse --verify HEAD~1 >/dev/null 2>&1; then
+        if rocm_base_changed_in_range \
+            "stable branch build; comparing against previous ${ROCM_BASE_STABLE_BRANCH:-main} commit" \
+            "HEAD~1..HEAD" \
+            "HEAD~1"; then
+            return 0
+        fi
+    else
+        git fetch --no-tags --depth=200 origin \
+            "+refs/heads/${base_branch}:${base_ref}" >/dev/null 2>&1 || true
+        merge_base=$(git merge-base HEAD "${base_ref}" 2>/dev/null || true)
+        if [[ -z "${merge_base}" ]]; then
+            echo "Unable to determine merge base with branch base ${base_ref}; skipping ROCm base refresh unless forced"
+            return 1
+        fi
+        if rocm_base_changed_in_range \
+            "branch build against ${base_ref}" \
+            "${merge_base}...HEAD" \
+            "${merge_base}"; then
+            return 0
+        fi
+    fi
+
+    return 1
+}
+
+should_push_stable_tag() {
+    if [[ "${BUILDKITE_PULL_REQUEST:-false}" != "false" ]]; then
+        return 1
+    fi
+
+    if [[ "${ROCM_BASE_PUSH_STABLE_TAG:-}" == "1" ]]; then
+        return 0
+    fi
+    if [[ "${ROCM_BASE_PUSH_STABLE_TAG:-}" == "0" ]]; then
+        return 1
+    fi
+
+    [[ "${BUILDKITE_PULL_REQUEST:-false}" == "false" \
+        && "${BUILDKITE_BRANCH:-}" == "${ROCM_BASE_STABLE_BRANCH:-main}" ]]
+}
+
+setup_builder() {
+    echo "--- :buildkite: Setting up buildx builder for ROCm base"
+    if docker buildx inspect "${BUILDER_NAME}" >/dev/null 2>&1; then
+        docker buildx use "${BUILDER_NAME}"
+    else
+        docker buildx create --name "${BUILDER_NAME}" --driver docker-container --use
+    fi
+    docker buildx inspect --bootstrap
+}
+
+compute_base_content_hash() {
+    local use_sccache="$1"
+    local base_image_digest="$2"
+    local content_files="${ROCM_BASE_CONTENT_FILES:-${DEFAULT_ROCM_BASE_CONTENT_FILES}}"
+    local content_args="${ROCM_BASE_CONTENT_ARGS:-${DEFAULT_ROCM_BASE_CONTENT_ARGS}}"
+    local -a content_paths=()
+    local -a content_arg_names=()
+
+    read -r -a content_paths <<< "${content_files}"
+    read -r -a content_arg_names <<< "${content_args}"
+
+    {
+        printf 'content-files-hash:%s\n' "$(compute_content_hash "${content_paths[@]}")"
+        printf 'dockerfile:%s\n' "${DOCKERFILE}"
+        printf 'resolved-build-args:\n'
+        hash_rocm_base_arg_values \
+            "${use_sccache}" "${base_image_digest}" "${content_arg_names[@]}"
+    } | sha256sum | cut -d' ' -f1
+}
+
+build_base_image() {
+    local use_sccache="${ROCM_BASE_USE_SCCACHE:-${USE_SCCACHE:-0}}"
+    local base_hash=""
+    local build_date=""
+    local build_suffix=""
+    local base_image_arg=""
+    local base_image_digest=""
+    local rocm_version=""
+    local triton_arg=""
+    local pytorch_arg=""
+    local pytorch_vision_arg=""
+    local pytorch_audio_arg=""
+    local fa_arg=""
+    local aiter_arg=""
+    local mori_arg=""
+    local python_version_arg=""
+    local pytorch_rocm_arch_arg=""
+    local pytorch_branch=""
+    local aiter_branch=""
+    local dependency_summary=""
+    local descriptor=""
+    local ci_descriptor=""
+    local descriptive_tag=""
+    local stable_tag="${BASE_REPO}:base"
+    local ci_descriptive_tag=""
+    local content_files="${ROCM_BASE_CONTENT_FILES:-${DEFAULT_ROCM_BASE_CONTENT_FILES}}"
+    local content_args="${ROCM_BASE_CONTENT_ARGS:-${DEFAULT_ROCM_BASE_CONTENT_ARGS}}"
+    local content_files_hash=""
+    local metadata_version="${ROCM_BASE_METADATA_VERSION:-${DEFAULT_ROCM_BASE_METADATA_VERSION}}"
+    local -a tags=()
+    local -a no_cache_args=()
+    local -a sccache_args=()
+    local -a content_paths=()
+
+    if [[ ! -f "${DOCKERFILE}" ]]; then
+        echo "Error: ROCm base Dockerfile not found: ${DOCKERFILE}" >&2
+        exit 1
+    fi
+
+    build_date="${ROCM_BASE_TAG_DATE:-$(date -u +%Y%m%d)}"
+    if [[ -n "${BUILDKITE_BUILD_NUMBER:-}" ]]; then
+        build_suffix="_bk_${BUILDKITE_BUILD_NUMBER}"
+    fi
+    base_image_arg="$(extract_arg_default BASE_IMAGE)"
+    base_image_digest="$(resolve_image_digest "${base_image_arg}")"
+    read -r -a content_paths <<< "${content_files}"
+    content_files_hash="$(compute_content_hash "${content_paths[@]}")"
+    base_hash=$(compute_base_content_hash "${use_sccache}" "${base_image_digest}")
+    rocm_version="$(rocm_version_from_base_image "${base_image_arg}")"
+    triton_arg="$(extract_arg_default TRITON_BRANCH)"
+    pytorch_arg="$(extract_arg_default PYTORCH_BRANCH)"
+    pytorch_vision_arg="$(extract_arg_default PYTORCH_VISION_BRANCH)"
+    pytorch_audio_arg="$(extract_arg_default PYTORCH_AUDIO_BRANCH)"
+    fa_arg="$(extract_arg_default FA_BRANCH)"
+    aiter_arg="$(extract_arg_default AITER_BRANCH)"
+    mori_arg="$(extract_arg_default MORI_BRANCH)"
+    python_version_arg="$(extract_arg_default PYTHON_VERSION)"
+    pytorch_rocm_arch_arg="$(extract_arg_default PYTORCH_ROCM_ARCH)"
+    pytorch_branch="$(tag_component "${pytorch_arg}" 16)"
+    aiter_branch="$(tag_component "${aiter_arg}" 24)"
+    dependency_summary="base=${base_image_arg},rocm=${rocm_version},python=${python_version_arg},pytorch=${pytorch_arg},torchvision=${pytorch_vision_arg},torchaudio=${pytorch_audio_arg},triton=${triton_arg},flash-attn=${fa_arg},aiter=${aiter_arg},mori=${mori_arg},pytorch-rocm-arch=${pytorch_rocm_arch_arg}"
+    descriptor="$(clean_docker_tag "base_custom_aiter_${aiter_branch}_torch_${pytorch_branch}_${build_date}${build_suffix}")"
+    ci_descriptor="$(clean_docker_tag "ci_custom_aiter_${aiter_branch}_torch_${pytorch_branch}_${build_date}${build_suffix}")"
+
+    descriptive_tag="${BASE_REPO}:${descriptor}"
+    ci_descriptive_tag="${CI_IMAGE_REPO}:${ci_descriptor}"
+
+    tags=(-t "${descriptive_tag}")
+    if should_push_stable_tag; then
+        tags+=(-t "${stable_tag}")
+        metadata_set "rocm-base-push-stable-tag" "1"
+    else
+        metadata_set "rocm-base-push-stable-tag" "0"
+    fi
+
+    if [[ "${ROCM_BASE_NO_CACHE:-1}" == "1" ]]; then
+        no_cache_args=(--no-cache)
+    fi
+
+    for env_name in \
+        SCCACHE_DOWNLOAD_URL \
+        SCCACHE_ENDPOINT \
+        SCCACHE_BUCKET_NAME \
+        SCCACHE_REGION_NAME \
+        SCCACHE_S3_NO_CREDENTIALS; do
+        if [[ -n "${!env_name:-}" ]]; then
+            sccache_args+=(--build-arg "${env_name}=${!env_name}")
+        fi
+    done
+
+    echo "--- :docker: Building ROCm base image"
+    echo "Dockerfile: ${DOCKERFILE}"
+    echo "Descriptive tag: ${descriptive_tag}"
+    echo "Stable tag: ${stable_tag} ($(should_push_stable_tag && echo enabled || echo disabled))"
+    echo "Content hash: ${base_hash}"
+    echo "Dependency summary: ${dependency_summary}"
+    echo "USE_SCCACHE: ${use_sccache}"
+
+    docker buildx build \
+        "${no_cache_args[@]}" \
+        --pull \
+        --progress "${BUILDKIT_PROGRESS:-plain}" \
+        --file "${DOCKERFILE}" \
+        --build-arg "USE_SCCACHE=${use_sccache}" \
+        "${sccache_args[@]}" \
+        --label "org.opencontainers.image.source=https://github.com/vllm-project/vllm" \
+        --label "org.opencontainers.image.vendor=vLLM" \
+        --label "org.opencontainers.image.title=vLLM ROCm base" \
+        --label "org.opencontainers.image.revision=${BUILDKITE_COMMIT:-}" \
+        --label "vllm.rocm_base.metadata_version=${metadata_version}" \
+        --label "vllm.rocm_base.content_hash=${base_hash}" \
+        --label "vllm.rocm_base.content_files_hash=${content_files_hash}" \
+        --label "vllm.rocm_base.dockerfile=${DOCKERFILE}" \
+        --label "vllm.rocm_base.image.descriptive=${descriptive_tag}" \
+        --label "vllm.rocm_base.image.stable=${stable_tag}" \
+        --label "vllm.rocm_base.git_commit=${BUILDKITE_COMMIT:-}" \
+        --label "vllm.rocm_base.stable_branch=${ROCM_BASE_STABLE_BRANCH:-main}" \
+        --label "vllm.rocm_base.descriptor=${descriptor}" \
+        --label "vllm.rocm_base.dependency_summary=${dependency_summary}" \
+        --label "vllm.rocm_base.base_image=${base_image_arg}" \
+        --label "vllm.rocm_base.base_image_digest=${base_image_digest}" \
+        --label "vllm.rocm_base.dependency.rocm=${rocm_version}" \
+        --label "vllm.rocm_base.dependency.python=${python_version_arg}" \
+        --label "vllm.rocm_base.dependency.pytorch=${pytorch_arg}" \
+        --label "vllm.rocm_base.dependency.torchvision=${pytorch_vision_arg}" \
+        --label "vllm.rocm_base.dependency.torchaudio=${pytorch_audio_arg}" \
+        --label "vllm.rocm_base.dependency.triton=${triton_arg}" \
+        --label "vllm.rocm_base.dependency.flash_attention=${fa_arg}" \
+        --label "vllm.rocm_base.dependency.aiter=${aiter_arg}" \
+        --label "vllm.rocm_base.dependency.mori=${mori_arg}" \
+        --label "vllm.rocm_base.pytorch_rocm_arch=${pytorch_rocm_arch_arg}" \
+        "${tags[@]}" \
+        --push \
+        .
+
+    docker buildx imagetools inspect "${descriptive_tag}" >/dev/null
+
+    metadata_set "rocm-base-refresh" "1"
+    metadata_set "rocm-base-image" "${descriptive_tag}"
+    metadata_set "rocm-base-image-descriptive" "${descriptive_tag}"
+    metadata_set "rocm-base-image-stable" "${stable_tag}"
+    metadata_set "rocm-base-image-ci-descriptive" "${ci_descriptive_tag}"
+    metadata_set "rocm-base-metadata-version" "${metadata_version}"
+    metadata_set "rocm-base-content-hash" "${base_hash}"
+    metadata_set "rocm-base-content-files-hash" "${content_files_hash}"
+    metadata_set "rocm-base-content-files" "${content_files}"
+    metadata_set "rocm-base-content-args" "${content_args}"
+    metadata_set "rocm-base-base-image-digest" "${base_image_digest}"
+    metadata_set "rocm-base-dockerfile" "${DOCKERFILE}"
+    metadata_set "rocm-base-descriptor" "${descriptor}"
+    metadata_set "rocm-base-dependency-summary" "${dependency_summary}"
+    metadata_set "rocm-base-dependency-rocm" "${rocm_version}"
+    metadata_set "rocm-base-dependency-python" "${python_version_arg}"
+    metadata_set "rocm-base-dependency-pytorch" "${pytorch_arg}"
+    metadata_set "rocm-base-dependency-torchvision" "${pytorch_vision_arg}"
+    metadata_set "rocm-base-dependency-torchaudio" "${pytorch_audio_arg}"
+    metadata_set "rocm-base-dependency-triton" "${triton_arg}"
+    metadata_set "rocm-base-dependency-flash-attention" "${fa_arg}"
+    metadata_set "rocm-base-dependency-aiter" "${aiter_arg}"
+    metadata_set "rocm-base-dependency-mori" "${mori_arg}"
+    metadata_set "rocm-base-pytorch-rocm-arch" "${pytorch_rocm_arch_arg}"
+    metadata_set "rocm-ci-image-descriptive" "${ci_descriptive_tag}"
+
+    echo "--- :white_check_mark: ROCm base image published"
+    echo "Use BASE_IMAGE=${descriptive_tag} for downstream ROCm CI builds"
+}
+
+main() {
+    metadata_set "rocm-base-refresh" "0"
+
+    if ! rocm_base_changed; then
+        echo "ROCm base Dockerfile did not change; skipping base image refresh"
+        return 0
+    fi
+
+    setup_builder
+    build_base_image
+}
+
+main "$@"

.buildkite/scripts/rocm/smoke-test-image.sh

@@ -0,0 +1,32 @@
+#!/usr/bin/env bash
+# Fast structural smoke test for the full ROCm CI image.
+
+set -euo pipefail
+
+image_ref="${VLLM_CI_SMOKE_IMAGE:-rocm/vllm-ci:${BUILDKITE_COMMIT:?BUILDKITE_COMMIT is required}}"
+
+docker run --rm --network=none --entrypoint /bin/bash "${image_ref}" -ec '
+  if [ ! -d /vllm-workspace ]; then echo Missing directory: /vllm-workspace >&2; exit 1; fi
+  if [ ! -d /vllm-workspace/tests ]; then echo Missing directory: /vllm-workspace/tests >&2; exit 1; fi
+  if [ ! -d /vllm-workspace/src/vllm ]; then echo Missing directory: /vllm-workspace/src/vllm >&2; exit 1; fi
+  if [ ! -x /vllm-workspace/src/vllm/vllm-rs ]; then echo Missing executable: /vllm-workspace/src/vllm/vllm-rs >&2; exit 1; fi
+
+  command -v python3
+  command -v uv
+  command -v pytest
+
+  if ! command -v amd-smi >/dev/null 2>&1 && ! command -v rocminfo >/dev/null 2>&1; then
+    echo No ROCm CLI found in image >&2
+    exit 1
+  fi
+
+  python3 - <<PY
+import torch
+import vllm
+
+print(torch.__version__)
+print(vllm.__version__)
+PY
+
+  echo AMD image smoke OK
+'

.buildkite/scripts/run-multi-node-test.sh

@@ -109,7 +109,9 @@ run_nodes() {
         if [ "$node" -ne 0 ]; then
             docker exec -d "node$node" /bin/bash -c "cd $WORKING_DIR ; ${COMMANDS[$node]}"
         else
-            docker exec "node$node" /bin/bash -c "cd $WORKING_DIR ; ${COMMANDS[$node]}"
+            # Allocate a TTY (-t -i) for the foreground head node so its output
+            # keeps ANSI color in the Buildkite log (see run-amd-test.sh).
+            docker exec -t -i "node$node" /bin/bash -c "cd $WORKING_DIR ; ${COMMANDS[$node]}"
         fi
     done
 }

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

@@ -8,7 +8,12 @@ if [[ "$MODE" != "style-clippy" && "$MODE" != "test" ]]; then
   exit 2
 fi
 
-ROOT_DIR="$(git rev-parse --show-toplevel)"
+if ROOT_DIR="$(git rev-parse --show-toplevel 2>/dev/null)"; then
+  :
+else
+  SCRIPT_DIR="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")" && pwd -P)"
+  ROOT_DIR="$(cd -- "${SCRIPT_DIR}/../.." && pwd -P)"
+fi
 cd "$ROOT_DIR"
 
 export CARGO_TERM_COLOR="${CARGO_TERM_COLOR:-always}"
@@ -163,8 +168,8 @@ run_style_clippy() {
   log_section "Checking Rust dependency bans"
   cargo deny \
     --manifest-path rust/Cargo.toml \
-    check \
     --config rust/deny.toml \
+    check \
     bans
 
   log_section "Running clippy"

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

@@ -23,6 +23,7 @@ NC='\033[0m' # No Color
 # Default configuration
 PIPELINE="ci"
 DRY_RUN=true
+TORCH_NIGHTLY=false
 
 usage() {
     cat <<EOF
@@ -34,12 +35,14 @@ Sets RUN_ALL=1 and NIGHTLY=1 environment variables.
 SAFETY: Dry-run by default. Use --execute to actually trigger a build.
 
 Options:
-    --execute       Actually trigger the build (default: dry-run)
-    --pipeline      Buildkite pipeline slug (default: ${PIPELINE})
-    --commit        Override commit SHA (default: current HEAD)
-    --branch        Override branch name (default: current branch)
-    --message       Custom build message (default: auto-generated)
-    --help          Show this help message
+    --execute        Actually trigger the build (default: dry-run)
+    --pipeline       Buildkite pipeline slug (default: ${PIPELINE})
+    --commit         Override commit SHA (default: current HEAD)
+    --branch         Override branch name (default: current branch)
+    --message        Custom build message (default: auto-generated)
+    --torch-nightly  Also build and run the full suite against torch nightly
+                     (sets TORCH_NIGHTLY=1)
+    --help           Show this help message
 
 Prerequisites:
     - bk CLI installed: brew tap buildkite/buildkite && brew install buildkite/buildkite/bk
@@ -49,6 +52,7 @@ Examples:
     $(basename "$0")                        # Dry-run, show what would happen
     $(basename "$0") --execute              # Actually trigger the build
     $(basename "$0") --pipeline ci-shadow   # Dry-run with different pipeline
+    $(basename "$0") --torch-nightly        # Dry-run a full torch-nightly run
 EOF
     exit 1
 }
@@ -96,6 +100,10 @@ while [[ $# -gt 0 ]]; do
             MESSAGE="$2"
             shift 2
             ;;
+        --torch-nightly)
+            TORCH_NIGHTLY=true
+            shift
+            ;;
         --help|-h)
             usage
             ;;
@@ -171,11 +179,17 @@ if [[ $(echo "$REMOTE_BRANCHES" | wc -l) -gt 5 ]]; then
 fi
 echo ""
 
+# Environment variables passed to the build.
+BUILD_ENV=("RUN_ALL=1" "NIGHTLY=1")
+if [[ "$TORCH_NIGHTLY" == true ]]; then
+    BUILD_ENV+=("TORCH_NIGHTLY=1")
+fi
+
 log_info "Pipeline: ${PIPELINE}"
 log_info "Branch: ${BRANCH}"
 log_info "Commit: ${COMMIT}"
 log_info "Message: ${MESSAGE}"
-log_info "Environment: RUN_ALL=1, NIGHTLY=1"
+log_info "Environment: ${BUILD_ENV[*]}"
 echo ""
 
 # Build the command
@@ -187,9 +201,10 @@ CMD=(bk build create
     --commit "${COMMIT}"
     --branch "${BRANCH}"
     --message "${MESSAGE}"
-    --env "RUN_ALL=1"
-    --env "NIGHTLY=1"
 )
+for env_var in "${BUILD_ENV[@]}"; do
+    CMD+=(--env "${env_var}")
+done
 
 if [[ "$DRY_RUN" == true ]]; then
     echo "=========================================="
@@ -210,8 +225,14 @@ if [[ "$DRY_RUN" == true ]]; then
     echo "    --commit '$(escape_for_shell "${COMMIT}")' \\"
     echo "    --branch '$(escape_for_shell "${BRANCH}")' \\"
     echo "    --message '$(escape_for_shell "${MESSAGE}")' \\"
-    echo "    --env 'RUN_ALL=1' \\"
-    echo "    --env 'NIGHTLY=1'"
+    last_idx=$(( ${#BUILD_ENV[@]} - 1 ))
+    for i in "${!BUILD_ENV[@]}"; do
+        if [[ $i -eq $last_idx ]]; then
+            echo "    --env '$(escape_for_shell "${BUILD_ENV[$i]}")'"
+        else
+            echo "    --env '$(escape_for_shell "${BUILD_ENV[$i]}")' \\"
+        fi
+    done
     echo ""
     echo "=========================================="
     echo -e "${YELLOW}To actually trigger this build, run:${NC}"

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

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

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

@@ -0,0 +1,21 @@
+#!/bin/bash
+
+set -ex
+
+ORIG_TAG_NAME="$BUILDKITE_COMMIT"
+REPO="vllm/vllm-openai-xpu"
+
+echo "Pushing original XPU tag ${ORIG_TAG_NAME}-xpu to nightly tags in ${REPO}"
+
+aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:"$ORIG_TAG_NAME"-x86_64-xpu
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:"$ORIG_TAG_NAME"-x86_64-xpu ${REPO}:nightly-x86_64
+docker push ${REPO}:nightly-x86_64
+
+docker manifest rm ${REPO}:nightly || true
+docker manifest rm ${REPO}:nightly-"$BUILDKITE_COMMIT" || true
+docker manifest create ${REPO}:nightly ${REPO}:nightly-x86_64 --amend
+docker manifest create ${REPO}:nightly-"$BUILDKITE_COMMIT" ${REPO}:nightly-x86_64 --amend
+docker manifest push ${REPO}:nightly
+docker manifest push ${REPO}:nightly-"$BUILDKITE_COMMIT"

.buildkite/test-amd.yaml

@@ -275,7 +275,7 @@ steps:
   - pytest -v -s v1/cudagraph/test_cudagraph_mode.py
 
 - label: e2e Core (1 GPU) # TBD
-  timeout_in_minutes: 35
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
   agent_pool: mi250_1
   optional: true
@@ -380,6 +380,7 @@ steps:
   - tests/test_outputs.py
   - tests/test_pooling_params.py
   - tests/test_ray_env.py
+  - tests/test_sampling_params.py
   - tests/multimodal
   - tests/renderers
   - tests/standalone_tests/lazy_imports.py
@@ -395,6 +396,7 @@ steps:
   - pytest -v -s test_outputs.py
   - pytest -v -s test_pooling_params.py
   - pytest -v -s test_ray_env.py
+  - pytest -v -s test_sampling_params.py
   - pytest -v -s -m 'cpu_test' multimodal
   - pytest -v -s renderers
   - pytest -v -s tokenizers_
@@ -404,10 +406,40 @@ steps:
   - pytest -v -s transformers_utils
   - pytest -v -s config
 
+#------------------------------------------------------------  mi250 · rust  -----------------------------------------------------------#
+
+- label: Rust Frontend Cargo Style + Clippy # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
+  agent_pool: mi250_1
+  no_gpu: true
+  working_dir: "/vllm-workspace"
+  source_file_dependencies:
+  - rust/
+  - rust-toolchain.toml
+  - .buildkite/test_areas/rust_frontend_cargo.yaml
+  - .buildkite/scripts/run-rust-frontend-cargo-ci.sh
+  commands:
+  - bash .buildkite/scripts/run-rust-frontend-cargo-ci.sh style-clippy
+
+- label: Rust Frontend Cargo Tests # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
+  agent_pool: mi250_1
+  no_gpu: true
+  working_dir: "/vllm-workspace"
+  source_file_dependencies:
+  - rust/
+  - rust-toolchain.toml
+  - .buildkite/test_areas/rust_frontend_cargo.yaml
+  - .buildkite/scripts/run-rust-frontend-cargo-ci.sh
+  commands:
+  - bash .buildkite/scripts/run-rust-frontend-cargo-ci.sh test
+
 #-----------------------------------------------------------  mi250 · docker  ----------------------------------------------------------#
 
 - label: Docker Build Metadata (ROCm) # TBD
-  timeout_in_minutes: 20
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
   agent_pool: mi250_1
   no_gpu: true
@@ -435,7 +467,7 @@ steps:
 #-----------------------------------------------------  mi300 · basic_correctness  -----------------------------------------------------#
 
 - label: Basic Correctness # TBD
-  timeout_in_minutes: 95
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   fast_check: true
@@ -452,7 +484,7 @@ steps:
   - pytest -v -s basic_correctness/test_cpu_offload.py
 
 - label: Distributed Model Tests (2 GPUs) # TBD
-  timeout_in_minutes: 110
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_2
   num_gpus: 2
@@ -472,7 +504,7 @@ steps:
   commands:
   - TARGET_TEST_SUITE=MI300 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)'
-  - pytest models/test_transformers.py -v -s -m 'distributed(num_gpus=2)'
+  - pytest models/transformers/test_backend.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 --ignore models/multimodal/generation/test_phi4siglip.py
   - pytest models/multimodal/generation/test_phi4siglip.py -v -s -m 'distributed(num_gpus=2)'
@@ -692,7 +724,7 @@ steps:
   - pytest -v -s distributed/test_eplb_spec_decode.py
 
 - label: Distributed Tests (2xH100-2xMI300) # TBD
-  timeout_in_minutes: 75
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_2
   num_gpus: 2
@@ -1025,6 +1057,23 @@ steps:
   commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-small.txt
 
+- label: MRCR Eval Small Models # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/model_executor/models/
+  - vllm/model_executor/model_loader/
+  - vllm/v1/attention/backends/
+  - vllm/v1/attention/selector.py
+  - vllm/_aiter_ops.py
+  - vllm/platforms/rocm.py
+  - tests/evals/mrcr/
+  commands:
+  - pytest -s -v evals/mrcr/test_mrcr_correctness.py --config-list-file=evals/mrcr/configs/models-small.txt
+
 - label: LM Eval Small Models (MI300) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
@@ -1248,7 +1297,7 @@ steps:
 #---------------------------------------------------------  mi300 · examples  ----------------------------------------------------------#
 
 - label: Examples # TBD
-  timeout_in_minutes: 90
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   optional: true
@@ -1283,8 +1332,23 @@ steps:
 
 #----------------------------------------------------------  mi300 · kernels  ----------------------------------------------------------#
 
+- label: vLLM IR Tests # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/"
+  source_file_dependencies:
+  - vllm/ir
+  - vllm/kernels
+  - vllm/_aiter_ops.py
+  - vllm/platforms/rocm.py
+  commands:
+  - pytest -v -s tests/ir
+  - pytest -v -s tests/kernels/ir
+
 - label: Kernels Attention Test %N # TBD
-  timeout_in_minutes: 100
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   optional: true
@@ -1317,8 +1381,23 @@ steps:
   commands:
   - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py  kernels/test_concat_mla_q.py kernels/test_top_k_per_row.py
 
+- label: Kernels KDA Test # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/model_executor/layers/fla/ops/kda.py
+  - vllm/model_executor/layers/fla/ops/chunk_delta_h.py
+  - vllm/model_executor/layers/fla/ops/l2norm.py
+  - tests/kernels/test_kda.py
+  - vllm/platforms/rocm.py
+  commands:
+  - pytest -v -s kernels/test_kda.py
+
 - label: Kernels MoE Test %N # TBD
-  timeout_in_minutes: 95
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   optional: true
@@ -1429,13 +1508,135 @@ steps:
   - pytest -v -s model_executor -m '(not slow_test)'
   - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py
 
+#----------------------------------------------------  mi300 · model_runner_v2  -------------------------------------------------------#
+
+- label: Model Runner V2 Core Tests # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/v1/worker/gpu/
+  - vllm/v1/worker/gpu_worker.py
+  - vllm/v1/core/sched/
+  - vllm/v1/attention/
+  - tests/v1/engine/test_llm_engine.py
+  - tests/v1/e2e/
+  - tests/entrypoints/llm/test_struct_output_generate.py
+  - vllm/platforms/rocm.py
+  commands:
+  - set -x
+  - export VLLM_USE_V2_MODEL_RUNNER=1
+  - pytest -v -s v1/engine/test_llm_engine.py -k "not test_engine_metrics"
+  - ENFORCE_EAGER=1 pytest -v -s v1/e2e/general/test_async_scheduling.py -k "not ngram"
+  - pytest -v -s v1/e2e/general/test_context_length.py
+  - pytest -v -s v1/e2e/general/test_min_tokens.py
+  - pytest -v -s entrypoints/llm/test_struct_output_generate.py -k "xgrammar and not speculative_config6 and not speculative_config7 and not speculative_config8 and not speculative_config0"
+
+- label: Model Runner V2 Examples # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/examples"
+  source_file_dependencies:
+  - vllm/v1/worker/gpu/
+  - vllm/v1/core/sched/
+  - vllm/v1/worker/gpu_worker.py
+  - examples/basic/offline_inference/
+  - examples/generate/multimodal/
+  - examples/features/
+  - examples/pooling/embed/vision_embedding_offline.py
+  - examples/features/tensorize_vllm_model.py
+  - examples/deployment/llm_engine_example.py
+  - vllm/platforms/rocm.py
+  commands:
+  - set -x
+  - export VLLM_USE_V2_MODEL_RUNNER=1
+  - pip install tensorizer
+  - python3 basic/offline_inference/chat.py
+  - python3 basic/offline_inference/generate.py --model facebook/opt-125m
+  - python3 generate/multimodal/audio_language_offline.py --seed 0
+  - python3 generate/multimodal/vision_language_offline.py --seed 0
+  - python3 generate/multimodal/vision_language_multi_image_offline.py --seed 0
+  - python3 generate/multimodal/encoder_decoder_multimodal_offline.py --model-type whisper --seed 0
+  - python3 pooling/embed/vision_embedding_offline.py --seed 0
+  - python3 features/automatic_prefix_caching/prefix_caching_offline.py
+  - python3 deployment/llm_engine_example.py
+  - python3 features/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 features/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
+  - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
+  - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
+
+- label: Model Runner V2 Distributed (2 GPUs) # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_2
+  num_gpus: 2
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/v1/worker/gpu/
+  - vllm/v1/worker/gpu_worker.py
+  - tests/basic_correctness/test_basic_correctness.py
+  - tests/v1/distributed/test_async_llm_dp.py
+  - tests/v1/distributed/test_eagle_dp.py
+  - vllm/platforms/rocm.py
+  commands:
+  - set -x
+  - export VLLM_USE_V2_MODEL_RUNNER=1
+  - TARGET_TEST_SUITE=MI300 pytest -v -s basic_correctness/test_basic_correctness.py -m 'distributed(num_gpus=2)' -k "not ray and not True"
+  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py -k "not ray"
+  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
+
+- label: Model Runner V2 Pipeline Parallelism (4 GPUs) # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_4
+  num_gpus: 4
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/v1/worker/gpu/
+  - vllm/v1/worker/gpu_worker.py
+  - tests/distributed/test_pipeline_parallel.py
+  - tests/distributed/test_pp_cudagraph.py
+  - vllm/platforms/rocm.py
+  commands:
+  - set -x
+  - export VLLM_USE_V2_MODEL_RUNNER=1
+  - pytest -v -s distributed/test_pipeline_parallel.py -k "not ray and not Jamba"
+  - pytest -v -s distributed/test_pp_cudagraph.py -k "not ray"
+
+- label: Model Runner V2 Spec Decode # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/v1/worker/gpu/
+  - vllm/v1/worker/gpu_worker.py
+  - tests/v1/spec_decode/test_max_len.py
+  - tests/v1/spec_decode/test_rejection_sampler_utils.py
+  - tests/v1/spec_decode/test_synthetic_rejection_sampler_utils.py
+  - tests/v1/e2e/spec_decode/test_spec_decode.py
+  - vllm/platforms/rocm.py
+  commands:
+  - set -x
+  - export VLLM_USE_V2_MODEL_RUNNER=1
+  - pytest -v -s v1/spec_decode/test_max_len.py -k "eagle or mtp"
+  - pytest -v -s v1/spec_decode/test_rejection_sampler_utils.py
+  - pytest -v -s v1/spec_decode/test_synthetic_rejection_sampler_utils.py
+  - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "eagle or mtp"
+
 #------------------------------------------------------  mi300 · models / basic  -------------------------------------------------------#
 
 - label: Basic Models Tests (Extra Initialization) %N # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
-  parallelism: 2
+  parallelism: 6
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/model_executor/models/
@@ -1460,7 +1661,7 @@ steps:
   - pytest -v -s models/test_initialization.py::test_can_initialize_small_subset
 
 - label: Basic Models Tests (Other) # TBD
-  timeout_in_minutes: 90
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   optional: true
@@ -1468,10 +1669,10 @@ steps:
   source_file_dependencies:
   - vllm/
   - tests/models/test_terratorch.py
-  - tests/models/test_transformers.py
+  - tests/models/transformers/test_backend.py
   - tests/models/test_registry.py
   commands:
-  - pytest -v -s models/test_terratorch.py models/test_transformers.py models/test_registry.py
+  - pytest -v -s models/test_terratorch.py models/transformers/test_backend.py models/test_registry.py
 
 #-----------------------------------------------------  mi300 · models / language  -----------------------------------------------------#
 
@@ -1605,7 +1806,7 @@ steps:
   - cd .. && VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s tests/models/multimodal/generation/test_whisper.py -m core_model
 
 - label: Multi-Modal Processor # 1h 42m
-  timeout_in_minutes: 138
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   optional: true
@@ -1691,7 +1892,7 @@ steps:
   - examples/
   commands:
   - pip install --upgrade git+https://github.com/huggingface/transformers
-  - pytest -v -s tests/models/test_transformers.py
+  - pytest -v -s tests/models/transformers/test_backend.py
   - pytest -v -s tests/models/multimodal/test_mapping.py
   - python3 examples/basic/offline_inference/chat.py
   - python3 examples/generate/multimodal/vision_language_offline.py --model-type qwen2_5_vl
@@ -1737,6 +1938,11 @@ steps:
   - pytest -v -s plugins_tests/test_stats_logger_plugins.py
   - pip uninstall dummy_stat_logger -y
   # END: `stat_logger` plugins test
+  # BEGIN: `endpoint` plugins test
+  - pip install -e ./plugins/vllm_add_dummy_endpoint_plugin
+  - pytest -v -s plugins_tests/test_endpoint_plugins.py
+  - pip uninstall vllm_add_dummy_endpoint_plugin -y
+  # END: `endpoint` plugins test
   # BEGIN: other tests
   - pytest -v -s plugins_tests/test_scheduler_plugins.py
   - pip install -e ./plugins/vllm_add_dummy_model
@@ -1745,6 +1951,144 @@ steps:
   - pytest -v -s plugins_tests/test_oot_registration_offline.py # it needs a clean process
   - pytest -v -s plugins_tests/lora_resolvers # unit tests for in-tree lora resolver plugins
 
+- label: GGUF Plugin # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  optional: true
+  soft_fail: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/model_executor/layers/quantization
+  - tests/plugins_tests/test_gguf_plugin.py
+  - tests/plugins_tests/gguf
+  - vllm/platforms/rocm.py
+  commands:
+  - pip install "vllm-gguf-plugin >= 0.0.2"
+  - pytest -v -s plugins_tests/gguf
+
+#-------------------------------------------------------  mi300 · rust_frontend  -------------------------------------------------------#
+
+- label: Rust Frontend OpenAI Coverage # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - rust/
+  - vllm/benchmarks/
+  - vllm/entrypoints/openai/
+  - vllm/entrypoints/serve/
+  - vllm/v1/sample/
+  - tests/utils.py
+  - tests/benchmarks/test_serve_cli.py
+  - tests/entrypoints/openai/chat_completion/test_chat_completion.py
+  - tests/entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py
+  - tests/entrypoints/openai/completion/test_shutdown.py
+  - tests/entrypoints/openai/test_return_token_ids.py
+  - tests/entrypoints/openai/test_uds.py
+  - tests/v1/sample/test_logprobs_e2e.py
+  - vllm/platforms/rocm.py
+  commands:
+  - export VLLM_USE_RUST_FRONTEND=1
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s benchmarks/test_serve_cli.py -k "not insecure and not (test_bench_serve and not test_bench_serve_chat)"
+  - pytest -v -s entrypoints/openai/chat_completion/test_chat_completion.py -k "not test_invalid_json_schema and not test_invalid_regex"
+  - pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not multiple"
+  - pytest -v -s entrypoints/openai/completion/test_shutdown.py -k "not engine_failure and not test_abort_timeout_exits_quickly"
+  - pytest -v -s entrypoints/openai/test_return_token_ids.py -k "not test_comparison"
+  - pytest -v -s entrypoints/openai/test_uds.py
+  - pytest -v -s v1/sample/test_logprobs_e2e.py -k "test_prompt_logprobs_e2e_server"
+
+- label: Rust Frontend Serve Admin Coverage # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - rust/
+  - vllm/entrypoints/openai/
+  - vllm/entrypoints/serve/
+  - vllm/v1/engine/
+  - tests/utils.py
+  - tests/entrypoints/serve/dev/rpc/test_collective_rpc.py
+  - tests/entrypoints/scale_out/token_in_token_out/test_serving_tokens.py
+  - tests/entrypoints/serve/instrumentator/test_basic.py
+  - tests/entrypoints/serve/instrumentator/test_metrics.py
+  - tests/entrypoints/serve/tokenize/test_tokenization.py
+  - vllm/platforms/rocm.py
+  commands:
+  - export VLLM_USE_RUST_FRONTEND=1
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
+  - pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not show_version and not server_load"
+  - pytest -v -s entrypoints/scale_out/token_in_token_out/test_serving_tokens.py -k "not stream and not lora and not test_generate_logprobs and not stop_string_workflow"
+  - pytest -v -s entrypoints/serve/instrumentator/test_metrics.py -k "text and not show and not run_batch and not test_metrics_counts and not test_metrics_exist"
+  - pytest -v -s entrypoints/serve/tokenize/test_tokenization.py -k "not tokenizer_info"
+
+- label: Rust Frontend Core Correctness # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - rust/
+  - vllm/entrypoints/openai/
+  - tests/utils.py
+  - tests/entrypoints/openai/correctness/test_lmeval.py
+  - vllm/platforms/rocm.py
+  commands:
+  - export VLLM_USE_RUST_FRONTEND=1
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
+
+- label: Rust Frontend Tool Use # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - rust/
+  - vllm/entrypoints/openai/
+  - vllm/tool_parsers/
+  - tests/utils.py
+  - tests/tool_use/
+  - vllm/platforms/rocm.py
+  commands:
+  - export VLLM_USE_RUST_FRONTEND=1
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s tool_use --ignore=tool_use/mistral --models llama3.2 -k "not test_response_format_with_tool_choice_required and not test_parallel_tool_calls_false and not test_tool_call_and_choice"
+
+- label: Rust Frontend Distributed # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_4
+  num_gpus: 4
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - rust/
+  - vllm/distributed/
+  - vllm/engine/
+  - vllm/executor/
+  - vllm/v1/engine/
+  - vllm/v1/worker/
+  - tests/utils.py
+  - tests/v1/distributed/test_external_lb_dp.py
+  - tests/v1/distributed/test_hybrid_lb_dp.py
+  - tests/v1/distributed/test_internal_lb_dp.py
+  - vllm/platforms/rocm.py
+  commands:
+  - export VLLM_USE_RUST_FRONTEND=1
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_internal_lb_dp.py -k "not 4 and not server_info"
+  - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py -k "not 4 and not server_info"
+  - TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_hybrid_lb_dp.py -k "not 4 and not server_info"
+
 #-------------------------------------------------------  mi300 · quantization  --------------------------------------------------------#
 
 - label: Quantization # TBD
@@ -1911,7 +2255,7 @@ steps:
   - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
 
 - label: Spec Decode Eagle # TBD
-  timeout_in_minutes: 90
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   optional: true
@@ -2063,7 +2407,7 @@ steps:
   - pytest -v -s v1/worker
   - pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
   - pytest -v -s -m 'not cpu_test' v1/metrics
-  - pip install -U git+https://github.com/robertgshaw2-redhat/lm-evaluation-harness.git@streaming-api
+  - pip install -U git+https://github.com/vllm-project/lm-evaluation-harness.git@streaming-api
   # - export HSA_NO_SCRATCH_RECLAIM=1
   - pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
 
@@ -2128,7 +2472,7 @@ steps:
   - DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py
 
 - label: Metrics, Tracing (2 GPUs) # TBD
-  timeout_in_minutes: 65
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_2
   optional: true
@@ -2254,6 +2598,59 @@ steps:
   - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
   - HYBRID_SSM=1 ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
+- label: Hybrid SSM NixlConnector PD prefix cache test (2 GPUs) # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_2
+  num_gpus: 2
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+  - vllm/v1/core/sched/
+  - vllm/v1/core/kv_cache_coordinator.py
+  - tests/v1/kv_connector/nixl_integration/
+  - vllm/platforms/rocm.py
+  commands:
+  - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
+  - ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/run_mamba_prefix_cache_test.sh
+
+- label: MultiConnector (Nixl+Offloading) PD accuracy (2 GPUs) # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_2
+  num_gpus: 2
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+  - vllm/distributed/kv_transfer/kv_connector/v1/multi_connector.py
+  - vllm/distributed/kv_transfer/kv_connector/v1/offloading_connector.py
+  - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
+  - tests/v1/kv_connector/nixl_integration/
+  - vllm/platforms/rocm.py
+  commands:
+  - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
+  - ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/run_multi_connector_accuracy_test.sh
+
+- label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs) # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_2
+  num_gpus: 2
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+  - vllm/distributed/kv_transfer/kv_connector/v1/multi_connector.py
+  - vllm/distributed/kv_transfer/kv_connector/v1/offloading_connector.py
+  - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
+  - tests/v1/kv_connector/nixl_integration/
+  - vllm/platforms/rocm.py
+  commands:
+  - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
+  - ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh
+
 - label: V1 e2e (4 GPUs) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
@@ -2281,7 +2678,7 @@ steps:
 #------------------------------------------------------  mi300 · weight_loading  -------------------------------------------------------#
 
 - label: Weight Loading Multiple GPU # TBD
-  timeout_in_minutes: 75
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_2
   num_gpus: 2
@@ -2293,7 +2690,7 @@ steps:
   - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-amd.txt
 
 - label: Weight Loading Multiple GPU - Large Models # TBD
-  timeout_in_minutes: 75
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_2
   num_gpus: 2
@@ -2544,15 +2941,19 @@ steps:
   - vllm/envs.py
   - examples/offline_inference/data_parallel.py
   - tests/distributed/test_context_parallel.py
+  - tests/distributed/test_rocm_aiter_custom_ar.py
   - tests/distributed/test_rocm_quick_reduce.py
   - tests/distributed/test_quick_all_reduce.py
+  - tests/v1/e2e/general/test_rocm_aiter_custom_ar.py
   - tests/v1/distributed/test_dbo.py
   - tests/utils.py
   commands:
   - pytest -v -s tests/distributed/test_context_parallel.py
-  - pytest -v -s tests/v1/distributed/test_dbo.py
+  - pytest -v -s tests/distributed/test_rocm_aiter_custom_ar.py
+  - pytest -v -s tests/v1/e2e/general/test_rocm_aiter_custom_ar.py
   - pytest -v -s tests/distributed/test_rocm_quick_reduce.py
   - pytest -v -s tests/distributed/test_quick_all_reduce.py
+  - pytest -v -s tests/v1/distributed/test_dbo.py
 
 #--------------------------------------------------------  mi355 · entrypoints  --------------------------------------------------------#
 
@@ -2688,7 +3089,7 @@ steps:
     - pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-gfx950.txt
 
 - label: LM Eval Qwen3-5 Models (B200-MI355) # TBD
-  timeout_in_minutes: 120
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_2
   num_gpus: 2
@@ -2770,7 +3171,7 @@ steps:
 #---------------------------------------------------------  mi355 · examples  ----------------------------------------------------------#
 
 - label: Examples # TBD
-  timeout_in_minutes: 45
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
   working_dir: "/vllm-workspace/examples"
@@ -2805,7 +3206,7 @@ steps:
 #----------------------------------------------------------  mi355 · kernels  ----------------------------------------------------------#
 
 - label: Kernels (B200-MI355) # TBD
-  timeout_in_minutes: 15
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
   working_dir: "/vllm-workspace/"
@@ -2830,7 +3231,7 @@ steps:
   - pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py
 
 - label: Kernels Attention Test %N # TBD
-  timeout_in_minutes: 100
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
   parallelism: 2
@@ -2847,7 +3248,7 @@ steps:
   - pytest -v -s kernels/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
 
 - label: Kernels MoE Test %N # TBD
-  timeout_in_minutes: 50
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
   parallelism: 5
@@ -3111,7 +3512,7 @@ steps:
   - pytest -v -s v1/attention
 
 - label: V1 Core + KV + Metrics # TBD
-  timeout_in_minutes: 60
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
   optional: true
@@ -3133,11 +3534,11 @@ steps:
   - pytest -v -s v1/worker
   - pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
   - pytest -v -s -m 'not cpu_test' v1/metrics
-  - pip install -U git+https://github.com/robertgshaw2-redhat/lm-evaluation-harness.git@streaming-api
+  - pip install -U git+https://github.com/vllm-project/lm-evaluation-harness.git@streaming-api
   - pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
 
 - label: V1 Sample + Logits # TBD
-  timeout_in_minutes: 60
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
   optional: true
@@ -3157,7 +3558,7 @@ steps:
   - pytest -v -s v1/test_outputs.py
 
 - label: V1 Spec Decode # TBD
-  timeout_in_minutes: 60
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
   working_dir: "/vllm-workspace/tests"
@@ -3170,7 +3571,7 @@ steps:
 #------------------------------------------------------  mi355 · weight_loading  -------------------------------------------------------#
 
 - label: Weight Loading Multiple GPU # TBD
-  timeout_in_minutes: 75
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_2
   num_gpus: 2
@@ -3182,7 +3583,7 @@ steps:
   - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-amd.txt
 
 - label: Weight Loading Multiple GPU - Large Models # TBD
-  timeout_in_minutes: 75
+  timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_2
   working_dir: "/vllm-workspace/tests"

.buildkite/test_areas/attention.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: V1 attention (H100-MI300)
   key: v1-attention-h100-mi300
-  timeout_in_minutes: 30
+  timeout_in_minutes: 85
   device: h100
   source_file_dependencies:
     - vllm/config/attention.py
@@ -16,7 +16,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 70
+      timeout_in_minutes: 95
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -30,7 +30,7 @@ steps:
 
 - label: V1 attention (B200)
   key: v1-attention-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 80
   device: b200-k8s
   source_file_dependencies:
     - vllm/config/attention.py

.buildkite/test_areas/basic_correctness.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Basic Correctness
   key: basic-correctness
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -19,6 +19,6 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 50
+      timeout_in_minutes: 70
       depends_on:
       - image-build-amd

.buildkite/test_areas/benchmarks.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Benchmarks CLI Test
   key: benchmarks-cli-test
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -23,7 +23,7 @@ steps:
   num_gpus: 2
   optional: true
   working_dir: "/vllm-workspace/"
-  timeout_in_minutes: 10
+  timeout_in_minutes: 20
   source_file_dependencies:
   - benchmarks/attention_benchmarks/
   - vllm/v1/attention/

.buildkite/test_areas/compile.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Sequence Parallel Correctness Tests (2 GPUs)
   key: sequence-parallel-correctness-tests-2-gpus
-  timeout_in_minutes: 50
+  timeout_in_minutes: 80
   working_dir: "/vllm-workspace/"
   num_devices: 2
   source_file_dependencies:
@@ -19,7 +19,7 @@ steps:
 
 - label: Sequence Parallel Correctness Tests (2xH100)
   key: sequence-parallel-correctness-tests-2xh100
-  timeout_in_minutes: 50
+  timeout_in_minutes: 75
   working_dir: "/vllm-workspace/"
   device: h100
   optional: true
@@ -30,7 +30,7 @@ steps:
 
 - label: AsyncTP Correctness Tests (2xH100)
   key: asynctp-correctness-tests-2xh100
-  timeout_in_minutes: 50
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/"
   device: h100
   optional: true
@@ -41,7 +41,7 @@ steps:
 
 - label: AsyncTP Correctness Tests (B200)
   key: asynctp-correctness-tests-b200
-  timeout_in_minutes: 50
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/"
   device: b200-k8s
   optional: true
@@ -52,7 +52,7 @@ steps:
 
 - label: Distributed Compile Unit Tests (2xH100)
   key: distributed-compile-unit-tests-2xh100
-  timeout_in_minutes: 20
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/"
   device: h100
   num_devices: 2
@@ -66,7 +66,7 @@ steps:
 
 - label: Fusion and Compile Unit Tests (2xB200)
   key: fusion-and-compile-unit-tests-2xb200
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/"
   device: b200-k8s
   source_file_dependencies:
@@ -96,7 +96,7 @@ steps:
 
 - label: Fusion E2E Quick (H100)
   key: fusion-e2e-quick-h100
-  timeout_in_minutes: 15
+  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/"
   device: h100
   num_devices: 1
@@ -115,7 +115,7 @@ steps:
 
 - label: Fusion E2E Config Sweep (H100)
   key: fusion-e2e-config-sweep-h100
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/"
   device: h100
   num_devices: 1
@@ -149,7 +149,7 @@ steps:
 
 - label: Fusion E2E TP2 Quick (H100)
   key: fusion-e2e-tp2-quick-h100
-  timeout_in_minutes: 20
+  timeout_in_minutes: 35
   working_dir: "/vllm-workspace/"
   device: h100
   num_devices: 2
@@ -167,7 +167,7 @@ steps:
 
 - label: Fusion E2E TP2 AR-RMS Config Sweep (H100)
   key: fusion-e2e-tp2-ar-rms-config-sweep-h100
-  timeout_in_minutes: 40
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/"
   device: h100
   num_devices: 2
@@ -207,7 +207,7 @@ steps:
 
 - label: Fusion E2E TP2 (B200)
   key: fusion-e2e-tp2-b200
-  timeout_in_minutes: 20
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/"
   device: b200-k8s
   num_devices: 2

.buildkite/test_areas/cuda.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Platform Tests
   key: platform-tests
-  timeout_in_minutes: 15
+  timeout_in_minutes: 20
   device: h200_18gb
   source_file_dependencies:
   - vllm/envs.py
@@ -19,7 +19,7 @@ steps:
 
 - label: Cudagraph
   key: cudagraph
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   source_file_dependencies:
   - tests/v1/cudagraph
   - vllm/v1/cudagraph_dispatcher.py

.buildkite/test_areas/disaggregated.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed NixlConnector PD accuracy (4 GPUs)
   key: distributed-nixlconnector-pd-accuracy-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 55
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -16,7 +16,7 @@ steps:
   mirror:
     amd:
       device: mi300_4
-      timeout_in_minutes: 110
+      timeout_in_minutes: 85
       depends_on:
         - image-build-amd
       source_file_dependencies:
@@ -29,7 +29,7 @@ steps:
 
 - label: Distributed FlashInfer NixlConnector PD accuracy (4 GPUs)
   key: distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 55
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -39,6 +39,19 @@ steps:
     - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
     - FLASHINFER=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
+- label: Push NixlConnector PP prefill PD accuracy (4 GPUs)
+  key: push-nixlconnector-pp-prefill-pd-accuracy-4-gpus
+  timeout_in_minutes: 30
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 4
+  source_file_dependencies:
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - tests/v1/kv_connector/nixl_integration/
+    - tests/v1/kv_connector/nixl_push_integration/
+  commands:
+    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+    - bash v1/kv_connector/nixl_push_integration/config_sweep_accuracy_test.sh
+
 - label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs)
   key: dp-ep-distributed-nixlconnector-pd-accuracy-tests-4-gpus
   timeout_in_minutes: 30
@@ -53,7 +66,7 @@ steps:
   mirror:
     amd:
       device: mi300_4
-      timeout_in_minutes: 50
+      timeout_in_minutes: 60
       depends_on:
         - image-build-amd
       source_file_dependencies:
@@ -66,7 +79,7 @@ steps:
 
 - label: CrossLayer KV layout Distributed NixlConnector PD accuracy tests (4 GPUs)
   key: crosslayer-kv-layout-distributed-nixlconnector-pd-accuracy-tests-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 55
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -78,7 +91,7 @@ steps:
   mirror:
     amd:
       device: mi300_4
-      timeout_in_minutes: 110
+      timeout_in_minutes: 85
       depends_on:
         - image-build-amd
       source_file_dependencies:
@@ -91,7 +104,7 @@ steps:
 
 - label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs)
   key: hybrid-ssm-nixlconnector-pd-accuracy-tests-4-gpus
-  timeout_in_minutes: 25
+  timeout_in_minutes: 60
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -103,7 +116,7 @@ steps:
   mirror:
     amd:
       device: mi300_4
-      timeout_in_minutes: 60
+      timeout_in_minutes: 80
       depends_on:
         - image-build-amd
       source_file_dependencies:
@@ -130,7 +143,7 @@ steps:
 
 - label: MultiConnector (Nixl+Offloading) PD accuracy (2 GPUs)
   key: multiconnector-nixl-offloading-pd-accuracy-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -145,7 +158,7 @@ steps:
 
 - label: NixlConnector PD + Spec Decode acceptance (2 GPUs)
   key: nixlconnector-pd-spec-decode-acceptance-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   device: a100
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
@@ -159,7 +172,7 @@ steps:
   mirror:
     amd:
       device: mi300_2
-      timeout_in_minutes: 60
+      timeout_in_minutes: 70
       depends_on:
         - image-build-amd
       source_file_dependencies:
@@ -173,7 +186,7 @@ steps:
 
 - label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs)
   key: multiconnector-nixl-offloading-pd-edge-cases-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:

.buildkite/test_areas/distributed.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed Comm Ops
   key: distributed-comm-ops
-  timeout_in_minutes: 20
+  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -18,7 +18,7 @@ steps:
 
 - label: Distributed DP Tests (2 GPUs)
   key: distributed-dp-tests-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 35
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -55,7 +55,7 @@ steps:
 
 - label: Distributed Compile + RPC Tests (2 GPUs)
   key: distributed-compile-rpc-tests-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 65
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -78,7 +78,7 @@ steps:
 
 - label: Distributed Torchrun + Shutdown Tests (2 GPUs)
   key: distributed-torchrun-shutdown-tests-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -133,7 +133,7 @@ steps:
 
 - label: Distributed DP Tests (4 GPUs)
   key: distributed-dp-tests-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -154,7 +154,7 @@ steps:
 
 - label: Distributed Compile + Comm (4 GPUs)
   key: distributed-compile-comm-4-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 70
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -176,7 +176,7 @@ steps:
 
 - label: Distributed Tests (8xH100)
   key: distributed-tests-8xh100
-  timeout_in_minutes: 10
+  timeout_in_minutes: 20
   device: h100
   num_devices: 8
   working_dir: "/vllm-workspace/tests"
@@ -212,7 +212,7 @@ steps:
 
 - label: Distributed Tests (2xH100-2xMI300)
   key: distributed-tests-2xh100-2xmi300
-  timeout_in_minutes: 15
+  timeout_in_minutes: 30
   device: h100
   optional: true
   working_dir: "/vllm-workspace/"
@@ -259,7 +259,7 @@ steps:
 
 - label: Pipeline + Context Parallelism (4 GPUs)
   key: pipeline-context-parallelism-4-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 55
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -274,7 +274,7 @@ steps:
 
 - label: RayExecutorV2 (4 GPUs)
   key: rayexecutorv2-4-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:

.buildkite/test_areas/docker.yaml

@@ -3,7 +3,7 @@ depends_on:
   - image-build-cpu
 steps:
 - label: Docker Build Metadata
-  timeout_in_minutes: 10
+  timeout_in_minutes: 20
   device: cpu-small
   source_file_dependencies:
     - .buildkite/release-pipeline.yaml

.buildkite/test_areas/e2e_integration.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: DeepSeek V2-Lite Sync EPLB Accuracy (4xH100)
   key: deepseek-v2-lite-sync-eplb-accuracy-4xh100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
   device: h100
   optional: true
   num_devices: 4
@@ -14,7 +14,7 @@ steps:
 
 - label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (4xH100)
   key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-4xh100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
   device: h100
   optional: true
   num_devices: 4
@@ -24,7 +24,7 @@ steps:
 
 - label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (2xB200)
   key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-2xb200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 20
   device: b200-k8s
   optional: true
   num_devices: 2
@@ -34,7 +34,7 @@ steps:
 
 - label: Qwen3-30B-A3B-FP8 DP4 Async EPLB Accuracy
   key: qwen3-30b-a3b-fp8-dp4-async-eplb-accuracy
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
   device: h100
   optional: true
   num_devices: 4
@@ -44,7 +44,7 @@ steps:
 
 - label: DeepSeek V2-Lite Prefetch Offload Accuracy (H100)
   key: deepseek-v2-lite-prefetch-offload-accuracy-h100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 20
   device: h100
   optional: true
   num_devices: 1

.buildkite/test_areas/engine.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Engine
   key: engine
-  timeout_in_minutes: 15
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
   - vllm/compilation/
@@ -29,13 +29,13 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 50
       depends_on:
       - image-build-amd
 
 - label: Engine (1 GPU)
   key: engine-1-gpu
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   source_file_dependencies:
     - vllm/v1/engine/
     - tests/v1/engine/
@@ -45,13 +45,13 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 40
+      timeout_in_minutes: 55
       depends_on:
       - image-build-amd
 
 - label: e2e Scheduling (1 GPU)
   key: e2e-scheduling-1-gpu
-  timeout_in_minutes: 30
+  timeout_in_minutes: 35
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/
@@ -60,15 +60,15 @@ steps:
     - pytest -v -s v1/e2e/general/test_async_scheduling.py
   mirror:
     amd:
-      device: mi250_1
-      timeout_in_minutes: 60
+      device: mi325_1
+      timeout_in_minutes: 70
       depends_on:
       - image-build-amd
 
 - label: e2e Core (1 GPU)
   device: h200_35gb
   key: e2e-core-1-gpu
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   source_file_dependencies:
     - vllm/v1/
     - tests/v1/e2e/general/
@@ -76,8 +76,8 @@ steps:
     - pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py
   mirror:
     amd:
-      device: mi250_1
-      timeout_in_minutes: 35
+      device: mi325_1
+      timeout_in_minutes: 60
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -87,7 +87,7 @@ steps:
 
 - label: V1 e2e (2 GPUs)
   key: v1-e2e-2-gpus
-  timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
+  timeout_in_minutes: 25 # TODO: Fix timeout after we have more confidence in the test stability
   optional: true
   num_devices: 2
   source_file_dependencies:
@@ -120,7 +120,7 @@ steps:
 
 - label: V1 e2e (4 GPUs)
   key: v1-e2e-4-gpus
-  timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
+  timeout_in_minutes: 20 # TODO: Fix timeout after we have more confidence in the test stability
   optional: true
   num_devices: 4
   source_file_dependencies:
@@ -148,7 +148,7 @@ steps:
 
 - label: V1 e2e (4xH100)
   key: v1-e2e-4xh100
-  timeout_in_minutes: 60
+  timeout_in_minutes: 35
   device: h100
   num_devices: 4
   optional: true

.buildkite/test_areas/entrypoints.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Entrypoints Unit Tests
   key: entrypoints-unit-tests
-  timeout_in_minutes: 10
+  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/entrypoints
@@ -16,7 +16,7 @@ steps:
 
 - label: Entrypoints Integration (LLM)
   key: entrypoints-integration-llm
-  timeout_in_minutes: 40
+  timeout_in_minutes: 60
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -37,7 +37,7 @@ steps:
 - label: Entrypoints Integration (API Server)
   key: entrypoints-integration-api-server
   device: h200_35gb
-  timeout_in_minutes: 130
+  timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -56,7 +56,7 @@ steps:
 
 - label: Entrypoints Integration (API Server OpenAI - Part 1)
   key: entrypoints-integration-api-server-openai-part-1
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -68,13 +68,13 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 80
+      timeout_in_minutes: 65
       depends_on:
       - image-build-amd
 
 - label: Entrypoints Integration (API Server OpenAI - Part 2)
   key: entrypoints-integration-api-server-openai-part-2
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -109,7 +109,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 65
       depends_on:
       - image-build-amd
 
@@ -126,7 +126,7 @@ steps:
 - label: Entrypoints Integration (Speech to Text)
   device: h200_35gb
   key: entrypoints-integration-speech_to_text
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -138,7 +138,7 @@ steps:
 - label: Entrypoints Integration (Multimodal)
   device: h200_35gb
   key: entrypoints-integration-multimodal
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -160,7 +160,7 @@ steps:
 
 - label: OpenAI API Correctness
   key: openai-api-correctness
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   device: h200_18gb
   source_file_dependencies:
   - csrc/

.buildkite/test_areas/expert_parallelism.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: EPLB Algorithm
   key: eplb-algorithm
-  timeout_in_minutes: 15
+  timeout_in_minutes: 20
   device: h200_18gb
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -27,7 +27,7 @@ steps:
 
 - label: EPLB Execution # 17min
   key: eplb-execution
-  timeout_in_minutes: 27
+  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -39,7 +39,7 @@ steps:
 
 - label: Elastic EP Scaling Test
   key: elastic-ep-scaling-test
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   device: h100
   working_dir: "/vllm-workspace/tests"
   num_devices: 4

.buildkite/test_areas/kernels.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: vLLM IR Tests
   key: vllm-ir-tests
-  timeout_in_minutes: 10
+  timeout_in_minutes: 35
   device: h200_18gb
   working_dir: "/vllm-workspace/"
   source_file_dependencies:
@@ -16,7 +16,7 @@ steps:
 
 - label: Kernels Core Operation Test
   key: kernels-core-operation-test
-  timeout_in_minutes: 75
+  timeout_in_minutes: 120
   source_file_dependencies:
   - csrc/
   - tests/kernels/core
@@ -27,7 +27,7 @@ steps:
 
 - label: Kernels MiniMax Reduce RMS Test (2 GPUs)
   key: kernels-minimax-reduce-rms-test-2-gpus
-  timeout_in_minutes: 15
+  timeout_in_minutes: 20
   num_devices: 2
   device: h100
   source_file_dependencies:
@@ -41,7 +41,7 @@ steps:
 
 - label: Deepseek V4 Kernel Test (H100)
   key: deepseek-v4-kernel-test-h100
-  timeout_in_minutes: 15
+  timeout_in_minutes: 30
   device: h100
   source_file_dependencies:
   - csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
@@ -54,7 +54,7 @@ steps:
 
 - label: Deepseek V4 Kernel Test (B200)
   key: deepseek-v4-kernel-test-b200
-  timeout_in_minutes: 15
+  timeout_in_minutes: 20
   device: b200-k8s
   source_file_dependencies:
   - csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
@@ -65,7 +65,7 @@ steps:
 
 - label: Kernels Attention Test %N
   key: kernels-attention-test
-  timeout_in_minutes: 35
+  timeout_in_minutes: 65
   source_file_dependencies:
   - csrc/attention/
   - vllm/v1/attention
@@ -79,7 +79,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 55
+      timeout_in_minutes: 90
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -106,7 +106,7 @@ steps:
 
 - label: Kernels Quantization Test %N
   key: kernels-quantization-test
-  timeout_in_minutes: 90
+  timeout_in_minutes: 60
   source_file_dependencies:
   - csrc/quantization/
   - vllm/model_executor/layers/quantization
@@ -131,7 +131,7 @@ steps:
 
 - label: Kernels MoE Test %N
   key: kernels-moe-test
-  timeout_in_minutes: 25
+  timeout_in_minutes: 50
   source_file_dependencies:
   - csrc/quantization/cutlass_w8a8/moe/
   - csrc/moe/
@@ -147,7 +147,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 50
+      timeout_in_minutes: 65
       source_file_dependencies:
       - csrc/quantization/cutlass_w8a8/moe/
       - csrc/moe/
@@ -163,7 +163,7 @@ steps:
 
 - label: Kernels Mamba Test
   key: kernels-mamba-test
-  timeout_in_minutes: 45
+  timeout_in_minutes: 40
   source_file_dependencies:
   - csrc/mamba/
   - tests/kernels/mamba
@@ -172,7 +172,7 @@ steps:
     - pytest -v -s kernels/mamba
 
 - label: Kernels KDA Test
-  timeout_in_minutes: 20
+  timeout_in_minutes: 25
   device: h200_18gb
   source_file_dependencies:
   - vllm/model_executor/layers/fla/ops/kda.py
@@ -184,7 +184,7 @@ steps:
 
 - label: Kernels DeepGEMM Test (H100)
   key: kernels-deepgemm-test-h100
-  timeout_in_minutes: 45
+  timeout_in_minutes: 35
   device: h100
   num_devices: 1
   source_file_dependencies:
@@ -211,7 +211,7 @@ steps:
 
 - label: Kernels (B200)
   key: kernels-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 80
   working_dir: "/vllm-workspace/"
   device: b200-k8s
   # optional: true
@@ -264,7 +264,7 @@ steps:
 
 - label: Kernels Helion Test
   key: kernels-helion-test
-  timeout_in_minutes: 30
+  timeout_in_minutes: 115
   device: h100
   source_file_dependencies:
   - vllm/utils/import_utils.py
@@ -276,7 +276,7 @@ steps:
  
 - label: Kernels FP8 MoE Test (1xH100)
   key: kernels-fp8-moe-test-1xh100
-  timeout_in_minutes: 90
+  timeout_in_minutes: 40
   device: h100
   num_devices: 1
   optional: true
@@ -293,7 +293,7 @@ steps:
 
 - label: Kernels FP8 MoE Test (2xH100)
   key: kernels-fp8-moe-test-2xh100
-  timeout_in_minutes: 90
+  timeout_in_minutes: 45
   device: h100
   num_devices: 2
   optional: true
@@ -303,7 +303,7 @@ steps:
 
 - label: Kernels Fp4 MoE Test (B200)
   key: kernels-fp4-moe-test-b200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
   device: b200-k8s
   num_devices: 1
   optional: true
@@ -316,7 +316,7 @@ steps:
 
 - label: Kernels FusedMoE Layer Test (2 H100s)
   key: kernels-fusedmoe-layer-test-2-h100s
-  timeout_in_minutes: 90
+  timeout_in_minutes: 30
   device: h100
   num_devices: 2
   source_file_dependencies:

.buildkite/test_areas/lm_eval.yaml

@@ -5,7 +5,7 @@ steps:
 - label: LM Eval Small Models
   device: h200_35gb
   key: lm-eval-small-models
-  timeout_in_minutes: 75
+  timeout_in_minutes: 45
   source_file_dependencies:
   - csrc/
   - vllm/model_executor/layers/quantization
@@ -56,7 +56,7 @@ steps:
 
 - label: LM Eval Small Models (1xB200)
   key: lm-eval-small-models-1xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 50
   device: b200-k8s
   optional: true
   source_file_dependencies:
@@ -80,7 +80,7 @@ steps:
 
 - label: LM Eval Large Models EP (2xB200)
   key: lm-eval-large-models-ep-2xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 60
   device: b200-k8s
   optional: true
   num_devices: 2
@@ -92,7 +92,7 @@ steps:
 
 - label: LM Eval Qwen3.5 Models (2xB200)
   key: lm-eval-qwen3-5-models-2xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 45
   device: b200-k8s
   optional: true
   num_devices: 2
@@ -109,7 +109,7 @@ steps:
 
 - label: LM Eval Large Models (8xH200)
   key: lm-eval-large-models-8xh200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 50
   device: h200
   optional: true
   num_devices: 8
@@ -118,7 +118,7 @@ steps:
   mirror:
     amd:
       device: mi300_8
-      timeout_in_minutes: 180
+      timeout_in_minutes: 60
       depends_on:
       - image-build-amd
       commands:
@@ -150,9 +150,9 @@ steps:
   commands:
     - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/moe-refactor-dp-ep/config-b200.txt
 
-- label: LM Eval Humming (A100 - TEMPORARY)
-  key: lm-eval-humming-a100
-  timeout_in_minutes: 30
+- label: LM Eval Humming f16 (A100 - TEMPORARY)
+  key: lm-eval-humming-f16-a100
+  timeout_in_minutes: 75
   device: a100
   optional: true
   num_devices: 1
@@ -160,13 +160,29 @@ steps:
   - vllm/model_executor/layers/quantization/humming.py
   - vllm/model_executor/layers/quantization/utils/humming_utils.py
   - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
-  - vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
+  - vllm/model_executor/layers/fused_moe/oracle/
+  - vllm/model_executor/kernels/linear/
   commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config.txt
 
-- label: LM Eval Humming (H100 - TEMPORARY)
-  key: lm-eval-humming-h100
-  timeout_in_minutes: 30
+- label: LM Eval Humming Act int8 (A100 - TEMPORARY)
+  key: lm-eval-humming-act-a100
+  timeout_in_minutes: 45
+  device: a100
+  optional: true
+  num_devices: 1
+  source_file_dependencies:
+  - vllm/model_executor/layers/quantization/humming.py
+  - vllm/model_executor/layers/quantization/utils/humming_utils.py
+  - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
+  - vllm/model_executor/layers/fused_moe/oracle/
+  - vllm/model_executor/kernels/linear/
+  commands:
+  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-int8.txt
+
+- label: LM Eval Humming f16 (H100 - TEMPORARY)
+  key: lm-eval-humming-f16-h100
+  timeout_in_minutes: 70
   device: h100
   optional: true
   num_devices: 1
@@ -174,14 +190,30 @@ steps:
   - vllm/model_executor/layers/quantization/humming.py
   - vllm/model_executor/layers/quantization/utils/humming_utils.py
   - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
-  - vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
+  - vllm/model_executor/layers/fused_moe/oracle/
+  - vllm/model_executor/kernels/linear/
   commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config.txt
-  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-fp8.txt
 
-- label: LM Eval Humming (B200 - TEMPORARY)
-  key: lm-eval-humming-b200
-  timeout_in_minutes: 30
+- label: LM Eval Humming Act fp8/int8 (H100 - TEMPORARY)
+  key: lm-eval-humming-act-h100
+  timeout_in_minutes: 70
+  device: h100
+  optional: true
+  num_devices: 1
+  source_file_dependencies:
+  - vllm/model_executor/layers/quantization/humming.py
+  - vllm/model_executor/layers/quantization/utils/humming_utils.py
+  - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
+  - vllm/model_executor/layers/fused_moe/oracle/
+  - vllm/model_executor/kernels/linear/
+  commands:
+  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-fp8.txt
+  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-int8.txt
+
+- label: LM Eval Humming f16 (B200 - TEMPORARY)
+  key: lm-eval-humming-f16-b200
+  timeout_in_minutes: 50
   device: b200-k8s
   optional: true
   num_devices: 1
@@ -189,14 +221,30 @@ steps:
   - vllm/model_executor/layers/quantization/humming.py
   - vllm/model_executor/layers/quantization/utils/humming_utils.py
   - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
-  - vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
+  - vllm/model_executor/layers/fused_moe/oracle/
+  - vllm/model_executor/kernels/linear/
   commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config.txt
+
+- label: LM Eval Humming Act fp8/int8 (B200 - TEMPORARY)
+  key: lm-eval-humming-act-b200
+  timeout_in_minutes: 50
+  device: b200-k8s
+  optional: true
+  num_devices: 1
+  source_file_dependencies:
+  - vllm/model_executor/layers/quantization/humming.py
+  - vllm/model_executor/layers/quantization/utils/humming_utils.py
+  - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
+  - vllm/model_executor/layers/fused_moe/oracle/
+  - vllm/model_executor/kernels/linear/
+  commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-fp8.txt
+  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-int8.txt
 
 - label: LM Eval TurboQuant KV Cache
   key: lm-eval-turboquant-kv-cache
-  timeout_in_minutes: 75
+  timeout_in_minutes: 55
   device: h200_18gb
   source_file_dependencies:
   - vllm/model_executor/layers/quantization/turboquant/
@@ -208,7 +256,7 @@ steps:
 
 - label: GPQA Eval (GPT-OSS) (2xH100)
   key: gpqa-eval-gpt-oss-2xh100
-  timeout_in_minutes: 120
+  timeout_in_minutes: 35
   device: h100
   optional: true
   num_devices: 2
@@ -222,7 +270,7 @@ steps:
 
 - label: GPQA Eval (GPT-OSS) (2xB200)
   key: gpqa-eval-gpt-oss-2xb200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 30
   device: b200-k8s
   optional: true
   num_devices: 2
@@ -236,7 +284,7 @@ steps:
 
 - label: GPQA Eval (GPT-OSS) (DGX Spark)
   key: gpqa-eval-gpt-oss-spark
-  timeout_in_minutes: 120
+  timeout_in_minutes: 35
   device: dgx-spark
   optional: true
   num_devices: 1
@@ -250,9 +298,50 @@ steps:
     - uv pip install --system 'gpt-oss[eval]==0.0.5'
     - pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-spark.txt
 
+- label: LM Eval KV-Offload (1xH200)
+  key: kv-offload-small
+  timeout_in_minutes: 30
+  device: h200_35gb
+  source_file_dependencies:
+  - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
+  - vllm/distributed/kv_transfer/kv_connector/v1/simple_cpu_offload_connector.py
+  - vllm/v1/kv_offload/
+  - vllm/v1/simple_kv_offload/
+  - tests/evals/gsm8k/test_gsm8k_offloading.py
+  commands:
+  - pytest -s -v evals/gsm8k/test_gsm8k_offloading.py -k "nemotron-h-8b or gemma-4-e4b-it"
+
+- label: LM Eval KV-Offload (2xH100)
+  key: kv-offload-medium
+  timeout_in_minutes: 30
+  device: h100
+  num_devices: 2
+  source_file_dependencies:
+  - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
+  - vllm/distributed/kv_transfer/kv_connector/v1/simple_cpu_offload_connector.py
+  - vllm/v1/kv_offload/
+  - vllm/v1/simple_kv_offload/
+  - tests/evals/gsm8k/test_gsm8k_offloading.py
+  commands:
+  - pytest -s -v evals/gsm8k/test_gsm8k_offloading.py -k "qwen3.5-35b"
+
+- label: LM Eval KV-Offload (4xH100)
+  key: kv-offload-large
+  timeout_in_minutes: 40
+  device: h100
+  num_devices: 4
+  source_file_dependencies:
+  - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
+  - vllm/distributed/kv_transfer/kv_connector/v1/simple_cpu_offload_connector.py
+  - vllm/v1/kv_offload/
+  - vllm/v1/simple_kv_offload/
+  - tests/evals/gsm8k/test_gsm8k_offloading.py
+  commands:
+  - pytest -s -v evals/gsm8k/test_gsm8k_offloading.py -k "deepseek-v4-flash"
+
 - label: MRCR Eval Small Models
   device: h200_35gb
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
   source_file_dependencies:
   - tests/evals/mrcr/
   commands:

.buildkite/test_areas/lora.yaml

@@ -5,7 +5,7 @@ steps:
 - label: LoRA %N
   device: h200_35gb
   key: lora
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   source_file_dependencies:
   - vllm/lora
   - tests/lora
@@ -16,7 +16,7 @@ steps:
     amd:
       device: mi325_1
       working_dir: "/vllm-workspace/tests"
-      timeout_in_minutes: 60
+      timeout_in_minutes: 65
       source_file_dependencies:
       - vllm/lora
       - tests/lora
@@ -27,7 +27,7 @@ steps:
 
 - label: LoRA TP (Distributed)
   key: lora-tp-distributed
-  timeout_in_minutes: 30
+  timeout_in_minutes: 60
   num_devices: 4
   source_file_dependencies:
   - vllm/lora

.buildkite/test_areas/misc.yaml

@@ -5,7 +5,7 @@ steps:
 - label: V1 Spec Decode
   device: h200_35gb
   key: v1-spec-decode
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   source_file_dependencies:
     - vllm/config/
     - vllm/distributed/
@@ -24,13 +24,13 @@ steps:
   mirror:
     amd:
       device: mi300_1
-      timeout_in_minutes: 65
+      timeout_in_minutes: 75
       depends_on:
       - image-build-amd
 
 - label: V1 Sample + Logits
   key: v1-sample-logits
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   device: h200_18gb
   source_file_dependencies:
     - vllm/config/
@@ -64,7 +64,7 @@ steps:
 
 - label: V1 Core + KV + Metrics
   key: v1-core-kv-metrics
-  timeout_in_minutes: 30
+  timeout_in_minutes: 60
   source_file_dependencies:
     - vllm/config/
     - vllm/distributed/
@@ -103,12 +103,12 @@ steps:
     - pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
     - pytest -v -s -m 'not cpu_test' v1/metrics
     # Integration test for streaming correctness (requires special branch).
-    - pip install -U git+https://github.com/robertgshaw2-redhat/lm-evaluation-harness.git@streaming-api
+    - pip install -U git+https://github.com/vllm-project/lm-evaluation-harness.git@streaming-api
     - pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 75
       depends_on:
       - image-build-amd
 
@@ -172,7 +172,7 @@ steps:
 
 - label: Regression
   key: regression
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
   - vllm/config/
@@ -195,7 +195,7 @@ steps:
 - label: Examples
   device: h200_35gb
   key: examples
-  timeout_in_minutes: 45
+  timeout_in_minutes: 40
   working_dir: "/vllm-workspace/examples"
   source_file_dependencies:
   - vllm/entrypoints
@@ -237,7 +237,7 @@ steps:
 
 - label: Metrics, Tracing (2 GPUs)
   key: metrics-tracing-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 25
   num_devices: 2
   source_file_dependencies:
   - vllm/config/
@@ -281,7 +281,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 20
+      timeout_in_minutes: 45
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -292,7 +292,7 @@ steps:
 - label: Async Engine, Inputs, Utils, Worker
   device: h200_35gb
   key: async-engine-inputs-utils-worker
-  timeout_in_minutes: 50
+  timeout_in_minutes: 25
   source_file_dependencies:
   - vllm/assets/
   - vllm/config/
@@ -319,7 +319,7 @@ steps:
   key: async-engine-inputs-utils-worker-config-cpu
   depends_on:
   - image-build-cpu
-  timeout_in_minutes: 30
+  timeout_in_minutes: 65
   source_file_dependencies:
   - vllm/assets/
   - vllm/config/
@@ -351,6 +351,7 @@ steps:
   - tests/test_outputs.py
   - tests/test_pooling_params.py
   - tests/test_ray_env.py
+  - tests/test_sampling_params.py
   - tests/multimodal
   - tests/renderers
   - tests/standalone_tests/lazy_imports.py
@@ -368,6 +369,7 @@ steps:
   - pytest -v -s test_outputs.py
   - pytest -v -s test_pooling_params.py
   - pytest -v -s test_ray_env.py
+  - pytest -v -s test_sampling_params.py
   - pytest -v -s -m 'cpu_test' multimodal
   - pytest -v -s renderers
   - pytest -v -s reasoning
@@ -379,7 +381,7 @@ steps:
 
 - label: Batch Invariance (A100)
   key: batch-invariance-a100
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   device: a100
   source_file_dependencies:
     - vllm/v1/attention
@@ -393,7 +395,7 @@ steps:
 
 - label: Batch Invariance (H100)
   key: batch-invariance-h100
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   device: h100
   source_file_dependencies:
     - vllm/v1/attention
@@ -409,7 +411,7 @@ steps:
 
 - label: Batch Invariance (B200)
   key: batch-invariance-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 35
   device: b200-k8s
   source_file_dependencies:
     - vllm/v1/attention
@@ -428,7 +430,7 @@ steps:
 - label: Acceptance Length Test (Large Models) # optional
   device: h200_35gb
   key: acceptance-length-test-large-models
-  timeout_in_minutes: 25
+  timeout_in_minutes: 20
   gpu: h100
   optional: true
   num_gpus: 1

.buildkite/test_areas/model_executor.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Model Executor
   key: model-executor
-  timeout_in_minutes: 35
+  timeout_in_minutes: 45
   source_file_dependencies:
   - vllm/engine/arg_utils.py
   - vllm/config/model.py

.buildkite/test_areas/model_runner_v2.yaml

@@ -5,7 +5,7 @@ steps:
 - label: Model Runner V2 Core Tests
   device: h200_35gb
   key: model-runner-v2-core-tests
-  timeout_in_minutes: 45
+  timeout_in_minutes: 35
   source_file_dependencies:
   - vllm/v1/worker/gpu/
   - vllm/v1/worker/gpu_worker.py
@@ -27,7 +27,7 @@ steps:
 - label: Model Runner V2 Examples
   device: h200_35gb
   key: model-runner-v2-examples
-  timeout_in_minutes: 45
+  timeout_in_minutes: 35
   working_dir: "/vllm-workspace/examples"
   source_file_dependencies:
     - vllm/v1/worker/gpu/
@@ -63,7 +63,7 @@ steps:
 
 - label: Model Runner V2 Distributed (2 GPUs)
   key: model-runner-v2-distributed-2-gpus
-  timeout_in_minutes: 45
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -84,7 +84,7 @@ steps:
 
 - label: Model Runner V2 Pipeline Parallelism (4 GPUs)
   key: model-runner-v2-pipeline-parallelism-4-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:

.buildkite/test_areas/models_basic.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Basic Models Tests (Initialization)
   key: basic-models-tests-initialization
-  timeout_in_minutes: 45
+  timeout_in_minutes: 25
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -17,7 +17,7 @@ steps:
 - label: Basic Models Tests (Extra Initialization) %N
   device: h200_35gb
   key: basic-models-tests-extra-initialization
-  timeout_in_minutes: 45
+  timeout_in_minutes: 100
   source_file_dependencies:
   - vllm/model_executor/models/
   - tests/models/test_initialization.py
@@ -32,14 +32,14 @@ steps:
 - label: Basic Models Tests (Other)
   device: h200_35gb
   key: basic-models-tests-other
-  timeout_in_minutes: 45
+  timeout_in_minutes: 35
   source_file_dependencies:
   - vllm/
   - tests/models/test_terratorch.py
-  - tests/models/test_transformers.py
+  - tests/models/transformers/test_backend.py
   - tests/models/test_registry.py
   commands:
-    - pytest -v -s models/test_terratorch.py models/test_transformers.py models/test_registry.py
+    - pytest -v -s models/test_terratorch.py models/transformers/test_backend.py models/test_registry.py
   mirror:
     amd:
       device: mi325_1
@@ -50,11 +50,12 @@ steps:
   key: basic-models-test-other-cpu
   depends_on:
   - image-build-cpu
-  timeout_in_minutes: 10
+  timeout_in_minutes: 20
   source_file_dependencies:
   - vllm/
   - tests/models/test_utils.py
   - tests/models/test_vision.py
+  - tests/models/transformers/fusers/
   device: cpu-small
   commands:
-    - pytest -v -s models/test_utils.py models/test_vision.py
+    - pytest -v -s models/test_utils.py models/test_vision.py models/transformers/fusers/

.buildkite/test_areas/models_distributed.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed Model Tests (2 GPUs)
   key: distributed-model-tests-2-gpus
-  timeout_in_minutes: 50
+  timeout_in_minutes: 60
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -17,7 +17,7 @@ steps:
   - 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)'
   # Avoid importing model tests that cause CUDA reinitialization error
-  - pytest models/test_transformers.py -v -s -m 'distributed(num_gpus=2)'
+  - pytest models/transformers/test_backend.py -v -s -m 'distributed(num_gpus=2)'
   - pytest models/language -v -s -m 'distributed(num_gpus=2)'
   - pytest models/multimodal/generation/test_phi4siglip.py -v -s -m 'distributed(num_gpus=2)'
   - pytest models/multimodal -v -s -m 'distributed(num_gpus=2)' --ignore models/multimodal/generation/test_whisper.py --ignore models/multimodal/generation/test_phi4siglip.py

.buildkite/test_areas/models_language.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Language Models Tests (Standard)
   key: language-models-tests-standard
-  timeout_in_minutes: 25
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -21,7 +21,7 @@ steps:
 
 - label: Language Models Tests (Extra Standard) %N
   key: language-models-tests-extra-standard
-  timeout_in_minutes: 45
+  timeout_in_minutes: 40
   source_file_dependencies:
   - vllm/model_executor/models/
   - tests/models/language/pooling/test_embedding.py
@@ -52,7 +52,7 @@ steps:
 
 - label: Language Models Tests (Hybrid) %N
   key: language-models-tests-hybrid
-  timeout_in_minutes: 75
+  timeout_in_minutes: 65
   source_file_dependencies:
   - vllm/
   - tests/models/language/generation
@@ -67,7 +67,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 90
+      timeout_in_minutes: 70
       depends_on:
       - image-build-amd
       commands:
@@ -78,7 +78,7 @@ steps:
 - label: Language Models Test (Extended Generation) # 80min
   device: h200_35gb
   key: language-models-test-extended-generation
-  timeout_in_minutes: 110
+  timeout_in_minutes: 65
   optional: true
   source_file_dependencies:
   - vllm/
@@ -92,7 +92,7 @@ steps:
 
 - label: Language Models Test (PPL)
   key: language-models-test-ppl
-  timeout_in_minutes: 110
+  timeout_in_minutes: 30
   device: h200_18gb
   optional: true
   source_file_dependencies:
@@ -104,7 +104,7 @@ steps:
 - label: Language Models Test (Extended Pooling)  # 36min
   device: h200_35gb
   key: language-models-test-extended-pooling
-  timeout_in_minutes: 50
+  timeout_in_minutes: 70
   optional: true
   source_file_dependencies:
   - vllm/
@@ -120,7 +120,7 @@ steps:
 
 - label: Language Models Test (MTEB)
   key: language-models-test-mteb
-  timeout_in_minutes: 110
+  timeout_in_minutes: 45
   device: h200_18gb
   optional: true
   source_file_dependencies:

.buildkite/test_areas/models_multimodal.yaml

@@ -20,13 +20,14 @@ steps:
 
 - label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
   key: multi-modal-models-standard-2-qwen3-gemma
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
   device: h200_18gb
   source_file_dependencies:
   - vllm/
   - tests/models/multimodal
   commands:
     - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen3 or gemma"
+    - pytest -v -s models/multimodal/generation/test_mm_prefix_lm.py -m core_model
     - pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model
   mirror:
     amd:
@@ -37,7 +38,7 @@ steps:
 - label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
   device: h200_35gb
   key: multi-modal-models-standard-3-llava-qwen2-vl
-  timeout_in_minutes: 45
+  timeout_in_minutes: 40
   source_file_dependencies:
   - vllm/
   - tests/models/multimodal
@@ -53,12 +54,12 @@ steps:
 - label: "Multi-Modal Models (Standard) 4: other + whisper"
   device: h200_35gb
   key: multi-modal-models-standard-4-other-whisper
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
   source_file_dependencies:
   - vllm/
   - tests/models/multimodal
   commands:
-    - pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py  --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/generation/test_vit_cudagraph.py --ignore models/multimodal/processing
+    - pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py --ignore models/multimodal/generation/test_mm_prefix_lm.py --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/generation/test_vit_cudagraph.py --ignore models/multimodal/processing
     - pytest -v -s models/multimodal/generation/test_vit_cudagraph.py -m core_model
     - pytest models/multimodal/generation/test_memory_leak.py -m core_model
     - cd .. && VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s tests/models/multimodal/generation/test_whisper.py -m core_model  # Otherwise, mp_method="spawn" doesn't work
@@ -72,7 +73,7 @@ steps:
   key: multi-modal-processor-cpu
   depends_on:
   - image-build-cpu
-  timeout_in_minutes: 60
+  timeout_in_minutes: 125
   source_file_dependencies:
   - vllm/
   - tests/models/multimodal
@@ -83,7 +84,7 @@ steps:
 
 - label: Multi-Modal Processor # 44min
   key: multi-modal-processor
-  timeout_in_minutes: 60
+  timeout_in_minutes: 65
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -95,7 +96,7 @@ steps:
 - label: Multi-Modal Accuracy Eval (Small Models) # 50min
   device: h200_35gb
   key: multi-modal-accuracy-eval-small-models
-  timeout_in_minutes: 70
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
   source_file_dependencies:
   - vllm/multimodal/
@@ -163,7 +164,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 75
       depends_on:
       - image-build-amd
       source_file_dependencies:

.buildkite/test_areas/plugins.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Plugin Tests (2 GPUs)
   key: plugin-tests-2-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 35
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -37,6 +37,11 @@ steps:
   - pytest -v -s plugins_tests/test_stats_logger_plugins.py
   - pip uninstall dummy_stat_logger -y
   # end stat_logger plugins test
+  # begin endpoint plugins test
+  - pip install -e ./plugins/vllm_add_dummy_endpoint_plugin
+  - pytest -v -s plugins_tests/test_endpoint_plugins.py
+  - pip uninstall vllm_add_dummy_endpoint_plugin -y
+  # end endpoint plugins test
   # other tests continue here:
   - pytest -v -s plugins_tests/test_scheduler_plugins.py
   - pip install -e ./plugins/vllm_add_dummy_model

.buildkite/test_areas/pytorch.yaml

@@ -5,7 +5,7 @@ steps:
 - label: PyTorch Compilation Unit Tests
   device: h200_35gb
   key: pytorch-compilation-unit-tests
-  timeout_in_minutes: 10
+  timeout_in_minutes: 90
   source_file_dependencies:
     - vllm/__init__.py
     - vllm/_aiter_ops.py
@@ -78,7 +78,7 @@ steps:
 
 - label: PyTorch Compilation Passes Unit Tests
   key: pytorch-compilation-passes-unit-tests
-  timeout_in_minutes: 20
+  timeout_in_minutes: 45
   source_file_dependencies:
     - vllm/__init__.py
     - vllm/_aiter_ops.py
@@ -110,13 +110,13 @@ steps:
   mirror:
     amd:
       device: mi300_1
-      timeout_in_minutes: 180
+      timeout_in_minutes: 65
       depends_on:
       - image-build-amd
 
 - label: PyTorch Fullgraph Smoke Test
   key: pytorch-fullgraph-smoke-test
-  timeout_in_minutes: 35
+  timeout_in_minutes: 60
   source_file_dependencies:
   - vllm/__init__.py
   - vllm/_aiter_ops.py
@@ -152,7 +152,7 @@ steps:
 
 - label: PyTorch Fullgraph
   key: pytorch-fullgraph
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   device: h200_18gb
   source_file_dependencies:
   - vllm/__init__.py

.buildkite/test_areas/quantization.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Quantization
   key: quantization
-  timeout_in_minutes: 90
+  timeout_in_minutes: 60
   source_file_dependencies:
   - csrc/
   - vllm/model_executor/layers/quantization
@@ -23,7 +23,7 @@ steps:
 
 - label: Quantized Fusions
   key: quantized-fusions
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   source_file_dependencies:
   - tests/fusion
   - vllm/model_executor/layers/fusion
@@ -35,7 +35,7 @@ steps:
 
 - label: Quantized MoE Test (B200)
   key: quantized-moe-test-b200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 120
   working_dir: "/vllm-workspace/"
   device: b200-k8s
   source_file_dependencies:
@@ -53,7 +53,7 @@ steps:
 
 - label: Quantized Models Test
   key: quantized-models-test
-  timeout_in_minutes: 60
+  timeout_in_minutes: 50
   source_file_dependencies:
   - vllm/model_executor/layers/quantization
   - tests/models/quantization

.buildkite/test_areas/rust_frontend.yaml

@@ -3,7 +3,7 @@ depends_on:
   - image-build
 steps:
 - label: Rust Frontend OpenAI Coverage
-  timeout_in_minutes: 90
+  timeout_in_minutes: 30
   device: h200_18gb
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -15,28 +15,30 @@ steps:
   - tests/utils.py
   - tests/benchmarks/test_serve_cli.py
   - tests/entrypoints/openai/chat_completion/test_chat_completion.py
-  # - tests/entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py
+  - tests/entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py
 
   # - tests/entrypoints/openai/completion/test_prompt_validation.py
   - tests/entrypoints/openai/completion/test_shutdown.py
-  # - tests/entrypoints/openai/test_return_token_ids.py
-  # - tests/entrypoints/openai/test_uds.py
+  - tests/entrypoints/openai/test_return_token_ids.py
+  - tests/entrypoints/openai/test_uds.py
   - tests/v1/sample/test_logprobs_e2e.py
   commands:
   - export VLLM_USE_RUST_FRONTEND=1
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - pytest -v -s benchmarks/test_serve_cli.py -k "not insecure and not (test_bench_serve and not test_bench_serve_chat)"
   - pytest -v -s entrypoints/openai/chat_completion/test_chat_completion.py -k "not test_invalid_json_schema and not test_invalid_regex"
-  # - pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not invalid"
+  - pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not multiple"
 
   # - pytest -v -s entrypoints/openai/completion/test_prompt_validation.py -k "not prompt_embeds"
   - pytest -v -s entrypoints/openai/completion/test_shutdown.py -k "not engine_failure and not test_abort_timeout_exits_quickly"
-  # - pytest -v -s entrypoints/openai/test_return_token_ids.py
-  # - pytest -v -s entrypoints/openai/test_uds.py
+  # test_comparison streams differently: Rust emits a separate first (prompt_token_ids) chunk and
+  # finish chunk without logprobs, while the test reads `logprobs.tokens` on every chunk.
+  - pytest -v -s entrypoints/openai/test_return_token_ids.py -k "not test_comparison"
+  - pytest -v -s entrypoints/openai/test_uds.py
   - pytest -v -s v1/sample/test_logprobs_e2e.py -k "test_prompt_logprobs_e2e_server"
 
 - label: Rust Frontend Serve/Admin Coverage
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
   device: h200_18gb
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -45,22 +47,27 @@ steps:
   - vllm/entrypoints/serve/
   - vllm/v1/engine/
   - tests/utils.py
-  # - tests/entrypoints/serve/dev/rpc/test_collective_rpc.py
+  - tests/entrypoints/serve/dev/rpc/test_collective_rpc.py
   - tests/entrypoints/scale_out/token_in_token_out/test_serving_tokens.py
   - tests/entrypoints/serve/instrumentator/test_basic.py
   - tests/entrypoints/serve/instrumentator/test_metrics.py
   # - tests/entrypoints/serve/dev/test_sleep.py
+  - tests/entrypoints/serve/tokenize/test_tokenization.py
   commands:
   - export VLLM_USE_RUST_FRONTEND=1
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  # - pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
+  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
+  # server_load can be flaky under the Rust frontend; keep it excluded for now.
   - pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not show_version and not server_load"
+  # test_generate_logprobs expects Python-style top_logprobs truncation (dedup sampled + cap at max(k, 1)).
   - pytest -v -s entrypoints/scale_out/token_in_token_out/test_serving_tokens.py -k "not stream and not lora and not test_generate_logprobs and not stop_string_workflow"
   - pytest -v -s entrypoints/serve/instrumentator/test_metrics.py -k "text and not show and not run_batch and not test_metrics_counts and not test_metrics_exist"
   # - pytest -v -s entrypoints/serve/dev/test_sleep.py
+  # /tokenizer_info is not implemented in the Rust frontend (the CLI flag is accepted as a no-op).
+  - pytest -v -s entrypoints/serve/tokenize/test_tokenization.py -k "not tokenizer_info"
 
 - label: Rust Frontend Core Correctness
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   device: h200_18gb
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -74,7 +81,7 @@ steps:
   - pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
 
 - label: Rust Frontend Tool Use
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - rust/
@@ -88,7 +95,7 @@ steps:
   - pytest -v -s tool_use --ignore=tool_use/mistral --models llama3.2 -k "not test_response_format_with_tool_choice_required and not test_parallel_tool_calls_false and not test_tool_call_and_choice"
 
 - label: Rust Frontend Distributed
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
   num_devices: 4
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:

.buildkite/test_areas/rust_frontend_cargo.yaml

@@ -4,7 +4,7 @@ steps:
 - label: Rust Frontend Cargo Style + Clippy
   key: rust-frontend-cargo-style-clippy
   depends_on: []
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   device: cpu-medium
   no_plugin: true
   source_file_dependencies:
@@ -18,7 +18,7 @@ steps:
 - label: Rust Frontend Cargo Tests
   key: rust-frontend-cargo-tests
   depends_on: []
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   device: cpu-medium
   no_plugin: true
   source_file_dependencies:

.buildkite/test_areas/samplers.yaml

@@ -5,7 +5,7 @@ steps:
 - label: Samplers Test
   device: h200_35gb
   key: samplers-test
-  timeout_in_minutes: 75
+  timeout_in_minutes: 40
   source_file_dependencies:
   - vllm/model_executor/layers
   - vllm/sampling_metadata.py
@@ -19,7 +19,7 @@ steps:
     - VLLM_USE_FLASHINFER_SAMPLER=1 pytest -v -s samplers
   mirror:
     amd:
-      device: mi250_1
+      device: mi325_1
       depends_on:
       - image-build-amd
       commands:

.buildkite/test_areas/spec_decode.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Spec Decode Eagle
   key: spec-decode-eagle
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -15,7 +15,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 45
+      timeout_in_minutes: 60
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -29,7 +29,7 @@ steps:
 
 - label: Spec Decode Eagle Nightly B200
   key: spec-decode-eagle-nightly-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
   device: b200-k8s
   optional: true
   source_file_dependencies:
@@ -41,7 +41,7 @@ steps:
 
 - label: Spec Decode Speculators + MTP
   key: spec-decode-speculators-mtp
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -82,7 +82,7 @@ steps:
   
 - label: Spec Decode Ngram + Suffix
   key: spec-decode-ngram-suffix
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -93,7 +93,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 65
+      timeout_in_minutes: 55
       # TODO(akaratza): Test after Torch >= 2.12 bump
       soft_fail: true
       depends_on:
@@ -109,7 +109,7 @@ steps:
 
 - label: Spec Decode Draft Model
   key: spec-decode-draft-model
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -120,7 +120,7 @@ steps:
   mirror:
     amd:
       device: mi325_1
-      timeout_in_minutes: 50
+      timeout_in_minutes: 55
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -134,7 +134,7 @@ steps:
 
 - label: Spec Decode Draft Model Nightly B200
   key: spec-decode-draft-model-nightly-b200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 40
   device: b200-k8s
   optional: true
   source_file_dependencies:
@@ -146,7 +146,7 @@ steps:
 
 - label: Speculators Correctness
   key: speculators-correctness
-  timeout_in_minutes: 60
+  timeout_in_minutes: 30
   device: h100
   optional: true
   num_devices: 1
@@ -159,7 +159,7 @@ steps:
     - pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test
 
 - label: Spec Decode MTP hybrid (B200)
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   device: b200-k8s
   optional: true
   source_file_dependencies:

.buildkite/test_areas/weight_loading.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Weight Loading Multiple GPU  # 33min
   key: weight-loading-multiple-gpu
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   optional: true

.github/CODEOWNERS

@@ -119,7 +119,7 @@
 
 # Transformers modeling backend
 /vllm/model_executor/models/transformers @hmellor
-/tests/models/test_transformers.py @hmellor
+/tests/models/transformers @hmellor
 
 # Docs
 /docs/mkdocs @hmellor

.github/workflows/pre-commit.yml

@@ -28,7 +28,8 @@ jobs:
             pull_number: context.payload.pull_request.number,
           });
 
-          const hasReadyLabel = pr.labels.some(l => l.name === 'ready');
+          const readyLabels = ['ready', 'ready-run-all-tests'];
+          const hasReadyLabel = pr.labels.some(l => readyLabels.includes(l.name));
           const hasVerifiedLabel = pr.labels.some(l => l.name === 'verified');
 
           const { data: mergedPRs } = await github.rest.search.issuesAndPullRequests({
@@ -40,7 +41,7 @@ jobs:
           if (hasReadyLabel || hasVerifiedLabel || mergedCount >= 4) {
             core.info(`Check passed: verified label=${hasVerifiedLabel}, ready label=${hasReadyLabel}, 4+ merged PRs=${mergedCount >= 4}`);
           } else {
-            core.setFailed(`PR must have the 'verified' or 'ready' (which also triggers tests) label or the author must have at least 4 merged PRs (found ${mergedCount}).`);
+            core.setFailed(`PR must have the 'verified', 'ready', or 'ready-run-all-tests' label (the ready labels also trigger tests) or the author must have at least 4 merged PRs (found ${mergedCount}).`);
           }
 
   pre-commit:

AGENTS.md

@@ -29,6 +29,7 @@ Do not open one-off PRs for tiny edits (single typo, isolated style change, one
 - PR descriptions for AI-assisted work **must** include:
     - Why this is not duplicating an existing PR.
     - Test commands run and results.
+    - Model evaluation results when the change affects output, accuracy, or serving.
     - Clear statement that AI assistance was used.
 
 ### Fail-closed behavior
@@ -66,23 +67,38 @@ VLLM_USE_PRECOMPILED=1 uv pip install -e . --torch-backend=auto
 uv pip install -e . --torch-backend=auto
 ```
 
-### Running tests
+### Tests
 
 > Requires [Environment setup](#environment-setup) and [Installing dependencies](#installing-dependencies).
 
 ```bash
-# Install test dependencies.
-# requirements/test/cuda.txt is pinned to x86_64; on other platforms, use the
-# unpinned source file instead:
-uv pip install -r requirements/test/cuda.in    # resolves for current platform
-# Or on x86_64:
-uv pip install -r requirements/test/cuda.txt
+# Install test dependencies (use cuda.in on non-x86_64):
+uv pip install -r requirements/test/cuda.in
 
-# Run a specific test file (use .venv/bin/python directly;
-# `source activate` does not persist in non-interactive shells):
+# Run a specific test file:
 .venv/bin/python -m pytest tests/path/to/test_file.py -v
 ```
 
+When adding tests:
+
+- **Design before you write.** Answer four questions first: what is the module
+  for, what is its I/O contract, what failure am I guarding against, and what is
+  the cheapest level that catches it (unit over integration over e2e)?
+- **Reuse before create.** Extend existing test files, `conftest.py` fixtures, and
+  helpers; add a new file only when no nearby suite fits.
+- **Test behavior with intent.** Assert observable outcomes through public APIs;
+  state why in the name or docstring. Skip trivial wiring; flaky tests are worse
+  than no tests.
+- **Keep it minimal.** One behavior per test and the smallest setup that
+  triggers it; if the test diff dwarfs the code change, cut scope.
+- **No one-off kernel benchmarks in `tests/`.** Put kernel perf work in
+  `benchmarks/kernels/`; prove correctness in existing pytest suites.
+- **Run model evals for model-affecting changes.** Search `tests/evals/` or use
+  `vllm bench` and include results in the PR — do not wait for reviewers to ask.
+
+For model-specific requirements, see
+[`docs/contributing/model/tests.md`](docs/contributing/model/tests.md).
+
 ### Running linters
 
 > Requires [Environment setup](#environment-setup).
@@ -107,23 +123,18 @@ Use [Google-style docstrings](https://google.github.io/styleguide/pyguide.html#3
 
 ### Coding style guidelines
 
-Follow these rules for all code changes in this repository:
-
-- Try to match existing code style.
-- Code should be self-documenting and self-explanatory.
-- Keep comments and docstrings minimal and concise.
+- Match existing code style
+- Minimize use of comments. Eliminate comments which are redundant, preferring legible and self-documenting code. When used, keep docstrings and comments brief and direct.
 - Assume the reader is familiar with vLLM.
 
 ### Commit messages
 
-Add attribution using commit trailers such as `Co-authored-by:` (other projects use `Assisted-by:` or `Generated-by:`). For example:
+Add attribution using commit trailers such as `Co-authored-by:` (other projects use `Assisted-by:` or `Generated-by:`):
 
 ```text
 Your commit message here
 
-Co-authored-by: GitHub Copilot
-Co-authored-by: Claude
-Co-authored-by: gemini-code-assist
+Co-authored-by: Agent Name Here
 Signed-off-by: Your Name <your.email@example.com>
 ```
 

CMakeLists.txt

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

benchmarks/backend_request_func.py

@@ -12,7 +12,7 @@ from dataclasses import dataclass, field
 import aiohttp
 import huggingface_hub.constants
 from tqdm.asyncio import tqdm
-from transformers import AutoTokenizer, PreTrainedTokenizer, PreTrainedTokenizerFast
+from transformers import AutoTokenizer, PythonBackend, TokenizersBackend
 
 # NOTE(simon): do not import vLLM here so the benchmark script
 # can run without vLLM installed.
@@ -609,7 +609,7 @@ def get_tokenizer(
     tokenizer_mode: str = "auto",
     trust_remote_code: bool = False,
     **kwargs,
-) -> PreTrainedTokenizer | PreTrainedTokenizerFast:
+) -> PythonBackend | TokenizersBackend:
     if pretrained_model_name_or_path is not None and not os.path.exists(
         pretrained_model_name_or_path
     ):

benchmarks/benchmark_topk_topp.py

@@ -132,8 +132,10 @@ def benchmark_function(
     reset_memory_stats()
 
     # Benchmark
-    start_events = [torch.Event(enable_timing=True) for _ in range(benchmark_iters)]
-    end_events = [torch.Event(enable_timing=True) for _ in range(benchmark_iters)]
+    start_events = [
+        torch.cuda.Event(enable_timing=True) for _ in range(benchmark_iters)
+    ]
+    end_events = [torch.cuda.Event(enable_timing=True) for _ in range(benchmark_iters)]
 
     for i in range(benchmark_iters):
         logits_copy = logits.clone()

benchmarks/kernels/benchmark_flydsl_moe_w4a16.py

@@ -17,7 +17,7 @@ from vllm.model_executor.layers.fused_moe.fused_flydsl_moe import fused_flydsl_m
 from vllm.model_executor.layers.quantization.compressed_tensors.compressed_tensors_moe import (  # noqa: E501
     compressed_tensors_moe_w4a16_flydsl,
 )
-from vllm.platforms import current_platform
+from vllm.utils.platform_utils import get_device_name_as_file_name
 
 RoutingBuffers = tuple[
     torch.Tensor,  # sorted_token_ids
@@ -259,7 +259,7 @@ def tune_flydsl_moe_w4a16(
                             )
                             us_best = us
                             tuned_config[str(num_tokens)] = tile_config
-            device_name = current_platform.get_device_name().replace(" ", "_")
+            device_name = get_device_name_as_file_name()
             tuned_config_file_name = (
                 f"E={num_experts},N={inter_dim},device_name={device_name},"
                 f"dtype=int4_w4a16,backend=flydsl.json"

benchmarks/kernels/benchmark_moe_defaults.py

@@ -134,8 +134,8 @@ def benchmark_config(
     torch.accelerator.synchronize()
 
     # Benchmark
-    start = torch.Event(enable_timing=True)
-    end = torch.Event(enable_timing=True)
+    start = torch.cuda.Event(enable_timing=True)
+    end = torch.cuda.Event(enable_timing=True)
     start.record()
     for _ in range(num_iters):
         with override_config(config):

benchmarks/kernels/benchmark_selective_state_update.py

@@ -170,8 +170,8 @@ def benchmark_config(
                 graph.replay()
             torch.accelerator.synchronize()
 
-            start = torch.Event(enable_timing=True)
-            end = torch.Event(enable_timing=True)
+            start = torch.cuda.Event(enable_timing=True)
+            end = torch.cuda.Event(enable_timing=True)
             latencies: list[float] = []
             for _ in range(num_iters):
                 start.record()

benchmarks/kernels/benchmark_w8a8_block_fp8.py

@@ -19,6 +19,7 @@ from vllm.model_executor.layers.quantization.utils.fp8_utils import (
 from vllm.platforms import current_platform
 from vllm.triton_utils import triton
 from vllm.utils.argparse_utils import FlexibleArgumentParser
+from vllm.utils.platform_utils import get_device_name_as_file_name
 
 mp.set_start_method("spawn", force=True)
 
@@ -264,7 +265,7 @@ def save_configs(
     input_type="fp8",
 ) -> None:
     os.makedirs(save_path, exist_ok=True)
-    device_name = current_platform.get_device_name().replace(" ", "_")
+    device_name = get_device_name_as_file_name()
     json_file_name = (
         f"N={N},K={K},device_name={device_name},dtype={input_type}_w8a8,"
         f"block_shape=[{block_n},{block_k}].json"

build_vllm_ppc64le.sh

@@ -0,0 +1,241 @@
+#!/bin/bash
+set -eoux pipefail
+
+########################################
+# Resolve repo root (IMPORTANT)
+########################################
+REPO_ROOT="$(pwd)"
+
+cd "$REPO_ROOT"
+
+########################################
+# DevPI configuration
+########################################
+
+IBM_DEVPI_URL=${IBM_DEVPI_URL:-"https://wheels.developerfirst.ibm.com/ppc64le/linux/+simple/"}
+RHOAI_INDEX_URL=${RHOAI_INDEX_URL:-"https://console.redhat.com/api/pypi/public-rhai/rhoai/3.4/cpu-ubi9/simple/"}
+
+########################################
+# wheel dir
+########################################
+
+WHEEL_DIR=${WHEEL_DIR:-"/tmp/wheels"}
+mkdir -p "$WHEEL_DIR"
+
+########################################
+# Helpers
+########################################
+try_install_from_devpi() {
+    local pkg=$1
+    uv pip install \
+        --extra-index-url "${IBM_DEVPI_URL}" \
+        --index-strategy unsafe-best-match \
+        --no-build-isolation \
+        "${pkg}"
+}
+
+########################################
+# Package Versions
+########################################
+cd "$REPO_ROOT"
+TORCH_VERSION=${TORCH_VERSION:-$(grep -E '^torch==.+==\s*"ppc64le"' requirements/cpu.txt | grep -Eo '\b[0-9\.]+\b' || true)}
+TORCH_VERSION=${TORCH_VERSION:-2.11.0}
+
+TORCHVISION_VERSION=${TORCHVISION_VERSION:-0.26.0}
+TORCHAUDIO_VERSION=${TORCHAUDIO_VERSION:-${TORCH_VERSION}}
+
+export TORCH_VERSION
+export TORCHVISION_VERSION
+export TORCHAUDIO_VERSION
+export OPENCV_VERSION=${OPENCV_VERSION:-4.13.0.92}
+export XGRAMMAR_VERSION=${XGRAMMAR_VERSION:-0.2.1}
+
+########################################
+# install system dependencies
+########################################
+
+rpm -ivh https://dl.fedoraproject.org/pub/epel/epel-release-latest-9.noarch.rpm || true
+
+microdnf install -y \
+    python3.12 python3.12-devel python3.12-pip gcc \
+    git jq gcc-toolset-14 gcc-toolset-14-libatomic-devel \
+    automake libtool clang-devel openssl-devel \
+    harfbuzz-devel kmod lcms2-devel libimagequant-devel libjpeg-turbo-devel \
+    llvm15-devel libraqm-devel libtiff-devel libwebp-devel libxcb-devel \
+    ninja-build openjpeg2-devel pkgconfig \
+    tcl-devel tk-devel xsimd-devel zeromq-devel zlib-devel patchelf file openblas openblas-devel protobuf numactl numactl-devel openmpi openmpi-devel
+    
+rpm -ivh --nodeps \
+    https://mirror.stream.centos.org/9-stream/CRB/ppc64le/os/Packages/protobuf-lite-devel-3.14.0-17.el9.ppc64le.rpm
+
+rpm -ivh --nodeps \
+    https://mirror.stream.centos.org/9-stream/CRB/ppc64le/os/Packages/protobuf-devel-3.14.0-17.el9.ppc64le.rpm
+
+rpm -ivh --nodeps \
+    https://mirror.stream.centos.org/9-stream/CRB/ppc64le/os/Packages/protobuf-compiler-3.14.0-17.el9.ppc64le.rpm
+
+########################################
+# Python 3.12 virtual environment
+########################################
+
+python3.12 -m venv /opt/vllm
+source /opt/vllm/bin/activate
+
+export PATH=/opt/vllm/bin:$PATH
+
+python --version
+
+########################################
+# install build tools (stable uv)
+########################################
+
+pip install -U pip setuptools-rust
+pip install uv
+pip install "setuptools<70" build wheel cmake auditwheel
+uv pip install "setuptools<70" cython meson-python pybind11 "sympy>=1.13.3" --no-build-isolation
+
+########################################
+# Rust
+########################################
+
+curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y
+source /root/.cargo/env
+
+########################################
+# Compiler env
+########################################
+
+source /opt/rh/gcc-toolset-14/enable
+
+export PATH=/usr/lib64/llvm15/bin:$PATH
+export LLVM_CONFIG=/usr/lib64/llvm15/bin/llvm-config
+export CMAKE_ARGS="-DPython3_EXECUTABLE=python"
+
+export MAX_JOBS=${MAX_JOBS:-$(nproc)}
+export GRPC_PYTHON_BUILD_SYSTEM_OPENSSL=1
+
+########################################
+# Install Packages From Devpi
+########################################
+uv pip install numpy==2.3.5 pillow==12.2.0 --extra-index-url "$IBM_DEVPI_URL"
+try_install_from_devpi "opencv-python-headless==${OPENCV_VERSION}"
+try_install_from_devpi "torch==${TORCH_VERSION}"
+try_install_from_devpi "torchvision==${TORCHVISION_VERSION}"
+
+########################################
+# torch audio
+########################################
+
+TEMP_BUILD_DIR=$(mktemp -d)
+cd "${TEMP_BUILD_DIR}"
+export BUILD_SOX=1 BUILD_KALDI=1 BUILD_RNNT=1 USE_FFMPEG=0 USE_ROCM=0 USE_CUDA=0
+export TORCHAUDIO_TEST_ALLOW_SKIP_IF_NO_FFMPEG=1
+git clone --recursive https://github.com/pytorch/audio.git -b v${TORCHAUDIO_VERSION}
+cd audio
+#patching 
+sed -i '
+s|_CSRC_DIR / "_torchaudio.cpp"|str(_CSRC_DIR / "_torchaudio.cpp")|;
+s|_CSRC_DIR / "utils.cpp"|str(_CSRC_DIR / "utils.cpp")|;
+s|sources=\[_CSRC_DIR / s for s in sources\]|sources=[str(_CSRC_DIR / s) for s in sources]|;
+' tools/setup_helpers/extension.py
+MAX_JOBS=${MAX_JOBS:-$(nproc)} \
+BUILD_VERSION=${TORCHAUDIO_VERSION} \
+uv build --wheel --out-dir "${WHEEL_DIR}" --no-build-isolation
+uv pip install "${WHEEL_DIR}"/torchaudio*.whl
+cd "${REPO_ROOT}"
+rm -rf "${TEMP_BUILD_DIR}"
+
+########################################
+# Xgrammar
+########################################
+uv pip install \
+   "scikit-build-core==0.11.6" \
+   "pyproject-metadata<0.8" \
+    pathspec \
+    packaging \
+    distro \
+   "setuptools<70" \
+    setuptools_scm \
+    cmake \
+    ninja \
+    pybind11 \
+    nanobind
+uv pip install apache-tvm-ffi==0.1.12 \
+  --no-build-isolation \
+  --no-cache
+
+TEMP_BUILD_DIR=$(mktemp -d)
+
+pushd "${TEMP_BUILD_DIR}"
+
+export CFLAGS="-fno-lto -mcpu=power9"
+export CXXFLAGS="-fno-lto -mcpu=power9"
+export LDFLAGS="-fno-lto"
+export PATH=/opt/vllm/bin:$PATH
+
+export Python_EXECUTABLE=/opt/vllm/bin/python3
+export Python3_EXECUTABLE=/opt/vllm/bin/python3
+export PYTHON_EXECUTABLE=/opt/vllm/bin/python3
+
+export Python_ROOT_DIR=/opt/vllm
+export Python3_ROOT_DIR=/opt/vllm
+
+git clone \
+    --recursive \
+    https://github.com/mlc-ai/xgrammar \
+    -b "v${XGRAMMAR_VERSION}"
+
+cd xgrammar
+
+cp cmake/config.cmake .
+export PYTHONPATH=/opt/vllm/lib64/python3.12/site-packages:/opt/vllm/lib/python3.12/site-packages:${PYTHONPATH:-}
+
+uv build \
+    --wheel \
+    --out-dir "${WHEEL_DIR}" \
+    --no-build-isolation
+
+uv pip install "${WHEEL_DIR}"/xgrammar*.whl -v
+
+popd
+
+rm -rf "${TEMP_BUILD_DIR}"
+cd "${REPO_ROOT}"
+
+########################################
+# RHOAI Binary Downloads
+########################################
+pip download \
+    --index-url "${RHOAI_INDEX_URL}" \
+    --only-binary=:all: \
+    --no-deps \
+    llvmlite==0.47.0 \
+    -d "${WHEEL_DIR}"
+
+pip download \
+    --index-url "${RHOAI_INDEX_URL}" \
+    --only-binary=:all: \
+    --no-deps \
+    Numba==0.65.0 \
+    -d "${WHEEL_DIR}"
+
+########################################
+# install built wheels
+########################################
+uv pip install setuptools_scm maturin setuptools-rust ninja scikit-build-core pybind11 nanobind \
+    --no-build-isolation
+uv pip install "${WHEEL_DIR}"/*.whl
+
+########################################
+# install remaining deps
+########################################
+
+sed -i.bak -e 's/.*torch.*//g' pyproject.toml requirements/*.txt
+
+uv pip install "setuptools>=78.1.1" --no-build-isolation
+
+export PKG_CONFIG_PATH=/usr/local/lib/pkgconfig:/usr/local/lib64/pkgconfig:/usr/lib64/pkgconfig
+
+uv pip install -r requirements/common.txt \
+               -r requirements/cpu.txt \
+               -r requirements/build/cpu.txt --index-strategy unsafe-best-match

cmake/cpu_extension.cmake

@@ -15,6 +15,7 @@ endif()
 #
 set(ENABLE_X86_ISA $ENV{VLLM_CPU_X86})
 set(ENABLE_ARM_BF16 $ENV{VLLM_CPU_ARM_BF16})
+set(ENABLE_RVV_BF16 $ENV{VLLM_CPU_RVV_BF16})
 
 include_directories("${CMAKE_SOURCE_DIR}/csrc")
 
@@ -110,6 +111,13 @@ else()
         set(ARM_BF16_FOUND ON)
         message(STATUS "ARM BF16 support enabled via VLLM_CPU_ARM_BF16 environment variable")
     endif()
+    # Some kernels (e.g. Bianbu on Spacemit X100) do not report zvfbfmin
+    # in /proc/cpuinfo despite hardware support. VLLM_CPU_RVV_BF16=1
+    # overrides the detection result.
+    if (ENABLE_RVV_BF16)
+        set(RVV_BF16_FOUND ON)
+        message(STATUS "RVV BF16 support enabled via VLLM_CPU_RVV_BF16 environment variable")
+    endif()
 endif()
 
 if (CMAKE_SYSTEM_PROCESSOR MATCHES "x86_64|amd64" OR ENABLE_X86_ISA)
@@ -178,7 +186,10 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
     # Override with -DVLLM_RVV_VLEN=128 or -DVLLM_RVV_VLEN=256 for RVV.
     if(NOT DEFINED VLLM_RVV_VLEN)
         # Auto-detect: find the largest zvl<N>b in /proc/cpuinfo isa line.
-        if(EXISTS /proc/cpuinfo)
+        # Skip when cross-compiling — /proc/cpuinfo describes the build host.
+        if(CMAKE_CROSSCOMPILING)
+            message(STATUS "Cross-compiling: skipping VLEN auto-detection from /proc/cpuinfo")
+        elseif(EXISTS /proc/cpuinfo)
             file(READ /proc/cpuinfo _cpuinfo)
             set(_best 0)
             foreach(_n IN ITEMS 128 256 512 1024)
@@ -186,6 +197,13 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
                     set(_best ${_n})
                 endif()
             endforeach()
+            # Only VLEN=128 and VLEN=256 are supported by the RVV kernels.
+            if(_best GREATER 256)
+                message(WARNING
+                    "Detected VLEN=${_best} but only 128/256 are supported; "
+                    "clamping to 256")
+                set(_best 256)
+            endif()
             if(_best GREATER 0)
                 set(VLLM_RVV_VLEN ${_best})
             endif()
@@ -195,9 +213,9 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
         if(NOT DEFINED VLLM_RVV_VLEN AND (RVV_FP16_FOUND OR RVV_BF16_FOUND))
             message(FATAL_ERROR
                 "RISC-V RVV is available but VLEN could not be auto-detected. "
-                "Please specify VLEN explicitly:\n"
-                "  -DVLLM_RVV_VLEN=128   (for VLEN=128 hardware)\n"
-                "  -DVLLM_RVV_VLEN=256   (for VLEN=256 hardware, e.g. Spacemit X100)")
+                "Please specify VLEN explicitly via CMAKE_ARGS:\n"
+                "  CMAKE_ARGS='-DVLLM_RVV_VLEN=128'   (for VLEN=128 hardware)\n"
+                "  CMAKE_ARGS='-DVLLM_RVV_VLEN=256'   (for VLEN=256 hardware, e.g. Spacemit X100)")
         endif()
     endif()
     if(VLLM_RVV_VLEN AND VLLM_RVV_VLEN GREATER 0)
@@ -209,7 +227,7 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
             message(STATUS "BF16 extension detected")
             set(MARCH_FLAGS -march=rv64gcv_zvfh_zfbfmin_zvfbfmin_zvl${VLLM_RVV_VLEN}b -mrvv-vector-bits=zvl -mabi=lp64d)
         elseif(RVV_FP16_FOUND)
-            message(WARNING "BF16 functionality is not available")
+            message(WARNING "BF16 functionality is not available.")
             set(MARCH_FLAGS -march=rv64gcv_zvfh_zvl${VLLM_RVV_VLEN}b -mrvv-vector-bits=zvl -mabi=lp64d)
         else()
             message(STATUS "compile riscv with scalar (no FP16/BF16)")

cmake/external_projects/fmha_sm100.cmake

@@ -17,7 +17,7 @@ else()
   FetchContent_Declare(
     fmha_sm100
     GIT_REPOSITORY https://github.com/vllm-project/MSA.git
-    GIT_TAG fee783153f3efe57e3e933c5cb7e267a7cebcfb5
+    GIT_TAG 2e63ec37a0fc29bc20f39cd1a52e0f5affc33a73
     GIT_PROGRESS TRUE
     CONFIGURE_COMMAND ""
     BUILD_COMMAND ""

cmake/external_projects/vllm_flash_attn.cmake

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

csrc/cpu/cpu_attn.cpp

@@ -13,7 +13,8 @@ static inline cpu_attention::Fp8KVCacheDataType parse_fp8_kv_dtype(
 
 bool cpu_attn_has_isa(const std::string& isa) {
   if (isa == "rvv") {
-#if defined(__riscv) && defined(__riscv_v_min_vlen) && __riscv_v_min_vlen == 128
+#if defined(__riscv) && defined(__riscv_v_min_vlen) && \
+    (__riscv_v_min_vlen == 128 || __riscv_v_min_vlen == 256)
     return true;
 #else
     return false;

csrc/cpu/cpu_attn_impl.hpp

@@ -417,8 +417,10 @@ class AttentionScheduler {
       has_decode_request = has_decode_request || (q_token_num == 1);
       decode_only_batch = decode_only_batch && (q_token_num == 1);
     }
-    int32_t q_head_per_kv = input.num_heads_q / input.num_heads_kv;
-    const bool supports_gqa = q_head_per_kv <= max_num_q_per_iter;
+    const int32_t original_q_head_per_kv =
+        input.num_heads_q / input.num_heads_kv;
+    int32_t q_head_per_kv = original_q_head_per_kv;
+    const bool supports_gqa = original_q_head_per_kv <= max_num_q_per_iter;
     const bool use_gqa_fast_path = supports_gqa && decode_only_batch;
     const bool use_gqa_scratchpad = supports_gqa && has_decode_request;
     if (!use_gqa_scratchpad) {
@@ -671,22 +673,62 @@ class AttentionScheduler {
     metadata_ptr->effective_thread_num = effective_thread_num;
 
     {
-      // when q_tile_size = max_num_q_per_iter, requires max
-      // attention_scratchpad_size
       AttentionScratchPad sc(0, *metadata_ptr, 0x0);
-      int64_t n = AttentionScheduler::calcu_tile_size_with_constant_q(
-          cache_size, input.head_dim, input.elem_size, input.q_buffer_elem_size,
-          input.logits_buffer_elem_size, input.output_buffer_elem_size,
-          max_num_q_per_iter, kv_len_alignment, max_num_q_per_iter, true);
-      sc.update(input.head_dim, input.q_buffer_elem_size,
-                input.logits_buffer_elem_size, input.output_buffer_elem_size,
-                max_num_q_per_iter, max_num_q_per_iter, n);
+      int64_t max_attention_scratchpad_size = 0;
+
+      for (const AttentionWorkItemGroup& item : workitems) {
+        const bool curr_use_gqa =
+            use_gqa_fast_path || (supports_gqa && item.q_token_num == 1);
+        const int32_t curr_q_heads_per_kv =
+            curr_use_gqa ? original_q_head_per_kv : 1;
+        const int32_t curr_default_q_tile_token_num =
+            default_tile_size / curr_q_heads_per_kv;
+
+        for (int32_t q_token_offset = 0; q_token_offset < item.q_token_num;
+             q_token_offset += curr_default_q_tile_token_num) {
+          const int32_t actual_q_token_num = std::min(
+              curr_default_q_tile_token_num, item.q_token_num - q_token_offset);
+          const int32_t q_head_tile_size =
+              actual_q_token_num * curr_q_heads_per_kv;
+          const int32_t rounded_q_head_tile_size =
+              ((q_head_tile_size + max_num_q_per_iter - 1) /
+               max_num_q_per_iter) *
+              max_num_q_per_iter;
+
+          const int64_t n = AttentionScheduler::calcu_tile_size_with_constant_q(
+              cache_size, input.head_dim, input.elem_size,
+              input.q_buffer_elem_size, input.logits_buffer_elem_size,
+              input.output_buffer_elem_size, max_num_q_per_iter,
+              kv_len_alignment, rounded_q_head_tile_size,
+              rounded_q_head_tile_size <= max_num_q_per_iter);
+
+          sc.update(input.head_dim, input.q_buffer_elem_size,
+                    input.logits_buffer_elem_size,
+                    input.output_buffer_elem_size, max_num_q_per_iter,
+                    rounded_q_head_tile_size, n);
+
+          max_attention_scratchpad_size = std::max(
+              max_attention_scratchpad_size, sc.get_thread_scratchpad_size());
+        }
+      }
+
       metadata_ptr->attention_scratchpad_size_per_thread =
-          ((sc.get_thread_scratchpad_size() + 63) / 64) * 64;
+          ((max_attention_scratchpad_size + 63) / 64) * 64;
+
+      int32_t max_reduction_q_head_tile_size = 0;
+      for (const ReductionWorkItemGroup& item : reduce_workitems) {
+        const bool curr_use_gqa =
+            use_gqa_fast_path || (supports_gqa && item.q_token_id_num == 1);
+        const int32_t curr_q_heads_per_kv =
+            curr_use_gqa ? original_q_head_per_kv : 1;
+
+        max_reduction_q_head_tile_size =
+            std::max(max_reduction_q_head_tile_size,
+                     item.q_token_id_num * curr_q_heads_per_kv);
+      }
 
       sc.update(0, metadata_ptr->reduction_split_num, input.head_dim,
-                q_head_per_kv * split_kv_q_token_num_threshold,
-                input.output_buffer_elem_size);
+                max_reduction_q_head_tile_size, input.output_buffer_elem_size);
       metadata_ptr->reduction_scratchpad_size_per_kv_head =
           ((sc.get_reduction_scratchpad_size() + 63) / 64) * 64;
     }

csrc/cpu/cpu_attn_vsx.hpp

@@ -323,8 +323,6 @@ class AttentionImpl<ISA::VSX, scalar_t, head_dim> {
       const int64_t num_blocks_stride, const int64_t cache_head_num_stride,
       const int64_t block_size, const int64_t block_size_stride,
       const float k_inv = 0.0f, const float v_inv = 0.0f) {
-    // k_inv and v_inv are unused on VSX: FP8 KV cache is not supported on
-    // PowerPC. The parameters are present to match the common interface.
 #pragma omp parallel for collapse(2)
     for (int64_t token_idx = 0; token_idx < token_num; ++token_idx) {
       for (int64_t head_idx = 0; head_idx < head_num; ++head_idx) {

csrc/cpu/cpu_types.hpp

@@ -4,7 +4,7 @@
 #if defined(__x86_64__)
   // x86 implementation
   #include "cpu_types_x86.hpp"
-#elif defined(__POWER9_VECTOR__)
+#elif defined(__powerpc__)
   // ppc implementation
   #include "cpu_types_vsx.hpp"
 #elif defined(__s390x__)
@@ -41,4 +41,4 @@ inline int get_max_threads() {
 }
 }  // namespace cpu_utils
 
-#endif
+#endif

csrc/cpu/cpu_types_riscv_impl.hpp

@@ -214,11 +214,18 @@ struct BF16Vec32 : public Vec<BF16Vec32> {
 
   explicit BF16Vec32(const BF16Vec8& v) {
     fixed_u16x8_t u16_val = bf16_to_u16(v.reg);
-    fixed_u16x32_t u16_combined =
-        RVVI4(__riscv_vcreate_v_u16, LMUL_128, _u16, LMUL_512)(
-            u16_val, u16_val, u16_val, u16_val);
-    reg = RVVI4(__riscv_vreinterpret_v_u16, LMUL_512, _bf16,
-                LMUL_512)(u16_combined);
+    // Widen LMUL_128 → LMUL_256 so vslideup operands share a type.
+    // At VLEN=256 this is mf2→m1 (both integer); at VLEN=128 it is m1→m2.
+    fixed_u16x16_t ext =
+        RVVI4(__riscv_vlmul_ext_v_u16, LMUL_128, _u16, LMUL_256)(u16_val);
+    // Build 16-element half: place the 8 elements at offsets 0 and 8.
+    fixed_u16x16_t half = RVVI(__riscv_vmv_v_x_u16, LMUL_256)(0, 16);
+    half = RVVI(__riscv_vslideup_vx_u16, LMUL_256)(half, ext, 0, 8);
+    half = RVVI(__riscv_vslideup_vx_u16, LMUL_256)(half, ext, 8, 16);
+    // Double to LMUL_512 (m1→m2 at VLEN=256, m2→m4 at VLEN=128).
+    fixed_u16x32_t dst =
+        RVVI4(__riscv_vcreate_v_u16, LMUL_256, _u16, LMUL_512)(half, half);
+    reg = RVVI4(__riscv_vreinterpret_v_u16, LMUL_512, _bf16, LMUL_512)(dst);
   };
 
   void save(void* ptr) const {
@@ -623,17 +630,29 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
             data.reg, data.reg)) {};
   explicit FP32Vec16(const FP32Vec16& data) : reg(data.reg) {};
   explicit FP32Vec16(int64_t value, const FP32Vec16& lut) {
-    const uint64_t q_values = static_cast<uint64_t>(value);
-    auto packed = RVVI(__riscv_vmv_v_x_u64, LMUL_1024)(q_values, VEC_ELEM_NUM);
-    auto lane_ids = RVVI(__riscv_vid_v_u64, LMUL_1024)(VEC_ELEM_NUM);
-    auto shifts =
-        RVVI(__riscv_vsll_vx_u64, LMUL_1024)(lane_ids, 2, VEC_ELEM_NUM);
-    auto shifted =
-        RVVI(__riscv_vsrl_vv_u64, LMUL_1024)(packed, shifts, VEC_ELEM_NUM);
-    auto idx64 =
-        RVVI(__riscv_vand_vx_u64, LMUL_1024)(shifted, 0xF, VEC_ELEM_NUM);
-    auto idx32 = RVVI(__riscv_vnsrl_wx_u32, LMUL_512)(idx64, 0, VEC_ELEM_NUM);
-    reg = RVVI(__riscv_vrgather_vv_f32, LMUL_512)(lut.reg, idx32, VEC_ELEM_NUM);
+    // Split into two 32-bit halves to avoid u64 @ LMUL_1024 (m8 on
+    // VLEN=128 / m4 on VLEN=256), which causes heavy register spilling.
+    constexpr int HALF = VEC_ELEM_NUM / 2;
+    const auto q = static_cast<uint64_t>(value);
+    const uint32_t lo = static_cast<uint32_t>(q);
+    const uint32_t hi = static_cast<uint32_t>(q >> 32);
+
+    auto lane_ids = RVVI(__riscv_vid_v_u32, LMUL_256)(HALF);
+    auto shifts = RVVI(__riscv_vsll_vx_u32, LMUL_256)(lane_ids, 2, HALF);
+
+    auto packed_lo = RVVI(__riscv_vmv_v_x_u32, LMUL_256)(lo, HALF);
+    auto idx_lo = RVVI(__riscv_vand_vx_u32, LMUL_256)(
+        RVVI(__riscv_vsrl_vv_u32, LMUL_256)(packed_lo, shifts, HALF), 0xF,
+        HALF);
+
+    auto packed_hi = RVVI(__riscv_vmv_v_x_u32, LMUL_256)(hi, HALF);
+    auto idx_hi = RVVI(__riscv_vand_vx_u32, LMUL_256)(
+        RVVI(__riscv_vsrl_vv_u32, LMUL_256)(packed_hi, shifts, HALF), 0xF,
+        HALF);
+
+    auto idx =
+        RVVI4(__riscv_vcreate_v_u32, LMUL_256, _u32, LMUL_512)(idx_lo, idx_hi);
+    reg = RVVI(__riscv_vrgather_vv_f32, LMUL_512)(lut.reg, idx, VEC_ELEM_NUM);
   }
   explicit FP32Vec16(const FP16Vec16& v);
 

csrc/cpu/cpu_types_vsx.hpp

@@ -344,53 +344,133 @@ struct FP32Vec8 : public Vec<FP32Vec8> {
 
     return result;
   }
-
   FP32Vec8 exp() const {
-    // TODO: Vectorize this
-    AliasReg ar;
-    ar.reg = reg;
-    f32x4x4_t ret;
-    ret.val[0][0] = std::exp(ar.values[0]);
-    ret.val[0][1] = std::exp(ar.values[1]);
-    ret.val[0][2] = std::exp(ar.values[2]);
-    ret.val[0][3] = std::exp(ar.values[3]);
-    ret.val[1][0] = std::exp(ar.values[4]);
-    ret.val[1][1] = std::exp(ar.values[5]);
-    ret.val[1][2] = std::exp(ar.values[6]);
-    ret.val[1][3] = std::exp(ar.values[7]);
-    return FP32Vec8(f32x4x2_t({ret.val[0], ret.val[1]}));
+    f32x4x2_t out;
+    const __vector float log2e = vec_splats(1.44269504088896341f);
+    const __vector float one = vec_splats(1.0f);
+    const __vector float min_x = vec_splats(-87.3f);
+    const __vector float max_x = vec_splats(88.7f);
+
+    // 5th-degree minimax polynomial for 2^r (r in [0,1))
+    const __vector float c1 = vec_splats(0.6931471805599453f);
+    const __vector float c2 = vec_splats(0.240226506959101f);
+    const __vector float c3 = vec_splats(0.05550410866482158f);
+    const __vector float c4 = vec_splats(0.009618129107628477f);
+    const __vector float c5 = vec_splats(0.0013333558146428443f);
+
+    for (int i = 0; i < 2; i++) {
+      __vector float x = reg.val[i];
+      x = vec_max(x, min_x);
+      x = vec_min(x, max_x);
+
+      __vector float y = vec_mul(x, log2e);
+
+      __vector float kf = vec_floor(y);
+      __vector float r = vec_sub(y, kf);
+
+      // Convert float to signed integer. Use vec_cts for PowerPC AltiVec
+      // compatibility.
+      __vector signed int k = vec_cts(kf, 0);
+      const __vector signed int min_k = vec_splats((signed int)-126);
+      const __vector signed int max_k = vec_splats((signed int)127);
+      k = vec_min(vec_max(k, min_k), max_k);
+
+      // Build 2^k from exponent bits
+      __vector signed int exp_int = vec_add(k, vec_splats((signed int)127));
+      __vector unsigned int bits = (__vector unsigned int)exp_int;
+      bits = vec_sl(bits, vec_splats((unsigned int)23));
+      __vector float pow2k = (__vector float)bits;
+
+      // Improved minimax polynomial
+      __vector float poly = vec_madd(c5, r, c4);
+      poly = vec_madd(poly, r, c3);
+      poly = vec_madd(poly, r, c2);
+      poly = vec_madd(poly, r, c1);
+      poly = vec_madd(poly, r, one);
+
+      out.val[i] = vec_mul(pow2k, poly);
+    }
+    return FP32Vec8(out);
   }
 
   FP32Vec8 tanh() const {
-    // TODO: Vectorize this
-    AliasReg ar;
-    ar.reg = reg;
-    f32x4x4_t ret;
-    ret.val[0][0] = std::tanh(ar.values[0]);
-    ret.val[0][1] = std::tanh(ar.values[1]);
-    ret.val[0][2] = std::tanh(ar.values[2]);
-    ret.val[0][3] = std::tanh(ar.values[3]);
-    ret.val[1][0] = std::tanh(ar.values[4]);
-    ret.val[1][1] = std::tanh(ar.values[5]);
-    ret.val[1][2] = std::tanh(ar.values[6]);
-    ret.val[1][3] = std::tanh(ar.values[7]);
-    return FP32Vec8(f32x4x2_t({ret.val[0], ret.val[1]}));
+    const __vector float one = vec_splats(1.0f);
+    const __vector float two = vec_splats(2.0f);
+    const __vector float zero = vec_splats(0.0f);
+    const __vector float sat = vec_splats(9.0f);
+
+    f32x4x2_t out;
+
+    for (int i = 0; i < 2; i++) {
+      __vector float x = reg.val[i];
+      __vector float ax = vec_abs(x);
+
+      __vector bool int mask = vec_cmpge(x, zero);
+      __vector float sign = vec_sel(vec_splats(-1.0f), one, mask);
+
+      __vector bool int saturated = vec_cmpge(ax, sat);
+
+      __vector float two_x = vec_mul(x, two);
+      f32x4x2_t tmp;
+      tmp.val[0] = two_x;
+      tmp.val[1] = two_x;
+      FP32Vec8 temp_vec(tmp);
+      vector float e = temp_vec.exp().reg.val[0];
+
+      vector float num = vec_sub(e, one);
+      vector float den = vec_add(e, one);
+      vector float t = vec_div(num, den);
+
+      out.val[i] = vec_sel(t, sign, saturated);
+    }
+    return FP32Vec8(out);
   }
 
   FP32Vec8 er() const {
-    // TODO: Vectorize this
-    AliasReg ar;
-    ar.reg = reg;
-    f32x4x4_t ret;
-    ret.val[0][0] = std::erf(ar.values[0]);
-    ret.val[0][1] = std::erf(ar.values[1]);
-    ret.val[0][2] = std::erf(ar.values[2]);
-    ret.val[0][3] = std::erf(ar.values[3]);
-    ret.val[1][0] = std::erf(ar.values[4]);
-    ret.val[1][1] = std::erf(ar.values[5]);
-    ret.val[1][2] = std::erf(ar.values[6]);
-    ret.val[1][3] = std::erf(ar.values[7]);
-    return FP32Vec8(f32x4x2_t({ret.val[0], ret.val[1]}));
+    const vector float a1 = vec_splats(0.254829592f);
+    const vector float a2 = vec_splats(-0.284496736f);
+    const vector float a3 = vec_splats(1.421413741f);
+    const vector float a4 = vec_splats(-1.453152027f);
+    const vector float a5 = vec_splats(1.061405429f);
+    const vector float p = vec_splats(0.3275911f);
+    const vector float one = vec_splats(1.0f);
+    const vector float zero = vec_splats(0.0f);
+    const vector float sat = vec_splats(6.0f);
+
+    f32x4x2_t ret;
+
+    for (int i = 0; i < 2; i++) {
+      vector float x = reg.val[i];
+      vector float ax = vec_abs(x);
+
+      vector bool int mask = vec_cmpge(x, zero);
+      vector float sign = vec_sel(vec_splats(-1.0f), one, mask);
+
+      vector bool int saturated = vec_cmpge(ax, sat);
+
+      vector float t = vec_div(one, vec_madd(p, ax, one));
+
+      vector float poly = a5;
+      poly = vec_madd(poly, t, a4);
+      poly = vec_madd(poly, t, a3);
+      poly = vec_madd(poly, t, a2);
+      poly = vec_madd(poly, t, a1);
+      poly = vec_mul(poly, t);
+
+      vector float x_squared = vec_mul(x, x);
+      vector float neg_x_squared = vec_mul(vec_splats(-1.0f), x_squared);
+      f32x4x2_t tmp;
+      tmp.val[0] = neg_x_squared;
+      tmp.val[1] = neg_x_squared;
+      FP32Vec8 exp_input(tmp);
+      vector float exp_term = exp_input.exp().reg.val[0];
+
+      vector float y = vec_nmsub(poly, exp_term, one);
+      vector float erf_val = vec_mul(sign, y);
+
+      ret.val[i] = vec_sel(erf_val, sign, saturated);
+    }
+    return FP32Vec8(ret);
   }
 
   FP32Vec8 operator*(const FP32Vec8& b) const {

csrc/cpu/torch_bindings.cpp

@@ -278,7 +278,8 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
   ops.def(
       "dynamic_4bit_int_moe("
       "Tensor x, Tensor topk_ids, Tensor topk_weights,"
-      "Tensor w13_packed, Tensor w2_packed, int H, int I, int I2,"
+      "Tensor w13_packed, Tensor w2_packed,"
+      "int hidden_size, int intermediate_size,"
       "int group_size, bool apply_router_weight_on_input, int activation_kind"
       ") -> Tensor");
 

csrc/cpu/utils.hpp

@@ -76,14 +76,14 @@ inline int64_t get_available_l2_size() {
     if (l2_cache_size == 0) {
       l2_cache_size = 256 * 1024;
     }
-    return static_cast<int64_t>(l2_cache_size) >> 1;  // use 50% of L2 cache
+    return static_cast<int64_t>(l2_cache_size) >> 1;
   }();
   return size;
 #else
   static int64_t size = []() {
     auto caps = at::cpu::get_cpu_capabilities();
     const uint32_t l2_cache_size = caps.at("l2_cache_size").toInt();
-    return l2_cache_size >> 1;  // use 50% of L2 cache
+    return l2_cache_size >> 1;
   }();
   return size;
 #endif

csrc/libtorch_stable/moe/moe_ops.h

@@ -15,7 +15,8 @@ void topk_sigmoid(torch::stable::Tensor& topk_weights,
                   torch::stable::Tensor& topk_indices,
                   torch::stable::Tensor& token_expert_indices,
                   torch::stable::Tensor& gating_output, bool renormalize,
-                  std::optional<torch::stable::Tensor> bias);
+                  std::optional<torch::stable::Tensor> bias,
+                  double routed_scaling_factor);
 
 void topk_softplus_sqrt(
     torch::stable::Tensor& topk_weights, torch::stable::Tensor& topk_indices,

csrc/libtorch_stable/moe/topk_softmax_kernels.cu

@@ -173,7 +173,8 @@ __launch_bounds__(TPB) __global__ void moeTopK(
     const int start_expert,
     const int end_expert,
     const bool renormalize,
-    const float* bias)
+    const float* bias,
+    const double routed_scaling_factor)
 {
 
     using cub_kvp = cub::KeyValuePair<int, float>;
@@ -241,14 +242,16 @@ __launch_bounds__(TPB) __global__ void moeTopK(
         __syncthreads();
     }
 
-    // Renormalize the k weights for this row to sum to 1, if requested.
-    if (renormalize) {
-        if (threadIdx.x == 0) {
+    // Apply renormalization and routed scaling factor to final weights.
+    if (threadIdx.x == 0) {
+        float scale = static_cast<float>(routed_scaling_factor);
+        if (renormalize) {
             const float denom = selected_sum > 0.f ? selected_sum : 1.f;
-            for (int k_idx = 0; k_idx < k; ++k_idx) {
-                const int idx = k * block_row + k_idx;
-                output[idx] = output[idx] / denom;
-            }
+            scale /= denom;
+        }
+        for (int k_idx = 0; k_idx < k; ++k_idx) {
+            const int idx = k * block_row + k_idx;
+            output[idx] = output[idx] * scale;
         }
     }
 }
@@ -274,7 +277,7 @@ template <int VPT, int NUM_EXPERTS, int WARPS_PER_CTA, int BYTES_PER_LDG, int WA
 __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
     void topkGating(const InputType* input, const bool* finished, float* output, const int num_rows, IndType* indices,
         int* source_rows, const int k, const int start_expert, const int end_expert, const bool renormalize,
-        const float* bias)
+        const float* bias, const double routed_scaling_factor)
 {
     static_assert(std::is_same_v<InputType, float> || std::is_same_v<InputType, __nv_bfloat16> ||
                       std::is_same_v<InputType, __half>,
@@ -570,17 +573,17 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
         }
     }
 
-    // Renormalize the k weights for this row to sum to 1, if requested.
-    if (renormalize) {
-        if (thread_group_idx == 0)
-        {
-            const float denom = selected_sum > 0.f ? selected_sum : 1.f;
-            for (int k_idx = 0; k_idx < k; ++k_idx)
-            {
-                const int idx = k * thread_row + k_idx;
-                output[idx] = output[idx] / denom;
-            }
-        }
+    // Apply renormalization and routed scaling factor to final weights.
+    if (thread_group_idx == 0) {
+      float scale = static_cast<float>(routed_scaling_factor);
+      if (renormalize) {
+        const float denom = selected_sum > 0.f ? selected_sum : 1.f;
+        scale /= denom;
+      }
+      for (int k_idx = 0; k_idx < k; ++k_idx) {
+        const int idx = k * thread_row + k_idx;
+        output[idx] = output[idx] * scale;
+      }
     }
 }
 
@@ -602,7 +605,7 @@ struct TopkConstants
 template <int EXPERTS, int WARPS_PER_TB, int WARP_SIZE_PARAM, int MAX_BYTES_PER_LDG, typename IndType, typename InputType, ScoringFunc SF>
 void topkGatingLauncherHelper(const InputType* input, const bool* finished, float* output, IndType* indices,
     int* source_row, const int num_rows, const int k, const int start_expert, const int end_expert, const bool renormalize,
-    const float* bias, cudaStream_t stream)
+    const float* bias, const double routed_scaling_factor, cudaStream_t stream)
 {
     static constexpr int BYTES_PER_LDG = MIN(MAX_BYTES_PER_LDG, sizeof(InputType) * EXPERTS);
     using Constants = detail::TopkConstants<EXPERTS, BYTES_PER_LDG, WARP_SIZE_PARAM, InputType>;
@@ -613,7 +616,7 @@ void topkGatingLauncherHelper(const InputType* input, const bool* finished, floa
 
     dim3 block_dim(WARP_SIZE_PARAM, WARPS_PER_TB);
     topkGating<VPT, EXPERTS, WARPS_PER_TB, BYTES_PER_LDG, WARP_SIZE_PARAM, IndType, InputType, SF><<<num_blocks, block_dim, 0, stream>>>(
-        input, finished, output, num_rows, indices, source_row, k, start_expert, end_expert, renormalize, bias);
+        input, finished, output, num_rows, indices, source_row, k, start_expert, end_expert, renormalize, bias, routed_scaling_factor);
 }
 
 #ifndef USE_ROCM
@@ -624,7 +627,7 @@ void topkGatingLauncherHelper(const InputType* input, const bool* finished, floa
                              IndType, InputType, SF>(                         \
         gating_output, nullptr, topk_weights, topk_indices,                   \
         token_expert_indices, num_tokens, topk, 0, num_experts, renormalize,  \
-        bias, stream);
+        bias, routed_scaling_factor, stream);
 #else
   #define LAUNCH_TOPK(NUM_EXPERTS, WARPS_PER_TB, MAX_BYTES)                    \
     if (WARP_SIZE == 64) {                                                     \
@@ -632,13 +635,13 @@ void topkGatingLauncherHelper(const InputType* input, const bool* finished, floa
                                IndType, InputType, SF>(                        \
           gating_output, nullptr, topk_weights, topk_indices,                  \
           token_expert_indices, num_tokens, topk, 0, num_experts, renormalize, \
-          bias, stream);                                                       \
+          bias, routed_scaling_factor, stream);                                                       \
     } else if (WARP_SIZE == 32) {                                              \
       topkGatingLauncherHelper<NUM_EXPERTS, WARPS_PER_TB, 32, MAX_BYTES,       \
                                IndType, InputType, SF>(                        \
           gating_output, nullptr, topk_weights, topk_indices,                  \
           token_expert_indices, num_tokens, topk, 0, num_experts, renormalize, \
-          bias, stream);                                                       \
+          bias, routed_scaling_factor, stream);                                                       \
     } else {                                                                   \
       assert(false &&                                                          \
              "Unsupported warp size. Only 32 and 64 are supported for ROCm");  \
@@ -657,6 +660,7 @@ void topkGatingKernelLauncher(
     const int topk,
     const bool renormalize,
     const float* bias,
+    const double routed_scaling_factor,
     cudaStream_t stream) {
     static constexpr int WARPS_PER_TB = 4;
     static constexpr int BYTES_PER_LDG_POWER_OF_2 = 16;
@@ -732,7 +736,7 @@ void topkGatingKernelLauncher(
             }
             moeTopK<TPB><<<num_tokens, TPB, 0, stream>>>(
                 workspace, nullptr, topk_weights, topk_indices, token_expert_indices,
-                num_experts, topk, 0, num_experts, renormalize, bias);
+                num_experts, topk, 0, num_experts, renormalize, bias, routed_scaling_factor);
         }
     }
 }
@@ -750,6 +754,7 @@ void dispatch_topk_launch(
     torch::stable::Tensor& softmax_workspace,
     int num_tokens, int num_experts, int topk, bool renormalize,
     std::optional<torch::stable::Tensor> bias,
+    double routed_scaling_factor,
     cudaStream_t stream)
  {
     const float* bias_ptr = nullptr;
@@ -772,7 +777,7 @@ void dispatch_topk_launch(
             token_expert_indices.mutable_data_ptr<int>(),
             softmax_workspace.mutable_data_ptr<float>(),
             num_tokens, num_experts, topk, renormalize,
-            bias_ptr, stream);
+            bias_ptr, routed_scaling_factor, stream);
     } else if (topk_indices.scalar_type() == torch::headeronly::ScalarType::UInt32) {
         vllm::moe::topkGatingKernelLauncher<uint32_t, ComputeType, SF>(
             reinterpret_cast<const ComputeType*>(gating_output.const_data_ptr()),
@@ -781,7 +786,7 @@ void dispatch_topk_launch(
             token_expert_indices.mutable_data_ptr<int>(),
             softmax_workspace.mutable_data_ptr<float>(),
             num_tokens, num_experts, topk, renormalize,
-            bias_ptr, stream);
+            bias_ptr, routed_scaling_factor, stream);
     } else {
         STD_TORCH_CHECK(topk_indices.scalar_type() == torch::headeronly::ScalarType::Long);
         vllm::moe::topkGatingKernelLauncher<int64_t, ComputeType, SF>(
@@ -791,7 +796,7 @@ void dispatch_topk_launch(
             token_expert_indices.mutable_data_ptr<int>(),
             softmax_workspace.mutable_data_ptr<float>(),
             num_tokens, num_experts, topk, renormalize,
-            bias_ptr, stream);
+            bias_ptr, routed_scaling_factor, stream);
     }
 }
 
@@ -820,15 +825,15 @@ void topk_softmax(
     if (gating_output.scalar_type() == torch::headeronly::ScalarType::Float) {
         dispatch_topk_launch<float, vllm::moe::SCORING_SOFTMAX>(gating_output, topk_weights, topk_indices,
             token_expert_indices, softmax_workspace, num_tokens, num_experts, topk, renormalize,
-            bias, stream);
+            bias, 1.0, stream);
     } else if (gating_output.scalar_type() == torch::headeronly::ScalarType::Half) {
         dispatch_topk_launch<__half, vllm::moe::SCORING_SOFTMAX>(gating_output, topk_weights, topk_indices,
             token_expert_indices, softmax_workspace, num_tokens, num_experts, topk, renormalize,
-            bias, stream);
+            bias, 1.0, stream);
     } else if (gating_output.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
         dispatch_topk_launch<__nv_bfloat16, vllm::moe::SCORING_SOFTMAX>(gating_output, topk_weights, topk_indices,
             token_expert_indices, softmax_workspace, num_tokens, num_experts, topk, renormalize,
-            bias, stream);
+            bias, 1.0, stream);
     } else {
         STD_TORCH_CHECK(false, "Unsupported gating_output data type: ", gating_output.scalar_type());
     }
@@ -840,7 +845,8 @@ void topk_sigmoid(
     torch::stable::Tensor& token_expert_indices,        // [num_tokens, topk]
     torch::stable::Tensor& gating_output,               // [num_tokens, num_experts]
     bool renormalize,
-    std::optional<torch::stable::Tensor> bias)
+    std::optional<torch::stable::Tensor> bias,
+    double routed_scaling_factor)
 {
     const int num_experts = gating_output.size(-1);
     const auto num_tokens = gating_output.numel() / num_experts;
@@ -859,15 +865,15 @@ void topk_sigmoid(
     if (gating_output.scalar_type() == torch::headeronly::ScalarType::Float) {
         dispatch_topk_launch<float, vllm::moe::SCORING_SIGMOID>(gating_output, topk_weights, topk_indices,
             token_expert_indices, workspace, num_tokens, num_experts, topk, renormalize,
-            bias, stream);
+            bias, routed_scaling_factor, stream);
     } else if (gating_output.scalar_type() == torch::headeronly::ScalarType::Half) {
         dispatch_topk_launch<__half, vllm::moe::SCORING_SIGMOID>(gating_output, topk_weights, topk_indices,
             token_expert_indices, workspace, num_tokens, num_experts, topk, renormalize,
-            bias, stream);
+            bias, routed_scaling_factor, stream);
     } else if (gating_output.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
         dispatch_topk_launch<__nv_bfloat16, vllm::moe::SCORING_SIGMOID>(gating_output, topk_weights, topk_indices,
             token_expert_indices, workspace, num_tokens, num_experts, topk, renormalize,
-            bias, stream);
+            bias, routed_scaling_factor, stream);
     } else {
         STD_TORCH_CHECK(false, "Unsupported gating_output data type: ", gating_output.scalar_type());
     }

csrc/libtorch_stable/moe/torch_bindings.cpp

@@ -13,8 +13,8 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_moe_C, m) {
   // Apply topk sigmoid to the gating outputs.
   m.def(
       "topk_sigmoid(Tensor! topk_weights, Tensor! topk_indices, Tensor! "
-      "token_expert_indices, Tensor gating_output, bool renormalize, Tensor? "
-      "bias) -> ()");
+      "token_expert_indices, Tensor gating_output, bool renormalize, "
+      "Tensor? bias, float routed_scaling_factor) -> ()");
 
   m.def(
       "topk_softplus_sqrt(Tensor! topk_weights, Tensor! topk_indices, Tensor! "

csrc/libtorch_stable/ngram_embedding_kernels.cu

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

csrc/libtorch_stable/ops.h

@@ -554,3 +554,12 @@ void cp_gather_indexer_k_quant_cache(
                                                 // quant_block_size * 4]
     const torch::stable::Tensor& block_table,   // [batch_size, num_blocks]
     const torch::stable::Tensor& cu_seq_lens);  // [batch_size + 1]
+
+// LongCat n-gram embedding index kernel (see ngram_embedding_kernels.cu).
+void ngram_compute_n_gram_ids(
+    int64_t ne_n, int64_t ne_k, torch::stable::Tensor& ne_weights,
+    torch::stable::Tensor& ne_mods,
+    torch::stable::Tensor& exclusive_ne_embedder_size_sums,
+    torch::stable::Tensor& exclusive_req_len_sums,
+    torch::stable::Tensor& ne_token_table, torch::stable::Tensor& row_indices,
+    torch::stable::Tensor& column_starts, torch::stable::Tensor& n_gram_ids);

csrc/libtorch_stable/torch_bindings.cpp

@@ -598,9 +598,22 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
       "Tensor? initial_state_idx,"
       "Tensor? cu_chunk_seqlen,"
       "Tensor? last_chunk_indices) -> ()");
+
+  // LongCat n-gram embedding index kernel. All tensor args are marked mutable
+  // to match the (non-const) stable-Tensor& C++ signature; only ne_token_table
+  // and n_gram_ids are actually written in place.
+  ops.def(
+      "ngram_compute_n_gram_ids(int ne_n, int ne_k, Tensor(a!) ne_weights, "
+      "Tensor(b!) ne_mods, Tensor(c!) exclusive_ne_embedder_size_sums, "
+      "Tensor(d!) exclusive_req_len_sums, Tensor(e!) ne_token_table, "
+      "Tensor(f!) row_indices, Tensor(g!) column_starts, "
+      "Tensor(h!) n_gram_ids) -> ()");
 }
 
 STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
+  // LongCat n-gram embedding index kernel.
+  ops.impl("ngram_compute_n_gram_ids", TORCH_BOX(&ngram_compute_n_gram_ids));
+
   // Per-token group quantization
   ops.impl("per_token_group_fp8_quant", TORCH_BOX(&per_token_group_quant_fp8));
   ops.impl("per_token_group_fp8_quant_packed",

csrc/moe/dynamic_4bit_int_moe_cpu.cpp

@@ -29,25 +29,37 @@ enum ActivationKind : int64_t {
 
 torch::Tensor dynamic_4bit_int_moe_cpu(
     torch::Tensor x, torch::Tensor topk_ids, torch::Tensor topk_weights,
-    torch::Tensor w13_packed, torch::Tensor w2_packed, int64_t H, int64_t I,
-    int64_t I2, int64_t group_size, bool apply_router_weight_on_input,
-    int64_t activation_kind) {
+    torch::Tensor w13_packed, torch::Tensor w2_packed, int64_t hidden_size,
+    int64_t intermediate_size, int64_t group_size,
+    bool apply_router_weight_on_input, int64_t activation_kind) {
   TORCH_CHECK(x.dim() == 2, "x must be 2D");
   TORCH_CHECK(topk_ids.dim() == 2 && topk_weights.dim() == 2,
               "topk tensors must be [T, K]");
   TORCH_CHECK(
       w13_packed.size(0) == w2_packed.size(0),
       "w13_packed and w2_packed must have same number of experts in dim 0");
-  TORCH_CHECK(I2 == 2 * I, "I2 must equal 2*I");
 
   const int64_t T = x.size(0);
   const int64_t K = topk_ids.size(1);
   const int64_t E = w13_packed.size(0);
   const int64_t N = T * K;
+  const int64_t w13_out_features = 2 * intermediate_size;
 
   auto x_c = x.contiguous();
+  // _dyn_quant_matmul_4bit kernel natively supports these pre-quant activation
+  // dtypes:
+  // - fp32: with channelwise and groupwise
+  // - bf16: with channelwise -> upcast to fp32 for groupwise
+  // - fp16: not supported -> upcast to fp32 for groupwise & channelwise
+  const auto output_dtype = x_c.scalar_type();
+  const bool should_cast_input =
+      ((group_size != -1) && output_dtype == at::kBFloat16) ||
+      output_dtype == at::kHalf;
+  if (should_cast_input) {
+    x_c = x_c.to(at::kFloat);
+  }
   auto ids_c = topk_ids.contiguous();
-  auto gates_c = topk_weights.to(at::kFloat).contiguous();
+  auto gates_c = topk_weights.to(x_c.scalar_type()).contiguous();
 
   // bucketing tokens -> experts
   c10::SmallVector<int64_t, 64> counts(
@@ -63,35 +75,42 @@ torch::Tensor dynamic_4bit_int_moe_cpu(
   c10::SmallVector<int64_t, 65> offsets(E + 1, 0);  // ( E +1 )
   for (int64_t e = 0; e < E; ++e) offsets[e + 1] = offsets[e] + counts[e];
 
+  // expert_tokens = [tokens indices for expert 0, ...]
+  // expert_gates = [router weights for tokens assigned to expert 0, ...]
   auto expert_tokens = at::empty({offsets[E]}, ids_c.options());
   auto expert_gates = at::empty({offsets[E]}, gates_c.options());
   {
     c10::SmallVector<int64_t, 64> cursor(E, 0);
-    const auto* ids_ptr = ids_c.data_ptr<int64_t>();
-    const auto* gts_ptr = gates_c.data_ptr<float>();
-    auto* tok_ptr = expert_tokens.data_ptr<int64_t>();
-    auto* gate_ptr = expert_gates.data_ptr<float>();
+    AT_DISPATCH_FLOATING_TYPES_AND2(
+        at::ScalarType::BFloat16, at::ScalarType::Half, gates_c.scalar_type(),
+        "bucket_expert_tokens_and_gates", [&] {
+          const auto* ids_ptr = ids_c.data_ptr<int64_t>();
+          const auto* gts_ptr = gates_c.data_ptr<scalar_t>();
+          auto* tok_ptr = expert_tokens.data_ptr<int64_t>();
+          auto* gate_ptr = expert_gates.data_ptr<scalar_t>();
 
-    for (int64_t t = 0; t < T; ++t) {
-      const int64_t base = t * K;
-      for (int64_t k = 0; k < K; ++k) {
-        const int64_t idx = base + k;
-        const int64_t e = ids_ptr[idx];
-        const int64_t p = offsets[e] + (cursor[e]++);
-        tok_ptr[p] = t;
-        gate_ptr[p] = gts_ptr[idx];
-      }
-    }
+          for (int64_t t = 0; t < T; ++t) {
+            const int64_t base = t * K;
+            for (int64_t k = 0; k < K; ++k) {
+              const int64_t idx = base + k;
+              const int64_t e = ids_ptr[idx];
+              const int64_t p = offsets[e] + (cursor[e]++);
+              tok_ptr[p] = t;
+              gate_ptr[p] = gts_ptr[idx];
+            }
+          }
+        });
   }
 
-  const int64_t g_eff_13 = (group_size != -1) ? group_size : H;
-  const int64_t g_eff_2 = (group_size != -1) ? group_size : I;
+  const int64_t g_eff_13 = (group_size != -1) ? group_size : hidden_size;
+  const int64_t g_eff_2 = (group_size != -1) ? group_size : intermediate_size;
 
+  // X_all [num_tokens * K, hidden_size]
   auto X_all = x_c.index_select(/*dim=*/0, expert_tokens);
   if (apply_router_weight_on_input) {
     X_all = X_all.mul(expert_gates.unsqueeze(1));
   }
-  auto Y_all = at::empty({offsets[E], H}, x_c.options());
+  auto Y_all = at::empty({offsets[E], hidden_size}, x_c.options());
 
   at::parallel_for(0, offsets[E], 0, [&](int64_t idx_begin, int64_t idx_end) {
     c10::InferenceMode guard;
@@ -109,11 +128,13 @@ torch::Tensor dynamic_4bit_int_moe_cpu(
       auto w2_e = w2_packed.select(/*dim=*/0, e);
 
       // W13
-      auto y13 =
-          mm(x_e, w13_e, g_eff_13, /*in_features=*/H, /*out_features=*/I2);
+      auto y13 = mm(x_e, w13_e, g_eff_13, /*in_features=*/hidden_size,
+                    /*out_features=*/w13_out_features);
 
-      auto g_part = y13.narrow(/*dim=*/1, /*start=*/0, /*length=*/I);
-      auto u_part = y13.narrow(/*dim=*/1, /*start=*/I, /*length=*/I);
+      auto g_part =
+          y13.narrow(/*dim=*/1, /*start=*/0, /*length=*/intermediate_size);
+      auto u_part = y13.narrow(/*dim=*/1, /*start=*/intermediate_size,
+                               /*length=*/intermediate_size);
 
       torch::Tensor act;
       if (activation_kind == ActivationKind::SwiGLUOAI) {  // SwiGLUOAI
@@ -128,7 +149,8 @@ torch::Tensor dynamic_4bit_int_moe_cpu(
       }
 
       // W2
-      auto y = mm(act, w2_e, g_eff_2, /*in_features=*/I, /*out_features=*/H);
+      auto y = mm(act, w2_e, g_eff_2, /*in_features=*/intermediate_size,
+                  /*out_features=*/hidden_size);
 
       // Store per-expert result
       Y_all.narrow(/*dim=*/0, /*start=*/start, /*length=*/te).copy_(y);
@@ -138,8 +160,11 @@ torch::Tensor dynamic_4bit_int_moe_cpu(
   if (!apply_router_weight_on_input) {
     Y_all = Y_all.mul(expert_gates.unsqueeze(1));
   }
+  if (Y_all.scalar_type() != output_dtype) {
+    Y_all = Y_all.to(output_dtype);
+  }
 
-  auto out = at::zeros({T, H}, x.options());
+  auto out = at::zeros({T, hidden_size}, x.options());
   out =
       at::index_add(out, /*dim=*/0, /*index=*/expert_tokens, /*source=*/Y_all);
 

csrc/ops.h

@@ -53,9 +53,9 @@ void dynamic_scaled_int8_quant(torch::Tensor& out, torch::Tensor const& input,
 
 torch::Tensor dynamic_4bit_int_moe_cpu(
     torch::Tensor x, torch::Tensor topk_ids, torch::Tensor topk_weights,
-    torch::Tensor w13_packed, torch::Tensor w2_packed, int64_t H, int64_t I,
-    int64_t I2, int64_t group_size, bool apply_router_weight_on_input,
-    int64_t activation_kind);
+    torch::Tensor w13_packed, torch::Tensor w2_packed, int64_t hidden_size,
+    int64_t intermediate_size, int64_t group_size,
+    bool apply_router_weight_on_input, int64_t activation_kind);
 
 using fptr_t = int64_t;
 #ifdef USE_ROCM

docker/Dockerfile.cpu

@@ -25,7 +25,6 @@ FROM ubuntu:22.04 AS base-common
 WORKDIR /workspace
 
 ARG PYTHON_VERSION=3.12
-ARG PIP_EXTRA_INDEX_URL="https://download.pytorch.org/whl/cpu"
 
 ARG max_jobs=32
 ENV MAX_JOBS=${max_jobs}
@@ -53,8 +52,6 @@ ENV PATH="$VIRTUAL_ENV/bin:$PATH"
 ENV UV_HTTP_TIMEOUT=500
 
 # Install Python dependencies
-ENV PIP_EXTRA_INDEX_URL=${PIP_EXTRA_INDEX_URL}
-ENV UV_EXTRA_INDEX_URL=${PIP_EXTRA_INDEX_URL}
 ENV UV_INDEX_STRATEGY="unsafe-best-match"
 ENV UV_LINK_MODE="copy"
 
@@ -64,7 +61,7 @@ COPY requirements/cpu.txt requirements/cpu.txt
 
 RUN --mount=type=cache,target=/root/.cache/uv \
     uv pip install --upgrade pip && \
-    uv pip install -r requirements/cpu.txt
+    uv pip install -r requirements/cpu.txt --torch-backend cpu
 
 ARG TARGETARCH
 ENV TARGETARCH=${TARGETARCH}
@@ -149,7 +146,7 @@ RUN if [ "$TARGETARCH" = "arm64" ] && [ "$VLLM_CPU_X86" != "0" ]; then \
 COPY requirements/build/cpu.txt requirements/build/cpu.txt
 
 RUN --mount=type=cache,target=/root/.cache/uv \
-    uv pip install -r requirements/build/cpu.txt
+    uv pip install -r requirements/build/cpu.txt --torch-backend cpu
 
 COPY . .
 
@@ -205,7 +202,7 @@ RUN case "$(uname -m)" in \
     esac
 
 RUN --mount=type=cache,target=/root/.cache/uv \
-    uv pip install -r requirements/test/cpu.txt
+    uv pip install -r requirements/test/cpu.txt --torch-backend cpu
 
 ######################### DEV IMAGE #########################
 FROM vllm-build AS vllm-dev
@@ -231,7 +228,7 @@ COPY --from=vllm-test-deps /vllm-workspace/requirements/test/cpu.txt requirement
 
 RUN --mount=type=cache,target=/root/.cache/uv \
     uv pip install -r requirements/lint.txt && \
-    uv pip install -r requirements/test/cpu.txt && \
+    uv pip install -r requirements/test/cpu.txt --torch-backend cpu && \
     pre-commit install --hook-type pre-commit --hook-type commit-msg
 
 ENTRYPOINT ["bash"]
@@ -301,6 +298,12 @@ LABEL ai.vllm.build.cpu-x86="${VLLM_CPU_X86:-false}"
 LABEL ai.vllm.build.cpu-arm-bf16="${VLLM_CPU_ARM_BF16:-false}"
 LABEL ai.vllm.build.python-version="${PYTHON_VERSION:-3.12}"
 
+# Copy the examples directory (including the chat/tool templates) so it is
+# present in the released image, as the CUDA image ships it too. The vllm-test
+# stage above adds examples/ for testing only, so without this the published
+# vllm-openai-cpu image would not ship examples/*.jinja.
+COPY examples examples
+
 ENTRYPOINT ["vllm", "serve"]
 
 

docker/Dockerfile.ppc64le

@@ -1,275 +1,80 @@
+# Base UBI image
 ARG BASE_UBI_IMAGE_TAG=9.6-1754584681
 
 ###############################################################
-# Stage to build openblas
+#                       BUILDER  STAGE                        #
 ###############################################################
 
-FROM registry.access.redhat.com/ubi9/ubi-minimal:${BASE_UBI_IMAGE_TAG} AS openblas-builder
-
-ARG MAX_JOBS
-ARG OPENBLAS_VERSION=0.3.30
-RUN microdnf install -y dnf && dnf install -y gcc-toolset-14 make wget unzip \
-    && source /opt/rh/gcc-toolset-14/enable \
-    && wget https://github.com/OpenMathLib/OpenBLAS/releases/download/v$OPENBLAS_VERSION/OpenBLAS-$OPENBLAS_VERSION.zip \
-    && unzip OpenBLAS-$OPENBLAS_VERSION.zip \
-    && cd OpenBLAS-$OPENBLAS_VERSION \
-    &&  make -j${MAX_JOBS} TARGET=POWER9 BINARY=64 USE_OPENMP=1 USE_THREAD=1 NUM_THREADS=120 DYNAMIC_ARCH=1 INTERFACE64=0 \
-    && cd /tmp && touch control
-
-
-###############################################################
-# base stage with dependencies coming from centos mirrors
-###############################################################
-FROM registry.access.redhat.com/ubi9/ubi-minimal:${BASE_UBI_IMAGE_TAG} AS centos-deps-builder
-RUN  microdnf install -y dnf && \ 
-     dnf install -y https://mirror.stream.centos.org/9-stream/BaseOS/`arch`/os/Packages/centos-gpg-keys-9.0-26.el9.noarch.rpm \
-        https://mirror.stream.centos.org/9-stream/BaseOS/`arch`/os/Packages/centos-stream-repos-9.0-26.el9.noarch.rpm \
-        https://dl.fedoraproject.org/pub/epel/epel-release-latest-9.noarch.rpm && \
-        dnf config-manager --set-enabled crb
-
-RUN dnf install -y openjpeg2-devel lcms2-devel tcl-devel tk-devel fribidi-devel yajl-devel && \
-    dnf remove -y centos-gpg-keys-9.0-24.el9.noarch centos-stream-repos-9.0-26.el9.noarch 
-
-
-###############################################################
-# base stage with basic dependencies
-###############################################################
-
-FROM centos-deps-builder AS base-builder
+FROM registry.access.redhat.com/ubi9/ubi-minimal:${BASE_UBI_IMAGE_TAG} AS builder-base
 
+ARG VLLM_VERSION="0.22.1"
 ARG PYTHON_VERSION=3.12
-ARG OPENBLAS_VERSION=0.3.30
-
-# Set Environment Variables for venv, cargo & openblas
-ENV VIRTUAL_ENV=/opt/vllm
-ENV PATH=${VIRTUAL_ENV}/bin:/root/.cargo/bin:$PATH
-ENV PKG_CONFIG_PATH=/usr/local/lib/pkgconfig/
-ENV LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/lib64:/usr/local/lib:/usr/lib64:/usr/lib
-ENV UV_LINK_MODE=copy
-
-# install gcc-13, python, rust, openblas
-# Note: A symlink for libatomic.so is created for gcc-13 (linker fails to find libatomic otherwise - reqd. for sentencepiece)
-# Note: A dummy file 'control' is created in /tmp/ to artificially create dependencies between stages when building stages in parallel
-#       when `--jobs=<N>` is passed with podman build command
-
-COPY --from=openblas-builder /tmp/control /dev/null
-
-RUN --mount=type=bind,from=openblas-builder,source=/OpenBLAS-$OPENBLAS_VERSION/,target=/openblas/,rw \
-    dnf install -y openssl-devel \
-    && dnf install -y \
-       git tar gcc-toolset-14 automake libtool \
-       pkgconfig xsimd zeromq-devel kmod findutils protobuf* \
-       libtiff-devel libjpeg-devel zlib-devel freetype-devel libwebp-devel \
-       harfbuzz-devel libraqm-devel libimagequant-devel libxcb-devel \
-       python${PYTHON_VERSION}-devel python${PYTHON_VERSION}-pip clang-devel \
-    && dnf clean all \
-    && PREFIX=/usr/local make -C /openblas install \
-    && ln -sf /usr/lib64/libatomic.so.1 /usr/lib64/libatomic.so \
-    && python${PYTHON_VERSION} -m venv ${VIRTUAL_ENV} \
-    && python -m pip install -U pip uv \
-    && uv pip install wheel build "setuptools<70" setuptools_scm setuptools_rust meson-python 'cmake<4' ninja cython scikit_build_core scikit_build \
-    && curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y \
-    && cd /tmp && touch control
-
-
-###############################################################
-# Stage to build torch family
-###############################################################
-
-FROM base-builder AS torch-builder
-
-ARG MAX_JOBS
-ARG TORCH_VERSION=2.7.0
-ARG _GLIBCXX_USE_CXX11_ABI=1
-ARG OPENBLAS_VERSION=0.3.30
-
-RUN --mount=type=cache,target=/root/.cache/uv \
-    source /opt/rh/gcc-toolset-14/enable &&  \
-    git clone --recursive https://github.com/pytorch/pytorch.git -b v${TORCH_VERSION} && \
-    cd pytorch && \
-    uv pip install -r requirements.txt && \
-    python setup.py develop && \
-    rm -f dist/torch*+git*whl && \
-    MAX_JOBS=${MAX_JOBS:-$(nproc)} \
-    PYTORCH_BUILD_VERSION=${TORCH_VERSION} PYTORCH_BUILD_NUMBER=1 uv build --wheel --out-dir /torchwheels/
-
-ARG TORCHVISION_VERSION=0.22.0
-ARG TORCHVISION_USE_NVJPEG=0
-ARG TORCHVISION_USE_FFMPEG=0
-RUN --mount=type=cache,target=/root/.cache/uv \
-    source /opt/rh/gcc-toolset-14/enable && \
-    git clone --recursive https://github.com/pytorch/vision.git -b v${TORCHVISION_VERSION} && \
-    cd vision && \
-    MAX_JOBS=${MAX_JOBS:-$(nproc)} \
-    BUILD_VERSION=${TORCHVISION_VERSION} \
-    uv build --wheel --out-dir /torchwheels/ --no-build-isolation
-
-ARG TORCHAUDIO_VERSION=2.7.0
-ARG BUILD_SOX=1
-ARG BUILD_KALDI=1
-ARG BUILD_RNNT=1
-ARG USE_FFMPEG=0
-ARG USE_ROCM=0
-ARG USE_CUDA=0
-ARG TORCHAUDIO_TEST_ALLOW_SKIP_IF_NO_FFMPEG=1
-RUN --mount=type=cache,target=/root/.cache/uv \
-    source /opt/rh/gcc-toolset-14/enable && \
-    git clone --recursive https://github.com/pytorch/audio.git -b v${TORCHAUDIO_VERSION} && \
-    cd audio && \
-    MAX_JOBS=${MAX_JOBS:-$(nproc)} \
-    BUILD_VERSION=${TORCHAUDIO_VERSION} \
-    uv build --wheel --out-dir /torchwheels/ --no-build-isolation
-
-###############################################################
-# Stage to build pyarrow
-###############################################################
-
-FROM base-builder AS arrow-builder
-
-ARG MAX_JOBS
-ARG PYARROW_PARALLEL
-ARG PYARROW_VERSION=21.0.0
-RUN --mount=type=cache,target=/root/.cache/uv \
-    source /opt/rh/gcc-toolset-14/enable && \
-    git clone --recursive https://github.com/apache/arrow.git -b apache-arrow-${PYARROW_VERSION} && \
-    cd arrow/cpp && \
-    mkdir build && cd build && \
-    cmake -DCMAKE_BUILD_TYPE=release \
-        -DCMAKE_INSTALL_PREFIX=/usr/local \
-        -DARROW_PYTHON=ON \
-        -DARROW_BUILD_TESTS=OFF \
-        -DARROW_JEMALLOC=ON \
-        -DARROW_BUILD_STATIC="OFF" \
-        -DARROW_PARQUET=ON \
-        .. && \
-    make install -j ${MAX_JOBS:-$(nproc)} && \
-    cd ../../python/ && \
-    uv pip install -v -r requirements-build.txt && uv pip install numpy==2.1.3 && \
-    PYARROW_PARALLEL=${PYARROW_PARALLEL:-$(nproc)} \
-    python setup.py build_ext \
-    --build-type=release --bundle-arrow-cpp \
-    bdist_wheel --dist-dir /arrowwheels/
-
-###############################################################
-# Stage to build opencv
-###############################################################
-
-FROM base-builder AS cv-builder
-
-ARG MAX_JOBS
-ARG OPENCV_VERSION=86
-# patch for version 4.11.0.86
-ARG OPENCV_PATCH=97f3f39
-ARG ENABLE_HEADLESS=1
-RUN --mount=type=cache,target=/root/.cache/uv \
-    source /opt/rh/gcc-toolset-14/enable && \
-    git clone --recursive https://github.com/opencv/opencv-python.git -b ${OPENCV_VERSION} && \
-    cd opencv-python && \
-    sed -i -E -e 's/"setuptools.+",/"setuptools",/g' pyproject.toml && \
-    cd opencv && git cherry-pick --no-commit $OPENCV_PATCH && cd .. && \
-    uv pip install scikit-build && \    
-    python -m build --wheel --installer=uv --outdir /opencvwheels/
-
-###############################################################
-# Stage to build numactl
-###############################################################
-
-FROM base-builder AS numa-builder
-
-# Note: Building numactl with gcc-11. Compiling with gcc-13 in this builder stage will
-# trigger recompilation with gcc-11 (and require libtool) in the final stage where we do not have gcc-13
-ARG MAX_JOBS
-ARG NUMACTL_VERSION=2.0.19
-RUN git clone --recursive https://github.com/numactl/numactl.git -b v${NUMACTL_VERSION} \
-    && cd numactl \
-    && autoreconf -i && ./configure \
-    && make -j ${MAX_JOBS:-$(nproc)}
-
-
-###############################################################
-# Stage to build numba 
-###############################################################
-
-FROM base-builder AS numba-builder
-
-ARG MAX_JOBS
-ARG NUMBA_VERSION=0.61.2
-
-# Clone all required dependencies
-RUN dnf install ninja-build llvm15 llvm15-devel -y && source /opt/rh/gcc-toolset-14/enable && export PATH=$PATH:/usr/lib64/llvm15/bin && \
-    git clone --recursive https://github.com/numba/numba.git -b ${NUMBA_VERSION} && \
-    cd ./numba && \
-    if ! grep '#include "dynamic_annotations.h"' numba/_dispatcher.cpp; then \
-       sed -i '/#include "internal\/pycore_atomic.h"/i\#include "dynamic_annotations.h"' numba/_dispatcher.cpp; \
-    fi && python -m build --wheel --installer=uv --outdir /numbawheels/
-
-###############################################################
-# Stage to build vllm - this stage builds and installs
-# vllm, tensorizer and vllm-tgis-adapter and builds uv cache
-# for transitive dependencies - eg. grpcio
-###############################################################
-
-FROM base-builder AS vllmcache-builder
-
-ENV LLVM_CONFIG=/usr/lib64/llvm15/bin/llvm-config
-ENV PATH=/usr/lib64/llvm15/bin:$PATH
-
-COPY --from=torch-builder /tmp/control /dev/null
-COPY --from=arrow-builder /tmp/control /dev/null
-COPY --from=cv-builder /tmp/control /dev/null
-COPY --from=numa-builder /tmp/control /dev/null
-COPY --from=numba-builder /tmp/control /dev/null
-
 ARG VLLM_TARGET_DEVICE=cpu
-ARG GRPC_PYTHON_BUILD_SYSTEM_OPENSSL=1
 
-# this step installs vllm and populates uv cache
-# with all the transitive dependencies
+USER root
+WORKDIR /root
+
+ENV HOME=/root \
+    WHEEL_DIR=/wheelsdir \
+    VIRTUAL_ENV=/opt/vllm \
+    GRPC_PYTHON_BUILD_SYSTEM_OPENSSL=1 \
+    CARGO_HOME=/root/.cargo \
+    RUSTUP_HOME=/root/.rustup \
+    UV_CACHE_DIR=$HOME/.cache/uv \
+    PATH=/root/.cargo/bin:/root/.rustup/bin:${VIRTUAL_ENV}/bin:$PATH
+
+RUN echo "DEBUG: VLLM_VERSION=${VLLM_VERSION}"
+RUN --mount=type=cache,target=/var/cache/dnf \
+    microdnf install -y \
+    python${PYTHON_VERSION}-devel python${PYTHON_VERSION}-pip \
+    && python${PYTHON_VERSION} -m venv ${VIRTUAL_ENV} \
+    && python${PYTHON_VERSION} -m pip install -U pip uv --no-cache
+
+# Important: Copy only bare minimum required for the script to run
+COPY requirements/ requirements/
+COPY pyproject.toml ./
+
+# The script is expected to install whatever python dependencies are missing
+# as well as whatever system libraries need to be installed from source
+COPY  build_vllm_*.sh ./
+
 RUN --mount=type=cache,target=/root/.cache/uv \
-    dnf install llvm15 llvm15-devel -y && \
-    rpm -ivh --nodeps https://mirror.stream.centos.org/9-stream/CRB/ppc64le/os/Packages/protobuf-lite-devel-3.14.0-16.el9.ppc64le.rpm && \
-    source /opt/rh/gcc-toolset-14/enable && \
-    git clone https://github.com/huggingface/xet-core.git && cd xet-core/hf_xet/ && \
-    uv pip install maturin && \
-    uv build --wheel --out-dir /hf_wheels/
+    sh ./build_vllm_$(uname -m).sh
+
+# copy vllm source code to build cache
+COPY . .
 
-ENV CXXFLAGS="-fno-lto -Wno-error=free-nonheap-object" \
-    CFLAGS="-fno-lto"
 RUN --mount=type=cache,target=/root/.cache/uv \
-    --mount=type=bind,from=torch-builder,source=/torchwheels/,target=/torchwheels/,ro \
-    --mount=type=bind,from=arrow-builder,source=/arrowwheels/,target=/arrowwheels/,ro \
-    --mount=type=bind,from=cv-builder,source=/opencvwheels/,target=/opencvwheels/,ro \
-    --mount=type=bind,from=numa-builder,source=/numactl/,target=/numactl/,rw \
-    --mount=type=bind,from=numba-builder,source=/numbawheels/,target=/numbawheels/,ro \
-    --mount=type=bind,src=.,dst=/src/,rw \
     source /opt/rh/gcc-toolset-14/enable && \
-    export PATH=$PATH:/usr/lib64/llvm15/bin && \
-    uv pip install /opencvwheels/*.whl /arrowwheels/*.whl /torchwheels/*.whl /numbawheels/*.whl && \
-    sed -i -e 's/.*torch.*//g' /src/pyproject.toml /src/requirements/*.txt && \
-    sed -i -e 's/.*sentencepiece.*//g' /src/pyproject.toml /src/requirements/*.txt && \
-    uv pip install sentencepiece==0.2.0 pandas pythran nanobind pybind11 /hf_wheels/*.whl && \
-    make -C /numactl install && \
-    # sentencepiece.pc is in some pkgconfig inside uv cache
-    export PKG_CONFIG_PATH=$(find / -type d -name "pkgconfig" 2>/dev/null | tr '\n' ':') && \
-    nanobind_DIR=$(uv pip show nanobind | grep Location | sed 's/^Location: //;s/$/\/nanobind\/cmake/') && uv pip install -r /src/requirements/common.txt -r /src/requirements/cpu.txt -r /src/requirements/build/cuda.txt --no-build-isolation && \
-    cd /src/ && \
-    uv build --wheel --out-dir /vllmwheel/ --no-build-isolation && \
-    uv pip install /vllmwheel/*.whl
+    pip install -U uv
 
+# build & install vLLM so that all transitive dependencies are build/downloaded into the uv cache
+RUN --mount=type=cache,target=/root/.cache/uv \
+    source /opt/rh/gcc-toolset-14/enable && \
+    export PATH=/opt/rh/gcc-toolset-14/root/usr/bin:$PATH && \
+    export CC=/opt/rh/gcc-toolset-14/root/usr/bin/gcc && \
+    export CXX=/opt/rh/gcc-toolset-14/root/usr/bin/g++ && \
+    export PKG_CONFIG_PATH=/usr/local/lib/pkgconfig:/usr/lib64/pkgconfig:$PKG_CONFIG_PATH && \
+    export CMAKE_PREFIX_PATH=/usr/local:/usr:$CMAKE_PREFIX_PATH && \
+    export Protobuf_PROTOC_EXECUTABLE=/usr/bin/protoc && \
+    export CFLAGS="-mcpu=power10 -mtune=power10" && \
+    export CXXFLAGS="-mcpu=power10 -mtune=power10" && \
+    export DNNL_ARCH_OPT_FLAGS="-mcpu=power10 -mtune=power10" && \
+    export C_INCLUDE_PATH=/usr/local/include:$C_INCLUDE_PATH && \
+    export CPLUS_INCLUDE_PATH=/usr/local/include:$CPLUS_INCLUDE_PATH && \
+    uv pip install 'setuptools>=78.1.1' && \
+    export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/opt/OpenBLAS/lib/:/usr/local/lib64:/usr/local/lib && \
+    export LIBGOMP=/opt/rh/gcc-toolset-14/root/usr/lib/gcc/ppc64le-redhat-linux/14/libgomp.so && \
 
-###############################################################
-# Stage to build lapack
-###############################################################
-
-FROM base-builder AS lapack-builder
-
-ARG MAX_JOBS
-ARG LAPACK_VERSION=3.12.1
-RUN git clone --recursive https://github.com/Reference-LAPACK/lapack.git -b v${LAPACK_VERSION} \
-    && cd lapack && source /opt/rh/gcc-toolset-14/enable \
-    && cmake -B build -S . \
-    && cmake --build build -j ${MAX_JOBS:-$(nproc)}
+    export CMAKE_LIBRARY_PATH=$(dirname $LIBGOMP):${CMAKE_LIBRARY_PATH} && \
+    export LIBRARY_PATH=$(dirname $LIBGOMP):${LIBRARY_PATH} && \
+    export LD_LIBRARY_PATH=$(dirname $LIBGOMP):${LD_LIBRARY_PATH} && \
 
+    echo "LIBGOMP=${LIBGOMP}" && \
+    find /root/.cache/uv -name "*.whl" && \
+    SETUPTOOLS_SCM_PRETEND_VERSION="$VLLM_VERSION" uv build \
+      --wheel --out-dir ${WHEEL_DIR} --no-build-isolation && \
+    uv pip install "$(echo ${WHEEL_DIR}/vllm*.whl)[tensorizer]" --refresh
 
 ###############################################################
 #                   FINAL VLLM IMAGE STAGE                    #
@@ -278,72 +83,74 @@ RUN git clone --recursive https://github.com/Reference-LAPACK/lapack.git -b v${L
 FROM registry.access.redhat.com/ubi9/ubi-minimal:${BASE_UBI_IMAGE_TAG} AS vllm-openai
 
 ARG PYTHON_VERSION=3.12
-ARG OPENBLAS_VERSION=0.3.30
+ENV VLLM_NO_USAGE_STATS=1
 
 # Set Environment Variables for venv & openblas
 ENV VIRTUAL_ENV=/opt/vllm
-ENV PATH=${VIRTUAL_ENV}/bin:$PATH
-ENV PKG_CONFIG_PATH=/usr/local/lib/pkgconfig/
-ENV LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/lib64:/usr/local/lib:/usr/lib64:/usr/lib
+ENV PCP_DIR=/opt/rh/gcc-toolset-14/root
+ENV PATH=${VIRTUAL_ENV}/bin:${PCP_DIR}/usr/bin:/usr/local/bin:$PATH
+ENV PKG_CONFIG_PATH=${PCP_DIR}/usr/lib64/pkgconfig:/usr/local/lib/pkgconfig/
+ENV C_INCLUDE_PATH="/usr/local/include:$C_INCLUDE_PATH"
+ENV LD_LIBRARY_PATH=${PCP_DIR}/usr/lib64:${PCP_DIR}/usr/lib:${VIRTUAL_ENV}/lib64/python${PYTHON_VERSION}/site-packages/torch/lib:/usr/local/lib:$LD_LIBRARY_PATH:/usr/local/lib64:/usr/lib64:/usr/lib
 ENV UV_LINK_MODE=copy
-ENV OMP_NUM_THREADS=16
+ARG VLLM_VERSION="0.22.1"
+ARG UV_EXTRA_INDEX_URL="https://wheels.developerfirst.ibm.com/ppc64le/linux/+simple/"
+ENV UV_EXTRA_INDEX_URL=${UV_EXTRA_INDEX_URL}
+ENV UV_INDEX_STRATEGY=first-match
 
-# create artificial dependencies between stages for independent stages to build in parallel
-COPY --from=torch-builder /tmp/control /dev/null
-COPY --from=arrow-builder /tmp/control /dev/null
-COPY --from=cv-builder /tmp/control /dev/null
-COPY --from=vllmcache-builder /tmp/control /dev/null
-COPY --from=numa-builder /tmp/control /dev/null
-COPY --from=lapack-builder /tmp/control /dev/null
-COPY --from=openblas-builder /tmp/control /dev/null
-COPY --from=numba-builder /tmp/control /dev/null
 
-# install gcc-11, python, openblas, numactl, lapack
 RUN --mount=type=cache,target=/root/.cache/uv \
-    --mount=type=bind,from=numa-builder,source=/numactl/,target=/numactl/,rw \
-    --mount=type=bind,from=lapack-builder,source=/lapack/,target=/lapack/,rw \
-    --mount=type=bind,from=openblas-builder,source=/OpenBLAS-$OPENBLAS_VERSION/,target=/openblas/,rw \
     rpm -ivh https://dl.fedoraproject.org/pub/epel/epel-release-latest-9.noarch.rpm && \
     microdnf install --nodocs -y \
-    libomp libicu tar findutils openssl llvm15 llvm15-devel \
-    pkgconfig xsimd g++ gcc-fortran libsndfile \
+    libomp libicu tar autoconf automake libtool findutils openssl numactl numactl-devel \
+    pkgconfig xsimd gcc-toolset-14 libsndfile \
     libtiff libjpeg openjpeg2 zlib zeromq \
     freetype lcms2 libwebp tcl tk utf8proc \
-    harfbuzz fribidi libraqm libimagequant libxcb util-linux \
+    harfbuzz fribidi libraqm libimagequant libxcb util-linux gperftools-libs \
     python${PYTHON_VERSION}-devel python${PYTHON_VERSION}-pip \
-    && export PATH=$PATH:/usr/lib64/llvm15/bin && microdnf clean all \
-    && python${PYTHON_VERSION} -m venv ${VIRTUAL_ENV} \
-    && python -m pip install -U pip uv --no-cache \
-    && make -C /numactl install \
-    && PREFIX=/usr/local make -C /openblas install \
-    && uv pip install 'cmake<4' \
-    && cmake --install /lapack/build \
-    && uv pip uninstall cmake
+    && source /opt/rh/gcc-toolset-14/enable \
+    && microdnf update -y \
+    && microdnf clean all
 
-# consume previously built wheels (including vllm)
-RUN --mount=type=cache,target=/root/.cache/uv \
-    --mount=type=bind,from=torch-builder,source=/torchwheels/,target=/torchwheels/,ro \
-    --mount=type=bind,from=arrow-builder,source=/arrowwheels/,target=/arrowwheels/,ro \
-    --mount=type=bind,from=cv-builder,source=/opencvwheels/,target=/opencvwheels/,ro \
-    --mount=type=bind,from=vllmcache-builder,source=/hf_wheels/,target=/hf_wheels/,ro \
-    --mount=type=bind,from=vllmcache-builder,source=/vllmwheel/,target=/vllmwheel/,ro \
-    --mount=type=bind,from=numba-builder,source=/numbawheels/,target=/numbawheels/,ro \
-    export PKG_CONFIG_PATH=$(find / -type d -name "pkgconfig" 2>/dev/null | tr '\n' ':') && uv pip install sentencepiece==0.2.0 && \
-    HOME=/root uv pip install /opencvwheels/*.whl /arrowwheels/*.whl /torchwheels/*.whl /numbawheels/*.whl /hf_wheels/*.whl /vllmwheel/*.whl
+# The `lscpu` command was added as a requirement in part of https://github.com/vllm-project/vllm/pull/21032, so installing it.
+RUN microdnf install --nodocs -y util-linux && \
+    microdnf clean all
 
+COPY --from=builder-base /usr/lib64/libprotobuf.so.25 /usr/lib64/
+COPY --from=builder-base /usr/lib64/libprotobuf.so.25.0.0 /usr/lib64/
 
-COPY ./ /workspace/vllm
-WORKDIR /workspace/vllm
-ARG GIT_REPO_CHECK=0
-RUN --mount=type=bind,source=.git,target=.git \
-    if [ "$GIT_REPO_CHECK" != 0 ]; then bash tools/check_repo.sh; fi
+# Use builder venv in final stage instead of wheel reinstallation
+COPY --from=builder-base /opt/vllm /opt/vllm
 
-# install development dependencies (for testing)
-RUN --mount=type=cache,target=/root/.cache/uv \
-    uv pip install -e tests/vllm_test_utils
+ENV LD_PRELOAD=/usr/lib64/libtcmalloc.so.4
 
-WORKDIR /workspace/
+WORKDIR /home/vllm
 
-RUN ln -s /workspace/vllm/tests && ln -s /workspace/vllm/examples && ln -s /workspace/vllm/benchmarks
+# setup non-root user for OpenShift
+RUN umask 002 && \
+    useradd --uid 2000 --gid 0 vllm && \
+    mkdir -p /home/vllm && \
+    chmod g+rwx /home/vllm
+
+ENV HOME=/home/vllm
+
+# Add labels to document build configuration
+LABEL org.opencontainers.image.title="vLLM CPU"
+LABEL org.opencontainers.image.description="vLLM inference engine for CPU platforms"
+LABEL org.opencontainers.image.vendor="vLLM Project"
+LABEL org.opencontainers.image.source="https://github.com/vllm-project/vllm"
+
+# Build configuration labels
+ARG TARGETARCH
+ARG VLLM_CPU_PPC64LE
+ARG PYTHON_VERSION
+
+LABEL ai.vllm.build.target-arch="${TARGETARCH}"
+LABEL ai.vllm.build.cpu-ppc64le="${VLLM_CPU_PPC64LE:-false}"
+LABEL ai.vllm.build.python-version="${PYTHON_VERSION:-3.12}"
+
+USER 2000
 
 ENTRYPOINT ["vllm", "serve"]
+
+ 

docker/Dockerfile.rocm

@@ -252,6 +252,8 @@ COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/docker/docker-bake.hcl /docker/doc
 COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/docker/docker-bake-rocm.hcl /docker/docker-bake-rocm.hcl
 COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/.buildkite /.buildkite
 COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/pyproject.toml /pyproject.toml
+COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/rust /rust
+COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/rust-toolchain.toml /rust-toolchain.toml
 COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/vllm/v1 /vllm_v1
 
 # RIXL/UCX build stages
@@ -543,6 +545,8 @@ COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/docker/docker-bake.h
 COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/docker/docker-bake-rocm.hcl /docker/docker-bake-rocm.hcl
 COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/.buildkite /.buildkite
 COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/pyproject.toml /pyproject.toml
+COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/rust /rust
+COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/rust-toolchain.toml /rust-toolchain.toml
 COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/vllm/v1 /vllm_v1
 
 # -----------------------
@@ -576,6 +580,7 @@ RUN apt-get update -q -y && apt-get install -q -y --no-install-recommends \
     libibverbs1 \
     ibverbs-providers \
     ibverbs-utils \
+    unzip \
     pkg-config ffmpeg libavcodec-dev libavformat-dev libavutil-dev \
     libswscale-dev libavdevice-dev libavfilter-dev libswresample-dev \
     && rm -rf /var/lib/apt/lists/*

docker/Dockerfile.rocm_base

@@ -9,7 +9,7 @@ ARG PYTORCH_AUDIO_BRANCH="v2.9.0"
 ARG PYTORCH_AUDIO_REPO="https://github.com/pytorch/audio.git"
 ARG FA_BRANCH="0e60e394"
 ARG FA_REPO="https://github.com/Dao-AILab/flash-attention.git"
-ARG AITER_BRANCH="v0.1.16.post2"
+ARG AITER_BRANCH="v0.1.16.post3"
 ARG AITER_REPO="https://github.com/ROCm/aiter.git"
 ARG MORI_BRANCH="v1.1.0"
 ARG MORI_REPO="https://github.com/ROCm/mori.git"

docker/Dockerfile.s390x

@@ -249,7 +249,7 @@ RUN --mount=type=cache,target=/root/.cache/uv \
      NUMBA_WHL_FILE=$(ls /tmp/numba-wheels/*.whl) && \
      OPENCV_WHL_FILE=$(ls /tmp/opencv-wheels/*.whl) && \
      GUIDANCE_WHL_FILE=$(ls /tmp/guidance-wheels/*.whl) && \
-     uv pip install -v \    
+     uv pip install -v \
         $ARROW_WHL_FILE \
         $VISION_WHL_FILE \
         $HF_XET_WHL_FILE \
@@ -257,6 +257,7 @@ RUN --mount=type=cache,target=/root/.cache/uv \
         $NUMBA_WHL_FILE \
         $OPENCV_WHL_FILE \
         $GUIDANCE_WHL_FILE \
+        --torch-backend cpu \
         --index-strategy unsafe-best-match \
         -r requirements/build/cpu.txt \
         -r requirements/cpu.txt

docs/cli/README.md

@@ -50,6 +50,21 @@ vllm serve --help=max-num-seqs
 vllm serve --help=max
 ```
 
+!!! tip "Human-readable integer arguments"
+    Many integer arguments accept human-readable suffixes for convenience. For example:
+
+    - `1k` = 1,000 (decimal kilo)
+    - `1K` = 1,024 (binary kibibyte)
+    - `1m` = 1,000,000 (decimal mega)
+    - `1M` = 1,048,576 (binary mebibyte)
+    - `1g` / `1G` = 1 billion / 1 gibibyte
+    - `1t` / `1T` = 1 trillion / 1 tebibyte
+    
+    Decimal suffixes (`k`, `m`, `g`, `t`) also accept floating point: `25.6k` = 25,600.
+    Binary suffixes (`K`, `M`, `G`, `T`) require integers: `32K` = 32,768.
+    
+    Supported arguments include: `--max-model-len`, `--max-num-batched-tokens`, `--max-num-scheduled-tokens`, `--kv-cache-memory-bytes`, `--safetensors-prefetch-block-size`.
+
 See [vllm serve](./serve.md) for the full reference of all available arguments.
 
 ## launch

docs/configuration/conserving_memory.md

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

docs/configuration/optimization.md

@@ -16,6 +16,14 @@ vLLM provides 4 optimization levels (`-O0`, `-O1`, `-O2`, `-O3`) that allow user
 
 For more information, see the [optimization level documentation](../design/optimization_levels.md).
 
+## Faster Startup
+
+Beyond the optimization levels, three mechanisms reduce time-to-first-token on repeated boots of the same (model, config, hardware) combination:
+
+- **Reuse the compile cache.** vLLM persists `torch.compile` artifacts under `VLLM_CACHE_ROOT` (default `~/.cache/vllm`), and the cache directory can be copied between machines or baked into a container image; see the [torch.compile design doc](../design/torch_compile.md). Set `VLLM_FORCE_AOT_LOAD=1` to fail loudly instead of silently recompiling when the cache misses (any change to the model, config, relevant `VLLM_*` environment variables, torch build, or GPU model invalidates it).
+- **Skip memory profiling with `--kv-cache-memory`.** On startup, vLLM logs the exact `--kv-cache-memory` value that reproduces the current allocation. Passing it back on the next boot skips the memory-profiling measurement and the CUDA-graph memory estimation pass. Note that this has performance implications: the KV cache is sized to exactly the given value instead of being measured, so a conservative value caps batch concurrency (and therefore throughput), while an optimistic one fails at allocation time. The value is only valid on the same GPU with the same initial free memory; if a boot OOMs after hardware or co-tenant changes, remove the flag to re-profile.
+- **Serve without CUDA graphs using `--enforce-eager`.** Skips both compilation and CUDA-graph capture for the fastest possible startup, at the cost of steady-state decode performance. Useful for development loops and for measuring how much of a boot is compile/capture.
+
 ## Preemption
 
 Due to the autoregressive nature of transformer architecture, there are times when KV cache space is insufficient to handle all batched requests.
@@ -276,8 +284,9 @@ By default vLLM uses the standard Hugging Face `tokenizers` library to power
 the fast tokenizer. For BPE tokenizers (Qwen, Llama, DeepSeek, GPT-OSS, etc.)
 you can switch to the [fastokens](https://github.com/crusoecloud/fastokens)
 Rust backend, a drop-in replacement that's substantially faster on
-encode/decode and on streaming detokenization. Enable it by setting
-`VLLM_USE_FASTOKENS=1`:
+encode/decode and on streaming detokenization. `VLLM_USE_FASTOKENS` is
+available in vLLM v0.23.0 and later. If your installed vLLM version does not
+recognize the environment variable, upgrade vLLM before enabling the override:
 
 ```console
 VLLM_USE_FASTOKENS=1 vllm serve Qwen/Qwen3-8B