[Core] Enable extensible KV cache for all attention backends and Mamba

Port of internal D110967544. The extensible KV cache flow (reserve KV virtual address space up front, capture CUDA graphs first, then size and commit the KV cache from post-capture free memory) previously required a block-major attention backend and rejected Mamba models. This enables it for every backend layout and for Mamba / linear attention: - ExtensibleTensor gains num_segments: the reservation is divided into equal segments that grow in lockstep, with committed bytes forming a prefix of each segment. Physical pages are mapped at allocation-granularity granules and deduped across overlapping ranges, so a granule straddling a segment boundary is mapped exactly once. resize_per_segment_(bytes, zero_new=True) zeroes only the newly committed logical range of each segment. - Each KV cache buffer keeps its layers' physical layout and is committed as one prefix per layout segment. The segment count is derived from the backend's get_kv_cache_shape / get_kv_cache_block_dim / stride order: K/V-split layouts (e.g. FlashAttention) get one prefix per half, block-major layouts (e.g. FlashInfer, MLA) a single prefix. Mamba state pages are block-major per layer, and hybrid-model attention caches are re-strided to block-major, so both use a single segment. - Removed the supports_extensible_kv_cache gate plumbing from EngineCore, Executor, Worker, WorkerBase and GPUModelRunner; a CUDA platform check remains in EngineCore. - enable_extensible_kv_cache is reported as unsupported by the V2 model runner so V2-default models fall back to the V1 runner (which implements the flow); also fixed initialize_kv_cache being called with the extensible kwarg on runners that do not accept it, which broke every default V2-runner boot on this branch. Tested on H100: - tests/utils_/test_extensible_tensor.py (5 passed, incl. new segmented lockstep-grow/zero, granule-dedup and invalid-usage tests) - tests/v1/worker/test_extensible_kv_cache.py (new, 6 passed: segment derivation, split grows both halves, block-major, legacy full commit, Mamba per-layer growth, hybrid attention+Mamba) - E2E Qwen3-0.6B greedy with VLLM_ATTENTION_BACKEND=FLASH_ATTN (a K/V-split backend the old gate rejected): extensible generations byte-identical to the legacy path; log shows reserve then "Extended KV cache to 34663 blocks". V2->V1 auto-fallback path verified as well.
add new files
2026-07-10 17:08:34 -07:00 · 2026-07-01 17:32:19 -07:00 · 2026-07-01 17:30:53 -07:00
1051 changed files with 14981 additions and 51686 deletions
@@ -3,7 +3,6 @@ job_dirs:
  - ".buildkite/intel_jobs"
 run_all_patterns:
  - ".buildkite/ci_config_intel.yaml"
-  - ".buildkite/scripts/hardware_ci/run-intel-test.sh"
  - "docker/Dockerfile"
  - "docker/Dockerfile.xpu"
  - "CMakeLists.txt"
@@ -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,28 +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  # 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
@@ -30,16 +7,13 @@ 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"
@@ -59,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"
@@ -17,14 +17,12 @@ steps:
  - tests/kernels/test_awq_int4_to_int8.py
  - tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
  - tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
-  - tests/kernels/mamba/test_cpu_short_conv.py
  commands:
    - |
      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
      pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
      pytest -x -v -s tests/kernels/moe/test_cpu_fused_moe.py
      pytest -x -v -s tests/kernels/moe/test_cpu_quant_fused_moe.py
-      pytest -x -v -s tests/kernels/mamba/test_cpu_short_conv.py
      pytest -x -v -s tests/kernels/test_onednn.py
      pytest -x -v -s tests/kernels/test_awq_int4_to_int8.py
      pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
@@ -141,22 +139,9 @@ steps:
  commands:
    - |
      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 45m "
-      pytest -x -v -s tests/models/multimodal/generation --ignore=tests/models/multimodal/generation/test_pixtral.py --ignore=tests/models/multimodal/generation/test_qwen2_5_vl.py -m cpu_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB"
+      pytest -x -v -s tests/models/multimodal/generation --ignore=tests/models/multimodal/generation/test_pixtral.py -m cpu_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB"
  parallelism: 4

- label: CPU-Qwen2.5-VL Multimodal Tests
-  depends_on: []
-  device: intel_cpu
-  no_plugin: true
-  source_file_dependencies:
-  # - vllm/
-  - vllm/model_executor/layers/rotary_embedding
-  - tests/models/multimodal/generation/
-  commands:
-    - |
-      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 40m "
-      VLLM_CI_ENV=0 pytest -x -v -s tests/models/multimodal/generation/test_qwen2_5_vl.py"
-
 - label: "Arm CPU Test"
  depends_on: []
  soft_fail: false
@@ -45,7 +45,7 @@ steps:
    agent_tags:
      label: production
      gpu: 1+
-      mem: 24+
+      mem: 16+
    no_plugin: true
    env:
      REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -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
@@ -81,9 +81,7 @@ steps:
      'cd tests &&
      export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
      set -o pipefail &&
-      pytest -v -s lora/test_punica_ops.py::test_kernels &&
-      pytest -v -s lora/test_punica_ops.py::test_kernels_hidden_size &&
-      pytest -v -s lora/test_punica_ops.py::test_add_lora_fused_moe_early_exit'
+      pytest -v -s lora/test_punica_ops.py --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-3-43264-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype1-1-2049-64-128-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-1-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-1-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-8-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype0-3-2049-128-8-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-8-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype1-1-2049-256-128-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-64256-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-2-29696-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-3-49408-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-2-16384-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-2-51328-32-4-4]" --deselect="tests/lora/test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-32000-32-4-4]" --deselect="tests/lora/test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-32000-32-4-4]"'

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

@@ -9,7 +9,7 @@ steps:
  agent_tags:
    label: production
    gpu: 2+
-    mem: 16+
+    mem: 24+
  no_plugin: true
  working_dir: "."
  env:
@@ -9,7 +9,7 @@ steps:
  agent_tags:
    label: production
    gpu: 1+
-    mem: 24+
+    mem: 16+
  no_plugin: true
  working_dir: "."
  env:
@@ -81,7 +81,7 @@ steps:
  agent_tags:
    label: production
    gpu: 1+
-    mem: 24+
+    mem: 16+
  no_plugin: true
  working_dir: "."
  env:
@@ -105,7 +105,7 @@ steps:
  agent_tags:
    label: production
    gpu: 1+
-    mem: 24+
+    mem: 16+
  no_plugin: true
  working_dir: "."
  env:
@@ -1,28 +0,0 @@
-group: Quantization
-depends_on:
-  - image-build-xpu
-steps:
- label: Quantization
-  key: quantization
-  timeout_in_minutes: 30
-  env:
-    REGISTRY: "public.ecr.aws/q9t5s3a7"
-    REPO: "vllm-ci-test-repo"
-    VLLM_TEST_DEVICE: "xpu"
-  no_plugin: true
-  working_dir: "."
-  device: intel_gpu
-  agent_tags:
-    label: production
-    gpu: 1+
-    mem: 16+
-  source_file_dependencies:
-  - csrc/
-  - vllm/model_executor/layers/quantization
-  - tests/quantization
-  commands:
-  # - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py
-  - >-
-    bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-    'VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s tests/quantization/test_per_token_kv_cache.py --deselect="tests/quantization/test_per_token_kv_cache.py::test_triton_unified_attention_per_token_head_scale[int4-16-128-num_heads0-seq_lens1]"'
-
@@ -42,37 +42,12 @@ steps:
        python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --kv-cache-dtype fp8 &&
        python3 examples/basic/offline_inference/generate.py --model nvidia/Llama-3.1-8B-Instruct-FP8 --block-size 64 --enforce-eager --quantization modelopt --kv-cache-dtype fp8 --attention-backend TRITON_ATTN --max-model-len 4096 &&
        python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --block-size 64 --enforce-eager --max-model-len 8192 &&
-        python3 examples/basic/offline_inference/generate.py --model TheBloke/TinyLlama-1.1B-Chat-v0.3-AWQ --block-size 64 --enforce-eager &&
        python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 &&
        python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 --enable-expert-parallel &&
        python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --max-model-len 8192 &&
        VLLM_XPU_FUSED_MOE_USE_REF=1 python3 examples/basic/offline_inference/generate.py --model Qwen/Qwen3-30B-A3B-Instruct-2507-FP8 --enforce-eager -tp 2 --max-model-len 8192 &&
        python3 examples/basic/offline_inference/generate.py --model INCModel/Qwen3-30B-A3B-Instruct-2507-MXFP4-LLMC --enforce-eager -tp 2 --max-model-len 8192
        '
-  - label: "XPU W8A8 FP8 Linear Examples"
-    depends_on:
-      - image-build-xpu
-    timeout_in_minutes: 60
-    device: intel_gpu
-    agent_tags:
-      label: production
-      gpu: 1+
-      mem: 24+
-    no_plugin: true
-    env:
-      REGISTRY: "public.ecr.aws/q9t5s3a7"
-      REPO: "vllm-ci-test-repo"
-      VLLM_TEST_DEVICE: "xpu"
-    source_file_dependencies:
-      - vllm/
-      - .buildkite/intel_jobs/test-intel.yaml
-    commands:
-      - >-
-        bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-        'python3 examples/basic/offline_inference/generate.py --linear-backend xpu --model RedHatAI/Meta-Llama-3.1-8B-Instruct-FP8 --enforce-eager --max-model-len 4096 &&
-        python3 examples/basic/offline_inference/generate.py --linear-backend xpu --model neuralmagic/Llama-3.2-1B-Instruct-FP8-dynamic --enforce-eager --max-model-len 4096 &&
-        python3 examples/basic/offline_inference/generate.py --linear-backend xpu --model meta-llama/Llama-3.2-1B-Instruct --quantization fp8 --enforce-eager --max-model-len 4096
-        '
  - label: "XPU V1 test"
    depends_on:
      - image-build-xpu
