Compare commits

..
Author SHA1 Message Date
Lucas WilkinsonandOpenAI Codex 79b560e29a Fix DeepSeek V4 SM10 sparse MLA default
Co-authored-by: OpenAI Codex <codex@openai.com>

Signed-off-by: Lucas Wilkinson <lwilkins@redhat.com>
2026-07-05 01:59:07 +00:00
731 changed files with 7784 additions and 33277 deletions
-1
View File
@@ -8,7 +8,6 @@ run_all_patterns:
- "docker/docker-bake-rocm.hcl"
- ".buildkite/hardware_tests/amd.yaml"
- ".buildkite/scripts/ci-bake-rocm.sh"
- ".buildkite/scripts/rocm/"
- ".buildkite/scripts/hardware_ci/run-amd-test.py"
- ".buildkite/scripts/hardware_ci/run-amd-test.sh"
- "CMakeLists.txt"
+29 -32
View File
@@ -1,28 +1,5 @@
group: Hardware - AMD Build
# ROCm image flow:
# 1. Refresh the long-lived ROCm base image only when Dockerfile.rocm_base changes.
# 2. Build ci_base from either the stable base or the freshly refreshed base.
# 3. Build the per-commit ROCm CI image and smoke-test it before GPU jobs run.
steps:
- label: "AMD: :docker: refresh ROCm base"
key: refresh-rocm-base-amd
depends_on: []
device: amd_cpu
no_plugin: true
commands:
- bash .buildkite/scripts/rocm/refresh-base-image.sh
env:
DOCKER_BUILDKIT: "1"
BUILDKIT_PROGRESS: "tty"
TERM: "xterm-256color"
retry:
automatic:
- exit_status: -1 # Agent was lost
limit: 1
- exit_status: -10 # Agent was lost
limit: 1
# Ensure ci_base is up-to-date before building the test image.
# Compares a content hash of ci_base-affecting files against the remote
# image label. If hashes match the build is skipped (< 30 s); if they
@@ -30,16 +7,13 @@ steps:
- label: "AMD: :docker: ensure ci_base"
key: ensure-ci-base-amd
soft_fail: false
depends_on:
- refresh-rocm-base-amd
depends_on: []
device: amd_cpu
no_plugin: true
commands:
- bash .buildkite/scripts/rocm/build-ci-base.sh
- bash .buildkite/scripts/ci-bake-rocm.sh ci-base-rocm-ci-with-deps
env:
DOCKER_BUILDKIT: "1"
BUILDKIT_PROGRESS: "tty"
TERM: "xterm-256color"
VLLM_BAKE_FILE: "docker/docker-bake-rocm.hcl"
PYTORCH_ROCM_ARCH: "gfx90a;gfx942;gfx950"
REMOTE_VLLM: "1"
@@ -59,12 +33,35 @@ steps:
device: amd_cpu
no_plugin: true
commands:
- bash .buildkite/scripts/rocm/build-test-image.sh
- bash .buildkite/scripts/rocm/smoke-test-image.sh
- |
if [[ "${ROCM_CI_ARTIFACT_ONLY:-0}" == "1" ]]; then
echo "ROCM_CI_ARTIFACT_ONLY=1; building ROCm wheel artifact only"
IMAGE_TAG="" bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-artifacts
else
bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-wheel
fi
- |
docker run --rm --network=none --entrypoint /bin/bash "rocm/vllm-ci:${BUILDKITE_COMMIT}" -ec '
if [ ! -d /vllm-workspace ]; then echo Missing directory: /vllm-workspace >&2; exit 1; fi
if [ ! -d /vllm-workspace/tests ]; then echo Missing directory: /vllm-workspace/tests >&2; exit 1; fi
if [ ! -d /vllm-workspace/src/vllm ]; then echo Missing directory: /vllm-workspace/src/vllm >&2; exit 1; fi
if [ ! -x /vllm-workspace/src/vllm/vllm-rs ]; then echo Missing executable: /vllm-workspace/src/vllm/vllm-rs >&2; exit 1; fi
command -v python3
command -v uv
command -v pytest
if ! command -v amd-smi >/dev/null 2>&1 && ! command -v rocminfo >/dev/null 2>&1; then
echo No ROCm CLI found in image >&2
exit 1
fi
python3 - <<PY
import torch, vllm
print(torch.__version__)
print(vllm.__version__)
PY
echo AMD image smoke OK
'
env:
DOCKER_BUILDKIT: "1"
BUILDKIT_PROGRESS: "tty"
TERM: "xterm-256color"
VLLM_BAKE_FILE: "docker/docker-bake-rocm.hcl"
PYTORCH_ROCM_ARCH: "gfx90a;gfx942;gfx950"
IMAGE_TAG: "rocm/vllm-ci:$BUILDKITE_COMMIT"
-2
View File
@@ -17,14 +17,12 @@ steps:
- tests/kernels/test_awq_int4_to_int8.py
- tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
- tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
- tests/kernels/mamba/test_cpu_short_conv.py
commands:
- |
bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 30m "
pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
pytest -x -v -s tests/kernels/moe/test_cpu_fused_moe.py
pytest -x -v -s tests/kernels/moe/test_cpu_quant_fused_moe.py
pytest -x -v -s tests/kernels/mamba/test_cpu_short_conv.py
pytest -x -v -s tests/kernels/test_onednn.py
pytest -x -v -s tests/kernels/test_awq_int4_to_int8.py
pytest -x -v -s tests/kernels/quantization/test_cpu_fp8_scaled_mm.py
@@ -45,7 +45,7 @@ steps:
agent_tags:
label: production
gpu: 1+
mem: 24+
mem: 16+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
-20
View File
@@ -79,18 +79,12 @@ setup_buildx_builder() {
docker buildx ls | grep -E '^\*|^NAME' || docker buildx ls
}
annotate_image_tags() {
.buildkite/scripts/annotate-image-build.sh \
"${IMAGE_TAG:-}" "${IMAGE_TAG_LATEST:-}"
}
check_and_skip_if_image_exists() {
if [[ -n "${IMAGE_TAG:-}" ]]; then
echo "--- :mag: Checking if image exists"
if docker manifest inspect "${IMAGE_TAG}" >/dev/null 2>&1; then
echo "Image already exists: ${IMAGE_TAG}"
echo "Skipping build"
annotate_image_tags
exit 0
fi
echo "Image not found, proceeding with build"
@@ -177,18 +171,6 @@ BRANCH=$4
IMAGE_TAG=$5
IMAGE_TAG_LATEST=${6:-} # only used for main branch, optional
# When TORCH_NIGHTLY=1, build the base CI image against PyTorch nightly so the
# entire existing pipeline runs on nightly torch (CUDA/GPU lane only). Delegate
# to the dedicated nightly build (PYTORCH_NIGHTLY=1, CUDA 13.0) and tag it at the
# normal IMAGE_TAG that every test step already pulls -- no separate image tag,
# no duplicate "vLLM Against PyTorch Nightly" pipeline section.
if [[ "${TORCH_NIGHTLY:-0}" == "1" ]]; then
echo "--- :warning: TORCH_NIGHTLY=1 -- building base image on PyTorch nightly"
SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
exec "${SCRIPT_DIR}/image_build_torch_nightly.sh" \
"${REGISTRY}" "${REPO}" "${BUILDKITE_COMMIT}" "${BRANCH}" "${IMAGE_TAG}"
fi
# build config
TARGET="test-ci"
VLLM_BAKE_FILE_PATH="${VLLM_BAKE_FILE_PATH:-docker/docker-bake.hcl}"
@@ -272,5 +254,3 @@ echo "--- :docker: Building ${TARGET}"
docker --debug buildx bake -f "${VLLM_BAKE_FILE_PATH}" -f "${CI_HCL_PATH}" --progress plain "${TARGET}"
echo "--- :white_check_mark: Build complete"
annotate_image_tags
+19 -20
View File
@@ -9,31 +9,30 @@ fi
REGISTRY=$1
REPO=$2
BUILDKITE_COMMIT=$3
IMAGE="$REGISTRY/$REPO:$BUILDKITE_COMMIT-arm64"
# authenticate with AWS ECR
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
# skip build if image already exists
if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
echo "Image found"
else
if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64) ]]; then
echo "Image not found, proceeding with build..."
# build for arm64 GPU targets: Grace/GH200 (sm_90),
# Blackwell/Thor (sm_100/sm_103/sm_110), and DGX Spark/GB10
# (sm_121, family-covered by 12.0 under CUDA 13)
docker build --file docker/Dockerfile \
--platform linux/arm64 \
--build-arg max_jobs=16 \
--build-arg nvcc_threads=4 \
--build-arg torch_cuda_arch_list="9.0 10.0 11.0 12.0" \
--build-arg USE_SCCACHE=1 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$IMAGE" \
--target test \
--progress plain .
# push
docker push "$IMAGE"
else
echo "Image found"
exit 0
fi
.buildkite/scripts/annotate-image-build.sh "$IMAGE"
# build for arm64 GPU targets: Grace/GH200 (sm_90) and DGX Spark/GB10
# (sm_121, family-covered by 12.0 under CUDA 13)
docker build --file docker/Dockerfile \
--platform linux/arm64 \
--build-arg max_jobs=16 \
--build-arg nvcc_threads=4 \
--build-arg torch_cuda_arch_list="9.0 12.0" \
--build-arg USE_SCCACHE=1 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64 \
--target test \
--progress plain .
# push
docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64
+15 -15
View File
@@ -9,26 +9,26 @@ fi
REGISTRY=$1
REPO=$2
BUILDKITE_COMMIT=$3
IMAGE="$REGISTRY/$REPO:$BUILDKITE_COMMIT-cpu"
# authenticate with AWS ECR
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
# skip build if image already exists
if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
echo "Image found"
else
if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-cpu) ]]; then
echo "Image not found, proceeding with build..."
# build
docker build --file docker/Dockerfile.cpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--build-arg VLLM_CPU_X86=true \
--tag "$IMAGE" \
--target vllm-test \
--progress plain .
# push
docker push "$IMAGE"
else
echo "Image found"
exit 0
fi
.buildkite/scripts/annotate-image-build.sh "$IMAGE"
# build
docker build --file docker/Dockerfile.cpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--build-arg VLLM_CPU_X86=true \
--tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-cpu \
--target vllm-test \
--progress plain .
# push
docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-cpu
+14 -14
View File
@@ -9,25 +9,25 @@ fi
REGISTRY=$1
REPO=$2
BUILDKITE_COMMIT=$3
IMAGE="$REGISTRY/$REPO:$BUILDKITE_COMMIT-arm64-cpu"
# authenticate with AWS ECR
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
# skip build if image already exists
if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
echo "Image found"
else
if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64-cpu) ]]; then
echo "Image not found, proceeding with build..."
# build
docker build --file docker/Dockerfile.cpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$IMAGE" \
--target vllm-test \
--progress plain .
# push
docker push "$IMAGE"
else
echo "Image found"
exit 0
fi
.buildkite/scripts/annotate-image-build.sh "$IMAGE"
# build
docker build --file docker/Dockerfile.cpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64-cpu \
--target vllm-test \
--progress plain .
# push
docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-arm64-cpu
+15 -15
View File
@@ -9,26 +9,26 @@ fi
REGISTRY=$1
REPO=$2
BUILDKITE_COMMIT=$3
IMAGE="$REGISTRY/$REPO:$BUILDKITE_COMMIT-hpu"
# authenticate with AWS ECR
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
# skip build if image already exists
if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
echo "Image found"
else
if [[ -z $(docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-hpu) ]]; then
echo "Image not found, proceeding with build..."
# build
docker build \
--file tests/pytorch_ci_hud_benchmark/Dockerfile.hpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$IMAGE" \
--progress plain \
https://github.com/vllm-project/vllm-gaudi.git
# push
docker push "$IMAGE"
else
echo "Image found"
exit 0
fi
.buildkite/scripts/annotate-image-build.sh "$IMAGE"
# build
docker build \
--file tests/pytorch_ci_hud_benchmark/Dockerfile.hpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-hpu \
--progress plain \
https://github.com/vllm-project/vllm-gaudi.git
# push
docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-hpu
@@ -40,7 +40,6 @@ docker buildx ls
echo "--- :mag: Checking if image already exists"
if docker manifest inspect "$IMAGE_TAG" >/dev/null 2>&1; then
echo "Image found: $IMAGE_TAG — skipping build"
.buildkite/scripts/annotate-image-build.sh "$IMAGE_TAG"
exit 0
fi
echo "Image not found, proceeding with build..."
@@ -67,5 +66,3 @@ docker buildx build --file docker/Dockerfile \
--progress plain .
echo "--- :white_check_mark: Torch nightly image build complete: $IMAGE_TAG"
.buildkite/scripts/annotate-image-build.sh "$IMAGE_TAG"
+14 -14
View File
@@ -9,26 +9,26 @@ fi
REGISTRY=$1
REPO=$2
BUILDKITE_COMMIT=$3
IMAGE="$REGISTRY/$REPO:$BUILDKITE_COMMIT-xpu"
# authenticate with AWS ECR
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin "$REGISTRY" || true
aws ecr get-login-password --region us-east-1 | docker login --username AWS --password-stdin 936637512419.dkr.ecr.us-east-1.amazonaws.com || true
# skip build if image already exists
if docker manifest inspect "$IMAGE" &> /dev/null; then
echo "Image found"
else
if ! docker manifest inspect "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu &> /dev/null; then
echo "Image not found, proceeding with build..."
# build
docker build \
--file docker/Dockerfile.xpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$IMAGE" \
--progress plain .
# push
docker push "$IMAGE"
else
echo "Image found"
exit 0
fi
.buildkite/scripts/annotate-image-build.sh "$IMAGE"
# build
docker build \
--file docker/Dockerfile.xpu \
--build-arg max_jobs=16 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu \
--progress plain .
# push
docker push "$REGISTRY"/"$REPO":"$BUILDKITE_COMMIT"-xpu
+5 -27
View File
@@ -38,7 +38,7 @@ steps:
agent_tags:
label: production
gpu: 1+
mem: 24+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -72,31 +72,9 @@ steps:
pytest -v -s v1/test_oracle.py &&
pytest -v -s v1/test_request.py &&
pytest -v -s v1/test_outputs.py &&
pytest -v -s v1/sample'
- label: Basic Models Tests (Initialization)
timeout_in_minutes: 60
device: intel_gpu
agent_tags:
label: production
gpu: 1+
mem: 16+
no_plugin: true
working_dir: "."
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
REPO: "vllm-ci-test-repo"
VLLM_TEST_DEVICE: "xpu"
source_file_dependencies:
- vllm/
- tests/models/test_initialization.py
- tests/models/registry.py
commands:
- >-
bash .buildkite/scripts/hardware_ci/run-intel-test.sh
'export VLLM_XPU_FUSED_MOE_USE_REF=1 &&
cd tests &&
pytest -v -s models/test_initialization.py::test_can_initialize_large_subset[Eagle3MiniMaxM2ForCausalLM]'
pytest -v -s v1/sample/test_topk_topp_sampler.py &&
pytest -v -s v1/sample/test_logprobs.py &&
pytest -v -s v1/sample/test_logprobs_e2e.py'
- label: XPU CPU Offload
timeout_in_minutes: 60
@@ -132,7 +110,7 @@ steps:
agent_tags:
label: production
gpu: 2+
mem: 16+
mem: 24+
no_plugin: true
working_dir: "."
env:
@@ -9,7 +9,7 @@ steps:
agent_tags:
label: production
gpu: 2+
mem: 16+
mem: 24+
no_plugin: true
working_dir: "."
env:
@@ -9,7 +9,7 @@ steps:
agent_tags:
label: production
gpu: 1+
mem: 24+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -81,7 +81,7 @@ steps:
agent_tags:
label: production
gpu: 1+
mem: 24+
mem: 16+
no_plugin: true
working_dir: "."
env:
@@ -105,7 +105,7 @@ steps:
agent_tags:
label: production
gpu: 1+
mem: 24+
mem: 16+
no_plugin: true
working_dir: "."
env:
+1 -1
View File
@@ -81,7 +81,7 @@ steps:
agent_tags:
label: production
gpu: 1+
mem: 24+
mem: 16+
no_plugin: true
env:
REGISTRY: "public.ecr.aws/q9t5s3a7"
-47
View File
@@ -350,25 +350,6 @@ steps:
- "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"
- 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-cu129-ubuntu2404"'
- label: ":docker: Build release image - x86_64 - XPU"
depends_on: ~
id: build-xpu-release-image
agents:
queue: cpu_queue_release
commands:
- "aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7"
- |
DOCKER_BUILDKIT=1 docker build \
$(bash .buildkite/scripts/docker-build-metadata-args.sh xpu) \
--build-arg GIT_REPO_CHECK=1 \
--target vllm-openai \
--progress plain \
-f docker/Dockerfile.xpu .
- "docker push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-xpu"
- 'bash .buildkite/scripts/annotate-build-artifact.sh "$$BUILDKITE_LABEL" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-$(uname -m)-xpu"'
env:
DOCKER_BUILDKIT: "1"
- block: "Build release image for x86_64 CPU"
key: block-cpu-release-image-build
depends_on: ~
@@ -464,16 +445,6 @@ steps:
- "docker manifest push public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"
- 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: CUDA 12.9 Ubuntu 24.04" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-cu129-ubuntu2404"'
- label: "Create manifest - XPU"
depends_on:
- build-xpu-release-image
id: create-manifest-xpu
agents:
queue: small_cpu_queue_release
commands:
- "bash .buildkite/scripts/xpu/create-xpu-ecr-manifest.sh"
- 'bash .buildkite/scripts/annotate-build-artifact.sh "Manifest: XPU" "public.ecr.aws/q9t5s3a7/vllm-release-repo:$BUILDKITE_COMMIT-xpu"'
- label: "Publish nightly multi-arch image to DockerHub"
depends_on:
- create-multi-arch-manifest
@@ -848,23 +819,6 @@ steps:
DOCKER_BUILDKIT: "1"
S3_BUCKET: "vllm-wheels"
- label: "Publish nightly XPU image to DockerHub"
depends_on:
- create-manifest-xpu
if: build.env("NIGHTLY") == "1"
agents:
queue: small_cpu_queue_release
commands:
- "bash .buildkite/scripts/xpu/push-nightly-builds-xpu.sh"
- "bash .buildkite/scripts/cleanup-nightly-builds.sh nightly- vllm/vllm-openai-xpu"
plugins:
- docker-login#v3.0.0:
username: vllmbot
password-env: DOCKERHUB_TOKEN
env:
DOCKER_BUILDKIT: "1"
DOCKERHUB_USERNAME: "vllmbot"
- label: "Publish nightly ROCm image to DockerHub"
depends_on:
- build-rocm-release-image
@@ -895,7 +849,6 @@ steps:
- create-multi-arch-manifest-cuda-12-9
- create-multi-arch-manifest-ubuntu2404
- create-multi-arch-manifest-cuda-12-9-ubuntu2404
- create-manifest-xpu
- build-rocm-release-image
- input-release-version
# Wait for CPU builds if their block steps were unblocked, so publish
@@ -1,36 +0,0 @@
#!/bin/bash
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
#
# Append the Docker image tag(s) an image-build step pushed to a Buildkite
# annotation, so the built image tags show up on the build page instead of
# being buried in the job logs.
#
# Usage: annotate-image-build.sh <image_tag> [<image_tag> ...]
set -euo pipefail
# buildkite-agent only exists on Buildkite agents; no-op elsewhere so the
# image build scripts stay runnable locally.
if ! command -v buildkite-agent >/dev/null 2>&1; then
echo "buildkite-agent not found; skipping image tag annotation"
exit 0
fi
label="${BUILDKITE_LABEL:-Image build}"
content=""
for image in "$@"; do
[[ -n "$image" ]] || continue
content+="- **${label}**: \`${image}\`"$'\n'
done
if [[ -z "$content" ]]; then
echo "No image tags provided; nothing to annotate"
exit 0
fi
# Best-effort: a flaky annotation must never fail an otherwise successful
# (and expensive) image build.
if ! printf '%s' "$content" | \
buildkite-agent annotate --append --style 'info' --context 'docker-images'; then
echo "warning: failed to annotate build with image tags"
fi
@@ -29,11 +29,7 @@ if python3 -c "import torch; assert torch.version.hip" 2>/dev/null; then
TORCH_INDEX_URL=""
fi
else
if [ "${TORCH_NIGHTLY:-0}" = "1" ]; then
TORCH_INDEX_URL="https://download.pytorch.org/whl/nightly/cu130"
else
TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
fi
TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
fi
echo ">>> Using PyTorch index: ${TORCH_INDEX_URL:-PyPI default}"
+3 -25
View File
@@ -15,7 +15,7 @@ set -euo pipefail
DEFAULT_REPO_SLUG="vllm-project/vllm"
DEFAULT_CI_HCL_SOURCE="docker/ci-rocm.hcl"
DEFAULT_CI_BASE_CONTENT_FILES="requirements/common.txt requirements/rocm.txt requirements/test/rocm.txt docker/Dockerfile.rocm_base docker/ci-rocm.hcl docker/docker-bake-rocm.hcl tools/install_torchcodec_rocm.sh tests/vllm_test_utils .buildkite/scripts/ci-bake-rocm.sh .buildkite/scripts/rocm/build-ci-base.sh"
DEFAULT_CI_BASE_CONTENT_FILES="requirements/common.txt requirements/rocm.txt requirements/test/rocm.txt docker/Dockerfile.rocm_base docker/ci-rocm.hcl docker/docker-bake-rocm.hcl tools/install_torchcodec_rocm.sh tests/vllm_test_utils .buildkite/scripts/ci-bake-rocm.sh"
DEFAULT_CI_BASE_DOCKERFILE="docker/Dockerfile.rocm"
DEFAULT_CI_BASE_DOCKERFILE_STAGES="base build_rixl build_rocshmem build_deepep mori_base ci_base"
DEFAULT_CI_BASE_METADATA_VERSION="1"
@@ -393,16 +393,6 @@ should_upload_wheel_artifacts() {
|| "${TARGET}" == *"artifact"* ]]
}
set_buildkite_metadata() {
local key="$1"
local value="$2"
[[ -n "${value}" ]] || return 0
if command -v buildkite-agent >/dev/null 2>&1; then
buildkite-agent meta-data set "${key}" "${value}" || true
fi
}
get_remote_image_label() {
local image_ref="$1"
local label_key="$2"
@@ -743,8 +733,6 @@ configure_ci_base_image_refs() {
if is_ci_base_target; then
IMAGE_TAG="${primary_tag}"
CI_BASE_IMAGE="${primary_tag}"
export CI_BASE_IMAGE
export IMAGE_TAG
echo "ci_base primary image tag: ${CI_BASE_IMAGE_TAG}"
@@ -762,10 +750,6 @@ configure_ci_base_image_refs() {
echo "ci_base stable alias will not be pushed for this build"
echo "Set NIGHTLY=1 on ${CI_BASE_STABLE_BRANCH:-main} to refresh ${stable_tag}"
fi
set_buildkite_metadata "rocm-ci-base-image" "${CI_BASE_IMAGE_TAG}"
set_buildkite_metadata "rocm-ci-base-image-content" "${content_tag}"
set_buildkite_metadata "rocm-ci-base-image-commit" "${CI_BASE_IMAGE_TAG_COMMIT:-}"
set_buildkite_metadata "rocm-ci-base-image-stable" "${CI_BASE_IMAGE_TAG_STABLE:-}"
return 0
fi
@@ -1795,10 +1779,7 @@ seed_dependency_caches_if_needed() {
echo "--- :docker: Seeding ${target}"
echo "Expected cache ref: ${cache_ref}"
docker buildx bake \
"${BAKE_FILES[@]}" \
--progress "${BUILDKIT_PROGRESS:-plain}" \
"${target}"
docker buildx bake "${BAKE_FILES[@]}" --progress plain "${target}"
verify_dependency_cache_ref "${cache_ref}"
done
}
@@ -1826,10 +1807,7 @@ run_bake() {
local build_rc=0
echo "--- :docker: Building ${TARGET}"
docker buildx bake \
"${BAKE_FILES[@]}" \
--progress "${BUILDKIT_PROGRESS:-plain}" \
"${BAKE_TARGETS[@]}" || build_rc=$?
docker buildx bake "${BAKE_FILES[@]}" --progress plain "${BAKE_TARGETS[@]}" || build_rc=$?
if [[ ${build_rc} -eq 0 ]]; then
echo "--- :white_check_mark: Build complete"
@@ -28,17 +28,6 @@
###############################################################################
set -o pipefail
: "${BUILDKIT_PROGRESS:=plain}"
: "${TERM:=xterm-256color}"
: "${FORCE_COLOR:=1}"
: "${CLICOLOR_FORCE:=1}"
: "${PY_COLORS:=1}"
: "${ROCM_DOCKER_TTY:=1}"
if [[ " ${PYTEST_ADDOPTS:-} " != *" --color"* ]]; then
PYTEST_ADDOPTS="${PYTEST_ADDOPTS:+${PYTEST_ADDOPTS} }--color=yes"
fi
export BUILDKIT_PROGRESS TERM FORCE_COLOR CLICOLOR_FORCE PY_COLORS PYTEST_ADDOPTS ROCM_DOCKER_TTY
# Export Python path for commands that run directly on the host. Containerized
# tests set this to /vllm-workspace below so spawned Python processes do not
# depend on their current working directory.
@@ -160,7 +149,6 @@ EOF
echo "--- Building local ROCm test image"
docker build \
--pull=false \
--progress "${BUILDKIT_PROGRESS}" \
--build-arg "BASE_IMAGE=${base_image}" \
-t "${artifact_image}" \
"${context_dir}" || return 1
@@ -547,13 +535,6 @@ if is_multi_node "$commands"; then
else
echo "--- Single-node job"
echo "Render devices: $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES"
docker_run_terminal_args=(-i)
if [[ "${ROCM_DOCKER_TTY}" == "1" ]]; then
docker_run_terminal_args+=(-t)
echo "Docker interactive stdin: enabled; TTY allocation: enabled"
else
echo "Docker interactive stdin: enabled; TTY allocation: disabled"
fi
ulimit_core_hard=$(ulimit -H -c)
if [[ "$ulimit_core_hard" == "unlimited" ]]; then
@@ -570,7 +551,6 @@ else
fi
docker run \
"${docker_run_terminal_args[@]}" \
--device /dev/kfd $BUILDKITE_AGENT_META_DATA_RENDER_DEVICES \
$RDMA_FLAGS \
--network=host \
@@ -585,11 +565,6 @@ else
-e AWS_SECRET_ACCESS_KEY \
-e BUILDKITE_PARALLEL_JOB \
-e BUILDKITE_PARALLEL_JOB_COUNT \
-e TERM \
-e FORCE_COLOR \
-e CLICOLOR_FORCE \
-e PY_COLORS \
-e PYTEST_ADDOPTS \
-v "${HF_CACHE}:${HF_MOUNT}" \
-e "HF_HOME=${HF_MOUNT}" \
-e "PYTHONPATH=${MYPYTHONPATH}" \
@@ -38,9 +38,7 @@ function cpu_tests() {
pytest -x -v -s tests/kernels/attention/test_cpu_attn.py
pytest -x -v -s tests/kernels/core/test_cpu_activation.py
pytest -x -v -s tests/kernels/moe/test_cpu_fused_moe.py
pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py
pytest -x -v -s tests/kernels/moe/test_cpu_int4_moe.py
pytest -x -v -s tests/kernels/mamba/test_cpu_short_conv.py"
pytest -x -v -s tests/kernels/mamba/cpu/test_cpu_gdn_ops.py"
# skip tests requiring model downloads if HF_TOKEN is not set
# due to rate-limits
@@ -64,6 +62,7 @@ function cpu_tests() {
set -e
pytest -x -v -s tests/quantization/test_compressed_tensors.py::test_compressed_tensors_w8a8_logprobs"
# basic online serving
docker exec cpu-test bash -c '
set -e
@@ -130,22 +130,6 @@ docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm
docker push vllm/vllm-openai-rocm:latest-base
docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
# ---- XPU ----
docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu
docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu vllm/vllm-openai-xpu:latest-x86_64
docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64
docker push vllm/vllm-openai-xpu:latest-x86_64
docker push vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64
docker manifest rm vllm/vllm-openai-xpu:latest || true
docker manifest rm vllm/vllm-openai-xpu:v${RELEASE_VERSION} || true
docker manifest create vllm/vllm-openai-xpu:latest vllm/vllm-openai-xpu:latest-x86_64 --amend
docker manifest create vllm/vllm-openai-xpu:v${RELEASE_VERSION} vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64 --amend
docker manifest push vllm/vllm-openai-xpu:latest
docker manifest push vllm/vllm-openai-xpu:v${RELEASE_VERSION}
# ---- CPU ----
# CPU images are behind separate block steps and may not have been built.
# All-or-nothing: inspect both arches first, then either publish everything
-32
View File
@@ -1,32 +0,0 @@
#!/usr/bin/env bash
# Build the ROCm ci_base image, optionally from a freshly rebuilt ROCm base.
set -euo pipefail
metadata_get() {
local key="$1"
if command -v buildkite-agent >/dev/null 2>&1; then
buildkite-agent meta-data get "${key}" 2>/dev/null || true
fi
}
main() {
local base_refreshed=""
base_refreshed="$(metadata_get rocm-base-refresh)"
if [[ "${base_refreshed}" == "1" ]]; then
export BASE_IMAGE
export CI_BASE_PUSH_STABLE_TAG
BASE_IMAGE="$(metadata_get rocm-base-image)"
CI_BASE_PUSH_STABLE_TAG="$(metadata_get rocm-base-push-stable-tag)"
CI_BASE_PUSH_STABLE_TAG="${CI_BASE_PUSH_STABLE_TAG:-0}"
echo "Using refreshed ROCm base image for ci_base: ${BASE_IMAGE}"
echo "Push stable ci_base tag: ${CI_BASE_PUSH_STABLE_TAG}"
fi
bash .buildkite/scripts/ci-bake-rocm.sh ci-base-rocm-ci-with-deps
}
main "$@"
@@ -1,57 +0,0 @@
#!/usr/bin/env bash
# Build the ROCm CI test image or wheel artifact.
#
# When Dockerfile.rocm_base changes, always build the full image so downstream
# ROCm tests can validate the freshly rebuilt base -> ci_base -> ci image chain.
set -euo pipefail
metadata_get() {
local key="$1"
if command -v buildkite-agent >/dev/null 2>&1; then
buildkite-agent meta-data get "${key}" 2>/dev/null || true
fi
}
use_refreshed_base_if_present() {
local base_refreshed=""
base_refreshed="$(metadata_get rocm-base-refresh)"
if [[ "${base_refreshed}" != "1" ]]; then
return 1
fi
export BASE_IMAGE
export CI_BASE_IMAGE
export IMAGE_TAG_LATEST
BASE_IMAGE="$(metadata_get rocm-base-image)"
CI_BASE_IMAGE="$(metadata_get rocm-ci-base-image)"
IMAGE_TAG_LATEST="$(metadata_get rocm-ci-image-descriptive)"
echo "Using refreshed ROCm base image for test image: ${BASE_IMAGE}"
echo "Using refreshed ROCm ci_base image for test image: ${CI_BASE_IMAGE}"
if [[ -n "${IMAGE_TAG_LATEST}" ]]; then
echo "Also tagging full ROCm CI image as: ${IMAGE_TAG_LATEST}"
fi
return 0
}
main() {
local base_refreshed=0
if use_refreshed_base_if_present; then
base_refreshed=1
fi
if [[ "${ROCM_CI_ARTIFACT_ONLY:-0}" == "1" && "${base_refreshed}" != "1" ]]; then
echo "ROCM_CI_ARTIFACT_ONLY=1; building ROCm wheel artifact only"
IMAGE_TAG="" bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-artifacts
return
fi
bash .buildkite/scripts/ci-bake-rocm.sh test-rocm-ci-with-wheel
}
main "$@"
@@ -1,513 +0,0 @@
#!/usr/bin/env bash
# Build and publish a fresh ROCm base image when Dockerfile.rocm_base changes.
#
# Normal AMD CI builds should not pay for this path. The script no-ops unless
# docker/Dockerfile.rocm_base changed relative to the branch base, the previous
# main commit, or ROCM_BASE_REFRESH_FORCE=1 is set.
set -euo pipefail
DOCKERFILE="${ROCM_BASE_DOCKERFILE:-docker/Dockerfile.rocm_base}"
BASE_REPO="${ROCM_BASE_IMAGE_REPO:-rocm/vllm-dev}"
CI_IMAGE_REPO="${ROCM_CI_IMAGE_REPO:-rocm/vllm-ci}"
BUILDER_NAME="${ROCM_BASE_BUILDER_NAME:-vllm-rocm-base-builder}"
DEFAULT_ROCM_BASE_METADATA_VERSION="1"
DEFAULT_ROCM_BASE_CONTENT_FILES="${DOCKERFILE}"
DEFAULT_ROCM_BASE_CONTENT_ARGS="BASE_IMAGE TRITON_BRANCH TRITON_REPO PYTORCH_BRANCH PYTORCH_REPO PYTORCH_VISION_BRANCH PYTORCH_VISION_REPO PYTORCH_AUDIO_BRANCH PYTORCH_AUDIO_REPO FA_BRANCH FA_REPO AITER_BRANCH AITER_REPO MORI_BRANCH MORI_REPO PYTORCH_ROCM_ARCH PYTHON_VERSION USE_SCCACHE"
metadata_set() {
local key="$1"
local value="$2"
[[ -n "${value}" ]] || return 0
if command -v buildkite-agent >/dev/null 2>&1; then
buildkite-agent meta-data set "${key}" "${value}" || true
fi
}
compute_content_hash() {
local path=""
local file=""
for path in "$@"; do
if [[ -d "${path}" ]]; then
while IFS= read -r -d '' file; do
printf 'file:%s\n' "${file}"
sha256sum "${file}"
done < <(find "${path}" -type f -print0 | sort -z)
elif [[ -f "${path}" ]]; then
printf 'file:%s\n' "${path}"
sha256sum "${path}"
else
printf 'missing:%s\n' "${path}"
fi
done | sha256sum | cut -d' ' -f1
}
clean_docker_tag() {
local input="$1"
echo "${input}" | sed 's/[^a-zA-Z0-9._-]/_/g' | cut -c1-128
}
tag_component() {
local input="$1"
local max_chars="${2:-24}"
clean_docker_tag "${input:-unknown}" | cut -c1-"${max_chars}"
}
extract_arg_default() {
local arg_name="$1"
sed -n -E "s/^[[:space:]]*ARG[[:space:]]+${arg_name}=\"?([^\"[:space:]]+)\"?.*/\\1/p" \
"${DOCKERFILE}" | head -1
}
resolve_image_digest() {
local image_ref="$1"
docker buildx imagetools inspect "${image_ref}" 2>/dev/null \
| sed -n -E 's/^Digest:[[:space:]]+//p' \
| head -1 || true
}
resolve_rocm_base_arg_value() {
local arg_name="$1"
local use_sccache="$2"
case "${arg_name}" in
USE_SCCACHE)
printf '%s\n' "${use_sccache}"
;;
*)
extract_arg_default "${arg_name}"
;;
esac
}
hash_rocm_base_arg_values() {
local use_sccache="$1"
local base_image_digest="$2"
local arg_name=""
local arg_value=""
shift 2 || true
for arg_name in "$@"; do
[[ -n "${arg_name}" ]] || continue
arg_value=$(resolve_rocm_base_arg_value "${arg_name}" "${use_sccache}")
printf 'arg:%s=%s\n' "${arg_name}" "${arg_value:-<empty>}"
if [[ "${arg_name}" == "BASE_IMAGE" && -n "${arg_value}" ]]; then
printf 'arg:%s.digest=%s\n' "${arg_name}" "${base_image_digest:-unknown}"
fi
done
}
rocm_version_from_base_image() {
local base_image="$1"
local version=""
version="$(sed -n -E 's/.*:([0-9]+\.[0-9]+(\.[0-9]+)?)-.*/\1/p' <<<"${base_image}")"
tag_component "${version:-${base_image}}" 16
}
git_diff_changed_base() {
local range="$1"
[[ -n "$(git diff --name-only "${range}" -- "${DOCKERFILE}" 2>/dev/null)" ]]
}
short_git_ref() {
local ref="$1"
git rev-parse --short "${ref}" 2>/dev/null || printf '%s\n' "${ref}"
}
extract_arg_default_from_ref() {
local ref="$1"
local arg_name="$2"
local content=""
content="$(git show "${ref}:${DOCKERFILE}" 2>/dev/null || true)"
sed -n -E "s/^[[:space:]]*ARG[[:space:]]+${arg_name}=\"?([^\"[:space:]]+)\"?.*/\\1/p" \
<<<"${content}" | head -1
}
log_arg_default_changes() {
local old_ref="$1"
local new_ref="$2"
local content_args="${ROCM_BASE_CONTENT_ARGS:-${DEFAULT_ROCM_BASE_CONTENT_ARGS}}"
local arg_name=""
local old_value=""
local new_value=""
local changed=0
echo "Changed ROCm base ARG defaults:"
for arg_name in ${content_args}; do
old_value="$(extract_arg_default_from_ref "${old_ref}" "${arg_name}")"
new_value="$(extract_arg_default_from_ref "${new_ref}" "${arg_name}")"
if [[ "${old_value}" != "${new_value}" ]]; then
echo " - ${arg_name}: ${old_value:-<unset>} -> ${new_value:-<unset>}"
changed=1
fi
done
if [[ "${changed}" == "0" ]]; then
echo " - none detected; Dockerfile instructions changed outside tracked ARG defaults"
fi
}
log_arg_line_diff() {
local range="$1"
local arg_diff=""
arg_diff="$(
git diff --unified=0 "${range}" -- "${DOCKERFILE}" 2>/dev/null \
| awk '/^[+-][[:space:]]*ARG[[:space:]]/ && $0 !~ /^(---|\+\+\+)/ { print " " $0 }' \
|| true
)"
if [[ -n "${arg_diff}" ]]; then
echo "Changed Dockerfile ARG lines:"
printf '%s\n' "${arg_diff}"
fi
}
log_rocm_base_change_check() {
local context="$1"
local range="$2"
local old_ref="$3"
local old_short=""
local head_short=""
old_short="$(short_git_ref "${old_ref}")"
head_short="$(short_git_ref HEAD)"
echo "--- :mag: ROCm base refresh check"
echo "Context: ${context}"
echo "Dockerfile: ${DOCKERFILE}"
echo "Base revision: ${old_short}"
echo "Head revision: ${head_short}"
echo "Git diff range: ${range}"
}
log_rocm_base_rebuild_reason() {
local context="$1"
local range="$2"
local old_ref="$3"
local changed_files=""
log_rocm_base_change_check "${context}" "${range}" "${old_ref}"
changed_files="$(git diff --name-only "${range}" -- "${DOCKERFILE}" 2>/dev/null || true)"
echo "Changed files:"
if [[ -n "${changed_files}" ]]; then
sed 's/^/ - /' <<<"${changed_files}"
else
echo " - ${DOCKERFILE}"
fi
log_arg_default_changes "${old_ref}" HEAD
log_arg_line_diff "${range}"
echo "Decision: rebuilding ROCm base image because ${DOCKERFILE} changed."
}
rocm_base_changed_in_range() {
local context="$1"
local range="$2"
local old_ref="$3"
if git_diff_changed_base "${range}"; then
log_rocm_base_rebuild_reason "${context}" "${range}" "${old_ref}"
return 0
fi
log_rocm_base_change_check "${context}" "${range}" "${old_ref}"
echo "Decision: ROCm base refresh not required; ${DOCKERFILE} is unchanged."
return 1
}
rocm_base_changed() {
local base_branch="${BUILDKITE_PULL_REQUEST_BASE_BRANCH:-main}"
local base_ref="refs/remotes/origin/${base_branch}"
local merge_base=""
if [[ "${ROCM_BASE_REFRESH_SKIP:-0}" == "1" ]]; then
echo "ROCM_BASE_REFRESH_SKIP=1 set; skipping ROCm base refresh"
return 1
fi
if [[ "${ROCM_BASE_REFRESH_FORCE:-0}" == "1" ]]; then
echo "ROCM_BASE_REFRESH_FORCE=1 set; refreshing ROCm base image"
return 0
fi
if ! git rev-parse --is-inside-work-tree >/dev/null 2>&1; then
echo "Not in a git checkout; skipping ROCm base refresh unless forced"
return 1
fi
if [[ "${BUILDKITE_PULL_REQUEST:-false}" != "false" ]]; then
git fetch --no-tags --depth=200 origin \
"+refs/heads/${base_branch}:${base_ref}" >/dev/null 2>&1 || true
merge_base=$(git merge-base HEAD "${base_ref}" 2>/dev/null || true)
if [[ -z "${merge_base}" ]]; then
echo "Unable to determine merge base with PR base ${base_ref}; skipping ROCm base refresh unless forced"
return 1
fi
if rocm_base_changed_in_range \
"pull request build against ${base_ref}" \
"${merge_base}...HEAD" \
"${merge_base}"; then
return 0
fi
elif [[ "${BUILDKITE_BRANCH:-}" == "${ROCM_BASE_STABLE_BRANCH:-main}" ]] \
&& git rev-parse --verify HEAD~1 >/dev/null 2>&1; then
if rocm_base_changed_in_range \
"stable branch build; comparing against previous ${ROCM_BASE_STABLE_BRANCH:-main} commit" \
"HEAD~1..HEAD" \
"HEAD~1"; then
return 0
fi
else
git fetch --no-tags --depth=200 origin \
"+refs/heads/${base_branch}:${base_ref}" >/dev/null 2>&1 || true
merge_base=$(git merge-base HEAD "${base_ref}" 2>/dev/null || true)
if [[ -z "${merge_base}" ]]; then
echo "Unable to determine merge base with branch base ${base_ref}; skipping ROCm base refresh unless forced"
return 1
fi
if rocm_base_changed_in_range \
"branch build against ${base_ref}" \
"${merge_base}...HEAD" \
"${merge_base}"; then
return 0
fi
fi
return 1
}
should_push_stable_tag() {
if [[ "${BUILDKITE_PULL_REQUEST:-false}" != "false" ]]; then
return 1
fi
if [[ "${ROCM_BASE_PUSH_STABLE_TAG:-}" == "1" ]]; then
return 0
fi
if [[ "${ROCM_BASE_PUSH_STABLE_TAG:-}" == "0" ]]; then
return 1
fi
[[ "${BUILDKITE_PULL_REQUEST:-false}" == "false" \
&& "${BUILDKITE_BRANCH:-}" == "${ROCM_BASE_STABLE_BRANCH:-main}" ]]
}
setup_builder() {
echo "--- :buildkite: Setting up buildx builder for ROCm base"
if docker buildx inspect "${BUILDER_NAME}" >/dev/null 2>&1; then
docker buildx use "${BUILDER_NAME}"
else
docker buildx create --name "${BUILDER_NAME}" --driver docker-container --use
fi
docker buildx inspect --bootstrap
}
compute_base_content_hash() {
local use_sccache="$1"
local base_image_digest="$2"
local content_files="${ROCM_BASE_CONTENT_FILES:-${DEFAULT_ROCM_BASE_CONTENT_FILES}}"
local content_args="${ROCM_BASE_CONTENT_ARGS:-${DEFAULT_ROCM_BASE_CONTENT_ARGS}}"
local -a content_paths=()
local -a content_arg_names=()
read -r -a content_paths <<< "${content_files}"
read -r -a content_arg_names <<< "${content_args}"
{
printf 'content-files-hash:%s\n' "$(compute_content_hash "${content_paths[@]}")"
printf 'dockerfile:%s\n' "${DOCKERFILE}"
printf 'resolved-build-args:\n'
hash_rocm_base_arg_values \
"${use_sccache}" "${base_image_digest}" "${content_arg_names[@]}"
} | sha256sum | cut -d' ' -f1
}
build_base_image() {
local use_sccache="${ROCM_BASE_USE_SCCACHE:-${USE_SCCACHE:-0}}"
local base_hash=""
local build_date=""
local build_suffix=""
local base_image_arg=""
local base_image_digest=""
local rocm_version=""
local triton_arg=""
local pytorch_arg=""
local pytorch_vision_arg=""
local pytorch_audio_arg=""
local fa_arg=""
local aiter_arg=""
local mori_arg=""
local python_version_arg=""
local pytorch_rocm_arch_arg=""
local pytorch_branch=""
local aiter_branch=""
local dependency_summary=""
local descriptor=""
local ci_descriptor=""
local descriptive_tag=""
local stable_tag="${BASE_REPO}:base"
local ci_descriptive_tag=""
local content_files="${ROCM_BASE_CONTENT_FILES:-${DEFAULT_ROCM_BASE_CONTENT_FILES}}"
local content_args="${ROCM_BASE_CONTENT_ARGS:-${DEFAULT_ROCM_BASE_CONTENT_ARGS}}"
local content_files_hash=""
local metadata_version="${ROCM_BASE_METADATA_VERSION:-${DEFAULT_ROCM_BASE_METADATA_VERSION}}"
local -a tags=()
local -a no_cache_args=()
local -a sccache_args=()
local -a content_paths=()
if [[ ! -f "${DOCKERFILE}" ]]; then
echo "Error: ROCm base Dockerfile not found: ${DOCKERFILE}" >&2
exit 1
fi
build_date="${ROCM_BASE_TAG_DATE:-$(date -u +%Y%m%d)}"
if [[ -n "${BUILDKITE_BUILD_NUMBER:-}" ]]; then
build_suffix="_bk_${BUILDKITE_BUILD_NUMBER}"
fi
base_image_arg="$(extract_arg_default BASE_IMAGE)"
base_image_digest="$(resolve_image_digest "${base_image_arg}")"
read -r -a content_paths <<< "${content_files}"
content_files_hash="$(compute_content_hash "${content_paths[@]}")"
base_hash=$(compute_base_content_hash "${use_sccache}" "${base_image_digest}")
rocm_version="$(rocm_version_from_base_image "${base_image_arg}")"
triton_arg="$(extract_arg_default TRITON_BRANCH)"
pytorch_arg="$(extract_arg_default PYTORCH_BRANCH)"
pytorch_vision_arg="$(extract_arg_default PYTORCH_VISION_BRANCH)"
pytorch_audio_arg="$(extract_arg_default PYTORCH_AUDIO_BRANCH)"
fa_arg="$(extract_arg_default FA_BRANCH)"
aiter_arg="$(extract_arg_default AITER_BRANCH)"
mori_arg="$(extract_arg_default MORI_BRANCH)"
python_version_arg="$(extract_arg_default PYTHON_VERSION)"
pytorch_rocm_arch_arg="$(extract_arg_default PYTORCH_ROCM_ARCH)"
pytorch_branch="$(tag_component "${pytorch_arg}" 16)"
aiter_branch="$(tag_component "${aiter_arg}" 24)"
dependency_summary="base=${base_image_arg},rocm=${rocm_version},python=${python_version_arg},pytorch=${pytorch_arg},torchvision=${pytorch_vision_arg},torchaudio=${pytorch_audio_arg},triton=${triton_arg},flash-attn=${fa_arg},aiter=${aiter_arg},mori=${mori_arg},pytorch-rocm-arch=${pytorch_rocm_arch_arg}"
descriptor="$(clean_docker_tag "base_custom_aiter_${aiter_branch}_torch_${pytorch_branch}_${build_date}${build_suffix}")"
ci_descriptor="$(clean_docker_tag "ci_custom_aiter_${aiter_branch}_torch_${pytorch_branch}_${build_date}${build_suffix}")"
descriptive_tag="${BASE_REPO}:${descriptor}"
ci_descriptive_tag="${CI_IMAGE_REPO}:${ci_descriptor}"
tags=(-t "${descriptive_tag}")
if should_push_stable_tag; then
tags+=(-t "${stable_tag}")
metadata_set "rocm-base-push-stable-tag" "1"
else
metadata_set "rocm-base-push-stable-tag" "0"
fi
if [[ "${ROCM_BASE_NO_CACHE:-1}" == "1" ]]; then
no_cache_args=(--no-cache)
fi
for env_name in \
SCCACHE_DOWNLOAD_URL \
SCCACHE_ENDPOINT \
SCCACHE_BUCKET_NAME \
SCCACHE_REGION_NAME \
SCCACHE_S3_NO_CREDENTIALS; do
if [[ -n "${!env_name:-}" ]]; then
sccache_args+=(--build-arg "${env_name}=${!env_name}")
fi
done
echo "--- :docker: Building ROCm base image"
echo "Dockerfile: ${DOCKERFILE}"
echo "Descriptive tag: ${descriptive_tag}"
echo "Stable tag: ${stable_tag} ($(should_push_stable_tag && echo enabled || echo disabled))"
echo "Content hash: ${base_hash}"
echo "Dependency summary: ${dependency_summary}"
echo "USE_SCCACHE: ${use_sccache}"
docker buildx build \
"${no_cache_args[@]}" \
--pull \
--progress "${BUILDKIT_PROGRESS:-plain}" \
--file "${DOCKERFILE}" \
--build-arg "USE_SCCACHE=${use_sccache}" \
"${sccache_args[@]}" \
--label "org.opencontainers.image.source=https://github.com/vllm-project/vllm" \
--label "org.opencontainers.image.vendor=vLLM" \
--label "org.opencontainers.image.title=vLLM ROCm base" \
--label "org.opencontainers.image.revision=${BUILDKITE_COMMIT:-}" \
--label "vllm.rocm_base.metadata_version=${metadata_version}" \
--label "vllm.rocm_base.content_hash=${base_hash}" \
--label "vllm.rocm_base.content_files_hash=${content_files_hash}" \
--label "vllm.rocm_base.dockerfile=${DOCKERFILE}" \
--label "vllm.rocm_base.image.descriptive=${descriptive_tag}" \
--label "vllm.rocm_base.image.stable=${stable_tag}" \
--label "vllm.rocm_base.git_commit=${BUILDKITE_COMMIT:-}" \
--label "vllm.rocm_base.stable_branch=${ROCM_BASE_STABLE_BRANCH:-main}" \
--label "vllm.rocm_base.descriptor=${descriptor}" \
--label "vllm.rocm_base.dependency_summary=${dependency_summary}" \
--label "vllm.rocm_base.base_image=${base_image_arg}" \
--label "vllm.rocm_base.base_image_digest=${base_image_digest}" \
--label "vllm.rocm_base.dependency.rocm=${rocm_version}" \
--label "vllm.rocm_base.dependency.python=${python_version_arg}" \
--label "vllm.rocm_base.dependency.pytorch=${pytorch_arg}" \
--label "vllm.rocm_base.dependency.torchvision=${pytorch_vision_arg}" \
--label "vllm.rocm_base.dependency.torchaudio=${pytorch_audio_arg}" \
--label "vllm.rocm_base.dependency.triton=${triton_arg}" \
--label "vllm.rocm_base.dependency.flash_attention=${fa_arg}" \
--label "vllm.rocm_base.dependency.aiter=${aiter_arg}" \
--label "vllm.rocm_base.dependency.mori=${mori_arg}" \
--label "vllm.rocm_base.pytorch_rocm_arch=${pytorch_rocm_arch_arg}" \
"${tags[@]}" \
--push \
.
docker buildx imagetools inspect "${descriptive_tag}" >/dev/null
metadata_set "rocm-base-refresh" "1"
metadata_set "rocm-base-image" "${descriptive_tag}"
metadata_set "rocm-base-image-descriptive" "${descriptive_tag}"
metadata_set "rocm-base-image-stable" "${stable_tag}"
metadata_set "rocm-base-image-ci-descriptive" "${ci_descriptive_tag}"
metadata_set "rocm-base-metadata-version" "${metadata_version}"
metadata_set "rocm-base-content-hash" "${base_hash}"
metadata_set "rocm-base-content-files-hash" "${content_files_hash}"
metadata_set "rocm-base-content-files" "${content_files}"
metadata_set "rocm-base-content-args" "${content_args}"
metadata_set "rocm-base-base-image-digest" "${base_image_digest}"
metadata_set "rocm-base-dockerfile" "${DOCKERFILE}"
metadata_set "rocm-base-descriptor" "${descriptor}"
metadata_set "rocm-base-dependency-summary" "${dependency_summary}"
metadata_set "rocm-base-dependency-rocm" "${rocm_version}"
metadata_set "rocm-base-dependency-python" "${python_version_arg}"
metadata_set "rocm-base-dependency-pytorch" "${pytorch_arg}"
metadata_set "rocm-base-dependency-torchvision" "${pytorch_vision_arg}"
metadata_set "rocm-base-dependency-torchaudio" "${pytorch_audio_arg}"
metadata_set "rocm-base-dependency-triton" "${triton_arg}"
metadata_set "rocm-base-dependency-flash-attention" "${fa_arg}"
metadata_set "rocm-base-dependency-aiter" "${aiter_arg}"
metadata_set "rocm-base-dependency-mori" "${mori_arg}"
metadata_set "rocm-base-pytorch-rocm-arch" "${pytorch_rocm_arch_arg}"
metadata_set "rocm-ci-image-descriptive" "${ci_descriptive_tag}"
echo "--- :white_check_mark: ROCm base image published"
echo "Use BASE_IMAGE=${descriptive_tag} for downstream ROCm CI builds"
}
main() {
metadata_set "rocm-base-refresh" "0"
if ! rocm_base_changed; then
echo "ROCm base Dockerfile did not change; skipping base image refresh"
return 0
fi
setup_builder
build_base_image
}
main "$@"
@@ -1,32 +0,0 @@
#!/usr/bin/env bash
# Fast structural smoke test for the full ROCm CI image.
set -euo pipefail
image_ref="${VLLM_CI_SMOKE_IMAGE:-rocm/vllm-ci:${BUILDKITE_COMMIT:?BUILDKITE_COMMIT is required}}"
docker run --rm --network=none --entrypoint /bin/bash "${image_ref}" -ec '
if [ ! -d /vllm-workspace ]; then echo Missing directory: /vllm-workspace >&2; exit 1; fi
if [ ! -d /vllm-workspace/tests ]; then echo Missing directory: /vllm-workspace/tests >&2; exit 1; fi
if [ ! -d /vllm-workspace/src/vllm ]; then echo Missing directory: /vllm-workspace/src/vllm >&2; exit 1; fi
if [ ! -x /vllm-workspace/src/vllm/vllm-rs ]; then echo Missing executable: /vllm-workspace/src/vllm/vllm-rs >&2; exit 1; fi
command -v python3
command -v uv
command -v pytest
if ! command -v amd-smi >/dev/null 2>&1 && ! command -v rocminfo >/dev/null 2>&1; then
echo No ROCm CLI found in image >&2
exit 1
fi
python3 - <<PY
import torch
import vllm
print(torch.__version__)
print(vllm.__version__)
PY
echo AMD image smoke OK
'
+1 -3
View File
@@ -109,9 +109,7 @@ run_nodes() {
if [ "$node" -ne 0 ]; then
docker exec -d "node$node" /bin/bash -c "cd $WORKING_DIR ; ${COMMANDS[$node]}"
else
# Allocate a TTY (-t -i) for the foreground head node so its output
# keeps ANSI color in the Buildkite log (see run-amd-test.sh).
docker exec -t -i "node$node" /bin/bash -c "cd $WORKING_DIR ; ${COMMANDS[$node]}"
docker exec "node$node" /bin/bash -c "cd $WORKING_DIR ; ${COMMANDS[$node]}"
fi
done
}
@@ -8,12 +8,7 @@ if [[ "$MODE" != "style-clippy" && "$MODE" != "test" ]]; then
exit 2
fi
if ROOT_DIR="$(git rev-parse --show-toplevel 2>/dev/null)"; then
:
else
SCRIPT_DIR="$(cd -- "$(dirname -- "${BASH_SOURCE[0]}")" && pwd -P)"
ROOT_DIR="$(cd -- "${SCRIPT_DIR}/../.." && pwd -P)"
fi
ROOT_DIR="$(git rev-parse --show-toplevel)"
cd "$ROOT_DIR"
export CARGO_TERM_COLOR="${CARGO_TERM_COLOR:-always}"
@@ -168,8 +163,8 @@ run_style_clippy() {
log_section "Checking Rust dependency bans"
cargo deny \
--manifest-path rust/Cargo.toml \
--config rust/deny.toml \
check \
--config rust/deny.toml \
bans
log_section "Running clippy"
+11 -32
View File
@@ -23,7 +23,6 @@ NC='\033[0m' # No Color
# Default configuration
PIPELINE="ci"
DRY_RUN=true
TORCH_NIGHTLY=false
usage() {
cat <<EOF
@@ -35,14 +34,12 @@ Sets RUN_ALL=1 and NIGHTLY=1 environment variables.
SAFETY: Dry-run by default. Use --execute to actually trigger a build.
Options:
--execute Actually trigger the build (default: dry-run)
--pipeline Buildkite pipeline slug (default: ${PIPELINE})
--commit Override commit SHA (default: current HEAD)
--branch Override branch name (default: current branch)
--message Custom build message (default: auto-generated)
--torch-nightly Also build and run the full suite against torch nightly
(sets TORCH_NIGHTLY=1)
--help Show this help message
--execute Actually trigger the build (default: dry-run)
--pipeline Buildkite pipeline slug (default: ${PIPELINE})
--commit Override commit SHA (default: current HEAD)
--branch Override branch name (default: current branch)
--message Custom build message (default: auto-generated)
--help Show this help message
Prerequisites:
- bk CLI installed: brew tap buildkite/buildkite && brew install buildkite/buildkite/bk
@@ -52,7 +49,6 @@ Examples:
$(basename "$0") # Dry-run, show what would happen
$(basename "$0") --execute # Actually trigger the build
$(basename "$0") --pipeline ci-shadow # Dry-run with different pipeline
$(basename "$0") --torch-nightly # Dry-run a full torch-nightly run
EOF
exit 1
}
@@ -100,10 +96,6 @@ while [[ $# -gt 0 ]]; do
MESSAGE="$2"
shift 2
;;
--torch-nightly)
TORCH_NIGHTLY=true
shift
;;
--help|-h)
usage
;;
@@ -179,17 +171,11 @@ if [[ $(echo "$REMOTE_BRANCHES" | wc -l) -gt 5 ]]; then
fi
echo ""
# Environment variables passed to the build.
BUILD_ENV=("RUN_ALL=1" "NIGHTLY=1")
if [[ "$TORCH_NIGHTLY" == true ]]; then
BUILD_ENV+=("TORCH_NIGHTLY=1")
fi
log_info "Pipeline: ${PIPELINE}"
log_info "Branch: ${BRANCH}"
log_info "Commit: ${COMMIT}"
log_info "Message: ${MESSAGE}"
log_info "Environment: ${BUILD_ENV[*]}"
log_info "Environment: RUN_ALL=1, NIGHTLY=1"
echo ""
# Build the command
@@ -201,10 +187,9 @@ CMD=(bk build create
--commit "${COMMIT}"
--branch "${BRANCH}"
--message "${MESSAGE}"
--env "RUN_ALL=1"
--env "NIGHTLY=1"
)
for env_var in "${BUILD_ENV[@]}"; do
CMD+=(--env "${env_var}")
done
if [[ "$DRY_RUN" == true ]]; then
echo "=========================================="
@@ -225,14 +210,8 @@ if [[ "$DRY_RUN" == true ]]; then
echo " --commit '$(escape_for_shell "${COMMIT}")' \\"
echo " --branch '$(escape_for_shell "${BRANCH}")' \\"
echo " --message '$(escape_for_shell "${MESSAGE}")' \\"
last_idx=$(( ${#BUILD_ENV[@]} - 1 ))
for i in "${!BUILD_ENV[@]}"; do
if [[ $i -eq $last_idx ]]; then
echo " --env '$(escape_for_shell "${BUILD_ENV[$i]}")'"
else
echo " --env '$(escape_for_shell "${BUILD_ENV[$i]}")' \\"
fi
done
echo " --env 'RUN_ALL=1' \\"
echo " --env 'NIGHTLY=1'"
echo ""
echo "=========================================="
echo -e "${YELLOW}To actually trigger this build, run:${NC}"
@@ -1,13 +0,0 @@
#!/bin/bash
set -euo pipefail
REGISTRY="public.ecr.aws/q9t5s3a7"
REPO="vllm-release-repo"
ARCH_TAG="${BUILDKITE_COMMIT}-$(uname -m)-xpu"
PLATFORM_TAG="${BUILDKITE_COMMIT}-xpu"
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin ${REGISTRY}
docker manifest rm ${REGISTRY}/${REPO}:${PLATFORM_TAG} || true
docker manifest create ${REGISTRY}/${REPO}:${PLATFORM_TAG} ${REGISTRY}/${REPO}:${ARCH_TAG} --amend
docker manifest push ${REGISTRY}/${REPO}:${PLATFORM_TAG}
@@ -1,21 +0,0 @@
#!/bin/bash
set -ex
ORIG_TAG_NAME="$BUILDKITE_COMMIT"
REPO="vllm/vllm-openai-xpu"
echo "Pushing original XPU tag ${ORIG_TAG_NAME}-xpu to nightly tags in ${REPO}"
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:"$ORIG_TAG_NAME"-x86_64-xpu
docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:"$ORIG_TAG_NAME"-x86_64-xpu ${REPO}:nightly-x86_64
docker push ${REPO}:nightly-x86_64
docker manifest rm ${REPO}:nightly || true
docker manifest rm ${REPO}:nightly-"$BUILDKITE_COMMIT" || true
docker manifest create ${REPO}:nightly ${REPO}:nightly-x86_64 --amend
docker manifest create ${REPO}:nightly-"$BUILDKITE_COMMIT" ${REPO}:nightly-x86_64 --amend
docker manifest push ${REPO}:nightly
docker manifest push ${REPO}:nightly-"$BUILDKITE_COMMIT"
+38 -254
View File
@@ -275,7 +275,7 @@ steps:
- pytest -v -s v1/cudagraph/test_cudagraph_mode.py
- label: e2e Core (1 GPU) # TBD
timeout_in_minutes: 180
timeout_in_minutes: 35
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
optional: true
@@ -380,7 +380,6 @@ steps:
- tests/test_outputs.py
- tests/test_pooling_params.py
- tests/test_ray_env.py
- tests/test_sampling_params.py
- tests/multimodal
- tests/renderers
- tests/standalone_tests/lazy_imports.py
@@ -396,7 +395,6 @@ steps:
- pytest -v -s test_outputs.py
- pytest -v -s test_pooling_params.py
- pytest -v -s test_ray_env.py
- pytest -v -s test_sampling_params.py
- pytest -v -s -m 'cpu_test' multimodal
- pytest -v -s renderers
- pytest -v -s tokenizers_
@@ -406,40 +404,10 @@ steps:
- pytest -v -s transformers_utils
- pytest -v -s config
#------------------------------------------------------------ mi250 · rust -----------------------------------------------------------#
- label: Rust Frontend Cargo Style + Clippy # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
no_gpu: true
working_dir: "/vllm-workspace"
source_file_dependencies:
- rust/
- rust-toolchain.toml
- .buildkite/test_areas/rust_frontend_cargo.yaml
- .buildkite/scripts/run-rust-frontend-cargo-ci.sh
commands:
- bash .buildkite/scripts/run-rust-frontend-cargo-ci.sh style-clippy
- label: Rust Frontend Cargo Tests # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
no_gpu: true
working_dir: "/vllm-workspace"
source_file_dependencies:
- rust/
- rust-toolchain.toml
- .buildkite/test_areas/rust_frontend_cargo.yaml
- .buildkite/scripts/run-rust-frontend-cargo-ci.sh
commands:
- bash .buildkite/scripts/run-rust-frontend-cargo-ci.sh test
#----------------------------------------------------------- mi250 · docker ----------------------------------------------------------#
- label: Docker Build Metadata (ROCm) # TBD
timeout_in_minutes: 180
timeout_in_minutes: 20
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
no_gpu: true
@@ -467,7 +435,7 @@ steps:
#----------------------------------------------------- mi300 · basic_correctness -----------------------------------------------------#
- label: Basic Correctness # TBD
timeout_in_minutes: 180
timeout_in_minutes: 95
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
fast_check: true
@@ -484,7 +452,7 @@ steps:
- pytest -v -s basic_correctness/test_cpu_offload.py
- label: Distributed Model Tests (2 GPUs) # TBD
timeout_in_minutes: 180
timeout_in_minutes: 110
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_2
num_gpus: 2
@@ -504,7 +472,7 @@ steps:
commands:
- TARGET_TEST_SUITE=MI300 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/transformers/test_backend.py -v -s -m 'distributed(num_gpus=2)'
- pytest models/test_transformers.py -v -s -m 'distributed(num_gpus=2)'
- pytest models/language -v -s -m 'distributed(num_gpus=2)'
- pytest models/multimodal -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)'
@@ -724,7 +692,7 @@ steps:
- pytest -v -s distributed/test_eplb_spec_decode.py
- label: Distributed Tests (2xH100-2xMI300) # TBD
timeout_in_minutes: 180
timeout_in_minutes: 75
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_2
num_gpus: 2
@@ -1297,7 +1265,7 @@ steps:
#--------------------------------------------------------- mi300 · examples ----------------------------------------------------------#
- label: Examples # TBD
timeout_in_minutes: 180
timeout_in_minutes: 90
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
@@ -1333,7 +1301,7 @@ steps:
#---------------------------------------------------------- mi300 · kernels ----------------------------------------------------------#
- label: vLLM IR Tests # TBD
timeout_in_minutes: 180
timeout_in_minutes: 30
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
@@ -1348,7 +1316,7 @@ steps:
- pytest -v -s tests/kernels/ir
- label: Kernels Attention Test %N # TBD
timeout_in_minutes: 180
timeout_in_minutes: 100
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
@@ -1382,7 +1350,7 @@ steps:
- 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 KDA Test # TBD
timeout_in_minutes: 180
timeout_in_minutes: 30
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
@@ -1397,7 +1365,7 @@ steps:
- pytest -v -s kernels/test_kda.py
- label: Kernels MoE Test %N # TBD
timeout_in_minutes: 180
timeout_in_minutes: 95
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
@@ -1511,7 +1479,7 @@ steps:
#---------------------------------------------------- mi300 · model_runner_v2 -------------------------------------------------------#
- label: Model Runner V2 Core Tests # TBD
timeout_in_minutes: 180
timeout_in_minutes: 60
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
@@ -1535,7 +1503,7 @@ 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 # TBD
timeout_in_minutes: 180
timeout_in_minutes: 60
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
@@ -1569,7 +1537,7 @@ steps:
- 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) # TBD
timeout_in_minutes: 180
timeout_in_minutes: 60
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_2
num_gpus: 2
@@ -1590,7 +1558,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) # TBD
timeout_in_minutes: 180
timeout_in_minutes: 60
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_4
num_gpus: 4
@@ -1609,7 +1577,7 @@ steps:
- pytest -v -s distributed/test_pp_cudagraph.py -k "not ray"
- label: Model Runner V2 Spec Decode # TBD
timeout_in_minutes: 180
timeout_in_minutes: 45
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
@@ -1636,7 +1604,7 @@ steps:
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
parallelism: 6
parallelism: 2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/model_executor/models/
@@ -1661,7 +1629,7 @@ steps:
- pytest -v -s models/test_initialization.py::test_can_initialize_small_subset
- label: Basic Models Tests (Other) # TBD
timeout_in_minutes: 180
timeout_in_minutes: 90
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
@@ -1669,10 +1637,10 @@ steps:
source_file_dependencies:
- vllm/
- tests/models/test_terratorch.py
- tests/models/transformers/test_backend.py
- tests/models/test_transformers.py
- tests/models/test_registry.py
commands:
- pytest -v -s models/test_terratorch.py models/transformers/test_backend.py models/test_registry.py
- pytest -v -s models/test_terratorch.py models/test_transformers.py models/test_registry.py
#----------------------------------------------------- mi300 · models / language -----------------------------------------------------#
@@ -1806,7 +1774,7 @@ steps:
- 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: 180
timeout_in_minutes: 138
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
@@ -1892,7 +1860,7 @@ steps:
- examples/
commands:
- pip install --upgrade git+https://github.com/huggingface/transformers
- pytest -v -s tests/models/transformers/test_backend.py
- pytest -v -s tests/models/test_transformers.py
- 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
@@ -1938,11 +1906,6 @@ steps:
- pytest -v -s plugins_tests/test_stats_logger_plugins.py
- pip uninstall dummy_stat_logger -y
# END: `stat_logger` plugins test
# BEGIN: `endpoint` plugins test
- pip install -e ./plugins/vllm_add_dummy_endpoint_plugin
- pytest -v -s plugins_tests/test_endpoint_plugins.py
- pip uninstall vllm_add_dummy_endpoint_plugin -y
# END: `endpoint` plugins test
# BEGIN: other tests
- pytest -v -s plugins_tests/test_scheduler_plugins.py
- pip install -e ./plugins/vllm_add_dummy_model
@@ -1952,7 +1915,7 @@ steps:
- pytest -v -s plugins_tests/lora_resolvers # unit tests for in-tree lora resolver plugins
- label: GGUF Plugin # TBD
timeout_in_minutes: 180
timeout_in_minutes: 30
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
@@ -1967,128 +1930,6 @@ steps:
- pip install "vllm-gguf-plugin >= 0.0.2"
- pytest -v -s plugins_tests/gguf
#------------------------------------------------------- mi300 · rust_frontend -------------------------------------------------------#
- label: Rust Frontend OpenAI Coverage # 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:
- rust/
- vllm/benchmarks/
- vllm/entrypoints/openai/
- vllm/entrypoints/serve/
- vllm/v1/sample/
- tests/utils.py
- tests/benchmarks/test_serve_cli.py
- tests/entrypoints/openai/chat_completion/test_chat_completion.py
- tests/entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py
- tests/entrypoints/openai/completion/test_shutdown.py
- tests/entrypoints/openai/test_return_token_ids.py
- tests/entrypoints/openai/test_uds.py
- tests/v1/sample/test_logprobs_e2e.py
- vllm/platforms/rocm.py
commands:
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s benchmarks/test_serve_cli.py -k "not insecure and not (test_bench_serve and not test_bench_serve_chat)"
- pytest -v -s entrypoints/openai/chat_completion/test_chat_completion.py -k "not test_invalid_json_schema and not test_invalid_regex"
- pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not multiple"
- pytest -v -s entrypoints/openai/completion/test_shutdown.py -k "not engine_failure and not test_abort_timeout_exits_quickly"
- pytest -v -s entrypoints/openai/test_return_token_ids.py -k "not test_comparison"
- pytest -v -s entrypoints/openai/test_uds.py
- pytest -v -s v1/sample/test_logprobs_e2e.py -k "test_prompt_logprobs_e2e_server"
- label: Rust Frontend Serve Admin Coverage # 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:
- rust/
- vllm/entrypoints/openai/
- vllm/entrypoints/serve/
- vllm/v1/engine/
- tests/utils.py
- tests/entrypoints/serve/dev/rpc/test_collective_rpc.py
- tests/entrypoints/scale_out/token_in_token_out/test_serving_tokens.py
- tests/entrypoints/serve/instrumentator/test_basic.py
- tests/entrypoints/serve/instrumentator/test_metrics.py
- tests/entrypoints/serve/tokenize/test_tokenization.py
- vllm/platforms/rocm.py
commands:
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
- pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not show_version and not server_load"
- pytest -v -s entrypoints/scale_out/token_in_token_out/test_serving_tokens.py -k "not stream and not lora and not test_generate_logprobs and not stop_string_workflow"
- pytest -v -s entrypoints/serve/instrumentator/test_metrics.py -k "text and not show and not run_batch and not test_metrics_counts and not test_metrics_exist"
- pytest -v -s entrypoints/serve/tokenize/test_tokenization.py -k "not tokenizer_info"
- label: Rust Frontend Core Correctness # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- rust/
- vllm/entrypoints/openai/
- tests/utils.py
- tests/entrypoints/openai/correctness/test_lmeval.py
- vllm/platforms/rocm.py
commands:
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
- label: Rust Frontend Tool Use # 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:
- rust/
- vllm/entrypoints/openai/
- vllm/tool_parsers/
- tests/utils.py
- tests/tool_use/
- vllm/platforms/rocm.py
commands:
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s tool_use --ignore=tool_use/mistral --models llama3.2 -k "not test_response_format_with_tool_choice_required and not test_parallel_tool_calls_false and not test_tool_call_and_choice"
- label: Rust Frontend Distributed # 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:
- rust/
- vllm/distributed/
- vllm/engine/
- vllm/executor/
- vllm/v1/engine/
- vllm/v1/worker/
- tests/utils.py
- tests/v1/distributed/test_external_lb_dp.py
- tests/v1/distributed/test_hybrid_lb_dp.py
- tests/v1/distributed/test_internal_lb_dp.py
- vllm/platforms/rocm.py
commands:
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_internal_lb_dp.py -k "not 4 and not server_info"
- TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py -k "not 4 and not server_info"
- TP_SIZE=1 DP_SIZE=4 pytest -v -s v1/distributed/test_hybrid_lb_dp.py -k "not 4 and not server_info"
#------------------------------------------------------- mi300 · quantization --------------------------------------------------------#
- label: Quantization # TBD
@@ -2255,7 +2096,7 @@ steps:
- 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
timeout_in_minutes: 90
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
@@ -2472,7 +2313,7 @@ steps:
- DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py
- label: Metrics, Tracing (2 GPUs) # TBD
timeout_in_minutes: 180
timeout_in_minutes: 65
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_2
optional: true
@@ -2598,59 +2439,6 @@ steps:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
- HYBRID_SSM=1 ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
- label: Hybrid SSM NixlConnector PD prefix cache test (2 GPUs) # 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:
- vllm/distributed/kv_transfer/kv_connector/v1/nixl/
- vllm/v1/core/sched/
- vllm/v1/core/kv_cache_coordinator.py
- tests/v1/kv_connector/nixl_integration/
- vllm/platforms/rocm.py
commands:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
- ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/run_mamba_prefix_cache_test.sh
- label: MultiConnector (Nixl+Offloading) PD accuracy (2 GPUs) # 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:
- 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/
- vllm/platforms/rocm.py
commands:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
- ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/run_multi_connector_accuracy_test.sh
- label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs) # 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:
- 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/
- vllm/platforms/rocm.py
commands:
- uv pip install --system -r /vllm-workspace/requirements/kv_connectors_rocm.txt
- ATTENTION_BACKEND=TRITON_ATTN bash v1/kv_connector/nixl_integration/run_multi_connector_edge_case_test.sh
- label: V1 e2e (4 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
@@ -2678,7 +2466,7 @@ steps:
#------------------------------------------------------ mi300 · weight_loading -------------------------------------------------------#
- label: Weight Loading Multiple GPU # TBD
timeout_in_minutes: 180
timeout_in_minutes: 75
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_2
num_gpus: 2
@@ -2690,7 +2478,7 @@ steps:
- bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-amd.txt
- label: Weight Loading Multiple GPU - Large Models # TBD
timeout_in_minutes: 180
timeout_in_minutes: 75
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_2
num_gpus: 2
@@ -2941,19 +2729,15 @@ steps:
- vllm/envs.py
- examples/offline_inference/data_parallel.py
- tests/distributed/test_context_parallel.py
- tests/distributed/test_rocm_aiter_custom_ar.py
- tests/distributed/test_rocm_quick_reduce.py
- tests/distributed/test_quick_all_reduce.py
- tests/v1/e2e/general/test_rocm_aiter_custom_ar.py
- tests/v1/distributed/test_dbo.py
- tests/utils.py
commands:
- pytest -v -s tests/distributed/test_context_parallel.py
- pytest -v -s tests/distributed/test_rocm_aiter_custom_ar.py
- pytest -v -s tests/v1/e2e/general/test_rocm_aiter_custom_ar.py
- pytest -v -s tests/v1/distributed/test_dbo.py
- pytest -v -s tests/distributed/test_rocm_quick_reduce.py
- pytest -v -s tests/distributed/test_quick_all_reduce.py
- pytest -v -s tests/v1/distributed/test_dbo.py
#-------------------------------------------------------- mi355 · entrypoints --------------------------------------------------------#
@@ -3089,7 +2873,7 @@ steps:
- pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-gfx950.txt
- label: LM Eval Qwen3-5 Models (B200-MI355) # TBD
timeout_in_minutes: 180
timeout_in_minutes: 120
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_2
num_gpus: 2
@@ -3171,7 +2955,7 @@ steps:
#--------------------------------------------------------- mi355 · examples ----------------------------------------------------------#
- label: Examples # TBD
timeout_in_minutes: 180
timeout_in_minutes: 100
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
working_dir: "/vllm-workspace/examples"
@@ -3206,7 +2990,7 @@ steps:
#---------------------------------------------------------- mi355 · kernels ----------------------------------------------------------#
- label: Kernels (B200-MI355) # TBD
timeout_in_minutes: 180
timeout_in_minutes: 15
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
working_dir: "/vllm-workspace/"
@@ -3231,7 +3015,7 @@ steps:
- pytest -v -s tests/kernels/attention/test_rocm_aiter_mla_decode_metadata.py
- label: Kernels Attention Test %N # TBD
timeout_in_minutes: 180
timeout_in_minutes: 100
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
parallelism: 2
@@ -3248,7 +3032,7 @@ steps:
- pytest -v -s kernels/attention --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
- label: Kernels MoE Test %N # TBD
timeout_in_minutes: 180
timeout_in_minutes: 50
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
parallelism: 5
@@ -3512,7 +3296,7 @@ steps:
- pytest -v -s v1/attention
- label: V1 Core + KV + Metrics # TBD
timeout_in_minutes: 180
timeout_in_minutes: 60
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
optional: true
@@ -3538,7 +3322,7 @@ steps:
- pytest -v -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
- label: V1 Sample + Logits # TBD
timeout_in_minutes: 180
timeout_in_minutes: 60
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
optional: true
@@ -3558,7 +3342,7 @@ steps:
- pytest -v -s v1/test_outputs.py
- label: V1 Spec Decode # TBD
timeout_in_minutes: 180
timeout_in_minutes: 60
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
working_dir: "/vllm-workspace/tests"
@@ -3571,7 +3355,7 @@ steps:
#------------------------------------------------------ mi355 · weight_loading -------------------------------------------------------#
- label: Weight Loading Multiple GPU # TBD
timeout_in_minutes: 180
timeout_in_minutes: 75
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_2
num_gpus: 2
@@ -3583,7 +3367,7 @@ steps:
- bash weight_loading/run_model_weight_loading_test.sh -c weight_loading/models-amd.txt
- label: Weight Loading Multiple GPU - Large Models # TBD
timeout_in_minutes: 180
timeout_in_minutes: 75
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_2
working_dir: "/vllm-workspace/tests"
+3 -3
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: V1 attention (H100-MI300)
key: v1-attention-h100-mi300
timeout_in_minutes: 85
timeout_in_minutes: 30
device: h100
source_file_dependencies:
- vllm/config/attention.py
@@ -16,7 +16,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 95
timeout_in_minutes: 70
depends_on:
- image-build-amd
source_file_dependencies:
@@ -30,7 +30,7 @@ steps:
- label: V1 attention (B200)
key: v1-attention-b200
timeout_in_minutes: 80
timeout_in_minutes: 30
device: b200-k8s
source_file_dependencies:
- vllm/config/attention.py
+2 -2
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Basic Correctness
key: basic-correctness
timeout_in_minutes: 45
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/
@@ -19,6 +19,6 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 70
timeout_in_minutes: 50
depends_on:
- image-build-amd
+2 -2
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Benchmarks CLI Test
key: benchmarks-cli-test
timeout_in_minutes: 30
timeout_in_minutes: 20
device: h200_18gb
source_file_dependencies:
- vllm/
@@ -23,7 +23,7 @@ steps:
num_gpus: 2
optional: true
working_dir: "/vllm-workspace/"
timeout_in_minutes: 20
timeout_in_minutes: 10
source_file_dependencies:
- benchmarks/attention_benchmarks/
- vllm/v1/attention/
+11 -11
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Sequence Parallel Correctness Tests (2 GPUs)
key: sequence-parallel-correctness-tests-2-gpus
timeout_in_minutes: 80
timeout_in_minutes: 50
working_dir: "/vllm-workspace/"
num_devices: 2
source_file_dependencies:
@@ -19,7 +19,7 @@ steps:
- label: Sequence Parallel Correctness Tests (2xH100)
key: sequence-parallel-correctness-tests-2xh100
timeout_in_minutes: 75
timeout_in_minutes: 50
working_dir: "/vllm-workspace/"
device: h100
optional: true
@@ -30,7 +30,7 @@ steps:
- label: AsyncTP Correctness Tests (2xH100)
key: asynctp-correctness-tests-2xh100
timeout_in_minutes: 30
timeout_in_minutes: 50
working_dir: "/vllm-workspace/"
device: h100
optional: true
@@ -41,7 +41,7 @@ steps:
- label: AsyncTP Correctness Tests (B200)
key: asynctp-correctness-tests-b200
timeout_in_minutes: 30
timeout_in_minutes: 50
working_dir: "/vllm-workspace/"
device: b200-k8s
optional: true
@@ -52,7 +52,7 @@ steps:
- label: Distributed Compile Unit Tests (2xH100)
key: distributed-compile-unit-tests-2xh100
timeout_in_minutes: 45
timeout_in_minutes: 20
working_dir: "/vllm-workspace/"
device: h100
num_devices: 2
@@ -66,7 +66,7 @@ steps:
- label: Fusion and Compile Unit Tests (2xB200)
key: fusion-and-compile-unit-tests-2xb200
timeout_in_minutes: 30
timeout_in_minutes: 20
working_dir: "/vllm-workspace/"
device: b200-k8s
source_file_dependencies:
@@ -96,7 +96,7 @@ steps:
- label: Fusion E2E Quick (H100)
key: fusion-e2e-quick-h100
timeout_in_minutes: 25
timeout_in_minutes: 15
working_dir: "/vllm-workspace/"
device: h100
num_devices: 1
@@ -115,7 +115,7 @@ steps:
- label: Fusion E2E Config Sweep (H100)
key: fusion-e2e-config-sweep-h100
timeout_in_minutes: 25
timeout_in_minutes: 30
working_dir: "/vllm-workspace/"
device: h100
num_devices: 1
@@ -149,7 +149,7 @@ steps:
- label: Fusion E2E TP2 Quick (H100)
key: fusion-e2e-tp2-quick-h100
timeout_in_minutes: 35
timeout_in_minutes: 20
working_dir: "/vllm-workspace/"
device: h100
num_devices: 2
@@ -167,7 +167,7 @@ steps:
- label: Fusion E2E TP2 AR-RMS Config Sweep (H100)
key: fusion-e2e-tp2-ar-rms-config-sweep-h100
timeout_in_minutes: 30
timeout_in_minutes: 40
working_dir: "/vllm-workspace/"
device: h100
num_devices: 2
@@ -207,7 +207,7 @@ steps:
- label: Fusion E2E TP2 (B200)
key: fusion-e2e-tp2-b200
timeout_in_minutes: 45
timeout_in_minutes: 20
working_dir: "/vllm-workspace/"
device: b200-k8s
num_devices: 2
+2 -2
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Platform Tests
key: platform-tests
timeout_in_minutes: 20
timeout_in_minutes: 15
device: h200_18gb
source_file_dependencies:
- vllm/envs.py
@@ -19,7 +19,7 @@ steps:
- label: Cudagraph
key: cudagraph
timeout_in_minutes: 30
timeout_in_minutes: 20
source_file_dependencies:
- tests/v1/cudagraph
- vllm/v1/cudagraph_dispatcher.py
+12 -25
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Distributed NixlConnector PD accuracy (4 GPUs)
key: distributed-nixlconnector-pd-accuracy-4-gpus
timeout_in_minutes: 55
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -16,7 +16,7 @@ steps:
mirror:
amd:
device: mi300_4
timeout_in_minutes: 85
timeout_in_minutes: 110
depends_on:
- image-build-amd
source_file_dependencies:
@@ -29,7 +29,7 @@ steps:
- label: Distributed FlashInfer NixlConnector PD accuracy (4 GPUs)
key: distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus
timeout_in_minutes: 55
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -39,19 +39,6 @@ steps:
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- FLASHINFER=1 bash v1/kv_connector/nixl_integration/config_sweep_accuracy_test.sh
- label: Push NixlConnector PP prefill PD accuracy (4 GPUs)
key: push-nixlconnector-pp-prefill-pd-accuracy-4-gpus
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
- vllm/distributed/kv_transfer/kv_connector/v1/nixl/
- tests/v1/kv_connector/nixl_integration/
- tests/v1/kv_connector/nixl_push_integration/
commands:
- bash /vllm-workspace/.buildkite/scripts/install-kv-connectors.sh
- bash v1/kv_connector/nixl_push_integration/config_sweep_accuracy_test.sh
- label: DP EP Distributed NixlConnector PD accuracy tests (4 GPUs)
key: dp-ep-distributed-nixlconnector-pd-accuracy-tests-4-gpus
timeout_in_minutes: 30
@@ -66,7 +53,7 @@ steps:
mirror:
amd:
device: mi300_4
timeout_in_minutes: 60
timeout_in_minutes: 50
depends_on:
- image-build-amd
source_file_dependencies:
@@ -79,7 +66,7 @@ steps:
- label: CrossLayer KV layout Distributed NixlConnector PD accuracy tests (4 GPUs)
key: crosslayer-kv-layout-distributed-nixlconnector-pd-accuracy-tests-4-gpus
timeout_in_minutes: 55
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -91,7 +78,7 @@ steps:
mirror:
amd:
device: mi300_4
timeout_in_minutes: 85
timeout_in_minutes: 110
depends_on:
- image-build-amd
source_file_dependencies:
@@ -104,7 +91,7 @@ steps:
- label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs)
key: hybrid-ssm-nixlconnector-pd-accuracy-tests-4-gpus
timeout_in_minutes: 60
timeout_in_minutes: 25
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -116,7 +103,7 @@ steps:
mirror:
amd:
device: mi300_4
timeout_in_minutes: 80
timeout_in_minutes: 60
depends_on:
- image-build-amd
source_file_dependencies:
@@ -143,7 +130,7 @@ steps:
- label: MultiConnector (Nixl+Offloading) PD accuracy (2 GPUs)
key: multiconnector-nixl-offloading-pd-accuracy-2-gpus
timeout_in_minutes: 40
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -158,7 +145,7 @@ steps:
- label: NixlConnector PD + Spec Decode acceptance (2 GPUs)
key: nixlconnector-pd-spec-decode-acceptance-2-gpus
timeout_in_minutes: 45
timeout_in_minutes: 30
device: a100
working_dir: "/vllm-workspace/tests"
num_devices: 2
@@ -172,7 +159,7 @@ steps:
mirror:
amd:
device: mi300_2
timeout_in_minutes: 70
timeout_in_minutes: 60
depends_on:
- image-build-amd
source_file_dependencies:
@@ -186,7 +173,7 @@ steps:
- label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs)
key: multiconnector-nixl-offloading-pd-edge-cases-2-gpus
timeout_in_minutes: 25
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
+10 -10
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Distributed Comm Ops
key: distributed-comm-ops
timeout_in_minutes: 25
timeout_in_minutes: 20
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -18,7 +18,7 @@ steps:
- label: Distributed DP Tests (2 GPUs)
key: distributed-dp-tests-2-gpus
timeout_in_minutes: 35
timeout_in_minutes: 20
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -55,7 +55,7 @@ steps:
- label: Distributed Compile + RPC Tests (2 GPUs)
key: distributed-compile-rpc-tests-2-gpus
timeout_in_minutes: 65
timeout_in_minutes: 20
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -78,7 +78,7 @@ steps:
- label: Distributed Torchrun + Shutdown Tests (2 GPUs)
key: distributed-torchrun-shutdown-tests-2-gpus
timeout_in_minutes: 30
timeout_in_minutes: 20
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -133,7 +133,7 @@ steps:
- label: Distributed DP Tests (4 GPUs)
key: distributed-dp-tests-4-gpus
timeout_in_minutes: 45
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -154,7 +154,7 @@ steps:
- label: Distributed Compile + Comm (4 GPUs)
key: distributed-compile-comm-4-gpus
timeout_in_minutes: 70
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -176,7 +176,7 @@ steps:
- label: Distributed Tests (8xH100)
key: distributed-tests-8xh100
timeout_in_minutes: 20
timeout_in_minutes: 10
device: h100
num_devices: 8
working_dir: "/vllm-workspace/tests"
@@ -212,7 +212,7 @@ steps:
- label: Distributed Tests (2xH100-2xMI300)
key: distributed-tests-2xh100-2xmi300
timeout_in_minutes: 30
timeout_in_minutes: 15
device: h100
optional: true
working_dir: "/vllm-workspace/"
@@ -259,7 +259,7 @@ steps:
- label: Pipeline + Context Parallelism (4 GPUs)
key: pipeline-context-parallelism-4-gpus
timeout_in_minutes: 55
timeout_in_minutes: 60
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -274,7 +274,7 @@ steps:
- label: RayExecutorV2 (4 GPUs)
key: rayexecutorv2-4-gpus
timeout_in_minutes: 45
timeout_in_minutes: 60
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
+1 -1
View File
@@ -3,7 +3,7 @@ depends_on:
- image-build-cpu
steps:
- label: Docker Build Metadata
timeout_in_minutes: 20
timeout_in_minutes: 10
device: cpu-small
source_file_dependencies:
- .buildkite/release-pipeline.yaml
+5 -5
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: DeepSeek V2-Lite Sync EPLB Accuracy (4xH100)
key: deepseek-v2-lite-sync-eplb-accuracy-4xh100
timeout_in_minutes: 25
timeout_in_minutes: 60
device: h100
optional: true
num_devices: 4
@@ -14,7 +14,7 @@ steps:
- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (4xH100)
key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-4xh100
timeout_in_minutes: 25
timeout_in_minutes: 60
device: h100
optional: true
num_devices: 4
@@ -24,7 +24,7 @@ steps:
- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (2xB200)
key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-2xb200
timeout_in_minutes: 20
timeout_in_minutes: 60
device: b200-k8s
optional: true
num_devices: 2
@@ -34,7 +34,7 @@ steps:
- label: Qwen3-30B-A3B-FP8 DP4 Async EPLB Accuracy
key: qwen3-30b-a3b-fp8-dp4-async-eplb-accuracy
timeout_in_minutes: 25
timeout_in_minutes: 60
device: h100
optional: true
num_devices: 4
@@ -44,7 +44,7 @@ steps:
- label: DeepSeek V2-Lite Prefetch Offload Accuracy (H100)
key: deepseek-v2-lite-prefetch-offload-accuracy-h100
timeout_in_minutes: 20
timeout_in_minutes: 60
device: h100
optional: true
num_devices: 1
+13 -13
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Engine
key: engine
timeout_in_minutes: 30
timeout_in_minutes: 15
device: h200_18gb
source_file_dependencies:
- vllm/compilation/
@@ -29,13 +29,13 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 50
timeout_in_minutes: 60
depends_on:
- image-build-amd
- label: Engine (1 GPU)
key: engine-1-gpu
timeout_in_minutes: 45
timeout_in_minutes: 30
source_file_dependencies:
- vllm/v1/engine/
- tests/v1/engine/
@@ -45,13 +45,13 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 55
timeout_in_minutes: 40
depends_on:
- image-build-amd
- label: e2e Scheduling (1 GPU)
key: e2e-scheduling-1-gpu
timeout_in_minutes: 35
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/v1/
@@ -60,15 +60,15 @@ steps:
- pytest -v -s v1/e2e/general/test_async_scheduling.py
mirror:
amd:
device: mi325_1
timeout_in_minutes: 70
device: mi250_1
timeout_in_minutes: 60
depends_on:
- image-build-amd
- label: e2e Core (1 GPU)
device: h200_35gb
key: e2e-core-1-gpu
timeout_in_minutes: 40
timeout_in_minutes: 30
source_file_dependencies:
- vllm/v1/
- tests/v1/e2e/general/
@@ -76,8 +76,8 @@ steps:
- pytest -v -s v1/e2e/general --ignore v1/e2e/general/test_async_scheduling.py
mirror:
amd:
device: mi325_1
timeout_in_minutes: 60
device: mi250_1
timeout_in_minutes: 35
depends_on:
- image-build-amd
source_file_dependencies:
@@ -87,7 +87,7 @@ steps:
- label: V1 e2e (2 GPUs)
key: v1-e2e-2-gpus
timeout_in_minutes: 25 # TODO: Fix timeout after we have more confidence in the test stability
timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
optional: true
num_devices: 2
source_file_dependencies:
@@ -120,7 +120,7 @@ steps:
- label: V1 e2e (4 GPUs)
key: v1-e2e-4-gpus
timeout_in_minutes: 20 # TODO: Fix timeout after we have more confidence in the test stability
timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
optional: true
num_devices: 4
source_file_dependencies:
@@ -148,7 +148,7 @@ steps:
- label: V1 e2e (4xH100)
key: v1-e2e-4xh100
timeout_in_minutes: 35
timeout_in_minutes: 60
device: h100
num_devices: 4
optional: true
+10 -10
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Entrypoints Unit Tests
key: entrypoints-unit-tests
timeout_in_minutes: 25
timeout_in_minutes: 10
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/entrypoints
@@ -16,7 +16,7 @@ steps:
- label: Entrypoints Integration (LLM)
key: entrypoints-integration-llm
timeout_in_minutes: 60
timeout_in_minutes: 40
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -37,7 +37,7 @@ steps:
- label: Entrypoints Integration (API Server)
key: entrypoints-integration-api-server
device: h200_35gb
timeout_in_minutes: 50
timeout_in_minutes: 130
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -56,7 +56,7 @@ steps:
- label: Entrypoints Integration (API Server OpenAI - Part 1)
key: entrypoints-integration-api-server-openai-part-1
timeout_in_minutes: 45
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -68,13 +68,13 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 65
timeout_in_minutes: 80
depends_on:
- image-build-amd
- label: Entrypoints Integration (API Server OpenAI - Part 2)
key: entrypoints-integration-api-server-openai-part-2
timeout_in_minutes: 45
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -109,7 +109,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 65
timeout_in_minutes: 60
depends_on:
- image-build-amd
@@ -126,7 +126,7 @@ steps:
- label: Entrypoints Integration (Speech to Text)
device: h200_35gb
key: entrypoints-integration-speech_to_text
timeout_in_minutes: 45
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -138,7 +138,7 @@ steps:
- label: Entrypoints Integration (Multimodal)
device: h200_35gb
key: entrypoints-integration-multimodal
timeout_in_minutes: 45
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -160,7 +160,7 @@ steps:
- label: OpenAI API Correctness
key: openai-api-correctness
timeout_in_minutes: 20
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- csrc/
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: EPLB Algorithm
key: eplb-algorithm
timeout_in_minutes: 20
timeout_in_minutes: 15
device: h200_18gb
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -27,7 +27,7 @@ steps:
- label: EPLB Execution # 17min
key: eplb-execution
timeout_in_minutes: 25
timeout_in_minutes: 27
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -39,7 +39,7 @@ steps:
- label: Elastic EP Scaling Test
key: elastic-ep-scaling-test
timeout_in_minutes: 30
timeout_in_minutes: 20
device: h100
working_dir: "/vllm-workspace/tests"
num_devices: 4
+19 -19
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: vLLM IR Tests
key: vllm-ir-tests
timeout_in_minutes: 35
timeout_in_minutes: 10
device: h200_18gb
working_dir: "/vllm-workspace/"
source_file_dependencies:
@@ -16,7 +16,7 @@ steps:
- label: Kernels Core Operation Test
key: kernels-core-operation-test
timeout_in_minutes: 120
timeout_in_minutes: 75
source_file_dependencies:
- csrc/
- tests/kernels/core
@@ -27,7 +27,7 @@ steps:
- label: Kernels MiniMax Reduce RMS Test (2 GPUs)
key: kernels-minimax-reduce-rms-test-2-gpus
timeout_in_minutes: 20
timeout_in_minutes: 15
num_devices: 2
device: h100
source_file_dependencies:
@@ -41,7 +41,7 @@ steps:
- label: Deepseek V4 Kernel Test (H100)
key: deepseek-v4-kernel-test-h100
timeout_in_minutes: 30
timeout_in_minutes: 15
device: h100
source_file_dependencies:
- csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
@@ -54,7 +54,7 @@ steps:
- label: Deepseek V4 Kernel Test (B200)
key: deepseek-v4-kernel-test-b200
timeout_in_minutes: 20
timeout_in_minutes: 15
device: b200-k8s
source_file_dependencies:
- csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
@@ -65,7 +65,7 @@ steps:
- label: Kernels Attention Test %N
key: kernels-attention-test
timeout_in_minutes: 65
timeout_in_minutes: 35
source_file_dependencies:
- csrc/attention/
- vllm/v1/attention
@@ -79,7 +79,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 90
timeout_in_minutes: 55
depends_on:
- image-build-amd
source_file_dependencies:
@@ -106,7 +106,7 @@ steps:
- label: Kernels Quantization Test %N
key: kernels-quantization-test
timeout_in_minutes: 60
timeout_in_minutes: 90
source_file_dependencies:
- csrc/quantization/
- vllm/model_executor/layers/quantization
@@ -131,7 +131,7 @@ steps:
- label: Kernels MoE Test %N
key: kernels-moe-test
timeout_in_minutes: 50
timeout_in_minutes: 25
source_file_dependencies:
- csrc/quantization/cutlass_w8a8/moe/
- csrc/moe/
@@ -147,7 +147,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 65
timeout_in_minutes: 50
source_file_dependencies:
- csrc/quantization/cutlass_w8a8/moe/
- csrc/moe/
@@ -163,7 +163,7 @@ steps:
- label: Kernels Mamba Test
key: kernels-mamba-test
timeout_in_minutes: 40
timeout_in_minutes: 45
source_file_dependencies:
- csrc/mamba/
- tests/kernels/mamba
@@ -172,7 +172,7 @@ steps:
- pytest -v -s kernels/mamba
- label: Kernels KDA Test
timeout_in_minutes: 25
timeout_in_minutes: 20
device: h200_18gb
source_file_dependencies:
- vllm/model_executor/layers/fla/ops/kda.py
@@ -184,7 +184,7 @@ steps:
- label: Kernels DeepGEMM Test (H100)
key: kernels-deepgemm-test-h100
timeout_in_minutes: 35
timeout_in_minutes: 45
device: h100
num_devices: 1
source_file_dependencies:
@@ -211,7 +211,7 @@ steps:
- label: Kernels (B200)
key: kernels-b200
timeout_in_minutes: 80
timeout_in_minutes: 30
working_dir: "/vllm-workspace/"
device: b200-k8s
# optional: true
@@ -264,7 +264,7 @@ steps:
- label: Kernels Helion Test
key: kernels-helion-test
timeout_in_minutes: 115
timeout_in_minutes: 30
device: h100
source_file_dependencies:
- vllm/utils/import_utils.py
@@ -276,7 +276,7 @@ steps:
- label: Kernels FP8 MoE Test (1xH100)
key: kernels-fp8-moe-test-1xh100
timeout_in_minutes: 40
timeout_in_minutes: 90
device: h100
num_devices: 1
optional: true
@@ -293,7 +293,7 @@ steps:
- label: Kernels FP8 MoE Test (2xH100)
key: kernels-fp8-moe-test-2xh100
timeout_in_minutes: 45
timeout_in_minutes: 90
device: h100
num_devices: 2
optional: true
@@ -303,7 +303,7 @@ steps:
- label: Kernels Fp4 MoE Test (B200)
key: kernels-fp4-moe-test-b200
timeout_in_minutes: 25
timeout_in_minutes: 60
device: b200-k8s
num_devices: 1
optional: true
@@ -316,7 +316,7 @@ steps:
- label: Kernels FusedMoE Layer Test (2 H100s)
key: kernels-fusedmoe-layer-test-2-h100s
timeout_in_minutes: 30
timeout_in_minutes: 90
device: h100
num_devices: 2
source_file_dependencies:
+23 -112
View File
@@ -5,7 +5,7 @@ steps:
- label: LM Eval Small Models
device: h200_35gb
key: lm-eval-small-models
timeout_in_minutes: 45
timeout_in_minutes: 75
source_file_dependencies:
- csrc/
- vllm/model_executor/layers/quantization
@@ -56,7 +56,7 @@ steps:
- label: LM Eval Small Models (1xB200)
key: lm-eval-small-models-1xb200
timeout_in_minutes: 50
timeout_in_minutes: 120
device: b200-k8s
optional: true
source_file_dependencies:
@@ -80,7 +80,7 @@ steps:
- label: LM Eval Large Models EP (2xB200)
key: lm-eval-large-models-ep-2xb200
timeout_in_minutes: 60
timeout_in_minutes: 120
device: b200-k8s
optional: true
num_devices: 2
@@ -92,7 +92,7 @@ steps:
- label: LM Eval Qwen3.5 Models (2xB200)
key: lm-eval-qwen3-5-models-2xb200
timeout_in_minutes: 45
timeout_in_minutes: 120
device: b200-k8s
optional: true
num_devices: 2
@@ -109,7 +109,7 @@ steps:
- label: LM Eval Large Models (8xH200)
key: lm-eval-large-models-8xh200
timeout_in_minutes: 50
timeout_in_minutes: 60
device: h200
optional: true
num_devices: 8
@@ -118,7 +118,7 @@ steps:
mirror:
amd:
device: mi300_8
timeout_in_minutes: 60
timeout_in_minutes: 180
depends_on:
- image-build-amd
commands:
@@ -150,9 +150,9 @@ steps:
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/moe-refactor-dp-ep/config-b200.txt
- label: LM Eval Humming f16 (A100 - TEMPORARY)
key: lm-eval-humming-f16-a100
timeout_in_minutes: 75
- label: LM Eval Humming (A100 - TEMPORARY)
key: lm-eval-humming-a100
timeout_in_minutes: 30
device: a100
optional: true
num_devices: 1
@@ -160,29 +160,13 @@ steps:
- vllm/model_executor/layers/quantization/humming.py
- vllm/model_executor/layers/quantization/utils/humming_utils.py
- vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
- vllm/model_executor/layers/fused_moe/oracle/
- vllm/model_executor/kernels/linear/
- vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config.txt
- label: LM Eval Humming Act int8 (A100 - TEMPORARY)
key: lm-eval-humming-act-a100
timeout_in_minutes: 45
device: a100
optional: true
num_devices: 1
source_file_dependencies:
- vllm/model_executor/layers/quantization/humming.py
- vllm/model_executor/layers/quantization/utils/humming_utils.py
- vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
- vllm/model_executor/layers/fused_moe/oracle/
- vllm/model_executor/kernels/linear/
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-int8.txt
- label: LM Eval Humming f16 (H100 - TEMPORARY)
key: lm-eval-humming-f16-h100
timeout_in_minutes: 70
- label: LM Eval Humming (H100 - TEMPORARY)
key: lm-eval-humming-h100
timeout_in_minutes: 30
device: h100
optional: true
num_devices: 1
@@ -190,30 +174,14 @@ steps:
- vllm/model_executor/layers/quantization/humming.py
- vllm/model_executor/layers/quantization/utils/humming_utils.py
- vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
- vllm/model_executor/layers/fused_moe/oracle/
- vllm/model_executor/kernels/linear/
- vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config.txt
- label: LM Eval Humming Act fp8/int8 (H100 - TEMPORARY)
key: lm-eval-humming-act-h100
timeout_in_minutes: 70
device: h100
optional: true
num_devices: 1
source_file_dependencies:
- vllm/model_executor/layers/quantization/humming.py
- vllm/model_executor/layers/quantization/utils/humming_utils.py
- vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
- vllm/model_executor/layers/fused_moe/oracle/
- vllm/model_executor/kernels/linear/
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-fp8.txt
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-int8.txt
- label: LM Eval Humming f16 (B200 - TEMPORARY)
key: lm-eval-humming-f16-b200
timeout_in_minutes: 50
- label: LM Eval Humming (B200 - TEMPORARY)
key: lm-eval-humming-b200
timeout_in_minutes: 30
device: b200-k8s
optional: true
num_devices: 1
@@ -221,30 +189,14 @@ steps:
- vllm/model_executor/layers/quantization/humming.py
- vllm/model_executor/layers/quantization/utils/humming_utils.py
- vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
- vllm/model_executor/layers/fused_moe/oracle/
- vllm/model_executor/kernels/linear/
- vllm/model_executor/layers/fused_moe/oracle/mxfp4.py
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config.txt
- label: LM Eval Humming Act fp8/int8 (B200 - TEMPORARY)
key: lm-eval-humming-act-b200
timeout_in_minutes: 50
device: b200-k8s
optional: true
num_devices: 1
source_file_dependencies:
- vllm/model_executor/layers/quantization/humming.py
- vllm/model_executor/layers/quantization/utils/humming_utils.py
- vllm/model_executor/layers/fused_moe/experts/fused_humming_moe.py
- vllm/model_executor/layers/fused_moe/oracle/
- vllm/model_executor/kernels/linear/
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-fp8.txt
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/humming/config-act-int8.txt
- label: LM Eval TurboQuant KV Cache
key: lm-eval-turboquant-kv-cache
timeout_in_minutes: 55
timeout_in_minutes: 75
device: h200_18gb
source_file_dependencies:
- vllm/model_executor/layers/quantization/turboquant/
@@ -256,7 +208,7 @@ steps:
- label: GPQA Eval (GPT-OSS) (2xH100)
key: gpqa-eval-gpt-oss-2xh100
timeout_in_minutes: 35
timeout_in_minutes: 120
device: h100
optional: true
num_devices: 2
@@ -270,7 +222,7 @@ steps:
- label: GPQA Eval (GPT-OSS) (2xB200)
key: gpqa-eval-gpt-oss-2xb200
timeout_in_minutes: 30
timeout_in_minutes: 120
device: b200-k8s
optional: true
num_devices: 2
@@ -284,7 +236,7 @@ steps:
- label: GPQA Eval (GPT-OSS) (DGX Spark)
key: gpqa-eval-gpt-oss-spark
timeout_in_minutes: 35
timeout_in_minutes: 120
device: dgx-spark
optional: true
num_devices: 1
@@ -298,50 +250,9 @@ steps:
- uv pip install --system 'gpt-oss[eval]==0.0.5'
- pytest -s -v evals/gpt_oss/test_gpqa_correctness.py --config-list-file=configs/models-spark.txt
- label: LM Eval KV-Offload (1xH200)
key: kv-offload-small
timeout_in_minutes: 30
device: h200_35gb
source_file_dependencies:
- vllm/distributed/kv_transfer/kv_connector/v1/offloading/
- vllm/distributed/kv_transfer/kv_connector/v1/simple_cpu_offload_connector.py
- vllm/v1/kv_offload/
- vllm/v1/simple_kv_offload/
- tests/evals/gsm8k/test_gsm8k_offloading.py
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_offloading.py -k "nemotron-h-8b or gemma-4-e4b-it"
- label: LM Eval KV-Offload (2xH100)
key: kv-offload-medium
timeout_in_minutes: 30
device: h100
num_devices: 2
source_file_dependencies:
- vllm/distributed/kv_transfer/kv_connector/v1/offloading/
- vllm/distributed/kv_transfer/kv_connector/v1/simple_cpu_offload_connector.py
- vllm/v1/kv_offload/
- vllm/v1/simple_kv_offload/
- tests/evals/gsm8k/test_gsm8k_offloading.py
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_offloading.py -k "qwen3.5-35b"
- label: LM Eval KV-Offload (4xH100)
key: kv-offload-large
timeout_in_minutes: 40
device: h100
num_devices: 4
source_file_dependencies:
- vllm/distributed/kv_transfer/kv_connector/v1/offloading/
- vllm/distributed/kv_transfer/kv_connector/v1/simple_cpu_offload_connector.py
- vllm/v1/kv_offload/
- vllm/v1/simple_kv_offload/
- tests/evals/gsm8k/test_gsm8k_offloading.py
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_offloading.py -k "deepseek-v4-flash"
- label: MRCR Eval Small Models
device: h200_35gb
timeout_in_minutes: 25
timeout_in_minutes: 30
source_file_dependencies:
- tests/evals/mrcr/
commands:
+3 -3
View File
@@ -5,7 +5,7 @@ steps:
- label: LoRA %N
device: h200_35gb
key: lora
timeout_in_minutes: 40
timeout_in_minutes: 30
source_file_dependencies:
- vllm/lora
- tests/lora
@@ -16,7 +16,7 @@ steps:
amd:
device: mi325_1
working_dir: "/vllm-workspace/tests"
timeout_in_minutes: 65
timeout_in_minutes: 60
source_file_dependencies:
- vllm/lora
- tests/lora
@@ -27,7 +27,7 @@ steps:
- label: LoRA TP (Distributed)
key: lora-tp-distributed
timeout_in_minutes: 60
timeout_in_minutes: 30
num_devices: 4
source_file_dependencies:
- vllm/lora
+15 -17
View File
@@ -5,7 +5,7 @@ steps:
- label: V1 Spec Decode
device: h200_35gb
key: v1-spec-decode
timeout_in_minutes: 40
timeout_in_minutes: 30
source_file_dependencies:
- vllm/config/
- vllm/distributed/
@@ -24,13 +24,13 @@ steps:
mirror:
amd:
device: mi300_1
timeout_in_minutes: 75
timeout_in_minutes: 65
depends_on:
- image-build-amd
- label: V1 Sample + Logits
key: v1-sample-logits
timeout_in_minutes: 45
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/config/
@@ -64,7 +64,7 @@ steps:
- label: V1 Core + KV + Metrics
key: v1-core-kv-metrics
timeout_in_minutes: 60
timeout_in_minutes: 30
source_file_dependencies:
- vllm/config/
- vllm/distributed/
@@ -108,7 +108,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 75
timeout_in_minutes: 60
depends_on:
- image-build-amd
@@ -172,7 +172,7 @@ steps:
- label: Regression
key: regression
timeout_in_minutes: 30
timeout_in_minutes: 20
device: h200_18gb
source_file_dependencies:
- vllm/config/
@@ -195,7 +195,7 @@ steps:
- label: Examples
device: h200_35gb
key: examples
timeout_in_minutes: 40
timeout_in_minutes: 45
working_dir: "/vllm-workspace/examples"
source_file_dependencies:
- vllm/entrypoints
@@ -237,7 +237,7 @@ steps:
- label: Metrics, Tracing (2 GPUs)
key: metrics-tracing-2-gpus
timeout_in_minutes: 25
timeout_in_minutes: 20
num_devices: 2
source_file_dependencies:
- vllm/config/
@@ -281,7 +281,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 45
timeout_in_minutes: 20
depends_on:
- image-build-amd
source_file_dependencies:
@@ -292,7 +292,7 @@ steps:
- label: Async Engine, Inputs, Utils, Worker
device: h200_35gb
key: async-engine-inputs-utils-worker
timeout_in_minutes: 25
timeout_in_minutes: 50
source_file_dependencies:
- vllm/assets/
- vllm/config/
@@ -319,7 +319,7 @@ steps:
key: async-engine-inputs-utils-worker-config-cpu
depends_on:
- image-build-cpu
timeout_in_minutes: 65
timeout_in_minutes: 30
source_file_dependencies:
- vllm/assets/
- vllm/config/
@@ -351,7 +351,6 @@ steps:
- tests/test_outputs.py
- tests/test_pooling_params.py
- tests/test_ray_env.py
- tests/test_sampling_params.py
- tests/multimodal
- tests/renderers
- tests/standalone_tests/lazy_imports.py
@@ -369,7 +368,6 @@ steps:
- pytest -v -s test_outputs.py
- pytest -v -s test_pooling_params.py
- pytest -v -s test_ray_env.py
- pytest -v -s test_sampling_params.py
- pytest -v -s -m 'cpu_test' multimodal
- pytest -v -s renderers
- pytest -v -s reasoning
@@ -381,7 +379,7 @@ steps:
- label: Batch Invariance (A100)
key: batch-invariance-a100
timeout_in_minutes: 40
timeout_in_minutes: 30
device: a100
source_file_dependencies:
- vllm/v1/attention
@@ -395,7 +393,7 @@ steps:
- label: Batch Invariance (H100)
key: batch-invariance-h100
timeout_in_minutes: 40
timeout_in_minutes: 30
device: h100
source_file_dependencies:
- vllm/v1/attention
@@ -411,7 +409,7 @@ steps:
- label: Batch Invariance (B200)
key: batch-invariance-b200
timeout_in_minutes: 35
timeout_in_minutes: 30
device: b200-k8s
source_file_dependencies:
- vllm/v1/attention
@@ -430,7 +428,7 @@ steps:
- label: Acceptance Length Test (Large Models) # optional
device: h200_35gb
key: acceptance-length-test-large-models
timeout_in_minutes: 20
timeout_in_minutes: 25
gpu: h100
optional: true
num_gpus: 1
+1 -1
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Model Executor
key: model-executor
timeout_in_minutes: 45
timeout_in_minutes: 35
source_file_dependencies:
- vllm/engine/arg_utils.py
- vllm/config/model.py
+4 -4
View File
@@ -5,7 +5,7 @@ steps:
- label: Model Runner V2 Core Tests
device: h200_35gb
key: model-runner-v2-core-tests
timeout_in_minutes: 35
timeout_in_minutes: 45
source_file_dependencies:
- vllm/v1/worker/gpu/
- vllm/v1/worker/gpu_worker.py
@@ -27,7 +27,7 @@ steps:
- label: Model Runner V2 Examples
device: h200_35gb
key: model-runner-v2-examples
timeout_in_minutes: 35
timeout_in_minutes: 45
working_dir: "/vllm-workspace/examples"
source_file_dependencies:
- vllm/v1/worker/gpu/
@@ -63,7 +63,7 @@ steps:
- label: Model Runner V2 Distributed (2 GPUs)
key: model-runner-v2-distributed-2-gpus
timeout_in_minutes: 30
timeout_in_minutes: 45
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -84,7 +84,7 @@ steps:
- label: Model Runner V2 Pipeline Parallelism (4 GPUs)
key: model-runner-v2-pipeline-parallelism-4-gpus
timeout_in_minutes: 50
timeout_in_minutes: 60
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
+7 -8
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Basic Models Tests (Initialization)
key: basic-models-tests-initialization
timeout_in_minutes: 25
timeout_in_minutes: 45
device: h200_18gb
source_file_dependencies:
- vllm/
@@ -17,7 +17,7 @@ steps:
- label: Basic Models Tests (Extra Initialization) %N
device: h200_35gb
key: basic-models-tests-extra-initialization
timeout_in_minutes: 100
timeout_in_minutes: 45
source_file_dependencies:
- vllm/model_executor/models/
- tests/models/test_initialization.py
@@ -32,14 +32,14 @@ steps:
- label: Basic Models Tests (Other)
device: h200_35gb
key: basic-models-tests-other
timeout_in_minutes: 35
timeout_in_minutes: 45
source_file_dependencies:
- vllm/
- tests/models/test_terratorch.py
- tests/models/transformers/test_backend.py
- tests/models/test_transformers.py
- tests/models/test_registry.py
commands:
- pytest -v -s models/test_terratorch.py models/transformers/test_backend.py models/test_registry.py
- pytest -v -s models/test_terratorch.py models/test_transformers.py models/test_registry.py
mirror:
amd:
device: mi325_1
@@ -50,12 +50,11 @@ steps:
key: basic-models-test-other-cpu
depends_on:
- image-build-cpu
timeout_in_minutes: 20
timeout_in_minutes: 10
source_file_dependencies:
- vllm/
- tests/models/test_utils.py
- tests/models/test_vision.py
- tests/models/transformers/fusers/
device: cpu-small
commands:
- pytest -v -s models/test_utils.py models/test_vision.py models/transformers/fusers/
- pytest -v -s models/test_utils.py models/test_vision.py
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Distributed Model Tests (2 GPUs)
key: distributed-model-tests-2-gpus
timeout_in_minutes: 60
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -17,7 +17,7 @@ steps:
- TARGET_TEST_SUITE=L4 pytest basic_correctness/ -v -s -m 'distributed(num_gpus=2)'
- CUDA_VISIBLE_DEVICES=0,1 pytest -v -s model_executor/model_loader/test_sharded_state_loader.py -m '(not slow_test)'
# Avoid importing model tests that cause CUDA reinitialization error
- pytest models/transformers/test_backend.py -v -s -m 'distributed(num_gpus=2)'
- pytest models/test_transformers.py -v -s -m 'distributed(num_gpus=2)'
- pytest models/language -v -s -m 'distributed(num_gpus=2)'
- pytest models/multimodal/generation/test_phi4siglip.py -v -s -m 'distributed(num_gpus=2)'
- pytest models/multimodal -v -s -m 'distributed(num_gpus=2)' --ignore models/multimodal/generation/test_whisper.py --ignore models/multimodal/generation/test_phi4siglip.py
+8 -8
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Language Models Tests (Standard)
key: language-models-tests-standard
timeout_in_minutes: 30
timeout_in_minutes: 25
device: h200_18gb
source_file_dependencies:
- vllm/
@@ -21,7 +21,7 @@ steps:
- label: Language Models Tests (Extra Standard) %N
key: language-models-tests-extra-standard
timeout_in_minutes: 40
timeout_in_minutes: 45
source_file_dependencies:
- vllm/model_executor/models/
- tests/models/language/pooling/test_embedding.py
@@ -52,7 +52,7 @@ steps:
- label: Language Models Tests (Hybrid) %N
key: language-models-tests-hybrid
timeout_in_minutes: 65
timeout_in_minutes: 75
source_file_dependencies:
- vllm/
- tests/models/language/generation
@@ -67,7 +67,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 70
timeout_in_minutes: 90
depends_on:
- image-build-amd
commands:
@@ -78,7 +78,7 @@ steps:
- label: Language Models Test (Extended Generation) # 80min
device: h200_35gb
key: language-models-test-extended-generation
timeout_in_minutes: 65
timeout_in_minutes: 110
optional: true
source_file_dependencies:
- vllm/
@@ -92,7 +92,7 @@ steps:
- label: Language Models Test (PPL)
key: language-models-test-ppl
timeout_in_minutes: 30
timeout_in_minutes: 110
device: h200_18gb
optional: true
source_file_dependencies:
@@ -104,7 +104,7 @@ steps:
- label: Language Models Test (Extended Pooling) # 36min
device: h200_35gb
key: language-models-test-extended-pooling
timeout_in_minutes: 70
timeout_in_minutes: 50
optional: true
source_file_dependencies:
- vllm/
@@ -120,7 +120,7 @@ steps:
- label: Language Models Test (MTEB)
key: language-models-test-mteb
timeout_in_minutes: 45
timeout_in_minutes: 110
device: h200_18gb
optional: true
source_file_dependencies:
+8 -9
View File
@@ -20,14 +20,13 @@ steps:
- label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
key: multi-modal-models-standard-2-qwen3-gemma
timeout_in_minutes: 50
timeout_in_minutes: 45
device: h200_18gb
source_file_dependencies:
- vllm/
- tests/models/multimodal
commands:
- pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "qwen3 or gemma"
- pytest -v -s models/multimodal/generation/test_mm_prefix_lm.py -m core_model
- pytest -v -s models/multimodal/generation/test_qwen2_5_vl.py -m core_model
mirror:
amd:
@@ -38,7 +37,7 @@ steps:
- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
device: h200_35gb
key: multi-modal-models-standard-3-llava-qwen2-vl
timeout_in_minutes: 40
timeout_in_minutes: 45
source_file_dependencies:
- vllm/
- tests/models/multimodal
@@ -54,12 +53,12 @@ steps:
- label: "Multi-Modal Models (Standard) 4: other + whisper"
device: h200_35gb
key: multi-modal-models-standard-4-other-whisper
timeout_in_minutes: 50
timeout_in_minutes: 45
source_file_dependencies:
- vllm/
- tests/models/multimodal
commands:
- pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py --ignore models/multimodal/generation/test_mm_prefix_lm.py --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/generation/test_vit_cudagraph.py --ignore models/multimodal/processing
- pytest -v -s models/multimodal -m core_model --ignore models/multimodal/generation/test_common.py --ignore models/multimodal/generation/test_ultravox.py --ignore models/multimodal/generation/test_qwen2_5_vl.py --ignore models/multimodal/generation/test_qwen2_vl.py --ignore models/multimodal/generation/test_whisper.py --ignore models/multimodal/generation/test_memory_leak.py --ignore models/multimodal/generation/test_vit_cudagraph.py --ignore models/multimodal/processing
- pytest -v -s models/multimodal/generation/test_vit_cudagraph.py -m core_model
- pytest models/multimodal/generation/test_memory_leak.py -m core_model
- cd .. && VLLM_WORKER_MULTIPROC_METHOD=spawn pytest -v -s tests/models/multimodal/generation/test_whisper.py -m core_model # Otherwise, mp_method="spawn" doesn't work
@@ -73,7 +72,7 @@ steps:
key: multi-modal-processor-cpu
depends_on:
- image-build-cpu
timeout_in_minutes: 125
timeout_in_minutes: 60
source_file_dependencies:
- vllm/
- tests/models/multimodal
@@ -84,7 +83,7 @@ steps:
- label: Multi-Modal Processor # 44min
key: multi-modal-processor
timeout_in_minutes: 65
timeout_in_minutes: 60
device: h200_18gb
source_file_dependencies:
- vllm/
@@ -96,7 +95,7 @@ steps:
- label: Multi-Modal Accuracy Eval (Small Models) # 50min
device: h200_35gb
key: multi-modal-accuracy-eval-small-models
timeout_in_minutes: 30
timeout_in_minutes: 70
working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
source_file_dependencies:
- vllm/multimodal/
@@ -164,7 +163,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 75
timeout_in_minutes: 60
depends_on:
- image-build-amd
source_file_dependencies:
+1 -6
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Plugin Tests (2 GPUs)
key: plugin-tests-2-gpus
timeout_in_minutes: 35
timeout_in_minutes: 60
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -37,11 +37,6 @@ steps:
- pytest -v -s plugins_tests/test_stats_logger_plugins.py
- pip uninstall dummy_stat_logger -y
# end stat_logger plugins test
# begin endpoint plugins test
- pip install -e ./plugins/vllm_add_dummy_endpoint_plugin
- pytest -v -s plugins_tests/test_endpoint_plugins.py
- pip uninstall vllm_add_dummy_endpoint_plugin -y
# end endpoint plugins test
# other tests continue here:
- pytest -v -s plugins_tests/test_scheduler_plugins.py
- pip install -e ./plugins/vllm_add_dummy_model
+5 -5
View File
@@ -5,7 +5,7 @@ steps:
- label: PyTorch Compilation Unit Tests
device: h200_35gb
key: pytorch-compilation-unit-tests
timeout_in_minutes: 90
timeout_in_minutes: 10
source_file_dependencies:
- vllm/__init__.py
- vllm/_aiter_ops.py
@@ -78,7 +78,7 @@ steps:
- label: PyTorch Compilation Passes Unit Tests
key: pytorch-compilation-passes-unit-tests
timeout_in_minutes: 45
timeout_in_minutes: 20
source_file_dependencies:
- vllm/__init__.py
- vllm/_aiter_ops.py
@@ -110,13 +110,13 @@ steps:
mirror:
amd:
device: mi300_1
timeout_in_minutes: 65
timeout_in_minutes: 180
depends_on:
- image-build-amd
- label: PyTorch Fullgraph Smoke Test
key: pytorch-fullgraph-smoke-test
timeout_in_minutes: 60
timeout_in_minutes: 35
source_file_dependencies:
- vllm/__init__.py
- vllm/_aiter_ops.py
@@ -152,7 +152,7 @@ steps:
- label: PyTorch Fullgraph
key: pytorch-fullgraph
timeout_in_minutes: 40
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/__init__.py
+4 -4
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Quantization
key: quantization
timeout_in_minutes: 60
timeout_in_minutes: 90
source_file_dependencies:
- csrc/
- vllm/model_executor/layers/quantization
@@ -23,7 +23,7 @@ steps:
- label: Quantized Fusions
key: quantized-fusions
timeout_in_minutes: 20
timeout_in_minutes: 30
source_file_dependencies:
- tests/fusion
- vllm/model_executor/layers/fusion
@@ -35,7 +35,7 @@ steps:
- label: Quantized MoE Test (B200)
key: quantized-moe-test-b200
timeout_in_minutes: 120
timeout_in_minutes: 60
working_dir: "/vllm-workspace/"
device: b200-k8s
source_file_dependencies:
@@ -53,7 +53,7 @@ steps:
- label: Quantized Models Test
key: quantized-models-test
timeout_in_minutes: 50
timeout_in_minutes: 60
source_file_dependencies:
- vllm/model_executor/layers/quantization
- tests/models/quantization
+13 -20
View File
@@ -3,7 +3,7 @@ depends_on:
- image-build
steps:
- label: Rust Frontend OpenAI Coverage
timeout_in_minutes: 30
timeout_in_minutes: 90
device: h200_18gb
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -15,30 +15,28 @@ steps:
- tests/utils.py
- tests/benchmarks/test_serve_cli.py
- tests/entrypoints/openai/chat_completion/test_chat_completion.py
- tests/entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py
# - tests/entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py
# - tests/entrypoints/openai/completion/test_prompt_validation.py
- tests/entrypoints/openai/completion/test_shutdown.py
- tests/entrypoints/openai/test_return_token_ids.py
- tests/entrypoints/openai/test_uds.py
# - tests/entrypoints/openai/test_return_token_ids.py
# - tests/entrypoints/openai/test_uds.py
- tests/v1/sample/test_logprobs_e2e.py
commands:
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s benchmarks/test_serve_cli.py -k "not insecure and not (test_bench_serve and not test_bench_serve_chat)"
- pytest -v -s entrypoints/openai/chat_completion/test_chat_completion.py -k "not test_invalid_json_schema and not test_invalid_regex"
- pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not multiple"
# - pytest -v -s entrypoints/openai/chat_completion/test_chat_logit_bias_validation.py -k "not invalid"
# - pytest -v -s entrypoints/openai/completion/test_prompt_validation.py -k "not prompt_embeds"
- pytest -v -s entrypoints/openai/completion/test_shutdown.py -k "not engine_failure and not test_abort_timeout_exits_quickly"
# test_comparison streams differently: Rust emits a separate first (prompt_token_ids) chunk and
# finish chunk without logprobs, while the test reads `logprobs.tokens` on every chunk.
- pytest -v -s entrypoints/openai/test_return_token_ids.py -k "not test_comparison"
- pytest -v -s entrypoints/openai/test_uds.py
# - pytest -v -s entrypoints/openai/test_return_token_ids.py
# - pytest -v -s entrypoints/openai/test_uds.py
- pytest -v -s v1/sample/test_logprobs_e2e.py -k "test_prompt_logprobs_e2e_server"
- label: Rust Frontend Serve/Admin Coverage
timeout_in_minutes: 25
timeout_in_minutes: 60
device: h200_18gb
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -47,27 +45,22 @@ steps:
- vllm/entrypoints/serve/
- vllm/v1/engine/
- tests/utils.py
- tests/entrypoints/serve/dev/rpc/test_collective_rpc.py
# - tests/entrypoints/serve/dev/rpc/test_collective_rpc.py
- tests/entrypoints/scale_out/token_in_token_out/test_serving_tokens.py
- tests/entrypoints/serve/instrumentator/test_basic.py
- tests/entrypoints/serve/instrumentator/test_metrics.py
# - tests/entrypoints/serve/dev/test_sleep.py
- tests/entrypoints/serve/tokenize/test_tokenization.py
commands:
- export VLLM_USE_RUST_FRONTEND=1
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
# server_load can be flaky under the Rust frontend; keep it excluded for now.
# - pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
- pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not show_version and not server_load"
# test_generate_logprobs expects Python-style top_logprobs truncation (dedup sampled + cap at max(k, 1)).
- pytest -v -s entrypoints/scale_out/token_in_token_out/test_serving_tokens.py -k "not stream and not lora and not test_generate_logprobs and not stop_string_workflow"
- pytest -v -s entrypoints/serve/instrumentator/test_metrics.py -k "text and not show and not run_batch and not test_metrics_counts and not test_metrics_exist"
# - pytest -v -s entrypoints/serve/dev/test_sleep.py
# /tokenizer_info is not implemented in the Rust frontend (the CLI flag is accepted as a no-op).
- pytest -v -s entrypoints/serve/tokenize/test_tokenization.py -k "not tokenizer_info"
- label: Rust Frontend Core Correctness
timeout_in_minutes: 20
timeout_in_minutes: 30
device: h200_18gb
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -81,7 +74,7 @@ steps:
- pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
- label: Rust Frontend Tool Use
timeout_in_minutes: 25
timeout_in_minutes: 60
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- rust/
@@ -95,7 +88,7 @@ steps:
- pytest -v -s tool_use --ignore=tool_use/mistral --models llama3.2 -k "not test_response_format_with_tool_choice_required and not test_parallel_tool_calls_false and not test_tool_call_and_choice"
- label: Rust Frontend Distributed
timeout_in_minutes: 25
timeout_in_minutes: 30
num_devices: 4
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -4,7 +4,7 @@ steps:
- label: Rust Frontend Cargo Style + Clippy
key: rust-frontend-cargo-style-clippy
depends_on: []
timeout_in_minutes: 20
timeout_in_minutes: 30
device: cpu-medium
no_plugin: true
source_file_dependencies:
@@ -18,7 +18,7 @@ steps:
- label: Rust Frontend Cargo Tests
key: rust-frontend-cargo-tests
depends_on: []
timeout_in_minutes: 20
timeout_in_minutes: 30
device: cpu-medium
no_plugin: true
source_file_dependencies:
+2 -2
View File
@@ -5,7 +5,7 @@ steps:
- label: Samplers Test
device: h200_35gb
key: samplers-test
timeout_in_minutes: 40
timeout_in_minutes: 75
source_file_dependencies:
- vllm/model_executor/layers
- vllm/sampling_metadata.py
@@ -19,7 +19,7 @@ steps:
- VLLM_USE_FLASHINFER_SAMPLER=1 pytest -v -s samplers
mirror:
amd:
device: mi325_1
device: mi250_1
depends_on:
- image-build-amd
commands:
+11 -11
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Spec Decode Eagle
key: spec-decode-eagle
timeout_in_minutes: 25
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/v1/spec_decode/
@@ -15,7 +15,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 60
timeout_in_minutes: 45
depends_on:
- image-build-amd
source_file_dependencies:
@@ -29,7 +29,7 @@ steps:
- label: Spec Decode Eagle Nightly B200
key: spec-decode-eagle-nightly-b200
timeout_in_minutes: 25
timeout_in_minutes: 30
device: b200-k8s
optional: true
source_file_dependencies:
@@ -41,7 +41,7 @@ steps:
- label: Spec Decode Speculators + MTP
key: spec-decode-speculators-mtp
timeout_in_minutes: 20
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/v1/spec_decode/
@@ -82,7 +82,7 @@ steps:
- label: Spec Decode Ngram + Suffix
key: spec-decode-ngram-suffix
timeout_in_minutes: 20
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/v1/spec_decode/
@@ -93,7 +93,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 55
timeout_in_minutes: 65
# TODO(akaratza): Test after Torch >= 2.12 bump
soft_fail: true
depends_on:
@@ -109,7 +109,7 @@ steps:
- label: Spec Decode Draft Model
key: spec-decode-draft-model
timeout_in_minutes: 45
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/v1/spec_decode/
@@ -120,7 +120,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 55
timeout_in_minutes: 50
depends_on:
- image-build-amd
source_file_dependencies:
@@ -134,7 +134,7 @@ steps:
- label: Spec Decode Draft Model Nightly B200
key: spec-decode-draft-model-nightly-b200
timeout_in_minutes: 40
timeout_in_minutes: 30
device: b200-k8s
optional: true
source_file_dependencies:
@@ -146,7 +146,7 @@ steps:
- label: Speculators Correctness
key: speculators-correctness
timeout_in_minutes: 30
timeout_in_minutes: 60
device: h100
optional: true
num_devices: 1
@@ -159,7 +159,7 @@ steps:
- pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test
- label: Spec Decode MTP hybrid (B200)
timeout_in_minutes: 20
timeout_in_minutes: 30
device: b200-k8s
optional: true
source_file_dependencies:
+1 -1
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Weight Loading Multiple GPU # 33min
key: weight-loading-multiple-gpu
timeout_in_minutes: 50
timeout_in_minutes: 45
working_dir: "/vllm-workspace/tests"
num_devices: 2
optional: true
+1 -1
View File
@@ -119,7 +119,7 @@
# Transformers modeling backend
/vllm/model_executor/models/transformers @hmellor
/tests/models/transformers @hmellor
/tests/models/test_transformers.py @hmellor
# Docs
/docs/mkdocs @hmellor
+2 -3
View File
@@ -28,8 +28,7 @@ jobs:
pull_number: context.payload.pull_request.number,
});
const readyLabels = ['ready', 'ready-run-all-tests'];
const hasReadyLabel = pr.labels.some(l => readyLabels.includes(l.name));
const hasReadyLabel = pr.labels.some(l => l.name === 'ready');
const hasVerifiedLabel = pr.labels.some(l => l.name === 'verified');
const { data: mergedPRs } = await github.rest.search.issuesAndPullRequests({
@@ -41,7 +40,7 @@ jobs:
if (hasReadyLabel || hasVerifiedLabel || mergedCount >= 4) {
core.info(`Check passed: verified label=${hasVerifiedLabel}, ready label=${hasReadyLabel}, 4+ merged PRs=${mergedCount >= 4}`);
} else {
core.setFailed(`PR must have the 'verified', 'ready', or 'ready-run-all-tests' label (the ready labels also trigger tests) or the author must have at least 4 merged PRs (found ${mergedCount}).`);
core.setFailed(`PR must have the 'verified' or 'ready' (which also triggers tests) label or the author must have at least 4 merged PRs (found ${mergedCount}).`);
}
pre-commit:
+18 -29
View File
@@ -29,7 +29,6 @@ Do not open one-off PRs for tiny edits (single typo, isolated style change, one
- PR descriptions for AI-assisted work **must** include:
- Why this is not duplicating an existing PR.
- Test commands run and results.
- Model evaluation results when the change affects output, accuracy, or serving.
- Clear statement that AI assistance was used.
### Fail-closed behavior
@@ -67,38 +66,23 @@ VLLM_USE_PRECOMPILED=1 uv pip install -e . --torch-backend=auto
uv pip install -e . --torch-backend=auto
```
### Tests
### Running tests
> Requires [Environment setup](#environment-setup) and [Installing dependencies](#installing-dependencies).
```bash
# Install test dependencies (use cuda.in on non-x86_64):
uv pip install -r requirements/test/cuda.in
# Install test dependencies.
# requirements/test/cuda.txt is pinned to x86_64; on other platforms, use the
# unpinned source file instead:
uv pip install -r requirements/test/cuda.in # resolves for current platform
# Or on x86_64:
uv pip install -r requirements/test/cuda.txt
# Run a specific test file:
# Run a specific test file (use .venv/bin/python directly;
# `source activate` does not persist in non-interactive shells):
.venv/bin/python -m pytest tests/path/to/test_file.py -v
```
When adding tests:
- **Design before you write.** Answer four questions first: what is the module
for, what is its I/O contract, what failure am I guarding against, and what is
the cheapest level that catches it (unit over integration over e2e)?
- **Reuse before create.** Extend existing test files, `conftest.py` fixtures, and
helpers; add a new file only when no nearby suite fits.
- **Test behavior with intent.** Assert observable outcomes through public APIs;
state why in the name or docstring. Skip trivial wiring; flaky tests are worse
than no tests.
- **Keep it minimal.** One behavior per test and the smallest setup that
triggers it; if the test diff dwarfs the code change, cut scope.
- **No one-off kernel benchmarks in `tests/`.** Put kernel perf work in
`benchmarks/kernels/`; prove correctness in existing pytest suites.
- **Run model evals for model-affecting changes.** Search `tests/evals/` or use
`vllm bench` and include results in the PR — do not wait for reviewers to ask.
For model-specific requirements, see
[`docs/contributing/model/tests.md`](docs/contributing/model/tests.md).
### Running linters
> Requires [Environment setup](#environment-setup).
@@ -123,18 +107,23 @@ Use [Google-style docstrings](https://google.github.io/styleguide/pyguide.html#3
### Coding style guidelines
- Match existing code style
- Minimize use of comments. Eliminate comments which are redundant, preferring legible and self-documenting code. When used, keep docstrings and comments brief and direct.
Follow these rules for all code changes in this repository:
- Try to match existing code style.
- Code should be self-documenting and self-explanatory.
- Keep comments and docstrings minimal and concise.
- Assume the reader is familiar with vLLM.
### Commit messages
Add attribution using commit trailers such as `Co-authored-by:` (other projects use `Assisted-by:` or `Generated-by:`):
Add attribution using commit trailers such as `Co-authored-by:` (other projects use `Assisted-by:` or `Generated-by:`). For example:
```text
Your commit message here
Co-authored-by: Agent Name Here
Co-authored-by: GitHub Copilot
Co-authored-by: Claude
Co-authored-by: gemini-code-assist
Signed-off-by: Your Name <your.email@example.com>
```
+3 -12
View File
@@ -70,15 +70,6 @@ endif()
#
set(TORCH_SUPPORTED_VERSION_CUDA "2.11.0")
set(TORCH_SUPPORTED_VERSION_ROCM "2.11.0")
# TORCH_NIGHTLY=1 builds run against unpinned nightly wheels, so the supported-
# version check would always warn. Only treat it as a nightly build when the
# value is exactly "1" (the bootstrap exports TORCH_NIGHTLY=0 by default, which
# must NOT suppress the warning for normal builds).
if (DEFINED ENV{TORCH_NIGHTLY} AND "$ENV{TORCH_NIGHTLY}" STREQUAL "1")
set(TORCH_NIGHTLY_BUILD TRUE)
else()
set(TORCH_NIGHTLY_BUILD FALSE)
endif()
#
# Try to find python package with an executable that exactly matches
@@ -184,7 +175,7 @@ endif()
if (NOT HIP_FOUND AND NOT PYTORCH_FOUND_HIP AND CUDA_FOUND)
set(VLLM_GPU_LANG "CUDA")
if (NOT TORCH_NIGHTLY_BUILD AND NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
if (NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
message(WARNING "Pytorch version ${TORCH_SUPPORTED_VERSION_CUDA} "
"expected for CUDA build, saw ${Torch_VERSION} instead.")
endif()
@@ -197,7 +188,7 @@ elseif(HIP_FOUND OR PYTORCH_FOUND_HIP)
enable_language(HIP)
# ROCm 5.X and 6.X
if (NOT TORCH_NIGHTLY_BUILD AND ROCM_VERSION_DEV_MAJOR GREATER_EQUAL 5 AND
if (ROCM_VERSION_DEV_MAJOR GREATER_EQUAL 5 AND
Torch_VERSION VERSION_LESS ${TORCH_SUPPORTED_VERSION_ROCM})
message(WARNING "Pytorch version >= ${TORCH_SUPPORTED_VERSION_ROCM} "
"expected for ROCm build, saw ${Torch_VERSION} instead.")
@@ -390,7 +381,6 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
"csrc/libtorch_stable/cuda_view.cu"
"csrc/libtorch_stable/cuda_utils_kernels.cu"
"csrc/libtorch_stable/activation_kernels.cu"
"csrc/libtorch_stable/ngram_embedding_kernels.cu"
"csrc/libtorch_stable/quantization/activation_kernels.cu"
"csrc/libtorch_stable/quantization/w8a8/int8/scaled_quant.cu"
"csrc/libtorch_stable/quantization/w8a8/fp8/common.cu"
@@ -410,6 +400,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
"csrc/libtorch_stable/topk.cu"
"csrc/libtorch_stable/mamba/selective_scan_fwd.cu"
"csrc/libtorch_stable/cache_kernels.cu"
"csrc/libtorch_stable/cache_kernels.cu"
"csrc/libtorch_stable/cache_kernels_fused.cu"
"csrc/libtorch_stable/custom_all_reduce.cu"
"csrc/libtorch_stable/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu")
+2 -2
View File
@@ -12,7 +12,7 @@ from dataclasses import dataclass, field
import aiohttp
import huggingface_hub.constants
from tqdm.asyncio import tqdm
from transformers import AutoTokenizer, PythonBackend, TokenizersBackend
from transformers import AutoTokenizer, PreTrainedTokenizer, PreTrainedTokenizerFast
# NOTE(simon): do not import vLLM here so the benchmark script
# can run without vLLM installed.
@@ -609,7 +609,7 @@ def get_tokenizer(
tokenizer_mode: str = "auto",
trust_remote_code: bool = False,
**kwargs,
) -> PythonBackend | TokenizersBackend:
) -> PreTrainedTokenizer | PreTrainedTokenizerFast:
if pretrained_model_name_or_path is not None and not os.path.exists(
pretrained_model_name_or_path
):
+2 -4
View File
@@ -132,10 +132,8 @@ def benchmark_function(
reset_memory_stats()
# Benchmark
start_events = [
torch.cuda.Event(enable_timing=True) for _ in range(benchmark_iters)
]
end_events = [torch.cuda.Event(enable_timing=True) for _ in range(benchmark_iters)]
start_events = [torch.Event(enable_timing=True) for _ in range(benchmark_iters)]
end_events = [torch.Event(enable_timing=True) for _ in range(benchmark_iters)]
for i in range(benchmark_iters):
logits_copy = logits.clone()
@@ -17,7 +17,7 @@ from vllm.model_executor.layers.fused_moe.fused_flydsl_moe import fused_flydsl_m
from vllm.model_executor.layers.quantization.compressed_tensors.compressed_tensors_moe import ( # noqa: E501
compressed_tensors_moe_w4a16_flydsl,
)
from vllm.utils.platform_utils import get_device_name_as_file_name
from vllm.platforms import current_platform
RoutingBuffers = tuple[
torch.Tensor, # sorted_token_ids
@@ -259,7 +259,7 @@ def tune_flydsl_moe_w4a16(
)
us_best = us
tuned_config[str(num_tokens)] = tile_config
device_name = get_device_name_as_file_name()
device_name = current_platform.get_device_name().replace(" ", "_")
tuned_config_file_name = (
f"E={num_experts},N={inter_dim},device_name={device_name},"
f"dtype=int4_w4a16,backend=flydsl.json"
+2 -2
View File
@@ -134,8 +134,8 @@ def benchmark_config(
torch.accelerator.synchronize()
# Benchmark
start = torch.cuda.Event(enable_timing=True)
end = torch.cuda.Event(enable_timing=True)
start = torch.Event(enable_timing=True)
end = torch.Event(enable_timing=True)
start.record()
for _ in range(num_iters):
with override_config(config):
@@ -170,8 +170,8 @@ def benchmark_config(
graph.replay()
torch.accelerator.synchronize()
start = torch.cuda.Event(enable_timing=True)
end = torch.cuda.Event(enable_timing=True)
start = torch.Event(enable_timing=True)
end = torch.Event(enable_timing=True)
latencies: list[float] = []
for _ in range(num_iters):
start.record()
@@ -19,7 +19,6 @@ from vllm.model_executor.layers.quantization.utils.fp8_utils import (
from vllm.platforms import current_platform
from vllm.triton_utils import triton
from vllm.utils.argparse_utils import FlexibleArgumentParser
from vllm.utils.platform_utils import get_device_name_as_file_name
mp.set_start_method("spawn", force=True)
@@ -265,7 +264,7 @@ def save_configs(
input_type="fp8",
) -> None:
os.makedirs(save_path, exist_ok=True)
device_name = get_device_name_as_file_name()
device_name = current_platform.get_device_name().replace(" ", "_")
json_file_name = (
f"N={N},K={K},device_name={device_name},dtype={input_type}_w8a8,"
f"block_shape=[{block_n},{block_k}].json"
-241
View File
@@ -1,241 +0,0 @@
#!/bin/bash
set -eoux pipefail
########################################
# Resolve repo root (IMPORTANT)
########################################
REPO_ROOT="$(pwd)"
cd "$REPO_ROOT"
########################################
# DevPI configuration
########################################
IBM_DEVPI_URL=${IBM_DEVPI_URL:-"https://wheels.developerfirst.ibm.com/ppc64le/linux/+simple/"}
RHOAI_INDEX_URL=${RHOAI_INDEX_URL:-"https://console.redhat.com/api/pypi/public-rhai/rhoai/3.4/cpu-ubi9/simple/"}
########################################
# wheel dir
########################################
WHEEL_DIR=${WHEEL_DIR:-"/tmp/wheels"}
mkdir -p "$WHEEL_DIR"
########################################
# Helpers
########################################
try_install_from_devpi() {
local pkg=$1
uv pip install \
--extra-index-url "${IBM_DEVPI_URL}" \
--index-strategy unsafe-best-match \
--no-build-isolation \
"${pkg}"
}
########################################
# Package Versions
########################################
cd "$REPO_ROOT"
TORCH_VERSION=${TORCH_VERSION:-$(grep -E '^torch==.+==\s*"ppc64le"' requirements/cpu.txt | grep -Eo '\b[0-9\.]+\b' || true)}
TORCH_VERSION=${TORCH_VERSION:-2.11.0}
TORCHVISION_VERSION=${TORCHVISION_VERSION:-0.26.0}
TORCHAUDIO_VERSION=${TORCHAUDIO_VERSION:-${TORCH_VERSION}}
export TORCH_VERSION
export TORCHVISION_VERSION
export TORCHAUDIO_VERSION
export OPENCV_VERSION=${OPENCV_VERSION:-4.13.0.92}
export XGRAMMAR_VERSION=${XGRAMMAR_VERSION:-0.2.1}
########################################
# install system dependencies
########################################
rpm -ivh https://dl.fedoraproject.org/pub/epel/epel-release-latest-9.noarch.rpm || true
microdnf install -y \
python3.12 python3.12-devel python3.12-pip gcc \
git jq gcc-toolset-14 gcc-toolset-14-libatomic-devel \
automake libtool clang-devel openssl-devel \
harfbuzz-devel kmod lcms2-devel libimagequant-devel libjpeg-turbo-devel \
llvm15-devel libraqm-devel libtiff-devel libwebp-devel libxcb-devel \
ninja-build openjpeg2-devel pkgconfig \
tcl-devel tk-devel xsimd-devel zeromq-devel zlib-devel patchelf file openblas openblas-devel protobuf numactl numactl-devel openmpi openmpi-devel
rpm -ivh --nodeps \
https://mirror.stream.centos.org/9-stream/CRB/ppc64le/os/Packages/protobuf-lite-devel-3.14.0-17.el9.ppc64le.rpm
rpm -ivh --nodeps \
https://mirror.stream.centos.org/9-stream/CRB/ppc64le/os/Packages/protobuf-devel-3.14.0-17.el9.ppc64le.rpm
rpm -ivh --nodeps \
https://mirror.stream.centos.org/9-stream/CRB/ppc64le/os/Packages/protobuf-compiler-3.14.0-17.el9.ppc64le.rpm
########################################
# Python 3.12 virtual environment
########################################
python3.12 -m venv /opt/vllm
source /opt/vllm/bin/activate
export PATH=/opt/vllm/bin:$PATH
python --version
########################################
# install build tools (stable uv)
########################################
pip install -U pip setuptools-rust
pip install uv
pip install "setuptools<70" build wheel cmake auditwheel
uv pip install "setuptools<70" cython meson-python pybind11 "sympy>=1.13.3" --no-build-isolation
########################################
# Rust
########################################
curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y
source /root/.cargo/env
########################################
# Compiler env
########################################
source /opt/rh/gcc-toolset-14/enable
export PATH=/usr/lib64/llvm15/bin:$PATH
export LLVM_CONFIG=/usr/lib64/llvm15/bin/llvm-config
export CMAKE_ARGS="-DPython3_EXECUTABLE=python"
export MAX_JOBS=${MAX_JOBS:-$(nproc)}
export GRPC_PYTHON_BUILD_SYSTEM_OPENSSL=1
########################################
# Install Packages From Devpi
########################################
uv pip install numpy==2.3.5 pillow==12.2.0 --extra-index-url "$IBM_DEVPI_URL"
try_install_from_devpi "opencv-python-headless==${OPENCV_VERSION}"
try_install_from_devpi "torch==${TORCH_VERSION}"
try_install_from_devpi "torchvision==${TORCHVISION_VERSION}"
########################################
# torch audio
########################################
TEMP_BUILD_DIR=$(mktemp -d)
cd "${TEMP_BUILD_DIR}"
export BUILD_SOX=1 BUILD_KALDI=1 BUILD_RNNT=1 USE_FFMPEG=0 USE_ROCM=0 USE_CUDA=0
export TORCHAUDIO_TEST_ALLOW_SKIP_IF_NO_FFMPEG=1
git clone --recursive https://github.com/pytorch/audio.git -b v${TORCHAUDIO_VERSION}
cd audio
#patching
sed -i '
s|_CSRC_DIR / "_torchaudio.cpp"|str(_CSRC_DIR / "_torchaudio.cpp")|;
s|_CSRC_DIR / "utils.cpp"|str(_CSRC_DIR / "utils.cpp")|;
s|sources=\[_CSRC_DIR / s for s in sources\]|sources=[str(_CSRC_DIR / s) for s in sources]|;
' tools/setup_helpers/extension.py
MAX_JOBS=${MAX_JOBS:-$(nproc)} \
BUILD_VERSION=${TORCHAUDIO_VERSION} \
uv build --wheel --out-dir "${WHEEL_DIR}" --no-build-isolation
uv pip install "${WHEEL_DIR}"/torchaudio*.whl
cd "${REPO_ROOT}"
rm -rf "${TEMP_BUILD_DIR}"
########################################
# Xgrammar
########################################
uv pip install \
"scikit-build-core==0.11.6" \
"pyproject-metadata<0.8" \
pathspec \
packaging \
distro \
"setuptools<70" \
setuptools_scm \
cmake \
ninja \
pybind11 \
nanobind
uv pip install apache-tvm-ffi==0.1.12 \
--no-build-isolation \
--no-cache
TEMP_BUILD_DIR=$(mktemp -d)
pushd "${TEMP_BUILD_DIR}"
export CFLAGS="-fno-lto -mcpu=power9"
export CXXFLAGS="-fno-lto -mcpu=power9"
export LDFLAGS="-fno-lto"
export PATH=/opt/vllm/bin:$PATH
export Python_EXECUTABLE=/opt/vllm/bin/python3
export Python3_EXECUTABLE=/opt/vllm/bin/python3
export PYTHON_EXECUTABLE=/opt/vllm/bin/python3
export Python_ROOT_DIR=/opt/vllm
export Python3_ROOT_DIR=/opt/vllm
git clone \
--recursive \
https://github.com/mlc-ai/xgrammar \
-b "v${XGRAMMAR_VERSION}"
cd xgrammar
cp cmake/config.cmake .
export PYTHONPATH=/opt/vllm/lib64/python3.12/site-packages:/opt/vllm/lib/python3.12/site-packages:${PYTHONPATH:-}
uv build \
--wheel \
--out-dir "${WHEEL_DIR}" \
--no-build-isolation
uv pip install "${WHEEL_DIR}"/xgrammar*.whl -v
popd
rm -rf "${TEMP_BUILD_DIR}"
cd "${REPO_ROOT}"
########################################
# RHOAI Binary Downloads
########################################
pip download \
--index-url "${RHOAI_INDEX_URL}" \
--only-binary=:all: \
--no-deps \
llvmlite==0.47.0 \
-d "${WHEEL_DIR}"
pip download \
--index-url "${RHOAI_INDEX_URL}" \
--only-binary=:all: \
--no-deps \
Numba==0.65.0 \
-d "${WHEEL_DIR}"
########################################
# install built wheels
########################################
uv pip install setuptools_scm maturin setuptools-rust ninja scikit-build-core pybind11 nanobind \
--no-build-isolation
uv pip install "${WHEEL_DIR}"/*.whl
########################################
# install remaining deps
########################################
sed -i.bak -e 's/.*torch.*//g' pyproject.toml requirements/*.txt
uv pip install "setuptools>=78.1.1" --no-build-isolation
export PKG_CONFIG_PATH=/usr/local/lib/pkgconfig:/usr/local/lib64/pkgconfig:/usr/lib64/pkgconfig
uv pip install -r requirements/common.txt \
-r requirements/cpu.txt \
-r requirements/build/cpu.txt --index-strategy unsafe-best-match
+5 -23
View File
@@ -15,7 +15,6 @@ endif()
#
set(ENABLE_X86_ISA $ENV{VLLM_CPU_X86})
set(ENABLE_ARM_BF16 $ENV{VLLM_CPU_ARM_BF16})
set(ENABLE_RVV_BF16 $ENV{VLLM_CPU_RVV_BF16})
include_directories("${CMAKE_SOURCE_DIR}/csrc")
@@ -111,13 +110,6 @@ else()
set(ARM_BF16_FOUND ON)
message(STATUS "ARM BF16 support enabled via VLLM_CPU_ARM_BF16 environment variable")
endif()
# Some kernels (e.g. Bianbu on Spacemit X100) do not report zvfbfmin
# in /proc/cpuinfo despite hardware support. VLLM_CPU_RVV_BF16=1
# overrides the detection result.
if (ENABLE_RVV_BF16)
set(RVV_BF16_FOUND ON)
message(STATUS "RVV BF16 support enabled via VLLM_CPU_RVV_BF16 environment variable")
endif()
endif()
if (CMAKE_SYSTEM_PROCESSOR MATCHES "x86_64|amd64" OR ENABLE_X86_ISA)
@@ -186,10 +178,7 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
# Override with -DVLLM_RVV_VLEN=128 or -DVLLM_RVV_VLEN=256 for RVV.
if(NOT DEFINED VLLM_RVV_VLEN)
# Auto-detect: find the largest zvl<N>b in /proc/cpuinfo isa line.
# Skip when cross-compiling — /proc/cpuinfo describes the build host.
if(CMAKE_CROSSCOMPILING)
message(STATUS "Cross-compiling: skipping VLEN auto-detection from /proc/cpuinfo")
elseif(EXISTS /proc/cpuinfo)
if(EXISTS /proc/cpuinfo)
file(READ /proc/cpuinfo _cpuinfo)
set(_best 0)
foreach(_n IN ITEMS 128 256 512 1024)
@@ -197,13 +186,6 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
set(_best ${_n})
endif()
endforeach()
# Only VLEN=128 and VLEN=256 are supported by the RVV kernels.
if(_best GREATER 256)
message(WARNING
"Detected VLEN=${_best} but only 128/256 are supported; "
"clamping to 256")
set(_best 256)
endif()
if(_best GREATER 0)
set(VLLM_RVV_VLEN ${_best})
endif()
@@ -213,9 +195,9 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
if(NOT DEFINED VLLM_RVV_VLEN AND (RVV_FP16_FOUND OR RVV_BF16_FOUND))
message(FATAL_ERROR
"RISC-V RVV is available but VLEN could not be auto-detected. "
"Please specify VLEN explicitly via CMAKE_ARGS:\n"
" CMAKE_ARGS='-DVLLM_RVV_VLEN=128' (for VLEN=128 hardware)\n"
" CMAKE_ARGS='-DVLLM_RVV_VLEN=256' (for VLEN=256 hardware, e.g. Spacemit X100)")
"Please specify VLEN explicitly:\n"
" -DVLLM_RVV_VLEN=128 (for VLEN=128 hardware)\n"
" -DVLLM_RVV_VLEN=256 (for VLEN=256 hardware, e.g. Spacemit X100)")
endif()
endif()
if(VLLM_RVV_VLEN AND VLLM_RVV_VLEN GREATER 0)
@@ -227,7 +209,7 @@ elseif (CMAKE_SYSTEM_PROCESSOR MATCHES "riscv64")
message(STATUS "BF16 extension detected")
set(MARCH_FLAGS -march=rv64gcv_zvfh_zfbfmin_zvfbfmin_zvl${VLLM_RVV_VLEN}b -mrvv-vector-bits=zvl -mabi=lp64d)
elseif(RVV_FP16_FOUND)
message(WARNING "BF16 functionality is not available.")
message(WARNING "BF16 functionality is not available")
set(MARCH_FLAGS -march=rv64gcv_zvfh_zvl${VLLM_RVV_VLEN}b -mrvv-vector-bits=zvl -mabi=lp64d)
else()
message(STATUS "compile riscv with scalar (no FP16/BF16)")
+1 -1
View File
@@ -17,7 +17,7 @@ else()
FetchContent_Declare(
fmha_sm100
GIT_REPOSITORY https://github.com/vllm-project/MSA.git
GIT_TAG 2e63ec37a0fc29bc20f39cd1a52e0f5affc33a73
GIT_TAG fee783153f3efe57e3e933c5cb7e267a7cebcfb5
GIT_PROGRESS TRUE
CONFIGURE_COMMAND ""
BUILD_COMMAND ""
@@ -39,7 +39,7 @@ else()
FetchContent_Declare(
vllm-flash-attn
GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
GIT_TAG b3964b1d8b95d8e8447435668ab169a2700bab65
GIT_TAG 2c839c33742309ec41e620bf837495ec9926c56e
GIT_PROGRESS TRUE
# Don't share the vllm-flash-attn build between build types
BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn
+1 -2
View File
@@ -13,8 +13,7 @@ static inline cpu_attention::Fp8KVCacheDataType parse_fp8_kv_dtype(
bool cpu_attn_has_isa(const std::string& isa) {
if (isa == "rvv") {
#if defined(__riscv) && defined(__riscv_v_min_vlen) && \
(__riscv_v_min_vlen == 128 || __riscv_v_min_vlen == 256)
#if defined(__riscv) && defined(__riscv_v_min_vlen) && __riscv_v_min_vlen == 128
return true;
#else
return false;
+2
View File
@@ -323,6 +323,8 @@ class AttentionImpl<ISA::VSX, scalar_t, head_dim> {
const int64_t num_blocks_stride, const int64_t cache_head_num_stride,
const int64_t block_size, const int64_t block_size_stride,
const float k_inv = 0.0f, const float v_inv = 0.0f) {
// k_inv and v_inv are unused on VSX: FP8 KV cache is not supported on
// PowerPC. The parameters are present to match the common interface.
#pragma omp parallel for collapse(2)
for (int64_t token_idx = 0; token_idx < token_num; ++token_idx) {
for (int64_t head_idx = 0; head_idx < head_num; ++head_idx) {
+2 -2
View File
@@ -4,7 +4,7 @@
#if defined(__x86_64__)
// x86 implementation
#include "cpu_types_x86.hpp"
#elif defined(__powerpc__)
#elif defined(__POWER9_VECTOR__)
// ppc implementation
#include "cpu_types_vsx.hpp"
#elif defined(__s390x__)
@@ -41,4 +41,4 @@ inline int get_max_threads() {
}
} // namespace cpu_utils
#endif
#endif
+16 -35
View File
@@ -214,18 +214,11 @@ struct BF16Vec32 : public Vec<BF16Vec32> {
explicit BF16Vec32(const BF16Vec8& v) {
fixed_u16x8_t u16_val = bf16_to_u16(v.reg);
// Widen LMUL_128 → LMUL_256 so vslideup operands share a type.
// At VLEN=256 this is mf2→m1 (both integer); at VLEN=128 it is m1→m2.
fixed_u16x16_t ext =
RVVI4(__riscv_vlmul_ext_v_u16, LMUL_128, _u16, LMUL_256)(u16_val);
// Build 16-element half: place the 8 elements at offsets 0 and 8.
fixed_u16x16_t half = RVVI(__riscv_vmv_v_x_u16, LMUL_256)(0, 16);
half = RVVI(__riscv_vslideup_vx_u16, LMUL_256)(half, ext, 0, 8);
half = RVVI(__riscv_vslideup_vx_u16, LMUL_256)(half, ext, 8, 16);
// Double to LMUL_512 (m1→m2 at VLEN=256, m2→m4 at VLEN=128).
fixed_u16x32_t dst =
RVVI4(__riscv_vcreate_v_u16, LMUL_256, _u16, LMUL_512)(half, half);
reg = RVVI4(__riscv_vreinterpret_v_u16, LMUL_512, _bf16, LMUL_512)(dst);
fixed_u16x32_t u16_combined =
RVVI4(__riscv_vcreate_v_u16, LMUL_128, _u16, LMUL_512)(
u16_val, u16_val, u16_val, u16_val);
reg = RVVI4(__riscv_vreinterpret_v_u16, LMUL_512, _bf16,
LMUL_512)(u16_combined);
};
void save(void* ptr) const {
@@ -630,29 +623,17 @@ struct FP32Vec16 : public Vec<FP32Vec16> {
data.reg, data.reg)) {};
explicit FP32Vec16(const FP32Vec16& data) : reg(data.reg) {};
explicit FP32Vec16(int64_t value, const FP32Vec16& lut) {
// Split into two 32-bit halves to avoid u64 @ LMUL_1024 (m8 on
// VLEN=128 / m4 on VLEN=256), which causes heavy register spilling.
constexpr int HALF = VEC_ELEM_NUM / 2;
const auto q = static_cast<uint64_t>(value);
const uint32_t lo = static_cast<uint32_t>(q);
const uint32_t hi = static_cast<uint32_t>(q >> 32);
auto lane_ids = RVVI(__riscv_vid_v_u32, LMUL_256)(HALF);
auto shifts = RVVI(__riscv_vsll_vx_u32, LMUL_256)(lane_ids, 2, HALF);
auto packed_lo = RVVI(__riscv_vmv_v_x_u32, LMUL_256)(lo, HALF);
auto idx_lo = RVVI(__riscv_vand_vx_u32, LMUL_256)(
RVVI(__riscv_vsrl_vv_u32, LMUL_256)(packed_lo, shifts, HALF), 0xF,
HALF);
auto packed_hi = RVVI(__riscv_vmv_v_x_u32, LMUL_256)(hi, HALF);
auto idx_hi = RVVI(__riscv_vand_vx_u32, LMUL_256)(
RVVI(__riscv_vsrl_vv_u32, LMUL_256)(packed_hi, shifts, HALF), 0xF,
HALF);
auto idx =
RVVI4(__riscv_vcreate_v_u32, LMUL_256, _u32, LMUL_512)(idx_lo, idx_hi);
reg = RVVI(__riscv_vrgather_vv_f32, LMUL_512)(lut.reg, idx, VEC_ELEM_NUM);
const uint64_t q_values = static_cast<uint64_t>(value);
auto packed = RVVI(__riscv_vmv_v_x_u64, LMUL_1024)(q_values, VEC_ELEM_NUM);
auto lane_ids = RVVI(__riscv_vid_v_u64, LMUL_1024)(VEC_ELEM_NUM);
auto shifts =
RVVI(__riscv_vsll_vx_u64, LMUL_1024)(lane_ids, 2, VEC_ELEM_NUM);
auto shifted =
RVVI(__riscv_vsrl_vv_u64, LMUL_1024)(packed, shifts, VEC_ELEM_NUM);
auto idx64 =
RVVI(__riscv_vand_vx_u64, LMUL_1024)(shifted, 0xF, VEC_ELEM_NUM);
auto idx32 = RVVI(__riscv_vnsrl_wx_u32, LMUL_512)(idx64, 0, VEC_ELEM_NUM);
reg = RVVI(__riscv_vrgather_vv_f32, LMUL_512)(lut.reg, idx32, VEC_ELEM_NUM);
}
explicit FP32Vec16(const FP16Vec16& v);
+40 -120
View File
@@ -344,133 +344,53 @@ struct FP32Vec8 : public Vec<FP32Vec8> {
return result;
}
FP32Vec8 exp() const {
f32x4x2_t out;
const __vector float log2e = vec_splats(1.44269504088896341f);
const __vector float one = vec_splats(1.0f);
const __vector float min_x = vec_splats(-87.3f);
const __vector float max_x = vec_splats(88.7f);
// 5th-degree minimax polynomial for 2^r (r in [0,1))
const __vector float c1 = vec_splats(0.6931471805599453f);
const __vector float c2 = vec_splats(0.240226506959101f);
const __vector float c3 = vec_splats(0.05550410866482158f);
const __vector float c4 = vec_splats(0.009618129107628477f);
const __vector float c5 = vec_splats(0.0013333558146428443f);
for (int i = 0; i < 2; i++) {
__vector float x = reg.val[i];
x = vec_max(x, min_x);
x = vec_min(x, max_x);
__vector float y = vec_mul(x, log2e);
__vector float kf = vec_floor(y);
__vector float r = vec_sub(y, kf);
// Convert float to signed integer. Use vec_cts for PowerPC AltiVec
// compatibility.
__vector signed int k = vec_cts(kf, 0);
const __vector signed int min_k = vec_splats((signed int)-126);
const __vector signed int max_k = vec_splats((signed int)127);
k = vec_min(vec_max(k, min_k), max_k);
// Build 2^k from exponent bits
__vector signed int exp_int = vec_add(k, vec_splats((signed int)127));
__vector unsigned int bits = (__vector unsigned int)exp_int;
bits = vec_sl(bits, vec_splats((unsigned int)23));
__vector float pow2k = (__vector float)bits;
// Improved minimax polynomial
__vector float poly = vec_madd(c5, r, c4);
poly = vec_madd(poly, r, c3);
poly = vec_madd(poly, r, c2);
poly = vec_madd(poly, r, c1);
poly = vec_madd(poly, r, one);
out.val[i] = vec_mul(pow2k, poly);
}
return FP32Vec8(out);
// TODO: Vectorize this
AliasReg ar;
ar.reg = reg;
f32x4x4_t ret;
ret.val[0][0] = std::exp(ar.values[0]);
ret.val[0][1] = std::exp(ar.values[1]);
ret.val[0][2] = std::exp(ar.values[2]);
ret.val[0][3] = std::exp(ar.values[3]);
ret.val[1][0] = std::exp(ar.values[4]);
ret.val[1][1] = std::exp(ar.values[5]);
ret.val[1][2] = std::exp(ar.values[6]);
ret.val[1][3] = std::exp(ar.values[7]);
return FP32Vec8(f32x4x2_t({ret.val[0], ret.val[1]}));
}
FP32Vec8 tanh() const {
const __vector float one = vec_splats(1.0f);
const __vector float two = vec_splats(2.0f);
const __vector float zero = vec_splats(0.0f);
const __vector float sat = vec_splats(9.0f);
f32x4x2_t out;
for (int i = 0; i < 2; i++) {
__vector float x = reg.val[i];
__vector float ax = vec_abs(x);
__vector bool int mask = vec_cmpge(x, zero);
__vector float sign = vec_sel(vec_splats(-1.0f), one, mask);
__vector bool int saturated = vec_cmpge(ax, sat);
__vector float two_x = vec_mul(x, two);
f32x4x2_t tmp;
tmp.val[0] = two_x;
tmp.val[1] = two_x;
FP32Vec8 temp_vec(tmp);
vector float e = temp_vec.exp().reg.val[0];
vector float num = vec_sub(e, one);
vector float den = vec_add(e, one);
vector float t = vec_div(num, den);
out.val[i] = vec_sel(t, sign, saturated);
}
return FP32Vec8(out);
// TODO: Vectorize this
AliasReg ar;
ar.reg = reg;
f32x4x4_t ret;
ret.val[0][0] = std::tanh(ar.values[0]);
ret.val[0][1] = std::tanh(ar.values[1]);
ret.val[0][2] = std::tanh(ar.values[2]);
ret.val[0][3] = std::tanh(ar.values[3]);
ret.val[1][0] = std::tanh(ar.values[4]);
ret.val[1][1] = std::tanh(ar.values[5]);
ret.val[1][2] = std::tanh(ar.values[6]);
ret.val[1][3] = std::tanh(ar.values[7]);
return FP32Vec8(f32x4x2_t({ret.val[0], ret.val[1]}));
}
FP32Vec8 er() const {
const vector float a1 = vec_splats(0.254829592f);
const vector float a2 = vec_splats(-0.284496736f);
const vector float a3 = vec_splats(1.421413741f);
const vector float a4 = vec_splats(-1.453152027f);
const vector float a5 = vec_splats(1.061405429f);
const vector float p = vec_splats(0.3275911f);
const vector float one = vec_splats(1.0f);
const vector float zero = vec_splats(0.0f);
const vector float sat = vec_splats(6.0f);
f32x4x2_t ret;
for (int i = 0; i < 2; i++) {
vector float x = reg.val[i];
vector float ax = vec_abs(x);
vector bool int mask = vec_cmpge(x, zero);
vector float sign = vec_sel(vec_splats(-1.0f), one, mask);
vector bool int saturated = vec_cmpge(ax, sat);
vector float t = vec_div(one, vec_madd(p, ax, one));
vector float poly = a5;
poly = vec_madd(poly, t, a4);
poly = vec_madd(poly, t, a3);
poly = vec_madd(poly, t, a2);
poly = vec_madd(poly, t, a1);
poly = vec_mul(poly, t);
vector float x_squared = vec_mul(x, x);
vector float neg_x_squared = vec_mul(vec_splats(-1.0f), x_squared);
f32x4x2_t tmp;
tmp.val[0] = neg_x_squared;
tmp.val[1] = neg_x_squared;
FP32Vec8 exp_input(tmp);
vector float exp_term = exp_input.exp().reg.val[0];
vector float y = vec_nmsub(poly, exp_term, one);
vector float erf_val = vec_mul(sign, y);
ret.val[i] = vec_sel(erf_val, sign, saturated);
}
return FP32Vec8(ret);
// TODO: Vectorize this
AliasReg ar;
ar.reg = reg;
f32x4x4_t ret;
ret.val[0][0] = std::erf(ar.values[0]);
ret.val[0][1] = std::erf(ar.values[1]);
ret.val[0][2] = std::erf(ar.values[2]);
ret.val[0][3] = std::erf(ar.values[3]);
ret.val[1][0] = std::erf(ar.values[4]);
ret.val[1][1] = std::erf(ar.values[5]);
ret.val[1][2] = std::erf(ar.values[6]);
ret.val[1][3] = std::erf(ar.values[7]);
return FP32Vec8(f32x4x2_t({ret.val[0], ret.val[1]}));
}
FP32Vec8 operator*(const FP32Vec8& b) const {
+1 -2
View File
@@ -278,8 +278,7 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
ops.def(
"dynamic_4bit_int_moe("
"Tensor x, Tensor topk_ids, Tensor topk_weights,"
"Tensor w13_packed, Tensor w2_packed,"
"int hidden_size, int intermediate_size,"
"Tensor w13_packed, Tensor w2_packed, int H, int I, int I2,"
"int group_size, bool apply_router_weight_on_input, int activation_kind"
") -> Tensor");
+2 -2
View File
@@ -76,14 +76,14 @@ inline int64_t get_available_l2_size() {
if (l2_cache_size == 0) {
l2_cache_size = 256 * 1024;
}
return static_cast<int64_t>(l2_cache_size) >> 1;
return static_cast<int64_t>(l2_cache_size) >> 1; // use 50% of L2 cache
}();
return size;
#else
static int64_t size = []() {
auto caps = at::cpu::get_cpu_capabilities();
const uint32_t l2_cache_size = caps.at("l2_cache_size").toInt();
return l2_cache_size >> 1;
return l2_cache_size >> 1; // use 50% of L2 cache
}();
return size;
#endif
+1 -2
View File
@@ -15,8 +15,7 @@ void topk_sigmoid(torch::stable::Tensor& topk_weights,
torch::stable::Tensor& topk_indices,
torch::stable::Tensor& token_expert_indices,
torch::stable::Tensor& gating_output, bool renormalize,
std::optional<torch::stable::Tensor> bias,
double routed_scaling_factor);
std::optional<torch::stable::Tensor> bias);
void topk_softplus_sqrt(
torch::stable::Tensor& topk_weights, torch::stable::Tensor& topk_indices,
@@ -173,8 +173,7 @@ __launch_bounds__(TPB) __global__ void moeTopK(
const int start_expert,
const int end_expert,
const bool renormalize,
const float* bias,
const double routed_scaling_factor)
const float* bias)
{
using cub_kvp = cub::KeyValuePair<int, float>;
@@ -242,16 +241,14 @@ __launch_bounds__(TPB) __global__ void moeTopK(
__syncthreads();
}
// Apply renormalization and routed scaling factor to final weights.
if (threadIdx.x == 0) {
float scale = static_cast<float>(routed_scaling_factor);
if (renormalize) {
// Renormalize the k weights for this row to sum to 1, if requested.
if (renormalize) {
if (threadIdx.x == 0) {
const float denom = selected_sum > 0.f ? selected_sum : 1.f;
scale /= denom;
}
for (int k_idx = 0; k_idx < k; ++k_idx) {
const int idx = k * block_row + k_idx;
output[idx] = output[idx] * scale;
for (int k_idx = 0; k_idx < k; ++k_idx) {
const int idx = k * block_row + k_idx;
output[idx] = output[idx] / denom;
}
}
}
}
@@ -277,7 +274,7 @@ template <int VPT, int NUM_EXPERTS, int WARPS_PER_CTA, int BYTES_PER_LDG, int WA
__launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
void topkGating(const InputType* input, const bool* finished, float* output, const int num_rows, IndType* indices,
int* source_rows, const int k, const int start_expert, const int end_expert, const bool renormalize,
const float* bias, const double routed_scaling_factor)
const float* bias)
{
static_assert(std::is_same_v<InputType, float> || std::is_same_v<InputType, __nv_bfloat16> ||
std::is_same_v<InputType, __half>,
@@ -573,17 +570,17 @@ __launch_bounds__(WARPS_PER_CTA* WARP_SIZE_PARAM) __global__
}
}
// Apply renormalization and routed scaling factor to final weights.
if (thread_group_idx == 0) {
float scale = static_cast<float>(routed_scaling_factor);
if (renormalize) {
const float denom = selected_sum > 0.f ? selected_sum : 1.f;
scale /= denom;
}
for (int k_idx = 0; k_idx < k; ++k_idx) {
const int idx = k * thread_row + k_idx;
output[idx] = output[idx] * scale;
}
// Renormalize the k weights for this row to sum to 1, if requested.
if (renormalize) {
if (thread_group_idx == 0)
{
const float denom = selected_sum > 0.f ? selected_sum : 1.f;
for (int k_idx = 0; k_idx < k; ++k_idx)
{
const int idx = k * thread_row + k_idx;
output[idx] = output[idx] / denom;
}
}
}
}
@@ -605,7 +602,7 @@ struct TopkConstants
template <int EXPERTS, int WARPS_PER_TB, int WARP_SIZE_PARAM, int MAX_BYTES_PER_LDG, typename IndType, typename InputType, ScoringFunc SF>
void topkGatingLauncherHelper(const InputType* input, const bool* finished, float* output, IndType* indices,
int* source_row, const int num_rows, const int k, const int start_expert, const int end_expert, const bool renormalize,
const float* bias, const double routed_scaling_factor, cudaStream_t stream)
const float* bias, cudaStream_t stream)
{
static constexpr int BYTES_PER_LDG = MIN(MAX_BYTES_PER_LDG, sizeof(InputType) * EXPERTS);
using Constants = detail::TopkConstants<EXPERTS, BYTES_PER_LDG, WARP_SIZE_PARAM, InputType>;
@@ -616,7 +613,7 @@ void topkGatingLauncherHelper(const InputType* input, const bool* finished, floa
dim3 block_dim(WARP_SIZE_PARAM, WARPS_PER_TB);
topkGating<VPT, EXPERTS, WARPS_PER_TB, BYTES_PER_LDG, WARP_SIZE_PARAM, IndType, InputType, SF><<<num_blocks, block_dim, 0, stream>>>(
input, finished, output, num_rows, indices, source_row, k, start_expert, end_expert, renormalize, bias, routed_scaling_factor);
input, finished, output, num_rows, indices, source_row, k, start_expert, end_expert, renormalize, bias);
}
#ifndef USE_ROCM
@@ -627,7 +624,7 @@ void topkGatingLauncherHelper(const InputType* input, const bool* finished, floa
IndType, InputType, SF>( \
gating_output, nullptr, topk_weights, topk_indices, \
token_expert_indices, num_tokens, topk, 0, num_experts, renormalize, \
bias, routed_scaling_factor, stream);
bias, stream);
#else
#define LAUNCH_TOPK(NUM_EXPERTS, WARPS_PER_TB, MAX_BYTES) \
if (WARP_SIZE == 64) { \
@@ -635,13 +632,13 @@ void topkGatingLauncherHelper(const InputType* input, const bool* finished, floa
IndType, InputType, SF>( \
gating_output, nullptr, topk_weights, topk_indices, \
token_expert_indices, num_tokens, topk, 0, num_experts, renormalize, \
bias, routed_scaling_factor, stream); \
bias, stream); \
} else if (WARP_SIZE == 32) { \
topkGatingLauncherHelper<NUM_EXPERTS, WARPS_PER_TB, 32, MAX_BYTES, \
IndType, InputType, SF>( \
gating_output, nullptr, topk_weights, topk_indices, \
token_expert_indices, num_tokens, topk, 0, num_experts, renormalize, \
bias, routed_scaling_factor, stream); \
bias, stream); \
} else { \
assert(false && \
"Unsupported warp size. Only 32 and 64 are supported for ROCm"); \
@@ -660,7 +657,6 @@ void topkGatingKernelLauncher(
const int topk,
const bool renormalize,
const float* bias,
const double routed_scaling_factor,
cudaStream_t stream) {
static constexpr int WARPS_PER_TB = 4;
static constexpr int BYTES_PER_LDG_POWER_OF_2 = 16;
@@ -736,7 +732,7 @@ void topkGatingKernelLauncher(
}
moeTopK<TPB><<<num_tokens, TPB, 0, stream>>>(
workspace, nullptr, topk_weights, topk_indices, token_expert_indices,
num_experts, topk, 0, num_experts, renormalize, bias, routed_scaling_factor);
num_experts, topk, 0, num_experts, renormalize, bias);
}
}
}
@@ -754,7 +750,6 @@ void dispatch_topk_launch(
torch::stable::Tensor& softmax_workspace,
int num_tokens, int num_experts, int topk, bool renormalize,
std::optional<torch::stable::Tensor> bias,
double routed_scaling_factor,
cudaStream_t stream)
{
const float* bias_ptr = nullptr;
@@ -777,7 +772,7 @@ void dispatch_topk_launch(
token_expert_indices.mutable_data_ptr<int>(),
softmax_workspace.mutable_data_ptr<float>(),
num_tokens, num_experts, topk, renormalize,
bias_ptr, routed_scaling_factor, stream);
bias_ptr, stream);
} else if (topk_indices.scalar_type() == torch::headeronly::ScalarType::UInt32) {
vllm::moe::topkGatingKernelLauncher<uint32_t, ComputeType, SF>(
reinterpret_cast<const ComputeType*>(gating_output.const_data_ptr()),
@@ -786,7 +781,7 @@ void dispatch_topk_launch(
token_expert_indices.mutable_data_ptr<int>(),
softmax_workspace.mutable_data_ptr<float>(),
num_tokens, num_experts, topk, renormalize,
bias_ptr, routed_scaling_factor, stream);
bias_ptr, stream);
} else {
STD_TORCH_CHECK(topk_indices.scalar_type() == torch::headeronly::ScalarType::Long);
vllm::moe::topkGatingKernelLauncher<int64_t, ComputeType, SF>(
@@ -796,7 +791,7 @@ void dispatch_topk_launch(
token_expert_indices.mutable_data_ptr<int>(),
softmax_workspace.mutable_data_ptr<float>(),
num_tokens, num_experts, topk, renormalize,
bias_ptr, routed_scaling_factor, stream);
bias_ptr, stream);
}
}
@@ -825,15 +820,15 @@ void topk_softmax(
if (gating_output.scalar_type() == torch::headeronly::ScalarType::Float) {
dispatch_topk_launch<float, vllm::moe::SCORING_SOFTMAX>(gating_output, topk_weights, topk_indices,
token_expert_indices, softmax_workspace, num_tokens, num_experts, topk, renormalize,
bias, 1.0, stream);
bias, stream);
} else if (gating_output.scalar_type() == torch::headeronly::ScalarType::Half) {
dispatch_topk_launch<__half, vllm::moe::SCORING_SOFTMAX>(gating_output, topk_weights, topk_indices,
token_expert_indices, softmax_workspace, num_tokens, num_experts, topk, renormalize,
bias, 1.0, stream);
bias, stream);
} else if (gating_output.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
dispatch_topk_launch<__nv_bfloat16, vllm::moe::SCORING_SOFTMAX>(gating_output, topk_weights, topk_indices,
token_expert_indices, softmax_workspace, num_tokens, num_experts, topk, renormalize,
bias, 1.0, stream);
bias, stream);
} else {
STD_TORCH_CHECK(false, "Unsupported gating_output data type: ", gating_output.scalar_type());
}
@@ -845,8 +840,7 @@ void topk_sigmoid(
torch::stable::Tensor& token_expert_indices, // [num_tokens, topk]
torch::stable::Tensor& gating_output, // [num_tokens, num_experts]
bool renormalize,
std::optional<torch::stable::Tensor> bias,
double routed_scaling_factor)
std::optional<torch::stable::Tensor> bias)
{
const int num_experts = gating_output.size(-1);
const auto num_tokens = gating_output.numel() / num_experts;
@@ -865,15 +859,15 @@ void topk_sigmoid(
if (gating_output.scalar_type() == torch::headeronly::ScalarType::Float) {
dispatch_topk_launch<float, vllm::moe::SCORING_SIGMOID>(gating_output, topk_weights, topk_indices,
token_expert_indices, workspace, num_tokens, num_experts, topk, renormalize,
bias, routed_scaling_factor, stream);
bias, stream);
} else if (gating_output.scalar_type() == torch::headeronly::ScalarType::Half) {
dispatch_topk_launch<__half, vllm::moe::SCORING_SIGMOID>(gating_output, topk_weights, topk_indices,
token_expert_indices, workspace, num_tokens, num_experts, topk, renormalize,
bias, routed_scaling_factor, stream);
bias, stream);
} else if (gating_output.scalar_type() == torch::headeronly::ScalarType::BFloat16) {
dispatch_topk_launch<__nv_bfloat16, vllm::moe::SCORING_SIGMOID>(gating_output, topk_weights, topk_indices,
token_expert_indices, workspace, num_tokens, num_experts, topk, renormalize,
bias, routed_scaling_factor, stream);
bias, stream);
} else {
STD_TORCH_CHECK(false, "Unsupported gating_output data type: ", gating_output.scalar_type());
}
+2 -2
View File
@@ -13,8 +13,8 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_moe_C, m) {
// Apply topk sigmoid to the gating outputs.
m.def(
"topk_sigmoid(Tensor! topk_weights, Tensor! topk_indices, Tensor! "
"token_expert_indices, Tensor gating_output, bool renormalize, "
"Tensor? bias, float routed_scaling_factor) -> ()");
"token_expert_indices, Tensor gating_output, bool renormalize, Tensor? "
"bias) -> ()");
m.def(
"topk_softplus_sqrt(Tensor! topk_weights, Tensor! topk_indices, Tensor! "
@@ -1,96 +0,0 @@
// N-gram embedding index kernel for LongCat-Flash (n-gram embedding variant).
//
// Adapted from SGLang:
// https://github.com/sgl-project/sglang/blob/main/python/sglang/jit_kernel/csrc/ngram_embedding.cuh
//
// For each position, computes the hashed n-gram embedding ids that index the
// concatenated embedder table. Integer tensors are int32 except ``row_indices``
// (int64); the token table is ``[max_running_reqs, max_context_len]`` int32,
// where a negative entry marks an ignored token (e.g. an EOS boundary).
#include "torch_utils.h"
#include "ops.h"
#include <cstdint>
namespace vllm::ngram_embedding {
constexpr int kBlockThreads = 256;
__global__ void ComputeNGramIdsKernel(
int batch_size, int ne_n, int ne_k,
int* ne_weights, // [ne_n-1, ne_k, ne_n]
int* ne_mods, // [ne_n-1, ne_k]
int* exclusive_ne_embedder_size_sums, // [(ne_n-1)*ne_k + 1]
int* exclusive_req_len_sums, // [batch_size + 1]
int* ne_token_table, // [max_running_reqs, max_context_len]
int max_context_len,
const int64_t* __restrict__ row_indices, // [batch_size]
int* column_starts, // [batch_size]
int* n_gram_ids // [token_num, (ne_n-1)*ne_k]
) {
const int req_id = blockIdx.x % batch_size;
const int config_id = (blockIdx.x - req_id) / batch_size;
// n and k are offset from their physical meaning: n = real_n - 2, k = real_k
// - 1 (they index into ne_weights / ne_mods).
const int k = config_id % ne_k;
const int n = (config_id - config_id % ne_k) / ne_k;
const int ne_weight_base_idx = n * ne_k * ne_n + k * ne_n;
const int ne_mod = ne_mods[n * ne_k + k];
for (int i = exclusive_req_len_sums[req_id] + threadIdx.x;
i < exclusive_req_len_sums[req_id + 1]; i += blockDim.x) {
uint64_t n_gram_id = 0;
const int64_t current_token_offset = i - exclusive_req_len_sums[req_id];
const int64_t req_token_table_index =
row_indices[req_id] * static_cast<int64_t>(max_context_len);
const int64_t current_token_table_index =
req_token_table_index + column_starts[req_id] + current_token_offset;
for (int j = 0; j < n + 2; j++) {
if (current_token_table_index - j < req_token_table_index) {
break; // outside this request's range
}
if (ne_token_table[current_token_table_index - j] < 0) {
break; // ignored token
}
const uint64_t term =
(uint64_t)ne_token_table[current_token_table_index - j] *
(uint64_t)ne_weights[ne_weight_base_idx + j];
n_gram_id += term % ne_mod;
}
n_gram_id %= ne_mod;
n_gram_id += exclusive_ne_embedder_size_sums[n * ne_k + k];
n_gram_ids[i * (ne_n - 1) * ne_k + n * ne_k + k] = (int)(n_gram_id);
}
}
} // namespace vllm::ngram_embedding
void ngram_compute_n_gram_ids(
int64_t ne_n, int64_t ne_k, torch::stable::Tensor& ne_weights,
torch::stable::Tensor& ne_mods,
torch::stable::Tensor& exclusive_ne_embedder_size_sums,
torch::stable::Tensor& exclusive_req_len_sums,
torch::stable::Tensor& ne_token_table, torch::stable::Tensor& row_indices,
torch::stable::Tensor& column_starts, torch::stable::Tensor& n_gram_ids) {
const int batch_size = static_cast<int>(exclusive_req_len_sums.size(0) - 1);
const int max_context_len = static_cast<int>(ne_token_table.size(1));
const int num_configs = (static_cast<int>(ne_n) - 1) * static_cast<int>(ne_k);
const int grid_size = num_configs * batch_size;
if (grid_size <= 0) return;
const torch::stable::accelerator::DeviceGuard device_guard(
ne_weights.get_device_index());
const cudaStream_t stream = get_current_cuda_stream();
vllm::ngram_embedding::ComputeNGramIdsKernel<<<
grid_size, vllm::ngram_embedding::kBlockThreads, 0, stream>>>(
batch_size, static_cast<int>(ne_n), static_cast<int>(ne_k),
ne_weights.mutable_data_ptr<int32_t>(),
ne_mods.mutable_data_ptr<int32_t>(),
exclusive_ne_embedder_size_sums.mutable_data_ptr<int32_t>(),
exclusive_req_len_sums.mutable_data_ptr<int32_t>(),
ne_token_table.mutable_data_ptr<int32_t>(), max_context_len,
row_indices.const_data_ptr<int64_t>(),
column_starts.mutable_data_ptr<int32_t>(),
n_gram_ids.mutable_data_ptr<int32_t>());
}
-9
View File
@@ -554,12 +554,3 @@ void cp_gather_indexer_k_quant_cache(
// quant_block_size * 4]
const torch::stable::Tensor& block_table, // [batch_size, num_blocks]
const torch::stable::Tensor& cu_seq_lens); // [batch_size + 1]
// LongCat n-gram embedding index kernel (see ngram_embedding_kernels.cu).
void ngram_compute_n_gram_ids(
int64_t ne_n, int64_t ne_k, torch::stable::Tensor& ne_weights,
torch::stable::Tensor& ne_mods,
torch::stable::Tensor& exclusive_ne_embedder_size_sums,
torch::stable::Tensor& exclusive_req_len_sums,
torch::stable::Tensor& ne_token_table, torch::stable::Tensor& row_indices,
torch::stable::Tensor& column_starts, torch::stable::Tensor& n_gram_ids);
-13
View File
@@ -598,22 +598,9 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
"Tensor? initial_state_idx,"
"Tensor? cu_chunk_seqlen,"
"Tensor? last_chunk_indices) -> ()");
// LongCat n-gram embedding index kernel. All tensor args are marked mutable
// to match the (non-const) stable-Tensor& C++ signature; only ne_token_table
// and n_gram_ids are actually written in place.
ops.def(
"ngram_compute_n_gram_ids(int ne_n, int ne_k, Tensor(a!) ne_weights, "
"Tensor(b!) ne_mods, Tensor(c!) exclusive_ne_embedder_size_sums, "
"Tensor(d!) exclusive_req_len_sums, Tensor(e!) ne_token_table, "
"Tensor(f!) row_indices, Tensor(g!) column_starts, "
"Tensor(h!) n_gram_ids) -> ()");
}
STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
// LongCat n-gram embedding index kernel.
ops.impl("ngram_compute_n_gram_ids", TORCH_BOX(&ngram_compute_n_gram_ids));
// Per-token group quantization
ops.impl("per_token_group_fp8_quant", TORCH_BOX(&per_token_group_quant_fp8));
ops.impl("per_token_group_fp8_quant_packed",
+28 -53
View File
@@ -29,37 +29,25 @@ enum ActivationKind : int64_t {
torch::Tensor dynamic_4bit_int_moe_cpu(
torch::Tensor x, torch::Tensor topk_ids, torch::Tensor topk_weights,
torch::Tensor w13_packed, torch::Tensor w2_packed, int64_t hidden_size,
int64_t intermediate_size, int64_t group_size,
bool apply_router_weight_on_input, int64_t activation_kind) {
torch::Tensor w13_packed, torch::Tensor w2_packed, int64_t H, int64_t I,
int64_t I2, int64_t group_size, bool apply_router_weight_on_input,
int64_t activation_kind) {
TORCH_CHECK(x.dim() == 2, "x must be 2D");
TORCH_CHECK(topk_ids.dim() == 2 && topk_weights.dim() == 2,
"topk tensors must be [T, K]");
TORCH_CHECK(
w13_packed.size(0) == w2_packed.size(0),
"w13_packed and w2_packed must have same number of experts in dim 0");
TORCH_CHECK(I2 == 2 * I, "I2 must equal 2*I");
const int64_t T = x.size(0);
const int64_t K = topk_ids.size(1);
const int64_t E = w13_packed.size(0);
const int64_t N = T * K;
const int64_t w13_out_features = 2 * intermediate_size;
auto x_c = x.contiguous();
// _dyn_quant_matmul_4bit kernel natively supports these pre-quant activation
// dtypes:
// - fp32: with channelwise and groupwise
// - bf16: with channelwise -> upcast to fp32 for groupwise
// - fp16: not supported -> upcast to fp32 for groupwise & channelwise
const auto output_dtype = x_c.scalar_type();
const bool should_cast_input =
((group_size != -1) && output_dtype == at::kBFloat16) ||
output_dtype == at::kHalf;
if (should_cast_input) {
x_c = x_c.to(at::kFloat);
}
auto ids_c = topk_ids.contiguous();
auto gates_c = topk_weights.to(x_c.scalar_type()).contiguous();
auto gates_c = topk_weights.to(at::kFloat).contiguous();
// bucketing tokens -> experts
c10::SmallVector<int64_t, 64> counts(
@@ -75,42 +63,35 @@ torch::Tensor dynamic_4bit_int_moe_cpu(
c10::SmallVector<int64_t, 65> offsets(E + 1, 0); // ( E +1 )
for (int64_t e = 0; e < E; ++e) offsets[e + 1] = offsets[e] + counts[e];
// expert_tokens = [tokens indices for expert 0, ...]
// expert_gates = [router weights for tokens assigned to expert 0, ...]
auto expert_tokens = at::empty({offsets[E]}, ids_c.options());
auto expert_gates = at::empty({offsets[E]}, gates_c.options());
{
c10::SmallVector<int64_t, 64> cursor(E, 0);
AT_DISPATCH_FLOATING_TYPES_AND2(
at::ScalarType::BFloat16, at::ScalarType::Half, gates_c.scalar_type(),
"bucket_expert_tokens_and_gates", [&] {
const auto* ids_ptr = ids_c.data_ptr<int64_t>();
const auto* gts_ptr = gates_c.data_ptr<scalar_t>();
auto* tok_ptr = expert_tokens.data_ptr<int64_t>();
auto* gate_ptr = expert_gates.data_ptr<scalar_t>();
const auto* ids_ptr = ids_c.data_ptr<int64_t>();
const auto* gts_ptr = gates_c.data_ptr<float>();
auto* tok_ptr = expert_tokens.data_ptr<int64_t>();
auto* gate_ptr = expert_gates.data_ptr<float>();
for (int64_t t = 0; t < T; ++t) {
const int64_t base = t * K;
for (int64_t k = 0; k < K; ++k) {
const int64_t idx = base + k;
const int64_t e = ids_ptr[idx];
const int64_t p = offsets[e] + (cursor[e]++);
tok_ptr[p] = t;
gate_ptr[p] = gts_ptr[idx];
}
}
});
for (int64_t t = 0; t < T; ++t) {
const int64_t base = t * K;
for (int64_t k = 0; k < K; ++k) {
const int64_t idx = base + k;
const int64_t e = ids_ptr[idx];
const int64_t p = offsets[e] + (cursor[e]++);
tok_ptr[p] = t;
gate_ptr[p] = gts_ptr[idx];
}
}
}
const int64_t g_eff_13 = (group_size != -1) ? group_size : hidden_size;
const int64_t g_eff_2 = (group_size != -1) ? group_size : intermediate_size;
const int64_t g_eff_13 = (group_size != -1) ? group_size : H;
const int64_t g_eff_2 = (group_size != -1) ? group_size : I;
// X_all [num_tokens * K, hidden_size]
auto X_all = x_c.index_select(/*dim=*/0, expert_tokens);
if (apply_router_weight_on_input) {
X_all = X_all.mul(expert_gates.unsqueeze(1));
}
auto Y_all = at::empty({offsets[E], hidden_size}, x_c.options());
auto Y_all = at::empty({offsets[E], H}, x_c.options());
at::parallel_for(0, offsets[E], 0, [&](int64_t idx_begin, int64_t idx_end) {
c10::InferenceMode guard;
@@ -128,13 +109,11 @@ torch::Tensor dynamic_4bit_int_moe_cpu(
auto w2_e = w2_packed.select(/*dim=*/0, e);
// W13
auto y13 = mm(x_e, w13_e, g_eff_13, /*in_features=*/hidden_size,
/*out_features=*/w13_out_features);
auto y13 =
mm(x_e, w13_e, g_eff_13, /*in_features=*/H, /*out_features=*/I2);
auto g_part =
y13.narrow(/*dim=*/1, /*start=*/0, /*length=*/intermediate_size);
auto u_part = y13.narrow(/*dim=*/1, /*start=*/intermediate_size,
/*length=*/intermediate_size);
auto g_part = y13.narrow(/*dim=*/1, /*start=*/0, /*length=*/I);
auto u_part = y13.narrow(/*dim=*/1, /*start=*/I, /*length=*/I);
torch::Tensor act;
if (activation_kind == ActivationKind::SwiGLUOAI) { // SwiGLUOAI
@@ -149,8 +128,7 @@ torch::Tensor dynamic_4bit_int_moe_cpu(
}
// W2
auto y = mm(act, w2_e, g_eff_2, /*in_features=*/intermediate_size,
/*out_features=*/hidden_size);
auto y = mm(act, w2_e, g_eff_2, /*in_features=*/I, /*out_features=*/H);
// Store per-expert result
Y_all.narrow(/*dim=*/0, /*start=*/start, /*length=*/te).copy_(y);
@@ -160,11 +138,8 @@ torch::Tensor dynamic_4bit_int_moe_cpu(
if (!apply_router_weight_on_input) {
Y_all = Y_all.mul(expert_gates.unsqueeze(1));
}
if (Y_all.scalar_type() != output_dtype) {
Y_all = Y_all.to(output_dtype);
}
auto out = at::zeros({T, hidden_size}, x.options());
auto out = at::zeros({T, H}, x.options());
out =
at::index_add(out, /*dim=*/0, /*index=*/expert_tokens, /*source=*/Y_all);
+3 -3
View File
@@ -53,9 +53,9 @@ void dynamic_scaled_int8_quant(torch::Tensor& out, torch::Tensor const& input,
torch::Tensor dynamic_4bit_int_moe_cpu(
torch::Tensor x, torch::Tensor topk_ids, torch::Tensor topk_weights,
torch::Tensor w13_packed, torch::Tensor w2_packed, int64_t hidden_size,
int64_t intermediate_size, int64_t group_size,
bool apply_router_weight_on_input, int64_t activation_kind);
torch::Tensor w13_packed, torch::Tensor w2_packed, int64_t H, int64_t I,
int64_t I2, int64_t group_size, bool apply_router_weight_on_input,
int64_t activation_kind);
using fptr_t = int64_t;
#ifdef USE_ROCM
+7 -4
View File
@@ -25,6 +25,7 @@ FROM ubuntu:22.04 AS base-common
WORKDIR /workspace
ARG PYTHON_VERSION=3.12
ARG PIP_EXTRA_INDEX_URL="https://download.pytorch.org/whl/cpu"
ARG max_jobs=32
ENV MAX_JOBS=${max_jobs}
@@ -52,6 +53,8 @@ ENV PATH="$VIRTUAL_ENV/bin:$PATH"
ENV UV_HTTP_TIMEOUT=500
# Install Python dependencies
ENV PIP_EXTRA_INDEX_URL=${PIP_EXTRA_INDEX_URL}
ENV UV_EXTRA_INDEX_URL=${PIP_EXTRA_INDEX_URL}
ENV UV_INDEX_STRATEGY="unsafe-best-match"
ENV UV_LINK_MODE="copy"
@@ -61,7 +64,7 @@ COPY requirements/cpu.txt requirements/cpu.txt
RUN --mount=type=cache,target=/root/.cache/uv \
uv pip install --upgrade pip && \
uv pip install -r requirements/cpu.txt --torch-backend cpu
uv pip install -r requirements/cpu.txt
ARG TARGETARCH
ENV TARGETARCH=${TARGETARCH}
@@ -146,7 +149,7 @@ RUN if [ "$TARGETARCH" = "arm64" ] && [ "$VLLM_CPU_X86" != "0" ]; then \
COPY requirements/build/cpu.txt requirements/build/cpu.txt
RUN --mount=type=cache,target=/root/.cache/uv \
uv pip install -r requirements/build/cpu.txt --torch-backend cpu
uv pip install -r requirements/build/cpu.txt
COPY . .
@@ -202,7 +205,7 @@ RUN case "$(uname -m)" in \
esac
RUN --mount=type=cache,target=/root/.cache/uv \
uv pip install -r requirements/test/cpu.txt --torch-backend cpu
uv pip install -r requirements/test/cpu.txt
######################### DEV IMAGE #########################
FROM vllm-build AS vllm-dev
@@ -228,7 +231,7 @@ COPY --from=vllm-test-deps /vllm-workspace/requirements/test/cpu.txt requirement
RUN --mount=type=cache,target=/root/.cache/uv \
uv pip install -r requirements/lint.txt && \
uv pip install -r requirements/test/cpu.txt --torch-backend cpu && \
uv pip install -r requirements/test/cpu.txt && \
pre-commit install --hook-type pre-commit --hook-type commit-msg
ENTRYPOINT ["bash"]
+298 -105
View File
@@ -1,80 +1,275 @@
# Base UBI image
ARG BASE_UBI_IMAGE_TAG=9.6-1754584681
###############################################################
# BUILDER STAGE #
# Stage to build openblas
###############################################################
FROM registry.access.redhat.com/ubi9/ubi-minimal:${BASE_UBI_IMAGE_TAG} AS builder-base
FROM registry.access.redhat.com/ubi9/ubi-minimal:${BASE_UBI_IMAGE_TAG} AS openblas-builder
ARG MAX_JOBS
ARG OPENBLAS_VERSION=0.3.30
RUN microdnf install -y dnf && dnf install -y gcc-toolset-14 make wget unzip \
&& source /opt/rh/gcc-toolset-14/enable \
&& wget https://github.com/OpenMathLib/OpenBLAS/releases/download/v$OPENBLAS_VERSION/OpenBLAS-$OPENBLAS_VERSION.zip \
&& unzip OpenBLAS-$OPENBLAS_VERSION.zip \
&& cd OpenBLAS-$OPENBLAS_VERSION \
&& make -j${MAX_JOBS} TARGET=POWER9 BINARY=64 USE_OPENMP=1 USE_THREAD=1 NUM_THREADS=120 DYNAMIC_ARCH=1 INTERFACE64=0 \
&& cd /tmp && touch control
###############################################################
# base stage with dependencies coming from centos mirrors
###############################################################
FROM registry.access.redhat.com/ubi9/ubi-minimal:${BASE_UBI_IMAGE_TAG} AS centos-deps-builder
RUN microdnf install -y dnf && \
dnf install -y https://mirror.stream.centos.org/9-stream/BaseOS/`arch`/os/Packages/centos-gpg-keys-9.0-26.el9.noarch.rpm \
https://mirror.stream.centos.org/9-stream/BaseOS/`arch`/os/Packages/centos-stream-repos-9.0-26.el9.noarch.rpm \
https://dl.fedoraproject.org/pub/epel/epel-release-latest-9.noarch.rpm && \
dnf config-manager --set-enabled crb
RUN dnf install -y openjpeg2-devel lcms2-devel tcl-devel tk-devel fribidi-devel yajl-devel && \
dnf remove -y centos-gpg-keys-9.0-24.el9.noarch centos-stream-repos-9.0-26.el9.noarch
###############################################################
# base stage with basic dependencies
###############################################################
FROM centos-deps-builder AS base-builder
ARG VLLM_VERSION="0.22.1"
ARG PYTHON_VERSION=3.12
ARG VLLM_TARGET_DEVICE=cpu
ARG OPENBLAS_VERSION=0.3.30
USER root
WORKDIR /root
# Set Environment Variables for venv, cargo & openblas
ENV VIRTUAL_ENV=/opt/vllm
ENV PATH=${VIRTUAL_ENV}/bin:/root/.cargo/bin:$PATH
ENV PKG_CONFIG_PATH=/usr/local/lib/pkgconfig/
ENV LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/lib64:/usr/local/lib:/usr/lib64:/usr/lib
ENV UV_LINK_MODE=copy
ENV HOME=/root \
WHEEL_DIR=/wheelsdir \
VIRTUAL_ENV=/opt/vllm \
GRPC_PYTHON_BUILD_SYSTEM_OPENSSL=1 \
CARGO_HOME=/root/.cargo \
RUSTUP_HOME=/root/.rustup \
UV_CACHE_DIR=$HOME/.cache/uv \
PATH=/root/.cargo/bin:/root/.rustup/bin:${VIRTUAL_ENV}/bin:$PATH
# install gcc-13, python, rust, openblas
# Note: A symlink for libatomic.so is created for gcc-13 (linker fails to find libatomic otherwise - reqd. for sentencepiece)
# Note: A dummy file 'control' is created in /tmp/ to artificially create dependencies between stages when building stages in parallel
# when `--jobs=<N>` is passed with podman build command
RUN echo "DEBUG: VLLM_VERSION=${VLLM_VERSION}"
RUN --mount=type=cache,target=/var/cache/dnf \
microdnf install -y \
python${PYTHON_VERSION}-devel python${PYTHON_VERSION}-pip \
COPY --from=openblas-builder /tmp/control /dev/null
RUN --mount=type=bind,from=openblas-builder,source=/OpenBLAS-$OPENBLAS_VERSION/,target=/openblas/,rw \
dnf install -y openssl-devel \
&& dnf install -y \
git tar gcc-toolset-14 automake libtool \
pkgconfig xsimd zeromq-devel kmod findutils protobuf* \
libtiff-devel libjpeg-devel zlib-devel freetype-devel libwebp-devel \
harfbuzz-devel libraqm-devel libimagequant-devel libxcb-devel \
python${PYTHON_VERSION}-devel python${PYTHON_VERSION}-pip clang-devel \
&& dnf clean all \
&& PREFIX=/usr/local make -C /openblas install \
&& ln -sf /usr/lib64/libatomic.so.1 /usr/lib64/libatomic.so \
&& python${PYTHON_VERSION} -m venv ${VIRTUAL_ENV} \
&& python${PYTHON_VERSION} -m pip install -U pip uv --no-cache
&& python -m pip install -U pip uv \
&& uv pip install wheel build "setuptools<70" setuptools_scm setuptools_rust meson-python 'cmake<4' ninja cython scikit_build_core scikit_build \
&& curl --proto '=https' --tlsv1.2 -sSf https://sh.rustup.rs | sh -s -- -y \
&& cd /tmp && touch control
# Important: Copy only bare minimum required for the script to run
COPY requirements/ requirements/
COPY pyproject.toml ./
# The script is expected to install whatever python dependencies are missing
# as well as whatever system libraries need to be installed from source
COPY build_vllm_*.sh ./
###############################################################
# Stage to build torch family
###############################################################
FROM base-builder AS torch-builder
ARG MAX_JOBS
ARG TORCH_VERSION=2.7.0
ARG _GLIBCXX_USE_CXX11_ABI=1
ARG OPENBLAS_VERSION=0.3.30
RUN --mount=type=cache,target=/root/.cache/uv \
sh ./build_vllm_$(uname -m).sh
# copy vllm source code to build cache
COPY . .
source /opt/rh/gcc-toolset-14/enable && \
git clone --recursive https://github.com/pytorch/pytorch.git -b v${TORCH_VERSION} && \
cd pytorch && \
uv pip install -r requirements.txt && \
python setup.py develop && \
rm -f dist/torch*+git*whl && \
MAX_JOBS=${MAX_JOBS:-$(nproc)} \
PYTORCH_BUILD_VERSION=${TORCH_VERSION} PYTORCH_BUILD_NUMBER=1 uv build --wheel --out-dir /torchwheels/
ARG TORCHVISION_VERSION=0.22.0
ARG TORCHVISION_USE_NVJPEG=0
ARG TORCHVISION_USE_FFMPEG=0
RUN --mount=type=cache,target=/root/.cache/uv \
source /opt/rh/gcc-toolset-14/enable && \
pip install -U uv
git clone --recursive https://github.com/pytorch/vision.git -b v${TORCHVISION_VERSION} && \
cd vision && \
MAX_JOBS=${MAX_JOBS:-$(nproc)} \
BUILD_VERSION=${TORCHVISION_VERSION} \
uv build --wheel --out-dir /torchwheels/ --no-build-isolation
# build & install vLLM so that all transitive dependencies are build/downloaded into the uv cache
ARG TORCHAUDIO_VERSION=2.7.0
ARG BUILD_SOX=1
ARG BUILD_KALDI=1
ARG BUILD_RNNT=1
ARG USE_FFMPEG=0
ARG USE_ROCM=0
ARG USE_CUDA=0
ARG TORCHAUDIO_TEST_ALLOW_SKIP_IF_NO_FFMPEG=1
RUN --mount=type=cache,target=/root/.cache/uv \
source /opt/rh/gcc-toolset-14/enable && \
export PATH=/opt/rh/gcc-toolset-14/root/usr/bin:$PATH && \
export CC=/opt/rh/gcc-toolset-14/root/usr/bin/gcc && \
export CXX=/opt/rh/gcc-toolset-14/root/usr/bin/g++ && \
export PKG_CONFIG_PATH=/usr/local/lib/pkgconfig:/usr/lib64/pkgconfig:$PKG_CONFIG_PATH && \
export CMAKE_PREFIX_PATH=/usr/local:/usr:$CMAKE_PREFIX_PATH && \
export Protobuf_PROTOC_EXECUTABLE=/usr/bin/protoc && \
export CFLAGS="-mcpu=power10 -mtune=power10" && \
export CXXFLAGS="-mcpu=power10 -mtune=power10" && \
export DNNL_ARCH_OPT_FLAGS="-mcpu=power10 -mtune=power10" && \
export C_INCLUDE_PATH=/usr/local/include:$C_INCLUDE_PATH && \
export CPLUS_INCLUDE_PATH=/usr/local/include:$CPLUS_INCLUDE_PATH && \
uv pip install 'setuptools>=78.1.1' && \
export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/opt/OpenBLAS/lib/:/usr/local/lib64:/usr/local/lib && \
export LIBGOMP=/opt/rh/gcc-toolset-14/root/usr/lib/gcc/ppc64le-redhat-linux/14/libgomp.so && \
git clone --recursive https://github.com/pytorch/audio.git -b v${TORCHAUDIO_VERSION} && \
cd audio && \
MAX_JOBS=${MAX_JOBS:-$(nproc)} \
BUILD_VERSION=${TORCHAUDIO_VERSION} \
uv build --wheel --out-dir /torchwheels/ --no-build-isolation
export CMAKE_LIBRARY_PATH=$(dirname $LIBGOMP):${CMAKE_LIBRARY_PATH} && \
export LIBRARY_PATH=$(dirname $LIBGOMP):${LIBRARY_PATH} && \
export LD_LIBRARY_PATH=$(dirname $LIBGOMP):${LD_LIBRARY_PATH} && \
###############################################################
# Stage to build pyarrow
###############################################################
FROM base-builder AS arrow-builder
ARG MAX_JOBS
ARG PYARROW_PARALLEL
ARG PYARROW_VERSION=21.0.0
RUN --mount=type=cache,target=/root/.cache/uv \
source /opt/rh/gcc-toolset-14/enable && \
git clone --recursive https://github.com/apache/arrow.git -b apache-arrow-${PYARROW_VERSION} && \
cd arrow/cpp && \
mkdir build && cd build && \
cmake -DCMAKE_BUILD_TYPE=release \
-DCMAKE_INSTALL_PREFIX=/usr/local \
-DARROW_PYTHON=ON \
-DARROW_BUILD_TESTS=OFF \
-DARROW_JEMALLOC=ON \
-DARROW_BUILD_STATIC="OFF" \
-DARROW_PARQUET=ON \
.. && \
make install -j ${MAX_JOBS:-$(nproc)} && \
cd ../../python/ && \
uv pip install -v -r requirements-build.txt && uv pip install numpy==2.1.3 && \
PYARROW_PARALLEL=${PYARROW_PARALLEL:-$(nproc)} \
python setup.py build_ext \
--build-type=release --bundle-arrow-cpp \
bdist_wheel --dist-dir /arrowwheels/
###############################################################
# Stage to build opencv
###############################################################
FROM base-builder AS cv-builder
ARG MAX_JOBS
ARG OPENCV_VERSION=86
# patch for version 4.11.0.86
ARG OPENCV_PATCH=97f3f39
ARG ENABLE_HEADLESS=1
RUN --mount=type=cache,target=/root/.cache/uv \
source /opt/rh/gcc-toolset-14/enable && \
git clone --recursive https://github.com/opencv/opencv-python.git -b ${OPENCV_VERSION} && \
cd opencv-python && \
sed -i -E -e 's/"setuptools.+",/"setuptools",/g' pyproject.toml && \
cd opencv && git cherry-pick --no-commit $OPENCV_PATCH && cd .. && \
uv pip install scikit-build && \
python -m build --wheel --installer=uv --outdir /opencvwheels/
###############################################################
# Stage to build numactl
###############################################################
FROM base-builder AS numa-builder
# Note: Building numactl with gcc-11. Compiling with gcc-13 in this builder stage will
# trigger recompilation with gcc-11 (and require libtool) in the final stage where we do not have gcc-13
ARG MAX_JOBS
ARG NUMACTL_VERSION=2.0.19
RUN git clone --recursive https://github.com/numactl/numactl.git -b v${NUMACTL_VERSION} \
&& cd numactl \
&& autoreconf -i && ./configure \
&& make -j ${MAX_JOBS:-$(nproc)}
###############################################################
# Stage to build numba
###############################################################
FROM base-builder AS numba-builder
ARG MAX_JOBS
ARG NUMBA_VERSION=0.61.2
# Clone all required dependencies
RUN dnf install ninja-build llvm15 llvm15-devel -y && source /opt/rh/gcc-toolset-14/enable && export PATH=$PATH:/usr/lib64/llvm15/bin && \
git clone --recursive https://github.com/numba/numba.git -b ${NUMBA_VERSION} && \
cd ./numba && \
if ! grep '#include "dynamic_annotations.h"' numba/_dispatcher.cpp; then \
sed -i '/#include "internal\/pycore_atomic.h"/i\#include "dynamic_annotations.h"' numba/_dispatcher.cpp; \
fi && python -m build --wheel --installer=uv --outdir /numbawheels/
###############################################################
# Stage to build vllm - this stage builds and installs
# vllm, tensorizer and vllm-tgis-adapter and builds uv cache
# for transitive dependencies - eg. grpcio
###############################################################
FROM base-builder AS vllmcache-builder
ENV LLVM_CONFIG=/usr/lib64/llvm15/bin/llvm-config
ENV PATH=/usr/lib64/llvm15/bin:$PATH
COPY --from=torch-builder /tmp/control /dev/null
COPY --from=arrow-builder /tmp/control /dev/null
COPY --from=cv-builder /tmp/control /dev/null
COPY --from=numa-builder /tmp/control /dev/null
COPY --from=numba-builder /tmp/control /dev/null
ARG VLLM_TARGET_DEVICE=cpu
ARG GRPC_PYTHON_BUILD_SYSTEM_OPENSSL=1
# this step installs vllm and populates uv cache
# with all the transitive dependencies
RUN --mount=type=cache,target=/root/.cache/uv \
dnf install llvm15 llvm15-devel -y && \
rpm -ivh --nodeps https://mirror.stream.centos.org/9-stream/CRB/ppc64le/os/Packages/protobuf-lite-devel-3.14.0-16.el9.ppc64le.rpm && \
source /opt/rh/gcc-toolset-14/enable && \
git clone https://github.com/huggingface/xet-core.git && cd xet-core/hf_xet/ && \
uv pip install maturin && \
uv build --wheel --out-dir /hf_wheels/
ENV CXXFLAGS="-fno-lto -Wno-error=free-nonheap-object" \
CFLAGS="-fno-lto"
RUN --mount=type=cache,target=/root/.cache/uv \
--mount=type=bind,from=torch-builder,source=/torchwheels/,target=/torchwheels/,ro \
--mount=type=bind,from=arrow-builder,source=/arrowwheels/,target=/arrowwheels/,ro \
--mount=type=bind,from=cv-builder,source=/opencvwheels/,target=/opencvwheels/,ro \
--mount=type=bind,from=numa-builder,source=/numactl/,target=/numactl/,rw \
--mount=type=bind,from=numba-builder,source=/numbawheels/,target=/numbawheels/,ro \
--mount=type=bind,src=.,dst=/src/,rw \
source /opt/rh/gcc-toolset-14/enable && \
export PATH=$PATH:/usr/lib64/llvm15/bin && \
uv pip install /opencvwheels/*.whl /arrowwheels/*.whl /torchwheels/*.whl /numbawheels/*.whl && \
sed -i -e 's/.*torch.*//g' /src/pyproject.toml /src/requirements/*.txt && \
sed -i -e 's/.*sentencepiece.*//g' /src/pyproject.toml /src/requirements/*.txt && \
uv pip install sentencepiece==0.2.0 pandas pythran nanobind pybind11 /hf_wheels/*.whl && \
make -C /numactl install && \
# sentencepiece.pc is in some pkgconfig inside uv cache
export PKG_CONFIG_PATH=$(find / -type d -name "pkgconfig" 2>/dev/null | tr '\n' ':') && \
nanobind_DIR=$(uv pip show nanobind | grep Location | sed 's/^Location: //;s/$/\/nanobind\/cmake/') && uv pip install -r /src/requirements/common.txt -r /src/requirements/cpu.txt -r /src/requirements/build/cuda.txt --no-build-isolation && \
cd /src/ && \
uv build --wheel --out-dir /vllmwheel/ --no-build-isolation && \
uv pip install /vllmwheel/*.whl
###############################################################
# Stage to build lapack
###############################################################
FROM base-builder AS lapack-builder
ARG MAX_JOBS
ARG LAPACK_VERSION=3.12.1
RUN git clone --recursive https://github.com/Reference-LAPACK/lapack.git -b v${LAPACK_VERSION} \
&& cd lapack && source /opt/rh/gcc-toolset-14/enable \
&& cmake -B build -S . \
&& cmake --build build -j ${MAX_JOBS:-$(nproc)}
echo "LIBGOMP=${LIBGOMP}" && \
find /root/.cache/uv -name "*.whl" && \
SETUPTOOLS_SCM_PRETEND_VERSION="$VLLM_VERSION" uv build \
--wheel --out-dir ${WHEEL_DIR} --no-build-isolation && \
uv pip install "$(echo ${WHEEL_DIR}/vllm*.whl)[tensorizer]" --refresh
###############################################################
# FINAL VLLM IMAGE STAGE #
@@ -83,74 +278,72 @@ RUN --mount=type=cache,target=/root/.cache/uv \
FROM registry.access.redhat.com/ubi9/ubi-minimal:${BASE_UBI_IMAGE_TAG} AS vllm-openai
ARG PYTHON_VERSION=3.12
ENV VLLM_NO_USAGE_STATS=1
ARG OPENBLAS_VERSION=0.3.30
# Set Environment Variables for venv & openblas
ENV VIRTUAL_ENV=/opt/vllm
ENV PCP_DIR=/opt/rh/gcc-toolset-14/root
ENV PATH=${VIRTUAL_ENV}/bin:${PCP_DIR}/usr/bin:/usr/local/bin:$PATH
ENV PKG_CONFIG_PATH=${PCP_DIR}/usr/lib64/pkgconfig:/usr/local/lib/pkgconfig/
ENV C_INCLUDE_PATH="/usr/local/include:$C_INCLUDE_PATH"
ENV LD_LIBRARY_PATH=${PCP_DIR}/usr/lib64:${PCP_DIR}/usr/lib:${VIRTUAL_ENV}/lib64/python${PYTHON_VERSION}/site-packages/torch/lib:/usr/local/lib:$LD_LIBRARY_PATH:/usr/local/lib64:/usr/lib64:/usr/lib
ENV PATH=${VIRTUAL_ENV}/bin:$PATH
ENV PKG_CONFIG_PATH=/usr/local/lib/pkgconfig/
ENV LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/usr/local/lib64:/usr/local/lib:/usr/lib64:/usr/lib
ENV UV_LINK_MODE=copy
ARG VLLM_VERSION="0.22.1"
ARG UV_EXTRA_INDEX_URL="https://wheels.developerfirst.ibm.com/ppc64le/linux/+simple/"
ENV UV_EXTRA_INDEX_URL=${UV_EXTRA_INDEX_URL}
ENV UV_INDEX_STRATEGY=first-match
ENV OMP_NUM_THREADS=16
# create artificial dependencies between stages for independent stages to build in parallel
COPY --from=torch-builder /tmp/control /dev/null
COPY --from=arrow-builder /tmp/control /dev/null
COPY --from=cv-builder /tmp/control /dev/null
COPY --from=vllmcache-builder /tmp/control /dev/null
COPY --from=numa-builder /tmp/control /dev/null
COPY --from=lapack-builder /tmp/control /dev/null
COPY --from=openblas-builder /tmp/control /dev/null
COPY --from=numba-builder /tmp/control /dev/null
# install gcc-11, python, openblas, numactl, lapack
RUN --mount=type=cache,target=/root/.cache/uv \
--mount=type=bind,from=numa-builder,source=/numactl/,target=/numactl/,rw \
--mount=type=bind,from=lapack-builder,source=/lapack/,target=/lapack/,rw \
--mount=type=bind,from=openblas-builder,source=/OpenBLAS-$OPENBLAS_VERSION/,target=/openblas/,rw \
rpm -ivh https://dl.fedoraproject.org/pub/epel/epel-release-latest-9.noarch.rpm && \
microdnf install --nodocs -y \
libomp libicu tar autoconf automake libtool findutils openssl numactl numactl-devel \
pkgconfig xsimd gcc-toolset-14 libsndfile \
libomp libicu tar findutils openssl llvm15 llvm15-devel \
pkgconfig xsimd g++ gcc-fortran libsndfile \
libtiff libjpeg openjpeg2 zlib zeromq \
freetype lcms2 libwebp tcl tk utf8proc \
harfbuzz fribidi libraqm libimagequant libxcb util-linux gperftools-libs \
harfbuzz fribidi libraqm libimagequant libxcb util-linux \
python${PYTHON_VERSION}-devel python${PYTHON_VERSION}-pip \
&& source /opt/rh/gcc-toolset-14/enable \
&& microdnf update -y \
&& microdnf clean all
&& export PATH=$PATH:/usr/lib64/llvm15/bin && microdnf clean all \
&& python${PYTHON_VERSION} -m venv ${VIRTUAL_ENV} \
&& python -m pip install -U pip uv --no-cache \
&& make -C /numactl install \
&& PREFIX=/usr/local make -C /openblas install \
&& uv pip install 'cmake<4' \
&& cmake --install /lapack/build \
&& uv pip uninstall cmake
# The `lscpu` command was added as a requirement in part of https://github.com/vllm-project/vllm/pull/21032, so installing it.
RUN microdnf install --nodocs -y util-linux && \
microdnf clean all
# consume previously built wheels (including vllm)
RUN --mount=type=cache,target=/root/.cache/uv \
--mount=type=bind,from=torch-builder,source=/torchwheels/,target=/torchwheels/,ro \
--mount=type=bind,from=arrow-builder,source=/arrowwheels/,target=/arrowwheels/,ro \
--mount=type=bind,from=cv-builder,source=/opencvwheels/,target=/opencvwheels/,ro \
--mount=type=bind,from=vllmcache-builder,source=/hf_wheels/,target=/hf_wheels/,ro \
--mount=type=bind,from=vllmcache-builder,source=/vllmwheel/,target=/vllmwheel/,ro \
--mount=type=bind,from=numba-builder,source=/numbawheels/,target=/numbawheels/,ro \
export PKG_CONFIG_PATH=$(find / -type d -name "pkgconfig" 2>/dev/null | tr '\n' ':') && uv pip install sentencepiece==0.2.0 && \
HOME=/root uv pip install /opencvwheels/*.whl /arrowwheels/*.whl /torchwheels/*.whl /numbawheels/*.whl /hf_wheels/*.whl /vllmwheel/*.whl
COPY --from=builder-base /usr/lib64/libprotobuf.so.25 /usr/lib64/
COPY --from=builder-base /usr/lib64/libprotobuf.so.25.0.0 /usr/lib64/
# Use builder venv in final stage instead of wheel reinstallation
COPY --from=builder-base /opt/vllm /opt/vllm
COPY ./ /workspace/vllm
WORKDIR /workspace/vllm
ARG GIT_REPO_CHECK=0
RUN --mount=type=bind,source=.git,target=.git \
if [ "$GIT_REPO_CHECK" != 0 ]; then bash tools/check_repo.sh; fi
ENV LD_PRELOAD=/usr/lib64/libtcmalloc.so.4
# install development dependencies (for testing)
RUN --mount=type=cache,target=/root/.cache/uv \
uv pip install -e tests/vllm_test_utils
WORKDIR /home/vllm
WORKDIR /workspace/
# setup non-root user for OpenShift
RUN umask 002 && \
useradd --uid 2000 --gid 0 vllm && \
mkdir -p /home/vllm && \
chmod g+rwx /home/vllm
ENV HOME=/home/vllm
# Add labels to document build configuration
LABEL org.opencontainers.image.title="vLLM CPU"
LABEL org.opencontainers.image.description="vLLM inference engine for CPU platforms"
LABEL org.opencontainers.image.vendor="vLLM Project"
LABEL org.opencontainers.image.source="https://github.com/vllm-project/vllm"
# Build configuration labels
ARG TARGETARCH
ARG VLLM_CPU_PPC64LE
ARG PYTHON_VERSION
LABEL ai.vllm.build.target-arch="${TARGETARCH}"
LABEL ai.vllm.build.cpu-ppc64le="${VLLM_CPU_PPC64LE:-false}"
LABEL ai.vllm.build.python-version="${PYTHON_VERSION:-3.12}"
USER 2000
RUN ln -s /workspace/vllm/tests && ln -s /workspace/vllm/examples && ln -s /workspace/vllm/benchmarks
ENTRYPOINT ["vllm", "serve"]
-5
View File
@@ -252,8 +252,6 @@ COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/docker/docker-bake.hcl /docker/doc
COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/docker/docker-bake-rocm.hcl /docker/docker-bake-rocm.hcl
COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/.buildkite /.buildkite
COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/pyproject.toml /pyproject.toml
COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/rust /rust
COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/rust-toolchain.toml /rust-toolchain.toml
COPY --from=build_vllm ${COMMON_WORKDIR}/vllm/vllm/v1 /vllm_v1
# RIXL/UCX build stages
@@ -545,8 +543,6 @@ COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/docker/docker-bake.h
COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/docker/docker-bake-rocm.hcl /docker/docker-bake-rocm.hcl
COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/.buildkite /.buildkite
COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/pyproject.toml /pyproject.toml
COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/rust /rust
COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/rust-toolchain.toml /rust-toolchain.toml
COPY --from=build_vllm_wheel_release ${COMMON_WORKDIR}/vllm/vllm/v1 /vllm_v1
# -----------------------
@@ -580,7 +576,6 @@ RUN apt-get update -q -y && apt-get install -q -y --no-install-recommends \
libibverbs1 \
ibverbs-providers \
ibverbs-utils \
unzip \
pkg-config ffmpeg libavcodec-dev libavformat-dev libavutil-dev \
libswscale-dev libavdevice-dev libavfilter-dev libswresample-dev \
&& rm -rf /var/lib/apt/lists/*
+1 -1
View File
@@ -9,7 +9,7 @@ ARG PYTORCH_AUDIO_BRANCH="v2.9.0"
ARG PYTORCH_AUDIO_REPO="https://github.com/pytorch/audio.git"
ARG FA_BRANCH="0e60e394"
ARG FA_REPO="https://github.com/Dao-AILab/flash-attention.git"
ARG AITER_BRANCH="v0.1.16.post3"
ARG AITER_BRANCH="v0.1.16.post2"
ARG AITER_REPO="https://github.com/ROCm/aiter.git"
ARG MORI_BRANCH="v1.1.0"
ARG MORI_REPO="https://github.com/ROCm/mori.git"
+1 -2
View File
@@ -249,7 +249,7 @@ RUN --mount=type=cache,target=/root/.cache/uv \
NUMBA_WHL_FILE=$(ls /tmp/numba-wheels/*.whl) && \
OPENCV_WHL_FILE=$(ls /tmp/opencv-wheels/*.whl) && \
GUIDANCE_WHL_FILE=$(ls /tmp/guidance-wheels/*.whl) && \
uv pip install -v \
uv pip install -v \
$ARROW_WHL_FILE \
$VISION_WHL_FILE \
$HF_XET_WHL_FILE \
@@ -257,7 +257,6 @@ RUN --mount=type=cache,target=/root/.cache/uv \
$NUMBA_WHL_FILE \
$OPENCV_WHL_FILE \
$GUIDANCE_WHL_FILE \
--torch-backend cpu \
--index-strategy unsafe-best-match \
-r requirements/build/cpu.txt \
-r requirements/cpu.txt
-15
View File
@@ -50,21 +50,6 @@ vllm serve --help=max-num-seqs
vllm serve --help=max
```
!!! tip "Human-readable integer arguments"
Many integer arguments accept human-readable suffixes for convenience. For example:
- `1k` = 1,000 (decimal kilo)
- `1K` = 1,024 (binary kibibyte)
- `1m` = 1,000,000 (decimal mega)
- `1M` = 1,048,576 (binary mebibyte)
- `1g` / `1G` = 1 billion / 1 gibibyte
- `1t` / `1T` = 1 trillion / 1 tebibyte
Decimal suffixes (`k`, `m`, `g`, `t`) also accept floating point: `25.6k` = 25,600.
Binary suffixes (`K`, `M`, `G`, `T`) require integers: `32K` = 32,768.
Supported arguments include: `--max-model-len`, `--max-num-batched-tokens`, `--max-num-scheduled-tokens`, `--kv-cache-memory-bytes`, `--safetensors-prefetch-block-size`.
See [vllm serve](./serve.md) for the full reference of all available arguments.
## launch
+1 -1
View File
@@ -42,7 +42,7 @@ and the maximum batch size (`max_num_seqs` option).
```python
from vllm import LLM
llm = LLM(model="Qwen/Qwen2.5-VL-3B-Instruct", max_model_len=2048, max_num_seqs=2)
llm = LLM(model="adept/fuyu-8b", max_model_len=2048, max_num_seqs=2)
```
## Reduce CUDA Graphs
+2 -11
View File
@@ -16,14 +16,6 @@ vLLM provides 4 optimization levels (`-O0`, `-O1`, `-O2`, `-O3`) that allow user
For more information, see the [optimization level documentation](../design/optimization_levels.md).
## Faster Startup
Beyond the optimization levels, three mechanisms reduce time-to-first-token on repeated boots of the same (model, config, hardware) combination:
- **Reuse the compile cache.** vLLM persists `torch.compile` artifacts under `VLLM_CACHE_ROOT` (default `~/.cache/vllm`), and the cache directory can be copied between machines or baked into a container image; see the [torch.compile design doc](../design/torch_compile.md). Set `VLLM_FORCE_AOT_LOAD=1` to fail loudly instead of silently recompiling when the cache misses (any change to the model, config, relevant `VLLM_*` environment variables, torch build, or GPU model invalidates it).
- **Skip memory profiling with `--kv-cache-memory`.** On startup, vLLM logs the exact `--kv-cache-memory` value that reproduces the current allocation. Passing it back on the next boot skips the memory-profiling measurement and the CUDA-graph memory estimation pass. Note that this has performance implications: the KV cache is sized to exactly the given value instead of being measured, so a conservative value caps batch concurrency (and therefore throughput), while an optimistic one fails at allocation time. The value is only valid on the same GPU with the same initial free memory; if a boot OOMs after hardware or co-tenant changes, remove the flag to re-profile.
- **Serve without CUDA graphs using `--enforce-eager`.** Skips both compilation and CUDA-graph capture for the fastest possible startup, at the cost of steady-state decode performance. Useful for development loops and for measuring how much of a boot is compile/capture.
## Preemption
Due to the autoregressive nature of transformer architecture, there are times when KV cache space is insufficient to handle all batched requests.
@@ -284,9 +276,8 @@ By default vLLM uses the standard Hugging Face `tokenizers` library to power
the fast tokenizer. For BPE tokenizers (Qwen, Llama, DeepSeek, GPT-OSS, etc.)
you can switch to the [fastokens](https://github.com/crusoecloud/fastokens)
Rust backend, a drop-in replacement that's substantially faster on
encode/decode and on streaming detokenization. `VLLM_USE_FASTOKENS` is
available in vLLM v0.23.0 and later. If your installed vLLM version does not
recognize the environment variable, upgrade vLLM before enabling the override:
encode/decode and on streaming detokenization. Enable it by setting
`VLLM_USE_FASTOKENS=1`:
```console
VLLM_USE_FASTOKENS=1 vllm serve Qwen/Qwen3-8B
+1 -1
View File
@@ -14,7 +14,7 @@ Wheels are built in the `Release` pipeline (`.buildkite/release-pipeline.yaml`)
Each build step:
1. Builds the wheel in a Docker container.
2. Renames the wheel filename to use the correct manylinux tag (currently `manylinux_2_28`) for PEP 600 compliance.
2. Renames the wheel filename to use the correct manylinux tag (currently `manylinux_2_31`) for PEP 600 compliance.
3. Uploads the wheel to S3 bucket `vllm-wheels` under `/{commit_hash}/`.
### Index Generation

Some files were not shown because too many files have changed in this diff Show More