Bump actions/setup-python from 6.1.0 to 6.3.0

Bumps [actions/setup-python](https://github.com/actions/setup-python) from 6.1.0 to 6.3.0.
- [Release notes](https://github.com/actions/setup-python/releases)
- [Commits](https://github.com/actions/setup-python/compare/83679a892e2d95755f2dac6acb0bfd1e9ac5d548...ece7cb06caefa5fff74198d8649806c4678c61a1)

---
updated-dependencies:
- dependency-name: actions/setup-python
  dependency-version: 6.3.0
  dependency-type: direct:production
  update-type: version-update:semver-minor
...

Signed-off-by: dependabot[bot] <support@github.com>

.buildkite/ci_config_intel.yaml

@@ -3,7 +3,6 @@ job_dirs:
   - ".buildkite/intel_jobs"
 run_all_patterns:
   - ".buildkite/ci_config_intel.yaml"
-  - ".buildkite/scripts/hardware_ci/run-intel-test.sh"
   - "docker/Dockerfile"
   - "docker/Dockerfile.xpu"
   - "CMakeLists.txt"

.buildkite/ci_config_rocm.yaml

@@ -8,7 +8,6 @@ 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,30 +1,5 @@
 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  # Transient Docker/BuildKit failure
-          limit: 1
-        - 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
@@ -32,24 +7,19 @@ 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/rocm/build-ci-base.sh
+      - bash .buildkite/scripts/ci-bake-rocm.sh ci-base-rocm-ci-with-deps
     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"
       VLLM_BRANCH: "$BUILDKITE_COMMIT"
     retry:
       automatic:
-        - exit_status: 1  # Transient Docker/BuildKit failure
-          limit: 1
         - exit_status: -1  # Agent was lost
           limit: 1
         - exit_status: -10  # Agent was lost
@@ -63,12 +33,35 @@ steps:
     device: amd_cpu
     no_plugin: true
     commands:
-      - bash .buildkite/scripts/rocm/build-test-image.sh
+      - |
-      - bash .buildkite/scripts/rocm/smoke-test-image.sh
+        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
+        '
     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"
@@ -76,8 +69,6 @@ steps:
       VLLM_BRANCH: "$BUILDKITE_COMMIT"
     retry:
       automatic:
-        - exit_status: 1  # Transient Docker/BuildKit failure
-          limit: 1
         - exit_status: -1  # Agent was lost
           limit: 1
         - exit_status: -10  # Agent was lost

.buildkite/hardware_tests/cpu.yaml

@@ -17,22 +17,16 @@ 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
-  - tests/kernels/mamba/test_causal_conv1d.py
-  - tests/kernels/mamba/test_mamba_ssm.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
-      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/mamba/test_causal_conv1d.py
-      pytest -x -v -s tests/kernels/mamba/test_mamba_ssm.py"
 
 # Note: SDE can't be downloaded from CI host because of AWS WAF
 # - label: CPU-Compatibility Tests
@@ -145,22 +139,9 @@ steps:
   commands:
     - |
       bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 45m "
-      pytest -x -v -s tests/models/multimodal/generation --ignore=tests/models/multimodal/generation/test_pixtral.py --ignore=tests/models/multimodal/generation/test_qwen2_5_vl.py -m cpu_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB"
+      pytest -x -v -s tests/models/multimodal/generation --ignore=tests/models/multimodal/generation/test_pixtral.py -m cpu_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB"
   parallelism: 4
 
-- label: CPU-Qwen2.5-VL Multimodal Tests
-  depends_on: []
-  device: intel_cpu
-  no_plugin: true
-  source_file_dependencies:
-  # - vllm/
-  - vllm/model_executor/layers/rotary_embedding
-  - tests/models/multimodal/generation/
-  commands:
-    - |
-      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 40m "
-      VLLM_CI_ENV=0 pytest -x -v -s tests/models/multimodal/generation/test_qwen2_5_vl.py"
-
 - label: "Arm CPU Test"
   depends_on: []
   soft_fail: false

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

@@ -18,7 +18,7 @@ steps:
   - label: "XPU example Test"
     depends_on:
       - image-build-xpu
-    timeout_in_minutes: 50
+    timeout_in_minutes: 30
     optional: true
     device: intel_gpu
     agent_tags:
@@ -39,13 +39,13 @@ steps:
   - label: "XPU V1 test"
     depends_on:
       - image-build-xpu
-    timeout_in_minutes: 70
+    timeout_in_minutes: 30
     optional: true
     device: intel_gpu
     agent_tags:
       label: production
       gpu: 1+
-      mem: 24+
+      mem: 16+
     no_plugin: true
     env:
       REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -60,7 +60,7 @@ steps:
   - label: "XPU server test"
     depends_on:
       - image-build-xpu
-    timeout_in_minutes: 45
+    timeout_in_minutes: 30
     optional: true
     device: intel_gpu
     agent_tags:

.buildkite/image_build/image_build.sh

@@ -79,18 +79,12 @@ 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"
@@ -177,18 +171,6 @@ 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}"
@@ -272,5 +254,3 @@ 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,31 +9,30 @@ 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 docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
+if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64) ]]; then
-  echo "Image found"
-else
   echo "Image not found, proceeding with build..."
-  # build for arm64 GPU targets: Grace/GH200 (sm_90),
+else
-  # Blackwell/Thor (sm_100/sm_103/sm_110), and DGX Spark/GB10
+  echo "Image found"
-  # (sm_121, family-covered by 12.0 under CUDA 13)
+  exit 0
-  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
 
-.buildkite/scripts/annotate-image-build.sh "$IMAGE"
+# 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/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 docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
+if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-cpu) ]]; then
-  echo "Image found"
-else
   echo "Image not found, proceeding with build..."
-  # build
+else
-  docker build --file docker/Dockerfile.cpu \
+  echo "Image found"
-    --build-arg max_jobs=16 \
+  exit 0
-    --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
-    --build-arg VLLM_CPU_X86=true \
-    --tag "$IMAGE" \
-    --target vllm-test \
-    --progress plain .
-  # push
-  docker push "$IMAGE"
 fi
 
-.buildkite/scripts/annotate-image-build.sh "$IMAGE"
+# 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/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 docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
+if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64-cpu) ]]; then
-  echo "Image found"
-else
   echo "Image not found, proceeding with build..."
-  # build
+else
-  docker build --file docker/Dockerfile.cpu \
+  echo "Image found"
-    --build-arg max_jobs=16 \
+  exit 0
-    --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
-    --tag "$IMAGE" \
-    --target vllm-test \
-    --progress plain .
-  # push
-  docker push "$IMAGE"
 fi
 
-.buildkite/scripts/annotate-image-build.sh "$IMAGE"
+# 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/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 docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
+if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-hpu) ]]; then
-  echo "Image found"
-else
   echo "Image not found, proceeding with build..."
-  # build
+else
-  docker build \
+  echo "Image found"
-    --file tests/pytorch_ci_hud_benchmark/Dockerfile.hpu \
+  exit 0
-    --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
 
-.buildkite/scripts/annotate-image-build.sh "$IMAGE"
+# 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/image_build/image_build_torch_nightly.sh

@@ -40,7 +40,6 @@ 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..."
@@ -67,5 +66,3 @@ 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 "$IMAGE" &> /dev/null; then
+if ! docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu &> /dev/null; then
-  echo "Image found"
-else
   echo "Image not found, proceeding with build..."
-  # build
+else
-  docker build \
+  echo "Image found"
-    --file docker/Dockerfile.xpu \
+  exit 0
-    --build-arg max_jobs=16 \
-    --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
-    --tag "$IMAGE" \
-    --progress plain .
-  # push
-  docker push "$IMAGE"
 fi
 
-.buildkite/scripts/annotate-image-build.sh "$IMAGE"
+# build
+docker build \
+  --file docker/Dockerfile.xpu \
+  --build-arg max_jobs=16 \
+  --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
+  --tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu \
+  --progress plain .
+
+# push
+docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu

.buildkite/intel_jobs/basic_correctness.yaml

@@ -3,7 +3,7 @@ depends_on:
   - image-build-xpu
 steps:
 - label: XPU Sleep Mode
-  timeout_in_minutes: 45
+  timeout_in_minutes: 30
   device: intel_gpu
   agent_tags:
     label: production

.buildkite/intel_jobs/lora_intel.yaml

@@ -81,12 +81,10 @@ steps:
       'cd tests &&
       export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       set -o pipefail &&
-      pytest -v -s lora/test_punica_ops.py::test_kernels &&
+      pytest -v -s lora/test_punica_ops.py --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-3-43264-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype1-1-2049-64-128-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-1-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-1-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-8-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype0-3-2049-128-8-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-8-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype1-1-2049-256-128-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-64256-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-2-29696-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-3-49408-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-2-16384-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-2-51328-32-4-4]"'
-      pytest -v -s lora/test_punica_ops.py::test_kernels_hidden_size &&
-      pytest -v -s lora/test_punica_ops.py::test_add_lora_fused_moe_early_exit'
 
 - label: LoRA Punica FP8/XPU Ops
-  timeout_in_minutes: 60
+  timeout_in_minutes: 45
   device: intel_gpu
   agent_tags:
     label: production

.buildkite/intel_jobs/misc_intel.yaml

@@ -3,7 +3,7 @@ depends_on:
   - image-build-xpu
 steps:
 - label: V1 Core + KV + Metrics
-  timeout_in_minutes: 45
+  timeout_in_minutes: 30
   device: intel_gpu
   agent_tags:
     label: production
@@ -33,12 +33,12 @@ steps:
       pytest -v -s v1/executor'
 
 - label: V1 Sample + Logits
-  timeout_in_minutes: 90
+  timeout_in_minutes: 30
   device: intel_gpu
   agent_tags:
     label: production
     gpu: 1+
-    mem: 24+
+    mem: 16+
   no_plugin: true
   working_dir: "."
   env:
@@ -72,31 +72,9 @@ 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'
+      pytest -v -s v1/sample/test_topk_topp_sampler.py &&
-
+      pytest -v -s v1/sample/test_logprobs.py &&
-- label: Basic Models Tests (Initialization)
+      pytest -v -s v1/sample/test_logprobs_e2e.py'
-  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
@@ -125,34 +103,9 @@ steps:
       pytest -v -s v1/kv_offload &&
       pytest -v -s v1/kv_connector/unit/test_offloading_connector.py'
 
-- label: NixlConnector PD accuracy (2 GPUs)
-  timeout_in_minutes: 60
-  num_devices: 2
-  device: intel_gpu
-  agent_tags:
-    label: production
-    gpu: 2+
-    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/distributed/kv_transfer/kv_connector/v1/nixl/
-    - vllm/v1/worker/kv_connector_model_runner_mixin.py
-    - tests/v1/kv_connector/nixl_integration/
-    - vllm/platforms/xpu.py
-  commands:
-    - >-
-      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'cd tests &&
-      bash v1/kv_connector/nixl_integration/run_xpu_disagg_accuracy_test.sh'
-
 - label: Regression
   key: regression
-  timeout_in_minutes: 50
+  timeout_in_minutes: 30
   device: intel_gpu
   agent_tags:
     label: production
@@ -180,13 +133,13 @@ steps:
   commands:
   - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'pip install modelscope\<1.38 &&
+      'pip install modelscope &&
       cd tests &&
       pytest -v -s test_regression.py'
 
 - label: Metrics, Tracing (2 GPUs)
   key: metrics-tracing-2-gpus
-  timeout_in_minutes: 45
+  timeout_in_minutes: 30
   num_devices: 2
   device: intel_gpu
   agent_tags:
@@ -222,7 +175,7 @@ steps:
 
 - label: Async Engine, Inputs, Utils, Worker
   key: async-engine-inputs-utils-worker
-  timeout_in_minutes: 55
+  timeout_in_minutes: 30
   device: intel_gpu
   agent_tags:
     label: production

.buildkite/intel_jobs/models_distributed_intel.yaml

@@ -4,12 +4,12 @@ depends_on:
 steps:
 - label: Distributed Model Tests (2 GPUs)
   key: distributed-model-tests-2-gpus
-  timeout_in_minutes: 65
+  timeout_in_minutes: 50
   device: intel_gpu
   agent_tags:
     label: production
     gpu: 2+
-    mem: 16+
+    mem: 24+
   no_plugin: true
   working_dir: "."
   env:

.buildkite/intel_jobs/models_multimodal_intel.yaml

@@ -4,32 +4,7 @@ depends_on:
 steps:
 - label: "Multi-Modal Models (Standard) 1: qwen2"
   key: multi-modal-models-standard-1-qwen2
-  timeout_in_minutes: 70
+  timeout_in_minutes: 45
-  device: intel_gpu
-  agent_tags:
-    label: production
-    gpu: 1+
-    mem: 24+
-  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/multimodal
-  commands:
-    - >-
-      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'pip install av &&
-      cd tests &&
-      pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen2" &&
-      pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model'
-
-- label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
-  key: multi-modal-models-standard-2-qwen3-gemma
-  timeout_in_minutes: 70
   device: intel_gpu
   agent_tags:
     label: production
@@ -47,17 +22,19 @@ steps:
   commands:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'cd tests &&
+      'pip install av git+https://github.com/TIGER-AI-Lab/Mantis.git &&
-      pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model'
+      cd tests &&
+      pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen2" &&
+      pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model'
 
-- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
+- label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
-  key: multi-modal-models-standard-3-llava-qwen2-vl
+  key: multi-modal-models-standard-2-qwen3-gemma
-  timeout_in_minutes: 65
+  timeout_in_minutes: 45
   device: intel_gpu
   agent_tags:
     label: production
     gpu: 1+
-    mem: 24+
+    mem: 16+
   no_plugin: true
   working_dir: "."
   env:
@@ -70,12 +47,12 @@ steps:
   commands:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'cd tests &&
+      'pip install git+https://github.com/TIGER-AI-Lab/Mantis.git &&
-      pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "not qwen2 and not qwen3 and not gemma" &&
+      cd tests &&
-      pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model'
+      pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model'
 
-- label: "Multi-Modal Models (Standard) 4: other + whisper"
+- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
-  key: multi-modal-models-standard-4-other-whisper
+  key: multi-modal-models-standard-3-llava-qwen2-vl
   timeout_in_minutes: 45
   device: intel_gpu
   agent_tags:
@@ -94,18 +71,43 @@ steps:
   commands:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'pip install av &&
+      'pip install git+https://github.com/TIGER-AI-Lab/Mantis.git &&
+      cd tests &&
+      pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "not qwen2 and not qwen3 and not gemma" &&
+      pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model'
+
+- label: "Multi-Modal Models (Standard) 4: other + whisper"
+  key: multi-modal-models-standard-4-other-whisper
+  timeout_in_minutes: 45
+  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/multimodal
+  commands:
+    - >-
+      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+      'pip install av git+https://github.com/TIGER-AI-Lab/Mantis.git &&
       cd tests &&
       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/processing'
 
 - label: Multi-Modal Processor # 44min
   key: multi-modal-processor
-  timeout_in_minutes: 60
+  timeout_in_minutes: 45
   device: intel_gpu
   agent_tags:
     label: production
     gpu: 1+
-    mem: 24+
+    mem: 16+
   no_plugin: true
   working_dir: "."
   env:
@@ -119,7 +121,7 @@ steps:
   commands:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-      'pip install av matplotlib ftfy &&
+      'pip install av matplotlib ftfy git+https://github.com/TIGER-AI-Lab/Mantis.git &&
       pip install open-clip-torch --no-deps &&
       cd tests &&
       pytest -v -s models/multimodal/processing/test_tensor_schema.py

.buildkite/intel_jobs/quantization.yaml

@@ -1,28 +0,0 @@
-group: Quantization
-depends_on:
-  - image-build-xpu
-steps:
-- label: Quantization
-  key: quantization
-  timeout_in_minutes: 30
-  env:
-    REGISTRY: "public.ecr.aws/q9t5s3a7"
-    REPO: "vllm-ci-test-repo"
-    VLLM_TEST_DEVICE: "xpu"
-  no_plugin: true
-  working_dir: "."
-  device: intel_gpu
-  agent_tags:
-    label: production
-    gpu: 1+
-    mem: 16+
-  source_file_dependencies:
-  - csrc/
-  - vllm/model_executor/layers/quantization
-  - tests/quantization
-  commands:
-  # - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py
-  - >-
-    bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-    'VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s tests/quantization/test_per_token_kv_cache.py --deselect="tests/quantization/test_per_token_kv_cache.py::test_triton_unified_attention_per_token_head_scale[int4-16-128-num_heads0-seq_lens1]"'
-

.buildkite/intel_jobs/test-intel.yaml

@@ -17,7 +17,7 @@ steps:
   - label: "XPU example Test"
     depends_on:
       - image-build-xpu
-    timeout_in_minutes: 50
+    timeout_in_minutes: 30
     device: intel_gpu
     agent_tags:
       label: production
@@ -42,46 +42,21 @@ steps:
         python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --kv-cache-dtype fp8 &&
         python3 examples/basic/offline_inference/generate.py --model nvidia/Llama-3.1-8B-Instruct-FP8 --block-size 64 --enforce-eager --quantization modelopt --kv-cache-dtype fp8 --attention-backend TRITON_ATTN --max-model-len 4096 &&
         python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --block-size 64 --enforce-eager --max-model-len 8192 &&
-        python3 examples/basic/offline_inference/generate.py --model TheBloke/TinyLlama-1.1B-Chat-v0.3-AWQ --block-size 64 --enforce-eager &&
         python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 &&
         python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 --enable-expert-parallel &&
         python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --max-model-len 8192 &&
         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
-    timeout_in_minutes: 70
+    timeout_in_minutes: 30
     device: intel_gpu
     agent_tags:
       label: production
       gpu: 1+
-      mem: 24+
+      mem: 16+
     no_plugin: true
     env:
       REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -93,19 +68,19 @@ steps:
       - >-
         bash .buildkite/scripts/hardware_ci/run-intel-test.sh
         'cd tests &&
+        bash v1/kv_connector/nixl_integration/run_xpu_disagg_accuracy_test.sh &&
         pytest -v -s v1/core --ignore=v1/core/test_reset_prefix_cache_e2e.py --ignore=v1/core/test_scheduler_e2e.py &&
         pytest -v -s v1/engine --ignore=v1/engine/test_output_processor.py &&
         pytest -v -s v1/sample --ignore=v1/sample/test_logprobs.py --ignore=v1/sample/test_logprobs_e2e.py -k "not test_topk_only and not test_topp_only and not test_topk_and_topp" &&
         pytest -v -s v1/worker --ignore=v1/worker/test_gpu_model_runner.py --ignore=v1/worker/test_worker_memory_snapshot.py &&
         pytest -v -s v1/structured_output &&
         pytest -v -s v1/test_serial_utils.py &&
-        pytest -v -s v1/e2e/general/test_correctness_sliding_window.py --deselect="tests/v1/e2e/general/test_correctness_sliding_window.py::test_sliding_window_retrieval[True-1-5-google/gemma-3-1b-it]" &&
         pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py --ignore=v1/spec_decode/test_speculators_correctness.py &&
         pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --ignore=v1/kv_connector/unit/test_lmcache_integration.py --ignore=v1/kv_connector/unit/test_hf3fs_client.py --ignore=v1/kv_connector/unit/test_hf3fs_connector.py --ignore=v1/kv_connector/unit/test_hf3fs_metadata_server.py --ignore=v1/kv_connector/unit/test_offloading_connector.py'
   - label: "XPU server test"
     depends_on:
       - image-build-xpu
-    timeout_in_minutes: 45
+    timeout_in_minutes: 30
     device: intel_gpu
     agent_tags:
       label: production
@@ -145,27 +120,4 @@ 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/lm-eval-harness/configs/Meta-Llama-4-Maverick-17B-128E-Instruct-FP8.yaml

@@ -1,7 +1,6 @@
 # For hf script, without -t option (tensor parallel size).
 # bash .buildkite/lm-eval-harness/run-lm-eval-mmlupro-vllm-baseline.sh -m meta-llama/Llama-4-Maverick-17B-128E-Instruct-FP8 -l 250 -t 8 -f 5
 model_name: "meta-llama/Llama-4-Maverick-17B-128E-Instruct-FP8"
-rocm_safetensors_load_strategy: lazy
 required_gpu_arch:
   - gfx942
   - gfx950

.buildkite/lm-eval-harness/test_lm_eval_correctness.py

@@ -72,11 +72,6 @@ def launch_lm_eval(eval_config, tp_size):
     if moe_backend is not None:
         model_args += f"moe_backend={moe_backend},"
 
