diff --git a/.buildkite/hardware_tests/cpu.yaml b/.buildkite/hardware_tests/cpu.yaml
index e466e2a52444..92b0196ec2a2 100644
--- a/.buildkite/hardware_tests/cpu.yaml
+++ b/.buildkite/hardware_tests/cpu.yaml
@@ -12,15 +12,19 @@ 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 20m "
+ bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
pytest -x -v -s tests/kernels/moe/test_cpu_fused_moe.py
+ pytest -x -v -s tests/kernels/moe/test_cpu_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/test_awq_int4_to_int8.py
+ pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py"
- label: CPU-Compatibility Tests
depends_on: []
@@ -46,7 +50,7 @@ steps:
- tests/models/language/pooling/
commands:
- |
- bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
+ bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 40m "
pytest -x -v -s tests/models/language/generation -m cpu_model
pytest -x -v -s tests/models/language/pooling -m cpu_model"
@@ -61,6 +65,7 @@ steps:
- 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:
@@ -69,11 +74,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
+- label: CPU-Distributed Tests (PP+TP)
depends_on: []
device: intel_cpu
no_plugin: true
- source_file_dependencies:
+ source_file_dependencies: &cpu_distributed_deps
- csrc/cpu/shm.cpp
- vllm/v1/worker/cpu_worker.py
- vllm/v1/worker/gpu_worker.py
@@ -82,10 +87,21 @@ 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"
+ bash .buildkite/scripts/hardware_ci/run-cpu-distributed-smoke-test.sh dp_tp"
- label: CPU-Multi-Modal Model Tests %N
depends_on: []
@@ -99,7 +115,7 @@ steps:
- |
bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 45m "
pytest -x -v -s tests/models/multimodal/generation --ignore=tests/models/multimodal/generation/test_pixtral.py -m cpu_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB"
- parallelism: 2
+ parallelism: 3
- label: "Arm CPU Test"
depends_on: []
diff --git a/.buildkite/image_build/image_build.sh b/.buildkite/image_build/image_build.sh
index 9131dfc71a0a..10c03c3e1773 100755
--- a/.buildkite/image_build/image_build.sh
+++ b/.buildkite/image_build/image_build.sh
@@ -92,8 +92,8 @@ check_and_skip_if_image_exists() {
}
ecr_login() {
- aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY"
- aws ecr get-login-password --region us-east-1 | docker login --username AWS --password-stdin 936637512419.dkr.ecr.us-east-1.amazonaws.com
+ aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
+ aws ecr get-login-password --region us-east-1 | docker login --username AWS --password-stdin 936637512419.dkr.ecr.us-east-1.amazonaws.com || true
}
prepare_cache_tags() {
@@ -192,6 +192,7 @@ 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
diff --git a/.buildkite/image_build/image_build_cpu.sh b/.buildkite/image_build/image_build_cpu.sh
index ccfe155fa2b7..035f070ab891 100755
--- a/.buildkite/image_build/image_build_cpu.sh
+++ b/.buildkite/image_build/image_build_cpu.sh
@@ -11,7 +11,7 @@ REPO=$2
BUILDKITE_COMMIT=$3
# authenticate with AWS ECR
-aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY"
+aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
# skip build if image already exists
if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-cpu) ]]; then
diff --git a/.buildkite/image_build/image_build_cpu_arm64.sh b/.buildkite/image_build/image_build_cpu_arm64.sh
index ff3d11c8d599..b561e2c2e463 100755
--- a/.buildkite/image_build/image_build_cpu_arm64.sh
+++ b/.buildkite/image_build/image_build_cpu_arm64.sh
@@ -11,7 +11,7 @@ REPO=$2
BUILDKITE_COMMIT=$3
# authenticate with AWS ECR
-aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY"
+aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
# skip build if image already exists
if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64-cpu) ]]; then
diff --git a/.buildkite/image_build/image_build_torch_nightly.sh b/.buildkite/image_build/image_build_torch_nightly.sh
index a23c658d46b9..cbd08aa7bd0b 100755
--- a/.buildkite/image_build/image_build_torch_nightly.sh
+++ b/.buildkite/image_build/image_build_torch_nightly.sh
@@ -46,7 +46,7 @@ echo "Image not found, proceeding with build..."
# --- CUDA 13.0 for nightly builds ---
# Nightly CI uses CUDA 13.0 while regular CI stays on CUDA 12.9
-NIGHTLY_CUDA_VERSION="13.0.0"
+NIGHTLY_CUDA_VERSION="13.0.2"
NIGHTLY_BUILD_BASE_IMAGE="nvidia/cuda:${NIGHTLY_CUDA_VERSION}-devel-ubuntu22.04"
NIGHTLY_FINAL_BASE_IMAGE="nvidia/cuda:${NIGHTLY_CUDA_VERSION}-base-ubuntu22.04"
diff --git a/.buildkite/intel_jobs/engine_intel.yaml b/.buildkite/intel_jobs/engine_intel.yaml
new file mode 100644
index 000000000000..c66576d40991
--- /dev/null
+++ b/.buildkite/intel_jobs/engine_intel.yaml
@@ -0,0 +1,21 @@
+group: Engine Intel
+depends_on:
+ - image-build-xpu
+steps:
+- label: Engine (1 GPU)
+ timeout_in_minutes: 30
+ device: intel_gpu
+ no_plugin: true
+ working_dir: "."
+ env:
+ REGISTRY: "public.ecr.aws/q9t5s3a7"
+ REPO: "vllm-ci-test-repo"
+ VLLM_TEST_DEVICE: "xpu"
+ source_file_dependencies:
+ - vllm/v1/engine/
+ - tests/v1/engine/
+ commands:
+ - >-
+ bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+ 'cd tests &&
+ pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py'
diff --git a/.buildkite/intel_jobs/kernels_intel.yaml b/.buildkite/intel_jobs/kernels_intel.yaml
new file mode 100644
index 000000000000..66a8db25f02e
--- /dev/null
+++ b/.buildkite/intel_jobs/kernels_intel.yaml
@@ -0,0 +1,21 @@
+group: Kernels Intel
+depends_on:
+ - image-build-xpu
+steps:
+- label: vLLM IR Tests
+ timeout_in_minutes: 30
+ device: intel_gpu
+ no_plugin: true
+ working_dir: "."
+ env:
+ REGISTRY: "public.ecr.aws/q9t5s3a7"
+ REPO: "vllm-ci-test-repo"
+ VLLM_TEST_DEVICE: "xpu"
+ source_file_dependencies:
+ - vllm/ir
+ - vllm/kernels
+ commands:
+ - >-
+ bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+ 'cd tests &&
+ pytest -v -s kernels/ir'
diff --git a/.buildkite/intel_jobs/lora_intel.yaml b/.buildkite/intel_jobs/lora_intel.yaml
new file mode 100644
index 000000000000..6d5bddacf1bc
--- /dev/null
+++ b/.buildkite/intel_jobs/lora_intel.yaml
@@ -0,0 +1,135 @@
+group: LoRA Intel
+depends_on:
+ - image-build-xpu
+steps:
+- label: LoRA Runtime + Utils
+ timeout_in_minutes: 45
+ device: intel_gpu
+ no_plugin: true
+ working_dir: "."
+ env:
+ REGISTRY: "public.ecr.aws/q9t5s3a7"
+ REPO: "vllm-ci-test-repo"
+ VLLM_TEST_DEVICE: "xpu"
+ source_file_dependencies:
+ - vllm/lora
+ - tests/lora
+ commands:
+ - >-
+ 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_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'
+
+- label: LoRA Fused/MoE Kernels
+ timeout_in_minutes: 45
+ device: intel_gpu
+ no_plugin: true
+ working_dir: "."
+ env:
+ REGISTRY: "public.ecr.aws/q9t5s3a7"
+ REPO: "vllm-ci-test-repo"
+ VLLM_TEST_DEVICE: "xpu"
+ source_file_dependencies:
+ - vllm/lora
+ - tests/lora
+ commands:
+ - >-
+ 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'
+
+- label: LoRA Punica Kernels
+ timeout_in_minutes: 45
+ device: intel_gpu
+ no_plugin: true
+ working_dir: "."
+ env:
+ REGISTRY: "public.ecr.aws/q9t5s3a7"
+ REPO: "vllm-ci-test-repo"
+ VLLM_TEST_DEVICE: "xpu"
+ source_file_dependencies:
+ - vllm/lora
+ - tests/lora
+ commands:
+ - >-
+ 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]"'
+
+- label: LoRA Punica FP8/XPU Ops
+ timeout_in_minutes: 45
+ device: intel_gpu
+ no_plugin: true
+ working_dir: "."
+ env:
+ REGISTRY: "public.ecr.aws/q9t5s3a7"
+ REPO: "vllm-ci-test-repo"
+ VLLM_TEST_DEVICE: "xpu"
+ source_file_dependencies:
+ - vllm/lora
+ - tests/lora
+ commands:
+ - >-
+ 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'
+
+- label: LoRA Models
+ timeout_in_minutes: 45
+ device: intel_gpu
+ no_plugin: true
+ working_dir: "."
+ env:
+ REGISTRY: "public.ecr.aws/q9t5s3a7"
+ REPO: "vllm-ci-test-repo"
+ VLLM_TEST_DEVICE: "xpu"
+ source_file_dependencies:
+ - vllm/lora
+ - tests/lora
+ commands:
+ - >-
+ 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'
+
+- label: LoRA Multimodal
+ timeout_in_minutes: 45
+ device: intel_gpu
+ no_plugin: true
+ working_dir: "."
+ env:
+ REGISTRY: "public.ecr.aws/q9t5s3a7"
+ REPO: "vllm-ci-test-repo"
+ VLLM_TEST_DEVICE: "xpu"
+ source_file_dependencies:
+ - vllm/lora
+ - tests/lora
+ commands:
+ - >-
+ 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_whisper.py'
diff --git a/.buildkite/intel_jobs/misc_intel.yaml b/.buildkite/intel_jobs/misc_intel.yaml
new file mode 100644
index 000000000000..864128bb5338
--- /dev/null
+++ b/.buildkite/intel_jobs/misc_intel.yaml
@@ -0,0 +1,55 @@
+group: Miscellaneous Intel
+depends_on:
+ - image-build-xpu
+steps:
+- label: V1 Core + KV + Metrics
+ timeout_in_minutes: 30
+ device: intel_gpu
+ no_plugin: true
+ working_dir: "."
+ env:
+ REGISTRY: "public.ecr.aws/q9t5s3a7"
+ REPO: "vllm-ci-test-repo"
+ VLLM_TEST_DEVICE: "xpu"
+ source_file_dependencies:
+ - vllm/
+ - tests/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:
+ - >-
+ bash .buildkite/scripts/hardware_ci/run-intel-test.sh
+ 'pip install -r requirements/kv_connectors.txt &&
+ export VLLM_WORKER_MULTIPROC_METHOD=spawn &&
+ cd tests &&
+ pytest -v -s v1/executor'
+
+- label: V1 Sample + Logits
+ timeout_in_minutes: 30
+ device: intel_gpu
+ no_plugin: true
+ working_dir: "."
+ env:
+ REGISTRY: "public.ecr.aws/q9t5s3a7"
+ REPO: "vllm-ci-test-repo"
+ VLLM_TEST_DEVICE: "xpu"
+ source_file_dependencies:
+ - vllm/
+ - tests/v1/sample
+ - tests/v1/logits_processors
+ - tests/v1/test_oracle.py
+ - tests/v1/test_request.py
+ - tests/v1/test_outputs.py
+ commands:
+ - >-
+ 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/test_oracle.py &&
+ pytest -v -s v1/test_request.py &&
+ pytest -v -s v1/test_outputs.py'
diff --git a/.buildkite/intel_jobs/test-intel.yaml b/.buildkite/intel_jobs/test-intel.yaml
index c59be699502f..336f71fdbcda 100644
--- a/.buildkite/intel_jobs/test-intel.yaml
+++ b/.buildkite/intel_jobs/test-intel.yaml
@@ -61,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_tree_attention.py --ignore=v1/spec_decode/test_speculators_eagle3.py --ignore=v1/spec_decode/test_acceptance_length.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'
diff --git a/.buildkite/release-pipeline.yaml b/.buildkite/release-pipeline.yaml
index 45b2996f7ead..8a900c0bf862 100644
--- a/.buildkite/release-pipeline.yaml
+++ b/.buildkite/release-pipeline.yaml
@@ -1,3 +1,13 @@
+# CUDA architecture lists — following PyTorch RELEASE.md
+# (https://github.com/pytorch/pytorch/blob/main/RELEASE.md)
+# SM86 included for broader Ampere coverage; SM89 for marlin fp8 support
+env:
+ CUDA_ARCH_X86: "7.5 8.0 8.6 8.9 9.0 10.0 12.0+PTX"
+ # aarch64 only architectures: 8.7 for Orin, 11.0 for Thor (since CUDA 13)
+ CUDA_ARCH_AARCH64: "8.0 8.7 8.9 9.0 10.0 11.0 12.0+PTX"
+ CUDA_ARCH_X86_CU129: "7.5 8.0 8.6 8.9 9.0 10.0 12.0"
+ CUDA_ARCH_AARCH64_CU129: "8.0 8.7 8.9 9.0 10.0 12.0"
+
steps:
- input: "Provide Release version here"
id: input-release-version
@@ -14,9 +24,7 @@ steps:
agents:
queue: arm64_cpu_queue_release
commands:
- # #NOTE: torch_cuda_arch_list is derived from upstream PyTorch build files here:
- # https://github.com/pytorch/pytorch/blob/main/.ci/aarch64_linux/aarch64_ci_build.sh#L7
- - "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='8.7 8.9 9.0 10.0+PTX 12.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=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"
@@ -29,12 +37,10 @@ steps:
agents:
queue: arm64_cpu_queue_release
commands:
- # #NOTE: torch_cuda_arch_list is derived from upstream PyTorch build files here:
- # https://github.com/pytorch/pytorch/blob/main/.ci/aarch64_linux/aarch64_ci_build.sh#L7
- - "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.1 --build-arg torch_cuda_arch_list='8.7 8.9 9.0 10.0+PTX 12.0' --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu22.04 --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_OS=manylinux --build-arg BUILD_BASE_IMAGE=pytorch/manylinuxaarch64-builder:cuda13.0 --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 manylinux_2_35"
+ - "bash .buildkite/scripts/upload-nightly-wheels.sh"
env:
DOCKER_BUILDKIT: "1"
@@ -47,7 +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 manylinux_2_35"
+ - "bash .buildkite/scripts/upload-nightly-wheels.sh"
env:
DOCKER_BUILDKIT: "1"
@@ -57,10 +63,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=12.9.1 --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=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 manylinux_2_31"
+ - "bash .buildkite/scripts/upload-nightly-wheels.sh"
env:
DOCKER_BUILDKIT: "1"
@@ -70,10 +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.1 --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu22.04 --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_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 ."
- "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 manylinux_2_35"
+ - "bash .buildkite/scripts/upload-nightly-wheels.sh"
env:
DOCKER_BUILDKIT: "1"
@@ -86,7 +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 manylinux_2_35"
+ - "bash .buildkite/scripts/upload-nightly-wheels.sh"
env:
DOCKER_BUILDKIT: "1"
@@ -98,105 +104,212 @@ steps:
commands:
- "bash .buildkite/scripts/generate-and-upload-nightly-index.sh"
+ - block: "Unblock to build release Docker images"
+ depends_on: ~
+ key: block-build-release-images
+ if: build.env("NIGHTLY") != "1"
+
- group: "Build release Docker images"
key: "build-release-images"
+ depends_on: block-build-release-images
+ allow_dependency_failure: true
steps:
- - label: "Build release image - x86_64 - CUDA 12.9"
+ - label: "Build release image - x86_64 - CUDA 13.0"
depends_on: ~
id: build-release-image-x86
agents:
queue: cpu_queue_release
commands:
- "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
- - "DOCKER_BUILDKIT=1 docker build --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 FLASHINFER_AOT_COMPILE=true --build-arg INSTALL_KV_CONNECTORS=true --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m) --target vllm-openai --progress plain -f docker/Dockerfile ."
+ - |
+ 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 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"
- - label: "Build release image - aarch64 - CUDA 12.9"
+ - label: "Build release image - aarch64 - CUDA 13.0"
depends_on: ~
id: build-release-image-arm64
agents:
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 --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 FLASHINFER_AOT_COMPILE=true --build-arg torch_cuda_arch_list='8.7 8.9 9.0 10.0+PTX 12.0' --build-arg INSTALL_KV_CONNECTORS=true --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m) --target vllm-openai --progress plain -f docker/Dockerfile ."
+ - |
+ 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 push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)"
- - label: "Build release image - x86_64 - CUDA 13.0"
+ - label: "Build release image - x86_64 - CUDA 12.9"
depends_on: ~
- id: build-release-image-x86-cuda-13-0
+ id: build-release-image-x86-cuda-12-9
agents:
queue: cpu_queue_release
commands:
- "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
- - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.1 --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu22.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130 --target vllm-openai --progress plain -f docker/Dockerfile ."
- - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130"
+ - |
+ 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 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)-cu130 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130"
- - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130"
+ - "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"
- - label: "Build release image - aarch64 - CUDA 13.0"
+ - label: "Build release image - aarch64 - CUDA 12.9"
depends_on: ~
- id: build-release-image-arm64-cuda-13-0
+ id: build-release-image-arm64-cuda-12-9
agents:
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"
- # compute capability 12.0 for RTX-50 series / RTX PRO 6000 Blackwell, 12.1 for DGX Spark
- - "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.1 --build-arg torch_cuda_arch_list='8.7 8.9 9.0 10.0+PTX 12.0 12.1' --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu22.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130 --target vllm-openai --progress plain -f docker/Dockerfile ."
- - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130"
+ - |
+ 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 push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129"
- - label: "Build release image - x86_64 - CUDA 12.9 - Ubuntu 24.04"
+ - label: "Build release image - x86_64 - CUDA 13.0 - Ubuntu 24.04"
depends_on: ~
id: build-release-image-x86-ubuntu2404
agents:
queue: cpu_queue_release
commands:
- "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
- - "DOCKER_BUILDKIT=1 docker build --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 FLASHINFER_AOT_COMPILE=true --build-arg torch_cuda_arch_list='8.7 8.9 9.0 10.0+PTX 12.0' --build-arg INSTALL_KV_CONNECTORS=true --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
+ - |
+ 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 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"
- - label: "Build release image - aarch64 - CUDA 12.9 - Ubuntu 24.04"
+ - label: "Build release image - aarch64 - CUDA 13.0 - Ubuntu 24.04"
depends_on: ~
id: build-release-image-arm64-ubuntu2404
agents:
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 --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 FLASHINFER_AOT_COMPILE=true --build-arg torch_cuda_arch_list='8.7 8.9 9.0 10.0+PTX 12.0' --build-arg INSTALL_KV_CONNECTORS=true --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
+ - |
+ 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 push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-ubuntu2404"
- - label: "Build release image - x86_64 - CUDA 13.0 - Ubuntu 24.04"
+ - label: "Build release image - x86_64 - CUDA 12.9 - Ubuntu 24.04"
depends_on: ~
- id: build-release-image-x86-cuda-13-0-ubuntu2404
+ id: build-release-image-x86-cuda-12-9-ubuntu2404
agents:
queue: cpu_queue_release
commands:
- "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
- - "DOCKER_BUILDKIT=1 docker build --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.1 --build-arg UBUNTU_VERSION=24.04 --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 --build-arg FLASHINFER_AOT_COMPILE=true --build-arg torch_cuda_arch_list='8.7 8.9 9.0 10.0+PTX 12.0 12.1' --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu24.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
- - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130-ubuntu2404"
- - "docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130-ubuntu2404"
- - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130-ubuntu2404"
+ - |
+ 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 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"
- - label: "Build release image - aarch64 - CUDA 13.0 - Ubuntu 24.04"
+ - label: "Build release image - aarch64 - CUDA 12.9 - Ubuntu 24.04"
depends_on: ~
- id: build-release-image-arm64-cuda-13-0-ubuntu2404
+ id: build-release-image-arm64-cuda-12-9-ubuntu2404
agents:
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 --build-arg max_jobs=16 --build-arg USE_SCCACHE=1 --build-arg GIT_REPO_CHECK=1 --build-arg CUDA_VERSION=13.0.1 --build-arg UBUNTU_VERSION=24.04 --build-arg GDRCOPY_OS_VERSION=Ubuntu24_04 --build-arg FLASHINFER_AOT_COMPILE=true --build-arg torch_cuda_arch_list='8.7 8.9 9.0 10.0+PTX 12.0 12.1' --build-arg INSTALL_KV_CONNECTORS=true --build-arg BUILD_BASE_IMAGE=nvidia/cuda:13.0.1-devel-ubuntu24.04 --tag public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130-ubuntu2404 --target vllm-openai --progress plain -f docker/Dockerfile ."
- - "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu130-ubuntu2404"
+ - |
+ 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 push 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
@@ -215,7 +328,8 @@ steps:
depends_on: ~
- label: "Build release image - arm64 - CPU"
- depends_on:
+ key: build-cpu-release-image-arm64
+ depends_on:
- block-arm64-cpu-release-image-build
- input-release-version
agents:
@@ -231,7 +345,7 @@ steps:
- group: "Publish release images"
key: "publish-release-images"
steps:
- - label: "Create multi-arch manifest - CUDA 12.9"
+ - label: "Create multi-arch manifest - CUDA 13.0"
depends_on:
- build-release-image-x86
- build-release-image-arm64
@@ -243,7 +357,7 @@ steps:
- "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"
- - label: "Annotate release workflow - CUDA 12.9"
+ - label: "Annotate release workflow - CUDA 13.0"
depends_on:
- create-multi-arch-manifest
id: annotate-release-workflow
@@ -252,19 +366,19 @@ steps:
commands:
- "bash .buildkite/scripts/annotate-release.sh"
- - label: "Create multi-arch manifest - CUDA 13.0"
+ - label: "Create multi-arch manifest - CUDA 12.9"
depends_on:
- - build-release-image-x86-cuda-13-0
- - build-release-image-arm64-cuda-13-0
- id: create-multi-arch-manifest-cuda-13-0
+ - build-release-image-x86-cuda-12-9
+ - build-release-image-arm64-cuda-12-9
+ id: create-multi-arch-manifest-cuda-12-9
agents:
queue: small_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 manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-cu130 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-cu130 --amend"
- - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130"
+ - "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"
- - label: "Create multi-arch manifest - CUDA 12.9 - Ubuntu 24.04"
+ - label: "Create multi-arch manifest - CUDA 13.0 - Ubuntu 24.04"
depends_on:
- build-release-image-x86-ubuntu2404
- build-release-image-arm64-ubuntu2404
@@ -276,17 +390,17 @@ steps:
- "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"
- - label: "Create multi-arch manifest - CUDA 13.0 - Ubuntu 24.04"
+ - label: "Create multi-arch manifest - CUDA 12.9 - Ubuntu 24.04"
depends_on:
- - build-release-image-x86-cuda-13-0-ubuntu2404
- - build-release-image-arm64-cuda-13-0-ubuntu2404
- id: create-multi-arch-manifest-cuda-13-0-ubuntu2404
+ - build-release-image-x86-cuda-12-9-ubuntu2404
+ - build-release-image-arm64-cuda-12-9-ubuntu2404
+ id: create-multi-arch-manifest-cuda-12-9-ubuntu2404
agents:
queue: small_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 manifest create public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-x86_64-cu130-ubuntu2404 public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-aarch64-cu130-ubuntu2404 --amend"
- - "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu130-ubuntu2404"
+ - "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"
- label: "Publish nightly multi-arch image to DockerHub"
depends_on:
@@ -306,16 +420,51 @@ steps:
DOCKER_BUILDKIT: "1"
DOCKERHUB_USERNAME: "vllmbot"
- - label: "Publish nightly multi-arch image to DockerHub - CUDA 13.0"
+ - label: "Publish nightly multi-arch image to DockerHub - CUDA 12.9"
depends_on:
- - create-multi-arch-manifest-cuda-13-0
+ - create-multi-arch-manifest-cuda-12-9
if: build.env("NIGHTLY") == "1"
agents:
queue: small_cpu_queue_release
commands:
- - "bash .buildkite/scripts/push-nightly-builds.sh cu130"
+ - "bash .buildkite/scripts/push-nightly-builds.sh cu129"
# Clean up old nightly builds (keep only last 14)
- - "bash .buildkite/scripts/cleanup-nightly-builds.sh cu130-nightly-"
+ - "bash .buildkite/scripts/cleanup-nightly-builds.sh cu129-nightly-"
+ plugins:
+ - docker-login#v3.0.0:
+ username: vllmbot
+ password-env: DOCKERHUB_TOKEN
+ env:
+ DOCKER_BUILDKIT: "1"
+ DOCKERHUB_USERNAME: "vllmbot"
+
+ - 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
@@ -611,12 +760,14 @@ steps:
- "bash tools/vllm-rocm/generate-rocm-wheels-root-index.sh"
env:
S3_BUCKET: "vllm-wheels"
- VARIANT: "rocm721"
+ VARIANT: "rocm722"
# ROCm Job 6: Build ROCm Release Docker Image
- label: ":docker: Build release image - x86_64 - ROCm"
id: build-rocm-release-image
depends_on:
+ - step: block-build-release-images
+ allow_failure: true
- step: build-rocm-base-wheels
allow_failure: false
agents:
diff --git a/.buildkite/scripts/annotate-release.sh b/.buildkite/scripts/annotate-release.sh
index 2da9db2f2e5d..afa884fba46b 100755
--- a/.buildkite/scripts/annotate-release.sh
+++ b/.buildkite/scripts/annotate-release.sh
@@ -8,112 +8,20 @@ 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_31_x86_64.whl .
-aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}-cp38-abi3-manylinux_2_31_aarch64.whl .
+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 13.0:
-aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}+cu130-cp38-abi3-manylinux_2_35_x86_64.whl .
-aws s3 cp s3://vllm-wheels/${BUILDKITE_COMMIT}/vllm-${RELEASE_VERSION}+cu130-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-cu130
-docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${BUILDKITE_COMMIT}-aarch64-cu130
-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-cu130 vllm/vllm-openai:x86_64-cu130
-docker tag vllm/vllm-openai:x86_64-cu130 vllm/vllm-openai:latest-x86_64-cu130
-docker tag vllm/vllm-openai:x86_64-cu130 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu130
-docker push vllm/vllm-openai:latest-x86_64-cu130
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu130
-
-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-cu130 vllm/vllm-openai:aarch64-cu130
-docker tag vllm/vllm-openai:aarch64-cu130 vllm/vllm-openai:latest-aarch64-cu130
-docker tag vllm/vllm-openai:aarch64-cu130 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu130
-docker push vllm/vllm-openai:latest-aarch64-cu130
-docker push vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu130
-
-## 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-cu130
-docker manifest create vllm/vllm-openai:latest-cu130 vllm/vllm-openai:latest-x86_64-cu130 vllm/vllm-openai:latest-aarch64-cu130
-docker manifest create vllm/vllm-openai:v${RELEASE_VERSION}-cu130 vllm/vllm-openai:v${RELEASE_VERSION}-x86_64-cu130 vllm/vllm-openai:v${RELEASE_VERSION}-aarch64-cu130
-docker manifest push vllm/vllm-openai:latest-cu130
-docker manifest push vllm/vllm-openai:v${RELEASE_VERSION}-cu130
-
-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}
-\`\`\`
+Docker images are published automatically by the "Publish release images to DockerHub" pipeline step.
EOF
diff --git a/.buildkite/scripts/check-ray-compatibility.sh b/.buildkite/scripts/check-ray-compatibility.sh
index 1572fe94168d..b056d4403dbb 100644
--- a/.buildkite/scripts/check-ray-compatibility.sh
+++ b/.buildkite/scripts/check-ray-compatibility.sh
@@ -29,7 +29,7 @@ if python3 -c "import torch; assert torch.version.hip" 2>/dev/null; then
TORCH_INDEX_URL=""
fi
else
- TORCH_INDEX_URL="https://download.pytorch.org/whl/cu129"
+ TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
fi
echo ">>> Using PyTorch index: ${TORCH_INDEX_URL:-PyPI default}"
diff --git a/.buildkite/scripts/ci-fetch-log.sh b/.buildkite/scripts/ci-fetch-log.sh
new file mode 100755
index 000000000000..02798b56f4a9
--- /dev/null
+++ b/.buildkite/scripts/ci-fetch-log.sh
@@ -0,0 +1,55 @@
+#!/bin/bash
+# Usage: ./ci-fetch-log.sh [output_file]
+# ./ci-fetch-log.sh [output_file]
+#
+# Downloads the raw log for a Buildkite job from the public, unauthenticated
+# /organizations//pipelines//builds//jobs//download
+# endpoint, then strips ANSI/timestamps via ci-clean-log.sh.
+#
+# Find and via:
+# gh pr checks --repo vllm-project/vllm
+# Each failing row's URL is .../builds/#.
+
+set -euo pipefail
+
+ORG="vllm"
+PIPELINE="ci"
+
+usage() {
+ echo "Usage: $0 [output_file]"
+ echo " $0 [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"
diff --git a/.buildkite/scripts/detect-manylinux-tag.py b/.buildkite/scripts/detect-manylinux-tag.py
new file mode 100644
index 000000000000..40fa6c6ffbb7
--- /dev/null
+++ b/.buildkite/scripts/detect-manylinux-tag.py
@@ -0,0 +1,142 @@
+#!/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_`` platform
+tag, which installers like ``pip`` won't accept off PyPI/our index. We need to
+rewrite the platform tag to the appropriate ``manylinux___``
+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
+
+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_``, 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())
diff --git a/.buildkite/scripts/docker-build-metadata-args.sh b/.buildkite/scripts/docker-build-metadata-args.sh
new file mode 100644
index 000000000000..9aa6fa9314f7
--- /dev/null
+++ b/.buildkite/scripts/docker-build-metadata-args.sh
@@ -0,0 +1,54 @@
+#!/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}"
diff --git a/.buildkite/scripts/generate-and-upload-nightly-index.sh b/.buildkite/scripts/generate-and-upload-nightly-index.sh
index 7cef252c6078..502ed0609310 100755
--- a/.buildkite/scripts/generate-and-upload-nightly-index.sh
+++ b/.buildkite/scripts/generate-and-upload-nightly-index.sh
@@ -9,21 +9,14 @@ set -ex
BUCKET="vllm-wheels"
INDICES_OUTPUT_DIR="indices"
-DEFAULT_VARIANT_ALIAS="cu129" # align with vLLM_MAIN_CUDA_VERSION in vllm/envs.py
-PYTHON="${PYTHON_PROG:-python3}" # try to read from env var, otherwise use python3
+DEFAULT_VARIANT_ALIAS="cu130" # align with vLLM_MAIN_CUDA_VERSION in vllm/envs.py
SUBPATH=$BUILDKITE_COMMIT
S3_COMMIT_PREFIX="s3://$BUCKET/$SUBPATH/"
-# 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)"
+# 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
# ======== generate and upload indices ========
diff --git a/.buildkite/scripts/hardware_ci/run-amd-test.sh b/.buildkite/scripts/hardware_ci/run-amd-test.sh
index 703a7d753220..7e8ddb12ec98 100755
--- a/.buildkite/scripts/hardware_ci/run-amd-test.sh
+++ b/.buildkite/scripts/hardware_ci/run-amd-test.sh
@@ -378,9 +378,11 @@ 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
@@ -397,9 +399,15 @@ fi
echo "Raw commands: $commands"
-# Fix quoting before ROCm overrides (so overrides see correct structure)
-commands=$(re_quote_pytest_markers "$commands")
-echo "After re-quoting: $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
commands=$(apply_rocm_test_overrides "$commands")
echo "Final commands: $commands"
diff --git a/.buildkite/scripts/hardware_ci/run-cpu-distributed-smoke-test.sh b/.buildkite/scripts/hardware_ci/run-cpu-distributed-smoke-test.sh
index 4ff15067aa0f..8ac27ed6583a 100644
--- a/.buildkite/scripts/hardware_ci/run-cpu-distributed-smoke-test.sh
+++ b/.buildkite/scripts/hardware_ci/run-cpu-distributed-smoke-test.sh
@@ -3,42 +3,37 @@ set -euox pipefail
export VLLM_CPU_CI_ENV=0
export VLLM_CPU_KVCACHE_SPACE=1 # avoid OOM
-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
+MODE=${1:-all}
-#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
+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
diff --git a/.buildkite/scripts/hardware_ci/run-cpu-test-arm.sh b/.buildkite/scripts/hardware_ci/run-cpu-test-arm.sh
index c2509a07b2c4..9c13fa79fcb2 100755
--- a/.buildkite/scripts/hardware_ci/run-cpu-test-arm.sh
+++ b/.buildkite/scripts/hardware_ci/run-cpu-test-arm.sh
@@ -31,6 +31,21 @@ 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
@@ -46,12 +61,6 @@ 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/moe/test_moe.py -k test_cpu_fused_moe_basic"
# basic online serving
docker exec cpu-test bash -c '
@@ -66,6 +75,21 @@ 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.
diff --git a/.buildkite/scripts/hardware_ci/run-cpu-test.sh b/.buildkite/scripts/hardware_ci/run-cpu-test.sh
index db75ad3083b2..27ec0068668f 100644
--- a/.buildkite/scripts/hardware_ci/run-cpu-test.sh
+++ b/.buildkite/scripts/hardware_ci/run-cpu-test.sh
@@ -16,5 +16,5 @@ echo "--- :docker: Building Docker image"
docker build --progress plain --tag "$IMAGE_NAME" --target vllm-test -f docker/Dockerfile.cpu .
# Run the image, setting --shm-size=4g for tensor parallel.
-docker run --rm --cpuset-cpus="$CORE_RANGE" --cpuset-mems="$NUMA_NODE" -v ~/.cache/huggingface:/root/.cache/huggingface --privileged=true -e HF_TOKEN -e VLLM_CPU_KVCACHE_SPACE=16 -e VLLM_CPU_CI_ENV=1 -e VLLM_CPU_SIM_MULTI_NUMA=1 --shm-size=4g "$IMAGE_NAME" \
+docker run --rm --cpuset-cpus="$CORE_RANGE" --cpuset-mems="$NUMA_NODE" -v ~/.cache/huggingface:/root/.cache/huggingface --privileged=true -e HF_TOKEN -e VLLM_CPU_KVCACHE_SPACE=16 -e VLLM_CPU_CI_ENV=1 -e VLLM_CPU_SIM_MULTI_NUMA=1 -e VLLM_CPU_ATTN_SPLIT_KV=0 --shm-size=4g "$IMAGE_NAME" \
timeout "$TIMEOUT_VAL" bash -c "set -euox pipefail; echo \"--- Print packages\"; pip list; echo \"--- Running tests\"; ${TEST_COMMAND}"
diff --git a/.buildkite/scripts/hardware_ci/run-intel-test.sh b/.buildkite/scripts/hardware_ci/run-intel-test.sh
index 836f0b849573..eae3b231b4be 100755
--- a/.buildkite/scripts/hardware_ci/run-intel-test.sh
+++ b/.buildkite/scripts/hardware_ci/run-intel-test.sh
@@ -25,22 +25,100 @@ export PYTHONPATH=".."
###############################################################################
cleanup_docker() {
+ # Share the same lock with image pull to avoid cleanup/pull races on one node.
+ local docker_lock="/tmp/docker-pull.lock"
+ exec 9>"$docker_lock"
+ flock 9
+
docker_root=$(docker info -f '{{.DockerRootDir}}')
if [ -z "$docker_root" ]; then
echo "Failed to determine Docker root directory." >&2
- exit 1
+ flock -u 9
+ return 1
fi
echo "Docker root directory: $docker_root"
disk_usage=$(df "$docker_root" | tail -1 | awk '{print $5}' | sed 's/%//')
threshold=70
if [ "$disk_usage" -gt "$threshold" ]; then
- echo "Disk usage is above $threshold%. Cleaning up Docker images and volumes..."
- docker image prune -f
- docker volume prune -f && docker system prune --force --filter "until=72h" --all
- echo "Docker images and volumes cleanup completed."
+ echo "Disk usage is above $threshold%. Running aggressive CI image cleanup..."
+ cleanup_old_ci_images "${REGISTRY}/${REPO}" "${image_name}" "${DOCKER_IMAGE_CLEANUP_HOURS:-72}" 1
+ else
+ echo "Disk usage is below $threshold%. Checking old CI images anyway."
+ cleanup_old_ci_images "${REGISTRY}/${REPO}" "${image_name}" "${DOCKER_IMAGE_CLEANUP_HOURS:-72}" 0
+ fi
+ echo "Old CI image cleanup completed."
+
+ flock -u 9
+}
+
+cleanup_old_ci_images() {
+ local repo_prefix="$1"
+ local current_image_ref="$2"
+ local ttl_hours="$3"
+ local aggressive_cleanup="$4"
+
+ if [[ -z "$repo_prefix" || "$repo_prefix" == "/" ]]; then
+ echo "Skip old-image cleanup: invalid repo prefix '${repo_prefix}'"
+ return 0
+ fi
+
+ if ! [[ "$ttl_hours" =~ ^[0-9]+$ ]]; then
+ echo "Invalid DOCKER_IMAGE_CLEANUP_HOURS='${ttl_hours}', fallback to 72"
+ ttl_hours=72
+ fi
+
+ local now_epoch cutoff_epoch
+ now_epoch=$(date +%s)
+ cutoff_epoch=$((now_epoch - ttl_hours * 3600))
+
+ local -a used_image_ids
+ mapfile -t used_image_ids < <(docker ps -aq | xargs -r docker inspect --format '{{.Image}}' | sort -u)
+
+ local removed_count=0
+ local examined_count=0
+ declare -A seen_ids=()
+
+ while read -r image_ref image_id; do
+ [[ -z "$image_ref" || -z "$image_id" ]] && continue
+ ((examined_count++))
+
+ # Keep the image this job is going to use.
+ if [[ "$image_ref" == "$current_image_ref" ]]; then
+ continue
+ fi
+
+ # Avoid duplicate deletes when multiple tags point to same image id.
+ if [[ -n "${seen_ids[$image_id]:-}" ]]; then
+ continue
+ fi
+ seen_ids[$image_id]=1
+
+ # Never delete images that are used by any container on this node.
+ if printf '%s\n' "${used_image_ids[@]}" | grep -qx "$image_id"; then
+ continue
+ fi
+
+ local created created_epoch
+ created=$(docker image inspect -f '{{.Created}}' "$image_id" 2>/dev/null || true)
+ [[ -z "$created" ]] && continue
+ created_epoch=$(date -d "$created" +%s 2>/dev/null || true)
+ [[ -z "$created_epoch" ]] && continue
+
+ if (( created_epoch < cutoff_epoch )) || [[ "$aggressive_cleanup" == "1" ]]; then
+ if docker image rm -f "$image_id" >/dev/null 2>&1; then
+ ((removed_count++))
+ fi
+ fi
+ done < <(docker image ls --no-trunc "$repo_prefix" --format '{{.Repository}}:{{.Tag}} {{.ID}}')
+
+ # Also trim old dangling layers; this is safe and does not remove referenced images.
+ docker image prune -f --filter "until=${ttl_hours}h" >/dev/null 2>&1 || true
+
+ if [[ "$aggressive_cleanup" == "1" ]]; then
+ echo "Examined ${examined_count} images under ${repo_prefix}, removed ${removed_count} unused images under disk pressure."
else
- echo "Disk usage is below $threshold%. No cleanup needed."
+ echo "Examined ${examined_count} images under ${repo_prefix}, removed ${removed_count} old images (>${ttl_hours}h)."
fi
}
@@ -240,7 +318,6 @@ fi
cleanup_docker
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY"
-aws ecr get-login-password --region us-east-1 | docker login --username AWS --password-stdin 936637512419.dkr.ecr.us-east-1.amazonaws.com
# --- Build or pull test image ---
IMAGE="${IMAGE_TAG_XPU:-${image_name}}"
@@ -266,8 +343,6 @@ fi
remove_docker_container() {
docker rm -f "${container_name}" || true
- docker image rm -f "${image_name}" || true
- docker system prune -f || true
}
trap remove_docker_container EXIT
@@ -283,6 +358,7 @@ docker run \
--ipc=host \
--privileged \
-v /dev/dri/by-path:/dev/dri/by-path \
+ -v ${HOME}/.cache/huggingface:/root/.cache/huggingface \
--entrypoint="" \
-e "HF_TOKEN=${HF_TOKEN:-}" \
-e "ZE_AFFINITY_MASK=${ZE_AFFINITY_MASK:-}" \
diff --git a/.buildkite/scripts/hardware_ci/run-tpu-v1-test.sh b/.buildkite/scripts/hardware_ci/run-tpu-v1-test.sh
index feaf2b356267..61ebddf82e40 100755
--- a/.buildkite/scripts/hardware_ci/run-tpu-v1-test.sh
+++ b/.buildkite/scripts/hardware_ci/run-tpu-v1-test.sh
@@ -136,8 +136,6 @@ 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" \
diff --git a/.buildkite/scripts/hardware_ci/run-xpu-test.sh b/.buildkite/scripts/hardware_ci/run-xpu-test.sh
index 6579810e9826..14bd08cfc1c4 100644
--- a/.buildkite/scripts/hardware_ci/run-xpu-test.sh
+++ b/.buildkite/scripts/hardware_ci/run-xpu-test.sh
@@ -12,9 +12,7 @@ docker build -t "${image_name}" -f docker/Dockerfile.xpu .
# Setup cleanup
remove_docker_container() {
- docker rm -f "${container_name}" || true;
- docker image rm -f "${image_name}" || true;
- docker system prune -f || true;
+ docker rm -f "${container_name}" || true
}
trap remove_docker_container EXIT
diff --git a/.buildkite/scripts/lib/manylinux.sh b/.buildkite/scripts/lib/manylinux.sh
new file mode 100644
index 000000000000..bde2dfe0a3dc
--- /dev/null
+++ b/.buildkite/scripts/lib/manylinux.sh
@@ -0,0 +1,127 @@
+#!/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_`` to the correct ``manylinux___``.
+# After sourcing, call ``apply_manylinux_tag `` 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"
+}
diff --git a/.buildkite/scripts/lib/select-python.sh b/.buildkite/scripts/lib/select-python.sh
new file mode 100644
index 000000000000..bc53030a2b50
--- /dev/null
+++ b/.buildkite/scripts/lib/select-python.sh
@@ -0,0 +1,41 @@
+#!/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)"
+}
diff --git a/.buildkite/scripts/publish-release-images.sh b/.buildkite/scripts/publish-release-images.sh
new file mode 100755
index 000000000000..ec319aa76006
--- /dev/null
+++ b/.buildkite/scripts/publish-release-images.sh
@@ -0,0 +1,180 @@
+#!/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}"
diff --git a/.buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_prefetch_offload.sh b/.buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_prefetch_offload.sh
index de48eb282a65..0eadfa1f80b4 100755
--- a/.buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_prefetch_offload.sh
+++ b/.buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_prefetch_offload.sh
@@ -51,6 +51,7 @@ 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=$!
diff --git a/.buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_dp4_async_eplb.sh b/.buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_dp4_async_eplb.sh
new file mode 100755
index 000000000000..06743f16b687
--- /dev/null
+++ b/.buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_dp4_async_eplb.sh
@@ -0,0 +1,55 @@
+#!/usr/bin/env bash
+set -euxo pipefail
+
+# args: [THRESHOLD] [NUM_QUESTIONS] [START_PORT]
+THRESHOLD=${1:-0.8}
+NUM_Q=${2:-1319}
+PORT=${3:-8050}
+OUT_DIR=${OUT_DIR:-/tmp/vllm-scheduled}
+mkdir -p "${OUT_DIR}"
+
+wait_for_server() {
+ local port=$1
+ timeout 600 bash -c '
+ until curl -sf "http://127.0.0.1:'"$port"'/health" > /dev/null; do
+ sleep 1
+ done'
+}
+
+MODEL="Qwen/Qwen3-30B-A3B-FP8"
+BACK="allgather_reducescatter"
+
+cleanup() {
+ if [[ -n "${SERVER_PID:-}" ]] && kill -0 "${SERVER_PID}" 2>/dev/null; then
+ kill "${SERVER_PID}" 2>/dev/null || true
+ for _ in {1..20}; do
+ kill -0 "${SERVER_PID}" 2>/dev/null || break
+ sleep 0.5
+ done
+ kill -9 "${SERVER_PID}" 2>/dev/null || true
+ fi
+}
+trap cleanup EXIT
+
+VLLM_DEEP_GEMM_WARMUP=skip \
+vllm serve "$MODEL" \
+--enforce-eager \
+--data-parallel-size 4 \
+--enable-expert-parallel \
+--enable-eplb \
+--all2all-backend "$BACK" \
+--eplb-config '{"window_size":20, "step_interval":100, "use_async":true}' \
+--trust-remote-code \
+--max-model-len 2048 \
+--port "$PORT" &
+SERVER_PID=$!
+wait_for_server "$PORT"
+
+TAG=$(echo "$MODEL" | tr '/: \\n' '_____')
+OUT="${OUT_DIR}/${TAG}_${BACK}.json"
+python3 tests/evals/gsm8k/gsm8k_eval.py --host http://127.0.0.1 --port "$PORT" --num-questions "${NUM_Q}" --save-results "${OUT}"
+python3 - <= ${THRESHOLD}, f"${MODEL} ${BACK} accuracy {acc}"
+PY
diff --git a/.buildkite/scripts/tool_call/run-bfcl-eval.sh b/.buildkite/scripts/tool_call/run-bfcl-eval.sh
index f3e5009e6fe3..3748cab62c7c 100755
--- a/.buildkite/scripts/tool_call/run-bfcl-eval.sh
+++ b/.buildkite/scripts/tool_call/run-bfcl-eval.sh
@@ -28,6 +28,7 @@
# 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
@@ -43,6 +44,7 @@ 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
@@ -139,7 +141,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" "$OUTPUT_DIR" << 'PYEOF' || bfcl_exit_code=$?
+python3 - "$MODEL" "$TEST_CATEGORY" "$NUM_THREADS" "$PORT" "$API_TYPE" "$TEMPERATURE" "$OUTPUT_DIR" << 'PYEOF' || bfcl_exit_code=$?
import os
import sys
@@ -148,7 +150,8 @@ test_category = sys.argv[2]
num_threads = int(sys.argv[3])
port = sys.argv[4]
api_type = sys.argv[5]
-output_dir = sys.argv[6] if len(sys.argv) > 6 and sys.argv[6] else os.getcwd()
+temperature = float(sys.argv[6])
+output_dir = sys.argv[7] if len(sys.argv) > 7 and sys.argv[7] else os.getcwd()
os.environ["OPENAI_BASE_URL"] = f"http://localhost:{port}/v1"
os.environ["OPENAI_API_KEY"] = "dummy"
@@ -204,6 +207,7 @@ 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 ----
diff --git a/.buildkite/scripts/upload-nightly-wheels.sh b/.buildkite/scripts/upload-nightly-wheels.sh
index cc72cda7d505..8cef31908809 100644
--- a/.buildkite/scripts/upload-nightly-wheels.sh
+++ b/.buildkite/scripts/upload-nightly-wheels.sh
@@ -2,14 +2,18 @@
set -ex
-# Upload a single wheel to S3 (rename linux -> manylinux).
+# Upload a single wheel to S3, after detecting and applying the appropriate
+# manylinux platform tag with auditwheel.
# 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/"
-# ========= collect, rename & upload the wheel ==========
+# ========= locate the wheel ==========
# Assume wheels are in artifacts/dist/*.whl
wheel_files=(artifacts/dist/*.whl)
@@ -21,19 +25,9 @@ if [[ ${#wheel_files[@]} -ne 1 ]]; then
fi
wheel="${wheel_files[0]}"
-# 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}"
+# ========= detect manylinux tag and rename ==========
-# 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"
+wheel="$(apply_manylinux_tag "$wheel")"
echo "Renamed wheel to: $wheel"
# Extract the version from the wheel
diff --git a/.buildkite/scripts/upload-release-wheels-pypi.sh b/.buildkite/scripts/upload-release-wheels-pypi.sh
index 058e5bbe4f4c..7e2077a2692c 100644
--- a/.buildkite/scripts/upload-release-wheels-pypi.sh
+++ b/.buildkite/scripts/upload-release-wheels-pypi.sh
@@ -39,10 +39,11 @@ fi
set -x # avoid printing secrets above
-# install twine from pypi
+# install twine and sdist build prerequisites from pypi
python3 -m venv /tmp/vllm-release-env
source /tmp/vllm-release-env/bin/activate
pip install twine
+pip install -r requirements/build/cuda.txt
python3 -m twine --version
# copy release wheels to local directory
diff --git a/.buildkite/scripts/upload-rocm-wheels.sh b/.buildkite/scripts/upload-rocm-wheels.sh
index a42848a16ffe..1f3655631204 100755
--- a/.buildkite/scripts/upload-rocm-wheels.sh
+++ b/.buildkite/scripts/upload-rocm-wheels.sh
@@ -20,10 +20,6 @@ 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"
@@ -34,19 +30,21 @@ echo "Commit: $BUILDKITE_COMMIT"
echo "Branch: $BUILDKITE_BRANCH"
echo "========================================"
-# ======== Part 0: Setup Python ========
+# ======== Part 0: Setup Python and helpers ========
-# 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
+# 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
-echo "Using python interpreter: $PYTHON"
-echo "Python version: $($PYTHON --version)"
+# 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
# ======== Part 1: Collect and prepare wheels ========
@@ -63,11 +61,18 @@ if [ "$WHEEL_COUNT" -eq 0 ]; then
exit 1
fi
-# Rename linux to manylinux in wheel filenames
+# Detect the appropriate manylinux platform tag for any wheel that still
+# carries the generic ``linux_`` 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.
for wheel in all-rocm-wheels/*.whl; do
if [[ "$wheel" == *"linux"* ]] && [[ "$wheel" != *"manylinux"* ]]; then
- new_wheel="${wheel/linux/$MANYLINUX_VERSION}"
- mv -- "$wheel" "$new_wheel"
+ new_wheel="$(apply_manylinux_tag "$wheel")"
echo "Renamed: $(basename "$wheel") -> $(basename "$new_wheel")"
fi
done
diff --git a/.buildkite/test-amd.yaml b/.buildkite/test-amd.yaml
index dda5d4064c3d..c5369950fb30 100644
--- a/.buildkite/test-amd.yaml
+++ b/.buildkite/test-amd.yaml
@@ -1,8 +1,8 @@
# In this file, you can add more tests to run either by adding a new step or
# adding a new command to an existing step. See different options here for examples.
-# This script will be feed into Jinja template in `test-template-aws.j2` at
-# https://github.com/vllm-project/buildkite-ci/blob/main/scripts/test-template-aws.j2
+# This script will be feed into Jinja template in `test-template-amd.j2` at
+# https://github.com/vllm-project/buildkite-ci/blob/main/scripts/test-template-amd.j2
# to generate the final pipeline yaml file.
# Documentation
@@ -39,8 +39,8 @@
#####################################################################################################################################
# #
# IMPORTANT: #
-# * Currently AMD CI has MI250 agents, MI325 agents, and MI355 agents. All upcoming feature improvements are tracked in: #
-# https://github.com/vllm-project/vllm/issues/34994 #
+# * Currently AMD CI has MI250 agents, MI300 agents, MI325 agents, and MI355 agents. All upcoming feature improvements are #
+# tracked in: https://github.com/vllm-project/vllm/issues/34994 #
# #
#-----------------------------------------------------------------------------------------------------------------------------------#
# #
@@ -104,193 +104,193 @@
#####################################################################################################################################
-
-
steps:
+#########################################################################################################################################
+# #
+# MI250 (gfx90a) tests #
+# #
+#########################################################################################################################################
-#####################################################################################################################################
-# #
-# MI250 test definitions ( currently the test set is completely mirrored // TBD which tests are to be routed there ultimately) #
-# #
-#####################################################################################################################################
+#----------------------------------------------------- mi250 · basic_correctness -----------------------------------------------------#
-- label: Pytorch Nightly Dependency Override Check # TBD
+- label: Distributed Model Tests (2 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- optional: true
- soft_fail: true
+ agent_pool: mi250_2
+ num_gpus: 2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - requirements/test/nightly-torch.txt
+ - vllm/model_executor/model_loader/sharded_state_loader.py
+ - vllm/model_executor/models/
+ - vllm/model_executor/layers/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - tests/basic_correctness/
+ - tests/model_executor/model_loader/test_sharded_state_loader.py
+ - tests/models/
+ - vllm/_aiter_ops.py
- vllm/platforms/rocm.py
commands:
- - bash standalone_tests/pytorch_nightly_dependency.sh
+ - TARGET_TEST_SUITE=L4 pytest basic_correctness/ -v -s -m 'distributed(num_gpus=2)'
+ - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s model_executor/model_loader/test_sharded_state_loader.py -m '(not slow_test)'
+ - pytest models/test_transformers.py -v -s -m 'distributed(num_gpus=2)'
+ - pytest models/language -v -s -m 'distributed(num_gpus=2)'
+ - pytest models/multimodal -v -s -m 'distributed(num_gpus=2)' --ignore models/multimodal/generation/test_whisper.py
+ - VLLM_WORKER_MULTIPROC_METHOD=spawn pytest models/multimodal/generation/test_whisper.py -v -s -m 'distributed(num_gpus=2)'
+#-------------------------------------------------------- mi250 · benchmarks ---------------------------------------------------------#
-- label: Async Engine, Inputs, Utils, Worker # TBD
+- label: Benchmarks # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
- working_dir: "/vllm-workspace/tests"
+ working_dir: "/vllm-workspace/.buildkite"
source_file_dependencies:
- - vllm/
- - tests/detokenizer
- - tests/multimodal
- - tests/utils_
+ - benchmarks/
+ - vllm/platforms/rocm.py
commands:
- - pytest -v -s detokenizer
- - pytest -v -s -m 'not cpu_test' multimodal
- - pytest -v -s utils_
+ - bash scripts/run-benchmarks.sh
+#---------------------------------------------------------- mi250 · compile ----------------------------------------------------------#
-- label: Async Engine, Inputs, Utils, Worker, Config (CPU) # TBD
+- label: PyTorch Compilation Unit Tests # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
+ torch_nightly: true
optional: true
- no_gpu: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/test_inputs.py
- - tests/test_outputs.py
- - tests/test_pooling_params.py
- - tests/test_ray_env.py
- - tests/multimodal
- - tests/renderers
- - tests/standalone_tests/lazy_imports.py
- - tests/tokenizers_
- - tests/tool_parsers
- - tests/transformers_utils
- - tests/config
+ - vllm/compilation/
+ - vllm/model_executor/layers/
+ - vllm/v1/worker/
+ - vllm/v1/attention/
+ - vllm/v1/cudagraph_dispatcher.py
+ - vllm/config/compilation.py
+ - csrc/
+ - tests/compile
+ - vllm/platforms/rocm.py
commands:
- - python3 standalone_tests/lazy_imports.py
- - pytest -v -s test_inputs.py
- - pytest -v -s test_outputs.py
- - pytest -v -s test_pooling_params.py
- - pytest -v -s test_ray_env.py
- - pytest -v -s -m 'cpu_test' multimodal
- - pytest -v -s renderers
- - pytest -v -s tokenizers_
- - pytest -v -s tool_parsers
- - pytest -v -s transformers_utils
- - pytest -v -s config
-
+ - "find compile/ -maxdepth 1 -name 'test_*.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
-- label: Python-only Installation # TBD
+- label: PyTorch Fullgraph # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
optional: true
+ torch_nightly: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - tests/standalone_tests/python_only_compile.sh
- - setup.py
+ - vllm/compilation/
+ - vllm/model_executor/
+ - vllm/v1/attention/
+ - vllm/config/compilation.py
+ - csrc/
+ - tests/compile
- vllm/platforms/rocm.py
commands:
- - bash standalone_tests/python_only_compile.sh
-
+ - pytest -v -s compile/fullgraph/test_full_graph.py -k 'not test_fp8_kv_scale_compile'
-- label: Basic Correctness # TBD
+- label: PyTorch Fullgraph Smoke Test # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
- fast_check: true
+ optional: true
torch_nightly: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/basic_correctness/test_basic_correctness
- - tests/basic_correctness/test_cpu_offload
- - tests/basic_correctness/test_cumem.py
- commands:
- - export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - pytest -v -s basic_correctness/test_cumem.py
- - pytest -v -s basic_correctness/test_basic_correctness.py
- - pytest -v -s basic_correctness/test_cpu_offload.py
-
-
-- label: Entrypoints Unit Tests # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- fast_check: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/entrypoints
- - tests/entrypoints/
+ - vllm/compilation/
+ - vllm/model_executor/
+ - vllm/v1/attention/
+ - vllm/config/compilation.py
+ - csrc/
+ - tests/compile
- vllm/platforms/rocm.py
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
-
+ - "find compile/fullgraph/ -name 'test_*.py' -not -name 'test_full_graph.py' -exec pytest -s -v {} \\\\;"
-- label: Entrypoints Integration (LLM) # TBD
+- label: Distributed Compile + RPC Tests (2 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- fast_check: true
- torch_nightly: true
+ agent_pool: mi250_2
+ num_gpus: 2
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/entrypoints/llm
- - tests/entrypoints/offline_mode
+ - vllm/compilation/
+ - vllm/distributed/
+ - vllm/engine/
+ - vllm/executor/
+ - vllm/worker/worker_base.py
+ - vllm/v1/engine/
+ - vllm/v1/worker/
+ - tests/compile/fullgraph/test_basic_correctness.py
+ - tests/compile/test_wrapper.py
+ - tests/entrypoints/llm/test_collective_rpc.py
+ - vllm/platforms/rocm.py
commands:
- - export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - pytest -v -s entrypoints/llm --ignore=entrypoints/llm/test_generate.py --ignore=entrypoints/llm/test_collective_rpc.py
- - pytest -v -s entrypoints/llm/test_generate.py
- - pytest -v -s entrypoints/offline_mode
+ - pytest -v -s entrypoints/llm/test_collective_rpc.py
+ - pytest -v -s ./compile/fullgraph/test_basic_correctness.py
+ - pytest -v -s ./compile/test_wrapper.py
+#-------------------------------------------------------- mi250 · distributed --------------------------------------------------------#
-- label: Entrypoints Integration (API Server 2) # TBD
+- label: Distributed Comm Ops # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- fast_check: true
- torch_nightly: true
+ agent_pool: mi250_2
+ num_gpus: 2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/entrypoints/rpc
- - tests/entrypoints/serve/instrumentator
- - tests/tool_use
+ - vllm/distributed
+ - tests/distributed
+ - vllm/platforms/rocm.py
commands:
- - export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - pytest -v -s entrypoints/serve/instrumentator
- - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
- - pytest -v -s tool_use
-
+ - pytest -v -s distributed/test_comm_ops.py
+ - pytest -v -s distributed/test_shm_broadcast.py
+ - pytest -v -s distributed/test_shm_buffer.py
+ - pytest -v -s distributed/test_shm_storage.py
-- label: Entrypoints Integration (Responses API) # TBD
+- label: Distributed Torchrun + Shutdown Tests (2 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- fast_check: true
- torch_nightly: true
+ agent_pool: mi250_2
+ num_gpus: 2
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/entrypoints/openai/responses
+ - vllm/distributed/
+ - vllm/engine/
+ - vllm/executor/
+ - vllm/worker/worker_base.py
+ - vllm/v1/engine/
+ - vllm/v1/worker/
+ - tests/distributed/
+ - tests/v1/shutdown
+ - tests/v1/worker/test_worker_memory_snapshot.py
+ - vllm/platforms/rocm.py
commands:
- - export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - pytest -v -s entrypoints/openai/responses
-
+ - 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
+ - pytest -v -s v1/worker/test_worker_memory_snapshot.py
-- label: EPLB Algorithm # TBD
+- label: Elastic EP Scaling Test # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
+ agent_pool: mi250_4
+ num_gpus: 4
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/distributed/eplb
- - tests/distributed/test_eplb_algo.py
+ - vllm/distributed/
+ - vllm/engine/
+ - vllm/executor/
+ - vllm/compilation/
+ - tests/distributed/
- vllm/platforms/rocm.py
commands:
- - pytest -v -s distributed/test_eplb_algo.py
-
+ - pytest -v -s distributed/test_elastic_ep.py
- label: EPLB Execution # TBD
timeout_in_minutes: 180
@@ -307,8 +307,7 @@ steps:
- pytest -v -s distributed/test_eplb_execute.py
- pytest -v -s distributed/test_eplb_spec_decode.py
-
-- label: Elastic EP Scaling Test # TBD
+- label: Pipeline + Context Parallelism (4 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_4
@@ -318,338 +317,303 @@ steps:
- vllm/distributed/
- vllm/engine/
- vllm/executor/
- - vllm/compilation/
+ - vllm/model_executor/models/
+ - vllm/model_executor/layers/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
- tests/distributed/
+ - vllm/_aiter_ops.py
- vllm/platforms/rocm.py
commands:
- - pytest -v -s distributed/test_elastic_ep.py
+ - pytest -v -s distributed/test_pp_cudagraph.py
+ - pytest -v -s distributed/test_pipeline_parallel.py
+#---------------------------------------------------------- mi250 · engine -----------------------------------------------------------#
-- label: Metrics, Tracing (2 GPUs) # TBD
+- label: Engine # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
- num_gpus: 2
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/v1/tracing
+ - tests/engine
+ - tests/test_sequence
+ - tests/test_config
+ - tests/test_logger
+ - tests/test_vllm_port
commands:
- - "pip install \
- 'opentelemetry-sdk>=1.26.0' \
- 'opentelemetry-api>=1.26.0' \
- 'opentelemetry-exporter-otlp>=1.26.0' \
- 'opentelemetry-semantic-conventions-ai>=0.4.1'"
- - pytest -v -s v1/tracing
+ - pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py
+#----------------------------------------------------------- mi250 · evals -----------------------------------------------------------#
-- label: Regression # TBD
+- label: Multi-Modal Accuracy Eval (Small Models) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
- working_dir: "/vllm-workspace/tests"
+ optional: true
+ working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
source_file_dependencies:
- - vllm/
- - tests/test_regression
+ - vllm/multimodal/
+ - vllm/inputs/
+ - vllm/v1/core/
+ - vllm/platforms/rocm.py
+ - vllm/model_executor/model_loader/
commands:
- - pip install modelscope
- - pytest -v -s test_regression.py
+ - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-mm-small.txt --tp-size=1
+#--------------------------------------------------------- mi250 · examples ----------------------------------------------------------#
-- label: Engine # TBD
+- label: Examples # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
optional: true
- working_dir: "/vllm-workspace/tests"
+ working_dir: "/vllm-workspace/examples"
source_file_dependencies:
- - vllm/
- - tests/engine
- - tests/test_sequence
- - tests/test_config
- - tests/test_logger
- - tests/test_vllm_port
+ - vllm/entrypoints
+ - vllm/multimodal
+ - examples/
+ - vllm/platforms/rocm.py
commands:
- - pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py
+ - pip install tensorizer
+ # Basic
+ - python3 basic/offline_inference/chat.py --attention-backend TRITON_ATTN
+ - python3 basic/offline_inference/generate.py --model facebook/opt-125m
+ - python3 basic/offline_inference/generate.py --model meta-llama/Llama-2-13b-chat-hf --cpu-offload-gb 10
+ - python3 basic/offline_inference/classify.py
+ - 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
+ # 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
+#---------------------------------------------------------- mi250 · kernels ----------------------------------------------------------#
-- label: Engine (1 GPU) # TBD
+- label: Kernels Core Operation Test # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/v1/
- - tests/v1/engine/
+ - csrc/
+ - tests/kernels/core
+ - tests/kernels/test_top_k_per_row.py
+ - tests/kernels/test_concat_mla_q.py
+ - vllm/model_executor/layers/rotary_embedding/
+ - vllm/_aiter_ops.py
- vllm/platforms/rocm.py
commands:
- - pytest -v -s v1/engine/test_preprocess_error_handling.py
- - pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py
-
+ - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py kernels/test_concat_mla_q.py kernels/test_top_k_per_row.py
-- label: e2e Scheduling (1 GPU) # TBD
+- label: Kernels Helion Test # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/v1/
- - tests/v1/e2e/general/
+ - vllm/utils/import_utils.py
+ - tests/kernels/helion/
- vllm/platforms/rocm.py
commands:
- - pytest -v -s v1/e2e/general/test_async_scheduling.py
-
+ - pip install helion==1.0.0
+ - pytest -v -s kernels/helion/
-- label: e2e Core (1 GPU) # TBD
+- label: Kernels Mamba Test # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/v1/
- - tests/v1/e2e/general/
+ - csrc/mamba/
+ - tests/kernels/mamba
+ - vllm/model_executor/layers/mamba/ops
- vllm/platforms/rocm.py
commands:
- - pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py
+ - pytest -v -s kernels/mamba
+#----------------------------------------------------------- mi250 · lora ------------------------------------------------------------#
-- label: Spec Decode Speculators + MTP # TBD
+- label: LoRA %N # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
+ parallelism: 4
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/v1/spec_decode/
- - vllm/v1/worker/gpu/spec_decode/
- - vllm/model_executor/model_loader/
- - vllm/v1/sample/
- - vllm/model_executor/layers/
- - vllm/transformers_utils/configs/speculators/
- - tests/v1/e2e/spec_decode/
+ - vllm/lora
+ - tests/lora
- vllm/platforms/rocm.py
commands:
- - pytest -v -s v1/e2e/spec_decode -k "speculators or mtp_correctness"
+ - 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 -------------------------------------------------------#
-- label: Spec Decode Ngram + Suffix # TBD
+- label: Model Executor # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
optional: true
+ torch_nightly: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/v1/spec_decode/
- - vllm/v1/worker/gpu/spec_decode/
- - vllm/model_executor/model_loader/
- - vllm/v1/sample/
- - vllm/model_executor/layers/
- - tests/v1/e2e/spec_decode/
+ - vllm/engine/arg_utils.py
+ - vllm/config/model.py
+ - vllm/model_executor
+ - tests/model_executor
+ - tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
+ - vllm/_aiter_ops.py
- vllm/platforms/rocm.py
commands:
- - pytest -v -s v1/e2e/spec_decode -k "ngram or suffix"
+ - apt-get update && apt-get install -y curl libsodium23
+ - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+ - pytest -v -s model_executor -m '(not slow_test)'
+ - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py
+#---------------------------------------------------------- mi250 · models -----------------------------------------------------------#
-- label: Spec Decode Draft Model # TBD
+- label: Basic Models Test (Other CPU) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
+ no_gpu: true
optional: true
+ torch_nightly: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/v1/spec_decode/
- - vllm/v1/worker/gpu/spec_decode/
- - vllm/model_executor/model_loader/
- - vllm/v1/sample/
- - vllm/model_executor/layers/
- - tests/v1/e2e/spec_decode/
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/models/test_utils.py
+ - tests/models/test_vision.py
commands:
- - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
-
+ - pytest -v -s models/test_utils.py models/test_vision.py
-- label: V1 e2e (2 GPUs) # TBD
+- label: Basic Models Tests (Extra Initialization) %N # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
- optional: true
+ torch_nightly: true
+ parallelism: 2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/v1/e2e
+ - vllm/model_executor/models/
+ - vllm/model_executor/layers/
+ - tests/models/test_initialization.py
+ - tests/models/registry.py
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
commands:
- - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "tensor_parallelism"
-
+ - pytest -v -s models/test_initialization.py -k 'not test_can_initialize_small_subset' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
-- label: V1 Sample + Logits # TBD
- timeout_in_minutes: 60
+- label: Basic Models Tests (Initialization) # TBD
+ timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
- optional: true
+ torch_nightly: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/v1/sample
- - tests/v1/logits_processors
- - tests/v1/test_oracle.py
- - tests/v1/test_request.py
- - tests/v1/test_outputs.py
+ - tests/models/test_initialization.py
+ - tests/models/registry.py
commands:
- - pytest -v -s v1/sample
- - 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
-
+ - pytest -v -s models/test_initialization.py::test_can_initialize_small_subset
-- label: V1 Core + KV + Metrics # TBD
- timeout_in_minutes: 60
+- label: Basic Models Tests (Other) # TBD
+ timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
- optional: true
+ torch_nightly: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - 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
+ - tests/models/test_terratorch.py
+ - tests/models/test_transformers.py
+ - tests/models/test_registry.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
-
+ - pytest -v -s models/test_terratorch.py models/test_transformers.py models/test_registry.py
-- label: V1 attention (H100-MI250) # TBD
+- label: Language Models Test (MTEB) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/config/attention.py
- - vllm/model_executor/layers/attention
- - vllm/v1/attention
- - tests/v1/attention
- - vllm/_aiter_ops.py
- - vllm/envs.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/models/language/pooling_mteb_test
commands:
- - pytest -v -s v1/attention
-
+ - pytest -v -s models/language/pooling_mteb_test
-- label: V1 others (CPU) # TBD
+- label: Language Models Test (PPL) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
- no_gpu: true
- optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/v1
+ - tests/models/language/generation_ppl_test
commands:
- - pytest -v -s -m 'cpu_test' v1/core
- - pytest -v -s v1/structured_output
- - pytest -v -s v1/test_serial_utils.py
- - pytest -v -s -m 'cpu_test' v1/kv_connector/unit
- - pytest -v -s -m 'cpu_test' v1/metrics
-
+ - pytest -v -s models/language/generation_ppl_test
-- label: Examples # TBD
+- label: Language Models Tests (Extra Standard) %N # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
- optional: true
- working_dir: "/vllm-workspace/examples"
+ torch_nightly: true
+ parallelism: 2
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/entrypoints
- - vllm/multimodal
- - examples/
+ - vllm/model_executor/models/
+ - vllm/model_executor/model_loader/
+ - vllm/model_executor/layers/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - tests/models/language/pooling/test_embedding.py
+ - tests/models/language/generation/test_common.py
+ - tests/models/language/pooling/test_classification.py
+ - vllm/_aiter_ops.py
- vllm/platforms/rocm.py
commands:
- - pip install tensorizer
- # Basic
- - python3 basic/offline_inference/chat.py --attention-backend TRITON_ATTN
- - python3 basic/offline_inference/generate.py --model facebook/opt-125m
- - python3 basic/offline_inference/generate.py --model meta-llama/Llama-2-13b-chat-hf --cpu-offload-gb 10
- - python3 basic/offline_inference/classify.py
- - python3 basic/offline_inference/embed.py
- - python3 basic/offline_inference/score.py
- # Multi-modal models
- - 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 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
-
+ - pip freeze | grep -E 'torch'
+ - pytest -v -s models/language -m 'core_model and slow_test' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
-- label: Platform Tests (CUDA) # TBD
+- label: Multi-Modal Models (Extended Generation 2) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/cuda
+ - tests/models/multimodal/generation
commands:
- - pytest -v -s cuda/test_cuda_context.py
- - pytest -v -s cuda/test_platform_no_cuda_init.py
-
+ - 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: Samplers Test # TBD
+- label: Multi-Modal Models (Extended Pooling) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
- optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/model_executor/layers
- - vllm/sampling_metadata.py
- - vllm/v1/sample/
- - vllm/beam_search.py
- - tests/samplers
- - tests/conftest.py
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
- commands:
- - pytest -v -s samplers
-
-
-- label: LoRA %N # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- parallelism: 4
- optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/lora
- - tests/lora
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/models/multimodal/pooling
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_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
-
+ - pytest -v -s models/multimodal/pooling -m 'not core_model'
-- label: PyTorch Compilation Unit Tests # TBD
+- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl" # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
@@ -657,56 +621,75 @@ steps:
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/compilation/
- - vllm/model_executor/layers/
- - vllm/v1/worker/
- - vllm/v1/attention/
- - vllm/v1/cudagraph_dispatcher.py
- - vllm/config/compilation.py
- - csrc/
- - tests/compile
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/models/multimodal
commands:
- - "find compile/ -maxdepth 1 -name 'test_*.py' -print0 | xargs -0 -n1 -I{} pytest -s -v '{}'"
+ - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
+ - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "not qwen2 and not qwen3 and not gemma"
+ - pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model
+#---------------------------------------------------------- mi250 · plugins ----------------------------------------------------------#
-- label: PyTorch Fullgraph Smoke Test # TBD
+- label: Plugin Tests (2 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
+ agent_pool: mi250_2
+ num_gpus: 2
optional: true
- torch_nightly: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/compilation/
- - vllm/model_executor/
- - vllm/v1/attention/
- - vllm/config/compilation.py
- - csrc/
- - tests/compile
+ - vllm/plugins/
+ - tests/plugins/
- vllm/platforms/rocm.py
commands:
- - "find compile/fullgraph/ -name 'test_*.py' -not -name 'test_full_graph.py' -exec pytest -s -v {} \\\\;"
+ # BEGIN: platform plugin and general plugin tests, all the code in-between runs on dummy platform
+ - pip install -e ./plugins/vllm_add_dummy_platform
+ - pytest -v -s plugins_tests/test_platform_plugins.py
+ - pip uninstall vllm_add_dummy_platform -y
+ # END: platform plugin tests
+ # BEGIN: `io_processor` plugins test, all the code in between uses the `prithvi_io_processor` plugin
+ - pip install -e ./plugins/prithvi_io_processor_plugin
+ - pytest -v -s plugins_tests/test_io_processor_plugins.py
+ - pytest -v -s plugins_tests/test_terratorch_io_processor_plugins.py
+ - pip uninstall prithvi_io_processor_plugin -y
+ # END: `io_processor` plugins test
+ # BEGIN: `bge_m3_sparse io_processor` test
+ - pip install -e ./plugins/bge_m3_sparse_plugin
+ - pytest -v -s plugins_tests/test_bge_m3_sparse_io_processor_plugins.py
+ - pip uninstall bge_m3_sparse_plugin -y
+ # END: `bge_m3_sparse io_processor` test
+ # BEGIN: `stat_logger` plugins test
+ - pip install -e ./plugins/vllm_add_dummy_stat_logger
+ - pytest -v -s plugins_tests/test_stats_logger_plugins.py
+ - pip uninstall dummy_stat_logger -y
+ # END: `stat_logger` plugins test
+ # BEGIN: other tests
+ - pytest -v -s plugins_tests/test_scheduler_plugins.py
+ - pip install -e ./plugins/vllm_add_dummy_model
+ - pytest -v -s distributed/test_distributed_oot.py
+ - pytest -v -s entrypoints/openai/chat_completion/test_oot_registration.py
+ - pytest -v -s models/test_oot_registration.py
+ - pytest -v -s plugins/lora_resolvers
+#------------------------------------------------------------ mi250 · v1 -------------------------------------------------------------#
-- label: PyTorch Fullgraph # TBD
+- label: Batch Invariance (H100-MI250) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
optional: true
- torch_nightly: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/compilation/
- - vllm/model_executor/
- - vllm/v1/attention/
- - vllm/config/compilation.py
- - csrc/
- - tests/compile
+ - vllm/v1/attention
+ - vllm/model_executor/layers
+ - tests/v1/determinism/
+ - vllm/_aiter_ops.py
- vllm/platforms/rocm.py
commands:
- - pytest -v -s compile/fullgraph/test_full_graph.py -k 'not test_fp8_kv_scale_compile'
-
+ - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+ - pip install pytest-timeout pytest-forked
+ - pytest -v -s v1/determinism/test_batch_invariance.py
+ - pytest -v -s v1/determinism/test_rms_norm_batch_invariant.py
- label: Cudagraph # TBD
timeout_in_minutes: 180
@@ -724,1214 +707,923 @@ steps:
- pytest -v -s v1/cudagraph/test_cudagraph_dispatch.py
- pytest -v -s v1/cudagraph/test_cudagraph_mode.py
-
-- label: Kernels Core Operation Test # TBD
+- label: e2e Core (1 GPU) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - csrc/
- - tests/kernels/core
- - tests/kernels/test_top_k_per_row.py
- - tests/kernels/test_concat_mla_q.py
- - vllm/model_executor/layers/rotary_embedding/
- - vllm/_aiter_ops.py
+ - vllm/v1/
+ - tests/v1/e2e/general/
- vllm/platforms/rocm.py
commands:
- - pytest -v -s kernels/core kernels/test_top_k_per_row.py
-
+ - pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py
-- label: Kernels Mamba Test # TBD
+- label: e2e Scheduling (1 GPU) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - csrc/mamba/
- - tests/kernels/mamba
- - vllm/model_executor/layers/mamba/ops
+ - vllm/v1/
+ - tests/v1/e2e/general/
- vllm/platforms/rocm.py
commands:
- - pytest -v -s kernels/mamba
-
+ - pytest -v -s v1/e2e/general/test_async_scheduling.py
-- label: Kernels Helion Test # TBD
+- label: Engine (1 GPU) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/utils/import_utils.py
- - tests/kernels/helion/
+ - vllm/v1/
+ - tests/v1/engine/
- vllm/platforms/rocm.py
commands:
- - pip install helion==0.3.3
- - pytest -v -s kernels/helion/
-
+ - pytest -v -s v1/engine/test_preprocess_error_handling.py
+ - pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py
-- label: Model Executor # TBD
+- label: Spec Decode Draft Model # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
optional: true
- torch_nightly: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/engine/arg_utils.py
- - vllm/config/model.py
- - vllm/model_executor
- - tests/model_executor
- - tests/entrypoints/openai/completion/test_tensorizer_entrypoint.py
- - vllm/_aiter_ops.py
+ - vllm/v1/spec_decode/
+ - vllm/v1/worker/gpu/spec_decode/
+ - vllm/model_executor/model_loader/
+ - vllm/v1/sample/
+ - vllm/model_executor/layers/
+ - tests/v1/e2e/spec_decode/
- vllm/platforms/rocm.py
commands:
- - apt-get update && apt-get install -y curl libsodium23
- - export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - pytest -v -s model_executor -m '(not slow_test)'
- - pytest -v -s entrypoints/openai/completion/test_tensorizer_entrypoint.py
-
+ - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
-- label: Benchmarks # TBD
+- label: Spec Decode Speculators + MTP # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
- working_dir: "/vllm-workspace/.buildkite"
+ optional: true
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - benchmarks/
+ - vllm/v1/spec_decode/
+ - vllm/v1/worker/gpu/spec_decode/
+ - vllm/model_executor/model_loader/
+ - vllm/v1/sample/
+ - vllm/model_executor/layers/
+ - vllm/transformers_utils/configs/speculators/
+ - tests/v1/e2e/spec_decode/
- vllm/platforms/rocm.py
commands:
- - bash scripts/run-benchmarks.sh
-
+ - pytest -v -s v1/e2e/spec_decode -k "speculators or mtp_correctness"
-- label: Benchmarks CLI Test # TBD
+- label: V1 attention (H100-MI250) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/benchmarks/
+ - vllm/config/attention.py
+ - vllm/model_executor/layers/attention
+ - vllm/v1/attention
+ - tests/v1/attention
+ - vllm/_aiter_ops.py
+ - vllm/envs.py
+ - vllm/platforms/rocm.py
commands:
- - pytest -v -s benchmarks/
-
+ - pytest -v -s v1/attention
-- label: OpenAI API correctness # TBD
- timeout_in_minutes: 180
+- 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:
- - csrc/
- - vllm/entrypoints/openai/
- - vllm/model_executor/models/whisper.py
- - vllm/model_executor/layers/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
- - vllm/model_executor/model_loader/
+ - 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:
- - bash ../tools/install_torchcodec_rocm.sh || exit 1
- - pytest -s entrypoints/openai/correctness/
-
+ - 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: Basic Models Tests (Initialization) # TBD
- timeout_in_minutes: 180
+- label: V1 Sample + Logits # TBD
+ timeout_in_minutes: 60
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
- torch_nightly: true
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/models/test_initialization.py
- - tests/models/registry.py
+ - tests/v1/sample
+ - tests/v1/logits_processors
+ - tests/v1/test_oracle.py
+ - tests/v1/test_request.py
+ - tests/v1/test_outputs.py
commands:
- - pytest -v -s models/test_initialization.py::test_can_initialize_small_subset
-
+ - pytest -v -s v1/sample
+ - 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
-- label: Basic Models Tests (Extra Initialization) %N # TBD
+- label: Distributed DP Tests (2 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- torch_nightly: true
- parallelism: 2
+ agent_pool: mi250_2
+ num_gpus: 2
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/model_executor/models/
- - vllm/model_executor/layers/
- - tests/models/test_initialization.py
- - tests/models/registry.py
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/distributed/
+ - vllm/engine/
+ - vllm/executor/
+ - vllm/worker/worker_base.py
+ - vllm/v1/engine/
+ - vllm/v1/worker/
+ - tests/v1/distributed
+ - tests/entrypoints/openai/test_multi_api_servers.py
+ - vllm/platforms/rocm.py
commands:
- - pytest -v -s models/test_initialization.py -k 'not test_can_initialize_small_subset' --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
-
+ - 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
+ - DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py
-- label: Basic Models Tests (Other) # TBD
+- label: NixlConnector PD + Spec Decode acceptance (2 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- torch_nightly: true
+ agent_pool: mi250_2
+ num_gpus: 2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/models/test_terratorch.py
- - tests/models/test_transformers.py
- - tests/models/test_registry.py
+ - 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/
+ - vllm/platforms/rocm.py
commands:
- - pytest -v -s models/test_terratorch.py models/test_transformers.py models/test_registry.py
-
+ - 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
-- label: Basic Models Test (Other CPU) # TBD
+- label: V1 e2e (2 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- no_gpu: true
+ agent_pool: mi250_2
optional: true
- torch_nightly: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/models/test_utils.py
- - tests/models/test_vision.py
+ - tests/v1/e2e
commands:
- - pytest -v -s models/test_utils.py models/test_vision.py
-
+ - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "tensor_parallelism"
-- label: Language Models Tests (Extra Standard) %N # TBD
+- label: Distributed NixlConnector PD accuracy (4 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- torch_nightly: true
- parallelism: 2
+ agent_pool: mi250_4
+ num_gpus: 4
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/model_executor/layers/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - tests/models/language/pooling/test_embedding.py
- - tests/models/language/generation/test_common.py
- - tests/models/language/pooling/test_classification.py
- - vllm/_aiter_ops.py
+ - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
+ - tests/v1/kv_connector/nixl_integration/
- 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
+ - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
+ - ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
+#------------------------------------------------------------- mi250 · misc ------------------------------------------------------------#
-- label: Language Models Test (PPL) # TBD
+- label: Async Engine, Inputs, Utils, Worker, Config (CPU) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
+ optional: true
+ no_gpu: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/models/language/generation_ppl_test
+ - tests/test_inputs.py
+ - tests/test_outputs.py
+ - tests/test_pooling_params.py
+ - tests/test_ray_env.py
+ - tests/multimodal
+ - tests/renderers
+ - tests/standalone_tests/lazy_imports.py
+ - tests/tokenizers_
+ - tests/reasoning
+ - tests/tool_parsers
+ - tests/transformers_utils
+ - tests/config
commands:
- - pytest -v -s models/language/generation_ppl_test
+ - python3 standalone_tests/lazy_imports.py
+ - pytest -v -s test_inputs.py
+ - pytest -v -s test_outputs.py
+ - pytest -v -s test_pooling_params.py
+ - pytest -v -s test_ray_env.py
+ - pytest -v -s -m 'cpu_test' multimodal
+ - pytest -v -s renderers
+ - pytest -v -s tokenizers_
+ - pytest -v -s reasoning --ignore=reasoning/test_seedoss_reasoning_parser.py --ignore=reasoning/test_glm4_moe_reasoning_parser.py
+ - pytest -v -s tool_parsers
+ - pytest -v -s transformers_utils
+ - pytest -v -s config
+#########################################################################################################################################
+# #
+# MI300 (gfx942) tests #
+# #
+#########################################################################################################################################
-- label: Language Models Test (Extended Pooling) # TBD
+#----------------------------------------------------- mi300 · basic_correctness -----------------------------------------------------#
+
+- label: Basic Correctness # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
+ 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/models/language/pooling
+ - tests/basic_correctness/test_basic_correctness
+ - tests/basic_correctness/test_cpu_offload
+ - tests/basic_correctness/test_cumem.py
commands:
- - pytest -v -s models/language/pooling -m 'not core_model'
-
+ - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+ - pytest -v -s basic_correctness/test_cumem.py
+ - pytest -v -s basic_correctness/test_basic_correctness.py
+ - pytest -v -s basic_correctness/test_cpu_offload.py
-- label: Language Models Test (MTEB) # TBD
+- label: Distributed Model Tests (2 GPUs) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_2
+ num_gpus: 2
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/models/language/pooling_mteb_test
+ - vllm/model_executor/model_loader/sharded_state_loader.py
+ - vllm/model_executor/models/
+ - vllm/model_executor/layers/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
+ - tests/basic_correctness/
+ - tests/model_executor/model_loader/test_sharded_state_loader.py
+ - tests/models/
commands:
- - pytest -v -s models/language/pooling_mteb_test
+ - TARGET_TEST_SUITE=L4 pytest basic_correctness/ -v -s -m 'distributed(num_gpus=2)'
+ - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s model_executor/model_loader/test_sharded_state_loader.py -m '(not slow_test)'
+ - pytest models/test_transformers.py -v -s -m 'distributed(num_gpus=2)'
+ - pytest models/language -v -s -m 'distributed(num_gpus=2)'
+ - pytest models/multimodal -v -s -m 'distributed(num_gpus=2)' --ignore models/multimodal/generation/test_whisper.py --ignore models/multimodal/generation/test_phi4siglip.py
+ - pytest models/multimodal/generation/test_phi4siglip.py -v -s -m 'distributed(num_gpus=2)'
+ - VLLM_WORKER_MULTIPROC_METHOD=spawn pytest models/multimodal/generation/test_whisper.py -v -s -m 'distributed(num_gpus=2)'
+#-------------------------------------------------------- mi300 · benchmarks ---------------------------------------------------------#
-- label: Multi-Modal Processor (CPU) # TBD
+- label: Benchmarks # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- no_gpu: true
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
optional: true
+ working_dir: "/vllm-workspace/.buildkite"
+ source_file_dependencies:
+ - benchmarks/
+ - vllm/platforms/rocm.py
+ commands:
+ - bash scripts/run-benchmarks.sh
+
+- label: Benchmarks CLI Test # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/models/multimodal
- - tests/models/registry.py
+ - tests/benchmarks/
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py
+ - pytest -v -s benchmarks/
+#---------------------------------------------------------- mi300 · compile ----------------------------------------------------------#
-- label: Multi-Modal Accuracy Eval (Small Models) # TBD
+- label: Fusion E2E Config Sweep (H100-MI300) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- optional: true
- working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ num_gpus: 1
+ working_dir: "/vllm-workspace/"
source_file_dependencies:
- - vllm/multimodal/
- - vllm/inputs/
- - vllm/v1/core/
+ - csrc/quantization/
+ - vllm/compilation/
+ - vllm/model_executor/layers/layernorm.py
+ - vllm/model_executor/layers/activation.py
+ - vllm/model_executor/layers/attention/attention.py
+ - vllm/model_executor/layers/quantization/input_quant_fp8.py
+ - tests/compile/fusions_e2e/
+ - vllm/_aiter_ops.py
- vllm/platforms/rocm.py
- - vllm/model_executor/model_loader/
commands:
- - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-mm-small.txt --tp-size=1
+ - rocm-smi
+ - pytest -v -s tests/compile/fusions_e2e/test_tp1_quant.py -k "llama-3"
+- label: Fusion E2E Quick (H100-MI300) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ num_gpus: 1
+ working_dir: "/vllm-workspace/"
+ source_file_dependencies:
+ - csrc/quantization/
+ - vllm/model_executor/
+ - vllm/v1/attention/
+ - vllm/compilation/
+ - tests/compile/fusions_e2e/
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
+ commands:
+ - rocm-smi
+ # Run all models and attn backends but only Inductor partition and native custom ops
+ - "pytest -v -s tests/compile/fusions_e2e/test_tp1_quant.py -k 'inductor_partition and not +rms_norm and not +quant_fp8'"
+ # Different from CUDA, Qwen requires +rms_norm and +quant_fp8 as rms+quant fusion is only supported on AITER
+ - "pytest -v -s tests/compile/fusions_e2e/test_tp1_quant.py -k 'inductor_partition and +rms_norm and +quant_fp8 and qwen3'"
-- label: "Multi-Modal Models (Standard) 1: qwen2" # TBD
+- label: PyTorch Compilation Passes Unit Tests # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- torch_nightly: true
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/models/multimodal
+ - tests/compile/passes
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen2"
- - pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model
-
+ - pytest -s -v compile/passes --ignore compile/passes/distributed
-- label: "Multi-Modal Models (Standard) 2: qwen3 + gemma" # TBD
+- label: Pytorch Nightly Dependency Override Check # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- torch_nightly: true
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ optional: true
+ soft_fail: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/models/multimodal
+ - requirements/test/nightly-torch.txt
+ - vllm/platforms/rocm.py
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen3 or gemma"
- - pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model
-
+ - bash standalone_tests/pytorch_nightly_dependency.sh
-- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl" # TBD
+- label: Distributed Compile Unit Tests (2xH100-2xMI300) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- torch_nightly: true
- working_dir: "/vllm-workspace/tests"
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_2
+ num_gpus: 2
+ working_dir: "/vllm-workspace/"
source_file_dependencies:
- - vllm/
- - tests/models/multimodal
+ - vllm/compilation/
+ - vllm/model_executor/layers
+ - tests/compile/passes/distributed/
+ - tests/compile/fusions_e2e/
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "not qwen2 and not qwen3 and not gemma"
- - pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model
+ - 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
+#----------------------------------------------------------- mi300 · cuda ------------------------------------------------------------#
-- label: "Multi-Modal Models (Standard) 4: other + whisper" # TBD
+- label: Platform Tests (CUDA) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- torch_nightly: true
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/models/multimodal/generation
- - tests/models/multimodal/test_mapping.py
+ - tests/cuda
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/processing
- - pytest -v -s 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
+ - pytest -v -s cuda/test_cuda_context.py
+ - pytest -v -s cuda/test_platform_no_cuda_init.py
+#-------------------------------------------------------- mi300 · detokenizer --------------------------------------------------------#
-- label: Multi-Modal Models (Extended Generation 1) # TBD
+- label: Async Engine, Inputs, Utils, Worker # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/models/multimodal/generation
- - tests/models/multimodal/test_mapping.py
+ - tests/detokenizer
+ - tests/multimodal
+ - tests/utils_
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation -m 'not core_model' --ignore models/multimodal/generation/test_common.py
- - pytest -v -s models/multimodal/test_mapping.py
+ - pytest -v -s detokenizer
+ - pytest -v -s -m 'not cpu_test' multimodal
+ - pytest -v -s utils_
+#-------------------------------------------------------- mi300 · distributed --------------------------------------------------------#
-- label: Multi-Modal Models (Extended Generation 2) # TBD
+- label: EPLB Algorithm # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/models/multimodal/generation
+ - vllm/distributed/eplb
+ - tests/distributed/test_eplb_algo.py
+ - vllm/platforms/rocm.py
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=0) and not core_model'
-
+ - pytest -v -s distributed/test_eplb_algo.py
+ - pytest -v -s distributed/test_eplb_utils.py
-- label: Multi-Modal Models (Extended Generation 3) # TBD
+- label: Distributed Tests (2xH100-2xMI250) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- working_dir: "/vllm-workspace/tests"
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_2
+ num_gpus: 2
+ working_dir: "/vllm-workspace/"
source_file_dependencies:
- - vllm/
- - tests/models/multimodal/generation
+ - vllm/distributed/
+ - vllm/v1/distributed/
+ - vllm/model_executor/layers/fused_moe/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - tests/distributed/test_context_parallel.py
+ - examples/features/data_parallel/data_parallel_offline.py
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=1) and not core_model'
-
+ - 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
-- label: Multi-Modal Models (Extended Pooling) # TBD
+- label: Distributed Tests (4xA100-4xMI300) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
+ num_gpus: 4
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/models/multimodal/pooling
commands:
- - pytest -v -s models/multimodal/pooling -m 'not core_model'
-
+ - 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)'
+ - pytest -v -s -x lora/test_mixtral.py
-- label: Distributed Comm Ops # TBD
+- label: Distributed Torchrun + Examples (4 GPUs) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- num_gpus: 2
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
+ num_gpus: 4
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/distributed
- - tests/distributed
+ - vllm/distributed/
+ - tests/distributed/test_torchrun_example.py
+ - tests/distributed/test_torchrun_example_moe.py
+ - examples/rl/
+ - tests/examples/features/data_parallel/data_parallel_offline.py
- vllm/platforms/rocm.py
commands:
- - pytest -v -s distributed/test_comm_ops.py
- - pytest -v -s distributed/test_shm_broadcast.py
- - pytest -v -s distributed/test_shm_buffer.py
- - pytest -v -s distributed/test_shm_storage.py
-
+ - 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
+ # 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 (2 GPUs) # TBD
+- label: Elastic EP Scaling Test # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_2
- num_gpus: 2
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
+ num_gpus: 4
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/distributed/
- vllm/engine/
- vllm/executor/
- - vllm/worker/worker_base.py
- - vllm/v1/engine/
- - vllm/v1/worker/
- - tests/v1/distributed
- - tests/entrypoints/openai/test_multi_api_servers.py
+ - vllm/compilation/
+ - tests/distributed/
- 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
- - DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py
-
+ - pytest -v -s distributed/test_elastic_ep.py
-- label: Distributed Compile + RPC Tests (2 GPUs) # TBD
+- label: RayExecutorV2 (4 GPUs) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_2
- num_gpus: 2
- optional: true
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
+ num_gpus: 4
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/compilation/
- - vllm/distributed/
- - vllm/engine/
- - vllm/executor/
- - vllm/worker/worker_base.py
- - vllm/v1/engine/
- - vllm/v1/worker/
- - tests/compile/fullgraph/test_basic_correctness.py
- - tests/compile/test_wrapper.py
- - tests/entrypoints/llm/test_collective_rpc.py
+ - vllm/v1/executor/ray_executor_v2.py
+ - vllm/v1/executor/abstract.py
+ - vllm/v1/executor/multiproc_executor.py
+ - tests/distributed/test_ray_v2_executor.py
+ - tests/distributed/test_ray_v2_executor_e2e.py
+ - tests/distributed/test_pipeline_parallel.py
+ - tests/basic_correctness/test_basic_correctness.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
+ - export VLLM_USE_RAY_V2_EXECUTOR_BACKEND=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"
+ - TARGET_TEST_SUITE=L4 pytest -v -s basic_correctness/test_basic_correctness.py -k "ray"
-
-- label: Distributed Torchrun + Shutdown Tests (2 GPUs) # TBD
+- label: Distributed Tests (8xH100-8xMI300) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_2
- num_gpus: 2
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_8
+ num_gpus: 8
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
+ - examples/features/torchrun/torchrun_dp_example_offline.py
+ - vllm/config/parallel.py
- vllm/distributed/
- - vllm/engine/
- - vllm/executor/
- - vllm/worker/worker_base.py
- - vllm/v1/engine/
- - vllm/v1/worker/
- - tests/distributed/
- - tests/v1/shutdown
- - tests/v1/worker/test_worker_memory_snapshot.py
+ - vllm/v1/engine/llm_engine.py
+ - vllm/v1/executor/uniproc_executor.py
+ - vllm/v1/worker/gpu_worker.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
- - pytest -v -s v1/worker/test_worker_memory_snapshot.py
+ - torchrun --nproc-per-node=8 ../examples/features/torchrun/torchrun_dp_example_offline.py --tp-size=2 --pp-size=1 --dp-size=4 --enable-ep
+#-------------------------------------------------------- mi300 · entrypoints --------------------------------------------------------#
-- label: Distributed Model Tests (2 GPUs) # TBD
+- label: Entrypoints Integration (API Server 2) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_2
- num_gpus: 2
+ 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/model_executor/model_loader/sharded_state_loader.py
- - vllm/model_executor/models/
- - vllm/model_executor/layers/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - tests/basic_correctness/
- - tests/model_executor/model_loader/test_sharded_state_loader.py
- - tests/models/
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/entrypoints/rpc
+ - tests/entrypoints/serve/instrumentator
+ - tests/tool_use
commands:
- - TARGET_TEST_SUITE=L4 pytest basic_correctness/ -v -s -m 'distributed(num_gpus=2)'
- - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s model_executor/model_loader/test_sharded_state_loader.py -m '(not slow_test)'
- - pytest models/test_transformers.py -v -s -m 'distributed(num_gpus=2)'
- - pytest models/language -v -s -m 'distributed(num_gpus=2)'
- - pytest models/multimodal -v -s -m 'distributed(num_gpus=2)' --ignore models/multimodal/generation/test_whisper.py
- - VLLM_WORKER_MULTIPROC_METHOD=spawn pytest models/multimodal/generation/test_whisper.py -v -s -m 'distributed(num_gpus=2)'
-
+ - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+ - pytest -v -s entrypoints/serve/instrumentator
+ - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
+ - pytest -v -s tool_use
-- label: Plugin Tests (2 GPUs) # TBD
+- label: Entrypoints Integration (API Server openai - Part 1) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_2
- num_gpus: 2
- optional: true
+ 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/plugins/
- - tests/plugins/
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/entrypoints/openai
+ - tests/entrypoints/test_chat_utils
commands:
- # BEGIN: platform plugin and general plugin tests, all the code in-between runs on dummy platform
- - pip install -e ./plugins/vllm_add_dummy_platform
- - pytest -v -s plugins_tests/test_platform_plugins.py
- - pip uninstall vllm_add_dummy_platform -y
- # END: platform plugin tests
- # BEGIN: `io_processor` plugins test, all the code in between uses the `prithvi_io_processor` plugin
- - pip install -e ./plugins/prithvi_io_processor_plugin
- - pytest -v -s plugins_tests/test_io_processor_plugins.py
- - pip uninstall prithvi_io_processor_plugin -y
- # END: `io_processor` plugins test
- # BEGIN: `bge_m3_sparse io_processor` test
- - pip install -e ./plugins/bge_m3_sparse_plugin
- - pytest -v -s plugins_tests/test_bge_m3_sparse_io_processor_plugins.py
- - pip uninstall bge_m3_sparse_plugin -y
- # END: `bge_m3_sparse io_processor` test
- # BEGIN: `stat_logger` plugins test
- - pip install -e ./plugins/vllm_add_dummy_stat_logger
- - pytest -v -s plugins_tests/test_stats_logger_plugins.py
- - pip uninstall dummy_stat_logger -y
- # END: `stat_logger` plugins test
- # BEGIN: other tests
- - pytest -v -s plugins_tests/test_scheduler_plugins.py
- - pip install -e ./plugins/vllm_add_dummy_model
- - pytest -v -s distributed/test_distributed_oot.py
- - pytest -v -s entrypoints/openai/chat_completion/test_oot_registration.py
- - pytest -v -s models/test_oot_registration.py
- - pytest -v -s plugins/lora_resolvers
- # END: other tests
-
+ - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+ - pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
-- label: Pipeline + Context Parallelism (4 GPUs) # TBD
+- label: Entrypoints Integration (API Server openai - Part 2) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_4
- num_gpus: 4
+ 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/distributed/
- - vllm/engine/
- - vllm/executor/
- - vllm/model_executor/models/
- - vllm/model_executor/layers/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - tests/distributed/
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/entrypoints/openai
+ - tests/entrypoints/test_chat_utils
commands:
- - pytest -v -s distributed/test_pp_cudagraph.py
- - pytest -v -s distributed/test_pipeline_parallel.py
-
+ - 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: Ray Dependency Compatibility Check # TBD
+- label: Entrypoints Integration (API Server openai - Part 3) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_1
- working_dir: "/"
- source_file_dependencies:
- - requirements/
- - setup.py
- - vllm/platforms/rocm.py
+ 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/openai
+ - tests/entrypoints/test_chat_utils
commands:
- - bash /vllm-workspace/.buildkite/scripts/check-ray-compatibility.sh
-
+ - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+ - 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: Distributed NixlConnector PD accuracy (4 GPUs) # TBD
+- label: Entrypoints Integration (LLM) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_4
- num_gpus: 4
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ optional: true
+ fast_check: true
+ torch_nightly: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
- - tests/v1/kv_connector/nixl_integration/
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/entrypoints/llm
+ - tests/entrypoints/offline_mode
commands:
- - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
- - ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
-
+ - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+ - pytest -v -s entrypoints/llm --ignore=entrypoints/llm/test_generate.py --ignore=entrypoints/llm/test_collective_rpc.py
+ - pytest -v -s entrypoints/llm/test_generate.py
+ - pytest -v -s entrypoints/offline_mode
-- label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs) # TBD
+- label: Entrypoints Integration (Pooling) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_4
- num_gpus: 4
+ 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/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
- - tests/v1/kv_connector/nixl_integration/
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/entrypoints/pooling
commands:
- - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
- - DP_EP=1 ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
-
+ - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+ - pytest -v -s entrypoints/pooling
-- label: NixlConnector PD + Spec Decode acceptance (2 GPUs) # TBD
+- label: Entrypoints Integration (Responses API) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_2
- num_gpus: 2
+ 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/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
- - vllm/v1/worker/kv_connector_model_runner_mixin.py
- - tests/v1/kv_connector/nixl_integration/
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/entrypoints/openai/responses
commands:
- - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
- - ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
-
+ - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+ - pytest -v -s entrypoints/openai/responses
-- label: CrossLayer KV layout Distributed NixlConnector PD accuracy tests (4 GPUs) # TBD
+- label: Entrypoints Unit Tests # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_4
- num_gpus: 4
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ fast_check: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
- - tests/v1/kv_connector/nixl_integration/
+ - vllm/entrypoints
+ - tests/entrypoints/
- vllm/platforms/rocm.py
commands:
- - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
- - CROSS_LAYERS_BLOCKS=True ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
-
+ - 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
-- label: Hyrbid SSM NixlConnector PD accuracy tests (4 GPUs) # TBD
+- label: OpenAI API correctness # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi250_4
- num_gpus: 4
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
- - tests/v1/kv_connector/nixl_integration/
+ - csrc/
+ - vllm/entrypoints/openai/
+ - vllm/model_executor/models/whisper.py
+ - vllm/model_executor/layers/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - vllm/_aiter_ops.py
- vllm/platforms/rocm.py
+ - vllm/model_executor/model_loader/
commands:
- - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
- - HYBRID_SSM=1 ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
+ - bash ../tools/install_torchcodec_rocm.sh || exit 1
+ - pytest -s entrypoints/openai/correctness/
+#----------------------------------------------------------- mi300 · evals -----------------------------------------------------------#
-- label: Distributed Tests (2 GPUs)(H100-MI250) # TBD
+- label: DeepSeek V2-Lite Prefetch Offload Accuracy (H100-MI300) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
- agent_pool: mi325_2
- num_gpus: 2
- working_dir: "/vllm-workspace/"
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ num_gpus: 1
+ working_dir: "/vllm-workspace"
source_file_dependencies:
- - vllm/distributed/
- - vllm/v1/distributed/
+ - vllm/model_executor/models/
+ - vllm/model_executor/model_loader/
- vllm/model_executor/layers/fused_moe/
+ - vllm/model_executor/layers/quantization/
- vllm/v1/attention/backends/
+ - vllm/v1/attention/backends/mla/
- vllm/v1/attention/selector.py
- - tests/distributed/test_context_parallel.py
- - examples/offline_inference/data_parallel.py
+ - .buildkite/scripts/scheduled_integration_test/
- 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/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
-
-
-#####################################################################################################################################
-# #
-# gfx942 #
-# #
-#####################################################################################################################################
-
+ - bash .buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_prefetch_offload.sh 0.25 200 8030
-- label: Entrypoints Integration (LLM) # 13.1m
- timeout_in_minutes: 22
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- optional: true
- fast_check: true
- torch_nightly: true
+- label: LM Eval Small Models # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/entrypoints/llm
- - tests/entrypoints/offline_mode
+ - csrc/
+ - vllm/model_executor/layers/quantization
+ - vllm/model_executor/models/
+ - vllm/model_executor/model_loader/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
commands:
- - export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - pytest -v -s entrypoints/llm --ignore=entrypoints/llm/test_generate.py --ignore=entrypoints/llm/test_collective_rpc.py
- - pytest -v -s entrypoints/llm/test_generate.py
- - pytest -v -s entrypoints/offline_mode
-
+ - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-small.txt
-- label: Entrypoints Integration (API Server openai - Part 1) # TBD
+- label: LM Eval Small Models (MI300) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- fast_check: true
- torch_nightly: true
- working_dir: "/vllm-workspace/tests"
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
source_file_dependencies:
- - vllm/
- - tests/entrypoints/openai
- - tests/entrypoints/test_chat_utils
+ - csrc/
+ - vllm/model_executor/layers/quantization
+ - vllm/model_executor/models/
+ - vllm/model_executor/model_loader/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
commands:
- - export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
-
+ - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-small-rocm.txt
-- label: Entrypoints Integration (API Server openai - Part 2) # TBD
+- label: GPQA Eval (GPT-OSS) (2xH100-2xMI300) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- fast_check: true
- torch_nightly: true
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_2
+ num_gpus: 2
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/entrypoints/openai
- - tests/entrypoints/test_chat_utils
+ - csrc/
+ - vllm/model_executor/layers/quantization
+ - vllm/model_executor/models/
+ - vllm/model_executor/model_loader/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - vllm/model_executor/layers/fused_moe/
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
+ - tests/evals/gpt_oss/
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
-
+ - 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-gfx942.txt
-- label: Entrypoints Integration (API Server openai - Part 3) # TBD
+- label: LM Eval Small Models (2xB200-2xMI300) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- fast_check: true
- torch_nightly: true
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/entrypoints/openai
- - tests/entrypoints/test_chat_utils
+ - csrc/
+ - vllm/model_executor/layers/quantization
+ - vllm/model_executor/models/
+ - vllm/model_executor/model_loader/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
commands:
- - export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - 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
+ - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx-fp8-and-mixed.txt
+- label: DeepSeek V2-Lite Accuracy (4xH100-4xMI300) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
+ num_gpus: 4
+ optional: true
+ working_dir: "/vllm-workspace"
+ source_file_dependencies:
+ - vllm/model_executor/models/
+ - vllm/model_executor/model_loader/
+ - vllm/distributed/eplb
+ - vllm/model_executor/layers/fused_moe/
+ - vllm/model_executor/layers/quantization/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/backends/mla/
+ - vllm/v1/attention/selector.py
+ - .buildkite/scripts/scheduled_integration_test/
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
+ commands:
+ - bash .buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_ep_eplb.sh 0.25 200 8010
-- label: Entrypoints Integration (API Server 2) #26.9m
- timeout_in_minutes: 45
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+- label: LM Eval Large Models (4xA100-4xMI300) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
+ num_gpus: 4
optional: true
- fast_check: true
- torch_nightly: true
- working_dir: "/vllm-workspace/tests"
+ working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
source_file_dependencies:
- - vllm/
- - tests/entrypoints/rpc
- - tests/entrypoints/serve/instrumentator
- - tests/tool_use
+ - csrc/
+ - vllm/model_executor/layers/quantization
+ - vllm/model_executor/models/
+ - vllm/model_executor/model_loader/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - pytest -v -s entrypoints/serve/instrumentator
- - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
- - pytest -v -s tool_use
+ - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large.txt --tp-size=4
-
-- label: Entrypoints Integration (Pooling) # 22.8m
- timeout_in_minutes: 48
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- fast_check: true
- torch_nightly: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/
- - tests/entrypoints/pooling
- commands:
- - export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - pytest -v -s entrypoints/pooling
-
-
-- label: Distributed Torchrun + Examples (4 GPUs) # TBD
- timeout_in_minutes: 80
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_4
- num_gpus: 4
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/distributed/
- - tests/distributed/test_torchrun_example.py
- - tests/distributed/test_torchrun_example_moe.py
- - examples/rl/
- - 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/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) # TBD
- timeout_in_minutes: 60
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_4
- num_gpus: 4
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/distributed/
- - tests/v1/distributed
- - tests/v1/engine/test_engine_core_client.py
- - 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
- - TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_internal_lb_dp.py
- - TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_hybrid_lb_dp.py
- - pytest -v -s v1/engine/test_engine_core_client.py::test_kv_cache_events_dp
- - pytest -v -s distributed/test_utils.py
-
-
-- label: Distributed Compile + Comm (4 GPUs) # TBD
- timeout_in_minutes: 40
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_4
- num_gpus: 4
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/distributed/
- - tests/distributed/test_pynccl
- - tests/distributed/test_events
- - tests/compile/fullgraph/test_basic_correctness.py
- - tests/distributed/test_symm_mem_allreduce.py
- - 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
- - pytest -v -s distributed/test_symm_mem_allreduce.py
- - pytest -v -s distributed/test_multiproc_executor.py::test_multiproc_executor_multi_node
-
-
-- label: Distributed Tests (8 GPUs)(H100-MI325) # 6.4m
- timeout_in_minutes: 10
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_8
- num_gpus: 8
- optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - 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
- - vllm/platforms/rocm.py
- commands:
- - 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
-
-
-- label: Elastic EP Scaling Test # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_4
- num_gpus: 4
- optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/distributed/
- - vllm/engine/
- - vllm/executor/
- - vllm/compilation/
- - tests/distributed/
- - vllm/platforms/rocm.py
- commands:
- - pytest -v -s distributed/test_elastic_ep.py
-
-
-- label: Engine # 11.3m
- timeout_in_minutes: 35
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/
- - tests/engine
- - tests/test_sequence
- - tests/test_config
- - tests/test_logger
- - tests/test_vllm_port
- commands:
- - pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py
-
-
-- label: Engine (1 GPU) # TBD
+- label: Qwen3-30B-A3B-FP8-block Accuracy (4xH100-4xMI300) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/v1/engine/
- - tests/v1/engine/
- - vllm/platforms/rocm.py
- commands:
- - pytest -v -s v1/engine/test_preprocess_error_handling.py
- - pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py
-
-
-- label: e2e Scheduling (1 GPU) # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/v1/
- - tests/v1/e2e/general/
- - vllm/platforms/rocm.py
- commands:
- - pytest -v -s v1/e2e/general/test_async_scheduling.py
-
-
-- label: e2e Core (1 GPU) # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/v1/
- - tests/v1/e2e/
- - vllm/platforms/rocm.py
- commands:
- - pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py
-
-
-- label: Spec Decode Eagle # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/v1/spec_decode/
- - vllm/v1/worker/gpu/spec_decode/
- - vllm/model_executor/model_loader/
- - vllm/v1/sample/
- - vllm/model_executor/layers/
- - tests/v1/e2e/spec_decode/
- - vllm/platforms/rocm.py
- commands:
- - pytest -v -s v1/e2e/spec_decode -k "eagle_correctness"
-
-
-- label: Spec Decode Speculators + MTP # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/v1/spec_decode/
- - vllm/v1/worker/gpu/spec_decode/
- - vllm/model_executor/model_loader/
- - vllm/v1/sample/
- - vllm/model_executor/layers/
- - vllm/transformers_utils/configs/speculators/
- - tests/v1/e2e/spec_decode/
- - vllm/platforms/rocm.py
- commands:
- - pytest -v -s v1/e2e/spec_decode -k "speculators or mtp_correctness"
-
-
-- label: Spec Decode Ngram + Suffix # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/v1/spec_decode/
- - vllm/v1/worker/gpu/spec_decode/
- - vllm/model_executor/model_loader/
- - vllm/v1/sample/
- - vllm/model_executor/layers/
- - tests/v1/e2e/spec_decode/
- - vllm/platforms/rocm.py
- commands:
- - pytest -v -s v1/e2e/spec_decode -k "ngram or suffix"
-
-
-- label: Spec Decode Draft Model # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- optional: true
- working_dir: "/vllm-workspace/tests"
+ working_dir: "/vllm-workspace"
source_file_dependencies:
- - vllm/v1/spec_decode/
- - vllm/v1/worker/gpu/spec_decode/
+ - vllm/model_executor/models/
- vllm/model_executor/model_loader/
- - vllm/v1/sample/
- - vllm/model_executor/layers/
- - tests/v1/e2e/spec_decode/
- - vllm/platforms/rocm.py
- commands:
- - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
-
-
-- label: V1 e2e (2 GPUs) # 7.1m
- timeout_in_minutes: 12
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_2
- optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/
- - tests/v1/e2e
- commands:
- - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "tensor_parallelism"
-
-
-- label: V1 e2e (4 GPUs) # 52.6m
- timeout_in_minutes: 106
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_4
- optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/
- - tests/v1/e2e
- commands:
- - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "eagle_correctness_heavy"
-
-
-- label: V1 e2e (4xH100-4xMI325) # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_4
- optional: true
- source_file_dependencies:
- - vllm/v1/attention/backends/utils.py
- - vllm/v1/worker/gpu_model_runner.py
- - tests/v1/e2e/test_hybrid_chunked_prefill.py
- commands:
- - pytest -v -s v1/e2e/test_hybrid_chunked_prefill.py
-
-
-- label: V1 Spec Decode # TBD
- timeout_in_minutes: 40
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/
- - tests/v1/spec_decode
- commands:
- - pytest -v -s -m 'not slow_test' v1/spec_decode
-
-
-- label: V1 Sample + Logits # TBD
- timeout_in_minutes: 40
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/
- - tests/v1/sample
- - tests/v1/logits_processors
- - tests/v1/test_oracle.py
- - tests/v1/test_request.py
- - tests/v1/test_outputs.py
- commands:
- - pytest -v -s v1/sample
- - 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
-
-
-- label: V1 Core + KV + Metrics # TBD
- timeout_in_minutes: 40
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_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
- # - export HSA_NO_SCRATCH_RECLAIM=1
- - pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
-
-
-
-- label: Acceptance Length Test (Large Models) # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/v1/spec_decode/
- - vllm/model_executor/models/mlp_speculator.py
- - tests/v1/spec_decode/test_acceptance_length.py
- - vllm/platforms/rocm.py
- commands:
- - export VLLM_ALLOW_INSECURE_SERIALIZATION=1
- - pytest -v -s v1/spec_decode/test_acceptance_length.py
-
-
-- label: V1 attention (H100-MI325) # 14.5m
- timeout_in_minutes: 40
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/config/attention.py
- - vllm/model_executor/layers/attention
- - vllm/v1/attention
- - tests/v1/attention
+ - vllm/model_executor/layers/quantization/
+ - vllm/distributed/eplb
+ - vllm/model_executor/layers/fused_moe/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - .buildkite/scripts/scheduled_integration_test/
- vllm/_aiter_ops.py
- - vllm/envs.py
- vllm/platforms/rocm.py
commands:
- - pytest -v -s v1/attention
-
+ - bash .buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_block_ep_eplb.sh 0.8 200 8020
-- label: Batch Invariance (H100-MI325) # 5.2m
- timeout_in_minutes: 12
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+- label: Qwen3-Next-80B-A3B-Instruct MTP Async EPLB Accuracy # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
+ num_gpus: 4
optional: true
- working_dir: "/vllm-workspace/tests"
+ working_dir: "/vllm-workspace"
source_file_dependencies:
- - vllm/v1/attention
- - vllm/model_executor/layers
- - tests/v1/determinism/
+ - vllm/model_executor/models/
+ - vllm/model_executor/model_loader/
+ - vllm/v1/spec_decode/
+ - vllm/distributed/eplb
+ - vllm/model_executor/layers/fused_moe/
+ - vllm/model_executor/layers/quantization/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - .buildkite/scripts/scheduled_integration_test/
- vllm/_aiter_ops.py
- vllm/platforms/rocm.py
commands:
- - export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - pip install pytest-timeout pytest-forked
- - pytest -v -s v1/determinism/test_batch_invariance.py
- - pytest -v -s v1/determinism/test_rms_norm_batch_invariant.py
-
+ - bash .buildkite/scripts/scheduled_integration_test/qwen3_next_mtp_async_eplb.sh 0.8 1319 8040
-- label: V1 others (CPU) # 10.4m
- timeout_in_minutes: 28
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- no_gpu: true
+- label: LM Eval Large Models (8xH200-8xMI300) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_8
optional: true
+ num_gpus: 8
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/v1
+ - vllm/model_executor/models/
+ - vllm/model_executor/model_loader/
+ - vllm/model_executor/layers/quantization/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - vllm/model_executor/layers/layernorm.py
+ - csrc/
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
+ - tests/evals/
commands:
- - pytest -v -s -m 'cpu_test' v1/core
- - pytest -v -s v1/structured_output
- - pytest -v -s v1/test_serial_utils.py
- - pytest -v -s -m 'cpu_test' v1/kv_connector/unit
- - pytest -v -s -m 'cpu_test' v1/metrics
+ - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+ - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx.txt
+#--------------------------------------------------------- mi300 · examples ----------------------------------------------------------#
-- label: Examples # 24.5m
- timeout_in_minutes: 55
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+- label: Examples # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
optional: true
working_dir: "/vllm-workspace/examples"
source_file_dependencies:
@@ -1949,67 +1641,25 @@ steps:
- python3 basic/offline_inference/embed.py
- python3 basic/offline_inference/score.py
# Multi-modal models
- - 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
+ - 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
# Pooling models
- python3 pooling/embed/vision_embedding_offline.py --seed 0
# Features demo
- - 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
-
-
-- label: Platform Tests (CUDA) # 5.0m
- timeout_in_minutes: 9
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/
- - tests/cuda
- commands:
- - pytest -v -s cuda/test_cuda_context.py
- - pytest -v -s cuda/test_platform_no_cuda_init.py
+ - 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
+#---------------------------------------------------------- mi300 · kernels ----------------------------------------------------------#
-- label: PyTorch Compilation Passes Unit Tests # TBD
+- label: Kernels Attention Test %N # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/
- - tests/compile/passes
- commands:
- - pytest -s -v compile/passes --ignore compile/passes/distributed
-
-
-- label: Kernels Core Operation Test # 26.8m
- timeout_in_minutes: 38
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - csrc/
- - tests/kernels/core
- - tests/kernels/test_top_k_per_row.py
- - tests/kernels/test_concat_mla_q.py
- - vllm/model_executor/layers/rotary_embedding/
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
- commands:
- - pytest -v -s kernels/core kernels/test_top_k_per_row.py
-
-
-- label: Kernels Attention Test %N # 17.7m
- timeout_in_minutes: 28
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
parallelism: 2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -2023,29 +1673,26 @@ steps:
commands:
- pytest -v -s kernels/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
-
-- label: Kernels Quantization Test %N # 15.2m
- timeout_in_minutes: 24
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- parallelism: 2
+- label: Kernels Core Operation Test # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - csrc/quantization/
- - vllm/model_executor/layers/quantization
- - tests/kernels/quantization
- - tests/kernels/quantization/test_rocm_skinny_gemms.py
+ - csrc/
+ - tests/kernels/core
+ - tests/kernels/test_top_k_per_row.py
+ - tests/kernels/test_concat_mla_q.py
+ - vllm/model_executor/layers/rotary_embedding/
- vllm/_aiter_ops.py
- vllm/platforms/rocm.py
- - vllm/model_executor/kernels/
commands:
- - pytest -v -s kernels/quantization --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
-
+ - pytest -v -s kernels/core --ignore=kernels/core/test_minimax_reduce_rms.py kernels/test_concat_mla_q.py kernels/test_top_k_per_row.py
- label: Kernels MoE Test %N # TBD
- timeout_in_minutes: 19
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
parallelism: 4
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -2062,11 +1709,27 @@ steps:
- pytest -v -s kernels/moe --ignore=kernels/moe/test_modular_oai_triton_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
- pytest -v -s kernels/moe/test_modular_oai_triton_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
+- label: Kernels Quantization Test %N # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ parallelism: 2
+ working_dir: "/vllm-workspace/tests"
+ 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/
+ commands:
+ - pytest -v -s kernels/quantization --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
-- label: Kernels FP8 MoE Test # TBD
+- label: Kernels FP8 MoE Test (2xH100-2xMI300) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_2
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_2
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -2080,40 +1743,222 @@ steps:
commands:
- pytest -v -s kernels/moe/test_deepep_moe.py
+#----------------------------------------------------------- mi300 · lora ------------------------------------------------------------#
+
+- label: LoRA TP (Distributed) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
+ num_gpus: 4
+ working_dir: "/vllm-workspace/tests"
+ source_file_dependencies:
+ - vllm/lora
+ - tests/lora
+ - vllm/platforms/rocm.py
+ commands:
+ - export VLLM_WORKER_MULTIPROC_METHOD=spawn
+ - 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_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
+
+#---------------------------------------------------------- mi300 · models -----------------------------------------------------------#
+
+- label: Language Models Test (Extended Pooling) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ working_dir: "/vllm-workspace/tests"
+ source_file_dependencies:
+ - vllm/
+ - tests/models/language/pooling
+ commands:
+ - pytest -v -s models/language/pooling -m 'not core_model'
+
+- label: Language Models Tests (Standard) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ optional: true
+ torch_nightly: true
+ working_dir: "/vllm-workspace/tests"
+ source_file_dependencies:
+ - vllm/
+ - tests/models/language
+ commands:
+ - pip freeze | grep -E 'torch'
+ - pytest -v -s models/language -m 'core_model and (not slow_test)'
+
+- label: Multi-Modal Models (Extended Generation 1) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ optional: true
+ working_dir: "/vllm-workspace/tests"
+ source_file_dependencies:
+ - vllm/
+ - 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
+
+
+- label: Multi-Modal Models (Extended Generation 2) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ optional: true
+ working_dir: "/vllm-workspace/tests"
+ source_file_dependencies:
+ - vllm/
+ - 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)'
+
+
+- label: Multi-Modal Models (Extended Generation 3) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ optional: true
+ working_dir: "/vllm-workspace/tests"
+ source_file_dependencies:
+ - vllm/
+ - tests/models/multimodal/generation
+ commands:
+ - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
+ - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=1) and not core_model'
+
+- label: "Multi-Modal Models (Standard) 1: qwen2" # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ torch_nightly: true
+ optional: true
+ working_dir: "/vllm-workspace/tests"
+ source_file_dependencies:
+ - vllm/
+ - tests/models/multimodal
+ commands:
+ - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
+ - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen2"
+ - pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model
+
+- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl" # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ torch_nightly: true
+ optional: true
+ working_dir: "/vllm-workspace/tests"
+ source_file_dependencies:
+ - vllm/
+ - tests/models/multimodal/generation
+ - tests/models/multimodal/test_mapping.py
+ commands:
+ - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
+ - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "not qwen2 and not qwen3 and not gemma"
+ - pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model
+
+- label: "Multi-Modal Models (Standard) 4: other + whisper" # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ torch_nightly: true
+ working_dir: "/vllm-workspace/tests"
+ source_file_dependencies:
+ - vllm/
+ - tests/models/multimodal/generation
+ - tests/models/multimodal/test_mapping.py
+ commands:
+ - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
+ - pytest -v -s models/multimodal -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 -v -s 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
+
+- label: Multi-Modal Processor # 1h 42m
+ timeout_in_minutes: 138
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ optional: true
+ working_dir: "/vllm-workspace/tests"
+ source_file_dependencies:
+ - vllm/
+ - tests/models/multimodal
+ - tests/models/registry.py
+ commands:
+ - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
+ - pytest -v -s models/multimodal/processing/test_tensor_schema.py
+
+- label: Multi-Modal Processor (CPU) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ no_gpu: true
+ optional: true
+ working_dir: "/vllm-workspace/tests"
+ source_file_dependencies:
+ - vllm/
+ - tests/models/multimodal
+ - tests/models/registry.py
+ commands:
+ - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
+ - pytest -v -s models/multimodal/processing --ignore models/multimodal/processing/test_tensor_schema.py
-- label: ROCm AITER Ops Test # TBD
+- label: Quantized Models Test # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
+ - vllm/model_executor/layers/quantization
- vllm/_aiter_ops.py
- - vllm/envs.py
- vllm/platforms/rocm.py
- - tests/rocm/aiter/
- - vllm/v1/attention/backends/mla/rocm_aiter_mla.py
- - vllm/v1/attention/selector.py
+ - tests/models/quantization
+ - vllm/model_executor/model_loader/
commands:
- - pytest -v -s rocm/aiter/
-
+ - pytest -v -s models/quantization
-- label: Benchmarks # 8.2m
- timeout_in_minutes: 20
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+- label: Transformers Nightly Models # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
optional: true
- working_dir: "/vllm-workspace/.buildkite"
+ working_dir: "/vllm-workspace/"
source_file_dependencies:
- - benchmarks/
+ - vllm/model_executor/models/
+ - vllm/model_executor/model_loader/
+ - vllm/multimodal/
+ - vllm/model_executor/layers/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - vllm/_aiter_ops.py
- vllm/platforms/rocm.py
+ - tests/models/
+ - examples/
commands:
- - bash scripts/run-benchmarks.sh
+ - pip install --upgrade git+https://github.com/huggingface/transformers
+ - pytest -v -s tests/models/test_initialization.py
+ - 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
+ - VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/generate/multimodal/audio_language_offline.py --model-type whisper
+#------------------------------------------------------- mi300 · quantization --------------------------------------------------------#
-- label: Quantization # 36.1m
- timeout_in_minutes: 60
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+- label: Quantization # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- csrc/
@@ -2134,262 +1979,292 @@ steps:
- uv pip install --system conch-triton-kernels
- VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py
+#----------------------------------------------------------- mi300 · rocm ------------------------------------------------------------#
-- label: Language Models Tests (Standard) # 22.8m
- timeout_in_minutes: 38
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- optional: true
- torch_nightly: true
+- label: ROCm AITER Ops Test # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/models/language
+ - vllm/_aiter_ops.py
+ - vllm/envs.py
+ - vllm/platforms/rocm.py
+ - tests/rocm/aiter/
+ - vllm/v1/attention/backends/mla/rocm_aiter_mla.py
+ - vllm/v1/attention/selector.py
commands:
- - pip freeze | grep -E 'torch'
- - pytest -v -s models/language -m 'core_model and (not slow_test)'
+ - pytest -v -s rocm/aiter/
+#--------------------------------------------------------- mi300 · samplers ----------------------------------------------------------#
-- label: Language Models Tests (Hybrid) %N # 34.9m
- timeout_in_minutes: 55
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- torch_nightly: true
- parallelism: 2
+- label: Samplers Test # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/models/language/generation
+ - vllm/model_executor/layers
+ - vllm/sampling_metadata.py
+ - vllm/v1/sample/
+ - vllm/beam_search.py
+ - tests/samplers
+ - tests/conftest.py
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
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.5.2'
- - pytest -v -s models/language/generation -m hybrid_model --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB
+ - pytest -v -s samplers
+#------------------------------------------------------------ mi300 · misc ------------------------------------------------------------#
-- label: Language Models Test (Extended Generation) # 32.2m
- timeout_in_minutes: 55
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+- label: Python-only Installation # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/models/language/generation
+ - tests/standalone_tests/python_only_compile.sh
+ - setup.py
+ - vllm/platforms/rocm.py
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.5.2'
- - pytest -v -s models/language/generation -m '(not core_model) and (not hybrid_model)'
-
+ - bash standalone_tests/python_only_compile.sh
-- label: Multi-Modal Processor # 1h 42m
- timeout_in_minutes: 138
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- optional: true
+- label: Regression # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/models/multimodal
- - tests/models/registry.py
+ - tests/test_regression
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/processing/test_tensor_schema.py
+ - pip install modelscope
+ - pytest -v -s test_regression.py
+#--------------------------------------------------------- mi300 · ray_compat ---------------------------------------------------------#
-- label: "Multi-Modal Models (Standard) 1: qwen2" # TBD
+- label: Ray Dependency Compatibility Check # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- torch_nightly: true
- optional: true
- working_dir: "/vllm-workspace/tests"
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ working_dir: "/"
source_file_dependencies:
- - vllm/
- - tests/models/multimodal
+ - requirements/
+ - setup.py
+ - vllm/platforms/rocm.py
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen2"
- - pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model
+ - bash /vllm-workspace/.buildkite/scripts/check-ray-compatibility.sh
+#------------------------------------------------------------ mi300 · v1 -------------------------------------------------------------#
-- label: "Multi-Modal Models (Standard) 2: qwen3 + gemma" # TBD
+- label: Acceptance Length Test (Large Models) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- torch_nightly: true
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/models/multimodal
+ - vllm/v1/spec_decode/
+ - vllm/model_executor/models/mlp_speculator.py
+ - tests/v1/spec_decode/test_acceptance_length.py
+ - vllm/platforms/rocm.py
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen3 or gemma"
- - pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model
-
+ - export VLLM_ALLOW_INSECURE_SERIALIZATION=1
+ - pytest -v -s v1/spec_decode/test_acceptance_length.py -m slow_test
-- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl" # TBD
+- label: e2e Core (1 GPU) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- torch_nightly: true
+ 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
+ - vllm/v1/
+ - tests/v1/e2e/
+ - vllm/platforms/rocm.py
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "not qwen2 and not qwen3 and not gemma"
- - pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model
-
+ - pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py
-- label: "Multi-Modal Models (Standard) 4: other + whisper" # TBD
+- label: e2e Scheduling (1 GPU) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- torch_nightly: true
+ 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
+ - vllm/v1/
+ - tests/v1/e2e/general/
+ - vllm/platforms/rocm.py
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/processing
- - pytest -v -s 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
+ - pytest -v -s v1/e2e/general/test_async_scheduling.py
+- label: Engine (1 GPU) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ optional: true
+ working_dir: "/vllm-workspace/tests"
+ source_file_dependencies:
+ - vllm/v1/engine/
+ - tests/v1/engine/
+ - vllm/platforms/rocm.py
+ commands:
+ - pytest -v -s v1/engine/test_preprocess_error_handling.py
+ - pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py
-- label: Multi-Modal Models (Extended Generation 1) # 1h 2m
- timeout_in_minutes: 106
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+- label: Spec Decode Draft Model # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/models/multimodal/generation
- - tests/models/multimodal/test_mapping.py
+ - vllm/v1/spec_decode/
+ - vllm/v1/worker/gpu/spec_decode/
+ - vllm/model_executor/model_loader/
+ - vllm/v1/sample/
+ - vllm/model_executor/layers/
+ - tests/v1/e2e/spec_decode/
+ - vllm/platforms/rocm.py
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation -m 'not core_model' --ignore models/multimodal/generation/test_common.py
- - pytest -v -s models/multimodal/test_mapping.py
+ - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
+- label: Spec Decode Eagle # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ working_dir: "/vllm-workspace/tests"
+ source_file_dependencies:
+ - vllm/v1/spec_decode/
+ - vllm/v1/worker/gpu/spec_decode/
+ - vllm/model_executor/model_loader/
+ - vllm/v1/sample/
+ - vllm/model_executor/layers/
+ - tests/v1/e2e/spec_decode/
+ - vllm/platforms/rocm.py
+ commands:
+ - pytest -v -s v1/e2e/spec_decode -k "eagle_correctness"
-- label: Multi-Modal Models (Extended Generation 2) # TBD
+- label: Spec Decode Ngram + Suffix # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+ 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
+ - vllm/v1/spec_decode/
+ - vllm/v1/worker/gpu/spec_decode/
+ - vllm/model_executor/model_loader/
+ - vllm/v1/sample/
+ - vllm/model_executor/layers/
+ - tests/v1/e2e/spec_decode/
+ - vllm/platforms/rocm.py
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=0) and not core_model'
+ - pytest -v -s v1/e2e/spec_decode -k "ngram or suffix"
+- label: Spec Decode Speculators + MTP # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ optional: true
+ working_dir: "/vllm-workspace/tests"
+ source_file_dependencies:
+ - vllm/v1/spec_decode/
+ - vllm/v1/worker/gpu/spec_decode/
+ - vllm/model_executor/model_loader/
+ - vllm/v1/sample/
+ - vllm/model_executor/layers/
+ - vllm/transformers_utils/configs/speculators/
+ - tests/v1/e2e/spec_decode/
+ - vllm/platforms/rocm.py
+ commands:
+ - pytest -v -s v1/e2e/spec_decode -k "speculators or mtp_correctness"
-- label: Multi-Modal Models (Extended Generation 3) # TBD
+- label: V1 attention (H100-MI300) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- optional: true
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/models/multimodal/generation
+ - vllm/config/attention.py
+ - vllm/model_executor/layers/attention
+ - vllm/v1/attention
+ - tests/v1/attention
+ - vllm/_aiter_ops.py
+ - vllm/envs.py
+ - vllm/platforms/rocm.py
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=1) and not core_model'
-
+ - pytest -v -s v1/attention
-- label: Multi-Modal Models (Extended Pooling) # TBD
+- label: V1 Core + KV + Metrics # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/models/multimodal/pooling
+ - 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:
- - pytest -v -s models/multimodal/pooling -m 'not core_model'
-
+ - 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
+ # - export HSA_NO_SCRATCH_RECLAIM=1
+ - pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
-- label: Quantized Models Test # 21.4m
- timeout_in_minutes: 38
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+- label: V1 others (CPU) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
+ no_gpu: true
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/model_executor/layers/quantization
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
- - tests/models/quantization
- - vllm/model_executor/model_loader/
+ - vllm/
+ - tests/v1
commands:
- - pytest -v -s models/quantization
-
+ - pytest -v -s -m 'cpu_test' v1/core
+ - pytest -v -s v1/structured_output
+ - pytest -v -s v1/test_serial_utils.py
+ - pytest -v -s -m 'cpu_test' v1/kv_connector/unit
+ - pytest -v -s -m 'cpu_test' v1/metrics
-- label: Transformers Nightly Models # 50.9m
- timeout_in_minutes: 102
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+- label: V1 Sample + Logits # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_1
optional: true
- working_dir: "/vllm-workspace/"
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/multimodal/
- - vllm/model_executor/layers/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
- - tests/models/
- - examples/
+ - vllm/
+ - tests/v1/sample
+ - tests/v1/logits_processors
+ - tests/v1/test_oracle.py
+ - tests/v1/test_request.py
+ - tests/v1/test_outputs.py
commands:
- - pip install --upgrade git+https://github.com/huggingface/transformers
- - pytest -v -s tests/models/test_initialization.py
- - 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/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
-
+ - pytest -v -s v1/sample
+ - 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
-- label: Quantized MoE Test (B200-MI325) # TBD
+- label: Distributed DP Tests (2 GPUs) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- working_dir: "/vllm-workspace/"
- source_file_dependencies:
- - tests/quantization/test_gfx3xx_moe.py
- - vllm/model_executor/models/deepseek_v2.py
- - vllm/model_executor/models/gpt_oss.py
- - vllm/model_executor/models/llama4.py
- - vllm/model_executor/layers/fused_moe
- - vllm/model_executor/layers/quantization/compressed_tensors
- - vllm/model_executor/layers/quantization/modelopt.py
- - vllm/model_executor/layers/quantization/mxfp4.py
- - vllm/v1/attention/backends/triton_attn.py
- - vllm/v1/attention/backends/rocm_attn.py
- - vllm/v1/attention/backends/rocm_aiter_fa.py
- - vllm/v1/attention/backends/mla/
- - vllm/v1/attention/selector.py
- - vllm/model_executor/layers/layernorm.py
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
- - vllm/model_executor/model_loader/
- commands:
- - pytest -s -v tests/quantization/test_gfx3xx_moe.py
-
-
-- label: Distributed DP Tests (2 GPUs) # 56.1m
- timeout_in_minutes: 102
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_2
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_2
num_gpus: 2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -2403,158 +2278,115 @@ 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
- DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py
-
-- label: Distributed Compile + RPC Tests (2 GPUs) # 56.1m
- timeout_in_minutes: 102
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_2
+- label: Distributed Tests (2xH100-2xMI300) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_2
num_gpus: 2
- working_dir: "/vllm-workspace/tests"
+ working_dir: "/vllm-workspace/"
source_file_dependencies:
- - vllm/compilation/
- vllm/distributed/
- - vllm/engine/
- - vllm/executor/
- - vllm/worker/worker_base.py
- - vllm/v1/engine/
- - vllm/v1/worker/
- - tests/compile/fullgraph/test_basic_correctness.py
- - tests/compile/test_wrapper.py
- - tests/entrypoints/llm/test_collective_rpc.py
+ - vllm/v1/distributed/
+ - vllm/model_executor/layers/fused_moe/
+ - tests/v1/distributed/test_dbo.py
+ - vllm/_aiter_ops.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
-
+ - 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
+ - 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 Torchrun + Shutdown Tests (2 GPUs) # 56.1m
- timeout_in_minutes: 102
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_2
+- label: Metrics, Tracing (2 GPUs) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_2
num_gpus: 2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/distributed/
- - vllm/engine/
- - vllm/executor/
- - vllm/worker/worker_base.py
- - vllm/v1/engine/
- - vllm/v1/worker/
- - tests/distributed/
- - tests/v1/shutdown
- - tests/v1/worker/test_worker_memory_snapshot.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/v1/tracing
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
- - pytest -v -s v1/worker/test_worker_memory_snapshot.py
-
+ - "pip install \
+ 'opentelemetry-sdk>=1.26.0' \
+ 'opentelemetry-api>=1.26.0' \
+ 'opentelemetry-exporter-otlp>=1.26.0' \
+ 'opentelemetry-semantic-conventions-ai>=0.4.1'"
+ - pytest -v -s v1/tracing
-- label: Distributed Model Tests (2 GPUs) # 19.3m
- timeout_in_minutes: 38
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_2
- num_gpus: 2
+- label: V1 e2e (2 GPUs) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_2
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/model_executor/model_loader/sharded_state_loader.py
- - vllm/model_executor/models/
- - vllm/model_executor/layers/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
- - tests/basic_correctness/
- - tests/model_executor/model_loader/test_sharded_state_loader.py
- - tests/models/
+ - vllm/
+ - tests/v1/e2e
commands:
- - TARGET_TEST_SUITE=L4 pytest basic_correctness/ -v -s -m 'distributed(num_gpus=2)'
- - CUDA_VISIBLE_DEVICES=0,1 pytest -v -s model_executor/model_loader/test_sharded_state_loader.py -m '(not slow_test)'
- - pytest models/test_transformers.py -v -s -m 'distributed(num_gpus=2)'
- - pytest models/language -v -s -m 'distributed(num_gpus=2)'
- - pytest models/multimodal -v -s -m 'distributed(num_gpus=2)' --ignore models/multimodal/generation/test_whisper.py
- - VLLM_WORKER_MULTIPROC_METHOD=spawn pytest models/multimodal/generation/test_whisper.py -v -s -m 'distributed(num_gpus=2)'
-
+ - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "tensor_parallelism"
-- label: LoRA TP (Distributed) # 9.8m
- timeout_in_minutes: 18
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_4
+- label: CrossLayer KV layout Distributed NixlConnector PD accuracy tests (4 GPUs) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
num_gpus: 4
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/lora
- - tests/lora
+ - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
+ - tests/v1/kv_connector/nixl_integration/
- vllm/platforms/rocm.py
commands:
- - export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - 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_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
-
-
-- label: Weight Loading Multiple GPU # 7.5m
- timeout_in_minutes: 14
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_2
- num_gpus: 2
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/
- - tests/weight_loading
- commands:
- - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-amd.txt
-
+ - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
+ - CROSS_LAYERS_BLOCKS=True ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
-- label: Weight Loading Multiple GPU - Large Models # 12.6m
- timeout_in_minutes: 26
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_2
- num_gpus: 2
- optional: true
+- label: Distributed DP Tests (4 GPUs) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
+ num_gpus: 4
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/weight_loading
+ - vllm/distributed/
+ - tests/v1/distributed
+ - tests/v1/engine/test_engine_core_client.py
+ - tests/distributed/test_utils
+ - vllm/platforms/rocm.py
commands:
- - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-large-amd.txt
-
+ - 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
+ - TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_internal_lb_dp.py
+ - TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_hybrid_lb_dp.py
+ - pytest -v -s v1/engine/test_engine_core_client.py::test_kv_cache_events_dp
+ - pytest -v -s distributed/test_utils.py
-- label: Ray Dependency Compatibility Check # TBD
+- label: Distributed NixlConnector PD accuracy (4 GPUs) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
+ num_gpus: 4
optional: true
- working_dir: "/"
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - requirements/
- - setup.py
+ - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
+ - tests/v1/kv_connector/nixl_integration/
- vllm/platforms/rocm.py
commands:
- - bash /vllm-workspace/.buildkite/scripts/check-ray-compatibility.sh
-
+ - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
+ - ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
-- label: Distributed NixlConnector PD accuracy (4 GPUs) # 27.4m
- timeout_in_minutes: 44
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_4
+- label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
num_gpus: 4
- optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
@@ -2562,15 +2394,13 @@ steps:
- vllm/platforms/rocm.py
commands:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
- - ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
-
+ - DP_EP=1 ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
-- label: CrossLayer KV layout Distributed NixlConnector PD accuracy tests (4 GPUs) # TBD
+- label: Hyrbid SSM NixlConnector PD accuracy tests (4 GPUs) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_4
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
num_gpus: 4
- optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
@@ -2578,166 +2408,157 @@ steps:
- vllm/platforms/rocm.py
commands:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
- - CROSS_LAYERS_BLOCKS=True ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
-
+ - HYBRID_SSM=1 ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
-- label: Distributed Tests (4 GPUs)(A100-MI325) # 20.9m
- timeout_in_minutes: 37
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_4
- num_gpus: 4
+- label: V1 e2e (4 GPUs) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
+ - tests/v1/e2e
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)'
- - pytest -v -s -x lora/test_mixtral.py
-
+ - pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "eagle_correctness_heavy"
-- label: Distributed Tests (2 GPUs)(H100-MI325) # TBD
+- label: V1 e2e (4xH100-4xMI300) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_2
- num_gpus: 2
- working_dir: "/vllm-workspace/"
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_4
+ optional: true
source_file_dependencies:
- - vllm/distributed/
- - vllm/v1/distributed/
- - vllm/model_executor/layers/fused_moe/
- - tests/v1/distributed/test_dbo.py
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/v1/attention/backends/utils.py
+ - vllm/v1/worker/gpu_model_runner.py
+ - tests/v1/e2e/test_hybrid_chunked_prefill.py
commands:
- - export TORCH_NCCL_BLOCKING_WAIT=1
- - pytest -v -s tests/v1/distributed/test_dbo.py
+ - pytest -v -s v1/e2e/test_hybrid_chunked_prefill.py
+#------------------------------------------------------ mi300 · weight_loading -------------------------------------------------------#
-- label: Distributed Compile Unit Tests (2xH100-2xMI325) # 14.3m
- timeout_in_minutes: 32
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_2
+- label: Weight Loading Multiple GPU # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_2
num_gpus: 2
- working_dir: "/vllm-workspace/"
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/compilation/
- - vllm/model_executor/layers
- - tests/compile/passes/distributed/
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/weight_loading
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/passes/distributed/test_sequence_parallelism.py
- # TODO: this test is not supported on ROCm, there are aiter kernels for this.
- # - pytest -v -s tests/compile/passes/distributed/test_fusion_all_reduce.py
- # - pytest -v -s tests/compile/distributed/test_fusions_e2e.py::test_tp2_attn_quant_allreduce_rmsnorm
- # - "VLLM_TEST_CLEAN_GPU_MEMORY=1 pytest -v -s tests/compile/distributed/test_fusions_e2e.py -k 'not Llama-4'"
-
+ - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-amd.txt
-- label: LM Eval Small Models # 13.3m
- timeout_in_minutes: 23
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
+- label: Weight Loading Multiple GPU - Large Models # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
+ agent_pool: mi300_2
+ num_gpus: 2
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - csrc/
- - vllm/model_executor/layers/quantization
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/weight_loading
commands:
- - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-small.txt
+ - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-large-amd.txt
+#########################################################################################################################################
+# #
+# MI325 (gfx942) tests #
+# #
+#########################################################################################################################################
-- label: LM Eval Small Models (MI325) # TBD
+#---------------------------------------------------------- mi325 · compile ----------------------------------------------------------#
+
+- label: Distributed Compile + RPC Tests (2 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
+ agent_pool: mi325_2
+ num_gpus: 2
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - csrc/
- - vllm/model_executor/layers/quantization
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - vllm/_aiter_ops.py
+ - vllm/compilation/
+ - vllm/distributed/
+ - vllm/engine/
+ - vllm/executor/
+ - vllm/worker/worker_base.py
+ - vllm/v1/engine/
+ - vllm/v1/worker/
+ - tests/compile/fullgraph/test_basic_correctness.py
+ - tests/compile/test_wrapper.py
+ - tests/entrypoints/llm/test_collective_rpc.py
- vllm/platforms/rocm.py
commands:
- - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-small-rocm.txt
+ - 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
+#-------------------------------------------------------- mi325 · distributed --------------------------------------------------------#
-- label: LM Eval Small Models (B200-MI325) # TBD
+- label: Distributed Torchrun + Shutdown Tests (2 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
agent_pool: mi325_2
+ num_gpus: 2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - csrc/
- - vllm/model_executor/layers/quantization
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - vllm/_aiter_ops.py
+ - vllm/distributed/
+ - vllm/engine/
+ - vllm/executor/
+ - vllm/worker/worker_base.py
+ - vllm/v1/engine/
+ - vllm/v1/worker/
+ - tests/distributed/
+ - tests/v1/shutdown
+ - tests/v1/worker/test_worker_memory_snapshot.py
- vllm/platforms/rocm.py
commands:
- - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx-fp8-and-mixed.txt
-
+ - 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
+ - pytest -v -s v1/worker/test_worker_memory_snapshot.py
-- label: LM Eval Large Models (H200-MI325) # TBD
+- label: Distributed Compile + Comm (4 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_8
- optional: true
- num_gpus: 8
+ agent_pool: mi325_4
+ num_gpus: 4
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/model_executor/layers/quantization/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - vllm/model_executor/layers/layernorm.py
- - csrc/
- - vllm/_aiter_ops.py
+ - vllm/distributed/
+ - tests/distributed/test_pynccl
+ - tests/distributed/test_events
+ - tests/compile/fullgraph/test_basic_correctness.py
+ - tests/distributed/test_symm_mem_allreduce.py
+ - tests/distributed/test_multiproc_executor.py
- vllm/platforms/rocm.py
- - tests/evals/
commands:
- - export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx.txt
+ - pytest -v -s compile/fullgraph/test_basic_correctness.py
+ - pytest -v -s distributed/test_pynccl.py
+ - pytest -v -s distributed/test_events.py
+ - pytest -v -s distributed/test_symm_mem_allreduce.py
+ - pytest -v -s distributed/test_multiproc_executor.py::test_multiproc_executor_multi_node
+#---------------------------------------------------------- mi325 · engine -----------------------------------------------------------#
-- label: LM Eval Large Models (4 GPUs)(FP8) # 24.8m
- timeout_in_minutes: 42
+- label: Engine # TBD
+ timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_4
- num_gpus: 4
- optional: true
- working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
+ agent_pool: mi325_1
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - csrc/
- - vllm/model_executor/layers/quantization
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/engine
+ - tests/test_sequence
+ - tests/test_config
+ - tests/test_logger
+ - tests/test_vllm_port
commands:
- - export VLLM_USE_DEEP_GEMM=0
- - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-rocm-fp8.txt --tp-size=4
+ - pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py
+#----------------------------------------------------------- mi325 · evals -----------------------------------------------------------#
-- label: LM Eval Large Models (4 GPUs)(A100-MI325) # 17.3m
- timeout_in_minutes: 27
+- label: LM Eval Large Models (4xH100-4xMI325) # TBD
+ timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
agent_pool: mi325_4
num_gpus: 4
@@ -2753,11 +2574,10 @@ steps:
- vllm/_aiter_ops.py
- vllm/platforms/rocm.py
commands:
- - export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large.txt --tp-size=4
-
+ - export VLLM_USE_DEEP_GEMM=0
+ - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-rocm-fp8.txt --tp-size=4
-- label: ROCm LM Eval Large Models (8 Card) # TBD
+- label: ROCm LM Eval Large Models (8 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
agent_pool: mi325_8
@@ -2778,237 +2598,150 @@ steps:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-rocm.txt --tp-size=8
+#---------------------------------------------------------- mi325 · models -----------------------------------------------------------#
-- label: GPQA Eval (GPT-OSS) (H100-MI325) # TBD
+- label: Language Models Test (Extended Generation) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_2
- num_gpus: 2
- optional: true
+ agent_pool: mi325_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - csrc/
- - vllm/model_executor/layers/quantization
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - vllm/model_executor/layers/fused_moe/
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
- - tests/evals/gpt_oss/
+ - vllm/
+ - tests/models/language/generation
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-gfx942.txt
-
+ - 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: DeepSeek V2-Lite Accuracy # 6.7m
- timeout_in_minutes: 12
+- label: Language Models Tests (Hybrid) %N # TBD
+ timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_4
- num_gpus: 4
+ agent_pool: mi325_1
+ torch_nightly: true
+ parallelism: 2
optional: true
- working_dir: "/vllm-workspace"
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/distributed/eplb
- - vllm/model_executor/layers/fused_moe/
- - vllm/model_executor/layers/quantization/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/backends/mla/
- - vllm/v1/attention/selector.py
- - .buildkite/scripts/scheduled_integration_test/
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/models/language/generation
commands:
- - bash .buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_ep_eplb.sh 0.25 200 8010
-
+ - 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: DeepSeek V2-Lite Prefetch Offload Accuracy (H100-MI325) # TBD
+- label: Multi-Modal Models (Extended Pooling) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
agent_pool: mi325_1
- num_gpus: 1
- working_dir: "/vllm-workspace"
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/model_executor/layers/fused_moe/
- - vllm/model_executor/layers/quantization/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/backends/mla/
- - vllm/v1/attention/selector.py
- - .buildkite/scripts/scheduled_integration_test/
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/models/multimodal/pooling
commands:
- - bash .buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_prefetch_offload.sh 0.25 200 8030
-
+ - pytest -v -s models/multimodal/pooling -m 'not core_model'
-- label: Qwen3-30B-A3B-FP8-block Accuracy # 6.4m
- timeout_in_minutes: 11
+- label: "Multi-Modal Models (Standard) 2: qwen3 + gemma" # TBD
+ timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_4
+ agent_pool: mi325_1
+ torch_nightly: true
optional: true
- working_dir: "/vllm-workspace"
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/model_executor/layers/quantization/
- - vllm/distributed/eplb
- - vllm/model_executor/layers/fused_moe/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - .buildkite/scripts/scheduled_integration_test/
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/models/multimodal
commands:
- - bash .buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_block_ep_eplb.sh 0.8 200 8020
+ - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
+ - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen3 or gemma"
+ - pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model
+#------------------------------------------------------------ mi325 · v1 -------------------------------------------------------------#
-- label: Qwen3-Next-80B-A3B-Instruct MTP Async EPLB Accuracy # 10.9m
- timeout_in_minutes: 22
+- label: V1 Spec Decode # TBD
+ timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_4
- num_gpus: 4
- optional: true
- working_dir: "/vllm-workspace"
+ agent_pool: mi325_1
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/v1/spec_decode/
- - vllm/distributed/eplb
- - vllm/model_executor/layers/fused_moe/
- - vllm/model_executor/layers/quantization/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - .buildkite/scripts/scheduled_integration_test/
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/v1/spec_decode
commands:
- - bash .buildkite/scripts/scheduled_integration_test/qwen3_next_mtp_async_eplb.sh 0.8 1319 8040
-
-##### .buildkite/test_areas/compile.yaml #####
-# Slowly setting up the tests so that it is also easier for the
-# CI team to review and upstream to the pipelinev2.
-# The following tests are important for vLLM IR Ops refactoring,
-# which affects fusion passes on ROCm. So we have to
-# enable them as as soon as possible.
+ - pytest -v -s -m 'not slow_test' v1/spec_decode
-## TODO: Enable the test in this group
-# # corresponds to .buildkite/test_areas/compile.yaml
-# - label: Fusion and Compile Unit Tests (2xB200-2xMI325) # TBD
-# timeout_in_minutes: 180
-# mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325, tj]
-# agent_pool: mi325_1 # changed to 1 GPU until the fusion all reduce is enabled then only revert back to 2 GPUs
-# num_gpus: 1
-# working_dir: "/vllm-workspace/"
-# source_file_dependencies:
-# - csrc/quantization/fp4/
-# - vllm/model_executor/layers/quantization/
-# - vllm/model_executor/layers/layernorm.py
-# - vllm/model_executor/layers/activation.py
-# - vllm/model_executor/layers/attention/attention.py
-# - vllm/v1/attention/backends/flashinfer.py
-# - vllm/compilation/ # TODO(luka) limit to vllm/compilation/passes
-# - tests/compile/test_fusion_attn.py
-# - tests/compile/test_silu_mul_quant_fusion.py
-# - tests/compile/distributed/test_fusion_all_reduce.py
-# - tests/compile/fullgraph/test_full_graph.py
-# commands:
-# - rocm-smi
-# # we run all backend tests on ROCm
-# # These two tests are covered in "PyTorch Compilation Passes Unit Tests"
-# # - "pytest -v -s tests/compile/passes/test_fusion_attn.py"
-# # - "pytest -v -s tests/compile/passes/test_silu_mul_quant_fusion.py"
-# # TODO: this test is not supported on ROCm, there are aiter kernels for this.
-# # - pytest -v -s tests/compile/passes/distributed/test_fusion_all_reduce.py
-# # TODO: find out more details
-# # - pytest -v -s tests/compile/fullgraph/test_full_graph.py::test_fp8_kv_scale_compile
+#########################################################################################################################################
+# #
+# MI355 (gfx950) tests #
+# #
+#########################################################################################################################################
+#-------------------------------------------------------- mi355 · benchmarks ---------------------------------------------------------#
-- label: Fusion E2E Quick (H100-MI325) # TBD
+- label: Attention Benchmarks Smoke Test (B200-MI355) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- num_gpus: 1
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
+ agent_pool: mi355_2
+ num_gpus: 2
working_dir: "/vllm-workspace/"
source_file_dependencies:
- - csrc/quantization/
- - vllm/model_executor/
+ - benchmarks/attention_benchmarks/
- vllm/v1/attention/
- - vllm/compilation/
- - tests/compile/fusions_e2e/
- vllm/_aiter_ops.py
- vllm/platforms/rocm.py
commands:
- - rocm-smi
- # Run all models and attn backends but only Inductor partition and native custom ops
- - "pytest -v -s tests/compile/fusions_e2e/test_tp1_quant.py -k 'inductor_partition and not +rms_norm and not +quant_fp8'"
- # Different from CUDA, Qwen requires +rms_norm and +quant_fp8 as rms+quant fusion is only supported on AITER
- - "pytest -v -s tests/compile/fusions_e2e/test_tp1_quant.py -k 'inductor_partition and +rms_norm and +quant_fp8 and qwen3'"
+ - python3 benchmarks/attention_benchmarks/benchmark.py --backends ROCM_ATTN ROCM_AITER_FA ROCM_AITER_UNIFIED_ATTN --batch-specs "8q1s1k" --repeats 1 --warmup-iters 1
+#-------------------------------------------------------- mi355 · distributed --------------------------------------------------------#
-- label: Fusion E2E Config Sweep (H100-MI325) # TBD
+- label: Distributed Tests (2xH100-2xMI355) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi325]
- agent_pool: mi325_1
- num_gpus: 1
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
+ agent_pool: mi355_2
+ num_gpus: 2
+ optional: true
working_dir: "/vllm-workspace/"
source_file_dependencies:
- - csrc/quantization/
- - vllm/compilation/
- - vllm/model_executor/layers/layernorm.py
- - vllm/model_executor/layers/activation.py
- - vllm/model_executor/layers/attention/attention.py
- - vllm/model_executor/layers/quantization/input_quant_fp8.py
- - tests/compile/fusions_e2e/
+ - vllm/distributed/
+ - vllm/v1/distributed/
+ - vllm/model_executor/layers/fused_moe/
+ - vllm/v1/attention/backends/
+ - 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
- vllm/_aiter_ops.py
- vllm/platforms/rocm.py
commands:
- - rocm-smi
- - pytest -v -s tests/compile/fusions_e2e/test_tp1_quant.py -k "llama-3"
-
-## There are no ops on ROCm for these tests.
-## The test still passes but the logs are not useful.
-## fused ops just call torch.ops.symm_mem which
-## exists in ROCm even though they don't work
-# - label: AsyncTP Correctness Tests (2xH100-2xMI325)
-# - label: Fusion E2E TP2 Quick (H100-MI325)
-# - label: Fusion E2E TP2 AsyncTP Config Sweep (H100-MI325)
-# - label: Fusion E2E TP2 (B200-MI325)
-# - label: Sequence Parallel Correctness Tests (2xH100-2xMI325)
-
+ - pytest -v -s tests/distributed/test_context_parallel.py
+ - pytest -v -s tests/v1/distributed/test_dbo.py
-#####################################################################################################################################
-# #
-# gfx950 #
-# #
-#####################################################################################################################################
+#-------------------------------------------------------- mi355 · entrypoints --------------------------------------------------------#
-- label: Entrypoints Integration (API Server openai - Part 1) # TBD
+- label: Entrypoints Integration (API Server 2) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
+ optional: true
fast_check: true
torch_nightly: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/entrypoints/openai
- - tests/entrypoints/test_chat_utils
+ - tests/entrypoints/rpc
+ - tests/entrypoints/serve/instrumentator
+ - tests/tool_use
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
-
+ - pytest -v -s entrypoints/serve/instrumentator
+ - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
+ - pytest -v -s tool_use
-- label: Entrypoints Integration (API Server openai - Part 2) # TBD
+- label: Entrypoints Integration (API Server openai - Part 1) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
fast_check: true
torch_nightly: true
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -3016,17 +2749,15 @@ steps:
- tests/entrypoints/test_chat_utils
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - pytest -v -s entrypoints/openai/completion --ignore=entrypoints/openai/completion/test_tensorizer_entrypoint.py
- - pytest -v -s entrypoints/openai/speech_to_text/
- - pytest -v -s entrypoints/test_chat_utils.py
-
+ - pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
-- label: Entrypoints Integration (API Server openai - Part 3) # TBD
+- label: Entrypoints Integration (API Server openai - Part 2) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
fast_check: true
torch_nightly: true
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -3034,28 +2765,24 @@ 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/speech_to_text/ --ignore=entrypoints/openai/correctness/ --ignore=entrypoints/openai/tool_parsers/ --ignore=entrypoints/openai/responses
-
+ - 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 2) # TBD
+- label: Entrypoints Integration (API Server openai - Part 3) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
- optional: true
fast_check: true
torch_nightly: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/entrypoints/rpc
- - tests/entrypoints/serve/instrumentator
- - tests/tool_use
+ - tests/entrypoints/openai
+ - tests/entrypoints/test_chat_utils
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- - pytest -v -s entrypoints/serve/instrumentator
- - PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/rpc
- - pytest -v -s tool_use
-
+ - pytest -v -s entrypoints/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
@@ -3071,95 +2798,112 @@ steps:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/pooling
+#----------------------------------------------------------- mi355 · evals -----------------------------------------------------------#
-- label: Regression # TBD
+- label: GPQA Eval (GPT-OSS) (2xB200-2xMI355) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_1
+ agent_pool: mi355_2
+ num_gpus: 2
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/test_regression
+ - csrc/
+ - vllm/model_executor/layers/quantization
+ - vllm/model_executor/models/
+ - vllm/model_executor/model_loader/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - vllm/model_executor/layers/fused_moe/
+ - tests/evals/gpt_oss/
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
commands:
- - pip install modelscope
- - pytest -v -s test_regression.py
-
+ - 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-gfx950.txt
-- label: V1 Spec Decode # TBD
- timeout_in_minutes: 60
+- label: LM Eval Qwen3-5 Models (B200-MI355) # TBD
+ timeout_in_minutes: 120
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_1
+ agent_pool: mi355_2
+ num_gpus: 2
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/v1/spec_decode
+ - vllm/model_executor/models/qwen3_5.py
+ - vllm/model_executor/models/qwen3_5_mtp.py
+ - vllm/transformers_utils/configs/qwen3_5.py
+ - vllm/transformers_utils/configs/qwen3_5_moe.py
+ - vllm/model_executor/models/qwen.py
+ - vllm/model_executor/models/qwen2.py
+ - vllm/model_executor/models/qwen3.py
+ - vllm/model_executor/models/qwen3_next.py
+ - vllm/model_executor/models/qwen3_next_mtp.py
+ - vllm/model_executor/layers/fla/ops/
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
commands:
- - pytest -v -s -m 'not slow_test' v1/spec_decode
-
+ - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-qwen35-mi355.txt
-- label: V1 Sample + Logits # TBD
- timeout_in_minutes: 60
+- label: LM Eval Small Models (2xB200-2xMI355) # TBD
+ timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_1
+ agent_pool: mi355_2
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/v1/sample
- - tests/v1/logits_processors
- - tests/v1/test_oracle.py
- - tests/v1/test_request.py
- - tests/v1/test_outputs.py
+ - csrc/
+ - vllm/model_executor/layers/quantization
+ - vllm/model_executor/models/
+ - vllm/model_executor/model_loader/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
commands:
- - pytest -v -s v1/sample
- - 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
-
+ - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx-fp8-and-mixed.txt
-- label: V1 Core + KV + Metrics # TBD
- timeout_in_minutes: 60
+- label: Qwen3-30B-A3B-FP8-block Accuracy (B200-MI355) # TBD
+ timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_1
- working_dir: "/vllm-workspace/tests"
+ agent_pool: mi355_2
+ num_gpus: 2
+ working_dir: "/vllm-workspace"
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
+ - vllm/model_executor/models/
+ - vllm/model_executor/model_loader/
+ - vllm/model_executor/layers/quantization/
+ - vllm/model_executor/layers/fused_moe/
+ - vllm/distributed/eplb
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
+ - .buildkite/scripts/scheduled_integration_test/
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.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
-
+ - bash .buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_block_ep_eplb.sh 0.8 200 8020 2 1
-- label: V1 attention (B200-MI355) # TBD
+- label: LM Eval Large Models (4xH100-4xMI355) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_1
- working_dir: "/vllm-workspace/tests"
+ agent_pool: mi355_4
+ num_gpus: 4
+ optional: true
+ working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
source_file_dependencies:
- - vllm/config/attention.py
- - vllm/model_executor/layers/attention
- - vllm/v1/attention
- - tests/v1/attention
+ - csrc/
+ - vllm/model_executor/layers/quantization
+ - vllm/model_executor/models/
+ - vllm/model_executor/model_loader/
+ - vllm/v1/attention/backends/
+ - vllm/v1/attention/selector.py
- vllm/_aiter_ops.py
- - vllm/envs.py
- vllm/platforms/rocm.py
commands:
- - pytest -v -s v1/attention
+ - export VLLM_USE_DEEP_GEMM=0
+ - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-rocm-fp8.txt --tp-size=4
+#--------------------------------------------------------- mi355 · examples ----------------------------------------------------------#
- label: Examples # TBD
timeout_in_minutes: 180
@@ -3181,19 +2925,44 @@ steps:
- python3 basic/offline_inference/embed.py
- python3 basic/offline_inference/score.py
# Multi-modal models
- - 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
+ - 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
# Pooling models
- python3 pooling/embed/vision_embedding_offline.py --seed 0
# Features demo
- - 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
+ - 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
+#---------------------------------------------------------- mi355 · kernels ----------------------------------------------------------#
+
+- label: Kernels (B200-MI355) # TBD
+ timeout_in_minutes: 180
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
+ agent_pool: mi355_1
+ working_dir: "/vllm-workspace/"
+ source_file_dependencies:
+ - csrc/quantization/fp4/
+ - csrc/attention/mla/
+ - csrc/quantization/cutlass_w8a8/moe/
+ - vllm/model_executor/layers/fused_moe/cutlass_moe.py
+ - vllm/v1/attention/backends/triton_attn.py
+ - vllm/v1/attention/backends/rocm_attn.py
+ - vllm/v1/attention/backends/rocm_aiter_fa.py
+ - vllm/v1/attention/backends/rocm_aiter_unified_attn.py
+ - vllm/v1/attention/backends/mla/aiter_triton_mla.py
+ - vllm/v1/attention/backends/mla/rocm_aiter_mla.py
+ - vllm/v1/attention/selector.py
+ - vllm/platforms/rocm.py
+ - vllm/_aiter_ops.py
+ commands:
+ - rocm-smi
+ - python3 examples/basic/offline_inference/chat.py
+ - pytest -v -s tests/kernels/attention/test_attention_selector.py
- label: Kernels Attention Test %N # TBD
timeout_in_minutes: 180
@@ -3213,25 +2982,6 @@ steps:
commands:
- pytest -v -s kernels/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
-
-- label: Kernels Quantization Test %N # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_1
- parallelism: 2
- working_dir: "/vllm-workspace/tests"
- 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/
- commands:
- - pytest -v -s kernels/quantization --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
-
-
- label: Kernels MoE Test %N # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
@@ -3252,54 +3002,40 @@ steps:
- pytest -v -s kernels/moe --ignore=kernels/moe/test_modular_oai_triton_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
- pytest -v -s kernels/moe/test_modular_oai_triton_moe.py --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
-
-- label: Kernels FP8 MoE Test # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_2
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - csrc/moe/
- - csrc/quantization/w8a8/cutlass/moe/
- - vllm/model_executor/layers/fused_moe/
- - tests/kernels/moe/test_deepep_moe.py
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
- - vllm/envs.py
- commands:
- - pytest -v -s kernels/moe/test_deepep_moe.py
-
-
-- label: Quantization # TBD
+- label: Kernels Quantization Test %N # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
+ parallelism: 2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - csrc/
+ - csrc/quantization/
- vllm/model_executor/layers/quantization
- - tests/quantization
+ - tests/kernels/quantization
+ - tests/kernels/quantization/test_rocm_skinny_gemms.py
- vllm/_aiter_ops.py
- vllm/platforms/rocm.py
+ - vllm/model_executor/kernels/
commands:
- - uv pip install --system torchao==0.17.0
- - uv pip install --system conch-triton-kernels
- - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py
-
+ - pytest -v -s kernels/quantization --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
-- label: Language Models Tests (Standard) # TBD
+- label: Kernels FP8 MoE Test (2xH100-2xMI355) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_1
- torch_nightly: true
+ agent_pool: mi355_2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/models/language
+ - csrc/moe/
+ - csrc/quantization/w8a8/cutlass/moe/
+ - vllm/model_executor/layers/fused_moe/
+ - tests/kernels/moe/test_deepep_moe.py
+ - vllm/_aiter_ops.py
+ - vllm/platforms/rocm.py
+ - vllm/envs.py
commands:
- - pip freeze | grep -E 'torch'
- - pytest -v -s models/language -m 'core_model and (not slow_test)'
+ - pytest -v -s kernels/moe/test_deepep_moe.py
+#---------------------------------------------------------- mi355 · models -----------------------------------------------------------#
- label: Language Models Test (Extended Generation) # TBD
timeout_in_minutes: 180
@@ -3310,15 +3046,15 @@ steps:
- vllm/
- tests/models/language/generation
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.5.2'
+ - 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)'
-
- label: Language Models Test (Extended Pooling) # TBD
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/
@@ -3326,72 +3062,46 @@ steps:
commands:
- pytest -v -s models/language/pooling -m 'not core_model'
-
-- label: "Multi-Modal Models (Standard) 1: qwen2" # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_1
- torch_nightly: true
- optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/
- - tests/models/multimodal
- commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen2"
- - pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model
-
-
-- label: "Multi-Modal Models (Standard) 2: qwen3 + gemma" # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_1
- torch_nightly: true
- optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/
- - tests/models/multimodal
- commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen3 or gemma"
- - pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model
-
-
-- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl" # TBD
+- label: Language Models Test (PPL) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
- torch_nightly: true
- optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/
- - tests/models/multimodal/generation
- - tests/models/multimodal/test_mapping.py
+ - vllm/model_executor/models/qwen3_5.py
+ - vllm/model_executor/models/qwen3_5_mtp.py
+ - vllm/transformers_utils/configs/qwen3_5.py
+ - vllm/transformers_utils/configs/qwen3_5_moe.py
+ - vllm/model_executor/models/qwen.py
+ - vllm/model_executor/models/qwen2.py
+ - vllm/model_executor/models/qwen3.py
+ - vllm/model_executor/models/qwen3_next.py
+ - vllm/model_executor/models/qwen3_next_mtp.py
+ - vllm/model_executor/layers/fla/ops/
+ - vllm/_aiter_ops.py
+ - vllm/v1/attention/backends/triton_attn.py
+ - vllm/v1/attention/backends/rocm_attn.py
+ - vllm/v1/attention/backends/rocm_aiter_unified_attn.py
+ - vllm/v1/attention/backends/rocm_aiter_fa.py
+ - vllm/v1/attention/backends/flex_attention.py
+ - vllm/v1/attention/ops/
+ - vllm/platforms/rocm.py
+ - tests/models/language/generation_ppl_test
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "not qwen2 and not qwen3 and not gemma"
- - pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model
-
+ - pytest -v -s models/language/generation_ppl_test
-- label: "Multi-Modal Models (Standard) 4: other + whisper" # TBD
+- label: Language Models Tests (Standard) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
torch_nightly: true
- optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/models/multimodal/generation
+ - tests/models/language
commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/processing
- - pytest -v -s 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
-
+ - pip freeze | grep -E 'torch'
+ - pytest -v -s models/language -m 'core_model and (not slow_test)'
- label: Multi-Modal Models (Extended Generation 1) # TBD
timeout_in_minutes: 180
@@ -3408,21 +3118,6 @@ steps:
- 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
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_1
- optional: true
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/
- - tests/models/multimodal/generation
- commands:
- - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- - pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=0) and not core_model'
-
-
- label: Multi-Modal Models (Extended Generation 3) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
@@ -3436,7 +3131,6 @@ steps:
- pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- pytest -v -s models/multimodal/generation/test_common.py -m 'split(group=1) and not core_model'
-
- label: Multi-Modal Models (Extended Pooling) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
@@ -3449,303 +3143,254 @@ steps:
commands:
- pytest -v -s models/multimodal/pooling -m 'not core_model'
-
-- label: Quantized Models Test # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_1
- working_dir: "/vllm-workspace/tests"
- source_file_dependencies:
- - vllm/model_executor/layers/quantization
- - tests/models/quantization
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
- - vllm/model_executor/model_loader/
- commands:
- - pytest -v -s models/quantization
-
-
-- label: Kernels (B200-MI355) # TBD
+- label: "Multi-Modal Models (Standard) 1: qwen2" # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
- working_dir: "/vllm-workspace/"
- source_file_dependencies:
- - csrc/quantization/fp4/
- - csrc/attention/mla/
- - csrc/quantization/cutlass_w8a8/moe/
- - vllm/model_executor/layers/fused_moe/cutlass_moe.py
- - vllm/v1/attention/backends/triton_attn.py
- - vllm/v1/attention/backends/rocm_attn.py
- - vllm/v1/attention/backends/rocm_aiter_fa.py
- - vllm/v1/attention/backends/rocm_aiter_unified_attn.py
- - vllm/v1/attention/backends/mla/aiter_triton_mla.py
- - vllm/v1/attention/backends/mla/rocm_aiter_mla.py
- - vllm/v1/attention/selector.py
- - vllm/platforms/rocm.py
- - vllm/_aiter_ops.py
- commands:
- - rocm-smi
- - python3 examples/basic/offline_inference/chat.py
- - pytest -v -s tests/kernels/attention/test_attention_selector.py
-
-
-- label: Weight Loading Multiple GPU # TBD
- timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_2
- num_gpus: 2
+ torch_nightly: true
+ optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/weight_loading
+ - tests/models/multimodal
commands:
- - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-amd.txt
-
+ - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
+ - pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen2"
+ - pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model
-- label: Weight Loading Multiple GPU - Large Models # TBD
+- label: "Multi-Modal Models (Standard) 4: other + whisper" # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_2
- working_dir: "/vllm-workspace/tests"
- num_gpus: 2
+ agent_pool: mi355_1
+ torch_nightly: true
optional: true
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- - tests/weight_loading
+ - tests/models/multimodal/generation
commands:
- - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-large-amd.txt
-
+ - pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
+ - pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/processing
+ - pytest -v -s 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
-- label: Ray Dependency Compatibility Check # TBD
+- label: Quantized Models Test # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
- optional: true
- working_dir: "/"
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - requirements/
- - setup.py
+ - vllm/model_executor/layers/quantization
+ - tests/models/quantization
+ - vllm/_aiter_ops.py
- vllm/platforms/rocm.py
+ - vllm/model_executor/model_loader/
commands:
- - bash /vllm-workspace/.buildkite/scripts/check-ray-compatibility.sh
+ - pytest -v -s models/quantization
+#------------------------------------------------------- mi355 · quantization --------------------------------------------------------#
-- label: Distributed NixlConnector PD accuracy (4 GPUs) # TBD
+- label: Quantization # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_4
- num_gpus: 4
- optional: true
+ agent_pool: mi355_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
- - tests/v1/kv_connector/nixl_integration/
+ - csrc/
+ - vllm/model_executor/layers/quantization
+ - tests/quantization
+ - vllm/_aiter_ops.py
- vllm/platforms/rocm.py
commands:
- - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
- - ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
+ - uv pip install --system torchao==0.17.0
+ - uv pip install --system conch-triton-kernels
+ - VLLM_TEST_FORCE_LOAD_FORMAT=auto pytest -v -s quantization/ --ignore quantization/test_blackwell_moe.py
+
+# - label: Quantized MoE Test (B200-MI355) # TBD
+# timeout_in_minutes: 180
+# mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
+# agent_pool: mi355_1
+# working_dir: "/vllm-workspace/"
+# source_file_dependencies:
+# - tests/quantization/test_gfx950_moe.py
+# - vllm/model_executor/models/deepseek_v2.py
+# - vllm/model_executor/models/gpt_oss.py
+# - vllm/model_executor/models/llama4.py
+# - vllm/model_executor/layers/fused_moe
+# - vllm/model_executor/layers/quantization/compressed_tensors
+# - vllm/model_executor/layers/quantization/modelopt.py
+# - vllm/model_executor/layers/quantization/mxfp4.py
+# - vllm/v1/attention/backends/triton_attn.py
+# - vllm/v1/attention/backends/rocm_attn.py
+# - vllm/v1/attention/backends/rocm_aiter_fa.py
+# - vllm/v1/attention/backends/mla/
+# - vllm/v1/attention/selector.py
+# - vllm/model_executor/layers/layernorm.py
+# - vllm/_aiter_ops.py
+# - vllm/platforms/rocm.py
+# - vllm/model_executor/model_loader/
+# commands:
+# - pytest -s -v tests/quantization/test_gfx950_moe.py
+#------------------------------------------------------------ mi355 · v1 -------------------------------------------------------------#
-- label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs) # TBD
+- label: V1 attention (B200-MI355) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_4
- num_gpus: 4
- optional: true
+ agent_pool: mi355_1
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
- - tests/v1/kv_connector/nixl_integration/
+ - vllm/config/attention.py
+ - vllm/model_executor/layers/attention
+ - vllm/v1/attention
+ - tests/v1/attention
+ - vllm/_aiter_ops.py
+ - vllm/envs.py
- vllm/platforms/rocm.py
commands:
- - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
- - DP_EP=1 ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
-
+ - pytest -v -s v1/attention
-- label: NixlConnector PD + Spec Decode acceptance (2 GPUs) # TBD
- timeout_in_minutes: 180
+- label: V1 Core + KV + Metrics # TBD
+ timeout_in_minutes: 60
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_2
- num_gpus: 2
+ agent_pool: mi355_1
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - 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/
- - vllm/platforms/rocm.py
+ - 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
- - ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
-
+ - 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: Distributed Tests (2 GPUs)(H100-MI355) # TBD
- timeout_in_minutes: 180
+- label: V1 Sample + Logits # TBD
+ timeout_in_minutes: 60
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_2
- num_gpus: 2
+ agent_pool: mi355_1
optional: true
- working_dir: "/vllm-workspace/"
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/distributed/
- - vllm/v1/distributed/
- - vllm/model_executor/layers/fused_moe/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - tests/distributed/test_context_parallel.py
- - tests/v1/distributed/test_dbo.py
- - examples/offline_inference/data_parallel.py
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/v1/sample
+ - tests/v1/logits_processors
+ - tests/v1/test_oracle.py
+ - tests/v1/test_request.py
+ - tests/v1/test_outputs.py
commands:
- - export TORCH_NCCL_BLOCKING_WAIT=1
- - pytest -v -s tests/distributed/test_context_parallel.py
- - 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
-
+ - pytest -v -s v1/sample
+ - 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
-- label: Distributed Compile Unit Tests (2xH100-2xMI355) # TBD
- timeout_in_minutes: 180
+- label: V1 Spec Decode # TBD
+ timeout_in_minutes: 60
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_2
- num_gpus: 2
- optional: true
- working_dir: "/vllm-workspace/"
+ agent_pool: mi355_1
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/compilation/
- - vllm/model_executor/layers
- - tests/compile/passes/distributed/
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/v1/spec_decode
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/passes/distributed/test_sequence_parallelism.py
- # TODO: this test is not supported on ROCm, there are aiter kernels for this.
- # - pytest -v -s tests/compile/passes/distributed/test_fusion_all_reduce.py
- # - pytest -v -s tests/compile/distributed/test_fusions_e2e.py::test_tp2_attn_quant_allreduce_rmsnorm
- # - "VLLM_TEST_CLEAN_GPU_MEMORY=1 pytest -v -s tests/compile/distributed/test_fusions_e2e.py -k 'not Llama-4'"
-
+ - pytest -v -s -m 'not slow_test' v1/spec_decode
-- label: LM Eval Small Models (B200-MI355) # TBD
+- label: NixlConnector PD + Spec Decode acceptance (2 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_2
+ num_gpus: 2
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - csrc/
- - vllm/model_executor/layers/quantization
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - vllm/_aiter_ops.py
+ - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
+ - vllm/v1/worker/kv_connector_model_runner_mixin.py
+ - tests/v1/kv_connector/nixl_integration/
- vllm/platforms/rocm.py
commands:
- - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-mi3xx-fp8-and-mixed.txt
-
+ - 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
-- label: LM Eval Large Models (4 GPUs)(FP8) # TBD
+- label: Distributed NixlConnector PD accuracy (4 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_4
num_gpus: 4
optional: true
- working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - csrc/
- - vllm/model_executor/layers/quantization
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - vllm/_aiter_ops.py
+ - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
+ - tests/v1/kv_connector/nixl_integration/
- vllm/platforms/rocm.py
commands:
- - export VLLM_USE_DEEP_GEMM=0
- - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-rocm-fp8.txt --tp-size=4
-
+ - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
+ - ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
-- label: GPQA Eval (GPT-OSS) (B200-MI355) # TBD
+- label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs) # TBD
timeout_in_minutes: 180
- mirror_hardwares: [amdexperimental, amdproduction, amdgfx955nightly, amdmi355]
- agent_pool: mi355_2
- num_gpus: 2
+ mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
+ agent_pool: mi355_4
+ num_gpus: 4
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - csrc/
- - vllm/model_executor/layers/quantization
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - vllm/model_executor/layers/fused_moe/
- - tests/evals/gpt_oss/
- - vllm/_aiter_ops.py
+ - vllm/distributed/kv_transfer/kv_connector/v1/nixl_connector.py
+ - tests/v1/kv_connector/nixl_integration/
- vllm/platforms/rocm.py
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-gfx950.txt
+ - uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
+ - DP_EP=1 ROCM_ATTN=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
+#------------------------------------------------------ mi355 · weight_loading -------------------------------------------------------#
-- label: Qwen3-30B-A3B-FP8-block Accuracy (B200-MI355) # TBD
+- label: Weight Loading Multiple GPU # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_2
num_gpus: 2
- working_dir: "/vllm-workspace"
+ working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/model_executor/models/
- - vllm/model_executor/model_loader/
- - vllm/model_executor/layers/quantization/
- - vllm/model_executor/layers/fused_moe/
- - vllm/distributed/eplb
- - vllm/v1/attention/backends/
- - vllm/v1/attention/selector.py
- - .buildkite/scripts/scheduled_integration_test/
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/weight_loading
commands:
- - bash .buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_block_ep_eplb.sh 0.8 200 8020 2 1
-
+ - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-amd.txt
-- label: Attention Benchmarks Smoke Test (B200-MI355) # TBD
+- label: Weight Loading Multiple GPU - Large Models # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_2
+ working_dir: "/vllm-workspace/tests"
num_gpus: 2
- working_dir: "/vllm-workspace/"
+ optional: true
source_file_dependencies:
- - benchmarks/attention_benchmarks/
- - vllm/v1/attention/
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/weight_loading
commands:
- - python3 benchmarks/attention_benchmarks/benchmark.py --backends ROCM_ATTN ROCM_AITER_FA ROCM_AITER_UNIFIED_ATTN --batch-specs "8q1s1k" --repeats 1 --warmup-iters 1
+ - bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-large-amd.txt
+#----------------------------------------------------------- mi355 · misc ------------------------------------------------------------#
-- label: LM Eval Qwen3-5 Models (B200-MI355) # TBD
- timeout_in_minutes: 120
+- label: Regression # TBD
+ timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
- agent_pool: mi355_2
- num_gpus: 2
+ agent_pool: mi355_1
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- - vllm/model_executor/models/qwen3_5.py
- - vllm/model_executor/models/qwen3_5_mtp.py
- - vllm/transformers_utils/configs/qwen3_5.py
- - vllm/transformers_utils/configs/qwen3_5_moe.py
- - vllm/model_executor/models/qwen.py
- - vllm/model_executor/models/qwen2.py
- - vllm/model_executor/models/qwen3.py
- - vllm/model_executor/models/qwen3_next.py
- - vllm/model_executor/models/qwen3_next_mtp.py
- - vllm/model_executor/layers/fla/ops/
- - vllm/_aiter_ops.py
- - vllm/platforms/rocm.py
+ - vllm/
+ - tests/test_regression
commands:
- - pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-qwen35-mi355.txt
+ - pip install modelscope
+ - pytest -v -s test_regression.py
diff --git a/.buildkite/test_areas/attention.yaml b/.buildkite/test_areas/attention.yaml
index 4bcf116f2756..d3947a03162b 100644
--- a/.buildkite/test_areas/attention.yaml
+++ b/.buildkite/test_areas/attention.yaml
@@ -3,6 +3,7 @@ depends_on:
- image-build
steps:
- label: V1 attention (H100)
+ key: v1-attention-h100
timeout_in_minutes: 30
device: h100
source_file_dependencies:
@@ -14,8 +15,9 @@ steps:
- pytest -v -s v1/attention
- label: V1 attention (B200)
+ key: v1-attention-b200
timeout_in_minutes: 30
- device: b200
+ device: b200-k8s
source_file_dependencies:
- vllm/config/attention.py
- vllm/model_executor/layers/attention
diff --git a/.buildkite/test_areas/basic_correctness.yaml b/.buildkite/test_areas/basic_correctness.yaml
index 042734e8433b..5d547cd48637 100644
--- a/.buildkite/test_areas/basic_correctness.yaml
+++ b/.buildkite/test_areas/basic_correctness.yaml
@@ -3,6 +3,7 @@ depends_on:
- image-build
steps:
- label: Basic Correctness
+ key: basic-correctness
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
diff --git a/.buildkite/test_areas/benchmarks.yaml b/.buildkite/test_areas/benchmarks.yaml
index 4cda6fff1443..85f804780179 100644
--- a/.buildkite/test_areas/benchmarks.yaml
+++ b/.buildkite/test_areas/benchmarks.yaml
@@ -3,6 +3,7 @@ depends_on:
- image-build
steps:
- label: Benchmarks CLI Test
+ key: benchmarks-cli-test
timeout_in_minutes: 20
device: h200_18gb
source_file_dependencies:
@@ -12,7 +13,8 @@ steps:
- pytest -v -s benchmarks/
- label: Attention Benchmarks Smoke Test (B200)
- device: b200
+ key: attention-benchmarks-smoke-test-b200
+ device: b200-k8s
num_gpus: 2
optional: true
working_dir: "/vllm-workspace/"
diff --git a/.buildkite/test_areas/compile.yaml b/.buildkite/test_areas/compile.yaml
index aa46447c24af..01248738d519 100644
--- a/.buildkite/test_areas/compile.yaml
+++ b/.buildkite/test_areas/compile.yaml
@@ -3,6 +3,7 @@ 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
@@ -17,6 +18,7 @@ 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
@@ -27,6 +29,7 @@ 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
@@ -37,9 +40,10 @@ 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
+ device: b200-k8s
optional: true
num_devices: 2
commands:
@@ -47,6 +51,7 @@ 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
@@ -60,9 +65,10 @@ 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
+ device: b200-k8s
source_file_dependencies:
- csrc/quantization/fp4/
- vllm/model_executor/layers/quantization/
@@ -89,6 +95,7 @@ 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
@@ -107,6 +114,7 @@ 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
@@ -126,9 +134,10 @@ 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
+ device: b200-k8s
num_devices: 1
optional: true
commands:
@@ -139,6 +148,7 @@ 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
@@ -156,6 +166,7 @@ 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
@@ -175,6 +186,7 @@ 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
@@ -194,9 +206,10 @@ 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
+ device: b200-k8s
num_devices: 2
source_file_dependencies:
- csrc/quantization/
diff --git a/.buildkite/test_areas/cuda.yaml b/.buildkite/test_areas/cuda.yaml
index 4d1efdb13c88..6254a6ba3dd5 100644
--- a/.buildkite/test_areas/cuda.yaml
+++ b/.buildkite/test_areas/cuda.yaml
@@ -3,16 +3,22 @@ depends_on:
- image-build
steps:
- label: Platform Tests (CUDA)
+ key: platform-tests-cuda
timeout_in_minutes: 15
device: h200_18gb
source_file_dependencies:
- - vllm/
+ - vllm/envs.py
+ - vllm/logger.py
+ - vllm/platforms/
+ - vllm/plugins/
+ - vllm/utils/
- 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
diff --git a/.buildkite/test_areas/disaggregated.yaml b/.buildkite/test_areas/disaggregated.yaml
new file mode 100644
index 000000000000..e68b9e1add8b
--- /dev/null
+++ b/.buildkite/test_areas/disaggregated.yaml
@@ -0,0 +1,106 @@
+group: Disaggregated
+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/
+ - 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/
+ - 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/
+ - 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/
+ - 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: 20
+ working_dir: "/vllm-workspace/tests"
+ num_devices: 4
+ source_file_dependencies:
+ - vllm/distributed/kv_transfer/kv_connector/v1/nixl/
+ - tests/v1/kv_connector/nixl_integration/
+ 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/multi_connector.py
+ - vllm/distributed/kv_transfer/kv_connector/v1/offloading_connector.py
+ - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
+ - tests/v1/kv_connector/nixl_integration/
+ commands:
+ - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
+ - bash v1/kv_connector/nixl_integration/run_multi_connector_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/v1/worker/kv_connector_model_runner_mixin.py
+ - tests/v1/kv_connector/nixl_integration/
+ commands:
+ - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
+ - bash 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/multi_connector.py
+ - vllm/distributed/kv_transfer/kv_connector/v1/offloading_connector.py
+ - vllm/distributed/kv_transfer/kv_connector/v1/offloading/
+ - tests/v1/kv_connector/nixl_integration/
+ commands:
+ - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
+ - bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh
\ No newline at end of file
diff --git a/.buildkite/test_areas/distributed.yaml b/.buildkite/test_areas/distributed.yaml
index 1c9ba95183c1..8aa41a9a26ab 100644
--- a/.buildkite/test_areas/distributed.yaml
+++ b/.buildkite/test_areas/distributed.yaml
@@ -3,6 +3,7 @@ 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
@@ -16,6 +17,7 @@ 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
@@ -37,6 +39,7 @@ 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
@@ -59,6 +62,7 @@ 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
@@ -81,6 +85,7 @@ 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
@@ -88,9 +93,8 @@ 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/offline_inference/data_parallel.py
+ - tests/examples/features/data_parallel/data_parallel_offline.py
commands:
# https://github.com/NVIDIA/nccl/issues/1838
- export NCCL_CUMEM_HOST_ENABLE=0
@@ -107,12 +111,13 @@ 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/offline_inference/data_parallel.py --enforce-eager
+ - python3 examples/features/data_parallel/data_parallel_offline.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
@@ -133,6 +138,7 @@ 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
@@ -154,24 +160,28 @@ 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/offline_inference/torchrun_dp_example.py
+ - examples/features/torchrun/torchrun_dp_example_offline.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/offline_inference/torchrun_dp_example.py --tp-size=2 --pp-size=1 --dp-size=4 --enable-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
- label: Distributed Tests (4 GPUs)(A100)
+ key: distributed-tests-4-gpus-a100
device: a100
optional: true
num_devices: 4
@@ -186,6 +196,7 @@ 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
@@ -194,12 +205,14 @@ 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/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
+ - 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
- pytest -v -s tests/v1/distributed/test_dbo.py
- - TP_SIZE=1 DP_SIZE=2 pytest -v -s tests/v1/distributed/test_eagle_dp.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)
- device: b200
+ key: distributed-tests-2-gpus-b200
+ device: b200-k8s
optional: true
working_dir: "/vllm-workspace/"
num_devices: 2
@@ -207,8 +220,12 @@ 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
@@ -221,96 +238,12 @@ steps:
- vllm/executor/
- vllm/model_executor/models/
- tests/distributed/
- - 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/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: 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_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: 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_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)
- timeout_in_minutes: 30
- working_dir: "/vllm-workspace/tests"
- num_devices: 4
- source_file_dependencies:
- - 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: Hyrbid SSM NixlConnector PD accuracy tests (4 GPUs)
- timeout_in_minutes: 20
- working_dir: "/vllm-workspace/tests"
- num_devices: 4
- source_file_dependencies:
- - 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)
- timeout_in_minutes: 30
- working_dir: "/vllm-workspace/tests"
- num_devices: 2
- source_file_dependencies:
- - 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/
- - tests/v1/kv_connector/nixl_integration/
- commands:
- - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
- - bash v1/kv_connector/nixl_integration/run_multi_connector_accuracy_test.sh
-
-- label: 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_connector.py
- - vllm/v1/worker/kv_connector_model_runner_mixin.py
- - tests/v1/kv_connector/nixl_integration/
- commands:
- - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
- - bash v1/kv_connector/nixl_integration/spec_decode_acceptance_test.sh
-
-- label: 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_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/
- - tests/v1/kv_connector/nixl_integration/
+ - tests/examples/features/data_parallel/data_parallel_offline.py
commands:
- - uv pip install --system -r /vllm-workspace/requirements/kv_connectors.txt
- - bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh
+ - ./.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"
- label: Pipeline + Context Parallelism (4 GPUs)
+ key: pipeline-context-parallelism-4-gpus
timeout_in_minutes: 60
working_dir: "/vllm-workspace/tests"
num_devices: 4
@@ -325,6 +258,7 @@ 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
diff --git a/.buildkite/test_areas/docker.yaml b/.buildkite/test_areas/docker.yaml
new file mode 100644
index 000000000000..9bf96221abe0
--- /dev/null
+++ b/.buildkite/test_areas/docker.yaml
@@ -0,0 +1,16 @@
+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
diff --git a/.buildkite/test_areas/e2e_integration.yaml b/.buildkite/test_areas/e2e_integration.yaml
index 5b7f96bc7a26..bb8aa14eac18 100644
--- a/.buildkite/test_areas/e2e_integration.yaml
+++ b/.buildkite/test_areas/e2e_integration.yaml
@@ -3,6 +3,7 @@ depends_on:
- image-build
steps:
- label: DeepSeek V2-Lite Accuracy
+ key: deepseek-v2-lite-accuracy
timeout_in_minutes: 60
device: h100
optional: true
@@ -12,6 +13,7 @@ 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
@@ -21,15 +23,27 @@ 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
+ device: b200-k8s
optional: true
num_devices: 2
working_dir: "/vllm-workspace"
commands:
- 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
+ num_devices: 4
+ working_dir: "/vllm-workspace"
+ commands:
+ - 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
diff --git a/.buildkite/test_areas/engine.yaml b/.buildkite/test_areas/engine.yaml
index 5e4361ec9ad6..c8b2d66a413a 100644
--- a/.buildkite/test_areas/engine.yaml
+++ b/.buildkite/test_areas/engine.yaml
@@ -3,19 +3,32 @@ depends_on:
- image-build
steps:
- label: Engine
+ key: engine
timeout_in_minutes: 15
device: h200_18gb
source_file_dependencies:
- - vllm/
+ - 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/
- 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
+ - pytest -v -s engine test_sequence.py test_config.py test_logger.py test_vllm_port.py test_jit_monitor.py
- label: Engine (1 GPU)
+ key: engine-1-gpu
timeout_in_minutes: 30
source_file_dependencies:
- vllm/v1/engine/
@@ -25,6 +38,7 @@ steps:
- pytest -v -s v1/engine --ignore v1/engine/test_preprocess_error_handling.py
- label: e2e Scheduling (1 GPU)
+ key: e2e-scheduling-1-gpu
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
@@ -34,6 +48,7 @@ steps:
- pytest -v -s v1/e2e/general/test_async_scheduling.py
- label: e2e Core (1 GPU)
+ key: e2e-core-1-gpu
timeout_in_minutes: 30
source_file_dependencies:
- vllm/v1/
@@ -42,38 +57,63 @@ 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/
- - tests/v1/e2e
+ - 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
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/
- - tests/v1/e2e
+ - 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
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
diff --git a/.buildkite/test_areas/entrypoints.yaml b/.buildkite/test_areas/entrypoints.yaml
index 8c2b529a8068..ba92d3a3aec0 100644
--- a/.buildkite/test_areas/entrypoints.yaml
+++ b/.buildkite/test_areas/entrypoints.yaml
@@ -2,7 +2,8 @@ group: Entrypoints
depends_on:
- image-build
steps:
-- label: Entrypoints Unit Tests
+- label: Entrypoints Unit Tests
+ key: entrypoints-unit-tests
timeout_in_minutes: 10
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -13,6 +14,7 @@ steps:
- 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:
@@ -24,8 +26,14 @@ 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:
@@ -35,14 +43,10 @@ steps:
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/openai/chat_completion --ignore=entrypoints/openai/chat_completion/test_chat_with_tool_reasoning.py --ignore=entrypoints/openai/chat_completion/test_oot_registration.py
- mirror:
- amd:
- 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:
@@ -53,13 +57,9 @@ steps:
- 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: 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"
@@ -72,6 +72,7 @@ steps:
- 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:
@@ -86,6 +87,7 @@ steps:
- pytest -v -s tool_use
- label: Entrypoints Integration (Pooling)
+ key: entrypoints-integration-pooling
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -96,6 +98,7 @@ 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:
@@ -105,6 +108,7 @@ 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:
diff --git a/.buildkite/test_areas/expert_parallelism.yaml b/.buildkite/test_areas/expert_parallelism.yaml
index c2adf52a2d57..0f7ab0d7157c 100644
--- a/.buildkite/test_areas/expert_parallelism.yaml
+++ b/.buildkite/test_areas/expert_parallelism.yaml
@@ -3,6 +3,7 @@ depends_on:
- image-build
steps:
- label: EPLB Algorithm
+ key: eplb-algorithm
timeout_in_minutes: 15
device: h200_18gb
working_dir: "/vllm-workspace/tests"
@@ -15,6 +16,7 @@ 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
@@ -26,6 +28,7 @@ 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"
diff --git a/.buildkite/test_areas/kernels.yaml b/.buildkite/test_areas/kernels.yaml
index 8b9765130aee..34e1e4832d9d 100644
--- a/.buildkite/test_areas/kernels.yaml
+++ b/.buildkite/test_areas/kernels.yaml
@@ -3,6 +3,7 @@ depends_on:
- image-build
steps:
- label: vLLM IR Tests
+ key: vllm-ir-tests
timeout_in_minutes: 10
device: h200_18gb
working_dir: "/vllm-workspace/"
@@ -14,6 +15,7 @@ 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/
@@ -23,6 +25,7 @@ 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
@@ -36,6 +39,7 @@ 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/
@@ -49,6 +53,7 @@ steps:
parallelism: 2
- label: Kernels Quantization Test %N
+ key: kernels-quantization-test
timeout_in_minutes: 90
source_file_dependencies:
- csrc/quantization/
@@ -59,6 +64,7 @@ steps:
parallelism: 2
- label: Kernels MoE Test %N
+ key: kernels-moe-test
timeout_in_minutes: 25
source_file_dependencies:
- csrc/quantization/cutlass_w8a8/moe/
@@ -74,6 +80,7 @@ steps:
parallelism: 5
- label: Kernels Mamba Test
+ key: kernels-mamba-test
timeout_in_minutes: 45
source_file_dependencies:
- csrc/mamba/
@@ -82,7 +89,18 @@ steps:
commands:
- pytest -v -s kernels/mamba
+- label: Kernels KDA Test
+ timeout_in_minutes: 20
+ 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
@@ -95,16 +113,19 @@ 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:
- 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
+ device: b200-k8s
# optional: true
source_file_dependencies:
- csrc/quantization/fp4/
@@ -141,6 +162,7 @@ steps:
- pytest -v -s tests/kernels/quantization/test_nvfp4_qutlass.py
- pytest -v -s tests/kernels/quantization/test_mxfp4_qutlass.py
- pytest -v -s tests/kernels/moe/test_nvfp4_moe.py
+ - pytest -v -s tests/kernels/moe/test_mxfp4_moe.py
- pytest -v -s tests/kernels/moe/test_ocp_mx_moe.py
- pytest -v -s tests/kernels/moe/test_flashinfer.py
- pytest -v -s tests/kernels/moe/test_flashinfer_moe.py
@@ -149,17 +171,19 @@ 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:
- vllm/utils/import_utils.py
- tests/kernels/helion/
commands:
- - pip install helion==0.3.3
+ - pip install helion==1.0.0
- pytest -v -s kernels/helion/
- label: Kernels FP8 MoE Test (1 H100)
+ key: kernels-fp8-moe-test-1-h100
timeout_in_minutes: 90
device: h100
num_devices: 1
@@ -176,6 +200,7 @@ 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
@@ -185,8 +210,9 @@ 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
+ device: b200-k8s
num_devices: 1
optional: true
commands:
@@ -197,6 +223,7 @@ steps:
- label: Kernels FusedMoE Layer Test (2 H100s)
+ key: kernels-fusedmoe-layer-test-2-h100s
timeout_in_minutes: 90
device: h100
num_devices: 2
@@ -213,6 +240,7 @@ steps:
- label: Kernels FusedMoE Layer Test (2 B200s)
+ key: kernels-fusedmoe-layer-test-2-b200s
timeout_in_minutes: 90
device: b200
num_devices: 2
diff --git a/.buildkite/test_areas/lm_eval.yaml b/.buildkite/test_areas/lm_eval.yaml
index 39029efe9cd9..e5a163d17c7e 100644
--- a/.buildkite/test_areas/lm_eval.yaml
+++ b/.buildkite/test_areas/lm_eval.yaml
@@ -3,6 +3,7 @@ depends_on:
- image-build
steps:
- label: LM Eval Small Models
+ key: lm-eval-small-models
timeout_in_minutes: 75
source_file_dependencies:
- csrc/
@@ -24,6 +25,7 @@ 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
@@ -36,6 +38,7 @@ 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
optional: true
@@ -46,8 +49,9 @@ steps:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-blackwell.txt
- label: LM Eval Qwen3.5 Models (B200)
+ key: lm-eval-qwen3-5-models-b200
timeout_in_minutes: 120
- device: b200
+ device: b200-k8s
optional: true
num_devices: 2
source_file_dependencies:
@@ -62,6 +66,7 @@ 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
@@ -70,6 +75,7 @@ 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
@@ -77,13 +83,15 @@ 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)
- device: b200
+ key: moe-refactor-integration-test-b200-temporary
+ device: b200-k8s
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
optional: true
num_devices: 2
@@ -91,7 +99,19 @@ steps:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/moe-refactor-dp-ep/config-b200.txt
+- label: LM Eval TurboQuant KV Cache
+ key: lm-eval-turboquant-kv-cache
+ timeout_in_minutes: 75
+ source_file_dependencies:
+ - vllm/model_executor/layers/quantization/turboquant/
+ - vllm/v1/attention/backends/turboquant_attn.py
+ - vllm/v1/attention/ops/triton_turboquant_decode.py
+ - vllm/v1/attention/ops/triton_turboquant_store.py
+ commands:
+ - 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
@@ -105,6 +125,7 @@ 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
optional: true
@@ -116,3 +137,10 @@ 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
diff --git a/.buildkite/test_areas/lora.yaml b/.buildkite/test_areas/lora.yaml
index 21f392ff737b..ebffcc59fc3b 100644
--- a/.buildkite/test_areas/lora.yaml
+++ b/.buildkite/test_areas/lora.yaml
@@ -3,6 +3,7 @@ depends_on:
- image-build
steps:
- label: LoRA %N
+ key: lora
timeout_in_minutes: 30
source_file_dependencies:
- vllm/lora
@@ -13,10 +14,12 @@ steps:
- 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
diff --git a/.buildkite/test_areas/misc.yaml b/.buildkite/test_areas/misc.yaml
index 0cf9ec433925..da04c18017db 100644
--- a/.buildkite/test_areas/misc.yaml
+++ b/.buildkite/test_areas/misc.yaml
@@ -3,25 +3,40 @@ depends_on:
- image-build
steps:
- label: V1 Spec Decode
+ key: v1-spec-decode
timeout_in_minutes: 30
source_file_dependencies:
- - vllm/
+ - vllm/config/
+ - vllm/distributed/
+ - vllm/inputs/
+ - vllm/model_executor/
+ - vllm/platforms/
+ - vllm/sampling_params.py
+ - vllm/transformers_utils/
+ - vllm/utils/
+ - vllm/v1/
- 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/
+ - vllm/config/
+ - vllm/distributed/
+ - vllm/engine/
+ - vllm/inputs/
+ - vllm/logger.py
+ - vllm/model_executor/
+ - vllm/platforms/
+ - vllm/sampling_params.py
+ - vllm/transformers_utils/
+ - vllm/utils/
+ - vllm/v1/
- tests/v1/sample
- tests/v1/logits_processors
- tests/v1/test_oracle.py
@@ -36,14 +51,31 @@ steps:
- pytest -v -s v1/test_outputs.py
mirror:
amd:
- device: mi325_1
+ device: mi300_1
depends_on:
- image-build-amd
- label: V1 Core + KV + Metrics
+ key: v1-core-kv-metrics
timeout_in_minutes: 30
source_file_dependencies:
- - vllm/
+ - 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/
- tests/v1/core
- tests/v1/executor
- tests/v1/kv_offload
@@ -64,17 +96,27 @@ steps:
# 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/
+ - 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/
- tests/v1
device: cpu-small
commands:
@@ -86,10 +128,21 @@ steps:
- pytest -v -s -m 'cpu_test' v1/metrics
- label: Regression
+ key: regression
timeout_in_minutes: 20
device: h200_18gb
source_file_dependencies:
- - vllm/
+ - vllm/config/
+ - vllm/distributed/
+ - vllm/engine/
+ - vllm/inputs/
+ - vllm/model_executor/
+ - vllm/multimodal/
+ - vllm/platforms/
+ - vllm/sampling_params.py
+ - vllm/transformers_utils/
+ - vllm/utils/
+ - vllm/v1/
- tests/test_regression
commands:
- pip install modelscope
@@ -97,6 +150,7 @@ steps:
working_dir: "/vllm-workspace/tests" # optional
- label: Examples
+ key: examples
timeout_in_minutes: 45
working_dir: "/vllm-workspace/examples"
source_file_dependencies:
@@ -113,25 +167,37 @@ steps:
- python3 basic/offline_inference/embed.py
- python3 basic/offline_inference/score.py
# for multi-modal models
- - 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
+ - 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
# for pooling models
- python3 pooling/embed/vision_embedding_offline.py --seed 0
# for features demo
- - 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 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
# 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 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
+ - 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
- label: Metrics, Tracing (2 GPUs)
+ key: metrics-tracing-2-gpus
timeout_in_minutes: 20
num_devices: 2
source_file_dependencies:
- - vllm/
+ - vllm/config/
+ - vllm/distributed/
+ - vllm/engine/
+ - vllm/inputs/
+ - vllm/model_executor/
+ - vllm/multimodal/
+ - vllm/platforms/
+ - vllm/sampling_params.py
+ - vllm/tracing/
+ - vllm/transformers_utils/
+ - vllm/utils/
+ - vllm/v1/
- tests/v1/tracing
commands:
- "pip install \
@@ -142,6 +208,7 @@ steps:
- pytest -v -s v1/tracing
- label: Python-only Installation
+ key: python-only-installation
depends_on: ~
timeout_in_minutes: 20
source_file_dependencies:
@@ -151,9 +218,22 @@ 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/
+ - vllm/assets/
+ - vllm/config/
+ - vllm/distributed/
+ - vllm/engine/
+ - vllm/inputs/
+ - vllm/model_executor/
+ - vllm/multimodal/
+ - vllm/platforms/
+ - vllm/sampling_params.py
+ - vllm/tokenizers/
+ - vllm/transformers_utils/
+ - vllm/utils/
+ - vllm/v1/
- tests/detokenizer
- tests/multimodal
- tests/utils_
@@ -163,11 +243,35 @@ steps:
- pytest -v -s utils_
- label: Async Engine, Inputs, Utils, Worker, Config (CPU)
- depends_on:
+ key: async-engine-inputs-utils-worker-config-cpu
+ depends_on:
- image-build-cpu
timeout_in_minutes: 30
source_file_dependencies:
- - vllm/
+ - 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/
- tests/test_inputs.py
- tests/test_outputs.py
- tests/test_pooling_params.py
@@ -190,12 +294,13 @@ 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 --ignore=reasoning/test_gemma4_reasoning_parser.py
+ - pytest -v -s reasoning --ignore=reasoning/test_seedoss_reasoning_parser.py --ignore=reasoning/test_glm4_moe_reasoning_parser.py
- pytest -v -s tool_parsers
- pytest -v -s transformers_utils
- pytest -v -s config
- label: Batch Invariance (H100)
+ key: batch-invariance-h100
timeout_in_minutes: 30
device: h100
source_file_dependencies:
@@ -211,8 +316,9 @@ 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
+ device: b200-k8s
source_file_dependencies:
- vllm/v1/attention
- vllm/model_executor/layers
@@ -227,6 +333,7 @@ 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
diff --git a/.buildkite/test_areas/model_executor.yaml b/.buildkite/test_areas/model_executor.yaml
index 212abfdbb906..c41ef8a7110d 100644
--- a/.buildkite/test_areas/model_executor.yaml
+++ b/.buildkite/test_areas/model_executor.yaml
@@ -3,6 +3,7 @@ depends_on:
- image-build
steps:
- label: Model Executor
+ key: model-executor
timeout_in_minutes: 35
source_file_dependencies:
- vllm/engine/arg_utils.py
diff --git a/.buildkite/test_areas/model_runner_v2.yaml b/.buildkite/test_areas/model_runner_v2.yaml
index 7aa1870f0db5..9dfd046289e8 100644
--- a/.buildkite/test_areas/model_runner_v2.yaml
+++ b/.buildkite/test_areas/model_runner_v2.yaml
@@ -3,6 +3,7 @@ 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/
@@ -25,16 +26,18 @@ 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/others/tensorize_vllm_model.py
+ - examples/features/tensorize_vllm_model.py
commands:
- set -x
- export VLLM_USE_V2_MODEL_RUNNER=1
@@ -44,21 +47,22 @@ 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 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
+ - 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
# for pooling models
- python3 pooling/embed/vision_embedding_offline.py --seed 0
# for features demo
- - 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 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
# 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 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
+ - 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
- 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
@@ -79,6 +83,7 @@ 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
@@ -94,6 +99,7 @@ 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:
diff --git a/.buildkite/test_areas/models_basic.yaml b/.buildkite/test_areas/models_basic.yaml
index 10b038d8b8a8..8fca203de44f 100644
--- a/.buildkite/test_areas/models_basic.yaml
+++ b/.buildkite/test_areas/models_basic.yaml
@@ -3,8 +3,8 @@ 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/
@@ -17,6 +17,7 @@ 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/
@@ -32,6 +33,7 @@ steps:
torch_nightly: {}
- label: Basic Models Tests (Other)
+ key: basic-models-tests-other
timeout_in_minutes: 45
source_file_dependencies:
- vllm/
@@ -40,14 +42,9 @@ 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
@@ -60,6 +57,7 @@ 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
@@ -70,6 +68,22 @@ 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/offline_inference/vision_language.py --model-type qwen2_5_vl
+ - python3 examples/generate/multimodal/vision_language_offline.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
+
+- label: Transformers Backward Compatibility Models Test
+ key: transformers-backward-compatibility-models-test
+ working_dir: "/vllm-workspace/"
+ optional: true
+ soft_fail: true
+ commands:
+ - pip install transformers==4.57.5
+ - pytest -v -s tests/models/test_initialization.py
+ - 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
# Whisper needs spawn method to avoid deadlock
- - VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/offline_inference/audio_language.py --model-type whisper
+ - VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/generate/multimodal/audio_language_offline.py --model-type whisper
diff --git a/.buildkite/test_areas/models_distributed.yaml b/.buildkite/test_areas/models_distributed.yaml
index 55e7410b8af4..b5758c55affa 100644
--- a/.buildkite/test_areas/models_distributed.yaml
+++ b/.buildkite/test_areas/models_distributed.yaml
@@ -3,6 +3,7 @@ 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
diff --git a/.buildkite/test_areas/models_language.yaml b/.buildkite/test_areas/models_language.yaml
index c13371e25f1d..b560c5a4769a 100644
--- a/.buildkite/test_areas/models_language.yaml
+++ b/.buildkite/test_areas/models_language.yaml
@@ -3,6 +3,7 @@ depends_on:
- image-build
steps:
- label: Language Models Tests (Standard)
+ key: language-models-tests-standard
timeout_in_minutes: 25
source_file_dependencies:
- vllm/
@@ -15,6 +16,7 @@ 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/
@@ -31,6 +33,7 @@ steps:
torch_nightly: {}
- label: Language Models Tests (Hybrid) %N
+ key: language-models-tests-hybrid
timeout_in_minutes: 75
source_file_dependencies:
- vllm/
@@ -45,8 +48,17 @@ 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:
@@ -58,17 +70,9 @@ 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
@@ -79,6 +83,7 @@ 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:
@@ -86,13 +91,9 @@ steps:
- tests/models/language/pooling
commands:
- pytest -v -s models/language/pooling -m 'not core_model'
- mirror:
- amd:
- 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
diff --git a/.buildkite/test_areas/models_multimodal.yaml b/.buildkite/test_areas/models_multimodal.yaml
index ff0fd2e7a626..1f66393df818 100644
--- a/.buildkite/test_areas/models_multimodal.yaml
+++ b/.buildkite/test_areas/models_multimodal.yaml
@@ -3,6 +3,7 @@ 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:
@@ -14,11 +15,12 @@ steps:
- pytest -v -s models/multimodal/generation/test_ultravox.py -m core_model
mirror:
amd:
- device: mi325_1
+ device: mi300_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:
@@ -28,13 +30,15 @@ steps:
- pip install git+https://github.com/TIGER-AI-Lab/Mantis.git
- pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen3 or gemma"
- pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model
+ - pytest -v -s models/multimodal/generation/test_vit_cudagraph.py -m core_model
mirror:
amd:
- device: mi325_1
+ device: mi300_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/
@@ -45,11 +49,12 @@ steps:
- pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model
mirror:
amd:
- device: mi325_1
+ device: mi300_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/
@@ -59,14 +64,10 @@ 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)
- depends_on:
+ key: multi-modal-processor-cpu
+ depends_on:
- image-build-cpu
timeout_in_minutes: 60
source_file_dependencies:
@@ -79,6 +80,7 @@ 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:
@@ -90,6 +92,7 @@ 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:
@@ -100,6 +103,7 @@ 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/
@@ -111,11 +115,12 @@ steps:
- pytest -v -s models/multimodal/test_mapping.py
mirror:
amd:
- device: mi325_1
+ device: mi300_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/
@@ -125,6 +130,7 @@ 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/
@@ -134,6 +140,7 @@ 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:
diff --git a/.buildkite/test_areas/plugins.yaml b/.buildkite/test_areas/plugins.yaml
index 8e0eb0284019..0d23180f3ef7 100644
--- a/.buildkite/test_areas/plugins.yaml
+++ b/.buildkite/test_areas/plugins.yaml
@@ -3,6 +3,7 @@ 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
diff --git a/.buildkite/test_areas/pytorch.yaml b/.buildkite/test_areas/pytorch.yaml
index a3648219d89d..c58a9986beba 100644
--- a/.buildkite/test_areas/pytorch.yaml
+++ b/.buildkite/test_areas/pytorch.yaml
@@ -3,9 +3,33 @@ depends_on:
- image-build
steps:
- label: PyTorch Compilation Unit Tests
+ key: pytorch-compilation-unit-tests
timeout_in_minutes: 10
source_file_dependencies:
- - vllm/
+ - vllm/__init__.py
+ - vllm/_aiter_ops.py
+ - vllm/_custom_ops.py
+ - vllm/compilation/
+ - vllm/config/
+ - vllm/distributed/
+ - vllm/engine/
+ - vllm/env_override.py
+ - vllm/envs.py
+ - vllm/forward_context.py
+ - vllm/inputs/
+ - vllm/ir/
+ - vllm/kernels/
+ - vllm/logger.py
+ - vllm/model_executor/
+ - vllm/multimodal/
+ - vllm/platforms/
+ - vllm/plugins/
+ - vllm/sampling_params.py
+ - vllm/sequence.py
+ - vllm/transformers_utils/
+ - vllm/triton_utils/
+ - vllm/utils/
+ - vllm/v1/
- tests/compile
commands:
# Run unit tests defined directly under compile/,
@@ -18,27 +42,99 @@ 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/
+ - vllm/__init__.py
+ - vllm/_aiter_ops.py
+ - vllm/_custom_ops.py
+ - vllm/compilation/
+ - vllm/config/
+ - vllm/distributed/
+ - vllm/engine/
+ - vllm/env_override.py
+ - vllm/envs.py
+ - vllm/forward_context.py
+ - vllm/inputs/
+ - vllm/ir/
+ - vllm/kernels/
+ - vllm/logger.py
+ - vllm/model_executor/
+ - vllm/multimodal/
+ - vllm/platforms/
+ - vllm/plugins/
+ - vllm/sampling_params.py
+ - vllm/sequence.py
+ - vllm/transformers_utils/
+ - vllm/triton_utils/
+ - vllm/utils/
+ - vllm/v1/
- tests/compile/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/
+ - vllm/__init__.py
+ - vllm/_aiter_ops.py
+ - vllm/_custom_ops.py
+ - vllm/compilation/
+ - vllm/config/
+ - vllm/distributed/
+ - vllm/engine/
+ - vllm/env_override.py
+ - vllm/envs.py
+ - vllm/forward_context.py
+ - vllm/inputs/
+ - vllm/ir/
+ - vllm/kernels/
+ - vllm/logger.py
+ - vllm/model_executor/
+ - vllm/multimodal/
+ - vllm/platforms/
+ - vllm/plugins/
+ - vllm/sampling_params.py
+ - vllm/sequence.py
+ - vllm/transformers_utils/
+ - vllm/triton_utils/
+ - vllm/utils/
+ - vllm/v1/
- tests/compile/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/
+ - vllm/__init__.py
+ - vllm/_aiter_ops.py
+ - vllm/_custom_ops.py
+ - vllm/compilation/
+ - vllm/config/
+ - vllm/distributed/
+ - vllm/engine/
+ - vllm/env_override.py
+ - vllm/envs.py
+ - vllm/forward_context.py
+ - vllm/inputs/
+ - vllm/ir/
+ - vllm/kernels/
+ - vllm/logger.py
+ - vllm/model_executor/
+ - vllm/multimodal/
+ - vllm/platforms/
+ - vllm/plugins/
+ - vllm/sampling_params.py
+ - vllm/sequence.py
+ - vllm/transformers_utils/
+ - vllm/triton_utils/
+ - vllm/utils/
+ - vllm/v1/
- tests/compile
commands:
# Run smoke tests under fullgraph directory, except test_full_graph.py
@@ -48,16 +144,41 @@ 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/
+ - vllm/__init__.py
+ - vllm/_aiter_ops.py
+ - vllm/_custom_ops.py
+ - vllm/compilation/
+ - vllm/config/
+ - vllm/distributed/
+ - vllm/engine/
+ - vllm/env_override.py
+ - vllm/envs.py
+ - vllm/forward_context.py
+ - vllm/inputs/
+ - vllm/ir/
+ - vllm/kernels/
+ - vllm/logger.py
+ - vllm/model_executor/
+ - vllm/multimodal/
+ - vllm/platforms/
+ - vllm/plugins/
+ - vllm/sampling_params.py
+ - vllm/sequence.py
+ - vllm/transformers_utils/
+ - vllm/triton_utils/
+ - vllm/utils/
+ - vllm/v1/
- tests/compile
commands:
# 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
diff --git a/.buildkite/test_areas/quantization.yaml b/.buildkite/test_areas/quantization.yaml
index a42d59b021c6..8a9a36da4481 100644
--- a/.buildkite/test_areas/quantization.yaml
+++ b/.buildkite/test_areas/quantization.yaml
@@ -3,6 +3,7 @@ depends_on:
- image-build
steps:
- label: Quantization
+ key: quantization
timeout_in_minutes: 90
source_file_dependencies:
- csrc/
@@ -21,9 +22,10 @@ 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
+ device: b200-k8s
source_file_dependencies:
- tests/quantization/test_blackwell_moe.py
- vllm/model_executor/models/deepseek_v2.py
@@ -38,6 +40,7 @@ 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
diff --git a/.buildkite/test_areas/ray_compat.yaml b/.buildkite/test_areas/ray_compat.yaml
index 3485e346532c..9207621a5830 100644
--- a/.buildkite/test_areas/ray_compat.yaml
+++ b/.buildkite/test_areas/ray_compat.yaml
@@ -3,6 +3,7 @@ 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.
diff --git a/.buildkite/test_areas/samplers.yaml b/.buildkite/test_areas/samplers.yaml
index 2052a379827a..48e9f55571e4 100644
--- a/.buildkite/test_areas/samplers.yaml
+++ b/.buildkite/test_areas/samplers.yaml
@@ -3,6 +3,7 @@ depends_on:
- image-build
steps:
- label: Samplers Test
+ key: samplers-test
timeout_in_minutes: 75
source_file_dependencies:
- vllm/model_executor/layers
@@ -10,11 +11,13 @@ steps:
- tests/samplers
- tests/conftest.py
commands:
- - pytest -v -s samplers
+ # 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
- VLLM_USE_FLASHINFER_SAMPLER=1 pytest -v -s samplers
mirror:
amd:
- device: mi325_1
+ device: mi250_1
depends_on:
- image-build-amd
commands:
diff --git a/.buildkite/test_areas/spec_decode.yaml b/.buildkite/test_areas/spec_decode.yaml
index a0b730968675..5253f54735aa 100644
--- a/.buildkite/test_areas/spec_decode.yaml
+++ b/.buildkite/test_areas/spec_decode.yaml
@@ -3,6 +3,7 @@ depends_on:
- image-build
steps:
- label: Spec Decode Eagle
+ key: spec-decode-eagle
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
@@ -12,7 +13,20 @@ steps:
commands:
- 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
+ optional: true
+ source_file_dependencies:
+ - vllm/v1/spec_decode/
+ - vllm/v1/worker/gpu/spec_decode/
+ - tests/v1/e2e/spec_decode/
+ commands:
+ - 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:
@@ -23,7 +37,21 @@ steps:
commands:
- 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
+ optional: true
+ source_file_dependencies:
+ - vllm/v1/spec_decode/
+ - vllm/v1/worker/gpu/spec_decode/
+ - vllm/transformers_utils/configs/speculators/
+ - tests/v1/e2e/spec_decode/
+ commands:
+ - 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:
@@ -34,6 +62,7 @@ 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:
@@ -42,3 +71,42 @@ steps:
- tests/v1/e2e/spec_decode/
commands:
- 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
+ optional: true
+ source_file_dependencies:
+ - vllm/v1/spec_decode/
+ - vllm/v1/worker/gpu/spec_decode/
+ - tests/v1/e2e/spec_decode/
+ commands:
+ - pytest -v -s v1/e2e/spec_decode -k "draft_model or no_sync or batch_inference"
+
+- label: DFlash Speculators Correctness
+ key: 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_dflash.py
+ commands:
+ - export VLLM_ALLOW_INSECURE_SERIALIZATION=1
+ - pytest -v -s v1/spec_decode/test_speculators_dflash.py -m slow_test
+
+- label: Spec Decode MTP hybrid (B200)
+ timeout_in_minutes: 30
+ device: b200
+ 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"
diff --git a/.buildkite/test_areas/weight_loading.yaml b/.buildkite/test_areas/weight_loading.yaml
index 8e86374a8ad0..01c6bb7809bc 100644
--- a/.buildkite/test_areas/weight_loading.yaml
+++ b/.buildkite/test_areas/weight_loading.yaml
@@ -3,6 +3,7 @@ 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
diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS
index e1a2a5820839..44cf10076ee7 100644
--- a/.github/CODEOWNERS
+++ b/.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
-/vllm/model_executor/layers/quantization @mgoin @robertgshaw2-redhat @tlrmchlsmth @yewentao256 @pavanimajety
+/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/mamba @tdoublep @tomeras91
/vllm/model_executor/layers/mamba/gdn_linear_attn.py @tdoublep @ZJY0516 @vadiklyutiy
/vllm/model_executor/layers/rotary_embedding.py @vadiklyutiy
@@ -18,7 +18,8 @@
/vllm/kernels/helion @ProExpertProg @zou3519
/vllm/multimodal @DarkLight1337 @ywang96 @NickLucche @tjtanaa
/vllm/vllm_flash_attn @LucasWilkinson @MatthewBonanni
-CMakeLists.txt @tlrmchlsmth @LucasWilkinson
+/CMakeLists.txt @tlrmchlsmth @LucasWilkinson @Harry-Chen
+/cmake @tlrmchlsmth @LucasWilkinson @Harry-Chen
# Any change to the VllmConfig changes can have a large user-facing impact,
# so spam a lot of people
@@ -44,8 +45,9 @@ CMakeLists.txt @tlrmchlsmth @LucasWilkinson
/vllm/pooling_params.py @noooop @DarkLight1337
/vllm/tokenizers @DarkLight1337 @njhill
/vllm/renderers @DarkLight1337 @njhill
-/vllm/reasoning @aarnphm @chaunceyjiang
-/vllm/tool_parsers @aarnphm @chaunceyjiang
+/vllm/reasoning @aarnphm @chaunceyjiang @sfeng33 @bbrowning
+/vllm/tool_parsers @aarnphm @chaunceyjiang @sfeng33 @bbrowning
+/vllm/parser @aarnphm @chaunceyjiang @sfeng33 @bbrowning
# vLLM V1
/vllm/v1/attention @LucasWilkinson @MatthewBonanni
@@ -69,6 +71,10 @@ CMakeLists.txt @tlrmchlsmth @LucasWilkinson
/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
@@ -76,11 +82,11 @@ CMakeLists.txt @tlrmchlsmth @LucasWilkinson
/tests/distributed/test_same_node.py @youkaichao
/tests/entrypoints @DarkLight1337 @robertgshaw2-redhat @aarnphm @NickLucche
/tests/evals @mgoin @vadiklyutiy
-/tests/kernels @mgoin @tlrmchlsmth @WoosukKwon @yewentao256
+/tests/kernels @mgoin @tlrmchlsmth @WoosukKwon @yewentao256 @zyongye
/tests/kernels/ir @ProExpertProg @tjtanaa
/tests/models @DarkLight1337 @ywang96
/tests/multimodal @DarkLight1337 @ywang96 @NickLucche
-/tests/quantization @mgoin @robertgshaw2-redhat @yewentao256 @pavanimajety
+/tests/quantization @mgoin @robertgshaw2-redhat @yewentao256 @pavanimajety @zyongye
/tests/test_inputs.py @DarkLight1337 @ywang96
/tests/entrypoints/llm/test_struct_output_generate.py @mgoin @russellb @aarnphm
/tests/v1/structured_output @mgoin @russellb @aarnphm
@@ -91,7 +97,10 @@ CMakeLists.txt @tlrmchlsmth @LucasWilkinson
/tests/v1/kv_connector/nixl_integration @NickLucche
/tests/v1/kv_connector @ApostaC @orozery
/tests/v1/kv_offload @ApostaC @orozery
-/tests/v1/determinism @yewentao256
+/tests/v1/determinism @yewentao256
+/tests/reasoning @aarnphm @chaunceyjiang @sfeng33 @bbrowning
+/tests/tool_parsers @aarnphm @chaunceyjiang @sfeng33 @bbrowning
+/tests/tool_use @aarnphm @chaunceyjiang @sfeng33 @bbrowning
# Transformers modeling backend
/vllm/model_executor/models/transformers @hmellor
@@ -143,6 +152,12 @@ 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
diff --git a/.github/PULL_REQUEST_TEMPLATE.md b/.github/PULL_REQUEST_TEMPLATE.md
index 8043df65d558..8a3934670e44 100644
--- a/.github/PULL_REQUEST_TEMPLATE.md
+++ b/.github/PULL_REQUEST_TEMPLATE.md
@@ -15,7 +15,6 @@ PLEASE FILL IN THE PR DESCRIPTION HERE ENSURING ALL CHECKLIST ITEMS (AT THE BOTT
- [ ] The test plan, such as providing test command.
- [ ] The test results, such as pasting the results comparison before and after, or e2e results
- [ ] (Optional) The necessary documentation update, such as updating `supported_models.md` and `examples` for a new model.
-- [ ] (Optional) Release notes update. If your change is user facing, please update the release notes draft in the [Google Doc](https://docs.google.com/document/d/1YyVqrgX4gHTtrstbq8oWUImOyPCKSGnJ7xtTpmXzlRs/edit?tab=t.0).
**BEFORE SUBMITTING, PLEASE READ ** (anything written below this line will be removed by GitHub Actions)
diff --git a/.github/mergify.yml b/.github/mergify.yml
index baf65e14a882..2d36e3507028 100644
--- a/.github/mergify.yml
+++ b/.github/mergify.yml
@@ -262,7 +262,7 @@ pull_request_rules:
- files~=^docker/Dockerfile.xpu
- files~=^\\.buildkite/intel_jobs/
- files=\.buildkite/ci_config_intel.yaml
- - files=vllm/model_executor/layers/fused_moe/xpu_fused_moe.py
+ - files=vllm/model_executor/layers/fused_moe/experts/xpu_moe.py
- files=vllm/model_executor/kernels/linear/mixed_precision/xpu.py
- files=vllm/model_executor/kernels/linear/mxfp8/xpu.py
- files=vllm/model_executor/kernels/linear/scaled_mm/xpu.py
@@ -308,8 +308,7 @@ 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/offline_inference/structured_outputs.py
- - files=examples/online_serving/structured_outputs/structured_outputs.py
+ - files=^examples/features/structured_outputs/
- files~=^tests/v1/structured_output/
- files=tests/entrypoints/llm/test_struct_output_generate.py
- files~=^vllm/v1/structured_output/
@@ -325,7 +324,7 @@ pull_request_rules:
- or:
- files~=^vllm/v1/spec_decode/
- files~=^tests/v1/spec_decode/
- - files~=^examples/.*(spec_decode|mlpspeculator|eagle|speculation).*\.py
+ - files=^examples/features/speculative_decoding/
- 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
@@ -389,11 +388,7 @@ pull_request_rules:
- files~=^tests/entrypoints/anthropic/.*tool.*
- files~=^vllm/tool_parsers/
- files=docs/features/tool_calling.md
- - 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
+ - files~=^examples/tool_calling/
actions:
label:
add:
@@ -482,9 +477,7 @@ pull_request_rules:
conditions:
- label != stale
- or:
- - files~=^examples/online_serving/disaggregated[^/]*/.*
- - files~=^examples/offline_inference/disaggregated[^/]*/.*
- - files~=^examples/others/lmcache/
+ - files~=^examples/disaggregated/
- files~=^tests/v1/kv_connector/
- files~=^vllm/distributed/kv_transfer/
- title~=(?i)\bP/?D\b
diff --git a/.github/workflows/macos-smoke-test.yml b/.github/workflows/macos-smoke-test.yml
index b8e7cebf2dc1..ea1c8b0feac3 100644
--- a/.github/workflows/macos-smoke-test.yml
+++ b/.github/workflows/macos-smoke-test.yml
@@ -45,6 +45,7 @@ jobs:
- name: Smoke test vllm serve
run: |
# Start server in background
+ VLLM_CPU_KVCACHE_SPACE=1 \
vllm serve Qwen/Qwen3-0.6B \
--max-model-len=2K \
--load-format=dummy \
diff --git a/.github/workflows/new_pr_bot.yml b/.github/workflows/new_pr_bot.yml
index ef5e30952c62..27100f9f4da0 100644
--- a/.github/workflows/new_pr_bot.yml
+++ b/.github/workflows/new_pr_bot.yml
@@ -62,14 +62,14 @@ jobs:
const prAuthor = context.payload.pull_request.user.login;
const { data: searchResults } = await github.rest.search.issuesAndPullRequests({
- q: `repo:${owner}/${repo} type:pr author:${prAuthor}`,
+ q: `repo:${owner}/${repo} type:pr is:merged author:${prAuthor}`,
per_page: 1,
});
- const authorPRCount = searchResults.total_count;
- console.log(`Found ${authorPRCount} PRs by ${prAuthor}`);
+ const mergedPRCount = searchResults.total_count;
+ console.log(`Found ${mergedPRCount} merged PRs by ${prAuthor}`);
- if (authorPRCount === 1) {
+ if (mergedPRCount === 0) {
console.log(`Posting welcome comment for first-time contributor: ${prAuthor}`);
await github.rest.issues.createComment({
owner,
@@ -98,5 +98,5 @@ jobs:
].join('\n'),
});
} else {
- console.log(`Skipping comment for ${prAuthor} - not their first PR (${authorPRCount} PRs found)`);
+ console.log(`Skipping comment for ${prAuthor} - not a first-time contributor (${mergedPRCount} merged PRs)`);
}
diff --git a/.github/workflows/pre-commit.yml b/.github/workflows/pre-commit.yml
index 8ab8d3e7035f..1dd31b0e50f6 100644
--- a/.github/workflows/pre-commit.yml
+++ b/.github/workflows/pre-commit.yml
@@ -16,11 +16,7 @@ permissions:
jobs:
pre-run-check:
- if: >-
- github.event_name == 'pull_request' &&
- (github.event.action != 'labeled' ||
- github.event.label.name == 'ready' ||
- github.event.label.name == 'verified')
+ if: github.event_name == 'pull_request'
runs-on: ubuntu-latest
steps:
- name: Check PR label and author merge count
@@ -49,12 +45,7 @@ jobs:
pre-commit:
needs: pre-run-check
- 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')
+ if: always() && (needs.pre-run-check.result == 'success' || needs.pre-run-check.result == 'skipped')
runs-on: ubuntu-latest
steps:
- uses: actions/checkout@8e8c483db84b4bee98b60c0593521ed34d9990e8 # v6.0.1
diff --git a/.github/workflows/scripts/build.sh b/.github/workflows/scripts/build.sh
index 30c63ee0da91..5bb9f7196806 100644
--- a/.github/workflows/scripts/build.sh
+++ b/.github/workflows/scripts/build.sh
@@ -14,7 +14,7 @@ $python_executable -m pip install -r requirements/build/cuda.txt -r requirements
# Limit the number of parallel jobs to avoid OOM
export MAX_JOBS=1
# Make sure release wheels are built for the following architectures
-export TORCH_CUDA_ARCH_LIST="7.0 7.5 8.0 8.6 8.9 9.0+PTX"
+export TORCH_CUDA_ARCH_LIST="7.5 8.0 8.6 8.9 9.0 10.0 12.0+PTX"
bash tools/check_repo.sh
diff --git a/.gitignore b/.gitignore
index 134bbc5cc893..e53d19b35340 100644
--- a/.gitignore
+++ b/.gitignore
@@ -237,6 +237,7 @@ 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
diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml
index 33b1db69dec4..f1e0afebf213 100644
--- a/.pre-commit-config.yaml
+++ b/.pre-commit-config.yaml
@@ -131,6 +131,16 @@ 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
diff --git a/.readthedocs.yaml b/.readthedocs.yaml
index 1e479fd03d91..c1029af69437 100644
--- a/.readthedocs.yaml
+++ b/.readthedocs.yaml
@@ -9,7 +9,44 @@ build:
python: "3.12"
jobs:
post_checkout:
- # - bash docs/maybe_skip_pr_build.sh
+ - |
+ 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
- git fetch origin main --unshallow --no-tags --filter=blob:none || true
pre_create_environment:
- pip install uv
diff --git a/CMakeLists.txt b/CMakeLists.txt
index f24c12eff83c..38af64851ed9 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -13,8 +13,12 @@ cmake_minimum_required(VERSION 3.26)
# cmake --install . --component _C
project(vllm_extensions LANGUAGES CXX)
-set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_STANDARD 20)
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)
@@ -34,10 +38,10 @@ install(CODE "set(CMAKE_INSTALL_LOCAL_ONLY TRUE)" ALL_COMPONENTS)
# Supported python versions. These versions will be searched in order, the
# first match will be selected. These should be kept in sync with setup.py.
#
-set(PYTHON_SUPPORTED_VERSIONS "3.10" "3.11" "3.12" "3.13")
+set(PYTHON_SUPPORTED_VERSIONS "3.10" "3.11" "3.12" "3.13" "3.14")
# Supported AMD GPU architectures.
-set(HIP_SUPPORTED_ARCHS "gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201")
+set(HIP_SUPPORTED_ARCHS "gfx906;gfx908;gfx90a;gfx942;gfx950;gfx1030;gfx1100;gfx1101;gfx1102;gfx1103;gfx1150;gfx1151;gfx1152;gfx1153;gfx1200;gfx1201")
# ROCm installation prefix. Default to /opt/rocm but allow override via
# -DROCM_PATH=/your/rocm/path when invoking cmake.
@@ -94,12 +98,15 @@ find_package(Torch REQUIRED)
# This check must happen after find_package(Torch) because that's when CMAKE_CUDA_COMPILER_VERSION gets defined
if(DEFINED CMAKE_CUDA_COMPILER_VERSION AND
CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 13.0)
- set(CUDA_SUPPORTED_ARCHS "7.5;8.0;8.6;8.7;8.9;9.0;10.0;11.0;12.0;12.1")
+ # starting from CUDA 12.9 and Blackwell (10.0), we use family-specific targets (10.0f, 12.0f, etc)
+ # to support the whole generation without specifying all sub-architectures
+ # see: https://developer.nvidia.com/blog/nvidia-blackwell-and-nvidia-cuda-12-9-introduce-family-specific-architecture-features/
+ set(CUDA_SUPPORTED_ARCHS "7.5;8.0;8.6;8.7;8.9;9.0;10.0;11.0;12.0")
elseif(DEFINED CMAKE_CUDA_COMPILER_VERSION AND
CMAKE_CUDA_COMPILER_VERSION VERSION_GREATER_EQUAL 12.8)
- set(CUDA_SUPPORTED_ARCHS "7.0;7.2;7.5;8.0;8.6;8.7;8.9;9.0;10.0;10.1;12.0;12.1")
+ set(CUDA_SUPPORTED_ARCHS "7.5;8.0;8.6;8.7;8.9;9.0;10.0;10.1;10.3;12.0;12.1")
else()
- set(CUDA_SUPPORTED_ARCHS "7.0;7.2;7.5;8.0;8.6;8.7;8.9;9.0")
+ set(CUDA_SUPPORTED_ARCHS "7.0;7.5;8.0;8.6;8.7;8.9;9.0")
endif()
#
@@ -304,10 +311,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/torch_bindings.cpp"
+ "csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu")
if(VLLM_GPU_LANG STREQUAL "CUDA")
- list(APPEND VLLM_EXT_SRC "csrc/minimax_reduce_rms_kernel.cu")
+ list(APPEND VLLM_EXT_SRC
+ "csrc/minimax_reduce_rms_kernel.cu")
SET(CUTLASS_ENABLE_HEADERS_ONLY ON CACHE BOOL "Enable only the header library")
@@ -923,6 +932,14 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
SRCS "${SRCS}"
CUDA_ARCHS "${FP4_ARCHS}")
list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
+ # nvfp4_kv_cache_kernels uses non-stable torch API and is called directly
+ # from cache_kernels.cu, so it belongs in _C rather than _C_stable.
+ set(NVFP4_KV_SRC "csrc/nvfp4_kv_cache_kernels.cu")
+ set_gencode_flags_for_srcs(
+ SRCS "${NVFP4_KV_SRC}"
+ CUDA_ARCHS "${FP4_ARCHS}")
+ target_sources(_C PRIVATE ${NVFP4_KV_SRC})
+ target_compile_definitions(_C PRIVATE ENABLE_NVFP4_SM120=1)
list(APPEND VLLM_GPU_FLAGS "-DENABLE_NVFP4_SM120=1")
list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MOE_SM120=1")
message(STATUS "Building NVFP4 for archs: ${FP4_ARCHS}")
@@ -944,11 +961,19 @@ if(VLLM_GPU_LANG STREQUAL "CUDA")
"csrc/libtorch_stable/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu"
"csrc/libtorch_stable/quantization/fp4/nvfp4_experts_quant.cu"
"csrc/libtorch_stable/quantization/fp4/nvfp4_scaled_mm_kernels.cu"
- "csrc/libtorch_stable/quantization/fp4/nvfp4_blockwise_moe_kernel.cu")
+ "csrc/libtorch_stable/quantization/fp4/nvfp4_blockwise_moe_kernel.cu"
+ "csrc/libtorch_stable/quantization/fp4/mxfp4_experts_quant.cu"
+ "csrc/libtorch_stable/quantization/fp4/mxfp4_blockwise_moe_kernel.cu")
set_gencode_flags_for_srcs(
SRCS "${SRCS}"
CUDA_ARCHS "${FP4_ARCHS}")
list(APPEND VLLM_STABLE_EXT_SRC "${SRCS}")
+ set(NVFP4_KV_SRC "csrc/nvfp4_kv_cache_kernels.cu")
+ set_gencode_flags_for_srcs(
+ SRCS "${NVFP4_KV_SRC}"
+ CUDA_ARCHS "${FP4_ARCHS}")
+ target_sources(_C PRIVATE ${NVFP4_KV_SRC})
+ target_compile_definitions(_C PRIVATE ENABLE_NVFP4_SM100=1)
list(APPEND VLLM_GPU_FLAGS "-DENABLE_NVFP4_SM100=1")
list(APPEND VLLM_GPU_FLAGS "-DENABLE_CUTLASS_MOE_SM100=1")
message(STATUS "Building NVFP4 for archs: ${FP4_ARCHS}")
@@ -1026,13 +1051,13 @@ 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_softmax_kernels.cu"
+ "csrc/moe/topk_softplus_sqrt_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/router_gemm.cu")
+ "csrc/moe/grouped_topk_kernels.cu")
endif()
if(VLLM_GPU_LANG STREQUAL "CUDA")
diff --git a/README.md b/README.md
index d94e33ba8b0a..42777436c63d 100644
--- a/README.md
+++ b/README.md
@@ -14,7 +14,7 @@ Easy, fast, and cheap LLM serving for everyone
| Documentation | Blog | Paper | Twitter/X | User Forum | Developer Slack |
-🔥 We have built a vllm website to help you get started with vllm. Please visit [vllm.ai](https://vllm.ai) to learn more.
+🔥 We have built a vLLM website to help you get started with vLLM. Please visit [vllm.ai](https://vllm.ai) to learn more.
For events, please visit [vllm.ai/events](https://vllm.ai/events) to join us.
---
@@ -50,7 +50,7 @@ vLLM is flexible and easy to use with:
- Efficient multi-LoRA support for dense and MoE layers
- Support for NVIDIA GPUs, AMD GPUs, and x86/ARM/PowerPC CPUs. Additionally, diverse hardware plugins such as Google TPUs, Intel Gaudi, IBM Spyre, Huawei Ascend, Rebellions NPU, Apple Silicon, MetaX GPU, and more.
-vLLM seamlessly supports 200+ model architectures on HuggingFace, including:
+vLLM seamlessly supports 200+ model architectures on Hugging Face, including:
- Decoder-only LLMs (e.g., Llama, Qwen, Gemma)
- Mixture-of-Expert LLMs (e.g., Mixtral, DeepSeek-V3, Qwen-MoE, GPT-OSS)
diff --git a/tests/entrypoints/pooling/pooling/__init__.py b/benchmarks/__init__.py
similarity index 100%
rename from tests/entrypoints/pooling/pooling/__init__.py
rename to benchmarks/__init__.py
diff --git a/benchmarks/attention_benchmarks/mla_runner.py b/benchmarks/attention_benchmarks/mla_runner.py
index f8bc7b4a10ed..b58ddffcd0b1 100644
--- a/benchmarks/attention_benchmarks/mla_runner.py
+++ b/benchmarks/attention_benchmarks/mla_runner.py
@@ -404,6 +404,7 @@ def _build_attention_metadata(
query_start_loc=q_start_gpu,
query_start_loc_cpu=q_start_cpu,
seq_lens=seq_lens_gpu,
+ seq_lens_cpu_upper_bound=seq_lens_cpu,
_seq_lens_cpu=seq_lens_cpu,
_num_computed_tokens_cpu=num_computed_tokens_cpu,
slot_mapping=slot_mapping,
diff --git a/benchmarks/kernels/__init__.py b/benchmarks/kernels/__init__.py
new file mode 100644
index 000000000000..e69de29bb2d1
diff --git a/benchmarks/kernels/benchmark_cutlass_moe_fp8.py b/benchmarks/kernels/benchmark_cutlass_moe_fp8.py
index 3f80b024e108..03d7fb386f74 100644
--- a/benchmarks/kernels/benchmark_cutlass_moe_fp8.py
+++ b/benchmarks/kernels/benchmark_cutlass_moe_fp8.py
@@ -16,7 +16,7 @@
maybe_make_prepare_finalize,
)
from vllm.model_executor.layers.fused_moe.config import fp8_w8a8_moe_quant_config
-from vllm.model_executor.layers.fused_moe.cutlass_moe import CutlassExpertsFp8
+from vllm.model_executor.layers.fused_moe.experts.cutlass_moe import CutlassExpertsFp8
from vllm.model_executor.layers.fused_moe.fused_moe import fused_experts, fused_topk
from vllm.platforms import current_platform
from vllm.utils.argparse_utils import FlexibleArgumentParser
diff --git a/benchmarks/kernels/benchmark_cutlass_moe_nvfp4.py b/benchmarks/kernels/benchmark_cutlass_moe_nvfp4.py
index 2d4afd38c097..7379bf858889 100644
--- a/benchmarks/kernels/benchmark_cutlass_moe_nvfp4.py
+++ b/benchmarks/kernels/benchmark_cutlass_moe_nvfp4.py
@@ -22,7 +22,7 @@
fp8_w8a8_moe_quant_config,
nvfp4_moe_quant_config,
)
-from vllm.model_executor.layers.fused_moe.cutlass_moe import (
+from vllm.model_executor.layers.fused_moe.experts.cutlass_moe import (
CutlassExpertsFp4,
)
from vllm.model_executor.layers.fused_moe.fused_moe import fused_experts, fused_topk
diff --git a/benchmarks/kernels/benchmark_grouped_gemm_cutlass.py b/benchmarks/kernels/benchmark_grouped_gemm_cutlass.py
index dd4060bbdb94..04fc2960d1e4 100644
--- a/benchmarks/kernels/benchmark_grouped_gemm_cutlass.py
+++ b/benchmarks/kernels/benchmark_grouped_gemm_cutlass.py
@@ -13,7 +13,7 @@
maybe_make_prepare_finalize,
)
from vllm.model_executor.layers.fused_moe.config import fp8_w8a8_moe_quant_config
-from vllm.model_executor.layers.fused_moe.cutlass_moe import CutlassExpertsFp8
+from vllm.model_executor.layers.fused_moe.experts.cutlass_moe import CutlassExpertsFp8
from vllm.model_executor.layers.fused_moe.fused_moe import (
fused_experts,
fused_topk,
diff --git a/benchmarks/kernels/benchmark_vit_fp8_attn.py b/benchmarks/kernels/benchmark_vit_fp8_attn.py
new file mode 100644
index 000000000000..7d7a067dde9d
--- /dev/null
+++ b/benchmarks/kernels/benchmark_vit_fp8_attn.py
@@ -0,0 +1,324 @@
+# 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,
+ )
diff --git a/benchmarks/kernels/ir/__init__.py b/benchmarks/kernels/ir/__init__.py
new file mode 100644
index 000000000000..e69de29bb2d1
diff --git a/benchmarks/kernels/ir/bench_ir_ops.py b/benchmarks/kernels/ir/bench_ir_ops.py
new file mode 100644
index 000000000000..b23c4e8ae327
--- /dev/null
+++ b/benchmarks/kernels/ir/bench_ir_ops.py
@@ -0,0 +1,378 @@
+#!/usr/bin/env python3
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""
+Generic benchmark harness for vLLM IR ops.
+
+Usage:
+ python benchmarks/kernels/ir/bench_ir_ops.py
+ python benchmarks/kernels/ir/bench_ir_ops.py --ops rms_norm
+ python benchmarks/kernels/ir/bench_ir_ops.py --ops rms_norm,silu_mul
+ python benchmarks/kernels/ir/bench_ir_ops.py --no-cuda-graph
+ python benchmarks/kernels/ir/bench_ir_ops.py --ops rms_norm --save-path ./results/
+"""
+
+import argparse
+import contextlib
+import csv
+import dataclasses
+import datetime
+import math
+import os
+import subprocess
+import sys
+import tempfile
+
+# Ensure repo root is on sys.path so `benchmarks` is importable as a package.
+_REPO_ROOT = os.path.abspath(os.path.join(os.path.dirname(__file__), "../../.."))
+if _REPO_ROOT not in sys.path:
+ sys.path.insert(0, _REPO_ROOT)
+
+# Suppress noisy C++ warnings from vllm kernel registration (written to fd 2
+# directly by the dynamic linker, so Python-level sys.stderr redirect won't
+# catch them).
+_saved_fd = os.dup(2)
+try:
+ with open(os.devnull, "w") as _devnull:
+ os.dup2(_devnull.fileno(), 2)
+ import torch
+
+ import vllm.kernels # noqa: E402, F401
+finally:
+ os.dup2(_saved_fd, 2)
+ os.close(_saved_fd)
+
+from tqdm import tqdm # noqa: E402
+
+from benchmarks.kernels.ir.shapes import SHAPE_CONFIGS # noqa: E402 # isort: skip
+from vllm.ir.op import IrOp # noqa: E402
+from vllm.platforms import current_platform # noqa: E402
+from vllm.triton_utils import triton # noqa: E402
+
+
+@dataclasses.dataclass(frozen=True)
+class BenchConfig:
+ use_cuda_graph: bool = True
+ warmup: int = 25
+ rep: int = 100
+
+
+def _pkg_version(name: str) -> str:
+ from importlib.metadata import PackageNotFoundError, version
+
+ with contextlib.suppress(PackageNotFoundError):
+ return version(name)
+ return "not installed"
+
+
+_METADATA_LABELS = {
+ "timestamp": "Timestamp",
+ "git_commit": "Git commit",
+ "vllm": "vLLM",
+ "pytorch": "PyTorch",
+ "cuda_runtime": "CUDA runtime",
+ "triton": "Triton",
+ "cutlass": "CUTLASS",
+ "helion": "Helion",
+ "device": "Device",
+ "bench_mode": "Bench mode",
+ "warmup": "Warmup",
+ "rep": "Repetitions",
+}
+
+
+def collect_env_metadata(cfg: BenchConfig) -> dict[str, str]:
+ from vllm.collect_env import get_env_info
+
+ env = get_env_info()
+
+ git_sha = "unknown"
+ with contextlib.suppress(subprocess.CalledProcessError, FileNotFoundError):
+ git_sha = (
+ subprocess.check_output(
+ ["git", "rev-parse", "--short", "HEAD"], stderr=subprocess.DEVNULL
+ )
+ .decode()
+ .strip()
+ )
+
+ device_name = current_platform.get_device_name()
+
+ warmup_note = " ms" if not cfg.use_cuda_graph else " ms (ignored)"
+ rep_note = " replays" if cfg.use_cuda_graph else " ms"
+
+ return {
+ "timestamp": datetime.datetime.now().strftime("%Y-%m-%d %H:%M:%S"),
+ "git_commit": git_sha,
+ "vllm": str(env.vllm_version),
+ "pytorch": str(env.torch_version),
+ "cuda_runtime": str(env.cuda_runtime_version),
+ "triton": triton.__version__,
+ "cutlass": _pkg_version("nvidia-cutlass-dsl"),
+ "helion": _pkg_version("helion"),
+ "device": device_name,
+ "bench_mode": "cuda_graph" if cfg.use_cuda_graph else "eager",
+ "warmup": f"{cfg.warmup}{warmup_note}",
+ "rep": f"{cfg.rep}{rep_note}",
+ }
+
+
+def print_metadata(metadata: dict[str, str]):
+ print("=" * 60)
+ for key, val in metadata.items():
+ print(f"{_METADATA_LABELS.get(key, key) + ':':<16}{val}")
+ print("=" * 60)
+
+
+def _clone_args(args: tuple) -> tuple:
+ return tuple(a.clone() if isinstance(a, torch.Tensor) else a for a in args)
+
+
+# TODO(gmagogsfm): When the `maybe_inplace` PR lands, ops marked as
+# inplace=True will mutate bench_args across iterations. Both CUDA graph
+# and eager modes will accumulate drift from repeated in-place mutation.
+# We need to re-clone inputs per iteration for inplace ops.
+def _bench_one(fn, args, cfg: BenchConfig) -> float:
+ bench_args = _clone_args(args)
+ bench_fn = lambda: fn(*bench_args)
+
+ if cfg.use_cuda_graph:
+ ms = triton.testing.do_bench_cudagraph(bench_fn, rep=cfg.rep, quantiles=[0.5])
+ else:
+ ms = triton.testing.do_bench(
+ bench_fn, warmup=cfg.warmup, rep=cfg.rep, quantiles=[0.5]
+ )
+ return ms * 1000
+
+
+# TODO(gmagogsfm): Once compiled native implementation lands (#38775),
+# the benchmark baseline should be the compiled native (what vLLM runs by
+# default) rather than the uncompiled native implementation.
+def collect_timings(
+ op: IrOp, shape_configs: list[dict], cfg: BenchConfig
+) -> tuple[list[str], list[str], dict[str, dict[str, float]]]:
+ def fmt(v) -> str:
+ return str(v).split(".")[-1] if isinstance(v, torch.dtype) else str(v)
+
+ case_names = [
+ "_".join(f"{k}={fmt(v)}" for k, v in kwargs.items()) for kwargs in shape_configs
+ ]
+ providers = [n for n, impl in op.impls.items() if impl.supported]
+
+ results: dict[str, dict[str, float]] = {c: {} for c in case_names}
+ for provider in providers:
+ impl = op.impls[provider]
+ desc = f"{op.name} / {provider}"
+ for case_name, kwargs in tqdm(
+ zip(case_names, shape_configs),
+ desc=desc,
+ total=len(case_names),
+ unit=" cases",
+ ):
+ args = op.generate_inputs(**kwargs)
+ if impl.supports_args(*args):
+ results[case_name][provider] = _bench_one(impl.impl_fn, args, cfg)
+ else:
+ results[case_name][provider] = float("nan")
+
+ return case_names, providers, results
+
+
+def analyze_results(
+ op_name: str,
+ case_names: list[str],
+ providers: list[str],
+ results: dict[str, dict[str, float]],
+) -> tuple[list[dict[str, str]], list[dict[str, str]], list[str]]:
+ native_col = "native"
+ non_native = [p for p in providers if p != native_col]
+
+ header_cols = ["case"]
+ for p in providers:
+ header_cols.append(f"{p} (us)")
+ for p in non_native:
+ header_cols.append(f"{p} speedup")
+
+ detail_rows: list[dict[str, str]] = []
+ speedup_data: dict[str, list[tuple[float, str]]] = {p: [] for p in non_native}
+
+ for case_name in case_names:
+ timings = results[case_name]
+ row: dict[str, str] = {"case": case_name}
+
+ for p in providers:
+ val = timings.get(p, float("nan"))
+ row[f"{p} (us)"] = f"{val:.2f}" if not math.isnan(val) else "n/a"
+
+ native_us = timings.get(native_col, float("nan"))
+ for p in non_native:
+ p_us = timings.get(p, float("nan"))
+ if not math.isnan(native_us) and not math.isnan(p_us) and p_us > 0:
+ speedup = native_us / p_us
+ row[f"{p} speedup"] = f"{speedup:.2f}x"
+ speedup_data[p].append((speedup, case_name))
+ else:
+ row[f"{p} speedup"] = "n/a"
+
+ detail_rows.append(row)
+
+ summary_rows: list[dict[str, str]] = []
+ for p in non_native:
+ entries = speedup_data[p]
+ if not entries:
+ continue
+ speedups = [s for s, _ in entries]
+ geomean = math.exp(sum(math.log(s) for s in speedups) / len(speedups))
+ best_val, best_case = max(entries)
+ worst_val, worst_case = min(entries)
+ wins = sum(1 for s in speedups if s > 1.0)
+ losses = sum(1 for s in speedups if s < 1.0)
+ total = len(speedups)
+
+ print(f"\n{p} vs native ({wins}/{total} faster, {losses}/{total} slower):")
+ print(f" geomean speedup: {geomean:.2f}x")
+ print(f" best: {best_val:.2f}x ({best_case})")
+ print(f" worst: {worst_val:.2f}x ({worst_case})")
+
+ summary_rows.append(
+ {
+ "op": op_name,
+ "provider": p,
+ "geomean_speedup": f"{geomean:.2f}",
+ "best_speedup": f"{best_val:.2f}",
+ "best_case": best_case,
+ "worst_speedup": f"{worst_val:.2f}",
+ "worst_case": worst_case,
+ "wins": str(wins),
+ "losses": str(losses),
+ "total": str(total),
+ }
+ )
+
+ return detail_rows, summary_rows, header_cols
+
+
+def write_csv(path: str, rows: list[dict[str, str]], fieldnames: list[str]):
+ with open(path, "w", newline="") as f:
+ writer = csv.DictWriter(f, fieldnames=fieldnames)
+ writer.writeheader()
+ writer.writerows(rows)
+
+
+def save_results(
+ save_dir: str,
+ op_name: str,
+ detail_rows: list[dict[str, str]],
+ header_cols: list[str],
+ all_summary_rows: list[dict[str, str]],
+ metadata: dict[str, str],
+):
+ write_csv(
+ os.path.join(save_dir, f"{op_name}_detail.csv"),
+ detail_rows,
+ header_cols,
+ )
+ if all_summary_rows:
+ write_csv(
+ os.path.join(save_dir, "summary.csv"),
+ all_summary_rows,
+ list(all_summary_rows[0].keys()),
+ )
+ write_csv(
+ os.path.join(save_dir, "metadata.csv"),
+ [metadata],
+ list(metadata.keys()),
+ )
+
+
+def parse_args():
+ parser = argparse.ArgumentParser(description="Benchmark vLLM IR ops")
+ parser.add_argument(
+ "--ops",
+ type=str,
+ default=None,
+ help="Comma-separated list of op names to benchmark (substring match)",
+ )
+ parser.add_argument(
+ "--no-cuda-graph",
+ action="store_true",
+ help="Disable CUDA graph; use do_bench with L2 cache flushing instead",
+ )
+ parser.add_argument(
+ "--warmup",
+ type=int,
+ default=25,
+ help="Warmup time in ms (do_bench) or ignored with CUDA graph (default: 25)",
+ )
+ parser.add_argument(
+ "--rep",
+ type=int,
+ default=100,
+ help="Repetition time in ms (do_bench) or number of graph replays "
+ "(do_bench_cudagraph) (default: 100)",
+ )
+ parser.add_argument(
+ "--save-path",
+ type=str,
+ default=None,
+ help="Directory to save results (default: auto-created temp dir)",
+ )
+ return parser.parse_args()
+
+
+def main():
+ args = parse_args()
+ cfg = BenchConfig(
+ use_cuda_graph=not args.no_cuda_graph,
+ warmup=args.warmup,
+ rep=args.rep,
+ )
+
+ torch.set_default_device(current_platform.device_type)
+
+ metadata = collect_env_metadata(cfg)
+ print_metadata(metadata)
+
+ timestamp = datetime.datetime.now().strftime("%Y%m%d_%H%M%S")
+ save_dir = args.save_path or os.path.join(
+ tempfile.gettempdir(), f"vllm_ir_bench_{timestamp}"
+ )
+ os.makedirs(save_dir, exist_ok=True)
+
+ op_filters = [f.strip() for f in args.ops.split(",")] if args.ops else None
+ all_summary_rows: list[dict[str, str]] = []
+
+ for op in IrOp.registry.values():
+ if op_filters and not any(f in op.name for f in op_filters):
+ continue
+ if not op.has_input_generator:
+ print(f"Skipping op '{op.name}': no input generator registered")
+ continue
+ if op.name not in SHAPE_CONFIGS:
+ raise RuntimeError(
+ f"No benchmark shape config for op '{op.name}'. "
+ f"Add it to benchmarks/kernels/ir/shapes.py"
+ )
+
+ case_names, providers, results = collect_timings(
+ op, SHAPE_CONFIGS[op.name], cfg
+ )
+ detail_rows, summary_rows, header_cols = analyze_results(
+ op.name, case_names, providers, results
+ )
+ all_summary_rows.extend(summary_rows)
+
+ save_results(
+ save_dir,
+ op.name,
+ detail_rows,
+ header_cols,
+ all_summary_rows,
+ metadata,
+ )
+
+ print(f"\nResults saved to: {save_dir}")
+
+
+if __name__ == "__main__":
+ main()
diff --git a/benchmarks/kernels/ir/shapes.py b/benchmarks/kernels/ir/shapes.py
new file mode 100644
index 000000000000..6cc44cf6cec1
--- /dev/null
+++ b/benchmarks/kernels/ir/shapes.py
@@ -0,0 +1,29 @@
+# SPDX-License-Identifier: Apache-2.0
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+"""
+Shape configurations for IR op benchmarks.
+"""
+
+import torch
+
+NUM_TOKENS = [1, 2, 4, 16, 64, 256, 1024, 4096, 16384]
+COMMON_HIDDEN_SIZES = [
+ 2048, # Llama 3.2 1B, Qwen 3 MoE 30B-A3B, Gemma 3n
+ 3072, # Gemma 7B/9B
+ 4096, # Llama 3 8B, Qwen 3 8B, Mistral 7B
+ 5120, # Llama 4 Scout 17B-16E
+ 7168, # DeepSeek V3
+ 8192, # Llama 3 70B
+ 16384, # Llama 3 405B
+]
+
+# Each entry maps an op name to a list of kwarg dicts that will be passed
+# to that op's registered input generator via op.generate_inputs(**kwargs).
+SHAPE_CONFIGS: dict[str, list[dict]] = {
+ "rms_norm": [
+ {"num_tokens": n, "hidden_size": d, "dtype": dtype}
+ for dtype in [torch.float16, torch.bfloat16, torch.float32]
+ for d in COMMON_HIDDEN_SIZES
+ for n in NUM_TOKENS
+ ],
+}
diff --git a/benchmarks/multi_turn/benchmark_serving_multi_turn.py b/benchmarks/multi_turn/benchmark_serving_multi_turn.py
index 881039f43f07..2f56099c66fd 100644
--- a/benchmarks/multi_turn/benchmark_serving_multi_turn.py
+++ b/benchmarks/multi_turn/benchmark_serving_multi_turn.py
@@ -217,6 +217,7 @@ 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,
@@ -225,6 +226,9 @@ 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}
@@ -419,6 +423,7 @@ async def send_turn(
min_tokens,
max_tokens,
req_args.timeout_sec,
+ conversation_id=conv_id,
)
if response.valid is False:
@@ -1468,6 +1473,12 @@ 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)
@@ -1510,7 +1521,9 @@ async def main() -> None:
np.random.seed(args.seed)
logger.info("Loading tokenizer")
- tokenizer = AutoTokenizer.from_pretrained(args.model)
+ tokenizer = AutoTokenizer.from_pretrained(
+ args.model, trust_remote_code=args.trust_remote_code
+ )
await get_server_info(args.url)
diff --git a/cmake/cpu_extension.cmake b/cmake/cpu_extension.cmake
index 0389ce1299a8..361f08b51054 100644
--- a/cmake/cpu_extension.cmake
+++ b/cmake/cpu_extension.cmake
@@ -1,7 +1,7 @@
include(FetchContent)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
-set(CMAKE_CXX_STANDARD 17)
+set(CMAKE_CXX_STANDARD 20)
set(CMAKE_CXX_EXTENSIONS ON)
set(CMAKE_EXPORT_COMPILE_COMMANDS ON)
@@ -30,6 +30,26 @@ else()
list(APPEND CXX_COMPILE_FLAGS
"-fopenmp"
"-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.
+ 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
+ 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()
if (NOT MACOSX_FOUND)
@@ -146,16 +166,49 @@ elseif (S390_FOUND)
"-mtune=native")
elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
message(STATUS "RISC-V detected")
- if(RVV_BF16_FOUND)
- message(STATUS "BF16 extension detected")
- set(MARCH_FLAGS -march=rv64gcv_zvfh_zfbfmin_zvfbfmin_zvl128b -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_zvl128b -mrvv-vector-bits=zvl -mabi=lp64d)
+ # VLLM_RVV_VLEN selects the target VLEN. Auto-detected from /proc/cpuinfo
+ # by default; override with -DVLLM_RVV_VLEN=128 or -DVLLM_RVV_VLEN=256.
+ if(NOT DEFINED VLLM_RVV_VLEN)
+ # Auto-detect: find the largest zvlb in /proc/cpuinfo isa line.
+ if(EXISTS /proc/cpuinfo)
+ file(READ /proc/cpuinfo _cpuinfo)
+ set(_best 0)
+ foreach(_n IN ITEMS 128 256 512 1024)
+ if(_cpuinfo MATCHES "zvl${_n}b")
+ set(_best ${_n})
+ endif()
+ endforeach()
+ if(_best GREATER 0)
+ set(VLLM_RVV_VLEN ${_best})
+ endif()
+ endif()
+ # If auto-detect failed (no /proc/cpuinfo or no zvlb reported)
+ # but the compiler supports RVV, require explicit specification.
+ if(NOT DEFINED VLLM_RVV_VLEN AND (RVV_FP16_FOUND OR RVV_BF16_FOUND))
+ message(FATAL_ERROR
+ "RISC-V RVV is available but VLEN could not be auto-detected. "
+ "Please specify VLEN explicitly:\n"
+ " -DVLLM_RVV_VLEN=128 (for VLEN=128 hardware)\n"
+ " -DVLLM_RVV_VLEN=256 (for VLEN=256 hardware, e.g. Spacemit X100)\n"
+ " -DVLLM_RVV_VLEN=0 (force scalar, no RVV)")
+ endif()
+ endif()
+ if(VLLM_RVV_VLEN AND VLLM_RVV_VLEN GREATER 0)
+ message(STATUS "RISC-V RVV VLEN=${VLLM_RVV_VLEN}")
+ 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)
+ else()
+ message(STATUS "compile riscv with scalar (no FP16/BF16)")
+ set(MARCH_FLAGS -march=rv64gc)
+ endif()
else()
message(STATUS "compile riscv with scalar")
- list(APPEND CXX_COMPILE_FLAGS "-march=rv64gc")
+ set(MARCH_FLAGS -march=rv64gc)
endif()
list(APPEND CXX_COMPILE_FLAGS ${MARCH_FLAGS})
else()
@@ -175,20 +228,6 @@ if (ENABLE_X86_ISA OR (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND) OR POWER9_FOUND
if(NOT NPROC)
set(NPROC 4)
endif()
- # locate PyTorch's libgomp (e.g. site-packages/torch.libs/libgomp-947d5fa1.so.1.0.0)
- # and create a local shim dir with it
- vllm_prepare_torch_gomp_shim(VLLM_TORCH_GOMP_SHIM_DIR)
-
- 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}")
- endif()
# Fetch and populate ACL
if(DEFINED ENV{ACL_ROOT_DIR} AND IS_DIRECTORY "$ENV{ACL_ROOT_DIR}")
@@ -287,14 +326,6 @@ 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)
@@ -307,8 +338,14 @@ 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)
- target_compile_options(dnnl_ext PRIVATE ${DNNL_COMPILE_FLAGS} -fPIC)
+ 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()
list(APPEND LIBS dnnl_ext)
+
+
set(USE_ONEDNN ON)
else()
set(USE_ONEDNN OFF)
@@ -360,6 +397,7 @@ set(VLLM_EXT_SRC
if (ASIMD_FOUND AND NOT APPLE_SILICON_FOUND)
set(VLLM_EXT_SRC
"csrc/cpu/shm.cpp"
+ "csrc/cpu/activation_lut_bf16.cpp"
${VLLM_EXT_SRC})
endif()
@@ -371,12 +409,15 @@ 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
@@ -395,10 +436,11 @@ 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"
@@ -413,7 +455,7 @@ if (ENABLE_X86_ISA)
set(_C_LIBS numa dnnl_ext)
set(_C_AVX512_LIBS numa dnnl_ext)
- set(_C_AVX2_LIBS numa)
+ set(_C_AVX2_LIBS numa dnnl_ext)
# AMX + AVX512F + AVX512BF16 + AVX512VNNI
define_extension_target(
diff --git a/cmake/external_projects/deepgemm.cmake b/cmake/external_projects/deepgemm.cmake
index c3a48a64fc77..07328c271388 100644
--- a/cmake/external_projects/deepgemm.cmake
+++ b/cmake/external_projects/deepgemm.cmake
@@ -20,7 +20,7 @@ else()
FetchContent_Declare(
deepgemm
GIT_REPOSITORY https://github.com/deepseek-ai/DeepGEMM.git
- GIT_TAG 477618cd51baffca09c4b0b87e97c03fe827ef03
+ GIT_TAG 891d57b4db1071624b5c8fa0d1e51cb317fa709f
GIT_SUBMODULES "third-party/cutlass" "third-party/fmt"
GIT_PROGRESS TRUE
CONFIGURE_COMMAND ""
@@ -76,7 +76,6 @@ if(DEEPGEMM_ARCHS)
"${deepgemm_SOURCE_DIR}/third-party/fmt/include")
target_compile_options(_deep_gemm_C PRIVATE
- $<$:-std=c++17>
$<$:-O3>
$<$:-Wno-psabi>
$<$:-Wno-deprecated-declarations>)
@@ -120,6 +119,11 @@ if(DEEPGEMM_ARCHS)
COMPONENT _deep_gemm_C
FILES_MATCHING PATTERN "*.py")
+ install(DIRECTORY "${deepgemm_SOURCE_DIR}/deep_gemm/mega/"
+ DESTINATION vllm/third_party/deep_gemm/mega
+ COMPONENT _deep_gemm_C
+ FILES_MATCHING PATTERN "*.py")
+
# Generate envs.py (normally generated by DeepGEMM's setup.py build step)
file(WRITE "${CMAKE_CURRENT_BINARY_DIR}/deep_gemm_envs.py"
"# Pre-installed environment variables\npersistent_envs = dict()\n")
diff --git a/cmake/external_projects/flashmla.cmake b/cmake/external_projects/flashmla.cmake
index 0f16b9161fa3..65986df55012 100644
--- a/cmake/external_projects/flashmla.cmake
+++ b/cmake/external_projects/flashmla.cmake
@@ -19,7 +19,7 @@ else()
FetchContent_Declare(
flashmla
GIT_REPOSITORY https://github.com/vllm-project/FlashMLA
- GIT_TAG 692917b1cda61b93ac9ee2d846ec54e75afe87b1
+ GIT_TAG a6ec2ba7bd0a7dff98b3f4d3e6b52b159c48d78b
GIT_PROGRESS TRUE
CONFIGURE_COMMAND ""
BUILD_COMMAND ""
diff --git a/csrc/activation_kernels.cu b/csrc/activation_kernels.cu
index 758a77795553..303433392c32 100644
--- a/csrc/activation_kernels.cu
+++ b/csrc/activation_kernels.cu
@@ -11,29 +11,74 @@
namespace vllm {
template
+ bool act_first, bool HAS_CLAMP>
__device__ __forceinline__ scalar_t compute(const scalar_t& x,
- const scalar_t& y) {
- return act_first ? ACT_FN(x) * y : x * ACT_FN(y);
+ const scalar_t& y,
+ const float limit) {
+ if constexpr (act_first) {
+ scalar_t gate = x;
+ scalar_t up = y;
+ if constexpr (HAS_CLAMP) {
+ gate = (scalar_t)fminf((float)gate, limit);
+ up = (scalar_t)fmaxf(fminf((float)up, limit), -limit);
+ }
+ return ACT_FN(gate) * up;
+ } else {
+ scalar_t gate = x;
+ scalar_t up = y;
+ if constexpr (HAS_CLAMP) {
+ gate = (scalar_t)fmaxf(fminf((float)gate, limit), -limit);
+ up = (scalar_t)fminf((float)up, limit);
+ }
+ return gate * ACT_FN(up);
+ }
}
template
+ bool act_first, bool HAS_CLAMP>
__device__ __forceinline__ packed_t packed_compute(const packed_t& x,
- const packed_t& y) {
- return act_first ? packed_mul(PACKED_ACT_FN(x), y)
- : packed_mul(x, PACKED_ACT_FN(y));
+ const packed_t& y,
+ const float limit) {
+ if constexpr (act_first) {
+ packed_t gate = x;
+ packed_t up = y;
+ if constexpr (HAS_CLAMP) {
+ float2 g = cast_to_float2(gate);
+ float2 u = cast_to_float2(up);
+ g.x = fminf(g.x, limit);
+ g.y = fminf(g.y, limit);
+ u.x = fmaxf(fminf(u.x, limit), -limit);
+ u.y = fmaxf(fminf(u.y, limit), -limit);
+ gate = cast_to_packed(g);
+ up = cast_to_packed(u);
+ }
+ return packed_mul(PACKED_ACT_FN(gate), up);
+ } else {
+ packed_t gate = x;
+ packed_t up = y;
+ if constexpr (HAS_CLAMP) {
+ float2 g = cast_to_float2(gate);
+ float2 u = cast_to_float2(up);
+ g.x = fmaxf(fminf(g.x, limit), -limit);
+ g.y = fmaxf(fminf(g.y, limit), -limit);
+ u.x = fminf(u.x, limit);
+ u.y = fminf(u.y, limit);
+ gate = cast_to_packed(g);
+ up = cast_to_packed(u);
+ }
+ return packed_mul(gate, PACKED_ACT_FN(up));
+ }
}
// Activation and gating kernel template.
template
+ bool use_vec, bool HAS_CLAMP, bool use_256b = false>
__global__ void act_and_mul_kernel(
scalar_t* __restrict__ out, // [..., d]
const scalar_t* __restrict__ input, // [..., 2, d]
- const int d) {
+ const int d, const float limit) {
const scalar_t* x_ptr = input + blockIdx.x * 2 * d;
const scalar_t* y_ptr = x_ptr + d;
scalar_t* out_ptr = out + blockIdx.x * d;
@@ -58,8 +103,9 @@ __global__ void act_and_mul_kernel(
}
#pragma unroll
for (int j = 0; j < pvec_t::NUM_ELTS; j++) {
- x.elts[j] = packed_compute(
- x.elts[j], y.elts[j]);
+ x.elts[j] =
+ packed_compute(
+ x.elts[j], y.elts[j], limit);
}
if constexpr (use_256b) {
st256(x, &out_vec[i]);
@@ -72,7 +118,8 @@ __global__ void act_and_mul_kernel(
for (int64_t idx = threadIdx.x; idx < d; idx += blockDim.x) {
const scalar_t x = VLLM_LDG(&x_ptr[idx]);
const scalar_t y = VLLM_LDG(&y_ptr[idx]);
- out_ptr[idx] = compute(x, y);
+ out_ptr[idx] =
+ compute(x, y, limit);
}
}
}
@@ -151,8 +198,11 @@ packed_gelu_tanh_kernel(const packed_t& val) {
// Launch activation and gating kernel.
// Use ACT_FIRST (bool) indicating whether to apply the activation function
-// first.
-#define LAUNCH_ACTIVATION_GATE_KERNEL(KERNEL, PACKED_KERNEL, ACT_FIRST) \
+// first. HAS_CLAMP (bool) enables pre-activation clamping: gate input is
+// clamped (max only) and up input is clamped (both sides) before the
+// activation function is applied.
+#define LAUNCH_ACTIVATION_GATE_KERNEL(KERNEL, PACKED_KERNEL, ACT_FIRST, \
+ HAS_CLAMP, LIMIT) \
auto dtype = input.scalar_type(); \
int d = input.size(-1) / 2; \
int64_t num_tokens = input.numel() / input.size(-1); \
@@ -177,8 +227,8 @@ packed_gelu_tanh_kernel(const packed_t& val) {
scalar_t, typename vllm::PackedTypeConverter::Type, \
KERNEL, \
PACKED_KERNEL::Type>, \
- ACT_FIRST, true, true><<>>( \
- out.data_ptr(), input.data_ptr(), d); \
+ ACT_FIRST, true, HAS_CLAMP, true><<>>( \
+ out.data_ptr(), input.data_ptr(), d, LIMIT); \
}); \
} else { \
VLLM_DISPATCH_FLOATING_TYPES(dtype, "act_and_mul_kernel", [&] { \
@@ -186,8 +236,8 @@ packed_gelu_tanh_kernel(const packed_t& val) {
scalar_t, typename vllm::PackedTypeConverter::Type, \
KERNEL, \
PACKED_KERNEL::Type>, \
- ACT_FIRST, true, false><<>>( \
- out.data_ptr(), input.data_ptr(), d); \
+ ACT_FIRST, true, HAS_CLAMP, false><<>>( \
+ out.data_ptr(), input.data_ptr(), d, LIMIT); \
}); \
} \
} else { \
@@ -197,8 +247,8 @@ packed_gelu_tanh_kernel(const packed_t& val) {
scalar_t, typename vllm::PackedTypeConverter::Type, \
KERNEL, \
PACKED_KERNEL::Type>, \
- ACT_FIRST, false><<>>( \
- out.data_ptr(), input.data_ptr(), d); \
+ ACT_FIRST, false, HAS_CLAMP><<>>( \
+ out.data_ptr(), input.data_ptr(), d, LIMIT); \
}); \
}
@@ -206,7 +256,14 @@ void silu_and_mul(torch::Tensor& out, // [..., d]
torch::Tensor& input) // [..., 2 * d]
{
LAUNCH_ACTIVATION_GATE_KERNEL(vllm::silu_kernel, vllm::packed_silu_kernel,
- true);
+ true, false, 0.0f);
+}
+
+void silu_and_mul_clamp(torch::Tensor& out, // [..., d]
+ torch::Tensor& input, // [..., 2 * d]
+ double limit) {
+ LAUNCH_ACTIVATION_GATE_KERNEL(vllm::silu_kernel, vllm::packed_silu_kernel,
+ true, true, (float)limit);
}
void mul_and_silu(torch::Tensor& out, // [..., d]
@@ -215,21 +272,21 @@ void mul_and_silu(torch::Tensor& out, // [..., d]
// The difference between mul_and_silu and silu_and_mul is that mul_and_silu
// applies the silu to the latter half of the input.
LAUNCH_ACTIVATION_GATE_KERNEL(vllm::silu_kernel, vllm::packed_silu_kernel,
- false);
+ false, false, 0.0f);
}
void gelu_and_mul(torch::Tensor& out, // [..., d]
torch::Tensor& input) // [..., 2 * d]
{
LAUNCH_ACTIVATION_GATE_KERNEL(vllm::gelu_kernel, vllm::packed_gelu_kernel,
- true);
+ true, false, 0.0f);
}
void gelu_tanh_and_mul(torch::Tensor& out, // [..., d]
torch::Tensor& input) // [..., 2 * d]
{
- LAUNCH_ACTIVATION_GATE_KERNEL(vllm::gelu_tanh_kernel,
- vllm::packed_gelu_tanh_kernel, true);
+ LAUNCH_ACTIVATION_GATE_KERNEL(
+ vllm::gelu_tanh_kernel, vllm::packed_gelu_tanh_kernel, true, false, 0.0f);
}
namespace vllm {
diff --git a/csrc/cache.h b/csrc/cache.h
index 821d5e719a44..a9e74b0dc2df 100644
--- a/csrc/cache.h
+++ b/csrc/cache.h
@@ -12,7 +12,8 @@ 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);
+ const torch::Tensor& sizes,
+ bool is_src_access_order_any);
void reshape_and_cache(torch::Tensor& key, torch::Tensor& value,
torch::Tensor& key_cache, torch::Tensor& value_cache,
diff --git a/csrc/cache_kernels.cu b/csrc/cache_kernels.cu
index d1cfbaeb05d2..9130dd2ccae7 100644
--- a/csrc/cache_kernels.cu
+++ b/csrc/cache_kernels.cu
@@ -77,7 +77,8 @@ 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) {
+ const torch::Tensor& sizes,
+ bool is_src_access_order_any) {
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");
@@ -97,13 +98,13 @@ void swap_blocks_batch(const torch::Tensor& src_ptrs,
const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
- // 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));
+ // 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.
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).
@@ -124,7 +125,12 @@ void swap_blocks_batch(const torch::Tensor& src_ptrs,
if (batch_fn != nullptr) {
CUmemcpyAttributes attr = {};
- attr.srcAccessOrder = CU_MEMCPY_SRC_ACCESS_ORDER_STREAM;
+ // 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;
size_t attrs_idx = 0;
size_t fail_idx = 0;
CUresult result = batch_fn(reinterpret_cast(dst_data),
@@ -134,12 +140,30 @@ void swap_blocks_batch(const torch::Tensor& src_ptrs,
&fail_idx, static_cast(stream));
TORCH_CHECK(result == CUDA_SUCCESS, "cuMemcpyBatchAsync failed at index ",
fail_idx, " with error ", result);
- } else
+ 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(dst_data), reinterpret_cast(src_data),
+ reinterpret_cast(size_data), static_cast(n), &attr,
+ &attrs_idx, 0, &fail_idx, static_cast(stream));
+ TORCH_CHECK(result == hipSuccess, "hipMemcpyBatchAsync failed at index ",
+ fail_idx, " with error ", result);
+ return;
+ }
#endif
{
- // Fallback for CUDA < 12.8, older drivers, and ROCm:
- // individual async copies.
- // cudaMemcpyDefault lets the driver infer direction from pointer types.
+ // Fallback for CUDA < 12.8, older CUDA drivers, and ROCm < 7.1:
+ // individual async copies. cudaMemcpyDefault lets the driver infer
+ // direction from pointer types.
for (int64_t i = 0; i < n; i++) {
cudaMemcpyAsync(reinterpret_cast(dst_data[i]),
reinterpret_cast(src_data[i]),
@@ -599,6 +623,11 @@ __global__ void cp_gather_indexer_k_quant_cache_kernel(
const int head_idx = (blockIdx.y * blockDim.x + threadIdx.x) * VEC_SIZE;
// Find batch index within a block
__shared__ int batch_idx[BLOCK_Y_SIZE];
+ if (threadIdx.x == 0) {
+ batch_idx[threadIdx.y] = -1;
+ }
+ __syncthreads();
+
for (int iter = 0; iter < cuda_utils::ceil_div(batch_size, int(blockDim.x));
iter++) {
int tid = iter * blockDim.x + threadIdx.x;
@@ -611,16 +640,18 @@ __global__ void cp_gather_indexer_k_quant_cache_kernel(
}
}
-#ifndef USE_ROCM
- __syncwarp();
-#endif
+ __syncthreads();
- if (head_idx >= head_dim || token_idx >= num_tokens) {
+ // num_tokens may be an allocation upper bound when Python avoids a D2H sync.
+ // Only tokens covered by the exact device-side cu_seq_lens are valid to
+ // gather.
+ const int batch = batch_idx[threadIdx.y];
+ if (head_idx >= head_dim || token_idx >= num_tokens || batch < 0) {
return;
}
- const int inbatch_seq_idx = token_idx - cu_seq_lens[batch_idx[threadIdx.y]];
- const int block_idx = block_table[batch_idx[threadIdx.y] * num_blocks +
- inbatch_seq_idx / cache_block_size];
+ const int inbatch_seq_idx = token_idx - cu_seq_lens[batch];
+ const int block_idx =
+ block_table[batch * num_blocks + inbatch_seq_idx / cache_block_size];
const int64_t src_block_offset = block_idx * block_stride;
const int64_t cache_inblock_offset =
(inbatch_seq_idx % cache_block_size) * head_dim + head_idx;
@@ -724,6 +755,28 @@ void reshape_and_cache_flash(
int num_tokens = slot_mapping.size(0);
int num_heads = key.size(1);
int head_size = key.size(2);
+
+ const at::cuda::OptionalCUDAGuard device_guard(device_of(key));
+ const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
+
+ if (kv_cache_dtype == "nvfp4") {
+#if defined(ENABLE_NVFP4_SM100) || defined(ENABLE_NVFP4_SM120)
+ // NVFP4 dispatch is compiled separately for SM100+.
+ extern void reshape_and_cache_nvfp4_dispatch(
+ torch::Tensor & key, torch::Tensor & value, torch::Tensor & key_cache,
+ torch::Tensor & value_cache, torch::Tensor & slot_mapping,
+ torch::Tensor & k_scale, torch::Tensor & v_scale);
+ reshape_and_cache_nvfp4_dispatch(key, value, key_cache, value_cache,
+ slot_mapping, k_scale, v_scale);
+ return;
+#else
+ TORCH_CHECK(false,
+ "NVFP4 KV cache requires SM100+ (Blackwell). "
+ "Please rebuild vllm with a Blackwell-compatible CUDA target.");
+#endif
+ }
+
+ // Original FP8/auto path.
int block_size = key_cache.size(1);
int64_t key_stride = key.stride(0);
@@ -741,8 +794,6 @@ void reshape_and_cache_flash(
dim3 grid(num_tokens);
dim3 block(std::min(num_heads * head_size, 512));
- const at::cuda::OptionalCUDAGuard device_guard(device_of(key));
- const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
DISPATCH_BY_KV_CACHE_DTYPE(key.dtype(), kv_cache_dtype,
CALL_RESHAPE_AND_CACHE_FLASH);
@@ -1470,6 +1521,9 @@ void concat_mla_q(torch::Tensor& ql_nope, // [num_tokens, num_heads, nope_dim]
TORCH_CHECK(ql_nope.stride(2) == 1, "ql_nope must have stride 1 in dim 2");
TORCH_CHECK(q_pe.stride(2) == 1, "q_pe must have stride 1 in dim 2");
TORCH_CHECK(q_out.stride(2) == 1, "q_out must have stride 1 in dim 2");
+ TORCH_CHECK(ql_nope.scalar_type() == at::ScalarType::Half ||
+ ql_nope.scalar_type() == at::ScalarType::BFloat16,
+ "ql_nope must be float16 or bfloat16 dtype");
if (num_tokens == 0) return;
@@ -1481,7 +1535,7 @@ void concat_mla_q(torch::Tensor& ql_nope, // [num_tokens, num_heads, nope_dim]
const at::cuda::OptionalCUDAGuard device_guard(device_of(ql_nope));
const cudaStream_t stream = at::cuda::getCurrentCUDAStream();
- VLLM_DISPATCH_FLOATING_TYPES(ql_nope.scalar_type(), "concat_mla_q", [&] {
+ VLLM_DISPATCH_HALF_TYPES(ql_nope.scalar_type(), "concat_mla_q", [&] {
vllm::ConcatMLAQKernel<<>>(
q_out.data_ptr(), ql_nope.data_ptr(),
q_pe.data_ptr(), num_tokens, num_heads, q_out.stride(0),
diff --git a/csrc/cpu/activation_lut_bf16.cpp b/csrc/cpu/activation_lut_bf16.cpp
new file mode 100644
index 000000000000..0ff2567e1ee8
--- /dev/null
+++ b/csrc/cpu/activation_lut_bf16.cpp
@@ -0,0 +1,71 @@
+#include "cpu_types.hpp"
+
+#include
+#include
+#include
+#include
+
+#include
+#include
+#include
+
+constexpr uint32_t ActivationLutSize = 1u << 16;
+
+at::Tensor gelu_reference(const at::Tensor& x) { return at::gelu(x, "none"); }
+
+void maybe_init_activation_lut_bf16(
+ uint16_t* lut, std::once_flag& once,
+ at::Tensor (*activation)(const at::Tensor&)) {
+ std::call_once(once, [&]() {
+ auto lut_input =
+ at::empty({static_cast(ActivationLutSize)},
+ at::TensorOptions().device(at::kCPU).dtype(at::kFloat));
+ auto* lut_input_ptr = lut_input.data_ptr();
+#pragma omp parallel for
+ for (uint32_t i = 0; i < ActivationLutSize; ++i) {
+ lut_input_ptr[i] = c10::detail::f32_from_bits(static_cast(i));
+ }
+
+ auto lut_output = activation(lut_input);
+ const auto* lut_output_ptr = lut_output.data_ptr();
+#pragma omp parallel for
+ for (uint32_t i = 0; i < ActivationLutSize; ++i) {
+ lut[i] = c10::detail::round_to_nearest_even(lut_output_ptr[i]);
+ }
+ });
+}
+
+void activation_lut_bf16(torch::Tensor& out, torch::Tensor& input,
+ const uint16_t* lut, const char* op_name) {
+ TORCH_CHECK(input.scalar_type() == at::kBFloat16, op_name,
+ ": input must be bfloat16");
+ TORCH_CHECK(out.scalar_type() == at::kBFloat16, op_name,
+ ": out must be bfloat16");
+ TORCH_CHECK(input.is_contiguous(), op_name, ": input must be contiguous");
+ TORCH_CHECK(out.is_contiguous(), op_name, ": out must be contiguous");
+
+ const auto* src =
+ reinterpret_cast(input.data_ptr());
+ auto* dst = reinterpret_cast(out.data_ptr());
+ const int64_t n = input.numel();
+
+ CPU_KERNEL_GUARD_IN(activation_lut_bf16_impl)
+#pragma omp parallel for
+ for (int64_t i = 0; i < n; ++i) {
+ dst[i] = lut[src[i]];
+ }
+ CPU_KERNEL_GUARD_OUT(activation_lut_bf16_impl)
+}
+
+void activation_lut_bf16(torch::Tensor& out, torch::Tensor& input,
+ const std::string& activation) {
+ if (activation == "gelu") {
+ static std::array lut{};
+ static std::once_flag once;
+ maybe_init_activation_lut_bf16(lut.data(), once, gelu_reference);
+ activation_lut_bf16(out, input, lut.data(), "gelu_lut");
+ return;
+ }
+
+ TORCH_CHECK(false, "Unsupported activation: ", activation);
+}
diff --git a/csrc/cpu/cpu_arch_macros.h b/csrc/cpu/cpu_arch_macros.h
index c73b62ecdec9..9be45a3efcef 100644
--- a/csrc/cpu/cpu_arch_macros.h
+++ b/csrc/cpu/cpu_arch_macros.h
@@ -61,8 +61,23 @@
#endif
#ifdef __aarch64__
- // Implementation copied from Arm Optimized Routines (expf AdvSIMD)
+ // Implementation of neon_expf copied from Arm Optimized Routines (expf
+ // AdvSIMD)
// https://github.com/ARM-software/optimized-routines/blob/master/math/aarch64/advsimd/expf.c
+ //
+ // Additional fast exponential intended for cases where outputs will be
+ // downcasted to FP16 / BF16 (e.g. attention softmax). Accurate within 1 ULP
+ // for FP16 Accurate within 1 ULP for BF16 for inputs in [-87.683, 88.376] &
+ // clamps inputs outside this range to 0 / inf. Implementation is similar to
+ // exp_u20, but:
+ // - uses a third degree polynomial approximation for exp(r) instead of a
+ // fifth degree one, with coefficients re-tuned.
+ // - does not split natural log (ln) into high / low parts
+ // - clamps exp(x) to 0 for x < -87.683113f and inf for x > 88.3762589f
+ // exp(x) = 2^n (exp(r))
+ // r = x - n*ln2, with n = round(x/ln2)
+ // exp(r) ~ poly(r) = 1 + r + r^2 * (c3 + c2 * r)
+ // n = round(x / ln2), r = x - n*ln2
#include
#define DEFINE_FAST_EXP \
const float32x4_t inv_ln2 = vdupq_n_f32(0x1.715476p+0f); \
@@ -106,7 +121,38 @@
result.val[2] = neon_expf(vec.reg.val[2]); \
result.val[3] = neon_expf(vec.reg.val[3]); \
return vec_op::FP32Vec16(result); \
- };
+ }; \
+ const float32x4_t lower_bound = vdupq_n_f32(-0x1.5ebb82p+6f); \
+ const float32x4_t upper_bound = vdupq_n_f32(0x1.61814ap+6f); \
+ constexpr float ln2 = 0x1.62e43p-1f; \
+ constexpr float f_c2 = 0x1.5592ecp-3f; \
+ const float32x4_t f_c3 = vdupq_n_f32(0x1.017d34p-1f); \
+ auto neon_expf_f16 = [&](float32x4_t values) __attribute__(( \
+ always_inline)) { \
+ const uint32x4_t lt_lower = vcltq_f32(values, lower_bound); \
+ const uint32x4_t gt_upper = vcgtq_f32(values, upper_bound); \
+ float32x4_t n = vrndaq_f32(vmulq_f32(values, inv_ln2)); \
+ float32x4_t r = vfmsq_n_f32(values, n, ln2); \
+ uint32x4_t e = vshlq_n_u32(vreinterpretq_u32_s32(vcvtq_s32_f32(n)), 23); \
+ float32x4_t r2 = vmulq_f32(r, r); \
+ float32x4_t q = vfmaq_n_f32(f_c3, r, f_c2); \
+ float32x4_t s = vaddq_f32(vdupq_n_f32(1.0f), r); \
+ float32x4_t p = vfmaq_f32(s, q, r2); \
+ float32x4_t y = \
+ vreinterpretq_f32_u32(vaddq_u32(vreinterpretq_u32_f32(p), e)); \
+ y = vbslq_f32(lt_lower, vdupq_n_f32(0.0f), y); \
+ y = vbslq_f32(gt_upper, vdupq_n_f32(INFINITY), y); \
+ return y; \
+ }; \
+ auto fast_exp_f16 = [&](const vec_op::FP32Vec16& vec) \
+ __attribute__((always_inline)) { \
+ float32x4x4_t result; \
+ result.val[0] = neon_expf_f16(vec.reg.val[0]); \
+ result.val[1] = neon_expf_f16(vec.reg.val[1]); \
+ result.val[2] = neon_expf_f16(vec.reg.val[2]); \
+ result.val[3] = neon_expf_f16(vec.reg.val[3]); \
+ return vec_op::FP32Vec16(result); \
+ };
#endif // __aarch64__
diff --git a/csrc/cpu/cpu_attn.cpp b/csrc/cpu/cpu_attn.cpp
index a582b4b4d7cc..4750dd78838d 100644
--- a/csrc/cpu/cpu_attn.cpp
+++ b/csrc/cpu/cpu_attn.cpp
@@ -1,5 +1,16 @@
#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,
@@ -18,6 +29,8 @@ torch::Tensor get_scheduler_metadata(
isa = cpu_attention::ISA::NEON;
} else if (isa_hint == "vxe") {
isa = cpu_attention::ISA::VXE;
+ } else if (isa_hint == "vsx") {
+ isa = cpu_attention::ISA::VSX;
} else {
TORCH_CHECK(false, "Unsupported CPU attention ISA hint: " + isa_hint);
}
@@ -49,7 +62,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, [&]() {
+ CPU_ATTN_DISPATCH(head_dim, isa, 0, [&]() {
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);
@@ -72,7 +85,9 @@ 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 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") {
TORCH_CHECK_EQ(key.dim(), 3);
TORCH_CHECK_EQ(value.dim(), 3);
TORCH_CHECK_EQ(key_cache.dim(), 4);
@@ -80,18 +95,30 @@ 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(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(k_scale) : 0.0f;
+ const float v_inv = is_fp8 ? 1.0f / static_cast(v_scale) : 0.0f;
+
const int64_t token_num = key.size(0);
- 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 head_num = key.size(1);
+ const int64_t head_dim = key.size(2);
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") {
@@ -104,21 +131,31 @@ void cpu_attn_reshape_and_cache(
return cpu_attention::ISA::NEON;
} else if (isa == "vxe") {
return cpu_attention::ISA::VXE;
+ } 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, [&]() {
+ CPU_ATTN_DISPATCH(head_dim, isa_tag, kv_cache_idx, [&]() {
+ using kv_t = typename attn_impl::kv_cache_t;
attn_impl::reshape_and_cache(
key.data_ptr(), value.data_ptr(),
- key_cache.data_ptr(), value_cache.data_ptr(),
- slot_mapping.data_ptr(), 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);
+ reinterpret_cast(key_cache.data_ptr()),
+ reinterpret_cast(value_cache.data_ptr()),
+ slot_mapping.data_ptr(), 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);
});
});
}
@@ -137,13 +174,26 @@ 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& s_aux // [num_heads]
-) {
+ const std::optional& s_aux, // [num_heads]
+ const double k_scale = 1.0, const double v_scale = 1.0,
+ const std::string& kv_cache_dtype = "auto") {
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(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(
scheduler_metadata.data_ptr());
@@ -165,25 +215,32 @@ void cpu_attention_with_kv_cache(
input.block_table = block_table.data_ptr();
input.alibi_slopes =
alibi_slopes.has_value() ? alibi_slopes->data_ptr() : nullptr;
- // For now sink must be bf16
input.s_aux = s_aux.has_value() ? s_aux->data_ptr() : 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;
}
- float softcap_fp32 = softcap;
- input.softcap = softcap_fp32;
+ input.softcap = static_cast(softcap);
+
+ if (is_fp8) {
+ input.k_scale_fp8 = static_cast(k_scale);
+ input.v_scale_fp8 = static_cast(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");
+ }
VLLM_DISPATCH_FLOATING_TYPES(
query.scalar_type(), "cpu_attention_with_kv_cache", [&]() {
- CPU_ATTN_DISPATCH(query.size(2), input.metadata->isa, [&]() {
- TORCH_CHECK_EQ(input.block_size % attn_impl::BlockSizeAlignment, 0);
- cpu_attention::AttentionMainLoop mainloop;
- mainloop(&input);
- });
+ 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 mainloop;
+ mainloop(&input);
+ });
});
}
diff --git a/csrc/cpu/cpu_attn_amx.hpp b/csrc/cpu/cpu_attn_amx.hpp
index 1c8644d52329..6a0341085dce 100644
--- a/csrc/cpu/cpu_attn_amx.hpp
+++ b/csrc/cpu/cpu_attn_amx.hpp
@@ -1,6 +1,7 @@
#ifndef CPU_ATTN_AMX_HPP
#define CPU_ATTN_AMX_HPP
+#include "cpu_attn_fp8.hpp"
#include "cpu_attn_impl.hpp"
namespace cpu_attention {
@@ -21,9 +22,10 @@ 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
-template
+// 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
class TileGemm224 {
public:
template
@@ -42,13 +44,56 @@ class TileGemm224 {
}
};
-template <>
-class TileGemm224 {
+// Dequantize one FP8 tile (AMX_TILE_ROW_NUM rows x 32 cols) to BF16.
+template
+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) {
+ 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
+FORCE_INLINE const c10::BFloat16* prepare_b_tile(const kv_cache_t* src,
+ c10::BFloat16* scratch) {
+ if constexpr (std::is_same_v ||
+ std::is_same_v) {
+ deq_tile_amx(reinterpret_cast(src), scratch);
+ return scratch;
+ } else {
+ return reinterpret_cast(src);
+ }
+}
+
+// Handles both BF16 and FP8 KV cache (2-2-4 pattern).
+template
+class TileGemm224 {
+ static_assert(std::is_same_v ||
+ std::is_same_v ||
+ std::is_same_v,
+ "kv_cache_t must be BFloat16, Float8_e4m3fn, or Float8_e5m2");
+
+ static constexpr bool fp8_kv =
+ std::is_same_v ||
+ std::is_same_v;
+
+ 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;
+
public:
template
FORCE_INLINE static void gemm(const int32_t m_size,
c10::BFloat16* __restrict__ a_tile,
- c10::BFloat16* __restrict__ b_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,
@@ -56,6 +101,7 @@ class TileGemm224 {
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 = [&]() {
@@ -70,8 +116,8 @@ class TileGemm224 {
}
}();
- c10::BFloat16* __restrict__ b_tile_2 = b_tile;
- c10::BFloat16* __restrict__ b_tile_3 = [&]() {
+ kv_cache_t* __restrict__ b_tile_2 = b_tile;
+ kv_cache_t* __restrict__ b_tile_3 = [&]() {
if constexpr (phase == AttentionGemmPhase::QK) {
// k_cache is prepacked
return b_tile + (k_size * AMX_TILE_ROW_BYTES / 4);
@@ -106,11 +152,16 @@ class TileGemm224 {
_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, b_tile_2, b_tile_stride);
+ _tile_stream_loadd(2, const_cast(load_2), b_tile_stride);
_tile_dpbf16ps(4, 0, 2);
- _tile_stream_loadd(3, b_tile_3, b_tile_stride);
+ _tile_stream_loadd(3, const_cast(load_3), b_tile_stride);
_tile_dpbf16ps(5, 0, 3);
_tile_loadd(1, a_tile_1, a_tile_stride);
_tile_dpbf16ps(6, 1, 2);
@@ -154,13 +205,13 @@ class TileGemm224 {
};
// 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, (6, 7): store results C matrix, row num should be
-// m
-template
+// 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
class TileGemm122 {
public:
template
@@ -179,13 +230,26 @@ class TileGemm122 {
}
};
-template <>
-class TileGemm122 {
+// Handles both BF16 and FP8 KV cache (1-2-2 pattern).
+template
+class TileGemm122 {
+ static_assert(std::is_same_v ||
+ std::is_same_v ||
+ std::is_same_v,
+ "kv_cache_t must be BFloat16, Float8_e4m3fn, or Float8_e5m2");
+
+ static constexpr bool fp8_kv =
+ std::is_same_v ||
+ std::is_same_v;
+
+ static constexpr int64_t tile_elems = AMX_TILE_BYTES / sizeof(c10::BFloat16);
+ static constexpr int64_t scratch_elems = fp8_kv ? tile_elems : 1;
+
public:
template
FORCE_INLINE static void gemm(const int32_t m_size,
c10::BFloat16* __restrict__ a_tile,
- c10::BFloat16* __restrict__ b_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,
@@ -215,21 +279,19 @@ class TileGemm122 {
}
}();
- c10::BFloat16* __restrict__ b_tile_2 = b_tile;
- c10::BFloat16* __restrict__ b_tile_3 = [&]() {
+ kv_cache_t* __restrict__ b_tile_2 = b_tile;
+ kv_cache_t* __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");
}
}();
- c10::BFloat16* __restrict__ b_tile_4 =
+ kv_cache_t* __restrict__ b_tile_4 =
b_tile_2 + AMX_TILE_BYTES / sizeof(c10::BFloat16);
- c10::BFloat16* __restrict__ b_tile_5 =
+ kv_cache_t* __restrict__ b_tile_5 =
b_tile_3 + AMX_TILE_BYTES / sizeof(c10::BFloat16);
int64_t b_stride = AMX_TILE_ROW_BYTES;
@@ -250,16 +312,25 @@ class TileGemm122 {
_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, b_tile_2, b_stride);
+ _tile_stream_loadd(2, const_cast(load_2), b_stride);
_tile_dpbf16ps(6, 0, 2);
- _tile_stream_loadd(3, b_tile_3, b_stride);
+ _tile_stream_loadd(3, const_cast(load_3), b_stride);
_tile_dpbf16ps(7, 0, 3);
_tile_loadd(1, a_tile_1, a_tile_stride);
- _tile_stream_loadd(4, b_tile_4, b_stride);
+ _tile_stream_loadd(4, const_cast(load_4), b_stride);
_tile_dpbf16ps(6, 1, 4);
- _tile_stream_loadd(5, b_tile_5, b_stride);
+ _tile_stream_loadd(5, const_cast(load_5), b_stride);
_tile_dpbf16ps(7, 1, 5);
// update ptrs
@@ -279,10 +350,13 @@ class TileGemm122 {
}
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, b_tile_2, b_stride);
+ _tile_stream_loadd(2, const_cast(load_2), b_stride);
_tile_dpbf16ps(6, 0, 2);
- _tile_stream_loadd(3, b_tile_3, b_stride);
+ _tile_stream_loadd(3, const_cast(load_3), b_stride);
_tile_dpbf16ps(7, 0, 3);
}
@@ -302,21 +376,25 @@ class TileGemm122 {
_tile_loadconfig(&config);
}
};
+
} // namespace
-template
-class AttentionImpl {
+template
+class AttentionImpl {
+ static constexpr bool fp8_kv =
+ std::is_same_v ||
+ std::is_same_v;
+
public:
using query_t = scalar_t;
using q_buffer_t = scalar_t;
- using kv_cache_t = scalar_t;
+ using kv_cache_t = kv_cache_scalar_t;
using logits_buffer_t = float;
using partial_output_buffer_t = float;
using prob_buffer_t = scalar_t;
constexpr static int64_t BlockSizeAlignment =
- AMX_TILE_ROW_BYTES /
- sizeof(kv_cache_t); // KV token num unit of QK and PV phases
+ 32; // AMX_TILE_ROW_NUM = 16 tokens/tile; 32 = 2 tiles
constexpr static int64_t HeadDimAlignment =
2 * (AMX_TILE_ROW_BYTES / 4); // headdim num unit of PV phase
constexpr static int64_t MaxQHeadNumPerIteration = 32;
@@ -324,6 +402,9 @@ class AttentionImpl {
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
@@ -332,21 +413,50 @@ class AttentionImpl {
~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 ? 0x1p112f : 0x1p120f;
+ return v_scale * bias;
+ }
+ return 1.0f;
+ }
+
template 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 ? 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::init_tile_config(q_head_num, amx_tile_config_);
+ TileGemm224::init_tile_config(q_head_num,
+ amx_tile_config_);
}
- attention> attention_iteration;
+ attention> attention_iteration;
attention_iteration(CPU_ATTENTION_PARAMS);
} else {
if (q_head_num != current_q_head_num_) {
current_q_head_num_ = q_head_num;
- TileGemm122::init_tile_config(q_head_num, amx_tile_config_);
+ TileGemm122::init_tile_config(q_head_num,
+ amx_tile_config_);
}
- attention> attention_iteration;
+ attention> attention_iteration;
attention_iteration(CPU_ATTENTION_PARAMS);
}
}
@@ -411,13 +521,26 @@ class AttentionImpl {
// reshape KV to AMX friendly layout
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,
+ kv_cache_t* __restrict__ key_cache, kv_cache_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 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();
+ reshape_and_cache_fp8_amx_impl(
+ key, value, reinterpret_cast(key_cache),
+ reinterpret_cast(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;
+ }
+
// 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
@@ -426,6 +549,9 @@ class AttentionImpl {
// 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(key_cache);
+ scalar_t* __restrict__ vc = reinterpret_cast(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) {
@@ -453,8 +579,7 @@ class AttentionImpl {
constexpr int64_t quadword_num_per_group =
token_num_per_group * quadword_num;
int32_t* key_cache_start_ptr =
- reinterpret_cast(key_cache +
- block_idx * num_blocks_stride +
+ reinterpret_cast(kc + block_idx * num_blocks_stride +
head_idx * cache_head_num_stride) +
group_idx * quadword_num_per_group + group_offset;
@@ -483,7 +608,7 @@ class AttentionImpl {
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 +
+ vc + 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;
diff --git a/csrc/cpu/cpu_attn_fp8.hpp b/csrc/cpu/cpu_attn_fp8.hpp
new file mode 100644
index 000000000000..764b6ed7f84a
--- /dev/null
+++ b/csrc/cpu/cpu_attn_fp8.hpp
@@ -0,0 +1,214 @@
+// SPDX-License-Identifier: Apache-2.0
+// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
+#pragma once
+#include
+#include
+#include
+#include
+#include
+
+#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::quiet_NaN();
+ uint32_t b_u32 = static_cast(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(&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(&v_f32_unscaled);
+ uint8_t sign = static_cast((bits >> 24) & 0x80);
+ uint8_t payload = static_cast((bits >> 20) & 0x7F);
+ if (payload == 0) return sign;
+ payload = std::min(payload, 0x7E); // keep 0x7F as NaN encoding
+ return static_cast(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
+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(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(ksrc[j * 2]), k_inv);
+ uint8_t fp8_1 = quant_fn(static_cast(ksrc[j * 2 + 1]), k_inv);
+ uint8_t bytes[2] = {fp8_0, fp8_1};
+ uint16_t hw = *reinterpret_cast(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(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::quiet_NaN();
+ const uint32_t sign = static_cast(b & 0x80) << 24;
+ if (exp_bits == 0x1F)
+ return sign ? -std::numeric_limits::infinity()
+ : std::numeric_limits::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(exp_bits) - 15 + 127) << 23) |
+ (static_cast(mant_bits) << 21);
+ float val = *reinterpret_cast(&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(&v);
+ const uint8_t sign = static_cast((bits >> 24) & 0x80);
+ const int32_t exp_fp32 = static_cast((bits >> 23) & 0xFF) - 127;
+ const uint8_t mant2 = static_cast((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(std::min(m, 3u));
+ }
+ const uint8_t exp5 = static_cast(exp_fp32 + 15);
+ return sign | (exp5 << 2) | mant2;
+}
+
+// ---------------------------------------------------------------------------
+// Select the FP8 quant function at compile time based on kv_cache_t.
+// ---------------------------------------------------------------------------
+template
+constexpr auto select_fp8_quant_fn() {
+ if constexpr (std::is_same_v)
+ 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
+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(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(vsrc[i]), v_inv);
+ }
+ }
+}
diff --git a/csrc/cpu/cpu_attn_impl.hpp b/csrc/cpu/cpu_attn_impl.hpp
index c15799fa950d..b9987fb26c19 100644
--- a/csrc/cpu/cpu_attn_impl.hpp
+++ b/csrc/cpu/cpu_attn_impl.hpp
@@ -12,10 +12,24 @@
#include "cpu/utils.hpp"
namespace cpu_attention {
-enum class ISA { AMX, VEC, VEC16, NEON, VXE };
+enum class ISA { AMX, VEC, VEC16, NEON, VXE, VSX };
-template
-class AttentionImpl {};
+// Mirrors csrc/attention/dtype_fp8.cuh Fp8KVCacheDataType exactly.
+enum class Fp8KVCacheDataType {
+ kAuto = 0,
+ kFp8E4M3 = 1,
+ kFp8E5M2 = 2,
+};
+
+struct AttentionInput;
+
+template
+class AttentionImpl {
+ public:
+ void init_from_input(const AttentionInput*) {}
+ float get_output_v_scale() const noexcept { return 1.0f; }
+};
struct AttentionWorkItemGroup {
int32_t req_id;
@@ -147,6 +161,12 @@ struct AttentionMetadata {
case ISA::NEON:
ss << "NEON, ";
break;
+ case ISA::VXE:
+ ss << "VXE, ";
+ break;
+ case ISA::VSX:
+ ss << "VSX, ";
+ break;
}
ss << "workitem_group_num: " << workitem_group_num
<< ", reduction_item_num: " << reduction_item_num
@@ -777,6 +797,9 @@ 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 \
@@ -1149,7 +1172,11 @@ class AttentionMainLoop {
bool use_sink) {
#ifdef DEFINE_FAST_EXP
DEFINE_FAST_EXP
+ bool constexpr IsReducedPrecision =
+ std::is_same_v ||
+ std::is_same_v;
#endif
+
using prob_buffer_vec_t = typename VecTypeTrait::vec_t;
static_assert(sizeof(prob_buffer_t) <= sizeof(logits_buffer_t));
@@ -1198,8 +1225,17 @@ class AttentionMainLoop {
vec = vec - max_vec;
// compute exp
-#ifdef DEFINE_FAST_EXP
- vec = fast_exp(vec);
+
+#if defined(DEFINE_FAST_EXP)
+ #ifdef __aarch64__
+ if constexpr (IsReducedPrecision) {
+ vec = fast_exp_f16(vec);
+ } else
+ #endif
+ {
+ vec = fast_exp(vec);
+ }
+
prob_buffer_vec_t output_vec(vec);
output_vec.save(curr_prob_buffer_iter);
#else
@@ -1255,7 +1291,11 @@ class AttentionMainLoop {
int32_t kv_tile_token_num, float softcap_scale) {
#ifdef DEFINE_FAST_EXP
DEFINE_FAST_EXP
+ bool constexpr IsReducedPrecision =
+ std::is_same_v ||
+ std::is_same_v;
#endif
+
float inv_softcap_scale = 1.0 / softcap_scale;
vec_op::FP32Vec16 softcap_scale_vec(softcap_scale);
vec_op::FP32Vec16 inv_softcap_scale_vec(inv_softcap_scale);
@@ -1269,8 +1309,15 @@ class AttentionMainLoop {
vec_op::FP32Vec16 vec(curr_logits_buffer_iter);
vec = vec * inv_softcap_scale_vec;
-#ifdef DEFINE_FAST_EXP
- vec = fast_exp(vec);
+#if defined(DEFINE_FAST_EXP)
+ #ifdef __aarch64__
+ if constexpr (IsReducedPrecision) {
+ vec = fast_exp_f16(vec);
+ } else
+ #endif
+ {
+ vec = fast_exp(vec);
+ }
vec_op::FP32Vec16 inv_vec = ones_vec / vec;
vec = (vec - inv_vec) / (vec + inv_vec);
#else
@@ -1347,6 +1394,13 @@ class AttentionMainLoop {
}
attention_impl_t attn_impl;
+ constexpr bool fp8_kv = std::is_same_v ||
+ std::is_same_v;
+ 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;
@@ -1726,7 +1780,7 @@ class AttentionMainLoop {
reinterpret_cast(input->output) +
output_buffer_offset,
sum_buffer, actual_q_heads_per_kv,
- actual_q_token_num, q_head_num);
+ actual_q_token_num, q_head_num, output_v_scale);
} else {
const int32_t stride =
actual_q_heads_per_kv * split_kv_q_token_num_threshold;
@@ -1796,7 +1850,7 @@ class AttentionMainLoop {
split_output_buffer,
reinterpret_cast(input->output) + output_buffer_offset,
split_sum_buffer, actual_q_heads_per_kv, curr_output_token_num,
- q_head_num);
+ q_head_num, output_v_scale);
}
}
}
@@ -1920,8 +1974,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 int32_t actual_q_token_num, const int32_t q_head_num,
+ const float v_scale = 1.0f) {
// final output
using output_vec_t = typename VecTypeTrait::vec_t;
@@ -1935,7 +1989,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(1.0 / *curr_sum_buffer);
+ vec_op::FP32Vec16 inv_sum_scale_vec(v_scale / *curr_sum_buffer);
for (int32_t i = 0; i < group_num_per_head; ++i) {
vec_op::FP32Vec16 vec(curr_partial_output_buffer_iter);
diff --git a/csrc/cpu/cpu_attn_neon.hpp b/csrc/cpu/cpu_attn_neon.hpp
index 3523893c38c5..db4c5df2e88d 100644
--- a/csrc/cpu/cpu_attn_neon.hpp
+++ b/csrc/cpu/cpu_attn_neon.hpp
@@ -248,8 +248,8 @@ class TileGemmNeonFMLA {
} // namespace
// this is similar to "ISA::VEC" at the moment
-template
-class AttentionImpl {
+template
+class AttentionImpl {
public:
using query_t = scalar_t;
using q_buffer_t = float;
@@ -343,7 +343,8 @@ class AttentionImpl {
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 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) {
@@ -388,7 +389,7 @@ class AttentionImpl {
#ifdef ARM_BF16_SUPPORT
// For BF16 on Arm, reuse the BFMMLA kernels with 32-token alignment.
template
-class AttentionImpl
+class AttentionImpl
: public AttentionImplNEONBFMMLA {};
#endif
diff --git a/csrc/cpu/cpu_attn_neon_bfmmla.hpp b/csrc/cpu/cpu_attn_neon_bfmmla.hpp
index fb133aa13098..4e4578a74f5b 100644
--- a/csrc/cpu/cpu_attn_neon_bfmmla.hpp
+++ b/csrc/cpu/cpu_attn_neon_bfmmla.hpp
@@ -602,7 +602,8 @@ 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) {
+ [[maybe_unused]] const int64_t block_size_stride,
+ const float /*k_inv*/ = 0.0f, const float /*v_inv*/ = 0.0f) {
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;
diff --git a/csrc/cpu/cpu_attn_vec.hpp b/csrc/cpu/cpu_attn_vec.hpp
index f51a232ba955..c3983e0578a5 100644
--- a/csrc/cpu/cpu_attn_vec.hpp
+++ b/csrc/cpu/cpu_attn_vec.hpp
@@ -1,11 +1,37 @@
#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
+FORCE_INLINE std::pair load_b_pair_vec(
+ const kv_cache_t* ptr) {
+ if constexpr (std::is_same_v) {
+ // BF16 container, but values are in the FP16 exponent range (bias 15 not
+ // 127).
+ vec_op::BF16Vec32 bf16_b_reg(reinterpret_cast(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) {
+ vec_op::BF16Vec32 bf16_b_reg(reinterpret_cast(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::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
class TileGemm82 {
@@ -54,10 +80,7 @@ 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::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;
@@ -76,16 +99,14 @@ 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) {
- 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);
+ auto [fp32_b_0_reg, fp32_b_1_reg] = load_b_pair_vec(curr_b);
float* __restrict__ curr_m_a = curr_a;
vec_op::unroll_loop([&](int32_t i) {
- float v = *curr_m_a;
- vec_op::FP32Vec16 a_reg(v);
+ vec_op::FP32Vec16 a_reg(*curr_m_a);
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;
@@ -95,8 +116,7 @@ class TileGemm82 {
// update
curr_a += 1;
- curr_b_0 += ldb;
- curr_b_1 += ldb;
+ curr_b += ldb;
}
vec_op::unroll_loop([&](int32_t i) {
@@ -109,15 +129,20 @@ class TileGemm82 {
});
}
};
+
} // namespace
// This is a general but naive implementation based on vector instructions
-template
-class AttentionImpl {
+template
+class AttentionImpl {
+ static constexpr bool fp8_kv =
+ std::is_same_v ||
+ std::is_same_v