p

Signed-off-by: khluu <khluu000@gmail.com>

.buildkite/ci_config.yaml

@@ -8,7 +8,6 @@ run_all_patterns:
   - "CMakeLists.txt"
   - "requirements/common.txt"
   - "requirements/cuda.txt"
-  - "requirements/kv_connectors.txt"
   - "requirements/build/cuda.txt"
   - "requirements/test/cuda.txt"
   - "setup.py"

.buildkite/hardware_tests/cpu.yaml

@@ -12,19 +12,15 @@ steps:
   - vllm/_custom_ops.py
   - tests/kernels/attention/test_cpu_attn.py
   - tests/kernels/moe/test_cpu_fused_moe.py
-  - tests/kernels/moe/test_cpu_fp8_fused_moe.py
   - tests/kernels/test_onednn.py
   - tests/kernels/test_awq_int4_to_int8.py
-  - tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
   commands:
     - |
-      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
+      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 20m "
       pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
       pytest -x -v -s tests/kernels/moe/test_cpu_fused_moe.py
-      pytest -x -v -s tests/kernels/moe/test_cpu_fp8_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/test_awq_int4_to_int8.py"
 
 - label: CPU-Compatibility Tests
   depends_on: []
@@ -61,11 +57,10 @@ steps:
   source_file_dependencies:
   - csrc/cpu/
   - vllm/model_executor/layers/quantization/cpu_wna16.py
-  - vllm/model_executor/layers/quantization/auto_gptq.py
+  - vllm/model_executor/layers/quantization/gptq_marlin.py
   - vllm/model_executor/layers/quantization/compressed_tensors/schemes/compressed_tensors_w8a8_int8.py
   - vllm/model_executor/layers/quantization/kernels/scaled_mm/cpu.py
   - vllm/model_executor/layers/quantization/kernels/mixed_precision/cpu.py
-  - vllm/model_executor/layers/fused_moe/experts/cpu_moe.py
   - tests/quantization/test_compressed_tensors.py
   - tests/quantization/test_cpu_wna16.py
   commands:
@@ -74,11 +69,11 @@ steps:
       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"
       
-- label: CPU-Distributed Tests (PP+TP)
+- label: CPU-Distributed Tests
   depends_on: []
   device: intel_cpu
   no_plugin: true
-  source_file_dependencies: &cpu_distributed_deps
+  source_file_dependencies:
   - csrc/cpu/shm.cpp
   - vllm/v1/worker/cpu_worker.py
   - vllm/v1/worker/gpu_worker.py
@@ -87,21 +82,10 @@ steps:
   - vllm/platforms/cpu.py
   - vllm/distributed/parallel_state.py
   - vllm/distributed/device_communicators/cpu_communicator.py
-  - .buildkite/scripts/hardware_ci/run-cpu-distributed-smoke-test.sh
   commands:
     - |
       bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 10m "
-      bash .buildkite/scripts/hardware_ci/run-cpu-distributed-smoke-test.sh tp_pp"
-
-- label: CPU-Distributed Tests (DP+TP)
-  depends_on: []
-  device: intel_cpu
-  no_plugin: true
-  source_file_dependencies: *cpu_distributed_deps
-  commands:
-    - |
-      bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 10m "
-      bash .buildkite/scripts/hardware_ci/run-cpu-distributed-smoke-test.sh dp_tp"
+      bash .buildkite/scripts/hardware_ci/run-cpu-distributed-smoke-test.sh"
 
 - label: CPU-Multi-Modal Model Tests %N
   depends_on: []

.buildkite/image_build/image_build.sh

@@ -192,7 +192,6 @@ export BUILDKITE_COMMIT
 export PARENT_COMMIT
 export IMAGE_TAG
 export IMAGE_TAG_LATEST
-export COMMIT="${COMMIT:-${BUILDKITE_COMMIT}}"
 export CACHE_FROM
 export CACHE_FROM_BASE_BRANCH
 export CACHE_FROM_MAIN

.buildkite/intel_jobs/lora_intel.yaml

@@ -18,18 +18,17 @@ steps:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       'cd tests &&
-      export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       pytest -v -s lora/test_layers.py &&
       pytest -v -s lora/test_lora_checkpoints.py &&
-      pytest -v -s lora/test_lora_functions.py &&
+      (pytest -v -s lora/test_lora_functions.py --deselect="tests/lora/test_lora_functions.py::test_lora_functions_sync" --deselect="tests/lora/test_lora_functions.py::test_lora_functions_async" || true) &&
       pytest -v -s lora/test_lora_huggingface.py &&
       pytest -v -s lora/test_lora_manager.py &&
       pytest -v -s lora/test_lora_utils.py &&
       pytest -v -s lora/test_peft_helper.py &&
       pytest -v -s lora/test_resolver.py &&
       pytest -v -s lora/test_utils.py &&
-      pytest -v -s lora/test_add_lora.py  &&
-      pytest -v -s lora/test_worker.py'
+      (pytest -v -s lora/test_add_lora.py --deselect="tests/lora/test_add_lora.py::test_add_lora" || true) &&
+      (pytest -v -s lora/test_worker.py --deselect="tests/lora/test_worker.py::test_worker_apply_lora" || true)'
 
 - label: LoRA Fused/MoE Kernels
   timeout_in_minutes: 45
@@ -47,7 +46,6 @@ steps:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       'cd tests &&
-      export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       pytest -v -s lora/test_fused_moe_lora_kernel.py && 
       pytest -v -s lora/test_moe_lora_align_sum.py'
 
@@ -67,9 +65,8 @@ steps:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       'cd tests &&
-      export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       set -o pipefail &&
-      pytest -v -s lora/test_punica_ops.py --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-3-43264-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype1-1-2049-64-128-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-1-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-1-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-8-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype0-3-2049-128-8-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-8-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype1-1-2049-256-128-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-64256-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-2-29696-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-3-49408-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-2-16384-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-2-51328-32-4-4]"'
+      pytest -v -s lora/test_punica_ops.py --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-2-2049-64-32-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype1-2-64000-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-1-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-1-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-256-8-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype0-3-2049-128-8-16]" --deselect="tests/lora/test_punica_ops.py::test_kernels[shrink-0-xpu:0-dtype0-1-2049-128-8-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels[expand-0-xpu:0-dtype1-1-2049-256-128-32]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-3-64256-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-2-29696-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype1-3-49408-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[shrink-0-xpu:0-dtype0-2-16384-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype0-2-51328-32-4-4]" --deselect="tests/lora/test_punica_ops.py::test_kernels_hidden_size[expand-0-xpu:0-dtype1-1-102656-32-4-4]"'
 
 - label: LoRA Punica FP8/XPU Ops
   timeout_in_minutes: 45
@@ -87,7 +84,6 @@ steps:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       'cd tests &&
-      export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       pytest -v -s lora/test_punica_ops_fp8.py &&
       pytest -v -s lora/test_punica_xpu_ops.py'
 
@@ -107,12 +103,10 @@ steps:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       'cd tests &&
-      export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       (pytest -v -s lora/test_mixtral.py --deselect="tests/lora/test_mixtral.py::test_mixtral_lora[4]" || true) &&
       pytest -v -s lora/test_quant_model.py --deselect="tests/lora/test_quant_model.py::test_quant_model_lora[model0]" --deselect="tests/lora/test_quant_model.py::test_quant_model_lora[model1]" --deselect="tests/lora/test_quant_model.py::test_quant_model_tp_equality[model0]" &&
-      pytest -v -s lora/test_transformers_model.py &&
-      pytest -v -s lora/test_chatglm3_tp.py &&
-      pytest -s -v lora/test_minicpmv_tp.py'
+      pytest -v -s lora/test_qwen35_densemodel_lora.py &&
+      pytest -v -s lora/test_transformers_model.py'
 
 - label: LoRA Multimodal
   timeout_in_minutes: 45
@@ -130,6 +124,7 @@ steps:
     - >-
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       'cd tests &&
-      export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       pytest -v -s lora/test_default_mm_loras.py && 
+      (pytest -v -s lora/test_qwen3_unembed.py || true) &&
+      (pytest -v -s lora/test_qwenvl.py || true) &&
       pytest -v -s lora/test_whisper.py'

.buildkite/intel_jobs/misc_intel.yaml

@@ -49,7 +49,7 @@ steps:
       bash .buildkite/scripts/hardware_ci/run-intel-test.sh
       'export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
       cd tests &&
-      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/logits_processors &&
       pytest -v -s v1/test_oracle.py &&
       pytest -v -s v1/test_request.py &&
       pytest -v -s v1/test_outputs.py'

.buildkite/intel_jobs/test-intel.yaml

@@ -36,7 +36,6 @@ steps:
         python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --attention-backend=TRITON_ATTN &&
         python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --quantization fp8 &&
         python3 examples/basic/offline_inference/generate.py --model facebook/opt-125m --block-size 64 --enforce-eager --kv-cache-dtype fp8 &&
-        python3 examples/basic/offline_inference/generate.py --model nvidia/Llama-3.1-8B-Instruct-FP8 --block-size 64 --enforce-eager --quantization modelopt --kv-cache-dtype fp8 --attention-backend TRITON_ATTN --max-model-len 4096 &&
         python3 examples/basic/offline_inference/generate.py --model superjob/Qwen3-4B-Instruct-2507-GPTQ-Int4 --block-size 64 --enforce-eager --max-model-len 8192 &&
         python3 examples/basic/offline_inference/generate.py --model 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'
@@ -62,5 +61,5 @@ steps:
         pytest -v -s v1/worker --ignore=v1/worker/test_gpu_model_runner.py --ignore=v1/worker/test_worker_memory_snapshot.py &&
         pytest -v -s v1/structured_output &&
         pytest -v -s v1/test_serial_utils.py &&
-        pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py &&
-        pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --ignore=v1/kv_connector/unit/test_lmcache_integration.py --ignore=v1/kv_connector/unit/test_hf3fs_client.py --ignore=v1/kv_connector/unit/test_hf3fs_connector.py --ignore=v1/kv_connector/unit/test_hf3fs_metadata_server.py --ignore=v1/kv_connector/unit/test_offloading_connector.py'
+        pytest -v -s v1/spec_decode --ignore=v1/spec_decode/test_max_len.py --ignore=v1/spec_decode/test_tree_attention.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.py &&
+        pytest -v -s v1/kv_connector/unit --ignore=v1/kv_connector/unit/test_multi_connector.py --ignore=v1/kv_connector/unit/test_example_connector.py --ignore=v1/kv_connector/unit/test_lmcache_integration.py --ignore=v1/kv_connector/unit/test_hf3fs_client.py --ignore=v1/kv_connector/unit/test_hf3fs_connector.py --ignore=v1/kv_connector/unit/test_hf3fs_metadata_server.py'

.buildkite/lm-eval-harness/run-lm-eval-chartqa-vllm-vlm-baseline.sh

@@ -2,7 +2,7 @@
 # We can use this script to compute baseline accuracy on chartqa for vllm.
 #
 # Make sure you have lm-eval-harness installed:
-#   pip install "lm-eval[api]>=0.4.12"
+#   pip install "lm-eval[api]>=0.4.11"
 
 usage() {
     echo``

.buildkite/lm-eval-harness/run-lm-eval-gsm-hf-baseline.sh

@@ -2,7 +2,7 @@
 # We can use this script to compute baseline accuracy on GSM for transformers.
 #
 # Make sure you have lm-eval-harness installed:
-#   pip install "lm-eval[api]>=0.4.12"
+#   pip install "lm-eval[api]>=0.4.11"
 
 usage() {
     echo``

.buildkite/lm-eval-harness/run-lm-eval-gsm-vllm-baseline.sh

@@ -3,7 +3,7 @@
 # We use this for fp8, which HF does not support.
 #
 # Make sure you have lm-eval-harness installed:
-#   pip install "lm-eval[api]>=0.4.12"
+#   pip install "lm-eval[api]>=0.4.11"
 
 usage() {
     echo``

.buildkite/lm-eval-harness/run-lm-eval-mmlupro-vllm-baseline.sh

@@ -3,7 +3,7 @@
 # We use this for fp8, which HF does not support.
 #
 # Make sure you have lm-eval-harness installed:
-#   pip install "lm-eval[api]>=0.4.12"
+#   pip install "lm-eval[api]>=0.4.11"
 
 usage() {
     echo``

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

@@ -31,9 +31,30 @@
     }
   },
   "tests": [
+    {
+      "test_name": "serving_llama8B_tp1_sharegpt",
+      "server_parameters": {
+        "tensor_parallel_size": 1
+      },
+      "client_parameters": {
+        "dataset_name": "sharegpt",
+        "dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json"
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp2_sharegpt",
+      "server_parameters": {
+        "tensor_parallel_size": 2
+      },
+      "client_parameters": {
+        "dataset_name": "sharegpt",
+        "dataset_path": "./ShareGPT_V3_unfiltered_cleaned_split.json"
+      }
+    },
     {
       "test_name": "serving_llama8B_tp1_random_128_128",
       "server_parameters": {
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "dataset_name": "random",
@@ -41,10 +62,158 @@
         "random-output-len": 128
       }
     },
+    {
+      "test_name": "serving_llama8B_tp2_random_128_128",
+      "server_parameters": {
+        "tensor_parallel_size": 2
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp4_random_128_128",
+      "server_parameters": {
+        "tensor_parallel_size": 4
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp1_random_128_2048",
+      "server_parameters": {
+        "tensor_parallel_size": 1
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 2048
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp2_random_128_2048",
+      "server_parameters": {
+        "tensor_parallel_size": 2
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 2048
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp4_random_128_2048",
+      "server_parameters": {
+        "tensor_parallel_size": 4
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 2048
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp1_random_2048_128",
+      "server_parameters": {
+        "tensor_parallel_size": 1
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 2048,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp2_random_2048_128",
+      "server_parameters": {
+        "tensor_parallel_size": 2
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 2048,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp4_random_2048_128",
+      "server_parameters": {
+        "tensor_parallel_size": 4
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 2048,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp1_random_2048_2048",
+      "server_parameters": {
+        "tensor_parallel_size": 1
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 2048,
+        "random-output-len": 2048
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp2_random_2048_2048",
+      "server_parameters": {
+        "tensor_parallel_size": 2
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 2048,
+        "random-output-len": 2048
+      }
+    },
+    {
+      "test_name": "serving_llama8B_tp4_random_2048_2048",
+      "server_parameters": {
+        "tensor_parallel_size": 4
+      },
+      "client_parameters": {
+        "dataset_name": "random",
+        "random-input-len": 2048,
+        "random-output-len": 2048
+      }
+    },
     {
       "test_name": "serving_llama8B_int4_tp1_random_128_128",
       "server_parameters": {
-        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4"
+        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
+        "tensor_parallel_size": 1
+      },
+      "client_parameters": {
+        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_int4_tp2_random_128_128",
+      "server_parameters": {
+        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
+        "tensor_parallel_size": 2
+      },
+      "client_parameters": {
+        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_int4_tp4_random_128_128",
+      "server_parameters": {
+        "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
+        "tensor_parallel_size": 4
       },
       "client_parameters": {
         "model": "hugging-quants/Meta-Llama-3.1-8B-Instruct-AWQ-INT4",
@@ -56,7 +225,8 @@
     {
       "test_name": "serving_llama8B_int8_tp1_random_128_128",
       "server_parameters": {
-        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8"
+        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
@@ -66,12 +236,26 @@
       }
     },
     {
-      "test_name": "serving_llama1B_tp1_random_128_128",
+      "test_name": "serving_llama8B_int8_tp2_random_128_128",
       "server_parameters": {
-        "model": "meta-llama/Llama-3.2-1B"
+        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
+        "tensor_parallel_size": 2
       },
       "client_parameters": {
-        "model": "meta-llama/Llama-3.2-1B",
+        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
+        "dataset_name": "random",
+        "random-input-len": 128,
+        "random-output-len": 128
+      }
+    },
+    {
+      "test_name": "serving_llama8B_int8_tp4_random_128_128",
+      "server_parameters": {
+        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
+        "tensor_parallel_size": 4
+      },
+      "client_parameters": {
+        "model": "RedHatAI/Meta-Llama-3.1-8B-Instruct-quantized.w8a8",
         "dataset_name": "random",
         "random-input-len": 128,
         "random-output-len": 128
@@ -80,7 +264,8 @@
     {
       "test_name": "serving_llama3B_tp1_random_128_128",
       "server_parameters": {
-        "model": "meta-llama/Llama-3.2-3B-Instruct"
+        "model": "meta-llama/Llama-3.2-3B-Instruct",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "meta-llama/Llama-3.2-3B-Instruct",
@@ -89,22 +274,11 @@
         "random-output-len": 128
       }
     },
-    {
-      "test_name": "serving_llama70B_tp1_random_128_128",
-      "server_parameters": {
-        "model": "meta-llama/Llama-3.3-70B-Instruct"
-      },
-      "client_parameters": {
-        "model": "meta-llama/Llama-3.3-70B-Instruct",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
     {
       "test_name": "serving_granite2B_tp1_random_128_128",
       "server_parameters": {
-        "model": "ibm-granite/granite-3.2-2b-instruct"
+        "model": "ibm-granite/granite-3.2-2b-instruct",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "ibm-granite/granite-3.2-2b-instruct",
@@ -116,7 +290,8 @@
     {
       "test_name": "serving_qwen1.7B_tp1_random_128_128",
       "server_parameters": {
-        "model": "Qwen/Qwen3-1.7B"
+        "model": "Qwen/Qwen3-1.7B",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "Qwen/Qwen3-1.7B",
@@ -128,7 +303,8 @@
     {
       "test_name": "serving_qwen4B_tp1_random_128_128",
       "server_parameters": {
-        "model": "Qwen/Qwen3-4B"
+        "model": "Qwen/Qwen3-4B",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "Qwen/Qwen3-4B",
@@ -140,7 +316,8 @@
     {
       "test_name": "serving_qwen8B_tp1_random_128_128",
       "server_parameters": {
-        "model": "Qwen/Qwen3-8B"
+        "model": "Qwen/Qwen3-8B",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "Qwen/Qwen3-8B",
@@ -149,34 +326,11 @@
         "random-output-len": 128
       }
     },
-    {
-      "test_name": "serving_qwen14B_tp1_random_128_128",
-      "server_parameters": {
-        "model": "Qwen/Qwen3-14B"
-      },
-      "client_parameters": {
-        "model": "Qwen/Qwen3-14B",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_qwen30B_tp1_random_128_128",
-      "server_parameters": {
-        "model": "Qwen/Qwen3-30B-A3B"
-      },
-      "client_parameters": {
-        "model": "Qwen/Qwen3-30B-A3B",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
     {
       "test_name": "serving_glm9B_tp1_random_128_128",
       "server_parameters": {
-        "model": "zai-org/glm-4-9b-hf"
+        "model": "zai-org/glm-4-9b-hf",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "zai-org/glm-4-9b-hf",
@@ -188,7 +342,8 @@
     {
       "test_name": "serving_gemma7B_tp1_random_128_128",
       "server_parameters": {
-        "model": "google/gemma-7b"
+        "model": "google/gemma-7b",
+        "tensor_parallel_size": 1
       },
       "client_parameters": {
         "model": "google/gemma-7b",
@@ -196,94 +351,6 @@
         "random-input-len": 128,
         "random-output-len": 128
       }
-    },
-    {
-      "test_name": "serving_gemma3-4b_tp1_random_128_128",
-      "server_environment_variables": {
-	"VLLM_CPU_SGL_KERNEL": 0
-      },
-      "server_parameters": {
-        "model": "google/gemma-3-4b-it"
-      },
-      "client_parameters": {
-        "model": "google/gemma-3-4b-it",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_gemma3-12b_tp1_random_128_128",
-      "server_environment_variables": {
-	"VLLM_CPU_SGL_KERNEL": 0
-      },
-      "server_parameters": {
-        "model": "google/gemma-3-12b-it"
-      },
-      "client_parameters": {
-        "model": "google/gemma-3-12b-it",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_gemma4-4b_tp1_random_128_128",
-      "server_environment_variables": {
-	"VLLM_CPU_SGL_KERNEL": 0
-      },
-      "server_parameters": {
-        "model": "google/gemma-4-E4B-it"
-      },
-      "client_parameters": {
-        "model": "google/gemma-4-E4B-it",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_gemma4-2b_tp1_random_128_128",
-      "server_environment_variables": {
-	"VLLM_CPU_SGL_KERNEL": 0
-      },
-      "server_parameters": {
-        "model": "google/gemma-4-E2B-it"
-      },
-      "client_parameters": {
-        "model": "google/gemma-4-E2B-it",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_gemma4-26b_tp1_random_128_128",
-      "server_environment_variables": {
-	"VLLM_CPU_SGL_KERNEL": 0,
-        "VLLM_CPU_ATTN_SPLIT_KV": 0
-      },
-      "server_parameters": {
-        "model": "google/gemma-4-26B-A4B-it"
-      },
-      "client_parameters": {
-        "model": "google/gemma-4-26B-A4B-it",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
-    },
-    {
-      "test_name": "serving_phi4_tp1_random_128_128",
-      "server_parameters": {
-        "model": "microsoft/Phi-4-reasoning"
-      },
-      "client_parameters": {
-        "model": "microsoft/Phi-4-reasoning",
-        "dataset_name": "random",
-        "random-input-len": 128,
-        "random-output-len": 128
-      }
     }
   ]
 }

.buildkite/release-pipeline.yaml

@@ -27,8 +27,7 @@ steps:
           - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=12.9.1 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64_CU129}\" --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
           - "mkdir artifacts"
           - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
-          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
+          - "bash .buildkite/scripts/upload-nightly-wheels.sh manylinux_2_31"
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -38,11 +37,10 @@ steps:
         agents:
           queue: arm64_cpu_queue_release
         commands:
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg BUILD_OS=manylinux --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda13.0 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04  --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
           - "mkdir artifacts"
           - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
-          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
+          - "bash .buildkite/scripts/upload-nightly-wheels.sh manylinux_2_35"
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -55,8 +53,7 @@ steps:
           - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg GIT_REPO_CHECK=1 --build-arg VLLM_BUILD_ACL=ON --tag vllm-ci:build-image --target vllm-build --progress plain -f docker/Dockerfile.cpu ."
           - "mkdir artifacts"
           - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
-          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
+          - "bash .buildkite/scripts/upload-nightly-wheels.sh manylinux_2_35"
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -69,8 +66,7 @@ steps:
           - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=12.9.1 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86_CU129}\" --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
           - "mkdir artifacts"
           - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
-          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
+          - "bash .buildkite/scripts/upload-nightly-wheels.sh manylinux_2_31"
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -80,11 +76,10 @@ steps:
         agents:
           queue: cpu_queue_release
         commands:
-          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86}\" --build-arg BUILD_OS=manylinux --build-arg BUILD_BASE_IMAGE=pytorch/manylinux2_28-builder:cuda13.0 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86}\" --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04 --tag vllm-ci:build-image --target build --progress plain -f docker/Dockerfile ."
           - "mkdir artifacts"
           - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
-          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
+          - "bash .buildkite/scripts/upload-nightly-wheels.sh manylinux_2_35"
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -97,8 +92,7 @@ steps:
           - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg GIT_REPO_CHECK=1 --build-arg VLLM_CPU_X86=true --tag vllm-ci:build-image --target vllm-build --progress plain -f docker/Dockerfile.cpu ."
           - "mkdir artifacts"
           - "docker run --rm -v $(pwd)/artifacts:/artifacts_host vllm-ci:build-image bash -c 'cp -r dist /artifacts_host && chmod -R a+rw /artifacts_host'"
-          - "bash .buildkite/scripts/upload-nightly-wheels.sh"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "s3://vllm-wheels/$$BUILDKITE_COMMIT/$(cd artifacts/dist && echo *.whl)"'
+          - "bash .buildkite/scripts/upload-nightly-wheels.sh manylinux_2_35"
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -127,24 +121,11 @@ steps:
           queue: cpu_queue_release
         commands:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=13.0.2 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_X86}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04 \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86}\" --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m) --target vllm-openai --progress plain -f docker/Dockerfile ."
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)"
           # re-tag to default image tag and push, just in case arm64 build fails
           - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m) public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)"'
 
       - label: "Build release image - aarch64 - CUDA 13.0"
         depends_on: ~
@@ -153,21 +134,8 @@ steps:
           queue: arm64_cpu_queue_release
         commands:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=13.0.2 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_AARCH64}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04 \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu22.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m) --target vllm-openai --progress plain -f docker/Dockerfile ."
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)"'
 
       - label: "Build release image - x86_64 - CUDA 12.9"
         depends_on: ~
@@ -176,23 +144,11 @@ steps:
           queue: cpu_queue_release
         commands:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh cu129) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=12.9.1 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_X86_CU129}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=12.9.1 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86_CU129}\" --build-arg INSTALL_KV_CONNECTORS=true --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129 --target vllm-openai --progress plain -f docker/Dockerfile ."
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129"
           # re-tag to default image tag and push, just in case arm64 build fails
           - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129"
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129"'
 
       - label: "Build release image - aarch64 - CUDA 12.9"
         depends_on: ~
@@ -201,20 +157,8 @@ steps:
           queue: arm64_cpu_queue_release
         commands:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh cu129) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=12.9.1 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_AARCH64_CU129}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=12.9.1 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64_CU129}\" --build-arg INSTALL_KV_CONNECTORS=true --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129 --target vllm-openai --progress plain -f docker/Dockerfile ."
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129"'
 
       - label: "Build release image - x86_64 - CUDA 13.0 - Ubuntu 24.04"
         depends_on: ~
@@ -223,25 +167,10 @@ steps:
           queue: cpu_queue_release
         commands:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh ubuntu2404) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=13.0.2 \
-              --build-arg UBUNTU_VERSION=24.04 \
-              --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_X86}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu24.04 \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg UBUNTU_VERSION=24.04 --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86}\" --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu24.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"
           - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"'
 
       - label: "Build release image - aarch64 - CUDA 13.0 - Ubuntu 24.04"
         depends_on: ~
@@ -250,23 +179,8 @@ steps:
           queue: arm64_cpu_queue_release
         commands:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh ubuntu2404) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=13.0.2 \
-              --build-arg UBUNTU_VERSION=24.04 \
-              --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_AARCH64}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu24.04 \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.2 --build-arg UBUNTU_VERSION=24.04 --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64}\" --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.2-devel-ubuntu24.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"'
 
       - label: "Build release image - x86_64 - CUDA 12.9 - Ubuntu 24.04"
         depends_on: ~
@@ -275,24 +189,10 @@ steps:
           queue: cpu_queue_release
         commands:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh cu129-ubuntu2404) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=12.9.1 \
-              --build-arg UBUNTU_VERSION=24.04 \
-              --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_X86_CU129}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=12.9.1 --build-arg UBUNTU_VERSION=24.04 --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_X86_CU129}\" --build-arg INSTALL_KV_CONNECTORS=true --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"
           - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"'
 
       - label: "Build release image - aarch64 - CUDA 12.9 - Ubuntu 24.04"
         depends_on: ~
@@ -301,29 +201,14 @@ steps:
           queue: arm64_cpu_queue_release
         commands:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
-          - |
-            DOCKER_BUILDKIT=1 docker build \
-              $(bash .buildkite/scripts/docker-build-metadata-args.sh cu129-ubuntu2404) \
-              --build-arg max_jobs=16 \
-              --build-arg USE_SCCACHE=1 \
-              --build-arg GIT_REPO_CHECK=1 \
-              --build-arg CUDA_VERSION=12.9.1 \
-              --build-arg UBUNTU_VERSION=24.04 \
-              --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 \
-              --build-arg torch_cuda_arch_list="${CUDA_ARCH_AARCH64_CU129}" \
-              --build-arg INSTALL_KV_CONNECTORS=true \
-              --target vllm-openai \
-              --progress plain \
-              -f docker/Dockerfile .
+          - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=12.9.1 --build-arg UBUNTU_VERSION=24.04 --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 --build-arg torch_cuda_arch_list=\"${CUDA_ARCH_AARCH64_CU129}\" --build-arg INSTALL_KV_CONNECTORS=true --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
           - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"'
 
       - block: "Build release image for x86_64 CPU"
         key: block-cpu-release-image-build
         depends_on: ~
 
       - label: "Build release image - x86_64 - CPU"
-        key: build-cpu-release-image-x86
         depends_on:
           - block-cpu-release-image-build
           - input-release-version
@@ -334,7 +219,6 @@ steps:
           - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg GIT_REPO_CHECK=1 --build-arg VLLM_CPU_X86=true --tag public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:$(buildkite-agent meta-data get release-version) --tag public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:latest --progress plain --target vllm-openai -f docker/Dockerfile.cpu ."
           - "docker push public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:latest"
           - "docker push public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:$(buildkite-agent meta-data get release-version)"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:$(buildkite-agent meta-data get release-version)"'
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -343,8 +227,7 @@ steps:
         depends_on: ~
 
       - label: "Build release image - arm64 - CPU"
-        key: build-cpu-release-image-arm64
-        depends_on:
+        depends_on: 
           - block-arm64-cpu-release-image-build
           - input-release-version
         agents:
@@ -354,7 +237,6 @@ steps:
           - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg GIT_REPO_CHECK=1 --tag public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:$(buildkite-agent meta-data get release-version) --tag public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:latest --progress plain --target vllm-openai -f docker/Dockerfile.cpu ."
           - "docker push public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:latest"
           - "docker push public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:$(buildkite-agent meta-data get release-version)"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:$(buildkite-agent meta-data get release-version)"'
         env:
           DOCKER_BUILDKIT: "1"
 
@@ -372,7 +254,15 @@ steps:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
           - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64 --amend"
           - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 13.0" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT"'
+
+      - label: "Annotate release workflow - CUDA 13.0"
+        depends_on:
+          - create-multi-arch-manifest
+        id: annotate-release-workflow
+        agents:
+          queue: small_cpu_queue_release
+        commands:
+          - "bash .buildkite/scripts/annotate-release.sh"
 
       - label: "Create multi-arch manifest - CUDA 12.9"
         depends_on:
@@ -385,7 +275,6 @@ steps:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
           - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-cu129 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-cu129 --amend"
           - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 12.9" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129"'
 
       - label: "Create multi-arch manifest - CUDA 13.0 - Ubuntu 24.04"
         depends_on:
@@ -398,7 +287,6 @@ steps:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
           - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-ubuntu2404 --amend"
           - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 13.0 Ubuntu 24.04" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-ubuntu2404"'
 
       - label: "Create multi-arch manifest - CUDA 12.9 - Ubuntu 24.04"
         depends_on:
@@ -411,7 +299,6 @@ steps:
           - "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
           - "docker manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-cu129-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-cu129-ubuntu2404 --amend"
           - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"
-          - 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 12.9 Ubuntu 24.04" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"'
 
       - label: "Publish nightly multi-arch image to DockerHub"
         depends_on:
@@ -449,6 +336,24 @@ steps:
           DOCKER_BUILDKIT: "1"
           DOCKERHUB_USERNAME: "vllmbot"
 
+  - group: "Publish wheels"
+    key: "publish-wheels"
+    steps:
+      - block: "Confirm update release wheels to PyPI (experimental, use with caution)?"
+        key: block-upload-release-wheels
+        depends_on:
+          - input-release-version
+          - build-wheels
+
+      - label: "Upload release wheels to PyPI"
+        depends_on:
+          - block-upload-release-wheels
+        id: upload-release-wheels
+        agents:
+          queue: small_cpu_queue_release
+        commands:
+          - "bash .buildkite/scripts/upload-release-wheels-pypi.sh"
+
   # =============================================================================
   # ROCm Release Pipeline (x86_64 only)
   # =============================================================================
@@ -562,7 +467,7 @@ steps:
           echo ""
           echo " Build complete - Image and wheels cached"
         fi
-
+        
     artifact_paths:
       - "artifacts/rocm-base-wheels/*.whl"
     env:
@@ -718,7 +623,7 @@ steps:
       - "bash tools/vllm-rocm/generate-rocm-wheels-root-index.sh"
     env:
       S3_BUCKET: "vllm-wheels"
-      VARIANT: "rocm722"
+      VARIANT: "rocm721"
 
   # ROCm Job 6: Build ROCm Release Docker Image
   - label: ":docker: Build release image - x86_64 - ROCm"
@@ -778,7 +683,7 @@ steps:
         
         # Push to ECR
         docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rocm
-
+        
         echo ""
         echo " Successfully built and pushed ROCm release image"
         echo "   Image: public.ecr.aws/q9t5s3a7/vllm-release-repo:$${BUILDKITE_COMMIT}-rocm"
@@ -805,60 +710,3 @@ steps:
     env:
       DOCKER_BUILDKIT: "1"
       DOCKERHUB_USERNAME: "vllmbot"
-
-  # =============================================================================
-  # Publish to DockerHub and PyPI (at the end so all builds complete first)
-  # =============================================================================
-
-  - block: "Publish release images to DockerHub"
-    key: block-publish-release-images
-    depends_on:
-      - create-multi-arch-manifest
-      - create-multi-arch-manifest-cuda-12-9
-      - create-multi-arch-manifest-ubuntu2404
-      - create-multi-arch-manifest-cuda-12-9-ubuntu2404
-      - build-rocm-release-image
-      - input-release-version
-      # Wait for CPU builds if their block steps were unblocked, so publish
-      # doesn't race the in-progress CPU build. allow_failure lets publish
-      # proceed when the operator legitimately leaves the CPU block steps
-      # unblocked or the CPU build fails.
-      - step: build-cpu-release-image-x86
-        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:
-      - block-publish-release-images
-    key: publish-release-images-dockerhub
-    agents:
-      queue: small_cpu_queue_release
-    commands:
-      - "bash .buildkite/scripts/publish-release-images.sh"
-    plugins:
-      - docker-login#v3.0.0:
-          username: vllmbot
-          password-env: DOCKERHUB_TOKEN
-    env:
-      DOCKER_BUILDKIT: "1"
-      DOCKERHUB_USERNAME: "vllmbot"
-
-  - group: "Publish wheels"
-    key: "publish-wheels"
-    steps:
-      - block: "Confirm update release wheels to PyPI (experimental, use with caution)?"
-        key: block-upload-release-wheels
-        depends_on:
-          - input-release-version
-          - build-wheels
-
-      - label: "Upload release wheels to PyPI"
-        depends_on:
-          - block-upload-release-wheels
-        id: upload-release-wheels
-        agents:
-          queue: small_cpu_queue_release
-        commands:
-          - "bash .buildkite/scripts/upload-release-wheels-pypi.sh"

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

@@ -1,9 +0,0 @@
-#!/bin/bash
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-#
-# Append a build artifact line to the Buildkite annotation.
-# Usage: annotate-build-artifact.sh <label> <value>
-set -e
-echo "- **${1}**: \`${2}\`" | \
-  buildkite-agent annotate --append --style 'info' --context 'release-artifacts'

.buildkite/scripts/annotate-release.sh

@@ -0,0 +1,119 @@
+#!/bin/bash
+
+set -ex
+
+# Get release version, default to 1.0.0.dev for nightly/per-commit builds
+RELEASE_VERSION=$(buildkite-agent meta-data get release-version 2>/dev/null | sed 's/^v//')
+if [ -z "${RELEASE_VERSION}" ]; then
+  RELEASE_VERSION="1.0.0.dev"
+fi
+
+ROCM_BASE_CACHE_KEY=$(.buildkite/scripts/cache-rocm-base-wheels.sh key)
+
+buildkite-agent annotate --style 'info' --context 'release-workflow' << EOF
+To download the wheel (by commit):
+\`\`\`
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}-cp38-abi3-manylinux_2_35_x86_64.whl .
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}-cp38-abi3-manylinux_2_35_aarch64.whl .
+
+(Optional) For CUDA 12.9:
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}+cu129-cp38-abi3-manylinux_2_31_x86_64.whl .
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}+cu129-cp38-abi3-manylinux_2_31_aarch64.whl .
+
+(Optional) For CPU:
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}+cpu-cp38-abi3-manylinux_2_35_x86_64.whl .
+aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}+cpu-cp38-abi3-manylinux_2_35_aarch64.whl .
+\`\`\`
+
+
+To download and upload the image:
+
+\`\`\`
+# Download images:
+
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-x86_64
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-aarch64
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-x86_64-cu129
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-aarch64-cu129
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm-base
+docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-rocm
+docker pull public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:v${RELEASE_VERSION}
+docker pull public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:v${RELEASE_VERSION}
+
+# Tag and push images:
+
+## CUDA
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-x86_64 vllm/vllm-openai:x86_64
+docker tag vllm/vllm-openai:x86_64 vllm/vllm-openai:latest-x86_64
+docker tag vllm/vllm-openai:x86_64 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64
+docker push vllm/vllm-openai:latest-x86_64
+docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-x86_64-cu129 vllm/vllm-openai:x86_64-cu129
+docker tag vllm/vllm-openai:x86_64-cu129 vllm/vllm-openai:latest-x86_64-cu129
+docker tag vllm/vllm-openai:x86_64-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129
+docker push vllm/vllm-openai:latest-x86_64-cu129
+docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-aarch64 vllm/vllm-openai:aarch64
+docker tag vllm/vllm-openai:aarch64 vllm/vllm-openai:latest-aarch64
+docker tag vllm/vllm-openai:aarch64 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64
+docker push vllm/vllm-openai:latest-aarch64
+docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-aarch64-cu129 vllm/vllm-openai:aarch64-cu129
+docker tag vllm/vllm-openai:aarch64-cu129 vllm/vllm-openai:latest-aarch64-cu129
+docker tag vllm/vllm-openai:aarch64-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129
+docker push vllm/vllm-openai:latest-aarch64-cu129
+docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129
+
+## ROCm
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-rocm vllm/vllm-openai-rocm:${BUILDKITE_COMMIT}
+docker tag vllm/vllm-openai-rocm:${BUILDKITE_COMMIT} vllm/vllm-openai-rocm:latest
+docker tag vllm/vllm-openai-rocm:${BUILDKITE_COMMIT} vllm/vllm-openai-rocm:v${RELEASE_VERSION}
+docker push vllm/vllm-openai-rocm:latest
+docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm-base vllm/vllm-openai-rocm:${BUILDKITE_COMMIT}-base
+docker tag vllm/vllm-openai-rocm:${BUILDKITE_COMMIT}-base vllm/vllm-openai-rocm:latest-base
+docker tag vllm/vllm-openai-rocm:${BUILDKITE_COMMIT}-base vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
+docker push vllm/vllm-openai-rocm:latest-base
+docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
+
+## CPU
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:v${RELEASE_VERSION} vllm/vllm-openai-cpu:x86_64
+docker tag vllm/vllm-openai-cpu:x86_64 vllm/vllm-openai-cpu:latest-x86_64
+docker tag vllm/vllm-openai-cpu:x86_64 vllm/vllm-openai-cpu:v${RELEASE_VERSION}-x86_64
+docker push vllm/vllm-openai-cpu:latest-x86_64
+docker push vllm/vllm-openai-cpu:v${RELEASE_VERSION}-x86_64
+
+docker tag public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:v${RELEASE_VERSION} vllm/vllm-openai-cpu:arm64
+docker tag vllm/vllm-openai-cpu:arm64 vllm/vllm-openai-cpu:latest-arm64
+docker tag vllm/vllm-openai-cpu:arm64 vllm/vllm-openai-cpu:v${RELEASE_VERSION}-arm64
+docker push vllm/vllm-openai-cpu:latest-arm64
+docker push vllm/vllm-openai-cpu:v${RELEASE_VERSION}-arm64
+
+# Create multi-arch manifest:
+
+docker manifest rm vllm/vllm-openai:latest
+docker manifest create vllm/vllm-openai:latest vllm/vllm-openai:latest-x86_64 vllm/vllm-openai:latest-aarch64
+docker manifest create vllm/vllm-openai:v${RELEASE_VERSION} vllm/vllm-openai:v${RELEASE_VERSION}-x86_64 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64
+docker manifest push vllm/vllm-openai:latest
+docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}
+
+docker manifest rm vllm/vllm-openai:latest-cu129
+docker manifest create vllm/vllm-openai:latest-cu129 vllm/vllm-openai:latest-x86_64-cu129 vllm/vllm-openai:latest-aarch64-cu129
+docker manifest create vllm/vllm-openai:v${RELEASE_VERSION}-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129
+docker manifest push vllm/vllm-openai:latest-cu129
+docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}-cu129
+
+docker manifest rm vllm/vllm-openai-cpu:latest || true
+docker manifest create vllm/vllm-openai-cpu:latest vllm/vllm-openai-cpu:latest-x86_64 vllm/vllm-openai-cpu:latest-arm64
+docker manifest create vllm/vllm-openai-cpu:v${RELEASE_VERSION} vllm/vllm-openai-cpu:v${RELEASE_VERSION}-x86_64 vllm/vllm-openai-cpu:v${RELEASE_VERSION}-arm64
+docker manifest push vllm/vllm-openai-cpu:latest
+docker manifest push vllm/vllm-openai-cpu:v${RELEASE_VERSION}
+\`\`\`
+EOF

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

@@ -1,55 +0,0 @@
-#!/bin/bash
-# Usage: ./ci-fetch-log.sh <buildkite_job_url> [output_file]
-#        ./ci-fetch-log.sh <build_number> <job_uuid> [output_file]
-#
-# 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.
-#
-# 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>.
-
-set -euo pipefail
-
-ORG="vllm"
-PIPELINE="ci"
-
-usage() {
-    echo "Usage: $0 <buildkite_job_url> [output_file]"
-    echo "       $0 <build_number> <job_uuid> [output_file]"
-    exit 1
-}
-
-if [ $# -lt 1 ]; then usage; fi
-
-if [[ "$1" == https://* ]]; then
-    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
-    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
-    usage
-fi
-
-COOKIES=$(mktemp)
-trap 'rm -f "$COOKIES"' EXIT
-
-# Buildkite issues a session cookie on first hit; subsequent /download needs it.
-curl -fsSL -c "$COOKIES" -A "vllm-ci-fetch-log" \
-    "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"
-
-bash "$(dirname "$0")/ci-clean-log.sh" "$OUT"
-
-echo "$OUT"

.buildkite/scripts/detect-manylinux-tag.py

@@ -1,142 +0,0 @@
-#!/usr/bin/env python3
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""Detect the manylinux platform tag for a wheel and rename it in place.
-
-vLLM's build images produce wheels with the generic ``linux_<arch>`` platform
-tag, which installers like ``pip`` won't accept off PyPI/our index. We need to
-rewrite the platform tag to the appropriate ``manylinux_<major>_<minor>_<arch>``
-before uploading.
-
-Historically the tag was hard-coded per build (``manylinux_2_31`` for the
-Ubuntu 20.04-based image, ``manylinux_2_35`` for the Ubuntu 22.04-based
-images). That is brittle: bumping the base image silently produces wheels
-labelled with the wrong glibc requirement. This script asks ``auditwheel``
-to derive the tag from the symbol versions actually referenced by the
-binaries inside the wheel, so the label tracks reality.
-
-We can't simply call ``auditwheel repair`` -- it tries to graft external
-shared libraries into the wheel and fails on vLLM's CUDA/cuBLAS dependencies.
-Instead we use ``auditwheel.wheel_abi.analyze_wheel_abi`` directly, which is
-the same call that powers ``auditwheel show``, and read off
-``winfo.sym_policy.name``.
-
-Usage:
-    detect-manylinux-tag.py <wheel_path>
-
-The wheel is renamed in place; the new path is printed on stdout. All
-diagnostics go to stderr so callers can capture stdout safely.
-"""
-
-from __future__ import annotations
-
-import argparse
-import sys
-from pathlib import Path
-
-from auditwheel.error import (
-    AuditwheelError,
-    NonPlatformWheelError,
-    WheelToolsError,
-)
-from auditwheel.wheel_abi import analyze_wheel_abi
-from auditwheel.wheeltools import get_wheel_architecture, get_wheel_libc
-
-
-def detect_platform_tag(wheel_path: Path) -> str:
-    """Return the most precise platform tag the wheel is consistent with.
-
-    Mirrors ``auditwheel show`` but returns ``sym_policy`` rather than
-    ``overall_policy``: we only care about the glibc symbol versions used,
-    not about other policy axes (ISA extensions, blacklist, etc.) that
-    ``overall_policy`` folds in.
-    """
-    fn = wheel_path.name
-
-    try:
-        arch = get_wheel_architecture(fn)
-    except (WheelToolsError, NonPlatformWheelError):
-        # Architecture isn't deducible from the filename; let auditwheel
-        # infer it from the ELF binaries inside the wheel.
-        arch = None
-
-    try:
-        libc = get_wheel_libc(fn)
-    except WheelToolsError:
-        # An unrepaired wheel uses ``linux_<arch>``, which doesn't encode
-        # libc. Let auditwheel infer it from the ELF binaries.
-        libc = None
-
-    winfo = analyze_wheel_abi(
-        libc,
-        arch,
-        wheel_path,
-        frozenset(),
-        disable_isa_ext_check=False,
-        allow_graft=False,
-    )
-    return winfo.sym_policy.name
-
-
-def rename_wheel(wheel_path: Path, new_platform_tag: str) -> Path:
-    """Rename the wheel in place, replacing only its platform tag."""
-    # Wheel filename per PEP 427:
-    #   {distribution}-{version}(-{build})?-{python}-{abi}-{platform}.whl
-    # The platform tag is always the last ``-``-separated token before
-    # ``.whl``. Compound tags like ``manylinux_2_31_x86_64`` use ``_`` as the
-    # internal separator, so ``-``-splitting is unambiguous.
-    parts = wheel_path.stem.split("-")
-    if len(parts) < 5:
-        raise ValueError(f"Unrecognised wheel filename: {wheel_path.name}")
-    parts[-1] = new_platform_tag
-    new_path = wheel_path.with_name("-".join(parts) + ".whl")
-    if new_path != wheel_path:
-        wheel_path.rename(new_path)
-    return new_path
-
-
-def main() -> int:
-    parser = argparse.ArgumentParser(
-        description="Detect a wheel's manylinux platform tag with "
-        "auditwheel and rename the wheel in place."
-    )
-    parser.add_argument(
-        "wheel",
-        type=Path,
-        help="Path to the wheel to inspect and rename.",
-    )
-    args = parser.parse_args()
-
-    wheel_path: Path = args.wheel
-    if not wheel_path.is_file():
-        print(f"error: {wheel_path} is not a file", file=sys.stderr)
-        return 1
-
-    # Catch the things that ``analyze_wheel_abi`` and ``rename_wheel`` can
-    # raise: any subclass of ``AuditwheelError`` (pure-Python wheels,
-    # invalid libc, malformed wheels), filesystem errors, or our own
-    # ``ValueError`` for an unrecognised wheel filename. Print a single
-    # ``ERROR_TYPE: message`` line to stderr instead of a Python
-    # traceback, which is much friendlier in CI logs.
-    try:
-        new_tag = detect_platform_tag(wheel_path)
-        print(f"detected platform tag: {new_tag}", file=sys.stderr)
-        new_path = rename_wheel(wheel_path, new_tag)
-    except (AuditwheelError, ValueError, OSError) as e:
-        print(
-            f"error: failed to retag {wheel_path.name}: {type(e).__name__}: {e}",
-            file=sys.stderr,
-        )
-        return 2
-
-    if new_path != wheel_path:
-        print(f"renamed {wheel_path.name} -> {new_path.name}", file=sys.stderr)
-    else:
-        print(f"wheel already tagged {new_tag}", file=sys.stderr)
-
-    print(new_path)
-    return 0
-
-
-if __name__ == "__main__":
-    sys.exit(main())

.buildkite/scripts/docker-build-metadata-args.sh

@@ -1,54 +0,0 @@
-#!/bin/bash
-# Emit docker build flags for release image provenance metadata.
-# Keep this helper best-effort: missing Buildkite metadata should fall back to
-# local/default values instead of blocking the Docker build.
-
-# Variant examples: "", "cu129", "ubuntu2404", "cu129-ubuntu2404".
-variant="${1:-}"
-variant_suffix="${variant:+-${variant}}"
-
-image_name="${VLLM_DOCKER_IMAGE_NAME:-vllm/vllm-openai}"
-staging_repo="${VLLM_STAGING_IMAGE_REPO:-public.ecr.aws/q9t5s3a7/vllm-release-repo}"
-build_commit="${VLLM_BUILD_COMMIT:-${BUILDKITE_COMMIT:-unknown}}"
-build_pipeline="${VLLM_BUILD_PIPELINE:-${BUILDKITE_PIPELINE_ID:-${BUILDKITE_PIPELINE_SLUG:-local}}}"
-build_url="${VLLM_BUILD_URL:-${BUILDKITE_BUILD_URL:-}}"
-tag_commit="${BUILDKITE_COMMIT:-${build_commit}}"
-
-if [[ -n "${BUILDKITE:-}" || -n "${BUILDKITE_COMMIT:-}" ]]; then
-  release_version="${RELEASE_VERSION:-}"
-  if command -v buildkite-agent >/dev/null 2>&1; then
-    release_version="${release_version:-$(buildkite-agent meta-data get release-version 2>/dev/null)}"
-  fi
-  release_version="${release_version#v}"
-  release_version="${release_version:-${tag_commit}}"
-
-  staging_image_ref="${staging_repo}:${tag_commit}-$(uname -m)${variant_suffix}"
-
-  if [[ "${NIGHTLY:-}" == "1" ]]; then
-    if [[ -z "${variant}" ]]; then
-      image_tag="${image_name}:nightly-${tag_commit}"
-    elif [[ "${variant}" == cu* ]]; then
-      cuda_variant="${variant%%-*}"
-      remaining_variant="${variant#${cuda_variant}}"
-      image_tag="${image_name}:${cuda_variant}-nightly-${tag_commit}${remaining_variant}"
-    else
-      image_tag="${image_name}:nightly-${tag_commit}${variant_suffix}"
-    fi
-  else
-    image_tag="${image_name}:v${release_version}${variant_suffix}"
-  fi
-else
-  image_tag="${VLLM_IMAGE_TAG:-local/vllm-openai:dev}"
-  staging_image_ref="${image_tag}"
-fi
-
-emit_arg() {
-  printf -- "--build-arg %s=%s " "$1" "$2"
-}
-
-emit_arg VLLM_BUILD_COMMIT "${build_commit}"
-emit_arg VLLM_BUILD_PIPELINE "${build_pipeline}"
-emit_arg VLLM_BUILD_URL "${build_url}"
-# This is the intended public tag. The final digest is only known after push.
-emit_arg VLLM_IMAGE_TAG "${image_tag}"
-printf -- "--tag %s " "${staging_image_ref}"

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

@@ -10,13 +10,20 @@ set -ex
 BUCKET="vllm-wheels"
 INDICES_OUTPUT_DIR="indices"
 DEFAULT_VARIANT_ALIAS="cu130" # align with vLLM_MAIN_CUDA_VERSION in vllm/envs.py
+PYTHON="${PYTHON_PROG:-python3}" # try to read from env var, otherwise use python3
 SUBPATH=$BUILDKITE_COMMIT
 S3_COMMIT_PREFIX="s3://$BUCKET/$SUBPATH/"
 
-# Select python3 (>= 3.12) -- local if available, else a docker fallback.
-# shellcheck source=lib/select-python.sh
-source .buildkite/scripts/lib/select-python.sh
-select_python
+# detect if python3.12+ is available
+has_new_python=$($PYTHON -c "print(1 if __import__('sys').version_info >= (3,12) else 0)")
+if [[ "$has_new_python" -eq 0 ]]; then
+    # use new python from docker
+    docker pull python:3-slim
+    PYTHON="docker run --rm -u $(id -u):$(id -g) -v $(pwd):/app -w /app python:3-slim python3"
+fi
+
+echo "Using python interpreter: $PYTHON"
+echo "Python version: $($PYTHON --version)"
 
 # ======== generate and upload indices ========
 

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

@@ -378,11 +378,9 @@ HF_MOUNT="/root/.cache/huggingface"
 # double-quotes will have been stripped by the calling shell.
 if [[ -n "${VLLM_TEST_COMMANDS:-}" ]]; then
   commands="${VLLM_TEST_COMMANDS}"
-  commands_source="env"
   echo "Commands sourced from VLLM_TEST_COMMANDS (quoting preserved)"
 else
   commands="$*"
-  commands_source="argv"
   if [[ -z "$commands" ]]; then
     echo "Error: No test commands provided." >&2
     echo "Usage:" >&2
@@ -399,15 +397,9 @@ fi
 
 echo "Raw commands: $commands"
 
-# Only try to repair stripped pytest -m/-k quoting in legacy argv mode.
-# VLLM_TEST_COMMANDS preserves inner quoting already, and re-quoting that path
-# can corrupt embedded echo strings or otherwise well-formed shell fragments.
-if [[ "$commands_source" == "argv" ]]; then
-  commands=$(re_quote_pytest_markers "$commands")
-  echo "After re-quoting: $commands"
-else
-  echo "Skipping re-quoting for VLLM_TEST_COMMANDS input"
-fi
+# Fix quoting before ROCm overrides (so overrides see correct structure)
+commands=$(re_quote_pytest_markers "$commands")
+echo "After re-quoting: $commands"
 
 commands=$(apply_rocm_test_overrides "$commands")
 echo "Final commands: $commands"

.buildkite/scripts/hardware_ci/run-cpu-distributed-smoke-test.sh

@@ -3,37 +3,42 @@ set -euox pipefail
 export VLLM_CPU_CI_ENV=0
 export VLLM_CPU_KVCACHE_SPACE=1 # avoid OOM
 
-MODE=${1:-all}
+echo "--- PP+TP"
+vllm serve meta-llama/Llama-3.2-3B-Instruct -tp=2 -pp=2 --max-model-len=4096 &
+server_pid=$!
+timeout 600 bash -c "until curl localhost:8000/v1/models > /dev/null 2>&1; do sleep 1; done" || exit 1
+vllm bench serve \
+    --backend vllm \
+    --dataset-name random \
+    --model meta-llama/Llama-3.2-3B-Instruct \
+    --num-prompts 20 \
+    --result-dir ./test_results \
+    --result-filename tp_pp.json \
+    --save-result \
+    --endpoint /v1/completions
+kill -s SIGTERM $server_pid; wait $server_pid || true
+failed_req=$(jq '.failed' ./test_results/tp_pp.json)
+if [ "$failed_req" -ne 0 ]; then
+  echo "Some requests were failed!"
+  exit 1
+fi
 
-run_scenario() {
-    local label="$1" result_file="$2"
-    shift 2
-    echo "--- $label"
-    vllm serve meta-llama/Llama-3.2-3B-Instruct "$@" --max-model-len=4096 &
-    local server_pid=$!
-    timeout 600 bash -c "until curl localhost:8000/v1/models > /dev/null 2>&1; do sleep 1; done" || exit 1
-    vllm bench serve \
-        --backend vllm \
-        --dataset-name random \
-        --model meta-llama/Llama-3.2-3B-Instruct \
-        --num-prompts 20 \
-        --result-dir ./test_results \
-        --result-filename "$result_file" \
-        --save-result \
-        --endpoint /v1/completions
-    kill -s SIGTERM "$server_pid"; wait "$server_pid" || true
-    if [ "$(jq '.failed' "./test_results/$result_file")" -ne 0 ]; then
-        echo "Some requests were failed in $label!"
-        exit 1
-    fi
-}
-
-case "$MODE" in
-    tp_pp) run_scenario "PP+TP" tp_pp.json -tp=2 -pp=2 ;;
-    dp_tp) run_scenario "DP+TP" dp_tp.json -tp=2 -dp=2 ;;
-    all)
-        run_scenario "PP+TP" tp_pp.json -tp=2 -pp=2
-        run_scenario "DP+TP" dp_tp.json -tp=2 -dp=2
-        ;;
-    *) echo "ERROR: unknown mode '$MODE' (expected: tp_pp | dp_tp | all)" >&2; exit 1 ;;
-esac
+echo "--- DP+TP"
+vllm serve meta-llama/Llama-3.2-3B-Instruct -tp=2 -dp=2 --max-model-len=4096 &
+server_pid=$!
+timeout 600 bash -c "until curl localhost:8000/v1/models > /dev/null 2>&1; do sleep 1; done" || exit 1
+vllm bench serve \
+   --backend vllm \
+   --dataset-name random \
+   --model meta-llama/Llama-3.2-3B-Instruct \
+   --num-prompts 20 \
+   --result-dir ./test_results \
+   --result-filename dp_pp.json \
+   --save-result \
+   --endpoint /v1/completions
+kill -s SIGTERM $server_pid; wait $server_pid || true
+failed_req=$(jq '.failed' ./test_results/dp_pp.json)
+if [ "$failed_req" -ne 0 ]; then
+ echo "Some requests were failed!"
+ exit 1
+fi

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

@@ -31,21 +31,6 @@ function cpu_tests() {
     set -e
     pip list"
 
-  # Run kernel tests
-  docker exec cpu-test bash -c "
-    set -e
-    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"
-
-  # skip tests requiring model downloads if HF_TOKEN is not set
-  # due to rate-limits
-  if [ -z "$HF_TOKEN" ]; then
-    echo "Warning: HF_TOKEN is not set. Skipping tests that require model downloads."
-    return
-  fi
-
   # offline inference
   docker exec cpu-test bash -c "
     set -e
@@ -61,6 +46,13 @@ function cpu_tests() {
     set -e
     pytest -x -v -s tests/quantization/test_compressed_tensors.py::test_compressed_tensors_w8a8_logprobs"
 
+  # Run kernel tests
+  docker exec cpu-test bash -c "
+    set -e
+    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"
 
   # basic online serving
   docker exec cpu-test bash -c '
@@ -75,21 +67,6 @@ function cpu_tests() {
       --num-prompts 20 \
       --endpoint /v1/completions
     kill -s SIGTERM $server_pid &'
-
-  # smoke test for Gated DeltaNet
-  docker exec cpu-test bash -c '
-    set -e
-    VLLM_CPU_OMP_THREADS_BIND=$E2E_OMP_THREADS vllm serve Qwen/Qwen3.5-0.8B --max-model-len 2048 &
-    server_pid=$!
-    timeout 600 bash -c "until curl localhost:8000/v1/models; do sleep 1; done" || exit 1
-    vllm bench serve \
-      --backend vllm \
-      --dataset-name random \
-      --model Qwen/Qwen3.5-0.8B \
-      --num-prompts 20 \
-      --endpoint /v1/completions
-    kill -s SIGTERM $server_pid &'
-
 }
 
 # All of CPU tests are expected to be finished less than 40 mins.

.buildkite/scripts/hardware_ci/run-tpu-v1-test-part2.sh

@@ -61,7 +61,7 @@ echo "Results will be stored in: $RESULTS_DIR"
 echo "--- Installing Python dependencies ---"
 python3 -m pip install --progress-bar off git+https://github.com/thuml/depyf.git \
     && python3 -m pip install --progress-bar off pytest pytest-asyncio tpu-info \
-    && python3 -m pip install --progress-bar off "lm-eval[api]>=0.4.12" \
+    && python3 -m pip install --progress-bar off "lm-eval[api]>=0.4.11" \
     && python3 -m pip install --progress-bar off hf-transfer tblib==3.1.0
 echo "--- Python dependencies installed ---"
 

.buildkite/scripts/hardware_ci/run-tpu-v1-test.sh

@@ -61,7 +61,7 @@ echo "Results will be stored in: $RESULTS_DIR"
 echo "--- Installing Python dependencies ---"
 python3 -m pip install --progress-bar off git+https://github.com/thuml/depyf.git \
     && python3 -m pip install --progress-bar off pytest pytest-asyncio tpu-info \
-    && python3 -m pip install --progress-bar off "lm-eval[api]>=0.4.12" \
+    && python3 -m pip install --progress-bar off "lm-eval[api]>=0.4.11" \
     && python3 -m pip install --progress-bar off hf-transfer tblib==3.1.0
 echo "--- Python dependencies installed ---"
 
@@ -136,6 +136,8 @@ run_and_track_test 3 "test_accuracy.py::test_lm_eval_accuracy_v1_engine" \
     "python3 -m pytest -s -v /workspace/vllm/tests/entrypoints/llm/test_accuracy.py::test_lm_eval_accuracy_v1_engine"
 run_and_track_test 4 "test_quantization_accuracy.py" \
     "python3 -m pytest -s -v /workspace/vllm/tests/tpu/test_quantization_accuracy.py"
+run_and_track_test 5 "examples/offline_inference/tpu.py" \
+    "python3 /workspace/vllm/examples/offline_inference/tpu.py"
 run_and_track_test 6 "test_tpu_model_runner.py" \
     "python3 -m pytest -s -v /workspace/vllm/tests/v1/tpu/worker/test_tpu_model_runner.py"
 run_and_track_test 7 "test_sampler.py" \

.buildkite/scripts/lib/manylinux.sh

@@ -1,127 +0,0 @@
-#!/usr/bin/env bash
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-#
-# Shared helper for rewriting a wheel's platform tag from the generic
-# ``linux_<arch>`` to the correct ``manylinux_<major>_<minor>_<arch>``.
-# After sourcing, call ``apply_manylinux_tag <wheel>`` on each wheel
-# that still carries the generic tag; the renamed path is printed on
-# stdout (logs go to stderr).
-#
-# Why a pinned Docker container instead of using whatever Python
-# happens to be on the agent:
-#   - vLLM's release agents are heterogeneous -- they don't agree on
-#     a Python minor version, and we can't rely on a particular
-#     ``auditwheel`` being installed.
-#   - ``detect-manylinux-tag.py`` reads ``auditwheel.wheel_abi`` and
-#     ``Policy.sym_policy``, which are *internal* APIs without a
-#     stability promise. Pinning both Python and auditwheel makes the
-#     detected tag a function of the inputs alone, and shifts version
-#     bumps from "implicit drift" to "deliberate, retested change".
-#   - Other release scripts (``generate-and-upload-nightly-index.sh``,
-#     ``upload-rocm-wheels.sh``) already use the python:3-slim image
-#     when the agent's interpreter is too old; this is the same idea
-#     made stricter.
-#
-# To keep the per-wheel cost down (the ROCm upload retags ~10 wheels
-# each run), we install auditwheel into a long-lived helper container
-# once on source, then ``docker exec`` into it for each call.
-#
-# Trap behaviour:
-# - Sourcing installs an EXIT trap that calls ``manylinux_cleanup`` to
-#   tear down the helper container. Any EXIT trap that was already in
-#   place when this file was sourced is captured and run AFTER our
-#   cleanup, so we don't silently clobber it.
-# - If a caller sets a new EXIT trap *after* sourcing, that trap will
-#   replace ours; in that case the caller should call
-#   ``manylinux_cleanup`` from their own handler.
-
-if [[ -n "${_MANYLINUX_LIB_SOURCED:-}" ]]; then
-    return 0
-fi
-_MANYLINUX_LIB_SOURCED=1
-
-# Pin both sides. Bump these deliberately and re-run a representative
-# wheel from each build target through the detection.
-_MANYLINUX_PYTHON_IMAGE="python:3.12-slim"
-_MANYLINUX_AUDITWHEEL_VERSION="6.6.0"
-
-# Resolve our own directory (and the sibling detect script) using the
-# canonical, symlink-resolved path. The container mounts cwd at the
-# same absolute path on both sides, so all paths we hand to it -- the
-# script, the wheel -- must canonicalise to a location under cwd.
-_MANYLINUX_LIB_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd -P)"
-_MANYLINUX_DETECT_SCRIPT="$(cd "${_MANYLINUX_LIB_DIR}/.." && pwd -P)/detect-manylinux-tag.py"
-_MANYLINUX_CWD="$(pwd -P)"
-
-docker pull --quiet "$_MANYLINUX_PYTHON_IMAGE" >/dev/null
-
-# Spin up a long-lived helper container so we install auditwheel once
-# and then ``docker exec`` into it for each wheel.
-#
-# The container runs as root so ``pip install`` can write into the
-# system site-packages; individual ``docker exec`` calls below pin
-# themselves to the host UID so any file rename happens with host
-# ownership, not root.
-_MANYLINUX_CONTAINER="$(docker run -d --rm \
-    -v "$_MANYLINUX_CWD:$_MANYLINUX_CWD" \
-    -w "$_MANYLINUX_CWD" \
-    "$_MANYLINUX_PYTHON_IMAGE" \
-    sleep infinity)"
-docker exec "$_MANYLINUX_CONTAINER" \
-    pip install --quiet --disable-pip-version-check \
-    --root-user-action=ignore \
-    "auditwheel==${_MANYLINUX_AUDITWHEEL_VERSION}"
-
-# Public cleanup -- safe to call multiple times.
-manylinux_cleanup() {
-    if [[ -n "${_MANYLINUX_CONTAINER:-}" ]]; then
-        docker rm -f "$_MANYLINUX_CONTAINER" >/dev/null 2>&1 || true
-        _MANYLINUX_CONTAINER=""
-    fi
-}
-
-# Capture any EXIT trap that was already in place so we can chain to
-# it rather than overwrite it. ``trap -p EXIT`` prints the handler in
-# eval-able form (``trap -- 'CMD' EXIT``) or nothing if unset; we
-# strip the wrapper to recover ``CMD``. Handles the common case --
-# CMDs without embedded single quotes -- and degrades gracefully (we
-# still run our own cleanup) for the pathological case.
-_manylinux_prev_exit_trap_cmd=""
-_manylinux_existing_exit_trap="$(trap -p EXIT)"
-if [[ -n "$_manylinux_existing_exit_trap" ]]; then
-    _tmp="${_manylinux_existing_exit_trap#trap -- \'}"
-    _manylinux_prev_exit_trap_cmd="${_tmp%\' EXIT}"
-    unset _tmp
-fi
-unset _manylinux_existing_exit_trap
-
-_manylinux_run_exit_chain() {
-    manylinux_cleanup
-    if [[ -n "$_manylinux_prev_exit_trap_cmd" ]]; then
-        eval "$_manylinux_prev_exit_trap_cmd"
-    fi
-}
-trap _manylinux_run_exit_chain EXIT
-
-# Detect the manylinux platform tag for a single wheel and rename it
-# in place, printing the renamed wheel path on stdout. Returns
-# non-zero on failure (which under ``set -e`` propagates to caller).
-#
-# The wheel must be reachable via a path under the host cwd so it's
-# visible inside the helper container; in CI the wheels always live
-# under ``artifacts/`` so this is fine.
-apply_manylinux_tag() {
-    local wheel="$1"
-    local abs_wheel
-    abs_wheel="$(realpath "$wheel")"
-    local new_wheel
-    new_wheel="$(docker exec -u "$(id -u):$(id -g)" \
-        "$_MANYLINUX_CONTAINER" \
-        python "$_MANYLINUX_DETECT_SCRIPT" "$abs_wheel")"
-    if [[ -z "$new_wheel" || ! -f "$new_wheel" ]]; then
-        echo "apply_manylinux_tag: detect-manylinux-tag.py did not produce a valid wheel path for $wheel" >&2
-        return 1
-    fi
-    printf '%s\n' "$new_wheel"
-}

.buildkite/scripts/lib/select-python.sh

@@ -1,41 +0,0 @@
-#!/usr/bin/env bash
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-#
-# Pick a Python interpreter for buildkite scripts: prefer a local
-# ``python3`` if it is recent enough (>= 3.12), otherwise fall back to
-# a one-shot Docker container running ``python:3-slim``. After
-# ``select_python`` returns, ``$PYTHON`` is set in the caller's shell
-# and is safe to use as a command (e.g. ``$PYTHON some_script.py``).
-#
-# The 3.12 threshold matches what the existing nightly-index work
-# expects -- typing features used by ``generate-nightly-index.py``.
-# This helper does not pin the *minor* version; if you need stricter
-# reproducibility (e.g. relying on auditwheel internals), invoke
-# Docker yourself with a pinned tag rather than calling this.
-
-if [[ -n "${_SELECT_PYTHON_LIB_SOURCED:-}" ]]; then
-    return 0
-fi
-_SELECT_PYTHON_LIB_SOURCED=1
-
-# Sets ``PYTHON`` in the caller's shell and exports it. Idempotent --
-# calling twice is safe and the second call simply re-runs the probe.
-select_python() {
-    local py="${PYTHON_PROG:-python3}"
-    local has_new_python
-    has_new_python=$("$py" -c \
-        "print(1 if __import__('sys').version_info >= (3,12) else 0)" \
-        2>/dev/null || echo 0)
-    if [[ "$has_new_python" -eq 0 ]]; then
-        # ``-u $(id -u):$(id -g)`` so files created via the container
-        # end up owned by the host user, not root.
-        docker pull python:3-slim
-        PYTHON="docker run --rm -u $(id -u):$(id -g) -v $(pwd):/app -w /app python:3-slim python3"
-    else
-        PYTHON="$py"
-    fi
-    export PYTHON
-    echo "Using python interpreter: $PYTHON"
-    echo "Python version: $($PYTHON --version)"
-}

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

@@ -1,180 +0,0 @@
-#!/bin/bash
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-#
-# Publish release Docker images from ECR to DockerHub.
-# Pulls per-arch images, tags with latest and versioned tags, pushes them,
-# then creates and pushes multi-arch manifests.
-
-set -euo pipefail
-
-RELEASE_VERSION=$(buildkite-agent meta-data get release-version --default "" | sed 's/^v//')
-if [ -z "${RELEASE_VERSION}" ]; then
-  echo "ERROR: release-version metadata not set"
-  exit 1
-fi
-
-COMMIT="$BUILDKITE_COMMIT"
-ROCM_BASE_CACHE_KEY=$(.buildkite/scripts/cache-rocm-base-wheels.sh key)
-
-echo "========================================"
-echo "Publishing release images v${RELEASE_VERSION}"
-echo "  Commit: ${COMMIT}"
-echo "  ROCm base cache key: ${ROCM_BASE_CACHE_KEY}"
-echo "========================================"
-
-# Login to ECR to pull staging images
-aws ecr-public get-login-password --region us-east-1 | \
-  docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
-
-# ---- CUDA (default: 13.0) ----
-
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64 vllm/vllm-openai:latest-x86_64
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64
-docker push vllm/vllm-openai:latest-x86_64
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64 vllm/vllm-openai:latest-aarch64
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64
-docker push vllm/vllm-openai:latest-aarch64
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64
-
-docker manifest rm vllm/vllm-openai:latest || true
-docker manifest rm vllm/vllm-openai:v${RELEASE_VERSION} || true
-docker manifest create vllm/vllm-openai:latest vllm/vllm-openai:latest-x86_64 vllm/vllm-openai:latest-aarch64
-docker manifest create vllm/vllm-openai:v${RELEASE_VERSION} vllm/vllm-openai:v${RELEASE_VERSION}-x86_64 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64
-docker manifest push vllm/vllm-openai:latest
-docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}
-
-# ---- CUDA 12.9 ----
-
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-cu129
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-cu129
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-cu129 vllm/vllm-openai:latest-x86_64-cu129
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129
-docker push vllm/vllm-openai:latest-x86_64-cu129
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-cu129 vllm/vllm-openai:latest-aarch64-cu129
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129
-docker push vllm/vllm-openai:latest-aarch64-cu129
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129
-
-docker manifest rm vllm/vllm-openai:latest-cu129 || true
-docker manifest rm vllm/vllm-openai:v${RELEASE_VERSION}-cu129 || true
-docker manifest create vllm/vllm-openai:latest-cu129 vllm/vllm-openai:latest-x86_64-cu129 vllm/vllm-openai:latest-aarch64-cu129
-docker manifest create vllm/vllm-openai:v${RELEASE_VERSION}-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129
-docker manifest push vllm/vllm-openai:latest-cu129
-docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}-cu129
-
-# ---- Ubuntu 24.04 (CUDA 13.0) ----
-
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-ubuntu2404
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-ubuntu2404
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-ubuntu2404 vllm/vllm-openai:latest-x86_64-ubuntu2404
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-ubuntu2404
-docker push vllm/vllm-openai:latest-x86_64-ubuntu2404
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-ubuntu2404
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-ubuntu2404 vllm/vllm-openai:latest-aarch64-ubuntu2404
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-ubuntu2404
-docker push vllm/vllm-openai:latest-aarch64-ubuntu2404
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-ubuntu2404
-
-docker manifest rm vllm/vllm-openai:latest-ubuntu2404 || true
-docker manifest rm vllm/vllm-openai:v${RELEASE_VERSION}-ubuntu2404 || true
-docker manifest create vllm/vllm-openai:latest-ubuntu2404 vllm/vllm-openai:latest-x86_64-ubuntu2404 vllm/vllm-openai:latest-aarch64-ubuntu2404
-docker manifest create vllm/vllm-openai:v${RELEASE_VERSION}-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-ubuntu2404
-docker manifest push vllm/vllm-openai:latest-ubuntu2404
-docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}-ubuntu2404
-
-# ---- Ubuntu 24.04 (CUDA 12.9) ----
-
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-cu129-ubuntu2404
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-cu129-ubuntu2404
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-cu129-ubuntu2404 vllm/vllm-openai:latest-x86_64-cu129-ubuntu2404
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-cu129-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129-ubuntu2404
-docker push vllm/vllm-openai:latest-x86_64-cu129-ubuntu2404
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129-ubuntu2404
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-cu129-ubuntu2404 vllm/vllm-openai:latest-aarch64-cu129-ubuntu2404
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-aarch64-cu129-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129-ubuntu2404
-docker push vllm/vllm-openai:latest-aarch64-cu129-ubuntu2404
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129-ubuntu2404
-
-docker manifest rm vllm/vllm-openai:latest-cu129-ubuntu2404 || true
-docker manifest rm vllm/vllm-openai:v${RELEASE_VERSION}-cu129-ubuntu2404 || true
-docker manifest create vllm/vllm-openai:latest-cu129-ubuntu2404 vllm/vllm-openai:latest-x86_64-cu129-ubuntu2404 vllm/vllm-openai:latest-aarch64-cu129-ubuntu2404
-docker manifest create vllm/vllm-openai:v${RELEASE_VERSION}-cu129-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu129-ubuntu2404 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu129-ubuntu2404
-docker manifest push vllm/vllm-openai:latest-cu129-ubuntu2404
-docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}-cu129-ubuntu2404
-
-# ---- ROCm ----
-
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-rocm
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm-base
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-rocm vllm/vllm-openai-rocm:latest
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-rocm vllm/vllm-openai-rocm:v${RELEASE_VERSION}
-docker push vllm/vllm-openai-rocm:latest
-docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}
-
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm-base vllm/vllm-openai-rocm:latest-base
-docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm-base vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
-docker push vllm/vllm-openai-rocm:latest-base
-docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
-
-# ---- CPU ----
-# CPU images are behind separate block steps and may not have been built.
-# All-or-nothing: inspect both arches first, then either publish everything
-# (per-arch + multi-arch manifest) or skip everything. Publishing only one
-# arch would leave `:latest-x86_64` pointing at the new release while the
-# `:latest` multi-arch manifest still resolves to the previous release.
-
-CPU_X86_TAG=public.ecr.aws/q9t5s3a7/vllm-cpu-release-repo:v${RELEASE_VERSION}
-CPU_ARM_TAG=public.ecr.aws/q9t5s3a7/vllm-arm64-cpu-release-repo:v${RELEASE_VERSION}
-
-CPU_X86_AVAILABLE=false
-CPU_ARM_AVAILABLE=false
-docker manifest inspect "${CPU_X86_TAG}" >/dev/null 2>&1 && CPU_X86_AVAILABLE=true
-docker manifest inspect "${CPU_ARM_TAG}" >/dev/null 2>&1 && CPU_ARM_AVAILABLE=true
-
-if [ "$CPU_X86_AVAILABLE" = "true" ] && [ "$CPU_ARM_AVAILABLE" = "true" ]; then
-  docker pull "${CPU_X86_TAG}"
-  docker tag "${CPU_X86_TAG}" vllm/vllm-openai-cpu:latest-x86_64
-  docker tag "${CPU_X86_TAG}" vllm/vllm-openai-cpu:v${RELEASE_VERSION}-x86_64
-  docker push vllm/vllm-openai-cpu:latest-x86_64
-  docker push vllm/vllm-openai-cpu:v${RELEASE_VERSION}-x86_64
-
-  docker pull "${CPU_ARM_TAG}"
-  docker tag "${CPU_ARM_TAG}" vllm/vllm-openai-cpu:latest-arm64
-  docker tag "${CPU_ARM_TAG}" vllm/vllm-openai-cpu:v${RELEASE_VERSION}-arm64
-  docker push vllm/vllm-openai-cpu:latest-arm64
-  docker push vllm/vllm-openai-cpu:v${RELEASE_VERSION}-arm64
-
-  docker manifest rm vllm/vllm-openai-cpu:latest || true
-  docker manifest rm vllm/vllm-openai-cpu:v${RELEASE_VERSION} || true
-  docker manifest create vllm/vllm-openai-cpu:latest vllm/vllm-openai-cpu:latest-x86_64 vllm/vllm-openai-cpu:latest-arm64
-  docker manifest create vllm/vllm-openai-cpu:v${RELEASE_VERSION} vllm/vllm-openai-cpu:v${RELEASE_VERSION}-x86_64 vllm/vllm-openai-cpu:v${RELEASE_VERSION}-arm64
-  docker manifest push vllm/vllm-openai-cpu:latest
-  docker manifest push vllm/vllm-openai-cpu:v${RELEASE_VERSION}
-elif [ "$CPU_X86_AVAILABLE" = "false" ] && [ "$CPU_ARM_AVAILABLE" = "false" ]; then
-  echo "WARNING: Neither CPU image found in ECR, skipping CPU publish (ensure block-cpu-release-image-build and block-arm64-cpu-release-image-build were unblocked and the builds finished pushing)"
-else
-  # Partial state: one arch built, the other did not. Fail loudly rather than
-  # ship a Docker Hub state where `:latest-${arch}` and `:latest` (multi-arch)
-  # disagree on which release they point at.
-  echo "ERROR: Partial CPU build detected (x86_64=${CPU_X86_AVAILABLE}, arm64=${CPU_ARM_AVAILABLE})."
-  echo "       Refusing to publish to avoid split-tag drift between per-arch and multi-arch tags."
-  echo "       Re-run the missing CPU build and retry, or manually publish if a single-arch release is intended."
-  exit 1
-fi
-
-echo ""
-echo "Successfully published release images for v${RELEASE_VERSION}"

.buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_prefetch_offload.sh

@@ -51,7 +51,6 @@ vllm serve "$MODEL" \
   --offload-num-in-group 2 \
   --offload-prefetch-step 1 \
   --offload-params w13_weight w2_weight \
-  --generation-config vllm \
   --port "$PORT" \
   ${EXTRA_ARGS+"${EXTRA_ARGS[@]}"} &
 SERVER_PID=$!

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

@@ -28,7 +28,6 @@
 #   BFCL_MAX_MODEL_LEN  - Max model length (default: 4096)
 #   BFCL_PORT           - Server port (default: 8000)
 #   BFCL_REASONING_PARSER - Reasoning parser name (default: disabled)
-#   BFCL_TEMPERATURE    - Temperature (default: 0.0)
 #   BFCL_EXTRA_ARGS     - Additional vLLM server args
 
 set -euo pipefail
@@ -44,7 +43,6 @@ TP_SIZE="${BFCL_TP_SIZE:-1}"
 MAX_MODEL_LEN="${BFCL_MAX_MODEL_LEN:-4096}"
 PORT="${BFCL_PORT:-8000}"
 REASONING_PARSER="${BFCL_REASONING_PARSER:-}"
-TEMPERATURE="${BFCL_TEMPERATURE:-0.0}"
 EXTRA_ARGS="${BFCL_EXTRA_ARGS:-}"
 
 # Set up output directory
@@ -141,7 +139,7 @@ echo "vLLM server is ready. (started in ${SECONDS_WAITED}s)"
 # be patched in-process so BFCL knows to use the OpenAI-compatible handler
 # against our local vLLM server.
 bfcl_exit_code=0
-python3 - "$MODEL" "$TEST_CATEGORY" "$NUM_THREADS" "$PORT" "$API_TYPE" "$TEMPERATURE" "$OUTPUT_DIR" << 'PYEOF' || bfcl_exit_code=$?
+python3 - "$MODEL" "$TEST_CATEGORY" "$NUM_THREADS" "$PORT" "$API_TYPE" "$OUTPUT_DIR" << 'PYEOF' || bfcl_exit_code=$?
 import os
 import sys
 
@@ -150,8 +148,7 @@ test_category = sys.argv[2]
 num_threads = int(sys.argv[3])
 port = sys.argv[4]
 api_type = sys.argv[5]
-temperature = float(sys.argv[6])
-output_dir = sys.argv[7] if len(sys.argv) > 7 and sys.argv[7] else os.getcwd()
+output_dir = sys.argv[6] if len(sys.argv) > 6 and sys.argv[6] else os.getcwd()
 
 os.environ["OPENAI_BASE_URL"] = f"http://localhost:{port}/v1"
 os.environ["OPENAI_API_KEY"] = "dummy"
@@ -207,7 +204,6 @@ gen_kwargs["model"] = [model]
 gen_kwargs["test_category"] = [c.strip() for c in test_category.split(",")]
 gen_kwargs["skip_server_setup"] = True
 gen_kwargs["num_threads"] = num_threads
-gen_kwargs["temperature"] = temperature
 generate(**gen_kwargs)
 
 # ---- evaluate ----

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

@@ -2,18 +2,14 @@
 
 set -ex
 
-# Upload a single wheel to S3, after detecting and applying the appropriate
-# manylinux platform tag with auditwheel.
+# Upload a single wheel to S3 (rename linux -> manylinux).
 # Index generation is handled separately by generate-and-upload-nightly-index.sh.
 
-# shellcheck source=lib/manylinux.sh
-source .buildkite/scripts/lib/manylinux.sh
-
 BUCKET="vllm-wheels"
 SUBPATH=$BUILDKITE_COMMIT
 S3_COMMIT_PREFIX="s3://$BUCKET/$SUBPATH/"
 
-# ========= locate the wheel ==========
+# ========= collect, rename & upload the wheel ==========
 
 # Assume wheels are in artifacts/dist/*.whl
 wheel_files=(artifacts/dist/*.whl)
@@ -25,9 +21,19 @@ if [[ ${#wheel_files[@]} -ne 1 ]]; then
 fi
 wheel="${wheel_files[0]}"
 
-# ========= detect manylinux tag and rename ==========
+# default build image uses ubuntu 20.04, which corresponds to manylinux_2_31
+# we also accept params as manylinux tag
+# refer to https://github.com/mayeut/pep600_compliance?tab=readme-ov-file#acceptable-distros-to-build-wheels
+manylinux_version="${1:-manylinux_2_31}"
 
-wheel="$(apply_manylinux_tag "$wheel")"
+# Rename 'linux' to the appropriate manylinux version in the wheel filename
+if [[ "$wheel" != *"linux"* ]]; then
+  echo "Error: Wheel filename does not contain 'linux': $wheel"
+  exit 1
+fi
+new_wheel="${wheel/linux/$manylinux_version}"
+mv -- "$wheel" "$new_wheel"
+wheel="$new_wheel"
 echo "Renamed wheel to: $wheel"
 
 # Extract the version from the wheel

.buildkite/scripts/upload-release-wheels-pypi.sh

@@ -39,17 +39,10 @@ fi
 
 set -x # avoid printing secrets above
 
-# install uv if not already available
-if ! command -v uv &> /dev/null; then
-  curl -LsSf https://astral.sh/uv/install.sh | UV_VERSION=0.11.14 sh
-  export PATH="$HOME/.local/bin:$PATH"
-fi
-
-# install twine and sdist build prerequisites using uv with Python 3.12
-uv venv --python 3.12 /tmp/vllm-release-env
+# install twine from pypi
+python3 -m venv /tmp/vllm-release-env
 source /tmp/vllm-release-env/bin/activate
-uv pip install twine
-uv pip install -r requirements/build/cuda.txt
+pip install twine
 python3 -m twine --version
 
 # copy release wheels to local directory

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

@@ -20,6 +20,10 @@ BUCKET="${S3_BUCKET:-vllm-wheels}"
 ROCM_SUBPATH="rocm/${BUILDKITE_COMMIT}"
 S3_COMMIT_PREFIX="s3://$BUCKET/$ROCM_SUBPATH/"
 INDICES_OUTPUT_DIR="rocm-indices"
+PYTHON="${PYTHON_PROG:-python3}"
+
+# ROCm uses manylinux_2_35 (Ubuntu 22.04 based)
+MANYLINUX_VERSION="manylinux_2_35"
 
 echo "========================================"
 echo "ROCm Wheel Upload Configuration"
@@ -30,21 +34,19 @@ echo "Commit: $BUILDKITE_COMMIT"
 echo "Branch: $BUILDKITE_BRANCH"
 echo "========================================"
 
-# ======== Part 0: Setup Python and helpers ========
+# ======== Part 0: Setup Python ========
 
-# Pick a Python interpreter for index generation -- local if recent
-# enough, else a one-shot docker fallback.
-# shellcheck source=lib/select-python.sh
-source .buildkite/scripts/lib/select-python.sh
-select_python
+# Detect if python3.12+ is available
+has_new_python=$($PYTHON -c "print(1 if __import__('sys').version_info >= (3,12) else 0)" 2>/dev/null || echo 0)
+if [[ "$has_new_python" -eq 0 ]]; then
+    # Use new python from docker
+    # Use --user to ensure files are created with correct ownership (not root)
+    docker pull python:3-slim
+    PYTHON="docker run --rm --user $(id -u):$(id -g) -v $(pwd):/app -w /app python:3-slim python3"
+fi
 
-# Set up auditwheel-in-a-container for the manylinux retagging step.
-# Distinct from select_python: ``manylinux.sh`` deliberately pins both
-# the Python and auditwheel versions (the script reads auditwheel
-# internals) and so always runs in a known-good container regardless
-# of what's on the agent.
-# shellcheck source=lib/manylinux.sh
-source .buildkite/scripts/lib/manylinux.sh
+echo "Using python interpreter: $PYTHON"
+echo "Python version: $($PYTHON --version)"
 
 # ======== Part 1: Collect and prepare wheels ========
 
@@ -61,18 +63,11 @@ if [ "$WHEEL_COUNT" -eq 0 ]; then
     exit 1
 fi
 
-# Detect the appropriate manylinux platform tag for any wheel that still
-# carries the generic ``linux_<arch>`` tag, and rename it in place. We use
-# auditwheel via ``apply_manylinux_tag`` (see lib/manylinux.sh) rather than
-# a hard-coded ``manylinux_2_35`` string so that the label tracks the actual
-# glibc symbol versions used by the binaries (and stays correct if the
-# rocm_base image is rebased).
-#
-# The ``linux``/``manylinux`` filter below skips both pre-tagged wheels
-# (e.g. upstream torch) and pure-Python ``-any.whl`` wheels.
+# Rename linux to manylinux in wheel filenames
 for wheel in all-rocm-wheels/*.whl; do
     if [[ "$wheel" == *"linux"* ]] && [[ "$wheel" != *"manylinux"* ]]; then
-        new_wheel="$(apply_manylinux_tag "$wheel")"
+        new_wheel="${wheel/linux/$MANYLINUX_VERSION}"
+        mv -- "$wheel" "$new_wheel"
         echo "Renamed: $(basename "$wheel") -> $(basename "$new_wheel")"
     fi
 done

.buildkite/test-amd.yaml

@@ -230,6 +230,7 @@ steps:
   - tests/entrypoints/llm/test_collective_rpc.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s entrypoints/llm/test_collective_rpc.py
   - pytest -v -s ./compile/fullgraph/test_basic_correctness.py
   - pytest -v -s ./compile/test_wrapper.py
@@ -271,6 +272,7 @@ steps:
   - tests/v1/worker/test_worker_memory_snapshot.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - VLLM_TEST_SAME_HOST=1 torchrun --nproc-per-node=4 distributed/test_same_node.py | grep 'Same node test passed'
   - VLLM_TEST_SAME_HOST=1 VLLM_TEST_WITH_DEFAULT_DEVICE_SET=1 torchrun --nproc-per-node=4 distributed/test_same_node.py | grep 'Same node test passed'
   - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s v1/shutdown
@@ -386,18 +388,18 @@ steps:
     - python3 basic/offline_inference/embed.py
     - python3 basic/offline_inference/score.py
     # Multi-modal models
-    - python3 generate/multimodal/audio_language_offline.py --seed 0
-    - python3 generate/multimodal/vision_language_offline.py --seed 0
-    - python3 generate/multimodal/vision_language_multi_image_offline.py --seed 0
-    - python3 generate/multimodal/encoder_decoder_multimodal_offline.py --model-type whisper --seed 0
+    - python3 offline_inference/audio_language.py --seed 0
+    - python3 offline_inference/vision_language.py --seed 0
+    - python3 offline_inference/vision_language_multi_image.py --seed 0
+    - python3 offline_inference/encoder_decoder_multimodal.py --model-type whisper --seed 0
     # Pooling models
     - python3 pooling/embed/vision_embedding_offline.py --seed 0
     # Features demo
-    - python3 features/automatic_prefix_caching/prefix_caching_offline.py
-    - python3 deployment/llm_engine_example.py
-    - python3 features/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 features/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
-    - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
-    - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
+    - python3 offline_inference/prefix_caching.py
+    - python3 offline_inference/llm_engine_example.py
+    - python3 others/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 others/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
+    - python3 offline_inference/spec_decode.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
+    - python3 offline_inference/spec_decode.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
 
 #----------------------------------------------------------  mi250 · kernels  ----------------------------------------------------------#
 
@@ -460,7 +462,7 @@ steps:
   - tests/lora
   - vllm/platforms/rocm.py
   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
+  - 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_llm_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  -------------------------------------------------------#
 
@@ -588,6 +590,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - pip freeze | grep -E 'torch'
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s models/language -m 'core_model and slow_test' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
 
 - label: Multi-Modal Models (Extended Generation 2) # TBD
@@ -618,7 +621,6 @@ steps:
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
   agent_pool: mi250_1
   torch_nightly: true
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -798,6 +800,32 @@ steps:
   commands:
   - pytest -v -s v1/attention
 
+- label: V1 Core + KV + Metrics # TBD
+  timeout_in_minutes: 60
+  mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
+  agent_pool: mi250_1
+  optional: true
+  working_dir: "/vllm-workspace/tests"
+  source_file_dependencies:
+  - vllm/
+  - tests/v1/core
+  - tests/v1/executor
+  - tests/v1/kv_offload
+  - tests/v1/worker
+  - tests/v1/kv_connector/unit
+  - tests/v1/metrics
+  - tests/entrypoints/openai/correctness/test_lmeval.py
+  commands:
+  - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
+  - pytest -v -s -m 'not cpu_test' v1/core
+  - pytest -v -s v1/executor
+  - pytest -v -s v1/kv_offload
+  - pytest -v -s v1/worker
+  - pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
+  - pytest -v -s -m 'not cpu_test' v1/metrics
+  - pip install -U git+https://github.com/robertgshaw2-redhat/lm-evaluation-harness.git@streaming-api
+  - pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
+
 - label: V1 Sample + Logits # TBD
   timeout_in_minutes: 60
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
@@ -836,6 +864,7 @@ steps:
   - tests/entrypoints/openai/test_multi_api_servers.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
   - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
   - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
@@ -854,7 +883,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
+  - ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
 
 - label: V1 e2e (2 GPUs) # TBD
   timeout_in_minutes: 180
@@ -901,9 +930,7 @@ steps:
   - tests/renderers
   - tests/standalone_tests/lazy_imports.py
   - tests/tokenizers_
-  - tests/reasoning
   - tests/tool_parsers
-  - tests/parser
   - tests/transformers_utils
   - tests/config
   commands:
@@ -915,9 +942,8 @@ steps:
   - 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 reasoning --ignore=reasoning/test_seedoss_reasoning_parser.py --ignore=reasoning/test_glm4_moe_reasoning_parser.py --ignore=reasoning/test_gemma4_reasoning_parser.py
   - pytest -v -s tool_parsers
-  - pytest -v -s parser
   - pytest -v -s transformers_utils
   - pytest -v -s config
 
@@ -1076,13 +1102,12 @@ steps:
   - vllm/compilation/
   - vllm/model_executor/layers
   - tests/compile/passes/distributed/
-  - tests/compile/fusions_e2e/
   - vllm/_aiter_ops.py
   - vllm/platforms/rocm.py
   commands:
   - export VLLM_TEST_CLEAN_GPU_MEMORY=1
   - VLLM_TEST_CLEAN_GPU_MEMORY=1 pytest -v -s tests/compile/passes/distributed/test_async_tp.py
-  - pytest -v -s tests/compile/fusions_e2e/test_tp2_ar_rms.py::test_tp2_ar_rms_fusions
+  - pytest -v -s tests/compile/passes/distributed/test_sequence_parallelism.py
 
 #-----------------------------------------------------------  mi300 · cuda  ------------------------------------------------------------#
 
@@ -1143,12 +1168,13 @@ steps:
   - vllm/v1/attention/backends/
   - vllm/v1/attention/selector.py
   - tests/distributed/test_context_parallel.py
-  - examples/features/data_parallel/data_parallel_offline.py
+  - examples/offline_inference/data_parallel.py
   - vllm/_aiter_ops.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s tests/distributed/test_context_parallel.py
-  - VLLM_LOGGING_LEVEL=DEBUG python3 examples/features/data_parallel/data_parallel_offline.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=allgather_reducescatter --disable-nccl-for-dp-synchronization
+  - VLLM_LOGGING_LEVEL=DEBUG python3 examples/offline_inference/data_parallel.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=allgather_reducescatter --disable-nccl-for-dp-synchronization
 
 - label: Distributed Tests (4xA100-4xMI300) # TBD
   timeout_in_minutes: 180
@@ -1160,6 +1186,7 @@ steps:
   source_file_dependencies:
   - vllm/
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s distributed/test_custom_all_reduce.py
   - torchrun --nproc_per_node=2 distributed/test_ca_buffer_sharing.py
   - TARGET_TEST_SUITE=A100 pytest basic_correctness/ -v -s -m 'distributed(num_gpus=2)'
@@ -1176,16 +1203,17 @@ steps:
   - tests/distributed/test_torchrun_example.py
   - tests/distributed/test_torchrun_example_moe.py
   - examples/rl/
-  - tests/examples/features/data_parallel/data_parallel_offline.py
+  - tests/examples/offline_inference/data_parallel.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - torchrun --nproc-per-node=4 distributed/test_torchrun_example.py
   - PP_SIZE=2 torchrun --nproc-per-node=4 distributed/test_torchrun_example.py
   - TP_SIZE=4 torchrun --nproc-per-node=4 distributed/test_torchrun_example_moe.py
   - PP_SIZE=2 TP_SIZE=2 torchrun --nproc-per-node=4 distributed/test_torchrun_example_moe.py
   - DP_SIZE=4 ENABLE_EP=1 torchrun --nproc-per-node=4 distributed/test_torchrun_example_moe.py
   - TP_SIZE=2 DP_SIZE=2 ENABLE_EP=1 torchrun --nproc-per-node=4 distributed/test_torchrun_example_moe.py
-  - python3 ../examples/features/data_parallel/data_parallel_offline.py --enforce-eager
+  - python3 ../examples/offline_inference/data_parallel.py --enforce-eager
   # rlhf examples
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 ../examples/rl/rlhf_nccl.py
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 ../examples/rl/rlhf_ipc.py
@@ -1224,6 +1252,7 @@ steps:
   - vllm/platforms/rocm.py
   commands:
   - export VLLM_USE_RAY_V2_EXECUTOR_BACKEND=1
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s distributed/test_ray_v2_executor.py
   - pytest -v -s distributed/test_ray_v2_executor_e2e.py
   - pytest -v -s distributed/test_pipeline_parallel.py -k "ray"
@@ -1237,7 +1266,7 @@ steps:
   optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
-  - examples/features/torchrun/torchrun_dp_example_offline.py
+  - examples/offline_inference/torchrun_dp_example.py
   - vllm/config/parallel.py
   - vllm/distributed/
   - vllm/v1/engine/llm_engine.py
@@ -1245,7 +1274,8 @@ steps:
   - vllm/v1/worker/gpu_worker.py
   - vllm/platforms/rocm.py
   commands:
-  - torchrun --nproc-per-node=8 ../examples/features/torchrun/torchrun_dp_example_offline.py --tp-size=2 --pp-size=1 --dp-size=4 --enable-ep
+  - export TORCH_NCCL_BLOCKING_WAIT=1
+  - torchrun --nproc-per-node=8 ../examples/offline_inference/torchrun_dp_example.py --tp-size=2 --pp-size=1 --dp-size=4 --enable-ep
 
 #--------------------------------------------------------  mi300 · entrypoints  --------------------------------------------------------#
 
@@ -1253,7 +1283,6 @@ steps:
   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"
@@ -1274,7 +1303,6 @@ steps:
   agent_pool: mi300_1
   fast_check: true
   torch_nightly: true
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -1288,7 +1316,6 @@ steps:
   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"
@@ -1299,13 +1326,13 @@ steps:
   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/openai/speech_to_text/
   - 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"
@@ -1315,21 +1342,7 @@ steps:
   - 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/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/speech_to_text/ --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
 
 - label: Entrypoints Integration (LLM) # TBD
   timeout_in_minutes: 180
@@ -1634,18 +1647,18 @@ steps:
     - python3 basic/offline_inference/embed.py
     - python3 basic/offline_inference/score.py
     # Multi-modal models
-    - python3 generate/multimodal/audio_language_offline.py --seed 0
-    - python3 generate/multimodal/vision_language_offline.py --seed 0
-    - python3 generate/multimodal/vision_language_multi_image_offline.py --seed 0
-    - python3 generate/multimodal/encoder_decoder_multimodal_offline.py --model-type whisper --seed 0
+    - python3 offline_inference/audio_language.py --seed 0
+    - python3 offline_inference/vision_language.py --seed 0
+    - python3 offline_inference/vision_language_multi_image.py --seed 0
+    - python3 offline_inference/encoder_decoder_multimodal.py --model-type whisper --seed 0
     # Pooling models
     - python3 pooling/embed/vision_embedding_offline.py --seed 0
     # Features demo
-    - python3 features/automatic_prefix_caching/prefix_caching_offline.py
-    - python3 deployment/llm_engine_example.py
-    - python3 features/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 features/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
-    - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
-    - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
+    - python3 offline_inference/prefix_caching.py
+    - python3 offline_inference/llm_engine_example.py
+    - python3 others/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 others/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
+    - python3 offline_inference/spec_decode.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
+    - python3 offline_inference/spec_decode.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
 
 #----------------------------------------------------------  mi300 · kernels  ----------------------------------------------------------#
 
@@ -1706,7 +1719,6 @@ steps:
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
-  optional: true
   parallelism: 2
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -1754,7 +1766,7 @@ steps:
   - export PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True
   - pytest -v -s -x lora/test_chatglm3_tp.py
   - pytest -v -s -x lora/test_llama_tp.py
-  - pytest -v -s -x lora/test_qwen3_with_multi_loras.py
+  - pytest -v -s -x lora/test_llm_with_multi_loras.py
   - pytest -v -s -x lora/test_olmoe_tp.py
   - pytest -v -s -x lora/test_gptoss_tp.py
   - pytest -v -s -x lora/test_qwen35_densemodel_lora.py
@@ -1765,7 +1777,6 @@ steps:
   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/
@@ -1791,17 +1802,15 @@ steps:
   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/
   - tests/models/multimodal/generation
   - tests/models/multimodal/test_mapping.py
   commands:
-  - uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@rocm-7.0-v2.3.0'
-  - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
-  - pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
-
+  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
+  - pytest -v -s models/multimodal/generation -m 'not core_model' --ignore models/multimodal/generation/test_common.py
+  - pytest -v -s models/multimodal/test_mapping.py
 
 - label: Multi-Modal Models (Extended Generation 2) # TBD
   timeout_in_minutes: 180
@@ -1813,10 +1822,8 @@ steps:
   - vllm/
   - tests/models/multimodal/generation
   commands:
-  - uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@rocm-7.0-v2.3.0'
-  - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
-  - pytest -v -s models/language/generation -m '(not core_model) and (not hybrid_model)'
-
+  - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
+  - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=0) and not core_model'
 
 - label: Multi-Modal Models (Extended Generation 3) # TBD
   timeout_in_minutes: 180
@@ -1836,7 +1843,6 @@ steps:
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
   agent_pool: mi300_1
   torch_nightly: true
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -1945,8 +1951,8 @@ steps:
   - pytest -v -s tests/models/multimodal/processing/
   - pytest -v -s tests/models/multimodal/test_mapping.py
   - python3 examples/basic/offline_inference/chat.py
-  - python3 examples/generate/multimodal/vision_language_offline.py --model-type qwen2_5_vl
-  - VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/generate/multimodal/audio_language_offline.py --model-type whisper
+  - python3 examples/offline_inference/vision_language.py --model-type qwen2_5_vl
+  - VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/offline_inference/audio_language.py --model-type whisper
 
 #-------------------------------------------------------  mi300 · quantization  --------------------------------------------------------#
 
@@ -2197,6 +2203,7 @@ steps:
   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/
@@ -2273,6 +2280,7 @@ steps:
   - tests/entrypoints/openai/test_multi_api_servers.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
   - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
   - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
@@ -2292,8 +2300,9 @@ steps:
   - vllm/_aiter_ops.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_async_new_apis.py
-  - VLLM_LOGGING_LEVEL=DEBUG python3 examples/features/data_parallel/data_parallel_offline.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=deepep_high_throughput
+  - VLLM_LOGGING_LEVEL=DEBUG python3 examples/offline_inference/data_parallel.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=deepep_high_throughput
   - pytest -v -s tests/v1/distributed/test_dbo.py
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 pytest -v -s tests/distributed/test_weight_transfer.py
   - pytest -v -s tests/distributed/test_packed_tensor.py
@@ -2354,6 +2363,7 @@ steps:
   - tests/distributed/test_utils
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - TP_SIZE=2 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
   - TP_SIZE=2 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
   - TP_SIZE=2 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
@@ -2483,6 +2493,7 @@ steps:
   - tests/entrypoints/llm/test_collective_rpc.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s entrypoints/llm/test_collective_rpc.py
   - pytest -v -s ./compile/fullgraph/test_basic_correctness.py
   - pytest -v -s ./compile/test_wrapper.py
@@ -2507,6 +2518,7 @@ steps:
   - tests/v1/worker/test_worker_memory_snapshot.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - VLLM_TEST_SAME_HOST=1 torchrun --nproc-per-node=4 distributed/test_same_node.py | grep 'Same node test passed'
   - VLLM_TEST_SAME_HOST=1 VLLM_TEST_WITH_DEFAULT_DEVICE_SET=1 torchrun --nproc-per-node=4 distributed/test_same_node.py | grep 'Same node test passed'
   - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s v1/shutdown
@@ -2527,6 +2539,7 @@ steps:
   - tests/distributed/test_multiproc_executor.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s compile/fullgraph/test_basic_correctness.py
   - pytest -v -s distributed/test_pynccl.py
   - pytest -v -s distributed/test_events.py
@@ -2614,7 +2627,6 @@ steps:
   agent_pool: mi325_1
   torch_nightly: true
   parallelism: 2
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -2640,7 +2652,6 @@ steps:
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
   agent_pool: mi325_1
   torch_nightly: true
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -2702,10 +2713,11 @@ steps:
   - vllm/v1/attention/selector.py
   - tests/distributed/test_context_parallel.py
   - tests/v1/distributed/test_dbo.py
-  - examples/features/data_parallel/data_parallel_offline.py
+  - examples/offline_inference/data_parallel.py
   - vllm/_aiter_ops.py
   - vllm/platforms/rocm.py
   commands:
+  - export TORCH_NCCL_BLOCKING_WAIT=1
   - pytest -v -s tests/distributed/test_context_parallel.py
   - pytest -v -s tests/v1/distributed/test_dbo.py
 
@@ -2736,7 +2748,6 @@ steps:
   agent_pool: mi355_1
   fast_check: true
   torch_nightly: true
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -2752,7 +2763,6 @@ steps:
   agent_pool: mi355_1
   fast_check: true
   torch_nightly: true
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -2761,6 +2771,7 @@ steps:
   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/openai/speech_to_text/
   - pytest -v -s entrypoints/test_chat_utils.py
 
 - label: Entrypoints Integration (API Server openai - Part 3) # TBD
@@ -2776,21 +2787,7 @@ steps:
   - 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, amdmi355]
-  agent_pool: mi355_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/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/speech_to_text/ --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
 
 - label: Entrypoints Integration (Pooling) # TBD
   timeout_in_minutes: 180
@@ -2933,18 +2930,18 @@ steps:
   - python3 basic/offline_inference/embed.py
   - python3 basic/offline_inference/score.py
   # Multi-modal models
-  - python3 generate/multimodal/audio_language_offline.py --seed 0
-  - python3 generate/multimodal/vision_language_offline.py --seed 0
-  - python3 generate/multimodal/vision_language_multi_image_offline.py --seed 0
-  - python3 generate/multimodal/encoder_decoder_multimodal_offline.py --model-type whisper --seed 0
+  - python3 offline_inference/audio_language.py --seed 0
+  - python3 offline_inference/vision_language.py --seed 0
+  - python3 offline_inference/vision_language_multi_image.py --seed 0
+  - python3 offline_inference/encoder_decoder_multimodal.py --model-type whisper --seed 0
   # Pooling models
   - python3 pooling/embed/vision_embedding_offline.py --seed 0
   # Features demo
-  - python3 features/automatic_prefix_caching/prefix_caching_offline.py
-  - python3 deployment/llm_engine_example.py
-  - python3 features/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 features/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
-  - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
-  - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
+  - python3 offline_inference/prefix_caching.py
+  - python3 offline_inference/llm_engine_example.py
+  - python3 others/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 others/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
+  - python3 offline_inference/spec_decode.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
+  - python3 offline_inference/spec_decode.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
 
 #----------------------------------------------------------  mi355 · kernels  ----------------------------------------------------------#
 
@@ -3054,7 +3051,7 @@ steps:
   - vllm/
   - tests/models/language/generation
   commands:
-  - uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@rocm-7.0-v2.3.0'
+  - uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@fix-rocm-7.0-warp-size-constexpr'
   - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
   - pytest -v -s models/language/generation -m '(not core_model) and (not hybrid_model)'
 
@@ -3062,7 +3059,6 @@ steps:
   timeout_in_minutes: 180
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -3262,7 +3258,6 @@ steps:
   timeout_in_minutes: 60
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -3288,7 +3283,6 @@ steps:
   timeout_in_minutes: 60
   mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
   agent_pool: mi355_1
-  optional: true
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/
@@ -3329,7 +3323,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
+  - ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
 
 - label: Distributed NixlConnector PD accuracy (4 GPUs) # TBD
   timeout_in_minutes: 180

.buildkite/test_areas/attention.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: V1 attention (H100)
-  key: v1-attention-h100
   timeout_in_minutes: 30
   device: h100
   source_file_dependencies:
@@ -15,9 +14,8 @@ steps:
     - pytest -v -s v1/attention
 
 - label: V1 attention (B200)
-  key: v1-attention-b200
   timeout_in_minutes: 30
-  device: b200-k8s
+  device: b200
   source_file_dependencies:
     - vllm/config/attention.py
     - vllm/model_executor/layers/attention

.buildkite/test_areas/basic_correctness.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: Basic Correctness
-  key: basic-correctness
   timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:

.buildkite/test_areas/benchmarks.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: Benchmarks CLI Test
-  key: benchmarks-cli-test
   timeout_in_minutes: 20
   device: h200_18gb
   source_file_dependencies:
@@ -13,8 +12,7 @@ steps:
   - pytest -v -s benchmarks/
 
 - label: Attention Benchmarks Smoke Test (B200)
-  key: attention-benchmarks-smoke-test-b200
-  device: b200-k8s
+  device: b200
   num_gpus: 2
   optional: true
   working_dir: "/vllm-workspace/"

.buildkite/test_areas/compile.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: Sequence Parallel Correctness Tests (2 GPUs)
-  key: sequence-parallel-correctness-tests-2-gpus
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/"
   num_devices: 2
@@ -18,7 +17,6 @@ steps:
   - pytest -v -s tests/compile/correctness_e2e/test_sequence_parallel.py
 
 - label: Sequence Parallel Correctness Tests (2xH100)
-  key: sequence-parallel-correctness-tests-2xh100
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/"
   device: h100
@@ -29,7 +27,6 @@ steps:
   - pytest -v -s tests/compile/correctness_e2e/test_sequence_parallel.py
 
 - label: AsyncTP Correctness Tests (2xH100)
-  key: asynctp-correctness-tests-2xh100
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/"
   device: h100
@@ -40,10 +37,9 @@ steps:
   - pytest -v -s tests/compile/correctness_e2e/test_async_tp.py
 
 - label: AsyncTP Correctness Tests (B200)
-  key: asynctp-correctness-tests-b200
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/"
-  device: b200-k8s
+  device: b200
   optional: true
   num_devices: 2
   commands:
@@ -51,7 +47,6 @@ steps:
   - pytest -v -s tests/compile/correctness_e2e/test_async_tp.py
 
 - label: Distributed Compile Unit Tests (2xH100)
-  key: distributed-compile-unit-tests-2xh100
   timeout_in_minutes: 20
   working_dir: "/vllm-workspace/"
   device: h100
@@ -65,10 +60,9 @@ steps:
   - pytest -s -v tests/compile/passes/distributed
 
 - label: Fusion and Compile Unit Tests (2xB200)
-  key: fusion-and-compile-unit-tests-2xb200
   timeout_in_minutes: 20
   working_dir: "/vllm-workspace/"
-  device: b200-k8s
+  device: b200
   source_file_dependencies:
   - csrc/quantization/fp4/
   - vllm/model_executor/layers/quantization/
@@ -95,7 +89,6 @@ steps:
     - pytest -v -s tests/compile/fullgraph/test_full_graph.py::test_fp8_kv_scale_compile
 
 - label: Fusion E2E Quick (H100)
-  key: fusion-e2e-quick-h100
   timeout_in_minutes: 15
   working_dir: "/vllm-workspace/"
   device: h100
@@ -114,7 +107,6 @@ steps:
     - pytest -v -s tests/compile/fusions_e2e/test_tp1_quant.py -k "inductor_partition and not +rms_norm and +quant_fp8 and (qwen3 or deepseek)"
 
 - label: Fusion E2E Config Sweep (H100)
-  key: fusion-e2e-config-sweep-h100
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/"
   device: h100
@@ -134,10 +126,9 @@ steps:
     - pytest -v -s tests/compile/fusions_e2e/test_tp1_quant.py -k "llama-3"
 
 - label: Fusion E2E Config Sweep (B200)
-  key: fusion-e2e-config-sweep-b200
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/"
-  device: b200-k8s
+  device: b200
   num_devices: 1
   optional: true
   commands:
@@ -148,7 +139,6 @@ steps:
     - pytest -v -s tests/compile/fusions_e2e/test_tp1_quant.py -k "inductor_partition and (FLASHINFER and not +rms_norm and (not +quant_fp8 or +quant_fp8 and (qwen3 or deepseek)) or llama-3)"
 
 - label: Fusion E2E TP2 Quick (H100)
-  key: fusion-e2e-tp2-quick-h100
   timeout_in_minutes: 20
   working_dir: "/vllm-workspace/"
   device: h100
@@ -166,7 +156,6 @@ steps:
     - pytest -v -s tests/compile/fusions_e2e/test_tp2_async_tp.py -k "inductor_partition and not +rms_norm and (not +quant_fp8 or +quant_fp8 and (qwen3 or deepseek))"
 
 - label: Fusion E2E TP2 AR-RMS Config Sweep (H100)
-  key: fusion-e2e-tp2-ar-rms-config-sweep-h100
   timeout_in_minutes: 40
   working_dir: "/vllm-workspace/"
   device: h100
@@ -186,7 +175,6 @@ steps:
     - pytest -v -s tests/compile/fusions_e2e/test_tp2_ar_rms.py -k "llama-3"
 
 - label: Fusion E2E TP2 AsyncTP Config Sweep (H100)
-  key: fusion-e2e-tp2-asynctp-config-sweep-h100
   timeout_in_minutes: 40
   working_dir: "/vllm-workspace/"
   device: h100
@@ -206,10 +194,9 @@ steps:
     - pytest -v -s tests/compile/fusions_e2e/test_tp2_async_tp.py -k "llama-3"
 
 - label: Fusion E2E TP2 (B200)
-  key: fusion-e2e-tp2-b200
   timeout_in_minutes: 20
   working_dir: "/vllm-workspace/"
-  device: b200-k8s
+  device: b200
   num_devices: 2
   source_file_dependencies:
     - csrc/quantization/

.buildkite/test_areas/cuda.yaml

@@ -3,22 +3,16 @@ depends_on:
   - image-build
 steps:
 - label: Platform Tests (CUDA)
-  key: platform-tests-cuda
   timeout_in_minutes: 15
   device: h200_18gb
   source_file_dependencies:
-  - vllm/envs.py
-  - vllm/logger.py
-  - vllm/platforms/
-  - vllm/plugins/
-  - vllm/utils/
+  - vllm/
   - tests/cuda
   commands:
     - pytest -v -s cuda/test_cuda_context.py
     - pytest -v -s cuda/test_platform_no_cuda_init.py
 
 - label: Cudagraph
-  key: cudagraph
   timeout_in_minutes: 20
   source_file_dependencies:
   - tests/v1/cudagraph

.buildkite/test_areas/disaggregated.yaml

@@ -3,71 +3,65 @@ depends_on:
   - image-build
 steps:
 - label: Distributed NixlConnector PD accuracy (4 GPUs)
-  key: distributed-nixlconnector-pd-accuracy-4-gpus
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - tests/v1/kv_connector/nixl_integration/
   commands:
     - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
     - 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
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - tests/v1/kv_connector/nixl_integration/
   commands:
     - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
     - FLASHINFER=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
 - label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs)
-  key: dp-ep-distributed-nixlconnector-pd-accuracy-tests-4-gpus
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - tests/v1/kv_connector/nixl_integration/
   commands:
     - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
     - 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)
-  key: crosslayer-kv-layout-distributed-nixlconnector-pd-accuracy-tests-4-gpus
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - tests/v1/kv_connector/nixl_integration/
   commands:
     - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
     - CROSS_LAYERS_BLOCKS=True bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
 - label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs)
-  key: hybrid-ssm-nixlconnector-pd-accuracy-tests-4-gpus
-  timeout_in_minutes: 25
+  timeout_in_minutes: 20
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - tests/v1/kv_connector/nixl_integration/
   commands:
     - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
     - HYBRID_SSM=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
 
 - label: MultiConnector (Nixl+Offloading) PD accuracy (2 GPUs)
-  key: multiconnector-nixl-offloading-pd-accuracy-2-gpus
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - vllm/distributed/kv_transfer/kv_connector/v1/multi_connector.py
     - vllm/distributed/kv_transfer/kv_connector/v1/offloading_connector.py
     - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
@@ -77,13 +71,12 @@ steps:
     - bash v1/kv_connector/nixl_integration/run_multi_connector_accuracy_test.sh
 
 - label: NixlConnector PD + Spec Decode acceptance (2 GPUs)
-  key: nixlconnector-pd-spec-decode-acceptance-2-gpus
   timeout_in_minutes: 30
   device: a100
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - vllm/v1/worker/kv_connector_model_runner_mixin.py
     - tests/v1/kv_connector/nixl_integration/
   commands:
@@ -91,12 +84,11 @@ steps:
     - bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
 
 - label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs)
-  key: multiconnector-nixl-offloading-pd-edge-cases-2-gpus
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
   source_file_dependencies:
-    - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+    - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
     - vllm/distributed/kv_transfer/kv_connector/v1/multi_connector.py
     - vllm/distributed/kv_transfer/kv_connector/v1/offloading_connector.py
     - vllm/distributed/kv_transfer/kv_connector/v1/offloading/

.buildkite/test_areas/distributed.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: Distributed Comm Ops
-  key: distributed-comm-ops
   timeout_in_minutes: 20
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
@@ -17,7 +16,6 @@ steps:
   - pytest -v -s distributed/test_shm_storage.py
 
 - label: Distributed DP Tests (2 GPUs)
-  key: distributed-dp-tests-2-gpus
   timeout_in_minutes: 20
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
@@ -39,7 +37,6 @@ steps:
   - DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py
 
 - label: Distributed Compile + RPC Tests (2 GPUs)
-  key: distributed-compile-rpc-tests-2-gpus
   timeout_in_minutes: 20
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
@@ -62,7 +59,6 @@ steps:
   - pytest -v -s ./compile/test_wrapper.py
 
 - label: Distributed Torchrun + Shutdown Tests (2 GPUs)
-  key: distributed-torchrun-shutdown-tests-2-gpus
   timeout_in_minutes: 20
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
@@ -85,7 +81,6 @@ steps:
   - pytest -v -s v1/worker/test_worker_memory_snapshot.py
 
 - label: Distributed Torchrun + Examples (4 GPUs)
-  key: distributed-torchrun-examples-4-gpus
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace"
   num_devices: 4
@@ -93,8 +88,9 @@ steps:
   - vllm/distributed/
   - tests/distributed/test_torchrun_example.py
   - tests/distributed/test_torchrun_example_moe.py
+  - examples/offline_inference/rlhf_colocate.py
   - examples/rl/
-  - tests/examples/features/data_parallel/data_parallel_offline.py
+  - tests/examples/offline_inference/data_parallel.py
   commands:
   # https://github.com/NVIDIA/nccl/issues/1838
   - export NCCL_CUMEM_HOST_ENABLE=0
@@ -111,13 +107,12 @@ steps:
   # test with torchrun tp=2 and dp=2 with ep
   - TP_SIZE=2 DP_SIZE=2 ENABLE_EP=1 torchrun --nproc-per-node=4 tests/distributed/test_torchrun_example_moe.py
   # test with internal dp
-  - python3 examples/features/data_parallel/data_parallel_offline.py --enforce-eager
+  - python3 examples/offline_inference/data_parallel.py --enforce-eager
   # rlhf examples
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_nccl.py
   - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_ipc.py
 
 - label: Distributed DP Tests (4 GPUs)
-  key: distributed-dp-tests-4-gpus
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
@@ -138,7 +133,6 @@ steps:
   - pytest -v -s distributed/test_utils.py
 
 - label: Distributed Compile + Comm (4 GPUs)
-  key: distributed-compile-comm-4-gpus
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
@@ -160,28 +154,24 @@ steps:
   - pytest -v -s distributed/test_multiproc_executor.py::test_multiproc_executor_multi_node
 
 - label: Distributed Tests (8 GPUs)(H100)
-  key: distributed-tests-8-gpus-h100
   timeout_in_minutes: 10
   device: h100
   num_devices: 8
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
-  - examples/features/torchrun/torchrun_dp_example_offline.py
+  - examples/offline_inference/torchrun_dp_example.py
   - vllm/config/parallel.py
   - vllm/distributed/
   - vllm/v1/engine/llm_engine.py
   - vllm/v1/executor/uniproc_executor.py
   - vllm/v1/worker/gpu_worker.py
-  - tests/distributed/test_mnnvl_alltoall.py
-
   commands:
   # https://github.com/NVIDIA/nccl/issues/1838
   - export NCCL_CUMEM_HOST_ENABLE=0
   # test with torchrun tp=2 and dp=4 with ep
-  - torchrun --nproc-per-node=8 ../examples/features/torchrun/torchrun_dp_example_offline.py --tp-size=2 --pp-size=1 --dp-size=4 --enable-ep
+  - torchrun --nproc-per-node=8 ../examples/offline_inference/torchrun_dp_example.py --tp-size=2 --pp-size=1 --dp-size=4 --enable-ep
 
 - label: Distributed Tests (4 GPUs)(A100)
-  key: distributed-tests-4-gpus-a100
   device: a100
   optional: true
   num_devices: 4
@@ -196,7 +186,6 @@ steps:
   - pytest -v -s -x lora/test_mixtral.py
 
 - label: Distributed Tests (2 GPUs)(H100)
-  key: distributed-tests-2-gpus-h100
   timeout_in_minutes: 15
   device: h100
   optional: true
@@ -205,14 +194,13 @@ steps:
   commands:
     - pytest -v -s tests/distributed/test_context_parallel.py
     - VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_async_new_apis.py
-    - VLLM_USE_DEEP_GEMM=1 VLLM_LOGGING_LEVEL=DEBUG python3 examples/features/data_parallel/data_parallel_offline.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=deepep_high_throughput
+    - VLLM_USE_DEEP_GEMM=1 VLLM_LOGGING_LEVEL=DEBUG python3 examples/offline_inference/data_parallel.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=deepep_high_throughput
     - pytest -v -s tests/v1/distributed/test_dbo.py
     - VLLM_ALLOW_INSECURE_SERIALIZATION=1 pytest -v -s tests/distributed/test_weight_transfer.py
     - pytest -v -s tests/distributed/test_packed_tensor.py
 
 - label: Distributed Tests (2 GPUs)(B200)
-  key: distributed-tests-2-gpus-b200
-  device: b200-k8s
+  device: b200
   optional: true
   working_dir: "/vllm-workspace/"
   num_devices: 2
@@ -220,12 +208,8 @@ steps:
     - pytest -v -s tests/distributed/test_context_parallel.py
     - pytest -v -s tests/distributed/test_nccl_symm_mem_allreduce.py
     - pytest -v -s tests/v1/distributed/test_dbo.py
-    - pytest -v -s tests/distributed/test_mnnvl_alltoall.py
-
-    
 
 - label: 2 Node Test (4 GPUs)
-  key: 2-node-test-4-gpus
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
@@ -238,12 +222,11 @@ steps:
   - vllm/executor/
   - vllm/model_executor/models/
   - tests/distributed/
-  - tests/examples/features/data_parallel/data_parallel_offline.py
+  - tests/examples/offline_inference/data_parallel.py
   commands:
-    - ./.buildkite/scripts/run-multi-node-test.sh /vllm-workspace/tests 2 2 $IMAGE_TAG "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/features/data_parallel/data_parallel_offline.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=0 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code && VLLM_MULTI_NODE=1 pytest -v -s distributed/test_multi_node_assignment.py && VLLM_MULTI_NODE=1 pytest -v -s distributed/test_pipeline_parallel.py" "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/features/data_parallel/data_parallel_offline.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=1 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code"
+    - ./.buildkite/scripts/run-multi-node-test.sh /vllm-workspace/tests 2 2 $IMAGE_TAG "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/offline_inference/data_parallel.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=0 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code && VLLM_MULTI_NODE=1 pytest -v -s distributed/test_multi_node_assignment.py && VLLM_MULTI_NODE=1 pytest -v -s distributed/test_pipeline_parallel.py" "VLLM_TEST_SAME_HOST=0 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_same_node.py | grep 'Same node test passed' && NUM_NODES=2 torchrun --nnodes 2 --nproc-per-node=2 --rdzv_backend=c10d --rdzv_endpoint=192.168.10.10 distributed/test_node_count.py | grep 'Node count test passed' && python3 ../examples/offline_inference/data_parallel.py -dp=2 -tp=1 --dp-num-nodes=2 --dp-node-rank=1 --dp-master-addr=192.168.10.10 --dp-master-port=12345 --enforce-eager --trust-remote-code"
 
 - label: Pipeline + Context Parallelism (4 GPUs)
-  key: pipeline-context-parallelism-4-gpus
   timeout_in_minutes: 60
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
@@ -258,7 +241,6 @@ steps:
   - pytest -v -s distributed/test_pipeline_parallel.py
 
 - label: RayExecutorV2 (4 GPUs)
-  key: rayexecutorv2-4-gpus
   timeout_in_minutes: 60
   working_dir: "/vllm-workspace/tests"
   num_devices: 4

.buildkite/test_areas/docker.yaml

@@ -1,16 +0,0 @@
-group: Docker
-depends_on:
-  - image-build-cpu
-steps:
-- label: Docker Build Metadata
-  timeout_in_minutes: 10
-  device: cpu-small
-  source_file_dependencies:
-    - .buildkite/release-pipeline.yaml
-    - .buildkite/scripts/docker-build-metadata-args.sh
-    - docker/Dockerfile
-    - docker/Dockerfile.cpu
-    - docker/docker-bake.hcl
-    - tests/tools/test_docker_build_metadata_args.py
-  commands:
-    - pytest -v -s tools/test_docker_build_metadata_args.py

.buildkite/test_areas/e2e_integration.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: DeepSeek V2-Lite Accuracy
-  key: deepseek-v2-lite-accuracy
   timeout_in_minutes: 60
   device: h100
   optional: true
@@ -13,7 +12,6 @@ steps:
   - 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
   timeout_in_minutes: 60
   device: h100
   optional: true
@@ -23,9 +21,8 @@ steps:
   - 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
   timeout_in_minutes: 60
-  device: b200-k8s
+  device: b200
   optional: true
   num_devices: 2
   working_dir: "/vllm-workspace"
@@ -33,7 +30,6 @@ steps:
   - bash .buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_block_ep_eplb.sh 0.8 200 8020 2 1
 
 - label: Qwen3-30B-A3B-FP8 DP4 Async EPLB Accuracy
-  key: qwen3-30b-a3b-fp8-dp4-async-eplb-accuracy
   timeout_in_minutes: 60
   device: h100
   optional: true
@@ -43,7 +39,6 @@ steps:
   - bash .buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_dp4_async_eplb.sh 0.8 200 8050
 
 - label: DeepSeek V2-Lite Prefetch Offload Accuracy (H100)
-  key: deepseek-v2-lite-prefetch-offload-accuracy-h100
   timeout_in_minutes: 60
   device: h100
   optional: true

.buildkite/test_areas/engine.yaml

@@ -3,32 +3,19 @@ depends_on:
   - image-build
 steps:
 - label: Engine
-  key: engine
   timeout_in_minutes: 15
   device: h200_18gb
   source_file_dependencies:
-  - vllm/compilation/
-  - vllm/config/
-  - vllm/engine/
-  - vllm/entrypoints/logger.py
-  - vllm/envs.py
-  - vllm/logger.py
-  - vllm/logging_utils/
-  - vllm/platforms/
-  - vllm/sequence.py
-  - vllm/triton_utils/
-  - vllm/utils/
+  - vllm/
   - tests/engine
   - tests/test_sequence
   - tests/test_config
   - tests/test_logger
   - tests/test_vllm_port
-  - 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
+  - pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py
 
 - label: Engine (1 GPU)
-  key: engine-1-gpu
   timeout_in_minutes: 30
   source_file_dependencies:
     - vllm/v1/engine/
@@ -36,15 +23,8 @@ steps:
   commands:
     - pytest -v -s v1/engine/test_preprocess_error_handling.py
     - pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py
-  mirror:
-    amd:
-      device: mi300_1
-      timeout_in_minutes: 40
-      depends_on:
-      - image-build-amd
 
 - label: e2e Scheduling (1 GPU)
-  key: e2e-scheduling-1-gpu
   timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
@@ -52,15 +32,8 @@ steps:
     - tests/v1/e2e/general/
   commands:
     - pytest -v -s v1/e2e/general/test_async_scheduling.py
-  mirror:
-    amd:
-      device: mi250_1
-      timeout_in_minutes: 60
-      depends_on:
-      - image-build-amd
 
 - label: e2e Core (1 GPU)
-  key: e2e-core-1-gpu
   timeout_in_minutes: 30
   source_file_dependencies:
     - vllm/v1/
@@ -69,63 +42,38 @@ steps:
     - pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py
 
 - label: V1 e2e (2 GPUs)
-  key: v1-e2e-2-gpus
   timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
   optional: true
   num_devices: 2
   source_file_dependencies:
-    - vllm/compilation/
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/engine/
-    - vllm/envs.py
-    - vllm/forward_context.py
-    - vllm/inputs/
-    - vllm/logger.py
-    - vllm/logging_utils/
-    - vllm/model_executor/
-    - vllm/multimodal/
-    - vllm/platforms/
-    - vllm/sampling_params.py
-    - vllm/transformers_utils/
-    - vllm/triton_utils/
-    - vllm/utils/
-    - vllm/v1/
-    - tests/v1/e2e/spec_decode
+    - vllm/
+    - tests/v1/e2e
   commands:
     # Only run tests that need exactly 2 GPUs
     - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "tensor_parallelism"
+  mirror:
+    amd:
+      device: mi325_2
+      depends_on:
+      - image-build-amd
 
 - label: V1 e2e (4 GPUs)
-  key: v1-e2e-4-gpus
   timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
   optional: true
   num_devices: 4
   source_file_dependencies:
-    - vllm/compilation/
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/engine/
-    - vllm/envs.py
-    - vllm/forward_context.py
-    - vllm/inputs/
-    - vllm/logger.py
-    - vllm/logging_utils/
-    - vllm/model_executor/
-    - vllm/multimodal/
-    - vllm/platforms/
-    - vllm/sampling_params.py
-    - vllm/transformers_utils/
-    - vllm/triton_utils/
-    - vllm/utils/
-    - vllm/v1/
-    - tests/v1/e2e/spec_decode
+    - vllm/
+    - tests/v1/e2e
   commands:
     # Only run tests that need 4 GPUs
     - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "eagle_correctness_heavy"
+  mirror:
+    amd:
+      device: mi325_4
+      depends_on:
+      - image-build-amd
 
 - label: V1 e2e (4xH100)
-  key: v1-e2e-4xh100
   timeout_in_minutes: 60
   device: h100
   num_devices: 4

.buildkite/test_areas/entrypoints.yaml

@@ -2,8 +2,7 @@ group: Entrypoints
 depends_on: 
   - image-build
 steps:
-- label: Entrypoints Unit Tests
-  key: entrypoints-unit-tests
+- label: Entrypoints Unit Tests  
   timeout_in_minutes: 10
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -11,10 +10,9 @@ steps:
   - tests/entrypoints/
   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/ --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
 
 - label: Entrypoints Integration (LLM)
-  key: entrypoints-integration-llm
   timeout_in_minutes: 40
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -26,14 +24,8 @@ steps:
   - 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 # it needs a clean process
   - pytest -v -s entrypoints/offline_mode # Needs to avoid interference with other tests
-  mirror:
-    amd:
-      device: mi300_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"
   source_file_dependencies:
@@ -45,13 +37,12 @@ steps:
   - 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
   mirror:
     amd:
-      device: mi300_1
-      timeout_in_minutes: 80
+      device: mi325_1
       depends_on:
       - image-build-amd
 
+
 - label: Entrypoints Integration (API Server openai - Part 2)
-  key: entrypoints-integration-api-server-openai-part-2
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -60,16 +51,15 @@ steps:
   - tests/entrypoints/test_chat_utils
   commands:
   - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
+  - pytest -v -s entrypoints/openai/speech_to_text/
   - pytest -v -s entrypoints/test_chat_utils.py
   mirror:
     amd:
-      device: mi300_1
-      timeout_in_minutes: 60
+      device: mi325_1
       depends_on:
       - image-build-amd
 
 - label: Entrypoints Integration (API Server openai - Part 3)
-  key: entrypoints-integration-api-server-openai-part-3
   timeout_in_minutes: 50
   device: h200_18gb
   working_dir: "/vllm-workspace/tests"
@@ -79,16 +69,9 @@ steps:
   - 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
-  mirror:
-    amd:
-      device: mi300_1
-      timeout_in_minutes: 60
-      depends_on:
-      - image-build-amd
+  - pytest -v -s entrypoints/openai --ignore=entrypoints/openai/chat_completion --ignore=entrypoints/openai/completion --ignore=entrypoints/openai/speech_to_text/ --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses --ignore=entrypoints/openai/test_multi_api_servers.py
 
 - label: Entrypoints Integration (API Server 2)
-  key: entrypoints-integration-api-server-2
   timeout_in_minutes: 130
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -101,25 +84,8 @@ steps:
   - pytest -v -s entrypoints/serve/instrumentator
   - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
   - pytest -v -s tool_use
-  mirror:
-    amd:
-      device: mi300_1
-      depends_on:
-      - image-build-amd
-
-- label: Entrypoints Integration (Speech to Text)
-  key: entrypoints-integration-speech_to_text
-  timeout_in_minutes: 50
-  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
 
 - label: Entrypoints Integration (Pooling)
-  key: entrypoints-integration-pooling
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
@@ -130,7 +96,6 @@ steps:
   - 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:
@@ -140,12 +105,11 @@ steps:
   - pytest -v -s entrypoints/openai/responses
 
 - label: OpenAI API Correctness
-  key: openai-api-correctness
   timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
   - csrc/
   - vllm/entrypoints/openai/
   - vllm/model_executor/models/whisper.py
-  commands: # LMEval
+  commands: # LMEval+Transcription WER check
   - pytest -s entrypoints/openai/correctness/

.buildkite/test_areas/expert_parallelism.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: EPLB Algorithm
-  key: eplb-algorithm
   timeout_in_minutes: 15
   device: h200_18gb
   working_dir: "/vllm-workspace/tests"
@@ -16,7 +15,6 @@ steps:
   - pytest -v -s distributed/test_eplb_utils.py
 
 - label: EPLB Execution # 17min
-  key: eplb-execution
   timeout_in_minutes: 27
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
@@ -28,7 +26,6 @@ steps:
   - pytest -v -s distributed/test_eplb_spec_decode.py
 
 - label: Elastic EP Scaling Test
-  key: elastic-ep-scaling-test
   timeout_in_minutes: 20
   device: h100
   working_dir: "/vllm-workspace/tests"

.buildkite/test_areas/kernels.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: vLLM IR Tests
-  key: vllm-ir-tests
   timeout_in_minutes: 10
   device: h200_18gb
   working_dir: "/vllm-workspace/"
@@ -15,7 +14,6 @@ steps:
     - pytest -v -s tests/kernels/ir
 
 - label: Kernels Core Operation Test
-  key: kernels-core-operation-test
   timeout_in_minutes: 75
   source_file_dependencies:
   - csrc/
@@ -25,7 +23,6 @@ steps:
     - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py  kernels/test_concat_mla_q.py
 
 - label: Kernels MiniMax Reduce RMS Test (2 GPUs)
-  key: kernels-minimax-reduce-rms-test-2-gpus
   timeout_in_minutes: 15
   num_devices: 2
   device: h100
@@ -39,7 +36,6 @@ steps:
     - pytest -v -s kernels/core/test_minimax_reduce_rms.py
 
 - label: Kernels Attention Test %N
-  key: kernels-attention-test
   timeout_in_minutes: 35
   source_file_dependencies:
   - csrc/attention/
@@ -53,7 +49,6 @@ steps:
   parallelism: 2
 
 - label: Kernels Quantization Test %N
-  key: kernels-quantization-test
   timeout_in_minutes: 90
   source_file_dependencies:
   - csrc/quantization/
@@ -62,22 +57,8 @@ steps:
   commands:
     - pytest -v -s kernels/quantization --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
   parallelism: 2
-  mirror:
-    amd:
-      device: mi300_1
-      source_file_dependencies:
-      - csrc/quantization/
-      - vllm/model_executor/layers/quantization
-      - tests/kernels/quantization
-      - tests/kernels/quantization/test_rocm_skinny_gemms.py
-      - vllm/_aiter_ops.py
-      - vllm/platforms/rocm.py
-      - vllm/model_executor/kernels/
-      depends_on:
-      - image-build-amd
 
 - label: Kernels MoE Test %N
-  key: kernels-moe-test
   timeout_in_minutes: 25
   source_file_dependencies:
   - csrc/quantization/cutlass_w8a8/moe/
@@ -93,7 +74,6 @@ steps:
   parallelism: 5
 
 - label: Kernels Mamba Test
-  key: kernels-mamba-test
   timeout_in_minutes: 45
   source_file_dependencies:
   - csrc/mamba/
@@ -102,28 +82,12 @@ steps:
   commands:
     - pytest -v -s kernels/mamba
 
-- label: Kernels KDA Test
-  timeout_in_minutes: 20
-  device: h200_18gb
-  source_file_dependencies:
-  - vllm/model_executor/layers/fla/ops/kda.py
-  - vllm/model_executor/layers/fla/ops/chunk_delta_h.py
-  - vllm/model_executor/layers/fla/ops/l2norm.py
-  - tests/kernels/test_kda.py
-  commands:
-    - pytest -v -s kernels/test_kda.py
-
 - label: Kernels DeepGEMM Test (H100)
-  key: kernels-deepgemm-test-h100
   timeout_in_minutes: 45
   device: h100
   num_devices: 1
   source_file_dependencies:
-  - cmake/external_projects/deepgemm.cmake
   - tools/install_deepgemm.sh
-  - tools/build_deepgemm_C.py
-  - tools/setup_deepgemm_pythons.sh
-  - tools/check_wheel_deepgemm.py
   - vllm/utils/deep_gemm.py
   - vllm/model_executor/layers/fused_moe
   - vllm/model_executor/layers/quantization
@@ -131,20 +95,16 @@ steps:
   - tests/kernels/moe/test_deepgemm.py
   - tests/kernels/moe/test_batched_deepgemm.py
   - tests/kernels/attention/test_deepgemm_attention.py
-  - tests/quantization/test_cutlass_w4a16.py
   commands:
-    - python3 ../tools/check_wheel_deepgemm.py
     - pytest -v -s kernels/quantization/test_block_fp8.py
     - pytest -v -s kernels/moe/test_deepgemm.py
     - pytest -v -s kernels/moe/test_batched_deepgemm.py
     - pytest -v -s kernels/attention/test_deepgemm_attention.py
-    - pytest -v -s quantization/test_cutlass_w4a16.py
 
 - label: Kernels (B200)
-  key: kernels-b200
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/"
-  device: b200-k8s
+  device: b200
   # optional: true
   source_file_dependencies:
   - csrc/quantization/fp4/
@@ -190,7 +150,6 @@ steps:
     - pytest -v -s tests/models/quantization/test_nvfp4.py
 
 - label: Kernels Helion Test
-  key: kernels-helion-test
   timeout_in_minutes: 30
   device: h100
   source_file_dependencies:
@@ -202,7 +161,6 @@ steps:
 
  
 - label: Kernels FP8 MoE Test (1 H100)
-  key: kernels-fp8-moe-test-1-h100
   timeout_in_minutes: 90
   device: h100
   num_devices: 1
@@ -219,7 +177,6 @@ steps:
     - pytest -v -s kernels/moe/test_triton_moe_ptpc_fp8.py
 
 - label: Kernels FP8 MoE Test (2 H100s)
-  key: kernels-fp8-moe-test-2-h100s
   timeout_in_minutes: 90
   device: h100
   num_devices: 2
@@ -229,9 +186,8 @@ steps:
     - pytest -v -s kernels/moe/test_deepep_moe.py
 
 - label: Kernels Fp4 MoE Test (B200)
-  key: kernels-fp4-moe-test-b200
   timeout_in_minutes: 60
-  device: b200-k8s
+  device: b200
   num_devices: 1
   optional: true
   commands:
@@ -242,7 +198,6 @@ steps:
 
 
 - label: Kernels FusedMoE Layer Test (2 H100s)
-  key: kernels-fusedmoe-layer-test-2-h100s
   timeout_in_minutes: 90
   device: h100
   num_devices: 2
@@ -259,9 +214,8 @@ steps:
 
 
 - label: Kernels FusedMoE Layer Test (2 B200s)
-  key: kernels-fusedmoe-layer-test-2-b200s
   timeout_in_minutes: 90
-  device: b200-k8s
+  device: b200
   num_devices: 2
   source_file_dependencies:
   - csrc/quantization/cutlass_w8a8/moe/

.buildkite/test_areas/lm_eval.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: LM Eval Small Models
-  key: lm-eval-small-models
   timeout_in_minutes: 75
   source_file_dependencies:
   - csrc/
@@ -25,7 +24,6 @@ steps:
 #   - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large.txt --tp-size=4
 
 - label: LM Eval Large Models (4 GPUs)(H100)
-  key: lm-eval-large-models-4-gpus-h100
   device: h100
   optional: true
   num_devices: 4
@@ -38,9 +36,8 @@ steps:
     - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-hopper.txt --tp-size=4
 
 - label: LM Eval Small Models (B200)
-  key: lm-eval-small-models-b200
   timeout_in_minutes: 120
-  device: b200-k8s
+  device: b200
   optional: true
   source_file_dependencies:
   - csrc/
@@ -48,22 +45,9 @@ steps:
   commands:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-blackwell.txt
 
-- label: LM Eval Large Models (B200, EP)
-  key: lm-eval-large-models-b200-ep
-  timeout_in_minutes: 120
-  device: b200-k8s
-  optional: true
-  num_devices: 2
-  source_file_dependencies:
-  - csrc/
-  - vllm/model_executor/layers/quantization
-  commands:
-  - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-blackwell-ep.txt
-
 - label: LM Eval Qwen3.5 Models (B200)
-  key: lm-eval-qwen3-5-models-b200
   timeout_in_minutes: 120
-  device: b200-k8s
+  device: b200
   optional: true
   num_devices: 2
   source_file_dependencies:
@@ -78,7 +62,6 @@ steps:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-qwen35-blackwell.txt
 
 - label: LM Eval Large Models (H200)
-  key: lm-eval-large-models-h200
   timeout_in_minutes: 60
   device: h200
   optional: true
@@ -87,7 +70,6 @@ steps:
     - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-h200.txt
 
 - label: MoE Refactor Integration Test (H100 - TEMPORARY)
-  key: moe-refactor-integration-test-h100-temporary
   device: h100
   optional: true
   num_devices: 2
@@ -95,16 +77,14 @@ steps:
     - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/moe-refactor/config-h100.txt
   
 - label: MoE Refactor Integration Test (B200 - TEMPORARY)
-  key: moe-refactor-integration-test-b200-temporary
-  device: b200-k8s
+  device: b200
   optional: true
   num_devices: 2
   commands:
     - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/moe-refactor/config-b200.txt
 
 - label: MoE Refactor Integration Test (B200 DP - TEMPORARY)
-  key: moe-refactor-integration-test-b200-dp-temporary
-  device: b200-k8s
+  device: b200
   optional: true
   num_devices: 2
   commands:
@@ -112,9 +92,7 @@ steps:
 
 
 - label: LM Eval TurboQuant KV Cache
-  key: lm-eval-turboquant-kv-cache
   timeout_in_minutes: 75
-  device: h200_18gb
   source_file_dependencies:
   - vllm/model_executor/layers/quantization/turboquant/
   - vllm/v1/attention/backends/turboquant_attn.py
@@ -124,7 +102,6 @@ steps:
   - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/models-turboquant.txt
 
 - label: GPQA Eval (GPT-OSS) (H100)
-  key: gpqa-eval-gpt-oss-h100
   timeout_in_minutes: 120
   device: h100
   optional: true
@@ -138,9 +115,8 @@ steps:
     - pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-h100.txt
 
 - label: GPQA Eval (GPT-OSS) (B200)
-  key: gpqa-eval-gpt-oss-b200
   timeout_in_minutes: 120
-  device: b200-k8s
+  device: b200
   optional: true
   num_devices: 2
   source_file_dependencies:
@@ -150,10 +126,3 @@ steps:
   commands:
     - uv pip install --system 'gpt-oss[eval]==0.0.5'
     - pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-b200.txt
-
-- label: MRCR Eval Small Models
-  timeout_in_minutes: 30
-  source_file_dependencies:
-  - tests/evals/mrcr/
-  commands:
-    - pytest -s -v evals/mrcr/test_mrcr_correctness.py --config-list-file=evals/mrcr/configs/models-small.txt

.buildkite/test_areas/lora.yaml

@@ -3,23 +3,20 @@ depends_on:
   - image-build
 steps:
 - label: LoRA %N
-  key: lora
   timeout_in_minutes: 30
   source_file_dependencies:
   - vllm/lora
   - tests/lora
   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 
+    - 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_llm_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 
   parallelism: 4
 
 
 - label: LoRA TP (Distributed)
-  key: lora-tp-distributed
   timeout_in_minutes: 30
   num_devices: 4
   source_file_dependencies:
   - vllm/lora
-  - vllm/model_executor/layers/fused_moe/
   - tests/lora
   commands:
     # FIXIT: find out which code initialize cuda before running the test
@@ -31,7 +28,7 @@ steps:
     # requires multi-GPU testing for validation.
     - pytest -v -s -x lora/test_chatglm3_tp.py
     - pytest -v -s -x lora/test_llama_tp.py
-    - pytest -v -s -x lora/test_qwen3_with_multi_loras.py
+    - pytest -v -s -x lora/test_llm_with_multi_loras.py
     - pytest -v -s -x lora/test_olmoe_tp.py
     - pytest -v -s -x lora/test_gptoss_tp.py
     - pytest -v -s -x lora/test_qwen35_densemodel_lora.py

.buildkite/test_areas/misc.yaml

@@ -3,40 +3,25 @@ depends_on:
   - image-build
 steps:
 - label: V1 Spec Decode
-  key: v1-spec-decode
   timeout_in_minutes: 30
   source_file_dependencies:
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/inputs/
-    - vllm/model_executor/
-    - vllm/platforms/
-    - vllm/sampling_params.py
-    - vllm/transformers_utils/
-    - vllm/utils/
-    - vllm/v1/
+    - vllm/
     - tests/v1/spec_decode
   commands:
     - export VLLM_WORKER_MULTIPROC_METHOD=spawn
     # TODO: create another `optional` test group for slow tests
     - pytest -v -s -m 'not slow_test' v1/spec_decode
+  mirror:
+    amd:
+      device: mi325_1
+      depends_on:
+      - image-build-amd
 
 - label: V1 Sample + Logits
-  key: v1-sample-logits
   timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
-    - 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/
+    - vllm/
     - tests/v1/sample
     - tests/v1/logits_processors
     - tests/v1/test_oracle.py
@@ -51,35 +36,17 @@ steps:
     - pytest -v -s v1/test_outputs.py
   mirror:
     amd:
-      device: mi300_1
+      device: mi325_1
       depends_on:
       - image-build-amd
 
 - label: V1 Core + KV + Metrics
-  key: v1-core-kv-metrics
   timeout_in_minutes: 30
   source_file_dependencies:
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/engine/
-    - vllm/entrypoints/pooling/
-    - vllm/inputs/
-    - vllm/lora/
-    - vllm/model_executor/
-    - vllm/multimodal/
-    - vllm/outputs.py
-    - vllm/platforms/
-    - vllm/pooling_params.py
-    - vllm/profiler/
-    - vllm/sampling_params.py
-    - vllm/tokenizers/
-    - vllm/transformers_utils/
-    - vllm/utils/
-    - vllm/v1/
+    - vllm/
     - tests/v1/core
     - tests/v1/executor
     - tests/v1/kv_offload
-    - tests/v1/simple_kv_offload
     - tests/v1/worker
     - tests/v1/kv_connector/unit
     - tests/v1/metrics
@@ -91,34 +58,23 @@ steps:
     - pytest -v -s -m 'not cpu_test' v1/core
     - pytest -v -s v1/executor
     - pytest -v -s v1/kv_offload
-    - pytest -v -s v1/simple_kv_offload
     - pytest -v -s v1/worker
     - pytest -v -s -m 'not cpu_test' v1/kv_connector/unit
     - pytest -v -s -m 'not cpu_test' v1/metrics
     # Integration test for streaming correctness (requires special branch).
     - pip install -U git+https://github.com/robertgshaw2-redhat/lm-evaluation-harness.git@streaming-api
     - pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
+  mirror:
+    amd:
+      device: mi325_1
+      depends_on:
+      - image-build-amd
 
 - label: V1 Others (CPU)
-  key: v1-others-cpu
   depends_on:
     - image-build-cpu
   source_file_dependencies:
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/engine/
-    - vllm/inputs/
-    - vllm/lora/
-    - vllm/multimodal/
-    - vllm/outputs.py
-    - vllm/platforms/
-    - vllm/pooling_params.py
-    - vllm/profiler/
-    - vllm/sampling_params.py
-    - vllm/tokenizers/
-    - vllm/transformers_utils/
-    - vllm/utils/
-    - vllm/v1/
+    - vllm/
     - tests/v1
   device: cpu-small
   commands:
@@ -129,35 +85,11 @@ steps:
     - pytest -v -s -m 'cpu_test' v1/kv_connector/unit
     - pytest -v -s -m 'cpu_test' v1/metrics
 
-- label: Extract Hidden States Integration
-  key: extract-hidden-states-integration
-  timeout_in_minutes: 20
-  device: h200_18gb
-  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
-    - 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: Regression
-  key: regression
   timeout_in_minutes: 20
   device: h200_18gb
   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/
+  - vllm/
   - tests/test_regression
   commands:
   - pip install modelscope
@@ -165,7 +97,6 @@ steps:
   working_dir: "/vllm-workspace/tests" # optional
 
 - label: Examples
-  key: examples
   timeout_in_minutes: 45
   working_dir: "/vllm-workspace/examples"
   source_file_dependencies:
@@ -182,37 +113,25 @@ steps:
     - python3 basic/offline_inference/embed.py
     - python3 basic/offline_inference/score.py
     # for multi-modal models
-    - python3 generate/multimodal/audio_language_offline.py --seed 0
-    - python3 generate/multimodal/vision_language_offline.py --seed 0
-    - python3 generate/multimodal/vision_language_multi_image_offline.py --seed 0
-    - python3 generate/multimodal/encoder_decoder_multimodal_offline.py --model-type whisper --seed 0
+    - python3 offline_inference/audio_language.py --seed 0
+    - python3 offline_inference/vision_language.py --seed 0
+    - python3 offline_inference/vision_language_multi_image.py --seed 0
+    - python3 offline_inference/encoder_decoder_multimodal.py --model-type whisper --seed 0
      # for pooling models
     - python3 pooling/embed/vision_embedding_offline.py --seed 0
     # for features demo
-    - python3 features/automatic_prefix_caching/prefix_caching_offline.py
-    - python3 deployment/llm_engine_example.py
-    - python3 features/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 features/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
-    - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
+    - python3 offline_inference/prefix_caching.py
+    - python3 offline_inference/llm_engine_example.py
+    - python3 others/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 others/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
+    - python3 offline_inference/spec_decode.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
     # https://github.com/vllm-project/vllm/pull/26682 uses slightly more memory in PyTorch 2.9+ causing this test to OOM in 1xL4 GPU
-    - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
+    - python3 offline_inference/spec_decode.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
 
 - label: Metrics, Tracing (2 GPUs)
-  key: metrics-tracing-2-gpus
   timeout_in_minutes: 20
   num_devices: 2
   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/
+  - vllm/
   - tests/v1/tracing
   commands:
   - "pip install \
@@ -223,11 +142,8 @@ steps:
   - pytest -v -s v1/tracing
 
 - label: Python-only Installation
-  key: python-only-installation
   depends_on: ~
-  optional: true
   timeout_in_minutes: 20
-  device: h200_18gb
   source_file_dependencies:
   - tests/standalone_tests/python_only_compile.sh
   - setup.py
@@ -235,22 +151,9 @@ steps:
   - bash standalone_tests/python_only_compile.sh
 
 - label: Async Engine, Inputs, Utils, Worker
-  key: async-engine-inputs-utils-worker
   timeout_in_minutes: 50
   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/
+  - vllm/
   - tests/detokenizer
   - tests/multimodal
   - tests/utils_
@@ -260,35 +163,11 @@ steps:
   - pytest -v -s utils_
 
 - label: Async Engine, Inputs, Utils, Worker, Config (CPU)
-  key: async-engine-inputs-utils-worker-config-cpu
-  depends_on:
+  depends_on: 
   - image-build-cpu
   timeout_in_minutes: 30
   source_file_dependencies:
-  - vllm/assets/
-  - vllm/config/
-  - vllm/engine/arg_utils.py
-  - vllm/entrypoints/chat_utils.py
-  - vllm/entrypoints/mcp/
-  - vllm/entrypoints/openai/chat_completion/protocol.py
-  - vllm/entrypoints/openai/engine/protocol.py
-  - vllm/envs.py
-  - vllm/exceptions.py
-  - vllm/inputs/
-  - vllm/model_executor/layers/quantization/quark/
-  - vllm/multimodal/
-  - vllm/outputs.py
-  - vllm/platforms/
-  - vllm/pooling_params.py
-  - vllm/ray/
-  - vllm/reasoning/
-  - vllm/renderers/
-  - vllm/sampling_params.py
-  - vllm/tokenizers/
-  - vllm/tool_parsers/
-  - vllm/transformers_utils/
-  - vllm/utils/
-  - vllm/v1/
+  - vllm/
   - tests/test_inputs.py
   - tests/test_outputs.py
   - tests/test_pooling_params.py
@@ -299,7 +178,6 @@ steps:
   - tests/tokenizers_
   - tests/reasoning
   - tests/tool_parsers
-  - tests/parser
   - tests/transformers_utils
   - tests/config
   device: cpu-small
@@ -312,28 +190,12 @@ steps:
   - 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 reasoning --ignore=reasoning/test_seedoss_reasoning_parser.py --ignore=reasoning/test_glm4_moe_reasoning_parser.py --ignore=reasoning/test_gemma4_reasoning_parser.py
   - pytest -v -s tool_parsers
-  - pytest -v -s parser
   - pytest -v -s transformers_utils
   - pytest -v -s config
 
-- label: Batch Invariance (A100)
-  key: batch-invariance-a100
-  timeout_in_minutes: 30
-  device: a100
-  source_file_dependencies:
-    - vllm/v1/attention
-    - vllm/model_executor/layers
-    - tests/v1/determinism/
-  commands:
-    - export VLLM_WORKER_MULTIPROC_METHOD=spawn
-    - pip install pytest-timeout pytest-forked
-    - pytest -v -s v1/determinism/test_batch_invariance.py
-    - VLLM_TEST_MODEL=deepseek-ai/DeepSeek-V2-Lite-Chat pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[TRITON_MLA]
-
 - label: Batch Invariance (H100)
-  key: batch-invariance-h100
   timeout_in_minutes: 30
   device: h100
   source_file_dependencies:
@@ -349,9 +211,8 @@ steps:
     - VLLM_TEST_MODEL=Qwen/Qwen3-30B-A3B-Thinking-2507-FP8 pytest -v -s v1/determinism/test_batch_invariance.py::test_v1_generation_is_deterministic_across_batch_sizes_with_needle[FLASH_ATTN]
 
 - label: Batch Invariance (B200)
-  key: batch-invariance-b200
   timeout_in_minutes: 30
-  device: b200-k8s
+  device: b200
   source_file_dependencies:
     - vllm/v1/attention
     - vllm/model_executor/layers
@@ -366,7 +227,6 @@ steps:
     - pytest -v -s v1/determinism/test_nvfp4_batch_invariant.py
   
 - label: Acceptance Length Test (Large Models) # optional
-  key: acceptance-length-test-large-models
   timeout_in_minutes: 25
   gpu: h100
   optional: true

.buildkite/test_areas/model_executor.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: Model Executor
-  key: model-executor
   timeout_in_minutes: 35
   source_file_dependencies:
   - vllm/engine/arg_utils.py

.buildkite/test_areas/model_runner_v2.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: Model Runner V2 Core Tests
-  key: model-runner-v2-core-tests
   timeout_in_minutes: 45
   source_file_dependencies:
   - vllm/v1/worker/gpu/
@@ -26,18 +25,16 @@ steps:
   - pytest -v -s entrypoints/llm/test_struct_output_generate.py -k "xgrammar and not speculative_config6 and not speculative_config7 and not speculative_config8 and not speculative_config0"
 
 - label: Model Runner V2 Examples
-  key: model-runner-v2-examples
   timeout_in_minutes: 45
   working_dir: "/vllm-workspace/examples"
   source_file_dependencies:
     - vllm/v1/worker/gpu/
     - vllm/v1/core/sched/
     - vllm/v1/worker/gpu_worker.py
+    - examples/offline_inference/
     - examples/basic/offline_inference/
-    - examples/generate/multimodal/
-    - examples/features/
     - examples/pooling/embed/vision_embedding_offline.py
-    - examples/features/tensorize_vllm_model.py
+    - examples/others/tensorize_vllm_model.py
   commands:
     - set -x
     - export VLLM_USE_V2_MODEL_RUNNER=1
@@ -47,22 +44,21 @@ steps:
     #- python3 basic/offline_inference/generate.py --model meta-llama/Llama-2-13b-chat-hf --cpu-offload-gb 10  # TODO
     #- python3 basic/offline_inference/embed.py   # TODO
     # for multi-modal models
-    - python3 generate/multimodal/audio_language_offline.py --seed 0
-    - python3 generate/multimodal/vision_language_offline.py --seed 0
-    - python3 generate/multimodal/vision_language_multi_image_offline.py --seed 0
-    - python3 generate/multimodal/encoder_decoder_multimodal_offline.py --model-type whisper --seed 0
+    - python3 offline_inference/audio_language.py --seed 0
+    - python3 offline_inference/vision_language.py --seed 0
+    - python3 offline_inference/vision_language_multi_image.py --seed 0
+    - python3 offline_inference/encoder_decoder_multimodal.py --model-type whisper --seed 0
     # for pooling models
     - python3 pooling/embed/vision_embedding_offline.py --seed 0
     # for features demo
-    - python3 features/automatic_prefix_caching/prefix_caching_offline.py
-    - python3 deployment/llm_engine_example.py
-    - python3 features/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 features/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
-    - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
+    - python3 offline_inference/prefix_caching.py
+    - python3 offline_inference/llm_engine_example.py
+    - python3 others/tensorize_vllm_model.py --model facebook/opt-125m serialize --serialized-directory /tmp/ --suffix v1 && python3 others/tensorize_vllm_model.py --model facebook/opt-125m deserialize --path-to-tensors /tmp/vllm/facebook/opt-125m/v1/model.tensors
+    - python3 offline_inference/spec_decode.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
     # https://github.com/vllm-project/vllm/pull/26682 uses slightly more memory in PyTorch 2.9+ causing this test to OOM in 1xL4 GPU
-    - python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
+    - python3 offline_inference/spec_decode.py --test --method eagle3 --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 1536
 
 - label: Model Runner V2 Distributed (2 GPUs)
-  key: model-runner-v2-distributed-2-gpus
   timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   num_devices: 2
@@ -83,7 +79,6 @@ steps:
     - TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
 
 - label: Model Runner V2 Pipeline Parallelism (4 GPUs)
-  key: model-runner-v2-pipeline-parallelism-4-gpus
   timeout_in_minutes: 60
   working_dir: "/vllm-workspace/tests"
   num_devices: 4
@@ -99,19 +94,19 @@ steps:
     - pytest -v -s distributed/test_pp_cudagraph.py -k "not ray"
 
 - label: Model Runner V2 Spec Decode
-  key: model-runner-v2-spec-decode
   timeout_in_minutes: 30
   working_dir: "/vllm-workspace/tests"
   source_file_dependencies:
   - vllm/v1/worker/gpu/
   - vllm/v1/worker/gpu_worker.py
   - tests/v1/spec_decode/test_max_len.py
-  - tests/v1/spec_decode/test_rejection_sampler_utils.py
+  - tests/v1/spec_decode/test_probabilistic_rejection_sampler_utils.py
+  - tests/v1/spec_decode/test_synthetic_rejection_sampler_utils.py
   - tests/v1/e2e/spec_decode/test_spec_decode.py
   commands:
   - set -x
   - export VLLM_USE_V2_MODEL_RUNNER=1
   - pytest -v -s v1/spec_decode/test_max_len.py -k "eagle or mtp"
-  - pytest -v -s v1/spec_decode/test_rejection_sampler_utils.py
+  - pytest -v -s v1/spec_decode/test_probabilistic_rejection_sampler_utils.py
   - pytest -v -s v1/spec_decode/test_synthetic_rejection_sampler_utils.py
   - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "eagle or mtp"

.buildkite/test_areas/models_basic.yaml

@@ -3,9 +3,7 @@ depends_on:
   - image-build
 steps:
 - label: Basic Models Tests (Initialization)
-  key: basic-models-tests-initialization
   timeout_in_minutes: 45
-  device: h200_18gb
   torch_nightly: true
   source_file_dependencies:
   - vllm/
@@ -18,7 +16,6 @@ steps:
     torch_nightly: {}
 
 - label: Basic Models Tests (Extra Initialization) %N
-  key: basic-models-tests-extra-initialization
   timeout_in_minutes: 45
   source_file_dependencies:
   - vllm/model_executor/models/
@@ -34,7 +31,6 @@ steps:
     torch_nightly: {}
 
 - label: Basic Models Tests (Other)
-  key: basic-models-tests-other
   timeout_in_minutes: 45
   source_file_dependencies:
   - vllm/
@@ -43,9 +39,14 @@ steps:
   - tests/models/test_registry.py
   commands:
     - pytest -v -s models/test_terratorch.py models/test_transformers.py models/test_registry.py
+  mirror:
+    amd:
+      device: mi325_1
+      depends_on:
+      - image-build-amd
+
 
 - label: Basic Models Test (Other CPU) # 5min
-  key: basic-models-test-other-cpu
   depends_on:
   - image-build-cpu
   timeout_in_minutes: 10
@@ -58,7 +59,6 @@ steps:
     - pytest -v -s models/test_utils.py models/test_vision.py
 
 - label: Transformers Nightly Models
-  key: transformers-nightly-models
   working_dir: "/vllm-workspace/"
   optional: true
   soft_fail: true
@@ -69,12 +69,11 @@ steps:
     - pytest -v -s tests/models/multimodal/processing/
     - pytest -v -s tests/models/multimodal/test_mapping.py
     - python3 examples/basic/offline_inference/chat.py
-    - python3 examples/generate/multimodal/vision_language_offline.py --model-type qwen2_5_vl
+    - python3 examples/offline_inference/vision_language.py --model-type qwen2_5_vl
     # Whisper needs spawn method to avoid deadlock
-    - VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/generate/multimodal/audio_language_offline.py --model-type whisper
+    - VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/offline_inference/audio_language.py --model-type whisper
 
 - label: Transformers Backward Compatibility Models Test
-  key: transformers-backward-compatibility-models-test
   working_dir: "/vllm-workspace/"
   optional: true
   soft_fail: true
@@ -84,7 +83,7 @@ steps:
     - pytest -v -s tests/models/test_transformers.py
     - pytest -v -s tests/models/multimodal/processing/
     - pytest -v -s tests/models/multimodal/test_mapping.py
-    - python3 examples/basic/offline_inference/chat.py
-    - python3 examples/generate/multimodal/vision_language_offline.py --model-type qwen2_5_vl
+    - python3 examples/offline_inference/basic/chat.py
+    - python3 examples/offline_inference/vision_language.py --model-type qwen2_5_vl
     # Whisper needs spawn method to avoid deadlock
-    - VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/generate/multimodal/audio_language_offline.py --model-type whisper
+    - VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/offline_inference/audio_language.py --model-type whisper

.buildkite/test_areas/models_distributed.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: Distributed Model Tests (2 GPUs)
-  key: distributed-model-tests-2-gpus
   timeout_in_minutes: 50
   working_dir: "/vllm-workspace/tests"
   num_devices: 2

.buildkite/test_areas/models_language.yaml

@@ -3,9 +3,7 @@ depends_on:
   - image-build
 steps:
 - label: Language Models Tests (Standard)
-  key: language-models-tests-standard
   timeout_in_minutes: 25
-  device: h200_18gb
   source_file_dependencies:
   - vllm/
   - tests/models/language
@@ -17,7 +15,6 @@ steps:
     torch_nightly: {}
 
 - label: Language Models Tests (Extra Standard) %N
-  key: language-models-tests-extra-standard
   timeout_in_minutes: 45
   source_file_dependencies:
   - vllm/model_executor/models/
@@ -34,7 +31,6 @@ steps:
     torch_nightly: {}
 
 - label: Language Models Tests (Hybrid) %N
-  key: language-models-tests-hybrid
   timeout_in_minutes: 75
   source_file_dependencies:
   - vllm/
@@ -49,17 +45,8 @@ steps:
   parallelism: 2
   mirror:
     torch_nightly: {}
-    amd:
-      device: mi300_1
-      depends_on:
-      - image-build-amd
-      commands:
-      - uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@fix-rocm-7.0-warp-size-constexpr'
-      - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
-      - pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
 
 - label: Language Models Test (Extended Generation) # 80min
-  key: language-models-test-extended-generation
   timeout_in_minutes: 110
   optional: true
   source_file_dependencies:
@@ -71,9 +58,17 @@ steps:
     - uv pip install --system --no-build-isolation 'git+https://github.com/state-spaces/mamba@v2.3.0'
     - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
     - pytest -v -s models/language/generation -m '(not core_model) and (not hybrid_model)'
+  mirror:
+    amd:
+      device: mi325_1
+      depends_on:
+      - image-build-amd
+      commands:
+      - uv pip install --system --no-build-isolation 'git+https://github.com/AndreasKaratzas/mamba@fix-rocm-7.0-warp-size-constexpr'
+      - uv pip install --system --no-build-isolation 'git+https://github.com/Dao-AILab/causal-conv1d@v1.6.0'
+      - pytest -v -s models/language/generation -m '(not core_model) and (not hybrid_model)'
 
 - label: Language Models Test (PPL)
-  key: language-models-test-ppl
   timeout_in_minutes: 110
   device: h200_18gb
   optional: true
@@ -84,7 +79,6 @@ steps:
     - pytest -v -s models/language/generation_ppl_test
 
 - label: Language Models Test (Extended Pooling)  # 36min
-  key: language-models-test-extended-pooling
   timeout_in_minutes: 50
   optional: true
   source_file_dependencies:
@@ -94,13 +88,11 @@ steps:
     - pytest -v -s models/language/pooling -m 'not core_model'
   mirror:
     amd:
-      device: mi300_1
-      timeout_in_minutes: 100
+      device: mi325_1
       depends_on:
       - image-build-amd
 
 - label: Language Models Test (MTEB)
-  key: language-models-test-mteb
   timeout_in_minutes: 110
   device: h200_18gb
   optional: true

.buildkite/test_areas/models_multimodal.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: "Multi-Modal Models (Standard) 1: qwen2"
-  key: multi-modal-models-standard-1-qwen2
   timeout_in_minutes: 45
   device: h200_18gb
   source_file_dependencies:
@@ -15,12 +14,11 @@ steps:
     - pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model
   mirror:
     amd:
-      device: mi300_1
+      device: mi325_1
       depends_on:
       - image-build-amd
 
 - label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
-  key: multi-modal-models-standard-2-qwen3-gemma
   timeout_in_minutes: 45
   device: h200_18gb
   source_file_dependencies:
@@ -33,12 +31,11 @@ steps:
     - pytest -v -s models/multimodal/generation/test_vit_cudagraph.py -m core_model
   mirror:
     amd:
-      device: mi300_1
+      device: mi325_1
       depends_on:
       - image-build-amd
 
 - label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
-  key: multi-modal-models-standard-3-llava-qwen2-vl
   timeout_in_minutes: 45
   source_file_dependencies:
   - vllm/
@@ -49,12 +46,11 @@ steps:
     - pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model
   mirror:
     amd:
-      device: mi300_1
+      device: mi325_1
       depends_on:
       - image-build-amd
 
 - label: "Multi-Modal Models (Standard) 4: other + whisper"
-  key: multi-modal-models-standard-4-other-whisper
   timeout_in_minutes: 45
   source_file_dependencies:
   - vllm/
@@ -64,10 +60,14 @@ steps:
     - 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
     - pytest models/multimodal/generation/test_memory_leak.py -m core_model
     - cd .. && VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s tests/models/multimodal/generation/test_whisper.py -m core_model  # Otherwise, mp_method="spawn" doesn't work
+  mirror:
+    amd:
+      device: mi325_1
+      depends_on:
+      - image-build-amd
 
 - label: Multi-Modal Processor (CPU)
-  key: multi-modal-processor-cpu
-  depends_on:
+  depends_on: 
   - image-build-cpu
   timeout_in_minutes: 60
   source_file_dependencies:
@@ -80,7 +80,6 @@ steps:
     - pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py
 
 - label: Multi-Modal Processor # 44min
-  key: multi-modal-processor
   timeout_in_minutes: 60
   device: h200_18gb
   source_file_dependencies:
@@ -92,7 +91,6 @@ steps:
     - pytest -v -s models/multimodal/processing/test_tensor_schema.py
 
 - label: Multi-Modal Accuracy Eval (Small Models) # 50min
-  key: multi-modal-accuracy-eval-small-models
   timeout_in_minutes: 70
   working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
   source_file_dependencies:
@@ -103,7 +101,6 @@ steps:
   - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-mm-small.txt --tp-size=1
 
 - label: Multi-Modal Models (Extended Generation 1)
-  key: multi-modal-models-extended-generation-1
   optional: true
   source_file_dependencies:
   - vllm/
@@ -115,12 +112,11 @@ steps:
     - pytest -v -s models/multimodal/test_mapping.py
   mirror:
     amd:
-      device: mi300_1
+      device: mi325_1
       depends_on:
       - image-build-amd
 
 - label: Multi-Modal Models (Extended Generation 2)
-  key: multi-modal-models-extended-generation-2
   optional: true
   source_file_dependencies:
   - vllm/
@@ -130,7 +126,6 @@ steps:
     - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=0) and not core_model'
 
 - label: Multi-Modal Models (Extended Generation 3)
-  key: multi-modal-models-extended-generation-3
   optional: true
   source_file_dependencies:
   - vllm/
@@ -140,7 +135,6 @@ steps:
     - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=1) and not core_model'
 
 - label: Multi-Modal Models (Extended Pooling)
-  key: multi-modal-models-extended-pooling
   optional: true
   device: h200_18gb
   source_file_dependencies:

.buildkite/test_areas/plugins.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: Plugin Tests (2 GPUs)
-  key: plugin-tests-2-gpus
   timeout_in_minutes: 60
   working_dir: "/vllm-workspace/tests"
   num_devices: 2

.buildkite/test_areas/pytorch.yaml

@@ -3,33 +3,9 @@ depends_on:
   - image-build
 steps:
 - label: PyTorch Compilation Unit Tests
-  key: pytorch-compilation-unit-tests
   timeout_in_minutes: 10
   source_file_dependencies:
-    - vllm/__init__.py
-    - vllm/_aiter_ops.py
-    - vllm/_custom_ops.py
-    - vllm/compilation/
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/engine/
-    - vllm/env_override.py
-    - vllm/envs.py
-    - vllm/forward_context.py
-    - vllm/inputs/
-    - vllm/ir/
-    - vllm/kernels/
-    - vllm/logger.py
-    - vllm/model_executor/
-    - vllm/multimodal/
-    - vllm/platforms/
-    - vllm/plugins/
-    - vllm/sampling_params.py
-    - vllm/sequence.py
-    - vllm/transformers_utils/
-    - vllm/triton_utils/
-    - vllm/utils/
-    - vllm/v1/
+    - vllm/
     - tests/compile
   commands:
   # Run unit tests defined directly under compile/,
@@ -42,99 +18,27 @@ steps:
   - "find compile/ -maxdepth 1 -name 'test_*.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
 
 - label: PyTorch Compilation Unit Tests (H100)
-  key: pytorch-compilation-unit-tests-h100
   timeout_in_minutes: 30
   device: h100
   num_devices: 1
   source_file_dependencies:
-    - vllm/__init__.py
-    - vllm/_aiter_ops.py
-    - vllm/_custom_ops.py
-    - vllm/compilation/
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/engine/
-    - vllm/env_override.py
-    - vllm/envs.py
-    - vllm/forward_context.py
-    - vllm/inputs/
-    - vllm/ir/
-    - vllm/kernels/
-    - vllm/logger.py
-    - vllm/model_executor/
-    - vllm/multimodal/
-    - vllm/platforms/
-    - vllm/plugins/
-    - vllm/sampling_params.py
-    - vllm/sequence.py
-    - vllm/transformers_utils/
-    - vllm/triton_utils/
-    - vllm/utils/
-    - vllm/v1/
+    - vllm/
     - tests/compile/h100/
   commands:
   - "find compile/h100/ -name 'test_*.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
 
 - label: PyTorch Compilation Passes Unit Tests
-  key: pytorch-compilation-passes-unit-tests
   timeout_in_minutes: 20
   source_file_dependencies:
-    - vllm/__init__.py
-    - vllm/_aiter_ops.py
-    - vllm/_custom_ops.py
-    - vllm/compilation/
-    - vllm/config/
-    - vllm/distributed/
-    - vllm/engine/
-    - vllm/env_override.py
-    - vllm/envs.py
-    - vllm/forward_context.py
-    - vllm/inputs/
-    - vllm/ir/
-    - vllm/kernels/
-    - vllm/logger.py
-    - vllm/model_executor/
-    - vllm/multimodal/
-    - vllm/platforms/
-    - vllm/plugins/
-    - vllm/sampling_params.py
-    - vllm/sequence.py
-    - vllm/transformers_utils/
-    - vllm/triton_utils/
-    - vllm/utils/
-    - vllm/v1/
+    - vllm/
     - tests/compile/passes
   commands:
   - pytest -s -v compile/passes --ignore compile/passes/distributed
 
 - label: PyTorch Fullgraph Smoke Test
-  key: pytorch-fullgraph-smoke-test
   timeout_in_minutes: 35
   source_file_dependencies:
-  - vllm/__init__.py
-  - vllm/_aiter_ops.py
-  - vllm/_custom_ops.py
-  - vllm/compilation/
-  - vllm/config/
-  - vllm/distributed/
-  - vllm/engine/
-  - vllm/env_override.py
-  - vllm/envs.py
-  - vllm/forward_context.py
-  - vllm/inputs/
-  - vllm/ir/
-  - vllm/kernels/
-  - vllm/logger.py
-  - vllm/model_executor/
-  - vllm/multimodal/
-  - vllm/platforms/
-  - vllm/plugins/
-  - vllm/sampling_params.py
-  - vllm/sequence.py
-  - vllm/transformers_utils/
-  - vllm/triton_utils/
-  - vllm/utils/
-  - vllm/v1/
+  - vllm/
   - tests/compile
   commands:
   # Run smoke tests under fullgraph directory, except test_full_graph.py
@@ -144,41 +48,16 @@ steps:
   - "find compile/fullgraph/ -name 'test_*.py' -not -name 'test_full_graph.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
 
 - label: PyTorch Fullgraph
-  key: pytorch-fullgraph
   timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
-  - vllm/__init__.py
-  - vllm/_aiter_ops.py
-  - vllm/_custom_ops.py
-  - vllm/compilation/
-  - vllm/config/
-  - vllm/distributed/
-  - vllm/engine/
-  - vllm/env_override.py
-  - vllm/envs.py
-  - vllm/forward_context.py
-  - vllm/inputs/
-  - vllm/ir/
-  - vllm/kernels/
-  - vllm/logger.py
-  - vllm/model_executor/
-  - vllm/multimodal/
-  - vllm/platforms/
-  - vllm/plugins/
-  - vllm/sampling_params.py
-  - vllm/sequence.py
-  - vllm/transformers_utils/
-  - vllm/triton_utils/
-  - vllm/utils/
-  - vllm/v1/
+  - vllm/
   - tests/compile
   commands:
     # fp8 kv scales not supported on sm89, tested on Blackwell instead
   - pytest -v -s compile/fullgraph/test_full_graph.py -k 'not test_fp8_kv_scale_compile'
 
 - label: Pytorch Nightly Dependency Override Check # 2min
-  key: pytorch-nightly-dependency-override-check
   # if this test fails, it means the nightly torch version is not compatible with some
   # of the dependencies. Please check the error message and add the package to whitelist
   # in /vllm/tools/pre_commit/generate_nightly_torch_test.py

.buildkite/test_areas/quantization.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: Quantization
-  key: quantization
   timeout_in_minutes: 90
   source_file_dependencies:
   - csrc/
@@ -22,10 +21,9 @@ steps:
   - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py
 
 - label: Quantized MoE Test (B200)
-  key: quantized-moe-test-b200
   timeout_in_minutes: 60
   working_dir: "/vllm-workspace/"
-  device: b200-k8s
+  device: b200
   source_file_dependencies:
   - tests/quantization/test_blackwell_moe.py
   - vllm/model_executor/models/deepseek_v2.py
@@ -40,7 +38,6 @@ steps:
     - pytest -s -v tests/quantization/test_blackwell_moe.py
 
 - label: Quantized Models Test
-  key: quantized-models-test
   timeout_in_minutes: 60
   source_file_dependencies:
   - vllm/model_executor/layers/quantization

.buildkite/test_areas/ray_compat.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: Ray Dependency Compatibility Check
-  key: ray-dependency-compatibility-check
   # Informational only — does not block the pipeline.
   # If this fails, it means the PR introduces a dependency that
   # conflicts with Ray's dependency constraints.

.buildkite/test_areas/samplers.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: Samplers Test
-  key: samplers-test
   timeout_in_minutes: 75
   source_file_dependencies:
   - vllm/model_executor/layers
@@ -11,13 +10,11 @@ steps:
   - tests/samplers
   - tests/conftest.py
   commands:
-    # VLLM_USE_FLASHINFER_SAMPLER defaults to 1 now, so we need to pin both
-    # values explicitly to still cover the PyTorch-native (Triton) path.
-    - VLLM_USE_FLASHINFER_SAMPLER=0 pytest -v -s samplers
+    - pytest -v -s samplers
     - VLLM_USE_FLASHINFER_SAMPLER=1 pytest -v -s samplers
   mirror:
     amd:
-      device: mi250_1
+      device: mi325_1
       depends_on:
       - image-build-amd
       commands:

.buildkite/test_areas/spec_decode.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: Spec Decode Eagle
-  key: spec-decode-eagle
   timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
@@ -14,9 +13,8 @@ steps:
     - pytest -v -s v1/e2e/spec_decode -k "eagle_correctness"
 
 - label: Spec Decode Eagle Nightly B200
-  key: spec-decode-eagle-nightly-b200
   timeout_in_minutes: 30
-  device: b200-k8s
+  device: b200
   optional: true
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -26,7 +24,6 @@ steps:
     - pytest -v -s v1/e2e/spec_decode -k "eagle_correctness"
 
 - label: Spec Decode Speculators + MTP
-  key: spec-decode-speculators-mtp
   timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
@@ -38,9 +35,8 @@ steps:
     - pytest -v -s v1/e2e/spec_decode -k "speculators or mtp_correctness"
 
 - label: Spec Decode Speculators + MTP Nightly B200
-  key: spec-decode-speculators-mtp-nightly-b200
   timeout_in_minutes: 30
-  device: b200-k8s
+  device: b200
   optional: true
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -51,7 +47,6 @@ steps:
     - pytest -v -s v1/e2e/spec_decode -k "speculators or mtp_correctness"
   
 - label: Spec Decode Ngram + Suffix
-  key: spec-decode-ngram-suffix
   timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
@@ -62,7 +57,6 @@ steps:
     - pytest -v -s v1/e2e/spec_decode -k "ngram or suffix"
 
 - label: Spec Decode Draft Model
-  key: spec-decode-draft-model
   timeout_in_minutes: 30
   device: h200_18gb
   source_file_dependencies:
@@ -73,9 +67,8 @@ steps:
     - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
 
 - label: Spec Decode Draft Model Nightly B200
-  key: spec-decode-draft-model-nightly-b200
   timeout_in_minutes: 30
-  device: b200-k8s
+  device: b200
   optional: true
   source_file_dependencies:
     - vllm/v1/spec_decode/
@@ -84,29 +77,15 @@ steps:
   commands:
     - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
 
-- label: Speculators Correctness
-  key: speculators-correctness
-  timeout_in_minutes: 60
+- label: DFlash Speculators Correctness
+  timeout_in_minutes: 30
   device: h100
   optional: true
   num_devices: 1
   source_file_dependencies:
     - vllm/v1/spec_decode/
     - vllm/model_executor/models/qwen3_dflash.py
-    - tests/v1/spec_decode/test_speculators_correctness.py
+    - tests/v1/spec_decode/test_speculators_dflash.py
   commands:
     - export VLLM_ALLOW_INSECURE_SERIALIZATION=1
-    - pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test
-
-- label: Spec Decode MTP hybrid (B200)
-  timeout_in_minutes: 30
-  device: b200-k8s
-  optional: true
-  source_file_dependencies:
-    - vllm/v1/spec_decode/
-    - vllm/v1/worker/gpu/spec_decode/
-    - vllm/model_executor/models/qwen3_5.py
-    - vllm/model_executor/models/qwen3_5_mtp.py
-    - tests/v1/e2e/spec_decode/
-  commands:
-    - pytest -v -s v1/e2e/spec_decode -k "qwen3_5-hybrid"
+    - pytest -v -s v1/spec_decode/test_speculators_dflash.py -m slow_test

.buildkite/test_areas/weight_loading.yaml

@@ -3,7 +3,6 @@ depends_on:
   - image-build
 steps:
 - label: Weight Loading Multiple GPU  # 33min
-  key: weight-loading-multiple-gpu
   timeout_in_minutes: 45
   working_dir: "/vllm-workspace/tests"
   num_devices: 2

.github/CODEOWNERS

@@ -6,8 +6,8 @@
 /vllm/distributed/kv_transfer @NickLucche @ApostaC @orozery @xuechendi
 /vllm/lora @jeejeelee
 /vllm/model_executor/layers/attention @LucasWilkinson @MatthewBonanni
-/vllm/model_executor/layers/fused_moe @mgoin @pavanimajety @zyongye
-/vllm/model_executor/layers/quantization @mgoin @robertgshaw2-redhat @tlrmchlsmth @yewentao256 @pavanimajety @zyongye
+/vllm/model_executor/layers/fused_moe @mgoin @pavanimajety
+/vllm/model_executor/layers/quantization @mgoin @robertgshaw2-redhat @tlrmchlsmth @yewentao256 @pavanimajety
 /vllm/model_executor/layers/mamba @tdoublep @tomeras91
 /vllm/model_executor/layers/mamba/gdn_linear_attn.py @tdoublep @ZJY0516 @vadiklyutiy
 /vllm/model_executor/layers/rotary_embedding.py @vadiklyutiy
@@ -18,8 +18,7 @@
 /vllm/kernels/helion @ProExpertProg @zou3519
 /vllm/multimodal @DarkLight1337 @ywang96 @NickLucche @tjtanaa
 /vllm/vllm_flash_attn @LucasWilkinson @MatthewBonanni
-/CMakeLists.txt @tlrmchlsmth @LucasWilkinson @Harry-Chen
-/cmake @tlrmchlsmth @LucasWilkinson @Harry-Chen
+CMakeLists.txt @tlrmchlsmth @LucasWilkinson
 
 # Any change to the VllmConfig changes can have a large user-facing impact,
 # so spam a lot of people
@@ -31,8 +30,8 @@
 /vllm/entrypoints/cli @hmellor @mgoin @DarkLight1337 @russellb
 /vllm/entrypoints/mcp @heheda12345
 /vllm/entrypoints/openai @aarnphm @chaunceyjiang @DarkLight1337 @russellb
-/vllm/entrypoints/speech_to_text/realtime @njhill
-/vllm/entrypoints/speech_to_text @NickLucche
+/vllm/entrypoints/openai/realtime @njhill
+/vllm/entrypoints/openai/speech_to_text @NickLucche
 /vllm/entrypoints/pooling @noooop
 /vllm/entrypoints/sagemaker @DarkLight1337
 /vllm/entrypoints/serve @njhill
@@ -71,10 +70,6 @@
 /vllm/v1/worker/gpu @WoosukKwon @njhill
 /vllm/v1/worker/gpu/kv_connector.py @orozery
 
-# CI & building
-/.buildkite @Harry-Chen
-/docker/Dockerfile @Harry-Chen
-
 # Test ownership
 /.buildkite/lm-eval-harness @mgoin 
 /tests/distributed/test_multi_node_assignment.py @youkaichao
@@ -82,11 +77,11 @@
 /tests/distributed/test_same_node.py @youkaichao
 /tests/entrypoints @DarkLight1337 @robertgshaw2-redhat @aarnphm @NickLucche
 /tests/evals @mgoin @vadiklyutiy
-/tests/kernels @mgoin @tlrmchlsmth @WoosukKwon @yewentao256 @zyongye
+/tests/kernels @mgoin @tlrmchlsmth @WoosukKwon @yewentao256
 /tests/kernels/ir @ProExpertProg @tjtanaa
 /tests/models @DarkLight1337 @ywang96
 /tests/multimodal @DarkLight1337 @ywang96 @NickLucche
-/tests/quantization @mgoin @robertgshaw2-redhat @yewentao256 @pavanimajety @zyongye
+/tests/quantization @mgoin @robertgshaw2-redhat @yewentao256 @pavanimajety
 /tests/test_inputs.py @DarkLight1337 @ywang96
 /tests/entrypoints/llm/test_struct_output_generate.py @mgoin @russellb @aarnphm
 /tests/v1/structured_output @mgoin @russellb @aarnphm
@@ -152,12 +147,6 @@ mkdocs.yaml @hmellor
 # MTP-specific files
 /vllm/model_executor/models/deepseek_mtp.py @luccafong
 
-# DeepseekV4-specific files
-/vllm/v1/attention/ops/deepseek_v4_ops @zyongye
-/vllm/model_executor/layers/deepseek_compressor.py @zyongye
-/vllm/model_executor/layers/deepseek_v4_attention.py @zyongye
-/vllm/model_executor/layers/sparse_attn_indexer.py @zyongye
-
 # Mistral-specific files
 /vllm/model_executor/models/mistral*.py @patrickvonplaten
 /vllm/model_executor/models/mixtral*.py @patrickvonplaten

.github/mergify.yml

@@ -308,7 +308,8 @@ pull_request_rules:
       - files=benchmarks/benchmark_serving_structured_output.py
       - files=benchmarks/run_structured_output_benchmark.sh
       - files=docs/features/structured_outputs.md
-      - files=^examples/features/structured_outputs/
+      - files=examples/offline_inference/structured_outputs.py
+      - files=examples/online_serving/structured_outputs/structured_outputs.py
       - files~=^tests/v1/structured_output/
       - files=tests/entrypoints/llm/test_struct_output_generate.py
       - files~=^vllm/v1/structured_output/
@@ -324,7 +325,7 @@ pull_request_rules:
     - or:
       - files~=^vllm/v1/spec_decode/
       - files~=^tests/v1/spec_decode/
-      - files=^examples/features/speculative_decoding/
+      - files~=^examples/.*(spec_decode|mlpspeculator|eagle|speculation).*\.py
       - files~=^vllm/model_executor/models/.*eagle.*\.py
       - files=vllm/model_executor/models/mlp_speculator.py
       - files~=^vllm/transformers_utils/configs/(eagle|medusa|mlp_speculator)\.py
@@ -388,7 +389,11 @@ pull_request_rules:
       - files~=^tests/entrypoints/anthropic/.*tool.*
       - files~=^vllm/tool_parsers/
       - files=docs/features/tool_calling.md
-      - files~=^examples/tool_calling/
+      - files~=^examples/tool_chat_*
+      - files=examples/offline_inference/chat_with_tools.py
+      - files=examples/online_serving/openai_chat_completion_client_with_tools_required.py
+      - files=examples/online_serving/openai_chat_completion_tool_calls_with_reasoning.py
+      - files=examples/online_serving/openai_chat_completion_client_with_tools.py
   actions:
     label:
       add:
@@ -477,7 +482,9 @@ pull_request_rules:
   conditions:
     - label != stale
     - or:
-      - files~=^examples/disaggregated/
+      - files~=^examples/online_serving/disaggregated[^/]*/.*
+      - files~=^examples/offline_inference/disaggregated[^/]*/.*
+      - files~=^examples/others/lmcache/
       - files~=^tests/v1/kv_connector/
       - files~=^vllm/distributed/kv_transfer/
       - title~=(?i)\bP/?D\b

.github/workflows/pre-commit.yml

@@ -16,7 +16,11 @@ permissions:
 
 jobs:
   pre-run-check:
-    if: github.event_name == 'pull_request'
+    if: >-
+      github.event_name == 'pull_request' &&
+      (github.event.action != 'labeled' ||
+       github.event.label.name == 'ready' ||
+       github.event.label.name == 'verified')
     runs-on: ubuntu-latest
     steps:
     - name: Check PR label and author merge count
@@ -45,7 +49,12 @@ jobs:
 
   pre-commit:
     needs: pre-run-check
-    if: always() && (needs.pre-run-check.result == 'success' || needs.pre-run-check.result == 'skipped')
+    if: >-
+      always() &&
+      (github.event.action != 'labeled' ||
+       github.event.label.name == 'ready' ||
+       github.event.label.name == 'verified') &&
+      (needs.pre-run-check.result == 'success' || needs.pre-run-check.result == 'skipped')
     runs-on: ubuntu-latest
     steps:
     - uses: actions/checkout@8e8c483db84b4bee98b60c0593521ed34d9990e8 # v6.0.1

.github/workflows/scripts/build.sh

@@ -9,9 +9,6 @@ PATH=${cuda_home}/bin:$PATH
 LD_LIBRARY_PATH=${cuda_home}/lib64:$LD_LIBRARY_PATH
 
 # Install requirements
-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
 
 # Limit the number of parallel jobs to avoid OOM

.gitignore

@@ -237,7 +237,6 @@ ep_kernels_workspace/
 
 # Allow tracked library source folders under submodules (e.g., benchmarks/lib)
 !vllm/benchmarks/lib/
-!.buildkite/scripts/lib/
 
 # Generated gRPC protobuf files (compiled at build time from vllm_engine.proto)
 vllm/grpc/vllm_engine_pb2.py

.pre-commit-config.yaml

@@ -131,16 +131,6 @@ repos:
         --python-version, "3.12",
       ]
       files: ^requirements/(common|xpu|test/xpu)\.(in|txt)$
-    - id: pip-compile
-      alias: pip-compile-docs
-      name: pip-compile-docs
-      args: [
-        requirements/docs.in,
-        -o, requirements/docs.txt,
-        --python-platform, x86_64-manylinux_2_28,
-        --python-version, "3.12",
-      ]
-      files: ^requirements/docs\.(in|txt)$
 - repo: local
   hooks:
   - id: format-torch-nightly-test

.readthedocs.yaml

@@ -9,44 +9,7 @@ build:
     python: "3.12"
   jobs:
     post_checkout:
-      - |
-        if [ "$READTHEDOCS_VERSION_TYPE" = "external" ]; then
-          MAX_WAIT=300
-          INTERVAL=60
-          ELAPSED=0
-          while :; do
-            RAW=$(curl -sS -w "\n%{http_code}" "https://api.github.com/repos/vllm-project/vllm/commits/${READTHEDOCS_GIT_COMMIT_HASH}/check-runs?check_name=pre-run-check&filter=latest")
-            HTTP_CODE=$(printf %s "$RAW" | tail -n1)
-            BODY=$(printf %s "$RAW" | head -n -1)
-            if [ "$HTTP_CODE" != "200" ]; then
-              echo "GitHub API returned HTTP $HTTP_CODE (likely rate-limited); skipping pre-run-check gate."
-              break
-            fi
-            STATUS=$(printf %s "$BODY" | python3 -c "import sys, json; r=json.load(sys.stdin).get(\"check_runs\",[]); print((r[0].get(\"status\") or \"\") if r else \"none\")")
-            CONCLUSION=$(printf %s "$BODY" | python3 -c "import sys, json; r=json.load(sys.stdin).get(\"check_runs\",[]); print((r[0].get(\"conclusion\") or \"\") if r else \"\")")
-            if [ "$STATUS" = "none" ]; then
-              echo "no pre-run-check found for this commit; skipping gate."
-              break
-            fi
-            if [ -n "$CONCLUSION" ]; then
-              echo "pre-run-check conclusion: $CONCLUSION"
-              if [ "$CONCLUSION" = "failure" ] || [ "$CONCLUSION" = "cancelled" ] || [ "$CONCLUSION" = "timed_out" ]; then
-                echo "pre-run-check did not pass; failing docs build."
-                exit 1
-              fi
-              break
-            fi
-            if [ "$ELAPSED" -ge "$MAX_WAIT" ]; then
-              echo "pre-run-check status=$STATUS after ${MAX_WAIT}s; skipping gate."
-              break
-            fi
-            echo "pre-run-check status=$STATUS; waiting ${INTERVAL}s..."
-            sleep "$INTERVAL"
-            ELAPSED=$((ELAPSED + INTERVAL))
-          done
-        else
-          echo "Not a PR build (version type=$READTHEDOCS_VERSION_TYPE); skipping pre-run-check gate."
-        fi
+      # - bash docs/maybe_skip_pr_build.sh
       - git fetch origin main --unshallow --no-tags --filter=blob:none || true
     pre_create_environment:
       - pip install uv

CMakeLists.txt

@@ -13,12 +13,8 @@ cmake_minimum_required(VERSION 3.26)
 # cmake --install . --component _C
 project(vllm_extensions LANGUAGES CXX)
 
-set(CMAKE_CXX_STANDARD 20)
+set(CMAKE_CXX_STANDARD 17)
 set(CMAKE_CXX_STANDARD_REQUIRED ON)
-set(CMAKE_CUDA_STANDARD 20)
-set(CMAKE_CUDA_STANDARD_REQUIRED ON)
-set(CMAKE_HIP_STANDARD 20)
-set(CMAKE_HIP_STANDARD_REQUIRED ON)
 
 
 # CUDA by default, can be overridden by using -DVLLM_TARGET_DEVICE=... (used by setup.py)
@@ -109,24 +105,6 @@ else()
   set(CUDA_SUPPORTED_ARCHS "7.0;7.5;8.0;8.6;8.7;8.9;9.0")
 endif()
 
-#
-# spinloop extension (pure CXX; must stay above the non-CUDA device branch so
-# CPU builds define the target before the early return)
-#
-set(VLLM_SPINLOOP_EXT_SRC "csrc/spinloop.cpp")
-set(SPINLOOP_COMPILE_FLAGS "")
-if(CMAKE_SYSTEM_PROCESSOR MATCHES "x86_64|amd64")
-  list(APPEND SPINLOOP_COMPILE_FLAGS "-mmwaitx")
-endif()
-define_extension_target(
-  spinloop
-  DESTINATION vllm
-  LANGUAGE CXX
-  SOURCES ${VLLM_SPINLOOP_EXT_SRC}
-  COMPILE_FLAGS ${SPINLOOP_COMPILE_FLAGS}
-  USE_SABI 3.11
-  WITH_SOABI)
-
 #
 # Forward the non-CUDA device extensions to external CMake scripts.
 #
@@ -329,12 +307,12 @@ set(VLLM_EXT_SRC
   "csrc/quantization/activation_kernels.cu"
   "csrc/cuda_utils_kernels.cu"
   "csrc/custom_all_reduce.cu"
-  "csrc/torch_bindings.cpp"
-  "csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu")
+  "csrc/torch_bindings.cpp")
 
 if(VLLM_GPU_LANG STREQUAL "CUDA")
   list(APPEND VLLM_EXT_SRC
-    "csrc/minimax_reduce_rms_kernel.cu")
+    "csrc/minimax_reduce_rms_kernel.cu"
+    "csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu")
 
   SET(CUTLASS_ENABLE_HEADERS_ONLY ON CACHE BOOL "Enable only the header library")
 
@@ -369,6 +347,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
   FetchContent_MakeAvailable(cutlass)
 
   list(APPEND VLLM_EXT_SRC
+    "csrc/quantization/awq/gemm_kernels.cu"
     "csrc/cutlass_extensions/common.cpp")
 
   set_gencode_flags_for_srcs(
@@ -500,6 +479,46 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
                    " in CUDA target architectures")
   endif()
 
+  # Only build AllSpark kernels if we are building for at least some compatible archs.
+  cuda_archs_loose_intersection(ALLSPARK_ARCHS "8.0;8.6;8.7;8.9" "${CUDA_ARCHS}")
+  if (ALLSPARK_ARCHS)
+    set(ALLSPARK_SRCS
+       "csrc/quantization/gptq_allspark/allspark_repack.cu"
+       "csrc/quantization/gptq_allspark/allspark_qgemm_w8a16.cu")
+    set_gencode_flags_for_srcs(
+      SRCS "${ALLSPARK_SRCS}"
+      CUDA_ARCHS "${ALLSPARK_ARCHS}")
+    list(APPEND VLLM_EXT_SRC "${ALLSPARK_SRCS}")
+    message(STATUS "Building AllSpark kernels for archs: ${ALLSPARK_ARCHS}")
+  else()
+    message(STATUS "Not building AllSpark kernels as no compatible archs found"
+                   " in CUDA target architectures")
+  endif()
+
+  # CUTLASS MLA Archs and flags
+  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
+    cuda_archs_loose_intersection(MLA_ARCHS "10.0f;11.0f;12.0f" "${CUDA_ARCHS}")
+  else()
+    cuda_archs_loose_intersection(MLA_ARCHS "10.0a;10.1a;10.3a;12.0a;12.1a" "${CUDA_ARCHS}")
+  endif()
+  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND MLA_ARCHS)
+    set(SRCS
+      "csrc/attention/mla/sm100_cutlass_mla_kernel.cu")
+    set_gencode_flags_for_srcs(
+      SRCS "${SRCS}"
+      CUDA_ARCHS "${MLA_ARCHS}")
+    list(APPEND VLLM_EXT_SRC "${SRCS}")
+    list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MLA=1")
+    # Add MLA-specific include directories only to MLA source files
+    set_source_files_properties(${SRCS}
+      PROPERTIES INCLUDE_DIRECTORIES "${CUTLASS_DIR}/examples/77_blackwell_fmha;${CUTLASS_DIR}/examples/common")
+    message(STATUS "Building CUTLASS MLA for archs: ${MLA_ARCHS}")
+  else()
+    message(STATUS "Not building CUTLASS MLA as no compatible archs were found.")
+    # clear MLA_ARCHS
+    set(MLA_ARCHS)
+  endif()
+
   # Expert-specialization MXFP8 blockscaled grouped kernels (SM100+).
   if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
     cuda_archs_loose_intersection(ES_MXFP8_GROUPED_MM_ARCHS "10.0f;11.0f" "${CUDA_ARCHS}")
@@ -527,6 +546,24 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
     endif()
   endif()
 
+  # DeepSeek V3 fused A GEMM kernel (requires SM 9.0+, Hopper and later)
+  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
+    cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0f;11.0f" "${CUDA_ARCHS}")
+  else()
+    cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
+  endif()
+  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND DSV3_FUSED_A_GEMM_ARCHS)
+    set(DSV3_FUSED_A_GEMM_SRC "csrc/dsv3_fused_a_gemm.cu")
+    set_gencode_flags_for_srcs(
+      SRCS "${DSV3_FUSED_A_GEMM_SRC}"
+      CUDA_ARCHS "${DSV3_FUSED_A_GEMM_ARCHS}")
+    list(APPEND VLLM_EXT_SRC ${DSV3_FUSED_A_GEMM_SRC})
+    message(STATUS "Building dsv3_fused_a_gemm for archs: ${DSV3_FUSED_A_GEMM_ARCHS}")
+  else()
+    message(STATUS "Not building dsv3_fused_a_gemm as no compatible archs found "
+                   "in CUDA target architectures.")
+  endif()
+
   #
   # Machete kernels
 
@@ -598,6 +635,16 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
   endif()
 
 
+  # Hadacore kernels
+  cuda_archs_loose_intersection(HADACORE_ARCHS "8.0+PTX;9.0+PTX" "${CUDA_ARCHS}")
+  if(HADACORE_ARCHS)
+    set(SRCS "csrc/quantization/hadamard/hadacore/hadamard_transform_cuda.cu")
+    set_gencode_flags_for_srcs(
+      SRCS "${SRCS}"
+      CUDA_ARCHS "${HADACORE_ARCHS}")
+    list(APPEND VLLM_EXT_SRC "${SRCS}")
+    message(STATUS "Building hadacore")
+  endif()
 
 # if CUDA endif
 endif()
@@ -647,8 +694,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
     list(APPEND VLLM_STABLE_EXT_SRC
       "csrc/libtorch_stable/permute_cols.cu"
       "csrc/libtorch_stable/quantization/w8a8/fp8/per_token_group_quant.cu"
-      "csrc/libtorch_stable/quantization/w8a8/int8/per_token_group_quant.cu"
-      "csrc/libtorch_stable/quantization/awq/gemm_kernels.cu")
+      "csrc/libtorch_stable/quantization/w8a8/int8/per_token_group_quant.cu")
   endif()
 
   if(VLLM_GPU_LANG STREQUAL "CUDA")
@@ -657,40 +703,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
       CUDA_ARCHS "${CUDA_ARCHS}")
   endif()
 
-  # DeepSeek V3 fused A GEMM kernel (requires SM 9.0+, Hopper and later)
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0f;11.0f" "${CUDA_ARCHS}")
-  else()
-    cuda_archs_loose_intersection(DSV3_FUSED_A_GEMM_ARCHS "9.0a;10.0a;10.1a;10.3a" "${CUDA_ARCHS}")
-  endif()
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.0 AND DSV3_FUSED_A_GEMM_ARCHS)
-    set(SRCS "csrc/libtorch_stable/dsv3_fused_a_gemm.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${SRCS}"
-      CUDA_ARCHS "${DSV3_FUSED_A_GEMM_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    message(STATUS "Building dsv3_fused_a_gemm for archs: ${DSV3_FUSED_A_GEMM_ARCHS}")
-  else()
-    message(STATUS "Not building dsv3_fused_a_gemm as no compatible archs found "
-                   "in CUDA target architectures.")
-  endif()
-
-  # Only build AllSpark kernels if we are building for at least some compatible archs.
-  cuda_archs_loose_intersection(ALLSPARK_ARCHS "8.0;8.6;8.7;8.9" "${CUDA_ARCHS}")
-  if (ALLSPARK_ARCHS)
-    set(SRCS
-       "csrc/libtorch_stable/quantization/gptq_allspark/allspark_repack.cu"
-       "csrc/libtorch_stable/quantization/gptq_allspark/allspark_qgemm_w8a16.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${SRCS}"
-      CUDA_ARCHS "${ALLSPARK_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    message(STATUS "Building AllSpark kernels for archs: ${ALLSPARK_ARCHS}")
-  else()
-    message(STATUS "Not building AllSpark kernels as no compatible archs found"
-                   " in CUDA target architectures")
-  endif()
-
   #
   # CUTLASS scaled_mm kernels (moved from _C to _C_stable_libtorch)
   #
@@ -1000,41 +1012,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
     endif()
   endif()
 
-  # CUTLASS MLA Archs and flags
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 13.0)
-    cuda_archs_loose_intersection(MLA_ARCHS "10.0f;11.0f;12.0f" "${CUDA_ARCHS}")
-  else()
-    cuda_archs_loose_intersection(MLA_ARCHS "10.0a;10.1a;10.3a;12.0a;12.1a" "${CUDA_ARCHS}")
-  endif()
-  if(${CMAKE_CUDA_COMPILER_VERSION} VERSION_GREATER_EQUAL 12.8 AND MLA_ARCHS)
-    set(SRCS
-      "csrc/libtorch_stable/attention/mla/sm100_cutlass_mla_kernel.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${SRCS}"
-      CUDA_ARCHS "${MLA_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MLA=1")
-    # Add MLA-specific include directories only to MLA source files
-    set_source_files_properties(${SRCS}
-      PROPERTIES INCLUDE_DIRECTORIES "${CUTLASS_DIR}/examples/77_blackwell_fmha;${CUTLASS_DIR}/examples/common")
-    message(STATUS "Building CUTLASS MLA for archs: ${MLA_ARCHS}")
-  else()
-    message(STATUS "Not building CUTLASS MLA as no compatible archs were found.")
-    # clear MLA_ARCHS
-    set(MLA_ARCHS)
-  endif()
-
-  # Hadacore kernels
-  cuda_archs_loose_intersection(HADACORE_ARCHS "8.0+PTX;9.0+PTX" "${CUDA_ARCHS}")
-  if(HADACORE_ARCHS)
-    set(SRCS "csrc/libtorch_stable/quantization/hadamard/hadacore/hadamard_transform_cuda.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${SRCS}"
-      CUDA_ARCHS "${HADACORE_ARCHS}")
-    list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
-    message(STATUS "Building hadacore")
-  endif()
-
   message(STATUS "Enabling C_stable extension.")
   define_extension_target(
     _C_stable_libtorch
@@ -1070,13 +1047,14 @@ endif()
 set(VLLM_MOE_EXT_SRC
   "csrc/moe/torch_bindings.cpp"
   "csrc/moe/moe_align_sum_kernels.cu"
-  "csrc/moe/topk_softmax_kernels.cu"
-  "csrc/moe/topk_softplus_sqrt_kernels.cu")
+  "csrc/moe/topk_softmax_kernels.cu")
 
 if(VLLM_GPU_LANG STREQUAL "CUDA")
   list(APPEND VLLM_MOE_EXT_SRC
     "csrc/moe/moe_wna16.cu"
-    "csrc/moe/grouped_topk_kernels.cu")
+    "csrc/moe/grouped_topk_kernels.cu"
+    "csrc/moe/router_gemm.cu"
+    "csrc/moe/topk_softplus_sqrt_kernels.cu")
 endif()
 
 if(VLLM_GPU_LANG STREQUAL "CUDA")
@@ -1222,16 +1200,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
       CUDA_ARCHS "${DSV3_ROUTER_GEMM_ARCHS}")
     list(APPEND VLLM_MOE_EXT_SRC "${DSV3_ROUTER_GEMM_SRC}")
     message(STATUS "Building DSV3 router GEMM kernel for archs: ${DSV3_ROUTER_GEMM_ARCHS}")
-
-    # DeepSeek V4 fused RMSNorm + router GEMV - same arch gating as DSV3.
-    set(DSV4_NORM_ROUTER_GEMM_SRC
-      "csrc/moe/dsv4_norm_router_gemm_entry.cu"
-      "csrc/moe/dsv4_norm_router_gemm_kernel.cu")
-    set_gencode_flags_for_srcs(
-      SRCS "${DSV4_NORM_ROUTER_GEMM_SRC}"
-      CUDA_ARCHS "${DSV3_ROUTER_GEMM_ARCHS}")
-    list(APPEND VLLM_MOE_EXT_SRC "${DSV4_NORM_ROUTER_GEMM_SRC}")
-    message(STATUS "Building DSV4 norm+router GEMV kernel for archs: ${DSV3_ROUTER_GEMM_ARCHS}")
   else()
     message(STATUS "Not building DSV3 router GEMM kernel as no compatible archs found"
                    " (requires SM90+ and CUDA >= 12.0)")

benchmarks/attention_benchmarks/configs/mla_decode.yaml

@@ -53,7 +53,6 @@ backends:
   - FLASHINFER_MLA
   - FLASH_ATTN_MLA  # Hopper only
   - FLASHMLA        # Hopper only
-  - TOKENSPEED_MLA  # Blackwell + R1 dims + FP8 KV (use --kv-cache-dtype fp8)
 
 device: "cuda:0"
 repeats: 100

benchmarks/attention_benchmarks/configs/mla_prefill.yaml

@@ -3,7 +3,6 @@
 # Compares all available MLA prefill backends:
 #   FA backends:  fa2, fa3, fa4 (FlashAttention versions)
 #   Non-FA:       flashinfer, cudnn, trtllm (Blackwell-only, require flashinfer)
-#   CuTe DSL:     tokenspeed (Blackwell + R1 dims, requires tokenspeed_mla)
 #
 # Uses cutlass_mla as the decode backend for impl construction
 # (only the prefill path is exercised).
@@ -121,7 +120,6 @@ prefill_backends:
   - flashinfer
   - cudnn
   - trtllm
-  - tokenspeed
 
 device: "cuda:0"
 repeats: 20

benchmarks/attention_benchmarks/mla_runner.py

@@ -29,7 +29,6 @@ from vllm.config import (
     VllmConfig,
     set_current_vllm_config,
 )
-from vllm.v1.attention.backends.mla.prefill.registry import MLAPrefillBackendEnum
 
 # ============================================================================
 # VllmConfig Creation
@@ -80,8 +79,8 @@ def create_minimal_vllm_config(
         index_topk: Optional topk value for sparse MLA backends. If provided,
                     the config will include index_topk for sparse attention.
         prefill_backend: Prefill backend name (e.g., "fa3", "fa4", "flashinfer",
-                        "trtllm"). Configures the attention config to force
-                        the specified prefill backend.
+                        "cudnn", "trtllm"). Configures the attention config to
+                        force the specified prefill backend.
 
     Returns:
         VllmConfig for benchmarking
@@ -180,13 +179,19 @@ def create_minimal_vllm_config(
 
     if prefill_backend is not None:
         prefill_cfg = get_prefill_backend_config(prefill_backend)
-        vllm_config.attention_config.mla_prefill_backend = prefill_cfg[
-            "mla_prefill_backend"
-        ]
         if prefill_cfg["flash_attn_version"] is not None:
             vllm_config.attention_config.flash_attn_version = prefill_cfg[
                 "flash_attn_version"
             ]
+        vllm_config.attention_config.disable_flashinfer_prefill = prefill_cfg[
+            "disable_flashinfer_prefill"
+        ]
+        vllm_config.attention_config.use_cudnn_prefill = prefill_cfg[
+            "use_cudnn_prefill"
+        ]
+        vllm_config.attention_config.use_trtllm_ragged_deepseek_prefill = prefill_cfg[
+            "use_trtllm_ragged_deepseek_prefill"
+        ]
 
     return vllm_config
 
@@ -201,27 +206,39 @@ def create_minimal_vllm_config(
 _PREFILL_BACKEND_CONFIG: dict[str, dict] = {
     "fa2": {
         "flash_attn_version": 2,
-        "mla_prefill_backend": MLAPrefillBackendEnum.FLASH_ATTN,
+        "disable_flashinfer_prefill": True,
+        "use_cudnn_prefill": False,
+        "use_trtllm_ragged_deepseek_prefill": False,
     },
     "fa3": {
         "flash_attn_version": 3,
-        "mla_prefill_backend": MLAPrefillBackendEnum.FLASH_ATTN,
+        "disable_flashinfer_prefill": True,
+        "use_cudnn_prefill": False,
+        "use_trtllm_ragged_deepseek_prefill": False,
     },
     "fa4": {
         "flash_attn_version": 4,
-        "mla_prefill_backend": MLAPrefillBackendEnum.FLASH_ATTN,
+        "disable_flashinfer_prefill": True,
+        "use_cudnn_prefill": False,
+        "use_trtllm_ragged_deepseek_prefill": False,
     },
     "flashinfer": {
         "flash_attn_version": None,
-        "mla_prefill_backend": MLAPrefillBackendEnum.FLASHINFER,
+        "disable_flashinfer_prefill": False,
+        "use_cudnn_prefill": False,
+        "use_trtllm_ragged_deepseek_prefill": False,
+    },
+    "cudnn": {
+        "flash_attn_version": None,
+        "disable_flashinfer_prefill": True,
+        "use_cudnn_prefill": True,
+        "use_trtllm_ragged_deepseek_prefill": False,
     },
     "trtllm": {
         "flash_attn_version": None,
-        "mla_prefill_backend": MLAPrefillBackendEnum.TRTLLM_RAGGED,
-    },
-    "tokenspeed": {
-        "flash_attn_version": None,
-        "mla_prefill_backend": MLAPrefillBackendEnum.TOKENSPEED_MLA,
+        "disable_flashinfer_prefill": True,
+        "use_cudnn_prefill": False,
+        "use_trtllm_ragged_deepseek_prefill": True,
     },
 }
 
@@ -608,21 +625,6 @@ def _create_backend_impl(
     # Create mock layer
     layer = MockLayer(device, impl=impl, kv_cache_spec=kv_cache_spec)
 
-    # Attach a prefill backend (MLAAttention does this in __init__; the metadata
-    # builder reads layer.prefill_backend from static_forward_context).
-    from vllm.v1.attention.backends.mla.prefill import get_mla_prefill_backend
-
-    prefill_backend_cls = get_mla_prefill_backend(vllm_config)
-    layer.prefill_backend = prefill_backend_cls(
-        num_heads=mla_dims["num_q_heads"],
-        scale=(mla_dims["qk_nope_head_dim"] + mla_dims["qk_rope_head_dim"]) ** -0.5,
-        kv_lora_rank=mla_dims["kv_lora_rank"],
-        qk_nope_head_dim=mla_dims["qk_nope_head_dim"],
-        qk_rope_head_dim=mla_dims["qk_rope_head_dim"],
-        v_head_dim=mla_dims["v_head_dim"],
-        vllm_config=vllm_config,
-    )
-
     # Create builder instance if needed
     builder_instance = None
     if builder_class:
@@ -959,6 +961,19 @@ def _run_mla_benchmark_batched(
     results = []
 
     with set_current_vllm_config(vllm_config):
+        # Clear cached prefill backend detection functions so they re-evaluate
+        # with the current VllmConfig. These are @functools.cache decorated and
+        # would otherwise return stale results from a previous backend's config.
+        from vllm.model_executor.layers.attention.mla_attention import (
+            use_cudnn_prefill,
+            use_flashinfer_prefill,
+            use_trtllm_ragged_deepseek_prefill,
+        )
+
+        use_flashinfer_prefill.cache_clear()
+        use_cudnn_prefill.cache_clear()
+        use_trtllm_ragged_deepseek_prefill.cache_clear()
+
         # Create backend impl, layer, builder, and indexer (reused across benchmarks)
         impl, layer, builder_instance, indexer = _create_backend_impl(
             backend_cfg,
@@ -970,36 +985,38 @@ def _run_mla_benchmark_batched(
             kv_cache_dtype=kv_cache_dtype,
         )
 
-        # Verify the actual prefill backend matches what was requested. The
-        # selector + impl construction already raise on misuse; here we just
-        # check the resolved class against the requested name as a sanity guard.
+        # Verify the actual prefill backend matches what was requested
         if prefill_backend is not None:
-            expected_class = {
-                "fa2": "FlashAttnPrefillBackend",
-                "fa3": "FlashAttnPrefillBackend",
-                "fa4": "FlashAttnPrefillBackend",
-                "flashinfer": "FlashInferPrefillBackend",
-                "trtllm": "TrtllmRaggedPrefillBackend",
-                "tokenspeed": "TokenspeedMLAPrefillBackend",
-            }.get(prefill_backend)
-            actual_class = type(getattr(layer, "prefill_backend", None)).__name__
-            if expected_class and actual_class != expected_class:
-                raise RuntimeError(
-                    f"Prefill backend '{prefill_backend}' requested "
-                    f"{expected_class}, got {actual_class}. Check "
-                    f"attention_config plumbing or installed deps."
-                )
-            if prefill_backend in {"fa2", "fa3", "fa4"}:
-                fa_version = int(prefill_backend[2:])
-                actual_fa_version = getattr(
-                    layer.prefill_backend, "vllm_flash_attn_version", None
-                )
+            prefill_cfg = get_prefill_backend_config(prefill_backend)
+            fa_version = prefill_cfg["flash_attn_version"]
+
+            if fa_version is not None:
+                # FA backend: verify the impl's FA version
+                actual_fa_version = getattr(impl, "vllm_flash_attn_version", None)
                 if actual_fa_version != fa_version:
                     raise RuntimeError(
                         f"Prefill backend '{prefill_backend}' requested FA "
-                        f"version {fa_version}, got "
-                        f"{actual_fa_version} on {actual_class}."
+                        f"version {fa_version}, but the impl is using FA "
+                        f"version {actual_fa_version}. Check "
+                        f"vllm/v1/attention/backends/fa_utils.py."
                     )
+            else:
+                # Non-FA backend: verify the builder picked the right path
+                expected_flags = {
+                    "flashinfer": "_use_fi_prefill",
+                    "cudnn": "_use_cudnn_prefill",
+                    "trtllm": "_use_trtllm_ragged_prefill",
+                }
+                flag_name = expected_flags.get(prefill_backend)
+                if flag_name and not getattr(builder_instance, flag_name, False):
+                    raise RuntimeError(
+                        f"Prefill backend '{prefill_backend}' was requested "
+                        f"but the metadata builder did not enable it. This "
+                        f"usually means a dependency is missing (e.g., "
+                        f"flashinfer not installed) or the platform doesn't "
+                        f"support it."
+                    )
+
         # Run each benchmark with the shared impl
         for config, threshold, num_splits in configs_with_params:
             # Set threshold for this benchmark (FlashAttn/FlashMLA only)

benchmarks/kernels/benchmark_moe.py

@@ -27,10 +27,10 @@ from vllm.model_executor.layers.fused_moe.config import (
     RoutingMethodType,
     _get_config_dtype_str,
 )
-from vllm.model_executor.layers.fused_moe.experts.triton_deep_gemm_moe import (
+from vllm.model_executor.layers.fused_moe.fused_moe import *
+from vllm.model_executor.layers.fused_moe.triton_deep_gemm_moe import (
     TritonOrDeepGemmExperts,
 )
-from vllm.model_executor.layers.fused_moe.fused_moe import *
 from vllm.transformers_utils.config import get_config
 from vllm.triton_utils import triton
 from vllm.utils.argparse_utils import FlexibleArgumentParser

benchmarks/kernels/benchmark_norm_router_gemm.py

@@ -1,183 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-"""Benchmark and correctness check for ``ops.dsv4_norm_router_gemm``.
-
-Two implementations are compared:
-
-  1. ``unfused``   — ``vllm_ops.rms_norm`` then ``ops.dsv3_router_gemm``,
-                     i.e. the current vLLM hot path (two kernel launches).
-  2. ``fused``     — ``ops.dsv4_norm_router_gemm``, the new single-kernel
-                     fused path.
-
-Both produce ``(normed_x: bf16, router_logits: fp32)``.  The correctness
-check verifies that ``fused`` and ``unfused`` agree to within ~1 bf16
-ULP — that is the precision floor for this op.
-"""
-
-import argparse
-
-import torch
-
-from vllm import _custom_ops as vllm_ops
-from vllm.triton_utils import triton
-
-# The fused dsv4_norm_router_gemm kernel is templated only for DSV4-Pro
-# (hidden_size=7168, num_experts=384).  Other shapes fall back to the
-# unfused path on the Python side (NormGatedLinear), so benchmark only
-# the configuration that the fused kernel actually targets.
-HIDDEN_SIZE = 7168
-NUM_EXPERTS_CHOICES = (384,)
-RMS_EPS = 1e-6
-
-
-def unfused_norm_router_gemm(
-    x: torch.Tensor,
-    norm_weight: torch.Tensor,
-    gate_weight: torch.Tensor,
-    eps: float,
-) -> tuple[torch.Tensor, torch.Tensor]:
-    # Call ``_C::rms_norm`` directly (mirroring ``_dsv4_pro_norm_gate``'s
-    # fallback path) so the benchmarked baseline doesn't inherit any
-    # Python wrapper overhead or risk falling through to the native
-    # eager-primitive ``RMSNorm.forward_native`` path.
-    normed = torch.empty_like(x)
-    torch.ops._C.rms_norm(normed, x, norm_weight, eps)
-    logits = vllm_ops.dsv3_router_gemm(normed, gate_weight, torch.float32)
-    return normed, logits
-
-
-def fused_norm_router_gemm(
-    x: torch.Tensor,
-    norm_weight: torch.Tensor,
-    gate_weight: torch.Tensor,
-    eps: float,
-) -> tuple[torch.Tensor, torch.Tensor]:
-    return vllm_ops.dsv4_norm_router_gemm(x, norm_weight, gate_weight, eps)
-
-
-def _make_inputs(num_tokens: int, num_experts: int, hidden_size: int, seed: int = 0):
-    torch.manual_seed(seed)
-    device = "cuda"
-    x = torch.randn(num_tokens, hidden_size, dtype=torch.bfloat16, device=device)
-    norm_w = torch.randn(hidden_size, dtype=torch.bfloat16, device=device)
-    gate_w = torch.randn(num_experts, hidden_size, dtype=torch.bfloat16, device=device)
-    # Down-scale gate_w so the GEMV output stays in a representable range.
-    gate_w = gate_w / float(hidden_size) ** 0.5
-    norm_w = (norm_w * 0.1) + 1.0
-    return x, norm_w, gate_w
-
-
-def calculate_diff(
-    num_tokens: int,
-    num_experts: int,
-    hidden_size: int = HIDDEN_SIZE,
-    normed_atol: float = 2e-3,
-    logits_atol: float = 1e-2,
-    rtol: float = 1e-2,
-) -> None:
-    x, norm_w, gate_w = _make_inputs(num_tokens, num_experts, hidden_size)
-
-    normed_unfused, logits_unfused = unfused_norm_router_gemm(
-        x.clone(), norm_w, gate_w, RMS_EPS
-    )
-    normed_fused, logits_fused = fused_norm_router_gemm(
-        x.clone(), norm_w, gate_w, RMS_EPS
-    )
-
-    def _max_abs(a, b):
-        return (a.float() - b.float()).abs().max().item()
-
-    print(f"\n=== M={num_tokens} E={num_experts} H={hidden_size} ===")
-    print(f"normed_x  |fused - unfused| = {_max_abs(normed_fused, normed_unfused):.3e}")
-    print(f"logits    |fused - unfused| = {_max_abs(logits_fused, logits_unfused):.3e}")
-
-    ok_normed = torch.allclose(
-        normed_fused.float(),
-        normed_unfused.float(),
-        atol=normed_atol,
-        rtol=rtol,
-    )
-    ok_logits = torch.allclose(
-        logits_fused.float(),
-        logits_unfused.float(),
-        atol=logits_atol,
-        rtol=rtol,
-    )
-    if ok_normed and ok_logits:
-        print(
-            f"OK   fused vs unfused within "
-            f"normed_atol={normed_atol:.0e} logits_atol={logits_atol:.0e} "
-            f"rtol={rtol:.0e}"
-        )
-    else:
-        print(
-            f"FAIL normed_ok={ok_normed} logits_ok={ok_logits}; "
-            f"see max-abs values above"
-        )
-
-
-def get_benchmark():
-    # Only num_tokens varies (DSV4-Pro hard-codes E=384); single-axis
-    # sweep yields a clean line plot with M on the x-axis.
-    num_experts = NUM_EXPERTS_CHOICES[0]
-
-    @triton.testing.perf_report(
-        triton.testing.Benchmark(
-            x_names=["num_tokens"],
-            x_vals=list(range(1, 17)),
-            line_arg="provider",
-            line_vals=["unfused", "fused"],
-            line_names=["unfused (rms+dsv3)", "fused (dsv4)"],
-            styles=[("green", "-"), ("red", "-")],
-            ylabel="us",
-            plot_name=f"norm-router-gemm-E{num_experts}-H{HIDDEN_SIZE}",
-            args={},
-        )
-    )
-    def benchmark(num_tokens, provider):
-        x, norm_w, gate_w = _make_inputs(num_tokens, num_experts, HIDDEN_SIZE)
-
-        quantiles = [0.5, 0.2, 0.8]
-        if provider == "unfused":
-            fn = lambda: unfused_norm_router_gemm(  # noqa: E731
-                x, norm_w, gate_w, RMS_EPS
-            )
-        else:
-            fn = lambda: fused_norm_router_gemm(  # noqa: E731
-                x, norm_w, gate_w, RMS_EPS
-            )
-
-        ms, min_ms, max_ms = triton.testing.do_bench(fn, quantiles=quantiles)
-        return 1000 * ms, 1000 * max_ms, 1000 * min_ms
-
-    return benchmark
-
-
-def main() -> None:
-    parser = argparse.ArgumentParser()
-    parser.add_argument(
-        "--save-path",
-        type=str,
-        default="./configs/norm_router_gemm/",
-    )
-    parser.add_argument(
-        "--skip-bench",
-        action="store_true",
-        help="Run only the correctness check, not the perf sweep.",
-    )
-    args = parser.parse_args()
-
-    # Correctness sweep over the full fast-path range M=1..16.
-    for m in range(1, 17):
-        for e in NUM_EXPERTS_CHOICES:
-            calculate_diff(num_tokens=m, num_experts=e, hidden_size=HIDDEN_SIZE)
-
-    if args.skip_bench:
-        return
-
-    benchmark = get_benchmark()
-    benchmark.run(print_data=True, save_path=args.save_path)
-
-
-if __name__ == "__main__":
-    main()

benchmarks/kernels/benchmark_vit_fp8_attn.py

@@ -1,324 +0,0 @@
-# SPDX-License-Identifier: Apache-2.0
-# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-# Benchmarks FP8 vs BF16 ViT attention via FlashInfer cuDNN backend.
-#
-# == Usage Examples ==
-#
-# Benchmark mode (default, FlashInfer CUDAGraph Bench)
-#   python3 benchmark_vit_fp8_attn.py
-#
-# Profile mode (PyTorch profiler, saves TensorBoard traces):
-#   python3 benchmark_vit_fp8_attn.py --profile
-#   python3 benchmark_vit_fp8_attn.py --profile --profile-output-dir ./profile_traces
-#
-# Custom seq_lens:
-#   python3 benchmark_vit_fp8_attn.py --seq-lens 4096 8192 16384
-
-from functools import partial
-
-import numpy as np
-import torch
-from torch.profiler import ProfilerActivity, profile, record_function
-
-from vllm.utils.argparse_utils import FlexibleArgumentParser
-
-# Qwen3-VL defaults
-NUM_HEADS = 16
-HEAD_DIM = 72
-DEFAULT_SEQ_LENS = [2304, 4096, 8192, 16384]
-
-
-def _setup_fp8_attention(num_heads: int, head_dim: int) -> tuple:
-    """Create FP8 and BF16 attention modules + workspace."""
-    from types import SimpleNamespace
-    from unittest.mock import patch
-
-    from vllm.config import VllmConfig, set_current_vllm_config
-    from vllm.config.multimodal import MultiModalConfig
-    from vllm.model_executor.layers.attention.mm_encoder_attention import (
-        MMEncoderAttention,
-        _get_flashinfer_workspace_buffer,
-    )
-    from vllm.v1.attention.backends.registry import AttentionBackendEnum
-
-    old_dtype = torch.get_default_dtype()
-    torch.set_default_dtype(torch.bfloat16)
-
-    backend_patch = patch(
-        "vllm.model_executor.layers.attention.mm_encoder_attention"
-        ".get_vit_attn_backend",
-        return_value=AttentionBackendEnum.FLASHINFER,
-    )
-
-    # FP8 attention
-    mm_config_fp8 = MultiModalConfig(mm_encoder_attn_dtype="fp8")
-    vllm_config_fp8 = VllmConfig()
-    vllm_config_fp8.model_config = SimpleNamespace(multimodal_config=mm_config_fp8)
-    with set_current_vllm_config(vllm_config_fp8), backend_patch:
-        attn_fp8 = MMEncoderAttention(
-            num_heads=num_heads,
-            head_size=head_dim,
-            prefix="visual.blocks.0.attn",
-        ).to("cuda")
-
-    # BF16 attention (no FP8)
-    with set_current_vllm_config(VllmConfig()), backend_patch:
-        attn_bf16 = MMEncoderAttention(
-            num_heads=num_heads,
-            head_size=head_dim,
-            prefix="visual.blocks.0.attn",
-        ).to("cuda")
-
-    torch.set_default_dtype(old_dtype)
-
-    workspace = _get_flashinfer_workspace_buffer()
-    return attn_fp8, attn_bf16, workspace
-
-
-def _build_meta(
-    seq_len: int,
-    num_heads: int,
-    head_dim: int,
-    fp8: bool,
-):
-    """Build cu_seqlens, max_seqlen, sequence_lengths."""
-    from vllm.model_executor.layers.attention.mm_encoder_attention import (
-        MMEncoderAttention,
-    )
-    from vllm.utils.math_utils import round_up
-    from vllm.v1.attention.backends.registry import AttentionBackendEnum
-
-    cu_np = np.array([0, seq_len], dtype=np.int32)
-    fp8_padded = num_heads * round_up(head_dim, 16) if fp8 else None
-
-    seq_lengths = MMEncoderAttention.maybe_compute_seq_lens(
-        AttentionBackendEnum.FLASHINFER, cu_np, torch.device("cuda")
-    )
-    max_seqlen = torch.tensor(
-        MMEncoderAttention.compute_max_seqlen(AttentionBackendEnum.FLASHINFER, cu_np),
-        dtype=torch.int32,
-    )
-    cu_seqlens = MMEncoderAttention.maybe_recompute_cu_seqlens(
-        AttentionBackendEnum.FLASHINFER,
-        cu_np,
-        num_heads * head_dim,
-        1,
-        torch.device("cuda"),
-        fp8_padded_hidden_size=fp8_padded,
-    )
-    return cu_seqlens, max_seqlen, seq_lengths
-
-
-def run_benchmark(
-    seq_lens: list[int],
-    num_heads: int,
-    head_dim: int,
-    method: str,
-):
-    """Benchmark FP8 vs BF16 attention across seq_lens.
-
-    Uses FlashInfer GPU-level timing to measure pure kernel time,
-    excluding CPU launch overhead.
-    """
-    if method == "cupti":
-        from flashinfer.testing import bench_gpu_time_with_cupti as bench_fn
-
-        bench_fn = partial(bench_fn, use_cuda_graph=True, cold_l2_cache=False)
-    elif method == "cudagraph":
-        from flashinfer.testing import (
-            bench_gpu_time_with_cudagraph as bench_fn,
-        )
-
-        bench_fn = partial(bench_fn, cold_l2_cache=False)
-    else:
-        raise ValueError(f"Invalid method: {method}")
-
-    attn_fp8, attn_bf16, workspace = _setup_fp8_attention(num_heads, head_dim)
-
-    print(f"Timing method: {method}")
-    print(f"{'seq_len':>8} {'BF16 (us)':>12} {'FP8 (us)':>12} {'Speedup':>10}")
-    print("-" * 46)
-
-    for seq_len in seq_lens:
-        torch.manual_seed(42)
-
-        q = torch.randn(
-            seq_len,
-            num_heads,
-            head_dim,
-            device="cuda",
-            dtype=torch.bfloat16,
-        )
-        k = torch.randn_like(q)
-        v = torch.randn_like(q)
-
-        cu_fp8, max_s, seq_l = _build_meta(seq_len, num_heads, head_dim, fp8=True)
-        # we can reuse cu_fp8 for cu_bf16 since q, k, and v are contiguous
-        cu_bf16 = cu_fp8.clone()
-
-        def bf16_fn(q=q, k=k, v=v, cu=cu_bf16, ms=max_s, sl=seq_l):
-            attn_bf16._forward_flashinfer(q, k, v, cu, ms, sl)
-
-        def fp8_fn(q=q, k=k, v=v, cu=cu_fp8, ms=max_s, sl=seq_l):
-            attn_fp8._forward_flashinfer(q, k, v, cu, ms, sl)
-
-        # bench_fn returns List[float] of per-iteration times in ms
-        bf16_times = bench_fn(bf16_fn)
-        fp8_times = bench_fn(fp8_fn)
-
-        bf16_us = np.median(bf16_times) * 1e3  # ms -> us
-        fp8_us = np.median(fp8_times) * 1e3
-        speedup = bf16_us / fp8_us if fp8_us > 0 else float("inf")
-
-        print(f"{seq_len:>8} {bf16_us:>12.1f} {fp8_us:>12.1f} {speedup:>9.2f}x")
-
-
-def _make_trace_handler(output_dir: str, worker_name: str, label: str):
-    """Create a trace handler that saves to TensorBoard and prints summary."""
-
-    def handler(prof):
-        torch.profiler.tensorboard_trace_handler(output_dir, worker_name)(prof)
-        print(f"\n{'=' * 80}")
-        print(label)
-        print(f"{'=' * 80}")
-        print(prof.key_averages().table(sort_by="cuda_time_total", row_limit=20))
-
-    return handler
-
-
-def run_profile(
-    seq_len: int,
-    num_heads: int,
-    head_dim: int,
-    warmup: int,
-    output_dir: str,
-):
-    """Profile FP8 vs BF16 attention with PyTorch profiler."""
-    attn_fp8, attn_bf16, workspace = _setup_fp8_attention(num_heads, head_dim)
-
-    torch.manual_seed(42)
-    q = torch.randn(
-        seq_len,
-        num_heads,
-        head_dim,
-        device="cuda",
-        dtype=torch.bfloat16,
-    )
-    k = torch.randn_like(q)
-    v = torch.randn_like(q)
-
-    cu_fp8, max_s, seq_l = _build_meta(seq_len, num_heads, head_dim, fp8=True)
-    # we can reuse cu_fp8 for cu_bf16 since q, k, and v are contiguous
-    cu_bf16 = cu_fp8.clone()
-
-    sched = torch.profiler.schedule(wait=0, warmup=warmup, active=1)
-
-    # Profile BF16 (warmup handled by profiler schedule)
-    with profile(
-        activities=[ProfilerActivity.CPU, ProfilerActivity.CUDA],
-        schedule=sched,
-        on_trace_ready=_make_trace_handler(
-            output_dir,
-            f"bf16_h{head_dim}_s{seq_len}",
-            f"BF16 Attention (seq_len={seq_len}, heads={num_heads}, "
-            f"head_dim={head_dim})",
-        ),
-    ) as prof_bf16:
-        for _ in range(warmup + 1):
-            with record_function("bf16_attention"):
-                attn_bf16._forward_flashinfer(
-                    q.clone(), k.clone(), v.clone(), cu_bf16, max_s, seq_l
-                )
-                torch.accelerator.synchronize()
-            prof_bf16.step()
-
-    # Profile FP8 (warmup handled by profiler schedule)
-    with profile(
-        activities=[ProfilerActivity.CPU, ProfilerActivity.CUDA],
-        schedule=sched,
-        on_trace_ready=_make_trace_handler(
-            output_dir,
-            f"fp8_h{head_dim}_s{seq_len}",
-            f"FP8 Attention (seq_len={seq_len}, heads={num_heads}, "
-            f"head_dim={head_dim})",
-        ),
-    ) as prof_fp8:
-        for _ in range(warmup + 1):
-            with record_function("fp8_attention"):
-                attn_fp8._forward_flashinfer(
-                    q.clone(), k.clone(), v.clone(), cu_fp8, max_s, seq_l
-                )
-                torch.accelerator.synchronize()
-            prof_fp8.step()
-
-    print(f"\nTensorBoard traces saved to: {output_dir}")
-    print(f"View with: tensorboard --logdir={output_dir}")
-
-
-if __name__ == "__main__":
-    parser = FlexibleArgumentParser(description="Benchmark FP8 vs BF16 ViT attention.")
-    parser.add_argument(
-        "--seq-lens",
-        type=int,
-        nargs="+",
-        default=DEFAULT_SEQ_LENS,
-        help="Sequence lengths to benchmark",
-    )
-    parser.add_argument(
-        "--num-heads",
-        type=int,
-        default=NUM_HEADS,
-    )
-    parser.add_argument(
-        "--head-dim",
-        type=int,
-        default=HEAD_DIM,
-    )
-    parser.add_argument(
-        "--method",
-        choices=["cupti", "cudagraph"],
-        default="cudagraph",
-        help="GPU timing method: cupti (CUPTI kernel timing) or "
-        "cudagraph (CUDA graph capture/replay). Default: cudagraph",
-    )
-    parser.add_argument(
-        "--warmup",
-        type=int,
-        default=10,
-        help="Warmup iterations (profile mode only)",
-    )
-    parser.add_argument(
-        "--profile",
-        action="store_true",
-        help="Run PyTorch profiler instead of benchmark",
-    )
-    parser.add_argument(
-        "--profile-seq-len",
-        type=int,
-        default=8192,
-        help="Sequence length for profiling (default: 8192)",
-    )
-    parser.add_argument(
-        "--profile-output-dir",
-        type=str,
-        default="./profile_traces",
-        help="Output directory for TensorBoard traces (default: ./profile_traces)",
-    )
-    args = parser.parse_args()
-
-    if args.profile:
-        run_profile(
-            args.profile_seq_len,
-            args.num_heads,
-            args.head_dim,
-            args.warmup,
-            args.profile_output_dir,
-        )
-    else:
-        run_benchmark(
-            args.seq_lens,
-            args.num_heads,
-            args.head_dim,
-            args.method,
-        )

benchmarks/kernels/cpu/benchmark_cpu_attn.py

@@ -12,6 +12,7 @@ from vllm._custom_ops import (
     cpu_attn_get_scheduler_metadata,
     cpu_attn_reshape_and_cache,
 )
+from vllm.platforms import CpuArchEnum, current_platform
 from vllm.utils.argparse_utils import FlexibleArgumentParser
 from vllm.utils.torch_utils import STR_DTYPE_TO_TORCH_DTYPE, set_random_seed
 from vllm.v1.attention.backends.cpu_attn import CPUAttentionBackend, _get_attn_isa
@@ -21,14 +22,15 @@ def get_attn_isa(
     block_size: int | None = None,
     dtype: torch.dtype | None = None,
 ):
-    # Delegate to _get_attn_isa so the fallback path applies the same arch
-    # gating (e.g. RISC-V RVV is only chosen when the build's hardcoded
-    # VLEN=128 kernel is actually present; on VLEN=256 / scalar hosts it
-    # correctly falls through to vec/vec16).
-    return _get_attn_isa(
-        dtype if dtype is not None else torch.bfloat16,
-        block_size if block_size else 32,
-    )
+    if block_size and dtype:
+        return _get_attn_isa(dtype, block_size)
+    else:
+        if current_platform.get_cpu_architecture() == CpuArchEnum.ARM:
+            return "neon"
+        elif torch.cpu._is_amx_tile_supported():
+            return "amx"
+        else:
+            return "vec"
 
 
 # rand number generation takes too much time, cache rand tensors
@@ -233,7 +235,7 @@ if __name__ == "__main__":
     )
     parser.add_argument("--use-sink", action="store_true")
     parser.add_argument(
-        "--isa", type=str, choices=["vec", "neon", "amx", "vec16", "rvv"], default=None
+        "--isa", type=str, choices=["vec", "neon", "amx", "vec16"], default=None
     )
     parser.add_argument("--seed", type=int, default=0)
     parser.add_argument("--iters", type=int, default=20)

benchmarks/multi_turn/benchmark_serving_multi_turn.py

@@ -217,7 +217,6 @@ async def send_request(
     min_tokens: int | None = None,
     max_tokens: int | None = None,
     timeout_sec: int = 120,
-    conversation_id: str | None = None,
 ) -> ServerResponse:
     payload = {
         "model": model,
@@ -226,9 +225,6 @@ async def send_request(
         "temperature": 0.0,
     }
 
-    if conversation_id is not None:
-        payload["conversation_id"] = conversation_id
-
     if stream:
         payload["stream"] = True
         payload["stream_options"] = {"include_usage": False}
@@ -423,7 +419,6 @@ async def send_turn(
         min_tokens,
         max_tokens,
         req_args.timeout_sec,
-        conversation_id=conv_id,
     )
 
     if response.valid is False:
@@ -1473,12 +1468,6 @@ async def main() -> None:
         "(for example: --warmup-percentages=0%%,50%%)",
     )
 
-    parser.add_argument(
-        "--trust-remote-code",
-        action="store_true",
-        help="Trust remote code when loading the tokenizer.",
-    )
-
     args = parser.parse_args()
 
     logger.info(args)
@@ -1521,9 +1510,7 @@ async def main() -> None:
     np.random.seed(args.seed)
 
     logger.info("Loading tokenizer")
-    tokenizer = AutoTokenizer.from_pretrained(
-        args.model, trust_remote_code=args.trust_remote_code
-    )
+    tokenizer = AutoTokenizer.from_pretrained(args.model)
 
     await get_server_info(args.url)
 

cmake/cpu_extension.cmake

@@ -1,7 +1,7 @@
 include(FetchContent)
 
 set(CMAKE_CXX_STANDARD_REQUIRED ON)
-set(CMAKE_CXX_STANDARD 20)
+set(CMAKE_CXX_STANDARD 17)
 set(CMAKE_CXX_EXTENSIONS ON)
 set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
 
@@ -32,23 +32,18 @@ else()
         "-DVLLM_CPU_EXTENSION")
 
     # locate PyTorch's libgomp (e.g. site-packages/torch.libs/libgomp-947d5fa1.so.1.0.0)
-    # and create a local shim dir with it. When PyTorch is built from source or packaged
-    # by a distro (common on RISC-V, s390x, Fedora/RHEL aarch64), no vendored libgomp
-    # exists and the shim dir is empty; fall back to the system libgomp in that case.
+    # and create a local shim dir with it
     vllm_prepare_torch_gomp_shim(VLLM_TORCH_GOMP_SHIM_DIR)
 
-    if(VLLM_TORCH_GOMP_SHIM_DIR)
-        find_library(OPEN_MP
-            NAMES gomp
-            PATHS "${VLLM_TORCH_GOMP_SHIM_DIR}"
-            NO_DEFAULT_PATH
-            REQUIRED
-        )
-        # Use the same libgomp as PyTorch at runtime
+    find_library(OPEN_MP
+        NAMES gomp
+        PATHS ${VLLM_TORCH_GOMP_SHIM_DIR}
+        NO_DEFAULT_PATH
+        REQUIRED
+    )
+    # Set LD_LIBRARY_PATH to include the shim dir at build time to use the same libgomp as PyTorch
+    if (OPEN_MP)
         set(ENV{LD_LIBRARY_PATH} "${VLLM_TORCH_GOMP_SHIM_DIR}:$ENV{LD_LIBRARY_PATH}")
-    else()
-        # Fall back to system / toolchain libgomp
-        find_library(OPEN_MP NAMES gomp REQUIRED)
     endif()
 endif()
 
@@ -195,12 +190,10 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
     endif()
     if(VLLM_RVV_VLEN AND VLLM_RVV_VLEN GREATER 0)
         message(STATUS "RISC-V RVV VLEN=${VLLM_RVV_VLEN}")
-        # Sources gate FP16/BF16 paths on the compiler-provided
-        # __riscv_zvfh / __riscv_zvfbfmin macros, which GCC and clang
-        # define automatically when those extensions appear in -march.
         if(RVV_BF16_FOUND)
             message(STATUS "BF16 extension detected")
             set(MARCH_FLAGS -march=rv64gcv_zvfh_zfbfmin_zvfbfmin_zvl${VLLM_RVV_VLEN}b -mrvv-vector-bits=zvl -mabi=lp64d)
+            add_compile_definitions(RISCV_BF16_SUPPORT)
         elseif(RVV_FP16_FOUND)
             message(WARNING "BF16 functionality is not available")
             set(MARCH_FLAGS -march=rv64gcv_zvfh_zvl${VLLM_RVV_VLEN}b -mrvv-vector-bits=zvl -mabi=lp64d)
@@ -328,6 +321,14 @@ if (ENABLE_X86_ISA OR (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND) OR POWER9_FOUND
     set(ONEDNN_VERBOSE "ON")
     set(CMAKE_POLICY_DEFAULT_CMP0077 NEW)
 
+    # TODO: Refactor this
+    if (ENABLE_X86_ISA)
+        # Note: only enable oneDNN for AVX512
+        list(APPEND DNNL_COMPILE_FLAGS ${CXX_COMPILE_FLAGS_AVX512})
+    else()
+        list(APPEND DNNL_COMPILE_FLAGS ${CXX_COMPILE_FLAGS})
+    endif()
+
     set(VLLM_BUILD_TYPE ${CMAKE_BUILD_TYPE})
     set(CMAKE_BUILD_TYPE "Release") # remove oneDNN debug symbols to reduce size
     FetchContent_MakeAvailable(oneDNN)
@@ -340,14 +341,8 @@ if (ENABLE_X86_ISA OR (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND) OR POWER9_FOUND
         PRIVATE ${oneDNN_SOURCE_DIR}/src
     )
     target_link_libraries(dnnl_ext dnnl torch)
-    if (ENABLE_X86_ISA)
-        target_compile_options(dnnl_ext PRIVATE ${CXX_COMPILE_FLAGS_AVX2} -fPIC)
-    else()
-        target_compile_options(dnnl_ext PRIVATE ${CXX_COMPILE_FLAGS} -fPIC)
-    endif()
+    target_compile_options(dnnl_ext PRIVATE ${DNNL_COMPILE_FLAGS} -fPIC)
     list(APPEND LIBS dnnl_ext)
-
-
     set(USE_ONEDNN ON)
 else()
     set(USE_ONEDNN OFF)
@@ -411,15 +406,12 @@ 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"
         "csrc/cpu/sgl-kernels/gemm_fp8.cpp"
         "csrc/cpu/sgl-kernels/gemm_int4.cpp"
         "csrc/cpu/sgl-kernels/moe.cpp"
         "csrc/cpu/sgl-kernels/moe_int8.cpp"
-        "csrc/cpu/sgl-kernels/moe_int4.cpp"
         "csrc/cpu/sgl-kernels/moe_fp8.cpp")
 
     set(VLLM_EXT_SRC_AVX512
@@ -438,11 +430,10 @@ if (ENABLE_X86_ISA)
         "csrc/cpu/pos_encoding.cpp"
         "csrc/moe/dynamic_4bit_int_moe_cpu.cpp") 
 
-    set(VLLM_EXT_SRC_AVX2
+    set(VLLM_EXT_SRC_AVX2 
         "csrc/cpu/utils.cpp"
         "csrc/cpu/spec_decode_utils.cpp"
         "csrc/cpu/cpu_attn.cpp"
-        "csrc/cpu/dnnl_kernels.cpp"
         "csrc/cpu/torch_bindings.cpp"
         # TODO: Remove these files
         "csrc/cpu/activation.cpp"
@@ -457,7 +448,7 @@ if (ENABLE_X86_ISA)
 
     set(_C_LIBS numa dnnl_ext)
     set(_C_AVX512_LIBS numa dnnl_ext)
-    set(_C_AVX2_LIBS numa dnnl_ext)
+    set(_C_AVX2_LIBS numa)
 
     # AMX + AVX512F + AVX512BF16 + AVX512VNNI
     define_extension_target(

cmake/external_projects/deepgemm.cmake

@@ -53,80 +53,49 @@ cuda_archs_loose_intersection(DEEPGEMM_ARCHS
 if(DEEPGEMM_ARCHS)
   message(STATUS "DeepGEMM CUDA architectures: ${DEEPGEMM_ARCHS}")
 
-  #
-  # DeepGEMM integration notes
-  # --------------------------
-  # We vendor DeepGEMM into vllm/third_party/deep_gemm/ and bundle a
-  # `_C.cpython-X.Y-*.so` for every CPython in `requires-python`. The
-  # per-Python build is delegated to tools/build_deepgemm_C.py.
-  #
-  # Why per-Python: DeepGEMM's binding uses PYBIND11_MODULE, which links
-  # private CPython symbols — a single `_C.abi3.so` is not viable today
-  # (see #41476 / #41512 for the failed attempt).
-  #
-  # TODOs (tracked in vllm-project/vllm#42431):
-  #   - Replace DeepGEMM's pybind11 binding with a TORCH_LIBRARY + shim
-  #     binding (cf. vllm-flash-attention/csrc/common/pytorch_shim.h) to
-  #     collapse to one `_C.abi3.so`. Needs either an upstream change or
-  #     a maintained binding fork in vLLM.
-  #   - AOT-compile DeepGEMM's CUDA kernels instead of runtime JIT to drop
-  #     the vendored CUTLASS/CCCL headers and the CUDA-toolkit-at-runtime
-  #     requirement.
-  #
+  find_package(CUDAToolkit REQUIRED)
 
-  # DEEPGEMM_PYTHON_INTERPRETERS: ":"-separated target Python paths.
-  # Empty/unset → fall back to the build interpreter (editable installs).
-  # (Empty-but-set env vars test as DEFINED in cmake — treat as unset.)
-  if(NOT "$ENV{DEEPGEMM_PYTHON_INTERPRETERS}" STREQUAL "")
-    string(REPLACE ":" ";" _dg_pythons "$ENV{DEEPGEMM_PYTHON_INTERPRETERS}")
-  else()
-    set(_dg_pythons "${Python_EXECUTABLE}")
-  endif()
-  message(STATUS "DeepGEMM _C will be built for: ${_dg_pythons}")
+  #
+  # Build the _C pybind11 extension from DeepGEMM's C++ source.
+  # This is a CXX-only module — CUDA kernels are JIT-compiled at runtime.
+  #
+  Python_add_library(_deep_gemm_C MODULE WITH_SOABI
+    "${deepgemm_SOURCE_DIR}/csrc/python_api.cpp")
 
-  # add_custom_command does no implicit header scanning; glob explicitly so
-  # header-only edits in DeepGEMM/cutlass/fmt re-trigger the rebuild.
-  file(GLOB_RECURSE _dg_headers
-    "${deepgemm_SOURCE_DIR}/csrc/*.h"
-    "${deepgemm_SOURCE_DIR}/csrc/*.hpp"
-    "${deepgemm_SOURCE_DIR}/deep_gemm/include/*.h"
-    "${deepgemm_SOURCE_DIR}/deep_gemm/include/*.hpp"
-    "${deepgemm_SOURCE_DIR}/deep_gemm/include/*.cuh")
+  # The pybind11 module name must be _C to match DeepGEMM's Python imports.
+  set_target_properties(_deep_gemm_C PROPERTIES OUTPUT_NAME "_C")
 
-  set(_dg_markers)
-  set(_dg_seen_soabis)
-  foreach(_pybin IN LISTS _dg_pythons)
-    execute_process(
-      COMMAND "${_pybin}" -c
-        "import sysconfig; print(sysconfig.get_config_var('SOABI'))"
-      OUTPUT_VARIABLE _dg_soabi
-      OUTPUT_STRIP_TRAILING_WHITESPACE
-      COMMAND_ERROR_IS_FATAL ANY)
-    # Dedup interpreters that resolve to the same CPython.
-    if(_dg_soabi IN_LIST _dg_seen_soabis)
-      continue()
-    endif()
-    list(APPEND _dg_seen_soabis "${_dg_soabi}")
-    set(_dg_dir "${CMAKE_CURRENT_BINARY_DIR}/deepgemm_C_${_dg_soabi}")
-    set(_dg_marker "${_dg_dir}/.built")
-    add_custom_command(
-      OUTPUT "${_dg_marker}"
-      COMMAND "${Python_EXECUTABLE}"
-              "${CMAKE_SOURCE_DIR}/tools/build_deepgemm_C.py"
-              "${deepgemm_SOURCE_DIR}" "${_dg_dir}" "${_pybin}"
-      COMMAND "${CMAKE_COMMAND}" -E touch "${_dg_marker}"
-      DEPENDS "${CMAKE_SOURCE_DIR}/tools/build_deepgemm_C.py"
-              "${deepgemm_SOURCE_DIR}/csrc/python_api.cpp"
-              ${_dg_headers}
-      COMMENT "Building DeepGEMM _C for ${_pybin}"
-      VERBATIM)
-    list(APPEND _dg_markers "${_dg_marker}")
-    install(DIRECTORY "${_dg_dir}/"
-      DESTINATION vllm/third_party/deep_gemm
-      COMPONENT _deep_gemm_C
-      FILES_MATCHING PATTERN "_C.cpython-*.so")
-  endforeach()
-  add_custom_target(_deep_gemm_C ALL DEPENDS ${_dg_markers})
+  target_compile_definitions(_deep_gemm_C PRIVATE
+    "-DTORCH_EXTENSION_NAME=_C")
+
+  target_include_directories(_deep_gemm_C PRIVATE
+    "${deepgemm_SOURCE_DIR}/csrc"
+    "${deepgemm_SOURCE_DIR}/deep_gemm/include"
+    "${deepgemm_SOURCE_DIR}/third-party/cutlass/include"
+    "${deepgemm_SOURCE_DIR}/third-party/cutlass/tools/util/include"
+    "${deepgemm_SOURCE_DIR}/third-party/fmt/include")
+
+  target_compile_options(_deep_gemm_C PRIVATE
+    $<$<COMPILE_LANGUAGE:CXX>:-std=c++17>
+    $<$<COMPILE_LANGUAGE:CXX>:-O3>
+    $<$<COMPILE_LANGUAGE:CXX>:-Wno-psabi>
+    $<$<COMPILE_LANGUAGE:CXX>:-Wno-deprecated-declarations>)
+
+  # torch_python is required because DeepGEMM uses pybind11 type casters
+  # for at::Tensor (via PYBIND11_MODULE), unlike vLLM's own extensions which
+  # use torch::Library custom ops.
+  find_library(TORCH_PYTHON_LIBRARY torch_python
+    PATHS "${TORCH_INSTALL_PREFIX}/lib"
+    REQUIRED)
+
+  target_link_libraries(_deep_gemm_C PRIVATE
+    torch ${TORCH_LIBRARIES} "${TORCH_PYTHON_LIBRARY}"
+    CUDA::cudart CUDA::nvrtc)
+
+  # Install the shared library into the vendored package directory
+  install(TARGETS _deep_gemm_C
+    LIBRARY DESTINATION vllm/third_party/deep_gemm
+    COMPONENT _deep_gemm_C)
 
   #
   # Vendor DeepGEMM Python package files

cmake/external_projects/vllm_flash_attn.cmake

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

csrc/attention/mla/sm100_cutlass_mla_kernel.cu

@@ -18,12 +18,13 @@ limitations under the License.
  * Taken from SGLANG PR https://github.com/sgl-project/sglang/pull/6929
  * by Alcanderian JieXin Liang
  */
-#include "libtorch_stable/torch_utils.h"
-
-#include <torch/csrc/stable/library.h>
+#include "core/registration.h"
 
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAGuard.h>
 #include <cutlass/cutlass.h>
 #include <cutlass/kernel_hardware_info.h>
+#include <torch/all.h>
 
 #include <cute/tensor.hpp>
 #include <iostream>
@@ -34,27 +35,27 @@ limitations under the License.
 // clang-format off
 #if !defined(CUDA_VERSION) || CUDA_VERSION < 12040
 void sm100_cutlass_mla_decode(
-    torch::stable::Tensor const& out,
-    torch::stable::Tensor const& lse,
-    torch::stable::Tensor const& q_nope,
-    torch::stable::Tensor const& q_pe,
-    torch::stable::Tensor const& kv_c_and_k_pe_cache,
-    torch::stable::Tensor const& seq_lens,
-    torch::stable::Tensor const& page_table,
-    torch::stable::Tensor const& workspace,
+    torch::Tensor const& out,
+    torch::Tensor const& lse,
+    torch::Tensor const& q_nope,
+    torch::Tensor const& q_pe,
+    torch::Tensor const& kv_c_and_k_pe_cache,
+    torch::Tensor const& seq_lens,
+    torch::Tensor const& page_table,
+    torch::Tensor const& workspace,
     double sm_scale,
     int64_t num_kv_splits) {
-  STD_TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_decode");
+  TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_decode");
 }
 int64_t sm100_cutlass_mla_get_workspace_size(int64_t max_seq_len, int64_t num_batches, int64_t sm_count, int64_t num_kv_splits) {
-  STD_TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_get_workspace_size");
+  TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_get_workspace_size");
 }
 #else
 
 #define CUTLASS_CHECK(status)                                                       \
   {                                                                                 \
     cutlass::Status error = status;                                                 \
-    STD_TORCH_CHECK(error == cutlass::Status::kSuccess, cutlassGetStatusString(error)); \
+    TORCH_CHECK(error == cutlass::Status::kSuccess, cutlassGetStatusString(error)); \
   }
 
 using namespace cute;
@@ -99,23 +100,23 @@ struct MlaSm100 {
 
 template <typename T>
 typename T::Fmha::Arguments args_from_options(
-    torch::stable::Tensor const& out,
-    torch::stable::Tensor const& lse,
-    torch::stable::Tensor const& q_nope,
-    torch::stable::Tensor const& q_pe,
-    torch::stable::Tensor const& kv_c_and_k_pe_cache,
-    torch::stable::Tensor const& seq_lens,
-    torch::stable::Tensor const& page_table,
+    at::Tensor const& out,
+    at::Tensor const& lse,
+    at::Tensor const& q_nope,
+    at::Tensor const& q_pe,
+    at::Tensor const& kv_c_and_k_pe_cache,
+    at::Tensor const& seq_lens,
+    at::Tensor const& page_table,
     double sm_scale,
     int64_t num_kv_splits) {
   cutlass::KernelHardwareInfo hw_info;
-  hw_info.device_id = q_nope.get_device_index();
+  hw_info.device_id = q_nope.device().index();
   hw_info.sm_count = cutlass::KernelHardwareInfo::query_device_multiprocessor_count(hw_info.device_id);
 
-  int batches = q_nope.size(0);
-  int page_count_per_seq = page_table.size(1);
-  int page_count_total = kv_c_and_k_pe_cache.size(0);
-  int page_size = kv_c_and_k_pe_cache.size(1);
+  int batches = q_nope.sizes()[0];
+  int page_count_per_seq = page_table.sizes()[1];
+  int page_count_total = kv_c_and_k_pe_cache.sizes()[0];
+  int page_size = kv_c_and_k_pe_cache.sizes()[1];
   int max_seq_len = page_size * page_count_per_seq;
   using TileShapeH = typename T::TileShapeH;
   using TileShapeD = typename T::TileShapeD;
@@ -185,14 +186,14 @@ typename T::Fmha::Arguments args_from_options(
 
 template <typename Element, typename ElementOut, bool IsPaged128, typename PersistenceOption>
 void runMla(
-    torch::stable::Tensor const& out,
-    torch::stable::Tensor const& lse,
-    torch::stable::Tensor const& q_nope,
-    torch::stable::Tensor const& q_pe,
-    torch::stable::Tensor const& kv_c_and_k_pe_cache,
-    torch::stable::Tensor const& seq_lens,
-    torch::stable::Tensor const& page_table,
-    torch::stable::Tensor const& workspace,
+    at::Tensor const& out,
+    at::Tensor const& lse,
+    at::Tensor const& q_nope,
+    at::Tensor const& q_pe,
+    at::Tensor const& kv_c_and_k_pe_cache,
+    at::Tensor const& seq_lens,
+    at::Tensor const& page_table,
+    at::Tensor const& workspace,
     double sm_scale,
     int64_t num_kv_splits,
     cudaStream_t stream) {
@@ -219,37 +220,37 @@ void runMla(
   }()
 
 void sm100_cutlass_mla_decode(
-    torch::stable::Tensor const& out,
-    torch::stable::Tensor const& lse,
-    torch::stable::Tensor const& q_nope,
-    torch::stable::Tensor const& q_pe,
-    torch::stable::Tensor const& kv_c_and_k_pe_cache,
-    torch::stable::Tensor const& seq_lens,
-    torch::stable::Tensor const& page_table,
-    torch::stable::Tensor const& workspace,
+    torch::Tensor const& out,
+    torch::Tensor const& lse,
+    torch::Tensor const& q_nope,
+    torch::Tensor const& q_pe,
+    torch::Tensor const& kv_c_and_k_pe_cache,
+    torch::Tensor const& seq_lens,
+    torch::Tensor const& page_table,
+    torch::Tensor const& workspace,
     double sm_scale,
     int64_t num_kv_splits) {
-  auto in_dtype = q_nope.scalar_type();
-  torch::stable::accelerator::DeviceGuard device_guard(q_nope.get_device_index());
-  const cudaStream_t stream = get_current_cuda_stream(q_nope.get_device_index());
-  const int page_size = kv_c_and_k_pe_cache.size(1);
-
+  auto in_dtype = q_nope.dtype();
+  at::cuda::CUDAGuard device_guard{(char)q_nope.get_device()};
+  const cudaStream_t stream = at::cuda::getCurrentCUDAStream(q_nope.get_device());
+  const int page_size = kv_c_and_k_pe_cache.sizes()[1];
+  
   // NOTE(alcanderian): IsPersistent has bug with manual split_kv.
   // Kernel will hang if batch is too large with large num_kv_splits. (for example bs=8, num_kv_splits=8)
   // Maybe per batch split kv will fix this.
   DISPATCH_BOOL(page_size == 128, IsPaged128, [&] {
     DISPATCH_BOOL(num_kv_splits <= 1, NotManualSplitKV, [&] {
-      if (in_dtype == torch::headeronly::ScalarType::Half) {
+      if (in_dtype == at::ScalarType::Half) {
         runMla<cutlass::half_t, cutlass::half_t, IsPaged128, IsPersistent<NotManualSplitKV>>(
           out, lse, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, workspace, sm_scale, num_kv_splits, stream);
-      } else if (in_dtype == torch::headeronly::ScalarType::BFloat16) {
+      } else if (in_dtype == at::ScalarType::BFloat16) {
         runMla<cutlass::bfloat16_t, cutlass::bfloat16_t, IsPaged128, IsPersistent<NotManualSplitKV>>(
           out, lse, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, workspace, sm_scale, num_kv_splits, stream);
-      } else if (in_dtype == torch::headeronly::ScalarType::Float8_e4m3fn) {
+      } else if (in_dtype == at::ScalarType::Float8_e4m3fn) {
         runMla<cutlass::float_e4m3_t, cutlass::bfloat16_t, IsPaged128, IsPersistent<NotManualSplitKV>>(
           out, lse, q_nope, q_pe, kv_c_and_k_pe_cache, seq_lens, page_table, workspace, sm_scale, num_kv_splits, stream);
       } else {
-        STD_TORCH_CHECK(false, "Unsupported input data type of MLA");
+        TORCH_CHECK(false, "Unsupported input data type of MLA");
       }
       return true;
     });
@@ -279,12 +280,12 @@ int64_t sm100_cutlass_mla_get_workspace_size(int64_t max_seq_len, int64_t num_ba
 
 #endif
 
-STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, m) {
-  m.impl("sm100_cutlass_mla_decode", TORCH_BOX(&sm100_cutlass_mla_decode));
+TORCH_LIBRARY_IMPL_EXPAND(TORCH_EXTENSION_NAME, CUDA, m) {
+  m.impl("sm100_cutlass_mla_decode", &sm100_cutlass_mla_decode);
 }
 
-STABLE_TORCH_LIBRARY_IMPL(_C, CompositeExplicitAutograd, m) {
-  m.impl("sm100_cutlass_mla_get_workspace_size", TORCH_BOX(&sm100_cutlass_mla_get_workspace_size));
+TORCH_LIBRARY_IMPL_EXPAND(TORCH_EXTENSION_NAME, CatchAll, m) {
+  m.impl("sm100_cutlass_mla_get_workspace_size", &sm100_cutlass_mla_get_workspace_size);
 }
 
 // clang-format on

csrc/cache.h

@@ -12,8 +12,7 @@ void swap_blocks(torch::Tensor& src, torch::Tensor& dst,
 
 void swap_blocks_batch(const torch::Tensor& src_ptrs,
                        const torch::Tensor& dst_ptrs,
-                       const torch::Tensor& sizes,
-                       bool is_src_access_order_any);
+                       const torch::Tensor& sizes);
 
 void reshape_and_cache(torch::Tensor& key, torch::Tensor& value,
                        torch::Tensor& key_cache, torch::Tensor& value_cache,

csrc/cache_kernels.cu

@@ -77,8 +77,7 @@ void swap_blocks(torch::Tensor& src, torch::Tensor& dst,
 
 void swap_blocks_batch(const torch::Tensor& src_ptrs,
                        const torch::Tensor& dst_ptrs,
-                       const torch::Tensor& sizes,
-                       bool is_src_access_order_any) {
+                       const torch::Tensor& sizes) {
   TORCH_CHECK(src_ptrs.device().is_cpu(), "src_ptrs must be on CPU");
   TORCH_CHECK(dst_ptrs.device().is_cpu(), "dst_ptrs must be on CPU");
   TORCH_CHECK(sizes.device().is_cpu(), "sizes must be on CPU");
@@ -98,13 +97,13 @@ void swap_blocks_batch(const torch::Tensor& src_ptrs,
 
   const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
 
-  // Use cuMemcpyBatchAsync / hipMemcpyBatchAsync to submit all copies in a
-  // single driver call, amortizing per-copy submission overhead. int64_t
-  // and CUdeviceptr/void*/size_t are all 8 bytes on 64-bit platforms, so we
-  // reinterpret_cast the tensor data directly to avoid copies.
+  // Use cuMemcpyBatchAsync (CUDA 12.8+) to submit all copies in a single
+  // driver call, amortizing per-copy submission overhead.
+  // int64_t and CUdeviceptr/size_t are both 8 bytes on 64-bit platforms,
+  // so we reinterpret_cast the tensor data directly to avoid copies.
+  static_assert(sizeof(CUdeviceptr) == sizeof(int64_t));
   static_assert(sizeof(size_t) == sizeof(int64_t));
 #if !defined(USE_ROCM) && defined(CUDA_VERSION) && CUDA_VERSION >= 12080
-  static_assert(sizeof(CUdeviceptr) == sizeof(int64_t));
   // Resolve cuMemcpyBatchAsync at runtime via cuGetProcAddress so that
   // binaries compiled with CUDA 12.8+ still work on older drivers, and
   // we avoid the CUDA 13.0 header remapping (#define to _v2 signature).
@@ -125,12 +124,7 @@ void swap_blocks_batch(const torch::Tensor& src_ptrs,
 
   if (batch_fn != nullptr) {
     CUmemcpyAttributes attr = {};
-    // ANY lets the DMA engine prefetch source bytes out of stream order,
-    // which is only safe when no GPU stream is concurrently writing the
-    // source.
-    attr.srcAccessOrder = is_src_access_order_any
-                              ? CU_MEMCPY_SRC_ACCESS_ORDER_ANY
-                              : CU_MEMCPY_SRC_ACCESS_ORDER_STREAM;
+    attr.srcAccessOrder = CU_MEMCPY_SRC_ACCESS_ORDER_STREAM;
     size_t attrs_idx = 0;
     size_t fail_idx = 0;
     CUresult result = batch_fn(reinterpret_cast<CUdeviceptr*>(dst_data),
@@ -140,30 +134,12 @@ void swap_blocks_batch(const torch::Tensor& src_ptrs,
                                &fail_idx, static_cast<CUstream>(stream));
     TORCH_CHECK(result == CUDA_SUCCESS, "cuMemcpyBatchAsync failed at index ",
                 fail_idx, " with error ", result);
-    return;
-  }
-#elif defined(USE_ROCM) && defined(HIP_VERSION) && HIP_VERSION >= 70100000
-  // ROCm 7.1+ exposes hipMemcpyBatchAsync. The 7.2.1 implementation early-
-  // returns hipErrorNotSupported whenever numAttrs > 0 (see ROCm/clr @
-  // rocm-7.2.1 hipamd/src/hip_memory.cpp:2819-2822), so call with
-  // numAttrs=0.
-  {
-    hipMemcpyAttributes attr = {};
-    size_t attrs_idx = 0;
-    size_t fail_idx = 0;
-    hipError_t result = hipMemcpyBatchAsync(
-        reinterpret_cast<void**>(dst_data), reinterpret_cast<void**>(src_data),
-        reinterpret_cast<size_t*>(size_data), static_cast<size_t>(n), &attr,
-        &attrs_idx, 0, &fail_idx, static_cast<hipStream_t>(stream));
-    TORCH_CHECK(result == hipSuccess, "hipMemcpyBatchAsync failed at index ",
-                fail_idx, " with error ", result);
-    return;
-  }
+  } else
 #endif
   {
-    // Fallback for CUDA < 12.8, older CUDA drivers, and ROCm < 7.1:
-    // individual async copies. cudaMemcpyDefault lets the driver infer
-    // direction from pointer types.
+    // Fallback for CUDA < 12.8, older drivers, and ROCm:
+    // individual async copies.
+    // cudaMemcpyDefault lets the driver infer direction from pointer types.
     for (int64_t i = 0; i < n; i++) {
       cudaMemcpyAsync(reinterpret_cast<void*>(dst_data[i]),
                       reinterpret_cast<void*>(src_data[i]),

csrc/cache_kernels_fused.cu

@@ -21,33 +21,28 @@ namespace vllm {
 
 // NOTE Be EXTRA careful with raw_kv_scalar_t, for __half and __nv_bfloat16 it's
 // using u16 as the backing type.
-template <typename qk_t, typename cos_sin_t, bool IS_NEOX,
-          typename raw_kv_scalar_t, typename cache_t, Fp8KVCacheDataType kv_dt>
+template <typename qk_t, bool IS_NEOX, typename raw_kv_scalar_t,
+          typename cache_t, Fp8KVCacheDataType kv_dt>
 __global__ void concat_and_cache_mla_rope_fused_kernel(
     const int64_t* __restrict__ positions,  // [num_tokens]
     qk_t* __restrict__ q_pe,        // [num_tokens, num_q_heads, rot_dim]
     qk_t* __restrict__ k_pe,        // [num_tokens, rot_dim]
     const qk_t* __restrict__ kv_c,  // [num_tokens, kv_lora_rank]
-    const cos_sin_t* __restrict__ rope_cos_sin_cache,  // [max_position, 2,
-                                                       // rot_dim // 2]
+    const qk_t* __restrict__ rope_cos_sin_cache,  // [max_position, 2,
+                                                  // rot_dim // 2]
     const int rot_dim, const int64_t q_pe_stride_token,
     const int64_t q_pe_stride_head, const int64_t k_pe_stride,
     const int64_t kv_c_stride, const int num_q_heads,
     cache_t* __restrict__ kv_cache,  // [num_blocks, block_size, (kv_lora_rank +
                                      // rot_dim)]
-    const int64_t* __restrict__ slot_mapping,  // [num_tokens]
+    const int64_t* __restrict__ kv_cache_slot_mapping,  // [num_tokens]
     const int block_stride, const int entry_stride, const int kv_lora_rank,
     const int block_size, const float* kv_cache_quant_scale) {
   // Each thread block is responsible for one token.
   const int64_t token_idx = blockIdx.x;
-  const int64_t slot_idx = slot_mapping[token_idx];
-  // NOTE: slot_idx can be -1 if the token is padded
-  if (slot_idx < 0) {
-    return;
-  }
   const int64_t pos = positions[token_idx];
 
-  const cos_sin_t* cos_sin_ptr = rope_cos_sin_cache + pos * rot_dim;
+  const qk_t* cos_sin_ptr = rope_cos_sin_cache + pos * rot_dim;
 
   const int embed_dim = rot_dim / 2;
 
@@ -59,8 +54,8 @@ __global__ void concat_and_cache_mla_rope_fused_kernel(
 
     // NOTE: Would be nice to have interleaved sin/cos so we could just load
     // both at the same time.
-    qk_t cos = static_cast<qk_t>(VLLM_LDG(cos_sin_ptr + pair_idx));
-    qk_t sin = static_cast<qk_t>(VLLM_LDG(cos_sin_ptr + pair_idx + embed_dim));
+    qk_t cos = VLLM_LDG(cos_sin_ptr + pair_idx);
+    qk_t sin = VLLM_LDG(cos_sin_ptr + pair_idx + embed_dim);
 
     qk_t* q_pe_head_ptr =
         q_pe + token_idx * q_pe_stride_token + head_idx * q_pe_stride_head;
@@ -86,15 +81,21 @@ __global__ void concat_and_cache_mla_rope_fused_kernel(
     q_pe_head_ptr[pair_idx_y] = y_dst;
   }
 
+  const int64_t slot_idx = kv_cache_slot_mapping[token_idx];
   const int64_t block_idx = slot_idx / block_size;
   const int64_t entry_idx = slot_idx % block_size;
 
+  // NOTE: slot_idx can be -1 if the token is padded
+  if (slot_idx < 0) {
+    return;
+  }
+
   // K with 1 HEAD
   for (int i = threadIdx.x; i < embed_dim; i += blockDim.x) {
     int pair_idx = i;
 
-    qk_t cos = static_cast<qk_t>(VLLM_LDG(cos_sin_ptr + pair_idx));
-    qk_t sin = static_cast<qk_t>(VLLM_LDG(cos_sin_ptr + pair_idx + embed_dim));
+    qk_t cos = VLLM_LDG(cos_sin_ptr + pair_idx);
+    qk_t sin = VLLM_LDG(cos_sin_ptr + pair_idx + embed_dim);
 
     qk_t* k_pe_head_ptr = k_pe + token_idx * k_pe_stride;
 
@@ -164,43 +165,36 @@ __global__ void concat_and_cache_mla_rope_fused_kernel(
 
 }  // namespace vllm
 
-#define CALL_CONCAT_AND_CACHE_MLA_ROPE_FUSED(RAW_KV_T, CACHE_T, KV_DTYPE)     \
-  do {                                                                        \
-    VLLM_DISPATCH_FLOATING_TYPES(q_pe.scalar_type(), "qk_scalar_type", [&] {  \
-      using qk_t = scalar_t;                                                  \
-      VLLM_DISPATCH_FLOATING_TYPES(                                           \
-          rope_cos_sin_cache.scalar_type(), "rope_cos_sin_cache_scalar_type", \
-          [&] {                                                               \
-            using cos_sin_t = scalar_t;                                       \
-            if (rope_is_neox) {                                               \
-              vllm::concat_and_cache_mla_rope_fused_kernel<                   \
-                  qk_t, cos_sin_t, true, RAW_KV_T, CACHE_T, KV_DTYPE>         \
-                  <<<grid, block, 0, stream>>>(                               \
-                      positions.data_ptr<int64_t>(), q_pe.data_ptr<qk_t>(),   \
-                      k_pe.data_ptr<qk_t>(), kv_c.data_ptr<qk_t>(),           \
-                      rope_cos_sin_cache.data_ptr<cos_sin_t>(), rot_dim,      \
-                      q_pe_stride_token, q_pe_stride_head, k_pe_stride,       \
-                      kv_c_stride, num_q_heads,                               \
-                      reinterpret_cast<CACHE_T*>(kv_cache.data_ptr()),        \
-                      slot_mapping.data_ptr<int64_t>(), block_stride,         \
-                      entry_stride, kv_lora_rank, block_size,                 \
-                      kv_cache_quant_scale.data_ptr<float>());                \
-            } else {                                                          \
-              vllm::concat_and_cache_mla_rope_fused_kernel<                   \
-                  qk_t, cos_sin_t, false, RAW_KV_T, CACHE_T, KV_DTYPE>        \
-                  <<<grid, block, 0, stream>>>(                               \
-                      positions.data_ptr<int64_t>(), q_pe.data_ptr<qk_t>(),   \
-                      k_pe.data_ptr<qk_t>(), kv_c.data_ptr<qk_t>(),           \
-                      rope_cos_sin_cache.data_ptr<cos_sin_t>(), rot_dim,      \
-                      q_pe_stride_token, q_pe_stride_head, k_pe_stride,       \
-                      kv_c_stride, num_q_heads,                               \
-                      reinterpret_cast<CACHE_T*>(kv_cache.data_ptr()),        \
-                      slot_mapping.data_ptr<int64_t>(), block_stride,         \
-                      entry_stride, kv_lora_rank, block_size,                 \
-                      kv_cache_quant_scale.data_ptr<float>());                \
-            }                                                                 \
-          });                                                                 \
-    });                                                                       \
+#define CALL_CONCAT_AND_CACHE_MLA_ROPE_FUSED(RAW_KV_T, CACHE_T, KV_DTYPE)      \
+  do {                                                                         \
+    VLLM_DISPATCH_FLOATING_TYPES(q_pe.scalar_type(), "qk_scalar_type", [&] {   \
+      using qk_t = scalar_t;                                                   \
+      if (rope_is_neox) {                                                      \
+        vllm::concat_and_cache_mla_rope_fused_kernel<qk_t, true, RAW_KV_T,     \
+                                                     CACHE_T, KV_DTYPE>        \
+            <<<grid, block, 0, stream>>>(                                      \
+                positions.data_ptr<int64_t>(), q_pe.data_ptr<qk_t>(),          \
+                k_pe.data_ptr<qk_t>(), kv_c.data_ptr<qk_t>(),                  \
+                rope_cos_sin_cache.data_ptr<qk_t>(), rot_dim,                  \
+                q_pe_stride_token, q_pe_stride_head, k_pe_stride, kv_c_stride, \
+                num_q_heads, reinterpret_cast<CACHE_T*>(kv_cache.data_ptr()),  \
+                kv_cache_slot_mapping.data_ptr<int64_t>(), block_stride,       \
+                entry_stride, kv_lora_rank, block_size,                        \
+                kv_cache_quant_scale.data_ptr<float>());                       \
+      } else {                                                                 \
+        vllm::concat_and_cache_mla_rope_fused_kernel<qk_t, false, RAW_KV_T,    \
+                                                     CACHE_T, KV_DTYPE>        \
+            <<<grid, block, 0, stream>>>(                                      \
+                positions.data_ptr<int64_t>(), q_pe.data_ptr<qk_t>(),          \
+                k_pe.data_ptr<qk_t>(), kv_c.data_ptr<qk_t>(),                  \
+                rope_cos_sin_cache.data_ptr<qk_t>(), rot_dim,                  \
+                q_pe_stride_token, q_pe_stride_head, k_pe_stride, kv_c_stride, \
+                num_q_heads, reinterpret_cast<CACHE_T*>(kv_cache.data_ptr()),  \
+                kv_cache_slot_mapping.data_ptr<int64_t>(), block_stride,       \
+                entry_stride, kv_lora_rank, block_size,                        \
+                kv_cache_quant_scale.data_ptr<float>());                       \
+      }                                                                        \
+    });                                                                        \
   } while (false)
 
 // Executes RoPE on q_pe and k_pe, then writes k_pe and kv_c in the kv cache.
@@ -214,52 +208,43 @@ void concat_and_cache_mla_rope_fused(
     torch::Tensor& kv_c,                // [num_tokens, kv_lora_rank]
     torch::Tensor& rope_cos_sin_cache,  // [max_position, rot_dim]
     bool rope_is_neox,
-    torch::Tensor& slot_mapping,  // [num_tokens] or [num_actual_tokens]
+    torch::Tensor&
+        kv_cache_slot_mapping,  // [num_tokens] or [num_actual_tokens]
     torch::Tensor&
         kv_cache,  // [num_blocks, block_size, (kv_lora_rank + rot_dim)]
     const std::string& kv_cache_dtype, torch::Tensor& kv_cache_quant_scale) {
-  // NOTE(woosuk): In vLLM V1, query/key/position.size(0) can be different from
-  // slot_mapping.size(0) because of padding for CUDA graphs.
-  // In vLLM V0, key.size(0) is always equal to slot_mapping.size(0) because
-  // both include padding.
-  // In vLLM V1, however, key.size(0) can be larger than slot_mapping.size(0)
-  // since key includes padding for CUDA graphs, while slot_mapping does not.
-  // In this case, slot_mapping.size(0) represents the actual number of tokens
-  // before padding.
-  // For compatibility with both cases, we use slot_mapping.size(0) as the
-  // number of tokens.
-  int num_tokens = slot_mapping.size(0);
-  int num_padded_tokens = q_pe.size(0);
-  TORCH_CHECK_GE(num_padded_tokens, num_tokens);
+  const int64_t num_tokens = q_pe.size(0);
 
   const int num_q_heads = q_pe.size(1);
   const int rot_dim = q_pe.size(2);
   const int kv_lora_rank = kv_c.size(1);
 
-  TORCH_CHECK_EQ(positions.size(0), num_padded_tokens);
+  TORCH_CHECK(positions.size(0) >=
+              num_tokens);  // CUDA Graphs might pad this for us
   TORCH_CHECK_EQ(positions.dim(), 1);
   TORCH_CHECK_EQ(positions.scalar_type(), c10::ScalarType::Long);
 
-  TORCH_CHECK_EQ(q_pe.dim(), 3);
-  TORCH_CHECK_EQ(q_pe.size(0), num_padded_tokens);
+  TORCH_CHECK_EQ(q_pe.size(0), num_tokens);
   TORCH_CHECK_EQ(q_pe.size(1), num_q_heads);
   TORCH_CHECK_EQ(q_pe.size(2), rot_dim);
+  TORCH_CHECK_EQ(q_pe.dim(), 3);
 
-  TORCH_CHECK_EQ(k_pe.dim(), 2);
-  TORCH_CHECK_EQ(k_pe.size(0), num_padded_tokens);
+  TORCH_CHECK_EQ(k_pe.size(0), num_tokens);
   TORCH_CHECK_EQ(k_pe.size(1), rot_dim);
+  TORCH_CHECK_EQ(k_pe.dim(), 2);
   TORCH_CHECK_EQ(k_pe.scalar_type(), q_pe.scalar_type());
 
-  TORCH_CHECK_EQ(kv_c.dim(), 2);
-  TORCH_CHECK_EQ(kv_c.size(0), num_padded_tokens);
+  TORCH_CHECK_EQ(kv_c.size(0), num_tokens);
   TORCH_CHECK_EQ(kv_c.size(1), kv_lora_rank);
+  TORCH_CHECK_EQ(kv_c.dim(), 2);
   TORCH_CHECK_EQ(kv_c.scalar_type(), q_pe.scalar_type());
   TORCH_CHECK_EQ(kv_c.dtype(), q_pe.dtype());
 
   TORCH_CHECK_EQ(rope_cos_sin_cache.size(1), rot_dim);
+  TORCH_CHECK_EQ(rope_cos_sin_cache.scalar_type(), q_pe.scalar_type());
 
-  TORCH_CHECK_EQ(slot_mapping.size(0), num_tokens);
-  TORCH_CHECK_EQ(slot_mapping.scalar_type(), c10::ScalarType::Long);
+  TORCH_CHECK_EQ(kv_cache_slot_mapping.size(0), num_tokens);
+  TORCH_CHECK_EQ(kv_cache_slot_mapping.scalar_type(), c10::ScalarType::Long);
 
   TORCH_CHECK_EQ(kv_cache.size(2), kv_lora_rank + rot_dim);
   TORCH_CHECK_EQ(kv_cache.dim(), 3);

csrc/core/batch_invariant.hpp

@@ -1,6 +1,7 @@
 #pragma once
 #include <cstdlib>
 #include <string>
+#include <cctype>
 
 namespace vllm {
 

csrc/core/scalar_type.hpp

@@ -1,13 +1,7 @@
 #pragma once
 
-#include <cstdint>
-#include <string>
-#include <tuple>
-#include <utility>
-#include <variant>
-
-// For STD_TORCH_CHECK
-#include <torch/headeronly/util/Exception.h>
+// For TORCH_CHECK
+#include <torch/library.h>
 
 namespace vllm {
 
@@ -51,7 +45,7 @@ class ScalarType {
   // IEEE 754 compliant floating point type
   static constexpr ScalarType float_IEEE754(uint8_t exponent,
                                             uint8_t mantissa) {
-    STD_TORCH_CHECK(mantissa > 0 && exponent > 0);
+    TORCH_CHECK(mantissa > 0 && exponent > 0);
     return ScalarType(exponent, mantissa, true, 0, false, NAN_IEEE_754);
   }
 
@@ -59,12 +53,11 @@ class ScalarType {
   static constexpr ScalarType float_(uint8_t exponent, uint8_t mantissa,
                                      bool finite_values_only,
                                      NanRepr nan_repr) {
-    STD_TORCH_CHECK(nan_repr < NAN_REPR_ID_MAX, "Invalid NanRepr");
-    STD_TORCH_CHECK(mantissa > 0 && exponent > 0);
-    STD_TORCH_CHECK(
-        nan_repr != NAN_IEEE_754,
-        "use `float_IEEE754` constructor for floating point types that "
-        "follow IEEE 754 conventions");
+    TORCH_CHECK(nan_repr < NAN_REPR_ID_MAX, "Invalid NanRepr");
+    TORCH_CHECK(mantissa > 0 && exponent > 0);
+    TORCH_CHECK(nan_repr != NAN_IEEE_754,
+                "use `float_IEEE754` constructor for floating point types that "
+                "follow IEEE 754 conventions");
     return ScalarType(exponent, mantissa, true, 0, finite_values_only,
                       nan_repr);
   }
@@ -183,8 +176,8 @@ class ScalarType {
 
  private:
   double _floating_point_max() const {
-    STD_TORCH_CHECK(mantissa <= 52 && exponent <= 11,
-                    "Cannot represent max/min as a double for type ", str());
+    TORCH_CHECK(mantissa <= 52 && exponent <= 11,
+                "Cannot represent max/min as a double for type ", str());
 
     uint64_t max_mantissa = (uint64_t(1) << mantissa) - 1;
     if (nan_repr == NAN_EXTD_RANGE_MAX_MIN) {
@@ -193,8 +186,8 @@ class ScalarType {
 
     uint64_t max_exponent = (uint64_t(1) << exponent) - 2;
     if (nan_repr == NAN_EXTD_RANGE_MAX_MIN || nan_repr == NAN_NONE) {
-      STD_TORCH_CHECK(exponent < 11,
-                      "Cannot represent max/min as a double for type ", str());
+      TORCH_CHECK(exponent < 11,
+                  "Cannot represent max/min as a double for type ", str());
       max_exponent += 1;
     }
 
@@ -223,17 +216,16 @@ class ScalarType {
     if (is_floating_point()) {
       return {_floating_point_max()};
     } else {
-      STD_TORCH_CHECK(size_bits() < 64 || size_bits() == 64 && is_signed(),
-                      "Cannot represent max as a int64_t");
+      TORCH_CHECK(size_bits() < 64 || size_bits() == 64 && is_signed(),
+                  "Cannot represent max as a int64_t");
       return {(int64_t(1) << mantissa) - 1};
     }
   }
 
   constexpr std::variant<int64_t, double> _raw_min() const {
     if (is_floating_point()) {
-      STD_TORCH_CHECK(
-          is_signed(),
-          "We currently assume all floating point types are signed");
+      TORCH_CHECK(is_signed(),
+                  "We currently assume all floating point types are signed");
       constexpr uint64_t sign_bit_double = (uint64_t(1) << 63);
 
       double max = _floating_point_max();
@@ -241,8 +233,8 @@ class ScalarType {
       uint64_t min_raw = max_raw | sign_bit_double;
       return {*reinterpret_cast<double*>(&min_raw)};
     } else {
-      STD_TORCH_CHECK(!is_signed() || size_bits() <= 64,
-                      "Cannot represent min as a int64_t");
+      TORCH_CHECK(!is_signed() || size_bits() <= 64,
+                  "Cannot represent min as a int64_t");
       if (is_signed()) {
         // set the top bit to 1 (i.e. INT64_MIN) and the rest to 0
         // then perform an arithmetic shift right to set all the bits above

csrc/cpu/cpu_arch_macros.h

@@ -156,20 +156,4 @@
 
 #endif  // __aarch64__
 
-// RISC-V RVV
-#ifdef __riscv_v
-  #include <riscv_vector.h>
-
-  #ifdef __riscv_zihintpause
-    #define FAST_SPINNING __riscv_pause();
-  #endif
-
-  // FP32Vec16::exp() in cpu_types_riscv.hpp already implements the full
-  // polynomial approximation for RVV, so we simply delegate to it.
-  #define DEFINE_FAST_EXP                             \
-    auto fast_exp = [&](const vec_op::FP32Vec16& vec) \
-                        __attribute__((always_inline)) { return vec.exp(); };
-
-#endif  // __riscv_v
-
 #endif

csrc/cpu/cpu_attn.cpp

@@ -1,16 +1,5 @@
 #include "cpu_attn_dispatch_generated.h"
 
-// Maps kv_cache_dtype string to Fp8KVCacheDataType enum.
-// "auto" -> kAuto(0); "fp8"/"fp8_e4m3" -> kFp8E4M3; "fp8_e5m2" -> kFp8E5M2.
-static inline cpu_attention::Fp8KVCacheDataType parse_fp8_kv_dtype(
-    const std::string& kv_cache_dtype) {
-  if (kv_cache_dtype == "fp8_e5m2")
-    return cpu_attention::Fp8KVCacheDataType::kFp8E5M2;
-  if (kv_cache_dtype == "fp8_e4m3" || kv_cache_dtype == "fp8")
-    return cpu_attention::Fp8KVCacheDataType::kFp8E4M3;
-  return cpu_attention::Fp8KVCacheDataType::kAuto;
-}
-
 torch::Tensor get_scheduler_metadata(
     const int64_t num_req, const int64_t num_heads_q,
     const int64_t num_heads_kv, const int64_t head_dim,
@@ -29,10 +18,6 @@ torch::Tensor get_scheduler_metadata(
     isa = cpu_attention::ISA::NEON;
   } else if (isa_hint == "vxe") {
     isa = cpu_attention::ISA::VXE;
-  } else if (isa_hint == "rvv") {
-    isa = cpu_attention::ISA::RVV;
-  } else if (isa_hint == "vsx") {
-    isa = cpu_attention::ISA::VSX;
   } else {
     TORCH_CHECK(false, "Unsupported CPU attention ISA hint: " + isa_hint);
   }
@@ -64,7 +49,7 @@ torch::Tensor get_scheduler_metadata(
   input.enable_kv_split = enable_kv_split;
 
   VLLM_DISPATCH_FLOATING_TYPES(dtype, "get_scheduler_metadata", [&]() {
-    CPU_ATTN_DISPATCH(head_dim, isa, 0, [&]() {
+    CPU_ATTN_DISPATCH(head_dim, isa, [&]() {
       input.elem_size = sizeof(scalar_t);
       input.q_buffer_elem_size = sizeof(attn_impl::q_buffer_t);
       input.logits_buffer_elem_size = sizeof(attn_impl::logits_buffer_t);
@@ -87,9 +72,7 @@ void cpu_attn_reshape_and_cache(
         key_cache,  // [num_blocks, num_kv_heads, block_size, head_size]
     torch::Tensor&
         value_cache,  // [num_blocks, num_kv_heads, block_size, head_size]
-    const torch::Tensor& slot_mapping, const std::string& isa,
-    const double k_scale = 1.0, const double v_scale = 1.0,
-    const std::string& kv_cache_dtype = "auto") {
+    const torch::Tensor& slot_mapping, const std::string& isa) {
   TORCH_CHECK_EQ(key.dim(), 3);
   TORCH_CHECK_EQ(value.dim(), 3);
   TORCH_CHECK_EQ(key_cache.dim(), 4);
@@ -97,30 +80,18 @@ void cpu_attn_reshape_and_cache(
   TORCH_CHECK_EQ(key.stride(2), 1);
   TORCH_CHECK_EQ(value.stride(2), 1);
 
-  const int64_t kv_cache_idx =
-      static_cast<int64_t>(parse_fp8_kv_dtype(kv_cache_dtype));
-  const bool is_fp8 = (kv_cache_idx != 0);
-
-  if (is_fp8) {
-    TORCH_CHECK(key_cache.scalar_type() == at::ScalarType::Byte,
-                "key_cache must be uint8 for FP8 path");
-    TORCH_CHECK(value_cache.scalar_type() == at::ScalarType::Byte,
-                "value_cache must be uint8 for FP8 path");
-    TORCH_CHECK(k_scale > 0, "k_scale must be positive for FP8 path");
-    TORCH_CHECK(v_scale > 0, "v_scale must be positive for FP8 path");
-  }
-
-  const float k_inv = is_fp8 ? 1.0f / static_cast<float>(k_scale) : 0.0f;
-  const float v_inv = is_fp8 ? 1.0f / static_cast<float>(v_scale) : 0.0f;
-
   const int64_t token_num = key.size(0);
-  const int64_t head_num = key.size(1);
-  const int64_t head_dim = key.size(2);
+  const int64_t key_token_num_stride = key.stride(0);
+  const int64_t value_token_num_stride = value.stride(0);
+  const int64_t head_num = value.size(1);
+  const int64_t key_head_num_stride = key.stride(1);
+  const int64_t value_head_num_stride = value.stride(1);
   const int64_t num_blocks = key_cache.size(0);
   const int64_t num_blocks_stride = key_cache.stride(0);
   const int64_t cache_head_num_stride = key_cache.stride(1);
   const int64_t block_size = key_cache.size(2);
   const int64_t block_size_stride = key_cache.stride(2);
+  const int64_t head_dim = key.size(-1);
 
   cpu_attention::ISA isa_tag = [&]() {
     if (isa == "amx") {
@@ -133,33 +104,21 @@ void cpu_attn_reshape_and_cache(
       return cpu_attention::ISA::NEON;
     } else if (isa == "vxe") {
       return cpu_attention::ISA::VXE;
-    } else if (isa == "rvv") {
-      return cpu_attention::ISA::RVV;
-    } else if (isa == "vsx") {
-      return cpu_attention::ISA::VSX;
     } else {
       TORCH_CHECK(false, "Invalid ISA type: " + isa);
     }
   }();
 
-  if (is_fp8) {
-    TORCH_CHECK(isa_tag == cpu_attention::ISA::AMX ||
-                    isa_tag == cpu_attention::ISA::VEC,
-                "FP8 KV cache is only supported on x86 (AMX/VEC) ISA");
-  }
-
   VLLM_DISPATCH_FLOATING_TYPES(
       key.scalar_type(), "cpu_attn_reshape_and_cache", [&]() {
-        CPU_ATTN_DISPATCH(head_dim, isa_tag, kv_cache_idx, [&]() {
-          using kv_t = typename attn_impl::kv_cache_t;
+        CPU_ATTN_DISPATCH(head_dim, isa_tag, [&]() {
           attn_impl::reshape_and_cache(
               key.data_ptr<scalar_t>(), value.data_ptr<scalar_t>(),
-              reinterpret_cast<kv_t*>(key_cache.data_ptr()),
-              reinterpret_cast<kv_t*>(value_cache.data_ptr()),
-              slot_mapping.data_ptr<int64_t>(), token_num, key.stride(0),
-              value.stride(0), head_num, key.stride(1), value.stride(1),
-              num_blocks, num_blocks_stride, cache_head_num_stride, block_size,
-              block_size_stride, k_inv, v_inv);
+              key_cache.data_ptr<scalar_t>(), value_cache.data_ptr<scalar_t>(),
+              slot_mapping.data_ptr<int64_t>(), token_num, key_token_num_stride,
+              value_token_num_stride, head_num, key_head_num_stride,
+              value_head_num_stride, num_blocks, num_blocks_stride,
+              cache_head_num_stride, block_size, block_size_stride);
         });
       });
 }
@@ -178,26 +137,13 @@ void cpu_attention_with_kv_cache(
     const int64_t sliding_window_left, const int64_t sliding_window_right,
     const torch::Tensor& block_table,  // [num_tokens, max_block_num]
     const double softcap, const torch::Tensor& scheduler_metadata,
-    const std::optional<torch::Tensor>& s_aux,  // [num_heads]
-    const double k_scale = 1.0, const double v_scale = 1.0,
-    const std::string& kv_cache_dtype = "auto") {
+    const std::optional<torch::Tensor>& s_aux  // [num_heads]
+) {
   TORCH_CHECK_EQ(query.dim(), 3);
   TORCH_CHECK_EQ(query.stride(2), 1);
   TORCH_CHECK_EQ(key_cache.dim(), 4);
   TORCH_CHECK_EQ(value_cache.dim(), 4);
 
-  const int64_t kv_cache_idx =
-      static_cast<int64_t>(parse_fp8_kv_dtype(kv_cache_dtype));
-  const bool is_fp8 = (kv_cache_idx != 0);
-  if (is_fp8) {
-    TORCH_CHECK(key_cache.scalar_type() == at::ScalarType::Byte,
-                "key_cache must be uint8 for FP8 path");
-    TORCH_CHECK(value_cache.scalar_type() == at::ScalarType::Byte,
-                "value_cache must be uint8 for FP8 path");
-    TORCH_CHECK(k_scale > 0, "k_scale must be positive for FP8 path");
-    TORCH_CHECK(v_scale > 0, "v_scale must be positive for FP8 path");
-  }
-
   cpu_attention::AttentionInput input;
   input.metadata = reinterpret_cast<cpu_attention::AttentionMetadata*>(
       scheduler_metadata.data_ptr());
@@ -219,32 +165,25 @@ void cpu_attention_with_kv_cache(
   input.block_table = block_table.data_ptr<int32_t>();
   input.alibi_slopes =
       alibi_slopes.has_value() ? alibi_slopes->data_ptr<float>() : nullptr;
+  // For now sink must be bf16
   input.s_aux = s_aux.has_value() ? s_aux->data_ptr<c10::BFloat16>() : nullptr;
   input.scale = scale;
   input.causal = causal;
   input.sliding_window_left = sliding_window_left;
   input.sliding_window_right = sliding_window_right;
   if (input.causal) {
+    // to make boundary calculation easier
     input.sliding_window_right = 0;
   }
-  input.softcap = static_cast<float>(softcap);
-
-  if (is_fp8) {
-    input.k_scale_fp8 = static_cast<float>(k_scale);
-    input.v_scale_fp8 = static_cast<float>(v_scale);
-    TORCH_CHECK(input.metadata->isa == cpu_attention::ISA::AMX ||
-                    input.metadata->isa == cpu_attention::ISA::VEC,
-                "FP8 KV cache is only supported on x86 (AMX/VEC) ISA");
-  }
+  float softcap_fp32 = softcap;
+  input.softcap = softcap_fp32;
 
   VLLM_DISPATCH_FLOATING_TYPES(
       query.scalar_type(), "cpu_attention_with_kv_cache", [&]() {
-        CPU_ATTN_DISPATCH(
-            query.size(2), input.metadata->isa, kv_cache_idx, [&]() {
-              TORCH_CHECK_EQ(input.block_size % attn_impl::BlockSizeAlignment,
-                             0);
-              cpu_attention::AttentionMainLoop<attn_impl> mainloop;
-              mainloop(&input);
-            });
+        CPU_ATTN_DISPATCH(query.size(2), input.metadata->isa, [&]() {
+          TORCH_CHECK_EQ(input.block_size % attn_impl::BlockSizeAlignment, 0);
+          cpu_attention::AttentionMainLoop<attn_impl> mainloop;
+          mainloop(&input);
+        });
       });
 }

csrc/cpu/cpu_attn_amx.hpp

@@ -1,7 +1,6 @@
 #ifndef CPU_ATTN_AMX_HPP
 #define CPU_ATTN_AMX_HPP
 
-#include "cpu_attn_fp8.hpp"
 #include "cpu_attn_impl.hpp"
 
 namespace cpu_attention {
@@ -22,10 +21,9 @@ typedef struct __tile_config {
 // 2-2-4 pattern, for 16 < m <= 32
 // TILE 0, 1: load A matrix, row num should be 16, m - 16
 // TILE 2, 3: load B matrix, row num should be 16
-// TILE 4, 5, 6, 7: store results C matrix, row num should be 16, 16,
-// m - 16, m - 16
-// q_buffer_t: A (Q/P) tile type; kv_cache_t: B (K/V cache) tile type.
-template <typename q_buffer_t, typename kv_cache_t>
+// TILE 4, 5, 6, 7: store results C matrix, row num should be 16, 16, m - 16, m
+// - 16
+template <typename kv_cache_t>
 class TileGemm224 {
  public:
   template <AttentionGemmPhase phase, int32_t k_size>
@@ -44,56 +42,13 @@ class TileGemm224 {
   }
 };
 
-// Dequantize one FP8 tile (AMX_TILE_ROW_NUM rows x 32 cols) to BF16.
-template <typename kv_cache_t>
-FORCE_INLINE void deq_tile_amx(const uint8_t* src, c10::BFloat16* dst) {
-  for (int r = 0; r < AMX_TILE_ROW_NUM; ++r) {
-    if constexpr (std::is_same_v<kv_cache_t, c10::Float8_e4m3fn>) {
-      vec_op::BF16Vec32(src + r * 32, vec_op::fp8_bf16_e4m3_tag{})
-          .save(dst + r * 32);
-    } else {
-      vec_op::BF16Vec32(src + r * 32, vec_op::fp8_bf16_e5m2_tag{})
-          .save(dst + r * 32);
-    }
-  }
-}
-
-// For FP8: dequant src into scratch and return scratch.
-// For BF16: return src directly (scratch is unused; the compiler elides it).
-template <typename kv_cache_t>
-FORCE_INLINE const c10::BFloat16* prepare_b_tile(const kv_cache_t* src,
-                                                 c10::BFloat16* scratch) {
-  if constexpr (std::is_same_v<kv_cache_t, c10::Float8_e4m3fn> ||
-                std::is_same_v<kv_cache_t, c10::Float8_e5m2>) {
-    deq_tile_amx<kv_cache_t>(reinterpret_cast<const uint8_t*>(src), scratch);
-    return scratch;
-  } else {
-    return reinterpret_cast<const c10::BFloat16*>(src);
-  }
-}
-
-// Handles both BF16 and FP8 KV cache (2-2-4 pattern).
-template <typename kv_cache_t>
-class TileGemm224<c10::BFloat16, kv_cache_t> {
-  static_assert(std::is_same_v<kv_cache_t, c10::BFloat16> ||
-                    std::is_same_v<kv_cache_t, c10::Float8_e4m3fn> ||
-                    std::is_same_v<kv_cache_t, c10::Float8_e5m2>,
-                "kv_cache_t must be BFloat16, Float8_e4m3fn, or Float8_e5m2");
-
-  static constexpr bool fp8_kv =
-      std::is_same_v<kv_cache_t, c10::Float8_e4m3fn> ||
-      std::is_same_v<kv_cache_t, c10::Float8_e5m2>;
-
-  static constexpr int64_t tile_elems = AMX_TILE_BYTES / sizeof(c10::BFloat16);
-  // BF16 path: scratch_elems=1 so the scratch array is eliminated by the
-  // compiler.
-  static constexpr int64_t scratch_elems = fp8_kv ? tile_elems : 1;
-
+template <>
+class TileGemm224<c10::BFloat16> {
  public:
   template <AttentionGemmPhase phase, int32_t k_size>
   FORCE_INLINE static void gemm(const int32_t m_size,
                                 c10::BFloat16* __restrict__ a_tile,
-                                kv_cache_t* __restrict__ b_tile,
+                                c10::BFloat16* __restrict__ b_tile,
                                 float* __restrict__ c_tile, const int64_t lda,
                                 const int64_t ldb, const int64_t ldc,
                                 const int32_t block_size,
@@ -101,7 +56,6 @@ class TileGemm224<c10::BFloat16, kv_cache_t> {
                                 const bool accum_c) {
     const int32_t k_times =
         dynamic_k_size / (AMX_TILE_ROW_NUM * 4 / sizeof(c10::BFloat16));
-
     c10::BFloat16* __restrict__ a_tile_0 = a_tile;
     c10::BFloat16* __restrict__ a_tile_1 = a_tile + lda * AMX_TILE_ROW_NUM;
     const int64_t a_tile_stride = [&]() {
@@ -116,8 +70,8 @@ class TileGemm224<c10::BFloat16, kv_cache_t> {
       }
     }();
 
-    kv_cache_t* __restrict__ b_tile_2 = b_tile;
-    kv_cache_t* __restrict__ b_tile_3 = [&]() {
+    c10::BFloat16* __restrict__ b_tile_2 = b_tile;
+    c10::BFloat16* __restrict__ b_tile_3 = [&]() {
       if constexpr (phase == AttentionGemmPhase::QK) {
         // k_cache is prepacked
         return b_tile + (k_size * AMX_TILE_ROW_BYTES / 4);
@@ -152,16 +106,11 @@ class TileGemm224<c10::BFloat16, kv_cache_t> {
       _tile_zero(7);
     }
 
-    alignas(64) c10::BFloat16 scratch_2[scratch_elems];
-    alignas(64) c10::BFloat16 scratch_3[scratch_elems];
     for (int32_t k = 0; k < k_times; ++k) {
-      const c10::BFloat16* load_2 = prepare_b_tile(b_tile_2, scratch_2);
-      const c10::BFloat16* load_3 = prepare_b_tile(b_tile_3, scratch_3);
-
       _tile_loadd(0, a_tile_0, a_tile_stride);
-      _tile_stream_loadd(2, const_cast<c10::BFloat16*>(load_2), b_tile_stride);
+      _tile_stream_loadd(2, b_tile_2, b_tile_stride);
       _tile_dpbf16ps(4, 0, 2);
-      _tile_stream_loadd(3, const_cast<c10::BFloat16*>(load_3), b_tile_stride);
+      _tile_stream_loadd(3, b_tile_3, b_tile_stride);
       _tile_dpbf16ps(5, 0, 3);
       _tile_loadd(1, a_tile_1, a_tile_stride);
       _tile_dpbf16ps(6, 1, 2);
@@ -205,13 +154,13 @@ class TileGemm224<c10::BFloat16, kv_cache_t> {
 };
 
 // 1-2-2 pattern, for 0 < m <= 16
-// TILE 0, (1): load A matrix, use extra 1 tile for prefetch, row num should
-// be m, m
-// TILE 2, 3, (4, 5): load B matrix, use extra 2 tiles for prefetch, row num
-// should be 16
-// TILE 6, 7: store results C matrix, row num should be m
-// q_buffer_t: A (Q/P) tile type; kv_cache_t: B (K/V cache) tile type.
-template <typename q_buffer_t, typename kv_cache_t>
+// TILE 0, (1): load A matrix, use extra 1 tile for prefetch, row num should be
+// m, m
+// TILE 2, 3, (4, 5): load B matrix, use extra 2 tiles for prefetch, row
+// num should be 16
+// TILE 6, 7, (6, 7): store results C matrix, row num should be
+// m
+template <typename kv_cache_t>
 class TileGemm122 {
  public:
   template <AttentionGemmPhase phase, int32_t k_size>
@@ -230,26 +179,13 @@ class TileGemm122 {
   }
 };
 
-// Handles both BF16 and FP8 KV cache (1-2-2 pattern).
-template <typename kv_cache_t>
-class TileGemm122<c10::BFloat16, kv_cache_t> {
-  static_assert(std::is_same_v<kv_cache_t, c10::BFloat16> ||
-                    std::is_same_v<kv_cache_t, c10::Float8_e4m3fn> ||
-                    std::is_same_v<kv_cache_t, c10::Float8_e5m2>,
-                "kv_cache_t must be BFloat16, Float8_e4m3fn, or Float8_e5m2");
-
-  static constexpr bool fp8_kv =
-      std::is_same_v<kv_cache_t, c10::Float8_e4m3fn> ||
-      std::is_same_v<kv_cache_t, c10::Float8_e5m2>;
-
-  static constexpr int64_t tile_elems = AMX_TILE_BYTES / sizeof(c10::BFloat16);
-  static constexpr int64_t scratch_elems = fp8_kv ? tile_elems : 1;
-
+template <>
+class TileGemm122<c10::BFloat16> {
  public:
   template <AttentionGemmPhase phase, int32_t k_size>
   FORCE_INLINE static void gemm(const int32_t m_size,
                                 c10::BFloat16* __restrict__ a_tile,
-                                kv_cache_t* __restrict__ b_tile,
+                                c10::BFloat16* __restrict__ b_tile,
                                 float* __restrict__ c_tile, const int64_t lda,
                                 const int64_t ldb, const int64_t ldc,
                                 const int32_t block_size,
@@ -279,19 +215,21 @@ class TileGemm122<c10::BFloat16, kv_cache_t> {
       }
     }();
 
-    kv_cache_t* __restrict__ b_tile_2 = b_tile;
-    kv_cache_t* __restrict__ b_tile_3 = [&]() {
+    c10::BFloat16* __restrict__ b_tile_2 = b_tile;
+    c10::BFloat16* __restrict__ b_tile_3 = [&]() {
       if constexpr (phase == AttentionGemmPhase::QK) {
+        // k_cache is prepacked
         return b_tile + (k_size * AMX_TILE_ROW_BYTES / 4);
       } else if constexpr (phase == AttentionGemmPhase::PV) {
+        // v_cache is prepacked
         return b_tile + (block_size * AMX_TILE_ROW_BYTES / 4);
       } else {
         TORCH_CHECK(false, "Unreachable");
       }
     }();
-    kv_cache_t* __restrict__ b_tile_4 =
+    c10::BFloat16* __restrict__ b_tile_4 =
         b_tile_2 + AMX_TILE_BYTES / sizeof(c10::BFloat16);
-    kv_cache_t* __restrict__ b_tile_5 =
+    c10::BFloat16* __restrict__ b_tile_5 =
         b_tile_3 + AMX_TILE_BYTES / sizeof(c10::BFloat16);
     int64_t b_stride = AMX_TILE_ROW_BYTES;
 
@@ -312,25 +250,16 @@ class TileGemm122<c10::BFloat16, kv_cache_t> {
       _tile_zero(7);
     }
 
-    alignas(64) c10::BFloat16 scratch_2[scratch_elems];
-    alignas(64) c10::BFloat16 scratch_3[scratch_elems];
-    alignas(64) c10::BFloat16 scratch_4[scratch_elems];
-    alignas(64) c10::BFloat16 scratch_5[scratch_elems];
     for (int32_t k = 0; k < k_group_times; ++k) {
-      const c10::BFloat16* load_2 = prepare_b_tile(b_tile_2, scratch_2);
-      const c10::BFloat16* load_3 = prepare_b_tile(b_tile_3, scratch_3);
-      const c10::BFloat16* load_4 = prepare_b_tile(b_tile_4, scratch_4);
-      const c10::BFloat16* load_5 = prepare_b_tile(b_tile_5, scratch_5);
-
       _tile_loadd(0, a_tile_0, a_tile_stride);
-      _tile_stream_loadd(2, const_cast<c10::BFloat16*>(load_2), b_stride);
+      _tile_stream_loadd(2, b_tile_2, b_stride);
       _tile_dpbf16ps(6, 0, 2);
-      _tile_stream_loadd(3, const_cast<c10::BFloat16*>(load_3), b_stride);
+      _tile_stream_loadd(3, b_tile_3, b_stride);
       _tile_dpbf16ps(7, 0, 3);
       _tile_loadd(1, a_tile_1, a_tile_stride);
-      _tile_stream_loadd(4, const_cast<c10::BFloat16*>(load_4), b_stride);
+      _tile_stream_loadd(4, b_tile_4, b_stride);
       _tile_dpbf16ps(6, 1, 4);
-      _tile_stream_loadd(5, const_cast<c10::BFloat16*>(load_5), b_stride);
+      _tile_stream_loadd(5, b_tile_5, b_stride);
       _tile_dpbf16ps(7, 1, 5);
 
       // update ptrs
@@ -350,13 +279,10 @@ class TileGemm122<c10::BFloat16, kv_cache_t> {
     }
 
     if (has_tail) {
-      const c10::BFloat16* load_2 = prepare_b_tile(b_tile_2, scratch_2);
-      const c10::BFloat16* load_3 = prepare_b_tile(b_tile_3, scratch_3);
-
       _tile_loadd(0, a_tile_0, a_tile_stride);
-      _tile_stream_loadd(2, const_cast<c10::BFloat16*>(load_2), b_stride);
+      _tile_stream_loadd(2, b_tile_2, b_stride);
       _tile_dpbf16ps(6, 0, 2);
-      _tile_stream_loadd(3, const_cast<c10::BFloat16*>(load_3), b_stride);
+      _tile_stream_loadd(3, b_tile_3, b_stride);
       _tile_dpbf16ps(7, 0, 3);
     }
 
@@ -376,25 +302,21 @@ class TileGemm122<c10::BFloat16, kv_cache_t> {
     _tile_loadconfig(&config);
   }
 };
-
 }  // namespace
 
-template <typename scalar_t, int64_t head_dim, typename kv_cache_scalar_t>
-class AttentionImpl<ISA::AMX, scalar_t, head_dim, kv_cache_scalar_t> {
-  static constexpr bool fp8_kv =
-      std::is_same_v<kv_cache_scalar_t, c10::Float8_e4m3fn> ||
-      std::is_same_v<kv_cache_scalar_t, c10::Float8_e5m2>;
-
+template <typename scalar_t, int64_t head_dim>
+class AttentionImpl<ISA::AMX, scalar_t, head_dim> {
  public:
   using query_t = scalar_t;
   using q_buffer_t = scalar_t;
-  using kv_cache_t = kv_cache_scalar_t;
+  using kv_cache_t = scalar_t;
   using logits_buffer_t = float;
   using partial_output_buffer_t = float;
   using prob_buffer_t = scalar_t;
 
   constexpr static int64_t BlockSizeAlignment =
-      32;  // AMX_TILE_ROW_NUM = 16 tokens/tile; 32 = 2 tiles
+      AMX_TILE_ROW_BYTES /
+      sizeof(kv_cache_t);  // KV token num unit of QK and PV phases
   constexpr static int64_t HeadDimAlignment =
       2 * (AMX_TILE_ROW_BYTES / 4);  // headdim num unit of PV phase
   constexpr static int64_t MaxQHeadNumPerIteration = 32;
@@ -402,9 +324,6 @@ class AttentionImpl<ISA::AMX, scalar_t, head_dim, kv_cache_scalar_t> {
   constexpr static ISA ISAType = ISA::AMX;
   constexpr static bool scale_on_logits = true;
 
-  float k_scale = 1.0f;
-  float v_scale = 1.0f;
-
  public:
   AttentionImpl() : current_q_head_num_(0) {
     // Use all columns in AMX tiles
@@ -413,50 +332,21 @@ class AttentionImpl<ISA::AMX, scalar_t, head_dim, kv_cache_scalar_t> {
 
   ~AttentionImpl() { _tile_release(); }
 
-  void init_from_input(const AttentionInput* input) {
-    if constexpr (fp8_kv) {
-      k_scale = input->k_scale_fp8;
-      v_scale = input->v_scale_fp8;
-    }
-  }
-
-  float get_output_v_scale() const noexcept {
-    if constexpr (fp8_kv) {
-      // AMX dequant places FP8 payload into a BF16 field (exponent bias 127).
-      // Correction = 2^(127 - FP8_bias): E4M3 bias=7 → 2^120, E5M2 bias=15 →
-      // 2^112.
-      constexpr float bias =
-          std::is_same_v<kv_cache_t, c10::Float8_e5m2> ? 0x1p112f : 0x1p120f;
-      return v_scale * bias;
-    }
-    return 1.0f;
-  }
-
   template <template <typename tile_gemm_t> typename attention>
   FORCE_INLINE void execute_attention(DEFINE_CPU_ATTENTION_PARAMS) {
-    if constexpr (fp8_kv) {
-      // Same bias correction as get_output_v_scale: AMX FP8→BF16 dequant
-      // shifts the exponent bias from FP8 to BF16 (127), so we multiply by
-      // 2^(127-FP8_bias) to recover the true value. E4M3: 2^120, E5M2: 2^112.
-      const float bias =
-          std::is_same_v<kv_cache_t, c10::Float8_e5m2> ? 0x1p112f : 0x1p120f;
-      scale *= k_scale * bias;
-    }
     if (q_head_num > AMX_TILE_ROW_NUM) {
       if (q_head_num != current_q_head_num_) {
         current_q_head_num_ = q_head_num;
-        TileGemm224<q_buffer_t, kv_cache_t>::init_tile_config(q_head_num,
-                                                              amx_tile_config_);
+        TileGemm224<kv_cache_t>::init_tile_config(q_head_num, amx_tile_config_);
       }
-      attention<TileGemm224<q_buffer_t, kv_cache_t>> attention_iteration;
+      attention<TileGemm224<kv_cache_t>> attention_iteration;
       attention_iteration(CPU_ATTENTION_PARAMS);
     } else {
       if (q_head_num != current_q_head_num_) {
         current_q_head_num_ = q_head_num;
-        TileGemm122<q_buffer_t, kv_cache_t>::init_tile_config(q_head_num,
-                                                              amx_tile_config_);
+        TileGemm122<kv_cache_t>::init_tile_config(q_head_num, amx_tile_config_);
       }
-      attention<TileGemm122<q_buffer_t, kv_cache_t>> attention_iteration;
+      attention<TileGemm122<kv_cache_t>> attention_iteration;
       attention_iteration(CPU_ATTENTION_PARAMS);
     }
   }
@@ -521,26 +411,13 @@ class AttentionImpl<ISA::AMX, scalar_t, head_dim, kv_cache_scalar_t> {
   // reshape KV to AMX friendly layout
   static void reshape_and_cache(
       const scalar_t* __restrict__ key, const scalar_t* __restrict__ value,
-      kv_cache_t* __restrict__ key_cache, kv_cache_t* __restrict__ value_cache,
+      scalar_t* __restrict__ key_cache, scalar_t* __restrict__ value_cache,
       const int64_t* __restrict__ slot_mapping, const int64_t token_num,
       const int64_t key_token_num_stride, const int64_t value_token_num_stride,
       const int64_t head_num, const int64_t key_head_num_stride,
       const int64_t value_head_num_stride, const int64_t num_blocks,
       const int64_t num_blocks_stride, const int64_t cache_head_num_stride,
-      const int64_t block_size, const int64_t block_size_stride,
-      const float k_inv = 0.0f, const float v_inv = 0.0f) {
-    if constexpr (fp8_kv) {
-      constexpr auto qfn = select_fp8_quant_fn<kv_cache_t>();
-      reshape_and_cache_fp8_amx_impl<scalar_t, qfn>(
-          key, value, reinterpret_cast<uint8_t*>(key_cache),
-          reinterpret_cast<uint8_t*>(value_cache), slot_mapping, token_num,
-          head_num, head_dim, block_size, key_token_num_stride,
-          key_head_num_stride, value_token_num_stride, value_head_num_stride,
-          num_blocks_stride, cache_head_num_stride, num_blocks_stride,
-          cache_head_num_stride, k_inv, v_inv);
-      return;
-    }
-
+      const int64_t block_size, const int64_t block_size_stride) {
     // For AMX 2D tiles, size of each line is 64 bytes
     constexpr int64_t amx_tile_row_size = AMX_TILE_ROW_BYTES;
     // For AMX B matrix, N always is 16
@@ -549,9 +426,6 @@ class AttentionImpl<ISA::AMX, scalar_t, head_dim, kv_cache_scalar_t> {
     // For now suppose block_size is divisible by amx_tile_column_num
     TORCH_CHECK_EQ(block_size % amx_b_tile_k_size, 0);
 
-    scalar_t* __restrict__ kc = reinterpret_cast<scalar_t*>(key_cache);
-    scalar_t* __restrict__ vc = reinterpret_cast<scalar_t*>(value_cache);
-
 #pragma omp parallel for collapse(2)
     for (int64_t token_idx = 0; token_idx < token_num; ++token_idx) {
       for (int64_t head_idx = 0; head_idx < head_num; ++head_idx) {
@@ -579,7 +453,8 @@ class AttentionImpl<ISA::AMX, scalar_t, head_dim, kv_cache_scalar_t> {
           constexpr int64_t quadword_num_per_group =
               token_num_per_group * quadword_num;
           int32_t* key_cache_start_ptr =
-              reinterpret_cast<int32_t*>(kc + block_idx * num_blocks_stride +
+              reinterpret_cast<int32_t*>(key_cache +
+                                         block_idx * num_blocks_stride +
                                          head_idx * cache_head_num_stride) +
               group_idx * quadword_num_per_group + group_offset;
 
@@ -608,7 +483,7 @@ class AttentionImpl<ISA::AMX, scalar_t, head_dim, kv_cache_scalar_t> {
                                             token_idx * value_token_num_stride +
                                             head_idx * value_head_num_stride;
           scalar_t* value_cache_start_ptr =
-              vc + block_idx * num_blocks_stride +
+              value_cache + block_idx * num_blocks_stride +
               head_idx * cache_head_num_stride +
               sub_group_idx * token_num_per_sub_group * amx_b_tile_n_size +
               sub_group_offset;

csrc/cpu/cpu_attn_fp8.hpp

@@ -1,214 +0,0 @@
-// SPDX-License-Identifier: Apache-2.0
-// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-#pragma once
-#include <algorithm>
-#include <cmath>
-#include <cstdint>
-#include <limits>
-#include <type_traits>
-
-#include "cpu/utils.hpp"
-
-typedef uint32_t __attribute__((__may_alias__)) u32_alias_t;
-typedef uint16_t __attribute__((__may_alias__)) u16_alias_t;
-typedef float __attribute__((__may_alias__)) f32_alias_t;
-
-// Reference scalar dequant — used to verify vectorized AMX dequant.
-inline float fp8e4m3_to_float_scalar(uint8_t b, float scale) noexcept {
-  // NaN encoding in E4M3
-  if ((b & 0x7F) == 0x7F) return std::numeric_limits<float>::quiet_NaN();
-  uint32_t b_u32 = static_cast<uint32_t>(b);
-  uint32_t sign = (b_u32 & 0x80) << 24;
-  uint32_t payload = (b_u32 & 0x7F) << 20;
-  uint32_t bits = sign | payload;
-  float b_f32_unscaled = *reinterpret_cast<const f32_alias_t*>(&bits);
-  float b_f32_scaled = b_f32_unscaled * scale * 0x1p120f;
-  return b_f32_scaled;
-}
-
-inline uint8_t float_to_fp8e4m3_scalar(float v, float inv_scale) noexcept {
-  v *= inv_scale;
-  constexpr float fp8_max = 448.0f;
-  v = std::max(-fp8_max, std::min(fp8_max, v));
-  if (v == 0.0f) return 0;
-
-  // Inverse mapping of fp8e4m3_to_float_scalar: shift the effective exponent
-  // bias from fp32 (127) back to fp8 e4m3 (7), then pack sign|payload.
-  float v_f32_unscaled = v * 0x1p-120f;
-  uint32_t bits = *reinterpret_cast<const u32_alias_t*>(&v_f32_unscaled);
-  uint8_t sign = static_cast<uint8_t>((bits >> 24) & 0x80);
-  uint8_t payload = static_cast<uint8_t>((bits >> 20) & 0x7F);
-  if (payload == 0) return sign;
-  payload = std::min<uint8_t>(payload, 0x7E);  // keep 0x7F as NaN encoding
-  return static_cast<uint8_t>(sign | payload);
-}
-
-// ---------------------------------------------------------------------------
-// AMX reshape impl — parameterised on the quantisation function.
-// Writes key/value into uint8 FP8 KV cache using the AMX tile-friendly layout.
-// K: halfword-packed (2 FP8 per uint16, token_num_per_group=16).
-// V: sub-group packing (token_num_per_sub_group=2, head_elems_per_group=16).
-// block_size must be divisible by 32.
-// ---------------------------------------------------------------------------
-template <typename scalar_t, uint8_t (*quant_fn)(float, float)>
-inline void reshape_and_cache_fp8_amx_impl(
-    const scalar_t* key_ptr, const scalar_t* value_ptr, uint8_t* key_cache_ptr,
-    uint8_t* value_cache_ptr, const int64_t* slot_ptr, int64_t token_num,
-    int64_t head_num, int64_t head_dim, int64_t block_size, int64_t k_stride0,
-    int64_t k_stride1, int64_t v_stride0, int64_t v_stride1, int64_t kc_stride0,
-    int64_t kc_stride1, int64_t vc_stride0, int64_t vc_stride1, float k_inv,
-    float v_inv) {
-  constexpr int64_t token_num_per_group = 16;  // AMX_TILE_ROW_NUM
-  const int64_t halfword_num = head_dim / 2;   // 2 FP8 per uint16
-  const int64_t halfword_num_per_group = token_num_per_group * halfword_num;
-  constexpr int64_t head_elems_per_group = 16;
-  constexpr int64_t token_num_per_sub_group = 2;  // = 4 / sizeof(BF16)
-  const int64_t group_num = head_dim / head_elems_per_group;
-  const int64_t group_size = block_size * head_elems_per_group;
-
-#pragma omp parallel for collapse(2) schedule(static)
-  for (int64_t tok = 0; tok < token_num; ++tok) {
-    for (int64_t h = 0; h < head_num; ++h) {
-      const int64_t slot = slot_ptr[tok];
-      if (slot < 0) continue;
-      const int64_t block_idx = slot / block_size;
-      const int64_t block_offset = slot % block_size;
-
-      // Key: halfword-packed, 2 FP8 per uint16
-      {
-        const scalar_t* ksrc = key_ptr + tok * k_stride0 + h * k_stride1;
-        const int64_t group_idx = block_offset / token_num_per_group;
-        const int64_t group_offset = block_offset % token_num_per_group;
-        uint16_t* kdst =
-            reinterpret_cast<uint16_t*>(key_cache_ptr + block_idx * kc_stride0 +
-                                        h * kc_stride1) +
-            group_idx * halfword_num_per_group + group_offset;
-        for (int64_t j = 0; j < halfword_num; ++j) {
-          uint8_t fp8_0 = quant_fn(static_cast<float>(ksrc[j * 2]), k_inv);
-          uint8_t fp8_1 = quant_fn(static_cast<float>(ksrc[j * 2 + 1]), k_inv);
-          uint8_t bytes[2] = {fp8_0, fp8_1};
-          uint16_t hw = *reinterpret_cast<const u16_alias_t*>(bytes);
-          kdst[j * token_num_per_group] = hw;
-        }
-      }
-
-      // Value: sub-group packing (token_num_per_sub_group = 2)
-      {
-        const scalar_t* vsrc = value_ptr + tok * v_stride0 + h * v_stride1;
-        const int64_t sub_group_idx = block_offset / token_num_per_sub_group;
-        const int64_t sub_group_offset = block_offset % token_num_per_sub_group;
-        uint8_t* vdst =
-            value_cache_ptr + block_idx * vc_stride0 + h * vc_stride1 +
-            sub_group_idx * token_num_per_sub_group * head_elems_per_group +
-            sub_group_offset;
-        for (int64_t i = 0; i < group_num; ++i) {
-          for (int64_t j = 0; j < head_elems_per_group; ++j)
-            vdst[j * token_num_per_sub_group] =
-                quant_fn(static_cast<float>(vsrc[j]), v_inv);
-          vsrc += head_elems_per_group;
-          vdst += group_size;
-        }
-      }
-    }
-  }
-}
-
-// ---------------------------------------------------------------------------
-// FP8 E5M2 scalar helpers
-// ---------------------------------------------------------------------------
-
-// Reference scalar dequant — used to verify vectorized AMX dequant.
-// FP8 E5M2: s[7] e[6:2] m[1:0], exponent bias = 15 (same as FP16).
-// Byte b → FP16 bits = b << 8 (no bias correction needed).
-inline float fp8e5m2_to_float_scalar(uint8_t b, float scale) noexcept {
-  const uint8_t exp_bits = (b >> 2) & 0x1F;
-  const uint8_t mant_bits = b & 0x03;
-  // NaN: exp=11111, mant!=00
-  if (exp_bits == 0x1F && mant_bits != 0)
-    return std::numeric_limits<float>::quiet_NaN();
-  const uint32_t sign = static_cast<uint32_t>(b & 0x80) << 24;
-  if (exp_bits == 0x1F)
-    return sign ? -std::numeric_limits<float>::infinity()
-                : std::numeric_limits<float>::infinity();
-  if (exp_bits == 0) {  // subnormal: (-1)^s * 2^-14 * mant/4
-    if (mant_bits == 0) return 0.0f;
-    float v = mant_bits * 0x1p-16f;
-    return (sign ? -v : v) * scale;
-  }
-  // Normal: FP32 exp = exp5 - 15 + 127, mantissa top 2 bits
-  uint32_t fp32_bits = sign |
-                       ((static_cast<uint32_t>(exp_bits) - 15 + 127) << 23) |
-                       (static_cast<uint32_t>(mant_bits) << 21);
-  float val = *reinterpret_cast<const f32_alias_t*>(&fp32_bits);
-  return val * scale;
-}
-
-inline uint8_t float_to_fp8e5m2_scalar(float v, float inv_scale) noexcept {
-  v *= inv_scale;
-  constexpr float fp8_e5m2_max = 57344.0f;
-  v = std::max(-fp8_e5m2_max, std::min(fp8_e5m2_max, v));
-  if (v == 0.0f) return 0;
-  uint32_t bits = *reinterpret_cast<const u32_alias_t*>(&v);
-  const uint8_t sign = static_cast<uint8_t>((bits >> 24) & 0x80);
-  const int32_t exp_fp32 = static_cast<int32_t>((bits >> 23) & 0xFF) - 127;
-  const uint8_t mant2 = static_cast<uint8_t>((bits >> 21) & 0x03);
-  if (exp_fp32 < -14) {  // subnormal in E5M2
-    const int shift = -14 - exp_fp32;
-    if (shift + 21 >= 32)
-      return sign;  // underflow: too small for E5M2 subnormal
-    const uint32_t m = (0x800000u | (bits & 0x7FFFFFu)) >> (shift + 21);
-    return sign | static_cast<uint8_t>(std::min<uint32_t>(m, 3u));
-  }
-  const uint8_t exp5 = static_cast<uint8_t>(exp_fp32 + 15);
-  return sign | (exp5 << 2) | mant2;
-}
-
-// ---------------------------------------------------------------------------
-// Select the FP8 quant function at compile time based on kv_cache_t.
-// ---------------------------------------------------------------------------
-template <typename kv_cache_t>
-constexpr auto select_fp8_quant_fn() {
-  if constexpr (std::is_same_v<kv_cache_t, c10::Float8_e5m2>)
-    return float_to_fp8e5m2_scalar;
-  else
-    return float_to_fp8e4m3_scalar;
-}
-
-// ---------------------------------------------------------------------------
-// VEC reshape impl — parameterised on the quantisation function.
-// Writes key (column-major) and value (row-major) into uint8 FP8 KV cache.
-// The pragma omp must live outside VLLM_DISPATCH_FLOATING_TYPES because
-// #pragma cannot appear inside variadic macro arguments.
-// ---------------------------------------------------------------------------
-template <typename scalar_t, uint8_t (*quant_fn)(float, float)>
-inline void reshape_and_cache_fp8_vec_impl(
-    const scalar_t* key_ptr, const scalar_t* value_ptr, uint8_t* key_cache_ptr,
-    uint8_t* value_cache_ptr, const int64_t* slot_ptr, int64_t token_num,
-    int64_t head_num, int64_t head_dim, int64_t block_size, int64_t k_stride0,
-    int64_t k_stride1, int64_t v_stride0, int64_t v_stride1, int64_t kc_stride0,
-    int64_t kc_stride1, int64_t vc_stride0, int64_t vc_stride1, float k_inv,
-    float v_inv) {
-#pragma omp parallel for collapse(2) schedule(static)
-  for (int64_t tok = 0; tok < token_num; ++tok) {
-    for (int64_t h = 0; h < head_num; ++h) {
-      const int64_t slot = slot_ptr[tok];
-      if (slot < 0) continue;
-      const int64_t block_idx = slot / block_size;
-      const int64_t block_offset = slot % block_size;
-
-      // Key layout: column-major within block
-      const scalar_t* ksrc = key_ptr + tok * k_stride0 + h * k_stride1;
-      uint8_t* kdst = key_cache_ptr + block_idx * kc_stride0 + h * kc_stride1 +
-                      block_offset;
-      for (int64_t i = 0; i < head_dim; ++i)
-        kdst[i * block_size] = quant_fn(static_cast<float>(ksrc[i]), k_inv);
-
-      // Value layout: row-major within block (contiguous head_dim bytes)
-      const scalar_t* vsrc = value_ptr + tok * v_stride0 + h * v_stride1;
-      uint8_t* vdst = value_cache_ptr + block_idx * vc_stride0 +
-                      h * vc_stride1 + block_offset * head_dim;
-      for (int64_t i = 0; i < head_dim; ++i)
-        vdst[i] = quant_fn(static_cast<float>(vsrc[i]), v_inv);
-    }
-  }
-}

csrc/cpu/cpu_attn_impl.hpp

@@ -12,24 +12,10 @@
 #include "cpu/utils.hpp"
 
 namespace cpu_attention {
-enum class ISA { AMX, VEC, VEC16, NEON, VXE, RVV, VSX };
+enum class ISA { AMX, VEC, VEC16, NEON, VXE };
 
-// Mirrors csrc/attention/dtype_fp8.cuh Fp8KVCacheDataType exactly.
-enum class Fp8KVCacheDataType {
-  kAuto = 0,
-  kFp8E4M3 = 1,
-  kFp8E5M2 = 2,
-};
-
-struct AttentionInput;
-
-template <ISA isa, typename scalar_t, int64_t head_dim,
-          typename kv_cache_scalar_t = scalar_t>
-class AttentionImpl {
- public:
-  void init_from_input(const AttentionInput*) {}
-  float get_output_v_scale() const noexcept { return 1.0f; }
-};
+template <ISA isa, typename scalar_t, int64_t head_dim>
+class AttentionImpl {};
 
 struct AttentionWorkItemGroup {
   int32_t req_id;
@@ -164,12 +150,6 @@ struct AttentionMetadata {
       case ISA::VXE:
         ss << "VXE, ";
         break;
-      case ISA::RVV:
-        ss << "RVV, ";
-        break;
-      case ISA::VSX:
-        ss << "VSX, ";
-        break;
     }
     ss << "workitem_group_num: " << workitem_group_num
        << ", reduction_item_num: " << reduction_item_num
@@ -800,9 +780,6 @@ struct AttentionInput {
   int32_t sliding_window_left;
   int32_t sliding_window_right;
   float softcap;
-  // FP8 KV cache scales (used by FP8 attention implementations)
-  float k_scale_fp8 = 1.0f;
-  float v_scale_fp8 = 1.0f;
 };
 
 #define DEFINE_CPU_ATTENTION_PARAMS                                         \
@@ -1397,13 +1374,6 @@ class AttentionMainLoop {
       }
 
       attention_impl_t attn_impl;
-      constexpr bool fp8_kv = std::is_same_v<kv_cache_t, c10::Float8_e4m3fn> ||
-                              std::is_same_v<kv_cache_t, c10::Float8_e5m2>;
-      float output_v_scale = 1.0f;
-      if constexpr (fp8_kv) {
-        attn_impl.init_from_input(input);
-        output_v_scale = attn_impl.get_output_v_scale();
-      }
 
       // general information
       const int32_t q_head_num = input->num_heads;
@@ -1783,7 +1753,7 @@ class AttentionMainLoop {
                                reinterpret_cast<query_t*>(input->output) +
                                    output_buffer_offset,
                                sum_buffer, actual_q_heads_per_kv,
-                               actual_q_token_num, q_head_num, output_v_scale);
+                               actual_q_token_num, q_head_num);
                 } else {
                   const int32_t stride =
                       actual_q_heads_per_kv * split_kv_q_token_num_threshold;
@@ -1853,7 +1823,7 @@ class AttentionMainLoop {
               split_output_buffer,
               reinterpret_cast<query_t*>(input->output) + output_buffer_offset,
               split_sum_buffer, actual_q_heads_per_kv, curr_output_token_num,
-              q_head_num, output_v_scale);
+              q_head_num);
         }
       }
     }
@@ -1977,8 +1947,8 @@ class AttentionMainLoop {
                     query_t* __restrict__ curr_output_buffer,
                     float* __restrict__ sum_buffer,
                     const int32_t q_heads_per_kv,
-                    const int32_t actual_q_token_num, const int32_t q_head_num,
-                    const float v_scale = 1.0f) {
+                    const int32_t actual_q_token_num,
+                    const int32_t q_head_num) {
     // final output
     using output_vec_t = typename VecTypeTrait<query_t>::vec_t;
 
@@ -1992,7 +1962,7 @@ class AttentionMainLoop {
           curr_partial_output_buffer;
       query_t* __restrict__ curr_output_buffer_iter = curr_output_buffer;
       for (int32_t head_idx = 0; head_idx < q_heads_per_kv; ++head_idx) {
-        vec_op::FP32Vec16 inv_sum_scale_vec(v_scale / *curr_sum_buffer);
+        vec_op::FP32Vec16 inv_sum_scale_vec(1.0 / *curr_sum_buffer);
 
         for (int32_t i = 0; i < group_num_per_head; ++i) {
           vec_op::FP32Vec16 vec(curr_partial_output_buffer_iter);

csrc/cpu/cpu_attn_neon.hpp

@@ -248,8 +248,8 @@ class TileGemmNeonFMLA {
 }  // namespace
 
 // this is similar to "ISA::VEC" at the moment
-template <typename scalar_t, int64_t head_dim, typename kv_cache_scalar_t>
-class AttentionImpl<ISA::NEON, scalar_t, head_dim, kv_cache_scalar_t> {
+template <typename scalar_t, int64_t head_dim>
+class AttentionImpl<ISA::NEON, scalar_t, head_dim> {
  public:
   using query_t = scalar_t;
   using q_buffer_t = float;
@@ -343,8 +343,7 @@ class AttentionImpl<ISA::NEON, scalar_t, head_dim, kv_cache_scalar_t> {
       const int64_t head_num, const int64_t key_head_num_stride,
       const int64_t value_head_num_stride, const int64_t num_blocks,
       const int64_t num_blocks_stride, const int64_t cache_head_num_stride,
-      const int64_t block_size, const int64_t block_size_stride,
-      const float /*k_inv*/ = 0.0f, const float /*v_inv*/ = 0.0f) {
+      const int64_t block_size, const int64_t block_size_stride) {
 #pragma omp parallel for collapse(2)
     for (int64_t token_idx = 0; token_idx < token_num; ++token_idx) {
       for (int64_t head_idx = 0; head_idx < head_num; ++head_idx) {
@@ -389,7 +388,7 @@ class AttentionImpl<ISA::NEON, scalar_t, head_dim, kv_cache_scalar_t> {
 #ifdef ARM_BF16_SUPPORT
 // For BF16 on Arm, reuse the BFMMLA kernels with 32-token alignment.
 template <int64_t head_dim>
-class AttentionImpl<ISA::NEON, c10::BFloat16, head_dim, c10::BFloat16>
+class AttentionImpl<ISA::NEON, c10::BFloat16, head_dim>
     : public AttentionImplNEONBFMMLA<BLOCK_SIZE_ALIGNMENT, ISA::NEON,
                                      head_dim> {};
 #endif

csrc/cpu/cpu_attn_neon_bfmmla.hpp

@@ -602,8 +602,7 @@ class AttentionImplNEONBFMMLA {
       [[maybe_unused]] const int64_t num_blocks,
       const int64_t num_blocks_stride, const int64_t cache_head_num_stride,
       const int64_t block_size,
-      [[maybe_unused]] const int64_t block_size_stride,
-      const float /*k_inv*/ = 0.0f, const float /*v_inv*/ = 0.0f) {
+      [[maybe_unused]] const int64_t block_size_stride) {
     const int64_t k_block_stride = (head_dim / TILE_K) * K_INNER_STRIDE;
     const int64_t v_pair_stride =
         (block_size / V_TOKENS_PER_ROW_BLOCK) * V_INNER_STRIDE;

csrc/cpu/cpu_attn_rvv.hpp

@@ -1,445 +0,0 @@
-// SPDX-License-Identifier: Apache-2.0
-// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
-
-#ifndef CPU_ATTN_RVV_HPP
-#define CPU_ATTN_RVV_HPP
-
-// This kernel is currently hardcoded to VLEN=128 (m1/m2 intrinsics, vl=8).
-// The fixed-width typedefs below use `riscv_rvv_vector_bits(128)`, which
-// only matches `vfloat16m1_t`/`vuint16m1_t` register layout when VLEN==128;
-// at VLEN>=256 those typedefs fail to compile.  Scalar RISC-V builds
-// (-march=rv64gc) additionally don't have <riscv_vector.h>.  For both
-// cases we omit the file entirely and let the dispatcher fall back to the
-// scalar VEC / VEC16 implementations.  TODO: migrate to RVVI() macros +
-// semantic names in cpu_types_riscv_defs.hpp to support VLEN>=256 natively.
-#if defined(__riscv_v_min_vlen) && __riscv_v_min_vlen == 128
-
-  #include "cpu_attn_impl.hpp"
-  #include <riscv_vector.h>
-  #include <type_traits>
-
-namespace cpu_attention {
-
-namespace {
-
-// File-local concrete-LMUL typedefs.  The shared _defs.hpp exposes
-// VLEN-independent semantic names (fixed_fp32x8_t, fixed_fp16x8_t, ...),
-// but this kernel is currently hardcoded to VLEN=128 (m1/m2 intrinsics),
-// so keep the legacy concrete aliases scoped to this file.
-typedef vfloat16m1_t fixed_vfloat16m1_t
-    __attribute__((riscv_rvv_vector_bits(128)));
-typedef vfloat32m2_t fixed_vfloat32m2_t
-    __attribute__((riscv_rvv_vector_bits(256)));
-typedef vuint16m1_t fixed_vuint16m1_t
-    __attribute__((riscv_rvv_vector_bits(128)));
-typedef vuint32m2_t fixed_vuint32m2_t
-    __attribute__((riscv_rvv_vector_bits(256)));
-  #ifdef __riscv_zvfbfmin
-typedef vbfloat16m1_t fixed_vbfloat16m1_t
-    __attribute__((riscv_rvv_vector_bits(128)));
-  #endif
-
-  #define BLOCK_SIZE_ALIGNMENT 32
-  #define HEAD_SIZE_ALIGNMENT 32
-  #define MAX_Q_HEAD_NUM_PER_ITER 16
-
-// ============================================================================
-// B-matrix row loading: load 8 elements as FP32 (using m2 LMUL at VLEN=128)
-// ============================================================================
-
-template <typename kv_cache_t>
-FORCE_INLINE fixed_vfloat32m2_t load_row8_B_as_f32(const kv_cache_t* p);
-
-template <>
-FORCE_INLINE fixed_vfloat32m2_t load_row8_B_as_f32<float>(const float* p) {
-  return __riscv_vle32_v_f32m2(p, 8);
-}
-
-template <>
-FORCE_INLINE fixed_vfloat32m2_t
-load_row8_B_as_f32<c10::Half>(const c10::Half* p) {
-  #ifdef __riscv_zvfh
-  fixed_vfloat16m1_t h =
-      __riscv_vle16_v_f16m1(reinterpret_cast<const _Float16*>(p), 8);
-  return __riscv_vfwcvt_f_f_v_f32m2(h, 8);
-  #else
-  // Fallback for hardware without Zvfh: scalar half->float conversion.
-  // c10::Half provides operator float() so this is correct on any RVV CPU
-  // that has only the base V extension.  Slower than the Zvfh path, but
-  // keeps the kernel buildable on Zvfhmin-only / no-fp16 hardware.
-  alignas(16) float tmp[8];
-  for (int i = 0; i < 8; ++i) {
-    tmp[i] = static_cast<float>(p[i]);
-  }
-  return __riscv_vle32_v_f32m2(tmp, 8);
-  #endif
-}
-
-template <>
-FORCE_INLINE fixed_vfloat32m2_t
-load_row8_B_as_f32<c10::BFloat16>(const c10::BFloat16* p) {
-  #ifdef __riscv_zvfbfmin
-  fixed_vbfloat16m1_t bf =
-      __riscv_vle16_v_bf16m1(reinterpret_cast<const __bf16*>(p), 8);
-  return __riscv_vfwcvtbf16_f_f_v_f32m2(bf, 8);
-  #else
-  // Fallback: load as uint16, zero-extend to uint32, shift left by 16
-  fixed_vuint16m1_t raw =
-      __riscv_vle16_v_u16m1(reinterpret_cast<const uint16_t*>(p), 8);
-  fixed_vuint32m2_t wide = __riscv_vzext_vf2_u32m2(raw, 8);
-  fixed_vuint32m2_t shifted = __riscv_vsll_vx_u32m2(wide, 16, 8);
-  return __riscv_vreinterpret_v_u32m2_f32m2(shifted);
-  #endif
-}
-
-// ============================================================================
-// Micro kernel: Mx8 tile, K unrolled by 4, RVV scalar-broadcast FMA
-// ============================================================================
-//
-// NEON uses vfmaq_laneq_f32 (lane-indexed FMA from a preloaded A vector).
-// RVV has no lane-indexed FMA; instead we load A elements as scalars and
-// use __riscv_vfmacc_vf (scalar * vector + accumulator), which is equally
-// efficient and avoids the need for vrgather/vslidedown.
-//
-// At VLEN=128, m2 holds 8 x FP32, matching the 8-column tile width.
-// Register budget: M accumulators (m2 each) + 1 B temp = 2M+2 regs.
-// M=8 => 18 regs out of 32 available — no spills.
-
-template <int32_t M, typename kv_cache_t>
-FORCE_INLINE void gemm_micro_rvv_fma_Mx8_Ku4(
-    const float* __restrict A,       // [M x K]
-    const kv_cache_t* __restrict B,  // [K x 8]
-    float* __restrict C,             // [M x 8]
-    int64_t lda, int64_t ldb, int64_t ldc, int32_t K, bool accumulate) {
-  static_assert(1 <= M && M <= 8, "M must be in [1,8]");
-
-  constexpr size_t vl = 8;
-
-  // helpers for per-M codegen
-  #define ROWS_APPLY(OP) OP(0) OP(1) OP(2) OP(3) OP(4) OP(5) OP(6) OP(7)
-  #define IF_M(i) if constexpr (M > (i))
-
-    // A row base pointers
-  #define DECL_A(i) const float* a##i = A + (i) * lda;
-  ROWS_APPLY(DECL_A)
-  #undef DECL_A
-
-    // declare one m2 accumulator per row
-  #define DECL_ACC(i) fixed_vfloat32m2_t acc##i;
-  ROWS_APPLY(DECL_ACC)
-  #undef DECL_ACC
-
-    // initialize accumulators
-  #define INIT_ACC(i)                                      \
-    IF_M(i) {                                              \
-      if (accumulate) {                                    \
-        acc##i = __riscv_vle32_v_f32m2(C + (i) * ldc, vl); \
-      } else {                                             \
-        acc##i = __riscv_vfmv_v_f_f32m2(0.f, vl);          \
-      }                                                    \
-    }
-  ROWS_APPLY(INIT_ACC)
-  #undef INIT_ACC
-
-  int32_t k = 0;
-
-  // K unrolled by 4
-  for (; k + 3 < K; k += 4) {
-    // k + 0
-    {
-      fixed_vfloat32m2_t b =
-          load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 0) * ldb);
-  #define STEP_K0(i)                                                    \
-    IF_M(i) {                                                           \
-      acc##i = __riscv_vfmacc_vf_f32m2(acc##i, *(a##i + k + 0), b, vl); \
-    }
-      ROWS_APPLY(STEP_K0)
-  #undef STEP_K0
-    }
-    // k + 1
-    {
-      fixed_vfloat32m2_t b =
-          load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 1) * ldb);
-  #define STEP_K1(i)                                                    \
-    IF_M(i) {                                                           \
-      acc##i = __riscv_vfmacc_vf_f32m2(acc##i, *(a##i + k + 1), b, vl); \
-    }
-      ROWS_APPLY(STEP_K1)
-  #undef STEP_K1
-    }
-    // k + 2
-    {
-      fixed_vfloat32m2_t b =
-          load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 2) * ldb);
-  #define STEP_K2(i)                                                    \
-    IF_M(i) {                                                           \
-      acc##i = __riscv_vfmacc_vf_f32m2(acc##i, *(a##i + k + 2), b, vl); \
-    }
-      ROWS_APPLY(STEP_K2)
-  #undef STEP_K2
-    }
-    // k + 3
-    {
-      fixed_vfloat32m2_t b =
-          load_row8_B_as_f32<kv_cache_t>(B + (int64_t)(k + 3) * ldb);
-  #define STEP_K3(i)                                                    \
-    IF_M(i) {                                                           \
-      acc##i = __riscv_vfmacc_vf_f32m2(acc##i, *(a##i + k + 3), b, vl); \
-    }
-      ROWS_APPLY(STEP_K3)
-  #undef STEP_K3
-    }
-  }
-
-  // K tail
-  for (; k < K; ++k) {
-    fixed_vfloat32m2_t b = load_row8_B_as_f32<kv_cache_t>(B + (int64_t)k * ldb);
-  #define TAIL_ROW(i) \
-    IF_M(i) { acc##i = __riscv_vfmacc_vf_f32m2(acc##i, *(a##i + k), b, vl); }
-    ROWS_APPLY(TAIL_ROW)
-  #undef TAIL_ROW
-  }
-
-    // store accumulators to C
-  #define STORE_ROW(i) \
-    IF_M(i) { __riscv_vse32_v_f32m2(C + (i) * ldc, acc##i, vl); }
-  ROWS_APPLY(STORE_ROW)
-  #undef STORE_ROW
-
-  #undef ROWS_APPLY
-  #undef IF_M
-}
-
-// ============================================================================
-// Macro kernel: dispatch M tiles of {8,4,2,1}, step N by 8
-// ============================================================================
-
-template <int32_t N, typename kv_cache_t>
-FORCE_INLINE void gemm_macro_rvv_fma_Mx8_Ku4(const float* __restrict A,
-                                             const kv_cache_t* __restrict B,
-                                             float* __restrict C, int32_t M,
-                                             int32_t K, int64_t lda,
-                                             int64_t ldb, int64_t ldc,
-                                             bool accumulate) {
-  static_assert(N % 8 == 0, "N must be a multiple of 8");
-  for (int32_t m = 0; m < M;) {
-    int32_t mb = (M - m >= 8) ? 8 : (M - m >= 4) ? 4 : (M - m >= 2) ? 2 : 1;
-    const float* Ab = A + m * lda;
-    float* Cb = C + m * ldc;
-
-    for (int32_t n = 0; n < N; n += 8) {
-      const kv_cache_t* Bn = B + n;
-      float* Cn = Cb + n;
-      switch (mb) {
-        case 8:
-          gemm_micro_rvv_fma_Mx8_Ku4<8, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
-                                                    K, accumulate);
-          break;
-        case 4:
-          gemm_micro_rvv_fma_Mx8_Ku4<4, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
-                                                    K, accumulate);
-          break;
-        case 2:
-          gemm_micro_rvv_fma_Mx8_Ku4<2, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
-                                                    K, accumulate);
-          break;
-        default:
-          gemm_micro_rvv_fma_Mx8_Ku4<1, kv_cache_t>(Ab, Bn, Cn, lda, ldb, ldc,
-                                                    K, accumulate);
-          break;
-      }
-    }
-    m += mb;
-  }
-}
-
-// ============================================================================
-// TileGemm wrapper — plugs into AttentionMainLoop
-// ============================================================================
-
-template <typename kv_cache_t>
-class TileGemmRVV {
- public:
-  template <AttentionGemmPhase phase, int32_t k_size>
-  FORCE_INLINE static void gemm(const int32_t m_size,
-                                float* __restrict__ a_tile,
-                                kv_cache_t* __restrict__ b_tile,
-                                float* __restrict__ c_tile, const int64_t lda,
-                                const int64_t ldb, const int64_t ldc,
-                                const int32_t block_size,
-                                const int32_t dynamic_k_size,
-                                const bool accum_c) {
-    if constexpr (phase == AttentionGemmPhase::QK) {
-      gemm_macro_rvv_fma_Mx8_Ku4<BLOCK_SIZE_ALIGNMENT, kv_cache_t>(
-          a_tile, b_tile, c_tile, m_size, k_size, lda, ldb, ldc, accum_c);
-    } else {
-      gemm_macro_rvv_fma_Mx8_Ku4<HEAD_SIZE_ALIGNMENT, kv_cache_t>(
-          a_tile, b_tile, c_tile, m_size, dynamic_k_size, lda, ldb, ldc,
-          accum_c);
-    }
-  }
-};
-
-}  // namespace
-
-// ============================================================================
-// AttentionImpl<ISA::RVV> — mirrors ISA::NEON specialization
-// ============================================================================
-
-template <typename scalar_t, int64_t head_dim, typename kv_cache_scalar_t>
-class AttentionImpl<ISA::RVV, scalar_t, head_dim, kv_cache_scalar_t> {
- public:
-  using query_t = scalar_t;
-  using q_buffer_t = float;
-  using kv_cache_t = scalar_t;
-  using logits_buffer_t = float;
-  using partial_output_buffer_t = float;
-  using prob_buffer_t = float;
-
-  constexpr static int64_t BlockSizeAlignment = BLOCK_SIZE_ALIGNMENT;
-  constexpr static int64_t HeadDimAlignment = HEAD_SIZE_ALIGNMENT;
-  constexpr static int64_t MaxQHeadNumPerIteration = MAX_Q_HEAD_NUM_PER_ITER;
-  constexpr static int64_t HeadDim = head_dim;
-  constexpr static ISA ISAType = ISA::RVV;
-  constexpr static bool scale_on_logits = false;
-
-  static_assert(HeadDim % HeadDimAlignment == 0);
-  static_assert(HeadDimAlignment % 8 == 0);
-  static_assert(BlockSizeAlignment % 8 == 0);
-
- public:
-  template <template <typename tile_gemm_t> typename attention>
-  FORCE_INLINE void execute_attention(DEFINE_CPU_ATTENTION_PARAMS) {
-    attention<TileGemmRVV<kv_cache_t>> attention_iteration;
-    attention_iteration(CPU_ATTENTION_PARAMS);
-  }
-
-  constexpr static int64_t k_cache_token_group_stride(
-      const int32_t block_size) {
-    return BlockSizeAlignment;
-  }
-
-  constexpr static int64_t v_cache_token_group_stride(
-      const int32_t block_size) {
-    return head_dim * BlockSizeAlignment;
-  }
-
-  constexpr static int64_t v_cache_head_group_stride(const int32_t block_size) {
-    return HeadDimAlignment;
-  }
-
-  static void copy_q_heads_tile(scalar_t* __restrict__ src,
-                                float* __restrict__ q_buffer,
-                                const int32_t q_num,
-                                const int32_t q_heads_per_kv,
-                                const int64_t q_num_stride,
-                                const int64_t q_head_stride, float scale) {
-    static_assert(head_dim % 16 == 0);
-    constexpr int32_t unroll_size = head_dim / 16;
-    using load_vec_t = typename VecTypeTrait<scalar_t>::vec_t;
-
-    vec_op::FP32Vec16 scale_vec(scale);
-    for (int32_t q_num_idx = 0; q_num_idx < q_num; ++q_num_idx) {
-      for (int32_t q_head_idx = 0; q_head_idx < q_heads_per_kv; ++q_head_idx) {
-        scalar_t* __restrict__ curr_q =
-            src + q_num_idx * q_num_stride + q_head_idx * q_head_stride;
-        float* __restrict__ curr_q_buffer =
-            q_buffer + q_num_idx * q_heads_per_kv * head_dim +
-            q_head_idx * head_dim;
-
-        vec_op::unroll_loop<int32_t, unroll_size>([&](int32_t i) {
-          load_vec_t vec(curr_q);
-          vec_op::FP32Vec16 fp32_vec(vec);
-          fp32_vec = fp32_vec * scale_vec;
-          fp32_vec.save(curr_q_buffer);
-
-          curr_q += 16;
-          curr_q_buffer += 16;
-        });
-      }
-    }
-  }
-
-  static void reshape_and_cache(
-      const scalar_t* __restrict__ key, const scalar_t* __restrict__ value,
-      scalar_t* __restrict__ key_cache, scalar_t* __restrict__ value_cache,
-      const int64_t* __restrict__ slot_mapping, const int64_t token_num,
-      const int64_t key_token_num_stride, const int64_t value_token_num_stride,
-      const int64_t head_num, const int64_t key_head_num_stride,
-      const int64_t value_head_num_stride, const int64_t num_blocks,
-      const int64_t num_blocks_stride, const int64_t cache_head_num_stride,
-      const int64_t block_size, const int64_t block_size_stride,
-      const float /*k_inv*/ = 0.0f, const float /*v_inv*/ = 0.0f) {
-  #pragma omp parallel for collapse(2)
-    for (int64_t token_idx = 0; token_idx < token_num; ++token_idx) {
-      for (int64_t head_idx = 0; head_idx < head_num; ++head_idx) {
-        const int64_t pos = slot_mapping[token_idx];
-        if (pos < 0) {
-          continue;
-        }
-
-        const int64_t block_idx = pos / block_size;
-        const int64_t block_offset = pos % block_size;
-        {
-          // Write Key (transpose to column-major: [head_dim, block_size])
-          const scalar_t* key_start_ptr = key +
-                                          token_idx * key_token_num_stride +
-                                          head_idx * key_head_num_stride;
-          scalar_t* key_cache_start_ptr =
-              key_cache + block_idx * num_blocks_stride +
-              head_idx * cache_head_num_stride + block_offset;
-
-          // Strided vector store for efficient transpose.
-          // Load contiguous key elements, store with stride = block_size.
-          {
-            const ptrdiff_t byte_stride = block_size * sizeof(scalar_t);
-            int64_t i = 0;
-            for (; i < head_dim;) {
-              size_t vl;
-              if constexpr (std::is_same_v<scalar_t, float>) {
-                vl = __riscv_vsetvl_e32m2(head_dim - i);
-                vfloat32m2_t v = __riscv_vle32_v_f32m2(
-                    reinterpret_cast<const float*>(key_start_ptr + i), vl);
-                __riscv_vsse32_v_f32m2(
-                    reinterpret_cast<float*>(key_cache_start_ptr +
-                                             i * block_size),
-                    byte_stride, v, vl);
-              } else {
-                // Half and BFloat16 are both 16-bit types
-                vl = __riscv_vsetvl_e16m1(head_dim - i);
-                vuint16m1_t v = __riscv_vle16_v_u16m1(
-                    reinterpret_cast<const uint16_t*>(key_start_ptr + i), vl);
-                __riscv_vsse16_v_u16m1(
-                    reinterpret_cast<uint16_t*>(key_cache_start_ptr +
-                                                i * block_size),
-                    byte_stride, v, vl);
-              }
-              i += vl;
-            }
-          }
-        }
-        {
-          // Write Value (row-major: [block_size, head_dim])
-          const scalar_t* value_start_ptr = value +
-                                            token_idx * value_token_num_stride +
-                                            head_idx * value_head_num_stride;
-          scalar_t* value_cache_start_ptr =
-              value_cache + block_idx * num_blocks_stride +
-              head_idx * cache_head_num_stride + block_offset * head_dim;
-          std::memcpy(value_cache_start_ptr, value_start_ptr,
-                      sizeof(scalar_t) * head_dim);
-        }
-      }
-    }
-  }
-};
-
-}  // namespace cpu_attention
-
-  #undef BLOCK_SIZE_ALIGNMENT
-  #undef HEAD_SIZE_ALIGNMENT
-  #undef MAX_Q_HEAD_NUM_PER_ITER
-
-#endif  // __riscv_v_min_vlen == 128
-
-#endif  // CPU_ATTN_RVV_HPP

csrc/cpu/cpu_attn_vec.hpp

@@ -1,37 +1,11 @@
 #ifndef CPU_ATTN_VEC_HPP
 #define CPU_ATTN_VEC_HPP
 
-#include "cpu_attn_fp8.hpp"
 #include "cpu_attn_impl.hpp"
 
 namespace cpu_attention {
 
 namespace {
-
-// Load 32 kv_cache_t elements starting at ptr and return them as two FP32Vec16s
-// covering the lower 16 and upper 16 positions.
-// For FP8: both halves come from a single BF16Vec32 dequant of 32 bytes.
-// For BF16/FP16/FP32: two separate vector loads at ptr and ptr+16.
-template <typename kv_cache_t>
-FORCE_INLINE std::pair<vec_op::FP32Vec16, vec_op::FP32Vec16> load_b_pair_vec(
-    const kv_cache_t* ptr) {
-  if constexpr (std::is_same_v<kv_cache_t, c10::Float8_e4m3fn>) {
-    // BF16 container, but values are in the FP16 exponent range (bias 15 not
-    // 127).
-    vec_op::BF16Vec32 bf16_b_reg(reinterpret_cast<const uint8_t*>(ptr),
-                                 vec_op::fp8_e4m3_tag{});
-    return {vec_op::FP32Vec16(bf16_b_reg, 0), vec_op::FP32Vec16(bf16_b_reg, 1)};
-  } else if constexpr (std::is_same_v<kv_cache_t, c10::Float8_e5m2>) {
-    vec_op::BF16Vec32 bf16_b_reg(reinterpret_cast<const uint8_t*>(ptr),
-                                 vec_op::fp8_e5m2_tag{});
-    return {vec_op::FP32Vec16(bf16_b_reg, 0), vec_op::FP32Vec16(bf16_b_reg, 1)};
-  } else {
-    using load_vec_t = typename VecTypeTrait<kv_cache_t>::vec_t;
-    return std::make_pair(vec_op::FP32Vec16(load_vec_t(ptr)),
-                          vec_op::FP32Vec16(load_vec_t(ptr + 16)));
-  }
-}
-
 // 8-2-16 pattern, 8 regs for A, 2 regs for B, 16 regs for C, [8, K] @ [k, 32]
 template <typename kv_cache_t>
 class TileGemm82 {
@@ -80,7 +54,10 @@ class TileGemm82 {
                          const int32_t block_size, const int32_t dynamic_k_size,
                          const bool accum_c) {
     static_assert(0 < M && M <= 8);
+    using load_vec_t = typename VecTypeTrait<kv_cache_t>::vec_t;
 
+    kv_cache_t* __restrict__ curr_b_0 = b_tile;
+    kv_cache_t* __restrict__ curr_b_1 = b_tile + 16;
     float* __restrict__ curr_c_0 = c_tile;
     float* __restrict__ curr_c_1 = c_tile + 16;
 
@@ -99,14 +76,16 @@ class TileGemm82 {
     }
 
     float* __restrict__ curr_a = a_tile;
-    kv_cache_t* __restrict__ curr_b = b_tile;
-
     for (int32_t k = 0; k < dynamic_k_size; ++k) {
-      auto [fp32_b_0_reg, fp32_b_1_reg] = load_b_pair_vec(curr_b);
+      load_vec_t b_0_reg(curr_b_0);
+      vec_op::FP32Vec16 fp32_b_0_reg(b_0_reg);
+      load_vec_t b_1_reg(curr_b_1);
+      vec_op::FP32Vec16 fp32_b_1_reg(b_1_reg);
 
       float* __restrict__ curr_m_a = curr_a;
       vec_op::unroll_loop<int32_t, M>([&](int32_t i) {
-        vec_op::FP32Vec16 a_reg(*curr_m_a);
+        float v = *curr_m_a;
+        vec_op::FP32Vec16 a_reg(v);
         c_regs[i * 2] = c_regs[i * 2] + a_reg * fp32_b_0_reg;
         c_regs[i * 2 + 1] = c_regs[i * 2 + 1] + a_reg * fp32_b_1_reg;
 
@@ -116,7 +95,8 @@ class TileGemm82 {
 
       // update
       curr_a += 1;
-      curr_b += ldb;
+      curr_b_0 += ldb;
+      curr_b_1 += ldb;
     }
 
     vec_op::unroll_loop<int32_t, M>([&](int32_t i) {
@@ -129,20 +109,15 @@ class TileGemm82 {
     });
   }
 };
-
 }  // namespace
 
 // This is a general but naive implementation based on vector instructions
-template <typename scalar_t, int64_t head_dim, typename kv_cache_scalar_t>
-class AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
-  static constexpr bool fp8_kv =
-      std::is_same_v<kv_cache_scalar_t, c10::Float8_e4m3fn> ||
-      std::is_same_v<kv_cache_scalar_t, c10::Float8_e5m2>;
-
+template <typename scalar_t, int64_t head_dim>
+class AttentionImpl<ISA::VEC, scalar_t, head_dim> {
  public:
   using query_t = scalar_t;
   using q_buffer_t = float;
-  using kv_cache_t = kv_cache_scalar_t;
+  using kv_cache_t = scalar_t;
   using logits_buffer_t = float;
   using partial_output_buffer_t = float;
   using prob_buffer_t = float;
@@ -154,45 +129,11 @@ class AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
   constexpr static int64_t MaxQHeadNumPerIteration = 8;
   constexpr static int64_t HeadDim = head_dim;
   constexpr static ISA ISAType = ISA::VEC;
-  constexpr static bool scale_on_logits = fp8_kv;
-
-  float k_scale = 1.0f;
-  float v_scale = 1.0f;
+  constexpr static bool scale_on_logits = false;  // apply scale on q_buffer
 
  public:
-  void init_from_input(const AttentionInput* input) {
-    if constexpr (fp8_kv) {
-      k_scale = input->k_scale_fp8;
-      v_scale = input->v_scale_fp8;
-    }
-  }
-
-  float get_output_v_scale() const noexcept {
-    if constexpr (fp8_kv) {
-      // VEC dequant unpacks FP8 into a pseudo-FP16 layout (exponent bias 15).
-      // E4M3 (bias=7) needs correction 2^(15-7) = 2^8; E5M2 bias matches FP16
-      // so no correction.
-      if constexpr (std::is_same_v<kv_cache_t, c10::Float8_e5m2>) {
-        return v_scale;
-      } else {
-        return v_scale * 0x1p8f;
-      }
-    }
-    return 1.0f;
-  }
-
   template <template <typename tile_gemm_t> typename attention>
   FORCE_INLINE void execute_attention(DEFINE_CPU_ATTENTION_PARAMS) {
-    if constexpr (fp8_kv) {
-      // Same bias correction as get_output_v_scale: VEC FP8→pseudo-FP16 dequant
-      // uses bias 15; E4M3 (bias=7) needs ×2^8, E5M2 (bias=15) needs no
-      // correction.
-      if constexpr (std::is_same_v<kv_cache_t, c10::Float8_e5m2>) {
-        scale *= k_scale;
-      } else {
-        scale *= k_scale * 0x1p8f;
-      }
-    }
     attention<TileGemm82<kv_cache_t>> attention_iteration;
     attention_iteration(CPU_ATTENTION_PARAMS);
   }
@@ -220,19 +161,17 @@ class AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
                               // row-major
   }
 
-  // Copy q to q_buffer and cast it to fp32.
-  // FP8: QK scale is folded into execute_attention; copy Q unscaled here.
-  void copy_q_heads_tile(scalar_t* __restrict__ src,
-                         float* __restrict__ q_buffer, const int32_t q_num,
-                         const int32_t q_heads_per_kv,
-                         const int64_t q_num_stride,
-                         const int64_t q_head_stride, float scale) {
+  // Copy q to q_buffer and cast it to fp32
+  static void copy_q_heads_tile(
+      scalar_t* __restrict__ src,  // [q_num, q_heads_per_kv, head_size]
+      float* __restrict__ q_buffer, const int32_t q_num,
+      const int32_t q_heads_per_kv, const int64_t q_num_stride,
+      const int64_t q_head_stride, float scale) {
     static_assert(head_dim % 16 == 0);
     constexpr int32_t unroll_size = head_dim / 16;
     using load_vec_t = typename VecTypeTrait<scalar_t>::vec_t;
 
-    const float effective_scale = fp8_kv ? 1.0f : scale;
-    vec_op::FP32Vec16 scale_vec(effective_scale);
+    vec_op::FP32Vec16 scale_vec(scale);
     for (int32_t q_num_idx = 0; q_num_idx < q_num; ++q_num_idx) {
       for (int32_t q_head_idx = 0; q_head_idx < q_heads_per_kv; ++q_head_idx) {
         scalar_t* __restrict__ curr_q =
@@ -257,26 +196,13 @@ class AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
   // reshape K as column-major and V as row-major
   static void reshape_and_cache(
       const scalar_t* __restrict__ key, const scalar_t* __restrict__ value,
-      kv_cache_t* __restrict__ key_cache, kv_cache_t* __restrict__ value_cache,
+      scalar_t* __restrict__ key_cache, scalar_t* __restrict__ value_cache,
       const int64_t* __restrict__ slot_mapping, const int64_t token_num,
       const int64_t key_token_num_stride, const int64_t value_token_num_stride,
       const int64_t head_num, const int64_t key_head_num_stride,
       const int64_t value_head_num_stride, const int64_t num_blocks,
       const int64_t num_blocks_stride, const int64_t cache_head_num_stride,
-      const int64_t block_size, const int64_t block_size_stride,
-      const float k_inv = 0.0f, const float v_inv = 0.0f) {
-    if constexpr (fp8_kv) {
-      constexpr auto qfn = select_fp8_quant_fn<kv_cache_t>();
-      reshape_and_cache_fp8_vec_impl<scalar_t, qfn>(
-          key, value, reinterpret_cast<uint8_t*>(key_cache),
-          reinterpret_cast<uint8_t*>(value_cache), slot_mapping, token_num,
-          head_num, head_dim, block_size, key_token_num_stride,
-          key_head_num_stride, value_token_num_stride, value_head_num_stride,
-          num_blocks_stride, cache_head_num_stride, num_blocks_stride,
-          cache_head_num_stride, k_inv, v_inv);
-      return;
-    }
-
+      const int64_t block_size, const int64_t block_size_stride) {
 #pragma omp parallel for collapse(2)
     for (int64_t token_idx = 0; token_idx < token_num; ++token_idx) {
       for (int64_t head_idx = 0; head_idx < head_num; ++head_idx) {
@@ -294,9 +220,8 @@ class AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
                                           token_idx * key_token_num_stride +
                                           head_idx * key_head_num_stride;
           scalar_t* key_cache_start_ptr =
-              reinterpret_cast<scalar_t*>(key_cache) +
-              block_idx * num_blocks_stride + head_idx * cache_head_num_stride +
-              block_offset;
+              key_cache + block_idx * num_blocks_stride +
+              head_idx * cache_head_num_stride + block_offset;
 
 #pragma GCC unroll 8
           for (int64_t i = 0, j = 0; i < head_dim; ++i, j += block_size) {
@@ -309,9 +234,8 @@ class AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
                                             token_idx * value_token_num_stride +
                                             head_idx * value_head_num_stride;
           scalar_t* value_cache_start_ptr =
-              reinterpret_cast<scalar_t*>(value_cache) +
-              block_idx * num_blocks_stride + head_idx * cache_head_num_stride +
-              block_offset * head_dim;
+              value_cache + block_idx * num_blocks_stride +
+              head_idx * cache_head_num_stride + block_offset * head_dim;
           std::memcpy(value_cache_start_ptr, value_start_ptr,
                       sizeof(scalar_t) * head_dim);
         }
@@ -319,7 +243,6 @@ class AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
     }
   }
 };
-
 }  // namespace cpu_attention
 
 #endif

csrc/cpu/cpu_attn_vec16.hpp

@@ -116,9 +116,9 @@ class TileGemm161 {
 }  // namespace
 
 // This is a general but naive implementation based on vector instructions
-template <typename scalar_t, int64_t head_dim, typename kv_cache_scalar_t>
-class AttentionImpl<ISA::VEC16, scalar_t, head_dim, kv_cache_scalar_t>
-    : public AttentionImpl<ISA::VEC, scalar_t, head_dim, kv_cache_scalar_t> {
+template <typename scalar_t, int64_t head_dim>
+class AttentionImpl<ISA::VEC16, scalar_t, head_dim>
+    : public AttentionImpl<ISA::VEC, scalar_t, head_dim> {
  public:
   using query_t = scalar_t;
   using q_buffer_t = float;