-    if current_platform.is_rocm():
-        rocm_load_strategy = eval_config.get("rocm_safetensors_load_strategy")
-        if rocm_load_strategy is not None:
-            model_args += f"safetensors_load_strategy={rocm_load_strategy},"
-
     env_vars = eval_config.get("env_vars", None)
     with scoped_env_vars(env_vars):
         results = lm_eval.simple_evaluate(

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

@@ -3,8 +3,7 @@
 # SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 #
 # Append a build artifact line to the Buildkite annotation.
-# Usage: annotate-build-artifact.sh <label> <value> <context>
+# Usage: annotate-build-artifact.sh <label> <value>
 set -e
 echo "- **${1}**: \`${2}\`" | \
-  buildkite-agent annotate --append --style 'info' \
+  buildkite-agent annotate --append --style 'info' --context 'release-artifacts'
-    --context "${3:?context is required}"

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

@@ -1,36 +0,0 @@
-#!/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/build-macos-wheel.sh

@@ -1,32 +0,0 @@
-#!/usr/bin/env bash
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-#
-# Build the macOS arm64 CPU wheel natively on a macOS agent (the `macmini`
-# queue) into artifacts/dist/ for upload-nightly-wheels.sh.
-
-set -euo pipefail
-
-# The Rust frontend build needs protoc.
-if ! command -v protoc >/dev/null 2>&1; then
-  brew install protobuf
-fi
-
-# upload-nightly-wheels.sh expects exactly one wheel.
-rm -rf artifacts/dist
-mkdir -p artifacts/dist
-
-export VLLM_TARGET_DEVICE=cpu
-export VLLM_REQUIRE_RUST_FRONTEND=1
-export MACOSX_DEPLOYMENT_TARGET=11.0
-# uv's CPython is universal2; force an arm64-only build and tag so the wheel
-# isn't mislabelled universal2 and installed on Intel Macs where import fails.
-export ARCHFLAGS="-arch arm64"
-export _PYTHON_HOST_PLATFORM="macosx-11.0-arm64"
-export CMAKE_BUILD_PARALLEL_LEVEL="${CMAKE_BUILD_PARALLEL_LEVEL:-4}"
-
-uv venv --python 3.12
-uv pip install -r requirements/build/cpu.txt --index-strategy unsafe-best-match
-uv build --wheel --no-build-isolation -o artifacts/dist
-
-ls -l artifacts/dist/*.whl

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

@@ -29,11 +29,7 @@ if python3 -c "import torch; assert torch.version.hip" 2>/dev/null; then
         TORCH_INDEX_URL=""
     fi
 else
-    if [ "${TORCH_NIGHTLY:-0}" = "1" ]; then
+    TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
-        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,9 +15,9 @@ 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 tools/install_protoc.sh rust-toolchain.toml tests/vllm_test_utils .buildkite/scripts/ci-bake-rocm.sh .buildkite/scripts/rocm/build-ci-base.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"
 DEFAULT_CI_BASE_DOCKERFILE="docker/Dockerfile.rocm"
-DEFAULT_CI_BASE_DOCKERFILE_STAGES="base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust-toolchain build_nixl build_rocshmem build_deepep mori_base ci_base"
+DEFAULT_CI_BASE_DOCKERFILE_STAGES="base build_rixl build_rocshmem build_deepep mori_base ci_base"
 DEFAULT_CI_BASE_METADATA_VERSION="1"
 IMAGE_EXISTED_BEFORE_BUILD=0
 
@@ -285,7 +285,7 @@ get_content_arg_names() {
     fi | awk 'NF && !seen[$0]++'
 }
 
-compute_ci_base_content_hash_once() {
+compute_ci_base_content_hash() {
     local -a content_paths=()
     local -a content_args=()
     local dockerfile="${CI_BASE_DOCKERFILE:-}"
@@ -301,8 +301,7 @@ compute_ci_base_content_hash_once() {
         if [[ -n "${dockerfile}" ]]; then
             printf 'dockerfile:%s\n' "${dockerfile}"
             printf 'resolved-build-args:\n'
-            hash_dockerfile_arg_values "${dockerfile}" "${content_args[@]}" \
+            hash_dockerfile_arg_values "${dockerfile}" "${content_args[@]}"
-                || return 1
             if [[ -n "${stages}" ]]; then
                 printf 'dockerfile-stages:%s\n' "${stages}"
                 if [[ -f "${dockerfile}" ]]; then
@@ -315,53 +314,6 @@ compute_ci_base_content_hash_once() {
     } | sha256sum | cut -d' ' -f1
 }
 
-compute_ci_base_content_hash() {
-    local attempts="${CI_BASE_HASH_ATTEMPTS:-3}"
-    local delay_secs="${CI_BASE_HASH_RETRY_DELAY:-5}"
-    local attempt=0
-    local hash=""
-    local failed=0
-    local -a hashes=()
-
-    if [[ ! "${attempts}" =~ ^[1-9][0-9]*$ ]]; then
-        echo "Invalid CI_BASE_HASH_ATTEMPTS: ${attempts}" >&2
-        return 1
-    fi
-    if [[ ! "${delay_secs}" =~ ^[0-9]+$ ]]; then
-        echo "Invalid CI_BASE_HASH_RETRY_DELAY: ${delay_secs}" >&2
-        return 1
-    fi
-
-    for ((attempt = 1; attempt <= attempts; attempt++)); do
-        if ! hash=$(compute_ci_base_content_hash_once); then
-            echo "ci_base content hash calculation ${attempt}/${attempts} failed" >&2
-            failed=1
-        else
-            hashes+=("${hash}")
-            echo "ci_base content hash calculation ${attempt}/${attempts}: ${hash}" >&2
-        fi
-
-        if ((attempt < attempts)); then
-            sleep "${delay_secs}"
-        fi
-    done
-
-    if ((failed)) || ((${#hashes[@]} != attempts)); then
-        echo "Could not calculate a reliable ci_base content hash" >&2
-        return 1
-    fi
-
-    for hash in "${hashes[@]:1}"; do
-        if [[ "${hash}" != "${hashes[0]}" ]]; then
-            echo "ci_base content hash changed between calculations" >&2
-            printf '  observed: %s\n' "${hashes[@]}" >&2
-            return 1
-        fi
-    done
-
-    printf '%s\n' "${hashes[0]}"
-}
-
 extract_dockerfile_arg_default() {
     local dockerfile="$1"
     local arg_name="$2"
@@ -414,11 +366,7 @@ hash_dockerfile_arg_values() {
         printf 'arg:%s=%s\n' "${arg_name}" "${arg_value:-<empty>}"
         if [[ "${arg_name}" == "BASE_IMAGE" && -n "${arg_value}" ]]; then
             digest=$(resolve_image_digest "${arg_value}")
-            if [[ -z "${digest}" ]]; then
+            printf 'arg:%s.digest=%s\n' "${arg_name}" "${digest:-unknown}"
-                echo "Failed to resolve digest for BASE_IMAGE=${arg_value}" >&2
-                return 1
-            fi
-            printf 'arg:%s.digest=%s\n' "${arg_name}" "${digest}"
         fi
     done
 }
@@ -445,16 +393,6 @@ 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"
@@ -795,8 +733,6 @@ 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}"
@@ -814,10 +750,6 @@ 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_REF:-}"
-        set_buildkite_metadata "rocm-ci-base-image-stable" "${CI_BASE_IMAGE_TAG_STABLE:-}"
         return 0
     fi
 
@@ -1159,8 +1091,8 @@ ci_base_metadata_pairs() {
     metadata_pair "vllm.rocm.nic_backend" "$(resolve_dockerfile_arg_value "${dockerfile}" "NIC_BACKEND")"
     metadata_pair "vllm.rocm.ainic_version" "$(resolve_dockerfile_arg_value "${dockerfile}" "AINIC_VERSION")"
     metadata_pair "vllm.rocm.ubuntu_codename" "$(resolve_dockerfile_arg_value "${dockerfile}" "UBUNTU_CODENAME")"
-    metadata_pair "vllm.rocm.nixl_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "NIXL_REPO")"
+    metadata_pair "vllm.rocm.rixl_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "RIXL_REPO")"
-    metadata_pair "vllm.rocm.nixl_commit" "${NIXL_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "NIXL_BRANCH")}"
+    metadata_pair "vllm.rocm.rixl_commit" "${RIXL_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "RIXL_BRANCH")}"
     metadata_pair "vllm.rocm.ucx_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "UCX_REPO")"
     metadata_pair "vllm.rocm.ucx_commit" "${UCX_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "UCX_BRANCH")}"
     metadata_pair "vllm.rocm.rocshmem_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "ROCSHMEM_REPO")"
@@ -1169,7 +1101,7 @@ ci_base_metadata_pairs() {
     metadata_pair "vllm.rocm.deepep_commit" "${DEEPEP_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_BRANCH")}"
     metadata_pair "vllm.rocm.deepep_nic" "$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_NIC")"
     metadata_pair "vllm.rocm.deepep_rocm_arch" "$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_ROCM_ARCH")"
-    metadata_pair "vllm.rocm.nixl_cache_key" "${NIXL_CACHE_KEY:-}"
+    metadata_pair "vllm.rocm.rixl_cache_key" "${RIXL_CACHE_KEY:-}"
     metadata_pair "vllm.rocm.rocshmem_cache_key" "${ROCSHMEM_CACHE_KEY:-}"
     metadata_pair "vllm.rocm.deepep_cache_key" "${DEEPEP_CACHE_KEY:-}"
 
@@ -1263,24 +1195,12 @@ uses_rocm_csrc_cache() {
     esac
 }
 
-uses_rocm_rust_cache() {
-    case "${TARGET}" in
-        rust-rocm-ci|test-rocm-ci|test-rocm-ci-with-wheel|test-rocm-ci-with-artifacts|export-wheel-rocm)
-            return 0
-            ;;
-        *)
-            return 1
-            ;;
-    esac
-}
-
 compute_rocm_csrc_content_hash() {
     local bake_dir=""
     local dockerfile_rocm=""
     local -a content_paths=(
         "requirements/common.txt"
         "requirements/rocm.txt"
-        "pyproject.toml"
         "setup.py"
         "CMakeLists.txt"
         "cmake"
@@ -1324,56 +1244,6 @@ compute_rocm_csrc_content_hash_if_needed() {
     echo "ROCm csrc content cache ref: ${ROCM_CSRC_CONTENT_CACHE_REF}"
 }
 
-compute_rocm_rust_content_hash() {
-    local bake_dir=""
-    local dockerfile_rocm=""
-    local -a content_paths=(
-        "requirements/build/rust.txt"
-        "rust/Cargo.lock"
-        "rust/Cargo.toml"
-        "rust/proto"
-        "rust/src"
-        "rust-toolchain.toml"
-        "tools/build_rust.py"
-        "tools/install_protoc.sh"
-        "build_rust.sh"
-    )
-    local -a content_args=()
-
-    bake_dir=$(dirname "${VLLM_BAKE_FILE}")
-    dockerfile_rocm="${bake_dir}/Dockerfile.rocm"
-    mapfile -t content_args < <(
-        get_content_arg_names "${dockerfile_rocm}" "base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build" "${ROCM_RUST_CONTENT_ARGS:-}"
-    )
-
-    {
-        printf 'rust-input-files-hash:%s\n' "$(compute_content_hash "${content_paths[@]}")"
-        printf 'dockerfile:%s\n' "${dockerfile_rocm}"
-        printf 'resolved-build-args:\n'
-        hash_dockerfile_arg_values "${dockerfile_rocm}" "${content_args[@]}"
-        printf 'dockerfile-stages:base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build\n'
-        if [[ -f "${dockerfile_rocm}" ]]; then
-            hash_dockerfile_stages "${dockerfile_rocm}" "base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build"
-        else
-            printf 'missing:%s\n' "${dockerfile_rocm}"
-        fi
-    } | sha256sum | cut -d' ' -f1
-}
-
-compute_rocm_rust_content_hash_if_needed() {
-    local cache_repo="${DOCKERHUB_CACHE_REPO:-rocm/vllm-ci-cache}"
-
-    if [[ "${ROCM_RUST_CONTENT_CACHE:-1}" == "0" ]] || ! uses_rocm_rust_cache; then
-        return 0
-    fi
-
-    ROCM_RUST_CONTENT_HASH=$(compute_rocm_rust_content_hash)
-    ROCM_RUST_CONTENT_CACHE_REF="${cache_repo}:rust-rocm-input-${ROCM_RUST_CONTENT_HASH}"
-    export ROCM_RUST_CONTENT_HASH
-    export ROCM_RUST_CONTENT_CACHE_REF
-    echo "ROCm Rust content cache ref: ${ROCM_RUST_CONTENT_CACHE_REF}"
-}
-
 write_hcl_string_list_entries() {
     local indent="$1"
     local value=""
@@ -1431,7 +1301,6 @@ write_rocm_build_arg_override() {
                 "${CI_BASE_DOCKERFILE_STAGES:-${DEFAULT_CI_BASE_DOCKERFILE_STAGES}}" \
                 "${CI_BASE_CONTENT_ARGS:-}"
             get_content_arg_names "${dockerfile_rocm}" "base csrc-build" "${ROCM_CSRC_CONTENT_ARGS:-}"
-            get_content_arg_names "${dockerfile_rocm}" "base rust_toolchain_input_0 rust_toolchain_input_1 rust-toolchain-input rust_input_0 rust_input_1 rust-input rust-toolchain rust-build" "${ROCM_RUST_CONTENT_ARGS:-}"
         } | awk 'NF && !seen[$0]++'
     )
 
@@ -1480,133 +1349,46 @@ validate_cache_export_mode() {
     esac
 }
 
-validate_content_cache_export_mode() {
-    local mode="$1"
-    local env_name="$2"
-
-    case "${mode}" in
-        missing|always|never)
-            ;;
-        *)
-            echo "Error: ${env_name} must be one of: missing, always, never"
-            exit 1
-            ;;
-    esac
-}
-
-should_export_content_cache_ref() {
-    local cache_ref="$1"
-    local cache_name="$2"
-    local mode="${ROCM_CONTENT_CACHE_EXPORT_MODE:-missing}"
-
-    case "${mode}" in
-        always)
-            echo "${cache_name} content cache export mode is always; exporting ${cache_ref}"
-            return 0
-            ;;
-        never)
-            echo "${cache_name} content cache export mode is never; not exporting ${cache_ref}"
-            return 1
-            ;;
-        missing|"")
-            if docker buildx imagetools inspect "${cache_ref}" >/dev/null 2>&1; then
-                echo "${cache_name} content cache exists; not re-exporting ${cache_ref}"
-                return 1
-            fi
-            echo "${cache_name} content cache missing; will export ${cache_ref}"
-            return 0
-            ;;
-        *)
-            echo "Error: ROCM_CONTENT_CACHE_EXPORT_MODE must be one of: missing, always, never"
-            exit 1
-            ;;
-    esac
-}
-
 write_rocm_cache_override() {
     local cache_repo="${DOCKERHUB_CACHE_REPO:-rocm/vllm-ci-cache}"
-    local content_cache_export_mode="${ROCM_CONTENT_CACHE_EXPORT_MODE:-missing}"
     local csrc_cache_to_mode="${ROCM_CSRC_CACHE_TO_MODE:-max}"
-    local rust_cache_to_mode="${ROCM_RUST_CACHE_TO_MODE:-max}"
     local rocm_cache_to_mode="${ROCM_FINAL_CACHE_TO_MODE:-min}"
-    local -a csrc_content_cache_from=()
+    local -a content_cache_from=()
-    local -a rust_content_cache_from=()
-    local -a combined_content_cache_from=()
     local -a csrc_cache_to=()
-    local -a rust_cache_to=()
     local -a rocm_cache_to=()
     local -a export_wheel_cache_to=()
-    local export_csrc_cache=1
-    local export_rust_cache=1
 
-    if ! uses_rocm_csrc_cache && ! uses_rocm_rust_cache; then
+    if ! uses_rocm_csrc_cache; then
         return 0
     fi
 
-    validate_content_cache_export_mode \
-        "${content_cache_export_mode}" \
-        "ROCM_CONTENT_CACHE_EXPORT_MODE"
     validate_cache_export_mode "${csrc_cache_to_mode}" "ROCM_CSRC_CACHE_TO_MODE"
-    validate_cache_export_mode "${rust_cache_to_mode}" "ROCM_RUST_CACHE_TO_MODE"
     validate_cache_export_mode "${rocm_cache_to_mode}" "ROCM_FINAL_CACHE_TO_MODE"
-    echo "ROCm content cache export mode: ${content_cache_export_mode}"
     echo "ROCm csrc cache export mode: ${csrc_cache_to_mode}"
-    echo "ROCm Rust cache export mode: ${rust_cache_to_mode}"
     echo "ROCm final image cache export mode: ${rocm_cache_to_mode}"
 
     if [[ -n "${ROCM_CSRC_CONTENT_CACHE_REF:-}" ]]; then
-        csrc_content_cache_from+=("type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF}")
+        content_cache_from+=("type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF}")
-        if should_export_content_cache_ref "${ROCM_CSRC_CONTENT_CACHE_REF}" "ROCm csrc"; then
+        csrc_cache_to+=(
-            csrc_cache_to+=(
+            "type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF},mode=${csrc_cache_to_mode},ignore-error=true"
-                "type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF},mode=${csrc_cache_to_mode},ignore-error=true"
+        )
-            )
-        else
-            export_csrc_cache=0
-        fi
     fi
 
-    if [[ -n "${ROCM_RUST_CONTENT_CACHE_REF:-}" ]]; then
-        rust_content_cache_from+=("type=registry,ref=${ROCM_RUST_CONTENT_CACHE_REF}")
-        if should_export_content_cache_ref "${ROCM_RUST_CONTENT_CACHE_REF}" "ROCm Rust"; then
-            rust_cache_to+=(
-                "type=registry,ref=${ROCM_RUST_CONTENT_CACHE_REF},mode=${rust_cache_to_mode},ignore-error=true"
-            )
-        else
-            export_rust_cache=0
-        fi
-    fi
-
-    combined_content_cache_from=("${csrc_content_cache_from[@]}" "${rust_content_cache_from[@]}")
-
     # Docker Hub cache exports are best-effort. A cache-only target failure can
     # otherwise cancel the sibling image target before its manifest is pushed.
     if [[ -n "${BUILDKITE_COMMIT:-}" ]]; then
-        if [[ ${export_csrc_cache} -eq 1 ]]; then
+        csrc_cache_to+=(
-            csrc_cache_to+=(
+            "type=registry,ref=${cache_repo}:csrc-rocm-${BUILDKITE_COMMIT},mode=${csrc_cache_to_mode},ignore-error=true"
-                "type=registry,ref=${cache_repo}:csrc-rocm-${BUILDKITE_COMMIT},mode=${csrc_cache_to_mode},ignore-error=true"
+        )
-            )
-        fi
-        if [[ ${export_rust_cache} -eq 1 ]]; then
-            rust_cache_to+=(
-                "type=registry,ref=${cache_repo}:rust-rocm-${BUILDKITE_COMMIT},mode=${rust_cache_to_mode},ignore-error=true"
-            )
-        fi
         rocm_cache_to+=(
             "type=registry,ref=${cache_repo}:rocm-${BUILDKITE_COMMIT},mode=${rocm_cache_to_mode},ignore-error=true"
         )
     fi
 
     if [[ -n "${ROCM_CACHE_BRANCH_TAG:-}" ]]; then
-        if [[ ${export_csrc_cache} -eq 1 ]]; then
+        csrc_cache_to+=(
-            csrc_cache_to+=(
+            "type=registry,ref=${cache_repo}:csrc-rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${csrc_cache_to_mode},ignore-error=true"
-                "type=registry,ref=${cache_repo}:csrc-rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${csrc_cache_to_mode},ignore-error=true"
+        )
-            )
-        fi
-        if [[ ${export_rust_cache} -eq 1 ]]; then
-            rust_cache_to+=(
-                "type=registry,ref=${cache_repo}:rust-rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${rust_cache_to_mode},ignore-error=true"
-            )
-        fi
         rocm_cache_to+=(
             "type=registry,ref=${cache_repo}:rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${rocm_cache_to_mode},ignore-error=true"
         )
@@ -1624,7 +1406,7 @@ target "csrc-rocm-ci" {
   cache-from = concat(
     get_cache_from_rocm_csrc(),
 EOF
-        write_hcl_string_list "    " "${csrc_content_cache_from[@]}"
+        write_hcl_string_list "    " "${content_cache_from[@]}"
         cat <<EOF
   )
 EOF
@@ -1632,23 +1414,11 @@ EOF
         cat <<EOF
 }
 
-target "rust-rocm-ci" {
-  cache-from = concat(
-    get_cache_from_rocm_rust(),
-EOF
-        write_hcl_string_list "    " "${rust_content_cache_from[@]}"
-        cat <<EOF
-  )
-EOF
-        write_hcl_string_list_attr "  " "cache-to" "${rust_cache_to[@]}"
-        cat <<EOF
-}
-
 target "test-rocm-ci" {
   cache-from = concat(
     get_cache_from_rocm(),
 EOF
-        write_hcl_string_list "    " "${combined_content_cache_from[@]}"
+        write_hcl_string_list "    " "${content_cache_from[@]}"
         cat <<EOF
   )
 EOF
@@ -1660,7 +1430,7 @@ target "export-wheel-rocm" {
   cache-from = concat(
     get_cache_from_rocm(),
 EOF
-        write_hcl_string_list "    " "${combined_content_cache_from[@]}"
+        write_hcl_string_list "    " "${content_cache_from[@]}"
         cat <<EOF
   )
 EOF
@@ -1686,7 +1456,7 @@ extract_dependency_pins() {
         return 0
     fi
 
-    for var in NIXL_BRANCH UCX_BRANCH ROCSHMEM_BRANCH DEEPEP_BRANCH; do
+    for var in RIXL_BRANCH UCX_BRANCH ROCSHMEM_BRANCH DEEPEP_BRANCH; do
         if [[ -n "${!var:-}" ]]; then
             echo "Using provided ${var}: ${!var}"
             continue
@@ -1706,30 +1476,30 @@ extract_dependency_pins() {
 compute_dependency_cache_keys() {
     local bake_dir=""
     local dockerfile_rocm=""
-    local nixl_branch=""
+    local rixl_branch=""
     local ucx_branch=""
     local rocshmem_branch=""
     local deepep_branch=""
-    local nixl_material=""
+    local rixl_material=""
     local rocshmem_material=""
     local deepep_material=""
 
     bake_dir=$(dirname "${VLLM_BAKE_FILE}")
     dockerfile_rocm="${bake_dir}/Dockerfile.rocm"
-    nixl_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "NIXL_BRANCH")
+    rixl_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "RIXL_BRANCH")
     ucx_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "UCX_BRANCH")
     rocshmem_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "ROCSHMEM_BRANCH")
     deepep_branch=$(resolve_dockerfile_arg_value "${dockerfile_rocm}" "DEEPEP_BRANCH")
 
-    if [[ -n "${nixl_branch}" && -n "${ucx_branch}" ]]; then
+    if [[ -n "${rixl_branch}" && -n "${ucx_branch}" ]]; then
-        nixl_material=$(compose_stage_cache_material "${dockerfile_rocm}" "base build_nixl")
+        rixl_material=$(compose_stage_cache_material "${dockerfile_rocm}" "base build_rixl")
-        NIXL_CACHE_KEY=$(
+        RIXL_CACHE_KEY=$(
             compose_dependency_cache_key \
-                "${nixl_branch}-ucx-${ucx_branch}" \
+                "${rixl_branch}-ucx-${ucx_branch}" \
-                "${nixl_material}"
+                "${rixl_material}"
         )
-        export NIXL_CACHE_KEY
+        export RIXL_CACHE_KEY
-        echo "NIXL dependency cache key: ${NIXL_CACHE_KEY}"
+        echo "RIXL dependency cache key: ${RIXL_CACHE_KEY}"
     fi
 
     if [[ -n "${rocshmem_branch}" ]]; then
@@ -1780,11 +1550,11 @@ dependency_cache_ref_for_target() {
     local cache_repo="${DOCKERHUB_CACHE_REPO:-rocm/vllm-ci-cache}"
 
     case "${target}" in
-        nixl-rocm-ci)
+        rixl-rocm-ci)
-            if [[ -n "${NIXL_CACHE_KEY:-}" ]]; then
+            if [[ -n "${RIXL_CACHE_KEY:-}" ]]; then
-                printf '%s\n' "${cache_repo}:nixl-rocm-${NIXL_CACHE_KEY}"
+                printf '%s\n' "${cache_repo}:rixl-rocm-${RIXL_CACHE_KEY}"
-            elif [[ -n "${NIXL_BRANCH:-}" ]]; then
+            elif [[ -n "${RIXL_BRANCH:-}" ]]; then
-                printf '%s\n' "${cache_repo}:nixl-rocm-${NIXL_BRANCH}-ucx-${UCX_BRANCH:-}"
+                printf '%s\n' "${cache_repo}:rixl-rocm-${RIXL_BRANCH}-ucx-${UCX_BRANCH:-}"
             fi
             ;;
         rocshmem-rocm-ci)
@@ -1815,7 +1585,7 @@ add_dependency_cache_target() {
 
 resolve_ci_base_dependency_targets() {
     local mode="${ROCM_DEP_CACHE_EXPORT_MODE:-missing}"
-    local nixl_ref=""
+    local rixl_ref=""
     local rocshmem_ref=""
     local deepep_ref=""
 
@@ -1824,7 +1594,7 @@ resolve_ci_base_dependency_targets() {
     case "${mode}" in
         always)
             echo "ROCM_DEP_CACHE_EXPORT_MODE=always; exporting all dependency caches serially"
-            for target in nixl-rocm-ci rocshmem-rocm-ci deepep-rocm-ci; do
+            for target in rixl-rocm-ci rocshmem-rocm-ci deepep-rocm-ci; do
                 if [[ -n "$(dependency_cache_ref_for_target "${target}")" ]]; then
                     add_dependency_cache_target "${target}"
                 fi
@@ -1844,13 +1614,13 @@ resolve_ci_base_dependency_targets() {
             ;;
     esac
 
-    if [[ "${mode}" != "always" && -n "${NIXL_CACHE_KEY:-}" ]]; then
+    if [[ "${mode}" != "always" && -n "${RIXL_CACHE_KEY:-}" ]]; then
-        nixl_ref=$(dependency_cache_ref_for_target "nixl-rocm-ci")
+        rixl_ref=$(dependency_cache_ref_for_target "rixl-rocm-ci")
-        if dependency_cache_ref_exists "${nixl_ref}"; then
+        if dependency_cache_ref_exists "${rixl_ref}"; then
-            echo "NIXL dependency cache exists: ${nixl_ref}"
+            echo "RIXL dependency cache exists: ${rixl_ref}"
         else
-            echo "NIXL dependency cache missing; will seed: ${nixl_ref}"
+            echo "RIXL dependency cache missing; will seed: ${rixl_ref}"
-            add_dependency_cache_target "nixl-rocm-ci"
+            add_dependency_cache_target "rixl-rocm-ci"
         fi
     fi
 
@@ -1950,8 +1720,8 @@ confirm_remote_image_push() {
         fi
 
         if [[ -z "${remote_revision}" \
-              && ${IMAGE_EXISTED_BEFORE_BUILD} -eq 0 ]] \
+              && ${IMAGE_EXISTED_BEFORE_BUILD} -eq 0 \
-            && image_tag_is_commit_scoped; then
+              && image_tag_is_commit_scoped ]]; then
             echo "Remote image exists under a commit-scoped tag; accepting push despite missing revision label."
             return 0
         fi
@@ -2009,10 +1779,7 @@ seed_dependency_caches_if_needed() {
 
         echo "--- :docker: Seeding ${target}"
         echo "Expected cache ref: ${cache_ref}"
-        docker buildx bake \
+        docker buildx bake "${BAKE_FILES[@]}" --progress plain "${target}"
-            "${BAKE_FILES[@]}" \
-            --progress "${BUILDKIT_PROGRESS:-plain}" \
-            "${target}"
         verify_dependency_cache_ref "${cache_ref}"
     done
 }
@@ -2040,10 +1807,7 @@ run_bake() {
     local build_rc=0
 
     echo "--- :docker: Building ${TARGET}"
-    docker buildx bake \
+    docker buildx bake "${BAKE_FILES[@]}" --progress plain "${BAKE_TARGETS[@]}" || build_rc=$?
-        "${BAKE_FILES[@]}" \
-        --progress "${BUILDKIT_PROGRESS:-plain}" \
-        "${BAKE_TARGETS[@]}" || build_rc=$?
 
     if [[ ${build_rc} -eq 0 ]]; then
         echo "--- :white_check_mark: Build complete"
@@ -2081,57 +1845,36 @@ upload_wheel_artifacts_if_present() {
     local wheel_dir="./wheel-export"
     local artifact_dir="artifacts/vllm-rocm-install"
     local archive_name="vllm-rocm-install.tar.gz"
-    local metadata_dir="${wheel_dir}/.vllm-ci-artifact"
-    local native_base_image=""
     local whl=""
     local whl_name=""
-    local -a wheels=()
 
     if ! should_upload_wheel_artifacts; then
         return 0
     fi
 
-    if [[ -d "${wheel_dir}" ]]; then
+    if [[ ! -d "${wheel_dir}" ]] || ! ls "${wheel_dir}"/*.whl >/dev/null 2>&1; then
-        mapfile -t wheels < <(find "${wheel_dir}" -maxdepth 1 -type f -name '*.whl' -print)
+        echo "No ROCm wheel artifacts found in ${wheel_dir}"
-    fi
+        return 0
-    if [[ ${#wheels[@]} -ne 1 ]]; then
-        echo "Expected exactly one ROCm wheel in ${wheel_dir}; found ${#wheels[@]}" >&2
-        return 1
-    fi
-    whl="${wheels[0]}"
-    whl_name=$(basename "${whl}")
-    native_base_image="${CI_BASE_IMAGE_TAG_COMMIT_REF:-${CI_BASE_IMAGE:-}}"
-    if [[ -z "${native_base_image}" ]]; then
-        echo "Native ROCm artifact requires a ci_base image reference" >&2
-        return 1
     fi
 
     echo "--- :package: Uploading ROCm vLLM install artifact"
-    rm -rf "${artifact_dir}" "${metadata_dir}"
+    mkdir -p "${artifact_dir}"
-    mkdir -p "${artifact_dir}" "${metadata_dir}"
-
-    printf '%s\n' "${BUILDKITE_COMMIT:-local}" > "${metadata_dir}/commit.txt"
-    printf '%s\n' "${native_base_image}" > "${metadata_dir}/native-base-image.txt"
-    printf '%s\n' "${CI_BASE_IMAGE:-}" > "${metadata_dir}/ci-base-image.txt"
-    printf '%s\n' "${IMAGE_TAG:-}" > "${metadata_dir}/fallback-image.txt"
-    printf '%s\n' "${whl_name}" > "${metadata_dir}/wheel-filename.txt"
 
     tar -C "${wheel_dir}" -czf "${artifact_dir}/${archive_name}" .
-    (
-        cd "${artifact_dir}"
-        sha256sum "${archive_name}" > "${archive_name}.sha256"
-    )
     echo "Created ${archive_name}: $(du -sh "${artifact_dir}/${archive_name}" | cut -f1)"
-    cp "${metadata_dir}"/*.txt "${artifact_dir}/"
+    printf '%s\n' "${CI_BASE_IMAGE:-}" > "${artifact_dir}/ci-base-image.txt"
-    cp "${whl}" "${artifact_dir}/${whl_name}"
+    printf '%s\n' "${IMAGE_TAG:-}" > "${artifact_dir}/fallback-image.txt"
-    echo "Copied ${whl_name}: $(du -sh "${artifact_dir}/${whl_name}" | cut -f1)"
+
+    for whl in "${wheel_dir}"/*.whl; do
+        [[ -f "${whl}" ]] || continue
+        whl_name=$(basename "${whl}")
+        cp "${whl}" "${artifact_dir}/${whl_name}"
+        echo "Copied ${whl_name}: $(du -sh "${artifact_dir}/${whl_name}" | cut -f1)"
+    done
 
     if command -v buildkite-agent >/dev/null 2>&1; then
-        buildkite-agent artifact upload "${artifact_dir}/*" || return 1
+        buildkite-agent artifact upload "${artifact_dir}/*"
         echo "ROCm vLLM install artifacts uploaded to ${artifact_dir}/"
-    elif [[ "${BUILDKITE:-false}" == "true" ]]; then
-        echo "buildkite-agent not found; cannot upload required ROCm artifacts" >&2
-        return 1
     else
         echo "Not in Buildkite, skipping artifact upload"
     fi
@@ -2155,7 +1898,6 @@ main() {
     compute_dependency_cache_keys
     write_ci_base_label_override
     compute_rocm_csrc_content_hash_if_needed
-    compute_rocm_rust_content_hash_if_needed
     write_rocm_cache_override
     resolve_ci_base_dependency_targets
     print_bake_config
@@ -2163,11 +1905,6 @@ main() {
         echo "BAKE_PRINT_ONLY=1 set; skipping build"
         return 0
     fi
-    if should_upload_wheel_artifacts; then
-        # wheel-export is an output directory, not a BuildKit cache. Starting
-        # clean prevents a failed/retried export from packaging a stale wheel.
-        rm -rf ./wheel-export
-    fi
     seed_dependency_caches_if_needed
     run_bake
     upload_wheel_artifacts_if_present

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

@@ -45,10 +45,8 @@ $PYTHON .buildkite/scripts/generate-nightly-index.py --version "$SUBPATH" --curr
 echo "Uploading indices to $S3_COMMIT_PREFIX"
 aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "$S3_COMMIT_PREFIX"
 
-# copy to /nightly/ only when enabled for a main branch build that is not a PR
+# copy to /nightly/ only if it is on the main branch and not a PR
-if [[ "${UPDATE_NIGHTLY_INDEX:-1}" == "1" && \
+if [[ "$BUILDKITE_BRANCH" == "main" && "$BUILDKITE_PULL_REQUEST" == "false" ]]; then
-      "$BUILDKITE_BRANCH" == "main" && \
-      "$BUILDKITE_PULL_REQUEST" == "false" ]]; then
     echo "Uploading indices to overwrite /nightly/"
     aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "s3://$BUCKET/nightly/"
 fi
@@ -69,7 +67,7 @@ pure_version="${version%%+*}"
 echo "Pure version (without variant): $pure_version"
 
 # re-generate and copy to /<pure_version>/ only if it does not have "dev" in the version
-if [[ "${UPDATE_VERSION_INDEX:-1}" == "1" && "$version" != *"dev"* ]]; then
+if [[ "$version" != *"dev"* ]]; then
     echo "Re-generating indices for /$pure_version/"
     rm -rf "${INDICES_OUTPUT_DIR:?}"
     mkdir -p "$INDICES_OUTPUT_DIR"

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

@@ -1,7 +1,7 @@
 #!/bin/bash
 
-# This script runs ROCm tests either directly in a native CI pod or inside the
+# This script runs tests inside the corresponding ROCm docker container.
-# corresponding Docker container. Multi-node tests continue to use Docker.
+# It handles both single-node and multi-node test configurations.
 #
 # Multi-node detection: Instead of matching on fragile group names, we detect
 # multi-node jobs structurally by looking for the bracket command syntax
@@ -28,34 +28,6 @@
 ###############################################################################
 set -o pipefail
 
-: "${BUILDKIT_PROGRESS:=plain}"
-: "${TERM:=xterm-256color}"
-: "${FORCE_COLOR:=1}"
-: "${CLICOLOR_FORCE:=1}"
-: "${PY_COLORS:=1}"
-: "${ROCM_DOCKER_TTY:=1}"
-: "${PYTHONFAULTHANDLER:=1}"
-: "${PYTEST_TIMEOUT:=2400}"
-if [[ " ${PYTEST_ADDOPTS:-} " != *" --color"* ]]; then
-  PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }--color=yes"
-fi
-if [[ " ${PYTEST_ADDOPTS:-} " != *" --durations="* ]]; then
-  PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }--durations=25"
-fi
-if [[ " ${PYTEST_ADDOPTS:-} " != *" --durations-min="* ]]; then
-  PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }--durations-min=1.0"
-fi
-# Dump stacks after 25 minutes, then stop an individual test after 40 minutes.
-if [[ " ${PYTEST_ADDOPTS:-} " != *" faulthandler_timeout="* ]]; then
-  PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }-o faulthandler_timeout=1500"
-fi
-if [[ " ${PYTEST_ADDOPTS:-} " != *" --timeout-method="* &&
-  " ${PYTEST_ADDOPTS:-} " != *" --timeout-method "* ]]; then
-  PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }--timeout-method=thread"
-fi
-export BUILDKIT_PROGRESS TERM FORCE_COLOR CLICOLOR_FORCE PY_COLORS PYTEST_ADDOPTS PYTEST_TIMEOUT ROCM_DOCKER_TTY
-export PYTHONFAULTHANDLER
-
 # 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.
@@ -70,28 +42,6 @@ report_docker_usage() {
   docker system df || true
 }
 
-clear_ci_orchestration_env() {
-  unset -v \
-    VLLM_TEST_GROUP_NAME \
-    VLLM_CI_REQUIRE_PERSISTENT_HF_CACHE \
-    VLLM_CI_ARTIFACT_STEP \
-    VLLM_TEST_CACHE \
-    VLLM_CI_EXECUTION_MODE \
-    VLLM_CI_WORKSPACE \
-    VLLM_CI_REQUIRE_WORKSPACE_MOUNT \
-    VLLM_TEST_COMMANDS \
-    VLLM_CI_BRANCH \
-    VLLM_CI_BASE_IMAGE \
-    VLLM_CI_FALLBACK_IMAGE \
-    VLLM_CI_DOCKER_DISABLED \
-    VLLM_CI_ARTIFACT_GLOB \
-    VLLM_CI_ARTIFACT_CHECKSUM_GLOB \
-    VLLM_CI_EXPECTED_GPU_COUNT \
-    VLLM_CI_USE_ARTIFACTS \
-    VLLM_CI_RESULTS_ROOT \
-    VLLM_ALLOW_DEPRECATED_BEAM_SEARCH
-}
-
 cleanup_network() {
   local max_nodes=${NUM_NODES:-2}
   for node in $(seq 0 $((max_nodes - 1))); do
@@ -184,11 +134,7 @@ prepare_artifact_image() {
   fi
 
   cp "${wheel_dir}"/*.whl "${context_dir}/wheels/" || return 1
-  tar -C "${wheel_dir}" \
+  tar -C "${wheel_dir}" --exclude='*.whl' -cf - . \
-    --exclude='*.whl' \
-    --exclude='.vllm-ci-artifact' \
-    --exclude='./.vllm-ci-artifact' \
-    -cf - . \
     | tar -C "${workspace_dir}" -xf - || return 1
   cat > "${context_dir}/Dockerfile" <<'EOF'
 ARG BASE_IMAGE
@@ -203,7 +149,6 @@ 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
@@ -211,259 +156,6 @@ EOF
   return 0
 }
 
-is_native_runtime() {
-  [[ "${AMD_CI_RUNTIME:-}" == "native" || "${NATIVE_CI:-}" == "true" ]]
-}
-
-validate_native_workspace() {
-  local workspace_dir="${VLLM_CI_WORKSPACE:-/vllm-workspace}"
-  local workspace_real=""
-  local checkout_real=""
-  local workspace_mount=""
-
-  mkdir -p "${workspace_dir}" || return 1
-  workspace_real=$(readlink -m "${workspace_dir}") || return 1
-  if [[ -n "${BUILDKITE_BUILD_CHECKOUT_PATH:-}" ]]; then
-    checkout_real=$(readlink -m "${BUILDKITE_BUILD_CHECKOUT_PATH}") || return 1
-    if [[ "${checkout_real}" == "${workspace_real}" \
-      || "${checkout_real}" == "${workspace_real}/"* \
-      || "${workspace_real}" == "${checkout_real}/"* ]]; then
-      echo "Refusing to replace ${workspace_real}; it overlaps the Buildkite checkout ${checkout_real}" >&2
-      return 1
-    fi
-  fi
-  if [[ "${VLLM_CI_REQUIRE_WORKSPACE_MOUNT:-1}" == "1" ]]; then
-    if ! command -v findmnt >/dev/null 2>&1; then
-      echo "findmnt is required to verify the native workspace mount" >&2
-      return 1
-    fi
-    workspace_mount=$(findmnt -n -T "${workspace_real}" -o TARGET 2>/dev/null || true)
-    if [[ "$(readlink -m "${workspace_mount:-/}")" != "${workspace_real}" ]]; then
-      echo "Native CI requires a dedicated volume mounted at ${workspace_real}" >&2
-      return 1
-    fi
-  fi
-}
-
-prepare_native_workspace() {
-  if [[ "${VLLM_CI_USE_ARTIFACTS:-0}" != "1" ]]; then
-    echo "Native CI requires VLLM_CI_USE_ARTIFACTS=1"
-    return 1
-  fi
-  if ! command -v buildkite-agent >/dev/null 2>&1; then
-    echo "buildkite-agent not found; cannot download ROCm wheel artifact"
-    return 1
-  fi
-  validate_native_workspace || return 1
-
-  local artifact_glob="${VLLM_CI_ARTIFACT_GLOB:-artifacts/vllm-rocm-install/vllm-rocm-install.tar.gz}"
-  local artifact_checksum_glob="${VLLM_CI_ARTIFACT_CHECKSUM_GLOB:-${artifact_glob}.sha256}"
-  local artifact_step="${VLLM_CI_ARTIFACT_STEP:-image-build-amd}"
-  local archive=""
-  local checksum=""
-  local download_dir=""
-  local metadata_dir=""
-  local recorded_base=""
-  local recorded_commit=""
-  local recorded_wheel=""
-  local workspace_dir="${VLLM_CI_WORKSPACE:-/vllm-workspace}"
-  local wheel_dir=""
-  local attempt=0
-  local attempt_dir=""
-  local -a archives=()
-  local -a checksums=()
-  local -a wheels=()
-
-  artifact_work_dir=$(mktemp -d -t vllm-rocm-artifact.XXXXXX) || return 1
-  wheel_dir="${artifact_work_dir}/wheels"
-  mkdir -p "${wheel_dir}" || return 1
-
-  echo "--- Downloading ROCm wheel artifact from ${artifact_step} (native in-pod)"
-  for attempt in 1 2 3; do
-    attempt_dir="${artifact_work_dir}/download-${attempt}"
-    rm -rf "${attempt_dir}" || return 1
-    mkdir -p "${attempt_dir}" || return 1
-    if buildkite-agent artifact download \
-      "${artifact_glob}" "${attempt_dir}" --step "${artifact_step}" \
-      && buildkite-agent artifact download \
-        "${artifact_checksum_glob}" "${attempt_dir}" --step "${artifact_step}"; then
-      download_dir="${attempt_dir}"
-      break
-    fi
-    echo "Artifact download attempt ${attempt}/3 failed"
-    if [[ "${attempt}" -lt 3 ]]; then
-      sleep $((attempt * 2))
-    fi
-  done
-  if [[ -z "${download_dir}" ]]; then
-    echo "Failed to download ${artifact_glob} and ${artifact_checksum_glob} from ${artifact_step}"
-    return 1
-  fi
-
-  mapfile -t archives < <(
-    find "${download_dir}" -name "vllm-rocm-install.tar.gz" -type f -print
-  )
-  mapfile -t checksums < <(
-    find "${download_dir}" -name "vllm-rocm-install.tar.gz.sha256" -type f -print
-  )
-  if [[ ${#archives[@]} -ne 1 || ${#checksums[@]} -ne 1 ]]; then
-    echo "Expected exactly one ROCm archive and checksum; found ${#archives[@]} archive(s) and ${#checksums[@]} checksum(s)" >&2
-    return 1
-  fi
-  archive="${archives[0]}"
-  checksum="${checksums[0]}"
-  if [[ "$(dirname "${archive}")" != "$(dirname "${checksum}")" ]]; then
-    echo "ROCm archive and checksum were downloaded to different directories" >&2
-    return 1
-  fi
-  (
-    cd "$(dirname "${archive}")"
-    sha256sum -c "$(basename "${checksum}")"
-  ) || return 1
-
-  tar --no-same-owner -xzf "${archive}" -C "${wheel_dir}" || return 1
-  mapfile -t wheels < <(
-    find "${wheel_dir}" -maxdepth 1 -type f -name '*.whl' -print
-  )
-  if [[ ${#wheels[@]} -ne 1 ]]; then
-    echo "ROCm artifact must contain exactly one top-level wheel; found ${#wheels[@]}" >&2
-    return 1
-  fi
-  metadata_dir="${wheel_dir}/.vllm-ci-artifact"
-  for metadata_file in commit.txt native-base-image.txt wheel-filename.txt; do
-    if [[ ! -s "${metadata_dir}/${metadata_file}" ]]; then
-      echo "ROCm artifact metadata is missing ${metadata_file}" >&2
-      return 1
-    fi
-  done
-  for metadata_file in ci-base-image.txt fallback-image.txt; do
-    if [[ ! -f "${metadata_dir}/${metadata_file}" ]]; then
-      echo "ROCm artifact metadata is missing ${metadata_file}" >&2
-      return 1
-    fi
-  done
-
-  recorded_commit=$(tr -d '\r\n' < "${metadata_dir}/commit.txt")
-  recorded_base=$(tr -d '\r\n' < "${metadata_dir}/native-base-image.txt")
-  recorded_wheel=$(tr -d '\r\n' < "${metadata_dir}/wheel-filename.txt")
-  if [[ -z "${BUILDKITE_COMMIT:-}" || "${recorded_commit}" != "${BUILDKITE_COMMIT}" ]]; then
-    echo "ROCm artifact commit ${recorded_commit} does not match ${BUILDKITE_COMMIT:-unset}" >&2
-    return 1
-  fi
-  if [[ -z "${VLLM_CI_BASE_IMAGE:-}" || "${recorded_base}" != "${VLLM_CI_BASE_IMAGE}" ]]; then
-    echo "ROCm artifact base ${recorded_base} does not match ${VLLM_CI_BASE_IMAGE:-unset}" >&2
-    return 1
-  fi
-  if [[ "${recorded_wheel}" != "$(basename "${wheels[0]}")" ]]; then
-    echo "ROCm artifact wheel manifest ${recorded_wheel} does not match $(basename "${wheels[0]}")" >&2
-    return 1
-  fi
-  for required_dir in tests .buildkite requirements; do
-    if [[ ! -d "${wheel_dir}/${required_dir}" ]]; then
-      echo "ROCm wheel artifact did not contain ${required_dir}/" >&2
-      return 1
-    fi
-  done
-
-  echo "--- Installing ROCm wheel into pod environment"
-  python3 -m pip install --no-deps --force-reinstall "${wheels[0]}" || return 1
-
-  echo "--- Preparing ${workspace_dir} from artifact"
-  find "${workspace_dir}" -mindepth 1 -maxdepth 1 -exec rm -rf -- {} + || return 1
-  tar -C "${wheel_dir}" \
-    --exclude='*.whl' \
-    --exclude='.vllm-ci-artifact' \
-    --exclude='./.vllm-ci-artifact' \
-    -cf - . | tar --no-same-owner -C "${workspace_dir}" -xf - || return 1
-  if [[ ! -d "${workspace_dir}/tests" ]]; then
-    echo "Failed to stage the native test workspace" >&2
-    return 1
-  fi
-
-  return 0
-}
-
-initialize_native_environment() {
-  local job_id="${BUILDKITE_JOB_ID:-${BUILDKITE_PARALLEL_JOB:-local}}"
-  local job_id_suffix=""
-  local native_root=""
-  local hf_mount=""
-
-  if [[ "$(id -u)" -ne 0 ]]; then
-    echo "Native ROCm CI currently requires the ci_base container to run as root" >&2
-    return 1
-  fi
-
-  job_id="${job_id//[^A-Za-z0-9_.-]/_}"
-  job_id_suffix="${job_id##*-}"
-  job_id_suffix="${job_id_suffix:0:12}"
-  native_root="/tmp/vllm-native-${job_id}"
-  TMPDIR="/tmp/vllm-${job_id_suffix}/tmp"
-  VLLM_RPC_BASE_PATH="/tmp"
-  TORCHINDUCTOR_CACHE_DIR="${native_root}/cache/torchinductor"
-  TRITON_CACHE_DIR="${native_root}/cache/triton"
-  VLLM_CACHE_ROOT="${native_root}/cache/vllm"
-  XDG_CACHE_HOME="${native_root}/cache/xdg"
-  : "${HF_HOME:=/home/buildkite-agent/huggingface}"
-  : "${HF_HUB_DOWNLOAD_TIMEOUT:=300}"
-  : "${HF_HUB_ETAG_TIMEOUT:=60}"
-  export TMPDIR VLLM_RPC_BASE_PATH
-  export TORCHINDUCTOR_CACHE_DIR TRITON_CACHE_DIR VLLM_CACHE_ROOT XDG_CACHE_HOME
-  export HF_HOME HF_HUB_DOWNLOAD_TIMEOUT HF_HUB_ETAG_TIMEOUT
-  export PYTORCH_ROCM_ARCH=""
-
-  mkdir -p "${TMPDIR}" \
-    "${TORCHINDUCTOR_CACHE_DIR}" \
-    "${TRITON_CACHE_DIR}" \
-    "${VLLM_CACHE_ROOT}" \
-    "${XDG_CACHE_HOME}" \
-    "${HF_HOME}" || return 1
-
-  echo "Native compile caches: VLLM_CACHE_ROOT=${VLLM_CACHE_ROOT} TORCHINDUCTOR_CACHE_DIR=${TORCHINDUCTOR_CACHE_DIR}"
-
-  if [[ "${VLLM_CI_REQUIRE_PERSISTENT_HF_CACHE:-0}" == "1" ]]; then
-    if ! command -v findmnt >/dev/null 2>&1; then
-      echo "findmnt is required to verify the native Hugging Face cache mount" >&2
-      return 1
-    fi
-    hf_mount=$(findmnt -n -T "${HF_HOME}" -o TARGET 2>/dev/null || true)
-    if [[ -z "${hf_mount}" || "${hf_mount}" == "/" ]]; then
-      echo "Native CI requires a persistent volume mounted at or above ${HF_HOME}" >&2
-      return 1
-    fi
-  fi
-}
-
-run_native_preflight() {
-  local expected_gpus="${VLLM_CI_EXPECTED_GPU_COUNT:-1}"
-
-  if [[ ! "${expected_gpus}" =~ ^[0-9]+$ ]]; then
-    echo "Invalid VLLM_CI_EXPECTED_GPU_COUNT=${expected_gpus}" >&2
-    return 1
-  fi
-
-  python3 -c "import encodings, importlib.metadata as im, importlib.util as iu; [im.version(d) for d in ('transformers', 'torch', 'ray', 'sympy', 'markupsafe', 'vllm')]; missing=[m for m in ('torch.utils.model_zoo', 'transformers.models.nomic_bert', 'ray.dag', 'sympy.physics', 'markupsafe._speedups') if iu.find_spec(m) is None]; assert not missing, missing" || return 1
-
-  if [[ "${expected_gpus}" == "0" ]]; then
-    echo "Native CPU-only AMD job: skipping ROCm device validation"
-    return 0
-  fi
-
-  echo "--- ROCm info"
-  rocminfo || return 1
-  VLLM_CI_EXPECTED_GPU_COUNT="${expected_gpus}" python3 - <<'PY'
-import os
-
-import torch
-
-expected = int(os.environ["VLLM_CI_EXPECTED_GPU_COUNT"])
-assert torch.version.hip, "PyTorch is not a ROCm build"
-assert torch.cuda.is_available(), "ROCm GPU is not available to PyTorch"
-actual = torch.cuda.device_count()
-assert actual == expected, f"Expected {expected} ROCm GPU(s), found {actual}"
-PY
-}
-
 is_multi_node() {
   local cmds="$1"
   # Primary signal: NUM_NODES environment variable set by the pipeline
@@ -646,58 +338,7 @@ re_quote_pytest_markers() {
 # Main
 ###############################################################################
 
-if is_native_runtime; then
+# --- GPU initialization ---
-  echo "--- Native in-pod ROCm CI (AMD_CI_RUNTIME=${AMD_CI_RUNTIME:-unset}, NATIVE_CI=${NATIVE_CI:-unset})"
-  artifact_work_dir=""
-
-  cleanup_native_workspace() {
-    if [[ -n "${artifact_work_dir}" ]]; then
-      rm -rf "${artifact_work_dir}"
-    fi
-  }
-  trap cleanup_native_workspace EXIT
-
-  if [[ -n "${VLLM_TEST_COMMANDS:-}" ]]; then
-    commands="${VLLM_TEST_COMMANDS}"
-    commands_source="env"
-  else
-    commands="$*"
-    commands_source="argv"
-    if [[ -z "$commands" ]]; then
-      echo "Error: No test commands provided for native CI." >&2
-      exit 1
-    fi
-  fi
-
-  if [[ "$commands_source" == "argv" ]]; then
-    commands=$(re_quote_pytest_markers "$commands")
-  fi
-
-  if is_multi_node "$commands"; then
-    echo "Native CI does not support multi-node jobs yet."
-    exit 1
-  fi
-
-  if ! initialize_native_environment; then
-    echo "Failed to initialize the native test environment"
-    exit 1
-  fi
-  if ! prepare_native_workspace; then
-    echo "Failed to prepare native test workspace"
-    exit 1
-  fi
-
-  export PYTHONPATH="${VLLM_CI_WORKSPACE:-/vllm-workspace}"
-
-  echo "Native test commands: $commands"
-  run_native_preflight || exit 1
-  # Keep AMD CI orchestration variables out of vLLM's runtime environment.
-  clear_ci_orchestration_env
-  /bin/bash -o pipefail -c "${commands}"
-  handle_pytest_exit "$?"
-fi
-
-# --- GPU initialization for legacy Docker execution ---
 echo "--- ROCm info"
 rocminfo
 
@@ -799,27 +440,23 @@ fi
 
 echo "Final commands: $commands"
 
-standalone_merge_base_env=()
+# The ROCm test image often ships /vllm-workspace without .git (artifact tarball unpack).
-if [[ "$commands" == *python_only_compile.sh* ]]; then
+# tests/standalone_tests/python_only_compile.sh uses merge-base(HEAD, origin/main) for
-  # The ROCm test image often ships /vllm-workspace without .git. Resolve the
+# wheels.vllm.ai; compute on the agent (full git checkout) and pass into the container.
-  # wheels.vllm.ai commit from the agent checkout for this test only.
+vllm_standalone_merge_base=""
-  vllm_standalone_merge_base=""
+checkout="${BUILDKITE_BUILD_CHECKOUT_PATH:-}"
-  checkout="${BUILDKITE_BUILD_CHECKOUT_PATH:-}"
+if [[ -z "${checkout}" || ! -d "${checkout}" ]]; then
-  if [[ -z "${checkout}" || ! -d "${checkout}" ]]; then
+  checkout="."
-    checkout="."
-  fi
-  # Pass safe.directory per-command because Buildkite uses mixed user IDs.
-  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 "safe.directory=${checkout}" -C "${checkout}" merge-base HEAD origin/main 2>/dev/null || true
-    )"
-  fi
-  if [[ -z "${vllm_standalone_merge_base}" ]]; then
-    vllm_standalone_merge_base="${BUILDKITE_COMMIT:-}"
-  fi
-  echo "INFO: passing CI_STANDALONE_MERGE_BASE into container: ${vllm_standalone_merge_base}"
-  standalone_merge_base_env=(-e "CI_STANDALONE_MERGE_BASE=${vllm_standalone_merge_base}")
 fi
+if git -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
+  )"
+fi
+if [[ -z "${vllm_standalone_merge_base}" ]]; then
+  vllm_standalone_merge_base="${BUILDKITE_COMMIT:-}"
+fi
+echo "INFO: passing VLLM_STANDALONE_MERGE_BASE into container: ${vllm_standalone_merge_base}"
 
 MYPYTHONPATH="/vllm-workspace"
 
@@ -850,7 +487,6 @@ else
 fi
 
 # --- Route: multi-node vs single-node ---
-clear_ci_orchestration_env
 if is_multi_node "$commands"; then
   echo "--- Multi-node job detected"
   export DCKR_VER=$(docker --version | sed 's/Docker version \(.*\), build .*/\1/')
@@ -897,37 +533,14 @@ 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
-    # docker run can't pass "unlimited" to --ulimit
-    ulimit_core_hard="-1"
-  fi
-   # Disable core dumps in the ROCm test container unless the ROCm debug agent is enabled
-  coredump_flags="--ulimit core=0:$ulimit_core_hard"
-  if [[ "$commands" == *"ROCm debug agent enabled"* ]]; then
-    # Works around https://github.com/rocm/rocm-systems/issues/6206
-    coredump_flags='-e HSA_COREDUMP_PATTERN="/tmp/gpucore.%p"'
-  else
-    echo "ROCm debug agent not enabled, coredumps are disabled in the test container."
-  fi
 
   docker run \
-    "${docker_run_terminal_args[@]}" \
     --device /dev/kfd $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES \
     $RDMA_FLAGS \
     --network=host \
     --shm-size=16gb \
     --group-add "$render_gid" \
     --rm \
-    $coredump_flags \
     -e HF_TOKEN \
     -e "HF_HUB_DOWNLOAD_TIMEOUT=${HF_HUB_DOWNLOAD_TIMEOUT}" \
     -e "HF_HUB_ETAG_TIMEOUT=${HF_HUB_ETAG_TIMEOUT}" \
@@ -935,13 +548,6 @@ 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 PYTHONFAULTHANDLER \
-    -e PYTEST_ADDOPTS \
-    -e PYTEST_TIMEOUT \
     -v "${HF_CACHE}:${HF_MOUNT}" \
     -e "HF_HOME=${HF_MOUNT}" \
     -e "PYTHONPATH=${MYPYTHONPATH}" \
@@ -951,7 +557,7 @@ else
     -e "VLLM_CACHE_ROOT=${CONTAINER_CACHE_ROOT}/vllm" \
     -e "XDG_CACHE_HOME=${CONTAINER_CACHE_ROOT}/xdg" \
     -e "PYTORCH_ROCM_ARCH=" \
-    "${standalone_merge_base_env[@]}" \
+    -e "VLLM_STANDALONE_MERGE_BASE=${vllm_standalone_merge_base}" \
     --name "${container_name}" \
     "${image_name}" \
     /bin/bash -c "${CONTAINER_PREFLIGHT} && ${commands}"

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

@@ -1,11 +1,10 @@
 #!/bin/bash
 set -euox pipefail
 
-export VLLM_CPU_KVCACHE_SPACE=1
+export VLLM_CPU_KVCACHE_SPACE=1 
 export VLLM_CPU_CI_ENV=1
-# Skip torch.compile via vLLM's --enforce-eager flag (passed below) instead of
+# Reduce sub-processes for acceleration
-# TORCH_COMPILE_DISABLE=1, which torch 2.12 no longer treats as a silent no-op
+export TORCH_COMPILE_DISABLE=1 
-# when callers specify fullgraph=True.
 export VLLM_ENABLE_V1_MULTIPROCESSING=0
 
 SDE_ARCHIVE="sde-external-10.7.0-2026-02-18-lin.tar.xz"
@@ -50,15 +49,15 @@ wait_for_pid_and_check_log() {
 }
 
 # Test Sky Lake (AVX512F)
-./sde/sde64 -skl -- python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --dtype bfloat16 --enforce-eager > test_0.log 2>&1 &
+./sde/sde64 -skl -- python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --dtype bfloat16 > test_0.log 2>&1 &
 PID_TEST_0=$!
 
 # Test Cascade Lake (AVX512F + VNNI)
-./sde/sde64 -clx -- python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --dtype bfloat16 --enforce-eager > test_1.log 2>&1 &
+./sde/sde64 -clx -- python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --dtype bfloat16 > test_1.log 2>&1 &
 PID_TEST_1=$!
 
 # Test Cooper Lake (AVX512F + VNNI + BF16)
-./sde/sde64 -cpx -- python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --dtype bfloat16 --enforce-eager > test_2.log 2>&1 &
+./sde/sde64 -cpx -- python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --dtype bfloat16 > test_2.log 2>&1 &
 PID_TEST_2=$!
 
 wait_for_pid_and_check_log $PID_TEST_0 test_0.log

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

@@ -38,11 +38,7 @@ 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
-    pytest -x -v -s tests/kernels/mamba/test_causal_conv1d.py
-    pytest -x -v -s tests/kernels/mamba/test_mamba_ssm.py"
 
   # skip tests requiring model downloads if HF_TOKEN is not set
   # due to rate-limits
@@ -66,6 +62,7 @@ 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
@@ -99,4 +96,3 @@ function cpu_tests() {
 # All of CPU tests are expected to be finished less than 40 mins.
 export -f cpu_tests
 timeout 2h bash -c cpu_tests
-

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

@@ -21,7 +21,6 @@ case "${test_suite}" in
     python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --kv-cache-dtype fp8
     python3 examples/basic/offline_inference/generate.py --model nvidia/Llama-3.1-8B-Instruct-FP8 --block-size 64 --enforce-eager --quantization modelopt --kv-cache-dtype fp8 --attention-backend TRITON_ATTN --max-model-len 4096
     python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --block-size 64 --enforce-eager --max-model-len 8192
-    python3 examples/basic/offline_inference/generate.py --model TheBloke/TinyLlama-1.1B-Chat-v0.3-AWQ --block-size 64 --enforce-eager 
     python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2
     python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 --enable-expert-parallel
     python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --max-model-len 8192
@@ -35,7 +34,6 @@ case "${test_suite}" in
     pytest -v -s v1/worker --ignore=v1/worker/test_gpu_model_runner.py --ignore=v1/worker/test_worker_memory_snapshot.py
     pytest -v -s v1/structured_output
     pytest -v -s v1/test_serial_utils.py
-    pytest -v -s v1/e2e/general/test_correctness_sliding_window.py --deselect="tests/v1/e2e/general/test_correctness_sliding_window.py::test_sliding_window_retrieval[True-1-5-google/gemma-3-1b-it]"
     pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py --ignore=v1/spec_decode/test_speculators_correctness.py
     pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --ignore=v1/kv_connector/unit/test_lmcache_integration.py --ignore=v1/kv_connector/unit/test_hf3fs_client.py --ignore=v1/kv_connector/unit/test_hf3fs_connector.py --ignore=v1/kv_connector/unit/test_hf3fs_metadata_server.py --ignore=v1/kv_connector/unit/test_offloading_connector.py
     ;;

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

@@ -360,7 +360,7 @@ export HF_TOKEN ZE_AFFINITY_MASK
     --ipc=host \
     --privileged \
     -v /dev/dri/by-path:/dev/dri/by-path \
-    -v "/data/huggingface:/root/.cache/huggingface" \
+    -v "${HOME}/.cache/huggingface:/root/.cache/huggingface" \
     --entrypoint='' \
     -e HF_TOKEN \
     -e ZE_AFFINITY_MASK \

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

@@ -85,7 +85,7 @@ RUN pip config set global.index-url http://cache-service-vllm.nginx-pypi-cache.s
 
 # Install for pytest to make the docker build cache layer always valid
 RUN --mount=type=cache,target=/root/.cache/pip \
-    pip install pytest>=6.0  'modelscope<1.38'
+    pip install pytest>=6.0  modelscope
 
 WORKDIR /workspace/vllm
 

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

@@ -130,22 +130,6 @@ 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

@@ -1,32 +0,0 @@
-#!/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

@@ -1,68 +0,0 @@
-#!/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_ci_base_if_present() {
-    local ci_base_image=""
-
-    ci_base_image="$(metadata_get rocm-ci-base-image)"
-    if [[ -z "${ci_base_image}" ]]; then
-        return 1
-    fi
-
-    export CI_BASE_IMAGE="${ci_base_image}"
-    echo "Using ROCm ci_base image selected by the preceding build step: ${CI_BASE_IMAGE}"
-}
-
-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 IMAGE_TAG_LATEST
-
-    BASE_IMAGE="$(metadata_get rocm-base-image)"
-    IMAGE_TAG_LATEST="$(metadata_get rocm-ci-image-descriptive)"
-
-    echo "Using refreshed ROCm base image for test image: ${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
-
-    use_ci_base_if_present || true
-
-    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

@@ -1,513 +0,0 @@
-#!/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

@@ -1,32 +0,0 @@
-#!/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,9 +109,7 @@ run_nodes() {
         if [ "$node" -ne 0 ]; then
             docker exec -d "node$node" /bin/bash -c "cd $WORKING_DIR ; ${COMMANDS[$node]}"
         else
-            # Allocate a TTY (-t -i) for the foreground head node so its output
+            docker exec "node$node" /bin/bash -c "cd $WORKING_DIR ; ${COMMANDS[$node]}"
-            # 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,12 +8,7 @@ if [[ "$MODE" != "style-clippy" && "$MODE" != "test" ]]; then
   exit 2
 fi
 
-if ROOT_DIR="$(git rev-parse --show-toplevel 2>/dev/null)"; then
+ROOT_DIR="$(git rev-parse --show-toplevel)"
-  :
-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}"
@@ -21,20 +16,16 @@ export CARGO_HOME="${CARGO_HOME:-$HOME/.cargo}"
 export RUSTUP_HOME="${RUSTUP_HOME:-$HOME/.rustup}"
 export PATH="$CARGO_HOME/bin:$PATH"
 