@@ -81,7 +56,7 @@ steps:
    agent_tags:
      label: production
      gpu: 1+
-      mem: 24+
+      mem: 16+
    no_plugin: true
    env:
      REGISTRY: "public.ecr.aws/q9t5s3a7"
@@ -144,27 +119,4 @@ steps:
      - >-
        bash .buildkite/scripts/hardware_ci/run-intel-test.sh
        'cd tests &&
-        pytest -v -s quantization/test_auto_round.py'
-  - label: "XPU compressed tensors FP8 test"
-    depends_on:
-      - image-build-xpu
-    timeout_in_minutes: 60
-    device: intel_gpu
-    agent_tags:
-      label: production
-      gpu: 1+
-      mem: 16+
-    no_plugin: true
-    env:
-      REGISTRY: "public.ecr.aws/q9t5s3a7"
-      REPO: "vllm-ci-test-repo"
-      VLLM_TEST_DEVICE: "xpu"
-    source_file_dependencies:
-      - vllm/
-      - tests/quantization/test_compressed_tensors.py
-      - .buildkite/intel_jobs/test-intel.yaml
-    commands:
-      - >-
-        bash .buildkite/scripts/hardware_ci/run-intel-test.sh
-        'cd tests &&
-        pytest -v -s quantization/test_compressed_tensors.py::test_compressed_tensors_fp8'
+        pytest -v -s quantization/test_auto_round.py'
@@ -350,25 +350,6 @@ steps:
          - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"
          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"'

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

-      - label: "Create manifest - XPU"
-        depends_on:
-          - build-xpu-release-image
-        id: create-manifest-xpu
-        agents:
-          queue: small_cpu_queue_release
-        commands:
-          - "bash .buildkite/scripts/xpu/create-xpu-ecr-manifest.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: XPU" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-xpu"'
-
      - label: "Publish nightly multi-arch image to DockerHub"
        depends_on:
          - create-multi-arch-manifest
@@ -848,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
@@ -895,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
@@ -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
@@ -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,7 +15,7 @@ set -euo pipefail

 DEFAULT_REPO_SLUG="vllm-project/vllm"
 DEFAULT_CI_HCL_SOURCE="docker/ci-rocm.hcl"
-DEFAULT_CI_BASE_CONTENT_FILES="requirements/common.txt requirements/rocm.txt requirements/test/rocm.txt docker/Dockerfile.rocm_base docker/ci-rocm.hcl docker/docker-bake-rocm.hcl tools/install_torchcodec_rocm.sh tests/vllm_test_utils .buildkite/scripts/ci-bake-rocm.sh .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 build_rixl build_rocshmem build_deepep mori_base ci_base"
 DEFAULT_CI_BASE_METADATA_VERSION="1"
@@ -393,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"
@@ -743,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}"
@@ -762,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:-}"
-        set_buildkite_metadata "rocm-ci-base-image-stable" "${CI_BASE_IMAGE_TAG_STABLE:-}"
        return 0
    fi

@@ -1795,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
 }
@@ -1826,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"
@@ -28,17 +28,6 @@
 ###############################################################################
 set -o pipefail

-: "${BUILDKIT_PROGRESS:=plain}"
-: "${TERM:=xterm-256color}"
-: "${FORCE_COLOR:=1}"
-: "${CLICOLOR_FORCE:=1}"
-: "${PY_COLORS:=1}"
-: "${ROCM_DOCKER_TTY:=1}"
-if [[ " ${PYTEST_ADDOPTS:-} " != *" --color"* ]]; then
-  PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }--color=yes"
-fi
-export BUILDKIT_PROGRESS TERM FORCE_COLOR CLICOLOR_FORCE PY_COLORS PYTEST_ADDOPTS ROCM_DOCKER_TTY
-
 # Export Python path for commands that run directly on the host. Containerized
 # tests set this to /vllm-workspace below so spawned Python processes do not
 # depend on their current working directory.
@@ -160,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
@@ -460,11 +448,9 @@ 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
+if git -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
+    git -C "${checkout}" merge-base HEAD origin/main 2>/dev/null || true
  )"
 fi
 if [[ -z "${vllm_standalone_merge_base}" ]]; then
@@ -547,13 +533,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
@@ -570,7 +549,6 @@ else
  fi

  docker run \
-    "${docker_run_terminal_args[@]}" \
    --device /dev/kfd $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES \
    $RDMA_FLAGS \
    --network=host \
@@ -585,11 +563,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 PYTEST_ADDOPTS \
    -v "${HF_CACHE}:${HF_MOUNT}" \
    -e "HF_HOME=${HF_MOUNT}" \
    -e "PYTHONPATH=${MYPYTHONPATH}" \
