[docs] Fix the docs build (#48008 )

Signed-off-by: Martin Hickey <martin.hickey@ie.ibm.com> Co-authored-by: Harry Mellor <19981378+hmellor@users.noreply.github.com> (cherry picked from commit 6cf7b26bd4)
Pin PyNvVideoCodec to tested 2.0.4 wheel (#48056 )
2026-07-11 11:25:04 +01:00 · 2026-07-09 12:43:21 -07:00 · 2026-07-09 03:24:39 -07:00 · 2026-07-09 03:24:39 -07:00 · 2026-07-08 21:12:19 -07:00 · 2026-07-08 21:12:19 -07:00
1161 changed files with 10986 additions and 65120 deletions
@@ -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"
@@ -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:
-      - refresh-rocm-base-amd
+    depends_on: []
    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
@@ -141,22 +141,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
@@ -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,7 +39,7 @@ 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:
@@ -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:
@@ -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
@@ -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
-  echo "Image found"
-else
+if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64) ]]; then
  echo "Image not found, proceeding with build..."
-  # build for arm64 GPU targets: Grace/GH200 (sm_90),
-  # Blackwell/Thor (sm_100/sm_103/sm_110), and DGX Spark/GB10
-  # (sm_121, family-covered by 12.0 under CUDA 13)
-  docker build --file docker/Dockerfile \
-    --platform linux/arm64 \
-    --build-arg max_jobs=16 \
-    --build-arg nvcc_threads=4 \
-    --build-arg torch_cuda_arch_list="9.0 10.0 11.0 12.0" \
-    --build-arg USE_SCCACHE=1 \
-    --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
-    --tag "$IMAGE" \
-    --target test \
-    --progress plain .
-  # push
-  docker push "$IMAGE"
+else
+  echo "Image found"
+  exit 0
 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
@@ -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
-  echo "Image found"
-else
+if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-cpu) ]]; then
  echo "Image not found, proceeding with build..."
-  # build
-  docker build --file docker/Dockerfile.cpu \
-    --build-arg max_jobs=16 \
-    --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
-    --build-arg VLLM_CPU_X86=true \
-    --tag "$IMAGE" \
-    --target vllm-test \
-    --progress plain .
-  # push
-  docker push "$IMAGE"
+else
+  echo "Image found"
+  exit 0
 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
@@ -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
-  echo "Image found"
-else
+if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64-cpu) ]]; then
  echo "Image not found, proceeding with build..."
-  # build
-  docker build --file docker/Dockerfile.cpu \
-    --build-arg max_jobs=16 \
-    --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
-    --tag "$IMAGE" \
-    --target vllm-test \
-    --progress plain .
-  # push
-  docker push "$IMAGE"
+else
+  echo "Image found"
+  exit 0
 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
@@ -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
-  echo "Image found"
-else
+if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-hpu) ]]; then
  echo "Image not found, proceeding with build..."
-  # build
-  docker build \
-    --file tests/pytorch_ci_hud_benchmark/Dockerfile.hpu \
-    --build-arg max_jobs=16 \
-    --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
-    --tag "$IMAGE" \
-    --progress plain \
-    https://github.com/vllm-project/vllm-gaudi.git
-  # push
-  docker push "$IMAGE"
+else
+  echo "Image found"
+  exit 0
 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
@@ -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"
@@ -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
-  echo "Image found"
-else
+if ! docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu &> /dev/null; then
  echo "Image not found, proceeding with build..."
-  # build
-  docker build \
-    --file docker/Dockerfile.xpu \
-    --build-arg max_jobs=16 \
-    --build-arg buildkite_commit="$BUILDKITE_COMMIT" \
-    --tag "$IMAGE" \
-    --progress plain .
-  # push
-  docker push "$IMAGE"
+else
+  echo "Image found"
+  exit 0
 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
@@ -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
@@ -86,7 +86,7 @@ steps:
      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
@@ -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,7 +33,7 @@ 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
@@ -72,7 +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 &&
+      pytest -v -s v1/sample/test_logprobs_e2e.py'

 - label: Basic Models Tests (Initialization)
  timeout_in_minutes: 60
@@ -152,7 +154,7 @@ steps:

 - label: Regression
  key: regression
-  timeout_in_minutes: 50
+  timeout_in_minutes: 30
  device: intel_gpu
  agent_tags:
    label: production
@@ -186,7 +188,7 @@ steps:

 - 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 +224,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
@@ -4,7 +4,7 @@ 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
@@ -4,7 +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
@@ -29,7 +29,7 @@ steps:

 - label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
  key: multi-modal-models-standard-2-qwen3-gemma
-  timeout_in_minutes: 70
+  timeout_in_minutes: 45
  device: intel_gpu
  agent_tags:
    label: production
@@ -52,7 +52,7 @@ steps:

 - label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
  key: multi-modal-models-standard-3-llava-qwen2-vl
-  timeout_in_minutes: 65
+  timeout_in_minutes: 45
  device: intel_gpu
  agent_tags:
    label: production
@@ -100,7 +100,7 @@ steps:

 - label: Multi-Modal Processor # 44min
  key: multi-modal-processor
-  timeout_in_minutes: 60
+  timeout_in_minutes: 45
  device: intel_gpu
  agent_tags:
    label: production
@@ -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
@@ -76,7 +76,7 @@ steps:
  - label: "XPU V1 test"
    depends_on:
      - image-build-xpu
-    timeout_in_minutes: 70
+    timeout_in_minutes: 30
    device: intel_gpu
    agent_tags:
      label: production
@@ -99,13 +99,12 @@ steps:
        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
@@ -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
@@ -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(
@@ -44,7 +44,7 @@ steps:
          - "mkdir artifacts"
          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)" release-wheels'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
        env:
          DOCKER_BUILDKIT: "1"

@@ -58,7 +58,7 @@ steps:
          - "mkdir artifacts"
          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)" release-wheels'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
        env:
          DOCKER_BUILDKIT: "1"

@@ -72,53 +72,10 @@ steps:
          - "mkdir artifacts"
          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)" release-wheels'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
        env:
          DOCKER_BUILDKIT: "1"

-      - label: "Build wheel - macOS arm64 - CPU"
-        depends_on: ~
-        id: build-wheel-macos-arm64-cpu
-        agents:
-          queue: macmini
-        timeout_in_minutes: 90
-        commands:
-          - "bash .buildkite/scripts/build-macos-wheel.sh"
-          - "mkdir -p artifacts/transfer"
-          - "basename artifacts/dist/*.whl > artifacts/transfer/wheel.name"
-          - "shasum -a 256 artifacts/dist/*.whl > artifacts/transfer/wheel.sha256"
-          - "split -b 1m artifacts/dist/*.whl artifacts/transfer/wheel.part."
-        env:
-          BUILDKITE_ARTIFACT_UPLOAD_CONCURRENCY: "1"
-          BUILDKITE_NO_HTTP2: "true"
-        artifact_paths:
-          - "artifacts/transfer/*"
-
-      - label: "Upload wheel - macOS arm64 - CPU"
-        depends_on: build-wheel-macos-arm64-cpu
-        id: upload-wheel-macos-arm64-cpu
-        agents:
-          queue: small_cpu_queue_release
-        timeout_in_minutes: 15
-        commands:
-          - 'buildkite-agent artifact download "artifacts/transfer/*" .'
-          - "mkdir -p artifacts/dist artifacts/reassembled"
-          - 'wheel_name=$$(cat artifacts/transfer/wheel.name)'
-          - '[[ "$$wheel_name" =~ ^[A-Za-z0-9][A-Za-z0-9._+-]*\.whl$$ && "$$wheel_name" != *..* ]]'
-          - 'expected_sha=$$(awk ''{print $$1}'' artifacts/transfer/wheel.sha256)'
-          - '[[ "$$expected_sha" =~ ^[0-9a-f]{64}$$ ]]'
-          - 'cat artifacts/transfer/wheel.part.* > artifacts/reassembled/wheel'
-          - 'actual_sha=$$(shasum -a 256 artifacts/reassembled/wheel | awk ''{print $$1}'')'
-          - '[[ "$$actual_sha" == "$$expected_sha" ]]'
-          - 'mv artifacts/reassembled/wheel "artifacts/dist/$$wheel_name"'
-          - "aws sts get-caller-identity"
-          - "VLLM_WHEEL_PLATFORM=macos bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
-        plugins:
-          - aws-assume-role-with-web-identity#v1.6.0:
-              role-arn: arn:aws:iam::936637512419:role/vllm-release-macos-wheel-uploader
-              region: us-west-2
-
      - label: "Build wheel - x86_64 - CUDA 12.9"
        depends_on: ~
        id: build-wheel-x86-cuda-12-9
@@ -129,7 +86,7 @@ steps:
          - "mkdir artifacts"
          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)" release-wheels'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
        env:
          DOCKER_BUILDKIT: "1"

@@ -143,7 +100,7 @@ steps:
          - "mkdir artifacts"
          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)" release-wheels'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
        env:
          DOCKER_BUILDKIT: "1"

@@ -157,7 +114,7 @@ steps:
          - "mkdir artifacts"
          - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)" release-wheels'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
        env:
          DOCKER_BUILDKIT: "1"

@@ -205,7 +162,7 @@ steps:
          # re-tag to default image tag and push, just in case arm64 build fails
          - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m) public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)" release-images'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)"'

      - label: "Build release image - aarch64 - CUDA 13.0"
        depends_on: ~
@@ -230,7 +187,7 @@ steps:
              --progress plain \
              -f docker/Dockerfile .
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)" release-images'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)"'

      - label: "Build release image - x86_64 - CUDA 12.9"
        depends_on: ~
@@ -257,7 +214,7 @@ steps:
          # re-tag to default image tag and push, just in case arm64 build fails
          - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129"
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129" release-images'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129"'

      - label: "Build release image - aarch64 - CUDA 12.9"
        depends_on: ~
@@ -281,7 +238,7 @@ steps:
              --progress plain \
              -f docker/Dockerfile .
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129" release-images'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129"'

      - label: "Build release image - x86_64 - CUDA 13.0 - Ubuntu 24.04"
        depends_on: ~
@@ -310,7 +267,7 @@ steps:
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"
          - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404" release-images'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"'

      - label: "Build release image - aarch64 - CUDA 13.0 - Ubuntu 24.04"
        depends_on: ~
@@ -337,7 +294,7 @@ steps:
              --progress plain \
              -f docker/Dockerfile .
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404" release-images'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"'

      - label: "Build release image - x86_64 - CUDA 12.9 - Ubuntu 24.04"
        depends_on: ~
@@ -365,7 +322,7 @@ steps:
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"
          - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404" release-images'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"'

      - label: "Build release image - aarch64 - CUDA 12.9 - Ubuntu 24.04"
        depends_on: ~
@@ -391,26 +348,7 @@ steps:
              --progress plain \
              -f docker/Dockerfile .
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404" release-images'
-
-      - label: ":docker: Build release image - x86_64 - XPU"
-        depends_on: ~
-        id: build-xpu-release-image
-        agents:
-          queue: cpu_queue_release
-        commands:
-          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh xpu) \
-              --build-arg GIT_REPO_CHECK=1 \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile.xpu .
-          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-xpu"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-xpu" release-images'
-        env:
-          DOCKER_BUILDKIT: "1"
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"'

      - block: "Build release image for x86_64 CPU"
        key: block-cpu-release-image-build
@@ -428,7 +366,7 @@ steps:
          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg GIT_REPO_CHECK=1 --build-arg VLLM_CPU_X86=true --tag public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:$(buildkite-agent meta-data get release-version) --tag public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:latest --progress plain --target vllm-openai -f docker/Dockerfile.cpu ."
          - "docker push public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:latest"
          - "docker push public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:$(buildkite-agent meta-data get release-version)"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:$(buildkite-agent meta-data get release-version)" release-images'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:$(buildkite-agent meta-data get release-version)"'
        env:
          DOCKER_BUILDKIT: "1"

@@ -448,7 +386,7 @@ steps:
          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg GIT_REPO_CHECK=1 --tag public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:$(buildkite-agent meta-data get release-version) --tag public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:latest --progress plain --target vllm-openai -f docker/Dockerfile.cpu ."
          - "docker push public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:latest"
          - "docker push public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:$(buildkite-agent meta-data get release-version)"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:$(buildkite-agent meta-data get release-version)" release-images'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:$(buildkite-agent meta-data get release-version)"'
        env:
          DOCKER_BUILDKIT: "1"

@@ -466,7 +404,7 @@ steps:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
          - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64 --amend"
          - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 13.0" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT" release-manifests'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 13.0" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"'

      - label: "Create multi-arch manifest - CUDA 12.9"
        depends_on:
@@ -479,7 +417,7 @@ steps:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
          - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-cu129 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-cu129 --amend"
          - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 12.9" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129" release-manifests'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 12.9" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129"'

      - label: "Create multi-arch manifest - CUDA 13.0 - Ubuntu 24.04"
        depends_on:
@@ -492,7 +430,7 @@ steps:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
          - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-ubuntu2404 --amend"
          - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 13.0 Ubuntu 24.04" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404" release-manifests'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 13.0 Ubuntu 24.04" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"'

      - label: "Create multi-arch manifest - CUDA 12.9 - Ubuntu 24.04"
        depends_on:
@@ -505,17 +443,7 @@ steps:
          - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
          - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-cu129-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-cu129-ubuntu2404 --amend"
          - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 12.9 Ubuntu 24.04" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404" release-manifests'
-
-      - label: "Create manifest - XPU"
-        depends_on:
-          - build-xpu-release-image
-        id: create-manifest-xpu
-        agents:
-          queue: small_cpu_queue_release
-        commands:
-          - "bash .buildkite/scripts/xpu/create-xpu-ecr-manifest.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: XPU" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-xpu" release-manifests'
+          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 12.9 Ubuntu 24.04" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"'

      - label: "Publish nightly multi-arch image to DockerHub"
        depends_on:
@@ -891,23 +819,6 @@ steps:
      DOCKER_BUILDKIT: "1"
      S3_BUCKET: "vllm-wheels"

-  - label: "Publish nightly XPU image to DockerHub"
-    depends_on:
-      - create-manifest-xpu
-    if: build.env("NIGHTLY") == "1"
-    agents:
-      queue: small_cpu_queue_release
-    commands:
-      - "bash .buildkite/scripts/xpu/push-nightly-builds-xpu.sh"
-      - "bash .buildkite/scripts/cleanup-nightly-builds.sh nightly- vllm/vllm-openai-xpu"
-    plugins:
-      - docker-login#v3.0.0:
-          username: vllmbot
-          password-env: DOCKERHUB_TOKEN
-    env:
-      DOCKER_BUILDKIT: "1"
-      DOCKERHUB_USERNAME: "vllmbot"
-
  - label: "Publish nightly ROCm image to DockerHub"
    depends_on:
      - build-rocm-release-image
@@ -938,7 +849,6 @@ steps:
      - create-multi-arch-manifest-cuda-12-9
      - create-multi-arch-manifest-ubuntu2404
      - create-multi-arch-manifest-cuda-12-9-ubuntu2404
-      - create-manifest-xpu
      - build-rocm-release-image
      - input-release-version
      # Wait for CPU builds if their block steps were unblocked, so publish
@@ -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' \
-    --context "${3:?context is required}"
+  buildkite-agent annotate --append --style 'info' --context 'release-artifacts'
@@ -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
@@ -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
@@ -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/nightly/cu130"
-    else
-        TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
-    fi
+    TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
 fi
 echo ">>> Using PyTorch index: ${TORCH_INDEX_URL:-PyPI default}"

@@ -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_rixl 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[@]}" \
-                || return 1
+            hash_dockerfile_arg_values "${dockerfile}" "${content_args[@]}"
            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
-                echo "Failed to resolve digest for BASE_IMAGE=${arg_value}" >&2
-                return 1
-            fi
-            printf 'arg:%s.digest=%s\n' "${arg_name}" "${digest}"
+            printf 'arg:%s.digest=%s\n' "${arg_name}" "${digest:-unknown}"
        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