-PROTOC_VERSION="${PROTOC_VERSION:-31.1}"
-CARGO_BINSTALL_VERSION="${CARGO_BINSTALL_VERSION:-1.20.1}"
-UV_VERSION="${UV_VERSION:-0.11.28}"
-PYO3_PYTHON_VERSION="${PYO3_PYTHON_VERSION:-3.12}"
-
-CARGO_SORT_VERSION_REQ="${CARGO_SORT_VERSION_REQ:-2}"
-CARGO_DENY_VERSION_REQ="${CARGO_DENY_VERSION_REQ:-0.20}"
-CARGO_NEXTEST_VERSION_REQ="${CARGO_NEXTEST_VERSION_REQ:-0.9}"
-
 log_section() {
   echo "--- $*"
 }
 
 install_protoc() {
+  if command -v protoc >/dev/null 2>&1; then
+    return
+  fi
+
+  local version="${PROTOC_VERSION:-31.1}"
   local arch
   case "$(uname -m)" in
     x86_64)
@@ -49,17 +40,16 @@ install_protoc() {
       ;;
   esac
 
-  local url="https://github.com/protocolbuffers/protobuf/releases/download/v${PROTOC_VERSION}/protoc-${PROTOC_VERSION}-linux-${arch}.zip"
+  local url="https://github.com/protocolbuffers/protobuf/releases/download/v${version}/protoc-${version}-linux-${arch}.zip"
   local tmp_dir
   tmp_dir="$(mktemp -d)"
 
-  log_section "Installing protoc ${PROTOC_VERSION}"
+  log_section "Installing protoc ${version}"
   curl -L --proto '=https' --tlsv1.2 -sSf "$url" -o "$tmp_dir/protoc.zip"
   mkdir -p "$CARGO_HOME/bin"
   unzip -q "$tmp_dir/protoc.zip" bin/protoc 'include/*' -d "$CARGO_HOME"
   chmod +x "$CARGO_HOME/bin/protoc"
   rm -rf "$tmp_dir"
-  protoc --version
 }
 
 rust_toolchain() {
@@ -80,48 +70,66 @@ install_rust_toolchain() {
 }
 
 install_cargo_binstall() {
-  log_section "Installing cargo-binstall ${CARGO_BINSTALL_VERSION}"
+  if command -v cargo-binstall >/dev/null 2>&1; then
+    return
+  fi
+
+  log_section "Installing cargo-binstall"
   curl -L --proto '=https' --tlsv1.2 -sSf \
-    "https://raw.githubusercontent.com/cargo-bins/cargo-binstall/v${CARGO_BINSTALL_VERSION}/install-from-binstall-release.sh" \
+    https://raw.githubusercontent.com/cargo-bins/cargo-binstall/main/install-from-binstall-release.sh \
-    | env BINSTALL_VERSION="$CARGO_BINSTALL_VERSION" bash
+    | bash
-  cargo-binstall -V
 }
 
 install_cargo_sort() {
-  log_section "Installing cargo-sort ${CARGO_SORT_VERSION_REQ}"
+  if command -v cargo-sort >/dev/null 2>&1; then
-  cargo binstall --no-confirm --force "cargo-sort@${CARGO_SORT_VERSION_REQ}"
+    return
+  fi
+
+  log_section "Installing cargo-sort"
+  install_cargo_binstall
+  cargo binstall --no-confirm cargo-sort
 }
 
 install_cargo_deny() {
-  log_section "Installing cargo-deny ${CARGO_DENY_VERSION_REQ}"
+  if command -v cargo-deny >/dev/null 2>&1; then
-  cargo binstall --no-confirm --force "cargo-deny@${CARGO_DENY_VERSION_REQ}"
+    return
+  fi
+
+  log_section "Installing cargo-deny"
+  install_cargo_binstall
+  cargo binstall --no-confirm cargo-deny
 }
 
 install_cargo_nextest() {
-  log_section "Installing cargo-nextest ${CARGO_NEXTEST_VERSION_REQ}"
+  if command -v cargo-nextest >/dev/null 2>&1; then
-  cargo binstall \
+    return
-    --no-confirm \
+  fi
-    --force \
+
-    --secure \
+  log_section "Installing cargo-nextest"
-    "cargo-nextest@${CARGO_NEXTEST_VERSION_REQ}"
+  install_cargo_binstall
+  cargo binstall --no-confirm --secure cargo-nextest
 }
 
 install_uv() {
-  log_section "Installing uv ${UV_VERSION}"
+  if command -v uv >/dev/null 2>&1; then
-  curl -L --proto '=https' --tlsv1.2 -sSf \
+    return
-    "https://github.com/astral-sh/uv/releases/download/${UV_VERSION}/uv-installer.sh" \
+  fi
+
+  log_section "Installing uv"
+  curl -LsSf --proto '=https' --tlsv1.2 https://astral.sh/uv/install.sh \
     | env UV_INSTALL_DIR="$CARGO_HOME/bin" sh
-  uv --version
 }
 
 setup_pyo3_python() {
-  log_section "Installing Python ${PYO3_PYTHON_VERSION} for PyO3 tests"
+  local python_version="${PYO3_PYTHON_VERSION:-3.12}"
-  uv python install "$PYO3_PYTHON_VERSION"
+
+  log_section "Installing Python ${python_version} for PyO3 tests"
+  uv python install "$python_version"
   PYO3_PYTHON="$(uv python find \
     --managed-python \
     --no-project \
     --resolve-links \
-    "$PYO3_PYTHON_VERSION")"
+    "$python_version")"
   export PYO3_PYTHON
 
   local python_libdir
@@ -143,7 +151,6 @@ PY
 }
 
 run_style_clippy() {
-  install_cargo_binstall
   install_cargo_sort
   install_cargo_deny
 
@@ -156,8 +163,8 @@ run_style_clippy() {
   log_section "Checking Rust dependency bans"
   cargo deny \
     --manifest-path rust/Cargo.toml \
-    --config rust/deny.toml \
     check \
+    --config rust/deny.toml \
     bans
 
   log_section "Running clippy"
@@ -174,7 +181,6 @@ run_style_clippy() {
 run_tests() {
   install_uv
   setup_pyo3_python
-  install_cargo_binstall
   install_cargo_nextest
 
   log_section "Running cargo nextest"

.buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_dp4_async_eplb.sh

@@ -18,10 +18,6 @@ wait_for_server() {
 
 MODEL="Qwen/Qwen3-30B-A3B-FP8"
 BACK="allgather_reducescatter"
-if command -v rocm-smi &> /dev/null || [[ -d /opt/rocm ]] || [[ -n "${ROCM_PATH:-}" ]]; then
-  # Disable MOE padding for ROCm since it is causing eplb to fail.
-  export VLLM_ROCM_MOE_PADDING=0
-fi
 
 cleanup() {
   if [[ -n "${SERVER_PID:-}" ]] && kill -0 "${SERVER_PID}" 2>/dev/null; then

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

@@ -23,7 +23,6 @@ NC='\033[0m' # No Color
 # Default configuration
 PIPELINE="ci"
 DRY_RUN=true
-TORCH_NIGHTLY=false
 
 usage() {
     cat <<EOF
@@ -35,14 +34,12 @@ Sets RUN_ALL=1 and NIGHTLY=1 environment variables.
 SAFETY: Dry-run by default. Use --execute to actually trigger a build.
 
 Options:
-    --execute        Actually trigger the build (default: dry-run)
+    --execute       Actually trigger the build (default: dry-run)
-    --pipeline       Buildkite pipeline slug (default: ${PIPELINE})
+    --pipeline      Buildkite pipeline slug (default: ${PIPELINE})
-    --commit         Override commit SHA (default: current HEAD)
+    --commit        Override commit SHA (default: current HEAD)
-    --branch         Override branch name (default: current branch)
+    --branch        Override branch name (default: current branch)
-    --message        Custom build message (default: auto-generated)
+    --message       Custom build message (default: auto-generated)
-    --torch-nightly  Also build and run the full suite against torch nightly
+    --help          Show this help message
-                     (sets TORCH_NIGHTLY=1)
-    --help           Show this help message
 
 Prerequisites:
     - bk CLI installed: brew tap buildkite/buildkite && brew install buildkite/buildkite/bk
@@ -52,7 +49,6 @@ 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
 }
@@ -100,10 +96,6 @@ while [[ $# -gt 0 ]]; do
             MESSAGE="$2"
             shift 2
             ;;
-        --torch-nightly)
-            TORCH_NIGHTLY=true
-            shift
-            ;;
         --help|-h)
             usage
             ;;
@@ -179,17 +171,11 @@ 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: ${BUILD_ENV[*]}"
+log_info "Environment: RUN_ALL=1, NIGHTLY=1"
 echo ""
 
 # Build the command
@@ -201,10 +187,9 @@ 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 "=========================================="
@@ -225,14 +210,8 @@ if [[ "$DRY_RUN" == true ]]; then
     echo "    --commit '$(escape_for_shell "${COMMIT}")' \\"
     echo "    --branch '$(escape_for_shell "${BRANCH}")' \\"
     echo "    --message '$(escape_for_shell "${MESSAGE}")' \\"
-    last_idx=$(( ${#BUILD_ENV[@]} - 1 ))
+    echo "    --env 'RUN_ALL=1' \\"
-    for i in "${!BUILD_ENV[@]}"; do
+    echo "    --env 'NIGHTLY=1'"
-        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/upload-nightly-wheels.sh

@@ -6,14 +6,8 @@ set -ex
 # manylinux platform tag with auditwheel.
 # Index generation is handled separately by generate-and-upload-nightly-index.sh.
 
-# auditwheel is Linux-only; macOS wheels already carry a valid tag, so skip the
+# shellcheck source=lib/manylinux.sh
-# manylinux retag for them.
+source .buildkite/scripts/lib/manylinux.sh
-WHEEL_PLATFORM="${VLLM_WHEEL_PLATFORM:-linux}"
-
-if [[ "$WHEEL_PLATFORM" == "linux" ]]; then
-  # shellcheck source=lib/manylinux.sh
-  source .buildkite/scripts/lib/manylinux.sh
-fi
 
 BUCKET="vllm-wheels"
 SUBPATH=$BUILDKITE_COMMIT
@@ -33,10 +27,8 @@ wheel="${wheel_files[0]}"
 
 # ========= detect manylinux tag and rename ==========
 
-if [[ "$WHEEL_PLATFORM" == "linux" ]]; then
+wheel="$(apply_manylinux_tag "$wheel")"
-  wheel="$(apply_manylinux_tag "$wheel")"
+echo "Renamed wheel to: $wheel"
-  echo "Renamed wheel to: $wheel"
-fi
 
 # Extract the version from the wheel
 version=$(unzip -p "$wheel" '**/METADATA' | grep '^Version: ' | cut -d' ' -f2)

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

@@ -113,8 +113,8 @@ $PYTHON .buildkite/scripts/generate-nightly-index.py \
 echo "Uploading indices to $S3_COMMIT_PREFIX"
 aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "$S3_COMMIT_PREFIX"
 
-# Only scheduled nightly builds should update the moving nightly index.
+# Update rocm/nightly/ if on main branch and not a PR
-if [[ "${NIGHTLY:-0}" == "1" ]]; then
+if [[ "$BUILDKITE_BRANCH" == "main" && "$BUILDKITE_PULL_REQUEST" == "false" ]] || [[ "$NIGHTLY" == "1" ]]; then
     echo "Updating rocm/nightly/ index..."
     aws s3 cp --recursive "$INDICES_OUTPUT_DIR/" "s3://$BUCKET/rocm/nightly/"
 fi
@@ -147,7 +147,7 @@ echo ""
 echo "Install command (by commit):"
 echo "  pip install vllm --extra-index-url https://${BUCKET}.s3.amazonaws.com/$ROCM_SUBPATH/"
 echo ""
-if [[ "${NIGHTLY:-0}" == "1" ]]; then
+if [[ "$BUILDKITE_BRANCH" == "main" ]] || [[ "$NIGHTLY" == "1" ]]; then
     echo "Install command (nightly):"
     echo "  pip install vllm --extra-index-url https://${BUCKET}.s3.amazonaws.com/rocm/nightly/"
 fi

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

@@ -1,13 +0,0 @@
-#!/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

@@ -1,21 +0,0 @@
-#!/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_areas/attention.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: V1 attention (H100-MI300)
   key: v1-attention-h100-mi300
-  timeout_in_minutes: 85
+  timeout_in_minutes: 30
   device: h100
   source_file_dependencies:
     - vllm/config/attention.py
@@ -12,13 +12,11 @@ steps:
     - vllm/v1/attention
     - tests/v1/attention
   commands:
-    - pytest -v -s v1/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+    - pytest -v -s v1/attention
-  parallelism: 2
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 70
-      timeout_in_minutes: 125
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -32,7 +30,7 @@ steps:
 
 - label: V1 attention (B200)
   key: v1-attention-b200
-  timeout_in_minutes: 80
+  timeout_in_minutes: 30
   device: b200-k8s
   source_file_dependencies:
     - vllm/config/attention.py
@@ -40,5 +38,4 @@ steps:
     - vllm/v1/attention
     - tests/v1/attention
   commands:
-    - pytest -v -s v1/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+    - pytest -v -s v1/attention
-  parallelism: 2

.buildkite/test_areas/basic_correctness.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Basic Correctness
   key: basic-correctness
-  timeout_in_minutes: 68
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -18,8 +18,7 @@ steps:
   - pytest -v -s basic_correctness/test_cpu_offload.py
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 50
-      timeout_in_minutes: 60
       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: 45
+  timeout_in_minutes: 20
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -13,9 +13,7 @@ steps:
   - pytest -v -s benchmarks/
   mirror:
     amd:
-      dind: false
       device: mi300_1
-      timeout_in_minutes: 40
       depends_on:
       - image-build-amd
 
@@ -25,7 +23,7 @@ steps:
   num_gpus: 2
   optional: true
   working_dir: "/vllm-workspace/"
-  timeout_in_minutes: 20
+  timeout_in_minutes: 10
   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: 80
+  timeout_in_minutes: 50
   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: 75
+  timeout_in_minutes: 50
   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: 30
+  timeout_in_minutes: 50
   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: 30
+  timeout_in_minutes: 50
   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: 45
+  timeout_in_minutes: 20
   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: 30
+  timeout_in_minutes: 20
   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: 25
+  timeout_in_minutes: 15
   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: 25
+  timeout_in_minutes: 30
   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: 35
+  timeout_in_minutes: 20
   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: 30
+  timeout_in_minutes: 40
   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: 45
+  timeout_in_minutes: 20
   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: 20
+  timeout_in_minutes: 15
   device: h200_18gb
   source_file_dependencies:
   - vllm/envs.py
@@ -18,18 +18,14 @@ steps:
     - pytest -v -s cuda/test_platform_no_cuda_init.py
 
 - label: Cudagraph
-  device: h200_35gb
   key: cudagraph
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   source_file_dependencies:
   - tests/v1/cudagraph
   - vllm/v1/cudagraph_dispatcher.py
   - vllm/config/compilation.py
   - vllm/compilation
-  - vllm/v1/worker/encoder_cudagraph.py
-  - vllm/v1/worker/encoder_cudagraph_defs.py
   commands:
     - pytest -v -s v1/cudagraph/test_cudagraph_dispatch.py
     - pytest -v -s v1/cudagraph/test_cudagraph_mode.py
-    - pytest -v -s v1/cudagraph/test_breakable_cudagraph.py
+    - pytest -v -s v1/cudagraph/test_breakable_cudagraph.py
-    - pytest -v -s v1/cudagraph/test_encoder_cudagraph.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: 55
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -15,9 +15,8 @@ steps:
     - bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
   mirror:
     amd:
-      dind: false
       device: mi300_4
-      timeout_in_minutes: 60
+      timeout_in_minutes: 110
       depends_on:
         - image-build-amd
       source_file_dependencies:
@@ -30,7 +29,7 @@ steps:
 
 - label: Distributed FlashInfer NixlConnector PD accuracy (4 GPUs)
   key: distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus
-  timeout_in_minutes: 55
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -40,19 +39,6 @@ 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
@@ -66,9 +52,8 @@ steps:
     - DP_EP=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
   mirror:
     amd:
-      dind: false
       device: mi300_4
-      timeout_in_minutes: 40
+      timeout_in_minutes: 50
       depends_on:
         - image-build-amd
       source_file_dependencies:
@@ -81,7 +66,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: 55
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -92,9 +77,8 @@ steps:
     - CROSS_LAYERS_BLOCKS=True bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
   mirror:
     amd:
-      dind: false
       device: mi300_4
-      timeout_in_minutes: 60
+      timeout_in_minutes: 110
       depends_on:
         - image-build-amd
       source_file_dependencies:
@@ -107,7 +91,7 @@ steps:
 
 - label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs)
   key: hybrid-ssm-nixlconnector-pd-accuracy-tests-4-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -118,9 +102,8 @@ steps:
     - HYBRID_SSM=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
   mirror:
     amd:
-      dind: false
       device: mi300_4
-      timeout_in_minutes: 55
+      timeout_in_minutes: 60
       depends_on:
         - image-build-amd
       source_file_dependencies:
@@ -147,7 +130,7 @@ steps:
 
 - label: MultiConnector (Nixl+Offloading) PD accuracy (2 GPUs)
   key: multiconnector-nixl-offloading-pd-accuracy-2-gpus
-  timeout_in_minutes: 40
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -162,7 +145,7 @@ steps:
 
 - label: NixlConnector PD + Spec Decode acceptance (2 GPUs)
   key: nixlconnector-pd-spec-decode-acceptance-2-gpus
-  timeout_in_minutes: 45
+  timeout_in_minutes: 30
   device: a100
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
@@ -175,9 +158,8 @@ steps:
     - bash v1/kv_connector/nixl_integration/config_sweep_spec_decode_test.sh
   mirror:
     amd:
-      dind: false
       device: mi300_2
-      timeout_in_minutes: 45
+      timeout_in_minutes: 60
       depends_on:
         - image-build-amd
       source_file_dependencies:
@@ -187,11 +169,11 @@ steps:
         - vllm/platforms/rocm.py
       commands:
         - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
-        - KV_CACHE_MEMORY_BYTES=8G ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_spec_decode_test.sh
+        - ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_spec_decode_test.sh
 
 - label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs)
   key: multiconnector-nixl-offloading-pd-edge-cases-2-gpus
-  timeout_in_minutes: 25
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -203,25 +185,3 @@ steps:
   commands:
     - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
     - bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh
-
-# P TP 4 - D DPEP 4 test case for DSv4-Flash
-- label: DSv4-Flash Disaggregated DP EP
-  key: dsv4-flash-disaggregated
-  timeout_in_minutes: 60
-  device: h200
-  optional: true
-  working_dir: "/vllm-workspace/tests"
-  num_devices: 8
-  env:
-    ENABLE_HMA_FLAG: "1"
-    DP_EP: "1"
-    GPU_MEMORY_UTILIZATION: "0.85"
-    PREFILLER_TP_SIZE: "4"
-    DECODER_TP_SIZE: "4"
-    PREFILL_BLOCK_SIZE: "256"
-    DECODE_BLOCK_SIZE: "256"
-    MODEL_NAMES: "deepseek-ai/DeepSeek-V4-Flash"
-    VLLM_SERVE_EXTRA_ARGS: "--trust-remote-code,--kv-cache-dtype,fp8"
-  commands:
-    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
-    - bash v1/kv_connector/nixl_integration/run_accuracy_test.sh

.buildkite/test_areas/distributed.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed Comm Ops
   key: distributed-comm-ops
-  timeout_in_minutes: 25
+  timeout_in_minutes: 20
   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: 35
+  timeout_in_minutes: 20
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -39,9 +39,7 @@ steps:
   - DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py
   mirror:
     amd:
-      dind: false
       device: mi300_2
-      timeout_in_minutes: 45
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -57,7 +55,7 @@ steps:
 
 - label: Distributed Compile + RPC Tests (2 GPUs)
   key: distributed-compile-rpc-tests-2-gpus
-  timeout_in_minutes: 65
+  timeout_in_minutes: 20
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -80,7 +78,7 @@ steps:
 
 - label: Distributed Torchrun + Shutdown Tests (2 GPUs)
   key: distributed-torchrun-shutdown-tests-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -135,7 +133,7 @@ steps:
 
 - label: Distributed DP Tests (4 GPUs)
   key: distributed-dp-tests-4-gpus
-  timeout_in_minutes: 45
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -156,7 +154,7 @@ steps:
 
 - label: Distributed Compile + Comm (4 GPUs)
   key: distributed-compile-comm-4-gpus
-  timeout_in_minutes: 70
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -178,7 +176,7 @@ steps:
 
 - label: Distributed Tests (8xH100)
   key: distributed-tests-8xh100
-  timeout_in_minutes: 20
+  timeout_in_minutes: 10
   device: h100
   num_devices: 8
   working_dir: "/vllm-workspace/tests"
@@ -214,7 +212,7 @@ steps:
 
 - label: Distributed Tests (2xH100-2xMI300)
   key: distributed-tests-2xh100-2xmi300
-  timeout_in_minutes: 30
+  timeout_in_minutes: 15
   device: h100
   optional: true
   working_dir: "/vllm-workspace/"
@@ -235,7 +233,7 @@ steps:
   num_devices: 2
   commands:
     - pytest -v -s tests/distributed/test_context_parallel.py
-    - pytest -v -s tests/distributed/test_nccl_symm_mem.py
+    - pytest -v -s tests/distributed/test_nccl_symm_mem_allreduce.py
     - pytest -v -s tests/v1/distributed/test_dbo.py
     - pytest -v -s tests/distributed/test_mnnvl_alltoall.py
 
@@ -261,7 +259,7 @@ steps:
 
 - label: Pipeline + Context Parallelism (4 GPUs)
   key: pipeline-context-parallelism-4-gpus
-  timeout_in_minutes: 55
+  timeout_in_minutes: 60
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -276,7 +274,7 @@ steps:
 
 - label: RayExecutorV2 (4 GPUs)
   key: rayexecutorv2-4-gpus
-  timeout_in_minutes: 45
+  timeout_in_minutes: 60
   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: 20
+  timeout_in_minutes: 10
   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: 25
+  timeout_in_minutes: 60
   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: 25
+  timeout_in_minutes: 60
   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: 20
+  timeout_in_minutes: 60
   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: 25
+  timeout_in_minutes: 60
   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: 20
+  timeout_in_minutes: 60
   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: 30
+  timeout_in_minutes: 15
   device: h200_18gb
   source_file_dependencies:
   - vllm/compilation/
@@ -28,15 +28,14 @@ steps:
   - pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py test_jit_monitor.py
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 60
-      timeout_in_minutes: 40
       depends_on:
       - image-build-amd
 
 - label: Engine (1 GPU)
   key: engine-1-gpu
-  timeout_in_minutes: 45
+  timeout_in_minutes: 30
   source_file_dependencies:
     - vllm/v1/engine/
     - tests/v1/engine/
@@ -45,14 +44,14 @@ steps:
     - pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py
   mirror:
     amd:
-      device: mi250_1
+      device: mi325_1
-      timeout_in_minutes: 45
+      timeout_in_minutes: 40
       depends_on:
       - image-build-amd
 
 - label: e2e Scheduling (1 GPU)
   key: e2e-scheduling-1-gpu
-  timeout_in_minutes: 53
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/
@@ -62,14 +61,14 @@ steps:
   mirror:
     amd:
       device: mi250_1
-      timeout_in_minutes: 55
+      timeout_in_minutes: 60
       depends_on:
       - image-build-amd
 
 - label: e2e Core (1 GPU)
   device: h200_35gb
   key: e2e-core-1-gpu
-  timeout_in_minutes: 40
+  timeout_in_minutes: 30
   source_file_dependencies:
     - vllm/v1/
     - tests/v1/e2e/general/
@@ -78,7 +77,7 @@ steps:
   mirror:
     amd:
       device: mi250_1
-      timeout_in_minutes: 50
+      timeout_in_minutes: 35
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -88,7 +87,7 @@ steps:
 
 - label: V1 e2e (2 GPUs)
   key: v1-e2e-2-gpus
-  timeout_in_minutes: 25 # TODO: Fix timeout after we have more confidence in the test stability
+  timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
   optional: true
   num_devices: 2
   source_file_dependencies:
@@ -115,15 +114,13 @@ steps:
     - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "tensor_parallelism"
   mirror:
     amd:
-      dind: false
       device: mi300_2
-      timeout_in_minutes: 30
       depends_on:
       - image-build-amd
 
 - label: V1 e2e (4 GPUs)
   key: v1-e2e-4-gpus
-  timeout_in_minutes: 20 # TODO: Fix timeout after we have more confidence in the test stability
+  timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
   optional: true
   num_devices: 4
   source_file_dependencies:
@@ -151,7 +148,7 @@ steps:
 
 - label: V1 e2e (4xH100)
   key: v1-e2e-4xh100
-  timeout_in_minutes: 35
+  timeout_in_minutes: 60
   device: h100
   num_devices: 4
   optional: true

.buildkite/test_areas/entrypoints.yaml

@@ -3,9 +3,8 @@ depends_on:
   - image-build
 steps:
 - label: Entrypoints Unit Tests
-  device: h200_35gb
   key: entrypoints-unit-tests
-  timeout_in_minutes: 25
+  timeout_in_minutes: 10
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/entrypoints
@@ -16,9 +15,8 @@ steps:
   - pytest -v -s entrypoints/weight_transfer
 
 - label: Entrypoints Integration (LLM)
-  device: h200_35gb
   key: entrypoints-integration-llm
-  timeout_in_minutes: 60
+  timeout_in_minutes: 40
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -30,16 +28,16 @@ steps:
   - pytest -v -s entrypoints/llm/offline_mode # Needs to avoid interference with other tests
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      # TODO(akaratza): Test after Torch >= 2.12 bump
-      timeout_in_minutes: 55
+      soft_fail: true
       depends_on:
       - image-build-amd
 
 - label: Entrypoints Integration (API Server)
   key: entrypoints-integration-api-server
   device: h200_35gb
-  timeout_in_minutes: 75
+  timeout_in_minutes: 130
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -52,16 +50,13 @@ steps:
   - pytest -v -s entrypoints/scale_out
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
-      timeout_in_minutes: 65
       depends_on:
       - image-build-amd
 
 - label: Entrypoints Integration (API Server OpenAI - Part 1)
-  device: h200_35gb
   key: entrypoints-integration-api-server-openai-part-1
-  timeout_in_minutes: 68
+  timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -72,16 +67,14 @@ steps:
   - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/correctness
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 80
-      timeout_in_minutes: 65
       depends_on:
       - image-build-amd
 
 - label: Entrypoints Integration (API Server OpenAI - Part 2)
-  device: h200_35gb
   key: entrypoints-integration-api-server-openai-part-2
-  timeout_in_minutes: 83
+  timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -93,14 +86,12 @@ steps:
   - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 80
-      timeout_in_minutes: 70
       depends_on:
       - image-build-amd
 
 - label: Entrypoints Integration (API Server Generate)
-  device: h200_35gb
   key: entrypoints-integration-api-server-generate
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
@@ -117,14 +108,12 @@ steps:
   - pytest -v -s entrypoints/anthropic
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 60
-      timeout_in_minutes: 65
       depends_on:
       - image-build-amd
 
 - label: Entrypoints Integration (Responses API)
-  device: h200_35gb
   key: entrypoints-integration-responses-api
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
@@ -137,7 +126,7 @@ steps:
 - label: Entrypoints Integration (Speech to Text)
   device: h200_35gb
   key: entrypoints-integration-speech_to_text
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -149,7 +138,7 @@ steps:
 - label: Entrypoints Integration (Multimodal)
   device: h200_35gb
   key: entrypoints-integration-multimodal
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -159,9 +148,8 @@ steps:
   - pytest -v -s entrypoints/multimodal
 
 - label: Entrypoints Integration (Pooling)
-  device: h200_35gb
   key: entrypoints-integration-pooling
-  timeout_in_minutes: 75
+  timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -172,7 +160,7 @@ steps:
 
 - label: OpenAI API Correctness
   key: openai-api-correctness
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
   - csrc/
@@ -181,9 +169,7 @@ steps:
   - pytest -s entrypoints/openai/correctness/
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
-      timeout_in_minutes: 30
       depends_on:
       - image-build-amd
       source_file_dependencies:

.buildkite/test_areas/expert_parallelism.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: EPLB Algorithm
   key: eplb-algorithm
-  timeout_in_minutes: 20
+  timeout_in_minutes: 15
   device: h200_18gb
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -16,9 +16,7 @@ steps:
   - pytest -v -s distributed/test_eplb_utils.py
   mirror:
     amd:
-      dind: false
       device: mi300_1
-      timeout_in_minutes: 30
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -29,7 +27,7 @@ steps:
 
 - label: EPLB Execution # 17min
   key: eplb-execution
-  timeout_in_minutes: 25
+  timeout_in_minutes: 27
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
@@ -41,7 +39,7 @@ steps:
 
 - label: Elastic EP Scaling Test
   key: elastic-ep-scaling-test
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   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: 35
+  timeout_in_minutes: 10
   device: h200_18gb
   working_dir: "/vllm-workspace/"
   source_file_dependencies:
@@ -15,21 +15,19 @@ steps:
     - pytest -v -s tests/kernels/ir
 
 - label: Kernels Core Operation Test
-  device: h200_35gb
   key: kernels-core-operation-test
-  timeout_in_minutes: 120
+  timeout_in_minutes: 75
   source_file_dependencies:
   - csrc/
   - tests/kernels/core
   - tests/kernels/test_concat_mla_q.py
   - tests/kernels/test_fused_qk_norm_rope_gate.py
   commands:
-    - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py kernels/test_concat_mla_q.py kernels/test_fused_qk_norm_rope_gate.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+    - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py  kernels/test_concat_mla_q.py kernels/test_fused_qk_norm_rope_gate.py
-  parallelism: 3
 
 - label: Kernels MiniMax Reduce RMS Test (2 GPUs)
   key: kernels-minimax-reduce-rms-test-2-gpus
-  timeout_in_minutes: 20
+  timeout_in_minutes: 15
   num_devices: 2
   device: h100
   source_file_dependencies:
@@ -43,7 +41,7 @@ steps:
 
 - label: Deepseek V4 Kernel Test (H100)
   key: deepseek-v4-kernel-test-h100
-  timeout_in_minutes: 30
+  timeout_in_minutes: 15
   device: h100
   source_file_dependencies:
   - csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
@@ -56,7 +54,7 @@ steps:
 
 - label: Deepseek V4 Kernel Test (B200)
   key: deepseek-v4-kernel-test-b200
-  timeout_in_minutes: 20
+  timeout_in_minutes: 15
   device: b200-k8s
   source_file_dependencies:
   - csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
@@ -67,7 +65,7 @@ steps:
 
 - label: Kernels Attention Test %N
   key: kernels-attention-test
-  timeout_in_minutes: 65
+  timeout_in_minutes: 35
   source_file_dependencies:
   - csrc/attention/
   - vllm/v1/attention
@@ -80,9 +78,8 @@ steps:
   parallelism: 2
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 55
-      timeout_in_minutes: 90
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -109,7 +106,7 @@ steps:
 
 - label: Kernels Quantization Test %N
   key: kernels-quantization-test
-  timeout_in_minutes: 60
+  timeout_in_minutes: 90
   source_file_dependencies:
   - csrc/quantization/
   - vllm/model_executor/layers/quantization
@@ -119,9 +116,7 @@ steps:
   parallelism: 2
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
-      timeout_in_minutes: 120
       source_file_dependencies:
       - csrc/quantization/
       - vllm/model_executor/layers/quantization
@@ -136,7 +131,7 @@ steps:
 
 - label: Kernels MoE Test %N
   key: kernels-moe-test
-  timeout_in_minutes: 50
+  timeout_in_minutes: 25
   source_file_dependencies:
   - csrc/quantization/cutlass_w8a8/moe/
   - csrc/moe/
@@ -151,9 +146,8 @@ steps:
   parallelism: 5
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 50
-      timeout_in_minutes: 55
       source_file_dependencies:
       - csrc/quantization/cutlass_w8a8/moe/
       - csrc/moe/
@@ -168,9 +162,8 @@ steps:
       - image-build-amd
 
 - label: Kernels Mamba Test
-  device: h200_35gb
   key: kernels-mamba-test
-  timeout_in_minutes: 40
+  timeout_in_minutes: 45
   source_file_dependencies:
   - csrc/mamba/
   - tests/kernels/mamba
@@ -179,19 +172,19 @@ steps:
     - pytest -v -s kernels/mamba
 
 - label: Kernels KDA Test
-  timeout_in_minutes: 25
+  timeout_in_minutes: 20
   device: h200_18gb
   source_file_dependencies:
-  - vllm/third_party/flash_linear_attention/ops/kda.py
+  - vllm/model_executor/layers/fla/ops/kda.py
-  - vllm/third_party/flash_linear_attention/ops/chunk_delta_h.py
+  - vllm/model_executor/layers/fla/ops/chunk_delta_h.py
-  - vllm/third_party/flash_linear_attention/ops/l2norm.py
+  - vllm/model_executor/layers/fla/ops/l2norm.py
   - tests/kernels/test_kda.py
   commands:
     - pytest -v -s kernels/test_kda.py
 
 - label: Kernels DeepGEMM Test (H100)
   key: kernels-deepgemm-test-h100
-  timeout_in_minutes: 35
+  timeout_in_minutes: 45
   device: h100
   num_devices: 1
   source_file_dependencies:
@@ -218,7 +211,7 @@ steps:
 
 - label: Kernels (B200)
   key: kernels-b200
-  timeout_in_minutes: 80
+  timeout_in_minutes: 30
   working_dir: "/vllm-workspace/"
   device: b200-k8s
   # optional: true
@@ -238,15 +231,6 @@ steps:
   - vllm/v1/attention/backends/mla/flashinfer_mla.py
   - vllm/v1/attention/selector.py
   - vllm/platforms/cuda.py
-  - vllm/model_executor/kernels/linear/cute_dsl/ll_bf16.py
-  - vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_dotprod.py
-  - vllm/model_executor/kernels/linear/cute_dsl/_ll_bf16_splitk.py
-  - vllm/cute_utils/
-  - vllm/model_executor/layers/mamba/ops/gdn_chunk_cutedsl/
-  - vllm/model_executor/layers/fused_moe/router/bf16x3_router_gemm_cutedsl.py
-  - tests/kernels/mamba/test_gdn_prefill_cutedsl.py
-  - tests/kernels/test_bf16x3_router_gemm_cutedsl.py
-  - tests/kernels/test_ll_bf16_gemm.py
   - tests/kernels/test_top_k_per_row.py
   commands:
     - nvidia-smi
@@ -275,28 +259,24 @@ steps:
     - pytest -v -s tests/kernels/moe/test_flashinfer_moe.py
     - pytest -v -s tests/kernels/moe/test_trtllm_nvfp4_moe.py
     - pytest -v -s tests/kernels/moe/test_cutedsl_moe.py
-    - pytest -v -s tests/kernels/mamba/test_gdn_prefill_cutedsl.py
-    - pytest -v -s tests/kernels/test_bf16x3_router_gemm_cutedsl.py
-    - pytest -v -s tests/kernels/test_ll_bf16_gemm.py
     # e2e
     - pytest -v -s tests/models/quantization/test_nvfp4.py
 
 - label: Kernels Helion Test
   key: kernels-helion-test
-  timeout_in_minutes: 115
+  timeout_in_minutes: 30
   device: h100
   source_file_dependencies:
   - vllm/utils/import_utils.py
   - tests/kernels/helion/
   commands:
     - pip install helion==1.1.0
-    - pytest -v -s kernels/helion/ --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+    - pytest -v -s kernels/helion/
-  parallelism: 2
 
  
 - label: Kernels FP8 MoE Test (1xH100)
   key: kernels-fp8-moe-test-1xh100
-  timeout_in_minutes: 40
+  timeout_in_minutes: 90
   device: h100
   num_devices: 1
   optional: true
@@ -313,7 +293,7 @@ steps:
 
 - label: Kernels FP8 MoE Test (2xH100)
   key: kernels-fp8-moe-test-2xh100
-  timeout_in_minutes: 45
+  timeout_in_minutes: 90
   device: h100
   num_devices: 2
   optional: true
@@ -323,7 +303,7 @@ steps:
 
 - label: Kernels Fp4 MoE Test (B200)
   key: kernels-fp4-moe-test-b200
-  timeout_in_minutes: 25
+  timeout_in_minutes: 60
   device: b200-k8s
   num_devices: 1
   optional: true
@@ -336,7 +316,7 @@ steps:
 
 - label: Kernels FusedMoE Layer Test (2 H100s)
   key: kernels-fusedmoe-layer-test-2-h100s
-  timeout_in_minutes: 30
+  timeout_in_minutes: 90
   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: 45
+  timeout_in_minutes: 75
   source_file_dependencies:
   - csrc/
   - vllm/model_executor/layers/quantization
@@ -14,9 +14,8 @@ steps:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-small.txt
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 55
-      timeout_in_minutes: 45
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -55,9 +54,9 @@ steps:
     - export VLLM_USE_DEEP_GEMM=0  # We found Triton is faster than DeepGEMM for H100
     - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-hopper.txt --tp-size=4
 
-- label: LM Eval Small Models (1xB200)
+- label: LM Eval Small Models (2xB200)
-  key: lm-eval-small-models-1xb200
+  key: lm-eval-small-models-2xb200
-  timeout_in_minutes: 50
+  timeout_in_minutes: 120
   device: b200-k8s
   optional: true
   source_file_dependencies:
@@ -66,10 +65,9 @@ steps:
   commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-blackwell.txt
 
-- label: LM Eval Small Models Distributed (2xB200)
+- label: LM Eval Small Models (2xL4)
-  key: lm-eval-small-models-distributed-2xb200
+  key: lm-eval-small-models-tp
-  timeout_in_minutes: 120
+  timeout_in_minutes: 10
-  device: b200-k8s
   num_devices: 2
   optional: true
   source_file_dependencies:
@@ -79,31 +77,9 @@ steps:
   commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-small-tp.txt
 
-- label: LM Eval PCP (4xB200)
-  key: lm-eval-pcp-4xb200
-  timeout_in_minutes: 360
-  device: b200-k8s
-  num_devices: 4
-  optional: true
-  source_file_dependencies:
-  - csrc/
-  - tests/evals/gsm8k/configs/GLM-5.2-NVFP4-TP2-PCP2-EP.yaml
-  - tests/evals/gsm8k/configs/GLM-5.2-NVFP4-TP1-PCP4-EP.yaml
-  - tests/evals/gsm8k/configs/models-pcp.txt
-  - vllm/model_executor/layers/quantization
-  - vllm/config/parallel.py
-  - vllm/distributed/parallel_state.py
-  - vllm/model_executor/layers/attention/mla_attention.py
-  - vllm/model_executor/layers/attention/pcp.py
-  - vllm/v1/worker/gpu/model_runner.py
-  - vllm/v1/worker/gpu/pcp_manager.py
-  autorun_on_main: true
-  commands:
-  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-pcp.txt
-
 - label: LM Eval Large Models EP (2xB200)
   key: lm-eval-large-models-ep-2xb200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 120
   device: b200-k8s
   optional: true
   num_devices: 2
@@ -115,7 +91,7 @@ steps:
 
 - label: LM Eval Qwen3.5 Models (2xB200)
   key: lm-eval-qwen3-5-models-2xb200
-  timeout_in_minutes: 45
+  timeout_in_minutes: 120
   device: b200-k8s
   optional: true
   num_devices: 2
@@ -126,13 +102,13 @@ steps:
   - vllm/transformers_utils/configs/qwen3_5_moe.py
   - vllm/model_executor/models/qwen3_next.py
   - vllm/model_executor/models/qwen3_next_mtp.py
-  - vllm/third_party/flash_linear_attention/ops/
+  - vllm/model_executor/layers/fla/ops/
   commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-qwen35-blackwell.txt
 
 - label: LM Eval Large Models (8xH200)
   key: lm-eval-large-models-8xh200
-  timeout_in_minutes: 50
+  timeout_in_minutes: 60
   device: h200
   optional: true
   num_devices: 8
@@ -140,9 +116,8 @@ steps:
     - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-h200.txt
   mirror:
     amd:
-      dind: false
       device: mi300_8
-      timeout_in_minutes: 40
+      timeout_in_minutes: 180
       depends_on:
       - image-build-amd
       commands:
@@ -174,9 +149,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 f16 (A100 - TEMPORARY)
+- label: LM Eval Humming (A100 - TEMPORARY)
-  key: lm-eval-humming-f16-a100
+  key: lm-eval-humming-a100
-  timeout_in_minutes: 75
+  timeout_in_minutes: 30
   device: a100
   optional: true
   num_devices: 1
@@ -184,29 +159,13 @@ 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/
+  - vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
-  - 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 int8 (A100 - TEMPORARY)
+- label: LM Eval Humming (H100 - TEMPORARY)
-  key: lm-eval-humming-act-a100
+  key: lm-eval-humming-h100
-  timeout_in_minutes: 45
+  timeout_in_minutes: 30
-  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
@@ -214,30 +173,14 @@ 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/
+  - vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
-  - 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 (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)
+- label: LM Eval Humming (B200 - TEMPORARY)
-  key: lm-eval-humming-f16-b200
+  key: lm-eval-humming-b200
-  timeout_in_minutes: 50
+  timeout_in_minutes: 30
   device: b200-k8s
   optional: true
   num_devices: 1
@@ -245,30 +188,14 @@ 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/
+  - vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
-  - 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: 55
+  timeout_in_minutes: 75
   device: h200_18gb
   source_file_dependencies:
   - vllm/model_executor/layers/quantization/turboquant/
@@ -280,7 +207,7 @@ steps:
 
 - label: GPQA Eval (GPT-OSS) (2xH100)
   key: gpqa-eval-gpt-oss-2xh100
-  timeout_in_minutes: 35
+  timeout_in_minutes: 120
   device: h100
   optional: true
   num_devices: 2
@@ -294,7 +221,7 @@ steps:
 
 - label: GPQA Eval (GPT-OSS) (2xB200)
   key: gpqa-eval-gpt-oss-2xb200
-  timeout_in_minutes: 30
+  timeout_in_minutes: 120
   device: b200-k8s
   optional: true
   num_devices: 2
@@ -308,7 +235,7 @@ steps:
 
 - label: GPQA Eval (GPT-OSS) (DGX Spark)
   key: gpqa-eval-gpt-oss-spark
-  timeout_in_minutes: 35
+  timeout_in_minutes: 120
   device: dgx-spark
   optional: true
   num_devices: 1
@@ -322,50 +249,9 @@ 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: 25
+  timeout_in_minutes: 30
   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: 40
+  timeout_in_minutes: 30
   source_file_dependencies:
   - vllm/lora
   - tests/lora
@@ -14,10 +14,9 @@ steps:
   parallelism: 4
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
       working_dir: "/vllm-workspace/tests"
-      timeout_in_minutes: 85
+      timeout_in_minutes: 60
       source_file_dependencies:
       - vllm/lora
       - tests/lora
@@ -28,7 +27,7 @@ steps:
 
 - label: LoRA TP (Distributed)
   key: lora-tp-distributed
-  timeout_in_minutes: 60
+  timeout_in_minutes: 30
   num_devices: 4
   source_file_dependencies:
   - vllm/lora
@@ -47,4 +46,4 @@ steps:
     - pytest -v -s -x lora/test_qwen3_with_multi_loras.py
     - pytest -v -s -x lora/test_olmoe_tp.py
     - pytest -v -s -x lora/test_gptoss_tp.py
-    - pytest -v -s -x lora/test_qwen35_densemodel_lora.py
+    - pytest -v -s -x lora/test_qwen35_densemodel_lora.py

.buildkite/test_areas/misc.yaml

@@ -5,7 +5,7 @@ steps:
 - label: V1 Spec Decode
   device: h200_35gb
   key: v1-spec-decode
-  timeout_in_minutes: 40
+  timeout_in_minutes: 30
   source_file_dependencies:
     - vllm/config/
     - vllm/distributed/
@@ -23,15 +23,14 @@ steps:
     - pytest -v -s -m 'not slow_test' v1/spec_decode
   mirror:
     amd:
-      dind: false
       device: mi300_1
-      timeout_in_minutes: 50
+      timeout_in_minutes: 65
       depends_on:
       - image-build-amd
 
 - label: V1 Sample + Logits
   key: v1-sample-logits
-  timeout_in_minutes: 83
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
     - vllm/config/
@@ -59,16 +58,13 @@ steps:
     - pytest -v -s v1/test_outputs.py
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
-      timeout_in_minutes: 70
       depends_on:
       - image-build-amd
 
 - label: V1 Core + KV + Metrics
-  device: h200_35gb
   key: v1-core-kv-metrics
-  timeout_in_minutes: 80
+  timeout_in_minutes: 30
   source_file_dependencies:
     - vllm/config/
     - vllm/distributed/
@@ -92,9 +88,7 @@ steps:
     - tests/v1/kv_offload
     - tests/v1/simple_kv_offload
     - tests/v1/worker
-    - tests/v1/streaming_input
     - tests/v1/kv_connector/unit
-    - tests/v1/ec_connector/unit
     - tests/v1/metrics
     - tests/entrypoints/openai/correctness/test_lmeval.py
   commands:
@@ -106,18 +100,15 @@ steps:
     - pytest -v -s v1/kv_offload
     - pytest -v -s v1/simple_kv_offload
     - pytest -v -s v1/worker
-    - pytest -v -s v1/streaming_input
     - pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
-    - pytest -v -s -m 'not cpu_test' v1/ec_connector/unit
     - pytest -v -s -m 'not cpu_test' v1/metrics
     # Integration test for streaming correctness (requires special branch).
-    - pip install -U git+https://github.com/vllm-project/lm-evaluation-harness.git@streaming-api
+    - pip install -U git+https://github.com/robertgshaw2-redhat/lm-evaluation-harness.git@streaming-api
     - pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 60
-      timeout_in_minutes: 65
       depends_on:
       - image-build-amd
 
@@ -148,7 +139,6 @@ steps:
     - pytest -v -s -m 'cpu_test' v1/core
     - pytest -v -s v1/structured_output
     - pytest -v -s v1/test_serial_utils.py
-    - pytest -v -s v1/cudagraph/test_cudagraph_manager.py
     - pytest -v -s -m 'cpu_test' v1/kv_connector/unit
     - pytest -v -s -m 'cpu_test' v1/metrics
 
@@ -182,7 +172,7 @@ steps:
 
 - label: Regression
   key: regression
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
   device: h200_18gb
   source_file_dependencies:
   - vllm/config/
@@ -198,21 +188,21 @@ steps:
   - vllm/v1/
   - tests/test_regression
   commands:
-  - pip install 'modelscope<1.38'
+  - pip install modelscope
   - pytest -v -s test_regression.py
   working_dir: "/vllm-workspace/tests" # optional
 
 - label: Examples
   device: h200_35gb
   key: examples
-  timeout_in_minutes: 40
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/examples"
   source_file_dependencies:
   - vllm/entrypoints
   - vllm/multimodal
   - examples/
   commands:
-    - pip install --no-deps tensorizer # for tensorizer test
+    - pip install tensorizer # for tensorizer test
      # for basic
     - python3 basic/offline_inference/chat.py
     - python3 basic/offline_inference/generate.py --model facebook/opt-125m
@@ -236,9 +226,7 @@ steps:
     - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
-      timeout_in_minutes: 75
       source_file_dependencies:
       - vllm/entrypoints
       - vllm/multimodal
@@ -249,7 +237,7 @@ steps:
 
 - label: Metrics, Tracing (2 GPUs)
   key: metrics-tracing-2-gpus
-  timeout_in_minutes: 25
+  timeout_in_minutes: 20
   num_devices: 2
   source_file_dependencies:
   - vllm/config/
@@ -274,11 +262,10 @@ steps:
   - pytest -v -s v1/tracing
   mirror:
     amd:
-      dind: false
+      device: mi325_2
-      device: mi300_2
-      timeout_in_minutes: 30
       depends_on:
       - image-build-amd
+      optional: true
 
 - label: Python-only Installation
   key: python-only-installation
@@ -293,8 +280,8 @@ steps:
   - bash standalone_tests/python_only_compile.sh
   mirror:
     amd:
-      device: mi250_1
+      device: mi325_1
-      timeout_in_minutes: 55
+      timeout_in_minutes: 20
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -305,7 +292,7 @@ steps:
 - label: Async Engine, Inputs, Utils, Worker
   device: h200_35gb
   key: async-engine-inputs-utils-worker
-  timeout_in_minutes: 25
+  timeout_in_minutes: 50
   source_file_dependencies:
   - vllm/assets/
   - vllm/config/
@@ -332,7 +319,7 @@ steps:
   key: async-engine-inputs-utils-worker-config-cpu
   depends_on:
   - image-build-cpu
-  timeout_in_minutes: 65
+  timeout_in_minutes: 30
   source_file_dependencies:
   - vllm/assets/
   - vllm/config/
@@ -364,7 +351,6 @@ 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
@@ -382,7 +368,6 @@ 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
@@ -394,7 +379,7 @@ steps:
 
 - label: Batch Invariance (A100)
   key: batch-invariance-a100
-  timeout_in_minutes: 40
+  timeout_in_minutes: 30
   device: a100
   source_file_dependencies:
     - vllm/v1/attention
@@ -408,7 +393,7 @@ steps:
 
 - label: Batch Invariance (H100)
   key: batch-invariance-h100
-  timeout_in_minutes: 40
+  timeout_in_minutes: 30
   device: h100
   source_file_dependencies:
     - vllm/v1/attention
@@ -424,7 +409,7 @@ steps:
 
 - label: Batch Invariance (B200)
   key: batch-invariance-b200
-  timeout_in_minutes: 35
+  timeout_in_minutes: 30
   device: b200-k8s
   source_file_dependencies:
     - vllm/v1/attention
@@ -443,7 +428,7 @@ steps:
 - label: Acceptance Length Test (Large Models) # optional
   device: h200_35gb
   key: acceptance-length-test-large-models
-  timeout_in_minutes: 20
+  timeout_in_minutes: 25
   gpu: h100
   optional: true
   num_gpus: 1

.buildkite/test_areas/model_executor.yaml

@@ -3,16 +3,13 @@ depends_on:
   - image-build
 steps:
 - label: Model Executor
-  device: h200_35gb
   key: model-executor
-  timeout_in_minutes: 60
+  timeout_in_minutes: 35
   source_file_dependencies:
   - vllm/engine/arg_utils.py
   - vllm/config/model.py
   - vllm/model_executor
-  - vllm/model_executor/warmup
   - tests/model_executor
-  - tests/model_executor/test_jit_warmup.py
   - tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
   commands:
     - apt-get update && apt-get install -y curl libsodium23
@@ -28,18 +25,14 @@ steps:
     - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py --timeout=900 --timeout-method=thread
   mirror:
     amd:
-      dind: false
       device: mi300_1
-      timeout_in_minutes: 60
       depends_on:
       - image-build-amd
       source_file_dependencies:
       - vllm/engine/arg_utils.py
       - vllm/config/model.py
       - vllm/model_executor
-      - vllm/model_executor/warmup
       - tests/model_executor
-      - tests/model_executor/test_jit_warmup.py
       - tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
       - vllm/_aiter_ops.py
       - vllm/platforms/rocm.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: 35
+  timeout_in_minutes: 45
   source_file_dependencies:
   - vllm/v1/worker/gpu/
   - vllm/v1/worker/gpu_worker.py
@@ -18,7 +18,9 @@ steps:
   - set -x
   - export VLLM_USE_V2_MODEL_RUNNER=1
   - pytest -v -s v1/engine/test_llm_engine.py -k "not test_engine_metrics"
-  - pytest -v -s v1/e2e/general/test_async_scheduling.py -k "not ngram"
+  # This requires eager until we sort out CG correctness issues.
+  # TODO: remove ENFORCE_EAGER here after https://github.com/vllm-project/vllm/pull/32936 is merged.
+  - 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
   # Temporary hack filter to exclude ngram spec decoding based tests.
@@ -27,7 +29,7 @@ steps:
 - label: Model Runner V2 Examples
   device: h200_35gb
   key: model-runner-v2-examples
-  timeout_in_minutes: 35
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/examples"
   source_file_dependencies:
     - vllm/v1/worker/gpu/
@@ -41,7 +43,7 @@ steps:
   commands:
     - set -x
     - export VLLM_USE_V2_MODEL_RUNNER=1
-    - pip install --no-deps tensorizer # for tensorizer test
+    - pip install tensorizer # for tensorizer test
     - python3 basic/offline_inference/chat.py # for basic
     - python3 basic/offline_inference/generate.py --model facebook/opt-125m
     #- python3 basic/offline_inference/generate.py --model meta-llama/Llama-2-13b-chat-hf --cpu-offload-gb 10  # TODO
@@ -63,7 +65,7 @@ steps:
 
 - label: Model Runner V2 Distributed (2 GPUs)
   key: model-runner-v2-distributed-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -84,7 +86,7 @@ steps:
 
 - label: Model Runner V2 Pipeline Parallelism (4 GPUs)
   key: model-runner-v2-pipeline-parallelism-4-gpus
-  timeout_in_minutes: 50
+  timeout_in_minutes: 60
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:

.buildkite/test_areas/models_basic.yaml

@@ -4,8 +4,9 @@ depends_on:
 steps:
 - label: Basic Models Tests (Initialization)
   key: basic-models-tests-initialization
-  timeout_in_minutes: 25
+  timeout_in_minutes: 45
   device: h200_18gb
+  torch_nightly: true
   source_file_dependencies:
   - vllm/
   - tests/models/test_initialization.py
@@ -13,11 +14,13 @@ steps:
   commands:
     # Run a subset of model initialization tests
     - pytest -v -s models/test_initialization.py::test_can_initialize_small_subset
+  mirror:
+    torch_nightly: {}
 
 - label: Basic Models Tests (Extra Initialization) %N
   device: h200_35gb
   key: basic-models-tests-extra-initialization
-  timeout_in_minutes: 100
+  timeout_in_minutes: 45
   source_file_dependencies:
   - vllm/model_executor/models/
   - tests/models/test_initialization.py
@@ -27,50 +30,36 @@ steps:
     # subset of supported models (the complement of the small subset in the above
     # test.) Also run if model initialization test file is modified
     - pytest -v -s models/test_initialization.py -k 'not test_can_initialize_small_subset' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
-  parallelism: 4
+  parallelism: 2
+  mirror:
+    torch_nightly: {}
 
 - label: Basic Models Tests (Other)
   device: h200_35gb
   key: basic-models-tests-other
-  timeout_in_minutes: 35
+  timeout_in_minutes: 45
   source_file_dependencies:
   - vllm/
   - tests/models/test_terratorch.py
-  - tests/models/transformers/test_backend.py
+  - tests/models/test_transformers.py
   - tests/models/test_registry.py
   commands:
-    - pytest -v -s models/test_terratorch.py models/transformers/test_backend.py models/test_registry.py
+    - pytest -v -s models/test_terratorch.py models/test_transformers.py models/test_registry.py
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
-      timeout_in_minutes: 50
       depends_on:
       - image-build-amd
 
-- label: Inkling Unit Tests (B200)
-  key: inkling-unit-tests-b200
-  timeout_in_minutes: 40
-  device: b200-k8s
-  source_file_dependencies:
-  - vllm/models/inkling/
-  - vllm/cute_utils/
-  - cmake/external_projects/tml_fa4.cmake
-  - tests/models/inkling/
-  commands:
-    # FA4 kernel tests require SM100; the suite skips them elsewhere.
-    - pytest -v -s models/inkling
-
 - label: Basic Models Test (Other CPU) # 5min
   key: basic-models-test-other-cpu
   depends_on:
   - image-build-cpu
-  timeout_in_minutes: 20
+  timeout_in_minutes: 10
   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 models/transformers/fusers/
+    - pytest -v -s models/test_utils.py models/test_vision.py

.buildkite/test_areas/models_distributed.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed Model Tests (2 GPUs)
   key: distributed-model-tests-2-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 50
   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/transformers/test_backend.py -v -s -m 'distributed(num_gpus=2)'
+  - pytest models/test_transformers.py -v -s -m 'distributed(num_gpus=2)'
   - pytest models/language -v -s -m 'distributed(num_gpus=2)'
   - pytest models/multimodal/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: 30
+  timeout_in_minutes: 25
   device: h200_18gb
   source_file_dependencies:
   - vllm/
@@ -14,17 +14,15 @@ steps:
     - pip freeze | grep -E 'torch'
     - pytest -v -s models/language -m 'core_model and (not slow_test)'
   mirror:
+    torch_nightly: {}
     amd:
-      dind: false
       device: mi300_1
-      timeout_in_minutes: 45
       depends_on:
       - image-build-amd
 
 - label: Language Models Tests (Extra Standard) %N
-  device: h200_35gb
   key: language-models-tests-extra-standard
-  timeout_in_minutes: 40
+  timeout_in_minutes: 45
   source_file_dependencies:
   - vllm/model_executor/models/
   - tests/models/language/pooling/test_embedding.py
@@ -37,10 +35,9 @@ steps:
     - pytest -v -s models/language -m 'core_model and slow_test' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
   parallelism: 2
   mirror:
+    torch_nightly: {}
     amd:
-      dind: false
       device: mi300_1
-      timeout_in_minutes: 40
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -54,25 +51,26 @@ steps:
       - tests/models/language/pooling/test_classification.py
       - vllm/_aiter_ops.py
       - vllm/platforms/rocm.py
+
 - label: Language Models Tests (Hybrid) %N
-  device: h200_35gb
   key: language-models-tests-hybrid
-  timeout_in_minutes: 65
+  timeout_in_minutes: 75
   source_file_dependencies:
   - vllm/
   - tests/models/language/generation
   commands:
     # Install fast path packages for testing against transformers
+    # Note: also needed to run plamo2 model in vLLM
     - uv pip install --system --no-build-isolation 'git+https://github.com/state-spaces/mamba@v2.3.0'
     - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
-    # Shard the hybrid language model tests that are numerically stable on Hopper.
+    # Shard hybrid language model tests
-    - pytest -v -s models/language/generation -m hybrid_model -k 'not granite-4.0-tiny-preview' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
+    - pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
   parallelism: 2
   mirror:
+    torch_nightly: {}
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 90
-      timeout_in_minutes: 60
       depends_on:
       - image-build-amd
       commands:
@@ -80,37 +78,24 @@ steps:
       - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
       - pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
 
-# Granite 4 hybrid generation is sensitive to hardware-specific Triton SSD
-# autotuning (https://github.com/vllm-project/vllm/issues/25194). Keep this one
-# correctness test on L4 until its H200 output matches the Transformers reference.
-- label: Language Models Tests (Granite L4 Compatibility)
-  key: language-models-tests-granite-l4-compatibility
-  timeout_in_minutes: 65
-  source_file_dependencies:
-  - vllm/
-  - tests/models/language/generation
-  commands:
-    - uv pip install --system --no-build-isolation 'git+https://github.com/state-spaces/mamba@v2.3.0'
-    - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
-    - pytest -v -s models/language/generation -m hybrid_model -k 'granite-4.0-tiny-preview'
-
 - label: Language Models Test (Extended Generation) # 80min
   device: h200_35gb
   key: language-models-test-extended-generation
-  timeout_in_minutes: 65
+  timeout_in_minutes: 110
   optional: true
   source_file_dependencies:
   - vllm/
   - tests/models/language/generation
   commands:
     # Install fast path packages for testing against transformers
+    # Note: also needed to run plamo2 model in vLLM
     - uv pip install --system --no-build-isolation 'git+https://github.com/state-spaces/mamba@v2.3.0'
     - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
     - pytest -v -s models/language/generation -m '(not core_model) and (not hybrid_model)'
 
 - label: Language Models Test (PPL)
   key: language-models-test-ppl
-  timeout_in_minutes: 30
+  timeout_in_minutes: 110
   device: h200_18gb
   optional: true
   source_file_dependencies:
@@ -119,10 +104,10 @@ steps:
   commands:
     - pytest -v -s models/language/generation_ppl_test
 
-- label: Language Models Test (Extended Pooling)
+- label: Language Models Test (Extended Pooling)  # 36min
   device: h200_35gb
   key: language-models-test-extended-pooling
-  timeout_in_minutes: 120
+  timeout_in_minutes: 50
   optional: true
   source_file_dependencies:
   - vllm/
@@ -131,15 +116,14 @@ steps:
     - pytest -v -s models/language/pooling -m 'not core_model'
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 100
-      timeout_in_minutes: 95
       depends_on:
       - image-build-amd
 
 - label: Language Models Test (MTEB)
   key: language-models-test-mteb
-  timeout_in_minutes: 68
+  timeout_in_minutes: 110
   device: h200_18gb
   optional: true
   source_file_dependencies:

.buildkite/test_areas/models_multimodal.yaml

@@ -4,107 +4,104 @@ depends_on:
 steps:
 - label: "Multi-Modal Models (Standard) 1: qwen2"
   key: multi-modal-models-standard-1-qwen2
-  timeout_in_minutes: 68
+  timeout_in_minutes: 45
   device: h200_18gb
   source_file_dependencies:
   - vllm/
   - tests/models/multimodal
   commands:
+    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
     - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen2"
     - pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
-      timeout_in_minutes: 65
       depends_on:
       - image-build-amd
 
 - label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
   key: multi-modal-models-standard-2-qwen3-gemma
-  timeout_in_minutes: 75
+  timeout_in_minutes: 45
   device: h200_18gb
   source_file_dependencies:
   - vllm/
   - tests/models/multimodal
   commands:
+    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
     - 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:
-      dind: false
+      device: mi325_1
-      device: mi300_1
-      timeout_in_minutes: 55
       depends_on:
       - image-build-amd
 
 - label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
   device: h200_35gb
   key: multi-modal-models-standard-3-llava-qwen2-vl
-  timeout_in_minutes: 40
+  timeout_in_minutes: 45
   source_file_dependencies:
   - vllm/
   - tests/models/multimodal
   commands:
+    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
     - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "not qwen2 and not qwen3 and not gemma"
     - pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model
   mirror:
     amd:
-      device: mi250_1
+      device: mi325_1
-      timeout_in_minutes: 55
       depends_on:
       - image-build-amd
 
 - label: "Multi-Modal Models (Standard) 4: other + whisper"
   device: h200_35gb
   key: multi-modal-models-standard-4-other-whisper
-  timeout_in_minutes: 75
+  timeout_in_minutes: 45
   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_mm_prefix_lm.py --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/generation/test_vit_cudagraph.py --ignore models/multimodal/processing
+    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
+    - 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/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
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
-      timeout_in_minutes: 50
       depends_on:
       - image-build-amd
 
-- label: Multi-Modal Processor (CPU) %N
+- label: Multi-Modal Processor (CPU)
   key: multi-modal-processor-cpu
   depends_on:
   - image-build-cpu
-  timeout_in_minutes: 125
+  timeout_in_minutes: 60
   source_file_dependencies:
   - vllm/
   - tests/models/multimodal
   - tests/models/registry.py
   device: cpu-medium
   commands:
-    - pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
+    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
-  parallelism: 4
+    - pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py
 
 - label: Multi-Modal Processor # 44min
   key: multi-modal-processor
-  timeout_in_minutes: 98
+  timeout_in_minutes: 60
   device: h200_18gb
   source_file_dependencies:
   - vllm/
   - tests/models/multimodal
   - tests/models/registry.py
   commands:
+    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
     - pytest -v -s models/multimodal/processing/test_tensor_schema.py
 
 - label: Multi-Modal Accuracy Eval (Small Models) # 50min
   device: h200_35gb
   key: multi-modal-accuracy-eval-small-models
-  timeout_in_minutes: 30
+  timeout_in_minutes: 70
   working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
   source_file_dependencies:
   - vllm/multimodal/
@@ -114,9 +111,7 @@ steps:
   - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-mm-small.txt --tp-size=1
   mirror:
     amd:
-      dind: false
       device: mi300_1
-      timeout_in_minutes: 35
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -127,7 +122,6 @@ steps:
       - vllm/model_executor/model_loader/
 
 - label: Multi-Modal Models (Extended Generation 1)
-  device: h200_35gb
   key: multi-modal-models-extended-generation-1
   optional: true
   source_file_dependencies:
@@ -135,13 +129,12 @@ steps:
   - tests/models/multimodal/generation
   - tests/models/multimodal/test_mapping.py
   commands:
+    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
     - pytest -v -s models/multimodal/generation -m 'not core_model' --ignore models/multimodal/generation/test_common.py
     - pytest -v -s models/multimodal/test_mapping.py
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
-      timeout_in_minutes: 90
       depends_on:
       - image-build-amd
 
@@ -153,6 +146,7 @@ steps:
   - vllm/
   - tests/models/multimodal/generation
   commands:
+    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
     - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=0) and not core_model'
 
 - label: Multi-Modal Models (Extended Generation 3)
@@ -163,6 +157,7 @@ steps:
   - vllm/
   - tests/models/multimodal/generation
   commands:
+    - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
     - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=1) and not core_model'
 
 - label: Multi-Modal Models (Extended Pooling)
@@ -176,8 +171,7 @@ steps:
     - pytest -v -s models/multimodal/pooling -m 'not core_model'
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
       timeout_in_minutes: 60
       depends_on:
       - image-build-amd

.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: 35
+  timeout_in_minutes: 60
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
@@ -37,11 +37,6 @@ 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: 150
+  timeout_in_minutes: 10
   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: 45
+  timeout_in_minutes: 20
   source_file_dependencies:
     - vllm/__init__.py
     - vllm/_aiter_ops.py
@@ -107,11 +107,16 @@ steps:
     - tests/compile/passes
   commands:
   - pytest -s -v compile/passes --ignore compile/passes/distributed
+  mirror:
+    amd:
+      device: mi300_1
+      timeout_in_minutes: 180
+      depends_on:
+      - image-build-amd
 
 - label: PyTorch Fullgraph Smoke Test
-  device: h200_35gb
   key: pytorch-fullgraph-smoke-test
-  timeout_in_minutes: 90
+  timeout_in_minutes: 35
   source_file_dependencies:
   - vllm/__init__.py
   - vllm/_aiter_ops.py
@@ -143,46 +148,11 @@ steps:
   # as it is a heavy test that is covered in other steps.
   # Use `find` to launch multiple instances of pytest so that
   # they do not suffer from https://github.com/vllm-project/vllm/issues/28965
-  - "find compile/fullgraph/ -name 'test_*.py' -not -name 'test_full_cudagraph.py' -not -name 'test_full_graph.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
+  - "find compile/fullgraph/ -name 'test_*.py' -not -name 'test_full_graph.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
-
-# Hopper-only DeepSeek-V2-Lite cases in this file require two 29.3-GiB model
-# instances and cannot fit a 35GB MIG slice. L4 retains the original coverage:
-# those SM90 cases skip while the architecture-compatible cases still run.
-- label: PyTorch Fullgraph CUDAGraph (L4 Compatibility)
-  key: pytorch-fullgraph-cudagraph-l4-compatibility
-  timeout_in_minutes: 60
-  source_file_dependencies:
-  - vllm/__init__.py
-  - vllm/_aiter_ops.py
-  - vllm/_custom_ops.py
-  - vllm/compilation/
-  - vllm/config/
-  - vllm/distributed/
-  - vllm/engine/
-  - vllm/env_override.py
-  - vllm/envs.py
-  - vllm/forward_context.py
-  - vllm/inputs/
-  - vllm/ir/
-  - vllm/kernels/
-  - vllm/logger.py
-  - vllm/model_executor/
-  - vllm/multimodal/
-  - vllm/platforms/
-  - vllm/plugins/
-  - vllm/sampling_params.py
-  - vllm/sequence.py
-  - vllm/transformers_utils/
-  - vllm/triton_utils/
-  - vllm/utils/
-  - vllm/v1/
-  - tests/compile
-  commands:
-  - pytest -s -v compile/fullgraph/test_full_cudagraph.py
 
 - label: PyTorch Fullgraph
   key: pytorch-fullgraph
-  timeout_in_minutes: 40
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
   - vllm/__init__.py
@@ -227,9 +197,7 @@ steps:
   - bash standalone_tests/pytorch_nightly_dependency.sh
   mirror:
     amd:
-      dind: false
       device: mi300_1
-      timeout_in_minutes: 30
       depends_on:
       - image-build-amd
       source_file_dependencies:

.buildkite/test_areas/quantization.yaml

@@ -3,11 +3,8 @@ depends_on:
   - image-build
 steps:
 - label: Quantization
-  device: h200_35gb
   key: quantization
-  timeout_in_minutes: 75
+  timeout_in_minutes: 90
-  env:
-    VLLM_USE_V2_MODEL_RUNNER: "0"
   source_file_dependencies:
   - csrc/
   - vllm/model_executor/layers/quantization
@@ -22,14 +19,11 @@ steps:
   # TODO(jerryzh168): resolve the above comment
   - uv pip install --system torchao==0.17.0 --index-url https://download.pytorch.org/whl/cu130
   - uv pip install --system conch-triton-kernels
-  # The SM90-only checkpoint currently contains a removed weight_chan_scale
+  - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py
-  # parameter. It was not exercised by the previous L4 job.
-  - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py -k 'not test_compressed_tensors_w4a8_fp8'
 
 - label: Quantized Fusions
-  device: h200_35gb
   key: quantized-fusions
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   source_file_dependencies:
   - tests/fusion
   - vllm/model_executor/layers/fusion
@@ -41,7 +35,7 @@ steps:
 
 - label: Quantized MoE Test (B200)
   key: quantized-moe-test-b200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 60
   working_dir: "/vllm-workspace/"
   device: b200-k8s
   source_file_dependencies:
@@ -58,11 +52,8 @@ steps:
     - pytest -s -v tests/quantization/test_blackwell_moe.py
 
 - label: Quantized Models Test
-  device: h200_35gb
   key: quantized-models-test
-  timeout_in_minutes: 65
+  timeout_in_minutes: 60
-  env:
-    VLLM_USE_V2_MODEL_RUNNER: "0"
   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: 30
+  timeout_in_minutes: 90
   device: h200_18gb
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -15,30 +15,28 @@ 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_return_token_ids.py
-  - tests/entrypoints/openai/test_uds.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 multiple"
+  # - pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not invalid"
 
   # - 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"
-  # test_comparison streams differently: Rust emits a separate first (prompt_token_ids) chunk and
+  # - pytest -v -s entrypoints/openai/test_return_token_ids.py
-  # finish chunk without logprobs, while the test reads `logprobs.tokens` on every chunk.
+  # - pytest -v -s entrypoints/openai/test_uds.py
-  - 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: 25
+  timeout_in_minutes: 60
   device: h200_18gb
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -47,27 +45,22 @@ 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
-  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
+  # - 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: 20
+  timeout_in_minutes: 30
   device: h200_18gb
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -81,8 +74,7 @@ steps:
   - pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
 
 - label: Rust Frontend Tool Use
-  device: h200_35gb
+  timeout_in_minutes: 60
-  timeout_in_minutes: 25
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - rust/
@@ -96,7 +88,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: 25
+  timeout_in_minutes: 30
   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: 20
+  timeout_in_minutes: 30
   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: 20
+  timeout_in_minutes: 30
   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: 40
+  timeout_in_minutes: 75
   source_file_dependencies:
   - vllm/model_executor/layers
   - vllm/sampling_metadata.py
@@ -20,17 +20,7 @@ steps:
   mirror:
     amd:
       device: mi250_1
-      timeout_in_minutes: 40
       depends_on:
       - image-build-amd
-      source_file_dependencies:
-      - vllm/model_executor/layers
-      - vllm/sampling_metadata.py
-      - vllm/v1/sample/
-      - vllm/entrypoints/generate/beam_search/
-      - tests/samplers
-      - tests/conftest.py
-      - vllm/_aiter_ops.py
-      - vllm/platforms/rocm.py
       commands:
       - pytest -v -s samplers

.buildkite/test_areas/spec_decode.yaml

@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Spec Decode Eagle
   key: spec-decode-eagle
-  timeout_in_minutes: 25
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -14,9 +14,8 @@ steps:
     - pytest -v -s v1/e2e/spec_decode -k "eagle_correctness"
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 45
-      timeout_in_minutes: 55
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -30,7 +29,7 @@ steps:
 
 - label: Spec Decode Eagle Nightly B200
   key: spec-decode-eagle-nightly-b200
-  timeout_in_minutes: 25
+  timeout_in_minutes: 30
   device: b200-k8s
   optional: true
   source_file_dependencies:
@@ -42,7 +41,7 @@ steps:
 
 - label: Spec Decode Speculators + MTP
   key: spec-decode-speculators-mtp
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -54,9 +53,8 @@ steps:
     - pytest -v -s v1/e2e/spec_decode -k "speculators or mtp_correctness"
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 65
-      timeout_in_minutes: 75
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -84,7 +82,7 @@ steps:
   
 - label: Spec Decode Ngram + Suffix
   key: spec-decode-ngram-suffix
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -94,9 +92,10 @@ steps:
     - pytest -v -s v1/e2e/spec_decode -k "ngram or suffix"
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 65
-      timeout_in_minutes: 35
+      # TODO(akaratza): Test after Torch >= 2.12 bump
+      soft_fail: true
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -110,7 +109,7 @@ steps:
 
 - label: Spec Decode Draft Model
   key: spec-decode-draft-model
-  timeout_in_minutes: 45
+  timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -120,9 +119,8 @@ steps:
     - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
   mirror:
     amd:
-      dind: false
+      device: mi325_1
-      device: mi300_1
+      timeout_in_minutes: 50
-      timeout_in_minutes: 55
       depends_on:
       - image-build-amd
       source_file_dependencies:
@@ -136,7 +134,7 @@ steps:
 
 - label: Spec Decode Draft Model Nightly B200
   key: spec-decode-draft-model-nightly-b200
-  timeout_in_minutes: 40
+  timeout_in_minutes: 30
   device: b200-k8s
   optional: true
   source_file_dependencies:
@@ -148,7 +146,7 @@ steps:
 
 - label: Speculators Correctness
   key: speculators-correctness
-  timeout_in_minutes: 30
+  timeout_in_minutes: 60
   device: h100
   optional: true
   num_devices: 1
@@ -161,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: 20
+  timeout_in_minutes: 30
   device: b200-k8s
   optional: true
   source_file_dependencies:
@@ -172,19 +170,3 @@ steps:
     - tests/v1/e2e/spec_decode/
   commands:
     - pytest -v -s v1/e2e/spec_decode -k "qwen3_5-hybrid"
-
-- label: Spec Decode DeepSeek MTP Parallel Load (B200)
-  key: spec-decode-deepseek-mtp-parallel-load-b200
-  timeout_in_minutes: 30
-  device: b200-k8s
-  optional: true
-  num_devices: 2
-  source_file_dependencies:
-    - vllm/v1/spec_decode/llm_base_proposer.py
-    - vllm/v1/spec_decode/eagle.py
-    - vllm/v1/worker/gpu/spec_decode/eagle/
-    - vllm/model_executor/models/deepseek_mtp.py
-    - vllm/model_executor/models/deepseek_v2.py
-    - tests/v1/e2e/spec_decode/test_mtp_parallel_load.py
-  commands:
-    - pytest -v -s v1/e2e/spec_decode/test_mtp_parallel_load.py

.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: 50
+  timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   optional: true
@@ -15,9 +15,7 @@ steps:
     - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models.txt
   mirror:
     amd:
-      dind: false
       device: mi300_2
-      timeout_in_minutes: 35
       depends_on:
       - image-build-amd
       commands:

.dockerignore

@@ -3,7 +3,6 @@
 dist
 vllm/*.so
 vllm/vllm-rs
-.git
 
 # Byte-compiled / optimized / DLL files
 __pycache__/

.github/CODEOWNERS

@@ -47,7 +47,6 @@
 
 # Rust Frontend
 /rust/ @BugenZhao @njhill
-/rust/src/bench @esmeetu
 /build_rust.sh @BugenZhao @njhill
 /rust-toolchain.toml @BugenZhao @njhill
 /.buildkite/test_areas/rust* @BugenZhao @njhill
@@ -120,7 +119,7 @@
 
 # Transformers modeling backend
 /vllm/model_executor/models/transformers @hmellor
-/tests/models/transformers @hmellor
+/tests/models/test_transformers.py @hmellor
 
 # Docs
 /docs/mkdocs @hmellor
@@ -173,7 +172,7 @@ mkdocs.yaml @hmellor
 # Kernels
 /vllm/v1/attention/ops/chunked_prefill_paged_decode.py @tdoublep
 /vllm/v1/attention/ops/triton_unified_attention.py @tdoublep
-/vllm/third_party/flash_linear_attention @ZJY0516 @vadiklyutiy
+/vllm/model_executor/layers/fla @ZJY0516 @vadiklyutiy
 
 # ROCm related: specify owner with write access to notify AMD folks for careful code review
 /vllm/**/*rocm* @tjtanaa @dllehr-amd

.github/mergify.yml

@@ -181,18 +181,6 @@ pull_request_rules:
       add:
         - performance
 
-- name: label-quantization
-  description: Automatically apply quantization label
-  conditions:
-    - label != stale
-    - or:
-      - files~=^vllm/model_executor/layers/quantization/
-      - title~=(?i)quant
-  actions:
-    label:
-      add:
-        - quantization
-
 - name: label-qwen
   description: Automatically apply qwen label
   conditions:

.github/workflows/issue_autolabel.yml

@@ -130,47 +130,6 @@ jobs:
                   },
                 ],
               },
-              quantization: {
-                keywords: [
-                  {
-                    term: "quantization",
-                    searchIn: "both"
-                  },
-                  {
-                    term: "quantized",
-                    searchIn: "both"
-                  },
-                ],
-              },
-              "intel-gpu": {
-                // Keyword search - matches whole words only (with word boundaries)
-                keywords: [
-                  {
-                    term: "B50",
-                    searchIn: "both"
-                  },
-                  {
-                    term: "B60",
-                    searchIn: "both"
-                  },
-                  {
-                    term: "B70",
-                    searchIn: "both"
-                  },
-                  {
-                    term: "intel gpu",
-                    searchIn: "both"
-                  },
-                  {
-                    term: "Arc GPU",
-                    searchIn: "both"
-                  },
-                  {
-                    term: "BMG",
-                    searchIn: "both"
-                  },
-                ],
-              },
               // Add more label configurations here as needed
               // example: {
               //   keywords: [...],
@@ -364,7 +323,7 @@ jobs:
             // {users} will be replaced with @mentions
             const ccConfig = {
               rocm: {
-                users: ['hongxiayang', 'tjtanaa', 'vllmellm', 'giuseppegrossi'],
+                users: ['hongxiayang', 'tjtanaa', 'vllmellm'],
                 message: 'CC {users} for ROCm-related issue',
               },
               mistral: {
@@ -532,4 +491,4 @@ jobs:
               issue_number: context.issue.number,
               body: message,
             });
-            core.notice(`Requested missing ROCm info from @${author}: ${missing.map(m => m.name).join(', ')}`);
+            core.notice(`Requested missing ROCm info from @${author}: ${missing.map(m => m.name).join(', ')}`);

.github/workflows/pre-commit.yml

@@ -28,8 +28,7 @@ jobs:
             pull_number: context.payload.pull_request.number,
           });
 
