[Bugfix][CI] Update Dockerfile dependency graph PNG (#45602)

Signed-off-by: sfeng33 <4florafeng@gmail.com>

.buildkite/ci_config_rocm.yaml

@@ -0,0 +1,23 @@
+name: vllm_rocm_ci
+job_dirs:
+  - ".buildkite/hardware_tests"
+run_all_patterns:
+  - "docker/Dockerfile.rocm"
+  - "docker/Dockerfile.rocm_base"
+  - "docker/ci-rocm.hcl"
+  - "docker/docker-bake-rocm.hcl"
+  - ".buildkite/hardware_tests/amd.yaml"
+  - ".buildkite/scripts/ci-bake-rocm.sh"
+  - ".buildkite/scripts/hardware_ci/run-amd-test.py"
+  - ".buildkite/scripts/hardware_ci/run-amd-test.sh"
+  - "CMakeLists.txt"
+  - "requirements/common.txt"
+  - "requirements/rocm.txt"
+  - "requirements/build/rocm.txt"
+  - "requirements/test/rocm.txt"
+  - "setup.py"
+  - "csrc/"
+  - "cmake/"
+run_all_exclude_patterns:
+  - "csrc/cpu/"
+  - "cmake/cpu_extension.cmake"

.buildkite/hardware_tests/amd.yaml

@@ -1,42 +1,73 @@
-group: Hardware - AMD Build 
+group: Hardware - AMD Build
 steps:
-  - label: "AMD: :docker: build image"
-    key: image-build-amd
+  # 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
+  # differ ci_base is rebuilt and pushed automatically.
+  - label: "AMD: :docker: ensure ci_base"
+    key: ensure-ci-base-amd
     depends_on: []
     device: amd_cpu
     no_plugin: true
     commands:
-    - >
-      docker build
-      --build-arg max_jobs=16
-      --build-arg REMOTE_VLLM=1
-      --build-arg ARG_PYTORCH_ROCM_ARCH='gfx90a;gfx942;gfx950'
-      --build-arg VLLM_BRANCH=$BUILDKITE_COMMIT
-      --tag "rocm/vllm-ci:${BUILDKITE_COMMIT}"
-      -f docker/Dockerfile.rocm
-      --target test
-      --no-cache
-      --progress plain .
-    - |
-      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
-      '
-    - docker push "rocm/vllm-ci:${BUILDKITE_COMMIT}"
+      - bash .buildkite/scripts/ci-bake-rocm.sh ci-base-rocm-ci-with-deps
     env:
       DOCKER_BUILDKIT: "1"
+      VLLM_BAKE_FILE: "docker/docker-bake-rocm.hcl"
+      PYTORCH_ROCM_ARCH: "gfx90a;gfx942;gfx950"
+      REMOTE_VLLM: "1"
+      VLLM_BRANCH: "$BUILDKITE_COMMIT"
+    retry:
+      automatic:
+        - exit_status: -1  # Agent was lost
+          limit: 1
+        - exit_status: -10  # Agent was lost
+          limit: 1
+
+  - label: "AMD: :docker: build test image and artifacts"
+    key: image-build-amd
+    depends_on:
+      - ensure-ci-base-amd
+    device: amd_cpu
+    no_plugin: true
+    commands:
+      - |
+        if [[ "${ROCM_CI_ARTIFACT_ONLY:-0}" == "1" ]]; then
+          echo "ROCM_CI_ARTIFACT_ONLY=1; building ROCm wheel artifact only"
+          IMAGE_TAG="" bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-artifacts
+        else
+          bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-wheel
+        fi
+      - |
+        docker run --rm --network=none --entrypoint /bin/bash "rocm/vllm-ci:${BUILDKITE_COMMIT}" -ec '
+          if [ ! -d /vllm-workspace ]; then echo Missing directory: /vllm-workspace >&2; exit 1; fi
+          if [ ! -d /vllm-workspace/tests ]; then echo Missing directory: /vllm-workspace/tests >&2; exit 1; fi
+          if [ ! -d /vllm-workspace/src/vllm ]; then echo Missing directory: /vllm-workspace/src/vllm >&2; exit 1; fi
+          if [ ! -x /vllm-workspace/src/vllm/vllm-rs ]; then echo Missing executable: /vllm-workspace/src/vllm/vllm-rs >&2; exit 1; fi
+          command -v python3
+          command -v uv
+          command -v pytest
+          if ! command -v amd-smi >/dev/null 2>&1 && ! command -v rocminfo >/dev/null 2>&1; then
+            echo No ROCm CLI found in image >&2
+            exit 1
+          fi
+          python3 - <<PY
+        import torch, vllm
+        print(torch.__version__)
+        print(vllm.__version__)
+        PY
+          echo AMD image smoke OK
+        '
+    env:
+      DOCKER_BUILDKIT: "1"
+      VLLM_BAKE_FILE: "docker/docker-bake-rocm.hcl"
+      PYTORCH_ROCM_ARCH: "gfx90a;gfx942;gfx950"
+      IMAGE_TAG: "rocm/vllm-ci:$BUILDKITE_COMMIT"
+      REMOTE_VLLM: "1"
+      VLLM_BRANCH: "$BUILDKITE_COMMIT"
+    retry:
+      automatic:
+        - exit_status: -1  # Agent was lost
+          limit: 1
+        - exit_status: -10  # Agent was lost
+          limit: 1

.buildkite/hardware_tests/cpu.yaml

@@ -16,6 +16,7 @@ steps:
   - tests/kernels/test_onednn.py
   - 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
   commands:
     - |
       bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
@@ -24,20 +25,22 @@ steps:
       pytest -x -v -s tests/kernels/moe/test_cpu_quant_fused_moe.py
       pytest -x -v -s tests/kernels/test_onednn.py
       pytest -x -v -s tests/kernels/test_awq_int4_to_int8.py
-      pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py"
+      pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
+      pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py"
 
-- label: CPU-Compatibility Tests
-  depends_on: []
-  device: intel_cpu
-  no_plugin: true
-  source_file_dependencies:
-  - cmake/cpu_extension.cmake
-  - setup.py
-  - vllm/platforms/cpu.py
-  commands:
-    - |
-      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 20m "
-      bash .buildkite/scripts/hardware_ci/run-cpu-compatibility-test.sh"
+# Note: SDE can't be downloaded from CI host because of AWS WAF
+# - label: CPU-Compatibility Tests
+#   depends_on: []
+#   device: intel_cpu
+#   no_plugin: true
+#   source_file_dependencies:
+#   - cmake/cpu_extension.cmake
+#   - setup.py
+#   - vllm/platforms/cpu.py
+#   commands:
+#     - |
+#       bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 20m "
+#       bash .buildkite/scripts/hardware_ci/run-cpu-compatibility-test.sh"
 
 - label: CPU-Language Generation and Pooling Model Tests
   depends_on: []
@@ -62,11 +65,16 @@ steps:
   source_file_dependencies:
   - vllm/v1/worker/cpu/
   - vllm/v1/worker/gpu/
+  - vllm/v1/sample/ops/topk_topp_triton.py
+  - vllm/v1/sample/ops/topk_topp_sampler.py
+  - tests/v1/sample/test_topk_topp_sampler.py
   commands:
     - |
       bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 45m "
       uv pip install git+https://github.com/triton-lang/triton-cpu.git@270e696d
-      VLLM_USE_V2_MODEL_RUNNER=1 pytest -x -v -s tests/models/language/generation/test_granite.py -m cpu_model"
+      VLLM_USE_V2_MODEL_RUNNER=1 pytest -x -v -s tests/models/language/generation/test_granite.py -m cpu_model
+      # TODO: move to CPU-Kernel Tests once triton-cpu has a pre-built wheel
+      pytest -x -v -s tests/v1/sample/test_topk_topp_sampler.py::TestTritonTopkTopp"
 
 - label: CPU-Quantization Model Tests
   depends_on: []
@@ -83,7 +91,7 @@ steps:
   - tests/quantization/test_cpu_wna16.py
   commands:
     - |
-      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
+      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 45m "
       pytest -x -v -s tests/quantization/test_compressed_tensors.py::test_compressed_tensors_w8a8_logprobs
       pytest -x -v -s tests/quantization/test_cpu_wna16.py"
       

.buildkite/image_build/image_build.yaml

@@ -6,14 +6,26 @@ steps:
     timeout_in_minutes: 600
     commands:
     - if [[ "$BUILDKITE_BRANCH" == "main" ]]; then .buildkite/image_build/image_build.sh $REGISTRY $REPO $BUILDKITE_COMMIT $BRANCH $IMAGE_TAG $IMAGE_TAG_LATEST; else .buildkite/image_build/image_build.sh $REGISTRY $REPO $BUILDKITE_COMMIT $BRANCH $IMAGE_TAG; fi
-    # Non-root smoke 1: the default (root) image must still be importable
+    retry:
+      automatic:
+        - exit_status: -1  # Agent was lost
+          limit: 2
+        - exit_status: -10  # Agent was lost
+          limit: 2
+
+  - label: ":docker: :smoking: Non-root smoke tests"
+    key: image-build-smoke-test
+    depends_on:
+      - image-build
+    commands:
+    # Smoke 1: the default (root) image must still be importable
     # under a non-root UID via `--user 2000:0`. Validates the `vllm` passwd
     # entry + group-0-writable /home/vllm + uv path cleanup from #31959.
     # Uses `import vllm` rather than `vllm serve --help` because the latter
     # instantiates `VllmConfig` which requires a GPU attached to the
     # container.
     - docker run --rm --user 2000:0 --entrypoint python3 "$IMAGE_TAG" -c "import vllm; print(vllm.__version__)"
-    # Non-root smoke 2: assert the non-root enabling invariants are baked
+    # Smoke 2: assert the non-root enabling invariants are baked
     # into the image. Runs as UID 2000:0 via a shell so we can verify
     # filesystem perms + passwd/group file state + wrapper presence without
     # triggering vLLM's GPU-requiring config-init path. The opt-in

.buildkite/intel_jobs/basic_correctness.yaml

@@ -0,0 +1,22 @@
+group: Basic Correctness
+depends_on:
+  - image-build-xpu
+steps:
+- label: XPU Sleep Mode
+  timeout_in_minutes: 30
+  device: intel_gpu
+  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/basic_correctness/test_cumem.py
+  commands:
+    - >-
+      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+      'cd tests &&
+      export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
+      pytest -v -s basic_correctness/test_mem.py::test_end_to_end'

.buildkite/intel_jobs/expert_parallelism_intel.yaml

@@ -0,0 +1,23 @@
+group: Expert Parallelism
+depends_on: 
+  - image-build-xpu
+steps:
+- label: EPLB Algorithm
+  key: eplb-algorithm
+  timeout_in_minutes: 45
+  device: intel_gpu
+  no_plugin: true
+  working_dir: "."
+  env:
+    REGISTRY: "public.ecr.aws/q9t5s3a7"
+    REPO: "vllm-ci-test-repo"
+    VLLM_TEST_DEVICE: "xpu"
+  source_file_dependencies:
+  - vllm/distributed/eplb
+  - tests/distributed/test_eplb_algo.py
+  - tests/distributed/test_eplb_utils.py
+  commands:
+    - >-
+      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+      'cd tests &&
+      pytest -v -s distributed/test_eplb_algo.py'

.buildkite/intel_jobs/misc_intel.yaml

@@ -38,7 +38,17 @@ steps:
     REPO: "vllm-ci-test-repo"
     VLLM_TEST_DEVICE: "xpu"
   source_file_dependencies:
-    - vllm/
+    - vllm/config/
+    - vllm/distributed/
+    - vllm/engine/
+    - vllm/inputs/
+    - vllm/logger.py
+    - vllm/model_executor/
+    - vllm/platforms/
+    - vllm/sampling_params.py
+    - vllm/transformers_utils/
+    - vllm/utils/
+    - vllm/v1/
     - tests/v1/sample
     - tests/v1/logits_processors
     - tests/v1/test_oracle.py
@@ -52,4 +62,126 @@ steps:
       pytest -v -s v1/logits_processors --ignore=v1/logits_processors/test_custom_online.py --ignore=v1/logits_processors/test_custom_offline.py &&
       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/test_outputs.py &&
+      pytest -v -s v1/sample/test_topk_topp_sampler.py'
+
+- label: XPU CPU Offload
+  timeout_in_minutes: 60
+  device: intel_gpu
+  no_plugin: true
+  working_dir: "."
+  env:
+    REGISTRY: "public.ecr.aws/q9t5s3a7"
+    REPO: "vllm-ci-test-repo"
+    VLLM_TEST_DEVICE: "xpu"
+  source_file_dependencies:
+    - vllm/
+    - vllm/v1/kv_offload/
+    - vllm/v1/kv_connector/
+    - tests/v1/kv_offload/
+    - tests/v1/kv_connector/unit/test_offloading_connector.py
+  commands:
+    - >-
+      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+      'export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
+      cd tests &&
+      pytest -v -s v1/kv_offload &&
+      pytest -v -s v1/kv_connector/unit/test_offloading_connector.py'
+
+- label: Regression
+  key: regression
+  timeout_in_minutes: 30
+  device: intel_gpu
+  no_plugin: true
+  working_dir: "."
+  env:
+    REGISTRY: "public.ecr.aws/q9t5s3a7"
+    REPO: "vllm-ci-test-repo"
+    VLLM_TEST_DEVICE: "xpu"
+  source_file_dependencies:
+  - vllm/config/
+  - vllm/distributed/
+  - vllm/engine/
+  - vllm/inputs/
+  - vllm/model_executor/
+  - vllm/multimodal/
+  - vllm/platforms/
+  - vllm/sampling_params.py
+  - vllm/transformers_utils/
+  - vllm/utils/
+  - vllm/v1/
+  - tests/test_regression
+  commands:
+  - >-
+      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+      'pip install modelscope &&
+      cd tests &&
+      pytest -v -s test_regression.py'
+
+- label: Metrics, Tracing (2 GPUs)
+  key: metrics-tracing-2-gpus
+  timeout_in_minutes: 30
+  num_devices: 2
+  device: intel_gpu
+  no_plugin: true
+  working_dir: "."
+  env:
+    REGISTRY: "public.ecr.aws/q9t5s3a7"
+    REPO: "vllm-ci-test-repo"
+    VLLM_TEST_DEVICE: "xpu"
+  source_file_dependencies:
+  - vllm/config/
+  - vllm/distributed/
+  - vllm/engine/
+  - vllm/inputs/
+  - vllm/model_executor/
+  - vllm/multimodal/
+  - vllm/platforms/
+  - vllm/sampling_params.py
+  - vllm/tracing/
+  - vllm/transformers_utils/
+  - vllm/utils/
+  - vllm/v1/
+  - tests/v1/tracing
+  commands:
+  - >-
+    bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+    'pip install opentelemetry-sdk\>=1.26.0 opentelemetry-api\>=1.26.0 opentelemetry-exporter-otlp\>=1.26.0 opentelemetry-semantic-conventions-ai\>=0.4.1 &&
+    cd tests &&
+    pytest -v -s v1/tracing'
+
+- label: Async Engine, Inputs, Utils, Worker
+  key: async-engine-inputs-utils-worker
+  timeout_in_minutes: 30
+  device: intel_gpu
+  no_plugin: true
+  working_dir: "."
+  env:
+    REGISTRY: "public.ecr.aws/q9t5s3a7"
+    REPO: "vllm-ci-test-repo"
+    VLLM_TEST_DEVICE: "xpu"
+  source_file_dependencies:
+  - vllm/assets/
+  - vllm/config/
+  - vllm/distributed/
+  - vllm/engine/
+  - vllm/inputs/
+  - vllm/model_executor/
+  - vllm/multimodal/
+  - vllm/platforms/
+  - vllm/sampling_params.py
+  - vllm/tokenizers/
+  - vllm/transformers_utils/
+  - vllm/utils/
+  - vllm/v1/
+  - tests/detokenizer
+  - tests/multimodal
+  - tests/utils_
+  commands:
+  - >-
+    bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+    'cd tests &&
+    pip install av &&
+    pytest -v -s detokenizer &&
+    pytest -v -s -m "not cpu_test" ./multimodal &&
+    pytest -v -s utils_ --ignore=utils_/test_mem_utils.py'

.buildkite/intel_jobs/models_multimodal_intel.yaml

@@ -0,0 +1,111 @@
+group: Models - Multimodal
+depends_on: 
+  - image-build-xpu
+steps:
+- label: "Multi-Modal Models (Standard) 1: qwen2"
+  key: multi-modal-models-standard-1-qwen2
+  timeout_in_minutes: 45
+  device: intel_gpu
+  no_plugin: true
+  working_dir: "."
+  env:
+    REGISTRY: "public.ecr.aws/q9t5s3a7"
+    REPO: "vllm-ci-test-repo"
+    VLLM_TEST_DEVICE: "xpu"
+  source_file_dependencies:
+  - vllm/
+  - tests/models/multimodal
+  commands:
+    - >-
+      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+      'pip install av git+https://github.com/TIGER-AI-Lab/Mantis.git &&
+      cd tests &&
+      pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen2" &&
+      pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model'
+
+- label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
+  key: multi-modal-models-standard-2-qwen3-gemma
+  timeout_in_minutes: 45
+  device: intel_gpu
+  no_plugin: true
+  working_dir: "."
+  env:
+    REGISTRY: "public.ecr.aws/q9t5s3a7"
+    REPO: "vllm-ci-test-repo"
+    VLLM_TEST_DEVICE: "xpu"
+  source_file_dependencies:
+  - vllm/
+  - tests/models/multimodal
+  commands:
+    - >-
+      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+      'pip install git+https://github.com/TIGER-AI-Lab/Mantis.git &&
+      cd tests &&
+      pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model'
+
+- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
+  key: multi-modal-models-standard-3-llava-qwen2-vl
+  timeout_in_minutes: 45
+  device: intel_gpu
+  no_plugin: true
+  working_dir: "."
+  env:
+    REGISTRY: "public.ecr.aws/q9t5s3a7"
+    REPO: "vllm-ci-test-repo"
+    VLLM_TEST_DEVICE: "xpu"
+  source_file_dependencies:
+  - vllm/
+  - tests/models/multimodal
+  commands:
+    - >-
+      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+      'pip install git+https://github.com/TIGER-AI-Lab/Mantis.git &&
+      cd tests &&
+      pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "not qwen2 and not qwen3 and not gemma" &&
+      pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model'
+
+- label: "Multi-Modal Models (Standard) 4: other + whisper"
+  key: multi-modal-models-standard-4-other-whisper
+  timeout_in_minutes: 45
+  device: intel_gpu
+  no_plugin: true
+  working_dir: "."
+  env:
+    REGISTRY: "public.ecr.aws/q9t5s3a7"
+    REPO: "vllm-ci-test-repo"
+    VLLM_TEST_DEVICE: "xpu"
+  source_file_dependencies:
+  - vllm/
+  - tests/models/multimodal
+  commands:
+    - >-
+      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+      'pip install av git+https://github.com/TIGER-AI-Lab/Mantis.git &&
+      cd tests &&
+      pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py  --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/processing'
+
+- label: Multi-Modal Processor # 44min
+  key: multi-modal-processor
+  timeout_in_minutes: 45
+  device: intel_gpu
+  no_plugin: true
+  working_dir: "."
+  env:
+    REGISTRY: "public.ecr.aws/q9t5s3a7"
+    REPO: "vllm-ci-test-repo"
+    VLLM_TEST_DEVICE: "xpu"
+  source_file_dependencies:
+  - vllm/
+  - tests/models/multimodal
+  - tests/models/registry.py
+  commands:
+    - >-
+      bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+      'pip install av matplotlib ftfy git+https://github.com/TIGER-AI-Lab/Mantis.git &&
+      pip install open-clip-torch --no-deps &&
+      cd tests &&
+      pytest -v -s models/multimodal/processing/test_tensor_schema.py
+      --deselect "tests/models/multimodal/processing/test_tensor_schema.py::test_model_tensor_schema[mistralai/Mistral-Large-3-675B-Instruct-2512-NVFP4]"
+      --deselect "tests/models/multimodal/processing/test_tensor_schema.py::test_model_tensor_schema[Qwen/Qwen2.5-Omni-7B-AWQ]"
+      --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB'
+  parallelism: 4

.buildkite/intel_jobs/test-intel.yaml

@@ -40,7 +40,9 @@ steps:
         python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --block-size 64 --enforce-eager --max-model-len 8192 &&
         python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 &&
         python3 examples/basic/offline_inference/generate.py --model ibm-research/PowerMoE-3b --block-size 64 --enforce-eager -tp 2 --enable-expert-parallel &&
-        python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --max-model-len 8192
+        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 V1 test"
     depends_on:
@@ -83,5 +85,5 @@ steps:
         bash .buildkite/scripts/hardware_ci/run-intel-test.sh
         'pip install av &&
         cd tests &&
-        pytest -v -s entrypoints/openai/chat_completion/test_audio_in_video.py &&
+        pytest -v -s entrypoints/multimodal/openai/chat_completion/test_audio_in_video.py &&
         pytest -v -s benchmarks/test_serve_cli.py'

.buildkite/lm-eval-harness/configs/NVIDIA-Nemotron-3-Nano-30B-A3B-FP8.yaml

@@ -6,9 +6,7 @@ tasks:
     value: 0.7142
   - name: "exact_match,flexible-extract"
     value: 0.4579
-env_vars:
-  VLLM_USE_FLASHINFER_MOE_FP8: "1"
-  VLLM_FLASHINFER_MOE_BACKEND: "throughput"
+moe_backend: "flashinfer_cutlass"
 limit: 1319
 num_fewshot: 5
 max_model_len: 262144

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

@@ -68,6 +68,10 @@ def launch_lm_eval(eval_config, tp_size):
     if current_platform.is_rocm() and "Nemotron-3" in eval_config["model_name"]:
         model_args += "attention_backend=TRITON_ATTN"
 
+    moe_backend = eval_config.get("moe_backend", None)
+    if moe_backend is not None:
+        model_args += f"moe_backend={moe_backend},"
+
     env_vars = eval_config.get("env_vars", None)
     with scoped_env_vars(env_vars):
         results = lm_eval.simple_evaluate(

.buildkite/performance-benchmarks/tests/latency-tests-arm64-cpu.json

@@ -2,7 +2,6 @@
     {
         "test_name": "latency_llama8B_tp1",
         "environment_variables": {
-            "VLLM_RPC_TIMEOUT": 100000,
             "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
             "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
             "VLLM_CPU_KVCACHE_SPACE": 40

.buildkite/performance-benchmarks/tests/latency-tests-cpu.json

@@ -2,7 +2,6 @@
     {
         "test_name": "latency_llama8B_tp2",
         "environment_variables": {
-            "VLLM_RPC_TIMEOUT": 100000,
 	    "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
 	    "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
 	    "VLLM_CPU_SGL_KERNEL": 1,

.buildkite/performance-benchmarks/tests/serving-tests-arm64-cpu.json

@@ -13,7 +13,6 @@
       200
     ],
     "server_environment_variables": {
-      "VLLM_RPC_TIMEOUT": 100000,
       "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
       "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
       "VLLM_CPU_SGL_KERNEL": 1,

.buildkite/performance-benchmarks/tests/serving-tests-cpu-asr.json

@@ -5,7 +5,6 @@
     ],
     "max_concurrency_list": [12, 16, 24, 32, 64, 128, 200],
     "server_environment_variables": {
-      "VLLM_RPC_TIMEOUT": 100000,
       "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120
     },
     "server_parameters": {

.buildkite/performance-benchmarks/tests/serving-tests-cpu-embed.json

@@ -9,7 +9,6 @@
       128
     ],
     "server_environment_variables": {
-      "VLLM_RPC_TIMEOUT": 100000,
       "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
       "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
       "VLLM_CPU_SGL_KERNEL": 1,

.buildkite/performance-benchmarks/tests/serving-tests-cpu-text.json

@@ -5,7 +5,6 @@
     ],
     "max_concurrency_list": [12, 16, 24, 32, 64, 128, 200],
     "server_environment_variables": {
-      "VLLM_RPC_TIMEOUT": 100000,
       "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
       "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
       "VLLM_CPU_SGL_KERNEL": 1,

.buildkite/performance-benchmarks/tests/serving-tests-cpu.json

@@ -5,7 +5,6 @@
     ],
     "max_concurrency_list": [12, 16, 24, 32, 64, 128, 200],
     "server_environment_variables": {
-      "VLLM_RPC_TIMEOUT": 100000,
       "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
       "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
       "VLLM_CPU_SGL_KERNEL": 1,

.buildkite/performance-benchmarks/tests/throughput-tests-arm64-cpu.json

@@ -2,7 +2,6 @@
     {
         "test_name": "throughput_llama8B_tp1",
         "environment_variables": {
-            "VLLM_RPC_TIMEOUT": 100000,
             "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
             "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
             "VLLM_CPU_KVCACHE_SPACE": 40

.buildkite/performance-benchmarks/tests/throughput-tests-cpu.json

@@ -2,7 +2,6 @@
     {
         "test_name": "throughput_llama8B_tp2",
         "environment_variables": {
-            "VLLM_RPC_TIMEOUT": 100000,
 	    "VLLM_ALLOW_LONG_MAX_MODEL_LEN": 1,
 	    "VLLM_ENGINE_ITERATION_TIMEOUT_S": 120,
 	    "VLLM_CPU_SGL_KERNEL": 1,

.buildkite/release-pipeline.yaml

@@ -1,12 +1,25 @@
 # CUDA architecture lists — following PyTorch RELEASE.md
 # (https://github.com/pytorch/pytorch/blob/main/RELEASE.md)
 # SM86 included for broader Ampere coverage; SM89 for marlin fp8 support
+# These requested arches are filtered by CMake's CUDA_SUPPORTED_ARCHS before
+# per-kernel arch selection. Do not add +PTX here: top-level +PTX is stripped
+# during that filtering, so kernels that need PTX must request it locally.
 env:
-  CUDA_ARCH_X86: "7.5 8.0 8.6 8.9 9.0 10.0 12.0+PTX"
-  # aarch64 only architectures: 8.7 for Orin, 11.0 for Thor (since CUDA 13)
-  CUDA_ARCH_AARCH64: "8.0 8.7 8.9 9.0 10.0 11.0 12.0+PTX"
+  # for CUDA >=13, sm_100+ targets have family specifiers (see CMakeLists.txt)
+  # so targets like 10.3 and 12.1 are automatically supported with this list
+  CUDA_ARCH_X86: "7.5 8.0 8.6 8.9 9.0 10.0 12.0"
+  # aarch64-only targets: Orin (8.7), Thor (11.0, CUDA 13+)
+  CUDA_ARCH_AARCH64: "8.0 8.7 8.9 9.0 10.0 11.0 12.0"
+
+  # for CUDA <13, we need to specify all needed targets
+  # some targets (10.3, 12.1) are skipped to limit the wheel size (< 500MB)
+  # please use CUDA 13 wheels or compile yourself on these new devices
   CUDA_ARCH_X86_CU129: "7.5 8.0 8.6 8.9 9.0 10.0 12.0"
   CUDA_ARCH_AARCH64_CU129: "8.0 8.7 8.9 9.0 10.0 12.0"
+  
+  # pre-built mooncake wheels
+  # the manylinux_2_35 wheel has compatibility issue on Ubuntu 24.04
+  # so we use different wheels for the time being
   MOONCAKE_WHEEL_AARCH64_2_35: "https://vllm-wheels.s3.amazonaws.com/mooncake/mooncake_transfer_engine-0.3.10.post2-0da9dfea3-cp312-cp312-manylinux_2_35_aarch64.whl"
   MOONCAKE_WHEEL_AARCH64_2_39: "https://vllm-wheels.s3.amazonaws.com/mooncake/mooncake_transfer_engine-0.3.10.post2-0da9dfea3-cp312-cp312-manylinux_2_39_aarch64.whl"
   MOONCAKE_WHEEL_X86_64: "https://vllm-wheels.s3.amazonaws.com/mooncake/mooncake_transfer_engine-0.3.10.post2-0da9dfea3-cp312-cp312-manylinux_2_35_x86_64.whl"
@@ -846,7 +859,6 @@ steps:
         allow_failure: true
       - step: build-cpu-release-image-arm64
         allow_failure: true
-    if: build.env("NIGHTLY") != "1"
 
   - label: "Publish release images to DockerHub"
     depends_on:

.buildkite/scripts/ci-clean-log.sh

@@ -13,5 +13,8 @@ INPUT_FILE="$1"
 # Strip timestamps
 sed -i 's/^\[[0-9]\{4\}-[0-9]\{2\}-[0-9]\{2\}T[0-9]\{2\}:[0-9]\{2\}:[0-9]\{2\}Z\] //' "$INPUT_FILE"
 
+# Strip Buildkite inline timestamp markers (ESC _bk;t=<ms> BEL)
+sed -i 's/\x1B_bk;t=[0-9]*\x07//g' "$INPUT_FILE"
+
 # Strip colorization
 sed -i -r 's/\x1B\[[0-9;]*[mK]//g' "$INPUT_FILE"

.buildkite/scripts/ci-fetch-log.sh

@@ -1,55 +1,178 @@
 #!/bin/bash
-# Usage: ./ci-fetch-log.sh <buildkite_job_url> [output_file]
-#        ./ci-fetch-log.sh <build_number> <job_uuid> [output_file]
+# Fetch vLLM Buildkite CI logs (public; no login required).
 #
-# Downloads the raw log for a Buildkite job from the public, unauthenticated
-# /organizations/<org>/pipelines/<pipeline>/builds/<n>/jobs/<uuid>/download
-# endpoint, then strips ANSI/timestamps via ci-clean-log.sh.
+# Usage:
+#   ci-fetch-log.sh [--soft|--all] --pr [<PR>]  failed jobs in the PR's latest
+#                                               build (current branch if omitted)
+#   ci-fetch-log.sh [--soft|--all] <build_url>  failed jobs in that build
+#   ci-fetch-log.sh <job_url> [output]          one job; both #<job_uuid> and
+#                                               ?sid=<id> URL forms work
+#   ci-fetch-log.sh <build> <job_uuid> [output]
 #
-# Find <build_number> and <job_uuid> via:
-#   gh pr checks <PR> --repo vllm-project/vllm
-# Each failing row's URL is .../builds/<build_number>#<job_uuid>.
+# --soft also fetches soft-failed jobs; --all fetches every finished job.
+# Saves each log as ci-<build>-<job-name>.log (ANSI/timestamps stripped) and
+# prints "<file>\t<job name>" per job. [output] is single-job only; "-"
+# streams to stdout. Existing files are kept; CI_FETCH_LOG_FORCE=1 refetches.
 
 set -euo pipefail
 
 ORG="vllm"
 PIPELINE="ci"
+UA="vllm-ci-fetch-log"
+UUID_RE='[0-9a-f]{8}-[0-9a-f]{4}-[0-9a-f]{4}-[0-9a-f]{4}-[0-9a-f]{12}'
 
 usage() {
-    echo "Usage: $0 <buildkite_job_url> [output_file]"
-    echo "       $0 <build_number> <job_uuid> [output_file]"
+    sed -n '2,15p' "$0" | sed 's/^# \{0,1\}//'
     exit 1
 }
 
-if [ $# -lt 1 ]; then usage; fi
+die() {
+    echo "$1" >&2
+    exit 1
+}
 
-if [[ "$1" == https://* ]]; then
+BUILD="" JOB="" SID="" OUT=""
+SCOPE="failed"
+
+while :; do
+    case "${1:-}" in
+    --soft) SCOPE="soft" ;;
+    --all) SCOPE="all" ;;
+    *) break ;;
+    esac
+    shift
+done
+
+case "${1:-}" in
+--pr)
+    PR="${2:-}"
+    # gh pr checks exits non-zero when checks are failing; that is the
+    # expected case here.
+    URL=$(gh pr checks ${PR:+"$PR"} --repo vllm-project/vllm 2>/dev/null |
+        grep -oE "https://buildkite.com/${ORG}/${PIPELINE}/builds/[0-9]+" |
+        sort -t/ -k7 -n | tail -1 || true)
+    [ -n "$URL" ] || die "No Buildkite build found via: gh pr checks ${PR:-<current branch>}"
+    BUILD="${URL##*/}"
+    ;;
+https://*)
     BUILD=$(echo "$1" | sed -nE 's#.*/builds/([0-9]+).*#\1#p')
-    JOB=$(echo "$1" | grep -oE '[0-9a-f]{8}-[0-9a-f-]+' | head -n 1)
-    OUT="${2:-ci-${BUILD}-${JOB:0:8}.log}"
-else
-    if [ $# -lt 2 ]; then usage; fi
+    JOB=$(echo "$1" | grep -oE "#${UUID_RE}" | head -n 1 | cut -c2- || true)
+    SID=$(echo "$1" | grep -oE "[?&]sid=${UUID_RE}" | head -n 1 | sed 's/.*sid=//' || true)
+    OUT="${2:-}"
+    [ -n "$BUILD" ] || die "Could not parse build number from: $1"
+    ;;
+[0-9]*)
+    [ $# -ge 2 ] || usage
     BUILD="$1"
     JOB="$2"
-    OUT="${3:-ci-${BUILD}-${JOB:0:8}.log}"
-fi
-
-if [ -z "$BUILD" ] || [ -z "$JOB" ]; then
-    echo "Could not parse build number or job UUID from: $1" >&2
+    OUT="${3:-}"
+    ;;
+*)
     usage
-fi
+    ;;
+esac
 
 COOKIES=$(mktemp)
-trap 'rm -f "$COOKIES"' EXIT
+JOBS_TSV=$(mktemp)
+trap 'rm -f "$COOKIES" "$JOBS_TSV"' EXIT
 
-# Buildkite issues a session cookie on first hit; subsequent /download needs it.
-curl -fsSL -c "$COOKIES" -A "vllm-ci-fetch-log" \
+# Buildkite issues a session cookie on first hit; later requests need it.
+curl -fsSL -c "$COOKIES" -A "$UA" \
     "https://buildkite.com/${ORG}/${PIPELINE}/builds/${BUILD}" -o /dev/null
 
-curl -fsSL -b "$COOKIES" -A "vllm-ci-fetch-log" \
-    "https://buildkite.com/organizations/${ORG}/pipelines/${PIPELINE}/builds/${BUILD}/jobs/${JOB}/download" \
-    -o "$OUT"
+# The build's job list (id, step uuid, state, name) is served as JSON from
+# the user-facing /data/jobs endpoint. Flatten it to TSV for easy filtering:
+#   job_id  step_uuid  failed  soft_failed  finished  slug  name
+curl -fsSL -b "$COOKIES" -A "$UA" \
+    "https://buildkite.com/${ORG}/${PIPELINE}/builds/${BUILD}/data/jobs" |
+    python3 -c '
+import json, re, sys
 
-bash "$(dirname "$0")/ci-clean-log.sh" "$OUT"
+data = json.load(sys.stdin)
+if data.get("has_next_page"):
+    print("warning: job list is paginated; some jobs not shown", file=sys.stderr)
+for r in data["records"]:
+    if r.get("type") != "script":
+        continue
+    name = (r.get("name") or "").replace("\t", " ").replace("\n", " ")
+    slug = re.sub(r"[^a-z0-9]+", "-", name.lower()).strip("-")[:60]
+    print("\t".join([
+        r["id"],
+        r.get("step_uuid") or "",
+        str(r.get("passed") is False),
+        str(bool(r.get("soft_failed"))),
+        str(bool(r.get("finished_at"))),
+        slug,
+        name,
+    ]))
+' >"$JOBS_TSV" || die "Could not list jobs for build ${BUILD}"
 
-echo "$OUT"
+if [ -n "$SID" ] && [ -z "$JOB" ]; then
+    # The ?sid= in builds/<N>/list URLs is the *step* uuid, not the job uuid.
+    JOB=$(awk -F'\t' -v s="$SID" '$1 == s || $2 == s {print $1; exit}' "$JOBS_TSV")
+    [ -n "$JOB" ] || die "No job matching sid=${SID} in build ${BUILD}"
+fi
+
+fetch_job() { # <job_uuid> <output_file>
+    curl -fsSL -b "$COOKIES" -A "$UA" \
+        "https://buildkite.com/organizations/${ORG}/pipelines/${PIPELINE}/builds/${BUILD}/jobs/$1/download" \
+        -o "$2"
+    bash "$(dirname "$0")/ci-clean-log.sh" "$2"
+}
+
+if [ -n "$JOB" ]; then
+    # Single-job mode.
+    NAME=$(awk -F'\t' -v j="$JOB" '$1 == j {print $7; exit}' "$JOBS_TSV")
+    SLUG=$(awk -F'\t' -v j="$JOB" '$1 == j {print $6; exit}' "$JOBS_TSV")
+    [ -n "$OUT" ] || OUT="ci-${BUILD}-${SLUG:-${JOB:0:13}}.log"
+    if [ "$OUT" = "-" ]; then
+        TMP=$(mktemp)
+        fetch_job "$JOB" "$TMP"
+        cat "$TMP"
+        rm -f "$TMP"
+        exit 0
+    fi
+    if [ -e "$OUT" ] && [ -z "${CI_FETCH_LOG_FORCE:-}" ]; then
+        die "Refusing to overwrite existing ${OUT} (set CI_FETCH_LOG_FORCE=1 or pass an output path)."
+    fi
+    fetch_job "$JOB" "$OUT"
+    printf '%s\t%s\n' "$OUT" "${NAME:-$JOB}"
+    exit 0
+fi
+
+# Build-wide mode: fetch finished jobs matching $SCOPE.
+[ -z "$OUT" ] || die "[output_file] is only valid when fetching a single job."
+
+case "$SCOPE" in
+failed) FILTER='$3 == "True" && $4 == "False" && $5 == "True"' ;;
+soft) FILTER='$3 == "True" && $5 == "True"' ;;
+all) FILTER='$5 == "True"' ;;
+esac
+
+if [ "$SCOPE" = "failed" ]; then
+    SOFT=$(awk -F'\t' '$3 == "True" && $4 == "True"' "$JOBS_TSV" | wc -l)
+    [ "$SOFT" -eq 0 ] || echo "Skipping ${SOFT} soft-failed job(s); use --soft to include them." >&2
+fi
+
+FOUND=0
+EMITTED=" "
+while IFS=$'\t' read -r job_id _ _ _ _ slug name; do
+    FOUND=$((FOUND + 1))
+    out="ci-${BUILD}-${slug:-${job_id:0:13}}.log"
+    # Retries share a name with the original job; disambiguate by uuid.
+    case "$EMITTED" in
+    *" $out "*) out="ci-${BUILD}-${slug:-job}-${job_id:0:13}.log" ;;
+    esac
+    EMITTED="${EMITTED}${out} "
+    if [ -e "$out" ] && [ -z "${CI_FETCH_LOG_FORCE:-}" ]; then
+        echo "Keeping existing ${out} (set CI_FETCH_LOG_FORCE=1 to refetch)." >&2
+    elif ! fetch_job "$job_id" "$out"; then
+        echo "Failed to download log for job ${job_id} (${name})." >&2
+        continue
+    fi
+    printf '%s\t%s\n' "$out" "$name"
+done < <(awk -F'\t' "$FILTER" "$JOBS_TSV")
+
+if [ "$FOUND" -eq 0 ]; then
+    echo "No matching jobs in build ${BUILD} (scope: ${SCOPE})." >&2
+fi

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

@@ -28,8 +28,10 @@
 ###############################################################################
 set -o pipefail
 
-# Export Python path
-export PYTHONPATH=".."
+# 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.
+export PYTHONPATH="${PYTHONPATH:-..}"
 
 ###############################################################################
 # Helper Functions
@@ -52,6 +54,108 @@ cleanup_network() {
   fi
 }
 
+prepare_artifact_image() {
+  if [[ "${VLLM_CI_USE_ARTIFACTS:-0}" != "1" ]]; then
+    return 1
+  fi
+  if ! command -v buildkite-agent >/dev/null 2>&1; then
+    echo "buildkite-agent not found; cannot download ROCm wheel artifact"
+    return 1
+  fi
+
+  local artifact_glob="${VLLM_CI_ARTIFACT_GLOB:-artifacts/vllm-rocm-install/vllm-rocm-install.tar.gz}"
+  local archive=""
+  local metadata_file=""
+  local base_image="${VLLM_CI_BASE_IMAGE:-rocm/vllm-dev:ci_base}"
+  local artifact_image=""
+  local artifact_key=""
+  local base_digest=""
+  local wheel_dir=""
+  local context_dir=""
+  local workspace_dir=""
+
+  artifact_work_dir=$(mktemp -d -t vllm-rocm-artifact.XXXXXX)
+  wheel_dir="${artifact_work_dir}/wheels"
+  context_dir="${artifact_work_dir}/context"
+  workspace_dir="${context_dir}/workspace"
+  mkdir -p "${wheel_dir}" "${context_dir}/wheels" "${workspace_dir}"
+
+  echo "--- Downloading ROCm wheel artifact"
+  if ! buildkite-agent artifact download "${artifact_glob}" "${artifact_work_dir}"; then
+    echo "Failed to download ${artifact_glob}"
+    return 1
+  fi
+  buildkite-agent artifact download \
+    "artifacts/vllm-rocm-install/ci-base-image.txt" \
+    "${artifact_work_dir}" >/dev/null 2>&1 || true
+
+  archive=$(find "${artifact_work_dir}" -name "vllm-rocm-install.tar.gz" -type f | head -1)
+  if [[ -z "${archive}" || ! -f "${archive}" ]]; then
+    echo "ROCm wheel artifact archive was not found"
+    return 1
+  fi
+
+  metadata_file=$(find "${artifact_work_dir}" -name "ci-base-image.txt" -type f | head -1)
+  if [[ -n "${metadata_file}" && -s "${metadata_file}" ]]; then
+    base_image=$(tr -d '[:space:]' < "${metadata_file}")
+  fi
+
+  echo "--- Preparing local ROCm test image"
+  echo "Base image: ${base_image}"
+  docker pull "${base_image}" || return 1
+  base_digest=$(
+    docker image inspect \
+      --format='{{if .RepoDigests}}{{index .RepoDigests 0}}{{else}}{{.Id}}{{end}}' \
+      "${base_image}" 2>/dev/null || printf '%s' "${base_image}"
+  )
+
+  artifact_key=$(
+    {
+      printf 'base-image:%s\n' "${base_digest}"
+      sha256sum "${archive}"
+    } | sha256sum | cut -c1-24
+  )
+  artifact_image="rocm/vllm-ci-artifact:${artifact_key}"
+
+  if docker image inspect "${artifact_image}" >/dev/null 2>&1; then
+    echo "Using existing local ROCm artifact image: ${artifact_image}"
+    image_name="${artifact_image}"
+    return 0
+  fi
+
+  tar -xzf "${archive}" -C "${wheel_dir}" || return 1
+  if ! ls "${wheel_dir}"/*.whl >/dev/null 2>&1; then
+    echo "ROCm wheel artifact did not contain a wheel"
+    return 1
+  fi
+  if [[ ! -d "${wheel_dir}/tests" ]]; then
+    echo "ROCm wheel artifact did not contain the test workspace"
+    return 1
+  fi
+
+  cp "${wheel_dir}"/*.whl "${context_dir}/wheels/" || return 1
+  tar -C "${wheel_dir}" --exclude='*.whl' -cf - . \
+    | tar -C "${workspace_dir}" -xf - || return 1
+  cat > "${context_dir}/Dockerfile" <<'EOF'
+ARG BASE_IMAGE
+FROM ${BASE_IMAGE}
+COPY wheels/ /tmp/vllm-wheels/
+COPY workspace/ /vllm-workspace/
+RUN python3 -m pip install --no-deps --force-reinstall /tmp/vllm-wheels/*.whl \
+    && rm -rf /tmp/vllm-wheels
+WORKDIR /vllm-workspace
+EOF
+
+  echo "--- Building local ROCm test image"
+  docker build \
+    --pull=false \
+    --build-arg "BASE_IMAGE=${base_image}" \
+    -t "${artifact_image}" \
+    "${context_dir}" || return 1
+  image_name="${artifact_image}"
+  return 0
+}
+
 is_multi_node() {
   local cmds="$1"
   # Primary signal: NUM_NODES environment variable set by the pipeline
@@ -243,22 +347,30 @@ report_docker_usage
 
 # --- Pull test image ---
 echo "--- Pulling container"
-image_name="rocm/vllm-ci:${BUILDKITE_COMMIT}"
+image_name="${VLLM_CI_FALLBACK_IMAGE:-rocm/vllm-ci:${BUILDKITE_COMMIT:-local}}"
+artifact_work_dir=""
 container_name="rocm_${BUILDKITE_COMMIT}_$(tr -dc A-Za-z0-9 < /dev/urandom | head -c 10; echo)"
-docker pull "${image_name}"
 
 remove_docker_container() {
-  # docker run uses --rm, so the container is normally already gone when the
-  # EXIT trap runs. Cleanup is best-effort and must not affect the test result.
-  docker rm -f "${container_name}" >/dev/null 2>&1 || true
+  if docker container inspect "${container_name}" >/dev/null 2>&1; then
+    docker rm -f "${container_name}" || true
+  fi
+  if [[ "${VLLM_CI_REMOVE_TEST_IMAGE:-0}" == "1" ]]; then
+    docker image rm -f "${image_name}" || true
+  else
+    # Keep images by default so later jobs on the same AMD node can reuse layers.
+    echo "Keeping ROCm test image locally: ${image_name}"
+  fi
+  if [[ -n "${artifact_work_dir}" ]]; then
+    rm -rf "${artifact_work_dir}"
+  fi
 }
+trap remove_docker_container EXIT
 
-on_exit() {
-  local exit_code=$?
-  remove_docker_container
-  exit "$exit_code"
-}
-trap on_exit EXIT
+if ! prepare_artifact_image; then
+  echo "Using full ROCm CI image: ${image_name}"
+  docker pull "${image_name}" || exit 1
+fi
 
 # --- Prepare commands ---
 echo "--- Running container"
@@ -267,6 +379,14 @@ HF_CACHE="$(realpath ~)/huggingface"
 mkdir -p "${HF_CACHE}"
 HF_MOUNT="/root/.cache/huggingface"
 
+# Hugging Face Hub defaults to 10s request/download timeouts, while the ROCm
+# CI image currently raises downloads to 60s. AMD model-test jobs routinely
+# start from a cold or partially-populated shared cache, and the 60s read cap
+# has still timed out before pytest reached the vLLM behavior under test.
+# Keep the CI default explicit and overridable from the Buildkite environment.
+: "${HF_HUB_DOWNLOAD_TIMEOUT:=300}"
+: "${HF_HUB_ETAG_TIMEOUT:=60}"
+
 # ---- Command source selection ----
 # Prefer VLLM_TEST_COMMANDS (preserves all inner quoting intact).
 # Fall back to $* for backward compatibility, but warn that inner
@@ -306,7 +426,14 @@ fi
 
 echo "Final commands: $commands"
 
-MYPYTHONPATH=".."
+MYPYTHONPATH="/vllm-workspace"
+
+container_job_id="${BUILDKITE_JOB_ID:-${BUILDKITE_PARALLEL_JOB:-0}}"
+container_job_id="${container_job_id//[^A-Za-z0-9_.-]/_}"
+container_job_id_short="${container_job_id:0:8}"
+CONTAINER_TMPDIR="/tmp/vllm-${container_job_id_short}"
+CONTAINER_CACHE_ROOT="/tmp/vllm-buildkite-${container_job_id}/cache"
+CONTAINER_PREFLIGHT="mkdir -p \"\$TMPDIR\" \"\$TORCHINDUCTOR_CACHE_DIR\" \"\$TRITON_CACHE_DIR\" \"\$VLLM_CACHE_ROOT\" \"\$XDG_CACHE_HOME\" && python -c \"import encodings, importlib.metadata as im, importlib.util as iu; [im.version(d) for d in ('transformers', 'torch', 'ray', 'sympy', 'markupsafe', 'vllm')]; missing=[m for m in ('torch.utils.model_zoo', 'transformers.models.nomic_bert', 'ray.dag', 'sympy.physics', 'markupsafe._speedups') if iu.find_spec(m) is None]; assert not missing, missing\""
 
 # Verify GPU access
 render_gid=$(getent group render | cut -d: -f3)
@@ -383,6 +510,8 @@ else
     --group-add "$render_gid" \
     --rm \
     -e HF_TOKEN \
+    -e "HF_HUB_DOWNLOAD_TIMEOUT=${HF_HUB_DOWNLOAD_TIMEOUT}" \
+    -e "HF_HUB_ETAG_TIMEOUT=${HF_HUB_ETAG_TIMEOUT}" \
     -e AWS_ACCESS_KEY_ID \
     -e AWS_SECRET_ACCESS_KEY \
     -e BUILDKITE_PARALLEL_JOB \
@@ -390,10 +519,15 @@ else
     -v "${HF_CACHE}:${HF_MOUNT}" \
     -e "HF_HOME=${HF_MOUNT}" \
     -e "PYTHONPATH=${MYPYTHONPATH}" \
+    -e "TMPDIR=${CONTAINER_TMPDIR}/tmp" \
+    -e "TORCHINDUCTOR_CACHE_DIR=${CONTAINER_CACHE_ROOT}/torchinductor" \
+    -e "TRITON_CACHE_DIR=${CONTAINER_CACHE_ROOT}/triton" \
+    -e "VLLM_CACHE_ROOT=${CONTAINER_CACHE_ROOT}/vllm" \
+    -e "XDG_CACHE_HOME=${CONTAINER_CACHE_ROOT}/xdg" \
     -e "PYTORCH_ROCM_ARCH=" \
     --name "${container_name}" \
     "${image_name}" \
-    /bin/bash -c "${commands}"
+    /bin/bash -c "${CONTAINER_PREFLIGHT} && ${commands}"
 
   exit_code=$?
   handle_pytest_exit "$exit_code"

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

@@ -37,7 +37,8 @@ function cpu_tests() {
     pytest -x -v -s tests/kernels/test_onednn.py
     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_moe.py -k test_cpu_fused_moe_basic"
+    pytest -x -v -s tests/kernels/moe/test_moe.py -k test_cpu_fused_moe_basic
+    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

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

@@ -324,23 +324,6 @@ IMAGE="${IMAGE_TAG_XPU:-${image_name}}"
 
 echo "Using image: ${IMAGE}"
 
-if docker image inspect "${IMAGE}" >/dev/null 2>&1; then
-  echo "Image already exists locally, skipping pull"
-else
-  echo "Image not found locally, waiting for lock..."
-
-  flock /tmp/docker-pull.lock bash -c "
-    if docker image inspect '${IMAGE}' >/dev/null 2>&1; then
-      echo 'Image already pulled by another runner'
-    else
-      echo 'Pulling image...'
-      timeout 900 docker pull '${IMAGE}'
-    fi
-  "
-
-  echo "Pull step completed"
-fi
-
 remove_docker_container() {
   docker rm -f "${container_name}" || true
 }
@@ -357,9 +340,12 @@ export HF_TOKEN ZE_AFFINITY_MASK
 
 {
   flock 9
-  if ! docker image inspect "${IMAGE}" >/dev/null 2>&1; then
-    echo 'Image missing before container creation, pulling again...'
+  if docker image inspect "${IMAGE}" >/dev/null 2>&1; then
+    echo "Image already exists locally, skipping pull"
+  else
+    echo "Image not found locally, pulling image..."
     timeout 900 docker pull "${IMAGE}"
+    echo "Pull step completed"
   fi
 
   docker create \
@@ -372,6 +358,8 @@ export HF_TOKEN ZE_AFFINITY_MASK
     --entrypoint='' \
     -e HF_TOKEN \
     -e ZE_AFFINITY_MASK \
+    -e BUILDKITE_PARALLEL_JOB \
+    -e BUILDKITE_PARALLEL_JOB_COUNT \
     -e CMDS \
     --name "${container_name}" \
     "${IMAGE}" \

.buildkite/scripts/install-kv-connectors.sh

@@ -0,0 +1,39 @@
+#!/bin/bash
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+set -euo pipefail
+
+REQUIREMENTS_FILE="${KV_CONNECTORS_REQUIREMENTS:-/vllm-workspace/requirements/kv_connectors.txt}"
+
+uv pip install --system -r "${REQUIREMENTS_FILE}"
+
+NIXL_METADATA=$(python3 - <<'PY'
+import importlib.metadata as metadata
+
+import torch
+
+cuda_version = torch.version.cuda
+if cuda_version is None:
+    raise SystemExit("torch.version.cuda is not set")
+
+print(cuda_version.split(".", 1)[0], metadata.version("nixl"))
+PY
+)
+read -r CUDA_MAJOR NIXL_VERSION <<<"${NIXL_METADATA}"
+
+# nixl>=1.1.0 can install multiple CUDA wheel variants. Keep only the variant
+# matching this CI image so nixl_ep_cpp links against the available libcudart.
+uv pip uninstall --system nixl-cu12 nixl-cu13 2>/dev/null || true
+uv pip install --system --no-deps "nixl-cu${CUDA_MAJOR}==${NIXL_VERSION}"
+
+python3 - <<'PY'
+import importlib.metadata as metadata
+
+for package_name in ("nixl", "nixl-cu12", "nixl-cu13"):
+    try:
+        version = metadata.version(package_name)
+    except metadata.PackageNotFoundError:
+        version = "not installed"
+    print(f"{package_name}: {version}")
+PY

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

@@ -110,6 +110,36 @@ install_uv() {
     | env UV_INSTALL_DIR="$CARGO_HOME/bin" sh
 }
 
+setup_pyo3_python() {
+  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 \
+    "$python_version")"
+  export PYO3_PYTHON
+
+  local python_libdir
+  python_libdir="$("$PYO3_PYTHON" - <<'PY'
+import pathlib
+import sysconfig
+
+libdir = pathlib.Path(sysconfig.get_config_var("LIBDIR"))
+ldlibrary = sysconfig.get_config_var("LDLIBRARY")
+assert sysconfig.get_config_var("Py_ENABLE_SHARED") == 1
+assert ldlibrary
+assert (libdir / ldlibrary).exists(), libdir / ldlibrary
+print(libdir)
+PY
+)"
+
+  export LD_LIBRARY_PATH="${python_libdir}:${LD_LIBRARY_PATH:-}"
+  export LIBRARY_PATH="${python_libdir}:${LIBRARY_PATH:-}"
+}
+
 run_style_clippy() {
   install_cargo_sort
 
@@ -132,6 +162,7 @@ run_style_clippy() {
 
 run_tests() {
   install_uv
+  setup_pyo3_python
   install_cargo_nextest
 
   log_section "Running cargo nextest"

.buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_ep_eplb.sh

@@ -49,6 +49,7 @@ for BACK in "${BACKENDS[@]}"; do
     --data-parallel-size 2 \
     --enable-expert-parallel \
     --enable-eplb \
+    --eplb-config '{"use_async": false}' \
     --trust-remote-code \
     --max-model-len 2048 \
     --all2all-backend "$BACK" \

.buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_block_ep_eplb.sh

@@ -48,7 +48,7 @@ for BACK in "${BACKENDS[@]}"; do
     --enforce-eager \
     --enable-eplb \
     --all2all-backend "$BACK" \
-    --eplb-config '{"window_size":10, "step_interval":100, "num_redundant_experts":0, "log_balancedness":true}' \
+    --eplb-config '{"window_size":10, "step_interval":100, "num_redundant_experts":0, "log_balancedness":true, "use_async":false}' \
     --tensor-parallel-size "${TENSOR_PARALLEL_SIZE}" \
     --data-parallel-size "${DATA_PARALLEL_SIZE}" \
     --enable-expert-parallel \

.buildkite/scripts/tool_call/run-bfcl-eval.sh

@@ -70,7 +70,7 @@ echo "============================================"
 # ---- Install bfcl-eval if missing ----
 if ! python3 -c "import bfcl_eval" 2>/dev/null; then
     echo "Installing bfcl-eval..."
-    pip install "bfcl-eval>=2025.10.20.1,<2026"
+    uv pip install "bfcl-eval>=2025.10.20.1,<2026"
 fi
 
 # ---- Cleanup handler ----
@@ -100,7 +100,7 @@ SERVE_ARGS=(
     --tensor-parallel-size "$TP_SIZE"
     --max-model-len "$MAX_MODEL_LEN"
     --enforce-eager
-    --no-enable-prefix-caching
+    --enable-prefix-caching
 )
 
 # Append reasoning parser if specified

.buildkite/test-amd.yaml

@@ -88,16 +88,16 @@
 #     - Do NOT remove `VLLM_WORKER_MULTIPROC_METHOD=spawn` setting as ROCm requires this for certain models to function.            #
 #   * [Transformers Nightly Models]: Whisper needs `VLLM_WORKER_MULTIPROC_METHOD=spawn` to avoid deadlock.                          #
 #   * [Plugin Tests (2 GPUs)]:                                                                                                      #
-#     - {`pytest -v -s entrypoints/openai/test_oot_registration.py`}: It needs a clean process                                      #
-#     - {`pytest -v -s models/test_oot_registration.py`}: It needs a clean process                                                  #
-#     - {`pytest -v -s plugins/lora_resolvers`}: Unit tests for in-tree lora resolver plugins                                       #
+#     - {`pytest -v -s plugins_tests/test_oot_registration_online.py`}: It needs a clean process                                    #
+#     - {`pytest -v -s plugins_tests/test_oot_registration_offline.py`}: It needs a clean process                                   #
+#     - {`pytest -v -s plugins_tests/lora_resolvers`}: Unit tests for in-tree lora resolver plugins                                 #
 #   * [LoRA TP (Distributed)]:                                                                                                      #
 #     - There is some Tensor Parallelism related processing logic in LoRA that requires multi-GPU testing for validation.           #
 #     - {`pytest -v -s -x lora/test_gptoss_tp.py`}: Disabled for now because MXFP4 backend on non-cuda platform doesn't support     #
 #                                                   LoRA yet.                                                                       #
 #   * [Distributed Tests (NxGPUs)(HW-TAG)]: Don't test llama model here, it seems hf implementation is buggy. See:                  #
 #                                           https://github.com/vllm-project/vllm/pull/5689                                          #
-#   * [Distributed Tests (NxGPUs)(HW-TAG)]: Some old E2E tests were removed in https://github.com/vllm-project/vllm/pull/33293     #
+#   * [Distributed Tests (NxGPUs)(HW-TAG)]: Some old E2E tests were removed in https://github.com/vllm-project/vllm/pull/33293      #
 #                                           in favor of new tests in fusions_e2e. We avoid replicating the new jobs in              #
 #                                           this file as it's deprecated.                                                           #
 #                                                                                                                                   #
@@ -315,24 +315,6 @@ steps:
   - pytest -v -s distributed/test_pp_cudagraph.py
   - pytest -v -s distributed/test_pipeline_parallel.py
 
-#----------------------------------------------------------  mi250 · engine  -----------------------------------------------------------#
-
-- label: Engine # TBD
-  timeout_in_minutes: 180
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
-  agent_pool: mi250_1
-  optional: true
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - vllm/
-  - tests/engine
-  - tests/test_sequence
-  - tests/test_config
-  - tests/test_logger
-  - tests/test_vllm_port
-  commands:
-  - pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py
-
 #-----------------------------------------------------------  mi250 · evals  -----------------------------------------------------------#
 
 - label: Multi-Modal Accuracy Eval (Small Models) # TBD
@@ -416,7 +398,7 @@ steps:
   - tests/kernels/helion/
   - vllm/platforms/rocm.py
   commands:
-  - pip install helion==1.0.0
+  - pip install helion==1.1.0
   - pytest -v -s kernels/helion/
 
 - label: Kernels Mamba Test # TBD
@@ -449,29 +431,6 @@ steps:
   commands:
   - pytest -v -s lora --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --ignore=lora/test_chatglm3_tp.py --ignore=lora/test_llama_tp.py --ignore=lora/test_qwen3_with_multi_loras.py --ignore=lora/test_olmoe_tp.py --ignore=lora/test_deepseekv2_tp.py --ignore=lora/test_gptoss_tp.py --ignore=lora/test_qwen3moe_tp.py --ignore=lora/test_qwen35_densemodel_lora.py
 
-#------------------------------------------------------  mi250 · model_executor  -------------------------------------------------------#
-
-- label: Model Executor # TBD
-  timeout_in_minutes: 180
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
-  agent_pool: mi250_1
-  optional: true
-  torch_nightly: true
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - vllm/engine/arg_utils.py
-  - vllm/config/model.py
-  - vllm/model_executor
-  - tests/model_executor
-  - tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
-  - vllm/_aiter_ops.py
-  - vllm/platforms/rocm.py
-  commands:
-  - apt-get update && apt-get install -y curl libsodium23
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s model_executor -m '(not slow_test)'
-  - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py
-
 #------------------------------------------------------  mi250 · models / basic  -------------------------------------------------------#
 
 - label: Basic Models Test (Other CPU) # TBD
@@ -658,9 +617,9 @@ steps:
   - pytest -v -s plugins_tests/test_scheduler_plugins.py
   - pip install -e ./plugins/vllm_add_dummy_model
   - pytest -v -s distributed/test_distributed_oot.py
-  - pytest -v -s entrypoints/openai/chat_completion/test_oot_registration.py
-  - pytest -v -s models/test_oot_registration.py
-  - pytest -v -s plugins/lora_resolvers
+  - pytest -v -s plugins_tests/test_oot_registration_online.py # it needs a clean process
+  - pytest -v -s plugins_tests/test_oot_registration_offline.py # it needs a clean process
+  - pytest -v -s plugins_tests/lora_resolvers # unit tests for in-tree lora resolver plugins
 
 #------------------------------------------------------------  mi250 · v1  -------------------------------------------------------------#
 
@@ -845,7 +804,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
-  - ATTENTION_BACKEND=ROCM_ATTN bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
+  - ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
 
 - label: V1 e2e (2 GPUs) # TBD
   timeout_in_minutes: 180
@@ -871,7 +830,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
-  - ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
+  - ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
 #-------------------------------------------------------------  mi250 · misc  ------------------------------------------------------------#
 
@@ -931,10 +890,10 @@ steps:
   - vllm/
   - tests/basic_correctness/test_basic_correctness
   - tests/basic_correctness/test_cpu_offload
-  - tests/basic_correctness/test_cumem.py
+  - tests/basic_correctness/test_mem.py
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s basic_correctness/test_cumem.py
+  - pytest -v -s basic_correctness/test_mem.py
   - pytest -v -s basic_correctness/test_basic_correctness.py
   - pytest -v -s basic_correctness/test_cpu_offload.py
 
@@ -1228,87 +1187,20 @@ steps:
 
 #--------------------------------------------------------  mi300 · entrypoints  --------------------------------------------------------#
 
-- label: Entrypoints Integration (API Server 2) # TBD
-  timeout_in_minutes: 180
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
-  agent_pool: mi300_1
-  optional: true
-  fast_check: true
-  torch_nightly: true
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - vllm/
-  - tests/entrypoints/rpc
-  - tests/entrypoints/serve/instrumentator
-  - tests/tool_use
-  commands:
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/serve/instrumentator
-  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
-  - pytest -v -s tool_use
-
-- label: Entrypoints Integration (API Server openai - Part 1) # TBD
+- label: Entrypoints Unit Tests # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   fast_check: true
-  torch_nightly: true
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
-  - vllm/
-  - tests/entrypoints/openai
-  - tests/entrypoints/test_chat_utils
+  - vllm/entrypoints
+  - tests/entrypoints/unit_tests
+  - tests/entrypoints/weight_transfer
+  - vllm/platforms/rocm.py
   commands:
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
-
-- label: Entrypoints Integration (API Server openai - Part 2) # TBD
-  timeout_in_minutes: 180
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
-  agent_pool: mi300_1
-  optional: true
-  fast_check: true
-  torch_nightly: true
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - vllm/
-  - tests/entrypoints/openai
-  - tests/entrypoints/test_chat_utils
-  commands:
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
-  - pytest -v -s entrypoints/test_chat_utils.py
-
-- label: Entrypoints Integration (API Server openai - Part 3) # TBD
-  timeout_in_minutes: 180
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
-  agent_pool: mi300_1
-  optional: true
-  fast_check: true
-  torch_nightly: true
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - vllm/
-  - tests/entrypoints/openai
-  - tests/entrypoints/test_chat_utils
-  commands:
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
-
-- label: Entrypoints Integration (Speech to Text) # TBD
-  timeout_in_minutes: 180
-  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
-  agent_pool: mi300_1
-  fast_check: true
-  torch_nightly: true
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - vllm/
-  - tests/entrypoints/speech_to_text
-  commands:
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/speech_to_text
+  - pytest -v -s entrypoints/unit_tests
+  - pytest -v -s entrypoints/weight_transfer
 
 - label: Entrypoints Integration (LLM) # TBD
   timeout_in_minutes: 180
@@ -1321,26 +1213,81 @@ steps:
   source_file_dependencies:
   - vllm/
   - tests/entrypoints/llm
-  - tests/entrypoints/offline_mode
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/llm --ignore=entrypoints/llm/test_generate.py --ignore=entrypoints/llm/test_collective_rpc.py
-  - pytest -v -s entrypoints/llm/test_generate.py
-  - pytest -v -s entrypoints/offline_mode
+  - pytest -v -s entrypoints/llm --ignore=entrypoints/llm/test_generate.py --ignore=entrypoints/llm/test_collective_rpc.py --ignore=entrypoints/llm/offline_mode
+  - pytest -v -s entrypoints/llm/test_generate.py # it needs a clean process
+  - pytest -v -s entrypoints/llm/offline_mode # Needs to avoid interference with other tests
 
-- label: Entrypoints Integration (Pooling) # TBD
+- label: Entrypoints Integration (API Server) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
+  optional: true
   fast_check: true
   torch_nightly: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
-  - tests/entrypoints/pooling
+  - tests/entrypoints/serve
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/pooling
+  - pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
+  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
+
+- label: Entrypoints Integration (API Server OpenAI - Part 1) # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  fast_check: true
+  torch_nightly: true
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/
+  - tests/entrypoints/openai
+  - tests/entrypoints/test_chat_utils
+  commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s entrypoints/openai/ --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/correctness
+
+- label: Entrypoints Integration (API Server OpenAI - Part 2) # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  fast_check: true
+  torch_nightly: true
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/
+  - tests/entrypoints/openai
+  - tests/entrypoints/test_chat_utils
+  commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s entrypoints/openai/chat_completion
+  - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
+
+- label: Entrypoints Integration (API Server Generate) # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  optional: true
+  fast_check: true
+  torch_nightly: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/
+  - tests/tool_use
+  - tests/entrypoints/tool_parsers
+  - tests/entrypoints/anthropic
+  - tests/entrypoints/generate
+  commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s tool_use
+  - pytest -v -s entrypoints/tool_parsers
+  - pytest -v -s entrypoints/generate
+  - pytest -v -s entrypoints/anthropic
 
 - label: Entrypoints Integration (Responses API) # TBD
   timeout_in_minutes: 180
@@ -1356,29 +1303,57 @@ steps:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - pytest -v -s entrypoints/openai/responses
 
-- label: Entrypoints Unit Tests # TBD
+- label: Entrypoints Integration (Speech to Text) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   fast_check: true
+  torch_nightly: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
-  - vllm/entrypoints
-  - tests/entrypoints/
-  - vllm/platforms/rocm.py
+  - vllm/
+  - tests/entrypoints/speech_to_text
   commands:
-  - pytest -v -s entrypoints/openai/tool_parsers
-  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s entrypoints/speech_to_text
+
+- label: Entrypoints Integration (Multimodal)
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  fast_check: true
+  torch_nightly: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/
+  - tests/entrypoints/multimodal
+  commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s entrypoints/multimodal
+
+- label: Entrypoints Integration (Pooling) # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  fast_check: true
+  torch_nightly: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/
+  - tests/entrypoints/pooling
+  commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s entrypoints/pooling
 
 - label: OpenAI API correctness # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
+  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - csrc/
   - vllm/entrypoints/openai/
-  - vllm/model_executor/models/whisper.py
   - vllm/model_executor/layers/
   - vllm/v1/attention/backends/
   - vllm/v1/attention/selector.py
@@ -1432,6 +1407,7 @@ steps:
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
+  optional: true
   working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
   source_file_dependencies:
   - csrc/
@@ -1484,7 +1460,7 @@ steps:
   commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx-fp8-and-mixed.txt
 
-- label: DeepSeek V2-Lite Accuracy (4xH100-4xMI300) # TBD
+- label: DeepSeek V2-Lite Sync EPLB Accuracy (4xH100-4xMI300) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_4
@@ -1526,7 +1502,7 @@ steps:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large.txt --tp-size=4
 
-- label: Qwen3-30B-A3B-FP8-block Accuracy (4xH100-4xMI300) # TBD
+- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (4xH100-4xMI300) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_4
@@ -1738,6 +1714,29 @@ steps:
   - pytest -v -s -x lora/test_gptoss_tp.py
   - pytest -v -s -x lora/test_qwen35_densemodel_lora.py
 
+#------------------------------------------------------  mi300 · model_executor  -------------------------------------------------------#
+
+- label: Model Executor # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  optional: true
+  torch_nightly: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/engine/arg_utils.py
+  - vllm/config/model.py
+  - vllm/model_executor
+  - tests/model_executor
+  - tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
+  - vllm/_aiter_ops.py
+  - vllm/platforms/rocm.py
+  commands:
+  - apt-get update && apt-get install -y curl libsodium23
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s model_executor -m '(not slow_test)'
+  - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py
+
 #-----------------------------------------------------  mi300 · models / language  -----------------------------------------------------#
 
 - label: Language Models Test (Extended Pooling)  # TBD
@@ -2183,10 +2182,72 @@ steps:
   commands:
   - pytest -v -s v1/e2e/spec_decode -k "speculators or mtp_correctness"
 
+- label: Speculators Correctness # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/config/speculative.py
+  - vllm/engine/arg_utils.py
+  - vllm/transformers_utils/config.py
+  - vllm/transformers_utils/configs/speculators/
+  - vllm/v1/spec_decode/
+  - vllm/v1/worker/gpu/spec_decode/
+  - vllm/v1/worker/gpu_model_runner.py
+  - vllm/v1/sample/
+  - vllm/v1/attention/backends/
+  - vllm/v1/attention/selector.py
+  - vllm/model_executor/model_loader/
+  - vllm/model_executor/layers/
+  - vllm/model_executor/models/llama_eagle3.py
+  - vllm/model_executor/models/qwen3.py
+  - vllm/model_executor/models/qwen3_dflash.py
+  - vllm/model_executor/models/registry.py
+  - vllm/_aiter_ops.py
+  - tests/evals/gsm8k/
+  - tests/v1/spec_decode/test_speculators_correctness.py
+  - vllm/platforms/rocm.py
+  commands:
+  - export VLLM_ALLOW_INSECURE_SERIALIZATION=1
+  - pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test
+
+- label: Extract Hidden States Integration # TBD
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+  agent_pool: mi300_1
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/config/speculative.py
+  - vllm/distributed/kv_transfer/kv_connector/
+  - vllm/model_executor/layers/attention/
+  - vllm/model_executor/layers/mamba/
+  - vllm/model_executor/model_loader/
+  - vllm/model_executor/models/extract_hidden_states.py
+  - vllm/model_executor/models/llama.py
+  - vllm/model_executor/models/qwen3_5.py
+  - vllm/model_executor/models/qwen3_next.py
+  - vllm/model_executor/models/registry.py
+  - vllm/transformers_utils/configs/extract_hidden_states.py
+  - vllm/transformers_utils/configs/qwen3_5.py
+  - vllm/v1/attention/backends/
+  - vllm/v1/attention/selector.py
+  - vllm/v1/kv_cache_interface.py
+  - vllm/v1/spec_decode/extract_hidden_states.py
+  - vllm/v1/worker/gpu_model_runner.py
+  - vllm/_aiter_ops.py
+  - tests/v1/kv_connector/extract_hidden_states_integration/
+  - vllm/platforms/rocm.py
+  commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s v1/kv_connector/extract_hidden_states_integration
+
 - label: V1 attention (H100-MI300) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
+  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/config/attention.py
@@ -2345,7 +2406,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
-  - CROSS_LAYERS_BLOCKS=True ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
+  - CROSS_LAYERS_BLOCKS=True ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
 - label: Distributed DP Tests (4 GPUs) # TBD
   timeout_in_minutes: 180
@@ -2381,7 +2442,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
-  - ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
+  - ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
 - label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs) # TBD
   timeout_in_minutes: 180
@@ -2395,7 +2456,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
-  - DP_EP=1 ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
+  - DP_EP=1 ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
 - label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs) # TBD
   timeout_in_minutes: 180
@@ -2409,7 +2470,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
-  - HYBRID_SSM=1 ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
+  - HYBRID_SSM=1 ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
 - label: V1 e2e (4 GPUs) # TBD
   timeout_in_minutes: 180
@@ -2545,6 +2606,7 @@ steps:
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
   agent_pool: mi325_1
+  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -2554,7 +2616,7 @@ steps:
   - tests/test_logger
   - tests/test_vllm_port
   commands:
-  - pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py
+  - pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py test_jit_monitor.py
 
 #-----------------------------------------------------------  mi325 · evals  -----------------------------------------------------------#
 
@@ -2636,6 +2698,7 @@ steps:
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
   agent_pool: mi325_1
+  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -2735,7 +2798,7 @@ steps:
 
 #--------------------------------------------------------  mi355 · entrypoints  --------------------------------------------------------#
 
-- label: Entrypoints Integration (API Server 2) # TBD
+- label: Entrypoints Integration (API Server) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
@@ -2745,16 +2808,13 @@ steps:
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
-  - tests/entrypoints/rpc
-  - tests/entrypoints/serve/instrumentator
-  - tests/tool_use
+  - tests/entrypoints/serve
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/serve/instrumentator
-  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
-  - pytest -v -s tool_use
+  - pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
+  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
 
-- label: Entrypoints Integration (API Server openai - Part 1) # TBD
+- label: Entrypoints Integration (API Server OpenAI - Part 1) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
@@ -2768,9 +2828,9 @@ steps:
   - tests/entrypoints/test_chat_utils
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
+  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/correctness
 
-- label: Entrypoints Integration (API Server openai - Part 2) # TBD
+- label: Entrypoints Integration (API Server OpenAI - Part 2) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
@@ -2784,23 +2844,29 @@ steps:
   - tests/entrypoints/test_chat_utils
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s entrypoints/openai/chat_completion
   - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
-  - pytest -v -s entrypoints/test_chat_utils.py
 
-- label: Entrypoints Integration (API Server openai - Part 3) # TBD
+- label: Entrypoints Integration (API Server Generate) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
   fast_check: true
   torch_nightly: true
+  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
-  - tests/entrypoints/openai
-  - tests/entrypoints/test_chat_utils
+  - tests/tool_use
+  - tests/entrypoints/tool_parsers
+  - tests/entrypoints/anthropic
+  - tests/entrypoints/generate
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
+  - pytest -v -s tool_use
+  - pytest -v -s entrypoints/tool_parsers
+  - pytest -v -s entrypoints/generate
+  - pytest -v -s entrypoints/anthropic
 
 - label: Entrypoints Integration (Speech to Text) # TBD
   timeout_in_minutes: 180
@@ -2816,6 +2882,20 @@ steps:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - pytest -v -s entrypoints/speech_to_text
 
+- label: Entrypoints Integration (Multimodal)
+  timeout_in_minutes: 180
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi355]
+  agent_pool: mi355_1
+  fast_check: true
+  torch_nightly: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/
+  - tests/entrypoints/multimodal
+  commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s entrypoints/multimodal
+
 - label: Entrypoints Integration (Pooling) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
@@ -2866,7 +2946,6 @@ steps:
   - vllm/model_executor/models/qwen3_5_mtp.py
   - vllm/transformers_utils/configs/qwen3_5.py
   - vllm/transformers_utils/configs/qwen3_5_moe.py
-  - vllm/model_executor/models/qwen.py
   - vllm/model_executor/models/qwen2.py
   - vllm/model_executor/models/qwen3.py
   - vllm/model_executor/models/qwen3_next.py
@@ -2895,7 +2974,7 @@ steps:
   commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx-fp8-and-mixed.txt
 
-- label: Qwen3-30B-A3B-FP8-block Accuracy (B200-MI355) # TBD
+- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (B200-MI355) # TBD
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_2
@@ -2993,7 +3072,7 @@ steps:
   - vllm/_aiter_ops.py
   commands:
   - rocm-smi
-  - python3 examples/basic/offline_inference/chat.py
+  - python3 examples/basic/offline_inference/chat.py --attention-backend TRITON_ATTN
   - pytest -v -s tests/kernels/attention/test_attention_selector.py
 
 - label: Kernels Attention Test %N # TBD
@@ -3104,7 +3183,6 @@ steps:
   - vllm/model_executor/models/qwen3_5_mtp.py
   - vllm/transformers_utils/configs/qwen3_5.py
   - vllm/transformers_utils/configs/qwen3_5_moe.py
-  - vllm/model_executor/models/qwen.py
   - vllm/model_executor/models/qwen2.py
   - vllm/model_executor/models/qwen3.py
   - vllm/model_executor/models/qwen3_next.py
@@ -3357,7 +3435,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
-  - ATTENTION_BACKEND=ROCM_ATTN bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
+  - ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
 
 - label: Distributed NixlConnector PD accuracy (4 GPUs) # TBD
   timeout_in_minutes: 180
@@ -3372,7 +3450,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
-  - ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
+  - ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
 - label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs) # TBD
   timeout_in_minutes: 180
@@ -3387,7 +3465,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
-  - DP_EP=1 ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
+  - DP_EP=1 ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
 #------------------------------------------------------  mi355 · weight_loading  -------------------------------------------------------#
 

.buildkite/test_areas/attention.yaml

@@ -2,8 +2,8 @@ group: Attention
 depends_on: 
   - image-build
 steps:
-- label: V1 attention (H100)
-  key: v1-attention-h100
+- label: V1 attention (H100-MI300)
+  key: v1-attention-h100-mi300
   timeout_in_minutes: 30
   device: h100
   source_file_dependencies:
@@ -13,6 +13,20 @@ steps:
     - tests/v1/attention
   commands:
     - pytest -v -s v1/attention
+  mirror:
+    amd:
+      device: mi325_1
+      timeout_in_minutes: 70
+      depends_on:
+      - image-build-amd
+      source_file_dependencies:
+      - vllm/config/attention.py
+      - vllm/model_executor/layers/attention
+      - vllm/v1/attention
+      - tests/v1/attention
+      - vllm/_aiter_ops.py
+      - vllm/envs.py
+      - vllm/platforms/rocm.py
 
 - label: V1 attention (B200)
   key: v1-attention-b200

.buildkite/test_areas/basic_correctness.yaml

@@ -10,9 +10,9 @@ steps:
   - vllm/
   - tests/basic_correctness/test_basic_correctness
   - tests/basic_correctness/test_cpu_offload
-  - tests/basic_correctness/test_cumem.py
+  - tests/basic_correctness/test_mem.py
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s basic_correctness/test_cumem.py
+  - pytest -v -s basic_correctness/test_mem.py
   - pytest -v -s basic_correctness/test_basic_correctness.py
   - pytest -v -s basic_correctness/test_cpu_offload.py

.buildkite/test_areas/disaggregated.yaml

@@ -11,7 +11,7 @@ steps:
     - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
     - tests/v1/kv_connector/nixl_integration/
   commands:
-    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
+    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
     - bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 - label: Distributed FlashInfer NixlConnector PD accuracy (4 GPUs)
   key: distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus
@@ -22,7 +22,7 @@ steps:
     - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
     - tests/v1/kv_connector/nixl_integration/
   commands:
-    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
+    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
     - FLASHINFER=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
 - label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs)
@@ -34,7 +34,7 @@ steps:
     - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
     - tests/v1/kv_connector/nixl_integration/
   commands:
-    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
+    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
     - DP_EP=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
 - label: CrossLayer KV layout Distributed NixlConnector PD accuracy tests (4 GPUs)
@@ -46,7 +46,7 @@ steps:
     - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
     - tests/v1/kv_connector/nixl_integration/
   commands:
-    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
+    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
     - CROSS_LAYERS_BLOCKS=True bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
 - label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs)
@@ -58,9 +58,23 @@ steps:
     - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
     - tests/v1/kv_connector/nixl_integration/
   commands:
-    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
+    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
     - HYBRID_SSM=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
+- label: Hybrid SSM NixlConnector PD prefix cache test (2 GPUs)
+  key: hybrid-ssm-nixlconnector-pd-prefix-cache-2-gpus
+  timeout_in_minutes: 25
+  working_dir: "/vllm-workspace/tests"
+  num_devices: 2
+  source_file_dependencies:
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/v1/core/sched/
+    - vllm/v1/core/kv_cache_coordinator.py
+    - tests/v1/kv_connector/nixl_integration/
+  commands:
+    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+    - bash v1/kv_connector/nixl_integration/run_mamba_prefix_cache_test.sh
+
 - label: MultiConnector (Nixl+Offloading) PD accuracy (2 GPUs)
   key: multiconnector-nixl-offloading-pd-accuracy-2-gpus
   timeout_in_minutes: 30
@@ -73,7 +87,7 @@ steps:
     - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
     - tests/v1/kv_connector/nixl_integration/
   commands:
-    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
+    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
     - bash v1/kv_connector/nixl_integration/run_multi_connector_accuracy_test.sh
 
 - label: NixlConnector PD + Spec Decode acceptance (2 GPUs)
@@ -87,7 +101,7 @@ steps:
     - vllm/v1/worker/kv_connector_model_runner_mixin.py
     - tests/v1/kv_connector/nixl_integration/
   commands:
-    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
+    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
     - bash v1/kv_connector/nixl_integration/config_sweep_spec_decode_test.sh
 
 - label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs)
@@ -102,5 +116,5 @@ steps:
     - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
     - tests/v1/kv_connector/nixl_integration/
   commands:
-    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
-    - bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh
+    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
+    - bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh

.buildkite/test_areas/e2e_integration.yaml

@@ -2,8 +2,8 @@ group: E2E Integration
 depends_on: 
   - image-build
 steps:
-- label: DeepSeek V2-Lite Accuracy
-  key: deepseek-v2-lite-accuracy
+- label: DeepSeek V2-Lite Sync EPLB Accuracy
+  key: deepseek-v2-lite-sync-eplb-accuracy
   timeout_in_minutes: 60
   device: h100
   optional: true
@@ -12,8 +12,8 @@ steps:
   commands:
   - bash .buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_ep_eplb.sh 0.25 200 8010
 
-- label: Qwen3-30B-A3B-FP8-block Accuracy
-  key: qwen3-30b-a3b-fp8-block-accuracy
+- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy
+  key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy
   timeout_in_minutes: 60
   device: h100
   optional: true
@@ -22,8 +22,8 @@ steps:
   commands:
   - bash .buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_block_ep_eplb.sh 0.8 200 8020
 
-- label: Qwen3-30B-A3B-FP8-block Accuracy (B200)
-  key: qwen3-30b-a3b-fp8-block-accuracy-b200
+- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (B200)
+  key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-b200
   timeout_in_minutes: 60
   device: b200-k8s
   optional: true

.buildkite/test_areas/engine.yaml

@@ -26,6 +26,12 @@ steps:
   - tests/test_jit_monitor.py
   commands:
   - pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py test_jit_monitor.py
+  mirror:
+    amd:
+      device: mi325_1
+      timeout_in_minutes: 60
+      depends_on:
+      - image-build-amd
 
 - label: Engine (1 GPU)
   key: engine-1-gpu

.buildkite/test_areas/entrypoints.yaml

@@ -8,10 +8,11 @@ steps:
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/entrypoints
-  - tests/entrypoints/
+  - tests/entrypoints/unit_tests
+  - tests/entrypoints/weight_transfer
   commands:
-  - pytest -v -s entrypoints/openai/tool_parsers
-  - pytest -v -s entrypoints/ --ignore=entrypoints/llm --ignore=entrypoints/rpc --ignore=entrypoints/sleep --ignore=entrypoints/serve/instrumentator --ignore=entrypoints/openai --ignore=entrypoints/offline_mode --ignore=entrypoints/test_chat_utils.py  --ignore=entrypoints/pooling --ignore=entrypoints/speech_to_text
+  - pytest -v -s entrypoints/unit_tests
+  - pytest -v -s entrypoints/weight_transfer
 
 - label: Entrypoints Integration (LLM)
   key: entrypoints-integration-llm
@@ -20,20 +21,36 @@ steps:
   source_file_dependencies:
   - vllm/
   - tests/entrypoints/llm
-  - tests/entrypoints/offline_mode
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/llm --ignore=entrypoints/llm/test_generate.py --ignore=entrypoints/llm/test_collective_rpc.py
+  - pytest -v -s entrypoints/llm --ignore=entrypoints/llm/test_generate.py --ignore=entrypoints/llm/test_collective_rpc.py --ignore=entrypoints/llm/offline_mode
   - pytest -v -s entrypoints/llm/test_generate.py # it needs a clean process
-  - pytest -v -s entrypoints/offline_mode # Needs to avoid interference with other tests
+  - pytest -v -s entrypoints/llm/offline_mode # Needs to avoid interference with other tests
   mirror:
     amd:
       device: mi325_1
-      soft_fail: true
       depends_on:
       - image-build-amd
 
-- label: Entrypoints Integration (API Server openai - Part 1)
+- label: Entrypoints Integration (API Server)
+  key: entrypoints-integration-api-server
+  device: h200_35gb
+  timeout_in_minutes: 130
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/
+  - tests/entrypoints/serve
+  commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
+  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
+  mirror:
+    amd:
+      device: mi325_1
+      depends_on:
+      - image-build-amd
+
+- label: Entrypoints Integration (API Server OpenAI - Part 1)
   key: entrypoints-integration-api-server-openai-part-1
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
@@ -43,16 +60,15 @@ steps:
   - tests/entrypoints/test_chat_utils
   commands:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
+  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/correctness
   mirror:
     amd:
       device: mi325_1
-      soft_fail: true
       timeout_in_minutes: 80
       depends_on:
       - image-build-amd
 
-- label: Entrypoints Integration (API Server openai - Part 2)
+- label: Entrypoints Integration (API Server OpenAI - Part 2)
   key: entrypoints-integration-api-server-openai-part-2
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
@@ -61,57 +77,47 @@ steps:
   - tests/entrypoints/openai
   - tests/entrypoints/test_chat_utils
   commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s entrypoints/openai/chat_completion
   - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
-  - pytest -v -s entrypoints/test_chat_utils.py
   mirror:
     amd:
       device: mi325_1
-      soft_fail: true
-      timeout_in_minutes: 60
+      timeout_in_minutes: 80
       depends_on:
       - image-build-amd
 
-- label: Entrypoints Integration (API Server openai - Part 3)
-  key: entrypoints-integration-api-server-openai-part-3
+- label: Entrypoints Integration (API Server Generate)
+  key: entrypoints-integration-api-server-generate
   timeout_in_minutes: 50
-  device: h200_18gb
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
-  - tests/entrypoints/openai
-  - tests/entrypoints/test_chat_utils
+  - tests/tool_use
+  - tests/entrypoints/tool_parsers
+  - tests/entrypoints/anthropic
+  - tests/entrypoints/generate
   commands:
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
+  - pytest -v -s tool_use
+  - pytest -v -s entrypoints/tool_parsers
+  - pytest -v -s entrypoints/generate
+  - pytest -v -s entrypoints/anthropic
   mirror:
     amd:
       device: mi325_1
-      soft_fail: true
       timeout_in_minutes: 60
       depends_on:
       - image-build-amd
 
-- label: Entrypoints Integration (API Server 2)
-  device: h200_35gb
-  key: entrypoints-integration-api-server-2
-  timeout_in_minutes: 130
+- label: Entrypoints Integration (Responses API)
+  key: entrypoints-integration-responses-api
+  timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
-  - tests/entrypoints/rpc
-  - tests/entrypoints/serve/instrumentator
-  - tests/tool_use
+  - tests/entrypoints/openai/responses
   commands:
-  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  - pytest -v -s entrypoints/serve/instrumentator
-  - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
-  - pytest -v -s tool_use
-  mirror:
-    amd:
-      device: mi325_1
-      soft_fail: true
-      depends_on:
-      - image-build-amd
+  - pytest -v -s entrypoints/openai/responses
 
 - label: Entrypoints Integration (Speech to Text)
   device: h200_35gb
@@ -125,6 +131,18 @@ steps:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - pytest -v -s entrypoints/speech_to_text
 
+- label: Entrypoints Integration (Multimodal)
+  device: h200_35gb
+  key: entrypoints-integration-multimodal
+  timeout_in_minutes: 50
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/
+  - tests/entrypoints/multimodal
+  commands:
+  - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+  - pytest -v -s entrypoints/multimodal
+
 - label: Entrypoints Integration (Pooling)
   key: entrypoints-integration-pooling
   timeout_in_minutes: 50
@@ -136,16 +154,6 @@ steps:
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
   - pytest -v -s entrypoints/pooling
 
-- label: Entrypoints Integration (Responses API)
-  key: entrypoints-integration-responses-api
-  timeout_in_minutes: 50
-  working_dir: "/vllm-workspace/tests"
-  source_file_dependencies:
-  - vllm/
-  - tests/entrypoints/openai/responses
-  commands:
-  - pytest -v -s entrypoints/openai/responses
-
 - label: OpenAI API Correctness
   key: openai-api-correctness
   timeout_in_minutes: 30
@@ -153,6 +161,22 @@ steps:
   source_file_dependencies:
   - csrc/
   - vllm/entrypoints/openai/
-  - vllm/model_executor/models/whisper.py
   commands: # LMEval
   - pytest -s entrypoints/openai/correctness/
+  mirror:
+    amd:
+      device: mi325_1
+      depends_on:
+      - image-build-amd
+      source_file_dependencies:
+      - csrc/
+      - vllm/entrypoints/openai/
+      - vllm/model_executor/layers/
+      - vllm/v1/attention/backends/
+      - vllm/v1/attention/selector.py
+      - vllm/_aiter_ops.py
+      - vllm/platforms/rocm.py
+      - vllm/model_executor/model_loader/
+      commands:
+      - bash ../tools/install_torchcodec_rocm.sh || exit 1
+      - pytest -s entrypoints/openai/correctness/

.buildkite/test_areas/kernels.yaml

@@ -21,8 +21,9 @@ steps:
   - 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
+    - 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
@@ -74,6 +75,19 @@ steps:
     - pytest -v -s kernels/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
   parallelism: 2
 
+- label: Kernels Attention DiffKV Test (H100)
+  key: kernels-attention-diffkv-test-h100
+  timeout_in_minutes: 20
+  device: h100
+  num_devices: 1
+  source_file_dependencies:
+  - vllm/v1/attention/ops/triton_unified_attention_diffkv.py
+  - vllm/v1/attention/backends/triton_attn_diffkv.py
+  - vllm/v1/attention/backends/flash_attn_diffkv.py
+  - tests/kernels/attention/test_triton_unified_attention_diffkv.py
+  commands:
+    - pytest -v -s kernels/attention/test_triton_unified_attention_diffkv.py
+
 - label: Kernels Quantization Test %N
   key: kernels-quantization-test
   timeout_in_minutes: 90
@@ -223,7 +237,7 @@ steps:
   - vllm/utils/import_utils.py
   - tests/kernels/helion/
   commands:
-    - pip install helion==1.0.0
+    - pip install helion==1.1.0
     - pytest -v -s kernels/helion/
 
  
@@ -299,3 +313,4 @@ steps:
   - vllm/config
   commands:
     - pytest -v -s kernels/moe/test_moe_layer.py
+    - pytest -v -s kernels/moe/test_deepep_v2_moe.py

.buildkite/test_areas/lm_eval.yaml

@@ -12,6 +12,21 @@ steps:
   autorun_on_main: true
   commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-small.txt
+  mirror:
+    amd:
+      device: mi325_1
+      timeout_in_minutes: 55
+      depends_on:
+      - image-build-amd
+      source_file_dependencies:
+      - csrc/
+      - vllm/model_executor/layers/quantization
+      - vllm/model_executor/models/
+      - vllm/model_executor/model_loader/
+      - vllm/v1/attention/backends/
+      - vllm/v1/attention/selector.py
+      - vllm/_aiter_ops.py
+      - vllm/platforms/rocm.py
 
 # - label: LM Eval Large Models (4 GPUs)(A100)
 #   device: a100

.buildkite/test_areas/misc.yaml

@@ -86,7 +86,7 @@ steps:
     - tests/v1/metrics
     - tests/entrypoints/openai/correctness/test_lmeval.py
   commands:
-    - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
+    - bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
     - export VLLM_WORKER_MULTIPROC_METHOD=spawn
     # split the test to avoid interference
     - pytest -v -s -m 'not cpu_test' v1/core
@@ -138,11 +138,26 @@ steps:
     - vllm/v1/spec_decode/extract_hidden_states.py
     - vllm/model_executor/models/extract_hidden_states.py
     - vllm/transformers_utils/configs/extract_hidden_states.py
+    - vllm/distributed/kv_transfer/kv_connector/v1/example_hidden_states_connector.py
     - tests/v1/kv_connector/extract_hidden_states_integration
   commands:
     - export VLLM_WORKER_MULTIPROC_METHOD=spawn
     - pytest -v -s v1/kv_connector/extract_hidden_states_integration
 
+- label: Extract Hidden States Integration (2 GPUs)
+  key: extract-hidden-states-integration-2-gpus
+  timeout_in_minutes: 20
+  num_devices: 2
+  source_file_dependencies:
+    - vllm/v1/spec_decode/extract_hidden_states.py
+    - vllm/model_executor/models/extract_hidden_states.py
+    - vllm/transformers_utils/configs/extract_hidden_states.py
+    - vllm/distributed/kv_transfer/kv_connector/v1/example_hidden_states_connector.py
+    - tests/v1/kv_connector/extract_hidden_states_integration
+  commands:
+    - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+    - pytest -v -s -m 'distributed' v1/kv_connector/extract_hidden_states_integration
+
 - label: Regression
   key: regression
   timeout_in_minutes: 20
@@ -281,6 +296,7 @@ steps:
   - vllm/model_executor/layers/quantization/quark/
   - vllm/multimodal/
   - vllm/outputs.py
+  - vllm/parser/
   - vllm/platforms/
   - vllm/pooling_params.py
   - vllm/ray/
@@ -292,6 +308,7 @@ steps:
   - vllm/transformers_utils/
   - vllm/utils/
   - vllm/v1/
+  - tests/test_envs.py
   - tests/test_inputs.py
   - tests/test_outputs.py
   - tests/test_pooling_params.py
@@ -299,24 +316,25 @@ steps:
   - tests/multimodal
   - tests/renderers
   - tests/standalone_tests/lazy_imports.py
-  - tests/tokenizers_
   - tests/reasoning
   - tests/tool_parsers
+  - tests/tokenizers_
   - tests/parser
   - tests/transformers_utils
   - tests/config
   device: cpu-small
   commands:
   - python3 standalone_tests/lazy_imports.py
+  - pytest -v -s test_envs.py
   - pytest -v -s test_inputs.py
   - pytest -v -s test_outputs.py
   - pytest -v -s test_pooling_params.py
   - pytest -v -s test_ray_env.py
   - pytest -v -s -m 'cpu_test' multimodal
   - pytest -v -s renderers
-  - pytest -v -s tokenizers_
   - pytest -v -s reasoning --ignore=reasoning/test_seedoss_reasoning_parser.py --ignore=reasoning/test_glm4_moe_reasoning_parser.py
   - pytest -v -s tool_parsers
+  - pytest -v -s tokenizers_
   - pytest -v -s parser
   - pytest -v -s transformers_utils
   - pytest -v -s config

.buildkite/test_areas/model_executor.yaml

@@ -14,5 +14,12 @@ steps:
   commands:
     - apt-get update && apt-get install -y curl libsodium23
     - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-    - pytest -v -s model_executor -m '(not slow_test)'
-    - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py
+    # Dump tracebacks of all threads if a test hangs, so a wedged GPU/CUDA
+    # init surfaces a stack instead of silently stalling.
+    - export PYTHONFAULTHANDLER=1
+    # Per-test watchdog: a single hung test (e.g. stuck during engine/CUDA
+    # init) fails fast with a traceback instead of running until the global
+    # build timeout. The `thread` method also handles hangs inside C/CUDA
+    # calls that the signal method cannot interrupt.
+    - pytest -v -s model_executor -m '(not slow_test)' --timeout=900 --timeout-method=thread
+    - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py --timeout=900 --timeout-method=thread

.buildkite/test_areas/model_runner_v2.yaml

@@ -94,11 +94,13 @@ steps:
     - vllm/v1/worker/gpu_worker.py
     - tests/distributed/test_pipeline_parallel.py
     - tests/distributed/test_pp_cudagraph.py
+    - tests/v1/distributed/test_pp_dp_v2.py
   commands:
     - set -x
     - export VLLM_USE_V2_MODEL_RUNNER=1
     - pytest -v -s distributed/test_pipeline_parallel.py -k "not ray and not Jamba"
     - pytest -v -s distributed/test_pp_cudagraph.py -k "not ray"
+    - pytest -v -s v1/distributed/test_pp_dp_v2.py
 
 - label: Model Runner V2 Spec Decode
   device: h200_35gb

.buildkite/test_areas/models_language.yaml

@@ -51,6 +51,7 @@ steps:
     torch_nightly: {}
     amd:
       device: mi325_1
+      timeout_in_minutes: 90
       depends_on:
       - image-build-amd
       commands:

.buildkite/test_areas/models_multimodal.yaml

@@ -153,3 +153,12 @@ steps:
   - tests/models/multimodal/pooling
   commands:
     - pytest -v -s models/multimodal/pooling -m 'not core_model'
+  mirror:
+    amd:
+      device: mi325_1
+      timeout_in_minutes: 60
+      depends_on:
+      - image-build-amd
+      source_file_dependencies:
+      - vllm/
+      - tests/models/multimodal/pooling

.buildkite/test_areas/plugins.yaml

@@ -37,6 +37,20 @@ steps:
   - pytest -v -s plugins_tests/test_scheduler_plugins.py
   - pip install -e ./plugins/vllm_add_dummy_model
   - pytest -v -s distributed/test_distributed_oot.py
-  - pytest -v -s entrypoints/openai/chat_completion/test_oot_registration.py # it needs a clean process
-  - pytest -v -s models/test_oot_registration.py # it needs a clean process
-  - pytest -v -s plugins/lora_resolvers # unit tests for in-tree lora resolver plugins
+  - pytest -v -s plugins_tests/test_oot_registration_online.py # it needs a clean process
+  - pytest -v -s plugins_tests/test_oot_registration_offline.py # it needs a clean process
+  - pytest -v -s plugins_tests/lora_resolvers # unit tests for in-tree lora resolver plugins
+
+
+- label: GGUF Plugin
+  key: gguf-plugin
+  device: h200_18gb
+  timeout_in_minutes: 30
+  soft_fail: true
+  optional: true
+  source_file_dependencies:
+  - vllm/model_executor/layers/quantization
+  - tests/plugins_tests/test_gguf_plugin.py
+  commands:
+    - pip install "vllm-gguf-plugin >= 0.0.2"
+    - pytest -v -s plugins_tests/gguf

.buildkite/test_areas/quantization.yaml

@@ -21,6 +21,18 @@ steps:
   - uv pip install --system conch-triton-kernels
   - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py
 
+- label: Quantized Fusions
+  key: quantized-fusions
+  timeout_in_minutes: 30
+  source_file_dependencies:
+  - tests/fusion
+  - vllm/model_executor/layers/fusion
+  - vllm/model_executor/kernels/linear
+  - vllm/model_executor/layers/quantization/compressed_tensors
+  - vllm/model_executor/layers/quantization/modelopt.py
+  commands:
+    - pytest -v -s fusion/
+
 - label: Quantized MoE Test (B200)
   key: quantized-moe-test-b200
   timeout_in_minutes: 60

.buildkite/test_areas/rust_frontend.yaml

@@ -16,7 +16,7 @@ steps:
   - 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_with_tool_reasoning.py
+
   # - tests/entrypoints/openai/completion/test_prompt_validation.py
   - tests/entrypoints/openai/completion/test_shutdown.py
   # - tests/entrypoints/openai/test_return_token_ids.py
@@ -28,7 +28,7 @@ steps:
   - 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
   # - pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not invalid"
-  # - pytest -v -s entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py
+
   # - pytest -v -s entrypoints/openai/completion/test_prompt_validation.py -k "not prompt_embeds"
   - pytest -v -s entrypoints/openai/completion/test_shutdown.py -k "not engine_failure and not test_abort_timeout_exits_quickly"
   # - pytest -v -s entrypoints/openai/test_return_token_ids.py
@@ -45,19 +45,19 @@ steps:
   - vllm/entrypoints/serve/
   - vllm/v1/engine/
   - tests/utils.py
-  # - tests/entrypoints/rpc/test_collective_rpc.py
+  # - tests/entrypoints/serve/dev/rpc/test_collective_rpc.py
   - tests/entrypoints/serve/disagg/test_serving_tokens.py
   - tests/entrypoints/serve/instrumentator/test_basic.py
   - tests/entrypoints/serve/instrumentator/test_metrics.py
-  # - tests/entrypoints/serve/instrumentator/test_sleep.py
+  # - tests/entrypoints/serve/dev/test_sleep.py
   commands:
   - export VLLM_USE_RUST_FRONTEND=1
   - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-  # - pytest -v -s entrypoints/rpc/test_collective_rpc.py
+  # - pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
   - pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not show_version and not server_load"
   - pytest -v -s entrypoints/serve/disagg/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/instrumentator/test_sleep.py
+  # - pytest -v -s entrypoints/serve/dev/test_sleep.py
 
 - label: Rust Frontend Core Correctness
   timeout_in_minutes: 30

.buildkite/test_areas/spec_decode.yaml

@@ -37,6 +37,21 @@ steps:
     - tests/v1/e2e/spec_decode/
   commands:
     - pytest -v -s v1/e2e/spec_decode -k "speculators or mtp_correctness"
+  mirror:
+    amd:
+      device: mi325_1
+      timeout_in_minutes: 65
+      depends_on:
+      - image-build-amd
+      source_file_dependencies:
+      - vllm/v1/spec_decode/
+      - vllm/v1/worker/gpu/spec_decode/
+      - vllm/model_executor/model_loader/
+      - vllm/v1/sample/
+      - vllm/model_executor/layers/
+      - vllm/transformers_utils/configs/speculators/
+      - tests/v1/e2e/spec_decode/
+      - vllm/platforms/rocm.py
 
 - label: Spec Decode Speculators + MTP Nightly B200
   key: spec-decode-speculators-mtp-nightly-b200
@@ -61,6 +76,20 @@ steps:
     - tests/v1/e2e/spec_decode/
   commands:
     - pytest -v -s v1/e2e/spec_decode -k "ngram or suffix"
+  mirror:
+    amd:
+      device: mi325_1
+      timeout_in_minutes: 65
+      depends_on:
+      - image-build-amd
+      source_file_dependencies:
+      - vllm/v1/spec_decode/
+      - vllm/v1/worker/gpu/spec_decode/
+      - vllm/model_executor/model_loader/
+      - vllm/v1/sample/
+      - vllm/model_executor/layers/
+      - tests/v1/e2e/spec_decode/
+      - vllm/platforms/rocm.py
 
 - label: Spec Decode Draft Model
   key: spec-decode-draft-model
@@ -72,6 +101,20 @@ steps:
     - tests/v1/e2e/spec_decode/
   commands:
     - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
+  mirror:
+    amd:
+      device: mi325_1
+      timeout_in_minutes: 50
+      depends_on:
+      - image-build-amd
+      source_file_dependencies:
+      - vllm/v1/spec_decode/
+      - vllm/v1/worker/gpu/spec_decode/
+      - vllm/model_executor/model_loader/
+      - vllm/v1/sample/
+      - vllm/model_executor/layers/
+      - tests/v1/e2e/spec_decode/
+      - vllm/platforms/rocm.py
 
 - label: Spec Decode Draft Model Nightly B200
   key: spec-decode-draft-model-nightly-b200

.github/CODEOWNERS

@@ -23,9 +23,14 @@
 
 # Any change to the VllmConfig changes can have a large user-facing impact,
 # so spam a lot of people
-/vllm/config @WoosukKwon @youkaichao @robertgshaw2-redhat @mgoin @tlrmchlsmth @houseroad @hmellor @yewentao256 @ProExpertProg
+/vllm/config @WoosukKwon @youkaichao @robertgshaw2-redhat @mgoin @tlrmchlsmth @houseroad @yewentao256 @ProExpertProg
 /vllm/config/cache.py @heheda12345
 
+# Config utils
+/vllm/config/utils.py @hmellor
+/vllm/engine/arg_utils.py @hmellor
+/vllm/utils/argparse_utils.py
+
 # Entrypoints
 /vllm/entrypoints/anthropic @mgoin @DarkLight1337
 /vllm/entrypoints/cli @hmellor @mgoin @DarkLight1337 @russellb
@@ -34,12 +39,19 @@
 /vllm/entrypoints/speech_to_text/realtime @njhill
 /vllm/entrypoints/speech_to_text @NickLucche
 /vllm/entrypoints/pooling @noooop
-/vllm/entrypoints/sagemaker @DarkLight1337
+/vllm/entrypoints/serve/sagemaker @DarkLight1337
 /vllm/entrypoints/serve @njhill
 /vllm/entrypoints/*.py @njhill
 /vllm/entrypoints/chat_utils.py @DarkLight1337
+/vllm/entrypoints/offline_utils.py @DarkLight1337
 /vllm/entrypoints/llm.py @DarkLight1337
 
+# Rust Frontend
+/rust/ @BugenZhao @njhill
+/build_rust.sh @BugenZhao @njhill
+/rust-toolchain.toml @BugenZhao @njhill
+/.buildkite/test_areas/rust* @BugenZhao @njhill
+
 # Input/Output Processing
 /vllm/sampling_params.py @njhill @NickLucche
 /vllm/pooling_params.py @noooop @DarkLight1337
@@ -68,15 +80,17 @@
 /vllm/v1/worker/kv_connector_model_runner_mixin.py @orozery @NickLucche
 
 # Model runner V2
-/vllm/v1/worker/gpu @WoosukKwon @njhill
+/vllm/v1/worker/gpu @WoosukKwon @njhill @yewentao256
 /vllm/v1/worker/gpu/kv_connector.py @orozery
 
 # CI & building
-/.buildkite @Harry-Chen
-/docker/Dockerfile @Harry-Chen
+/.buildkite @Harry-Chen @khluu
+/docker/Dockerfile @Harry-Chen @khluu
+/pyproject.toml @khluu
+/setup.py @khluu
 
 # Test ownership
-/.buildkite/lm-eval-harness @mgoin 
+/.buildkite/lm-eval-harness @mgoin
 /tests/distributed/test_multi_node_assignment.py @youkaichao
 /tests/distributed/test_pipeline_parallel.py @youkaichao
 /tests/distributed/test_same_node.py @youkaichao

.github/dependabot.yml

@@ -21,7 +21,6 @@ updates:
       - dependency-name: "torchvision"
       - dependency-name: "xformers"
       - dependency-name: "lm-format-enforcer"
-      - dependency-name: "gguf"
       - dependency-name: "compressed-tensors"
       - dependency-name: "ray[cgraph]" # Ray Compiled Graph
       - dependency-name: "lm-eval"

.github/mergify.yml

@@ -21,6 +21,9 @@ pull_request_rules:
     - check-failure=pre-commit
     - -closed
     - -draft
+    - or:
+      - label=ready
+      - label=verified
   actions:
     comment:
       message: |
@@ -36,18 +39,6 @@ pull_request_rules:
 
         For future commits, `pre-commit` will run automatically on changed files before each commit.
 
-        > [!TIP]
-        > <details>
-        > <summary>Is <code>mypy</code> failing?</summary>
-        > <br/>
-        > <code>mypy</code> is run differently in CI. If the failure is related to this check, please use the following command to run it locally:
-        >
-        > ```bash
-        > # For mypy (substitute "3.10" with the failing version if needed)
-        > pre-commit run --hook-stage manual mypy-3.10
-        > ```
-        > </details>
-
 - name: comment-dco-failure
   description: Comment on PR when DCO check fails
   conditions:
@@ -153,12 +144,12 @@ pull_request_rules:
     - label != stale
     - or:
       - files~=^examples/.*mistral.*\.py
-      - files~=^tests/.*mistral.*\.py
-      - files~=^vllm/model_executor/models/.*mistral.*\.py
+      - files~=^tests/.*(?:mistral|voxtral|mixtral|pixtral).*\.py
+      - files~=^vllm/model_executor/models/.*(?:mistral|voxtral|mixtral|pixtral).*\.py
       - files~=^vllm/reasoning/.*mistral.*\.py
       - files~=^vllm/tool_parsers/.*mistral.*\.py
-      - files~=^vllm/transformers_utils/.*mistral.*\.py
-      - title~=(?i)Mistral
+      - files~=^vllm/transformers_utils/.*(?:mistral|voxtral|pixtral).*\.py
+      - title~=(?i)(?:mistral|ministral|voxtral|mixtral|pixtral)
   actions:
     label:
       add:

.github/workflows/add_label_automerge.yml

@@ -10,7 +10,7 @@ jobs:
         runs-on: ubuntu-latest
         steps:
             -   name: Add label
-                uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
+                uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
                 with:
                     script: |
                         github.rest.issues.addLabels({

.github/workflows/issue_autolabel.yml

@@ -14,7 +14,7 @@ jobs:
     steps:
       - name: Label issues based on keywords
         id: label-step
-        uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd  # v8.0.0
+        uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3  # v9.0.0
         with:
           script: |
             // Configuration: Add new labels and keywords here
@@ -315,7 +315,7 @@ jobs:
 
       - name: CC users for labeled issues
         if: steps.label-step.outputs.labels_added != '[]'
-        uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd  # v8.0.0
+        uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3  # v9.0.0
         with:
           script: |
             // Configuration: Map labels to GitHub users to CC
@@ -392,7 +392,7 @@ jobs:
 
       - name: Request missing ROCm info from issue author
         if: contains(steps.label-step.outputs.labels_added, 'rocm') && contains(toJSON(github.event.issue.labels.*.name), 'bug')
-        uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd  # v8.0.0
+        uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3  # v9.0.0
         with:
           script: |
             const body = (context.payload.issue.body || '').toLowerCase();

.github/workflows/new_pr_bot.yml

@@ -12,7 +12,7 @@ jobs:
     runs-on: ubuntu-latest
     steps:
       - name: Update PR description
-        uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
+        uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
         with:
           script: |
             const { owner, repo } = context.repo;
@@ -55,7 +55,7 @@ jobs:
     runs-on: ubuntu-latest
     steps:
       - name: Post welcome comment for first-time contributors
-        uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
+        uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
         with:
           script: |
             const { owner, repo } = context.repo;

.github/workflows/pre-commit.yml

@@ -20,7 +20,7 @@ jobs:
     runs-on: ubuntu-latest
     steps:
     - name: Check PR label and author merge count
-      uses: actions/github-script@ed597411d8f924073f98dfc5c65a23a2325f34cd # v8.0.0
+      uses: actions/github-script@3a2844b7e9c422d3c10d287c895573f7108da1b3 # v9.0.0
       with:
         script: |
           const { data: pr } = await github.rest.pulls.get({

.github/workflows/scripts/build.sh

@@ -9,7 +9,7 @@ PATH=${cuda_home}/bin:$PATH
 LD_LIBRARY_PATH=${cuda_home}/lib64:$LD_LIBRARY_PATH
 
 # Install requirements
-if [ "$(echo $2 | cut -d. -f1)" = "12" ]; then
+if [ "$(echo "$2" | cut -d. -f1)" = "12" ]; then
     sed -i 's/^nvidia-cutlass-dsl\[cu13\]>=/nvidia-cutlass-dsl>=/' requirements/cuda.txt
 fi
 $python_executable -m pip install -r requirements/build/cuda.txt -r requirements/cuda.txt
@@ -17,7 +17,10 @@ $python_executable -m pip install -r requirements/build/cuda.txt -r requirements
 # Limit the number of parallel jobs to avoid OOM
 export MAX_JOBS=1
 # Make sure release wheels are built for the following architectures
-export TORCH_CUDA_ARCH_LIST="7.5 8.0 8.6 8.9 9.0 10.0 12.0+PTX"
+# Do not add +PTX here: vLLM filters torch's top-level PTX flag when it
+# converts global gencode flags into per-kernel arch lists. If a specific
+# kernel needs PTX, add +PTX to that kernel's CMake arch list instead.
+export TORCH_CUDA_ARCH_LIST="7.5 8.0 8.6 8.9 9.0 10.0 12.0"
 
 bash tools/check_repo.sh
 

.github/workflows/stale.yml

@@ -15,7 +15,7 @@ jobs:
       actions: write
     runs-on: ubuntu-latest
     steps:
-      - uses: actions/stale@997185467fa4f803885201cee163a9f38240193d # v10.1.1
+      - uses: actions/stale@eb5cf3af3ac0a1aa4c9c45633dd1ae542a27a899 # v10.3.0
         with:
           # Increasing this value ensures that changes to this workflow
           # propagate to all issues and PRs in days rather than months

.gitignore

@@ -233,7 +233,7 @@ actionlint
 shellcheck*/
 
 # Ignore moe/marlin_moe gen code
-csrc/moe/marlin_moe_wna16/kernel_*
+csrc/libtorch_stable/moe/marlin_moe_wna16/kernel_*
 
 # Ignore ep_kernels_workspace folder
 ep_kernels_workspace/

.pre-commit-config.yaml

@@ -21,7 +21,7 @@ repos:
   rev: v21.1.2
   hooks:
   - id: clang-format
-    exclude: 'csrc/(moe/topk_softmax_kernels.cu|quantization/gguf/(ggml-common.h|dequantize.cuh|vecdotq.cuh|mmq.cuh|mmvq.cuh))|vllm/third_party/.*'
+    exclude: 'csrc/libtorch_stable/moe/topk_softmax_kernels.cu|vllm/third_party/.*'
     types_or: [c++, cuda]
     args: [--style=file, --verbose]
 - repo: https://github.com/DavidAnson/markdownlint-cli2
@@ -148,33 +148,27 @@ repos:
     language: python
     entry: python tools/pre_commit/generate_nightly_torch_test.py
     files: ^requirements/test/cuda\.(in|txt)$
-  - id: mypy-local
-    name: Run mypy locally for lowest supported Python version
-    entry: python tools/pre_commit/mypy.py 0 "3.10"
-    stages: [pre-commit] # Don't run in CI
+  - id: mypy-3.10 # TODO: Use https://github.com/pre-commit/mirrors-mypy when mypy setup is less awkward
+    name: Run mypy for Python 3.10
+    entry: python tools/pre_commit/mypy.py "3.10"
     <<: &mypy_common
       language: python
       types_or: [python, pyi]
       require_serial: true
-      additional_dependencies: ["mypy[faster-cache]==1.19.1", regex, types-cachetools, types-setuptools, types-PyYAML, types-requests, types-torch, pydantic]
-  - id: mypy-3.10 # TODO: Use https://github.com/pre-commit/mirrors-mypy when mypy setup is less awkward
-    name: Run mypy for Python 3.10
-    entry: python tools/pre_commit/mypy.py 1 "3.10"
-    <<: *mypy_common
-    stages: [manual] # Only run in CI
+      additional_dependencies: ["mypy==1.20.2", regex, types-cachetools, types-setuptools, types-PyYAML, types-requests, types-torch, pydantic]
   - id: mypy-3.11 # TODO: Use https://github.com/pre-commit/mirrors-mypy when mypy setup is less awkward
     name: Run mypy for Python 3.11
-    entry: python tools/pre_commit/mypy.py 1 "3.11"
+    entry: python tools/pre_commit/mypy.py "3.11"
     <<: *mypy_common
     stages: [manual] # Only run in CI
   - id: mypy-3.12 # TODO: Use https://github.com/pre-commit/mirrors-mypy when mypy setup is less awkward
     name: Run mypy for Python 3.12
-    entry: python tools/pre_commit/mypy.py 1 "3.12"
+    entry: python tools/pre_commit/mypy.py "3.12"
     <<: *mypy_common
     stages: [manual] # Only run in CI
   - id: mypy-3.13 # TODO: Use https://github.com/pre-commit/mirrors-mypy when mypy setup is less awkward
     name: Run mypy for Python 3.13
-    entry: python tools/pre_commit/mypy.py 1 "3.13"
+    entry: python tools/pre_commit/mypy.py "3.13"
     <<: *mypy_common
     stages: [manual] # Only run in CI
   - id: shellcheck

.readthedocs.yaml

@@ -9,8 +9,8 @@ build:
     python: "3.12"
   jobs:
     post_checkout:
-      - bash docs/pre_run_check.sh
       - git fetch origin main --unshallow --no-tags --filter=blob:none || true
+      - bash docs/pre_run_check.sh
     pre_create_environment:
       - pip install uv
     create_environment:

AGENTS.md

@@ -98,11 +98,33 @@ pre-commit run --all-files
 pre-commit run ruff-check --all-files
 
 # Run mypy as it is in CI:
-pre-commit run mypy-3.10 --all-files --hook-stage manual
+pre-commit run mypy-3.12 --all-files --hook-stage manual
 ```
 
 The line length limit for Python code is 88 characters. If you are not sure, use pre-commit to check.
 
+Use [Google-style docstrings](https://google.github.io/styleguide/pyguide.html#38-comments-and-docstrings) (`Args:`/`Returns:`/`Raises:` sections), not reStructuredText/Sphinx fields (`:param:`, `:return:`, `:rtype:`).
+
+### Coding style guidelines
+
+Follow these rules for all code changes in this repository:
+
+- Try to match existing code style.
+- Code should be self-documenting and self-explanatory.
+- Keep comments and docstrings minimal and concise.
+- Assume the reader is familiar with vLLM.
+
+### Diagnosing CI failures
+
+Buildkite logs are public; no login needed. Details: [docs/contributing/ci/failures.md](docs/contributing/ci/failures.md).
+
+```bash
+# All failed-job logs for a PR's latest build (current branch's PR if omitted):
+.buildkite/scripts/ci-fetch-log.sh --pr <PR>
+# Any Buildkite build or job URL also works:
+.buildkite/scripts/ci-fetch-log.sh "<buildkite_url>"
+```
+
 ### Commit messages
 
 Add attribution using commit trailers such as `Co-authored-by:` (other projects use `Assisted-by:` or `Generated-by:`). For example:

MANIFEST.in

@@ -4,6 +4,7 @@ include requirements/cuda.txt
 include requirements/rocm.txt
 include requirements/cpu.txt
 include CMakeLists.txt
+include tools/build_rust.py
 
 recursive-include cmake *
 recursive-include csrc *

SECURITY.md

@@ -34,6 +34,15 @@ Vulnerabilities that cause denial of service or partial disruption, but do not a
 
 Minor issues such as informational disclosures, logging errors, non-exploitable flaws, or weaknesses that require local or high-privilege access and offer negligible impact. Examples include side channel attacks or hash collisions. These issues often have CVSS scores less than 4.0
 
+## Fix disclosure policy
+
+When a security report is accepted, the fix process depends on the severity:
+
+* **CRITICAL and HIGH severity**: Fixes are developed in a private security fork and coordinated with the prenotification group before public disclosure.
+* **MODERATE and LOW severity**: Fixes are developed and submitted as public pull requests. These issues do not require embargo since they do not enable arbitrary code execution or significant data breach, and public visibility accelerates community review and adoption of the fix.
+
+The vulnerability management team reserves the right to adjust the disclosure approach on a case-by-case basis, taking into account factors such as active exploitation, unusual attack surface, or coordination requirements with downstream vendors.
+
 ## Prenotification policy
 
 For certain security issues of CRITICAL, HIGH, or MODERATE severity level, we may prenotify certain organizations or vendors that ship vLLM. The purpose of this prenotification is to allow for a coordinated release of fixes for severe issues.

benchmarks/attention_benchmarks/README.md

@@ -108,7 +108,6 @@ python benchmark.py \
     --backends flash triton flashinfer \
     --batch-specs "q2k" "8q1s1k" "2q2k_32q1s1k" \
     --num-layers 10 \
-    --repeats 5 \
     --output-csv results.csv
 ```
 
@@ -164,14 +163,17 @@ python benchmark.py \
 # Model configuration
 --num-layers N                      # Number of layers
 --head-dim N                        # Head dimension
+--v-head-dim N                      # Value head dimension (defaults to --head-dim)
 --num-q-heads N                     # Query heads
 --num-kv-heads N                    # KV heads
 --block-size N                      # Block size
+--kv-lora-rank N                    # MLA KV LoRA rank
+--qk-nope-head-dim N                # MLA non-RoPE QK head dim
+--qk-rope-head-dim N                # MLA RoPE QK head dim
 
 # Benchmark settings
 --device DEVICE                     # Device (default: cuda:0)
---repeats N                         # Repetitions
---warmup-iters N                    # Warmup iterations
+--warmup-ms N                       # Warmup window in ms for triton do_bench
 --profile-memory                    # Profile memory usage
 
 # Parameter sweeps
@@ -211,8 +213,6 @@ config = BenchmarkConfig(
     num_kv_heads=1,
     block_size=128,
     device="cuda:0",
-    repeats=5,
-    warmup_iters=3,
 )
 
 # CUTLASS MLA with specific num_kv_splits
@@ -253,14 +253,10 @@ formatter.save_json(results, "output.json")
 
 ## Tips
 
-**1. Warmup matters** - Use `--warmup-iters 10` for stable results
+**1. Save results** - Always use `--output-csv` or `--output-json`
 
-**2. Multiple repeats** - Use `--repeats 20` for low variance
+**2. Test incrementally** - Start with `--num-layers 1`
 
-**3. Save results** - Always use `--output-csv` or `--output-json`
+**3. Extended grammar** - Leverage spec decode, chunked prefill patterns
 
-**4. Test incrementally** - Start with `--num-layers 1 --repeats 1`
-
-**5. Extended grammar** - Leverage spec decode, chunked prefill patterns
-
-**6. Parameter sweeps** - Use `--sweep-param` and `--sweep-values` to find optimal values
+**4. Parameter sweeps** - Use `--sweep-param` and `--sweep-values` to find optimal values

benchmarks/attention_benchmarks/benchmark.py

@@ -26,6 +26,9 @@ Examples:
 """
 
 import argparse
+import os
+import shutil
+import subprocess
 import sys
 from dataclasses import replace
 from pathlib import Path
@@ -83,13 +86,15 @@ def run_benchmark(config: BenchmarkConfig, **kwargs) -> BenchmarkResult:
         else:
             return run_standard_attention_benchmark(config)
     except Exception as e:
+        error_msg = str(e) or repr(e)
         return BenchmarkResult(
             config=config,
             mean_time=float("inf"),
+            median_time=float("inf"),
             std_time=0,
             min_time=float("inf"),
             max_time=float("inf"),
-            error=str(e),
+            error=error_msg,
         )
 
 
@@ -115,9 +120,12 @@ def run_model_parameter_sweep(
     """
     all_results = []
 
-    console.print(
-        f"[yellow]Model sweep mode: testing {sweep.param_name} = {sweep.values}[/]"
+    sweep_desc = (
+        f"{sweep.param_name} = {sweep.values}"
+        if sweep.param_name
+        else f"{len(sweep.values)} configurations"
     )
+    console.print(f"[yellow]Model sweep mode: testing {sweep_desc}[/]")
 
     total = len(backends) * len(batch_specs) * len(sweep.values)
 
@@ -125,9 +133,9 @@ def run_model_parameter_sweep(
         for backend in backends:
             for spec in batch_specs:
                 for value in sweep.values:
-                    # Create config with modified model parameter
+                    # Create config with modified model parameter(s)
                     config_args = base_config_args.copy()
-                    config_args[sweep.param_name] = value
+                    sweep.apply(config_args, value)
 
                     # Create config with original backend for running
                     clean_config = BenchmarkConfig(
@@ -144,13 +152,21 @@ def run_model_parameter_sweep(
                     all_results.append(result)
 
                     if not result.success:
+                        err_label = (
+                            f"{sweep.param_name}={value}"
+                            if sweep.param_name
+                            else f"{value}"
+                        )
                         console.print(
-                            f"[red]Error {backend} {spec} {sweep.param_name}="
-                            f"{value}: {result.error}[/]"
+                            f"[red]Error {backend} {spec} {err_label}"
+                            f": {result.error}[/]"
                         )
 
                     pbar.update(1)
 
+    if base_config_args.get("ncu_profile"):
+        return all_results
+
     # Display sweep results - create separate table for each parameter value
     console.print("\n[bold green]Model Parameter Sweep Results:[/]")
     formatter = ResultsFormatter(console)
@@ -184,7 +200,10 @@ def run_model_parameter_sweep(
     )
 
     for param_value in sorted_param_values:
-        console.print(f"\n[bold cyan]{sweep.param_name} = {param_value}[/]")
+        label = (
+            f"{sweep.param_name} = {param_value}" if sweep.param_name else param_value
+        )
+        console.print(f"\n[bold cyan]{label}[/]")
         param_results = by_param_value[param_value]
 
         # Create modified results with original backend names
@@ -200,8 +219,9 @@ def run_model_parameter_sweep(
         formatter.print_table(modified_results, backends, compare_to_fastest=True)
 
     # Show optimal backend for each (param_value, batch_spec) combination
+    sweep_name = sweep.param_name or "config"
     console.print(
-        f"\n[bold cyan]Optimal backend for each ({sweep.param_name}, batch_spec):[/]"
+        f"\n[bold cyan]Optimal backend for each ({sweep_name}, batch_spec):[/]"
     )
 
     # Group by (param_value, batch_spec)
@@ -236,7 +256,10 @@ def run_model_parameter_sweep(
     for param_value, spec in sorted_keys:
         # Print header when param value changes
         if param_value != current_param_value:
-            console.print(f"\n  [bold]{sweep.param_name}={param_value}:[/]")
+            header = (
+                f"{sweep.param_name}={param_value}" if sweep.param_name else param_value
+            )
+            console.print(f"\n  [bold]{header}:[/]")
             current_param_value = param_value
 
         results = by_param_and_spec[(param_value, spec)]
@@ -322,6 +345,9 @@ def run_parameter_sweep(
 
                     pbar.update(1)
 
+    if base_config_args.get("ncu_profile"):
+        return all_results
+
     # Display sweep results
     console.print("\n[bold green]Sweep Results:[/]")
     backend_labels = [sweep.get_label(b, v) for b in backends for v in sweep_values]
@@ -474,11 +500,35 @@ def main():
     parser.add_argument("--num-q-heads", type=int, default=32, help="Query heads")
     parser.add_argument("--num-kv-heads", type=int, default=8, help="KV heads")
     parser.add_argument("--block-size", type=int, default=16, help="Block size")
+    parser.add_argument(
+        "--v-head-dim",
+        type=int,
+        default=None,
+        help="Value head dimension (defaults to --head-dim if unset)",
+    )
+
+    # MLA-specific model dimensions
+    parser.add_argument(
+        "--kv-lora-rank", type=int, default=None, help="MLA KV LoRA rank"
+    )
+    parser.add_argument(
+        "--qk-nope-head-dim", type=int, default=None, help="MLA non-RoPE QK head dim"
+    )
+    parser.add_argument(
+        "--qk-rope-head-dim", type=int, default=None, help="MLA RoPE QK head dim"
+    )
 
     # Benchmark settings
     parser.add_argument("--device", default="cuda:0", help="Device")
-    parser.add_argument("--repeats", type=int, default=1, help="Repetitions")
-    parser.add_argument("--warmup-iters", type=int, default=3, help="Warmup iterations")
+    parser.add_argument(
+        "--warmup-ms",
+        type=int,
+        default=None,
+        help=(
+            "Warmup window in ms for triton's do_bench (default: triton's own). "
+            "Has no effect with CUDA graphs; pass --no-cuda-graphs to use it."
+        ),
+    )
     parser.add_argument("--profile-memory", action="store_true", help="Profile memory")
     parser.add_argument(
         "--kv-cache-dtype",
@@ -491,10 +541,33 @@ def main():
         action=argparse.BooleanOptionalAction,
         default=True,
         help=(
-            "Launch kernels with CUDA graphs to eliminate CPU overhead"
-            "in measurements (default: True)"
+            "Use triton do_bench_cudagraph (True) or do_bench (False) "
+            "for timing. CUDA graphs eliminate CPU launch overhead "
+            "(default: True)"
         ),
     )
+    parser.add_argument(
+        "--num-splits",
+        type=int,
+        default=None,
+        help="FlashAttention split-K factor (0=auto heuristic, 1=disabled, >1=force N)",
+    )
+    parser.add_argument(
+        "--ncu-profile",
+        action="store_true",
+        default=False,
+        help=(
+            "Enable Nsight Compute profiling mode. Automatically wraps the "
+            "script with ncu, capturing a profile with source correlation. "
+            "Use --ncu-output to set the output file name."
+        ),
+    )
+    parser.add_argument(
+        "--ncu-output",
+        type=str,
+        default="profile",
+        help="Output file name for ncu profile (default: 'profile').",
+    )
 
     # Parameter sweep (use YAML config for advanced sweeps)
     parser.add_argument(
@@ -576,23 +649,28 @@ def main():
             model = yaml_config["model"]
             args.num_layers = model.get("num_layers", args.num_layers)
             args.head_dim = model.get("head_dim", args.head_dim)
+            args.v_head_dim = model.get("v_head_dim", args.v_head_dim)
             args.num_q_heads = model.get("num_q_heads", args.num_q_heads)
             args.num_kv_heads = model.get("num_kv_heads", args.num_kv_heads)
             args.block_size = model.get("block_size", args.block_size)
+            # MLA-specific dimensions
+            args.kv_lora_rank = model.get("kv_lora_rank", args.kv_lora_rank)
+            args.qk_nope_head_dim = model.get("qk_nope_head_dim", args.qk_nope_head_dim)
+            args.qk_rope_head_dim = model.get("qk_rope_head_dim", args.qk_rope_head_dim)
 
         # Benchmark settings (top-level keys)
         if "device" in yaml_config:
             args.device = yaml_config["device"]
-        if "repeats" in yaml_config:
-            args.repeats = yaml_config["repeats"]
-        if "warmup_iters" in yaml_config:
-            args.warmup_iters = yaml_config["warmup_iters"]
+        if "warmup_ms" in yaml_config:
+            args.warmup_ms = yaml_config["warmup_ms"]
         if "profile_memory" in yaml_config:
             args.profile_memory = yaml_config["profile_memory"]
         if "kv_cache_dtype" in yaml_config:
             args.kv_cache_dtype = yaml_config["kv_cache_dtype"]
         if "cuda_graphs" in yaml_config:
             args.cuda_graphs = yaml_config["cuda_graphs"]
+        if "ncu_profile" in yaml_config:
+            args.ncu_profile = yaml_config["ncu_profile"]
 
         # Parameter sweep configuration
         if "parameter_sweep" in yaml_config:
@@ -612,7 +690,7 @@ def main():
         if "model_parameter_sweep" in yaml_config:
             sweep_config = yaml_config["model_parameter_sweep"]
             args.model_parameter_sweep = ModelParameterSweep(
-                param_name=sweep_config["param_name"],
+                param_name=sweep_config.get("param_name"),
                 values=sweep_config["values"],
                 label_format=sweep_config.get(
                     "label_format", "{backend}_{param_name}_{value}"
@@ -631,6 +709,32 @@ def main():
 
         console.print()
 
+    # Re-exec under ncu if --ncu-profile and not already inside ncu. This runs
+    # after YAML processing so ncu_profile set via config file is honored.
+    if args.ncu_profile and "_NCU_INNER" not in os.environ:
+        ncu = shutil.which("ncu")
+        if ncu is None:
+            print("Error: 'ncu' not found in PATH", file=sys.stderr)
+            sys.exit(1)
+        cmd = [
+            ncu,
+            "--profile-from-start",
+            "off",
+            "--set",
+            "full",
+            "--import-source",
+            "yes",
+            "-o",
+            args.ncu_output,
+            sys.executable,
+            *sys.argv,
+        ]
+        env = os.environ.copy()
+        env["CUTE_DSL_LINEINFO"] = "1"
+        env["_NCU_INNER"] = "1"
+        print(f"Launching: {' '.join(cmd)}")
+        sys.exit(subprocess.call(cmd, env=env))
+
     # Handle CLI-based parameter sweep (if not from YAML)
     if (
         (not hasattr(args, "parameter_sweep") or args.parameter_sweep is None)
@@ -655,6 +759,18 @@ def main():
     console.print(f"Batch specs: {', '.join(args.batch_specs)}")
     console.print(f"KV cache dtype: {args.kv_cache_dtype}")
     console.print(f"CUDA graphs: {args.cuda_graphs}")
+    if args.warmup_ms is not None and args.cuda_graphs:
+        console.print(
+            "[yellow]Warning: --warmup-ms is ignored with CUDA graphs "
+            "(do_bench_cudagraph warms up internally). Pass --no-cuda-graphs "
+            "to use it.[/]"
+        )
+    if args.num_splits == 0 and args.cuda_graphs:
+        console.print(
+            "[yellow]Warning: --num-splits 0 (FA3 heuristic) is not CUDA-graph "
+            "compatible and may fail or fall back. Pass --no-cuda-graphs or use "
+            "--num-splits >=1.[/]"
+        )
     console.print()
 
     init_workspace_manager(args.device)
@@ -662,6 +778,15 @@ def main():
     # Run benchmarks
     all_results = []
 
+    # Under ncu profiling the kernels run only to be captured by the profiler;
+    # timings are placeholder zeros, so the result tables and saved metrics are
+    # skipped. The Nsight Compute report (--ncu-output) holds the real data.
+    if args.ncu_profile:
+        console.print(
+            "[dim]ncu profiling enabled: result tables and saved metrics are "
+            "skipped (timings are placeholder zeros).[/]"
+        )
+
     # Handle special mode: decode_vs_prefill comparison
     if hasattr(args, "mode") and args.mode == "decode_vs_prefill":
         console.print("[yellow]Mode: Decode vs Prefill pipeline comparison[/]")
@@ -708,11 +833,11 @@ def main():
                         num_kv_heads=args.num_kv_heads,
                         block_size=args.block_size,
                         device=args.device,
-                        repeats=args.repeats,
-                        warmup_iters=args.warmup_iters,
                         profile_memory=args.profile_memory,
                         kv_cache_dtype=args.kv_cache_dtype,
                         use_cuda_graphs=args.cuda_graphs,
+                        ncu_profile=args.ncu_profile,
+                        warmup_ms=args.warmup_ms,
                     )
 
                     # Add decode pipeline config
@@ -749,6 +874,7 @@ def main():
                         result = BenchmarkResult(
                             config=config,
                             mean_time=timing["mean"],
+                            median_time=timing.get("median", timing["mean"]),
                             std_time=timing["std"],
                             min_time=timing["min"],
                             max_time=timing["max"],
@@ -770,6 +896,7 @@ def main():
                         result = BenchmarkResult(
                             config=config,
                             mean_time=float("inf"),
+                            median_time=float("inf"),
                             std_time=0,
                             min_time=float("inf"),
                             max_time=float("inf"),
@@ -779,6 +906,9 @@ def main():
 
                 pbar.update(1)
 
+        if args.ncu_profile:
+            return
+
         # Display decode vs prefill results
         console.print("\n[bold green]Decode vs Prefill Results:[/]")
 
@@ -858,15 +988,20 @@ def main():
         base_config_args = {
             "num_layers": args.num_layers,
             "head_dim": args.head_dim,
+            "v_head_dim": args.v_head_dim,
             "num_q_heads": args.num_q_heads,
             "num_kv_heads": args.num_kv_heads,
             "block_size": args.block_size,
             "device": args.device,
-            "repeats": args.repeats,
-            "warmup_iters": args.warmup_iters,
             "profile_memory": args.profile_memory,
             "kv_cache_dtype": args.kv_cache_dtype,
             "use_cuda_graphs": args.cuda_graphs,
+            "ncu_profile": args.ncu_profile,
+            "warmup_ms": args.warmup_ms,
+            "num_splits": args.num_splits,
+            "kv_lora_rank": args.kv_lora_rank,
+            "qk_nope_head_dim": args.qk_nope_head_dim,
+            "qk_rope_head_dim": args.qk_rope_head_dim,
         }
         all_results = run_model_parameter_sweep(
             backends,
@@ -882,15 +1017,17 @@ def main():
         base_config_args = {
             "num_layers": args.num_layers,
             "head_dim": args.head_dim,
+            "v_head_dim": args.v_head_dim,
             "num_q_heads": args.num_q_heads,
             "num_kv_heads": args.num_kv_heads,
             "block_size": args.block_size,
             "device": args.device,
-            "repeats": args.repeats,
-            "warmup_iters": args.warmup_iters,
             "profile_memory": args.profile_memory,
             "kv_cache_dtype": args.kv_cache_dtype,
             "use_cuda_graphs": args.cuda_graphs,
+            "ncu_profile": args.ncu_profile,
+            "warmup_ms": args.warmup_ms,
+            "num_splits": args.num_splits,
         }
         all_results = run_parameter_sweep(
             backends, args.batch_specs, base_config_args, args.parameter_sweep, console
@@ -914,15 +1051,17 @@ def main():
                             batch_spec=spec,
                             num_layers=args.num_layers,
                             head_dim=args.head_dim,
+                            v_head_dim=getattr(args, "v_head_dim", None),
                             num_q_heads=args.num_q_heads,
                             num_kv_heads=args.num_kv_heads,
                             block_size=args.block_size,
                             device=args.device,
-                            repeats=args.repeats,
-                            warmup_iters=args.warmup_iters,
                             profile_memory=args.profile_memory,
                             kv_cache_dtype=args.kv_cache_dtype,
                             use_cuda_graphs=args.cuda_graphs,
+                            ncu_profile=args.ncu_profile,
+                            warmup_ms=args.warmup_ms,
+                            num_splits=args.num_splits,
                         )
 
                         result = run_benchmark(config)
@@ -935,9 +1074,10 @@ def main():
 
                         pbar.update(1)
 
-            console.print("\n[bold green]Results:[/]")
-            formatter = ResultsFormatter(console)
-            formatter.print_table(decode_results, backends)
+            if not args.ncu_profile:
+                console.print("\n[bold green]Results:[/]")
+                formatter = ResultsFormatter(console)
+                formatter.print_table(decode_results, backends)
 
         # Run prefill backend comparison
         if prefill_backends:
@@ -962,9 +1102,8 @@ def main():
                             num_kv_heads=args.num_kv_heads,
                             block_size=args.block_size,
                             device=args.device,
-                            repeats=args.repeats,
-                            warmup_iters=args.warmup_iters,
                             profile_memory=args.profile_memory,
+                            warmup_ms=args.warmup_ms,
                             prefill_backend=pb,
                         )
 
@@ -980,16 +1119,17 @@ def main():
 
                         pbar.update(1)
 
-            console.print("\n[bold green]Prefill Backend Results:[/]")
-            formatter = ResultsFormatter(console)
-            formatter.print_table(
-                prefill_results, prefill_backends, compare_to_fastest=True
-            )
+            if not args.ncu_profile:
+                console.print("\n[bold green]Prefill Backend Results:[/]")
+                formatter = ResultsFormatter(console)
+                formatter.print_table(
+                    prefill_results, prefill_backends, compare_to_fastest=True
+                )
 
         all_results = decode_results + prefill_results
 
-    # Save results
-    if all_results:
+    # Save results (skip ncu profiling runs: timings are placeholder zeros)
+    if all_results and not args.ncu_profile:
         formatter = ResultsFormatter(console)
         if args.output_csv:
             formatter.save_csv(all_results, args.output_csv)

benchmarks/attention_benchmarks/common.py

@@ -15,6 +15,8 @@ from batch_spec import get_batch_type, parse_batch_spec
 from rich.console import Console
 from rich.table import Table
 
+from vllm.triton_utils import triton
+
 
 def batch_spec_sort_key(spec: str) -> tuple[int, int, int]:
     """
@@ -34,6 +36,30 @@ def batch_spec_sort_key(spec: str) -> tuple[int, int, int]:
         return (0, 0, 0)
 
 
+def run_do_bench(
+    benchmark_fn,
+    use_cuda_graphs: bool,
+    warmup_ms: int | None = None,
+) -> list[float]:
+    kwargs: dict[str, Any] = {"return_mode": "all"}
+    if use_cuda_graphs:
+        result = triton.testing.do_bench_cudagraph(benchmark_fn, **kwargs)
+    else:
+        if warmup_ms is not None:
+            kwargs["warmup"] = warmup_ms
+        result = triton.testing.do_bench(benchmark_fn, **kwargs)
+    return result
+
+
+def run_ncu_profile(benchmark_fn) -> None:
+    benchmark_fn()
+    torch.accelerator.synchronize()
+    torch.cuda.cudart().cudaProfilerStart()
+    benchmark_fn()
+    torch.accelerator.synchronize()
+    torch.cuda.cudart().cudaProfilerStop()
+
+
 # Mock classes for vLLM attention infrastructure
 
 
@@ -182,18 +208,37 @@ class ParameterSweep:
 
 @dataclass
 class ModelParameterSweep:
-    """Configuration for sweeping a model configuration parameter."""
+    """Configuration for sweeping model configuration parameter(s).
 
-    param_name: str  # Name of the model config parameter to sweep (e.g., "num_q_heads")
-    values: list[Any]  # List of values to test
-    label_format: str = "{backend}_{param_name}_{value}"  # Result label template
+    Supports two modes:
+    - Single param: param_name="head_dim", values=[128, 256, 512]
+    - Multi param: values=[{head_dim: 192, v_head_dim: 128}, {head_dim: 256}]
+      When values are dicts, each dict's keys are applied as config overrides.
+    """
+
+    param_name: str | None = None
+    values: list[Any] | None = None
+    label_format: str = "{backend}_{param_name}_{value}"
 
     def get_label(self, backend: str, value: Any) -> str:
         """Generate a label for a specific parameter value."""
+        if isinstance(value, dict):
+            return self.label_format.format(
+                backend=backend, param_name=self.param_name, value=value, **value
+            )
         return self.label_format.format(
             backend=backend, param_name=self.param_name, value=value
         )
 
+    def apply(self, config_args: dict, value: Any) -> None:
+        """Apply a sweep value to config args."""
+        if isinstance(value, dict):
+            config_args.update(value)
+        elif self.param_name is not None:
+            config_args[self.param_name] = value
+        else:
+            raise ValueError("param_name must be set if sweep values are not dicts")
+
 
 @dataclass
 class BenchmarkConfig:
@@ -208,10 +253,10 @@ class BenchmarkConfig:
     block_size: int
     device: str
     dtype: torch.dtype = torch.float16
-    repeats: int = 1
-    warmup_iters: int = 3
     profile_memory: bool = False
     use_cuda_graphs: bool = False
+    ncu_profile: bool = False
+    warmup_ms: int | None = None
 
     # "auto" or "fp8"
     kv_cache_dtype: str = "auto"
@@ -226,6 +271,7 @@ class BenchmarkConfig:
     # Backend-specific tuning
     num_kv_splits: int | None = None  # CUTLASS MLA
     reorder_batch_threshold: int | None = None  # FlashAttn MLA, FlashMLA
+    num_splits: int | None = None  # FlashAttention split-K (0=auto, 1=disabled)
 
 
 @dataclass
@@ -234,6 +280,7 @@ class BenchmarkResult:
 
     config: BenchmarkConfig
     mean_time: float  # seconds
+    median_time: float  # seconds
     std_time: float  # seconds
     min_time: float  # seconds
     max_time: float  # seconds
@@ -252,6 +299,7 @@ class BenchmarkResult:
         return {
             "config": asdict(self.config),
             "mean_time": self.mean_time,
+            "median_time": self.median_time,
             "std_time": self.std_time,
             "min_time": self.min_time,
             "max_time": self.max_time,

benchmarks/attention_benchmarks/configs/mla_decode.yaml

@@ -56,8 +56,6 @@ backends:
   - TOKENSPEED_MLA  # Blackwell + R1 dims + FP8 KV (use --kv-cache-dtype fp8)
 
 device: "cuda:0"
-repeats: 100
-warmup_iters: 10
 profile_memory: true
 
 # Backend-specific tuning

benchmarks/attention_benchmarks/configs/mla_mixed_batch.yaml

@@ -51,8 +51,6 @@ backends:
   - FLASHMLA         # Hopper only
 
 device: "cuda:0"
-repeats: 5
-warmup_iters: 3
 profile_memory: true
 
 # Analyze chunked prefill workspace size impact

benchmarks/attention_benchmarks/configs/mla_prefill.yaml

@@ -124,5 +124,3 @@ prefill_backends:
   - tokenspeed
 
 device: "cuda:0"
-repeats: 20
-warmup_iters: 5

benchmarks/attention_benchmarks/configs/mla_sparse_decode.yaml

@@ -53,6 +53,4 @@ backends:
   - FLASHINFER_MLA_SPARSE
 
 device: "cuda:0"
-repeats: 100
-warmup_iters: 10
 profile_memory: true

benchmarks/attention_benchmarks/configs/mla_sparse_prefill.yaml

@@ -57,6 +57,4 @@ backends:
   - FLASHINFER_MLA_SPARSE
 
 device: "cuda:0"
-repeats: 10
-warmup_iters: 3
 profile_memory: true

benchmarks/attention_benchmarks/configs/reorder_threshold.yaml

@@ -63,8 +63,6 @@ model:
 
 # Benchmark settings
 device: "cuda:0"
-repeats: 15          # More repeats for spec decode variance
-warmup_iters: 5
 profile_memory: false
 
 # Output

benchmarks/attention_benchmarks/configs/speculative_decode.yaml

@@ -49,8 +49,6 @@ backends:
 
 # Benchmark settings
 device: "cuda:0"
-repeats: 10  # More repeats for statistical significance
-warmup_iters: 5
 profile_memory: false
 
 # Test these threshold values for optimization

benchmarks/attention_benchmarks/configs/standard_attention.yaml

@@ -43,6 +43,4 @@ backends:
   - FLASHINFER
 
 device: "cuda:0"
-repeats: 5
-warmup_iters: 3
 profile_memory: false

benchmarks/attention_benchmarks/configs/standard_decode.yaml

@@ -0,0 +1,142 @@
+# Standard attention decode benchmark configuration
+# Sweeps num_q_heads and num_kv_heads to isolate effects of:
+#   1. GQA ratio (fixed num_q_heads=32, vary num_kv_heads)
+#   2. Absolute head count (fixed 4:1 ratio, vary scale)
+
+model:
+  num_layers: 32
+  num_q_heads: 32 # Base value, overridden by sweep
+  num_kv_heads: 8 # Base value, overridden by sweep
+  head_dim: 128
+  block_size: 16
+
+# Head count sweep: each entry overrides num_q_heads, num_kv_heads, and
+# head_dim where it differs from the base (128). Head counts are per-GPU
+# (i.e. after TP sharding).
+#
+# Group A — vary GQA ratio (fixed q=32, head_dim=128):
+#   32:32 (MHA), 32:8 (GQA 4:1), 32:4 (GQA 8:1), 32:1 (MQA)
+#
+# Groups B-E — real model configs at various TP degrees:
+#   Model              head_dim  Full    TP2      TP4      TP8
+#   Llama 3 8B           128    32:8   16:4      8:2      4:1
+#   Llama 3 70B          128    64:8   32:4     16:2      8:1
+#   GPT-OSS 120B          64    64:8   32:4     16:2      8:1
+#   Llama 3 405B         128   128:8   64:4     32:2     16:1
+model_parameter_sweep:
+  values:
+    # --- head_dim=128 (Llama 3 family) ---
+    - { num_q_heads: 32, num_kv_heads: 32, head_dim: 128 } # MHA 1:1
+    - { num_q_heads: 32, num_kv_heads: 1, head_dim: 128 } # MQA 32:1
+    - { num_q_heads: 4, num_kv_heads: 1, head_dim: 128 } # Llama 3 8B TP8
+    - { num_q_heads: 8, num_kv_heads: 2, head_dim: 128 } # Llama 3 8B TP4
+    - { num_q_heads: 16, num_kv_heads: 4, head_dim: 128 } # Llama 3 8B TP2
+    - { num_q_heads: 32, num_kv_heads: 8, head_dim: 128 } # Llama 3 8B TP1 / GQA 4:1
+    - { num_q_heads: 8, num_kv_heads: 1, head_dim: 128 } # Llama 3 70B TP8
+    - { num_q_heads: 16, num_kv_heads: 2, head_dim: 128 } # Llama 3 70B TP4
+    - { num_q_heads: 32, num_kv_heads: 4, head_dim: 128 } # Llama 3 70B TP2 / GQA 8:1
+    - { num_q_heads: 64, num_kv_heads: 8, head_dim: 128 } # Llama 3 70B TP1
+    - { num_q_heads: 16, num_kv_heads: 1, head_dim: 128 } # Llama 3 405B TP8
+    - { num_q_heads: 32, num_kv_heads: 2, head_dim: 128 } # Llama 3 405B TP4
+    - { num_q_heads: 64, num_kv_heads: 4, head_dim: 128 } # Llama 3 405B TP2
+    - { num_q_heads: 128, num_kv_heads: 8, head_dim: 128 } # Llama 3 405B TP1
+    # --- head_dim=64 (GPT-OSS 120B) ---
+    - { num_q_heads: 8, num_kv_heads: 1, head_dim: 64 } # GPT-OSS 120B TP8
+    - { num_q_heads: 16, num_kv_heads: 2, head_dim: 64 } # GPT-OSS 120B TP4
+    - { num_q_heads: 32, num_kv_heads: 4, head_dim: 64 } # GPT-OSS 120B TP2
+    - { num_q_heads: 64, num_kv_heads: 8, head_dim: 64 } # GPT-OSS 120B TP1
+  label_format: "{backend}_q{num_q_heads}kv{num_kv_heads}d{head_dim}"
+
+batch_specs:
+  # ---- batch_size x seq_len grid (decode: q_len=1) ----
+  # Small grid for quick iteration. Uncomment for full sweep.
+
+  # Batch size 1
+  - "q1s1k"
+  - "q1s512"
+  - "q1s2k"
+  - "q1s4k"
+  - "q1s8k"
+  - "q1s16k"
+  - "q1s32k"
+
+  # Batch size 2
+  - "2q1s512"
+  - "2q1s1k"
+  - "2q1s2k"
+  - "2q1s4k"
+  - "2q1s8k"
+  - "2q1s16k"
+  - "2q1s32k"
+
+  # Batch size 4
+  - "4q1s512"
+  - "4q1s1k"
+  - "4q1s2k"
+  - "4q1s4k"
+  - "4q1s8k"
+  - "4q1s16k"
+  - "4q1s32k"
+
+  # Batch size 8
+  - "8q1s1k"
+  - "8q1s512"
+  - "8q1s2k"
+  - "8q1s4k"
+  - "8q1s8k"
+  - "8q1s16k"
+  - "8q1s32k"
+
+  # Batch size 16
+  - "16q1s512"
+  - "16q1s1k"
+  - "16q1s2k"
+  - "16q1s4k"
+  - "16q1s8k"
+  - "16q1s16k"
+  - "16q1s32k"
+
+  # Batch size 32
+  - "32q1s512"
+  - "32q1s1k"
+  - "32q1s2k"
+  - "32q1s4k"
+  - "32q1s8k"
+  - "32q1s16k"
+  - "32q1s32k"
+
+  # Batch size 64
+  - "64q1s1k"
+  - "64q1s512"
+  - "64q1s2k"
+  - "64q1s4k"
+  - "64q1s8k"
+  - "64q1s16k"
+  - "64q1s32k"
+
+  # Batch size 128
+  - "128q1s512"
+  - "128q1s1k"
+  - "128q1s2k"
+  - "128q1s4k"
+  - "128q1s8k"
+  - "128q1s16k"
+  - "128q1s32k"
+
+  # Batch size 256
+  - "256q1s1k"
+  - "256q1s512"
+  - "256q1s2k"
+  - "256q1s4k"
+  - "256q1s8k"
+  - "256q1s16k"
+  - "256q1s32k"
+
+# Available backends: FLASH_ATTN, TRITON_ATTN, FLASHINFER
+backends:
+  - FLASH_ATTN
+  - TRITON_ATTN
+  - FLASHINFER
+
+device: "cuda:0"
+profile_memory: false

benchmarks/attention_benchmarks/configs/standard_prefill.yaml

@@ -0,0 +1,108 @@
+# Standard attention prefill benchmark configuration
+# Sweeps num_q_heads and num_kv_heads to isolate effects of:
+#   1. GQA ratio (fixed num_q_heads=32, vary num_kv_heads)
+#   2. Absolute head count (fixed 4:1 ratio, vary scale)
+
+model:
+  num_layers: 32
+  num_q_heads: 32 # Base value, overridden by sweep
+  num_kv_heads: 8 # Base value, overridden by sweep
+  head_dim: 128
+  block_size: 16
+
+# Head count sweep: each entry overrides num_q_heads, num_kv_heads, and
+# head_dim where it differs from the base (128). Head counts are per-GPU
+# (i.e. after TP sharding).
+#
+# Group A — vary GQA ratio (fixed q=32, head_dim=128):
+#   32:32 (MHA), 32:8 (GQA 4:1), 32:4 (GQA 8:1), 32:1 (MQA)
+#
+# Groups B-E — real model configs at various TP degrees:
+#   Model              head_dim  Full    TP2      TP4      TP8
+#   Llama 3 8B           128    32:8   16:4      8:2      4:1
+#   Llama 3 70B          128    64:8   32:4     16:2      8:1
+#   GPT-OSS 120B          64    64:8   32:4     16:2      8:1
+#   Llama 3 405B         128   128:8   64:4     32:2     16:1
+model_parameter_sweep:
+  values:
+    # --- head_dim=128 (Llama 3 family) ---
+    - { num_q_heads: 32, num_kv_heads: 32, head_dim: 128 } # MHA 1:1
+    - { num_q_heads: 32, num_kv_heads: 1, head_dim: 128 } # MQA 32:1
+    - { num_q_heads: 4, num_kv_heads: 1, head_dim: 128 } # Llama 3 8B TP8
+    - { num_q_heads: 8, num_kv_heads: 2, head_dim: 128 } # Llama 3 8B TP4
+    - { num_q_heads: 16, num_kv_heads: 4, head_dim: 128 } # Llama 3 8B TP2
+    - { num_q_heads: 32, num_kv_heads: 8, head_dim: 128 } # Llama 3 8B TP1 / GQA 4:1
+    - { num_q_heads: 8, num_kv_heads: 1, head_dim: 128 } # Llama 3 70B TP8
+    - { num_q_heads: 16, num_kv_heads: 2, head_dim: 128 } # Llama 3 70B TP4
+    - { num_q_heads: 32, num_kv_heads: 4, head_dim: 128 } # Llama 3 70B TP2 / GQA 8:1
+    - { num_q_heads: 64, num_kv_heads: 8, head_dim: 128 } # Llama 3 70B TP1
+    - { num_q_heads: 16, num_kv_heads: 1, head_dim: 128 } # Llama 3 405B TP8
+    - { num_q_heads: 32, num_kv_heads: 2, head_dim: 128 } # Llama 3 405B TP4
+    - { num_q_heads: 64, num_kv_heads: 4, head_dim: 128 } # Llama 3 405B TP2
+    - { num_q_heads: 128, num_kv_heads: 8, head_dim: 128 } # Llama 3 405B TP1
+    # --- head_dim=64 (GPT-OSS 120B) ---
+    - { num_q_heads: 8, num_kv_heads: 1, head_dim: 64 } # GPT-OSS 120B TP8
+    - { num_q_heads: 16, num_kv_heads: 2, head_dim: 64 } # GPT-OSS 120B TP4
+    - { num_q_heads: 32, num_kv_heads: 4, head_dim: 64 } # GPT-OSS 120B TP2
+    - { num_q_heads: 64, num_kv_heads: 8, head_dim: 64 } # GPT-OSS 120B TP1
+  label_format: "{backend}_q{num_q_heads}kv{num_kv_heads}d{head_dim}"
+
+batch_specs:
+  # ---- batch_size x prefill_len grid (prefill: q_len == seq_len) ----
+  # Total tokens = batch_size * prefill_len, and prefill compute scales with
+  # prefill_len^2, so the largest cells are expensive. Trim batch sizes or
+  # lengths for quick iteration.
+
+  # Batch size 1
+  - "q512"
+  - "q1k"
+  - "q2k"
+  - "q4k"
+  - "q8k"
+  - "q16k"
+  - "q32k"
+
+  # Batch size 2
+  - "2q512"
+  - "2q1k"
+  - "2q2k"
+  - "2q4k"
+  - "2q8k"
+  - "2q16k"
+  - "2q32k"
+
+  # Batch size 4
+  - "4q512"
+  - "4q1k"
+  - "4q2k"
+  - "4q4k"
+  - "4q8k"
+  - "4q16k"
+  - "4q32k"
+
+  # Batch size 8
+  - "8q512"
+  - "8q1k"
+  - "8q2k"
+  - "8q4k"
+  - "8q8k"
+  - "8q16k"
+  - "8q32k"
+
+  # Batch size 16
+  - "16q512"
+  - "16q1k"
+  - "16q2k"
+  - "16q4k"
+  - "16q8k"
+  - "16q16k"
+  - "16q32k"
+
+# Available backends: FLASH_ATTN, TRITON_ATTN, FLASHINFER
+backends:
+  - FLASH_ATTN
+  - TRITON_ATTN
+  - FLASHINFER
+
+device: "cuda:0"
+profile_memory: false

benchmarks/attention_benchmarks/mla_runner.py

@@ -8,6 +8,8 @@ This module provides helpers for running MLA backends without
 needing full VllmConfig integration.
 """
 
+import statistics
+
 import numpy as np
 import torch
 from batch_spec import parse_batch_spec
@@ -17,6 +19,8 @@ from common import (
     MockIndexer,
     MockKVBProj,
     MockLayer,
+    run_do_bench,
+    run_ncu_profile,
     setup_mla_dims,
 )
 
@@ -820,7 +824,7 @@ def _run_single_benchmark(
             num_prefill, mla_dims, query_fmt, device, torch.bfloat16
         )
 
-    # Build forward function
+    # Build forward function (runs a single decode/prefill pass)
     def forward_fn():
         results = []
         if has_decode:
@@ -839,44 +843,35 @@ def _run_single_benchmark(
             )
         return results[0] if len(results) == 1 else tuple(results)
 
-    # Warmup
-    for _ in range(config.warmup_iters):
-        forward_fn()
-    torch.accelerator.synchronize()
-
-    # Optionally capture a CUDA graph after warmup.
-    # Graph replay eliminates CPU launch overhead so timings reflect pure
-    # kernel time.
-    if config.use_cuda_graphs:
-        graph = torch.cuda.CUDAGraph()
-        with torch.cuda.graph(graph):
-            forward_fn()
-        benchmark_fn = graph.replay
-    else:
-        benchmark_fn = forward_fn
-
-    # Benchmark
-    times = []
-    for _ in range(config.repeats):
-        start = torch.cuda.Event(enable_timing=True)
-        end = torch.cuda.Event(enable_timing=True)
-
-        start.record()
+    def benchmark_fn():
         for _ in range(config.num_layers):
-            benchmark_fn()
-        end.record()
+            forward_fn()
 
-        torch.accelerator.synchronize()
-        elapsed_ms = start.elapsed_time(end)
-        times.append(elapsed_ms / 1000.0 / config.num_layers)
+    if config.ncu_profile:
+        run_ncu_profile(benchmark_fn)
+        return BenchmarkResult(
+            config=config,
+            mean_time=0.0,
+            median_time=0.0,
+            std_time=0.0,
+            min_time=0.0,
+            max_time=0.0,
+            throughput_tokens_per_sec=0.0,
+        )
+
+    all_ms = run_do_bench(benchmark_fn, config.use_cuda_graphs, config.warmup_ms)
+
+    # Convert ms to seconds per layer
+    times = [t / 1000.0 / config.num_layers for t in all_ms]
+    mean_time = statistics.mean(times)
 
-    mean_time = float(np.mean(times))
     return BenchmarkResult(
         config=config,
         mean_time=mean_time,
-        std_time=float(np.std(times)),
-        min_time=float(np.min(times)),
-        max_time=float(np.max(times)),
+        median_time=statistics.median(times),
+        std_time=statistics.stdev(times) if len(times) > 1 else 0.0,
+        min_time=min(times),
+        max_time=max(times),
         throughput_tokens_per_sec=total_q / mean_time if mean_time > 0 else 0,
     )
 

benchmarks/attention_benchmarks/runner.py

@@ -9,13 +9,20 @@ This module provides helpers for running standard attention backends
 """
 
 import logging
+import statistics
 import types
 from contextlib import contextmanager
 
-import numpy as np
 import torch
 from batch_spec import parse_batch_spec, reorder_for_flashinfer
-from common import BenchmarkConfig, BenchmarkResult, MockLayer, get_attention_scale
+from common import (
+    BenchmarkConfig,
+    BenchmarkResult,
+    MockLayer,
+    get_attention_scale,
+    run_do_bench,
+    run_ncu_profile,
+)
 
 from vllm.config import (
     CacheConfig,
@@ -208,6 +215,13 @@ def _create_backend_impl(
 
     scale = get_attention_scale(config.head_dim)
 
+    # Set v_head_dim for diff-headdim backends. Always reset (defaulting to
+    # head_dim) so a prior run's value doesn't leak into this one via the
+    # backend's class-level state.
+    if hasattr(backend_class, "set_head_size_v"):
+        v_dim = config.v_head_dim if config.v_head_dim is not None else config.head_dim
+        backend_class.set_head_size_v(v_dim)
+
     impl = backend_class.get_impl_cls()(
         num_heads=config.num_q_heads,
         head_size=config.head_dim,
@@ -300,6 +314,7 @@ def _create_input_tensors(
         from vllm.platforms import current_platform
 
         q_dtype = current_platform.fp8_dtype()
+    v_dim = config.v_head_dim if config.v_head_dim is not None else config.head_dim
     q_list = [
         torch.randn(
             total_q, config.num_q_heads, config.head_dim, device=device, dtype=dtype
@@ -313,9 +328,7 @@ def _create_input_tensors(
         for _ in range(config.num_layers)
     ]
     v_list = [
-        torch.randn(
-            total_q, config.num_kv_heads, config.head_dim, device=device, dtype=dtype
-        )
+        torch.randn(total_q, config.num_kv_heads, v_dim, device=device, dtype=dtype)
         for _ in range(config.num_layers)
     ]
     return q_list, k_list, v_list
@@ -389,14 +402,17 @@ def _run_single_benchmark(
     device: torch.device,
     dtype: torch.dtype,
 ) -> tuple:
-    """Run single benchmark iteration with warmup and timing loop."""
-    total_q = q_list[0].shape[0]
-    out = torch.empty(
-        total_q, config.num_q_heads, config.head_dim, device=device, dtype=dtype
-    )
+    """Run single benchmark using triton's do_bench_cudagraph/do_bench.
 
-    # Warmup
-    for _ in range(config.warmup_iters):
+    Returns:
+        (timing_stats, mem_stats) where timing_stats is a dict with
+        mean/std/min/max in seconds per layer.
+    """
+    total_q = q_list[0].shape[0]
+    v_dim = config.v_head_dim if config.v_head_dim is not None else config.head_dim
+    out = torch.empty(total_q, config.num_q_heads, v_dim, device=device, dtype=dtype)
+
+    def benchmark_fn():
         for i in range(config.num_layers):
             impl.forward(
                 layer,
@@ -407,52 +423,22 @@ def _run_single_benchmark(
                 attn_metadata,
                 output=out,
             )
-    torch.accelerator.synchronize()
 
-    # Optionally capture a CUDA graph after warmup.
-    # Graph replay eliminates CPU launch overhead so timings reflect pure
-    # kernel time.
-    if config.use_cuda_graphs:
-        graph = torch.cuda.CUDAGraph()
-        with torch.cuda.graph(graph):
-            for i in range(config.num_layers):
-                impl.forward(
-                    layer,
-                    q_list[i],
-                    k_list[i],
-                    v_list[i],
-                    cache_list[i],
-                    attn_metadata,
-                    output=out,
-                )
-        benchmark_fn = graph.replay
+    if config.ncu_profile:
+        run_ncu_profile(benchmark_fn)
+        timing_stats = dict.fromkeys(("mean", "median", "std", "min", "max"), 0.0)
     else:
+        all_ms = run_do_bench(benchmark_fn, config.use_cuda_graphs, config.warmup_ms)
 
-        def benchmark_fn():
-            for i in range(config.num_layers):
-                impl.forward(
-                    layer,
-                    q_list[i],
-                    k_list[i],
-                    v_list[i],
-                    cache_list[i],
-                    attn_metadata,
-                    output=out,
-                )
-
-    # Benchmark
-    times = []
-    for _ in range(config.repeats):
-        start = torch.cuda.Event(enable_timing=True)
-        end = torch.cuda.Event(enable_timing=True)
-
-        start.record()
-        benchmark_fn()
-        end.record()
-
-        torch.accelerator.synchronize()
-        elapsed_ms = start.elapsed_time(end)
-        times.append(elapsed_ms / 1000.0 / config.num_layers)  # seconds per layer
+        # Convert ms to seconds per layer
+        times = [t / 1000.0 / config.num_layers for t in all_ms]
+        timing_stats = {
+            "mean": statistics.mean(times),
+            "std": statistics.stdev(times) if len(times) > 1 else 0.0,
+            "min": min(times),
+            "max": max(times),
+            "median": statistics.median(times),
+        }
 
     mem_stats = {}
     if config.profile_memory:
@@ -461,7 +447,7 @@ def _run_single_benchmark(
             "reserved_mb": torch.accelerator.memory_reserved(device) / 1024**2,
         }
 
-    return times, mem_stats
+    return timing_stats, mem_stats
 
 
 # ============================================================================
@@ -541,6 +527,12 @@ def run_attention_benchmark(config: BenchmarkConfig) -> BenchmarkResult:
                 common_attn_metadata=common_metadata,
             )
 
+            # Override num_splits for split-K testing (FlashAttention only)
+            if config.num_splits is not None and hasattr(
+                attn_metadata, "max_num_splits"
+            ):
+                attn_metadata.max_num_splits = config.num_splits
+
             # Only quantize queries when the impl supports it
             quantize_query = config.kv_cache_dtype.startswith("fp8") and getattr(
                 impl, "supports_quant_query_input", False
@@ -553,7 +545,7 @@ def run_attention_benchmark(config: BenchmarkConfig) -> BenchmarkResult:
                 config, max_num_blocks, backend_class, device, dtype
             )
 
-            times, mem_stats = _run_single_benchmark(
+            timing_stats, mem_stats = _run_single_benchmark(
                 config,
                 impl,
                 layer,
@@ -566,15 +558,16 @@ def run_attention_benchmark(config: BenchmarkConfig) -> BenchmarkResult:
                 dtype,
             )
 
-    mean_time = np.mean(times)
+    mean_time = timing_stats["mean"]
     throughput = total_q / mean_time if mean_time > 0 else 0
 
     return BenchmarkResult(
         config=config,
         mean_time=mean_time,
-        std_time=np.std(times),
-        min_time=np.min(times),
-        max_time=np.max(times),
+        median_time=timing_stats["median"],
+        std_time=timing_stats["std"],
+        min_time=timing_stats["min"],
+        max_time=timing_stats["max"],
         throughput_tokens_per_sec=throughput,
         memory_allocated_mb=mem_stats.get("allocated_mb"),
         memory_reserved_mb=mem_stats.get("reserved_mb"),

benchmarks/benchmark_hidden_state_extraction.py

@@ -92,7 +92,6 @@ def run_baseline(
     llm = LLM(
         model=model,
         enable_prefix_caching=False,
-        enable_chunked_prefill=False,
         **extra_args,
     )
     sampling_params = SamplingParams(max_tokens=1)
@@ -194,7 +193,6 @@ async def _run_extraction_async(
     engine_args = AsyncEngineArgs(
         model=model,
         enable_prefix_caching=False,
-        enable_chunked_prefill=False,
         max_num_batched_tokens=40960,
         max_model_len=40960,
         speculative_config={

benchmarks/disagg_benchmarks/disagg_overhead_benchmark.sh

@@ -1,143 +0,0 @@
-#!/bin/bash
-
-# benchmark the overhead of disaggregated prefill.
-# methodology:
-# - send all request to prefill vLLM instance. It will buffer KV cache.
-# - then send all request to decode instance.
-# - The TTFT of decode instance is the overhead.
-
-set -ex
-
-kill_gpu_processes() {
-  # kill all processes on GPU.
-  pgrep pt_main_thread | xargs -r kill -9
-  pgrep python3 | xargs -r kill -9
-  # vLLM now names the process with VLLM prefix after https://github.com/vllm-project/vllm/pull/21445
-  pgrep VLLM | xargs -r kill -9
-  sleep 10
-
-  # remove vllm config file
-  rm -rf ~/.config/vllm
-
-  # Print the GPU memory usage
-  # so that we know if all GPU processes are killed.
-  gpu_memory_usage=$(nvidia-smi --query-gpu=memory.used --format=csv,noheader,nounits -i 0)
-  # The memory usage should be 0 MB.
-  echo "GPU 0 Memory Usage: $gpu_memory_usage MB"
-}
-
-wait_for_server() {
-  # wait for vllm server to start
-  # return 1 if vllm server crashes
-  local port=$1
-  timeout 1200 bash -c "
-    until curl -s localhost:${port}/v1/completions > /dev/null; do
-      sleep 1
-    done" && return 0 || return 1
-}
-
-
-benchmark() {
-
-  export VLLM_LOGGING_LEVEL=DEBUG
-  export VLLM_HOST_IP=$(hostname -I | awk '{print $1}')
-
-  # compare chunked prefill with disaggregated prefill
-
-  results_folder="./results"
-  model="meta-llama/Meta-Llama-3.1-8B-Instruct"
-  dataset_name="sonnet"
-  dataset_path="../sonnet_4x.txt"
-  num_prompts=10
-  qps=$1
-  prefix_len=50
-  input_len=2048
-  output_len=$2
-
-
-  CUDA_VISIBLE_DEVICES=0 vllm serve $model \
-    --port 8100 \
-    --max-model-len 10000 \
-    --gpu-memory-utilization 0.6 \
-    --kv-transfer-config \
-    '{"kv_connector":"P2pNcclConnector","kv_role":"kv_producer","kv_rank":0,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
-
-
-  CUDA_VISIBLE_DEVICES=1 vllm serve $model \
-    --port 8200 \
-    --max-model-len 10000 \
-    --gpu-memory-utilization 0.6 \
-    --kv-transfer-config \
-    '{"kv_connector":"P2pNcclConnector","kv_role":"kv_consumer","kv_rank":1,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
-
-  wait_for_server 8100
-  wait_for_server 8200
-
-  # let the prefill instance finish prefill
-  vllm bench serve \
-    --backend vllm \
-    --model $model \
-    --dataset-name $dataset_name \
-    --dataset-path $dataset_path \
-    --sonnet-input-len $input_len \
-    --sonnet-output-len "$output_len" \
-    --sonnet-prefix-len $prefix_len \
-    --num-prompts $num_prompts \
-    --port 8100 \
-    --save-result \
-    --result-dir $results_folder \
-    --result-filename disagg_prefill_tp1.json \
-    --request-rate "inf"
-
-
-  # send the request to decode.
-  # The TTFT of this command will be the overhead of disagg prefill impl.
-  vllm bench serve \
-    --backend vllm \
-    --model $model \
-    --dataset-name $dataset_name \
-    --dataset-path $dataset_path \
-    --sonnet-input-len $input_len \
-    --sonnet-output-len "$output_len" \
-    --sonnet-prefix-len $prefix_len \
-    --num-prompts $num_prompts \
-    --port 8200 \
-    --save-result \
-    --result-dir $results_folder \
-    --result-filename disagg_prefill_tp1_overhead.json \
-    --request-rate "$qps"
-  kill_gpu_processes
-
-}
-
-
-main() {
-
-  (which wget && which curl) || (apt-get update && apt-get install -y wget curl)
-  (which jq) || (apt-get -y install jq)
-  (which socat) || (apt-get -y install socat)
-
-  pip install quart httpx datasets
-
-  cd "$(dirname "$0")"
-
-  cd ..
-  # create sonnet-4x.txt
-  echo "" > sonnet_4x.txt
-  for _ in {1..4}
-  do
-    cat sonnet.txt >> sonnet_4x.txt
-  done
-  cd disagg_benchmarks
-
-  rm -rf results
-  mkdir results
-
-  default_qps=1
-  default_output_len=1
-  benchmark $default_qps $default_output_len
-
-}
-
-
-main "$@"

benchmarks/disagg_benchmarks/disagg_performance_benchmark.sh

@@ -1,157 +0,0 @@
-#!/bin/bash
-
-# Requirement: 2x GPUs.
-
-
-# Model: meta-llama/Meta-Llama-3.1-8B-Instruct
-# Query: 1024 input tokens, 6 output tokens, QPS 2/4/6/8, 100 requests
-# Resource: 2x GPU
-# Approaches:
-# 2. Chunked prefill: 2 vllm instance with tp=4, equivalent to 1 tp=4 instance with QPS 4
-# 3. Disaggregated prefill: 1 prefilling instance and 1 decoding instance
-# Prefilling instance: max_output_token=1
-# Decoding instance: force the input tokens be the same across requests to bypass prefilling
-
-set -ex
-
-kill_gpu_processes() {
-  # kill all processes on GPU.
-  pgrep pt_main_thread | xargs -r kill -9
-  pgrep python3 | xargs -r kill -9
-  # vLLM now names the process with VLLM prefix after https://github.com/vllm-project/vllm/pull/21445
-  pgrep VLLM | xargs -r kill -9
-  for port in 8000 8100 8200; do lsof -t -i:$port | xargs -r kill -9; done
-  sleep 1
-}
-
-wait_for_server() {
-  # wait for vllm server to start
-  # return 1 if vllm server crashes
-  local port=$1
-  timeout 1200 bash -c "
-    until curl -s localhost:${port}/v1/completions > /dev/null; do
-      sleep 1
-    done" && return 0 || return 1
-}
-
-
-launch_chunked_prefill() {
-  model="meta-llama/Meta-Llama-3.1-8B-Instruct"
-  # disagg prefill
-  CUDA_VISIBLE_DEVICES=0 vllm serve $model \
-    --port 8100 \
-    --max-model-len 10000 \
-    --enable-chunked-prefill \
-    --gpu-memory-utilization 0.6 &
-  CUDA_VISIBLE_DEVICES=1 vllm serve $model \
-    --port 8200 \
-    --max-model-len 10000 \
-    --enable-chunked-prefill \
-    --gpu-memory-utilization 0.6 &
-  wait_for_server 8100
-  wait_for_server 8200
-  python3 round_robin_proxy.py &
-  sleep 1
-}
-
-
-launch_disagg_prefill() {
-  model="meta-llama/Meta-Llama-3.1-8B-Instruct"
-  # disagg prefill
-  CUDA_VISIBLE_DEVICES=0 vllm serve $model \
-    --port 8100 \
-    --max-model-len 10000 \
-    --gpu-memory-utilization 0.6 \
-    --kv-transfer-config \
-    '{"kv_connector":"P2pNcclConnector","kv_role":"kv_producer","kv_rank":0,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
-
-  CUDA_VISIBLE_DEVICES=1 vllm serve $model \
-    --port 8200 \
-    --max-model-len 10000 \
-    --gpu-memory-utilization 0.6 \
-    --kv-transfer-config \
-    '{"kv_connector":"P2pNcclConnector","kv_role":"kv_consumer","kv_rank":1,"kv_parallel_size":2,"kv_buffer_size":5e9}' &
-
-  wait_for_server 8100
-  wait_for_server 8200
-  python3 disagg_prefill_proxy_server.py &
-  sleep 1
-}
-
-
-benchmark() {
-  results_folder="./results"
-  model="meta-llama/Meta-Llama-3.1-8B-Instruct"
-  dataset_name="sonnet"
-  dataset_path="../sonnet_4x.txt"
-  num_prompts=100
-  qps=$1
-  prefix_len=50
-  input_len=1024
-  output_len=$2
-  tag=$3
-
-  vllm bench serve \
-    --backend vllm \
-    --model $model \
-    --dataset-name $dataset_name \
-    --dataset-path $dataset_path \
-    --sonnet-input-len $input_len \
-    --sonnet-output-len "$output_len" \
-    --sonnet-prefix-len $prefix_len \
-    --num-prompts $num_prompts \
-    --port 8000 \
-    --save-result \
-    --result-dir $results_folder \
-    --result-filename "$tag"-qps-"$qps".json \
-    --request-rate "$qps"
-
-  sleep 2
-}
-
-
-main() {
-
-  (which wget && which curl) || (apt-get update && apt-get install -y wget curl)
-  (which jq) || (apt-get -y install jq)
-  (which socat) || (apt-get -y install socat)
-  (which lsof) || (apt-get -y install lsof)
-
-  pip install quart httpx matplotlib aiohttp datasets
-
-  cd "$(dirname "$0")"
-
-  cd ..
-  # create sonnet-4x.txt so that we can sample 2048 tokens for input
-  echo "" > sonnet_4x.txt
-  for _ in {1..4}
-  do
-    cat sonnet.txt >> sonnet_4x.txt
-  done
-  cd disagg_benchmarks
-
-  rm -rf results
-  mkdir results
-
-  default_output_len=6
-
-  export VLLM_HOST_IP=$(hostname -I | awk '{print $1}')
-
-  launch_chunked_prefill
-  for qps in 2 4 6 8; do
-  benchmark $qps $default_output_len chunked_prefill
-  done
-  kill_gpu_processes
-
-  launch_disagg_prefill
-  for qps in 2 4 6 8; do
-  benchmark $qps $default_output_len disagg_prefill
-  done
-  kill_gpu_processes
-
-  python3 visualize_benchmark_results.py
-
-}
-
-
-main "$@"

benchmarks/disagg_benchmarks/disagg_prefill_proxy_server.py

@@ -1,260 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-import argparse
-import asyncio
-import logging
-import os
-import time
-import uuid
-from urllib.parse import urlparse
-
-import aiohttp
-from quart import Quart, Response, make_response, request
-
-# Configure logging
-logging.basicConfig(level=logging.INFO)
-logger = logging.getLogger(__name__)
-
-
-def parse_args():
-    """parse command line arguments"""
-    parser = argparse.ArgumentParser(description="vLLM P/D disaggregation proxy server")
-
-    # Add args
-    parser.add_argument(
-        "--timeout",
-        type=float,
-        default=6 * 60 * 60,
-        help="Timeout for backend service requests in seconds (default: 21600)",
-    )
-    parser.add_argument(
-        "--port",
-        type=int,
-        default=8000,
-        help="Port to run the server on (default: 8000)",
-    )
-    parser.add_argument(
-        "--prefill-url",
-        type=str,
-        default="http://localhost:8100",
-        help="Prefill service base URL (protocol + host[:port])",
-    )
-    parser.add_argument(
-        "--decode-url",
-        type=str,
-        default="http://localhost:8200",
-        help="Decode service base URL (protocol + host[:port])",
-    )
-    parser.add_argument(
-        "--kv-host",
-        type=str,
-        default="localhost",
-        help="Hostname or IP used by KV transfer (default: localhost)",
-    )
-    parser.add_argument(
-        "--prefill-kv-port",
-        type=int,
-        default=14579,
-        help="Prefill KV port (default: 14579)",
-    )
-    parser.add_argument(
-        "--decode-kv-port",
-        type=int,
-        default=14580,
-        help="Decode KV port (default: 14580)",
-    )
-
-    return parser.parse_args()
-
-
-def main():
-    """parse command line arguments"""
-    args = parse_args()
-
-    # Initialize configuration using command line parameters
-    AIOHTTP_TIMEOUT = aiohttp.ClientTimeout(total=args.timeout)
-    PREFILL_SERVICE_URL = args.prefill_url
-    DECODE_SERVICE_URL = args.decode_url
-    PORT = args.port
-
-    PREFILL_KV_ADDR = f"{args.kv_host}:{args.prefill_kv_port}"
-    DECODE_KV_ADDR = f"{args.kv_host}:{args.decode_kv_port}"
-
-    logger.info(
-        "Proxy resolved KV addresses -> prefill: %s, decode: %s",
-        PREFILL_KV_ADDR,
-        DECODE_KV_ADDR,
-    )
-
-    app = Quart(__name__)
-
-    # Attach the configuration object to the application instance so helper
-    # coroutines can read the resolved backend URLs and timeouts without using
-    # globals.
-    app.config.update(
-        {
-            "AIOHTTP_TIMEOUT": AIOHTTP_TIMEOUT,
-            "PREFILL_SERVICE_URL": PREFILL_SERVICE_URL,
-            "DECODE_SERVICE_URL": DECODE_SERVICE_URL,
-            "PREFILL_KV_ADDR": PREFILL_KV_ADDR,
-            "DECODE_KV_ADDR": DECODE_KV_ADDR,
-        }
-    )
-
-    def _normalize_base_url(url: str) -> str:
-        """Remove any trailing slash so path joins behave predictably."""
-        return url.rstrip("/")
-
-    def _get_host_port(url: str) -> str:
-        """Return the hostname:port portion for logging and KV headers."""
-        parsed = urlparse(url)
-        host = parsed.hostname or "localhost"
-        port = parsed.port
-        if port is None:
-            port = 80 if parsed.scheme == "http" else 443
-        return f"{host}:{port}"
-
-    PREFILL_BASE = _normalize_base_url(PREFILL_SERVICE_URL)
-    DECODE_BASE = _normalize_base_url(DECODE_SERVICE_URL)
-    KV_TARGET = _get_host_port(DECODE_SERVICE_URL)
-
-    def _build_headers(request_id: str) -> dict[str, str]:
-        """Construct the headers expected by vLLM's P2P disagg connector."""
-        headers: dict[str, str] = {"X-Request-Id": request_id, "X-KV-Target": KV_TARGET}
-        api_key = os.environ.get("OPENAI_API_KEY")
-        if api_key:
-            headers["Authorization"] = f"Bearer {api_key}"
-        return headers
-
-    async def _run_prefill(
-        request_path: str,
-        payload: dict,
-        headers: dict[str, str],
-        request_id: str,
-    ):
-        url = f"{PREFILL_BASE}{request_path}"
-        start_ts = time.perf_counter()
-        logger.info("[prefill] start request_id=%s url=%s", request_id, url)
-        try:
-            async with (
-                aiohttp.ClientSession(timeout=AIOHTTP_TIMEOUT) as session,
-                session.post(url=url, json=payload, headers=headers) as resp,
-            ):
-                if resp.status != 200:
-                    error_text = await resp.text()
-                    raise RuntimeError(
-                        f"Prefill backend error {resp.status}: {error_text}"
-                    )
-                await resp.read()
-                logger.info(
-                    "[prefill] done request_id=%s status=%s elapsed=%.2fs",
-                    request_id,
-                    resp.status,
-                    time.perf_counter() - start_ts,
-                )
-        except asyncio.TimeoutError as exc:
-            raise RuntimeError(f"Prefill service timeout at {url}") from exc
-        except aiohttp.ClientError as exc:
-            raise RuntimeError(f"Prefill service unavailable at {url}") from exc
-
-    async def _stream_decode(
-        request_path: str,
-        payload: dict,
-        headers: dict[str, str],
-        request_id: str,
-    ):
-        url = f"{DECODE_BASE}{request_path}"
-        # Stream tokens from the decode service once the prefill stage has
-        # materialized KV caches on the target workers.
-        logger.info("[decode] start request_id=%s url=%s", request_id, url)
-        try:
-            async with (
-                aiohttp.ClientSession(timeout=AIOHTTP_TIMEOUT) as session,
-                session.post(url=url, json=payload, headers=headers) as resp,
-            ):
-                if resp.status != 200:
-                    error_text = await resp.text()
-                    logger.error(
-                        "Decode backend error %s - %s", resp.status, error_text
-                    )
-                    err_msg = (
-                        '{"error": "Decode backend error ' + str(resp.status) + '"}'
-                    )
-                    yield err_msg.encode()
-                    return
-                logger.info(
-                    "[decode] streaming response request_id=%s status=%s",
-                    request_id,
-                    resp.status,
-                )
-                async for chunk_bytes in resp.content.iter_chunked(1024):
-                    yield chunk_bytes
-                logger.info("[decode] finished streaming request_id=%s", request_id)
-        except asyncio.TimeoutError:
-            logger.error("Decode service timeout at %s", url)
-            yield b'{"error": "Decode service timeout"}'
-        except aiohttp.ClientError as exc:
-            logger.error("Decode service error at %s: %s", url, exc)
-            yield b'{"error": "Decode service unavailable"}'
-
-    async def process_request():
-        """Process a single request through prefill and decode stages"""
-        try:
-            original_request_data = await request.get_json()
-
-            # Create prefill request (max_tokens=1)
-            prefill_request = original_request_data.copy()
-            prefill_request["max_tokens"] = 1
-            if "max_completion_tokens" in prefill_request:
-                prefill_request["max_completion_tokens"] = 1
-
-            # Execute prefill stage
-            # The request id encodes both KV socket addresses so the backend can
-            # shuttle tensors directly via NCCL once the prefill response
-            # completes.
-            request_id = (
-                f"___prefill_addr_{PREFILL_KV_ADDR}___decode_addr_"
-                f"{DECODE_KV_ADDR}_{uuid.uuid4().hex}"
-            )
-
-            headers = _build_headers(request_id)
-            await _run_prefill(request.path, prefill_request, headers, request_id)
-
-            # Execute decode stage and stream response
-            # Pass the unmodified user request so the decode phase can continue
-            # sampling with the already-populated KV cache.
-            generator = _stream_decode(
-                request.path, original_request_data, headers, request_id
-            )
-            response = await make_response(generator)
-            response.timeout = None  # Disable timeout for streaming response
-            return response
-
-        except Exception:
-            logger.exception("Error processing request")
-            return Response(
-                response=b'{"error": "Internal server error"}',
-                status=500,
-                content_type="application/json",
-            )
-
-    @app.route("/v1/completions", methods=["POST"])
-    async def handle_request():
-        """Handle incoming API requests with concurrency and rate limiting"""
-        try:
-            return await process_request()
-        except asyncio.CancelledError:
-            logger.warning("Request cancelled")
-            return Response(
-                response=b'{"error": "Request cancelled"}',
-                status=503,
-                content_type="application/json",
-            )
-
-    # Start the Quart server with host can be set to 0.0.0.0
-    app.run(port=PORT)
-
-
-if __name__ == "__main__":
-    main()

benchmarks/disagg_benchmarks/round_robin_proxy.py

@@ -1,63 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-import asyncio
-import itertools
-
-import aiohttp
-from aiohttp import web
-
-
-class RoundRobinProxy:
-    def __init__(self, target_ports):
-        self.target_ports = target_ports
-        self.port_cycle = itertools.cycle(self.target_ports)
-
-    async def handle_request(self, request):
-        target_port = next(self.port_cycle)
-        target_url = f"http://localhost:{target_port}{request.path_qs}"
-
-        async with aiohttp.ClientSession() as session:
-            try:
-                # Forward the request
-                async with session.request(
-                    method=request.method,
-                    url=target_url,
-                    headers=request.headers,
-                    data=request.content,
-                ) as response:
-                    # Start sending the response
-                    resp = web.StreamResponse(
-                        status=response.status, headers=response.headers
-                    )
-                    await resp.prepare(request)
-
-                    # Stream the response content
-                    async for chunk in response.content.iter_any():
-                        await resp.write(chunk)
-
-                    await resp.write_eof()
-                    return resp
-
-            except Exception as e:
-                return web.Response(text=f"Error: {str(e)}", status=500)
-
-
-async def main():
-    proxy = RoundRobinProxy([8100, 8200])
-    app = web.Application()
-    app.router.add_route("*", "/{path:.*}", proxy.handle_request)
-
-    runner = web.AppRunner(app)
-    await runner.setup()
-    site = web.TCPSite(runner, "localhost", 8000)
-    await site.start()
-
-    print("Proxy server started on http://localhost:8000")
-
-    # Keep the server running
-    await asyncio.Event().wait()
-
-
-if __name__ == "__main__":
-    asyncio.run(main())

benchmarks/disagg_benchmarks/visualize_benchmark_results.py

@@ -1,47 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-import json
-
-import matplotlib.pyplot as plt
-import pandas as pd
-
-if __name__ == "__main__":
-    data = []
-    for name in ["disagg_prefill", "chunked_prefill"]:
-        for qps in [2, 4, 6, 8]:
-            with open(f"results/{name}-qps-{qps}.json") as f:
-                x = json.load(f)
-                x["name"] = name
-                x["qps"] = qps
-                data.append(x)
-
-    df = pd.DataFrame.from_dict(data)
-    dis_df = df[df["name"] == "disagg_prefill"]
-    chu_df = df[df["name"] == "chunked_prefill"]
-
-    plt.style.use("bmh")
-    plt.rcParams["font.size"] = 20
-
-    for key in [
-        "mean_ttft_ms",
-        "median_ttft_ms",
-        "p99_ttft_ms",
-        "mean_itl_ms",
-        "median_itl_ms",
-        "p99_itl_ms",
-    ]:
-        fig, ax = plt.subplots(figsize=(11, 7))
-        plt.plot(
-            dis_df["qps"], dis_df[key], label="disagg_prefill", marker="o", linewidth=4
-        )
-        plt.plot(
-            chu_df["qps"], chu_df[key], label="chunked_prefill", marker="o", linewidth=4
-        )
-        ax.legend()
-
-        ax.set_xlabel("QPS")
-        ax.set_ylabel(key)
-        ax.set_ylim(bottom=0)
-        fig.savefig(f"results/{key}.png")
-        plt.close(fig)

benchmarks/kernels/benchmark_flydsl_moe_w4a16.py

@@ -0,0 +1,277 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+# Copyright (c) 2025 FlyDSL Project Contributors
+
+import json
+import os
+
+import torch
+from aiter.test_common import run_perftest
+
+from vllm.model_executor.layers.fused_moe import fused_experts
+from vllm.model_executor.layers.fused_moe.activation import MoEActivation
+from vllm.model_executor.layers.fused_moe.config import (
+    int4_w4a16_moe_quant_config,
+)
+from vllm.model_executor.layers.fused_moe.fused_flydsl_moe import fused_flydsl_moe
+from vllm.model_executor.layers.quantization.compressed_tensors.compressed_tensors_moe import (  # noqa: E501
+    compressed_tensors_moe_w4a16_flydsl,
+)
+from vllm.platforms import current_platform
+
+RoutingBuffers = tuple[
+    torch.Tensor,  # sorted_token_ids
+    torch.Tensor,  # sorted_weights
+    torch.Tensor,  # sorted_expert_ids
+    torch.Tensor,  # num_valid_ids (shape [1], i32)
+    int,  # sorted_size
+    int,  # blocks
+]
+
+MODEL_PARAMS_TO_TUNE = [
+    # (num_experts, inter_dim, hidden_size, topk)
+    (384, 256, 7168, 8),  # Kimi K2.5 TP=8
+    (384, 512, 7168, 8),  # Kimi K2.5 TP=4
+]
+
+NUM_TOKENS_TO_TUNE = [
+    1,
+    2,
+    4,
+    8,
+    16,
+    24,
+    32,
+    48,
+    64,
+    128,
+    256,
+    512,
+    1024,
+    2048,
+    4096,
+    8192,
+]
+
+TILE_M_SEARCH_SPACE = [16, 32, 64, 128, 256]
+TILE_N_SEARCH_SPACE = [16, 32, 64, 128, 256]
+TILE_K_SEARCH_SPACE = [16, 32, 64, 128, 256, 512]
+TILE_N2_SEARCH_SPACE = [16, 32, 64, 128, 256]
+TILE_K2_SEARCH_SPACE = [16, 32, 64, 128, 256, 512]
+
+TILE_CONFIGS = []
+for tile_m in TILE_M_SEARCH_SPACE:
+    for tile_n in TILE_N_SEARCH_SPACE:
+        for tile_k in TILE_K_SEARCH_SPACE:
+            for tile_n2 in TILE_N2_SEARCH_SPACE:
+                for tile_k2 in TILE_K2_SEARCH_SPACE:
+                    TILE_CONFIGS.append(
+                        {
+                            "tile_m": tile_m,
+                            "tile_n": tile_n,
+                            "tile_k": tile_k,
+                            "tile_n2": tile_n2,
+                            "tile_k2": tile_k2,
+                        }
+                    )
+
+
+def tune_flydsl_moe_w4a16(
+    device: str = "cuda", num_iters: int = 100, num_warmup: int = 10
+):
+    packed_factor = 8
+    w13_num_shards = 2
+    params_dtype = torch.bfloat16
+    group_size = 32
+    scale_factor = 0.01
+
+    for model_params in MODEL_PARAMS_TO_TUNE:
+        num_experts = model_params[0]
+        inter_dim = model_params[1]
+        hidden_size = model_params[2]
+        topk = model_params[3]
+        print(
+            f"\nTuning: num_experts={num_experts}, inter_dim={inter_dim}, "
+            f"hidden_size={hidden_size}, topk={topk}...\n"
+        )
+
+        w2_scales_size = inter_dim
+        num_groups_w2 = w2_scales_size // group_size
+        num_groups_w13 = hidden_size // group_size
+
+        w13_weight = torch.randint(
+            0,
+            255,
+            (num_experts, hidden_size // packed_factor, w13_num_shards * inter_dim),
+            dtype=torch.int32,
+            device=device,
+        )
+
+        w2_weight = torch.randint(
+            0,
+            255,
+            (num_experts, inter_dim // packed_factor, hidden_size),
+            dtype=torch.int32,
+            device=device,
+        )
+        w13_scale = scale_factor * torch.randn(
+            num_experts,
+            num_groups_w13,
+            w13_num_shards * inter_dim,
+            dtype=params_dtype,
+            device=device,
+        )
+        w2_scale = scale_factor * torch.randn(
+            num_experts, num_groups_w2, hidden_size, dtype=params_dtype, device=device
+        )
+
+        w13 = w13_weight
+        w13 = compressed_tensors_moe_w4a16_flydsl._gptq_int32_to_flydsl_packed(w13)
+        w13 = w13.view(-1).contiguous()
+
+        w2 = w2_weight
+        w2 = compressed_tensors_moe_w4a16_flydsl._gptq_int32_to_flydsl_packed(w2)
+        w2 = w2.view(-1).contiguous()
+
+        w13_scale_flydsl = w13_scale
+        w2_scale_flydsl = w2_scale
+
+        if group_size > 0 and w13_scale.dim() == 3 and w13_scale.shape[1] > 1:
+            E, G, N = w13_scale.shape
+            w13_scale_flydsl = (
+                w13_scale_flydsl.view(E, G // 2, 2, N)
+                .permute(0, 1, 3, 2)
+                .contiguous()
+                .view(-1)
+                .contiguous()
+            )
+        elif w13_scale.dim() == 3 and w13_scale.shape[1] == 1:
+            w13_scale_flydsl = w13_scale_flydsl.squeeze(1)
+
+        if group_size > 0 and w2_scale.dim() == 3 and w2_scale.shape[1] > 1:
+            E, G, N = w2_scale.shape
+            w2_scale_flydsl = (
+                w2_scale_flydsl.view(E, G // 2, 2, N)
+                .permute(0, 1, 3, 2)
+                .contiguous()
+                .view(-1)
+                .contiguous()
+            )
+        elif w2_scale.dim() == 3 and w2_scale.shape[1] == 1:
+            w2_scale_flydsl = w2_scale_flydsl.squeeze(1)
+
+        w13_scale_flydsl = w13_scale_flydsl.contiguous()
+        w2_scale_flydsl = w2_scale_flydsl.contiguous()
+
+        w13.is_shuffled = True
+        w2.is_shuffled = True
+
+        w13_weight_scale = w13_scale.transpose(1, 2).contiguous()
+        w2_weight_scale = w2_scale.transpose(1, 2).contiguous()
+        w13_weight_packed = w13_weight.transpose(1, 2).contiguous().view(torch.uint8)
+        w2_weight_packed = w2_weight.transpose(1, 2).contiguous().view(torch.uint8)
+
+        moe_quant_config = int4_w4a16_moe_quant_config(
+            w1_scale=w13_weight_scale,
+            w2_scale=w2_weight_scale,
+            w1_zp=None,
+            w2_zp=None,
+            block_shape=[0, group_size],
+        )
+
+        tuned_config = {}
+
+        for num_tokens in NUM_TOKENS_TO_TUNE:
+            score = torch.rand(
+                (num_tokens, num_experts), device=device, dtype=torch.float32
+            )
+            topk_vals, topk_ids = torch.topk(score, k=topk, dim=1)
+            topk_weights = torch.softmax(topk_vals, dim=1).to(torch.float32)
+            x = torch.randn(
+                (num_tokens, hidden_size), dtype=torch.bfloat16, device=device
+            )
+            us_best = float("inf")
+            for tile_config in TILE_CONFIGS:
+                try:
+                    tile_m = tile_config["tile_m"]
+                    tile_n = tile_config["tile_n"]
+                    tile_k = tile_config["tile_k"]
+                    tile_n2 = tile_config["tile_n2"]
+                    tile_k2 = tile_config["tile_k2"]
+
+                    model_dim = x.shape[1]
+                    assert model_dim % 64 == 0
+                    assert model_dim % tile_k == 0
+                    assert inter_dim % tile_n == 0
+                    assert model_dim % tile_n2 == 0
+                    assert inter_dim % tile_k2 == 0
+                    assert ((tile_m * tile_k2) % 256) == 0
+                    bytes_per_thread_x = (tile_m * tile_k2) // 256
+                    assert (bytes_per_thread_x % 4) == 0
+
+                    out, _us = run_perftest(
+                        fused_flydsl_moe,
+                        x,
+                        w13,
+                        w2,
+                        num_experts,
+                        inter_dim,
+                        topk_weights,
+                        topk_ids,
+                        num_iters=num_iters,
+                        num_warmup=num_warmup,
+                        w1_scale=w13_scale_flydsl,
+                        w2_scale=w2_scale_flydsl,
+                        topk=topk_weights.shape[-1],
+                        group_size=group_size,
+                        doweight_stage1=False,
+                        scale_is_bf16=True,
+                        config=tile_config,
+                    )
+                    torch.accelerator.synchronize()
+                except Exception:
+                    torch.accelerator.synchronize()
+                    continue
+                else:
+                    us = _us.item()
+                    if us < us_best:
+                        out_ref = fused_experts(
+                            x,
+                            w13_weight_packed,
+                            w2_weight_packed,
+                            topk_weights=topk_weights,
+                            topk_ids=topk_ids,
+                            activation=MoEActivation.SILU,
+                            apply_router_weight_on_input=False,
+                            global_num_experts=num_experts,
+                            expert_map=None,
+                            quant_config=moe_quant_config,
+                        )
+                        try:
+                            assert torch.allclose(out, out_ref, atol=0.5, rtol=0.1)
+                        except Exception:
+                            continue
+                        else:
+                            print(
+                                f"For [num_tokens={num_tokens}, num_experts={num_experts}, "  # noqa: E501
+                                f"inter_dim={inter_dim}] found new best "  # noqa: E501
+                                f"config={tile_config}, us={us:0.3f}"
+                            )
+                            us_best = us
+                            tuned_config[str(num_tokens)] = tile_config
+            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"
+            )
+            tuner_dir_path = os.path.dirname(os.path.realpath(__file__))
+            store_path = os.path.join(tuner_dir_path, tuned_config_file_name)
+            with open(store_path, "w") as f:
+                json.dump(tuned_config, f, indent=4)
+            print(
+                f"\nTuned config for num_tokens={num_tokens} was stored at {store_path}\n"  # noqa: E501
+            )
+
+
+if __name__ == "__main__":
+    tune_flydsl_moe_w4a16(device="cuda")

benchmarks/kernels/benchmark_moe.py

@@ -250,7 +250,7 @@ def benchmark_config(
                     num_experts=num_experts,
                     experts_per_token=topk,
                     hidden_dim=hidden_size,
-                    intermediate_size_per_partition=shard_intermediate_size,
+                    intermediate_size=shard_intermediate_size,
                     num_local_experts=num_experts,
                     num_logical_experts=num_experts,
                     activation=MoEActivation.SILU,
@@ -271,7 +271,6 @@ def benchmark_config(
                     moe_config=moe_config,
                     quant_config=quant_config,
                 ),
-                inplace=not disable_inplace(),
             )
 
         with override_config(config):
@@ -279,7 +278,6 @@ def benchmark_config(
                 x, input_gating, topk, renormalize=not use_deep_gemm
             )
 
-            inplace = not disable_inplace()
             if use_deep_gemm:
                 return deep_gemm_experts.apply(
                     x,
@@ -298,7 +296,6 @@ def benchmark_config(
                 w2,
                 topk_weights,
                 topk_ids,
-                inplace=inplace,
                 quant_config=quant_config,
             )
 
@@ -795,6 +792,12 @@ def get_model_params(config):
         topk = text_config.num_experts_per_tok
         intermediate_size = text_config.moe_intermediate_size
         hidden_size = text_config.hidden_size
+    elif architecture == "DiffusionGemmaForBlockDiffusion":
+        text_config = config.get_text_config()
+        E = text_config.num_experts
+        topk = text_config.top_k_experts
+        intermediate_size = text_config.moe_intermediate_size
+        hidden_size = text_config.hidden_size
     elif architecture == "HunYuanMoEV1ForCausalLM":
         E = config.num_experts
         topk = config.moe_topk[0]

benchmarks/kernels/benchmark_router_gemm.py

@@ -0,0 +1,154 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+
+import torch
+import torch.nn.functional as F
+
+from vllm import _custom_ops as ops
+from vllm.platforms import current_platform
+from vllm.transformers_utils.config import get_config
+from vllm.triton_utils import triton
+from vllm.utils.argparse_utils import FlexibleArgumentParser
+
+# Dimensions supported by the DSV3 specialized kernel
+DSV3_SUPPORTED_NUM_EXPERTS = [256, 384]
+DSV3_SUPPORTED_HIDDEN_SIZES = [7168]
+
+# Dimensions supported by the gpt-oss specialized kernel
+GPT_OSS_SUPPORTED_NUM_EXPERTS = [32, 128]
+GPT_OSS_SUPPORTED_HIDDEN_SIZES = [2880]
+
+# Dimensions supported by the fp32 specialized kernel (MiniMax-M2)
+FP32_SUPPORTED_NUM_EXPERTS = [256]
+FP32_SUPPORTED_HIDDEN_SIZES = [3072]
+FP32_MAX_TOKENS = 32
+
+
+def get_batch_size_range(max_batch_size):
+    return [2**x for x in range(14) if 2**x <= max_batch_size]
+
+
+def get_model_params(config):
+    if config.architectures[0] in (
+        "DeepseekV2ForCausalLM",
+        "DeepseekV3ForCausalLM",
+        "DeepseekV32ForCausalLM",
+    ):
+        num_experts = config.n_routed_experts
+        hidden_size = config.hidden_size
+    elif config.architectures[0] in ("GptOssForCausalLM",) or config.architectures[
+        0
+    ] in ("MiniMaxM2ForCausalLM",):
+        num_experts = config.num_local_experts
+        hidden_size = config.hidden_size
+    else:
+        raise ValueError(f"Unsupported architecture: {config.architectures}")
+    return num_experts, hidden_size
+
+
+def get_benchmark(model, max_batch_size, trust_remote_code):
+    @triton.testing.perf_report(
+        triton.testing.Benchmark(
+            x_names=["batch_size"],
+            x_vals=get_batch_size_range(max_batch_size),
+            x_log=False,
+            line_arg="provider",
+            line_vals=[
+                "torch",
+                "vllm",
+            ],
+            line_names=["PyTorch", "vLLM"],
+            styles=([("blue", "-"), ("red", "-")]),
+            ylabel="TFLOPs",
+            plot_name=f"{model} router gemm throughput",
+            args={},
+        )
+    )
+    def benchmark(batch_size, provider):
+        config = get_config(model=model, trust_remote_code=trust_remote_code)
+        num_experts, hidden_size = get_model_params(config)
+
+        is_hopper_or_blackwell = current_platform.is_device_capability(
+            90
+        ) or current_platform.is_device_capability_family(100)
+        allow_dsv3_router_gemm = (
+            is_hopper_or_blackwell
+            and num_experts in DSV3_SUPPORTED_NUM_EXPERTS
+            and hidden_size in DSV3_SUPPORTED_HIDDEN_SIZES
+        )
+        allow_gpt_oss_router_gemm = (
+            is_hopper_or_blackwell
+            and num_experts in GPT_OSS_SUPPORTED_NUM_EXPERTS
+            and hidden_size in GPT_OSS_SUPPORTED_HIDDEN_SIZES
+        )
+        is_fp32_router_model = (
+            is_hopper_or_blackwell
+            and num_experts in FP32_SUPPORTED_NUM_EXPERTS
+            and hidden_size in FP32_SUPPORTED_HIDDEN_SIZES
+        )
+        allow_fp32_router_gemm = is_fp32_router_model and batch_size <= FP32_MAX_TOKENS
+
+        # Weight dtype: fp32 kernel requires fp32 weights; others use bf16.
+        weight_dtype = torch.float32 if is_fp32_router_model else torch.bfloat16
+        mat_a = torch.randn(
+            (batch_size, hidden_size), dtype=torch.bfloat16, device="cuda"
+        ).contiguous()
+        mat_b = torch.randn(
+            (num_experts, hidden_size), dtype=weight_dtype, device="cuda"
+        ).contiguous()
+        bias = torch.randn(
+            num_experts, dtype=torch.bfloat16, device="cuda"
+        ).contiguous()
+
+        has_bias = allow_gpt_oss_router_gemm
+
+        quantiles = [0.5, 0.2, 0.8]
+
+        if provider == "torch":
+
+            def runner():
+                if allow_fp32_router_gemm:
+                    F.linear(mat_a.float(), mat_b)
+                elif has_bias:
+                    F.linear(mat_a, mat_b, bias)
+                else:
+                    F.linear(mat_a, mat_b)
+        elif provider == "vllm":
+
+            def runner():
+                if allow_dsv3_router_gemm:
+                    ops.dsv3_router_gemm(mat_a, mat_b, torch.bfloat16)
+                elif allow_fp32_router_gemm:
+                    ops.fp32_router_gemm(mat_a, mat_b)
+                elif allow_gpt_oss_router_gemm:
+                    ops.gpt_oss_router_gemm(mat_a, mat_b, bias)
+                elif is_fp32_router_model:
+                    # batch_size > FP32_MAX_TOKENS: fall back to F.linear
+                    F.linear(mat_a.float(), mat_b)
+                else:
+                    F.linear(mat_a, mat_b)
+
+        ms, min_ms, max_ms = triton.testing.do_bench_cudagraph(
+            runner, quantiles=quantiles
+        )
+
+        def tflops(t_ms):
+            flops = 2 * batch_size * hidden_size * num_experts
+            return flops / (t_ms * 1e-3) / 1e12
+
+        return tflops(ms), tflops(max_ms), tflops(min_ms)
+
+    return benchmark
+
+
+if __name__ == "__main__":
+    parser = FlexibleArgumentParser()
+    parser.add_argument("--model", type=str, default="openai/gpt-oss-20b")
+    parser.add_argument("--max-batch-size", default=16, type=int)
+    parser.add_argument("--trust-remote-code", action="store_true")
+    args = parser.parse_args()
+
+    # Get the benchmark function
+    benchmark = get_benchmark(args.model, args.max_batch_size, args.trust_remote_code)
+    # Run performance benchmark
+    benchmark.run(print_data=True)

benchmarks/kv_cache_watermark.sh

@@ -0,0 +1,248 @@
+#!/bin/bash
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+#
+# Reproducible demonstration of the KV cache watermark (`--watermark`) for
+# reducing preemption thrashing.
+#
+# The watermark is the fraction of total KV cache blocks the scheduler keeps
+# free when admitting a waiting/preempted request into the running queue.
+#
+# Why this workload triggers thrashing:
+#   Requests are admitted based on the KV cache they need *at admission time*.
+#   With `--scheduler-reserve-full-isl` (default) the input length is reserved up
+#   front, but the *output* length is unknown and unreserved. A decode-heavy
+#   workload (output >> input) at high concurrency therefore over-admits while
+#   requests are short, then runs out of KV cache as they all grow during decode
+#   -> the scheduler preempts (recompute) recently-admitted requests, re-prefills
+#   them later, and repeats. The watermark keeps a block of KV cache free so
+#   running requests can grow into it instead of triggering this churn.
+#
+# This script launches `vllm serve` under a deliberately KV-constrained config
+# and a decode-heavy workload, sweeping the watermark across several values, and
+# reports the preemption count (scraped from /metrics), throughput, and latency
+# percentiles for each. It then plots the results.
+#
+# Default workload: concurrency 200, input ~300 tokens, output ~4000 tokens
+# (+/- 20% variance), sized to run each config for ~5 minutes.
+#
+# Usage:
+#   benchmarks/kv_cache_watermark.sh
+#   MODEL=Qwen/Qwen2.5-14B-Instruct TP=2 benchmarks/kv_cache_watermark.sh
+#
+# Run inside the vLLM virtualenv (so `vllm` and `python` resolve to it).
+set -euo pipefail
+
+# ---- Config (override via environment) -------------------------------------
+MODEL=${MODEL:-Qwen/Qwen2.5-7B-Instruct}
+TP=${TP:-1}
+PORT=${PORT:-8000}
+URL="http://127.0.0.1:${PORT}"
+# Constrain the KV cache to a *near-critical* size: large enough that the engine
+# can run stably, but small enough that greedy over-admission tips it into
+# preemption thrashing. (Independent of GPU size, so the demo is reproducible.)
+# At the default workload this fits ~1.5x the mean concurrent KV demand.
+KV_CACHE_MEMORY_GB=${KV_CACHE_MEMORY_GB:-16}
+MAX_MODEL_LEN=${MAX_MODEL_LEN:-8192}
+MAX_NUM_SEQS=${MAX_NUM_SEQS:-256}
+# Optional weight loader (e.g. fastsafetensors on the GCP cluster).
+LOAD_FORMAT=${LOAD_FORMAT:-auto}
+# Decode-heavy workload: moderate input, long output, with length variance. The
+# long output means preempted requests have generated a lot before eviction, so
+# resuming them re-prefills a long sequence (high recomputation cost).
+INPUT_LEN=${INPUT_LEN:-1000}
+OUTPUT_LEN=${OUTPUT_LEN:-5000}
+RANGE_RATIO=${RANGE_RATIO:-0.2}
+CONCURRENCY=${CONCURRENCY:-128}
+# Enough prompts to keep each config saturated for ~5+ minutes.
+NUM_PROMPTS=${NUM_PROMPTS:-450}
+OUTDIR=${OUTDIR:-./watermark_bench_results}
+# Watermark fractions compared. "label value" per line; value=0 disables it.
+CONFIGS=${CONFIGS:-"off    0
+w0.02  0.02
+w0.05  0.05
+w0.10  0.10
+w0.15  0.15"}
+
+KV_CACHE_MEMORY_BYTES=$((KV_CACHE_MEMORY_GB * 1024 * 1024 * 1024))
+mkdir -p "$OUTDIR"
+
+SERVER_PID=""
+cleanup() { [[ -n "$SERVER_PID" ]] && kill "$SERVER_PID" 2>/dev/null || true; }
+trap cleanup EXIT
+
+scrape_preemptions() {
+  # Sum the vllm:num_preemptions_total counter across engines.
+  python - "${URL}/metrics" <<'PY'
+import sys, urllib.request
+total = 0.0
+try:
+    body = urllib.request.urlopen(sys.argv[1], timeout=10).read().decode("utf-8", "replace")
+    for line in body.splitlines():
+        if line.startswith("vllm:num_preemptions_total"):
+            total += float(line.rsplit(" ", 1)[-1])
+except Exception as e:  # noqa: BLE001
+    print(f"scrape error: {e}", file=sys.stderr)
+print(int(total))
+PY
+}
+
+wait_for_server() {
+  for _ in $(seq 1 300); do
+    if curl -s "${URL}/health" >/dev/null 2>&1; then return 0; fi
+    if ! kill -0 "$SERVER_PID" 2>/dev/null; then
+      echo "ERROR: server process exited during startup" >&2; return 1
+    fi
+    sleep 5
+  done
+  echo "ERROR: server did not become ready" >&2; return 1
+}
+
+run_one() {
+  local label=$1 watermark=$2
+  echo
+  echo "==================== watermark: ${label} (${watermark}) ===================="
+  vllm serve "$MODEL" \
+    --tensor-parallel-size "$TP" \
+    --load-format "$LOAD_FORMAT" \
+    --kv-cache-memory-bytes "$KV_CACHE_MEMORY_BYTES" \
+    --max-model-len "$MAX_MODEL_LEN" \
+    --max-num-seqs "$MAX_NUM_SEQS" \
+    --no-enable-prefix-caching \
+    --watermark "$watermark" \
+    --port "$PORT" >"${OUTDIR}/serve_${label}.log" 2>&1 &
+  SERVER_PID=$!
+  wait_for_server
+  sleep 5
+
+  local pre post
+  pre=$(scrape_preemptions)
+  vllm bench serve \
+    --backend vllm \
+    --base-url "$URL" \
+    --model "$MODEL" \
+    --dataset-name random \
+    --random-input-len "$INPUT_LEN" \
+    --random-output-len "$OUTPUT_LEN" \
+    --random-range-ratio "$RANGE_RATIO" \
+    --ignore-eos \
+    --num-prompts "$NUM_PROMPTS" \
+    --max-concurrency "$CONCURRENCY" \
+    --percentile-metrics "ttft,tpot,itl,e2el" \
+    --metric-percentiles "50,90,99" \
+    --save-result \
+    --result-dir "$OUTDIR" \
+    --result-filename "bench_${label}.json"
+  post=$(scrape_preemptions)
+  echo "${label} ${watermark} $((post - pre))" >>"${OUTDIR}/preemptions.txt"
+
+  kill "$SERVER_PID" 2>/dev/null || true
+  for _ in $(seq 1 60); do curl -s "${URL}/health" >/dev/null 2>&1 || break; sleep 2; done
+  SERVER_PID=""
+  sleep 10
+}
+
+: >"${OUTDIR}/preemptions.txt"
+while read -r label watermark; do
+  [[ -z "${label:-}" ]] && continue
+  run_one "$label" "$watermark"
+done <<<"$CONFIGS"
+
+echo
+echo "==================== summary ===================="
+python - "$OUTDIR" <<'PY'
+import json, os, sys
+outdir = sys.argv[1]
+pre = {}
+order = []
+for line in open(os.path.join(outdir, "preemptions.txt")):
+    label, watermark, n = line.split()
+    pre[label] = (float(watermark), int(n))
+    order.append(label)
+
+def g(d, *names):
+    for n in names:
+        if d.get(n) is not None:
+            return d[n]
+    return float("nan")
+
+cols = ["watermark", "frac", "preempt", "out_tok/s", "req/s",
+        "TTFT_p50", "TTFT_p99", "ITL_p99", "E2EL_p50"]
+print("  ".join(f"{c:>10}" for c in cols))
+rows = []
+for label in order:
+    watermark, n = pre[label]
+    d = json.load(open(os.path.join(outdir, f"bench_{label}.json")))
+    rows.append(dict(
+        label=label, watermark=watermark, preempt=n,
+        out_tok_s=g(d, "output_throughput"),
+        req_s=g(d, "request_throughput"),
+        ttft_p50=g(d, "p50_ttft_ms", "median_ttft_ms"),
+        ttft_p99=g(d, "p99_ttft_ms"),
+        itl_p99=g(d, "p99_itl_ms"),
+        e2el_p50=g(d, "p50_e2el_ms", "median_e2el_ms"),
+    ))
+    print("  ".join(f"{str(v):>10}" for v in [
+        label, watermark, n,
+        f"{rows[-1]['out_tok_s']:.0f}",
+        f"{rows[-1]['req_s']:.3f}",
+        f"{rows[-1]['ttft_p50']/1000:.2f}",
+        f"{rows[-1]['ttft_p99']/1000:.2f}",
+        f"{rows[-1]['itl_p99']:.2f}",
+        f"{rows[-1]['e2el_p50']/1000:.1f}",
+    ]))
+print("\n(TTFT/E2EL in seconds; ITL in ms. Lower preempt is better.)")
+
+# ---- Plot -------------------------------------------------------------------
+try:
+    import matplotlib
+    matplotlib.use("Agg")
+    import matplotlib.pyplot as plt
+except Exception as e:  # noqa: BLE001
+    print(f"\n(skip plot: matplotlib unavailable: {e})")
+    sys.exit(0)
+
+x = [r["watermark"] for r in rows]
+xt = [f"{r['watermark']:g}\n({r['label']})" for r in rows]
+idx = list(range(len(rows)))
+
+fig, axes = plt.subplots(2, 2, figsize=(12, 8))
+fig.suptitle(
+    f"KV cache watermark sweep — {os.path.basename(os.path.abspath(outdir))}",
+    fontsize=12,
+)
+
+ax = axes[0][0]
+ax.bar(idx, [r["preempt"] for r in rows], color="tab:red")
+ax.set_title("Preemptions (lower is better)")
+ax.set_ylabel("preemptions")
+ax.set_xticks(idx); ax.set_xticklabels(xt)
+
+ax = axes[0][1]
+ax.plot(idx, [r["out_tok_s"] for r in rows], "o-", color="tab:green")
+ax.set_title("Output throughput (higher is better)")
+ax.set_ylabel("tokens/s")
+ax.set_xticks(idx); ax.set_xticklabels(xt)
+
+ax = axes[1][0]
+ax.plot(idx, [r["itl_p99"] for r in rows], "o-", color="tab:blue")
+ax.set_title("Inter-token latency p99 (lower is better)")
+ax.set_ylabel("ITL p99 (ms)")
+ax.set_xlabel("watermark fraction")
+ax.set_xticks(idx); ax.set_xticklabels(xt)
+
+ax = axes[1][1]
+ax.plot(idx, [r["ttft_p50"] / 1000 for r in rows], "o-", label="TTFT p50")
+ax.plot(idx, [r["ttft_p99"] / 1000 for r in rows], "o-", label="TTFT p99")
+ax.plot(idx, [r["e2el_p50"] / 1000 for r in rows], "o-", label="E2EL p50")
+ax.set_title("Latency (lower is better)")
+ax.set_ylabel("seconds")
+ax.set_xlabel("watermark fraction")
+ax.set_xticks(idx); ax.set_xticklabels(xt)
+ax.legend()
+
+fig.tight_layout(rect=(0, 0, 1, 0.95))
+out_png = os.path.join(outdir, "watermark_results.png")
+fig.savefig(out_png, dpi=120)
+print(f"\nWrote plot: {out_png}")
+PY

benchmarks/multi_turn/benchmark_serving_multi_turn.py

@@ -65,6 +65,32 @@ class RequestArgs(NamedTuple):
     limit_min_tokens: int  # Use negative value for no limit
     limit_max_tokens: int  # Use negative value for no limit
     timeout_sec: int
+    send_conversation_id: bool
+    headers: dict[str, str]
+
+
+def parse_custom_header(header: str) -> tuple[str, str]:
+    separators = (":", "=")
+    for separator in separators:
+        if separator in header:
+            key, value = header.split(separator, 1)
+            key = key.strip()
+            value = value.strip()
+            if key:
+                return key, value
+            break
+    raise argparse.ArgumentTypeError(
+        "Headers must be provided as 'Header-Name: value' or 'Header-Name=value'"
+    )
+
+
+def build_request_headers(
+    api_key: str | None, custom_headers: list[tuple[str, str]] | None
+) -> dict[str, str]:
+    headers = dict(custom_headers or [])
+    if api_key:
+        headers["Authorization"] = f"Bearer {api_key}"
+    return headers
 
 
 class BenchmarkArgs(NamedTuple):
@@ -218,12 +244,11 @@ async def send_request(
     max_tokens: int | None = None,
     timeout_sec: int = 120,
     conversation_id: str | None = None,
+    headers: dict[str, str] | None = None,
 ) -> ServerResponse:
     payload = {
         "model": model,
         "messages": messages,
-        "seed": 0,
-        "temperature": 0.0,
     }
 
     if conversation_id is not None:
@@ -233,13 +258,17 @@ async def send_request(
         payload["stream"] = True
         payload["stream_options"] = {"include_usage": False}
 
-    if min_tokens is not None:
-        payload["min_tokens"] = min_tokens
+    # if min_tokens is not None:
+    #     payload["min_tokens"] = min_tokens
 
     if max_tokens is not None:
         payload["max_tokens"] = max_tokens
 
-    headers = {"Content-Type": "application/json"}
+    request_headers = {"Content-Type": "application/json"}
+    if conversation_id is not None:
+        request_headers["X-Session-ID"] = str(conversation_id)
+    if headers is not None:
+        request_headers.update(headers)
 
     # Calculate the timeout for the request
     if max_tokens is not None:
@@ -265,7 +294,7 @@ async def send_request(
     most_recent_timestamp: int = start_time
 
     async with session.post(
-        url=chat_url, json=payload, headers=headers, timeout=timeout
+        url=chat_url, json=payload, headers=request_headers, timeout=timeout
     ) as response:
         http_status = HTTPStatus(response.status)
         if http_status == HTTPStatus.OK:
@@ -317,6 +346,8 @@ async def send_request(
             latency = time.perf_counter_ns() - start_time
 
     if ttft is None:
+        if stream:
+            valid_response = False
         # The response was a single chunk
         ttft = latency
 
@@ -423,7 +454,8 @@ async def send_turn(
         min_tokens,
         max_tokens,
         req_args.timeout_sec,
-        conversation_id=conv_id,
+        conversation_id=conv_id if req_args.send_conversation_id else None,
+        headers=req_args.headers,
     )
 
     if response.valid is False:
@@ -872,6 +904,7 @@ def get_client_config(
     # Arguments for API requests
     chat_url = f"{args.url}/v1/chat/completions"
     model_name = args.served_model_name if args.served_model_name else args.model
+    headers = build_request_headers(args.api_key, args.header)
 
     req_args = RequestArgs(
         chat_url=chat_url,
@@ -880,6 +913,8 @@ def get_client_config(
         limit_min_tokens=args.limit_min_tokens,
         limit_max_tokens=args.limit_max_tokens,
         timeout_sec=args.request_timeout_sec,
+        send_conversation_id=args.send_conversation_id,
+        headers=headers,
     )
 
     return client_args, req_args
@@ -1245,19 +1280,19 @@ def process_statistics(
         )
 
 
-async def get_server_info(url: str) -> None:
+async def get_server_info(url: str, headers: dict[str, str] | None = None) -> None:
     logger.info(f"{Color.BLUE}Collecting information from server: {url}{Color.RESET}")
     async with aiohttp.ClientSession() as session:
         # Get server version (not mandatory, "version" endpoint may not exist)
         url_version = f"{url}/version"
-        async with session.get(url_version) as response:
+        async with session.get(url_version, headers=headers) as response:
             if HTTPStatus(response.status) == HTTPStatus.OK:
                 text = await response.text()
                 logger.info(f"{Color.BLUE}Server version: {text}{Color.RESET}")
 
         # Get available models
         url_models = f"{url}/v1/models"
-        async with session.get(url_models) as response:
+        async with session.get(url_models, headers=headers) as response:
             if HTTPStatus(response.status) == HTTPStatus.OK:
                 text = await response.text()
                 logger.info(f"{Color.BLUE}Models:{Color.RESET}")
@@ -1323,6 +1358,22 @@ async def main() -> None:
         help="Base URL for the LLM API server",
     )
 
+    parser.add_argument(
+        "--api-key",
+        type=str,
+        default=None,
+        help="API key to send as an Authorization bearer token",
+    )
+    parser.add_argument(
+        "--header",
+        action="append",
+        type=parse_custom_header,
+        default=None,
+        metavar="KEY=VALUE",
+        help="Custom request header. Can be specified multiple times. "
+        "Accepts 'Header-Name: value' or 'Header-Name=value'.",
+    )
+
     parser.add_argument(
         "-p",
         "--num-clients",
@@ -1437,6 +1488,22 @@ async def main() -> None:
         help="Disable stream/streaming mode (set 'stream' to False in the API request)",
     )
 
+    parser.add_argument(
+        "--send-conversation-id",
+        default=False,
+        action="store_true",
+        help=(
+            "Inject a `conversation_id` field into each Chat Completions "
+            "payload. This is a non-standard OpenAI extension consumed by "
+            "vLLM's disaggregated multi-turn proxy "
+            "(examples/disaggregated/disaggregated_serving/"
+            "disagg_proxy_multiturn.py) to key cross-turn KV cache reuse. "
+            "Leave disabled (default) when targeting strict "
+            "OpenAI-compatible endpoints; enable when benchmarking the "
+            "disaggregated proxy."
+        ),
+    )
+
     parser.add_argument(
         "-e",
         "--excel-output",
@@ -1525,7 +1592,8 @@ async def main() -> None:
         args.model, trust_remote_code=args.trust_remote_code
     )
 
-    await get_server_info(args.url)
+    headers = build_request_headers(args.api_key, args.header)
+    await get_server_info(args.url, headers=headers)
 
     # Load the input file (either conversations of configuration file)
     logger.info(f"Reading input file: {args.input_file}")

build_rust.sh

@@ -1,5 +1,5 @@
 #!/bin/bash
-# Build the vllm-rs Rust frontend binary and install it into the vllm package.
+# Build vLLM Rust artifacts and install them into the vllm package.
 # Usage: ./build_rust.sh [--debug]
 #
 # By default builds in release mode. Pass --debug for faster compile times
@@ -8,8 +8,6 @@
 set -euo pipefail
 
 REPO_ROOT="$(cd "$(dirname "$0")" && pwd)"
-RUST_DIR="$REPO_ROOT/rust"
-TARGET_PATH="${VLLM_RS_TARGET_PATH:-$REPO_ROOT/vllm/vllm-rs}"
 
 # Read the required toolchain from rust-toolchain.toml.
 TOOLCHAIN=$(grep '^channel' "$REPO_ROOT/rust-toolchain.toml" | sed 's/.*= *"\(.*\)"/\1/')
@@ -27,18 +25,9 @@ if ! rustup run "$TOOLCHAIN" rustc --version &>/dev/null; then
 fi
 
 if [[ "${1:-}" == "--debug" ]]; then
-    PROFILE_ARGS=()
-    PROFILE_DIR="debug"
+    PROFILE_ARG="--debug"
 else
-    PROFILE_ARGS=(--release)
-    PROFILE_DIR="release"
+    PROFILE_ARG="--release"
 fi
 
-cargo +"$TOOLCHAIN" build "${PROFILE_ARGS[@]}" \
-    --manifest-path "$RUST_DIR/Cargo.toml" \
-    --bin vllm-rs \
-    --features native-tls-vendored
-
-mkdir -p "$(dirname "$TARGET_PATH")"
-cp "$RUST_DIR/target/$PROFILE_DIR/vllm-rs" "$TARGET_PATH"
-echo "Installed vllm-rs to $TARGET_PATH"
+python3 "$REPO_ROOT/tools/build_rust.py" "$PROFILE_ARG"

cmake/cpu_extension.cmake

@@ -166,6 +166,10 @@ elseif (S390_FOUND)
         "-mtune=native")
 elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
     message(STATUS "RISC-V detected")
+    if(DEFINED VLLM_RVV_VLEN AND NOT VLLM_RVV_VLEN GREATER 0)
+        message(FATAL_ERROR
+            "VLLM_RVV_VLEN must be a positive integer; got '${VLLM_RVV_VLEN}'")
+    endif()
     # VLLM_RVV_VLEN selects the target VLEN. Auto-detected from /proc/cpuinfo
     # by default; override with -DVLLM_RVV_VLEN=128 or -DVLLM_RVV_VLEN=256.
     if(NOT DEFINED VLLM_RVV_VLEN)
@@ -189,8 +193,7 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
                 "RISC-V RVV is available but VLEN could not be auto-detected. "
                 "Please specify VLEN explicitly:\n"
                 "  -DVLLM_RVV_VLEN=128   (for VLEN=128 hardware)\n"
-                "  -DVLLM_RVV_VLEN=256   (for VLEN=256 hardware, e.g. Spacemit X100)\n"
-                "  -DVLLM_RVV_VLEN=0     (force scalar, no RVV)")
+                "  -DVLLM_RVV_VLEN=256   (for VLEN=256 hardware, e.g. Spacemit X100)")
         endif()
     endif()
     if(VLLM_RVV_VLEN AND VLLM_RVV_VLEN GREATER 0)
@@ -219,7 +222,7 @@ endif()
 
 
 # Build oneDNN for GEMM kernels
-if (ENABLE_X86_ISA OR (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND) OR POWER9_FOUND OR POWER10_FOUND OR POWER11_FOUND)
+if (ENABLE_X86_ISA OR (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND) OR POWER9_FOUND OR POWER10_FOUND OR POWER11_FOUND OR RVV_FP16_FOUND OR RVV_BF16_FOUND)
     # Fetch and build Arm Compute Library (ACL) as oneDNN's backend for AArch64
     # TODO [fadara01]: remove this once ACL can be fetched and built automatically as a dependency of oneDNN
     set(ONEDNN_AARCH64_USE_ACL OFF CACHE BOOL "")
@@ -369,6 +372,18 @@ else()
     add_compile_definitions(-DVLLM_NUMA_DISABLED)
 endif()
 
+# check if the pytorch wheel ships libopenblas.so.
+set(VLLM_OPENBLAS_LIB "")
+if (NOT ENABLE_X86_ISA)
+    file(GLOB _VLLM_TORCH_OPENBLAS_LIBS
+        "${TORCH_INSTALL_PREFIX}/lib/libopenblas*.so*")
+    # Note: we don't link openblas directly to _C extension, as it's available through libtorch.so 
+    if (_VLLM_TORCH_OPENBLAS_LIBS)
+        list(GET _VLLM_TORCH_OPENBLAS_LIBS 0 VLLM_OPENBLAS_LIB)
+        message(STATUS "CPU OpenBLAS library: ${VLLM_OPENBLAS_LIB}")
+    endif()
+endif()
+
 #
 # Generate CPU attention dispatch header
 #
@@ -387,6 +402,7 @@ endif()
 #
 set(VLLM_EXT_SRC
     "csrc/cpu/activation.cpp"
+    "csrc/cpu/sgl-kernels/fla.cpp"
     "csrc/cpu/utils.cpp"
     "csrc/cpu/spec_decode_utils.cpp"
     "csrc/cpu/layernorm.cpp"
@@ -396,6 +412,13 @@ set(VLLM_EXT_SRC
     "csrc/cpu/cpu_attn.cpp"
     "csrc/cpu/torch_bindings.cpp")
 
+if (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64" AND VLLM_RVV_VLEN AND
+        VLLM_RVV_VLEN GREATER 0 AND (RVV_FP16_FOUND OR RVV_BF16_FOUND))
+    set(VLLM_EXT_SRC
+        "csrc/cpu/cpu_wna16.cpp"
+        ${VLLM_EXT_SRC})
+endif()
+
 if (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND)
     set(VLLM_EXT_SRC
         "csrc/cpu/shm.cpp"
@@ -403,15 +426,26 @@ if (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND)
         ${VLLM_EXT_SRC})
 endif()
 
+if (POWER9_FOUND OR POWER10_FOUND OR POWER11_FOUND)	
+    set(VLLM_EXT_SRC
+        "csrc/cpu/shm.cpp"
+        ${VLLM_EXT_SRC})
+endif()
+
 if(USE_ONEDNN)
     set(VLLM_EXT_SRC
         "csrc/cpu/dnnl_kernels.cpp"
         ${VLLM_EXT_SRC})
 endif()
 
+if (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
+    set(VLLM_EXT_SRC
+        "csrc/cpu/sgl-kernels/gemm_int4.cpp"
+        ${VLLM_EXT_SRC})
+endif()
+
 if (ENABLE_X86_ISA)
     set(VLLM_EXT_SRC_SGL
-        "csrc/cpu/sgl-kernels/fla.cpp"
         "csrc/cpu/sgl-kernels/conv.cpp"
         "csrc/cpu/sgl-kernels/gemm.cpp"
         "csrc/cpu/sgl-kernels/gemm_int8.cpp"
@@ -423,6 +457,7 @@ if (ENABLE_X86_ISA)
         "csrc/cpu/sgl-kernels/moe_fp8.cpp")
 
     set(VLLM_EXT_SRC_AVX512
+        "csrc/cpu/sgl-kernels/fla.cpp"
         "csrc/cpu/shm.cpp"
         "csrc/cpu/cpu_wna16.cpp"
         "csrc/cpu/cpu_fused_moe.cpp"
@@ -439,6 +474,7 @@ if (ENABLE_X86_ISA)
         "csrc/moe/dynamic_4bit_int_moe_cpu.cpp") 
 
     set(VLLM_EXT_SRC_AVX2
+        "csrc/cpu/sgl-kernels/fla.cpp"
         "csrc/cpu/utils.cpp"
         "csrc/cpu/spec_decode_utils.cpp"
         "csrc/cpu/cpu_attn.cpp"
@@ -512,6 +548,9 @@ else()
         USE_SABI 3
         WITH_SOABI
     )
+    if (VLLM_OPENBLAS_LIB)
+        target_compile_definitions(_C PRIVATE VLLM_HAS_OPENBLAS)
+    endif()
 endif()
 
 message(STATUS "Enabling C extension.")

cmake/external_projects/qutlass.cmake

@@ -32,21 +32,33 @@ endif()
 message(STATUS "[QUTLASS] QuTLASS is available at ${qutlass_SOURCE_DIR}")
 
 if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-  cuda_archs_loose_intersection(QUTLASS_ARCHS "10.0f;12.0f" "${CUDA_ARCHS}")
+  cuda_archs_loose_intersection(QUTLASS_SM120_ARCHS "12.0f" "${CUDA_ARCHS}")
+  cuda_archs_loose_intersection(QUTLASS_SM100_ARCHS "10.0f" "${CUDA_ARCHS}")
 else()
-  cuda_archs_loose_intersection(QUTLASS_ARCHS "12.0a;12.1a;10.0a;10.3a" "${CUDA_ARCHS}")
+  cuda_archs_loose_intersection(QUTLASS_SM120_ARCHS "12.0a;12.1a" "${CUDA_ARCHS}")
+  cuda_archs_loose_intersection(QUTLASS_SM100_ARCHS "10.0a;10.3a" "${CUDA_ARCHS}")
+endif()
+
+# QUTLASS uses TARGET_CUDA_ARCH as a single preprocessor selector for all its
+# sources. Do not compile a mixed SM100/SM120 arch list with one selector; prefer
+# SM100 when both families are requested because that is the primary deployed
+# target for this extension today.
+if(QUTLASS_SM100_ARCHS)
+  set(QUTLASS_ARCHS "${QUTLASS_SM100_ARCHS}")
+  set(QUTLASS_TARGET_CC 100)
+  if(QUTLASS_SM120_ARCHS)
+    message(WARNING
+      "[QUTLASS] Both SM100 and SM120 archs were requested; selecting SM100 "
+      "because TARGET_CUDA_ARCH is a single compile-time selector.")
+  endif()
+elseif(QUTLASS_SM120_ARCHS)
+  set(QUTLASS_ARCHS "${QUTLASS_SM120_ARCHS}")
+  set(QUTLASS_TARGET_CC 120)
+else()
+  set(QUTLASS_ARCHS)
 endif()
 
 if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND QUTLASS_ARCHS)
-
-  if(QUTLASS_ARCHS MATCHES "10\\.(0a|3a|0f)")
-    set(QUTLASS_TARGET_CC 100)
-  elseif(QUTLASS_ARCHS MATCHES "12\\.[01][af]?")
-    set(QUTLASS_TARGET_CC 120)
-  else()
-    message(FATAL_ERROR "[QUTLASS] internal error parsing CUDA_ARCHS='${QUTLASS_ARCHS}'.")
-  endif()
-
   set(QUTLASS_SOURCES
     ${qutlass_SOURCE_DIR}/qutlass/csrc/bindings.cpp
     ${qutlass_SOURCE_DIR}/qutlass/csrc/gemm.cu

cmake/external_projects/vllm_flash_attn.cmake

@@ -31,7 +31,7 @@ endif()
 
 if(VLLM_FLASH_ATTN_SRC_DIR)
   FetchContent_Declare(
-          vllm-flash-attn SOURCE_DIR 
+          vllm-flash-attn SOURCE_DIR
           ${VLLM_FLASH_ATTN_SRC_DIR}
           BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn
   )
@@ -39,7 +39,7 @@ else()
   FetchContent_Declare(
           vllm-flash-attn
           GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
-          GIT_TAG bce29425653ec0fbc579d329883030e832d15ada
+          GIT_TAG 803020a8fa15407871341d41eba4919ade2ee1ee
           GIT_PROGRESS TRUE
           # Don't share the vllm-flash-attn build between build types
           BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn

cmake/hipify.py

@@ -14,7 +14,18 @@ import argparse
 import os
 import shutil
 
-from torch.utils.hipify.hipify_python import hipify
+from torch.utils.hipify.hipify_python import get_hip_file_path, hipify
+
+
+def _expected_hip_build_path(source_abs: str, output_directory: str) -> str:
+    """Match torch.utils.hipify.hipify_python.preprocessor fout_path naming."""
+    rel = os.path.relpath(source_abs, output_directory)
+    return os.path.abspath(
+        os.path.join(
+            output_directory, get_hip_file_path(rel, is_pytorch_extension=True)
+        )
+    )
+
 
 if __name__ == "__main__":
     parser = argparse.ArgumentParser()
@@ -53,7 +64,11 @@ if __name__ == "__main__":
     hipify_result = hipify(
         project_directory=args.project_dir,
         output_directory=args.output_dir,
-        header_include_dirs=[],
+        # Hipify resolves quoted includes next to the including file first; vLLM
+        # uses paths relative to csrc/ (e.g. "libtorch_stable/torch_utils.h"
+        # from quantization/w8a8/fp8/*.cu). Without an include root here, those
+        # headers are never found and are not hipified or rewritten in dependents.
+        header_include_dirs=["."],
         includes=includes,
         extra_files=extra_files,
         show_detailed=True,
@@ -64,14 +79,20 @@ if __name__ == "__main__":
     hipified_sources = []
     for source in args.sources:
         s_abs = os.path.abspath(source)
-        hipified_s_abs = (
-            hipify_result[s_abs].hipified_path
-            if (
-                s_abs in hipify_result
-                and hipify_result[s_abs].hipified_path is not None
-            )
-            else s_abs
-        )
+        if s_abs in hipify_result and hipify_result[s_abs].hipified_path is not None:
+            path = hipify_result[s_abs].hipified_path
+            # PyTorch skips writing when is_pytorch_extension and text unchanged;
+            # hipified_path then stays *.cu. CMake expects *.hip under output_dir.
+            if s_abs.endswith(".cu") and path.endswith(".cu"):
+                dest = _expected_hip_build_path(s_abs, args.output_dir)
+                if os.path.normpath(path) != os.path.normpath(dest):
+                    os.makedirs(os.path.dirname(dest), exist_ok=True)
+                    shutil.copy2(path, dest)
+                hipified_s_abs = dest
+            else:
+                hipified_s_abs = path
+        else:
+            hipified_s_abs = s_abs
         hipified_sources.append(hipified_s_abs)
 
     assert len(hipified_sources) == len(args.sources)

cmake/utils.cmake

@@ -81,6 +81,14 @@ function (hipify_sources_target OUT_SRCS NAME ORIG_SRCS)
   set_property(GLOBAL APPEND PROPERTY VLLM_HIPIFY_ALL_SRCS ${SRCS})
   set_property(GLOBAL APPEND PROPERTY VLLM_HIPIFY_ALL_BYPRODUCTS ${HIP_SRCS})
 
+  # Chain hipify targets so they run sequentially. Parallel hipify
+  # invocations race on shutil.copytree, overwriting .hip files
+  # produced by another target back to .cu originals.
+  if (DEFINED _VLLM_LAST_HIPIFY_TARGET)
+    add_dependencies(hipify${NAME} ${_VLLM_LAST_HIPIFY_TARGET})
+  endif()
+  set(_VLLM_LAST_HIPIFY_TARGET "hipify${NAME}" PARENT_SCOPE)
+
   # Swap out original extension sources with hipified sources.
   list(APPEND HIP_SRCS ${CXX_SRCS})
   set(${OUT_SRCS} ${HIP_SRCS} PARENT_SCOPE)
@@ -476,6 +484,16 @@ function(cuda_archs_loose_intersection OUT_CUDA_ARCHS SRC_CUDA_ARCHS TGT_CUDA_AR
   set(${OUT_CUDA_ARCHS} ${_CUDA_ARCHS} PARENT_SCOPE)
 endfunction()
 
+
+function(cuda_archs_sm90plus OUT_CUDA_ARCHS TGT_CUDA_ARCHS)
+  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
+    cuda_archs_loose_intersection(_archs "9.0a;10.0f;11.0f;12.0f" "${TGT_CUDA_ARCHS}")
+  else()
+    cuda_archs_loose_intersection(_archs "9.0a;10.0a;10.1a;10.3a;12.0a;12.1a" "${TGT_CUDA_ARCHS}")
+  endif()
+  set(${OUT_CUDA_ARCHS} ${_archs} PARENT_SCOPE)
+endfunction()
+
 #
 # Override the GPU architectures detected by cmake/torch and filter them by
 # `GPU_SUPPORTED_ARCHES`. Sets the final set of architectures in

csrc/cpu/cpu_attn_impl.hpp

@@ -822,8 +822,8 @@ struct AttentionInput {
       logits_buffer_t *__restrict__ logits_buffer,                          \
       float *__restrict__ partial_q_buffer, float *__restrict__ max_buffer, \
       float *__restrict__ sum_buffer, int32_t *__restrict__ block_table,    \
-      const int32_t kv_tile_start_pos, const int32_t kv_tile_end_pos,       \
-      const int32_t kv_tile_token_num,                                      \
+      const int32_t kv_end_pos, const int32_t kv_tile_start_pos,            \
+      const int32_t kv_tile_end_pos, const int32_t kv_tile_token_num,       \
       const int64_t kv_cache_num_blocks_stride, const int32_t q_head_num,   \
       const int32_t q_token_num, const int32_t q_tile_start_pos,            \
       const int32_t q_heads_per_kv, const int32_t block_size,               \
@@ -834,7 +834,7 @@ struct AttentionInput {
 
 #define CPU_ATTENTION_PARAMS                                                  \
   q_heads_buffer, k_head_cache_ptr, v_head_cache_ptr, logits_buffer,          \
-      partial_q_buffer, max_buffer, sum_buffer, block_table,                  \
+      partial_q_buffer, max_buffer, sum_buffer, block_table, kv_end_pos,      \
       kv_tile_start_pos, kv_tile_end_pos, kv_tile_token_num,                  \
       kv_cache_num_blocks_stride, q_head_num, q_token_num, q_tile_start_pos,  \
       q_heads_per_kv, block_size, left_window_size, right_window_size, scale, \
@@ -917,6 +917,7 @@ class AttentionMainLoop {
     //  - max_buffer: [MaxQHeadNumPerIteration, 1], store max logits
     //  - sum_buffer: [MaxQHeadNumPerIteration, 1], store sum of exp
     //  - block_table
+    //  - kv_end_pos: un-aligned end position of KV cache
     //  - kv_tile_start_pos: start position of KV cache, aligned to
     //  BlockSizeAlignment
     //  - kv_tile_end_pos: end position of KV cache, aligned to
@@ -1043,7 +1044,7 @@ class AttentionMainLoop {
         }
 
         apply_mask(logits_buffer, kv_tile_token_num, q_tile_start_pos,
-                   kv_tile_start_pos, kv_tile_end_pos, q_token_num,
+                   kv_end_pos, kv_tile_start_pos, kv_tile_end_pos, q_token_num,
                    q_heads_per_kv, left_window_size, right_window_size);
 
         // if (debug_info){
@@ -1126,7 +1127,7 @@ class AttentionMainLoop {
 
     void apply_mask(logits_buffer_t* __restrict__ logits_buffer,
                     const int64_t logits_buffer_stride,
-                    const int32_t q_tile_start_pos,
+                    const int32_t q_tile_start_pos, const int32_t kv_end_pos,
                     const int32_t kv_tile_start_pos,
                     const int32_t kv_tile_end_pos, const int32_t q_token_num,
                     const int32_t q_heads_per_kv,
@@ -1154,7 +1155,7 @@ class AttentionMainLoop {
                            std::max(kv_tile_start_pos,
                                     curr_token_pos + sliding_window_right + 1));
           }
-          return pos;
+          return std::min(pos, kv_end_pos);
         }();
 
         int32_t left_invalid_token_num = left_kv_pos - kv_tile_start_pos;
@@ -1789,7 +1790,7 @@ class AttentionMainLoop {
                   attn_impl.template execute_attention<Attention>(
                       curr_q_heads_buffer, curr_k_cache, curr_v_cache,
                       logits_buffer, curr_partial_q_buffer, curr_max_buffer,
-                      curr_sum_buffer, curr_block_table,
+                      curr_sum_buffer, curr_block_table, kv_end_pos,
                       aligned_actual_kv_tile_pos_left,
                       aligned_actual_kv_tile_pos_right, actual_kv_token_num,
                       kv_cache_block_num_stride, q_tile_head_num,

csrc/cpu/cpu_fused_moe.cpp

@@ -30,7 +30,12 @@
   }()
 
 namespace {
-enum class FusedMOEAct { SiluAndMul, SwigluOAIAndMul, GeluAndMul };
+enum class FusedMOEAct {
+  SiluAndMul,
+  SwigluOAIAndMul,
+  GeluAndMul,
+  GeluTanhAndMul,
+};
 
 FusedMOEAct get_act_type(const std::string& act) {
   if (act == "silu") {
@@ -39,6 +44,8 @@ FusedMOEAct get_act_type(const std::string& act) {
     return FusedMOEAct::SwigluOAIAndMul;
   } else if (act == "gelu") {
     return FusedMOEAct::GeluAndMul;
+  } else if (act == "gelu_tanh") {
+    return FusedMOEAct::GeluTanhAndMul;
   } else {
     TORCH_CHECK(false, "Invalid act type: " + act);
   }
@@ -143,6 +150,44 @@ void gelu_and_mul(float* __restrict__ input, scalar_t* __restrict__ output,
   }
 }
 
+template <typename scalar_t>
+void gelu_tanh_and_mul(float* __restrict__ input, scalar_t* __restrict__ output,
+                       const int32_t m_size, const int32_t n_size,
+                       const int32_t input_stride,
+                       const int32_t output_stride) {
+  using scalar_vec_t = typename cpu_utils::VecTypeTrait<scalar_t>::vec_t;
+  const int32_t dim = n_size / 2;
+  float* __restrict__ gate = input;
+  float* __restrict__ up = input + dim;
+  vec_op::FP32Vec16 one_vec(1.0);
+  vec_op::FP32Vec16 w1_vec(0.7978845608028654);
+  vec_op::FP32Vec16 w2_vec(0.5);
+  vec_op::FP32Vec16 w3_vec(0.044715);
+  alignas(64) float temp[16];
+
+  for (int32_t m = 0; m < m_size; ++m) {
+    for (int32_t n = 0; n < dim; n += 16) {
+      vec_op::FP32Vec16 gate_vec(gate + n);
+      vec_op::FP32Vec16 up_vec(up + n);
+      auto gate_pow3_vec = gate_vec * gate_vec * gate_vec;
+      auto inner_vec = w1_vec * (gate_vec + w3_vec * gate_pow3_vec);
+
+      inner_vec.save(temp);
+      for (int32_t i = 0; i < 16; ++i) {
+        temp[i] = std::tanh(temp[i]);
+      }
+      vec_op::FP32Vec16 tanh_vec(temp);
+      auto gelu_tanh = gate_vec * w2_vec * (one_vec + tanh_vec);
+      auto gated_output_fp32 = up_vec * gelu_tanh;
+      scalar_vec_t gated_output = scalar_vec_t(gated_output_fp32);
+      gated_output.save(output + n);
+    }
+    gate += input_stride;
+    up += input_stride;
+    output += output_stride;
+  }
+}
+
 template <typename scalar_t>
 FORCE_INLINE void apply_gated_act(const FusedMOEAct act,
                                   float* __restrict__ input,
@@ -160,6 +205,9 @@ FORCE_INLINE void apply_gated_act(const FusedMOEAct act,
     case FusedMOEAct::GeluAndMul:
       gelu_and_mul(input, output, m, n, input_stride, output_stride);
       return;
+    case FusedMOEAct::GeluTanhAndMul:
+      gelu_tanh_and_mul(input, output, m, n, input_stride, output_stride);
+      return;
     default:
       TORCH_CHECK(false, "Unsupported act type.");
   }