@@ -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 rust_content_cache_from=()
-    local -a combined_content_cache_from=()
+    local -a 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}")
-        if should_export_content_cache_ref "${ROCM_CSRC_CONTENT_CACHE_REF}" "ROCm csrc"; then
-            csrc_cache_to+=(
-                "type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF},mode=${csrc_cache_to_mode},ignore-error=true"
-            )
-        else
-            export_csrc_cache=0
-        fi
+        content_cache_from+=("type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF}")
+        csrc_cache_to+=(
+            "type=registry,ref=${ROCM_CSRC_CONTENT_CACHE_REF},mode=${csrc_cache_to_mode},ignore-error=true"
+        )
    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+=(
-                "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
+        csrc_cache_to+=(
+            "type=registry,ref=${cache_repo}:csrc-rocm-${BUILDKITE_COMMIT},mode=${csrc_cache_to_mode},ignore-error=true"
+        )
        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+=(
-                "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
+        csrc_cache_to+=(
+            "type=registry,ref=${cache_repo}:csrc-rocm-branch-${ROCM_CACHE_BRANCH_TAG},mode=${csrc_cache_to_mode},ignore-error=true"
+        )
        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
@@ -1950,8 +1720,8 @@ confirm_remote_image_push() {
        fi

        if [[ -z "${remote_revision}" \
-              && ${IMAGE_EXISTED_BEFORE_BUILD} -eq 0 ]] \
-            && image_tag_is_commit_scoped; then
+              && ${IMAGE_EXISTED_BEFORE_BUILD} -eq 0 \
+              && 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 \
-            "${BAKE_FILES[@]}" \
-            --progress "${BUILDKIT_PROGRESS:-plain}" \
-            "${target}"
+        docker buildx bake "${BAKE_FILES[@]}" --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 \
-        "${BAKE_FILES[@]}" \
-        --progress "${BUILDKIT_PROGRESS:-plain}" \
-        "${BAKE_TARGETS[@]}" || build_rc=$?
+    docker buildx bake "${BAKE_FILES[@]}" --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
-        mapfile -t wheels < <(find "${wheel_dir}" -maxdepth 1 -type f -name '*.whl' -print)
-    fi
-    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
+    if [[ ! -d "${wheel_dir}" ]] || ! ls "${wheel_dir}"/*.whl >/dev/null 2>&1; then
+        echo "No ROCm wheel artifacts found in ${wheel_dir}"
+        return 0
    fi

    echo "--- :package: Uploading ROCm vLLM install artifact"
-    rm -rf "${artifact_dir}" "${metadata_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"
+    mkdir -p "${artifact_dir}"

    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}/"
-    cp "${whl}" "${artifact_dir}/${whl_name}"
-    echo "Copied ${whl_name}: $(du -sh "${artifact_dir}/${whl_name}" | cut -f1)"
+    printf '%s\n' "${CI_BASE_IMAGE:-}" > "${artifact_dir}/ci-base-image.txt"
+    printf '%s\n' "${IMAGE_TAG:-}" > "${artifact_dir}/fallback-image.txt"
+
+    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
@@ -1,7 +1,7 @@
 #!/bin/bash

-# This script runs ROCm tests either directly in a native CI pod or inside the
-# corresponding Docker container. Multi-node tests continue to use Docker.
+# This script runs tests inside the corresponding ROCm docker container.
+# 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:=2100}"
-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 15 minutes, then stop an individual test after 35 minutes.
-if [[ " ${PYTEST_ADDOPTS:-} " != *" faulthandler_timeout="* ]]; then
-  PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }-o faulthandler_timeout=900"
-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}" \
-    --exclude='*.whl' \
-    --exclude='.vllm-ci-artifact' \
-    --exclude='./.vllm-ci-artifact' \
-    -cf - . \
+  tar -C "${wheel_dir}" --exclude='*.whl' -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,257 +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
-
-  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
@@ -644,58 +338,7 @@ re_quote_pytest_markers() {
 # Main
 ###############################################################################

-if is_native_runtime; then
-  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 ---
+# --- GPU initialization ---
 echo "--- ROCm info"
 rocminfo

@@ -797,27 +440,25 @@ fi

 echo "Final commands: $commands"

-standalone_merge_base_env=()
-if [[ "$commands" == *python_only_compile.sh* ]]; then
-  # The ROCm test image often ships /vllm-workspace without .git. Resolve the
-  # wheels.vllm.ai commit from the agent checkout for this test only.
-  vllm_standalone_merge_base=""
-  checkout="${BUILDKITE_BUILD_CHECKOUT_PATH:-}"
-  if [[ -z "${checkout}" || ! -d "${checkout}" ]]; then
-    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}")
+# The ROCm test image often ships /vllm-workspace without .git (artifact tarball unpack).
+# tests/standalone_tests/python_only_compile.sh uses merge-base(HEAD, origin/main) for
+# wheels.vllm.ai; compute on the agent (full git checkout) and pass into the container.
+vllm_standalone_merge_base=""
+checkout="${BUILDKITE_BUILD_CHECKOUT_PATH:-}"
+if [[ -z "${checkout}" || ! -d "${checkout}" ]]; then
+  checkout="."
 fi
+# Pass safe.directory per-command (-c) because buildkite runs will always fail
+# the next check on git 2.35.2+ due to mixed uses of root and buildkite-agent/uids.
+if git -c "safe.directory=${checkout}" -C "${checkout}" rev-parse --is-inside-work-tree >/dev/null 2>&1; then
+  vllm_standalone_merge_base="$(
+    git -c "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 VLLM_STANDALONE_MERGE_BASE into container: ${vllm_standalone_merge_base}"

 MYPYTHONPATH="/vllm-workspace"

@@ -848,7 +489,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/')
@@ -895,13 +535,6 @@ 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
@@ -918,7 +551,7 @@ else
  fi

  docker run \
-    "${docker_run_terminal_args[@]}" \
+    -t -i \
    --device /dev/kfd $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES \
    $RDMA_FLAGS \
    --network=host \
@@ -933,13 +566,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}" \
@@ -949,7 +575,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}"
@@ -35,7 +35,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
    ;;
@@ -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
@@ -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 "$@"
@@ -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 "$@"
@@ -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 "$@"
@@ -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
-'
@@ -21,20 +21,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 +45,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 +75,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" \
-    | env BINSTALL_VERSION="$CARGO_BINSTALL_VERSION" bash
-  cargo-binstall -V
+    https://raw.githubusercontent.com/cargo-bins/cargo-binstall/main/install-from-binstall-release.sh \
+    | bash
 }

 install_cargo_sort() {
-  log_section "Installing cargo-sort ${CARGO_SORT_VERSION_REQ}"
-  cargo binstall --no-confirm --force "cargo-sort@${CARGO_SORT_VERSION_REQ}"
+  if command -v cargo-sort >/dev/null 2>&1; then
+    return
+  fi
+
+  log_section "Installing cargo-sort"
+  install_cargo_binstall
+  cargo binstall --no-confirm cargo-sort
 }

 install_cargo_deny() {
-  log_section "Installing cargo-deny ${CARGO_DENY_VERSION_REQ}"
-  cargo binstall --no-confirm --force "cargo-deny@${CARGO_DENY_VERSION_REQ}"
+  if command -v cargo-deny >/dev/null 2>&1; then
+    return
+  fi
+
+  log_section "Installing cargo-deny"
+  install_cargo_binstall
+  cargo binstall --no-confirm cargo-deny
 }

 install_cargo_nextest() {
-  log_section "Installing cargo-nextest ${CARGO_NEXTEST_VERSION_REQ}"
-  cargo binstall \
-    --no-confirm \
-    --force \
-    --secure \
-    "cargo-nextest@${CARGO_NEXTEST_VERSION_REQ}"
+  if command -v cargo-nextest >/dev/null 2>&1; then
+    return
+  fi
+
+  log_section "Installing cargo-nextest"
+  install_cargo_binstall
+  cargo binstall --no-confirm --secure cargo-nextest
 }

 install_uv() {
-  log_section "Installing uv ${UV_VERSION}"
-  curl -L --proto '=https' --tlsv1.2 -sSf \
-    "https://github.com/astral-sh/uv/releases/download/${UV_VERSION}/uv-installer.sh" \
+  if command -v uv >/dev/null 2>&1; then
+    return
+  fi
+
+  log_section "Installing uv"
+  curl -LsSf --proto '=https' --tlsv1.2 https://astral.sh/uv/install.sh \
    | env UV_INSTALL_DIR="$CARGO_HOME/bin" sh
-  uv --version
 }

 setup_pyo3_python() {
-  log_section "Installing Python ${PYO3_PYTHON_VERSION} for PyO3 tests"
-  uv python install "$PYO3_PYTHON_VERSION"
+  local python_version="${PYO3_PYTHON_VERSION:-3.12}"
+
+  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 +156,6 @@ PY
 }

 run_style_clippy() {
-  install_cargo_binstall
  install_cargo_sort
  install_cargo_deny

@@ -156,8 +168,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 +186,6 @@ run_style_clippy() {
 run_tests() {
  install_uv
  setup_pyo3_python
-  install_cargo_binstall
  install_cargo_nextest

  log_section "Running cargo nextest"
@@ -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)
-    --pipeline       Buildkite pipeline slug (default: ${PIPELINE})
-    --commit         Override commit SHA (default: current HEAD)
-    --branch         Override branch name (default: current branch)
-    --message        Custom build message (default: auto-generated)
-    --torch-nightly  Also build and run the full suite against torch nightly
-                     (sets TORCH_NIGHTLY=1)
-    --help           Show this help message
+    --execute       Actually trigger the build (default: dry-run)
+    --pipeline      Buildkite pipeline slug (default: ${PIPELINE})
+    --commit        Override commit SHA (default: current HEAD)
+    --branch        Override branch name (default: current branch)
+    --message       Custom build message (default: auto-generated)
+    --help          Show this help message

 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 ))
-    for i in "${!BUILD_ENV[@]}"; do
-        if [[ $i -eq $last_idx ]]; then
-            echo "    --env '$(escape_for_shell "${BUILD_ENV[$i]}")'"
-        else
-            echo "    --env '$(escape_for_shell "${BUILD_ENV[$i]}")' \\"
-        fi
-    done
+    echo "    --env 'RUN_ALL=1' \\"
+    echo "    --env 'NIGHTLY=1'"
    echo ""
    echo "=========================================="
    echo -e "${YELLOW}To actually trigger this build, run:${NC}"
@@ -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
-# manylinux retag for them.
-WHEEL_PLATFORM="${VLLM_WHEEL_PLATFORM:-linux}"
-
-if [[ "$WHEEL_PLATFORM" == "linux" ]]; then
-  # shellcheck source=lib/manylinux.sh
-  source .buildkite/scripts/lib/manylinux.sh
-fi
+# shellcheck source=lib/manylinux.sh
+source .buildkite/scripts/lib/manylinux.sh

 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")"
-  echo "Renamed wheel to: $wheel"
-fi
+wheel="$(apply_manylinux_tag "$wheel")"
+echo "Renamed wheel to: $wheel"

 # Extract the version from the wheel
 version=$(unzip -p "$wheel" '**/METADATA' | grep '^Version: ' | cut -d' ' -f2)
@@ -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}
@@ -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"
@@ -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,12 +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
-  parallelism: 2
+    - pytest -v -s v1/attention
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 95
+      timeout_in_minutes: 70
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -31,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
@@ -39,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
-  parallelism: 2
+    - pytest -v -s v1/attention
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Basic Correctness
  key: basic-correctness
-  timeout_in_minutes: 45
+  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -19,6 +19,6 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 70
+      timeout_in_minutes: 50
      depends_on:
      - image-build-amd
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Benchmarks CLI Test
  key: benchmarks-cli-test
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -13,7 +13,6 @@ steps:
  - pytest -v -s benchmarks/
  mirror:
    amd:
-      dind: false
      device: mi300_1
      depends_on:
      - image-build-amd
@@ -24,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/
@@ -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
@@ -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
@@ -19,7 +19,7 @@ steps:

 - label: Cudagraph
  key: cudagraph
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  source_file_dependencies:
  - tests/v1/cudagraph
  - vllm/v1/cudagraph_dispatcher.py
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed NixlConnector PD accuracy (4 GPUs)
  key: distributed-nixlconnector-pd-accuracy-4-gpus
-  timeout_in_minutes: 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: 85
+      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: 60
+      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: 85
+      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: 80
+      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: 70
+      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
@@ -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,7 +39,6 @@ steps:
  - DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py
  mirror:
    amd:
-      dind: false
      device: mi300_2
      depends_on:
      - image-build-amd
@@ -56,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:
@@ -79,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:
@@ -134,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:
@@ -155,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:
@@ -177,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"
@@ -213,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/"
@@ -260,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:
@@ -275,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:
@@ -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
@@ -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
@@ -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/
@@ -29,13 +29,13 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 50
+      timeout_in_minutes: 60
      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,13 +45,13 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 55
+      timeout_in_minutes: 40
      depends_on:
      - image-build-amd

 - label: e2e Scheduling (1 GPU)
  key: e2e-scheduling-1-gpu
-  timeout_in_minutes: 35
+  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/
@@ -60,15 +60,15 @@ steps:
    - pytest -v -s v1/e2e/general/test_async_scheduling.py
  mirror:
    amd:
-      device: mi325_1
-      timeout_in_minutes: 70
+      device: mi250_1
+      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/
@@ -76,8 +76,8 @@ steps:
    - pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py
  mirror:
    amd:
-      device: mi325_1
-      timeout_in_minutes: 60
+      device: mi250_1
+      timeout_in_minutes: 35
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -87,7 +87,7 @@ steps:

 - label: V1 e2e (2 GPUs)
  key: v1-e2e-2-gpus
-  timeout_in_minutes: 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:
@@ -114,14 +114,13 @@ steps:
    - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "tensor_parallelism"
  mirror:
    amd:
-      dind: false
      device: mi300_2
      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:
@@ -149,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
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Entrypoints Unit Tests
  key: entrypoints-unit-tests
-  timeout_in_minutes: 25
+  timeout_in_minutes: 10
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/entrypoints
@@ -16,7 +16,7 @@ steps:

 - label: Entrypoints Integration (LLM)
  key: entrypoints-integration-llm
-  timeout_in_minutes: 60
+  timeout_in_minutes: 40
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -37,7 +37,7 @@ steps:
 - label: Entrypoints Integration (API Server)
  key: entrypoints-integration-api-server
  device: h200_35gb
-  timeout_in_minutes: 50
+  timeout_in_minutes: 130
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -56,7 +56,7 @@ steps:

 - label: Entrypoints Integration (API Server OpenAI - Part 1)
  key: entrypoints-integration-api-server-openai-part-1
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -68,13 +68,13 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 65
+      timeout_in_minutes: 80
      depends_on:
      - image-build-amd

 - label: Entrypoints Integration (API Server OpenAI - Part 2)
  key: entrypoints-integration-api-server-openai-part-2
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -109,7 +109,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 65
+      timeout_in_minutes: 60
      depends_on:
      - image-build-amd

@@ -126,7 +126,7 @@ steps:
 - label: Entrypoints Integration (Speech to Text)
  device: h200_35gb
  key: entrypoints-integration-speech_to_text
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -138,7 +138,7 @@ steps:
 - label: Entrypoints Integration (Multimodal)
  device: h200_35gb
  key: entrypoints-integration-multimodal
-  timeout_in_minutes: 45
+  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - vllm/
@@ -160,7 +160,7 @@ steps:

 - label: OpenAI API Correctness
  key: openai-api-correctness
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
  - csrc/
@@ -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,7 +16,6 @@ steps:
  - pytest -v -s distributed/test_eplb_utils.py
  mirror:
    amd:
-      dind: false
      device: mi300_1
      depends_on:
      - image-build-amd
@@ -28,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:
@@ -40,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
@@ -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:
@@ -16,19 +16,18 @@ steps:

 - label: Kernels Core Operation Test
  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
-  parallelism: 3
+    - 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

 - 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:
@@ -42,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
@@ -55,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
@@ -66,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,7 +79,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 90
+      timeout_in_minutes: 55
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -107,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
@@ -132,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/
@@ -148,7 +147,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 65
+      timeout_in_minutes: 50
      source_file_dependencies:
      - csrc/quantization/cutlass_w8a8/moe/
      - csrc/moe/
@@ -164,7 +163,7 @@ steps:

 - label: Kernels Mamba Test
  key: kernels-mamba-test
-  timeout_in_minutes: 40
+  timeout_in_minutes: 45
  source_file_dependencies:
  - csrc/mamba/
  - tests/kernels/mamba
@@ -173,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/third_party/flash_linear_attention/ops/chunk_delta_h.py
-  - vllm/third_party/flash_linear_attention/ops/l2norm.py
+  - vllm/model_executor/layers/fla/ops/kda.py
+  - vllm/model_executor/layers/fla/ops/chunk_delta_h.py
+  - vllm/model_executor/layers/fla/ops/l2norm.py
  - tests/kernels/test_kda.py
  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:
@@ -212,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
@@ -232,10 +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
-  - tests/kernels/test_ll_bf16_gemm.py
  - tests/kernels/test_top_k_per_row.py
  commands:
    - nvidia-smi
@@ -264,26 +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/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
-  parallelism: 2
+    - pytest -v -s kernels/helion/

 
 - 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
@@ -300,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
@@ -310,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
@@ -323,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:
@@ -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
@@ -56,7 +56,7 @@ steps:

 - label: LM Eval Small Models (1xB200)
  key: lm-eval-small-models-1xb200
-  timeout_in_minutes: 50
+  timeout_in_minutes: 120
  device: b200-k8s
  optional: true
  source_file_dependencies:
@@ -80,7 +80,7 @@ steps:

 - label: LM Eval Large Models EP (2xB200)
  key: lm-eval-large-models-ep-2xb200
-  timeout_in_minutes: 60
+  timeout_in_minutes: 120
  device: b200-k8s
  optional: true
  num_devices: 2
@@ -92,7 +92,7 @@ steps:

 - label: LM Eval Qwen3.5 Models (2xB200)
  key: lm-eval-qwen3-5-models-2xb200
-  timeout_in_minutes: 45
+  timeout_in_minutes: 120
  device: b200-k8s
  optional: true
  num_devices: 2
@@ -103,13 +103,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
@@ -117,9 +117,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: 60
+      timeout_in_minutes: 180
      depends_on:
      - image-build-amd
      commands:
@@ -153,7 +152,7 @@ steps:

 - label: LM Eval Humming f16 (A100 - TEMPORARY)
  key: lm-eval-humming-f16-a100
-  timeout_in_minutes: 75
+  timeout_in_minutes: 120
  device: a100
  optional: true
  num_devices: 1
@@ -168,7 +167,7 @@ steps:

 - label: LM Eval Humming Act int8 (A100 - TEMPORARY)
  key: lm-eval-humming-act-a100
-  timeout_in_minutes: 45
+  timeout_in_minutes: 120
  device: a100
  optional: true
  num_devices: 1
@@ -183,7 +182,7 @@ steps:

 - label: LM Eval Humming f16 (H100 - TEMPORARY)
  key: lm-eval-humming-f16-h100
-  timeout_in_minutes: 70
+  timeout_in_minutes: 120
  device: h100
  optional: true
  num_devices: 1
@@ -198,7 +197,7 @@ steps:

 - label: LM Eval Humming Act fp8/int8 (H100 - TEMPORARY)
  key: lm-eval-humming-act-h100
-  timeout_in_minutes: 70
+  timeout_in_minutes: 120
  device: h100
  optional: true
  num_devices: 1
@@ -214,7 +213,7 @@ steps:

 - label: LM Eval Humming f16 (B200 - TEMPORARY)
  key: lm-eval-humming-f16-b200
-  timeout_in_minutes: 50
+  timeout_in_minutes: 120
  device: b200-k8s
  optional: true
  num_devices: 1
@@ -229,7 +228,7 @@ steps:

 - label: LM Eval Humming Act fp8/int8 (B200 - TEMPORARY)
  key: lm-eval-humming-act-b200
-  timeout_in_minutes: 50
+  timeout_in_minutes: 120
  device: b200-k8s
  optional: true
  num_devices: 1
@@ -245,7 +244,7 @@ steps:

 - 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/
@@ -257,7 +256,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
@@ -271,7 +270,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
@@ -285,7 +284,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
@@ -299,50 +298,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:
@@ -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
@@ -16,7 +16,7 @@ steps:
    amd:
      device: mi325_1
      working_dir: "/vllm-workspace/tests"
-      timeout_in_minutes: 65
+      timeout_in_minutes: 60
      source_file_dependencies:
      - vllm/lora
      - tests/lora
@@ -27,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
@@ -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: 75
+      timeout_in_minutes: 65
      depends_on:
      - image-build-amd

 - label: V1 Sample + Logits
  key: v1-sample-logits
-  timeout_in_minutes: 45
+  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
    - vllm/config/
@@ -65,7 +64,7 @@ steps:

 - label: V1 Core + KV + Metrics
  key: v1-core-kv-metrics
-  timeout_in_minutes: 60
+  timeout_in_minutes: 30
  source_file_dependencies:
    - vllm/config/
    - vllm/distributed/
@@ -90,7 +89,6 @@ steps:
    - tests/v1/simple_kv_offload
    - tests/v1/worker
    - tests/v1/kv_connector/unit
-    - tests/v1/ec_connector/unit
    - tests/v1/metrics
    - tests/entrypoints/openai/correctness/test_lmeval.py
  commands:
@@ -103,7 +101,6 @@ steps:
    - pytest -v -s v1/simple_kv_offload
    - pytest -v -s v1/worker
    - pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
-    - pytest -v -s -m 'not cpu_test' v1/ec_connector/unit
    - pytest -v -s -m 'not cpu_test' v1/metrics
    # Integration test for streaming correctness (requires special branch).
    - pip install -U git+https://github.com/vllm-project/lm-evaluation-harness.git@streaming-api
@@ -111,7 +108,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 75
+      timeout_in_minutes: 60
      depends_on:
      - image-build-amd

@@ -175,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,7 +195,7 @@ steps:
 - 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
@@ -240,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/
@@ -284,7 +281,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 45
+      timeout_in_minutes: 20
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -295,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/
@@ -322,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/
@@ -354,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
@@ -372,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
@@ -384,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
@@ -398,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
@@ -414,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
@@ -433,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
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Model Executor
  key: model-executor
-  timeout_in_minutes: 45
+  timeout_in_minutes: 35
  source_file_dependencies:
  - vllm/engine/arg_utils.py
  - vllm/config/model.py
@@ -25,7 +25,6 @@ steps:
    - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py --timeout=900 --timeout-method=thread
  mirror:
    amd:
-      dind: false
      device: mi300_1
      depends_on:
      - image-build-amd
@@ -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
@@ -27,7 +27,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/
@@ -63,7 +63,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 +84,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:
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Basic Models Tests (Initialization)
  key: basic-models-tests-initialization
-  timeout_in_minutes: 25
+  timeout_in_minutes: 45
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -17,7 +17,7 @@ steps:
 - label: Basic Models Tests (Extra Initialization) %N
  device: h200_35gb
  key: basic-models-tests-extra-initialization
-  timeout_in_minutes: 100
+  timeout_in_minutes: 45
  source_file_dependencies:
  - vllm/model_executor/models/
  - tests/models/test_initialization.py
@@ -27,12 +27,12 @@ steps:
    # subset of supported models (the complement of the small subset in the above
    # test.) Also run if model initialization test file is modified
    - pytest -v -s models/test_initialization.py -k 'not test_can_initialize_small_subset' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
-  parallelism: 4
+  parallelism: 2

 - 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
@@ -50,7 +50,7 @@ steps:
  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
@@ -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:
@@ -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/
@@ -15,14 +15,13 @@ steps:
    - pytest -v -s models/language -m 'core_model and (not slow_test)'
  mirror:
    amd:
-      dind: false
      device: mi300_1
      depends_on:
      - image-build-amd

 - label: Language Models Tests (Extra Standard) %N
  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
@@ -36,7 +35,6 @@ steps:
  parallelism: 2
  mirror:
    amd:
-      dind: false
      device: mi300_1
      depends_on:
      - image-build-amd
@@ -54,7 +52,7 @@ steps:

 - label: Language Models Tests (Hybrid) %N
  key: language-models-tests-hybrid
-  timeout_in_minutes: 65
+  timeout_in_minutes: 75
  source_file_dependencies:
  - vllm/
  - tests/models/language/generation
@@ -69,7 +67,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 70
+      timeout_in_minutes: 90
      depends_on:
      - image-build-amd
      commands:
@@ -80,7 +78,7 @@ steps:
 - 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/
@@ -94,7 +92,7 @@ steps:

 - 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:
@@ -103,10 +101,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/
@@ -116,13 +114,13 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 120
+      timeout_in_minutes: 100
      depends_on:
      - image-build-amd

 - label: Language Models Test (MTEB)
  key: language-models-test-mteb
-  timeout_in_minutes: 45
+  timeout_in_minutes: 110
  device: h200_18gb
  optional: true
  source_file_dependencies:
@@ -20,7 +20,7 @@ steps:

 - label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
  key: multi-modal-models-standard-2-qwen3-gemma
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -38,7 +38,7 @@ steps:
 - label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
  device: h200_35gb
  key: multi-modal-models-standard-3-llava-qwen2-vl
-  timeout_in_minutes: 40
+  timeout_in_minutes: 45
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal
@@ -54,7 +54,7 @@ steps:
 - label: "Multi-Modal Models (Standard) 4: other + whisper"
  device: h200_35gb
  key: multi-modal-models-standard-4-other-whisper
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
  source_file_dependencies:
  - vllm/
  - tests/models/multimodal
@@ -69,23 +69,22 @@ steps:
      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
-  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: 65
+  timeout_in_minutes: 60
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -97,7 +96,7 @@ steps:
 - label: Multi-Modal Accuracy Eval (Small Models) # 50min
  device: h200_35gb
  key: multi-modal-accuracy-eval-small-models
-  timeout_in_minutes: 30
+  timeout_in_minutes: 70
  working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
  source_file_dependencies:
  - vllm/multimodal/
@@ -107,7 +106,6 @@ 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
      depends_on:
      - image-build-amd
@@ -166,7 +164,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 75
+      timeout_in_minutes: 60
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Plugin Tests (2 GPUs)
  key: plugin-tests-2-gpus
-  timeout_in_minutes: 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
@@ -5,7 +5,7 @@ steps:
 - label: PyTorch Compilation Unit Tests
  device: h200_35gb
  key: pytorch-compilation-unit-tests
-  timeout_in_minutes: 90
+  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
@@ -109,15 +109,14 @@ steps:
  - pytest -s -v compile/passes --ignore compile/passes/distributed
  mirror:
    amd:
-      dind: false
      device: mi300_1
-      timeout_in_minutes: 65
+      timeout_in_minutes: 180
      depends_on:
      - image-build-amd

 - label: PyTorch Fullgraph Smoke Test
  key: pytorch-fullgraph-smoke-test
-  timeout_in_minutes: 60
+  timeout_in_minutes: 35
  source_file_dependencies:
  - vllm/__init__.py
  - vllm/_aiter_ops.py
@@ -153,7 +152,7 @@ steps:

 - label: PyTorch Fullgraph
  key: pytorch-fullgraph
-  timeout_in_minutes: 40
+  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
  - vllm/__init__.py
@@ -198,7 +197,6 @@ steps:
  - bash standalone_tests/pytorch_nightly_dependency.sh
  mirror:
    amd:
-      dind: false
      device: mi300_1
      depends_on:
      - image-build-amd
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Quantization
  key: quantization
-  timeout_in_minutes: 60
+  timeout_in_minutes: 90
  source_file_dependencies:
  - csrc/
  - vllm/model_executor/layers/quantization
@@ -23,7 +23,7 @@ steps:

 - label: Quantized Fusions
  key: quantized-fusions
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  source_file_dependencies:
  - tests/fusion
  - vllm/model_executor/layers/fusion
@@ -35,7 +35,7 @@ steps:

 - label: Quantized MoE Test (B200)
  key: quantized-moe-test-b200
-  timeout_in_minutes: 120
+  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/"
  device: b200-k8s
  source_file_dependencies:
@@ -53,7 +53,7 @@ steps:

 - label: Quantized Models Test
  key: quantized-models-test
-  timeout_in_minutes: 50
+  timeout_in_minutes: 60
  source_file_dependencies:
  - vllm/model_executor/layers/quantization
  - tests/models/quantization
@@ -3,7 +3,7 @@ depends_on:
  - image-build
 steps:
 - label: Rust Frontend OpenAI Coverage
-  timeout_in_minutes: 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_uds.py
+  # - tests/entrypoints/openai/test_return_token_ids.py
+  # - tests/entrypoints/openai/test_uds.py
  - tests/v1/sample/test_logprobs_e2e.py
  commands:
  - export VLLM_USE_RUST_FRONTEND=1
  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
  - pytest -v -s benchmarks/test_serve_cli.py -k "not insecure and not (test_bench_serve and not test_bench_serve_chat)"
  - pytest -v -s entrypoints/openai/chat_completion/test_chat_completion.py -k "not test_invalid_json_schema and not test_invalid_regex"
-  - pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not 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
-  # finish chunk without logprobs, while the test reads `logprobs.tokens` on every chunk.
-  - pytest -v -s entrypoints/openai/test_return_token_ids.py -k "not test_comparison"
-  - pytest -v -s entrypoints/openai/test_uds.py
+  # - pytest -v -s entrypoints/openai/test_return_token_ids.py
+  # - 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
-  # server_load can be flaky under the Rust frontend; keep it excluded for now.
+  # - pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
  - pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not show_version and not server_load"
-  # 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,7 +74,7 @@ steps:
  - pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine

 - label: Rust Frontend Tool Use
-  timeout_in_minutes: 25
+  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/tests"
  source_file_dependencies:
  - rust/
@@ -95,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:
@@ -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:
@@ -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
@@ -19,7 +19,7 @@ steps:
    - VLLM_USE_FLASHINFER_SAMPLER=1 pytest -v -s samplers
  mirror:
    amd:
-      device: mi325_1
+      device: mi250_1
      depends_on:
      - image-build-amd
      commands:
@@ -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/
@@ -15,7 +15,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 60
+      timeout_in_minutes: 45
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -29,7 +29,7 @@ steps:

 - label: Spec Decode Eagle Nightly B200
  key: spec-decode-eagle-nightly-b200
-  timeout_in_minutes: 25
+  timeout_in_minutes: 30
  device: b200-k8s
  optional: true
  source_file_dependencies:
@@ -41,7 +41,7 @@ steps:

 - label: Spec Decode Speculators + MTP
  key: spec-decode-speculators-mtp
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/spec_decode/
@@ -82,7 +82,7 @@ steps:
  
 - label: Spec Decode Ngram + Suffix
  key: spec-decode-ngram-suffix
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/spec_decode/
@@ -93,7 +93,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 55
+      timeout_in_minutes: 65
      # TODO(akaratza): Test after Torch >= 2.12 bump
      soft_fail: true
      depends_on:
@@ -109,7 +109,7 @@ steps:

 - label: Spec Decode Draft Model
  key: spec-decode-draft-model
-  timeout_in_minutes: 45
+  timeout_in_minutes: 30
  device: h200_18gb
  source_file_dependencies:
    - vllm/v1/spec_decode/
@@ -120,7 +120,7 @@ steps:
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 55
+      timeout_in_minutes: 50
      depends_on:
      - image-build-amd
      source_file_dependencies:
@@ -134,7 +134,7 @@ steps:

 - label: Spec Decode Draft Model Nightly B200
  key: spec-decode-draft-model-nightly-b200
-  timeout_in_minutes: 40
+  timeout_in_minutes: 30
  device: b200-k8s
  optional: true
  source_file_dependencies:
@@ -146,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
@@ -159,7 +159,7 @@ steps:
    - pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test

 - label: Spec Decode MTP hybrid (B200)
-  timeout_in_minutes: 20
+  timeout_in_minutes: 30
  device: b200-k8s
  optional: true
  source_file_dependencies:
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Weight Loading Multiple GPU  # 33min
  key: weight-loading-multiple-gpu
-  timeout_in_minutes: 50
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  optional: true
@@ -15,7 +15,6 @@ steps:
    - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models.txt
  mirror:
    amd:
-      dind: false
      device: mi300_2
      depends_on:
      - image-build-amd
@@ -172,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
@@ -323,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: {
@@ -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 => readyLabels.includes(l.name));
+          const hasReadyLabel = pr.labels.some(l => l.name === 'ready');
          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:
@@ -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

@@ -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
@@ -0,0 +1,2 @@
+collect_env.py
+vllm/model_executor/layers/fla/ops/*.py
@@ -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):
-uv pip install -r requirements/test/cuda.in
+# Install test dependencies.
+# requirements/test/cuda.txt is pinned to x86_64; on other platforms, use the
+# unpinned source file instead:
+uv pip install -r requirements/test/cuda.in    # resolves for current platform
+# Or on x86_64:
+uv pip install -r requirements/test/cuda.txt

-# 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
- Minimize use of comments. Eliminate comments which are redundant, preferring legible and self-documenting code. When used, keep docstrings and comments brief and direct.
+Follow these rules for all code changes in this repository:
+
+- Try to match existing code style.
+- Code should be self-documenting and self-explanatory.
+- Keep comments and docstrings minimal and concise.
 - 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>
 ```

@@ -70,15 +70,6 @@ endif()
 #
 set(TORCH_SUPPORTED_VERSION_CUDA "2.11.0")
 set(TORCH_SUPPORTED_VERSION_ROCM "2.11.0")
-# TORCH_NIGHTLY=1 builds run against unpinned nightly wheels, so the supported-
-# version check would always warn. Only treat it as a nightly build when the
-# value is exactly "1" (the bootstrap exports TORCH_NIGHTLY=0 by default, which
-# must NOT suppress the warning for normal builds).
-if (DEFINED ENV{TORCH_NIGHTLY} AND "$ENV{TORCH_NIGHTLY}" STREQUAL "1")
-  set(TORCH_NIGHTLY_BUILD TRUE)
-else()
-  set(TORCH_NIGHTLY_BUILD FALSE)
-endif()

 #
 # Try to find python package with an executable that exactly matches
@@ -184,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()
@@ -197,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.")
@@ -390,7 +381,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
    "csrc/libtorch_stable/cuda_view.cu"
    "csrc/libtorch_stable/cuda_utils_kernels.cu"
    "csrc/libtorch_stable/activation_kernels.cu"
-    "csrc/libtorch_stable/ngram_embedding_kernels.cu"
    "csrc/libtorch_stable/quantization/activation_kernels.cu"
    "csrc/libtorch_stable/quantization/w8a8/int8/scaled_quant.cu"
    "csrc/libtorch_stable/quantization/w8a8/fp8/common.cu"
@@ -1364,7 +1354,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)
@@ -1408,7 +1397,6 @@ if (VLLM_GPU_LANG STREQUAL "CUDA")
    include(cmake/external_projects/fmha_sm100.cmake)
    include(cmake/external_projects/flashmla.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)
@@ -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,
-        prefill_backend=config.prefill_backend,
-        sparse_mla_force_mqa=config.sparse_mla_force_mqa,
-        **kwargs,
+        config.backend, config, prefill_backend=config.prefill_backend, **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
@@ -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)


@@ -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"
@@ -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(
-            total_q,
-            max_kv_len,
-            getattr(config, "sparse_mla_topk_pattern", "random"),
-        )
+        indexer.fill_random_indices(total_q, max_kv_len)

    # Determine which forward methods to use based on metadata.
-    # Non-sparse backends use .decode/.prefill sub-objects.
-    # Sparse backends use num_decode_tokens/num_prefills directly.
-    #
-    # 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
+    # Sparse MLA backends always use forward_mqa
+    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
    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
-            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:
+            if 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,
    )
@@ -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
    ):
@@ -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"
@@ -1,201 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""
-Benchmark the RDNAHybridW4A16LinearKernel across decode and prefill shapes.
-
-Usage:
-    python benchmark_int4_gemm.py
-    python benchmark_int4_gemm.py --models Qwen/Qwen3-4B
-    python benchmark_int4_gemm.py --group-size 128
-"""
-
-import argparse
-import copy
-import itertools
-import os
-
-import torch
-
-from vllm.triton_utils import triton
-
-# ---------------------------------------------------------------------------
-# Weight shapes: [K, N], TP_SPLIT_DIM
-# ---------------------------------------------------------------------------
-WEIGHT_SHAPES = {
-    "Qwen/Qwen3-4B": [
-        ([2560, 3840], 1),  # qkv_proj
-        ([2560, 2560], 0),  # o_proj
-        ([2560, 19456], 1),  # gate_up_proj
-        ([9728, 2560], 0),  # down_proj
-    ],
-    "Qwen/Qwen2.5-7B-Instruct": [
-        ([3584, 4608], 1),
-        ([3584, 3584], 0),
-        ([3584, 37888], 1),
-        ([18944, 3584], 0),
-    ],
-    "trymirai/SmolLM2-1.7B-Instruct-AWQ": [
-        ([2048, 6144], 1),  # qkv_proj
-        ([2048, 2048], 0),  # o_proj
-        ([2048, 16384], 1),  # gate_up_proj
-        ([8192, 2048], 0),  # down_proj
-    ],
-    "RedHatAI/Qwen3-8B-quantized.w4a16": [
-        ([4096, 6144], 1),  # qkv_proj
-        ([4096, 4096], 0),  # o_proj
-        ([4096, 24576], 1),  # gate_up_proj
-        ([12288, 4096], 0),  # down_proj
-    ],
-}
-
-
-# ---------------------------------------------------------------------------
-# Weight packing
-# ---------------------------------------------------------------------------
-def prepare_hybrid_weights(K, N, group_size, device="cuda"):
-    """Create random weights for benchmarking.
-
-    Returns (w_q_skinny, w_s_skinny, w_fp16, w_zp). The triton path derives
-    its int32 view from w_q_skinny, so no separate int32 buffer is returned.
-    """
-    num_groups = K // group_size
-
-    # Random packed weights — actual values don't matter for throughput
-    w_q_skinny_i32 = torch.randint(
-        0, 2**31, (N, K // 8), dtype=torch.int32, device=device
-    )
-    w_q_skinny = w_q_skinny_i32.view(torch.int8).contiguous()
-    w_s_skinny = torch.randn(N, num_groups, dtype=torch.float16, device=device) * 0.01
-
-    # Raw per-group zero-points for asymmetric benchmarks
-    w_zp = torch.randint(0, 16, (N, num_groups), dtype=torch.int32, device=device).to(
-        torch.float16
-    )
-
-    # FP16 baseline for F.linear
-    w_fp16 = torch.randn(N, K, dtype=torch.float16, device=device) * 0.01
-
-    return w_q_skinny, w_s_skinny, w_fp16, w_zp
-
-
-# ---------------------------------------------------------------------------
-# Benchmark
-# ---------------------------------------------------------------------------
-PROVIDERS = ["torch-fp16", "hybrid-w4a16", "hybrid-w4a16-zp"]
-
-
-@triton.testing.perf_report(
-    triton.testing.Benchmark(
-        x_names=["batch_size"],
-        x_vals=[1, 2, 4, 8, 16, 32, 64, 128, 256, 512, 1024, 2048, 4096],
-        x_log=False,
-        line_arg="provider",
-        line_vals=PROVIDERS,
-        line_names=PROVIDERS,
-        ylabel="TFLOP/s (larger is better)",
-        plot_name="FP16 vs Hybrid W4A16",
-        args={},
-    )
-)
-def benchmark(batch_size, provider, N, K, group_size, weights):
-    M = batch_size
-    device = "cuda"
-    dtype = torch.float16
-    a = torch.randn((M, K), device=device, dtype=dtype)
-
-    quantiles = [0.5, 0.2, 0.8]
-
-    if provider == "torch-fp16":
-        w_fp16 = weights["w_fp16"]
-        ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
-            lambda: torch.nn.functional.linear(a, w_fp16),
-            quantiles=quantiles,
-        )
-    elif provider in ("hybrid-w4a16", "hybrid-w4a16-zp"):
-        from vllm.model_executor.kernels.linear.mixed_precision import (
-            rdna_hybrid_w4a16 as _k,
-        )
-
-        _rdna_hybrid_w4a16_apply_impl = _k._rdna_hybrid_w4a16_apply_impl
-        from vllm.utils.platform_utils import num_compute_units
-
-        w = weights
-        cu_count = num_compute_units()
-        use_zp = provider == "hybrid-w4a16-zp"
-
-        def run():
-            return _rdna_hybrid_w4a16_apply_impl(
-                a,
-                w["w_q_skinny"],
-                w["w_s_skinny"],
-                w["w_zp"] if use_zp else None,
-                None,  # bias
-                cu_count,
-                group_size,
-            )
-
-        ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
-            run,
-            quantiles=quantiles,
-        )
-    else:
-        return 0.0, 0.0, 0.0
-
-    to_tflops = lambda t_ms: (2 * M * N * K) * 1e-12 / (t_ms * 1e-3)
-    return to_tflops(ms), to_tflops(max_ms), to_tflops(min_ms)
-
-
-def prepare_shapes(args):
-    KN_model_names = []
-    for model, tp_size in itertools.product(args.models, args.tp_sizes):
-        for KN, tp_dim in copy.deepcopy(WEIGHT_SHAPES[model]):
-            KN[tp_dim] //= tp_size
-            KN.append(model)
-            KN_model_names.append(KN)
-    return KN_model_names
-
-
-if __name__ == "__main__":
-    parser = argparse.ArgumentParser(
-        description="Benchmark RDNAHybridW4A16LinearKernel"
-    )
-    parser.add_argument(
-        "--models",
-        nargs="+",
-        type=str,
-        default=["Qwen/Qwen3-4B"],
-        choices=list(WEIGHT_SHAPES.keys()),
-    )
-    parser.add_argument("--tp-sizes", nargs="+", type=int, default=[1])
-    parser.add_argument("--group-size", type=int, default=128)
-    parser.add_argument("--save-path", type=str, default=None)
-    args = parser.parse_args()
-
-    for K, N, model in prepare_shapes(args):
-        group_size = args.group_size
-        print(f"\n{'=' * 70}")
-        print(f"{model}, N={N} K={K}, group_size={group_size}")
-        print(f"{'=' * 70}")
-
-        w_q_skinny, w_s_skinny, w_fp16, w_zp = prepare_hybrid_weights(K, N, group_size)
-
-        weights = {
-            "w_q_skinny": w_q_skinny,
-            "w_s_skinny": w_s_skinny,
-            "w_fp16": w_fp16,
-            "w_zp": w_zp,
-        }
-
-        save_path = args.save_path or f"bench_int4_res_n{N}_k{K}"
-        os.makedirs(save_path, exist_ok=True)
-        benchmark.run(
-            print_data=True,
-            show_plots=False,
-            save_path=save_path,
-            N=N,
-            K=K,
-            group_size=group_size,
-            weights=weights,
-        )
-
-    print("\nBenchmark finished!")
@@ -1,108 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-# Benchmark ReLUSquaredActivation: custom CUDA kernel vs forward_native, both
-# eager and under torch.compile (Inductor fuses relu+square into one kernel).
-
-import itertools
-
-import torch
-import torch.nn.functional as F
-
-import vllm.model_executor.layers.activation  # noqa: F401
-from vllm.benchmarks.lib.utils import default_vllm_config
-from vllm.triton_utils import triton
-from vllm.utils.argparse_utils import FlexibleArgumentParser
-from vllm.utils.torch_utils import STR_DTYPE_TO_TORCH_DTYPE, set_random_seed
-
-# Capped so the largest tensor stays under 2**31 elements: the shared activation
-# kernel computes the per-token pointer offset (blockIdx.x * d) in 32-bit, which
-# overflows for tensors with >2**32 elements. Realistic token counts are well
-# below this; the kernel-vs-native gap is already clear at these sizes.
-batch_size_range = [1, 16, 128]
-seq_len_range = [1, 16, 64, 1024]
-intermediate_size = [3072, 9728, 12288]
-configs = list(itertools.product(batch_size_range, seq_len_range, intermediate_size))
-
-
-@default_vllm_config()
-def benchmark_relu_squared(
-    batch_size: int,
-    seq_len: int,
-    intermediate_size: int,
-    provider: str,
-    dtype: torch.dtype,
-):
-    device = "cuda"
-    num_tokens = batch_size * seq_len
-    set_random_seed(42)
-    torch.set_default_device(device)
-
-    x = torch.randn(num_tokens, intermediate_size, dtype=dtype, device=device)
-    out = torch.empty_like(x)
-
-    def native(x: torch.Tensor) -> torch.Tensor:
-        return torch.square(F.relu(x))
-
-    # Verify the custom kernel matches the native implementation before timing.
-    ref = native(x)
-    torch.ops._C.relu_squared(out, x)
-    torch.testing.assert_close(out, ref)
-
-    if provider == "custom":
-        # Custom CUDA kernel — single fused kernel.
-        fn = lambda: torch.ops._C.relu_squared(out, x)
-    elif provider == "native":
-        # forward_native, eager — relu and square as separate ops.
-        fn = lambda: native(x)
-    elif provider == "native_compiled":
-        # forward_native under torch.compile — Inductor fuses relu+square.
-        # This is the real production baseline (custom ops are off when
-        # Inductor is enabled), so it is the comparison reviewers care about.
-        compiled = torch.compile(native)
-        compiled(x)  # warm up / trigger compilation before timing
-        fn = lambda: compiled(x)
-
-    ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
-        fn, quantiles=[0.5, 0.2, 0.8]
-    )
-    return ms, max_ms, min_ms
-
-
-if __name__ == "__main__":
-    parser = FlexibleArgumentParser(
-        description="Benchmark ReLUSquaredActivation: custom kernel vs native."
-    )
-    parser.add_argument(
-        "--dtype",
-        type=str,
-        choices=["half", "bfloat16", "float"],
-        default="bfloat16",
-    )
-    args = parser.parse_args()
-
-    dtype = STR_DTYPE_TO_TORCH_DTYPE[args.dtype]
-
-    perf_report = triton.testing.perf_report(
-        triton.testing.Benchmark(
-            x_names=["batch_size", "seq_len", "intermediate_size"],
-            x_vals=configs,
-            line_arg="provider",
-            line_vals=["custom", "native_compiled", "native"],
-            line_names=[
-                "Custom Kernel",
-                "Native (torch.compile)",
-                "Native (eager)",
-            ],
-            styles=[("blue", "-"), ("green", "-"), ("red", "-")],
-            ylabel="ms",
-            plot_name="relu_squared-eager-performance",
-            args={},
-        )
-    )
-
-    perf_report(
-        lambda batch_size, seq_len, intermediate_size, provider: benchmark_relu_squared(
-            batch_size, seq_len, intermediate_size, provider, dtype
-        )
-    ).run(print_data=True)
@@ -19,7 +19,6 @@ from vllm.model_executor.layers.quantization.utils.fp8_utils import (
 from vllm.platforms import current_platform
 from vllm.triton_utils import triton
 from vllm.utils.argparse_utils import FlexibleArgumentParser
-from vllm.utils.platform_utils import get_device_name_as_file_name

 mp.set_start_method("spawn", force=True)

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

    set(FlashMLA_SOURCES
-        ${flashmla_SOURCE_DIR}/csrc/api/api.cpp
+        ${flashmla_SOURCE_DIR}/csrc/torch_api.cpp

        # Misc kernels for decoding
        ${flashmla_SOURCE_DIR}/csrc/smxx/decode/get_decoding_sched_meta/get_decoding_sched_meta.cu
@@ -128,7 +128,6 @@ if(FLASH_MLA_ARCHS)

    set(FlashMLA_Extension_INCLUDES
        ${flashmla_SOURCE_DIR}/csrc
-        ${flashmla_SOURCE_DIR}/csrc/kerutils/include
        ${flashmla_SOURCE_DIR}/csrc/extension/sm90/dense_fp8/
        ${flashmla_SOURCE_DIR}/csrc/cutlass/include
        ${flashmla_SOURCE_DIR}/csrc/cutlass/tools/util/include
@@ -153,18 +152,15 @@ if(FLASH_MLA_ARCHS)
        USE_SABI 3
        WITH_SOABI)

-    # Enable C++20 for the FlashMLA sources (required for std::span, requires, etc.)
+    # Keep Stable ABI for the module, but *not* for CUDA/C++ files.
+    # This prevents Py_LIMITED_API from affecting nvcc and C++ compiles.
+    # Also enable C++20 for the FlashMLA sources (required for std::span, requires, etc.)
    target_compile_options(_flashmla_C PRIVATE
+        $<$<COMPILE_LANGUAGE:CUDA>:-UPy_LIMITED_API>
+        $<$<COMPILE_LANGUAGE:CXX>:-UPy_LIMITED_API>
        $<$<COMPILE_LANGUAGE:CXX>:-std=c++20>
        $<$<COMPILE_LANGUAGE:CUDA>:-std=c++20>)

-    # _flashmla_C is now ABI-stable torch 2.11+
-    target_compile_definitions(_flashmla_C PRIVATE
-        TORCH_TARGET_VERSION=0x020B000000000000ULL)
-    if(VLLM_GPU_LANG STREQUAL "CUDA")
-        target_compile_definitions(_flashmla_C PRIVATE USE_CUDA)
-    endif()
-
    define_extension_target(
        _flashmla_extension_C
        DESTINATION vllm
@@ -176,12 +172,11 @@ if(FLASH_MLA_ARCHS)
        USE_SABI 3
        WITH_SOABI)

-    # _flashmla_extension_C is now ABI-stable w/ torch 2.11+
-    target_compile_definitions(_flashmla_extension_C PRIVATE
-        TORCH_TARGET_VERSION=0x020B000000000000ULL)
-    if(VLLM_GPU_LANG STREQUAL "CUDA")
-        target_compile_definitions(_flashmla_extension_C PRIVATE USE_CUDA)
-    endif()
+    # Keep Stable ABI for the module, but *not* for CUDA/C++ files.
+    # This prevents Py_LIMITED_API from affecting nvcc and C++ compiles.
+    target_compile_options(_flashmla_extension_C PRIVATE
+        $<$<COMPILE_LANGUAGE:CUDA>:-UPy_LIMITED_API>
+        $<$<COMPILE_LANGUAGE:CXX>:-UPy_LIMITED_API>)
 else()
    message(STATUS "FlashMLA will not compile: unsupported CUDA architecture ${CUDA_ARCHS}")
    # Create empty targets for setup.py on unsupported systems
@@ -1,50 +0,0 @@
-include(FetchContent)
-
-if(DEFINED ENV{TML_FA4_SRC_DIR})
-  set(TML_FA4_SRC_DIR $ENV{TML_FA4_SRC_DIR})
-endif()
-
-if(TML_FA4_SRC_DIR)
-  FetchContent_Declare(
-    tml_fa4
-    SOURCE_DIR ${TML_FA4_SRC_DIR}
-    CONFIGURE_COMMAND ""
-    BUILD_COMMAND "")
-else()
-  FetchContent_Declare(
-    tml_fa4
-    GIT_REPOSITORY https://github.com/vllm-project/tml-fa4.git
-    GIT_TAG 13374f0c855acc1add1bf30444bd67aebbc24a8e
-    GIT_PROGRESS TRUE
-    CONFIGURE_COMMAND ""
-    BUILD_COMMAND "")
-endif()
-
-FetchContent_GetProperties(tml_fa4)
-if(NOT tml_fa4_POPULATED)
-  FetchContent_Populate(tml_fa4)
-endif()
-message(STATUS "tml-fa4 is available at ${tml_fa4_SOURCE_DIR}")
-
-add_custom_target(tml_fa4)
-
-# Install into a private namespace so this implementation cannot shadow the
-# flash_attn package used by vLLM's standard attention backends.
-install(CODE "
-  file(GLOB_RECURSE TML_FA4_PY_FILES
-    \"${tml_fa4_SOURCE_DIR}/flash_attn/cute/*.py\")
-  foreach(SRC_FILE \${TML_FA4_PY_FILES})
-    file(RELATIVE_PATH REL_PATH
-      \"${tml_fa4_SOURCE_DIR}/flash_attn/cute\" \${SRC_FILE})
-    set(DST_FILE
-      \"\${CMAKE_INSTALL_PREFIX}/vllm/third_party/tml_fa4/\${REL_PATH}\")
-    get_filename_component(DST_DIR \${DST_FILE} DIRECTORY)
-    file(MAKE_DIRECTORY \${DST_DIR})
-    file(READ \${SRC_FILE} FILE_CONTENTS)
-    string(REPLACE
-      \"flash_attn.cute\"
-      \"vllm.third_party.tml_fa4\"
-      FILE_CONTENTS \"\${FILE_CONTENTS}\")
-    file(WRITE \${DST_FILE} \"\${FILE_CONTENTS}\")
-  endforeach()
-" COMPONENT tml_fa4)
@@ -39,7 +39,7 @@ else()
  FetchContent_Declare(
          vllm-flash-attn
          GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
-          GIT_TAG caaa4eb59845388a20b1f435ecaafb4bd9517ad8
+          GIT_TAG 2c839c33742309ec41e620bf837495ec9926c56e
          GIT_PROGRESS TRUE
          # Don't share the vllm-flash-attn build between build types
          BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn
@@ -323,6 +323,8 @@ class AttentionImpl<ISA::VSX, scalar_t, head_dim> {
      const int64_t num_blocks_stride, const int64_t cache_head_num_stride,
      const int64_t block_size, const int64_t block_size_stride,
      const float k_inv = 0.0f, const float v_inv = 0.0f) {
+    // k_inv and v_inv are unused on VSX: FP8 KV cache is not supported on
+    // PowerPC. The parameters are present to match the common interface.
 #pragma omp parallel for collapse(2)
    for (int64_t token_idx = 0; token_idx < token_num; ++token_idx) {
      for (int64_t head_idx = 0; head_idx < head_num; ++head_idx) {
@@ -4,7 +4,7 @@
 #if defined(__x86_64__)
  // x86 implementation
  #include "cpu_types_x86.hpp"
-#elif defined(__powerpc__)
+#elif defined(__POWER9_VECTOR__)
  // ppc implementation
  #include "cpu_types_vsx.hpp"
 #elif defined(__s390x__)
@@ -41,4 +41,4 @@ inline int get_max_threads() {
 }
 }  // namespace cpu_utils

-#endif
+#endif
@@ -344,133 +344,53 @@ struct FP32Vec8 : public Vec<FP32Vec8> {

    return result;
  }
+
  FP32Vec8 exp() const {
-    f32x4x2_t out;
-    const __vector float log2e = vec_splats(1.44269504088896341f);
-    const __vector float one = vec_splats(1.0f);
-    const __vector float min_x = vec_splats(-87.3f);
-    const __vector float max_x = vec_splats(88.7f);
-
-    // 5th-degree minimax polynomial for 2^r (r in [0,1))
-    const __vector float c1 = vec_splats(0.6931471805599453f);
-    const __vector float c2 = vec_splats(0.240226506959101f);
-    const __vector float c3 = vec_splats(0.05550410866482158f);
-    const __vector float c4 = vec_splats(0.009618129107628477f);
-    const __vector float c5 = vec_splats(0.0013333558146428443f);
-
-    for (int i = 0; i < 2; i++) {
-      __vector float x = reg.val[i];
-      x = vec_max(x, min_x);
-      x = vec_min(x, max_x);
-
-      __vector float y = vec_mul(x, log2e);
-
-      __vector float kf = vec_floor(y);
-      __vector float r = vec_sub(y, kf);
-
-      // Convert float to signed integer. Use vec_cts for PowerPC AltiVec
-      // compatibility.
-      __vector signed int k = vec_cts(kf, 0);
-      const __vector signed int min_k = vec_splats((signed int)-126);
-      const __vector signed int max_k = vec_splats((signed int)127);
-      k = vec_min(vec_max(k, min_k), max_k);
-
-      // Build 2^k from exponent bits
-      __vector signed int exp_int = vec_add(k, vec_splats((signed int)127));
-      __vector unsigned int bits = (__vector unsigned int)exp_int;
-      bits = vec_sl(bits, vec_splats((unsigned int)23));
-      __vector float pow2k = (__vector float)bits;
-
-      // Improved minimax polynomial
-      __vector float poly = vec_madd(c5, r, c4);
-      poly = vec_madd(poly, r, c3);
-      poly = vec_madd(poly, r, c2);
-      poly = vec_madd(poly, r, c1);
-      poly = vec_madd(poly, r, one);
-
-      out.val[i] = vec_mul(pow2k, poly);
-    }
-    return FP32Vec8(out);
+    // TODO: Vectorize this
+    AliasReg ar;
+    ar.reg = reg;
+    f32x4x4_t ret;
+    ret.val[0][0] = std::exp(ar.values[0]);
+    ret.val[0][1] = std::exp(ar.values[1]);
+    ret.val[0][2] = std::exp(ar.values[2]);
+    ret.val[0][3] = std::exp(ar.values[3]);
+    ret.val[1][0] = std::exp(ar.values[4]);
+    ret.val[1][1] = std::exp(ar.values[5]);
+    ret.val[1][2] = std::exp(ar.values[6]);
+    ret.val[1][3] = std::exp(ar.values[7]);
+    return FP32Vec8(f32x4x2_t({ret.val[0], ret.val[1]}));
  }

  FP32Vec8 tanh() const {
-    const __vector float one = vec_splats(1.0f);
-    const __vector float two = vec_splats(2.0f);
-    const __vector float zero = vec_splats(0.0f);
-    const __vector float sat = vec_splats(9.0f);
-
-    f32x4x2_t out;
-
-    for (int i = 0; i < 2; i++) {
-      __vector float x = reg.val[i];
-      __vector float ax = vec_abs(x);
-
-      __vector bool int mask = vec_cmpge(x, zero);
-      __vector float sign = vec_sel(vec_splats(-1.0f), one, mask);
-
-      __vector bool int saturated = vec_cmpge(ax, sat);
-
-      __vector float two_x = vec_mul(x, two);
-      f32x4x2_t tmp;
-      tmp.val[0] = two_x;
-      tmp.val[1] = two_x;
-      FP32Vec8 temp_vec(tmp);
-      vector float e = temp_vec.exp().reg.val[0];
-
-      vector float num = vec_sub(e, one);
-      vector float den = vec_add(e, one);
-      vector float t = vec_div(num, den);
-
-      out.val[i] = vec_sel(t, sign, saturated);
-    }
-    return FP32Vec8(out);
+    // TODO: Vectorize this
+    AliasReg ar;
+    ar.reg = reg;
+    f32x4x4_t ret;
+    ret.val[0][0] = std::tanh(ar.values[0]);
+    ret.val[0][1] = std::tanh(ar.values[1]);
+    ret.val[0][2] = std::tanh(ar.values[2]);
+    ret.val[0][3] = std::tanh(ar.values[3]);
+    ret.val[1][0] = std::tanh(ar.values[4]);
+    ret.val[1][1] = std::tanh(ar.values[5]);
+    ret.val[1][2] = std::tanh(ar.values[6]);
+    ret.val[1][3] = std::tanh(ar.values[7]);
+    return FP32Vec8(f32x4x2_t({ret.val[0], ret.val[1]}));
  }

  FP32Vec8 er() const {
-    const vector float a1 = vec_splats(0.254829592f);
-    const vector float a2 = vec_splats(-0.284496736f);
-    const vector float a3 = vec_splats(1.421413741f);
-    const vector float a4 = vec_splats(-1.453152027f);
-    const vector float a5 = vec_splats(1.061405429f);
-    const vector float p = vec_splats(0.3275911f);
-    const vector float one = vec_splats(1.0f);
-    const vector float zero = vec_splats(0.0f);
-    const vector float sat = vec_splats(6.0f);
-
-    f32x4x2_t ret;
-
-    for (int i = 0; i < 2; i++) {
-      vector float x = reg.val[i];
-      vector float ax = vec_abs(x);
-
-      vector bool int mask = vec_cmpge(x, zero);
-      vector float sign = vec_sel(vec_splats(-1.0f), one, mask);
-
-      vector bool int saturated = vec_cmpge(ax, sat);
-
-      vector float t = vec_div(one, vec_madd(p, ax, one));
-
-      vector float poly = a5;
-      poly = vec_madd(poly, t, a4);
-      poly = vec_madd(poly, t, a3);
-      poly = vec_madd(poly, t, a2);
-      poly = vec_madd(poly, t, a1);
-      poly = vec_mul(poly, t);
-
-      vector float x_squared = vec_mul(x, x);
-      vector float neg_x_squared = vec_mul(vec_splats(-1.0f), x_squared);
-      f32x4x2_t tmp;
-      tmp.val[0] = neg_x_squared;
-      tmp.val[1] = neg_x_squared;
-      FP32Vec8 exp_input(tmp);
-      vector float exp_term = exp_input.exp().reg.val[0];
-
-      vector float y = vec_nmsub(poly, exp_term, one);
-      vector float erf_val = vec_mul(sign, y);
-
-      ret.val[i] = vec_sel(erf_val, sign, saturated);
-    }
-    return FP32Vec8(ret);
+    // TODO: Vectorize this
+    AliasReg ar;
+    ar.reg = reg;
+    f32x4x4_t ret;
+    ret.val[0][0] = std::erf(ar.values[0]);
+    ret.val[0][1] = std::erf(ar.values[1]);
+    ret.val[0][2] = std::erf(ar.values[2]);
+    ret.val[0][3] = std::erf(ar.values[3]);
+    ret.val[1][0] = std::erf(ar.values[4]);
+    ret.val[1][1] = std::erf(ar.values[5]);
+    ret.val[1][2] = std::erf(ar.values[6]);
+    ret.val[1][3] = std::erf(ar.values[7]);
+    return FP32Vec8(f32x4x2_t({ret.val[0], ret.val[1]}));
  }

  FP32Vec8 operator*(const FP32Vec8& b) const {
@@ -1116,164 +1116,6 @@ void fused_sigmoid_gating_delta_rule_update_kernel_impl(
  });
 }

-// Speculative-decode variant: processes a varlen batch where each sequence has
-// ``q_len`` draft tokens, runs the recurrence sequentially over those tokens
-// (inside the kernel, so one dispatch handles the whole draft block), reads the
-// initial state from cache slot ``num_accepted-1`` and stores the state *after*
-// token ``t`` into cache slot ``t`` (multi-slot rollback, matching the GPU
-// kernel). Parallelized over (sequence, v_head); the per-sequence token loop is
-// sequential as required by the recurrence.
-template <typename scalar_t, typename param_t>
-void fused_sigmoid_gating_delta_rule_update_spec_kernel_impl(
-    const scalar_t* __restrict__ q_ptr,  // [T, HK, EK]
-    const scalar_t* __restrict__ k_ptr,  // [T, HK, EK]
-    const scalar_t* __restrict__ v_ptr,  // [T, HV, EV]
-    const param_t* __restrict__ A_log_ptr,
-    const scalar_t* __restrict__ a_ptr,  // [T, HV]
-    const scalar_t* __restrict__ dt_bias_ptr,
-    const scalar_t* __restrict__ b_ptr,  // [T, HV]
-    const int32_t* __restrict__ spec_indices_ptr,  // [N, S]
-    const int32_t* __restrict__ num_accepted_ptr,  // [N]
-    const int32_t* __restrict__ cu_seqlens_ptr,     // [N + 1]
-    float* __restrict__ state_ptr,
-    scalar_t* __restrict__ o_ptr,  // [T, HV, EV]
-    float* __restrict__ qk_scale_buf,  // [2, T, HK]
-    int64_t total_tokens,
-    int64_t batch_size,
-    int64_t spec_stride,
-    int64_t num_heads,
-    int64_t head_dim,
-    int64_t v_num_heads,
-    int64_t v_head_dim,
-    int64_t q_strideT,
-    int64_t q_strideH,
-    int64_t k_strideT,
-    int64_t k_strideH,
-    int64_t v_strideT,
-    int64_t v_strideH,
-    int64_t state_slot_stride,
-    bool use_qk_l2norm_in_kernel,
-    double softplus_threshold) {
-  using bVec = at::vec::Vectorized<scalar_t>;
-  using fVec = at::vec::Vectorized<float>;
-  constexpr int64_t VecSize = bVec::size();
-  constexpr int64_t fVecSize = fVec::size();
-  int64_t group_size = v_num_heads / num_heads;
-  double scale = 1 / std::sqrt((double)head_dim);
-  fVec scale_vec = fVec((float)scale);
-
-  if (use_qk_l2norm_in_kernel) {
-    float eps = 1e-5f;
-    at::parallel_for(0, total_tokens * num_heads, 0, [&](int64_t begin, int64_t end) {
-      for (int64_t i = begin; i < end; ++i) {
-        int64_t ti = i / num_heads;
-        int64_t ni = i % num_heads;
-        const scalar_t* qp = q_ptr + ti * q_strideT + ni * q_strideH;
-        const scalar_t* kp = k_ptr + ti * k_strideT + ni * k_strideH;
-        float sq = 0.f, sk = 0.f;
-        for (int64_t d = 0; d < head_dim; ++d) {
-          float qv = (float)qp[d];
-          sq += qv * qv;
-          float kv = (float)kp[d];
-          sk += kv * kv;
-        }
-        qk_scale_buf[ti * num_heads + ni] = 1.f / std::sqrt(sq + eps);
-        qk_scale_buf[total_tokens * num_heads + ti * num_heads + ni] = 1.f / std::sqrt(sk + eps);
-      }
-    });
-  }
-
-  at::parallel_for(0, batch_size * v_num_heads, 0, [&](int64_t begin, int64_t end) {
-    for (int64_t idx = begin; idx < end; ++idx) {
-      int64_t bi = idx / v_num_heads;
-      int64_t ni = idx % v_num_heads;
-      int64_t kh = ni / group_size;
-      int64_t q_start = cu_seqlens_ptr[bi];
-      int64_t q_len = cu_seqlens_ptr[bi + 1] - q_start;
-      if (q_len <= 0) {
-        continue;
-      }
-      int64_t acc = (int64_t)num_accepted_ptr[bi];
-      // Clamp acc-1 to >=0: when num_accepted is 0 the unclamped index reads
-      // out of bounds and yields an arbitrary prev_slot used to index the SSM
-      // state. Mirrors the GPU guard tl.maximum(num_accepted - 1, 0).
-      int64_t prev_slot =
-          (int64_t)spec_indices_ptr[bi * spec_stride + (acc > 0 ? acc - 1 : 0)];
-      for (int64_t t = 0; t < q_len; ++t) {
-        int64_t cur_slot = (int64_t)spec_indices_ptr[bi * spec_stride + t];
-        int64_t token = q_start + t;
-        const float* src = state_ptr + prev_slot * state_slot_stride + ni * head_dim * v_head_dim;
-        float* dst = state_ptr + cur_slot * state_slot_stride + ni * head_dim * v_head_dim;
-        float g_val = -std::exp((float)A_log_ptr[ni]) *
-            softplus((float)a_ptr[token * v_num_heads + ni] + (float)dt_bias_ptr[ni], softplus_threshold);
-        float g_val_exp = std::exp(g_val);
-        fVec g_val_exp_vec = fVec(g_val_exp);
-        float beta_val = 1.f / (1.f + std::exp(-(float)b_ptr[token * v_num_heads + ni]));
-        fVec beta_vec = fVec(beta_val);
-        int64_t q_offset = token * q_strideT + kh * q_strideH;
-        int64_t k_offset = token * k_strideT + kh * k_strideH;
-        float q_scale = use_qk_l2norm_in_kernel ? qk_scale_buf[token * num_heads + kh] : 1.f;
-        float k_scale =
-            use_qk_l2norm_in_kernel ? qk_scale_buf[total_tokens * num_heads + token * num_heads + kh] : 1.f;
-        int64_t v_offset = token * v_strideT + ni * v_strideH;
-        int64_t o_offset = (token * v_num_heads + ni) * v_head_dim;
-        int64_t dvi = 0;
-        for (; dvi <= v_head_dim - VecSize; dvi += VecSize) {
-          fVec kv_mem_vec0 = fVec(0.f);
-          fVec kv_mem_vec1 = fVec(0.f);
-          for (int di = 0; di < head_dim; ++di) {
-            fVec k_val_vec = fVec((float)k_ptr[k_offset + di] * k_scale);
-            fVec sv0 = fVec::loadu(src + di * v_head_dim + dvi);
-            fVec sv1 = fVec::loadu(src + di * v_head_dim + dvi + fVecSize);
-            kv_mem_vec0 = kv_mem_vec0 + sv0 * g_val_exp_vec * k_val_vec;
-            kv_mem_vec1 = kv_mem_vec1 + sv1 * g_val_exp_vec * k_val_vec;
-          }
-          bVec v_bvec = bVec::loadu(v_ptr + v_offset + dvi);
-          fVec v_vec0, v_vec1;
-          std::tie(v_vec0, v_vec1) = at::vec::convert_to_float(v_bvec);
-          fVec dt_vec0 = (v_vec0 - kv_mem_vec0) * beta_vec;
-          fVec dt_vec1 = (v_vec1 - kv_mem_vec1) * beta_vec;
-          fVec o_vec0 = fVec(0.f);
-          fVec o_vec1 = fVec(0.f);
-          for (int di = 0; di < head_dim; ++di) {
-            fVec q_vec = fVec((float)q_ptr[q_offset + di] * q_scale);
-            fVec k_vec = fVec((float)k_ptr[k_offset + di] * k_scale);
-            fVec sv0 = fVec::loadu(src + di * v_head_dim + dvi);
-            fVec sv1 = fVec::loadu(src + di * v_head_dim + dvi + fVecSize);
-            sv0 = sv0 * g_val_exp_vec + k_vec * dt_vec0;
-            sv1 = sv1 * g_val_exp_vec + k_vec * dt_vec1;
-            o_vec0 = o_vec0 + sv0 * q_vec * scale_vec;
-            o_vec1 = o_vec1 + sv1 * q_vec * scale_vec;
-            sv0.store(dst + di * v_head_dim + dvi);
-            sv1.store(dst + di * v_head_dim + dvi + fVecSize);
-          }
-          bVec o_vec = at::vec::convert_from_float<scalar_t>(o_vec0, o_vec1);
-          o_vec.store(o_ptr + o_offset + dvi);
-        }
-        for (; dvi < v_head_dim; ++dvi) {
-          float kv_mem_val = 0.f;
-          for (int di = 0; di < head_dim; ++di) {
-            float k_val = (float)k_ptr[k_offset + di] * k_scale;
-            kv_mem_val += src[di * v_head_dim + dvi] * g_val_exp * k_val;
-          }
-          float v_val = (float)v_ptr[v_offset + dvi];
-          float dt_val = (v_val - kv_mem_val) * beta_val;
-          float o_val = 0.f;
-          for (int di = 0; di < head_dim; ++di) {
-            float q_val = (float)q_ptr[q_offset + di] * q_scale;
-            float k_val = (float)k_ptr[k_offset + di] * k_scale;
-            float ns = src[di * v_head_dim + dvi] * g_val_exp + k_val * dt_val;
-            dst[di * v_head_dim + dvi] = ns;
-            o_val += ns * q_val * scale;
-          }
-          o_ptr[o_offset + dvi] = (scalar_t)o_val;
-        }
-        prev_slot = cur_slot;
-      }
-    }
-  });
-}
-
 template <typename scalar_t>
 void fused_gdn_gating_kernel_impl(
    float* __restrict__ A_log,
@@ -1658,103 +1500,6 @@ at::Tensor fused_sigmoid_gating_delta_rule_update_cpu(
  return core_attn_out;
 }

-// Speculative-decode update (multi-token, multi-slot rollback).
-// q: [T, HK, EK]  k: [T, HK, EK]  v: [T, HV, EV]
-// a: [T, HV]  b: [T, HV]
-// initial_state_source: [N_slots, HV, EK, EV] FP32 (updated in place)
-// spec_state_indices: [batch, S] INT32 (S = num_spec + 1)
-// num_accepted_tokens: [batch] INT32
-// cu_seqlens: [batch + 1] INT32
-// Returns output: [T, HV, EV]
-at::Tensor fused_sigmoid_gating_delta_rule_update_spec_cpu(
-    const at::Tensor& A_log,
-    const at::Tensor& dt_bias,
-    const at::Tensor& q,
-    const at::Tensor& k,
-    const at::Tensor& v,
-    const at::Tensor& a,
-    const at::Tensor& b,
-    at::Tensor& initial_state_source,
-    const at::Tensor& spec_state_indices,
-    const at::Tensor& num_accepted_tokens,
-    const at::Tensor& cu_seqlens,
-    bool use_qk_l2norm_in_kernel,
-    double softplus_beta = 1.0,
-    double softplus_threshold = 20.0) {
-  CHECK_DIM(3, q);
-  CHECK_DIM(3, v);
-  CHECK_LAST_DIM_CONTIGUOUS_INPUT(q);
-  int64_t total_tokens = q.size(0);
-  int64_t num_heads = q.size(1);
-  int64_t head_dim = q.size(2);
-  int64_t v_num_heads = v.size(1);
-  int64_t v_head_dim = v.size(2);
-  int64_t batch_size = cu_seqlens.size(0) - 1;
-  int64_t spec_stride = spec_state_indices.stride(0);
-  CHECK_INPUT_SHAPE_DTYPE<true>(k, {total_tokens, num_heads, head_dim}, q.scalar_type());
-  CHECK_INPUT_SHAPE_DTYPE<true>(v, {total_tokens, v_num_heads, v_head_dim}, q.scalar_type());
-  CHECK_INPUT_SHAPE_DTYPE<true>(a, {total_tokens, v_num_heads}, q.scalar_type());
-  CHECK_INPUT_SHAPE_DTYPE<true>(b, {total_tokens, v_num_heads}, q.scalar_type());
-  CHECK_INPUT_SHAPE_DTYPE<true>(dt_bias, {v_num_heads}, q.scalar_type());
-  CHECK_INPUT_SHAPE_DTYPE<true>(num_accepted_tokens, {batch_size}, at::kInt);
-  CHECK_INPUT_SHAPE_DTYPE<true>(cu_seqlens, {batch_size + 1}, at::kInt);
-  CHECK_EQ(v_num_heads % num_heads, 0);
-  TORCH_CHECK(A_log.sizes() == at::IntArrayRef({v_num_heads}));
-  CHECK_INPUT_SHAPE_DTYPE<true>(
-      initial_state_source,
-      {initial_state_source.size(0), v_num_heads, head_dim, v_head_dim},
-      at::kFloat);
-  TORCH_CHECK(initial_state_source.size(0) >= batch_size,
-      "initial_state_source capacity too small: size(0)=",
-      initial_state_source.size(0), ", batch_size=", batch_size);
-
-  int64_t q_strideT = q.stride(0);
-  int64_t q_strideH = q.stride(1);
-  int64_t k_strideT = k.stride(0);
-  int64_t k_strideH = k.stride(1);
-  int64_t v_strideT = v.stride(0);
-  int64_t v_strideH = v.stride(1);
-  int64_t state_slot_stride = initial_state_source.stride(0);
-
-  at::Tensor o = at::empty({total_tokens, v_num_heads, v_head_dim}, q.options());
-  at::Tensor qk_scale_buf = at::empty({2, total_tokens, num_heads}, at::kFloat);
-
-  CPU_DISPATCH_REDUCED_FLOATING_TYPES_EXT(
-      q.scalar_type(), A_log.scalar_type(), "fused_sigmoid_gating_delta_rule_update_spec_kernel_impl", [&] {
-        fused_sigmoid_gating_delta_rule_update_spec_kernel_impl<scalar_t, param_t>(
-            q.data_ptr<scalar_t>(),
-            k.data_ptr<scalar_t>(),
-            v.data_ptr<scalar_t>(),
-            A_log.data_ptr<param_t>(),
-            a.data_ptr<scalar_t>(),
-            dt_bias.data_ptr<scalar_t>(),
-            b.data_ptr<scalar_t>(),
-            spec_state_indices.data_ptr<int32_t>(),
-            num_accepted_tokens.data_ptr<int32_t>(),
-            cu_seqlens.data_ptr<int32_t>(),
-            initial_state_source.data_ptr<float>(),
-            o.data_ptr<scalar_t>(),
-            qk_scale_buf.data_ptr<float>(),
-            total_tokens,
-            batch_size,
-            spec_stride,
-            num_heads,
-            head_dim,
-            v_num_heads,
-            v_head_dim,
-            q_strideT,
-            q_strideH,
-            k_strideT,
-            k_strideH,
-            v_strideT,
-            v_strideH,
-            state_slot_stride,
-            use_qk_l2norm_in_kernel,
-            softplus_threshold);
-      });
-  return o;
-}
-
 // A_log: [num_v_heads]
 // a: [batch, num_v_heads]
 // b: [batch, num_v_heads]
@@ -120,14 +120,6 @@ at::Tensor fused_sigmoid_gating_delta_rule_update_cpu(
    bool use_qk_l2norm_in_kernel, double softplus_beta = 1.0,
    double softplus_threshold = 20.0);

-at::Tensor fused_sigmoid_gating_delta_rule_update_spec_cpu(
-    const at::Tensor& A_log, const at::Tensor& dt_bias, const at::Tensor& q,
-    const at::Tensor& k, const at::Tensor& v, const at::Tensor& a,
-    const at::Tensor& b, at::Tensor& initial_state_source,
-    const at::Tensor& spec_state_indices, const at::Tensor& num_accepted_tokens,
-    const at::Tensor& cu_seqlens, bool use_qk_l2norm_in_kernel,
-    double softplus_beta = 1.0, double softplus_threshold = 20.0);
-
 std::tuple<at::Tensor, at::Tensor> fused_gdn_gating_cpu(
    const at::Tensor& A_log, const at::Tensor& a, const at::Tensor& b,
    const at::Tensor& dt_bias);
@@ -516,15 +508,6 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
      "softplus_threshold=20.0) -> Tensor");
  ops.impl("fused_sigmoid_gating_delta_rule_update_cpu", torch::kCPU,
           &fused_sigmoid_gating_delta_rule_update_cpu);
-  ops.def(
-      "fused_sigmoid_gating_delta_rule_update_spec_cpu(Tensor A_log, Tensor "
-      "dt_bias, Tensor q, Tensor k, Tensor v, Tensor a, Tensor b, "
-      "Tensor(a!) initial_state_source, Tensor spec_state_indices, "
-      "Tensor num_accepted_tokens, Tensor cu_seqlens, bool "
-      "use_qk_l2norm_in_kernel, float softplus_beta=1.0, float "
-      "softplus_threshold=20.0) -> Tensor");
-  ops.impl("fused_sigmoid_gating_delta_rule_update_spec_cpu", torch::kCPU,
-           &fused_sigmoid_gating_delta_rule_update_spec_cpu);
  ops.def(
      "fused_gdn_gating_cpu(Tensor A_log, Tensor a, Tensor b, Tensor dt_bias) "
      "-> (Tensor, Tensor)");
@@ -570,8 +553,7 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
 #endif

  // fused moe
-#if defined(__AVX512F__) || \
-    (defined(__aarch64__) && !defined(__APPLE__) && defined(ARM_BF16_SUPPORT))
+#if defined(__AVX512F__) || (defined(ARM_BF16_SUPPORT))
  ops.def(
      "prepack_moe_weight(Tensor weight, Tensor(a1!) packed_weight, str isa) "
      "-> ()");
@@ -582,7 +564,7 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
      "bool skip_weighted, "
      "str act, str isa) -> ()");
  ops.impl("cpu_fused_moe", torch::kCPU, &cpu_fused_moe);
-#endif
+#endif  // #if defined(__AVX512F__) || (defined(ARM_BF16_SUPPORT))
  ops.def(
      "mla_decode_kvcache("
      "   Tensor! out, Tensor query, Tensor kv_cache,"
@@ -76,14 +76,14 @@ inline int64_t get_available_l2_size() {
    if (l2_cache_size == 0) {
      l2_cache_size = 256 * 1024;
    }
-    return static_cast<int64_t>(l2_cache_size) >> 1;
+    return static_cast<int64_t>(l2_cache_size) >> 1;  // use 50% of L2 cache
  }();
  return size;
 #else
  static int64_t size = []() {
    auto caps = at::cpu::get_cpu_capabilities();
    const uint32_t l2_cache_size = caps.at("l2_cache_size").toInt();
-    return l2_cache_size >> 1;
+    return l2_cache_size >> 1;  // use 50% of L2 cache
  }();
  return size;
 #endif
@@ -669,14 +669,6 @@ __device__ __forceinline__ T gelu_quick_kernel(const T& x) {
  return (T)(((float)x) / (1.0f + expf(-1.702f * (float)x)));
 }

-template <typename T>
-__device__ __forceinline__ T relu_squared_kernel(const T& x) {
-  // relu(x)^2 — introduced in https://arxiv.org/abs/2109.08668v2
-  const float f = (float)x;
-  const float val = f > 0.0f ? f : 0.0f;
-  return (T)(val * val);
-}
-
 }  // namespace vllm

 void gelu_new(torch::stable::Tensor& out,    // [..., d]
@@ -696,9 +688,3 @@ void gelu_quick(torch::stable::Tensor& out,    // [..., d]
 {
  LAUNCH_ACTIVATION_KERNEL(vllm::gelu_quick_kernel);
 }
-
-void relu_squared(torch::stable::Tensor& out,    // [..., d]
-                  torch::stable::Tensor& input)  // [..., d]
-{
-  LAUNCH_ACTIVATION_KERNEL(vllm::relu_squared_kernel);
-}
@@ -21,21 +21,6 @@
 #define VLLM_STABLE_DISPATCH_FP8_CASE(enum_type, ...) \
  THO_PRIVATE_CASE_TYPE_USING_HINT(enum_type, fp8_t, __VA_ARGS__)

-// Same idea, for dispatching on an int32/int64 index tensor (e.g. topk_ids)
-// nested inside a value-type dispatch. Named 'idx_t' instead of 'scalar_t'.
-#define VLLM_STABLE_DISPATCH_IDX_CASE(enum_type, ...) \
-  THO_PRIVATE_CASE_TYPE_USING_HINT(enum_type, idx_t, __VA_ARGS__)
-
-#define VLLM_STABLE_DISPATCH_CASE_IDX_TYPES(...)                     \
-  VLLM_STABLE_DISPATCH_IDX_CASE(torch::headeronly::ScalarType::Int,  \
-                                __VA_ARGS__)                         \
-  VLLM_STABLE_DISPATCH_IDX_CASE(torch::headeronly::ScalarType::Long, \
-                                __VA_ARGS__)
-
-#define VLLM_STABLE_DISPATCH_IDX_TYPES(TYPE, NAME, ...) \
-  THO_DISPATCH_SWITCH(TYPE, NAME,                       \
-                      VLLM_STABLE_DISPATCH_CASE_IDX_TYPES(__VA_ARGS__))
-
 #define VLLM_STABLE_DISPATCH_CASE_FLOATING_TYPES(...)                  \
  THO_DISPATCH_CASE(torch::headeronly::ScalarType::Float, __VA_ARGS__) \
  THO_DISPATCH_CASE(torch::headeronly::ScalarType::Half, __VA_ARGS__)  \
@@ -1,18 +1,13 @@
 // SPDX-License-Identifier: Apache-2.0
 // SPDX-FileCopyrightText: Copyright contributors to the vLLM project
 //
-// Router GEMM: activation(T) x weight(fp32) -> fp32, M<=32, for the
-// supported (E, H) pairs listed at the bottom of this file.
+// Router GEMM: activation(T) x weight(fp32) -> fp32, H=3072, E=256, M<=32.
 // Supports bf16 or fp32 activation; weight is always fp32.
 // Adapted from dsv3_router_gemm_float_out.cu.
-// (E=256, H=6144) bf16 uses a B300-tuned wide-block geometry; see
-// invokeFp32RouterGemm.

 #include <cuda_bf16.h>
 #include <cuda_runtime.h>

-#include <type_traits>
-
 // ---------------------------------------------------------------------------
 // Load helpers
 // ---------------------------------------------------------------------------
@@ -78,113 +73,94 @@ __device__ __forceinline__ void load_activation<__nv_bfloat16, 8>(
 // InputT : type of activation (float or __nv_bfloat16)
 // Weight is always fp32; output is always fp32.
 // VPT = 16 / sizeof(InputT):  4 for fp32, 8 for bf16
-// Each block computes kEPB expert columns; wider blocks / kEPB > 1 are
-// selected per (shape, M) in invokeFp32RouterGemm (B300-tuned, see below).
-// kTGroups > 1 splits the tokens across groups of kBlockSize threads within
-// the block: all groups scan the same weight K-slices (group 0 misses to
-// DRAM, later groups hit L1) so weight traffic stays 1x, while per-thread
-// accumulator registers drop by kTGroups (at M=16 the 32 fp32 accumulators
-// push the kernel to 128 regs/thread and 1 block/SM).
-template <typename InputT, int kBlockSize, int kNumTokens, int kEPB,
-          int kNumExperts, int kHiddenDim, int kTGroups = 1>
-__global__ __launch_bounds__(
-    kBlockSize* kTGroups, 1) void fp32_router_gemm_kernel(float* out,
-                                                          InputT const* mat_a,
-                                                          float const* mat_b) {
+template <typename InputT, int kBlockSize, int kNumTokens, int kNumExperts,
+          int kHiddenDim>
+__global__ __launch_bounds__(128, 1) void fp32_router_gemm_kernel(
+    float* out, InputT const* mat_a, float const* mat_b) {
  constexpr int VPT = 16 / sizeof(InputT);
  constexpr int k_elems_per_k_iteration = VPT * kBlockSize;
  constexpr int k_iterations = kHiddenDim / k_elems_per_k_iteration;
-  static_assert(kHiddenDim % k_elems_per_k_iteration == 0);
-  static_assert(kNumTokens % kTGroups == 0);
  constexpr int kWarpSize = 32;
-  constexpr int kNumWarps = kBlockSize / kWarpSize;  // per token group
-  constexpr int kMG = kNumTokens / kTGroups;         // tokens per group
+  constexpr int kNumWarps = kBlockSize / kWarpSize;

-  int const e_base = blockIdx.x * kEPB;
-  int const tid = threadIdx.x % kBlockSize;
-  int const m0 = (threadIdx.x / kBlockSize) * kMG;
+  int const n_idx = blockIdx.x;
+  int const tid = threadIdx.x;
  int const warpId = tid / kWarpSize;
  int const laneId = tid % kWarpSize;

-  float acc[kMG][kEPB] = {};
-  __shared__ float sm_reduction[kNumTokens][kEPB][kNumWarps];
+  float acc[kNumTokens] = {};
+  __shared__ float sm_reduction[kNumTokens][kNumWarps];
+
+  float const* b_col = mat_b + n_idx * kHiddenDim;
+
+  int k_bases[k_iterations];
+#pragma unroll
+  for (int ki = 0; ki < k_iterations; ki++) {
+    k_bases[ki] = ki * k_elems_per_k_iteration + tid * VPT;
+  }

 #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
  cudaGridDependencySynchronize();
-  // Fire the PDL trigger right after our own wait instead of at kernel end:
-  // a gridsync-ing consumer is unaffected (its wait always targets full grid
-  // completion), while a consumer that reads none of our outputs (e.g. the
-  // NVFP4 activation quant, which reads the same hidden_states) can launch
-  // now and fully overlap this kernel's body.
-  cudaTriggerProgrammaticLaunchCompletion();
 #endif

-#pragma unroll
  for (int ki = 0; ki < k_iterations; ki++) {
-    int const k_base = ki * k_elems_per_k_iteration + tid * VPT;
+    int const k_base = k_bases[ki];

-    float b_float[kEPB][VPT];
-#pragma unroll
-    for (int e = 0; e < kEPB; e++) {
-      load_weight<VPT>(mat_b + (e_base + e) * kHiddenDim + k_base, b_float[e]);
-    }
+    float b_float[VPT];
+    load_weight<VPT>(b_col + k_base, b_float);

 #pragma unroll
-    for (int m_idx = 0; m_idx < kMG; m_idx++) {
+    for (int m_idx = 0; m_idx < kNumTokens; m_idx++) {
      float a_float[VPT];
-      load_activation<InputT, VPT>(
-          mat_a + (size_t)(m0 + m_idx) * kHiddenDim + k_base, a_float);
+      load_activation<InputT, VPT>(mat_a + m_idx * kHiddenDim + k_base,
+                                   a_float);
 #pragma unroll
-      for (int e = 0; e < kEPB; e++) {
-#pragma unroll
-        for (int k = 0; k < VPT; k++) {
-          acc[m_idx][e] += a_float[k] * b_float[e][k];
-        }
+      for (int k = 0; k < VPT; k++) {
+        acc[m_idx] += a_float[k] * b_float[k];
      }
    }
  }

  // Warp-level butterfly reduction
 #pragma unroll
-  for (int m = 0; m < kMG; m++) {
-#pragma unroll
-    for (int e = 0; e < kEPB; e++) {
-      float sum = acc[m][e];
-      sum += __shfl_xor_sync(0xffffffff, sum, 16);
-      sum += __shfl_xor_sync(0xffffffff, sum, 8);
-      sum += __shfl_xor_sync(0xffffffff, sum, 4);
-      sum += __shfl_xor_sync(0xffffffff, sum, 2);
-      sum += __shfl_xor_sync(0xffffffff, sum, 1);
-      if (laneId == 0) sm_reduction[m0 + m][e][warpId] = sum;
-    }
+  for (int m = 0; m < kNumTokens; m++) {
+    float sum = acc[m];
+    sum += __shfl_xor_sync(0xffffffff, sum, 16);
+    sum += __shfl_xor_sync(0xffffffff, sum, 8);
+    sum += __shfl_xor_sync(0xffffffff, sum, 4);
+    sum += __shfl_xor_sync(0xffffffff, sum, 2);
+    sum += __shfl_xor_sync(0xffffffff, sum, 1);
+    if (laneId == 0) sm_reduction[m][warpId] = sum;
  }

  __syncthreads();

-  // Parallel finalize: one thread per (m, e) output.
-  for (int idx = threadIdx.x; idx < kNumTokens * kEPB;
-       idx += kBlockSize * kTGroups) {
-    int const m = idx / kEPB;
-    int const e = idx % kEPB;
-    float final_sum = 0.0f;
+  if (tid == 0) {
 #pragma unroll
-    for (int w = 0; w < kNumWarps; w++) final_sum += sm_reduction[m][e][w];
-    out[m * kNumExperts + e_base + e] = final_sum;
+    for (int m = 0; m < kNumTokens; m++) {
+      float final_sum = 0.0f;
+#pragma unroll
+      for (int w = 0; w < kNumWarps; w++) final_sum += sm_reduction[m][w];
+      out[m * kNumExperts + n_idx] = final_sum;
+    }
  }
+
+#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 900)
+  cudaTriggerProgrammaticLaunchCompletion();
+#endif
 }

 // ---------------------------------------------------------------------------
 // Launcher
 // ---------------------------------------------------------------------------

-template <typename InputT, int kBlockSize, int kEPB, int kNumTokens,
-          int kNumExperts, int kHiddenDim, int kTGroups = 1>
-static void launchFp32RouterGemm(float* output, InputT const* mat_a,
-                                 float const* mat_b, cudaStream_t stream) {
-  static_assert(kNumExperts % kEPB == 0);
+template <typename InputT, int kNumTokens, int kNumExperts, int kHiddenDim>
+void invokeFp32RouterGemm(float* output, InputT const* mat_a,
+                          float const* mat_b, cudaStream_t stream) {
+  constexpr int kBlockSize = 128;
  cudaLaunchConfig_t config;
-  config.gridDim = kNumExperts / kEPB;
-  config.blockDim = kBlockSize * kTGroups;
+  config.gridDim = kNumExperts;
+  config.blockDim = kBlockSize;
  config.dynamicSmemBytes = 0;
  config.stream = stream;
  cudaLaunchAttribute attrs[1];
@@ -192,112 +168,15 @@ static void launchFp32RouterGemm(float* output, InputT const* mat_a,
  attrs[0].val.programmaticStreamSerializationAllowed = 1;
  config.numAttrs = 1;
  config.attrs = attrs;
-  cudaLaunchKernelEx(
-      &config,
-      fp32_router_gemm_kernel<InputT, kBlockSize, kNumTokens, kEPB, kNumExperts,
-                              kHiddenDim, kTGroups>,
-      output, mat_a, mat_b);
-}
-
-static bool isBlackwellFamily() {
-  static int sm = []() {
-    int dev = 0, major = 0, minor = 0;
-    cudaGetDevice(&dev);
-    cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, dev);
-    cudaDeviceGetAttribute(&minor, cudaDevAttrComputeCapabilityMinor, dev);
-    return major * 10 + minor;
-  }();
-  return sm >= 100;
-}
-
-template <typename InputT, int kNumTokens, int kNumExperts, int kHiddenDim>
-void invokeFp32RouterGemm(float* output, InputT const* mat_a,
-                          float const* mat_b, cudaStream_t stream) {
-  // Geometry tuned on B300 per supported shape, bf16 activation, under a
-  // production-fidelity harness (CUDA-graph replay, per-layer cold weights).
-  // GLM-5.2 (E=256, H=6144):
-  //   M <= 4        : BS=768, EPB=1 (2.7us vs cast+cuBLAS 8.1us at M=1)
-  //   M in [5, 15]
-  //   or odd        : BS=384, EPB=2 (crossover vs BS=768 measured in (4, 8))
-  //   M >= 16, even : BS=192, EPB=2, 2 token groups (M=16 4.79us vs 5.04,
-  //                   M=24 5.71 vs 6.38, M=32 6.79 vs 7.72; M=12 loses at
-  //                   0.97x, so the boundary is 16).
-  // Only enabled on the Blackwell family where it was validated; Hopper and
-  // other shapes / fp32 activation keep the legacy geometry.
-  if constexpr (std::is_same_v<InputT, __nv_bfloat16> && kNumExperts == 256 &&
-                kHiddenDim == 6144) {
-    if (!isBlackwellFamily()) {
-      launchFp32RouterGemm<InputT, 128, 1, kNumTokens, kNumExperts, kHiddenDim>(
-          output, mat_a, mat_b, stream);
-      return;
-    }
-    if constexpr (kNumTokens <= 4) {
-      launchFp32RouterGemm<InputT, 768, 1, kNumTokens, kNumExperts, kHiddenDim>(
-          output, mat_a, mat_b, stream);
-    } else if constexpr (kNumTokens >= 16 && kNumTokens % 2 == 0) {
-      launchFp32RouterGemm<InputT, 192, 2, kNumTokens, kNumExperts, kHiddenDim,
-                           2>(output, mat_a, mat_b, stream);
-    } else {
-      launchFp32RouterGemm<InputT, 384, 2, kNumTokens, kNumExperts, kHiddenDim>(
-          output, mat_a, mat_b, stream);
-    }
-  } else if constexpr (std::is_same_v<InputT, __nv_bfloat16> &&
-                       kNumExperts == 128 && kHiddenDim == 6144) {
-    // MiniMax-M3. Legacy 128/1 only fills 128 blocks and pays the same
-    // accumulator register cliffs; B300 sweep:
-    //   even M in [6, 10] : BS=384, EPB=1, 2 token groups (1.26-1.43x)
-    //   even M >= 12      : BS=192, EPB=1, 2 token groups (1.59-1.66x at
-    //                       M >= 18; re-measured on B300+B200: 192 also wins
-    //                       M=12/14 by 5-11%% on both, ties 384 at 16)
-    //   M <= 5 / odd      : BS=384, EPB=1 (1.03-1.19x)
-    if (!isBlackwellFamily()) {
-      launchFp32RouterGemm<InputT, 128, 1, kNumTokens, kNumExperts, kHiddenDim>(
-          output, mat_a, mat_b, stream);
-      return;
-    }
-    if constexpr (kNumTokens >= 12 && kNumTokens % 2 == 0) {
-      launchFp32RouterGemm<InputT, 192, 1, kNumTokens, kNumExperts, kHiddenDim,
-                           2>(output, mat_a, mat_b, stream);
-    } else if constexpr (kNumTokens >= 6 && kNumTokens % 2 == 0) {
-      launchFp32RouterGemm<InputT, 384, 1, kNumTokens, kNumExperts, kHiddenDim,
-                           2>(output, mat_a, mat_b, stream);
-    } else {
-      launchFp32RouterGemm<InputT, 384, 1, kNumTokens, kNumExperts, kHiddenDim>(
-          output, mat_a, mat_b, stream);
-    }
-  } else if constexpr (std::is_same_v<InputT, __nv_bfloat16> &&
-                       kNumExperts == 256 && kHiddenDim == 3072) {
-    // MiniMax-M2/M2.5. The 3.1MB weight is latency-floor bound at small M
-    // (legacy already optimal); token groups win only at even M >= 8
-    // (1.05-1.17x). EPB crossover measured between 12 and 16.
-    if (!isBlackwellFamily()) {
-      launchFp32RouterGemm<InputT, 128, 1, kNumTokens, kNumExperts, kHiddenDim>(
-          output, mat_a, mat_b, stream);
-      return;
-    }
-    if constexpr (kNumTokens >= 14 && kNumTokens % 2 == 0) {
-      // M=14 originally measured 0.91x and stayed on legacy; two fresh
-      // sweeps (B300 dev1 + B200) both put 192/2/tg2 ahead by 3.5-4%%.
-      launchFp32RouterGemm<InputT, 192, 2, kNumTokens, kNumExperts, kHiddenDim,
-                           2>(output, mat_a, mat_b, stream);
-    } else if constexpr (kNumTokens >= 8 && kNumTokens <= 12 &&
-                         kNumTokens % 2 == 0) {
-      launchFp32RouterGemm<InputT, 192, 1, kNumTokens, kNumExperts, kHiddenDim,
-                           2>(output, mat_a, mat_b, stream);
-    } else {
-      launchFp32RouterGemm<InputT, 128, 1, kNumTokens, kNumExperts, kHiddenDim>(
-          output, mat_a, mat_b, stream);
-    }
-  } else {
-    launchFp32RouterGemm<InputT, 128, 1, kNumTokens, kNumExperts, kHiddenDim>(
-        output, mat_a, mat_b, stream);
-  }
+  cudaLaunchKernelEx(&config,
+                     fp32_router_gemm_kernel<InputT, kBlockSize, kNumTokens,
+                                             kNumExperts, kHiddenDim>,
+                     output, mat_a, mat_b);
 }

 // ---------------------------------------------------------------------------
 // Explicit instantiations: M=1..32, for both input types, for the supported
-// (E, H) pairs:  (256, 3072) [MiniMax-M2/M2.5],  (128, 6144) [MiniMax-M3]
-// and  (256, 6144) [GLM-5.2].
+// (E, H) pairs:  (256, 3072) [MiniMax-M2/M2.5]  and  (128, 6144) [MiniMax-M3].
 // ---------------------------------------------------------------------------

 #define INSTANTIATE(T, M, E, H)                                    \
@@ -342,8 +221,6 @@ INSTANTIATE_ALL(float, 256, 3072)
 INSTANTIATE_ALL(__nv_bfloat16, 256, 3072)
 INSTANTIATE_ALL(float, 128, 6144)
 INSTANTIATE_ALL(__nv_bfloat16, 128, 6144)
-INSTANTIATE_ALL(float, 256, 6144)
-INSTANTIATE_ALL(__nv_bfloat16, 256, 6144)

 #undef INSTANTIATE_ALL
 #undef INSTANTIATE
@@ -25,12 +25,10 @@ inline int getSMVersion() {
 static constexpr int FP32_MAX_TOKENS = 32;

 // Supported (hidden_dim, num_experts) pairs (must match the instantiations in
-// fp32_router_gemm.cu): (3072, 256) for MiniMax-M2/M2.5, (6144, 128) for M3,
-// (6144, 256) for GLM-5.2.
+// fp32_router_gemm.cu): (3072, 256) for MiniMax-M2/M2.5, (6144, 128) for M3.
 static inline bool fp32_router_gemm_supported(int hidden_dim, int num_experts) {
  return (hidden_dim == 3072 && num_experts == 256) ||
-         (hidden_dim == 6144 && num_experts == 128) ||
-         (hidden_dim == 6144 && num_experts == 256);
+         (hidden_dim == 6144 && num_experts == 128);
 }

 // Forward declarations — 4 template params must match fp32_router_gemm.cu
@@ -79,9 +77,6 @@ void dispatchFp32RouterGemm(int num_experts, int hidden_dim, int num_tokens,
  } else if (num_experts == 128 && hidden_dim == 6144) {
    Fp32LoopUnroller<InputT, 128, 6144, 1, FP32_MAX_TOKENS>::unroll(
        num_tokens, output, mat_a, mat_b, stream);
-  } else if (num_experts == 256 && hidden_dim == 6144) {
-    Fp32LoopUnroller<InputT, 256, 6144, 1, FP32_MAX_TOKENS>::unroll(
-        num_tokens, output, mat_a, mat_b, stream);
  } else {
    throw std::invalid_argument(
        "fp32_router_gemm: unsupported (hidden_dim, num_experts) pair");
@@ -116,7 +111,7 @@ void fp32_router_gemm(
  STD_TORCH_CHECK(
      fp32_router_gemm_supported(hidden_dim, num_experts),
      "fp32_router_gemm: supported (hidden_dim, num_experts) pairs are "
-      "(3072, 256), (6144, 128) and (6144, 256)");
+      "(3072, 256) and (6144, 128)");
  STD_TORCH_CHECK(num_tokens <= FP32_MAX_TOKENS,
                  "fp32_router_gemm: num_tokens must be in [0, 32]");
  STD_TORCH_CHECK(
@@ -39,10 +39,12 @@
 * The IQ/IK warps address the index_q/index_k sub-blocks *inside* qkv at the
 * fixed physical offsets (nq+2*nkv)*128 and (nq+2*nkv+niq)*128.
 *
- * Dense vs sparse row layout and index-branch processing are separate template
- * choices. Skip-index-topk reuse layers still have sparse rows and insert main
- * K/V cache entries, but compile away index_q/index_k work and index-cache
- * writes.
+ * Dense vs sparse is a compile-time choice via the ``kIsSparse``/``kInsertKV``
+ * template bools (3 instantiations: dense <false,false>, sparse-profiling
+ * <true,false>, sparse-serving <true,true>), so the index slots, the V slots
+ * and the cache inserts fold away entirely on paths that don't use them. The
+ * dense layer passes no caches/index: norm+RoPE happens in place and the
+ * generic ``Attention`` layer owns the cache write.
 *
 * Q/K and (sparse) index_q/index_k are all rewritten in place inside the fused
 * ``qkv`` tensor.  Caches (bf16) are scatter-written by slot.
@@ -221,25 +223,10 @@ __device__ __forceinline__ void storeCacheElems(
    // model dtype directly. FP8 cache dtypes use the conversion path below.
    storeElems<scalar_t>(reinterpret_cast<scalar_t*>(dst), elems);
  } else {
-#ifdef USE_ROCM
-    // Match ROCm's model-dtype materialization before FP8 cache conversion.
-    using Converter = vllm::_typeConvert<scalar_t>;
-    using rounded_t = typename Converter::hip_type;
-    rounded_t rounded[kElemsPerLane];
-  #pragma unroll
-    for (int i = 0; i < kElemsPerLane; i++) {
-      rounded[i] = Converter::convert(elems[i]);
-    }
-  #pragma unroll
-    for (int i = 0; i < kElemsPerLane; i++) {
-      dst[i] = fp8::scaled_convert<cache_t, rounded_t, kv_dt>(rounded[i], 1.0f);
-    }
-#else
-  #pragma unroll
+#pragma unroll
    for (int i = 0; i < kElemsPerLane; i++) {
      dst[i] = fp8::scaled_convert<cache_t, float, kv_dt>(elems[i], 1.0f);
    }
-#endif
  }
 }

@@ -275,25 +262,20 @@ __device__ __forceinline__ void storeElemsFp8(
 // Grid: 1D, ceil(num_tokens * slots_per_token / warps_per_block).
 // Each warp = one (token, slot).
 //
-// `kHasIndex`, `kProcessIndex`, and `kInsertKV` are compile-time template
-// bools, so branch decisions that distinguish the dense layer from the sparse
-// layer (index slots, KV/index inserts, V slots) fold away per instantiation.
-// Slots per token:
+// `kIsSparse` and `kInsertKV` are compile-time template bools, so all the
+// branch decisions that distinguish the dense layer from the sparse layer
+// (index slots, KV/index inserts, V slots) fold away per instantiation.
+// Three instantiations are built: dense <false,false>, sparse-profiling
+// <true,false> and sparse-serving <true,true>.  Slots per token:
 //     Q : nq                            (always — norm+RoPE)
 //     K : nkv                           (always — norm+RoPE; +K-cache insert)
 //     V : nkv  only if kInsertKV        (V-cache insert; no warps in dense)
-//     IQ: niq  only if kProcessIndex    (norm+RoPE)
-//     IK: 1    only if kProcessIndex    (norm+RoPE; +index-cache insert)
+//     IQ: niq  only if kIsSparse        (norm+RoPE)
+//     IK: 1    only if kIsSparse        (norm+RoPE; +index-cache insert)
 // cache_t/kv_dt: main attention KV-cache dtype (auto/fp8). out_idx_t/kFp8Idx:
 // indexer index-K cache + index-Q output dtype (scalar_t or e4m3 byte).
-// kHasIndex means the qkv row is laid out as sparse [q|k|v|index_q|index_k].
-// kProcessIndex controls whether this launch actually norms/ropes the index
-// branch and writes index_q/index_k outputs. Skip-index-topk reuse layers keep
-// kHasIndex=true but set kProcessIndex=false.
 template <typename scalar_t, typename cache_t, Fp8KVCacheDataType kv_dt,
-          typename out_idx_t, bool kHasIndex, bool kInsertKV,
-          bool kProcessIndex,
-          bool kFp8Idx>
+          typename out_idx_t, bool kIsSparse, bool kInsertKV, bool kFp8Idx>
 __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
    scalar_t* __restrict__ qkv,  // [N, qkv_row] in/out (packs index if sparse)
    scalar_t* __restrict__ q_out,         // [N, nq*128] contiguous, or nullptr
@@ -306,15 +288,16 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
    int64_t const* __restrict__ positions,       // [N] i64
    int64_t const* __restrict__ slot_mapping,    // main K/V slots or nullptr
    int64_t const* __restrict__ index_slot_mapping,  // index K slots/nullptr
-    cache_t* __restrict__ kv_cache,       // [nb,nkv,bs,2*128] or nullptr
+    cache_t* __restrict__ kv_cache,       // [nb,2,bs,nkv,128] or nullptr
    out_idx_t* __restrict__ index_cache,  // [nb*bs, 128]; scalar_t or e4m3 byte
    float const eps, int const rotary_dim, int const num_tokens, int const nq,
    int const nkv, int const niq, int const block_size,
-    // kv_cache strides (in elements) for logical shape [nb, nkv, bs, 2*128].
-    // The content (last) dim is always innermost-contiguous (stride 1), so the
-    // NHD/HND layout choice is captured by the head/token strides.
-    int64_t const kv_s_block, int64_t const kv_s_head, int64_t const kv_s_token,
-    int64_t const kv_s_dim) {
+    // kv_cache strides (in elements) for logical shape [nb, 2, bs, nkv, 128].
+    // The head_dim (last) dim is always innermost-contiguous (stride 1), so the
+    // NHD/HND layout choice is fully captured by these four strides: NHD keeps
+    // s_token < s_head, HND swaps them. dim_base addresses head_dim directly.
+    int64_t const kv_s_block, int64_t const kv_s_kv, int64_t const kv_s_token,
+    int64_t const kv_s_head) {
 #if (!defined(__CUDA_ARCH__) || __CUDA_ARCH__ < 800) && !defined(USE_ROCM)
  // _typeConvert<BFloat16> is unavailable on pre-Ampere; the M3 kernel only
  // runs with bf16/fp16 inputs in practice.  Discard the bf16 body there.
@@ -326,12 +309,9 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
    int const laneId = threadIdx.x % 32;
    int const globalWarpIdx = blockIdx.x * warpsPerBlock + (threadIdx.x / 32);

-    static_assert(!kProcessIndex || kHasIndex,
-                  "index processing requires sparse row layout");
-
    // Slot layout (compile-time gated: dense has neither V nor index slots).
    int const v_slots = kInsertKV ? nkv : 0;
-    int const idx_slots = kProcessIndex ? niq + 1 : 0;
+    int const idx_slots = kIsSparse ? niq + 1 : 0;
    int const slots_per_token = nq + nkv + v_slots + idx_slots;

    int const tokenIdx = globalWarpIdx / slots_per_token;
@@ -342,14 +322,14 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
    int const k_begin = nq;
    int const v_begin = nq + nkv;             // valid only when kInsertKV
    int const iq_begin = nq + nkv + v_slots;  // index block start
-    int const ik_slot = iq_begin + niq;       // valid only when kProcessIndex
+    int const ik_slot = iq_begin + niq;       // valid only when kIsSparse

    bool const isQ = slot < k_begin;
    bool const isK = slot >= k_begin && slot < v_begin;
    bool isV = false;
    if constexpr (kInsertKV) isV = slot >= v_begin && slot < v_begin + nkv;
    bool isIQ = false, isIK = false;
-    if constexpr (kProcessIndex) {
+    if constexpr (kIsSparse) {
      isIQ = slot >= iq_begin && slot < ik_slot;
      isIK = slot == ik_slot;
    }
@@ -357,7 +337,7 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
    int const dim_base = laneId * kElemsPerLane;
    // Physical row width of qkv: the dense layer packs [q|k|v]; the sparse
    // layer additionally packs [index_q (niq heads) | index_k (1 head)].
-    int const qkv_row = (nq + 2 * nkv + (kHasIndex ? (niq + 1) : 0)) * kHeadDim;
+    int const qkv_row = (nq + 2 * nkv + (kIsSparse ? (niq + 1) : 0)) * kHeadDim;

    // ── Resolve source pointer + per-branch parameters. ────────────────────
    scalar_t* row_ptr = nullptr;       // in-place output location
@@ -388,13 +368,10 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
      row_ptr = qkv + static_cast<int64_t>(tokenIdx) * qkv_row +
                (nq + 2 * nkv + ih) * kHeadDim;
      norm_w = iq_norm_w;
-    } else if (isIK) {
-      // Single shared index key at (nq+2*nkv+niq)*128.
+    } else {  // isIK -- single shared index key at (nq+2*nkv+niq)*128.
      row_ptr = qkv + static_cast<int64_t>(tokenIdx) * qkv_row +
                (nq + 2 * nkv + niq) * kHeadDim;
      norm_w = ik_norm_w;
-    } else {
-      return;
    }

    // Store destination.  Q and index_q are gathered into dedicated contiguous
@@ -450,12 +427,9 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
    // ── Cache inserts (sparse serving only). ───────────────────────────────
    if constexpr (kInsertKV) {
      // Guard (not early-return) so every thread reaches the PDL trigger below.
-      int64_t sm = -1;
-      if (isK || isV) {
-        sm = slot_mapping[tokenIdx];
-      } else if constexpr (kProcessIndex) {
-        if (isIK) sm = index_slot_mapping[tokenIdx];
-      }
+      int64_t const sm = (isK || isV)
+                             ? slot_mapping[tokenIdx]
+                             : (isIK ? index_slot_mapping[tokenIdx] : -1);
      if (sm >= 0) {  // skip padded / unscheduled tokens
        if (isIK) {
          if constexpr (kFp8Idx) {
@@ -464,16 +438,16 @@ __global__ void fusedMiniMaxM3QNormRopeKVInsertKernel(
            storeElems<scalar_t>(index_cache + sm * kHeadDim + dim_base, elems);
          }
        } else if (isK || isV) {
-          // kv_cache logical shape [num_blocks, nkv, block_size, 2*head_dim].
+          // kv_cache logical shape [num_blocks, 2, block_size, nkv, head_dim].
          // Paging is logical (block = sm/block_size, token = sm%block_size);
          // the physical NHD/HND layout is honoured via the passed strides.
          int64_t const b = sm / block_size;
          int64_t const t = sm % block_size;
          int const kv = isK ? 0 : 1;
-          int64_t const off = b * kv_s_block + head * kv_s_head +
-                              t * kv_s_token +
-                              (kv * kHeadDim + dim_base) * kv_s_dim;
-          storeCacheElems<scalar_t, cache_t, kv_dt>(kv_cache + off, elems);
+          int64_t const off =
+              b * kv_s_block + kv * kv_s_kv + t * kv_s_token + head * kv_s_head;
+          storeCacheElems<scalar_t, cache_t, kv_dt>(kv_cache + off + dim_base,
+                                                    elems);
        }
      }
    }
@@ -500,14 +474,14 @@ void launchFusedMiniMaxM3(
    int64_t const* index_slot_mapping, cache_t* kv_cache, void* index_cache,
    float const eps, int const rotary_dim, int const num_tokens, int const nq,
    int const nkv, int const niq, int const block_size,
-    int64_t const kv_s_block, int64_t const kv_s_head, int64_t const kv_s_token,
-    int64_t const kv_s_dim, bool const has_index, bool const insert_kv,
-    bool const process_index, bool const fp8_idx, cudaStream_t stream) {
+    int64_t const kv_s_block, int64_t const kv_s_kv, int64_t const kv_s_token,
+    int64_t const kv_s_head, bool const has_index, bool const insert_kv,
+    bool const fp8_idx, cudaStream_t stream) {
  // Index outputs are scalar_t (bf16) or e4m3 bytes (uint8_t); reinterpret the
  // void* pointers per instantiation in the LAUNCH macro.
  // Slot count must match the kernel's compile-time gating.
  int const v_slots = insert_kv ? nkv : 0;
-  int const idx_slots = process_index ? niq + 1 : 0;
+  int const idx_slots = has_index ? niq + 1 : 0;
  int const slots_per_token = nq + nkv + v_slots + idx_slots;

  constexpr int kBlockSize = 256;
@@ -534,59 +508,50 @@ void launchFusedMiniMaxM3(
  config.attrs = attrs;
  config.numAttrs = (sm_version >= 90) ? 1 : 0;

-  #define LAUNCH(HAS_INDEX, INSERT, PROCESS_INDEX, FP8, OUT_T)                 \
+  #define LAUNCH(IS_SPARSE, INSERT, FP8, OUT_T)                                \
    cudaLaunchKernelEx(                                                        \
        &config,                                                               \
        fusedMiniMaxM3QNormRopeKVInsertKernel<scalar_t, cache_t, kv_dt, OUT_T, \
-                                              HAS_INDEX, INSERT,               \
-                                              PROCESS_INDEX, FP8>,             \
+                                              IS_SPARSE, INSERT, FP8>,         \
        qkv, q_out, reinterpret_cast<OUT_T*>(index_q_out), q_norm_w, k_norm_w, \
        iq_norm_w, ik_norm_w, cos_sin_cache, positions, slot_mapping,          \
        index_slot_mapping, kv_cache, reinterpret_cast<OUT_T*>(index_cache),   \
        eps, rotary_dim, num_tokens, nq, nkv, niq, block_size, kv_s_block,     \
-        kv_s_head, kv_s_token, kv_s_dim)
+        kv_s_kv, kv_s_token, kv_s_head)
 #else
  // ROCm: standard kernel launch syntax (no PDL/stream serialization).
  // clang-format off
-  #define LAUNCH(HAS_INDEX, INSERT, PROCESS_INDEX, FP8, OUT_T)              \
-    fusedMiniMaxM3QNormRopeKVInsertKernel<                                  \
-        scalar_t, cache_t, kv_dt, OUT_T, HAS_INDEX, INSERT, PROCESS_INDEX,  \
-        FP8><<<grid, kBlockSize, 0, stream>>>(                              \
-        qkv, q_out, reinterpret_cast<OUT_T*>(index_q_out), q_norm_w,        \
-        k_norm_w, iq_norm_w, ik_norm_w, cos_sin_cache, positions,           \
-        slot_mapping, index_slot_mapping, kv_cache,                         \
-        reinterpret_cast<OUT_T*>(index_cache), eps, rotary_dim, num_tokens, \
-        nq, nkv, niq, block_size, kv_s_block, kv_s_head, kv_s_token,         \
-        kv_s_dim)
+  #define LAUNCH(IS_SPARSE, INSERT, FP8, OUT_T)                              \
+    fusedMiniMaxM3QNormRopeKVInsertKernel<scalar_t, cache_t, kv_dt, OUT_T,   \
+                                          IS_SPARSE, INSERT, FP8>            \
+        <<<grid, kBlockSize, 0, stream>>>(                                   \
+            qkv, q_out, reinterpret_cast<OUT_T*>(index_q_out), q_norm_w,     \
+            k_norm_w, iq_norm_w, ik_norm_w, cos_sin_cache, positions,        \
+            slot_mapping, index_slot_mapping, kv_cache,                      \
+            reinterpret_cast<OUT_T*>(index_cache), eps, rotary_dim,          \
+            num_tokens, nq, nkv, niq, block_size, kv_s_block, kv_s_kv,       \
+            kv_s_token, kv_s_head)
  // clang-format on
 #endif

  if (has_index) {
-    if (!process_index) {
-      if (insert_kv) {
-        LAUNCH(true, true, false, false, scalar_t);
-      } else {
-        LAUNCH(true, false, false, false, scalar_t);
-      }
-    } else if (insert_kv) {
+    if (insert_kv) {
      if (fp8_idx) {
-        LAUNCH(true, true, true, true,
-               uint8_t);  // sparse serving, fp8 index outputs
+        LAUNCH(true, true, true, uint8_t);  // sparse serving, fp8 index outputs
      } else {
-        LAUNCH(true, true, true, false, scalar_t);  // sparse serving, bf16
+        LAUNCH(true, true, false, scalar_t);  // sparse serving, bf16
      }
    } else {
      if (fp8_idx) {
-        LAUNCH(true, false, true, true,
-               uint8_t);  // sparse profiling, fp8 index_q
+        LAUNCH(true, false, true, uint8_t);  // sparse profiling, fp8 index_q
      } else {
-        LAUNCH(true, false, true, false, scalar_t);  // sparse profiling, bf16
+        LAUNCH(true, false, false, scalar_t);  // sparse profiling, bf16
      }
    }
  } else {
    // Dense layer: never has an index branch and never inserts here (the
    // generic Attention layer owns the KV insert).
-    LAUNCH(false, false, false, false, scalar_t);
+    LAUNCH(false, false, false, scalar_t);
  }
 #undef LAUNCH
 }
@@ -594,7 +559,6 @@ void launchFusedMiniMaxM3(
 }  // namespace minimax_m3_fused_ops
 }  // namespace vllm

-// clang-format off
 #define CALL_FUSED_MINIMAX_M3(_RAW_T, CACHE_T, KV_DTYPE)                       \
  vllm::minimax_m3_fused_ops::launchFusedMiniMaxM3<st, CACHE_T, KV_DTYPE>(     \
      reinterpret_cast<st*>(qkv.data_ptr()),                                   \
@@ -604,29 +568,24 @@ void launchFusedMiniMaxM3(
          : nullptr,                                                           \
      reinterpret_cast<st const*>(q_norm_weight.data_ptr()),                   \
      reinterpret_cast<st const*>(k_norm_weight.data_ptr()),                   \
-      process_index                                                            \
-          ? reinterpret_cast<st const*>(index_q_norm_weight->data_ptr())       \
-          : nullptr,                                                           \
-      process_index                                                            \
-          ? reinterpret_cast<st const*>(index_k_norm_weight->data_ptr())       \
-          : nullptr,                                                           \
+      has_index ? reinterpret_cast<st const*>(index_q_norm_weight->data_ptr()) \
+                : nullptr,                                                     \
+      has_index ? reinterpret_cast<st const*>(index_k_norm_weight->data_ptr()) \
+                : nullptr,                                                     \
      reinterpret_cast<st const*>(cos_sin_cache.data_ptr()),                   \
      reinterpret_cast<int64_t const*>(positions.data_ptr()),                  \
      insert_kv ? reinterpret_cast<int64_t const*>(slot_mapping->data_ptr())   \
                : nullptr,                                                     \
-      (insert_kv && process_index)                                             \
-          ? reinterpret_cast<int64_t const*>(                                  \
-                effective_index_slot_mapping->data_ptr())                      \
-          : nullptr,                                                           \
+      insert_kv ? reinterpret_cast<int64_t const*>(                            \
+                      effective_index_slot_mapping->data_ptr())                \
+                : nullptr,                                                     \
      insert_kv ? reinterpret_cast<CACHE_T*>(kv_cache->data_ptr()) : nullptr,  \
-      (insert_kv && process_index)                                             \
+      (insert_kv && has_index)                                                 \
          ? reinterpret_cast<void*>(index_cache->data_ptr())                   \
          : nullptr,                                                           \
      static_cast<float>(eps), static_cast<int>(rotary_dim), num_tokens, nq,   \
-      nkv, niq, static_cast<int>(block_size), kv_s_block, kv_s_head,           \
-      kv_s_token, kv_s_dim, has_index, insert_kv, process_index, fp8_idx,       \
-      stream)
-// clang-format on
+      nkv, niq, static_cast<int>(block_size), kv_s_block, kv_s_kv, kv_s_token, \
+      kv_s_head, has_index, insert_kv, fp8_idx, stream)

 // ────────────────────────────────────────────────────────────────────────────
 // Torch op wrapper
@@ -643,13 +602,13 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
    int64_t num_index_heads,                                  // niq; 0 => dense
    std::optional<torch::stable::Tensor> slot_mapping,        // [N] i64
    std::optional<torch::stable::Tensor> index_slot_mapping,  // [N] i64
-    std::optional<torch::stable::Tensor> kv_cache,     // [nb,nkv,bs,2*128]
+    std::optional<torch::stable::Tensor> kv_cache,     // [nb,2,bs,nkv,128]
    std::optional<torch::stable::Tensor> index_cache,  // [nb,bs,128]
    int64_t block_size,
    std::optional<torch::stable::Tensor> q_out,  // [N, nq*128] contiguous
    std::optional<torch::stable::Tensor>
        index_q_out,  // [N, niq*128] contiguous
-    const std::string& kv_cache_dtype, bool skip_index_branch) {
+    const std::string& kv_cache_dtype) {
  STD_TORCH_CHECK(qkv.is_cuda() && qkv.is_contiguous(),
                  "qkv must be contiguous CUDA");
  STD_TORCH_CHECK(
@@ -688,7 +647,6 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
  // (1 head)]) right after [q|k|v] in the same row; the dense layer does not.
  bool const has_index = niq > 0;
  bool const insert_kv = kv_cache.has_value();
-  bool const process_index = has_index && !skip_index_branch;
  vllm::Fp8KVCacheDataType const kv_dt =
      vllm::get_fp8_kv_cache_data_type(kv_cache_dtype);
  int const kHeadDim = vllm::minimax_m3_fused_ops::kHeadDim;
@@ -704,9 +662,7 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
  STD_TORCH_CHECK(
      !insert_kv || has_index,
      "insert mode (kv_cache) requires the index branch (sparse layer)");
-  STD_TORCH_CHECK(has_index || !skip_index_branch,
-                  "skip_index_branch requires sparse qkv rows");
-  if (process_index) {
+  if (has_index) {
    STD_TORCH_CHECK(
        index_q_norm_weight.has_value() && index_k_norm_weight.has_value(),
        "index branch requires both index norm weights");
@@ -717,26 +673,24 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
                        index_k_norm_weight->numel() == kHeadDim,
                    "index norm weights must have 128 elements");
  }
-  // kv_cache strides (logical shape [nb, nkv, bs, 2*head_dim]). Read straight
+  // kv_cache strides (logical shape [nb, 2, bs, nkv, head_dim]). Read straight
  // off the tensor so the kernel honours whatever physical layout the attention
-  // backend allocated (NHD: stride order (0,2,1,3); HND: (0,1,2,3)). No new
+  // backend allocated (NHD: stride order (0,1,2,3,4); HND: (0,1,3,2,4)). No new
  // op argument is needed -- the strides ride along with the tensor itself.
-  int64_t kv_s_block = 0, kv_s_head = 0, kv_s_token = 0, kv_s_dim = 0;
+  int64_t kv_s_block = 0, kv_s_kv = 0, kv_s_token = 0, kv_s_head = 0;
  torch::stable::Tensor const* effective_index_slot_mapping = nullptr;
  if (insert_kv) {
    STD_TORCH_CHECK(
        slot_mapping.has_value() && slot_mapping->is_cuda() &&
            slot_mapping->scalar_type() == torch::headeronly::ScalarType::Long,
        "insert mode requires int64 CUDA slot_mapping");
-    if (process_index) {
-      STD_TORCH_CHECK(
-          !index_slot_mapping.has_value() ||
-              (index_slot_mapping->is_cuda() &&
-               index_slot_mapping->scalar_type() ==
-                   torch::headeronly::ScalarType::Long &&
-               index_slot_mapping->numel() == slot_mapping->numel()),
-          "index_slot_mapping must be int64 CUDA with slot_mapping length");
-    }
+    STD_TORCH_CHECK(
+        !index_slot_mapping.has_value() ||
+            (index_slot_mapping->is_cuda() &&
+             index_slot_mapping->scalar_type() ==
+                 torch::headeronly::ScalarType::Long &&
+             index_slot_mapping->numel() == slot_mapping->numel()),
+        "index_slot_mapping must be int64 CUDA with slot_mapping length");
    // Main attention KV cache: auto matches qkv, fp8 uses uint8 storage.
    if (kv_dt == vllm::Fp8KVCacheDataType::kAuto) {
      STD_TORCH_CHECK(kv_cache->scalar_type() == qkv.scalar_type(),
@@ -747,26 +701,22 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
          "fp8 kv_cache must use uint8 storage");
    }
    // Indexer index-K cache: independent dtype -- qkv dtype or fp8 e4m3.
-    if (process_index) {
-      STD_TORCH_CHECK(
-          index_cache.has_value() &&
-              (index_cache->scalar_type() == qkv.scalar_type() ||
-               index_cache->scalar_type() ==
-                   torch::headeronly::ScalarType::Float8_e4m3fn),
-          "insert mode requires index_cache matching qkv dtype or fp8 e4m3");
-    }
-    STD_TORCH_CHECK(kv_cache->dim() == 4 && kv_cache->stride(3) == 1,
-                    "kv_cache must be [nb,nkv,bs,2*head_dim] with contiguous "
-                    "content dim (stride(3)==1)");
+    STD_TORCH_CHECK(
+        index_cache.has_value() &&
+            (index_cache->scalar_type() == qkv.scalar_type() ||
+             index_cache->scalar_type() ==
+                 torch::headeronly::ScalarType::Float8_e4m3fn),
+        "insert mode requires index_cache matching qkv dtype or fp8 e4m3");
+    STD_TORCH_CHECK(kv_cache->dim() == 5 && kv_cache->stride(4) == 1,
+                    "kv_cache must be [nb,2,bs,nkv,head_dim] with contiguous "
+                    "head_dim (stride(4)==1)");
    kv_s_block = kv_cache->stride(0);
-    kv_s_head = kv_cache->stride(1);
+    kv_s_kv = kv_cache->stride(1);
    kv_s_token = kv_cache->stride(2);
-    kv_s_dim = kv_cache->stride(3);
-    if (process_index) {
-      effective_index_slot_mapping = index_slot_mapping.has_value()
-                                         ? &index_slot_mapping.value()
-                                         : &slot_mapping.value();
-    }
+    kv_s_head = kv_cache->stride(3);
+    effective_index_slot_mapping = index_slot_mapping.has_value()
+                                       ? &index_slot_mapping.value()
+                                       : &slot_mapping.value();
  }
  // Optional contiguous gather targets: when given, the normed/roped q (and
  // index_q) are written here instead of in place, so callers avoid a separate
@@ -781,8 +731,9 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
        "q_out must have num_tokens * num_heads * 128 elements");
  }
  if (index_q_out.has_value()) {
-    STD_TORCH_CHECK(process_index,
-                    "index_q_out requires index branch processing");
+    STD_TORCH_CHECK(
+        has_index,
+        "index_q_out requires the index branch (num_index_heads > 0)");
    STD_TORCH_CHECK(
        index_q_out->is_cuda() && index_q_out->is_contiguous() &&
            (index_q_out->scalar_type() == qkv.scalar_type() ||
@@ -799,9 +750,8 @@ void fused_minimax_m3_qknorm_rope_kv_insert(
  // q/k/v + q_out stay qkv dtype. Both index outputs must agree.
  auto const kFp8 = torch::headeronly::ScalarType::Float8_e4m3fn;
  bool const fp8_idx =
-      process_index &&
-      ((index_cache.has_value() && index_cache->scalar_type() == kFp8) ||
-       (index_q_out.has_value() && index_q_out->scalar_type() == kFp8));
+      (index_cache.has_value() && index_cache->scalar_type() == kFp8) ||
+      (index_q_out.has_value() && index_q_out->scalar_type() == kFp8);
  if (fp8_idx) {
    STD_TORCH_CHECK(
        !index_cache.has_value() || index_cache->scalar_type() == kFp8,
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
Martin HickeyandHarry Mellor	702f4814fe	[docs] Fix the docs build (#48008 ) Signed-off-by: Martin Hickey <martin.hickey@ie.ibm.com> Co-authored-by: Harry Mellor <19981378+hmellor@users.noreply.github.com> (cherry picked from commit `6cf7b26bd4`)	2026-07-11 11:25:04 +01:00
Brandon Pelfreyandkhluu	dd10e03f95	Pin PyNvVideoCodec to tested 2.0.4 wheel (#48056 ) (cherry picked from commit `753c5039f0`)	2026-07-09 12:43:21 -07:00
Chaunceyandkhluu	3d1c21a6fc	[P/D][Bugfix] Fix PD async KV load lookahead handling for MTP spec decode (#46694 ) Signed-off-by: chaunceyjiang <chaunceyjiang@gmail.com> Co-authored-by: Nick Hill <nickhill123@gmail.com> (cherry picked from commit `ae6170f874`)	2026-07-09 03:24:39 -07:00
liuzhenweiandkhluu	e89b9c26d9	[XPU] Fix Event init failure w/ blocking (#47868 ) Signed-off-by: zhenwei-intel <zhenwei.liu@intel.com> (cherry picked from commit `bdaf27519f`)	2026-07-09 03:24:39 -07:00
Harry Mellorandkhluu	8bbc0062b1	Fix embed scaling + CUDA graphs in Transformers modelling backend (#48010 ) Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com> (cherry picked from commit `56da398dac`)	2026-07-08 21:12:19 -07:00
Thien Tranandkhluu	65168ef45c	Allow FlashInfer A2A backends for TRTLLM FP8 MoE Modular (#46661 ) Signed-off-by: Thien Tran <gau.nernst@yahoo.com.sg> (cherry picked from commit `0d2f4e7c9c`)	2026-07-08 21:12:19 -07:00
Robert Shawandkhluu	51e3572d30	[Bug] Fix Batched DeepGEMM (#47884 ) Signed-off-by: Robert Shaw <robertgshaw2-redhat@dgx-b200-02.mgmt.accl-001.lab.rdu2.dc.redhat.com> Co-authored-by: Robert Shaw <robertgshaw2-redhat@dgx-b200-02.mgmt.accl-001.lab.rdu2.dc.redhat.com> Co-authored-by: Claude <noreply@anthropic.com> (cherry picked from commit `572b25b03e`)	2026-07-08 21:12:19 -07:00
Canlin Guoandkhluu	654e9e5b90	[Model] Support MOSS-Transcribe-Diarize (#47729 ) Signed-off-by: gcanlin <canlinguosdu@gmail.com> (cherry picked from commit `285c08c036`)	2026-07-08 21:12:19 -07:00
Wei Zhaoandkhluu	15d7edbbcf	[Perf] Minimax M3 - Support cross-layer allreduce-norm fusion (#47631 ) Signed-off-by: wzhao18 <wzhao18.sz@gmail.com> Co-authored-by: Yongye Zhu <zyy1102000@gmail.com> (cherry picked from commit `2afa3f7e95`)	2026-07-08 21:12:19 -07:00
Jeff (Junze) Maandkhluu	b53923c9ef	[Bugfix] DSV4 TP16 garbage output (#47493 ) Signed-off-by: Jeff Ma <jeffjma@umich.edu> Co-authored-by: Lucas Wilkinson <LucasWilkinson@users.noreply.github.com> (cherry picked from commit `80eb01e93d`)	2026-07-08 21:12:19 -07:00
Yongye Zhuandkhluu	cf3df33687	[Minimax-M3] Using tok_sparse_select from MSA instead of triton kernels (#47502 ) Signed-off-by: Yongye Zhu <zyy1102000@gmail.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com> (cherry picked from commit `9021589498`)	2026-07-08 21:12:19 -07:00
Dakai Anandkhluu	ae3665a6ab	[Bugfix] Fix PD disagg + MTP correctness for Qwen3.5(GDN) (#47466 ) Signed-off-by: Dakai An <dakaian108@gmail.com> Signed-off-by: Nick Hill <nickhill123@gmail.com> Co-authored-by: Nick Hill <nickhill123@gmail.com> (cherry picked from commit `65dcde1695`)	2026-07-08 21:12:19 -07:00
Summer Yangandkhluu	f5ba5da988	[Perf] Use blocking CUDA events to avoid busy polling cuda driver lock (#47081 ) Signed-off-by: Jingyi Yang <girasoleyang@gmail.com> (cherry picked from commit `d3e69fd671`)	2026-07-08 21:12:19 -07:00
Wentao Yeandkhluu	7da76e4147	[CI Bug Fix] Temp fix for v3.2 accuracy (#47902 ) (cherry picked from commit `3f99883d97`)	2026-07-07 18:25:17 -07:00
Rahul Vishwakarmaandkhluu	b2dec4ac5a	[CPU][Bugfix] Fix flaky ShortConv prefill test on ARM (uninitialized weights) (#47848 ) Signed-off-by: Rahul Vishwakarma <Rahul.Vishwakarma2@ibm.com> (cherry picked from commit `f7efab58ec`)	2026-07-07 18:20:30 -07:00
Harry Mellorandkhluu	88550342ce	[CI] Fix Transformers modeling backend LoRA test (#47832 ) Signed-off-by: Harry Mellor <19981378+hmellor@users.noreply.github.com> (cherry picked from commit `0ed05b6f82`)	2026-07-07 18:07:57 -07:00
danielafrimiandkhluu	ad49bb65d9	[BugFix] Fix ModelOpt mixed-precision quantization for sparse `quantized_layers` configs. (#47318 ) Signed-off-by: Daniel Afrimi <dafrimi@nvidia.com> Signed-off-by: <dafrimi@nvidia.com> (cherry picked from commit `0a2965b1b3`)	2026-07-07 18:05:03 -07:00