-          const readyLabels = ['ready', 'ready-run-all-tests'];
+          const hasReadyLabel = pr.labels.some(l => l.name === 'ready');
-          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({
@@ -41,7 +40,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', '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}).`);
+            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}).`);
           }
 
   pre-commit:
@@ -50,7 +49,7 @@ jobs:
     runs-on: [self-hosted, linux, x64, vllm-runners]
     steps:
     - uses: actions/checkout@9c091bb21b7c1c1d1991bb908d89e4e9dddfe3e0 # v7.0.0
-    - uses: actions/setup-python@83679a892e2d95755f2dac6acb0bfd1e9ac5d548 # v6.1.0
+    - uses: actions/setup-python@ece7cb06caefa5fff74198d8649806c4678c61a1 # v6.3.0
       with:
         python-version: "3.12"
     # Provide shellcheck on PATH so tools/pre_commit/shellcheck.sh skips its

.gitignore

@@ -18,9 +18,6 @@ vllm/third_party/deep_gemm/
 # fmha_sm100 vendored package built from source
 vllm/third_party/fmha_sm100/
 
-# tml-fa4 vendored package built from source
-vllm/third_party/tml_fa4/
-
 # triton jit
 .triton
 

.pre-commit-config.yaml

@@ -4,7 +4,7 @@ default_install_hook_types:
 default_stages:
   - pre-commit # Run locally
   - manual # Run in CI