@@ -38,9 +38,7 @@ function cpu_tests() {
    pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
    pytest -x -v -s tests/kernels/core/test_cpu_activation.py
    pytest -x -v -s tests/kernels/moe/test_cpu_fused_moe.py
-    pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
-    pytest -x -v -s tests/kernels/moe/test_cpu_int4_moe.py
-    pytest -x -v -s tests/kernels/mamba/test_cpu_short_conv.py"
+    pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py"

  # skip tests requiring model downloads if HF_TOKEN is not set
  # due to rate-limits
@@ -64,6 +62,7 @@ function cpu_tests() {
    set -e
    pytest -x -v -s tests/quantization/test_compressed_tensors.py::test_compressed_tensors_w8a8_logprobs"

+
  # basic online serving
  docker exec cpu-test bash -c '
    set -e
@@ -21,7 +21,6 @@ case "${test_suite}" in
    python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --kv-cache-dtype fp8
    python3 examples/basic/offline_inference/generate.py --model nvidia/Llama-3.1-8B-Instruct-FP8 --block-size 64 --enforce-eager --quantization modelopt --kv-cache-dtype fp8 --attention-backend TRITON_ATTN --max-model-len 4096
    python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --block-size 64 --enforce-eager --max-model-len 8192
-    python3 examples/basic/offline_inference/generate.py --model TheBloke/TinyLlama-1.1B-Chat-v0.3-AWQ --block-size 64 --enforce-eager 
    python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2
    python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 --enable-expert-parallel
    python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --max-model-len 8192
@@ -360,7 +360,7 @@ export HF_TOKEN ZE_AFFINITY_MASK
    --ipc=host \
    --privileged \
    -v /dev/dri/by-path:/dev/dri/by-path \
-    -v "/data/huggingface:/root/.cache/huggingface" \
+    -v "${HOME}/.cache/huggingface:/root/.cache/huggingface" \
    --entrypoint='' \
    -e HF_TOKEN \
    -e ZE_AFFINITY_MASK \
@@ -85,7 +85,7 @@ RUN pip config set global.index-url http://cache-service-vllm.nginx-pypi-cache.s

 # Install for pytest to make the docker build cache layer always valid
 RUN --mount=type=cache,target=/root/.cache/pip \
-    pip install pytest>=6.0  'modelscope<1.38'
+    pip install pytest>=6.0  modelscope

 WORKDIR /workspace/vllm

@@ -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,57 +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_refreshed_base_if_present() {
-    local base_refreshed=""
-
-    base_refreshed="$(metadata_get rocm-base-refresh)"
-    if [[ "${base_refreshed}" != "1" ]]; then
-        return 1
-    fi
-
-    export BASE_IMAGE
-    export CI_BASE_IMAGE
-    export IMAGE_TAG_LATEST
-
-    BASE_IMAGE="$(metadata_get rocm-base-image)"
-    CI_BASE_IMAGE="$(metadata_get rocm-ci-base-image)"
-    IMAGE_TAG_LATEST="$(metadata_get rocm-ci-image-descriptive)"
-
-    echo "Using refreshed ROCm base image for test image: ${BASE_IMAGE}"
-    echo "Using refreshed ROCm ci_base image for test image: ${CI_BASE_IMAGE}"
-    if [[ -n "${IMAGE_TAG_LATEST}" ]]; then
-        echo "Also tagging full ROCm CI image as: ${IMAGE_TAG_LATEST}"
-    fi
-
-    return 0
-}
-
-main() {
-    local base_refreshed=0
-
-    if use_refreshed_base_if_present; then
-        base_refreshed=1
-    fi
-
-    if [[ "${ROCM_CI_ARTIFACT_ONLY:-0}" == "1" && "${base_refreshed}" != "1" ]]; then
-        echo "ROCM_CI_ARTIFACT_ONLY=1; building ROCm wheel artifact only"
-        IMAGE_TAG="" bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-artifacts
-        return
-    fi
-
-    bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-wheel
-}
-
-main "$@"
@@ -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
-'
@@ -109,9 +109,7 @@ run_nodes() {
        if [ "$node" -ne 0 ]; then
            docker exec -d "node$node" /bin/bash -c "cd $WORKING_DIR ; ${COMMANDS[$node]}"
        else
-            # Allocate a TTY (-t -i) for the foreground head node so its output
-            # keeps ANSI color in the Buildkite log (see run-amd-test.sh).
-            docker exec -t -i "node$node" /bin/bash -c "cd $WORKING_DIR ; ${COMMANDS[$node]}"
+            docker exec "node$node" /bin/bash -c "cd $WORKING_DIR ; ${COMMANDS[$node]}"
        fi
    done
 }
@@ -8,12 +8,7 @@ if [[ "$MODE" != "style-clippy" && "$MODE" != "test" ]]; then
  exit 2
 fi

-if ROOT_DIR="$(git rev-parse --show-toplevel 2>/dev/null)"; then
-  :
-else
-  SCRIPT_DIR="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")" && pwd -P)"
-  ROOT_DIR="$(cd -- "${SCRIPT_DIR}/../.." && pwd -P)"
-fi
+ROOT_DIR="$(git rev-parse --show-toplevel)"
 cd "$ROOT_DIR"

 export CARGO_TERM_COLOR="${CARGO_TERM_COLOR:-always}"
@@ -21,20 +16,16 @@ export CARGO_HOME="${CARGO_HOME:-$HOME/.cargo}"
 export RUSTUP_HOME="${RUSTUP_HOME:-$HOME/.rustup}"
 export PATH="$CARGO_HOME/bin:$PATH"

-PROTOC_VERSION="${PROTOC_VERSION:-31.1}"
-CARGO_BINSTALL_VERSION="${CARGO_BINSTALL_VERSION:-1.20.1}"
-UV_VERSION="${UV_VERSION:-0.11.28}"
-PYO3_PYTHON_VERSION="${PYO3_PYTHON_VERSION:-3.12}"
-
-CARGO_SORT_VERSION_REQ="${CARGO_SORT_VERSION_REQ:-2}"
-CARGO_DENY_VERSION_REQ="${CARGO_DENY_VERSION_REQ:-0.20}"
-CARGO_NEXTEST_VERSION_REQ="${CARGO_NEXTEST_VERSION_REQ:-0.9}"
-
 log_section() {
  echo "--- $*"
 }

 install_protoc() {
+  if command -v protoc >/dev/null 2>&1; then
+    return
+  fi
+
+  local version="${PROTOC_VERSION:-31.1}"
  local arch
  case "$(uname -m)" in
    x86_64)
@@ -49,17 +40,16 @@ install_protoc() {
      ;;
  esac

-  local url="https://github.com/protocolbuffers/protobuf/releases/download/v${PROTOC_VERSION}/protoc-${PROTOC_VERSION}-linux-${arch}.zip"
+  local url="https://github.com/protocolbuffers/protobuf/releases/download/v${version}/protoc-${version}-linux-${arch}.zip"
  local tmp_dir
  tmp_dir="$(mktemp -d)"

-  log_section "Installing protoc ${PROTOC_VERSION}"
+  log_section "Installing protoc ${version}"
  curl -L --proto '=https' --tlsv1.2 -sSf "$url" -o "$tmp_dir/protoc.zip"
  mkdir -p "$CARGO_HOME/bin"
  unzip -q "$tmp_dir/protoc.zip" bin/protoc 'include/*' -d "$CARGO_HOME"
  chmod +x "$CARGO_HOME/bin/protoc"
  rm -rf "$tmp_dir"
-  protoc --version
 }

 rust_toolchain() {
@@ -80,48 +70,66 @@ install_rust_toolchain() {
 }

 install_cargo_binstall() {
-  log_section "Installing cargo-binstall ${CARGO_BINSTALL_VERSION}"
+  if command -v cargo-binstall >/dev/null 2>&1; then
+    return
+  fi
+
+  log_section "Installing cargo-binstall"
  curl -L --proto '=https' --tlsv1.2 -sSf \
-    "https://raw.githubusercontent.com/cargo-bins/cargo-binstall/v${CARGO_BINSTALL_VERSION}/install-from-binstall-release.sh" \
-    | 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 +151,6 @@ PY
 }

 run_style_clippy() {
-  install_cargo_binstall
  install_cargo_sort
  install_cargo_deny

@@ -156,8 +163,8 @@ run_style_clippy() {
  log_section "Checking Rust dependency bans"
  cargo deny \
    --manifest-path rust/Cargo.toml \
-    --config rust/deny.toml \
    check \
+    --config rust/deny.toml \
    bans

  log_section "Running clippy"
@@ -174,7 +181,6 @@ run_style_clippy() {
 run_tests() {
  install_uv
  setup_pyo3_python
-  install_cargo_binstall
  install_cargo_nextest

  log_section "Running cargo nextest"
@@ -18,10 +18,6 @@ wait_for_server() {

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

 cleanup() {
  if [[ -n "${SERVER_PID:-}" ]] && kill -0 "${SERVER_PID}" 2>/dev/null; then
@@ -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}"
@@ -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/
@@ -23,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:
@@ -16,7 +16,7 @@ steps:
  mirror:
    amd:
      device: mi300_4
-      timeout_in_minutes: 85
+      timeout_in_minutes: 110
      depends_on:
        - image-build-amd
      source_file_dependencies:
@@ -29,7 +29,7 @@ steps:

 - label: Distributed FlashInfer NixlConnector PD accuracy (4 GPUs)
  key: distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus
-  timeout_in_minutes: 55
+  timeout_in_minutes: 30
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -39,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,7 +53,7 @@ steps:
  mirror:
    amd:
      device: mi300_4
-      timeout_in_minutes: 60
+      timeout_in_minutes: 50
      depends_on:
        - image-build-amd
      source_file_dependencies:
@@ -79,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:
@@ -91,7 +78,7 @@ steps:
  mirror:
    amd:
      device: mi300_4
-      timeout_in_minutes: 85
+      timeout_in_minutes: 110
      depends_on:
        - image-build-amd
      source_file_dependencies:
@@ -104,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:
@@ -116,7 +103,7 @@ steps:
  mirror:
    amd:
      device: mi300_4
-      timeout_in_minutes: 80
+      timeout_in_minutes: 60
      depends_on:
        - image-build-amd
      source_file_dependencies:
@@ -143,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:
@@ -158,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
@@ -172,7 +159,7 @@ steps:
  mirror:
    amd:
      device: mi300_2
-      timeout_in_minutes: 70
+      timeout_in_minutes: 60
      depends_on:
        - image-build-amd
      source_file_dependencies:
@@ -186,7 +173,7 @@ steps:

 - 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:
@@ -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:
@@ -55,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:
@@ -78,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:
@@ -133,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:
@@ -154,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:
@@ -176,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"
@@ -212,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/"
@@ -259,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:
@@ -274,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:
@@ -120,7 +120,7 @@ steps:

 - 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:
@@ -148,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:
@@ -27,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:
@@ -39,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,7 +172,7 @@ 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/model_executor/layers/fla/ops/kda.py
@@ -185,7 +184,7 @@ steps:

 - 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
@@ -265,20 +264,19 @@ steps:

 - 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
@@ -295,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
@@ -305,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
@@ -318,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
@@ -109,7 +109,7 @@ steps:

 - 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
@@ -118,7 +118,7 @@ steps:
  mirror:
    amd:
      device: mi300_8
-      timeout_in_minutes: 60
+      timeout_in_minutes: 180
      depends_on:
      - image-build-amd
      commands:
@@ -150,9 +150,9 @@ steps:
  commands:
    - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/moe-refactor-dp-ep/config-b200.txt

- label: LM Eval Humming f16 (A100 - TEMPORARY)
-  key: lm-eval-humming-f16-a100
-  timeout_in_minutes: 75
+- label: LM Eval Humming (A100 - TEMPORARY)
+  key: lm-eval-humming-a100
+  timeout_in_minutes: 30
  device: a100
  optional: true
  num_devices: 1
@@ -160,29 +160,13 @@ steps:
  - vllm/model_executor/layers/quantization/humming.py
  - vllm/model_executor/layers/quantization/utils/humming_utils.py
  - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
-  - vllm/model_executor/layers/fused_moe/oracle/
-  - vllm/model_executor/kernels/linear/
+  - vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
  commands:
  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config.txt

- label: LM Eval Humming Act int8 (A100 - TEMPORARY)
-  key: lm-eval-humming-act-a100
-  timeout_in_minutes: 45
-  device: a100
-  optional: true
-  num_devices: 1
-  source_file_dependencies:
-  - vllm/model_executor/layers/quantization/humming.py
-  - vllm/model_executor/layers/quantization/utils/humming_utils.py
-  - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
-  - vllm/model_executor/layers/fused_moe/oracle/
-  - vllm/model_executor/kernels/linear/
-  commands:
-  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-int8.txt
-
- label: LM Eval Humming f16 (H100 - TEMPORARY)
-  key: lm-eval-humming-f16-h100
-  timeout_in_minutes: 70
+- label: LM Eval Humming (H100 - TEMPORARY)
+  key: lm-eval-humming-h100
+  timeout_in_minutes: 30
  device: h100
  optional: true
  num_devices: 1
@@ -190,30 +174,14 @@ steps:
  - vllm/model_executor/layers/quantization/humming.py
  - vllm/model_executor/layers/quantization/utils/humming_utils.py
  - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
-  - vllm/model_executor/layers/fused_moe/oracle/
-  - vllm/model_executor/kernels/linear/
+  - vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
  commands:
  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config.txt
-
- label: LM Eval Humming Act fp8/int8 (H100 - TEMPORARY)
-  key: lm-eval-humming-act-h100
-  timeout_in_minutes: 70
-  device: h100
-  optional: true
-  num_devices: 1
-  source_file_dependencies:
-  - vllm/model_executor/layers/quantization/humming.py
-  - vllm/model_executor/layers/quantization/utils/humming_utils.py
-  - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
-  - vllm/model_executor/layers/fused_moe/oracle/
-  - vllm/model_executor/kernels/linear/
-  commands:
  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-fp8.txt
-  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-int8.txt

- label: LM Eval Humming f16 (B200 - TEMPORARY)
-  key: lm-eval-humming-f16-b200
-  timeout_in_minutes: 50
+- label: LM Eval Humming (B200 - TEMPORARY)
+  key: lm-eval-humming-b200
+  timeout_in_minutes: 30
  device: b200-k8s
  optional: true
  num_devices: 1
@@ -221,30 +189,14 @@ steps:
  - vllm/model_executor/layers/quantization/humming.py
  - vllm/model_executor/layers/quantization/utils/humming_utils.py
  - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
-  - vllm/model_executor/layers/fused_moe/oracle/
-  - vllm/model_executor/kernels/linear/
+  - vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
  commands:
  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config.txt
-
- label: LM Eval Humming Act fp8/int8 (B200 - TEMPORARY)
-  key: lm-eval-humming-act-b200
-  timeout_in_minutes: 50
-  device: b200-k8s
-  optional: true
-  num_devices: 1
-  source_file_dependencies:
-  - vllm/model_executor/layers/quantization/humming.py
-  - vllm/model_executor/layers/quantization/utils/humming_utils.py
-  - vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
-  - vllm/model_executor/layers/fused_moe/oracle/
-  - vllm/model_executor/kernels/linear/
-  commands:
  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-fp8.txt
-  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-int8.txt

 - label: LM Eval TurboQuant KV Cache
  key: lm-eval-turboquant-kv-cache
-  timeout_in_minutes: 55
+  timeout_in_minutes: 75
  device: h200_18gb
  source_file_dependencies:
  - vllm/model_executor/layers/quantization/turboquant/
@@ -256,7 +208,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
@@ -270,7 +222,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
@@ -284,7 +236,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
@@ -298,50 +250,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/
@@ -24,13 +24,13 @@ steps:
  mirror:
    amd:
      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/
@@ -64,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/
@@ -89,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:
@@ -102,15 +101,14 @@ 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
+    - pip install -U git+https://github.com/robertgshaw2-redhat/lm-evaluation-harness.git@streaming-api
    - pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
  mirror:
    amd:
      device: mi325_1
-      timeout_in_minutes: 75
+      timeout_in_minutes: 60
      depends_on:
      - image-build-amd

@@ -174,7 +172,7 @@ steps:

 - label: Regression
  key: regression
-  timeout_in_minutes: 30
+  timeout_in_minutes: 20
  device: h200_18gb
  source_file_dependencies:
  - vllm/config/
@@ -190,14 +188,14 @@ steps:
  - vllm/v1/
  - tests/test_regression
  commands:
-  - pip install 'modelscope<1.38'
+  - pip install modelscope
  - pytest -v -s test_regression.py
  working_dir: "/vllm-workspace/tests" # optional

 - label: Examples
  device: h200_35gb
  key: examples
-  timeout_in_minutes: 40
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/examples"
  source_file_dependencies:
  - vllm/entrypoints
@@ -239,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/
@@ -283,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:
@@ -294,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/
@@ -321,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/
@@ -353,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
@@ -371,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
@@ -383,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
@@ -397,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
@@ -413,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
@@ -432,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
@@ -5,7 +5,7 @@ steps:
 - label: Model Runner V2 Core Tests
  device: h200_35gb
  key: model-runner-v2-core-tests
-  timeout_in_minutes: 35
+  timeout_in_minutes: 45
  source_file_dependencies:
  - vllm/v1/worker/gpu/
  - vllm/v1/worker/gpu_worker.py
@@ -18,7 +18,9 @@ steps:
  - set -x
  - export VLLM_USE_V2_MODEL_RUNNER=1
  - pytest -v -s v1/engine/test_llm_engine.py -k "not test_engine_metrics"
-  - pytest -v -s v1/e2e/general/test_async_scheduling.py -k "not ngram"
+  # This requires eager until we sort out CG correctness issues.
+  # TODO: remove ENFORCE_EAGER here after https://github.com/vllm-project/vllm/pull/32936 is merged.
+  - ENFORCE_EAGER=1 pytest -v -s v1/e2e/general/test_async_scheduling.py -k "not ngram"
  - pytest -v -s v1/e2e/general/test_context_length.py
  - pytest -v -s v1/e2e/general/test_min_tokens.py
  # Temporary hack filter to exclude ngram spec decoding based tests.
@@ -27,7 +29,7 @@ steps:
 - label: Model Runner V2 Examples
  device: h200_35gb
  key: model-runner-v2-examples
-  timeout_in_minutes: 35
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/examples"
  source_file_dependencies:
    - vllm/v1/worker/gpu/
@@ -63,7 +65,7 @@ steps:

 - label: Model Runner V2 Distributed (2 GPUs)
  key: model-runner-v2-distributed-2-gpus
-  timeout_in_minutes: 30
+  timeout_in_minutes: 45
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -84,7 +86,7 @@ steps:

 - label: Model Runner V2 Pipeline Parallelism (4 GPUs)
  key: model-runner-v2-pipeline-parallelism-4-gpus
-  timeout_in_minutes: 50
+  timeout_in_minutes: 60
  working_dir: "/vllm-workspace/tests"
  num_devices: 4
  source_file_dependencies:
@@ -4,8 +4,9 @@ depends_on:
 steps:
 - label: Basic Models Tests (Initialization)
  key: basic-models-tests-initialization
-  timeout_in_minutes: 25
+  timeout_in_minutes: 45
  device: h200_18gb
+  torch_nightly: true
  source_file_dependencies:
  - vllm/
  - tests/models/test_initialization.py
@@ -13,11 +14,13 @@ steps:
  commands:
    # Run a subset of model initialization tests
    - pytest -v -s models/test_initialization.py::test_can_initialize_small_subset
+  mirror:
+    torch_nightly: {}

 - label: Basic Models Tests (Extra Initialization) %N
  device: h200_35gb
  key: basic-models-tests-extra-initialization
-  timeout_in_minutes: 100
+  timeout_in_minutes: 45
  source_file_dependencies:
  - vllm/model_executor/models/
  - tests/models/test_initialization.py
@@ -27,19 +30,21 @@ steps:
    # subset of supported models (the complement of the small subset in the above
    # test.) Also run if model initialization test file is modified
    - pytest -v -s models/test_initialization.py -k 'not test_can_initialize_small_subset' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
-  parallelism: 4
+  parallelism: 2
+  mirror:
+    torch_nightly: {}

 - label: Basic Models Tests (Other)
  device: h200_35gb
  key: basic-models-tests-other
-  timeout_in_minutes: 35
+  timeout_in_minutes: 45
  source_file_dependencies:
  - vllm/
  - tests/models/test_terratorch.py
-  - tests/models/transformers/test_backend.py
+  - tests/models/test_transformers.py
  - tests/models/test_registry.py
  commands:
-    - pytest -v -s models/test_terratorch.py models/transformers/test_backend.py models/test_registry.py
+    - pytest -v -s models/test_terratorch.py models/test_transformers.py models/test_registry.py
  mirror:
    amd:
      device: mi325_1
@@ -50,12 +55,11 @@ 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
  - tests/models/test_vision.py
-  - tests/models/transformers/fusers/
  device: cpu-small
  commands:
-    - pytest -v -s models/test_utils.py models/test_vision.py models/transformers/fusers/
+    - pytest -v -s models/test_utils.py models/test_vision.py
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Distributed Model Tests (2 GPUs)
  key: distributed-model-tests-2-gpus
-  timeout_in_minutes: 60
+  timeout_in_minutes: 50
  working_dir: "/vllm-workspace/tests"
  num_devices: 2
  source_file_dependencies:
@@ -17,7 +17,7 @@ steps:
  - TARGET_TEST_SUITE=L4 pytest basic_correctness/ -v -s -m 'distributed(num_gpus=2)'
  - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s model_executor/model_loader/test_sharded_state_loader.py -m '(not slow_test)'
  # Avoid importing model tests that cause CUDA reinitialization error
-  - pytest models/transformers/test_backend.py -v -s -m 'distributed(num_gpus=2)'
+  - pytest models/test_transformers.py -v -s -m 'distributed(num_gpus=2)'
  - pytest models/language -v -s -m 'distributed(num_gpus=2)'
  - pytest models/multimodal/generation/test_phi4siglip.py -v -s -m 'distributed(num_gpus=2)'
  - pytest models/multimodal -v -s -m 'distributed(num_gpus=2)' --ignore models/multimodal/generation/test_whisper.py --ignore models/multimodal/generation/test_phi4siglip.py
@@ -4,7 +4,7 @@ depends_on:
 steps:
 - label: Language Models Tests (Standard)
  key: language-models-tests-standard
-  timeout_in_minutes: 30
+  timeout_in_minutes: 25
  device: h200_18gb
  source_file_dependencies:
  - vllm/
@@ -14,6 +14,7 @@ steps:
    - pip freeze | grep -E 'torch'
    - pytest -v -s models/language -m 'core_model and (not slow_test)'
  mirror:
+    torch_nightly: {}
    amd:
      device: mi300_1
      depends_on:
@@ -21,7 +22,7 @@ steps:

 - 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
@@ -34,6 +35,7 @@ steps:
    - pytest -v -s models/language -m 'core_model and slow_test' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
  parallelism: 2
  mirror:
+    torch_nightly: {}
    amd:
      device: mi300_1
      depends_on:
@@ -52,7 +54,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
@@ -65,9 +67,10 @@ steps:
    - pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
  parallelism: 2
  mirror:
+    torch_nightly: {}
    amd:
      device: mi325_1
-      timeout_in_minutes: 70
+      timeout_in_minutes: 90
      depends_on:
      - image-build-amd
      commands:
@@ -78,7 +81,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/
@@ -92,7 +95,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:
@@ -104,7 +107,7 @@ steps:
 - label: Language Models Test (Extended Pooling)  # 36min
  device: h200_35gb
  key: language-models-test-extended-pooling
-  timeout_in_minutes: 70
+  timeout_in_minutes: 50
  optional: true
  source_file_dependencies:
  - vllm/
@@ -120,7 +123,7 @@ steps:

 - 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,14 +20,13 @@ 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/
  - tests/models/multimodal
  commands:
    - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen3 or gemma"
-    - pytest -v -s models/multimodal/generation/test_mm_prefix_lm.py -m core_model
    - pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model
  mirror:
    amd:
@@ -38,7 +37,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,12 +53,12 @@ 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
  commands:
-    - pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py --ignore models/multimodal/generation/test_mm_prefix_lm.py --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/generation/test_vit_cudagraph.py --ignore models/multimodal/processing
+    - pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py  --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/generation/test_vit_cudagraph.py --ignore models/multimodal/processing
    - pytest -v -s models/multimodal/generation/test_vit_cudagraph.py -m core_model
    - pytest models/multimodal/generation/test_memory_leak.py -m core_model
    - cd .. && VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s tests/models/multimodal/generation/test_whisper.py -m core_model  # Otherwise, mp_method="spawn" doesn't work
@@ -69,23 +68,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 +95,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/
@@ -165,7 +163,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
@@ -110,13 +110,13 @@ steps:
  mirror:
    amd:
      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
@@ -152,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
@@ -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
@@ -119,7 +119,7 @@

 # Transformers modeling backend
 /vllm/model_executor/models/transformers @hmellor
-/tests/models/transformers @hmellor
+/tests/models/test_transformers.py @hmellor

 # Docs
 /docs/mkdocs @hmellor
@@ -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:
@@ -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.")
@@ -359,7 +350,9 @@ endif()
 if(VLLM_GPU_LANG STREQUAL "HIP")
  set(VLLM_EXT_SRC
    "csrc/torch_bindings.cpp"
-    "csrc/custom_quickreduce.cu")
+    "csrc/custom_quickreduce.cu"
+    "csrc/cuda_view.cu"
+    "csrc/libtorch_stable/cuda_utils_kernels.cu")

  message(STATUS "Enabling C extension.")
  define_extension_target(
@@ -387,10 +380,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
  #
  set(VLLM_STABLE_EXT_SRC
    "csrc/libtorch_stable/torch_bindings.cpp"
-    "csrc/libtorch_stable/cuda_view.cu"
-    "csrc/libtorch_stable/cuda_utils_kernels.cu"
    "csrc/libtorch_stable/activation_kernels.cu"
-    "csrc/libtorch_stable/ngram_embedding_kernels.cu"
    "csrc/libtorch_stable/quantization/activation_kernels.cu"
    "csrc/libtorch_stable/quantization/w8a8/int8/scaled_quant.cu"
    "csrc/libtorch_stable/quantization/w8a8/fp8/common.cu"
@@ -409,6 +399,9 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
    "csrc/libtorch_stable/sampler.cu"
    "csrc/libtorch_stable/topk.cu"
    "csrc/libtorch_stable/mamba/selective_scan_fwd.cu"
+    "csrc/libtorch_stable/attention/paged_attention_v1.cu"
+    "csrc/libtorch_stable/attention/paged_attention_v2.cu"
+    "csrc/libtorch_stable/cache_kernels.cu"
    "csrc/libtorch_stable/cache_kernels.cu"
    "csrc/libtorch_stable/cache_kernels_fused.cu"
    "csrc/libtorch_stable/custom_all_reduce.cu"
@@ -466,6 +459,8 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
    FetchContent_MakeAvailable(cutlass)

    list(APPEND VLLM_STABLE_EXT_SRC
+      "csrc/libtorch_stable/cuda_view.cu"
+      "csrc/libtorch_stable/cuda_utils_kernels.cu"
      "csrc/libtorch_stable/cutlass_extensions/common.cpp"
      "csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_entry.cu"
      "csrc/libtorch_stable/quantization/fp4/nvfp4_quant_entry.cu"
@@ -1096,15 +1091,14 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
    USE_SABI 3
    WITH_SOABI)

-  # Set TORCH_TARGET_VERSION for stable ABI compatibility.
-  # This ensures we only use C-shim APIs available in PyTorch 2.11.
-  # _C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
-  # which is currently set to 2.11.
-  target_compile_definitions(_C_stable_libtorch PRIVATE
-    TORCH_TARGET_VERSION=0x020B000000000000ULL)
-
  # Needed to use cuda/hip APIs from C-shim
  if(VLLM_GPU_LANG STREQUAL "CUDA")
+    # Set TORCH_TARGET_VERSION for stable ABI compatibility.
+    # This ensures we only use C-shim APIs available in PyTorch 2.11.
+    # _C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
+    # which is currently set to 2.11.
+    target_compile_definitions(_C_stable_libtorch PRIVATE
+      TORCH_TARGET_VERSION=0x020B000000000000ULL)
    target_compile_definitions(_C_stable_libtorch PRIVATE USE_CUDA)
    if(COOPERATIVE_TOPK_ARCHS)
      target_compile_definitions(_C_stable_libtorch PRIVATE
@@ -1114,6 +1108,12 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
    target_compile_definitions(_C_stable_libtorch PRIVATE
      CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1)
  elseif(VLLM_GPU_LANG STREQUAL "HIP")
+    # Set TORCH_TARGET_VERSION for stable ABI compatibility.
+    # This ensures we only use C-shim APIs available in PyTorch 2.10.
+    # _C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
+    # which is currently set to 2.10.
+    target_compile_definitions(_C_stable_libtorch PRIVATE
+      TORCH_TARGET_VERSION=0x020A000000000000ULL)
    target_compile_definitions(_C_stable_libtorch PRIVATE USE_ROCM)
  endif()

@@ -1321,20 +1321,25 @@ define_extension_target(
  USE_SABI 3
  WITH_SOABI)

-# Set TORCH_TARGET_VERSION for stable ABI compatibility.
-# This ensures we only use C-shim APIs available in PyTorch 2.11.
-# _moe_C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
-# which is currently set to 2.11.
-target_compile_definitions(_moe_C_stable_libtorch PRIVATE
-  TORCH_TARGET_VERSION=0x020B000000000000ULL)
-
 # Needed to use cuda/hip APIs from C-shim
 if(VLLM_GPU_LANG STREQUAL "CUDA")
+  # Set TORCH_TARGET_VERSION for stable ABI compatibility.
+  # This ensures we only use C-shim APIs available in PyTorch 2.11.
+  # _moe_C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
+  # which is currently set to 2.11.
+  target_compile_definitions(_moe_C_stable_libtorch PRIVATE
+    TORCH_TARGET_VERSION=0x020B000000000000ULL)
  target_compile_definitions(_moe_C_stable_libtorch PRIVATE USE_CUDA)
  # Needed by CUTLASS kernels
  target_compile_definitions(_moe_C_stable_libtorch PRIVATE
    CUTLASS_ENABLE_DIRECT_CUDA_DRIVER_CALL=1)
 elseif(VLLM_GPU_LANG STREQUAL "HIP")
+  # Set TORCH_TARGET_VERSION for stable ABI compatibility.
+  # This ensures we only use C-shim APIs available in PyTorch 2.10.
+  # _moe_C_stable_libtorch is abi compatible with PyTorch >= TORCH_TARGET_VERSION
+  # which is currently set to 2.10.
+  target_compile_definitions(_moe_C_stable_libtorch PRIVATE
+    TORCH_TARGET_VERSION=0x020A000000000000ULL)
  target_compile_definitions(_moe_C_stable_libtorch PRIVATE USE_ROCM)
 endif()

@@ -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
    ):
@@ -132,10 +132,8 @@ def benchmark_function(
    reset_memory_stats()

    # Benchmark
-    start_events = [
-        torch.cuda.Event(enable_timing=True) for _ in range(benchmark_iters)
-    ]
-    end_events = [torch.cuda.Event(enable_timing=True) for _ in range(benchmark_iters)]
+    start_events = [torch.Event(enable_timing=True) for _ in range(benchmark_iters)]
+    end_events = [torch.Event(enable_timing=True) for _ in range(benchmark_iters)]

    for i in range(benchmark_iters):
        logits_copy = logits.clone()
@@ -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"
@@ -134,8 +134,8 @@ def benchmark_config(
    torch.accelerator.synchronize()

    # Benchmark
-    start = torch.cuda.Event(enable_timing=True)
-    end = torch.cuda.Event(enable_timing=True)
+    start = torch.Event(enable_timing=True)
+    end = torch.Event(enable_timing=True)
    start.record()
    for _ in range(num_iters):
        with override_config(config):
@@ -19,11 +19,13 @@ from vllm.utils.torch_utils import (
 logger = init_logger(__name__)

 NUM_BLOCKS = 128 * 1024
+PARTITION_SIZE = 512
 PARTITION_SIZE_ROCM = 256


@torch.inference_mode()
 def main(
+    version: str,
    num_seqs: int,
    seq_len: int,
    num_query_heads: int,
@@ -80,20 +82,27 @@ def main(

    # Prepare for the paged attention kernel.
    output = torch.empty_like(query)
-    num_partitions = (max_seq_len + PARTITION_SIZE_ROCM - 1) // PARTITION_SIZE_ROCM
-    tmp_output = torch.empty(
-        size=(num_seqs, num_query_heads, num_partitions, head_size),
-        dtype=output.dtype,
-        device=output.device,
-    )
-    exp_sums = torch.empty(
-        size=(num_seqs, num_query_heads, num_partitions),
-        dtype=torch.float32,
-        device=output.device,
-    )
-    max_logits = torch.empty_like(exp_sums)
+    if version == "v2":
+        if current_platform.is_rocm():
+            global PARTITION_SIZE
+            if not args.custom_paged_attn and not current_platform.is_navi():
+                PARTITION_SIZE = 1024
+            else:
+                PARTITION_SIZE = PARTITION_SIZE_ROCM
+        num_partitions = (max_seq_len + PARTITION_SIZE - 1) // PARTITION_SIZE
+        tmp_output = torch.empty(
+            size=(num_seqs, num_query_heads, num_partitions, head_size),
+            dtype=output.dtype,
+            device=output.device,
+        )
+        exp_sums = torch.empty(
+            size=(num_seqs, num_query_heads, num_partitions),
+            dtype=torch.float32,
+            device=output.device,
+        )
+        max_logits = torch.empty_like(exp_sums)

-    def run_benchmark(num_iters: int, profile: bool = False) -> float:
+    def run_cuda_benchmark(num_iters: int, profile: bool = False) -> float:
        torch.accelerator.synchronize()
        if profile:
            torch.cuda.cudart().cudaProfilerStart()
@@ -103,26 +112,67 @@ def main(
        k_scale = v_scale = torch.tensor(1.0, dtype=torch.float32, device=device)

        for _ in range(num_iters):
-            ops.paged_attention_rocm(
-                output,
-                exp_sums,
-                max_logits,
-                tmp_output,
-                query,
-                key_cache,
-                value_cache,
-                num_kv_heads,
-                scale,
-                block_tables,
-                seq_lens,
-                None,
-                block_size,
-                max_seq_len,
-                alibi_slopes,
-                kv_cache_dtype,
-                k_scale,
-                v_scale,
-            )
+            if version == "v1":
+                ops.paged_attention_v1(
+                    output,
+                    query,
+                    key_cache,
+                    value_cache,
+                    num_kv_heads,
+                    scale,
+                    block_tables,
+                    seq_lens,
+                    block_size,
+                    max_seq_len,
+                    alibi_slopes,
+                    kv_cache_dtype,
+                    k_scale,
+                    v_scale,
+                )
+            elif version == "v2":
+                if not args.custom_paged_attn:
+                    ops.paged_attention_v2(
+                        output,
+                        exp_sums,
+                        max_logits,
+                        tmp_output,
+                        query,
+                        key_cache,
+                        value_cache,
+                        num_kv_heads,
+                        scale,
+                        block_tables,
+                        seq_lens,
+                        block_size,
+                        max_seq_len,
+                        alibi_slopes,
+                        kv_cache_dtype,
+                        k_scale,
+                        v_scale,
+                    )
+                else:
+                    ops.paged_attention_rocm(
+                        output,
+                        exp_sums,
+                        max_logits,
+                        tmp_output,
+                        query,
+                        key_cache,
+                        value_cache,
+                        num_kv_heads,
+                        scale,
+                        block_tables,
+                        seq_lens,
+                        None,
+                        block_size,
+                        max_seq_len,
+                        alibi_slopes,
+                        kv_cache_dtype,
+                        k_scale,
+                        v_scale,
+                    )
+            else:
+                raise ValueError(f"Invalid version: {version}")
        torch.accelerator.synchronize()

        end_time = time.perf_counter()
@@ -132,6 +182,7 @@ def main(

    # Warmup.
    print("Warming up...")
+    run_benchmark = run_cuda_benchmark
    run_benchmark(num_iters=3, profile=False)

    # Benchmark.
@@ -144,13 +195,12 @@ def main(

 if __name__ == "__main__":
    logger.warning(
-        "This script benchmarks the ROCm paged attention kernel. "
+        "This script benchmarks the paged attention kernel. "
        "By default this is no longer used in vLLM inference."
    )
-    if not current_platform.is_rocm():
-        raise RuntimeError("This benchmark requires the ROCm platform.")

    parser = FlexibleArgumentParser(description="Benchmark the paged attention kernel.")
+    parser.add_argument("--version", type=str, choices=["v1", "v2"], default="v2")
    parser.add_argument("--batch-size", type=int, default=8)
    parser.add_argument("--seq-len", type=int, default=4096)
    parser.add_argument("--num-query-heads", type=int, default=64)
@@ -158,7 +208,7 @@ if __name__ == "__main__":
    parser.add_argument(
        "--head-size",
        type=int,
-        choices=[64, 128],
+        choices=[64, 80, 96, 112, 120, 128, 192, 256],
        default=128,
    )
    parser.add_argument("--block-size", type=int, choices=[16, 32], default=16)
@@ -174,7 +224,11 @@ if __name__ == "__main__":
        choices=["auto", "fp8", "fp8_e5m2", "fp8_e4m3"],
        default="auto",
        help="Data type for kv cache storage. If 'auto', will use model "
-        "data type. ROCm (AMD GPU) supports fp8 (=fp8_e4m3)",
+        "data type. CUDA 11.8+ supports fp8 (=fp8_e4m3) and fp8_e5m2. "
+        "ROCm (AMD GPU) supports fp8 (=fp8_e4m3)",
+    )
+    parser.add_argument(
+        "--custom-paged-attn", action="store_true", help="Use custom paged attention"
    )
    args = parser.parse_args()
    print(args)
@@ -182,6 +236,7 @@ if __name__ == "__main__":
    if args.num_query_heads % args.num_kv_heads != 0:
        raise ValueError("num_query_heads must be divisible by num_kv_heads")
    main(
+        version=args.version,
        num_seqs=args.batch_size,
        seq_len=args.seq_len,
        num_query_heads=args.num_query_heads,
@@ -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)
@@ -170,8 +170,8 @@ def benchmark_config(
                graph.replay()
            torch.accelerator.synchronize()

-            start = torch.cuda.Event(enable_timing=True)
-            end = torch.cuda.Event(enable_timing=True)
+            start = torch.Event(enable_timing=True)
+            end = torch.Event(enable_timing=True)
            latencies: list[float] = []
            for _ in range(num_iters):
                start.record()
@@ -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
@@ -15,7 +15,6 @@ endif()
 #
 set(ENABLE_X86_ISA $ENV{VLLM_CPU_X86})
 set(ENABLE_ARM_BF16 $ENV{VLLM_CPU_ARM_BF16})
-set(ENABLE_RVV_BF16 $ENV{VLLM_CPU_RVV_BF16})

 include_directories("${CMAKE_SOURCE_DIR}/csrc")

@@ -111,13 +110,6 @@ else()
        set(ARM_BF16_FOUND ON)
        message(STATUS "ARM BF16 support enabled via VLLM_CPU_ARM_BF16 environment variable")
    endif()
-    # Some kernels (e.g. Bianbu on Spacemit X100) do not report zvfbfmin
-    # in /proc/cpuinfo despite hardware support. VLLM_CPU_RVV_BF16=1
-    # overrides the detection result.
-    if (ENABLE_RVV_BF16)
-        set(RVV_BF16_FOUND ON)
-        message(STATUS "RVV BF16 support enabled via VLLM_CPU_RVV_BF16 environment variable")
-    endif()
 endif()

 if (CMAKE_SYSTEM_PROCESSOR MATCHES "x86_64|amd64" OR ENABLE_X86_ISA)
@@ -186,10 +178,7 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
    # Override with -DVLLM_RVV_VLEN=128 or -DVLLM_RVV_VLEN=256 for RVV.
    if(NOT DEFINED VLLM_RVV_VLEN)
        # Auto-detect: find the largest zvl<N>b in /proc/cpuinfo isa line.
-        # Skip when cross-compiling — /proc/cpuinfo describes the build host.
-        if(CMAKE_CROSSCOMPILING)
-            message(STATUS "Cross-compiling: skipping VLEN auto-detection from /proc/cpuinfo")
-        elseif(EXISTS /proc/cpuinfo)
+        if(EXISTS /proc/cpuinfo)
            file(READ /proc/cpuinfo _cpuinfo)
            set(_best 0)
            foreach(_n IN ITEMS 128 256 512 1024)
@@ -197,13 +186,6 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
                    set(_best ${_n})
                endif()
            endforeach()
-            # Only VLEN=128 and VLEN=256 are supported by the RVV kernels.
-            if(_best GREATER 256)
-                message(WARNING
-                    "Detected VLEN=${_best} but only 128/256 are supported; "
-                    "clamping to 256")
-                set(_best 256)
-            endif()
            if(_best GREATER 0)
                set(VLLM_RVV_VLEN ${_best})
            endif()
@@ -213,9 +195,9 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
        if(NOT DEFINED VLLM_RVV_VLEN AND (RVV_FP16_FOUND OR RVV_BF16_FOUND))
            message(FATAL_ERROR
                "RISC-V RVV is available but VLEN could not be auto-detected. "
-                "Please specify VLEN explicitly via CMAKE_ARGS:\n"
-                "  CMAKE_ARGS='-DVLLM_RVV_VLEN=128'   (for VLEN=128 hardware)\n"
-                "  CMAKE_ARGS='-DVLLM_RVV_VLEN=256'   (for VLEN=256 hardware, e.g. Spacemit X100)")
+                "Please specify VLEN explicitly:\n"
+                "  -DVLLM_RVV_VLEN=128   (for VLEN=128 hardware)\n"
+                "  -DVLLM_RVV_VLEN=256   (for VLEN=256 hardware, e.g. Spacemit X100)")
        endif()
    endif()
    if(VLLM_RVV_VLEN AND VLLM_RVV_VLEN GREATER 0)
@@ -227,7 +209,7 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
            message(STATUS "BF16 extension detected")
            set(MARCH_FLAGS -march=rv64gcv_zvfh_zfbfmin_zvfbfmin_zvl${VLLM_RVV_VLEN}b -mrvv-vector-bits=zvl -mabi=lp64d)
        elseif(RVV_FP16_FOUND)
-            message(WARNING "BF16 functionality is not available.")
+            message(WARNING "BF16 functionality is not available")
            set(MARCH_FLAGS -march=rv64gcv_zvfh_zvl${VLLM_RVV_VLEN}b -mrvv-vector-bits=zvl -mabi=lp64d)
        else()
            message(STATUS "compile riscv with scalar (no FP16/BF16)")
@@ -347,7 +329,7 @@ if (ENABLE_X86_ISA OR (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND) OR POWER9_FOUND
    set(ONEDNN_ENABLE_PRIMITIVE "MATMUL;REORDER")
    set(ONEDNN_BUILD_GRAPH "OFF")
    set(ONEDNN_ENABLE_JIT_PROFILING "ON")
-    set(ONEDNN_ENABLE_ITT_TASKS "ON")
+    set(ONEDNN_ENABLE_ITT_TASKS "OFF")
    set(ONEDNN_ENABLE_MAX_CPU_ISA "ON")
    set(ONEDNN_ENABLE_CPU_ISA_HINTS "ON")
    set(ONEDNN_VERBOSE "ON")
@@ -29,8 +29,7 @@ if(DEEPGEMM_SRC_DIR)
 else()
  # Keep in sync with tools/install_deepgemm.sh
  set(_DEEPGEMM_UPSTREAM_REPO "https://github.com/deepseek-ai/DeepGEMM.git")
-  # NOTE: This is currently targeting nv-dev branch due to sm120 support
-  set(_DEEPGEMM_UPSTREAM_TAG "a6b593d2826719dcf4892609af7b84ee23aaf32a")
+  set(_DEEPGEMM_UPSTREAM_TAG "891d57b4db1071624b5c8fa0d1e51cb317fa709f")

  set(_deepgemm_fc_root "${FETCHCONTENT_BASE_DIR}")
  if(NOT _deepgemm_fc_root)
@@ -19,7 +19,7 @@ else()
  FetchContent_Declare(
        flashmla
        GIT_REPOSITORY https://github.com/vllm-project/FlashMLA
-        GIT_TAG b70aff3d110a2b1a037e62eac295166b5143643a
+        GIT_TAG a6ec2ba7bd0a7dff98b3f4d3e6b52b159c48d78b
        GIT_PROGRESS TRUE
        CONFIGURE_COMMAND ""
        BUILD_COMMAND ""
@@ -17,7 +17,7 @@ else()
  FetchContent_Declare(
    fmha_sm100
    GIT_REPOSITORY https://github.com/vllm-project/MSA.git
-    GIT_TAG 2e63ec37a0fc29bc20f39cd1a52e0f5affc33a73
+    GIT_TAG fee783153f3efe57e3e933c5cb7e267a7cebcfb5
    GIT_PROGRESS TRUE
    CONFIGURE_COMMAND ""
    BUILD_COMMAND ""
@@ -39,7 +39,7 @@ else()
  FetchContent_Declare(
          vllm-flash-attn
          GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
-          GIT_TAG bb9a72e7dde0dc614ffc663e052cd6a19ce73a42
+          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
@@ -13,8 +13,7 @@ static inline cpu_attention::Fp8KVCacheDataType parse_fp8_kv_dtype(

 bool cpu_attn_has_isa(const std::string& isa) {
  if (isa == "rvv") {
-#if defined(__riscv) && defined(__riscv_v_min_vlen) && \
-    (__riscv_v_min_vlen == 128 || __riscv_v_min_vlen == 256)
+#if defined(__riscv) && defined(__riscv_v_min_vlen) && __riscv_v_min_vlen == 128
    return true;
 #else
    return false;
@@ -417,10 +417,8 @@ class AttentionScheduler {
      has_decode_request = has_decode_request || (q_token_num == 1);
      decode_only_batch = decode_only_batch && (q_token_num == 1);
    }
-    const int32_t original_q_head_per_kv =
-        input.num_heads_q / input.num_heads_kv;
-    int32_t q_head_per_kv = original_q_head_per_kv;
-    const bool supports_gqa = original_q_head_per_kv <= max_num_q_per_iter;
+    int32_t q_head_per_kv = input.num_heads_q / input.num_heads_kv;
+    const bool supports_gqa = q_head_per_kv <= max_num_q_per_iter;
    const bool use_gqa_fast_path = supports_gqa && decode_only_batch;
    const bool use_gqa_scratchpad = supports_gqa && has_decode_request;
    if (!use_gqa_scratchpad) {
@@ -673,62 +671,22 @@ class AttentionScheduler {
    metadata_ptr->effective_thread_num = effective_thread_num;

    {
+      // when q_tile_size = max_num_q_per_iter, requires max
+      // attention_scratchpad_size
      AttentionScratchPad sc(0, *metadata_ptr, 0x0);
-      int64_t max_attention_scratchpad_size = 0;
-
-      for (const AttentionWorkItemGroup& item : workitems) {
-        const bool curr_use_gqa =
-            use_gqa_fast_path || (supports_gqa && item.q_token_num == 1);
-        const int32_t curr_q_heads_per_kv =
-            curr_use_gqa ? original_q_head_per_kv : 1;
-        const int32_t curr_default_q_tile_token_num =
-            default_tile_size / curr_q_heads_per_kv;
-
-        for (int32_t q_token_offset = 0; q_token_offset < item.q_token_num;
-             q_token_offset += curr_default_q_tile_token_num) {
-          const int32_t actual_q_token_num = std::min(
-              curr_default_q_tile_token_num, item.q_token_num - q_token_offset);
-          const int32_t q_head_tile_size =
-              actual_q_token_num * curr_q_heads_per_kv;
-          const int32_t rounded_q_head_tile_size =
-              ((q_head_tile_size + max_num_q_per_iter - 1) /
-               max_num_q_per_iter) *
-              max_num_q_per_iter;
-
-          const int64_t n = AttentionScheduler::calcu_tile_size_with_constant_q(
-              cache_size, input.head_dim, input.elem_size,
-              input.q_buffer_elem_size, input.logits_buffer_elem_size,
-              input.output_buffer_elem_size, max_num_q_per_iter,
-              kv_len_alignment, rounded_q_head_tile_size,
-              rounded_q_head_tile_size <= max_num_q_per_iter);
-
-          sc.update(input.head_dim, input.q_buffer_elem_size,
-                    input.logits_buffer_elem_size,
-                    input.output_buffer_elem_size, max_num_q_per_iter,
-                    rounded_q_head_tile_size, n);
-
-          max_attention_scratchpad_size = std::max(
-              max_attention_scratchpad_size, sc.get_thread_scratchpad_size());
-        }
-      }
-
+      int64_t n = AttentionScheduler::calcu_tile_size_with_constant_q(
+          cache_size, input.head_dim, input.elem_size, input.q_buffer_elem_size,
+          input.logits_buffer_elem_size, input.output_buffer_elem_size,
+          max_num_q_per_iter, kv_len_alignment, max_num_q_per_iter, true);
+      sc.update(input.head_dim, input.q_buffer_elem_size,
+                input.logits_buffer_elem_size, input.output_buffer_elem_size,
+                max_num_q_per_iter, max_num_q_per_iter, n);
      metadata_ptr->attention_scratchpad_size_per_thread =
-          ((max_attention_scratchpad_size + 63) / 64) * 64;
-
-      int32_t max_reduction_q_head_tile_size = 0;
-      for (const ReductionWorkItemGroup& item : reduce_workitems) {
-        const bool curr_use_gqa =
-            use_gqa_fast_path || (supports_gqa && item.q_token_id_num == 1);
-        const int32_t curr_q_heads_per_kv =
-            curr_use_gqa ? original_q_head_per_kv : 1;
-
-        max_reduction_q_head_tile_size =
-            std::max(max_reduction_q_head_tile_size,
-                     item.q_token_id_num * curr_q_heads_per_kv);
-      }
+          ((sc.get_thread_scratchpad_size() + 63) / 64) * 64;

      sc.update(0, metadata_ptr->reduction_split_num, input.head_dim,
-                max_reduction_q_head_tile_size, input.output_buffer_elem_size);
+                q_head_per_kv * split_kv_q_token_num_threshold,
+                input.output_buffer_elem_size);
      metadata_ptr->reduction_scratchpad_size_per_kv_head =
          ((sc.get_reduction_scratchpad_size() + 63) / 64) * 64;
    }
@@ -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
@@ -214,18 +214,11 @@ struct BF16Vec32 : public Vec<BF16Vec32> {

  explicit BF16Vec32(const BF16Vec8& v) {
    fixed_u16x8_t u16_val = bf16_to_u16(v.reg);
-    // Widen LMUL_128 → LMUL_256 so vslideup operands share a type.
-    // At VLEN=256 this is mf2→m1 (both integer); at VLEN=128 it is m1→m2.
-    fixed_u16x16_t ext =
-        RVVI4(__riscv_vlmul_ext_v_u16, LMUL_128, _u16, LMUL_256)(u16_val);
-    // Build 16-element half: place the 8 elements at offsets 0 and 8.
-    fixed_u16x16_t half = RVVI(__riscv_vmv_v_x_u16, LMUL_256)(0, 16);
-    half = RVVI(__riscv_vslideup_vx_u16, LMUL_256)(half, ext, 0, 8);
-    half = RVVI(__riscv_vslideup_vx_u16, LMUL_256)(half, ext, 8, 16);
-    // Double to LMUL_512 (m1→m2 at VLEN=256, m2→m4 at VLEN=128).
-    fixed_u16x32_t dst =
-        RVVI4(__riscv_vcreate_v_u16, LMUL_256, _u16, LMUL_512)(half, half);
-    reg = RVVI4(__riscv_vreinterpret_v_u16, LMUL_512, _bf16, LMUL_512)(dst);
+    fixed_u16x32_t u16_combined =
+        RVVI4(__riscv_vcreate_v_u16, LMUL_128, _u16, LMUL_512)(
+            u16_val, u16_val, u16_val, u16_val);
+    reg = RVVI4(__riscv_vreinterpret_v_u16, LMUL_512, _bf16,
+                LMUL_512)(u16_combined);
  };

  void save(void* ptr) const {
@@ -630,29 +623,17 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
            data.reg, data.reg)) {};
  explicit FP32Vec16(const FP32Vec16& data) : reg(data.reg) {};
  explicit FP32Vec16(int64_t value, const FP32Vec16& lut) {
-    // Split into two 32-bit halves to avoid u64 @ LMUL_1024 (m8 on
-    // VLEN=128 / m4 on VLEN=256), which causes heavy register spilling.
-    constexpr int HALF = VEC_ELEM_NUM / 2;
-    const auto q = static_cast<uint64_t>(value);
-    const uint32_t lo = static_cast<uint32_t>(q);
-    const uint32_t hi = static_cast<uint32_t>(q >> 32);
-
-    auto lane_ids = RVVI(__riscv_vid_v_u32, LMUL_256)(HALF);
-    auto shifts = RVVI(__riscv_vsll_vx_u32, LMUL_256)(lane_ids, 2, HALF);
-
-    auto packed_lo = RVVI(__riscv_vmv_v_x_u32, LMUL_256)(lo, HALF);
-    auto idx_lo = RVVI(__riscv_vand_vx_u32, LMUL_256)(
-        RVVI(__riscv_vsrl_vv_u32, LMUL_256)(packed_lo, shifts, HALF), 0xF,
-        HALF);
-
-    auto packed_hi = RVVI(__riscv_vmv_v_x_u32, LMUL_256)(hi, HALF);
-    auto idx_hi = RVVI(__riscv_vand_vx_u32, LMUL_256)(
-        RVVI(__riscv_vsrl_vv_u32, LMUL_256)(packed_hi, shifts, HALF), 0xF,
-        HALF);
-
-    auto idx =
-        RVVI4(__riscv_vcreate_v_u32, LMUL_256, _u32, LMUL_512)(idx_lo, idx_hi);
-    reg = RVVI(__riscv_vrgather_vv_f32, LMUL_512)(lut.reg, idx, VEC_ELEM_NUM);
+    const uint64_t q_values = static_cast<uint64_t>(value);
+    auto packed = RVVI(__riscv_vmv_v_x_u64, LMUL_1024)(q_values, VEC_ELEM_NUM);
+    auto lane_ids = RVVI(__riscv_vid_v_u64, LMUL_1024)(VEC_ELEM_NUM);
+    auto shifts =
+        RVVI(__riscv_vsll_vx_u64, LMUL_1024)(lane_ids, 2, VEC_ELEM_NUM);
+    auto shifted =
+        RVVI(__riscv_vsrl_vv_u64, LMUL_1024)(packed, shifts, VEC_ELEM_NUM);
+    auto idx64 =
+        RVVI(__riscv_vand_vx_u64, LMUL_1024)(shifted, 0xF, VEC_ELEM_NUM);
+    auto idx32 = RVVI(__riscv_vnsrl_wx_u32, LMUL_512)(idx64, 0, VEC_ELEM_NUM);
+    reg = RVVI(__riscv_vrgather_vv_f32, LMUL_512)(lut.reg, idx32, VEC_ELEM_NUM);
  }
  explicit FP32Vec16(const FP16Vec16& v);

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

@@ -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
@@ -0,0 +1,60 @@
+// TODO: Remove this once ROCm upgrade to torch 2.11.
+#include <torch/all.h>
+#include <torch/cuda.h>
+#include <cuda_runtime.h>
+
+// This function assumes that `cpu_tensor` is a CPU tensor,
+// and that UVA (Unified Virtual Addressing) is enabled.
+torch::Tensor get_cuda_view_from_cpu_tensor(torch::Tensor& cpu_tensor) {
+  TORCH_CHECK(cpu_tensor.device().is_cpu(), "Input tensor must be on CPU");
+
+  // handle empty tensor
+  if (cpu_tensor.numel() == 0) {
+    return torch::empty(cpu_tensor.sizes(),
+                        cpu_tensor.options().device(torch::kCUDA));
+  }
+
+  if (cpu_tensor.is_pinned()) {
+    // If CPU tensor is pinned, directly get the device pointer.
+    void* host_ptr = const_cast<void*>(cpu_tensor.data_ptr());
+    void* device_ptr = nullptr;
+    cudaError_t err = cudaHostGetDevicePointer(&device_ptr, host_ptr, 0);
+    TORCH_CHECK(err == cudaSuccess,
+                "cudaHostGetDevicePointer failed: ", cudaGetErrorString(err));
+
+    return torch::from_blob(
+        device_ptr, cpu_tensor.sizes(), cpu_tensor.strides(),
+        [base = cpu_tensor](void*) {},  // keep cpu tensor alive
+        cpu_tensor.options().device(torch::kCUDA));
+  }
+
+  // If CPU tensor is not pinned, allocate a new pinned memory buffer.
+  torch::Tensor contiguous_cpu = cpu_tensor.contiguous();
+  size_t nbytes = contiguous_cpu.nbytes();
+
+  void* host_ptr = nullptr;
+  cudaError_t err = cudaHostAlloc(&host_ptr, nbytes, cudaHostAllocMapped);
+  if (err != cudaSuccess) {
+    AT_ERROR("cudaHostAlloc failed: ", cudaGetErrorString(err));
+  }
+
+  err = cudaMemcpy(host_ptr, contiguous_cpu.data_ptr(), nbytes,
+                   cudaMemcpyDefault);
+  if (err != cudaSuccess) {
+    cudaFreeHost(host_ptr);
+    AT_ERROR("cudaMemcpy failed: ", cudaGetErrorString(err));
+  }
+
+  void* device_ptr = nullptr;
+  err = cudaHostGetDevicePointer(&device_ptr, host_ptr, 0);
+  if (err != cudaSuccess) {
+    cudaFreeHost(host_ptr);
+    AT_ERROR("cudaHostGetDevicePointer failed: ", cudaGetErrorString(err));
+  }
+
+  auto deleter = [host_ptr](void*) { cudaFreeHost(host_ptr); };
+
+  return torch::from_blob(device_ptr, contiguous_cpu.sizes(),
+                          contiguous_cpu.strides(), deleter,
+                          contiguous_cpu.options().device(torch::kCUDA));
+}
--- a/Show More
+++ b/Show More