-exclude: 'vllm/third_party/.*|vllm/models/kimi_k3/nvidia/ops/third_party/.*|vllm/models/kimi_k3/amd/ops/third_party/.*'
+exclude: 'vllm/third_party/.*'
 repos:
 - repo: https://github.com/astral-sh/ruff-pre-commit
   rev: v0.14.0
@@ -30,7 +30,7 @@ repos:
   - id: markdownlint-cli2
     language_version: lts
     args: [--fix]
-    exclude: (^|/)CLAUDE\.md$
+    exclude: ^CLAUDE\.md$
 - repo: https://github.com/rhysd/actionlint
   rev: v1.7.7
   hooks:
@@ -210,7 +210,7 @@ repos:
     name: Check SPDX headers
     entry: python tools/pre_commit/check_spdx_header.py
     language: python
-    types_or: [python, rust, proto]
+    types: [python]
   - id: check-root-lazy-imports
     name: Check root lazy imports
     entry: python tools/pre_commit/check_init_lazy_imports.py

.yapfignore

@@ -0,0 +1,2 @@
+collect_env.py
+vllm/model_executor/layers/fla/ops/*.py

AGENTS.md

@@ -29,7 +29,6 @@ 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
@@ -67,38 +66,23 @@ VLLM_USE_PRECOMPILED=1 uv pip install -e . --torch-backend=auto
 uv pip install -e . --torch-backend=auto
 ```
 
-### Tests
+### Running tests
 
 > Requires [Environment setup](#environment-setup) and [Installing dependencies](#installing-dependencies).
 
 ```bash
-# Install test dependencies (use cuda.in on non-x86_64):
+# Install test dependencies.
-uv pip install -r requirements/test/cuda.in
+# 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
 
-# Run a specific test file:
+# Run a specific test file (use .venv/bin/python directly;
+# `source activate` does not persist in non-interactive shells):
 .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).
@@ -123,18 +107,23 @@ Use [Google-style docstrings](https://google.github.io/styleguide/pyguide.html#3
 
 ### Coding style guidelines
 
-- Match existing code style
+Follow these rules for all code changes in this repository:
-- Minimize use of comments. Eliminate comments which are redundant, preferring legible and self-documenting code. When used, keep docstrings and comments brief and direct.
+
+- Try to match existing code style.
+- Code should be self-documenting and self-explanatory.
+- Keep comments and docstrings minimal and concise.
 - 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:`):
+Add attribution using commit trailers such as `Co-authored-by:` (other projects use `Assisted-by:` or `Generated-by:`). For example:
 
 ```text
 Your commit message here
 
-Co-authored-by: Agent Name Here
+Co-authored-by: GitHub Copilot
+Co-authored-by: Claude
+Co-authored-by: gemini-code-assist
 Signed-off-by: Your Name <your.email@example.com>
 ```
 

CMakeLists.txt

@@ -68,17 +68,8 @@ endif()
 # requirements.txt files and should be kept consistent.  The ROCm torch
 # versions are derived from docker/Dockerfile.rocm
 #
-set(TORCH_SUPPORTED_VERSION_CUDA "2.13.0")
+set(TORCH_SUPPORTED_VERSION_CUDA "2.11.0")
-set(TORCH_SUPPORTED_VERSION_ROCM "2.13.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
@@ -114,11 +105,6 @@ find_package(Torch REQUIRED)
 # Supported NVIDIA architectures.
 # This check must happen after find_package(Torch) because that's when CMAKE_CUDA_COMPILER_VERSION gets defined
 if(DEFINED CMAKE_CUDA_COMPILER_VERSION AND
-   CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.4)
-  # Rubin (10.7) can run SM100 family code, but CUDA 13.4 also supports
-  # targeting it directly.
-  set(CUDA_SUPPORTED_ARCHS "7.5;8.0;8.6;8.7;8.9;9.0;10.0;10.7;11.0;12.0")
-elseif(DEFINED CMAKE_CUDA_COMPILER_VERSION AND
    CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.0)
    # starting from CUDA 12.9 and Blackwell (10.0), we use family-specific targets (10.0f, 12.0f, etc)
    # to support the whole generation without specifying all sub-architectures
@@ -189,7 +175,7 @@ endif()
 if (NOT HIP_FOUND AND NOT PYTORCH_FOUND_HIP AND CUDA_FOUND)
   set(VLLM_GPU_LANG "CUDA")
 
-  if (NOT TORCH_NIGHTLY_BUILD AND NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
+  if (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()
@@ -202,7 +188,7 @@ elseif(HIP_FOUND OR PYTORCH_FOUND_HIP)
   enable_language(HIP)
 
   # ROCm 5.X and 6.X
-  if (NOT TORCH_NIGHTLY_BUILD AND ROCM_VERSION_DEV_MAJOR GREATER_EQUAL 5 AND
+  if (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.")
@@ -364,7 +350,9 @@ endif()
 if(VLLM_GPU_LANG STREQUAL "HIP")
   set(VLLM_EXT_SRC
     "csrc/torch_bindings.cpp"
-    "csrc/custom_quickreduce.cu")
+    "csrc/custom_quickreduce.cu"
+    "csrc/cuda_view.cu"
+    "csrc/libtorch_stable/cuda_utils_kernels.cu")
 
   message(STATUS "Enabling C extension.")
   define_extension_target(
@@ -392,10 +380,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   #
   set(VLLM_STABLE_EXT_SRC
     "csrc/libtorch_stable/torch_bindings.cpp"
-    "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"
@@ -414,13 +399,13 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     "csrc/libtorch_stable/sampler.cu"
     "csrc/libtorch_stable/topk.cu"
     "csrc/libtorch_stable/mamba/selective_scan_fwd.cu"
+    "csrc/libtorch_stable/attention/paged_attention_v1.cu"
+    "csrc/libtorch_stable/attention/paged_attention_v2.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_gather_reduce_scatter.cu"
-    "csrc/libtorch_stable/custom_all_gather_reduce_scatter_ops.cpp"
     "csrc/libtorch_stable/custom_all_reduce.cu"
-    "csrc/libtorch_stable/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu"
+    "csrc/libtorch_stable/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu")
-    "csrc/libtorch_stable/fused_kimi_k3_mla_key_concat_kv_cache_kernel.cu")
 
   if(VLLM_GPU_LANG STREQUAL "CUDA" AND
      DEFINED CMAKE_CUDA_COMPILER_VERSION AND
@@ -428,7 +413,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
 
     if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
       cuda_archs_loose_intersection(COOPERATIVE_TOPK_ARCHS
-        "9.0a;10.0f;10.1f;10.3f;10.7f;11.0f;12.0f;12.1f" "${CUDA_ARCHS}")
+        "9.0a;10.0f;10.1f;10.3f;11.0f;12.0f;12.1f" "${CUDA_ARCHS}")
     else()
       cuda_archs_loose_intersection(COOPERATIVE_TOPK_ARCHS
         "9.0a;10.0a;10.1a;10.3a;12.0a;12.1a" "${CUDA_ARCHS}")
@@ -474,6 +459,8 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     FetchContent_MakeAvailable(cutlass)
 
     list(APPEND VLLM_STABLE_EXT_SRC
+      "csrc/libtorch_stable/cuda_view.cu"
+      "csrc/libtorch_stable/cuda_utils_kernels.cu"
       "csrc/libtorch_stable/cutlass_extensions/common.cpp"
       "csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_entry.cu"
       "csrc/libtorch_stable/quantization/fp4/nvfp4_quant_entry.cu"
@@ -703,7 +690,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
 
   # DeepSeek V3 fused A GEMM kernel (requires SM 9.0+, Hopper and later)
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0f;10.7f;11.0f;12.0f" "${CUDA_ARCHS}")
+    cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0f;11.0f;12.0f" "${CUDA_ARCHS}")
   else()
     cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0a;10.1a;10.3a;12.0a;12.1a" "${CUDA_ARCHS}")
   endif()
@@ -823,7 +810,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   # The cutlass_scaled_mm kernels for Blackwell SM100 (c3x, i.e. CUTLASS 3.x)
   # require CUDA 12.8 or later
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(SCALED_MM_ARCHS "10.0f;10.7f;11.0f" "${CUDA_ARCHS}")
+    cuda_archs_loose_intersection(SCALED_MM_ARCHS "10.0f;11.0f" "${CUDA_ARCHS}")
   else()
     cuda_archs_loose_intersection(SCALED_MM_ARCHS "10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
   endif()
@@ -907,7 +894,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   endif()
 
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(SCALED_MM_ARCHS "10.0f;10.7f;11.0f" "${CUDA_ARCHS}")
+    cuda_archs_loose_intersection(SCALED_MM_ARCHS "10.0f;11.0f" "${CUDA_ARCHS}")
   else()
     cuda_archs_loose_intersection(SCALED_MM_ARCHS "10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
   endif()
@@ -932,7 +919,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
 
   # moe_data.cu is used by all CUTLASS MoE kernels.
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(CUTLASS_MOE_DATA_ARCHS "9.0a;10.0f;10.7f;11.0f;12.0f" "${CUDA_ARCHS}")
+    cuda_archs_loose_intersection(CUTLASS_MOE_DATA_ARCHS "9.0a;10.0f;11.0f;12.0f" "${CUDA_ARCHS}")
   else()
     cuda_archs_loose_intersection(CUTLASS_MOE_DATA_ARCHS "9.0a;10.0a;10.1a;10.3a;12.0a;12.1a" "${CUDA_ARCHS}")
   endif()
@@ -989,7 +976,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   # SM10x/11x FP4 kernels. MXFP4 experts quantization is currently compiled
   # only in this block; SM12x has separate NVFP4 matmul/MoE kernels above.
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(FP4_SM100_ARCHS "10.0f;10.7f;11.0f" "${CUDA_ARCHS}")
+    cuda_archs_loose_intersection(FP4_SM100_ARCHS "10.0f;11.0f" "${CUDA_ARCHS}")
   else()
     cuda_archs_loose_intersection(FP4_SM100_ARCHS "10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
   endif()
@@ -1055,7 +1042,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
   # Runtime dispatch is gated in
   # vllm/v1/attention/backends/mla/cutlass_mla.py.
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(MLA_ARCHS "10.0f;10.7f;11.0f" "${CUDA_ARCHS}")
+    cuda_archs_loose_intersection(MLA_ARCHS "10.0f;11.0f" "${CUDA_ARCHS}")
   else()
     cuda_archs_loose_intersection(MLA_ARCHS "10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
   endif()
@@ -1077,41 +1064,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     set(MLA_ARCHS)
   endif()
 
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(FUSED_KDA_DECODE_ARCHS
-      "9.0a;10.0f;12.0f" "${CUDA_ARCHS}")
-  endif()
-  if(FUSED_KDA_DECODE_ARCHS)
-    set(FUSED_KDA_DECODE_SRC
-      "csrc/libtorch_stable/kimi_k3/fused_kda_decode_kernel.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${FUSED_KDA_DECODE_SRC}"
-      CUDA_ARCHS "${FUSED_KDA_DECODE_ARCHS}")
-    set_property(SOURCE ${FUSED_KDA_DECODE_SRC} APPEND PROPERTY
-      COMPILE_OPTIONS "$<$<COMPILE_LANGUAGE:CUDA>:--use_fast_math>")
-    list(APPEND VLLM_STABLE_EXT_SRC "${FUSED_KDA_DECODE_SRC}")
-    message(STATUS
-      "Building fused KDA decode for archs: ${FUSED_KDA_DECODE_ARCHS}")
-  endif()
-
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(KIMI_K3_ATTN_RES_ARCHS
-      "10.0f" "${CUDA_ARCHS}")
-  endif()
-  if(KIMI_K3_ATTN_RES_ARCHS)
-    set(KIMI_K3_ATTN_RES_SRC
-      "csrc/libtorch_stable/kimi_k3/attn_res_kernel.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${KIMI_K3_ATTN_RES_SRC}"
-      CUDA_ARCHS "${KIMI_K3_ATTN_RES_ARCHS}")
-    set_property(SOURCE ${KIMI_K3_ATTN_RES_SRC} APPEND PROPERTY
-      COMPILE_OPTIONS
-      "$<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr;--expt-extended-lambda;--use_fast_math>")
-    list(APPEND VLLM_STABLE_EXT_SRC "${KIMI_K3_ATTN_RES_SRC}")
-    message(STATUS
-      "Building Kimi K3 AttnRes for archs: ${KIMI_K3_ATTN_RES_ARCHS}")
-  endif()
-
   # Hadacore kernels
   cuda_archs_loose_intersection(HADACORE_ARCHS "8.0+PTX;9.0+PTX" "${CUDA_ARCHS}")
   if(HADACORE_ARCHS)
@@ -1139,32 +1091,29 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
     USE_SABI 3
     WITH_SOABI)
 
-  # Set TORCH_TARGET_VERSION for stable ABI compatibility.
-  # This ensures we only use C-shim APIs available in PyTorch 2.11.
-  # _C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
-  # which is currently set to 2.11.
-  target_compile_definitions(_C_stable_libtorch PRIVATE
-    TORCH_TARGET_VERSION=0x020B000000000000ULL)
-
   # Needed to use cuda/hip APIs from C-shim
   if(VLLM_GPU_LANG STREQUAL "CUDA")
+    # Set TORCH_TARGET_VERSION for stable ABI compatibility.
+    # This ensures we only use C-shim APIs available in PyTorch 2.11.
+    # _C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
+    # which is currently set to 2.11.
+    target_compile_definitions(_C_stable_libtorch PRIVATE
+      TORCH_TARGET_VERSION=0x020B000000000000ULL)
     target_compile_definitions(_C_stable_libtorch PRIVATE USE_CUDA)
     if(COOPERATIVE_TOPK_ARCHS)
       target_compile_definitions(_C_stable_libtorch PRIVATE
         VLLM_ENABLE_COOPERATIVE_TOPK=1)
     endif()
-    if(FUSED_KDA_DECODE_ARCHS)
-      target_compile_definitions(_C_stable_libtorch PRIVATE
-        VLLM_ENABLE_FUSED_KDA_DECODE=1)
-    endif()
-    if(KIMI_K3_ATTN_RES_ARCHS)
-      target_compile_definitions(_C_stable_libtorch PRIVATE
-        VLLM_ENABLE_KIMI_K3_ATTN_RES=1)
-    endif()
     # Needed by CUTLASS kernels
     target_compile_definitions(_C_stable_libtorch PRIVATE
       CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1)
   elseif(VLLM_GPU_LANG STREQUAL "HIP")
+    # Set TORCH_TARGET_VERSION for stable ABI compatibility.
+    # This ensures we only use C-shim APIs available in PyTorch 2.10.
+    # _C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
+    # which is currently set to 2.10.
+    target_compile_definitions(_C_stable_libtorch PRIVATE
+      TORCH_TARGET_VERSION=0x020A000000000000ULL)
     target_compile_definitions(_C_stable_libtorch PRIVATE USE_ROCM)
   endif()
 
@@ -1372,20 +1321,25 @@ define_extension_target(
   USE_SABI 3
   WITH_SOABI)
 
-# Set TORCH_TARGET_VERSION for stable ABI compatibility.
-# This ensures we only use C-shim APIs available in PyTorch 2.11.
-# _moe_C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
-# which is currently set to 2.11.
-target_compile_definitions(_moe_C_stable_libtorch PRIVATE
-  TORCH_TARGET_VERSION=0x020B000000000000ULL)
-
 # Needed to use cuda/hip APIs from C-shim
 if(VLLM_GPU_LANG STREQUAL "CUDA")
+  # Set TORCH_TARGET_VERSION for stable ABI compatibility.
+  # This ensures we only use C-shim APIs available in PyTorch 2.11.
+  # _moe_C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
+  # which is currently set to 2.11.
+  target_compile_definitions(_moe_C_stable_libtorch PRIVATE
+    TORCH_TARGET_VERSION=0x020B000000000000ULL)
   target_compile_definitions(_moe_C_stable_libtorch PRIVATE USE_CUDA)
   # Needed by CUTLASS kernels
   target_compile_definitions(_moe_C_stable_libtorch PRIVATE
     CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1)
 elseif(VLLM_GPU_LANG STREQUAL "HIP")
+  # Set TORCH_TARGET_VERSION for stable ABI compatibility.
+  # This ensures we only use C-shim APIs available in PyTorch 2.10.
+  # _moe_C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
+  # which is currently set to 2.10.
+  target_compile_definitions(_moe_C_stable_libtorch PRIVATE
+    TORCH_TARGET_VERSION=0x020A000000000000ULL)
   target_compile_definitions(_moe_C_stable_libtorch PRIVATE USE_ROCM)
 endif()
 
@@ -1415,7 +1369,6 @@ if(VLLM_GPU_LANG STREQUAL "HIP")
   set(VLLM_ROCM_EXT_SRC
     "csrc/rocm/torch_bindings.cpp"
     "csrc/rocm/skinny_gemms.cu"
-    "csrc/rocm/skinny_gemms_int4.cu"
     "csrc/rocm/attention.cu")
 
   set(VLLM_ROCM_HAS_GFX1100 OFF)
@@ -1458,9 +1411,7 @@ if (VLLM_GPU_LANG STREQUAL "CUDA")
     include(cmake/external_projects/deepgemm.cmake)
     include(cmake/external_projects/fmha_sm100.cmake)
     include(cmake/external_projects/flashmla.cmake)
-    include(cmake/external_projects/flashkda.cmake)
     include(cmake/external_projects/qutlass.cmake)
-    include(cmake/external_projects/tml_fa4.cmake)
 
     # vllm-flash-attn should be last as it overwrites some CMake functions
     include(cmake/external_projects/vllm_flash_attn.cmake)

README.md

@@ -48,7 +48,7 @@ vLLM is flexible and easy to use with:
 - Tool calling and reasoning parsers
 - OpenAI-compatible API server, plus Anthropic Messages API and gRPC support
 - Efficient multi-LoRA support for dense and MoE layers
-- Support for NVIDIA GPUs, AMD GPUs, Intel GPUs, and x86/ARM/PowerPC CPUs. Additionally, diverse hardware plugins such as Google TPUs, Intel Gaudi, IBM Spyre, Huawei Ascend, Rebellions NPU, Apple Silicon, MetaX GPU, and more.
+- Support for NVIDIA GPUs, AMD GPUs, and x86/ARM/PowerPC CPUs. Additionally, diverse hardware plugins such as Google TPUs, Intel Gaudi, IBM Spyre, Huawei Ascend, Rebellions NPU, Apple Silicon, MetaX GPU, and more.
 
 vLLM seamlessly supports 200+ model architectures on Hugging Face, including:
 

benchmarks/attention_benchmarks/benchmark.py

@@ -75,11 +75,7 @@ def run_mla_benchmark(config: BenchmarkConfig, **kwargs) -> BenchmarkResult:
     from mla_runner import run_mla_benchmark as run_mla
 
     return run_mla(
-        config.backend,
+        config.backend, config, prefill_backend=config.prefill_backend, **kwargs
-        config,
-        prefill_backend=config.prefill_backend,
-        sparse_mla_force_mqa=config.sparse_mla_force_mqa,
-        **kwargs,
     )
 
 
@@ -596,30 +592,6 @@ def main():
         default="profile",
         help="Output file name for ncu profile (default: 'profile').",
     )
-    parser.add_argument(
-        "--torch-profile",
-        action="store_true",
-        default=False,
-        help="Collect a PyTorch profiler Chrome trace for each benchmark run.",
-    )
-    parser.add_argument(
-        "--torch-profile-dir",
-        default=None,
-        help="Directory for PyTorch profiler traces.",
-    )
-    parser.add_argument(
-        "--torch-profile-iters",
-        type=int,
-        default=3,
-        help="Number of forward passes to record per PyTorch profiler trace.",
-    )
-    parser.add_argument(
-        "--sparse-mla-mha-variants",
-        nargs="+",
-        default=None,
-        choices=["dense_mha", "mqa"],
-        help="Sparse MLA variants to run in mha_vs_mqa mode. Defaults to both.",
-    )
 
     # Parameter sweep (use YAML config for advanced sweeps)
     parser.add_argument(
@@ -669,7 +641,6 @@ def main():
 
         # Prefill backends (e.g., ["fa3", "fa4"])
         args.prefill_backends = yaml_config.get("prefill_backends", None)
-        args.prefill_backend = yaml_config.get("prefill_backend", None)
 
         # FP8 output benchmark knobs; CLI wins.
         if args.fp8_output_scale is None:
@@ -712,9 +683,6 @@ def main():
             args.num_q_heads = model.get("num_q_heads", args.num_q_heads)
             args.num_kv_heads = model.get("num_kv_heads", args.num_kv_heads)
             args.block_size = model.get("block_size", args.block_size)
-            args.max_model_len = model.get(
-                "max_model_len", getattr(args, "max_model_len", None)
-            )
             # MLA-specific dimensions
             args.kv_lora_rank = model.get("kv_lora_rank", args.kv_lora_rank)
             args.qk_nope_head_dim = model.get("qk_nope_head_dim", args.qk_nope_head_dim)
@@ -733,21 +701,6 @@ def main():
             args.cuda_graphs = yaml_config["cuda_graphs"]
         if "ncu_profile" in yaml_config:
             args.ncu_profile = yaml_config["ncu_profile"]
-        if "torch_profile" in yaml_config:
-            args.torch_profile = yaml_config["torch_profile"]
-        if "torch_profile_dir" in yaml_config:
-            args.torch_profile_dir = yaml_config["torch_profile_dir"]
-        if "torch_profile_iters" in yaml_config:
-            args.torch_profile_iters = yaml_config["torch_profile_iters"]
-        args.sparse_mla_topk_pattern = yaml_config.get(
-            "sparse_mla_topk_pattern", "random"
-        )
-        args.sparse_mla_dense_mha_max_seq_len = yaml_config.get(
-            "sparse_mla_dense_mha_max_seq_len", None
-        )
-        args.sparse_mla_mha_variants = yaml_config.get(
-            "sparse_mla_mha_variants", args.sparse_mla_mha_variants
-        )
 
         # Parameter sweep configuration
         if "parameter_sweep" in yaml_config:
@@ -889,6 +842,8 @@ def main():
                         num_kv_heads=args.num_kv_heads,
                         block_size=args.block_size,
                         device=args.device,
+                        repeats=args.repeats,
+                        warmup_iters=args.warmup_iters,
                         profile_memory=args.profile_memory,
                         kv_cache_dtype=args.kv_cache_dtype,
                         use_cuda_graphs=args.cuda_graphs,
@@ -1108,133 +1063,6 @@ def main():
                     f"\n  [yellow]Prefill always faster for batch_size={bs}[/]"
                 )
 
-    # Handle MHA vs MQA comparison mode for sparse MLA
-    elif hasattr(args, "mode") and args.mode == "mha_vs_mqa":
-        console.print("[yellow]Mode: MHA vs MQA comparison for sparse MLA[/]")
-
-        sparse_mla_topk_pattern = getattr(args, "sparse_mla_topk_pattern", "random")
-        dense_mha_max_seq_len = getattr(args, "sparse_mla_dense_mha_max_seq_len", None)
-        prefill_backend = getattr(args, "prefill_backend", None)
-        if prefill_backend:
-            console.print(f"Prefill backend: {prefill_backend}")
-        available_variants = [
-            ("dense_mha", False, "dense"),
-            ("mqa", True, "auto"),
-        ]
-        requested_variants = getattr(args, "sparse_mla_mha_variants", None)
-        if requested_variants is not None:
-            valid_variants = {label for label, _, _ in available_variants}
-            invalid_variants = sorted(set(requested_variants) - valid_variants)
-            if invalid_variants:
-                raise ValueError(
-                    "Invalid sparse_mla_mha_variants entries: "
-                    f"{invalid_variants}. Valid variants are: "
-                    f"{sorted(valid_variants)}"
-                )
-            requested_variant_set = set(requested_variants)
-            variants = [
-                variant
-                for variant in available_variants
-                if variant[0] in requested_variant_set
-            ]
-        else:
-            variants = available_variants
-        formatter = ResultsFormatter(console)
-        total = 0
-        for spec in args.batch_specs:
-            q_len = max(request.q_len for request in parse_batch_spec(spec))
-            for variant_label, _, _ in variants:
-                if (
-                    variant_label == "dense_mha"
-                    and dense_mha_max_seq_len is not None
-                    and q_len > dense_mha_max_seq_len
-                ):
-                    continue
-                total += len(backends)
-
-        with tqdm(total=total, desc="Benchmarking") as pbar:
-            for spec in args.batch_specs:
-                q_len = max(request.q_len for request in parse_batch_spec(spec))
-                for backend in backends:
-                    for variant_label, force_mqa, mha_mode in variants:
-                        if (
-                            variant_label == "dense_mha"
-                            and dense_mha_max_seq_len is not None
-                            and q_len > dense_mha_max_seq_len
-                        ):
-                            continue
-                        config = BenchmarkConfig(
-                            backend=f"{backend}_{variant_label}",
-                            batch_spec=spec,
-                            num_layers=args.num_layers,
-                            head_dim=args.head_dim,
-                            num_q_heads=args.num_q_heads,
-                            num_kv_heads=args.num_kv_heads,
-                            block_size=args.block_size,
-                            device=args.device,
-                            max_model_len=getattr(args, "max_model_len", None),
-                            kv_cache_dtype=args.kv_cache_dtype,
-                            profile_memory=args.profile_memory,
-                            use_cuda_graphs=args.cuda_graphs,
-                            ncu_profile=args.ncu_profile,
-                            torch_profile=args.torch_profile,
-                            torch_profile_dir=args.torch_profile_dir,
-                            torch_profile_iters=args.torch_profile_iters,
-                            warmup_ms=args.warmup_ms,
-                            kv_lora_rank=getattr(args, "kv_lora_rank", None),
-                            qk_nope_head_dim=getattr(args, "qk_nope_head_dim", None),
-                            qk_rope_head_dim=getattr(args, "qk_rope_head_dim", None),
-                            v_head_dim=getattr(args, "v_head_dim", None),
-                            sparse_mla_force_mqa=force_mqa,
-                            sparse_mla_mha_mode=mha_mode,
-                            sparse_mla_dense_mha_max_seq_len=dense_mha_max_seq_len,
-                            sparse_mla_topk_pattern=sparse_mla_topk_pattern,
-                            prefill_backend=prefill_backend,
-                        )
-
-                        # run_mla_benchmark needs the real backend name
-                        from mla_runner import run_mla_benchmark as run_mla
-
-                        run_label = f"{backend}_{variant_label} {spec}"
-                        pbar.set_postfix_str(run_label)
-
-                        try:
-                            result = run_mla(
-                                backend,
-                                config,
-                                prefill_backend=prefill_backend,
-                                sparse_mla_force_mqa=force_mqa,
-                            )
-                        except Exception as e:
-                            result = BenchmarkResult(
-                                config=config,
-                                mean_time=float("inf"),
-                                median_time=float("inf"),
-                                std_time=0,
-                                min_time=float("inf"),
-                                max_time=float("inf"),
-                                error=str(e),
-                            )
-
-                        all_results.append(result)
-                        if args.output_csv:
-                            formatter.save_csv(all_results, args.output_csv)
-                        if args.output_json:
-                            formatter.save_json(all_results, args.output_json)
-
-                        if not result.success:
-                            console.print(
-                                f"[red]Error {backend}_{variant_label} "
-                                f"{spec}: {result.error}[/]"
-                            )
-
-                        pbar.update(1)
-
-        # Display results with variant labels as separate "backends"
-        console.print("\n[bold green]MHA vs MQA Results:[/]")
-        variant_backends = [f"{b}_{v}" for b in backends for v, _, _ in variants]
-        formatter.print_table(all_results, variant_backends)
-
     # Handle model parameter sweep mode
     elif hasattr(args, "model_parameter_sweep") and args.model_parameter_sweep:
         # Model parameter sweep

benchmarks/attention_benchmarks/common.py

@@ -4,10 +4,8 @@
 """Common utilities for attention benchmarking."""
 
 import csv
-import gc
 import json
 import math
-from collections.abc import Sequence
 from dataclasses import asdict, dataclass
 from pathlib import Path
 from typing import Any
@@ -46,13 +44,10 @@ def run_do_bench(
     kwargs: dict[str, Any] = {"return_mode": "all"}
     if use_cuda_graphs:
         result = triton.testing.do_bench_cudagraph(benchmark_fn, **kwargs)
-        gc.collect()
-        torch.accelerator.empty_cache()
     else:
         if warmup_ms is not None:
             kwargs["warmup"] = warmup_ms
         result = triton.testing.do_bench(benchmark_fn, **kwargs)
-    torch.accelerator.synchronize()
     return result
 
 
@@ -96,6 +91,42 @@ except ImportError:
     AttentionLayerBase = object  # Fallback
 
 
+class MockKVBProj:
+    """Mock KV projection layer for MLA prefill mode.
+
+    Mimics ColumnParallelLinear behavior for kv_b_proj in MLA backends.
+    Projects kv_c_normed to [qk_nope_head_dim + v_head_dim] per head.
+    """
+
+    def __init__(self, num_heads: int, qk_nope_head_dim: int, v_head_dim: int):
+        self.num_heads = num_heads
+        self.qk_nope_head_dim = qk_nope_head_dim
+        self.v_head_dim = v_head_dim
+        self.out_dim = qk_nope_head_dim + v_head_dim
+        self.weight = torch.empty(0, dtype=torch.bfloat16)
+
+    def __call__(self, x: torch.Tensor) -> tuple[torch.Tensor]:
+        """
+        Project kv_c_normed to output space.
+
+        Args:
+            x: Input tensor [num_tokens, kv_lora_rank]
+
+        Returns:
+            Tuple containing output tensor
+                [num_tokens, num_heads, qk_nope_head_dim + v_head_dim]
+        """
+        num_tokens = x.shape[0]
+        result = torch.randn(
+            num_tokens,
+            self.num_heads,
+            self.out_dim,
+            device=x.device,
+            dtype=x.dtype,
+        )
+        return (result,)  # Return as tuple to match ColumnParallelLinear API
+
+
 class MockIndexer:
     """Mock Indexer for sparse MLA backends.
 
@@ -127,60 +158,6 @@ class MockIndexer:
         )
         self.topk_indices_buffer[:num_tokens] = indices
 
-    def fill_indices(
-        self,
-        num_tokens: int,
-        max_kv_len: int,
-        pattern: str = "random",
-        requests: Sequence[Any] | None = None,
-    ):
-        if pattern == "random":
-            self.fill_random_indices(num_tokens, max_kv_len)
-            return
-        if pattern == "prefix":
-            indices = torch.arange(
-                self.topk_tokens,
-                dtype=torch.int32,
-                device=self.topk_indices_buffer.device,
-            )
-            indices = (indices % max_kv_len).expand(num_tokens, -1)
-            self.topk_indices_buffer[:num_tokens] = indices
-            return
-        if pattern == "sliding_window":
-            if requests is None:
-                start = max(max_kv_len - self.topk_tokens, 0)
-                indices = torch.arange(
-                    start,
-                    start + self.topk_tokens,
-                    dtype=torch.int32,
-                    device=self.topk_indices_buffer.device,
-                )
-                indices = indices.clamp(max=max_kv_len - 1).expand(num_tokens, -1)
-                self.topk_indices_buffer[:num_tokens] = indices
-                return
-
-            rows = []
-            offsets = torch.arange(
-                self.topk_tokens,
-                dtype=torch.int32,
-                device=self.topk_indices_buffer.device,
-            ) - (self.topk_tokens - 1)
-            for request in requests:
-                q_len = request.q_len
-                kv_len = request.kv_len
-                context_len = kv_len - q_len
-                positions = torch.arange(
-                    context_len,
-                    kv_len,
-                    dtype=torch.int32,
-                    device=self.topk_indices_buffer.device,
-                )
-                row_indices = positions[:, None] + offsets[None, :]
-                rows.append(row_indices.clamp(min=0, max=kv_len - 1))
-            self.topk_indices_buffer[:num_tokens] = torch.cat(rows, dim=0)
-            return
-        raise ValueError(f"Unknown sparse MLA topk pattern: {pattern}")
-
 
 class MockLayer(AttentionLayerBase):
     """Mock attention layer with scale parameters and impl.
@@ -275,14 +252,10 @@ class BenchmarkConfig:
     num_kv_heads: int
     block_size: int
     device: str
-    max_model_len: int | None = None
     dtype: torch.dtype = torch.float16
     profile_memory: bool = False
-    use_cuda_graphs: bool = True
+    use_cuda_graphs: bool = False
     ncu_profile: bool = False
-    torch_profile: bool = False
-    torch_profile_dir: str | None = None
-    torch_profile_iters: int = 3
     warmup_ms: int | None = None
 
     # "auto" or "fp8"
@@ -298,10 +271,6 @@ class BenchmarkConfig:
     # Backend-specific tuning
     num_kv_splits: int | None = None  # CUTLASS MLA
     reorder_batch_threshold: int | None = None  # FlashAttn MLA, FlashMLA
-    sparse_mla_force_mqa: bool = False  # Force MQA path for sparse MLA
-    sparse_mla_mha_mode: str = "auto"  # "auto" or "dense"
-    sparse_mla_dense_mha_max_seq_len: int | None = None
-    sparse_mla_topk_pattern: str = "random"  # "random", "prefix", "sliding_window"
     num_splits: int | None = None  # FlashAttention split-K (0=auto, 1=disabled)
 
 

benchmarks/attention_benchmarks/configs/mla_sparse_mha_vs_mqa.yaml

@@ -1,474 +0,0 @@
-# Sparse MLA benchmark: forward_mha vs forward_mqa
-#
-# Usage:
-#   python benchmark.py --config configs/mla_sparse_mha_vs_mqa.yaml
-#
-# Heatmap grid:
-#   - batch_size: 1, 2, 4, 8, 16, 32
-#   - seq_len: 32, 64, 128, 256, 512, 1024, 2048
-#   - q_len: powers of two through seq_len
-#
-# Specs with q_len < seq_len include context; the q_len == seq_len diagonal
-# covers pure prefill.
-# The model shape below is the DP case. For the TP8 run, manually change
-# model.num_q_heads from 128 to 16 before rerunning this benchmark.
-
-mode: mha_vs_mqa
-
-model:
-  name: "deepseek-v3"
-  num_layers: 60
-  num_q_heads: 128
-  num_kv_heads: 1
-  head_dim: 576
-  kv_lora_rank: 512
-  qk_nope_head_dim: 128
-  qk_rope_head_dim: 64
-  v_head_dim: 128
-  block_size: 128
-  max_model_len: 2048
-
-batch_specs:
-  # Batch size 1
-  # seq_len = 32
-  - "1q1s32"
-  - "1q2s32"
-  - "1q4s32"
-  - "1q8s32"
-  - "1q16s32"
-  - "1q32"
-  # seq_len = 64
-  - "1q1s64"
-  - "1q2s64"
-  - "1q4s64"
-  - "1q8s64"
-  - "1q16s64"
-  - "1q32s64"
-  - "1q64"
-  # seq_len = 128
-  - "1q1s128"
-  - "1q2s128"
-  - "1q4s128"
-  - "1q8s128"
-  - "1q16s128"
-  - "1q32s128"
-  - "1q64s128"
-  - "1q128"
-  # seq_len = 256
-  - "1q1s256"
-  - "1q2s256"
-  - "1q4s256"
-  - "1q8s256"
-  - "1q16s256"
-  - "1q32s256"
-  - "1q64s256"
-  - "1q128s256"
-  - "1q256"
-  # seq_len = 512
-  - "1q1s512"
-  - "1q2s512"
-  - "1q4s512"
-  - "1q8s512"
-  - "1q16s512"
-  - "1q32s512"
-  - "1q64s512"
-  - "1q128s512"
-  - "1q256s512"
-  - "1q512"
-  # seq_len = 1024
-  - "1q1s1024"
-  - "1q2s1024"
-  - "1q4s1024"
-  - "1q8s1024"
-  - "1q16s1024"
-  - "1q32s1024"
-  - "1q64s1024"
-  - "1q128s1024"
-  - "1q256s1024"
-  - "1q512s1024"
-  - "1q1024"
-  # seq_len = 2048
-  - "1q1s2048"
-  - "1q2s2048"
-  - "1q4s2048"
-  - "1q8s2048"
-  - "1q16s2048"
-  - "1q32s2048"
-  - "1q64s2048"
-  - "1q128s2048"
-  - "1q256s2048"
-  - "1q512s2048"
-  - "1q1024s2048"
-  - "1q2048"
-
-  # Batch size 2
-  # seq_len = 32
-  - "2q1s32"
-  - "2q2s32"
-  - "2q4s32"
-  - "2q8s32"
-  - "2q16s32"
-  - "2q32"
-  # seq_len = 64
-  - "2q1s64"
-  - "2q2s64"
-  - "2q4s64"
-  - "2q8s64"
-  - "2q16s64"
-  - "2q32s64"
-  - "2q64"
-  # seq_len = 128
-  - "2q1s128"
-  - "2q2s128"
-  - "2q4s128"
-  - "2q8s128"
-  - "2q16s128"
-  - "2q32s128"
-  - "2q64s128"
-  - "2q128"
-  # seq_len = 256
-  - "2q1s256"
-  - "2q2s256"
-  - "2q4s256"
-  - "2q8s256"
-  - "2q16s256"
-  - "2q32s256"
-  - "2q64s256"
-  - "2q128s256"
-  - "2q256"
-  # seq_len = 512
-  - "2q1s512"
-  - "2q2s512"
-  - "2q4s512"
-  - "2q8s512"
-  - "2q16s512"
-  - "2q32s512"
-  - "2q64s512"
-  - "2q128s512"
-  - "2q256s512"
-  - "2q512"
-  # seq_len = 1024
-  - "2q1s1024"
-  - "2q2s1024"
-  - "2q4s1024"
-  - "2q8s1024"
-  - "2q16s1024"
-  - "2q32s1024"
-  - "2q64s1024"
-  - "2q128s1024"
-  - "2q256s1024"
-  - "2q512s1024"
-  - "2q1024"
-  # seq_len = 2048
-  - "2q1s2048"
-  - "2q2s2048"
-  - "2q4s2048"
-  - "2q8s2048"
-  - "2q16s2048"
-  - "2q32s2048"
-  - "2q64s2048"
-  - "2q128s2048"
-  - "2q256s2048"
-  - "2q512s2048"
-  - "2q1024s2048"
-  - "2q2048"
-
-  # Batch size 4
-  # seq_len = 32
-  - "4q1s32"
-  - "4q2s32"
-  - "4q4s32"
-  - "4q8s32"
-  - "4q16s32"
-  - "4q32"
-  # seq_len = 64
-  - "4q1s64"
-  - "4q2s64"
-  - "4q4s64"
-  - "4q8s64"
-  - "4q16s64"
-  - "4q32s64"
-  - "4q64"
-  # seq_len = 128
-  - "4q1s128"
-  - "4q2s128"
-  - "4q4s128"
-  - "4q8s128"
-  - "4q16s128"
-  - "4q32s128"
-  - "4q64s128"
-  - "4q128"
-  # seq_len = 256
-  - "4q1s256"
-  - "4q2s256"
-  - "4q4s256"
-  - "4q8s256"
-  - "4q16s256"
-  - "4q32s256"
-  - "4q64s256"
-  - "4q128s256"
-  - "4q256"
-  # seq_len = 512
-  - "4q1s512"
-  - "4q2s512"
-  - "4q4s512"
-  - "4q8s512"
-  - "4q16s512"
-  - "4q32s512"
-  - "4q64s512"
-  - "4q128s512"
-  - "4q256s512"
-  - "4q512"
-  # seq_len = 1024
-  - "4q1s1024"
-  - "4q2s1024"
-  - "4q4s1024"
-  - "4q8s1024"
-  - "4q16s1024"
-  - "4q32s1024"
-  - "4q64s1024"
-  - "4q128s1024"
-  - "4q256s1024"
-  - "4q512s1024"
-  - "4q1024"
-  # seq_len = 2048
-  - "4q1s2048"
-  - "4q2s2048"
-  - "4q4s2048"
-  - "4q8s2048"
-  - "4q16s2048"
-  - "4q32s2048"
-  - "4q64s2048"
-  - "4q128s2048"
-  - "4q256s2048"
-  - "4q512s2048"
-  - "4q1024s2048"
-  - "4q2048"
-
-  # Batch size 8
-  # seq_len = 32
-  - "8q1s32"
-  - "8q2s32"
-  - "8q4s32"
-  - "8q8s32"
-  - "8q16s32"
-  - "8q32"
-  # seq_len = 64
-  - "8q1s64"
-  - "8q2s64"
-  - "8q4s64"
-  - "8q8s64"
-  - "8q16s64"
-  - "8q32s64"
-  - "8q64"
-  # seq_len = 128
-  - "8q1s128"
-  - "8q2s128"
-  - "8q4s128"
-  - "8q8s128"
-  - "8q16s128"
-  - "8q32s128"
-  - "8q64s128"
-  - "8q128"
-  # seq_len = 256
-  - "8q1s256"
-  - "8q2s256"
-  - "8q4s256"
-  - "8q8s256"
-  - "8q16s256"
-  - "8q32s256"
-  - "8q64s256"
-  - "8q128s256"
-  - "8q256"
-  # seq_len = 512
-  - "8q1s512"
-  - "8q2s512"
-  - "8q4s512"
-  - "8q8s512"
-  - "8q16s512"
-  - "8q32s512"
-  - "8q64s512"
-  - "8q128s512"
-  - "8q256s512"
-  - "8q512"
-  # seq_len = 1024
-  - "8q1s1024"
-  - "8q2s1024"
-  - "8q4s1024"
-  - "8q8s1024"
-  - "8q16s1024"
-  - "8q32s1024"
-  - "8q64s1024"
-  - "8q128s1024"
-  - "8q256s1024"
-  - "8q512s1024"
-  - "8q1024"
-  # seq_len = 2048
-  - "8q1s2048"
-  - "8q2s2048"
-  - "8q4s2048"
-  - "8q8s2048"
-  - "8q16s2048"
-  - "8q32s2048"
-  - "8q64s2048"
-  - "8q128s2048"
-  - "8q256s2048"
-  - "8q512s2048"
-  - "8q1024s2048"
-  - "8q2048"
-
-  # Batch size 16
-  # seq_len = 32
-  - "16q1s32"
-  - "16q2s32"
-  - "16q4s32"
-  - "16q8s32"
-  - "16q16s32"
-  - "16q32"
-  # seq_len = 64
-  - "16q1s64"
-  - "16q2s64"
-  - "16q4s64"
-  - "16q8s64"
-  - "16q16s64"
-  - "16q32s64"
-  - "16q64"
-  # seq_len = 128
-  - "16q1s128"
-  - "16q2s128"
-  - "16q4s128"
-  - "16q8s128"
-  - "16q16s128"
-  - "16q32s128"
-  - "16q64s128"
-  - "16q128"
-  # seq_len = 256
-  - "16q1s256"
-  - "16q2s256"
-  - "16q4s256"
-  - "16q8s256"
-  - "16q16s256"
-  - "16q32s256"
-  - "16q64s256"
-  - "16q128s256"
-  - "16q256"
-  # seq_len = 512
-  - "16q1s512"
-  - "16q2s512"
-  - "16q4s512"
-  - "16q8s512"
-  - "16q16s512"
-  - "16q32s512"
-  - "16q64s512"
-  - "16q128s512"
-  - "16q256s512"
-  - "16q512"
-  # seq_len = 1024
-  - "16q1s1024"
-  - "16q2s1024"
-  - "16q4s1024"
-  - "16q8s1024"
-  - "16q16s1024"
-  - "16q32s1024"
-  - "16q64s1024"
-  - "16q128s1024"
-  - "16q256s1024"
-  - "16q512s1024"
-  - "16q1024"
-  # seq_len = 2048
-  - "16q1s2048"
-  - "16q2s2048"
-  - "16q4s2048"
-  - "16q8s2048"
-  - "16q16s2048"
-  - "16q32s2048"
-  - "16q64s2048"
-  - "16q128s2048"
-  - "16q256s2048"
-  - "16q512s2048"
-  - "16q1024s2048"
-  - "16q2048"
-
-  # Batch size 32
-  # seq_len = 32
-  - "32q1s32"
-  - "32q2s32"
-  - "32q4s32"
-  - "32q8s32"
-  - "32q16s32"
-  - "32q32"
-  # seq_len = 64
-  - "32q1s64"
-  - "32q2s64"
-  - "32q4s64"
-  - "32q8s64"
-  - "32q16s64"
-  - "32q32s64"
-  - "32q64"
-  # seq_len = 128
-  - "32q1s128"
-  - "32q2s128"
-  - "32q4s128"
-  - "32q8s128"
-  - "32q16s128"
-  - "32q32s128"
-  - "32q64s128"
-  - "32q128"
-  # seq_len = 256
-  - "32q1s256"
-  - "32q2s256"
-  - "32q4s256"
-  - "32q8s256"
-  - "32q16s256"
-  - "32q32s256"
-  - "32q64s256"
-  - "32q128s256"
-  - "32q256"
-  # seq_len = 512
-  - "32q1s512"
-  - "32q2s512"
-  - "32q4s512"
-  - "32q8s512"
-  - "32q16s512"
-  - "32q32s512"
-  - "32q64s512"
-  - "32q128s512"
-  - "32q256s512"
-  - "32q512"
-  # seq_len = 1024
-  - "32q1s1024"
-  - "32q2s1024"
-  - "32q4s1024"
-  - "32q8s1024"
-  - "32q16s1024"
-  - "32q32s1024"
-  - "32q64s1024"
-  - "32q128s1024"
-  - "32q256s1024"
-  - "32q512s1024"
-  - "32q1024"
-  # seq_len = 2048
-  - "32q1s2048"
-  - "32q2s2048"
-  - "32q4s2048"
-  - "32q8s2048"
-  - "32q16s2048"
-  - "32q32s2048"
-  - "32q64s2048"
-  - "32q128s2048"
-  - "32q256s2048"
-  - "32q512s2048"
-  - "32q1024s2048"
-  - "32q2048"
-
-backends:
-  - FLASHMLA_SPARSE
-
-device: "cuda:0"
-profile_memory: false
-sparse_mla_dense_mha_max_seq_len: 2048
-sparse_mla_topk_pattern: "random"
-
-output:
-  csv: "benchmark_output/mla_sparse_mha_vs_mqa.csv"
-  json: "benchmark_output/mla_sparse_mha_vs_mqa.json"

benchmarks/attention_benchmarks/mla_runner.py

@@ -9,8 +9,6 @@ needing full VllmConfig integration.
 """
 
 import statistics
-import tempfile
-from pathlib import Path
 
 import numpy as np
 import torch
@@ -19,6 +17,7 @@ from common import (
     BenchmarkResult,
     MockHfConfig,
     MockIndexer,
+    MockKVBProj,
     MockLayer,
     run_do_bench,
     run_ncu_profile,
@@ -34,59 +33,8 @@ from vllm.config import (
     VllmConfig,
     set_current_vllm_config,
 )
-from vllm.model_executor.layers.linear import ColumnParallelLinear
 from vllm.v1.attention.backends.mla.prefill.registry import MLAPrefillBackendEnum
 
-
-def _safe_profile_name(value: str) -> str:
-    return "".join(c if c.isalnum() or c in "._-" else "_" for c in value)
-
-
-def _create_kv_b_proj(
-    mla_dims: dict,
-    device: torch.device,
-):
-    kv_b_proj = ColumnParallelLinear(
-        mla_dims["kv_lora_rank"],
-        mla_dims["num_q_heads"]
-        * (mla_dims["qk_nope_head_dim"] + mla_dims["v_head_dim"]),
-        bias=False,
-        params_dtype=torch.bfloat16,
-        quant_config=None,
-        prefix="benchmark.kv_b_proj",
-    ).to(device)
-    with torch.no_grad():
-        kv_b_proj.weight.copy_(torch.randn_like(kv_b_proj.weight))
-    return kv_b_proj
-
-
-def _ensure_single_rank_model_parallel() -> None:
-    import torch.distributed as dist
-
-    from vllm.distributed import (
-        ensure_model_parallel_initialized,
-        init_distributed_environment,
-        model_parallel_is_initialized,
-    )
-
-    if not dist.is_available():
-        return
-    if not dist.is_initialized():
-        with tempfile.NamedTemporaryFile(
-            prefix="vllm_bench_dist_", delete=False
-        ) as init_file:
-            distributed_init_method = f"file://{init_file.name}"
-        init_distributed_environment(
-            world_size=1,
-            rank=0,
-            distributed_init_method=distributed_init_method,
-            local_rank=0,
-            backend="nccl",
-        )
-    if not model_parallel_is_initialized():
-        ensure_model_parallel_initialized(1, 1)
-
-
 # ============================================================================
 # VllmConfig Creation
 # ============================================================================
@@ -118,12 +66,10 @@ def create_minimal_vllm_config(
     block_size: int = 128,
     max_num_seqs: int = 256,
     max_num_batched_tokens: int = 8192,
-    max_model_len: int = 32768,
     mla_dims: dict | None = None,
     index_topk: int | None = None,
     prefill_backend: str | None = None,
     kv_cache_dtype: str = "auto",
-    sparse_mla_force_mqa: bool = False,
 ) -> VllmConfig:
     """
     Create minimal VllmConfig for MLA benchmarks.
@@ -140,8 +86,6 @@ def create_minimal_vllm_config(
         prefill_backend: Prefill backend name (e.g., "fa3", "fa4", "flashinfer",
                         "trtllm"). Configures the attention config to force
                         the specified prefill backend.
-        sparse_mla_force_mqa: If True, forces all sparse MLA tokens through
-                    forward_mqa (even prefill tokens).
 
     Returns:
         VllmConfig for benchmarking
@@ -187,7 +131,7 @@ def create_minimal_vllm_config(
             trust_remote_code=True,
             dtype="bfloat16",
             seed=0,
-            max_model_len=max_model_len,
+            max_model_len=32768,
             quantization=None,
             enforce_eager=False,
             max_logprobs=20,
@@ -219,7 +163,7 @@ def create_minimal_vllm_config(
     scheduler_config = SchedulerConfig(
         max_num_seqs=max_num_seqs,
         max_num_batched_tokens=max(max_num_batched_tokens, max_num_seqs),
-        max_model_len=max_model_len,
+        max_model_len=32768,
         is_encoder_decoder=False,
         enable_chunked_prefill=True,
     )
@@ -248,9 +192,6 @@ def create_minimal_vllm_config(
                 "flash_attn_version"
             ]
 
-    if sparse_mla_force_mqa:
-        vllm_config.attention_config.sparse_mla_force_mqa = True
-
     return vllm_config
 
 
@@ -607,7 +548,12 @@ def _create_backend_impl(
     # Calculate scale
     scale = 1.0 / np.sqrt(mla_dims["qk_nope_head_dim"] + mla_dims["qk_rope_head_dim"])
 
-    kv_b_proj = _create_kv_b_proj(mla_dims, device)
+    # Create mock kv_b_proj layer for prefill mode
+    mock_kv_b_proj = MockKVBProj(
+        num_heads=mla_dims["num_q_heads"],
+        qk_nope_head_dim=mla_dims["qk_nope_head_dim"],
+        v_head_dim=mla_dims["v_head_dim"],
+    )
 
     # Create indexer for sparse backends
     indexer = None
@@ -638,7 +584,7 @@ def _create_backend_impl(
         "qk_rope_head_dim": mla_dims["qk_rope_head_dim"],
         "qk_head_dim": mla_dims["qk_nope_head_dim"] + mla_dims["qk_rope_head_dim"],
         "v_head_dim": mla_dims["v_head_dim"],
-        "kv_b_proj": kv_b_proj,
+        "kv_b_proj": mock_kv_b_proj,
     }
 
     # Add indexer for sparse backends
@@ -839,35 +785,14 @@ def _run_single_benchmark(
     # Fill indexer with random indices for sparse backends
     is_sparse = backend_cfg.get("is_sparse", False)
     if is_sparse and indexer is not None:
-        indexer.fill_indices(
+        indexer.fill_random_indices(total_q, max_kv_len)
-            total_q,
-            max_kv_len,
-            getattr(config, "sparse_mla_topk_pattern", "random"),
-        )
 
     # Determine which forward methods to use based on metadata.
-    # Non-sparse backends use .decode/.prefill sub-objects.
+    # Sparse MLA backends always use forward_mqa
-    # Sparse backends use num_decode_tokens/num_prefills directly.
+    has_decode = is_sparse or getattr(metadata, "decode", None) is not None
-    #
+    has_prefill = not is_sparse and getattr(metadata, "prefill", None) is not None
-    # sparse_mla_force_mqa overrides: even for prefill metadata, use MQA.
-    force_mqa = getattr(config, "sparse_mla_force_mqa", False)
-    force_dense_mha = getattr(config, "sparse_mla_mha_mode", "auto") == "dense"
-    if force_mqa:
-        has_decode = True
-        has_prefill = False
-    elif is_sparse:
-        has_decode = metadata.num_decode_tokens > 0
-        has_prefill = metadata.num_prefills > 0
-    else:
-        has_decode = metadata.decode is not None
-        has_prefill = metadata.prefill is not None
     if not has_decode and not has_prefill:
         raise RuntimeError("Metadata has neither decode nor prefill metadata")
-    if is_sparse and force_dense_mha and not has_prefill:
-        raise RuntimeError(
-            "Sparse MLA dense_mha benchmark did not produce prefill metadata. "
-            "Check reorder_batch_threshold/path forcing."
-        )
 
     num_decode = (
         metadata.num_decode_tokens
@@ -946,6 +871,7 @@ def _run_single_benchmark(
                 metadata,
                 prefill_inputs["k_scale"],
                 prefill_fp8_output if fused_output else prefill_inputs["output"],
+                prefill_output_scale if fused_output else None,
             )
             if fused_output:
                 out = prefill_fp8_output
@@ -972,48 +898,6 @@ def _run_single_benchmark(
             throughput_tokens_per_sec=0.0,
         )
 
-    if config.torch_profile:
-        profile_dir = Path(
-            config.torch_profile_dir or "benchmark_outputs/torch_profiles"
-        )
-        profile_dir.mkdir(parents=True, exist_ok=True)
-        trace_name = _safe_profile_name(f"{config.backend}_{config.batch_spec}")
-        trace_path = profile_dir / f"{trace_name}.json"
-        iters = max(config.torch_profile_iters, 1)
-
-        forward_fn()
-        torch.accelerator.synchronize()
-        with torch.profiler.profile(
-            activities=[
-                torch.profiler.ProfilerActivity.CPU,
-                torch.profiler.ProfilerActivity.CUDA,
-            ],
-            record_shapes=True,
-            profile_memory=True,
-            with_stack=False,
-        ) as prof:
-            for _ in range(iters):
-                forward_fn()
-                torch.accelerator.synchronize()
-                prof.step()
-        prof.export_chrome_trace(str(trace_path))
-        print(f"Saved PyTorch profiler trace to {trace_path}")
-        print(
-            prof.key_averages().table(
-                sort_by="cuda_time_total",
-                row_limit=25,
-            )
-        )
-        return BenchmarkResult(
-            config=config,
-            mean_time=0.0,
-            median_time=0.0,
-            std_time=0.0,
-            min_time=0.0,
-            max_time=0.0,
-            throughput_tokens_per_sec=0.0,
-        )
-
     all_ms = run_do_bench(benchmark_fn, config.use_cuda_graphs, config.warmup_ms)
 
     # Convert ms to seconds per layer
@@ -1036,7 +920,6 @@ def _run_mla_benchmark_batched(
     configs_with_params: list[tuple],  # [(config, threshold, num_splits), ...]
     index_topk: int = 2048,
     prefill_backend: str | None = None,
-    sparse_mla_force_mqa: bool = False,
     output_scale: float | None = None,
     fuse_quant_op: bool = False,
 ) -> list[BenchmarkResult]:
@@ -1057,8 +940,6 @@ def _run_mla_benchmark_batched(
         index_topk: Topk value for sparse MLA backends (default 2048)
         prefill_backend: Prefill backend name (e.g., "fa3", "fa4").
             When set, forces the specified FlashAttention version for prefill.
-        sparse_mla_force_mqa: If True, forces all sparse MLA tokens through
-            forward_mqa (even prefill tokens).
 
     Returns:
         List of BenchmarkResult objects
@@ -1099,41 +980,21 @@ def _run_mla_benchmark_batched(
         sum(r.q_len for r in parse_batch_spec(cfg.batch_spec))
         for cfg, *_ in configs_with_params
     )
-    max_model_len = max(
-        max_total_q,
-        max(
-            getattr(cfg, "max_model_len", None) or 32768
-            for cfg, *_ in configs_with_params
-        ),
-    )
 
     # Create and set vLLM config for MLA (reused across all benchmarks)
     vllm_config = create_minimal_vllm_config(
         model_name="deepseek-v3",  # Used only for model path
         block_size=block_size,
         max_num_batched_tokens=max_total_q,
-        max_model_len=max_model_len,
         mla_dims=mla_dims,  # Use custom dims from config or default
         index_topk=index_topk if is_sparse else None,
         prefill_backend=prefill_backend,
         kv_cache_dtype=kv_cache_dtype,
-        sparse_mla_force_mqa=sparse_mla_force_mqa,
     )
 
     results = []
 
-    # Initialize workspace manager (needed by metadata builders)
-    from vllm.v1.worker.workspace import (
-        init_workspace_manager,
-        is_workspace_manager_initialized,
-    )
-
-    if not is_workspace_manager_initialized():
-        init_workspace_manager(device)
-
     with set_current_vllm_config(vllm_config):
-        _ensure_single_rank_model_parallel()
-
         # Create backend impl, layer, builder, and indexer (reused across benchmarks)
         impl, layer, builder_instance, indexer = _create_backend_impl(
             backend_cfg,
@@ -1179,20 +1040,9 @@ def _run_mla_benchmark_batched(
         for config, threshold, num_splits in configs_with_params:
             # Set threshold for this benchmark (FlashAttn/FlashMLA only)
             original_threshold = None
-            effective_threshold = threshold
+            if threshold is not None and builder_instance:
-            force_dense_mha = (
-                is_sparse
-                and getattr(config, "sparse_mla_mha_mode", "auto") == "dense"
-                and not getattr(config, "sparse_mla_force_mqa", False)
-            )
-            if force_dense_mha:
-                # Sparse MLA normally treats q_len <= 1 as decode. Use an
-                # impossible threshold so dense_mha benchmarks actually run
-                # the prefill/MHA path, including q_len=1 short extends.
-                effective_threshold = -1
-            if effective_threshold is not None and builder_instance:
                 original_threshold = builder_instance.reorder_batch_threshold
-                builder_instance.reorder_batch_threshold = effective_threshold
+                builder_instance.reorder_batch_threshold = threshold
 
             # Set num_splits for CUTLASS
             original_num_splits = None
@@ -1240,7 +1090,6 @@ def run_mla_benchmark(
     num_kv_splits: int | None = None,
     index_topk: int = 2048,
     prefill_backend: str | None = None,
-    sparse_mla_force_mqa: bool = False,
     output_scale: float | None = None,
     fuse_quant_op: bool = False,
 ) -> BenchmarkResult | list[BenchmarkResult]:
@@ -1262,8 +1111,6 @@ def run_mla_benchmark(
         index_topk: Topk value for sparse MLA backends (default 2048)
         prefill_backend: Prefill backend name (e.g., "fa3", "fa4").
             When set, forces the specified FlashAttention version for prefill.
-        sparse_mla_force_mqa: If True, forces all sparse MLA tokens through
-            forward_mqa (even prefill tokens).
         output_scale: Static per-tensor FP8 scale for prefill output (None = bf16).
         fuse_quant_op: With output_scale set, fuse the FP8 write into the prefill
             kernel vs a standalone post-quant kernel. See _run_single_benchmark.
@@ -1295,7 +1142,6 @@ def run_mla_benchmark(
         configs_with_params,
         index_topk,
         prefill_backend=prefill_backend,
-        sparse_mla_force_mqa=sparse_mla_force_mqa,
         output_scale=output_scale,
         fuse_quant_op=fuse_quant_op,
     )

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, PythonBackend, TokenizersBackend
+from transformers import AutoTokenizer, PreTrainedTokenizer, PreTrainedTokenizerFast
 
 # 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,
-) -> PythonBackend | TokenizersBackend:
+) -> PreTrainedTokenizer | PreTrainedTokenizerFast:
     if pretrained_model_name_or_path is not None and not os.path.exists(
         pretrained_model_name_or_path
     ):

benchmarks/kernels/bench_cp_gather_fp8.py

@@ -69,11 +69,12 @@ def make_inputs(total_tokens, num_reqs, block_size):
     # Output workspace
     dst = torch.zeros(total_tokens, HEAD_DIM, dtype=torch.bfloat16, device="cuda")
 
+    seq_lens_t = torch.tensor(seq_lens, dtype=torch.int32, device="cuda")
     workspace_starts_t = torch.tensor(
         workspace_starts, dtype=torch.int32, device="cuda"
     )
 
-    return cache, dst, block_table, workspace_starts_t
+    return cache, dst, block_table, seq_lens_t, workspace_starts_t
 
 
 def bench_scenario(label, num_reqs, total_tokens_list, save_path):
@@ -93,7 +94,7 @@ def bench_scenario(label, num_reqs, total_tokens_list, save_path):
         )
     )
     def bench_fn(total_tokens, provider, num_reqs):
-        cache, dst, block_table, ws_starts = make_inputs(
+        cache, dst, block_table, seq_lens_t, ws_starts = make_inputs(
             total_tokens, num_reqs, BLOCK_SIZE
         )
 
@@ -101,7 +102,7 @@ def bench_scenario(label, num_reqs, total_tokens_list, save_path):
 
         ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
             lambda: ops.cp_gather_and_upconvert_fp8_kv_cache(
-                cache, dst, block_table, ws_starts, num_reqs
+                cache, dst, block_table, seq_lens_t, ws_starts, num_reqs
             ),
             quantiles=quantiles,
             rep=500,

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.utils.platform_utils import get_device_name_as_file_name
+from vllm.platforms import current_platform
 
 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 = get_device_name_as_file_name()
+            device_name = current_platform.get_device_name().replace(" ", "_")
             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_k3_cutedsl_residual.py

@@ -1,367 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""Benchmark the Kimi-K3 latent MoE addmm against CuTe residual GEMM.
-
-The benchmark covers ``BF16[M, 3584] @ BF16[7168, 3584].T + BF16[M, 7168]``
-with FP32 accumulation and BF16 output. Both backends execute through CUDA
-Graph replay. Weights and residuals rotate across buffers exceeding L2 so the
-comparison models the full latent MoE projection-and-add path.
-"""
-
-from __future__ import annotations
-
-import argparse
-import dataclasses
-import importlib.util
-import json
-import math
-import statistics
-from collections.abc import Callable, Sequence
-from pathlib import Path
-from typing import Any
-
-import cutlass
-import cutlass.cute as cute
-import torch
-from cuda.bindings import driver as cuda
-from cuda.bindings.driver import CUstream
-from quack.compile_utils import make_fake_tensor
-
-N = 7168
-K = 3584
-
-
-@dataclasses.dataclass(frozen=True, slots=True)
-class Config:
-    block_size: int
-    outputs_per_block: int
-    k_unroll: int
-    vector_width: int = 8
-
-
-def parse_config(value: str) -> Config:
-    try:
-        parts = [int(part) for part in value.split(",")]
-    except ValueError as error:
-        raise argparse.ArgumentTypeError(
-            "config must be BLOCK,OUTPUTS,K_UNROLL[,VECTOR_WIDTH]"
-        ) from error
-    if len(parts) == 3:
-        return Config(*parts)
-    if len(parts) == 4:
-        return Config(*parts)
-    raise argparse.ArgumentTypeError(
-        "config must be BLOCK,OUTPUTS,K_UNROLL[,VECTOR_WIDTH]"
-    )
-
-
-def production_residual_config(m: int) -> Config | None:
-    """The measured Latent-MoE residual config for M, from the K3 table."""
-    from vllm.models.kimi_k3.nvidia.low_latency_gemm import KIMI_K3_PROJECTIONS
-
-    spec = KIMI_K3_PROJECTIONS.get((N, K))
-    config = spec.residual_config(m) if spec is not None else None
-    if config is None:
-        return None
-    return Config(
-        config.block_size,
-        config.outputs_per_block,
-        config.k_unroll,
-        config.vector_width,
-    )
-
-
-def candidate_configs(mode: str, selected: Config | None, m: int) -> list[Config]:
-    if mode == "selected":
-        if selected is not None:
-            return [selected]
-        # No explicit --config: fall back to the production table for this M.
-        config = production_residual_config(m)
-        return [config] if config is not None else []
-    if mode == "baseline":
-        return [Config(224, 4, 2)]
-    return [
-        Config(block_size, outputs_per_block, k_unroll, vector_width)
-        for vector_width in (4, 8)
-        for block_size in (32, 64, 128, 224, 448)
-        if block_size % 32 == 0 and K % (block_size * vector_width) == 0
-        for outputs_per_block in (1, 2, 4, 7, 8)
-        if N % outputs_per_block == 0
-        for k_unroll in (1, 2, 4)
-    ]
-
-
-def load_kernel_class(path: Path):
-    spec = importlib.util.spec_from_file_location("cute_skinny_device", path)
-    if spec is None or spec.loader is None:
-        raise RuntimeError(f"cannot load CuTe kernel from {path}")
-    module = importlib.util.module_from_spec(spec)
-    spec.loader.exec_module(module)
-    return module.CuteSkinnyGemm
-
-
-def stream() -> CUstream:
-    return CUstream(torch.cuda.current_stream().cuda_stream)
-
-
-def compile_kernel(kernel_class, m: int, config: Config, max_registers: int):
-    element_type = cutlass.BFloat16
-    n = cute.sym_int(divisibility=config.outputs_per_block)
-    k = cute.sym_int(divisibility=config.block_size * config.vector_width)
-    a = make_fake_tensor(element_type, (m, k), divisibility=config.vector_width)
-    b = make_fake_tensor(element_type, (n, k), divisibility=config.vector_width)
-    residual = make_fake_tensor(element_type, (m, n), divisibility=1)
-    c = make_fake_tensor(element_type, (m, n), divisibility=1)
-    kernel = kernel_class(
-        element_type=element_type,
-        num_rows=m,
-        block_size=config.block_size,
-        outputs_per_block=config.outputs_per_block,
-        vector_width=config.vector_width,
-        k_unroll=config.k_unroll,
-        has_residual=True,
-        use_pdl=True,
-    )
-    return cute.compile(
-        kernel,
-        a,
-        b,
-        residual,
-        c,
-        stream(),
-        options=(
-            "--enable-tvm-ffi --keep-cubin "
-            f"--ptxas-options -maxrregcount={max_registers} "
-            "--ptxas-options -lineinfo"
-        ),
-    )
-
-
-def resource_usage(compiled) -> dict[str, Any]:
-    executor = getattr(compiled, "_default_executor", None)
-    context = getattr(executor, "exec_context", None)
-    functions = getattr(context, "kernel_functions", None)
-    if not functions:
-        return {"resource_metrics_available": False}
-
-    def attribute(name, function) -> int:
-        error, value = cuda.cuFuncGetAttribute(name, function)
-        if error != cuda.CUresult.CUDA_SUCCESS:
-            raise RuntimeError(f"cuFuncGetAttribute failed with {error}")
-        return int(value)
-
-    registers = [
-        attribute(cuda.CUfunction_attribute.CU_FUNC_ATTRIBUTE_NUM_REGS, function)
-        for function in functions
-    ]
-    local_bytes = [
-        attribute(
-            cuda.CUfunction_attribute.CU_FUNC_ATTRIBUTE_LOCAL_SIZE_BYTES,
-            function,
-        )
-        for function in functions
-    ]
-    return {
-        "resource_metrics_available": True,
-        "registers_per_thread": max(registers, default=0),
-        "spill_bytes": max(local_bytes, default=0),
-    }
-
-
-def rotating_buffer_count(m: int, multiplier: float, limit: int) -> int:
-    properties = torch.cuda.get_device_properties(0)
-    bytes_per_pair = (N * K + m * N) * 2
-    target = math.ceil(multiplier * properties.L2_cache_size)
-    return max(2, min(limit, math.ceil(target / bytes_per_pair)))
-
-
-def graph_samples(
-    launch: Callable[[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor], None],
-    activation: torch.Tensor,
-    weights: Sequence[torch.Tensor],
-    residuals: Sequence[torch.Tensor],
-    repeats: int,
-    replays: int,
-) -> tuple[list[float], list[torch.Tensor]]:
-    outputs = [torch.empty_like(residual) for residual in residuals]
-    for weight, residual, output in zip(weights, residuals, outputs):
-        launch(activation, weight, residual, output)
-    torch.cuda.synchronize()
-
-    graph = torch.cuda.CUDAGraph()
-    with torch.cuda.graph(graph):
-        for weight, residual, output in zip(weights, residuals, outputs):
-            launch(activation, weight, residual, output)
-    for _ in range(20):
-        graph.replay()
-    torch.cuda.synchronize()
-
-    samples = []
-    for _ in range(repeats):
-        start = torch.cuda.Event(enable_timing=True)
-        end = torch.cuda.Event(enable_timing=True)
-        start.record()
-        for _ in range(replays):
-            graph.replay()
-        end.record()
-        end.synchronize()
-        samples.append(start.elapsed_time(end) * 1000.0 / (replays * len(weights)))
-    return samples, outputs
-
-
-def summarize(samples: Sequence[float]) -> dict[str, Any]:
-    ordered = sorted(samples)
-
-    def percentile(fraction: float) -> float:
-        position = fraction * (len(ordered) - 1)
-        lower = math.floor(position)
-        upper = math.ceil(position)
-        if lower == upper:
-            return ordered[lower]
-        weight = position - lower
-        return ordered[lower] * (1.0 - weight) + ordered[upper] * weight
-
-    mean = statistics.mean(samples)
-    return {
-        "median_us": statistics.median(samples),
-        "p10_us": percentile(0.1),
-        "p90_us": percentile(0.9),
-        "mean_us": mean,
-        "cv_pct": statistics.pstdev(samples) / mean * 100.0,
-        "samples_us": list(samples),
-    }
-
-
-def correctness(
-    output: torch.Tensor,
-    activation: torch.Tensor,
-    weight: torch.Tensor,
-    residual: torch.Tensor,
-) -> dict[str, Any]:
-    actual = output.float()
-    reference = activation.float() @ weight.float().t() + residual.float()
-    error = (actual - reference).abs()
-    scaled_error = error / (reference.abs() + 1.0)
-    cosine = torch.nn.functional.cosine_similarity(
-        actual.flatten(), reference.flatten(), dim=0
-    ).item()
-    return {
-        "valid": cosine > 0.999,
-        "cosine": cosine,
-        "max_abs_error": error.max().item(),
-        "max_scaled_error": scaled_error.max().item(),
-    }
-
-
-def main() -> None:
-    parser = argparse.ArgumentParser(description=__doc__)
-    parser.add_argument("--kernel", type=Path, required=True)
-    parser.add_argument("--output", type=Path, required=True)
-    parser.add_argument(
-        "--mode", choices=("baseline", "sweep", "selected"), default="baseline"
-    )
-    parser.add_argument("--config", type=parse_config)
-    parser.add_argument("--m", type=int, action="append")
-    parser.add_argument("--config-shard", type=int, default=0)
-    parser.add_argument("--num-config-shards", type=int, default=1)
-    parser.add_argument("--repeats", type=int, default=21)
-    parser.add_argument("--replays", type=int, default=200)
-    parser.add_argument("--cache-multiplier", type=float, default=3.0)
-    parser.add_argument("--max-buffers", type=int, default=32)
-    parser.add_argument("--max-registers", type=int, default=64)
-    args = parser.parse_args()
-
-    token_counts = args.m or list(range(1, 17))
-    if any(not 1 <= m <= 16 for m in token_counts):
-        raise ValueError("expected 1 <= M <= 16")
-    if not 0 <= args.config_shard < args.num_config_shards:
-        raise ValueError("config shard must be in [0, num_config_shards)")
-    torch.cuda.set_device(0)
-    if torch.cuda.get_device_capability() != (10, 3):
-        raise RuntimeError("this benchmark requires SM103")
-
-    kernel_class = load_kernel_class(args.kernel)
-    properties = torch.cuda.get_device_properties(0)
-    metadata = {
-        "device": properties.name,
-        "compute_capability": list(torch.cuda.get_device_capability()),
-        "torch_version": torch.__version__,
-        "cuda_version": torch.version.cuda,
-    }
-    args.output.parent.mkdir(parents=True, exist_ok=True)
-    with args.output.open("w", encoding="utf-8") as output_file:
-        for m in token_counts:
-            configs = candidate_configs(args.mode, args.config, m)
-            torch.manual_seed(20260722 + m)
-            count = rotating_buffer_count(m, args.cache_multiplier, args.max_buffers)
-            activation = torch.randn((m, K), device="cuda", dtype=torch.bfloat16)
-            weights = [
-                torch.randn((N, K), device="cuda", dtype=torch.bfloat16)
-                for _ in range(count)
-            ]
-            residuals = [
-                torch.randn((m, N), device="cuda", dtype=torch.bfloat16)
-                for _ in range(count)
-            ]
-            candidates: list[tuple[str, Config | None]] = [("cublas_addmm", None)]
-            candidates.extend(
-                ("cute_residual", config)
-                for index, config in enumerate(configs)
-                if index % args.num_config_shards == args.config_shard
-            )
-            for backend, config in candidates:
-                row: dict[str, Any] = {
-                    "m": m,
-                    "n": N,
-                    "k": K,
-                    "backend": backend,
-                    "mode": args.mode,
-                    "config": dataclasses.asdict(config) if config else {},
-                    "num_buffers": count,
-                    "cache_multiplier": args.cache_multiplier,
-                    **metadata,
-                }
-                try:
-                    if backend == "cublas_addmm":
-                        launch = lambda a, b, residual, c: torch.addmm(
-                            residual, a, b.t(), out=c
-                        )
-                    else:
-                        if config is None:
-                            raise AssertionError("missing CuTe config")
-                        compiled = compile_kernel(
-                            kernel_class, m, config, args.max_registers
-                        )
-                        launch = lambda a, b, residual, c, fn=compiled: fn(
-                            a, b, residual, c, stream()
-                        )
-                        row.update(resource_usage(compiled))
-                    samples, outputs = graph_samples(
-                        launch,
-                        activation,
-                        weights,
-                        residuals,
-                        args.repeats,
-                        args.replays,
-                    )
-                    row.update(
-                        correctness(outputs[0], activation, weights[0], residuals[0])
-                    )
-                    row.update(summarize(samples))
-                except Exception as error:  # noqa: BLE001
-                    row.update(
-                        {
-                            "valid": False,
-                            "error": f"{type(error).__name__}: {error}",
-                        }
-                    )
-                output_file.write(json.dumps(row, sort_keys=True) + "\n")
-                output_file.flush()
-                print(json.dumps(row, sort_keys=True), flush=True)
-
-            del activation, weights, residuals
-            torch.cuda.empty_cache()
-
-
-if __name__ == "__main__":
-    main()

benchmarks/kernels/benchmark_kimi_k3_latent_moe_tail.py

@@ -1,806 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""Benchmark the Kimi K3 latent-MoE tail and its up-projection kernels.
-
-The ``up-projection`` subcommand isolates the TP-local dynamic and static-M
-skinny GEMMs. It rotates weights through a working set larger than L2 to model
-successive model layers.
-
-The ``whole-tail`` subcommand measures the distributed operator. Its reference
-path includes two AllReduces, RMSNorm, the replicated up-projection, and the
-final add. CUDA-event samples report the slowest rank so cross-rank skew is
-included.
-
-Examples:
-
-.. code-block:: console
-
-    .venv/bin/python \
-      benchmarks/kernels/benchmark_kimi_k3_latent_moe_tail.py up-projection
-
-    torchrun --nproc-per-node=8 \
-      benchmarks/kernels/benchmark_kimi_k3_latent_moe_tail.py whole-tail
-
-For multi-node runs, launch one ``torchrun`` agent per node and use a shared
-rendezvous endpoint.
-"""
-
-from __future__ import annotations
-
-import argparse
-import json
-import math
-import os
-import statistics
-from collections.abc import Callable, Sequence
-from dataclasses import asdict
-from pathlib import Path
-from typing import Any
-
-import cutlass
-import cutlass.utils as utils
-import torch
-import torch.distributed as dist
-import torch.nn.functional as F
-from cuda.bindings import driver as cuda
-
-from vllm.distributed import get_tp_group
-from vllm.distributed.parallel_state import (
-    init_distributed_environment,
-    initialize_model_parallel,
-    set_custom_all_reduce,
-)
-from vllm.model_executor.warmup.cutedsl_warmup import cutedsl_warmup
-from vllm.models.kimi_k3.nvidia.ops import latent_moe_tail
-from vllm.models.kimi_k3.nvidia.ops.cute_dsl.latent_moe_tail import (
-    fused_add_multicast_gemm,
-    fused_add_multicast_skinny_gemm,
-)
-
-HIDDEN_SIZE = 7168
-LATENT_SIZE = 3584
-RMS_EPS = 0.1
-MAX_NUM_TOKENS = 16
-MMA_TILER_MN = (64, 32)
-CLUSTER_SHAPE_MN = (1, 8)
-B_PRIME_STAGES = 2
-
-
-def parse_up_projection_config(
-    value: str,
-) -> fused_add_multicast_skinny_gemm.SkinnyConfig:
-    try:
-        values = [int(part) for part in value.split(",")]
-    except ValueError as error:
-        raise argparse.ArgumentTypeError(
-            "config must be BLOCK,OUTPUTS,K_UNROLL[,VECTOR_WIDTH[,PREFETCH_B]]"
-        ) from error
-    if len(values) in (3, 4):
-        return fused_add_multicast_skinny_gemm.SkinnyConfig(*values)
-    if len(values) == 5 and values[4] in (0, 1):
-        return fused_add_multicast_skinny_gemm.SkinnyConfig(
-            *values[:4],
-            prefetch_b_before_pdl=bool(values[4]),
-        )
-    raise argparse.ArgumentTypeError(
-        "config must be BLOCK,OUTPUTS,K_UNROLL"
-        "[,VECTOR_WIDTH[,PREFETCH_B]], where PREFETCH_B is 0 or 1"
-    )
-
-
-def parse_tail_skinny_config(
-    value: str,
-) -> tuple[int, fused_add_multicast_skinny_gemm.SkinnyConfig]:
-    try:
-        values = [int(part) for part in value.split(",")]
-    except ValueError as error:
-        raise argparse.ArgumentTypeError(
-            "config must be M,BLOCK,OUTPUTS,K_UNROLL[,VECTOR_WIDTH[,PREFETCH_B]]"
-        ) from error
-    if len(values) == 4:
-        num_tokens, *config = values
-        return num_tokens, fused_add_multicast_skinny_gemm.SkinnyConfig(*config)
-    if len(values) == 5:
-        num_tokens, *config = values
-        return num_tokens, fused_add_multicast_skinny_gemm.SkinnyConfig(*config)
-    if len(values) == 6 and values[5] in (0, 1):
-        num_tokens, block, outputs, unroll, vector_width, prefetch = values
-        return num_tokens, fused_add_multicast_skinny_gemm.SkinnyConfig(
-            block,
-            outputs,
-            unroll,
-            vector_width,
-            bool(prefetch),
-        )
-    raise argparse.ArgumentTypeError(
-        "config must be M,BLOCK,OUTPUTS,K_UNROLL"
-        "[,VECTOR_WIDTH[,PREFETCH_B]], where PREFETCH_B is 0 or 1"
-    )
-
-
-def parse_args() -> argparse.Namespace:
-    parser = argparse.ArgumentParser(description=__doc__)
-    subparsers = parser.add_subparsers(dest="scope", required=True)
-
-    up_projection = subparsers.add_parser(
-        "up-projection",
-        help="Benchmark the isolated TP-local up-projection kernels.",
-    )
-    up_projection.add_argument(
-        "--backend",
-        choices=("dynamic", "skinny", "both"),
-        default="both",
-    )
-    up_projection.add_argument("--tp-size", type=int, default=16)
-    up_projection.add_argument(
-        "--num-tokens",
-        type=int,
-        nargs="+",
-        default=[*range(1, 9), 16],
-    )
-    up_projection.add_argument(
-        "--skinny-config",
-        type=parse_up_projection_config,
-        action="append",
-        help="Benchmark a static-M config for every selected token count.",
-    )
-    up_projection.add_argument("--cache-multiplier", type=float, default=2.0)
-    up_projection.add_argument("--max-weights", type=int, default=64)
-    up_projection.add_argument("--warmup-replays", type=int, default=10)
-    up_projection.add_argument("--samples", type=int, default=31)
-    up_projection.add_argument("--output", type=Path)
-
-    whole_tail = subparsers.add_parser(
-        "whole-tail",
-        help="Benchmark the distributed latent-MoE tail operator.",
-    )
-    whole_tail.add_argument(
-        "--backend",
-        choices=("reference", "fused", "both"),
-        default="both",
-    )
-    whole_tail.add_argument(
-        "--num-tokens",
-        type=int,
-        nargs="+",
-        default=[1, 5, 8, 16],
-    )
-    whole_tail.add_argument("--warmup-replays", type=int, default=20)
-    whole_tail.add_argument("--samples", type=int, default=51)
-    whole_tail.add_argument(
-        "--skinny-max-num-tokens",
-        type=int,
-        nargs="+",
-        help="Override the fused operator's static-M cutoff; use 0 for dynamic-only.",
-    )
-    whole_tail.add_argument(
-        "--skinny-config",
-        type=parse_tail_skinny_config,
-        action="append",
-        help="Override one static-M config for tuning.",
-    )
-    whole_tail.add_argument("--output", type=Path)
-    return parser.parse_args()
-
-
-def percentile(samples: Sequence[float], fraction: float) -> float:
-    ordered = sorted(samples)
-    position = fraction * (len(ordered) - 1)
-    lower = math.floor(position)
-    upper = math.ceil(position)
-    if lower == upper:
-        return ordered[lower]
-    upper_weight = position - lower
-    return ordered[lower] * (1.0 - upper_weight) + ordered[upper] * upper_weight
-
-
-def summarize(samples_us: Sequence[float]) -> dict[str, Any]:
-    mean_us = statistics.mean(samples_us)
-    return {
-        "median_us": statistics.median(samples_us),
-        "p10_us": percentile(samples_us, 0.1),
-        "p90_us": percentile(samples_us, 0.9),
-        "mean_us": mean_us,
-        "cv_pct": statistics.pstdev(samples_us) / mean_us * 100.0,
-        "samples_us": list(samples_us),
-    }
-
-
-def rotating_weight_count(
-    shard_size: int,
-    cache_multiplier: float,
-    limit: int,
-) -> int:
-    properties = torch.cuda.get_device_properties(
-        torch.accelerator.current_device_index()
-    )
-    weight_bytes = shard_size * LATENT_SIZE * 2
-    target_bytes = math.ceil(properties.L2_cache_size * cache_multiplier)
-    return max(2, min(limit, math.ceil(target_bytes / weight_bytes)))
-
-
-def capture_up_projection_graph(
-    launches: Sequence[Callable[[], None]],
-) -> torch.cuda.CUDAGraph:
-    for launch in launches:
-        launch()
-    torch.accelerator.synchronize()
-
-    graph = torch.cuda.CUDAGraph()
-    with torch.cuda.graph(graph):
-        for launch in launches:
-            launch()
-    torch.accelerator.synchronize()
-    return graph
-
-
-def benchmark_up_projection_graph(
-    graph: torch.cuda.CUDAGraph,
-    *,
-    operations_per_replay: int,
-    warmup_replays: int,
-    samples: int,
-) -> dict[str, Any]:
-    for _ in range(warmup_replays):
-        graph.replay()
-    torch.accelerator.synchronize()
-
-    samples_us = []
-    start = torch.cuda.Event(enable_timing=True)
-    end = torch.cuda.Event(enable_timing=True)
-    for _ in range(samples):
-        start.record()
-        graph.replay()
-        end.record()
-        end.synchronize()
-        samples_us.append(start.elapsed_time(end) * 1000.0 / operations_per_replay)
-    return summarize(samples_us)
-
-
-class DynamicKernel:
-    def __init__(
-        self,
-        shard_size: int,
-        mailbox: torch.Tensor,
-        shared_shard: torch.Tensor,
-    ) -> None:
-        self.shard_size = shard_size
-        self.mailbox = mailbox
-        self.mailbox_c = fused_add_multicast_gemm._as_cute(mailbox)
-        compile_latent = torch.empty(
-            (1, MAX_NUM_TOKENS, LATENT_SIZE),
-            dtype=torch.bfloat16,
-            device=mailbox.device,
-        )
-        compile_weight = torch.empty(
-            (1, shard_size, LATENT_SIZE),
-            dtype=torch.bfloat16,
-            device=mailbox.device,
-        )
-        cluster_size = math.prod(CLUSTER_SHAPE_MN)
-        max_active_clusters = utils.HardwareInfo().get_max_active_clusters(cluster_size)
-        self.compiled = fused_add_multicast_gemm.compile_kernel(
-            (MAX_NUM_TOKENS, shard_size, LATENT_SIZE, 1),
-            fused_add_multicast_gemm._as_cute(
-                compile_latent,
-                dynamic_m=True,
-            ),
-            fused_add_multicast_gemm._as_cute(compile_weight),
-            self.mailbox_c,
-            fused_add_multicast_gemm._as_cute(shared_shard),
-            HIDDEN_SIZE,
-            shard_size,
-            MMA_TILER_MN,
-            CLUSTER_SHAPE_MN,
-            max_active_clusters,
-            B_PRIME_STAGES,
-        )
-
-    def launch(
-        self,
-        latent: torch.Tensor,
-        weight: torch.Tensor,
-        shared_shard: torch.Tensor,
-    ) -> None:
-        stream = cuda.CUstream(torch.cuda.current_stream().cuda_stream)
-        self.compiled(
-            fused_add_multicast_gemm._as_cute(
-                latent.unsqueeze(0),
-                dynamic_m=True,
-            ),
-            fused_add_multicast_gemm._as_cute(weight.unsqueeze(0)),
-            self.mailbox_c,
-            fused_add_multicast_gemm._as_cute(shared_shard),
-            cutlass.Int64(latent.shape[0]),
-            cutlass.Int64(self.mailbox.data_ptr()),
-            stream,
-        )
-
-
-class SkinnyKernel:
-    def __init__(
-        self,
-        num_tokens: int,
-        shard_size: int,
-        config: fused_add_multicast_skinny_gemm.SkinnyConfig,
-    ) -> None:
-        self.compiled = fused_add_multicast_skinny_gemm.compile_kernel(
-            num_rows=num_tokens,
-            latent_dim=LATENT_SIZE,
-            hidden_dim=HIDDEN_SIZE,
-            shard_dim=shard_size,
-            config=config,
-        )
-
-    def launch(
-        self,
-        latent: torch.Tensor,
-        weight: torch.Tensor,
-        shared_shard: torch.Tensor,
-        mailbox: torch.Tensor,
-    ) -> None:
-        self.compiled(
-            fused_add_multicast_skinny_gemm._as_cute(latent),
-            fused_add_multicast_skinny_gemm._as_cute(weight),
-            fused_add_multicast_skinny_gemm._as_cute(shared_shard),
-            cutlass.Int64(mailbox.data_ptr()),
-            cuda.CUstream(torch.cuda.current_stream().cuda_stream),
-        )
-
-
-def check_up_projection_output(
-    actual: torch.Tensor,
-    latent: torch.Tensor,
-    weight: torch.Tensor,
-    shared_shard: torch.Tensor,
-) -> None:
-    gemm = F.linear(latent.float(), weight.float()).to(torch.bfloat16)
-    expected = (gemm.float() + shared_shard.float()).to(torch.bfloat16)
-    torch.testing.assert_close(actual, expected, atol=8e-2, rtol=3e-2)
-
-
-def make_up_projection_launches(
-    launch: Callable[[torch.Tensor, torch.Tensor, torch.Tensor], None],
-    latent: torch.Tensor,
-    weights: Sequence[torch.Tensor],
-    shared_shard: torch.Tensor,
-) -> list[Callable[[], None]]:
-    return [
-        lambda weight=weight: launch(latent, weight, shared_shard) for weight in weights
-    ]
-
-
-def benchmark_up_projection(args: argparse.Namespace) -> None:
-    if args.tp_size <= 0 or HIDDEN_SIZE % args.tp_size:
-        raise ValueError("TP size must be positive and divide the hidden size")
-    if any(not 1 <= num_tokens <= MAX_NUM_TOKENS for num_tokens in args.num_tokens):
-        raise ValueError("--num-tokens values must be in [1, 16]")
-    if args.cache_multiplier <= 0 or args.max_weights <= 0:
-        raise ValueError("cache multiplier and max weights must be positive")
-    if args.warmup_replays < 0 or args.samples <= 0:
-        raise ValueError("warmup replays must be nonnegative and samples positive")
-
-    torch.accelerator.set_device_index(0)
-    device = torch.device("cuda", 0)
-    if torch.cuda.get_device_capability(device)[0] != 10:
-        raise RuntimeError("Kimi K3 latent-MoE tail requires SM100")
-
-    shard_size = HIDDEN_SIZE // args.tp_size
-    weight_count = rotating_weight_count(
-        shard_size,
-        args.cache_multiplier,
-        args.max_weights,
-    )
-    torch.manual_seed(20260726)
-    weights = [
-        torch.randn(
-            (shard_size, LATENT_SIZE),
-            dtype=torch.bfloat16,
-            device=device,
-        )
-        / LATENT_SIZE**0.5
-        for _ in range(weight_count)
-    ]
-    mailbox = torch.empty(
-        (1, MAX_NUM_TOKENS, HIDDEN_SIZE),
-        dtype=torch.bfloat16,
-        device=device,
-    )
-    shared = torch.randn(
-        (MAX_NUM_TOKENS, HIDDEN_SIZE),
-        dtype=torch.bfloat16,
-        device=device,
-    )
-    shared_shard = shared[:, :shard_size]
-    use_dynamic = args.backend in ("dynamic", "both")
-    use_skinny = args.backend in ("skinny", "both")
-    dynamic_kernel = (
-        DynamicKernel(shard_size, mailbox, shared_shard) if use_dynamic else None
-    )
-
-    results = []
-    for num_tokens in args.num_tokens:
-        latent = torch.randn(
-            (num_tokens, LATENT_SIZE),
-            dtype=torch.bfloat16,
-            device=device,
-        )
-        result: dict[str, Any] = {"num_tokens": num_tokens}
-        if dynamic_kernel is not None:
-            launches = make_up_projection_launches(
-                dynamic_kernel.launch,
-                latent,
-                weights,
-                shared_shard,
-            )
-            graph = capture_up_projection_graph(launches)
-            result["dynamic"] = benchmark_up_projection_graph(
-                graph,
-                operations_per_replay=len(launches),
-                warmup_replays=args.warmup_replays,
-                samples=args.samples,
-            )
-            check_up_projection_output(
-                mailbox[0, :num_tokens, :shard_size],
-                latent,
-                weights[-1],
-                shared_shard[:num_tokens],
-            )
-        if use_skinny:
-            configs = args.skinny_config or [
-                fused_add_multicast_skinny_gemm.config_for_m(
-                    num_tokens,
-                    shard_size,
-                )
-            ]
-            skinny_results = []
-            for config in configs:
-                skinny_kernel = SkinnyKernel(num_tokens, shard_size, config)
-
-                def launch_skinny(
-                    latent: torch.Tensor,
-                    weight: torch.Tensor,
-                    shared_shard: torch.Tensor,
-                    *,
-                    skinny_kernel: SkinnyKernel = skinny_kernel,
-                    num_tokens: int = num_tokens,
-                ) -> None:
-                    skinny_kernel.launch(
-                        latent,
-                        weight,
-                        shared_shard[:num_tokens],
-                        mailbox,
-                    )
-
-                launches = make_up_projection_launches(
-                    launch_skinny,
-                    latent,
-                    weights,
-                    shared_shard,
-                )
-                graph = capture_up_projection_graph(launches)
-                timing = benchmark_up_projection_graph(
-                    graph,
-                    operations_per_replay=len(launches),
-                    warmup_replays=args.warmup_replays,
-                    samples=args.samples,
-                )
-                check_up_projection_output(
-                    mailbox[0, :num_tokens, :shard_size],
-                    latent,
-                    weights[-1],
-                    shared_shard[:num_tokens],
-                )
-                skinny_results.append(
-                    {
-                        "config": asdict(config),
-                        **timing,
-                    }
-                )
-            result["skinny"] = skinny_results
-        results.append(result)
-
-    properties = torch.cuda.get_device_properties(device)
-    report = {
-        "scope": "up-projection",
-        "device": properties.name,
-        "compute_capability": list(torch.cuda.get_device_capability(device)),
-        "tp_size": args.tp_size,
-        "shard_size": shard_size,
-        "weight_count": weight_count,
-        "cache_multiplier": args.cache_multiplier,
-        "warmup_replays": args.warmup_replays,
-        "samples": args.samples,
-        "results": results,
-    }
-    rendered = json.dumps(report, indent=2)
-    print(rendered, flush=True)
-    if args.output is not None:
-        args.output.parent.mkdir(parents=True, exist_ok=True)
-        args.output.write_text(rendered + "\n", encoding="utf-8")
-
-
-def capture_tail_graph(
-    operation: Callable[[], torch.Tensor],
-    cpu_group: dist.ProcessGroup,
-) -> tuple[torch.cuda.CUDAGraph, torch.Tensor]:
-    for _ in range(3):
-        dist.barrier(group=cpu_group)
-        output = operation()
-    torch.accelerator.synchronize()
-
-    dist.barrier(group=cpu_group)
-    graph = torch.cuda.CUDAGraph()
-    with torch.cuda.graph(graph):
-        output = operation()
-    torch.accelerator.synchronize()
-    return graph, output
-
-
-def benchmark_tail_graph(
-    graph: torch.cuda.CUDAGraph,
-    *,
-    warmup_replays: int,
-    samples: int,
-    device_group: dist.ProcessGroup,
-    cpu_group: dist.ProcessGroup,
-) -> dict[str, Any]:
-    for _ in range(warmup_replays):
-        graph.replay()
-    torch.accelerator.synchronize()
-
-    dist.barrier(group=cpu_group)
-    starts = [torch.cuda.Event(enable_timing=True) for _ in range(samples + 1)]
-    ends = [torch.cuda.Event(enable_timing=True) for _ in range(samples + 1)]
-    for start, end in zip(starts, ends):
-        start.record()
-        graph.replay()
-        end.record()
-    torch.accelerator.synchronize()
-
-    samples_us = torch.tensor(
-        [start.elapsed_time(end) * 1000.0 for start, end in zip(starts, ends)],
-        dtype=torch.float64,
-        device=torch.accelerator.current_device_index(),
-    )
-    dist.all_reduce(samples_us, op=dist.ReduceOp.MAX, group=device_group)
-    return summarize(samples_us[1:].tolist())
-
-
-def make_inputs(
-    num_tokens: int,
-    rank: int,
-    device: torch.device,
-) -> tuple[torch.Tensor, torch.Tensor]:
-    torch.manual_seed(20260726 + 100 * num_tokens + rank)
-    routed = torch.randn(
-        (num_tokens, LATENT_SIZE),
-        dtype=torch.bfloat16,
-        device=device,
-    ).mul_(0.01)
-    shared = torch.randn(
-        (num_tokens, HIDDEN_SIZE),
-        dtype=torch.bfloat16,
-        device=device,
-    )
-    return routed, shared
-
-
-def make_reference(
-    routed: torch.Tensor,
-    shared: torch.Tensor,
-    rms_weight: torch.Tensor,
-    up_weight: torch.Tensor,
-    device_group: dist.ProcessGroup,
-) -> Callable[[], torch.Tensor]:
-    routed_workspace = torch.empty_like(routed)
-    shared_workspace = torch.empty_like(shared)
-
-    def reference() -> torch.Tensor:
-        routed_workspace.copy_(routed)
-        dist.all_reduce(routed_workspace, group=device_group)
-        normalized = F.rms_norm(
-            routed_workspace,
-            (LATENT_SIZE,),
-            rms_weight,
-            RMS_EPS,
-        )
-        projected = F.linear(normalized, up_weight)
-        shared_workspace.copy_(shared)
-        dist.all_reduce(shared_workspace, group=device_group)
-        return projected.add(shared_workspace)
-
-    return reference
-
-
-def check_fused_output(
-    fused_output: torch.Tensor,
-    reference: Callable[[], torch.Tensor],
-    cpu_group: dist.ProcessGroup,
-) -> None:
-    dist.barrier(group=cpu_group)
-    expected = reference()
-    torch.testing.assert_close(fused_output, expected, atol=8e-2, rtol=3e-2)
-
-
-def benchmark_whole_tail(args: argparse.Namespace) -> None:
-    if any(not 1 <= num_tokens <= 16 for num_tokens in args.num_tokens):
-        raise ValueError("--num-tokens values must be in [1, 16]")
-    if args.warmup_replays < 0 or args.samples <= 0:
-        raise ValueError("warmup replays must be nonnegative and samples positive")
-    if args.skinny_max_num_tokens is not None and any(
-        not 0 <= cutoff <= 8 for cutoff in args.skinny_max_num_tokens
-    ):
-        raise ValueError("--skinny-max-num-tokens must be in [0, 8]")
-    skinny_configs = dict(args.skinny_config or ())
-    if len(skinny_configs) != len(args.skinny_config or ()):
-        raise ValueError("--skinny-config must not repeat an M value")
-    if any(not 1 <= num_tokens <= 8 for num_tokens in skinny_configs):
-        raise ValueError("--skinny-config M values must be in [1, 8]")
-    if not {"RANK", "WORLD_SIZE", "LOCAL_RANK"} <= os.environ.keys():
-        raise RuntimeError("launch this benchmark with torchrun")
-
-    rank = int(os.environ["RANK"])
-    world_size = int(os.environ["WORLD_SIZE"])
-    local_rank = int(os.environ["LOCAL_RANK"])
-    device = torch.device("cuda", local_rank)
-    torch.accelerator.set_device_index(device)
-    init_distributed_environment()
-    if world_size > 8:
-        set_custom_all_reduce(False)
-    initialize_model_parallel(tensor_model_parallel_size=world_size)
-    device_group = get_tp_group().device_group
-    cpu_group = dist.new_group(backend="gloo")
-
-    if torch.cuda.get_device_capability(device)[0] != 10:
-        raise RuntimeError("Kimi K3 latent-MoE tail requires SM100")
-
-    torch.manual_seed(20260726)
-    rms_weight = 1 + 0.1 * torch.randn(
-        LATENT_SIZE,
-        dtype=torch.bfloat16,
-        device=device,
-    )
-    up_weight = (
-        torch.randn(
-            (HIDDEN_SIZE, LATENT_SIZE),
-            dtype=torch.bfloat16,
-            device=device,
-        )
-        / LATENT_SIZE**0.5
-    )
-
-    use_reference = args.backend in ("reference", "both")
-    use_fused = args.backend in ("fused", "both")
-    fused_ops = []
-    if use_fused:
-        production_config_for_m = fused_add_multicast_skinny_gemm.config_for_m
-
-        def config_for_m(
-            num_rows: int,
-            shard_dim: int = 896,
-        ) -> fused_add_multicast_skinny_gemm.SkinnyConfig:
-            config = skinny_configs.get(num_rows)
-            if config is not None:
-                return config
-            return production_config_for_m(num_rows, shard_dim)
-
-        fused_add_multicast_skinny_gemm.config_for_m = config_for_m
-        cutoffs = args.skinny_max_num_tokens or [latent_moe_tail._SKINNY_MAX_NUM_TOKENS]
-        for cutoff in cutoffs:
-            latent_moe_tail._SKINNY_MAX_NUM_TOKENS = cutoff
-            latent_moe_tail.KimiK3LatentMoETailOp._instances.clear()
-            fused_ops.append(
-                (
-                    cutoff,
-                    latent_moe_tail.KimiK3LatentMoETailOp.initialize(
-                        hidden_size=HIDDEN_SIZE,
-                        latent_size=LATENT_SIZE,
-                        dtype=torch.bfloat16,
-                        device=device,
-                        rms_eps=RMS_EPS,
-                    ),
-                )
-            )
-        cutedsl_warmup()
-
-    results = []
-    for num_tokens in args.num_tokens:
-        routed, shared = make_inputs(num_tokens, rank, device)
-        reference = make_reference(
-            routed,
-            shared,
-            rms_weight,
-            up_weight,
-            device_group,
-        )
-        result: dict[str, Any] = {"num_tokens": num_tokens}
-        if use_reference:
-            reference_graph, _ = capture_tail_graph(reference, cpu_group)
-            result["reference"] = benchmark_tail_graph(
-                reference_graph,
-                warmup_replays=args.warmup_replays,
-                samples=args.samples,
-                device_group=device_group,
-                cpu_group=cpu_group,
-            )
-        for cutoff, fused_op in fused_ops:
-
-            def fused(
-                routed: torch.Tensor = routed,
-                shared: torch.Tensor = shared,
-                fused_op: latent_moe_tail.KimiK3LatentMoETailOp = fused_op,
-            ) -> torch.Tensor:
-                return fused_op(routed, shared, rms_weight, up_weight)
-
-            fused_graph, fused_output = capture_tail_graph(fused, cpu_group)
-            fused_key = "fused" if len(fused_ops) == 1 else f"fused_skinny_max_{cutoff}"
-            result[fused_key] = benchmark_tail_graph(
-                fused_graph,
-                warmup_replays=args.warmup_replays,
-                samples=args.samples,
-                device_group=device_group,
-                cpu_group=cpu_group,
-            )
-            check_fused_output(fused_output, reference, cpu_group)
-            if "reference" in result:
-                speedup = (
-                    result["reference"]["median_us"] / result[fused_key]["median_us"]
-                )
-                if len(fused_ops) == 1:
-                    result["speedup"] = speedup
-                else:
-                    result[f"{fused_key}_speedup"] = speedup
-        results.append(result)
-
-    properties = torch.cuda.get_device_properties(device)
-    report = {
-        "scope": "whole-tail",
-        "device": properties.name,
-        "compute_capability": list(torch.cuda.get_device_capability(device)),
-        "world_size": world_size,
-        "torch_version": torch.__version__,
-        "cuda_version": torch.version.cuda,
-        "warmup_replays": args.warmup_replays,
-        "samples": args.samples,
-        "skinny_max_num_tokens": [cutoff for cutoff, _ in fused_ops],
-        "skinny_configs": {
-            str(num_tokens): asdict(config)
-            for num_tokens, config in skinny_configs.items()
-        },
-        "timing_scope": {
-            "reference": (
-                "two input copies, two AllReduces, RMSNorm, full replicated "
-                "up-projection GEMM, and final add"
-            ),
-            "fused": (
-                "routed AllReduce/RMSNorm plus shared ReduceScatter, sharded "
-                "up-projection/multicast, and Lamport copy"
-            ),
-        },
-        "results": results,
-    }
-    if rank == 0:
-        rendered = json.dumps(report, indent=2)
-        print(rendered, flush=True)
-        if args.output is not None:
-            args.output.parent.mkdir(parents=True, exist_ok=True)
-            args.output.write_text(rendered + "\n", encoding="utf-8")
-
-    dist.barrier(group=cpu_group)
-
-
-def main() -> None:
-    args = parse_args()
-    if args.scope == "up-projection":
-        benchmark_up_projection(args)
-        return
-
-    from vllm.config import VllmConfig, set_current_vllm_config
-
-    with set_current_vllm_config(VllmConfig()):
-        benchmark_whole_tail(args)
-
-
-if __name__ == "__main__":
-    main()

benchmarks/kernels/benchmark_kimi_k3_sp_collectives.py

@@ -1,239 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-import argparse
-import json
-import os
-import statistics
-from collections.abc import Callable
-
-import torch
-import torch.distributed as dist
-
-import vllm._custom_ops as ops
-from vllm.distributed.device_communicators.custom_all_reduce import CustomAllreduce
-
-
-def parse_args() -> argparse.Namespace:
-    parser = argparse.ArgumentParser()
-    parser.add_argument("--tokens", type=int, nargs="+", default=[8, 32, 128, 1024])
-    parser.add_argument("--hidden-size", type=int, default=7168)
-    parser.add_argument("--graph-repeats", type=int, default=20)
-    parser.add_argument("--warmup-replays", type=int, default=5)
-    parser.add_argument("--samples", type=int, default=15)
-    return parser.parse_args()
-
-
-def capture_graph(op: Callable[[], None], repeats: int) -> torch.cuda.CUDAGraph:
-    stream = torch.cuda.Stream()
-    stream.wait_stream(torch.cuda.current_stream())
-    with torch.cuda.stream(stream):
-        for _ in range(3):
-            op()
-    stream.synchronize()
-
-    graph = torch.cuda.CUDAGraph()
-    with torch.cuda.graph(graph, stream=stream):
-        for _ in range(repeats):
-            op()
-    torch.cuda.current_stream().wait_stream(stream)
-    return graph
-
-
-def max_rank_graph_time(
-    graph: torch.cuda.CUDAGraph,
-    repeats: int,
-    warmup_replays: int,
-    samples: int,
-    device_group: dist.ProcessGroup,
-    cpu_group: dist.ProcessGroup,
-) -> float:
-    for _ in range(warmup_replays):
-        graph.replay()
-    torch.accelerator.synchronize()
-
-    timings = []
-    start = torch.cuda.Event(enable_timing=True)
-    end = torch.cuda.Event(enable_timing=True)
-    for _ in range(samples):
-        dist.barrier(group=cpu_group)
-        start.record()
-        graph.replay()
-        end.record()
-        end.synchronize()
-        elapsed = torch.tensor(
-            start.elapsed_time(end) / repeats,
-            dtype=torch.float64,
-            device=torch.accelerator.current_device_index(),
-        )
-        dist.all_reduce(elapsed, op=dist.ReduceOp.MAX, group=device_group)
-        timings.append(elapsed.item())
-    return statistics.median(timings)
-
-
-def check_outputs(
-    comm: CustomAllreduce,
-    local: torch.Tensor,
-    reduce_input: torch.Tensor,
-    device_group: dist.ProcessGroup,
-) -> None:
-    expected_gather = torch.empty(
-        (local.shape[0] * dist.get_world_size(), local.shape[1]),
-        dtype=local.dtype,
-        device=local.device,
-    )
-    dist.all_gather_into_tensor(expected_gather, local, group=device_group)
-    gathered = comm.custom_all_gather(local)
-    assert gathered is not None
-    torch.testing.assert_close(gathered, expected_gather)
-
-    expected_scatter = torch.empty_like(local)
-    dist.reduce_scatter_tensor(
-        expected_scatter,
-        reduce_input.clone(),
-        group=device_group,
-    )
-    scattered = comm.custom_reduce_scatter(reduce_input)
-    assert scattered is not None
-    torch.testing.assert_close(scattered, expected_scatter)
-
-
-def benchmark_shape(
-    comm: CustomAllreduce,
-    global_tokens: int,
-    hidden_size: int,
-    graph_repeats: int,
-    warmup_replays: int,
-    samples: int,
-    device_group: dist.ProcessGroup,
-    cpu_group: dist.ProcessGroup,
-) -> dict[str, float | int]:
-    world_size = dist.get_world_size()
-    rank = dist.get_rank()
-    padded_tokens = (global_tokens + world_size - 1) // world_size * world_size
-    local_tokens = padded_tokens // world_size
-    local = torch.full(
-        (local_tokens, hidden_size),
-        rank + 1,
-        dtype=torch.bfloat16,
-        device=torch.accelerator.current_device_index(),
-    )
-    reduce_input = torch.full(
-        (padded_tokens, hidden_size),
-        rank + 1,
-        dtype=torch.bfloat16,
-        device=local.device,
-    )
-    check_outputs(comm, local, reduce_input, device_group)
-
-    custom_gather_out = torch.empty(
-        (padded_tokens, hidden_size),
-        dtype=local.dtype,
-        device=local.device,
-    )
-    custom_scatter_out = torch.empty_like(local)
-    nccl_gather_out = torch.empty_like(custom_gather_out)
-    nccl_scatter_out = torch.empty_like(local)
-
-    def custom_ag() -> None:
-        ops.mnnvl_lamport_all_gather(
-            comm._ptr,
-            local,
-            custom_gather_out,
-            comm.mnnvl_lamport_ag_local_ptr,
-            comm.mnnvl_lamport_ag_multicast_ptr,
-            comm.mnnvl_lamport_ag_epoch_ptr,
-            comm.mnnvl_buffer_size,
-        )
-
-    def custom_rs() -> None:
-        ops.mnnvl_lamport_reduce_scatter(
-            comm._ptr,
-            reduce_input,
-            custom_scatter_out,
-            comm.mnnvl_lamport_rs_local_ptr,
-            comm.mnnvl_lamport_rs_epoch_ptr,
-            comm.mnnvl_buffer_size,
-        )
-
-    def nccl_ag() -> None:
-        dist.all_gather_into_tensor(nccl_gather_out, local, group=device_group)
-
-    def nccl_rs() -> None:
-        dist.reduce_scatter_tensor(
-            nccl_scatter_out,
-            reduce_input,
-            group=device_group,
-        )
-
-    graphs = {
-        "custom_ag_us": capture_graph(custom_ag, graph_repeats),
-        "nccl_ag_us": capture_graph(nccl_ag, graph_repeats),
-        "custom_rs_us": capture_graph(custom_rs, graph_repeats),
-        "nccl_rs_us": capture_graph(nccl_rs, graph_repeats),
-    }
-    times = {
-        name: max_rank_graph_time(
-            graph,
-            graph_repeats,
-            warmup_replays,
-            samples,
-            device_group,
-            cpu_group,
-        )
-        * 1000
-        for name, graph in graphs.items()
-    }
-    torch.testing.assert_close(custom_gather_out, nccl_gather_out)
-    torch.testing.assert_close(custom_scatter_out, nccl_scatter_out)
-    return {
-        "global_tokens": global_tokens,
-        "padded_tokens": padded_tokens,
-        "local_bytes": local.nbytes,
-        "full_bytes": reduce_input.nbytes,
-        **times,
-        "ag_speedup": times["nccl_ag_us"] / times["custom_ag_us"],
-        "rs_speedup": times["nccl_rs_us"] / times["custom_rs_us"],
-    }
-
-
-def main() -> None:
-    args = parse_args()
-    local_rank = int(os.environ["LOCAL_RANK"])
-    torch.accelerator.set_device_index(local_rank)
-    dist.init_process_group("nccl")
-    device_group = dist.group.WORLD
-    cpu_group = dist.new_group(backend="gloo")
-
-    comm = CustomAllreduce(
-        group=cpu_group,
-        device=torch.device("cuda", local_rank),
-    )
-    assert not comm.disabled
-    assert comm.world_size == 16
-    assert comm.mnnvl_only
-    assert comm.mnnvl_multicast_ptr
-
-    results = [
-        benchmark_shape(
-            comm,
-            tokens,
-            args.hidden_size,
-            args.graph_repeats,
-            args.warmup_replays,
-            args.samples,
-            device_group,
-            cpu_group,
-        )
-        for tokens in args.tokens
-    ]
-    if dist.get_rank() == 0:
-        print(json.dumps(results, indent=2), flush=True)
-
-    comm.close()
-    dist.destroy_process_group(cpu_group)
-    dist.destroy_process_group()
-
-
-if __name__ == "__main__":
-    main()

benchmarks/kernels/benchmark_paged_attention.py

@@ -19,11 +19,13 @@ from vllm.utils.torch_utils import (
 logger = init_logger(__name__)
 
 NUM_BLOCKS = 128 * 1024
+PARTITION_SIZE = 512
 PARTITION_SIZE_ROCM = 256
 
 
 @torch.inference_mode()
 def main(
+    version: str,
     num_seqs: int,
     seq_len: int,
     num_query_heads: int,
@@ -80,20 +82,27 @@ def main(
 
     # Prepare for the paged attention kernel.
     output = torch.empty_like(query)
-    num_partitions = (max_seq_len + PARTITION_SIZE_ROCM - 1) // PARTITION_SIZE_ROCM
+    if version == "v2":
-    tmp_output = torch.empty(
+        if current_platform.is_rocm():
-        size=(num_seqs, num_query_heads, num_partitions, head_size),
+            global PARTITION_SIZE
-        dtype=output.dtype,
+            if not args.custom_paged_attn and not current_platform.is_navi():
-        device=output.device,
+                PARTITION_SIZE = 1024
-    )
+            else:
-    exp_sums = torch.empty(
+                PARTITION_SIZE = PARTITION_SIZE_ROCM
-        size=(num_seqs, num_query_heads, num_partitions),
+        num_partitions = (max_seq_len + PARTITION_SIZE - 1) // PARTITION_SIZE
-        dtype=torch.float32,
+        tmp_output = torch.empty(
-        device=output.device,
+            size=(num_seqs, num_query_heads, num_partitions, head_size),
-    )
+            dtype=output.dtype,
-    max_logits = torch.empty_like(exp_sums)
+            device=output.device,
+        )
+        exp_sums = torch.empty(
+            size=(num_seqs, num_query_heads, num_partitions),
+            dtype=torch.float32,
+            device=output.device,
+        )
+        max_logits = torch.empty_like(exp_sums)
 
-    def run_benchmark(num_iters: int, profile: bool = False) -> float:
+    def run_cuda_benchmark(num_iters: int, profile: bool = False) -> float:
         torch.accelerator.synchronize()
         if profile:
             torch.cuda.cudart().cudaProfilerStart()
@@ -103,26 +112,67 @@ def main(
         k_scale = v_scale = torch.tensor(1.0, dtype=torch.float32, device=device)
 
         for _ in range(num_iters):
-            ops.paged_attention_rocm(
+            if version == "v1":
-                output,
+                ops.paged_attention_v1(
-                exp_sums,
+                    output,
-                max_logits,
+                    query,
-                tmp_output,
+                    key_cache,
-                query,
+                    value_cache,
-                key_cache,
+                    num_kv_heads,
-                value_cache,
+                    scale,
-                num_kv_heads,
+                    block_tables,
-                scale,
+                    seq_lens,
-                block_tables,
+                    block_size,
-                seq_lens,
+                    max_seq_len,
-                None,
+                    alibi_slopes,
-                block_size,
+                    kv_cache_dtype,
-                max_seq_len,
+                    k_scale,
-                alibi_slopes,
+                    v_scale,
-                kv_cache_dtype,
+                )
-                k_scale,
+            elif version == "v2":
-                v_scale,
+                if not args.custom_paged_attn:
-            )
+                    ops.paged_attention_v2(
+                        output,
+                        exp_sums,
+                        max_logits,
+                        tmp_output,
+                        query,
+                        key_cache,
+                        value_cache,
+                        num_kv_heads,
+                        scale,
+                        block_tables,
+                        seq_lens,
+                        block_size,
+                        max_seq_len,
+                        alibi_slopes,
+                        kv_cache_dtype,
+                        k_scale,
+                        v_scale,
+                    )
+                else:
+                    ops.paged_attention_rocm(
+                        output,
+                        exp_sums,
+                        max_logits,
+                        tmp_output,
+                        query,
+                        key_cache,
+                        value_cache,
+                        num_kv_heads,
+                        scale,
+                        block_tables,
+                        seq_lens,
+                        None,
+                        block_size,
+                        max_seq_len,
+                        alibi_slopes,
+                        kv_cache_dtype,
+                        k_scale,
+                        v_scale,
+                    )
+            else:
+                raise ValueError(f"Invalid version: {version}")
         torch.accelerator.synchronize()
 
         end_time = time.perf_counter()
@@ -132,6 +182,7 @@ def main(
 
     # Warmup.
     print("Warming up...")
+    run_benchmark = run_cuda_benchmark
     run_benchmark(num_iters=3, profile=False)
 
     # Benchmark.
@@ -144,13 +195,12 @@ def main(
 
 if __name__ == "__main__":
     logger.warning(
-        "This script benchmarks the ROCm paged attention kernel. "
+        "This script benchmarks the paged attention kernel. "
         "By default this is no longer used in vLLM inference."
     )
-    if not current_platform.is_rocm():
-        raise RuntimeError("This benchmark requires the ROCm platform.")
 
     parser = FlexibleArgumentParser(description="Benchmark the paged attention kernel.")
+    parser.add_argument("--version", type=str, choices=["v1", "v2"], default="v2")
     parser.add_argument("--batch-size", type=int, default=8)
     parser.add_argument("--seq-len", type=int, default=4096)
     parser.add_argument("--num-query-heads", type=int, default=64)
@@ -158,7 +208,7 @@ if __name__ == "__main__":
     parser.add_argument(
         "--head-size",
         type=int,
-        choices=[64, 128],
+        choices=[64, 80, 96, 112, 120, 128, 192, 256],
         default=128,
     )
     parser.add_argument("--block-size", type=int, choices=[16, 32], default=16)
@@ -174,7 +224,11 @@ if __name__ == "__main__":
         choices=["auto", "fp8", "fp8_e5m2", "fp8_e4m3"],
         default="auto",
         help="Data type for kv cache storage. If 'auto', will use model "
-        "data type. ROCm (AMD GPU) supports fp8 (=fp8_e4m3)",
+        "data type. CUDA 11.8+ supports fp8 (=fp8_e4m3) and fp8_e5m2. "
+        "ROCm (AMD GPU) supports fp8 (=fp8_e4m3)",
+    )
+    parser.add_argument(
+        "--custom-paged-attn", action="store_true", help="Use custom paged attention"
     )
     args = parser.parse_args()
     print(args)
@@ -182,6 +236,7 @@ if __name__ == "__main__":
     if args.num_query_heads % args.num_kv_heads != 0:
         raise ValueError("num_query_heads must be divisible by num_kv_heads")
     main(
+        version=args.version,
         num_seqs=args.batch_size,
         seq_len=args.seq_len,
         num_query_heads=args.num_query_heads,