karylab_agents/vllm
1
0
Fork 0
forked from Karylab-cklius/vllm
Code Pull Requests 2 Actions 1 Packages Activity

Merge branch 'main' into wentao-optimize-per-token-group-quant

Browse Source
This commit is contained in:
Wentao Ye
2026-06-29 17:37:43 -04:00
committed by GitHub
parent 69655da6d9 8632c884dc
commit b41f7d8ffd
507 changed files with 27437 additions and 8385 deletions
Download Patch File Download Diff File Expand all files Collapse all files
.buildkite/hardware_tests/amd.yaml
+2
View File
@@ -6,6 +6,7 @@ steps:
# differ ci_base is rebuilt and pushed automatically.
- label: "AMD: :docker: ensure ci_base"
key: ensure-ci-base-amd
soft_fail: false
depends_on: []
device: amd_cpu
no_plugin: true
@@ -26,6 +27,7 @@ steps:
- label: "AMD: :docker: build test image and artifacts"
key: image-build-amd
soft_fail: false
depends_on:
- ensure-ci-base-amd
device: amd_cpu
.buildkite/hardware_tests/cpu.yaml
+3 -1
View File
@@ -91,11 +91,13 @@ steps:
- vllm/model_executor/layers/fused_moe/experts/cpu_moe.py
- tests/quantization/test_compressed_tensors.py
- tests/quantization/test_cpu_wna16.py
- tests/quantization/test_cpu_w8a8.py
commands:
- |
bash .buildkite/scripts/hardware_ci/run-cpu-test.sh 45m "
pytest -x -v -s tests/quantization/test_compressed_tensors.py::test_compressed_tensors_w8a8_logprobs
pytest -x -v -s tests/quantization/test_cpu_wna16.py"
pytest -x -v -s tests/quantization/test_cpu_wna16.py
pytest -x -v -s tests/quantization/test_cpu_w8a8.py"
- label: CPU-Distributed Tests (PP+TP)
depends_on: []
.buildkite/intel_jobs/models_multimodal_intel.yaml
-2
View File
@@ -125,7 +125,5 @@ steps:
pip install open-clip-torch --no-deps &&
cd tests &&
pytest -v -s models/multimodal/processing/test_tensor_schema.py
--deselect "tests/models/multimodal/processing/test_tensor_schema.py::test_model_tensor_schema[mistralai/Mistral-Large-3-675B-Instruct-2512-NVFP4]"
--deselect "tests/models/multimodal/processing/test_tensor_schema.py::test_model_tensor_schema[Qwen/Qwen2.5-Omni-7B-AWQ]"
--num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT --shard-id=$$BUILDKITE_PARALLEL_JOB'
parallelism: 4
.buildkite/scripts/ci-bake-rocm.sh
+180 -16
View File
@@ -15,9 +15,10 @@ 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 tools/install_torchcodec_rocm.sh tests/vllm_test_utils"
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"
IMAGE_EXISTED_BEFORE_BUILD=0
TARGET=""
@@ -525,6 +526,22 @@ get_remote_image_label_with_retry() {
return 0
}
remote_ci_base_metadata_is_current() {
local image_ref="$1"
local metadata_version=""
metadata_version=$(get_remote_image_label "${image_ref}" "vllm.ci_base.metadata_version")
[[ "${metadata_version}" == "${CI_BASE_METADATA_VERSION:-${DEFAULT_CI_BASE_METADATA_VERSION}}" ]]
}
remote_ci_base_metadata_is_current_with_retry() {
local image_ref="$1"
local metadata_version=""
metadata_version=$(get_remote_image_label_with_retry "${image_ref}" "vllm.ci_base.metadata_version")
[[ "${metadata_version}" == "${CI_BASE_METADATA_VERSION:-${DEFAULT_CI_BASE_METADATA_VERSION}}" ]]
}
remote_image_exists() {
local image_ref="$1"
docker manifest inspect "${image_ref}" >/dev/null 2>&1
@@ -581,6 +598,7 @@ init_config() {
CI_BASE_CONTENT_FILES="${CI_BASE_CONTENT_FILES:-${DEFAULT_CI_BASE_CONTENT_FILES}}"
CI_BASE_DOCKERFILE="${CI_BASE_DOCKERFILE:-${DEFAULT_CI_BASE_DOCKERFILE}}"
CI_BASE_DOCKERFILE_STAGES="${CI_BASE_DOCKERFILE_STAGES:-${DEFAULT_CI_BASE_DOCKERFILE_STAGES}}"
CI_BASE_METADATA_VERSION="${CI_BASE_METADATA_VERSION:-${DEFAULT_CI_BASE_METADATA_VERSION}}"
CI_BASE_IMAGE_TAG="${CI_BASE_IMAGE_TAG:-rocm/vllm-dev:ci_base}"
export PYTORCH_ROCM_ARCH
@@ -635,6 +653,10 @@ load_ci_hcl() {
echo "Copied ${CI_HCL_SOURCE} to ${CI_HCL_PATH}"
}
init_bake_files() {
BAKE_FILES=(-f "${VLLM_BAKE_FILE}" -f "${CI_HCL_PATH}")
}
compute_ci_base_hash_if_needed() {
if [[ -z "${CI_BASE_CONTENT_FILES:-}" ]]; then
return 0
@@ -676,12 +698,14 @@ configure_ci_base_image_refs() {
fi
content_tag=$(ci_base_tag_with_suffix "${stable_tag}" "${CI_BASE_CONTENT_HASH}")
CI_BASE_IMAGE_TAG_CONTENT_REF="${content_tag}"
if [[ -n "${BUILDKITE_COMMIT:-}" ]]; then
commit_tag=$(ci_base_tag_with_suffix "${stable_tag}" "${BUILDKITE_COMMIT}")
CI_BASE_IMAGE_TAG_COMMIT="${commit_tag}"
export CI_BASE_IMAGE_TAG_COMMIT
fi
CI_BASE_IMAGE_TAG_COMMIT_REF="${commit_tag}"
# *_REF is the logical tag recorded in metadata. *_EXTRA is only passed to
# bake when that tag is not already the primary tag, avoiding duplicates.
if should_push_stable_ci_base_tag; then
primary_tag="${content_tag}"
CI_BASE_IMAGE_TAG_STABLE="${stable_tag}"
@@ -691,19 +715,33 @@ configure_ci_base_image_refs() {
fi
CI_BASE_IMAGE_TAG="${primary_tag}"
if [[ "${primary_tag}" == "${content_tag}" ]]; then
CI_BASE_IMAGE_TAG_CONTENT=""
CI_BASE_IMAGE_TAG_CONTENT_EXTRA=""
else
CI_BASE_IMAGE_TAG_CONTENT="${content_tag}"
CI_BASE_IMAGE_TAG_CONTENT_EXTRA="${content_tag}"
fi
export CI_BASE_IMAGE_TAG CI_BASE_IMAGE_TAG_CONTENT CI_BASE_IMAGE_TAG_STABLE
if [[ -n "${commit_tag}" && "${commit_tag}" != "${primary_tag}" ]]; then
CI_BASE_IMAGE_TAG_COMMIT_EXTRA="${commit_tag}"
else
CI_BASE_IMAGE_TAG_COMMIT_EXTRA=""
fi
export CI_BASE_IMAGE_TAG
export CI_BASE_IMAGE_TAG_COMMIT_EXTRA
export CI_BASE_IMAGE_TAG_CONTENT_EXTRA
export CI_BASE_IMAGE_TAG_CONTENT_REF
export CI_BASE_IMAGE_TAG_COMMIT_REF
export CI_BASE_IMAGE_TAG_STABLE
if is_ci_base_target; then
IMAGE_TAG="${primary_tag}"
export IMAGE_TAG
echo "ci_base primary image tag: ${CI_BASE_IMAGE_TAG}"
if [[ -n "${CI_BASE_IMAGE_TAG_COMMIT:-}" ]]; then
echo "ci_base commit image tag: ${CI_BASE_IMAGE_TAG_COMMIT}"
if [[ -n "${commit_tag}" ]]; then
if [[ "${commit_tag}" == "${primary_tag}" ]]; then
echo "ci_base commit image tag: ${commit_tag} (primary)"
else
echo "ci_base commit image tag: ${commit_tag}"
fi
fi
echo "ci_base content image tag: ${content_tag}"
if [[ -n "${CI_BASE_IMAGE_TAG_STABLE}" ]]; then
@@ -728,8 +766,8 @@ ci_base_candidate_refs() {
printf '%s\n' \
"${IMAGE_TAG:-}" \
"${CI_BASE_IMAGE_TAG:-}" \
"${CI_BASE_IMAGE_TAG_COMMIT:-}" \
"${CI_BASE_IMAGE_TAG_CONTENT:-}" \
"${CI_BASE_IMAGE_TAG_COMMIT_EXTRA:-}" \
"${CI_BASE_IMAGE_TAG_CONTENT_EXTRA:-}" \
"${CI_BASE_IMAGE_TAG_STABLE:-}" \
| awk 'NF && !seen[$0]++'
}
@@ -743,6 +781,10 @@ find_matching_ci_base_ref() {
remote_image_exists "${candidate}" || continue
candidate_hash=$(get_remote_image_label "${candidate}" "vllm.ci_base.content_hash")
if [[ "${candidate_hash}" == "${CI_BASE_CONTENT_HASH}" ]]; then
if ! remote_ci_base_metadata_is_current "${candidate}"; then
echo "Found matching ci_base content hash but stale metadata: ${candidate}" >&2
continue
fi
printf '%s\n' "${candidate}"
return 0
fi
@@ -817,6 +859,10 @@ maybe_skip_existing_image() {
if [[ -n "${remote_hash}" ]]; then
echo "Remote ci_base content hash: ${remote_hash:0:16}..."
if [[ "${remote_hash}" == "${CI_BASE_CONTENT_HASH}" ]]; then
if ! remote_ci_base_metadata_is_current "${IMAGE_TAG}"; then
echo "Content hashes match but ci_base metadata is stale; rebuilding to refresh metadata"
return 0
fi
if ! refresh_ci_base_tags_from_ref "${IMAGE_TAG}"; then
echo "ci_base tag refresh failed; rebuilding to push expected tags"
return 0
@@ -998,12 +1044,104 @@ prepare_git_cache_metadata() {
fi
}
ci_base_metadata_pairs() {
local dockerfile="${CI_BASE_DOCKERFILE:-${DEFAULT_CI_BASE_DOCKERFILE}}"
local stages="${CI_BASE_DOCKERFILE_STAGES:-${DEFAULT_CI_BASE_DOCKERFILE_STAGES}}"
local content_files="${CI_BASE_CONTENT_FILES:-${DEFAULT_CI_BASE_CONTENT_FILES}}"
local content_files_hash=""
local base_image=""
local base_image_digest=""
local git_branch=""
local -a content_paths=()
local -a content_args=()
read -r -a content_paths <<< "${content_files}"
if [[ ${#content_paths[@]} -gt 0 ]]; then
content_files_hash=$(compute_content_hash "${content_paths[@]}")
fi
mapfile -t content_args < <(
get_content_arg_names "${dockerfile}" "${stages}" "${CI_BASE_CONTENT_ARGS:-}"
)
base_image=$(resolve_dockerfile_arg_value "${dockerfile}" "BASE_IMAGE")
if [[ -n "${base_image}" ]]; then
base_image_digest=$(resolve_image_digest "${base_image}")
fi
git_branch="${BUILDKITE_BRANCH:-${VLLM_BRANCH:-}}"
metadata_pair "vllm.ci_base.metadata_version" "${CI_BASE_METADATA_VERSION:-${DEFAULT_CI_BASE_METADATA_VERSION}}"
metadata_pair "vllm.ci_base.content_hash" "${CI_BASE_CONTENT_HASH:-}"
metadata_pair "vllm.ci_base.content_files_hash" "${content_files_hash}"
metadata_pair "vllm.ci_base.content_files" "${content_files}"
metadata_pair "vllm.ci_base.content_args" "$(join_words "${content_args[@]}")"
metadata_pair "vllm.ci_base.dockerfile" "${dockerfile}"
metadata_pair "vllm.ci_base.dockerfile_stages" "${stages}"
metadata_pair "vllm.ci_base.image.primary" "${CI_BASE_IMAGE_TAG:-}"
metadata_pair "vllm.ci_base.image.content" "${CI_BASE_IMAGE_TAG_CONTENT_REF:-${CI_BASE_IMAGE_TAG_CONTENT_EXTRA:-}}"
metadata_pair "vllm.ci_base.image.commit" "${CI_BASE_IMAGE_TAG_COMMIT_REF:-${CI_BASE_IMAGE_TAG_COMMIT_EXTRA:-}}"
metadata_pair "vllm.ci_base.image.stable" "${CI_BASE_IMAGE_TAG_STABLE:-}"
metadata_pair "vllm.ci_base.git_commit" "${BUILDKITE_COMMIT:-}"
metadata_pair "vllm.ci_base.git_branch" "${git_branch}"
metadata_pair "vllm.ci_base.vllm_branch" "${VLLM_BRANCH:-}"
metadata_pair "vllm.ci_base.stable_branch" "${CI_BASE_STABLE_BRANCH:-main}"
metadata_pair "vllm.rocm.base_image" "${base_image}"
metadata_pair "vllm.rocm.base_image_digest" "${base_image_digest}"
metadata_pair "vllm.rocm.pytorch_rocm_arch" "${PYTORCH_ROCM_ARCH:-}"
metadata_pair "vllm.rocm.nic_backend" "$(resolve_dockerfile_arg_value "${dockerfile}" "NIC_BACKEND")"
metadata_pair "vllm.rocm.ainic_version" "$(resolve_dockerfile_arg_value "${dockerfile}" "AINIC_VERSION")"
metadata_pair "vllm.rocm.ubuntu_codename" "$(resolve_dockerfile_arg_value "${dockerfile}" "UBUNTU_CODENAME")"
metadata_pair "vllm.rocm.rixl_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "RIXL_REPO")"
metadata_pair "vllm.rocm.rixl_commit" "${RIXL_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "RIXL_BRANCH")}"
metadata_pair "vllm.rocm.ucx_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "UCX_REPO")"
metadata_pair "vllm.rocm.ucx_commit" "${UCX_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "UCX_BRANCH")}"
metadata_pair "vllm.rocm.rocshmem_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "ROCSHMEM_REPO")"
metadata_pair "vllm.rocm.rocshmem_commit" "${ROCSHMEM_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "ROCSHMEM_BRANCH")}"
metadata_pair "vllm.rocm.deepep_repo" "$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_REPO")"
metadata_pair "vllm.rocm.deepep_commit" "${DEEPEP_BRANCH:-$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_BRANCH")}"
metadata_pair "vllm.rocm.deepep_nic" "$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_NIC")"
metadata_pair "vllm.rocm.deepep_rocm_arch" "$(resolve_dockerfile_arg_value "${dockerfile}" "DEEPEP_ROCM_ARCH")"
metadata_pair "vllm.rocm.rixl_cache_key" "${RIXL_CACHE_KEY:-}"
metadata_pair "vllm.rocm.rocshmem_cache_key" "${ROCSHMEM_CACHE_KEY:-}"
metadata_pair "vllm.rocm.deepep_cache_key" "${DEEPEP_CACHE_KEY:-}"
metadata_pair "vllm.buildkite.build_number" "${BUILDKITE_BUILD_NUMBER:-}"
metadata_pair "vllm.buildkite.build_id" "${BUILDKITE_BUILD_ID:-}"
}
write_ci_base_metadata_annotations() {
local metadata="$1"
local key=""
local value=""
local annotation=""
[[ -n "${metadata}" ]] || return 0
while IFS=$'\t' read -r key value; do
[[ -n "${key}" && -n "${value}" ]] || continue
annotation="manifest:${key}=${value}"
printf ' "%s",\n' "$(hcl_escape_string "${annotation}")"
done <<< "${metadata}"
}
write_ci_base_metadata_labels() {
local metadata="$1"
local key=""
local value=""
[[ -n "${metadata}" ]] || return 0
while IFS=$'\t' read -r key value; do
[[ -n "${key}" && -n "${value}" ]] || continue
printf ' "%s" = "%s"\n' \
"$(hcl_escape_string "${key}")" \
"$(hcl_escape_string "${value}")"
done <<< "${metadata}"
}
write_ci_base_label_override() {
local target_name=""
local metadata=""
local -a ci_base_targets=()
BAKE_FILES=(-f "${VLLM_BAKE_FILE}" -f "${CI_HCL_PATH}")
if [[ -z "${CI_BASE_CONTENT_HASH:-}" ]]; then
return 0
fi
@@ -1019,16 +1157,23 @@ write_ci_base_label_override() {
return 0
fi
metadata=$(ci_base_metadata_pairs)
: > "${CI_BASE_LABEL_OVERRIDE_PATH}"
for target_name in "${ci_base_targets[@]}"; do
cat >> "${CI_BASE_LABEL_OVERRIDE_PATH}" <<EOF
target "${target_name}" {
annotations = [
"manifest:org.opencontainers.image.revision=",
EOF
write_ci_base_metadata_annotations "${metadata}" >> "${CI_BASE_LABEL_OVERRIDE_PATH}"
cat >> "${CI_BASE_LABEL_OVERRIDE_PATH}" <<EOF
]
labels = {
"org.opencontainers.image.revision" = ""
"vllm.ci_base.content_hash" = "${CI_BASE_CONTENT_HASH}"
EOF
write_ci_base_metadata_labels "${metadata}" >> "${CI_BASE_LABEL_OVERRIDE_PATH}"
cat >> "${CI_BASE_LABEL_OVERRIDE_PATH}" <<EOF
}
}
@@ -1036,7 +1181,7 @@ EOF
done
BAKE_FILES+=(-f "${CI_BASE_LABEL_OVERRIDE_PATH}")
echo "Appended ci_base content-hash label override for targets: ${ci_base_targets[*]}"
echo "Appended ci_base metadata label override for targets: ${ci_base_targets[*]}"
}
uses_rocm_csrc_cache() {
@@ -1119,6 +1264,18 @@ hcl_escape_string() {
printf '%s' "${value}"
}
join_words() {
local IFS=" "
printf '%s' "$*"
}
metadata_pair() {
local key="$1"
local value="${2:-}"
printf '%s\t%s\n' "${key}" "${value}"
}
write_hcl_string_list() {
local indent="$1"
shift
@@ -1541,7 +1698,13 @@ confirm_remote_image_push() {
remote_hash=$(get_remote_image_label_with_retry "${image_ref}" "vllm.ci_base.content_hash")
if [[ -n "${remote_hash}" && "${remote_hash}" == "${CI_BASE_CONTENT_HASH}" ]]; then
return 0
if remote_ci_base_metadata_is_current_with_retry "${image_ref}"; then
return 0
fi
echo "Remote image exists with the expected ci_base content hash but stale metadata."
echo " expected metadata version: ${CI_BASE_METADATA_VERSION:-${DEFAULT_CI_BASE_METADATA_VERSION}}"
return 1
fi
echo "Remote image exists but does not have the expected ci_base content hash."
@@ -1724,15 +1887,16 @@ main() {
print_header
validate_inputs
load_ci_hcl
init_bake_files
compute_ci_base_hash_if_needed
configure_ci_base_image_refs
maybe_skip_existing_image
setup_builder
prepare_git_cache_metadata
write_ci_base_label_override
extract_dependency_pins
write_rocm_build_arg_override
compute_dependency_cache_keys
write_ci_base_label_override
compute_rocm_csrc_content_hash_if_needed
write_rocm_cache_override
resolve_ci_base_dependency_targets
.buildkite/scripts/hardware_ci/run-cpu-test.sh
+38
View File
@@ -12,6 +12,44 @@ IMAGE_NAME="cpu-test-${NUMA_NODE}${AGENT_SLOT:+-${AGENT_SLOT}}"
TIMEOUT_VAL=$1
TEST_COMMAND=$2
# Disk hygiene knobs. Reclaim space only once the Docker root filesystem crosses
# DISK_USAGE_THRESHOLD percent, and cap the shared BuildKit cache at
# BUILDKIT_CACHE_MAX so subsequent builds keep reusing the hottest layers.
DISK_USAGE_THRESHOLD=${DISK_USAGE_THRESHOLD:-70}
BUILDKIT_CACHE_MAX=${BUILDKIT_CACHE_MAX:-80GB}
# Reclaim disk only when the host is under pressure. We trim (not purge) the
# shared BuildKit cache so cross-job/cross-agent reuse stays intact, and only
# touch dangling images; other agents' uniquely tagged images are left alone.
prune_if_disk_pressure() {
local docker_root disk_usage
docker_root=$(docker info -f '{{.DockerRootDir}}' 2>/dev/null || true)
if [ -z "$docker_root" ]; then
return 0
fi
disk_usage=$(df "$docker_root" 2>/dev/null | tail -1 | awk '{print $5}' | tr -d '%')
if [ "${disk_usage:-0}" -gt "$DISK_USAGE_THRESHOLD" ]; then
echo "--- :broom: Disk usage ${disk_usage}% exceeds ${DISK_USAGE_THRESHOLD}%, reclaiming space"
docker image prune -f || true
docker builder prune -f --keep-storage="$BUILDKIT_CACHE_MAX" || true
else
echo "Disk usage ${disk_usage:-unknown}% within ${DISK_USAGE_THRESHOLD}% threshold; skipping prune"
fi
}
# Always drop this agent's image once the job ends (the default builder never
# uses it as a cache source, so removing it costs no rebuild speed), then
# reclaim space if needed. Guard every docker call with `|| true` so the trap
# never overrides the test's exit code.
cleanup() {
docker image rm -f "$IMAGE_NAME" || true
prune_if_disk_pressure
}
trap cleanup EXIT
# Free space up front so a nearly-full host doesn't fail the build.
prune_if_disk_pressure
# building the docker image
echo "--- :docker: Building Docker image"
docker build --progress plain --tag "$IMAGE_NAME" --target vllm-test -f docker/Dockerfile.cpu .
.buildkite/scripts/run-rust-frontend-cargo-ci.sh
+18
View File
@@ -90,6 +90,16 @@ install_cargo_sort() {
cargo binstall --no-confirm cargo-sort
}
install_cargo_deny() {
if command -v cargo-deny >/dev/null 2>&1; then
return
fi
log_section "Installing cargo-deny"
install_cargo_binstall
cargo binstall --no-confirm cargo-deny
}
install_cargo_nextest() {
if command -v cargo-nextest >/dev/null 2>&1; then
return
@@ -142,6 +152,7 @@ PY
run_style_clippy() {
install_cargo_sort
install_cargo_deny
log_section "Checking Rust formatting"
cargo fmt --manifest-path rust/Cargo.toml --all -- --check
@@ -149,6 +160,13 @@ run_style_clippy() {
log_section "Checking Cargo.toml ordering"
cargo sort --workspace --check rust
log_section "Checking Rust dependency bans"
cargo deny \
--manifest-path rust/Cargo.toml \
check \
--config rust/deny.toml \
bans
log_section "Running clippy"
cargo clippy \
--manifest-path rust/Cargo.toml \
.buildkite/test-amd.yaml
+86 -203
View File
@@ -112,33 +112,6 @@ steps:
# #
#########################################################################################################################################
#----------------------------------------------------- mi250 · basic_correctness -----------------------------------------------------#
- label: Distributed Model Tests (2 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_2
num_gpus: 2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/model_executor/model_loader/sharded_state_loader.py
- vllm/model_executor/models/
- vllm/model_executor/layers/
- vllm/v1/attention/backends/
- vllm/v1/attention/selector.py
- tests/basic_correctness/
- tests/model_executor/model_loader/test_sharded_state_loader.py
- tests/models/
- vllm/_aiter_ops.py
- vllm/platforms/rocm.py
commands:
- TARGET_TEST_SUITE=MI250 pytest basic_correctness/ -v -s -m 'distributed(num_gpus=2)'
- CUDA_VISIBLE_DEVICES=0,1 pytest -v -s model_executor/model_loader/test_sharded_state_loader.py -m '(not slow_test)'
- pytest models/test_transformers.py -v -s -m 'distributed(num_gpus=2)'
- pytest models/language -v -s -m 'distributed(num_gpus=2)'
- pytest models/multimodal -v -s -m 'distributed(num_gpus=2)' --ignore models/multimodal/generation/test_whisper.py
- VLLM_WORKER_MULTIPROC_METHOD=spawn pytest models/multimodal/generation/test_whisper.py -v -s -m 'distributed(num_gpus=2)'
#---------------------------------------------------------- mi250 · compile ----------------------------------------------------------#
- label: PyTorch Compilation Unit Tests # TBD
@@ -179,48 +152,8 @@ steps:
commands:
- "find compile/fullgraph/ -name 'test_*.py' -not -name 'test_full_graph.py' -exec pytest -s -v {} \\\\;"
- label: Distributed Compile + RPC Tests (2 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_2
num_gpus: 2
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/compilation/
- vllm/distributed/
- vllm/engine/
- vllm/executor/
- vllm/worker/worker_base.py
- vllm/v1/engine/
- vllm/v1/worker/
- tests/compile/fullgraph/test_basic_correctness.py
- tests/compile/test_wrapper.py
- tests/entrypoints/llm/test_collective_rpc.py
- vllm/platforms/rocm.py
commands:
- pytest -v -s entrypoints/llm/test_collective_rpc.py
- pytest -v -s ./compile/fullgraph/test_basic_correctness.py
- pytest -v -s ./compile/test_wrapper.py
#-------------------------------------------------------- mi250 · distributed --------------------------------------------------------#
- label: Distributed Comm Ops # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_2
num_gpus: 2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/distributed
- tests/distributed
- vllm/platforms/rocm.py
commands:
- pytest -v -s distributed/test_comm_ops.py
- pytest -v -s distributed/test_shm_broadcast.py
- pytest -v -s distributed/test_shm_buffer.py
- pytest -v -s distributed/test_shm_storage.py
- label: Pipeline + Context Parallelism (4 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
@@ -330,54 +263,6 @@ steps:
- pytest -v -s models/multimodal/generation/test_common.py -m core_model -k "not qwen2 and not qwen3 and not gemma"
- pytest -v -s models/multimodal/generation/test_qwen2_vl.py -m core_model
#---------------------------------------------------------- mi250 · plugins ----------------------------------------------------------#
- label: Plugin Tests (2 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_2
num_gpus: 2
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/plugins/
- tests/plugins/
- vllm/platforms/rocm.py
commands:
# BEGIN: platform plugin and general plugin tests, all the code in-between runs on dummy platform
- pip install -e ./plugins/vllm_add_dummy_platform
- pytest -v -s plugins_tests/test_platform_plugins.py
- pip uninstall vllm_add_dummy_platform -y
# END: platform plugin tests
# BEGIN: `io_processor` plugins test, all the code in between uses the `prithvi_io_processor` plugin
- pip install -e ./plugins/prithvi_io_processor_plugin
- pytest -v -s plugins_tests/test_io_processor_plugins.py
- pytest -v -s plugins_tests/test_terratorch_io_processor_plugins.py
- pip uninstall prithvi_io_processor_plugin -y
# END: `io_processor` plugins test
# BEGIN: `bge_m3_sparse io_processor` test
- pip install -e ./plugins/bge_m3_sparse_plugin
- pytest -v -s plugins_tests/test_bge_m3_sparse_io_processor_plugins.py
- pip uninstall bge_m3_sparse_plugin -y
# END: `bge_m3_sparse io_processor` test
# BEGIN: `colbert_query io_processor` test
- pip install -e ./plugins/colbert_query_plugin
- pytest -v -s plugins_tests/test_colbert_query_io_processor_plugins.py
- pip uninstall colbert_query_plugin -y
# END: `colbert_query io_processor` test
# BEGIN: `stat_logger` plugins test
- pip install -e ./plugins/vllm_add_dummy_stat_logger
- pytest -v -s plugins_tests/test_stats_logger_plugins.py
- pip uninstall dummy_stat_logger -y
# END: `stat_logger` plugins test
# BEGIN: other tests
- pytest -v -s plugins_tests/test_scheduler_plugins.py
- pip install -e ./plugins/vllm_add_dummy_model
- pytest -v -s distributed/test_distributed_oot.py
- pytest -v -s plugins_tests/test_oot_registration_online.py # it needs a clean process
- pytest -v -s plugins_tests/test_oot_registration_offline.py # it needs a clean process
- pytest -v -s plugins_tests/lora_resolvers # unit tests for in-tree lora resolver plugins
#------------------------------------------------------------ mi250 · v1 -------------------------------------------------------------#
- label: Batch Invariance (H100-MI250) # TBD
@@ -505,61 +390,6 @@ steps:
commands:
- pytest -v -s v1/attention
- label: V1 Sample + Logits # TBD
timeout_in_minutes: 60
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_1
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- tests/v1/sample
- tests/v1/logits_processors
- tests/v1/test_oracle.py
- tests/v1/test_request.py
- tests/v1/test_outputs.py
commands:
- pytest -v -s v1/sample
- pytest -v -s v1/logits_processors
- pytest -v -s v1/test_oracle.py
- pytest -v -s v1/test_request.py
- pytest -v -s v1/test_outputs.py
- label: Distributed DP Tests (2 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_2
num_gpus: 2
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/distributed/
- vllm/engine/
- vllm/executor/
- vllm/worker/worker_base.py
- vllm/v1/engine/
- vllm/v1/worker/
- tests/v1/distributed
- tests/entrypoints/openai/test_multi_api_servers.py
- vllm/platforms/rocm.py
commands:
- TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_async_llm_dp.py
- TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_eagle_dp.py
- TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
- DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py
- label: V1 e2e (2 GPUs) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx90anightly, amdmi250]
agent_pool: mi250_2
optional: true
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
- tests/v1/e2e
commands:
- pytest -v -s v1/e2e/spec_decode/test_spec_decode.py -k "tensor_parallelism"
#------------------------------------------------------------- mi250 · misc ------------------------------------------------------------#
- label: Async Engine, Inputs, Utils, Worker, Config (CPU) # TBD
@@ -608,7 +438,7 @@ steps:
#----------------------------------------------------- mi300 · basic_correctness -----------------------------------------------------#
- label: Basic Correctness # TBD
timeout_in_minutes: 40
timeout_in_minutes: 50
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
fast_check: true
@@ -626,7 +456,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: 65
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_2
num_gpus: 2
@@ -771,7 +601,7 @@ steps:
#----------------------------------------------------------- mi300 · cuda ------------------------------------------------------------#
- label: Platform Tests (CUDA) # TBD
- label: Platform Tests # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
@@ -802,6 +632,22 @@ steps:
#-------------------------------------------------------- mi300 · distributed --------------------------------------------------------#
- label: Distributed Comm Ops # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_2
num_gpus: 2
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/distributed
- tests/distributed
- vllm/platforms/rocm.py
commands:
- pytest -v -s distributed/test_comm_ops.py
- pytest -v -s distributed/test_shm_broadcast.py
- pytest -v -s distributed/test_shm_buffer.py
- pytest -v -s distributed/test_shm_storage.py
- label: EPLB Algorithm # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
@@ -831,8 +677,8 @@ steps:
- pytest -v -s distributed/test_eplb_execute.py
- pytest -v -s distributed/test_eplb_spec_decode.py
- label: Distributed Tests (2xH100-2xMI250) # TBD
timeout_in_minutes: 180
- label: Distributed Tests (2xH100-2xMI300) # TBD
timeout_in_minutes: 30
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_2
num_gpus: 2
@@ -843,13 +689,19 @@ steps:
- vllm/model_executor/layers/fused_moe/
- vllm/v1/attention/backends/
- vllm/v1/attention/selector.py
- tests/v1/distributed/test_dbo.py
- tests/distributed/test_context_parallel.py
- examples/features/data_parallel/data_parallel_offline.py
- vllm/_aiter_ops.py
- vllm/platforms/rocm.py
commands:
- pytest -v -s tests/distributed/test_context_parallel.py
- VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_async_new_apis.py
- VLLM_LOGGING_LEVEL=DEBUG python3 examples/features/data_parallel/data_parallel_offline.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=deepep_high_throughput
- VLLM_LOGGING_LEVEL=DEBUG python3 examples/features/data_parallel/data_parallel_offline.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=allgather_reducescatter --disable-nccl-for-dp-synchronization
- pytest -v -s tests/v1/distributed/test_dbo.py
- VLLM_ALLOW_INSECURE_SERIALIZATION=1 pytest -v -s tests/distributed/test_weight_transfer.py
- pytest -v -s tests/distributed/test_packed_tensor.py
- label: Distributed Tests (4xA100-4xMI300) # TBD
timeout_in_minutes: 180
@@ -993,10 +845,12 @@ steps:
source_file_dependencies:
- vllm/
- tests/entrypoints/serve
- tests/entrypoints/scale_out
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
- pytest -v -s entrypoints/scale_out
- label: Entrypoints Integration (API Server OpenAI - Part 1) # TBD
timeout_in_minutes: 180
@@ -1347,7 +1201,7 @@ steps:
#--------------------------------------------------------- mi300 · examples ----------------------------------------------------------#
- label: Examples # TBD
timeout_in_minutes: 180
timeout_in_minutes: 45
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_1
optional: true
@@ -1360,7 +1214,7 @@ steps:
commands:
- pip install tensorizer
# Basic
- python3 basic/offline_inference/chat.py --attention-backend TRITON_ATTN
- python3 basic/offline_inference/chat.py
- python3 basic/offline_inference/generate.py --model facebook/opt-125m
- python3 basic/offline_inference/generate.py --model meta-llama/Llama-2-13b-chat-hf --cpu-offload-gb 10
- python3 basic/offline_inference/classify.py
@@ -1814,6 +1668,54 @@ steps:
- python3 examples/generate/multimodal/vision_language_offline.py --model-type qwen2_5_vl
- VLLM_WORKER_MULTIPROC_METHOD=spawn python3 examples/generate/multimodal/audio_language_offline.py --model-type whisper
#---------------------------------------------------------- mi300 · plugins ----------------------------------------------------------#
- label: Plugin Tests (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/plugins/
- tests/plugins/
- vllm/platforms/rocm.py
commands:
# BEGIN: platform plugin and general plugin tests, all the code in-between runs on dummy platform
- pip install -e ./plugins/vllm_add_dummy_platform
- pytest -v -s plugins_tests/test_platform_plugins.py
- pip uninstall vllm_add_dummy_platform -y
# END: platform plugin tests
# BEGIN: `io_processor` plugins test, all the code in between uses the `prithvi_io_processor` plugin
- pip install -e ./plugins/prithvi_io_processor_plugin
- pytest -v -s plugins_tests/test_io_processor_plugins.py
- pytest -v -s plugins_tests/test_terratorch_io_processor_plugins.py
- pip uninstall prithvi_io_processor_plugin -y
# END: `io_processor` plugins test
# BEGIN: `bge_m3_sparse io_processor` test
- pip install -e ./plugins/bge_m3_sparse_plugin
- pytest -v -s plugins_tests/test_bge_m3_sparse_io_processor_plugins.py
- pip uninstall bge_m3_sparse_plugin -y
# END: `bge_m3_sparse io_processor` test
# BEGIN: `colbert_query io_processor` test
- pip install -e ./plugins/colbert_query_plugin
- pytest -v -s plugins_tests/test_colbert_query_io_processor_plugins.py
- pip uninstall colbert_query_plugin -y
# END: `colbert_query io_processor` test
# BEGIN: `stat_logger` plugins test
- pip install -e ./plugins/vllm_add_dummy_stat_logger
- pytest -v -s plugins_tests/test_stats_logger_plugins.py
- pip uninstall dummy_stat_logger -y
# END: `stat_logger` plugins test
# BEGIN: other tests
- pytest -v -s plugins_tests/test_scheduler_plugins.py
- pip install -e ./plugins/vllm_add_dummy_model
- pytest -v -s distributed/test_distributed_oot.py
- pytest -v -s plugins_tests/test_oot_registration_online.py # it needs a clean process
- pytest -v -s plugins_tests/test_oot_registration_offline.py # it needs a clean process
- pytest -v -s plugins_tests/lora_resolvers # unit tests for in-tree lora resolver plugins
#------------------------------------------------------- mi300 · quantization --------------------------------------------------------#
- label: Quantization # TBD
@@ -2195,26 +2097,6 @@ steps:
- TP_SIZE=1 DP_SIZE=2 pytest -v -s v1/distributed/test_external_lb_dp.py
- DP_SIZE=2 pytest -v -s entrypoints/openai/test_multi_api_servers.py
- label: Distributed Tests (2xH100-2xMI300) # TBD
timeout_in_minutes: 180
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
agent_pool: mi300_2
num_gpus: 2
working_dir: "/vllm-workspace/"
source_file_dependencies:
- vllm/distributed/
- vllm/v1/distributed/
- vllm/model_executor/layers/fused_moe/
- tests/v1/distributed/test_dbo.py
- vllm/_aiter_ops.py
- vllm/platforms/rocm.py
commands:
- VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_async_new_apis.py
- VLLM_LOGGING_LEVEL=DEBUG python3 examples/features/data_parallel/data_parallel_offline.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048 --all2all-backend=deepep_high_throughput
- pytest -v -s tests/v1/distributed/test_dbo.py
- VLLM_ALLOW_INSECURE_SERIALIZATION=1 pytest -v -s tests/distributed/test_weight_transfer.py
- pytest -v -s tests/distributed/test_packed_tensor.py
- label: Metrics, Tracing (2 GPUs) # TBD
timeout_in_minutes: 20
mirror_hardwares: [amdexperimental, amdproduction, amdgfx942nightly, amdmi300]
@@ -2679,10 +2561,12 @@ steps:
source_file_dependencies:
- vllm/
- tests/entrypoints/serve
- tests/entrypoints/scale_out
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
- pytest -v -s entrypoints/scale_out
- label: Entrypoints Integration (API Server OpenAI - Part 1) # TBD
timeout_in_minutes: 180
@@ -2804,13 +2688,12 @@ 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-gfx950.txt
- label: LM Eval Qwen3-5 Models (B200-MI355) %N # TBD
timeout_in_minutes: 180
- label: LM Eval Qwen3-5 Models (B200-MI355) # TBD
timeout_in_minutes: 120
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_2
num_gpus: 2
optional: true
parallelism: 4
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/model_executor/models/qwen3_5.py
@@ -2825,7 +2708,7 @@ steps:
- vllm/_aiter_ops.py
- vllm/platforms/rocm.py
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-qwen35-mi355.txt --shard-id=$$BUILDKITE_PARALLEL_JOB --num-shards=$$BUILDKITE_PARALLEL_JOB_COUNT
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-qwen35-mi355.txt
- label: LM Eval Small Models (2xB200-2xMI355) # TBD
timeout_in_minutes: 180
@@ -2888,7 +2771,7 @@ steps:
#--------------------------------------------------------- mi355 · examples ----------------------------------------------------------#
- label: Examples # TBD
timeout_in_minutes: 180
timeout_in_minutes: 45
mirror_hardwares: [amdexperimental, amdproduction, amdgfx950nightly, amdmi355]
agent_pool: mi355_1
working_dir: "/vllm-workspace/examples"
@@ -2900,7 +2783,7 @@ steps:
commands:
- pip install tensorizer
# Basic
- python3 basic/offline_inference/chat.py --attention-backend TRITON_ATTN
- python3 basic/offline_inference/chat.py
- python3 basic/offline_inference/generate.py --model facebook/opt-125m
- python3 basic/offline_inference/generate.py --model meta-llama/Llama-2-13b-chat-hf --cpu-offload-gb 10
- python3 basic/offline_inference/classify.py
.buildkite/test_areas/basic_correctness.yaml
+1 -1
View File
@@ -19,6 +19,6 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 40
timeout_in_minutes: 50
depends_on:
- image-build-amd
.buildkite/test_areas/cuda.yaml
+2 -2
View File
@@ -2,8 +2,8 @@ group: CUDA
depends_on:
- image-build
steps:
- label: Platform Tests (CUDA)
key: platform-tests-cuda
- label: Platform Tests
key: platform-tests
timeout_in_minutes: 15
device: h200_18gb
source_file_dependencies:
.buildkite/test_areas/distributed.yaml
+7 -19
View File
@@ -174,8 +174,8 @@ steps:
# test multi-node TP with multiproc executor (simulated on single node)
- pytest -v -s distributed/test_multiproc_executor.py::test_multiproc_executor_multi_node
- label: Distributed Tests (8 GPUs)(H100)
key: distributed-tests-8-gpus-h100
- label: Distributed Tests (8xH100)
key: distributed-tests-8xh100
timeout_in_minutes: 10
device: h100
num_devices: 8
@@ -195,8 +195,8 @@ steps:
# test with torchrun tp=2 and dp=4 with ep
- torchrun --nproc-per-node=8 ../examples/features/torchrun/torchrun_dp_example_offline.py --tp-size=2 --pp-size=1 --dp-size=4 --enable-ep
- label: Distributed Tests (4 GPUs)(A100)
key: distributed-tests-4-gpus-a100
- label: Distributed Tests (4xA100)
key: distributed-tests-4xa100
device: a100
optional: true
num_devices: 4
@@ -211,7 +211,7 @@ steps:
- pytest -v -s -x lora/test_mixtral.py
- label: Distributed Tests (2xH100-2xMI300)
key: distributed-tests-2-gpus-h100
key: distributed-tests-2xh100-2xmi300
timeout_in_minutes: 15
device: h100
optional: true
@@ -224,21 +224,9 @@ steps:
- pytest -v -s tests/v1/distributed/test_dbo.py
- VLLM_ALLOW_INSECURE_SERIALIZATION=1 pytest -v -s tests/distributed/test_weight_transfer.py
- pytest -v -s tests/distributed/test_packed_tensor.py
mirror:
amd:
device: mi300_2
timeout_in_minutes: 180
depends_on:
- image-build-amd
commands:
- VLLM_ALLOW_INSECURE_SERIALIZATION=1 python3 examples/rl/rlhf_async_new_apis.py
- VLLM_LOGGING_LEVEL=DEBUG python3 examples/features/data_parallel/data_parallel_offline.py --model=Qwen/Qwen1.5-MoE-A2.7B -tp=1 -dp=2 --max-model-len=2048
- pytest -v -s tests/v1/distributed/test_dbo.py
- VLLM_ALLOW_INSECURE_SERIALIZATION=1 pytest -v -s tests/distributed/test_weight_transfer.py
- pytest -v -s tests/distributed/test_packed_tensor.py
- label: Distributed Tests (2 GPUs)(B200)
key: distributed-tests-2-gpus-b200
- label: Distributed Tests (2xB200)
key: distributed-tests-2xb200
device: b200-k8s
optional: true
working_dir: "/vllm-workspace/"
.buildkite/test_areas/e2e_integration.yaml
+6 -6
View File
@@ -2,8 +2,8 @@ group: E2E Integration
depends_on:
- image-build
steps:
- label: DeepSeek V2-Lite Sync EPLB Accuracy
key: deepseek-v2-lite-sync-eplb-accuracy
- label: DeepSeek V2-Lite Sync EPLB Accuracy (4xH100)
key: deepseek-v2-lite-sync-eplb-accuracy-4xh100
timeout_in_minutes: 60
device: h100
optional: true
@@ -12,8 +12,8 @@ steps:
commands:
- bash .buildkite/scripts/scheduled_integration_test/deepseek_v2_lite_ep_eplb.sh 0.25 200 8010
- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy
key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy
- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (4xH100)
key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-4xh100
timeout_in_minutes: 60
device: h100
optional: true
@@ -22,8 +22,8 @@ steps:
commands:
- bash .buildkite/scripts/scheduled_integration_test/qwen30b_a3b_fp8_block_ep_eplb.sh 0.8 200 8020
- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (B200)
key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-b200
- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (2xB200)
key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-2xb200
timeout_in_minutes: 60
device: b200-k8s
optional: true
.buildkite/test_areas/entrypoints.yaml
+4
View File
@@ -29,6 +29,8 @@ steps:
mirror:
amd:
device: mi325_1
# TODO(akaratza): Test after Torch >= 2.12 bump
soft_fail: true
depends_on:
- image-build-amd
@@ -40,10 +42,12 @@ steps:
source_file_dependencies:
- vllm/
- tests/entrypoints/serve
- tests/entrypoints/scale_out
commands:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
- pytest -v -s entrypoints/serve --ignore=entrypoints/serve/dev/rpc
- PYTHONPATH=/vllm-workspace pytest -v -s entrypoints/serve/dev/rpc
- pytest -v -s entrypoints/scale_out
mirror:
amd:
device: mi325_1
.buildkite/test_areas/kernels.yaml
+4 -4
View File
@@ -274,8 +274,8 @@ steps:
- pytest -v -s kernels/helion/
- label: Kernels FP8 MoE Test (1 H100)
key: kernels-fp8-moe-test-1-h100
- label: Kernels FP8 MoE Test (1xH100)
key: kernels-fp8-moe-test-1xh100
timeout_in_minutes: 90
device: h100
num_devices: 1
@@ -291,8 +291,8 @@ steps:
- pytest -v -s kernels/moe/test_triton_moe_no_act_mul.py
- pytest -v -s kernels/moe/test_triton_moe_ptpc_fp8.py
- label: Kernels FP8 MoE Test (2 H100s)
key: kernels-fp8-moe-test-2-h100s
- label: Kernels FP8 MoE Test (2xH100)
key: kernels-fp8-moe-test-2xh100
timeout_in_minutes: 90
device: h100
num_devices: 2
.buildkite/test_areas/lm_eval.yaml
+30 -15
View File
@@ -28,7 +28,8 @@ steps:
- vllm/_aiter_ops.py
- vllm/platforms/rocm.py
# - label: LM Eval Large Models (4 GPUs)(A100)
# - label: LM Eval Large Models (4xA100)
# key: lm-eval-large-models-4xa100
# device: a100
# optional: true
# num_devices: 4
@@ -40,8 +41,8 @@ steps:
# - export VLLM_WORKER_MULTIPROC_METHOD=spawn
# - pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large.txt --tp-size=4
- label: LM Eval Large Models (4 GPUs)(H100)
key: lm-eval-large-models-4-gpus-h100
- label: LM Eval Large Models (4xH100)
key: lm-eval-large-models-4xh100
device: h100
optional: true
num_devices: 4
@@ -53,8 +54,8 @@ steps:
- export VLLM_USE_DEEP_GEMM=0 # We found Triton is faster than DeepGEMM for H100
- pytest -s -v test_lm_eval_correctness.py --config-list-file=configs/models-large-hopper.txt --tp-size=4
- label: LM Eval Small Models (B200)
key: lm-eval-small-models-b200
- label: LM Eval Small Models (2xB200)
key: lm-eval-small-models-2xb200
timeout_in_minutes: 120
device: b200-k8s
optional: true
@@ -64,8 +65,20 @@ steps:
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-blackwell.txt
- label: LM Eval Large Models (B200, EP)
key: lm-eval-large-models-b200-ep
- label: LM Eval Small Models (2xL4)
key: lm-eval-small-models-tp
timeout_in_minutes: 10
num_devices: 2
optional: true
source_file_dependencies:
- csrc/
- vllm/model_executor/layers/quantization
autorun_on_main: true
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-small-tp.txt
- label: LM Eval Large Models EP (2xB200)
key: lm-eval-large-models-ep-2xb200
timeout_in_minutes: 120
device: b200-k8s
optional: true
@@ -76,8 +89,8 @@ steps:
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-blackwell-ep.txt
- label: LM Eval Qwen3.5 Models (B200)
key: lm-eval-qwen3-5-models-b200
- label: LM Eval Qwen3.5 Models (2xB200)
key: lm-eval-qwen3-5-models-2xb200
timeout_in_minutes: 120
device: b200-k8s
optional: true
@@ -93,8 +106,8 @@ steps:
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=configs/models-qwen35-blackwell.txt
- label: LM Eval Large Models (H200)
key: lm-eval-large-models-h200
- label: LM Eval Large Models (8xH200)
key: lm-eval-large-models-8xh200
timeout_in_minutes: 60
device: h200
optional: true
@@ -192,8 +205,8 @@ steps:
commands:
- pytest -s -v evals/gsm8k/test_gsm8k_correctness.py --config-list-file=evals/gsm8k/configs/models-turboquant.txt
- label: GPQA Eval (GPT-OSS) (H100)
key: gpqa-eval-gpt-oss-h100
- label: GPQA Eval (GPT-OSS) (2xH100)
key: gpqa-eval-gpt-oss-2xh100
timeout_in_minutes: 120
device: h100
optional: true
@@ -206,8 +219,8 @@ 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-h100.txt
- label: GPQA Eval (GPT-OSS) (B200)
key: gpqa-eval-gpt-oss-b200
- label: GPQA Eval (GPT-OSS) (2xB200)
key: gpqa-eval-gpt-oss-2xb200
timeout_in_minutes: 120
device: b200-k8s
optional: true
@@ -226,6 +239,8 @@ steps:
device: dgx-spark
optional: true
num_devices: 1
depends_on:
- arm64-image-build
source_file_dependencies:
- csrc/
- vllm/model_executor/layers/quantization
.buildkite/test_areas/misc.yaml
+10
View File
@@ -224,6 +224,16 @@ steps:
- python3 features/speculative_decoding/spec_decode_offline.py --test --method eagle --num_spec_tokens 3 --dataset-name hf --dataset-path philschmid/mt-bench --num-prompts 80 --temp 0 --top-p 1.0 --top-k -1 --tp 1 --enable-chunked-prefill --max-model-len 2048
# https://github.com/vllm-project/vllm/pull/26682 uses slightly more memory in PyTorch 2.9+ causing this test to OOM in 1xL4 GPU
- python3 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
mirror:
amd:
device: mi325_1
source_file_dependencies:
- vllm/entrypoints
- vllm/multimodal
- examples/
- vllm/platforms/rocm.py
depends_on:
- image-build-amd
- label: Metrics, Tracing (2 GPUs)
key: metrics-tracing-2-gpus
.buildkite/test_areas/rust_frontend.yaml
+2 -2
View File
@@ -46,7 +46,7 @@ steps:
- vllm/v1/engine/
- tests/utils.py
# - tests/entrypoints/serve/dev/rpc/test_collective_rpc.py
- tests/entrypoints/serve/disagg/test_serving_tokens.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
@@ -55,7 +55,7 @@ steps:
- export VLLM_WORKER_MULTIPROC_METHOD=spawn
# - pytest -v -s entrypoints/serve/dev/rpc/test_collective_rpc.py
- pytest -v -s entrypoints/serve/instrumentator/test_basic.py -k "not show_version and not server_load"
- pytest -v -s entrypoints/serve/disagg/test_serving_tokens.py -k "not stream and not lora and not test_generate_logprobs and not stop_string_workflow"
- pytest -v -s entrypoints/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
.buildkite/test_areas/spec_decode.yaml
+2
View File
@@ -94,6 +94,8 @@ steps:
amd:
device: mi325_1
timeout_in_minutes: 65
# TODO(akaratza): Test after Torch >= 2.12 bump
soft_fail: true
depends_on:
- image-build-amd
source_file_dependencies:
.github/CODEOWNERS
+10 -8
View File
@@ -3,7 +3,7 @@
# This lists cover the "core" components of vLLM that require careful review
/vllm/compilation @zou3519 @youkaichao @ProExpertProg @BoyuanFeng
/vllm/distributed/kv_transfer @NickLucche @ApostaC @orozery @xuechendi
/vllm/distributed/kv_transfer @NickLucche @ApostaC @orozery @xuechendi @ivanium
/vllm/lora @jeejeelee
/vllm/model_executor/layers/attention @LucasWilkinson @MatthewBonanni
/vllm/model_executor/layers/fused_moe @mgoin @pavanimajety @zyongye
@@ -11,7 +11,7 @@
/vllm/model_executor/layers/mamba @tdoublep @tomeras91
/vllm/model_executor/layers/mamba/gdn/qwen_gdn_linear_attn.py @tdoublep @ZJY0516 @vadiklyutiy
/vllm/model_executor/model_loader @22quinn
/vllm/model_executor/layers/batch_invariant.py @yewentao256
/vllm/model_executor/layers/batch_invariant.py @yewentao256
/vllm/ir @ProExpertProg
/vllm/kernels/ @ProExpertProg @tjtanaa
/vllm/kernels/helion @ProExpertProg @zou3519
@@ -23,7 +23,7 @@
# Any change to the VllmConfig changes can have a large user-facing impact,
# so spam a lot of people
/vllm/config @WoosukKwon @youkaichao @robertgshaw2-redhat @mgoin @tlrmchlsmth @houseroad @yewentao256 @ProExpertProg
/vllm/config/cache.py @heheda12345
/vllm/config/cache.py @heheda12345 @ivanium
# Config utils
/vllm/config/utils.py @hmellor
@@ -67,16 +67,17 @@
/vllm/v1/attention/backends/flashinfer.py @mgoin @pavanimajety @vadiklyutiy
/vllm/v1/attention/backends/triton_attn.py @tdoublep
/vllm/v1/attention/backends/gdn_attn.py @ZJY0516 @vadiklyutiy
/vllm/v1/core @WoosukKwon @robertgshaw2-redhat @njhill @ywang96 @alexm-redhat @heheda12345 @ApostaC @orozery
/vllm/v1/core @WoosukKwon @robertgshaw2-redhat @njhill @ywang96 @alexm-redhat @heheda12345 @ApostaC @orozery @ivanium
/vllm/v1/sample @22quinn @houseroad @njhill
/vllm/v1/spec_decode @benchislett @luccafong @MatthewBonanni
/vllm/v1/structured_output @mgoin @russellb @aarnphm @benchislett
/vllm/v1/kv_cache_interface.py @heheda12345
/vllm/v1/kv_cache_interface.py @heheda12345 @ivanium
/vllm/v1/kv_offload @ApostaC @orozery
/vllm/v1/simple_kv_offload @ivanium
/vllm/v1/engine @njhill
/vllm/v1/executor @njhill
/vllm/v1/worker @njhill
/vllm/v1/worker/kv_connector_model_runner_mixin.py @orozery @NickLucche
/vllm/v1/worker/kv_connector_model_runner_mixin.py @orozery @NickLucche @ivanium
# Model runner V2
/vllm/v1/worker/gpu @WoosukKwon @njhill @yewentao256
@@ -103,13 +104,14 @@
/tests/test_inputs.py @DarkLight1337 @ywang96
/tests/entrypoints/llm/test_struct_output_generate.py @mgoin @russellb @aarnphm
/tests/v1/structured_output @mgoin @russellb @aarnphm
/tests/v1/core @WoosukKwon @robertgshaw2-redhat @njhill @ywang96 @alexm-redhat @heheda12345 @ApostaC @orozery
/tests/v1/core @WoosukKwon @robertgshaw2-redhat @njhill @ywang96 @alexm-redhat @heheda12345 @ApostaC @orozery @ivanium
/tests/weight_loading @mgoin @youkaichao @yewentao256
/tests/lora @jeejeelee
/tests/models/language/generation/test_hybrid.py @tdoublep @tomeras91
/tests/v1/kv_connector/nixl_integration @NickLucche
/tests/v1/kv_connector @ApostaC @orozery
/tests/v1/kv_connector @ApostaC @orozery @ivanium
/tests/v1/kv_offload @ApostaC @orozery
/tests/v1/simple_kv_offload @ivanium
/tests/v1/determinism @yewentao256
/tests/reasoning @aarnphm @chaunceyjiang @sfeng33 @bbrowning
/tests/tool_parsers @aarnphm @chaunceyjiang @sfeng33 @bbrowning
.github/actionlint.yaml
+2
View File
@@ -3,3 +3,5 @@
self-hosted-runner:
labels:
- vllm-runners
# Not yet in actionlint's known-label set.
- macos-26
.github/workflows/macos-smoke-test.yml
+17 -8
View File
@@ -11,7 +11,19 @@ permissions:
jobs:
macos-m1-smoke-test:
runs-on: macos-latest
# macos-26 (the supported target) is still a preview runner, so gate on GA
# macos-15 and keep macos-26 non-blocking.
strategy:
fail-fast: false
matrix:
include:
- os: macos-15
required: true
- os: macos-26
required: false
name: macos-m1-smoke-test (${{ matrix.os }})
runs-on: ${{ matrix.os }}
continue-on-error: ${{ !matrix.required }}
timeout-minutes: 30
steps:
@@ -72,14 +84,11 @@ jobs:
# Test health endpoint
curl -f http://localhost:8000/health
# Test completion
curl -f http://localhost:8000/v1/completions \
# Long prompt: hits the split-KV path that short prompts skip (#46769).
PAYLOAD=$(python -c "import json; print(json.dumps({'model': 'Qwen/Qwen3-0.6B', 'prompt': 'The quick brown fox jumps over the lazy dog. ' * 24, 'max_tokens': 16}))")
curl -f --max-time 120 http://localhost:8000/v1/completions \
-H "Content-Type: application/json" \
-d '{
"model": "Qwen/Qwen3-0.6B",
"prompt": "Hello",
"max_tokens": 5
}'
-d "$PAYLOAD"
# Cleanup
kill "$SERVER_PID"
CMakeLists.txt
+25
View File
@@ -140,6 +140,21 @@ if(Python_VERSION VERSION_GREATER_EQUAL "3.11")
WITH_SOABI)
endif()
#
# fs_io extension (pure CXX; must stay above the non-CUDA device branch
# so CPU builds define the target before the early return).
# GIL-releasing filesystem helpers for FileSystemTierManager.
#
if(Python_VERSION VERSION_GREATER_EQUAL "3.11")
define_extension_target(
fs_io_C
DESTINATION vllm
LANGUAGE CXX
SOURCES csrc/fs_io.cpp
USE_SABI 3.11
WITH_SOABI)
endif()
#
# Forward the non-CUDA device extensions to external CMake scripts.
#
@@ -270,6 +285,16 @@ if(VLLM_GPU_LANG STREQUAL "HIP")
#
set(CMAKE_${VLLM_GPU_LANG}_FLAGS "${CMAKE_${VLLM_GPU_LANG}_FLAGS} -Wno-unused-result -Wno-unused-value")
set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-unused-result -Wno-unused-value")
# When using LTO then *.cpp files must be compiled with same compiler as used linker
# So if HIP uses clang linker we also must use it
# Otherwise symbols will be missing from .so
if (CMAKE_CXX_FLAGS MATCHES "\-flto")
if(NOT CMAKE_CXX_COMPILER_ID STREQUAL CMAKE_HIP_COMPILER_ID)
message(FATAL_ERROR "LTO is enabled for ROCm build, but the C++ compiler (${CMAKE_CXX_COMPILER_ID}) and HIP compiler (${CMAKE_HIP_COMPILER_ID}) are different which is not supported. "
"Please ensure they are same by setting CXX=${CMAKE_HIP_COMPILER} environment variable. Or alternatively disable LTO.")
endif()
endif()
endif()
#
cmake/cpu_extension.cmake
+3
View File
@@ -24,7 +24,10 @@ set (ENABLE_NUMA TRUE)
# Check the compile flags
#
if(MACOSX_FOUND)
# Apple clang needs -Xpreprocessor to enable OpenMP. No runtime link is
# needed: _C is a dynamic_lookup bundle and resolves libomp from torch.
list(APPEND CXX_COMPILE_FLAGS
"-Xpreprocessor" "-fopenmp"
"-DVLLM_CPU_EXTENSION")
else()
list(APPEND CXX_COMPILE_FLAGS
cmake/external_projects/vllm_flash_attn.cmake
+1 -1
View File
@@ -39,7 +39,7 @@ else()
FetchContent_Declare(
vllm-flash-attn
GIT_REPOSITORY https://github.com/vllm-project/flash-attention.git
GIT_TAG 803020a8fa15407871341d41eba4919ade2ee1ee
GIT_TAG b3964b1d8b95d8e8447435668ab169a2700bab65
GIT_PROGRESS TRUE
# Don't share the vllm-flash-attn build between build types
BINARY_DIR ${CMAKE_BINARY_DIR}/vllm-flash-attn
csrc/cpu/cpu_attn_impl.hpp
+3 -3
View File
@@ -124,7 +124,7 @@ struct AttentionMetadata {
workitem_group_num(workitem_group_num),
reduction_item_num(reduction_item_num),
reduction_split_num(reduction_split_num),
thread_num(omp_get_max_threads()),
thread_num(cpu_utils::get_max_threads()),
effective_thread_num(thread_num),
split_kv_q_token_num_threshold(split_kv_q_token_num_threshold),
attention_scratchpad_size_per_thread(0),
@@ -405,7 +405,7 @@ class AttentionScheduler {
torch::Tensor schedule(const ScheduleInput& input) const {
const bool causal = input.causal;
const bool is_dynamic_causal = input.dynamic_causal != nullptr;
const int32_t thread_num = omp_get_max_threads();
const int32_t thread_num = cpu_utils::get_max_threads();
const int64_t cache_size = cpu_utils::get_available_l2_size();
const int32_t max_num_q_per_iter = input.max_num_q_per_iter;
const int32_t kv_len_alignment = input.kv_block_alignment;
@@ -1423,7 +1423,7 @@ class AttentionMainLoop {
public:
void operator()(const AttentionInput* input) {
const int thread_num = omp_get_max_threads();
const int thread_num = cpu_utils::get_max_threads();
TORCH_CHECK_EQ(input->metadata->thread_num, thread_num);
std::atomic<int32_t> guard_counter(0);
std::atomic<int32_t>* guard_counter_ptr = &guard_counter;
csrc/cpu/cpu_fused_moe.cpp
+1 -1
View File
@@ -267,7 +267,7 @@ void fused_moe_impl(scalar_t* __restrict__ output, scalar_t* __restrict__ input,
TORCH_CHECK_EQ(output_size_2 % gemm_n_tile_size, 0);
TORCH_CHECK_EQ(output_size_13 / 2, input_size_2);
const int32_t thread_num = omp_get_max_threads();
const int32_t thread_num = cpu_utils::get_max_threads();
const int32_t w13_input_buffer_size = cpu_utils::round_up<64>(
gemm_m_tile_size * input_size_13 * sizeof(scalar_t));
csrc/cpu/cpu_types.hpp
+16
View File
@@ -25,4 +25,20 @@
#include <omp.h>
#endif
#include <c10/util/Exception.h>
namespace cpu_utils {
// Without OpenMP the omp pragmas compile to serial loops, so report 1: kernels
// that barrier on the thread count would otherwise deadlock.
inline int get_max_threads() {
#ifdef _OPENMP
return omp_get_max_threads();
#else
TORCH_WARN_ONCE(
"vLLM CPU was built without OpenMP; running single-threaded.");
return 1;
#endif
}
} // namespace cpu_utils
#endif
csrc/cpu/cpu_wna16.cpp
+1 -1
View File
@@ -155,7 +155,7 @@ void cpu_gemm_wna16_impl(
constexpr int32_t gemm_m_tile_size = gemm_t::MaxMSize;
constexpr int32_t n_block_size = 16;
static_assert(gemm_n_tile_size % n_block_size == 0);
const int32_t thread_num = omp_get_max_threads();
const int32_t thread_num = cpu_utils::get_max_threads();
// a simple schedule policy, just to hold more B tiles in L2 and make sure
// each thread has tasks
csrc/cpu/dnnl_kernels.cpp
+1 -1
View File
@@ -202,7 +202,7 @@ void dynamic_quant_epilogue(const float* input, scalar_t* output,
using cvt_vec_t = typename KernelVecType<scalar_t>::cvt_vec_type;
constexpr int vec_elem_num = load_vec_t::VEC_ELEM_NUM;
const int64_t thread_num = omp_get_max_threads();
const int64_t thread_num = cpu_utils::get_max_threads();
if (num_tokens > thread_num) {
#pragma omp parallel for
for (int64_t i = 0; i < num_tokens; ++i) {
csrc/cpu/mla_decode.cpp
+1 -1
View File
@@ -251,7 +251,7 @@ void mla_decode_kvcache_cpu_impl(
constexpr int QK_NUM_ELEM = qk_vec_type::VEC_ELEM_NUM;
// shared across threads
const int max_threads = omp_get_max_threads();
const int max_threads = cpu_utils::get_max_threads();
const int acc_out_nbytes =
max_threads * num_heads * V_HEAD_DIM * sizeof(float);
float* acc_out = static_cast<float*>(std::aligned_alloc(64, acc_out_nbytes));
csrc/custom_all_reduce_test.cu
-361
View File
@@ -1,361 +0,0 @@
/**
* This is a standalone test for custom allreduce.
* To compile, make sure you have MPI and NCCL installed in your system.
* export MPI_HOME=XXX
* nvcc -O2 -arch=native -std=c++17 custom_all_reduce_test.cu -o
* custom_all_reduce_test -lnccl -I${MPI_HOME}/include -lmpi
*
* Warning: this C++ test is not designed to be very readable and was used
* during the rapid prototyping process.
*
* To run:
* mpirun --allow-run-as-root -np 8 ./custom_all_reduce_test
*/
#include <cuda.h>
#include <curand_kernel.h>
#include <stdio.h>
#include <stdlib.h>
#include <limits>
#include <vector>
#include "cuda_profiler_api.h"
#include "custom_all_reduce.cuh"
#include "mpi.h"
#ifdef USE_ROCM
#include <hip/hip_bf16.h>
typedef __hip_bfloat16 nv_bfloat16;
#include "rccl/rccl.h"
#include "custom_all_reduce_hip.cuh"
#else
#include "nccl.h"
#include "custom_all_reduce.cuh"
#endif
#define MPICHECK(cmd) \
do { \
int e = cmd; \
if (e != MPI_SUCCESS) { \
printf("Failed: MPI error %s:%d '%d'\n", __FILE__, __LINE__, e); \
exit(EXIT_FAILURE); \
} \
} while (0)
#define NCCLCHECK(cmd) \
do { \
ncclResult_t r = cmd; \
if (r != ncclSuccess) { \
printf("Failed, NCCL error %s:%d '%s'\n", __FILE__, __LINE__, \
ncclGetErrorString(r)); \
exit(EXIT_FAILURE); \
} \
} while (0)
#ifdef USE_ROCM
__global__ void dummy_kernel() {
for (int i = 0; i < 100; i++) {
uint64_t start = wall_clock64();
uint64_t cycles_elapsed;
do {
cycles_elapsed = wall_clock64() - start;
} while (cycles_elapsed < 100);
}
for (int i = 0; i < 100; i++) __nanosleep(1000000); // 100ms
}
#else
__global__ void dummy_kernel() {
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 700
for (int i = 0; i < 100; i++) __nanosleep(1000000); // 100ms
#else
for (int i = 0; i < 100; i++) {
long long int start = clock64();
while (clock64() - start < 150000000); // approximately 98.4ms on P40
}
#endif
}
#endif
template <typename T>
__global__ void set_data(T* data, int size, int myRank) {
for (int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < size;
idx += gridDim.x * blockDim.x) {
data[idx] = myRank * 0.11f;
}
}
template <typename T>
__global__ void convert_data(const T* data1, const T* data2, double* fdata1,
double* fdata2, int size) {
for (int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < size;
idx += gridDim.x * blockDim.x) {
fdata1[idx] = data1[idx];
fdata2[idx] = data2[idx];
}
}
__global__ void init_rand(curandState_t* state, int size, int nRanks) {
for (int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < size;
idx += gridDim.x * blockDim.x) {
for (int i = 0; i < nRanks; i++) {
curand_init(i + 1, idx, 0, &state[idx * nRanks + i]);
}
}
}
template <typename T>
__global__ void gen_data(curandState_t* state, T* data, double* ground_truth,
int myRank, int nRanks, int size) {
for (int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < size;
idx += gridDim.x * blockDim.x) {
double sum = 0.0;
for (int i = 0; i < nRanks; i++) {
double val = curand_uniform_double(&state[idx * nRanks + i]) * 4;
T hval = val; // downcast first
sum += static_cast<double>(hval);
if (i == myRank) data[idx] = hval;
}
ground_truth[idx] = sum;
}
}
template <typename T>
void run(int myRank, int nRanks, ncclComm_t& comm, int threads, int block_limit,
int data_size, bool performance_test) {
T* result;
cudaStream_t stream;
CUDACHECK(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking));
CUDACHECK(cudaMalloc(&result, data_size * sizeof(T)));
CUDACHECK(cudaMemset(result, 0, data_size * sizeof(T)));
cudaIpcMemHandle_t self_data_handle;
cudaIpcMemHandle_t data_handles[8];
vllm::Signal* buffer;
T* self_data_copy;
/**
* Allocate IPC buffer
*
* The first section is a temporary buffer for storing intermediate allreduce
* results, if a particular algorithm requires it. The second section is for
* the input to the allreduce. The actual API takes the input pointer as an
* argument (that is, they can and usually should be allocated separately).
* But since the input pointers and the temporary buffer all require IPC
* registration, they are allocated and registered together in the test for
* convenience.
*/
#ifdef USE_ROCM
CUDACHECK(hipExtMallocWithFlags(
(void**)&buffer, 2 * data_size * sizeof(T) + sizeof(vllm::Signal),
hipDeviceMallocUncached));
#else
CUDACHECK(
cudaMalloc(&buffer, 2 * data_size * sizeof(T) + sizeof(vllm::Signal)));
#endif
CUDACHECK(
cudaMemset(buffer, 0, 2 * data_size * sizeof(T) + sizeof(vllm::Signal)));
CUDACHECK(cudaMalloc(&self_data_copy, data_size * sizeof(T)));
CUDACHECK(cudaIpcGetMemHandle(&self_data_handle, buffer));
MPICHECK(MPI_Allgather(&self_data_handle, sizeof(cudaIpcMemHandle_t),
MPI_BYTE, data_handles, sizeof(cudaIpcMemHandle_t),
MPI_BYTE, MPI_COMM_WORLD));
void* rank_data;
size_t rank_data_sz = 16 * 1024 * 1024;
CUDACHECK(cudaMalloc(&rank_data, rank_data_sz));
vllm::Signal* ipc_ptrs[8];
for (int i = 0; i < nRanks; i++) {
if (i == myRank)
ipc_ptrs[i] = buffer;
else
CUDACHECK(cudaIpcOpenMemHandle((void**)&ipc_ptrs[i], data_handles[i],
cudaIpcMemLazyEnablePeerAccess));
}
vllm::CustomAllreduce fa(ipc_ptrs, rank_data, rank_data_sz, myRank, nRanks);
auto* self_data =
reinterpret_cast<T*>(reinterpret_cast<char*>(buffer) +
sizeof(vllm::Signal) + data_size * sizeof(T));
// hack buffer registration
{
void* data[8];
for (int i = 0; i < nRanks; i++) {
data[i] =
((char*)ipc_ptrs[i]) + sizeof(vllm::Signal) + data_size * sizeof(T);
}
fa.register_buffer(data);
}
double* ground_truth;
CUDACHECK(cudaMallocHost(&ground_truth, data_size * sizeof(double)));
curandState_t* states;
CUDACHECK(cudaMalloc(&states, sizeof(curandState_t) * nRanks * data_size));
init_rand<<<108, 1024, 0, stream>>>(states, data_size, nRanks);
gen_data<T><<<108, 1024, 0, stream>>>(states, self_data, ground_truth, myRank,
nRanks, data_size);
CUDACHECK(cudaMemcpyAsync(self_data_copy, self_data, data_size * sizeof(T),
cudaMemcpyDeviceToDevice, stream));
cudaEvent_t start, stop;
CUDACHECK(cudaEventCreate(&start));
CUDACHECK(cudaEventCreate(&stop));
ncclDataType_t ncclDtype;
if (std::is_same<T, half>::value) {
ncclDtype = ncclFloat16;
} else if (std::is_same<T, nv_bfloat16>::value) {
ncclDtype = ncclBfloat16;
} else {
ncclDtype = ncclFloat;
}
double *nccl_result, *my_result;
CUDACHECK(cudaMallocHost(&nccl_result, data_size * sizeof(double)));
CUDACHECK(cudaMallocHost(&my_result, data_size * sizeof(double)));
if (performance_test) {
dummy_kernel<<<1, 1, 0, stream>>>();
constexpr int warmup_iters = 5;
constexpr int num_iters = 100;
// warmup
for (int i = 0; i < warmup_iters; i++) {
NCCLCHECK(ncclAllReduce(result, result, data_size, ncclDtype, ncclSum,
comm, stream));
}
CUDACHECK(cudaEventRecord(start, stream));
for (int i = 0; i < num_iters; i++) {
NCCLCHECK(ncclAllReduce(result, result, data_size, ncclDtype, ncclSum,
comm, stream));
}
CUDACHECK(cudaEventRecord(stop, stream));
CUDACHECK(cudaStreamSynchronize(stream));
float allreduce_ms = 0;
cudaEventElapsedTime(&allreduce_ms, start, stop);
dummy_kernel<<<1, 1, 0, stream>>>();
// warm up
for (int i = 0; i < warmup_iters; i++) {
fa.allreduce<T>(stream, self_data, result, data_size, threads,
block_limit);
}
CUDACHECK(cudaEventRecord(start, stream));
for (int i = 0; i < num_iters; i++) {
fa.allreduce<T>(stream, self_data, result, data_size, threads,
block_limit);
}
CUDACHECK(cudaEventRecord(stop, stream));
CUDACHECK(cudaStreamSynchronize(stream));
float duration_ms = 0;
cudaEventElapsedTime(&duration_ms, start, stop);
if (myRank == 0)
printf(
"Rank %d done, nGPUs:%d, sz (kb): %d, %d, %d, my time:%.2fus, nccl "
"time:%.2fus\n",
myRank, nRanks, data_size * sizeof(T) / 1024, threads, block_limit,
duration_ms * 1e3 / num_iters, allreduce_ms * 1e3 / num_iters);
// And wait for all the queued up work to complete
CUDACHECK(cudaStreamSynchronize(stream));
NCCLCHECK(ncclAllReduce(self_data_copy, self_data, data_size, ncclDtype,
ncclSum, comm, stream));
convert_data<T><<<108, 1024, 0, stream>>>(self_data, result, nccl_result,
my_result, data_size);
CUDACHECK(cudaStreamSynchronize(stream));
for (unsigned long j = 0; j < data_size; j++) {
auto diff = abs(nccl_result[j] - my_result[j]);
if (diff >= 4e-2) {
printf("Rank %d: Verification mismatch at %lld: %f != (my) %f, gt=%f\n",
myRank, j, nccl_result[j], my_result[j], ground_truth[j]);
break;
}
}
long double nccl_diffs = 0.0;
long double my_diffs = 0.0;
for (int j = 0; j < data_size; j++) {
nccl_diffs += abs(nccl_result[j] - ground_truth[j]);
my_diffs += abs(my_result[j] - ground_truth[j]);
}
if (myRank == 0)
std::cout << "average abs diffs: nccl: " << nccl_diffs / data_size
<< " me: " << my_diffs / data_size << std::endl;
} else {
for (int i = 0; i < 100; i++) {
fa.allreduce<T>(stream, self_data, result, data_size, threads,
block_limit);
CUDACHECK(cudaStreamSynchronize(stream));
NCCLCHECK(ncclAllReduce(self_data, self_data_copy, data_size, ncclDtype,
ncclSum, comm, stream));
convert_data<T><<<108, 1024, 0, stream>>>(
self_data_copy, result, nccl_result, my_result, data_size);
CUDACHECK(cudaStreamSynchronize(stream));
for (unsigned long j = 0; j < data_size; j++) {
auto diff = abs(nccl_result[j] - my_result[j]);
if (diff >= 4e-2) {
printf(
"Rank %d: Verification mismatch at %lld: %f != (my) %f, gt=%f\n",
myRank, j, nccl_result[j], my_result[j], ground_truth[j]);
break;
}
}
}
if (myRank == 0)
printf("Test passed: nGPUs:%d, sz (kb): %d, %d, %d\n", nRanks,
data_size * sizeof(T) / 1024, threads, block_limit);
// long double nccl_diffs = 0.0;
// long double my_diffs = 0.0;
// for (int j = 0; j < data_size; j++) {
// nccl_diffs += abs(nccl_result[j] - ground_truth[j]);
// my_diffs += abs(my_result[j] - ground_truth[j]);
// }
// if (myRank == 0)
// std::cout << "average abs diffs: nccl: " << nccl_diffs / data_size
// << " me: " << my_diffs / data_size << std::endl;
}
CUDACHECK(cudaFree(result));
CUDACHECK(cudaFree(self_data_copy));
CUDACHECK(cudaFree(rank_data));
CUDACHECK(cudaFree(buffer));
CUDACHECK(cudaFree(states));
CUDACHECK(cudaFreeHost(ground_truth));
CUDACHECK(cudaFreeHost(nccl_result));
CUDACHECK(cudaFreeHost(my_result));
CUDACHECK(cudaStreamDestroy(stream));
}
int main(int argc, char** argv) {
int nRanks, myRank;
MPICHECK(MPI_Init(&argc, &argv));
MPICHECK(MPI_Comm_rank(MPI_COMM_WORLD, &myRank));
MPICHECK(MPI_Comm_size(MPI_COMM_WORLD, &nRanks));
CUDACHECK(cudaSetDevice(myRank));
ncclUniqueId id;
ncclComm_t comm;
if (myRank == 0) ncclGetUniqueId(&id);
MPICHECK(MPI_Bcast(static_cast<void*>(&id), sizeof(id), MPI_BYTE, 0,
MPI_COMM_WORLD));
NCCLCHECK(ncclCommInitRank(&comm, nRanks, id, myRank));
bool performance_test = true;
cudaProfilerStart();
// Uncomment to scan through different block size configs.
// for (int threads : {256, 512, 1024}) {
// for (int block_limit = 16; block_limit < 112; block_limit += 4) {
// run<half>(myRank, nRanks, comm, threads, block_limit, 1024 * 1024,
// performance_test);
// }
// }
#ifdef USE_ROCM
const int block_limit = 16;
#else
const int block_limit = 36;
#endif
// Scan through different sizes to test performance.
for (int sz = 512; sz <= (8 << 20); sz *= 2) {
run<half>(myRank, nRanks, comm, 512, 36, sz + 8 * 47, performance_test);
}
cudaProfilerStop();
MPICHECK(MPI_Finalize());
return EXIT_SUCCESS;
}
csrc/fs_io.cpp
+69
View File
@@ -0,0 +1,69 @@
// SPDX-License-Identifier: Apache-2.0
// SPDX-FileCopyrightText: Copyright contributors to the vLLM project
#include <Python.h>
#include <unistd.h>
#include <vector>
extern "C" {
static void _batch_lookup(const std::vector<const char*>& paths,
std::vector<int>& exists_flags) {
for (size_t i = 0; i < paths.size(); i++) {
exists_flags[i] = (access(paths[i], F_OK) == 0) ? 1 : 0;
}
}
/// @brief Check file existence for a batch of paths.
/// @param paths list[str] – absolute paths to check.
/// @return list[bool] – True if the corresponding path exists, False otherwise.
/// @note Releases the GIL for the entire batch. File existence via access(2).
static PyObject* batch_lookup(PyObject* /*self*/, PyObject* args) {
PyObject* path_list;
if (!PyArg_ParseTuple(args, "O!", &PyList_Type, &path_list)) {
return nullptr;
}
const Py_ssize_t n = PyList_Size(path_list);
std::vector<const char*> paths(n);
for (Py_ssize_t i = 0; i < n; i++) {
paths[i] = PyUnicode_AsUTF8AndSize(PyList_GetItem(path_list, i), nullptr);
if (paths[i] == nullptr) {
return nullptr;
}
}
std::vector<int> exists_flags(n);
{
Py_BEGIN_ALLOW_THREADS _batch_lookup(paths, exists_flags);
Py_END_ALLOW_THREADS
}
PyObject* result = PyList_New(n);
if (result == nullptr) {
return nullptr;
}
for (Py_ssize_t i = 0; i < n; i++) {
PyList_SetItem(result, i, PyBool_FromLong(exists_flags[i]));
}
return result;
}
static PyMethodDef fs_io_C_methods[] = {
{"batch_lookup", batch_lookup, METH_VARARGS,
"batch_lookup(paths: list[str]) -> list[bool]\n"
"\n"
"Check file existence for a batch of paths."},
{nullptr, nullptr, 0, nullptr},
};
static struct PyModuleDef fs_io_C_module = {
PyModuleDef_HEAD_INIT, "fs_io_C", "Filesystem helpers for KV offload", -1,
fs_io_C_methods,
};
PyMODINIT_FUNC PyInit_fs_io_C(void) { return PyModule_Create(&fs_io_C_module); }
} // extern "C"
csrc/libtorch_stable/layernorm_kernels.cu
+24 -8
View File
@@ -20,20 +20,26 @@ __global__ void rms_norm_kernel(
const int64_t input_stride_d4, // input.stride(-4)
const int64_t input_shape_d2, // input.size(-2)
const int64_t input_shape_d3, // input.size(-3)
const scalar_t* __restrict__ weight, // [hidden_size], null if !HasWeight
const scalar_t* __restrict__ weight, // [hidden_size] or
// [num_groups, hidden_size];
// null if !HasWeight
const int64_t weight_stride, // 0 or weight.stride(0)
const float epsilon, const int num_tokens, const int hidden_size) {
__shared__ float s_variance;
float variance = 0.0f;
const scalar_t* input_row;
const scalar_t* weight_row;
if constexpr (NUM_DIMS == 2) {
// 2D for layernorm normal case [batch_size, hidden]
input_row = input + blockIdx.x * input_stride_d2;
weight_row = weight + blockIdx.x * weight_stride;
} else if constexpr (NUM_DIMS == 3) {
// 3D for q/k norm [batch_size, num_heads, head_size]
int batch_idx = blockIdx.x / input_shape_d2;
int head_idx = blockIdx.x % input_shape_d2;
input_row =
input + batch_idx * input_stride_d3 + head_idx * input_stride_d2;
weight_row = weight + batch_idx * weight_stride;
} else if constexpr (NUM_DIMS == 4) {
// 4D for transformers model_impl qk norm [batch, seq, head, head_dim]
int batch_idx = blockIdx.x / (input_shape_d3 * input_shape_d2);
@@ -42,6 +48,7 @@ __global__ void rms_norm_kernel(
int head_idx = remaining % input_shape_d2;
input_row = input + batch_idx * input_stride_d4 +
seq_idx * input_stride_d3 + head_idx * input_stride_d2;
weight_row = weight + batch_idx * weight_stride;
}
auto vec_op = [&variance](const vec_n_t<scalar_t, VEC_SIZE>& vec) {
@@ -69,7 +76,7 @@ __global__ void rms_norm_kernel(
scalar_t* out_row = out + blockIdx.x * hidden_size;
auto* v_in = reinterpret_cast<const vec_n_t<scalar_t, VEC_SIZE>*>(input_row);
auto* v_w = reinterpret_cast<const vec_n_t<scalar_t, VEC_SIZE>*>(weight);
auto* v_w = reinterpret_cast<const vec_n_t<scalar_t, VEC_SIZE>*>(weight_row);
auto* v_out = reinterpret_cast<vec_n_t<scalar_t, VEC_SIZE>*>(out_row);
for (int i = threadIdx.x; i < hidden_size / VEC_SIZE; i += blockDim.x) {
vec_n_t<scalar_t, VEC_SIZE> dst;
@@ -211,15 +218,24 @@ fused_add_rms_norm_kernel(
void rms_norm(torch::stable::Tensor& out, // [..., hidden_size]
torch::stable::Tensor& input, // [..., hidden_size]
std::optional<torch::stable::Tensor> weight, // [hidden_size]
double epsilon) {
std::optional<torch::stable::Tensor> weight, double epsilon) {
STD_TORCH_CHECK(out.is_contiguous());
if (input.stride(-1) != 1) {
input = torch::stable::contiguous(input);
}
STD_TORCH_CHECK(input.stride(-1) == 1);
int64_t weight_stride = 0;
if (weight.has_value()) {
STD_TORCH_CHECK(weight->is_contiguous());
if (weight->dim() == 1) {
STD_TORCH_CHECK(weight->size(0) == input.size(-1));
} else if (weight->dim() == 2) {
STD_TORCH_CHECK(weight->size(0) == input.size(0));
STD_TORCH_CHECK(weight->size(-1) == input.size(-1));
weight_stride = weight->stride(0);
} else {
STD_TORCH_CHECK(false, "rms_norm weight must be 1D or 2D");
}
}
int hidden_size = input.size(-1);
@@ -256,16 +272,16 @@ void rms_norm(torch::stable::Tensor& out, // [..., hidden_size]
out.mutable_data_ptr<scalar_t>(),
input.const_data_ptr<scalar_t>(), input_stride_d2,
input_stride_d3, input_stride_d4, input_shape_d2,
input_shape_d3, weight_ptr, epsilon, num_tokens,
hidden_size);
input_shape_d3, weight_ptr, weight_stride, epsilon,
num_tokens, hidden_size);
} else {
vllm::rms_norm_kernel<scalar_t, vec_size, tensor_rank, false>
<<<grid, block, 0, stream>>>(
out.mutable_data_ptr<scalar_t>(),
input.const_data_ptr<scalar_t>(), input_stride_d2,
input_stride_d3, input_stride_d4, input_shape_d2,
input_shape_d3, weight_ptr, epsilon, num_tokens,
hidden_size);
input_shape_d3, weight_ptr, /*weight_stride=*/0, epsilon,
num_tokens, hidden_size);
}
});
});
csrc/libtorch_stable/moe/moe_align_sum_kernels.cu
+159 -45
View File
@@ -11,6 +11,7 @@
#include "../../cuda_compat.h"
#include "libtorch_stable/core/math.hpp"
#include "libtorch_stable/dispatch_utils.h"
#include "libtorch_stable/quantization/vectorization.cuh"
#include "libtorch_stable/torch_utils.h"
#define CEILDIV(x, y) (((x) + (y) - 1) / (y))
@@ -349,19 +350,102 @@ __global__ void count_and_sort_expert_tokens_kernel(
max_num_tokens_padded, nullptr, 0, topk_num, has_expert_map);
}
// Reduce the topk expert outputs per token (summed in fp32). The output is
// dense [num_tokens, d]; the input is addressed by its strides so non-
// contiguous inputs work without a copy. A 16B-vectorized path is used when
// the hidden dim is contiguous (innermost stride 1) and aligned; otherwise a
// scalar kernel reads via arbitrary strides. topk is a compile-time constant
// for common values and runtime otherwise.
// Elements per 16-byte vector (8 for bf16/fp16, 4 for fp32).
template <typename scalar_t>
constexpr int MOE_SUM_VEC = 16 / sizeof(scalar_t);
template <typename scalar_t, int TOPK>
__global__ void moe_sum_kernel(
scalar_t* __restrict__ out, // [..., d]
const scalar_t* __restrict__ input, // [..., topk, d]
const int d) {
const int64_t token_idx = blockIdx.x;
for (int64_t idx = threadIdx.x; idx < d; idx += blockDim.x) {
scalar_t x = 0.0;
__global__ void moe_sum_vec_kernel(
scalar_t* __restrict__ out, // [num_tokens, d], contiguous
const scalar_t* __restrict__ input, // [num_tokens, topk, d], d contiguous
const int64_t num_tokens, const int d, const int64_t stride_token,
const int64_t stride_topk) {
using vec_t = vllm::vec_n_t<scalar_t, MOE_SUM_VEC<scalar_t>>; // 16-byte pack
constexpr int VEC = MOE_SUM_VEC<scalar_t>;
const int64_t n_vec = d / VEC;
const int64_t total = num_tokens * n_vec;
for (int64_t i = blockIdx.x * blockDim.x + threadIdx.x; i < total;
i += (int64_t)gridDim.x * blockDim.x) {
const int64_t token = i / n_vec;
const int64_t v = i % n_vec;
const scalar_t* in_tok = input + token * stride_token + v * VEC;
float acc[VEC];
#pragma unroll
for (int j = 0; j < VEC; ++j) acc[j] = 0.f;
#pragma unroll
for (int k = 0; k < TOPK; ++k) {
x += VLLM_LDG(&input[token_idx * TOPK * d + k * d + idx]);
vec_t packed = *reinterpret_cast<const vec_t*>(in_tok + k * stride_topk);
#pragma unroll
for (int j = 0; j < VEC; ++j) acc[j] += static_cast<float>(packed.val[j]);
}
out[token_idx * d + idx] = x;
vec_t outp;
#pragma unroll
for (int j = 0; j < VEC; ++j) outp.val[j] = static_cast<scalar_t>(acc[j]);
*reinterpret_cast<vec_t*>(out + token * d + v * VEC) = outp;
}
}
// Runtime-topk variant of the above.
template <typename scalar_t>
__global__ void moe_sum_vec_dynamic_kernel(
scalar_t* __restrict__ out, // [num_tokens, d], contiguous
const scalar_t* __restrict__ input, // [num_tokens, topk, d], d contiguous
const int64_t num_tokens, const int d, const int topk,
const int64_t stride_token, const int64_t stride_topk) {
using vec_t = vllm::vec_n_t<scalar_t, MOE_SUM_VEC<scalar_t>>;
constexpr int VEC = MOE_SUM_VEC<scalar_t>;
const int64_t n_vec = d / VEC;
const int64_t total = num_tokens * n_vec;
for (int64_t i = blockIdx.x * blockDim.x + threadIdx.x; i < total;
i += (int64_t)gridDim.x * blockDim.x) {
const int64_t token = i / n_vec;
const int64_t v = i % n_vec;
const scalar_t* in_tok = input + token * stride_token + v * VEC;
float acc[VEC];
#pragma unroll
for (int j = 0; j < VEC; ++j) acc[j] = 0.f;
for (int k = 0; k < topk; ++k) {
vec_t packed = *reinterpret_cast<const vec_t*>(in_tok + k * stride_topk);
#pragma unroll
for (int j = 0; j < VEC; ++j) acc[j] += static_cast<float>(packed.val[j]);
}
vec_t outp;
#pragma unroll
for (int j = 0; j < VEC; ++j) outp.val[j] = static_cast<scalar_t>(acc[j]);
*reinterpret_cast<vec_t*>(out + token * d + v * VEC) = outp;
}
}
// Stride-aware scalar fallback: handles unaligned/non-vectorizable hidden dims
// (including a non-contiguous hidden stride) via per-element strided reads.
template <typename scalar_t>
__global__ void moe_sum_scalar_kernel(
scalar_t* __restrict__ out, // [num_tokens, d], contiguous
const scalar_t* __restrict__ input, // [num_tokens, topk, d]
const int d, const int topk, const int64_t stride_token,
const int64_t stride_topk, const int64_t stride_hidden) {
const int64_t token_idx = blockIdx.x;
const scalar_t* in_tok = input + token_idx * stride_token;
for (int64_t idx = threadIdx.x; idx < d; idx += blockDim.x) {
float x = 0.f;
for (int k = 0; k < topk; ++k) {
x += static_cast<float>(
VLLM_LDG(&in_tok[k * stride_topk + idx * stride_hidden]));
}
out[token_idx * d + idx] = static_cast<scalar_t>(x);
}
}
@@ -626,52 +710,82 @@ void batched_moe_align_block_size(int64_t max_tokens_per_batch,
void moe_sum(torch::stable::Tensor& input, // [num_tokens, topk, hidden_size]
torch::stable::Tensor& output) // [num_tokens, hidden_size]
{
// Output is dense and written in place, so it must be contiguous. The input
// is read by its strides (no copy); only the hidden dim needs to be
// contiguous to take the vectorized path.
STD_TORCH_CHECK(output.is_contiguous(),
"moe_sum expects a contiguous output");
const int hidden_size = input.size(-1);
const auto num_tokens = output.numel() / hidden_size;
const int64_t num_tokens = output.numel() / hidden_size;
const int topk = input.size(1);
const int64_t stride_token = input.stride(0);
const int64_t stride_topk = input.stride(1);
const int64_t stride_hidden = input.stride(2);
dim3 grid(num_tokens);
dim3 block(std::min(hidden_size, 1024));
const torch::stable::accelerator::DeviceGuard device_guard(
output.get_device_index());
const cudaStream_t stream =
get_current_cuda_stream(output.get_device_index());
switch (topk) {
case 2:
VLLM_STABLE_DISPATCH_FLOATING_TYPES(
input.scalar_type(), "moe_sum_kernel", [&] {
vllm::moe::moe_sum_kernel<scalar_t, 2><<<grid, block, 0, stream>>>(
reinterpret_cast<scalar_t*>(output.mutable_data_ptr()),
reinterpret_cast<const scalar_t*>(input.const_data_ptr()),
hidden_size);
});
break;
#define LAUNCH_MOE_SUM_VEC(TOPK) \
vllm::moe::moe_sum_vec_kernel<scalar_t, TOPK> \
<<<grid, dim3(block), 0, stream>>>( \
out_ptr, in_ptr, num_tokens, hidden_size, stride_token, stride_topk)
case 3:
VLLM_STABLE_DISPATCH_FLOATING_TYPES(
input.scalar_type(), "moe_sum_kernel", [&] {
vllm::moe::moe_sum_kernel<scalar_t, 3><<<grid, block, 0, stream>>>(
reinterpret_cast<scalar_t*>(output.mutable_data_ptr()),
reinterpret_cast<const scalar_t*>(input.const_data_ptr()),
hidden_size);
});
break;
VLLM_STABLE_DISPATCH_FLOATING_TYPES(input.scalar_type(), "moe_sum", [&] {
constexpr int VEC = vllm::moe::MOE_SUM_VEC<scalar_t>;
constexpr int WIDTH = VEC * sizeof(scalar_t); // 16 bytes
auto* out_ptr = reinterpret_cast<scalar_t*>(output.mutable_data_ptr());
auto* in_ptr = reinterpret_cast<const scalar_t*>(input.const_data_ptr());
case 4:
VLLM_STABLE_DISPATCH_FLOATING_TYPES(
input.scalar_type(), "moe_sum_kernel", [&] {
vllm::moe::moe_sum_kernel<scalar_t, 4><<<grid, block, 0, stream>>>(
reinterpret_cast<scalar_t*>(output.mutable_data_ptr()),
reinterpret_cast<const scalar_t*>(input.const_data_ptr()),
hidden_size);
});
break;
default:
torch::stable::sum_out(output, input, std::array<int64_t, 1>{1});
break;
}
// Vectorize along hidden only when it is contiguous (innermost stride 1),
// a whole number of vectors, and every row offset stays 16B-aligned.
const bool can_vec = (stride_hidden == 1) && (hidden_size % VEC == 0) &&
(stride_token % VEC == 0) &&
(stride_topk % VEC == 0) &&
(reinterpret_cast<uintptr_t>(in_ptr) % WIDTH == 0) &&
(reinterpret_cast<uintptr_t>(out_ptr) % WIDTH == 0);
if (can_vec) {
const int64_t n_vec = hidden_size / VEC;
const int64_t total = num_tokens * n_vec;
const int block = 256;
const dim3 grid(std::min<int64_t>((total + block - 1) / block, 65535));
switch (topk) {
case 1:
LAUNCH_MOE_SUM_VEC(1);
break;
case 2:
LAUNCH_MOE_SUM_VEC(2);
break;
case 4:
LAUNCH_MOE_SUM_VEC(4);
break;
case 6:
LAUNCH_MOE_SUM_VEC(6);
break;
case 8:
LAUNCH_MOE_SUM_VEC(8);
break;
case 9:
LAUNCH_MOE_SUM_VEC(9);
break;
default:
vllm::moe::moe_sum_vec_dynamic_kernel<scalar_t>
<<<grid, dim3(block), 0, stream>>>(out_ptr, in_ptr, num_tokens,
hidden_size, topk,
stride_token, stride_topk);
break;
}
} else {
dim3 grid(num_tokens);
dim3 block(std::min(hidden_size, 1024));
vllm::moe::moe_sum_scalar_kernel<scalar_t><<<grid, block, 0, stream>>>(
out_ptr, in_ptr, hidden_size, topk, stride_token, stride_topk,
stride_hidden);
}
});
#undef LAUNCH_MOE_SUM_VEC
}
void moe_lora_align_block_size(
csrc/ops.h
-6
View File
@@ -41,12 +41,6 @@ void gelu_fast(torch::Tensor& out, torch::Tensor& input);
void gelu_quick(torch::Tensor& out, torch::Tensor& input);
void cutlass_mla_decode(torch::Tensor const& out, torch::Tensor const& q_nope,
torch::Tensor const& q_pe,
torch::Tensor const& kv_c_and_k_pe_cache,
torch::Tensor const& seq_lens,
torch::Tensor const& page_table, double scale);
void static_scaled_int8_quant(torch::Tensor& out, torch::Tensor const& input,
torch::Tensor const& scale,
std::optional<torch::Tensor> const& azp);
docker/Dockerfile
+4 -3
View File
@@ -311,9 +311,10 @@ ENV CARGO_BUILD_JOBS=4
# Build the release artifacts. Cache cargo registry/git, but not target/,
# because stale target metadata can outlive source updates across BuildKit
# cache reuse.
RUN --mount=type=cache,target=/root/.cargo/registry \
--mount=type=cache,target=/root/.cargo/git \
RUN --mount=type=cache,target=/root/.cargo/registry,sharing=locked \
--mount=type=cache,target=/root/.cargo/git,sharing=locked \
bash build_rust.sh
#################### RUST BUILD IMAGE ####################
#################### CSRC BUILD IMAGE ####################
@@ -792,7 +793,7 @@ RUN --mount=type=cache,target=/opt/uv/cache \
# Install FlashInfer JIT cache (requires CUDA-version-specific index URL)
# https://docs.flashinfer.ai/installation.html
# From versions.json: .flashinfer.version
ARG FLASHINFER_VERSION=0.6.12
ARG FLASHINFER_VERSION=0.6.13
RUN --mount=type=cache,target=/opt/uv/cache \
uv pip install --system flashinfer-jit-cache==${FLASHINFER_VERSION} \
--index-url https://flashinfer.ai/whl/cu$(echo $CUDA_VERSION | cut -d. -f1,2 | tr -d '.')
docker/Dockerfile.cpu
+1
View File
@@ -194,6 +194,7 @@ FROM base AS vllm-test-deps
WORKDIR /vllm-workspace
# Copy test requirements
COPY requirements/common.txt requirements/common.txt
COPY requirements/test/cuda.in requirements/test/cpu.in
RUN \
docker/Dockerfile.nightly_torch
+5 -4
View File
@@ -123,9 +123,10 @@ COPY build_rust.sh build_rust.sh
# (rustc spawns enough concurrent processes to hit RLIMIT_NOFILE otherwise).
ENV CARGO_BUILD_JOBS=4
RUN --mount=type=cache,target=/root/.cargo/registry \
--mount=type=cache,target=/root/.cargo/git \
RUN --mount=type=cache,target=/root/.cargo/registry,sharing=locked \
--mount=type=cache,target=/root/.cargo/git,sharing=locked \
bash build_rust.sh
#################### RUST BUILD IMAGE ####################
#################### WHEEL BUILD IMAGE ####################
@@ -256,13 +257,13 @@ RUN pip install setuptools==75.6.0 packaging==23.2 ninja==1.11.1.3 build==1.2.2.
# build flashinfer for torch nightly from source around 10 mins
# release version: v0.6.12
# release version: v0.6.13
# todo(elainewy): cache flashinfer build result for faster build
ENV CCACHE_DIR=/root/.cache/ccache
RUN --mount=type=cache,target=/root/.cache/ccache \
--mount=type=cache,target=/root/.cache/uv \
echo "git clone flashinfer..." \
&& git clone --depth 1 --branch v0.6.12 --recursive https://github.com/flashinfer-ai/flashinfer.git \
&& git clone --depth 1 --branch v0.6.13 --recursive https://github.com/flashinfer-ai/flashinfer.git \
&& cd flashinfer \
&& git submodule update --init --recursive \
&& echo "finish git clone flashinfer..." \
docker/Dockerfile.rocm
+6
View File
@@ -575,6 +575,9 @@ RUN --mount=type=cache,target=/root/.cache/uv \
ENV HF_XET_HIGH_PERFORMANCE=1
ENV HF_HUB_DOWNLOAD_TIMEOUT=60
# Keep torch.cuda.is_available() fork-safe (see vllm/env_override.py).
ENV PYTORCH_NVML_BASED_CUDA_CHECK=1
# Pre-install vLLM test dependencies.
COPY requirements/test/rocm.txt /tmp/rocm-test-reqs.txt
RUN --mount=type=cache,target=/root/.cache/uv \
@@ -695,6 +698,9 @@ ENV SAFETENSORS_FAST_GPU=1
# Performance environment variable.
ENV HIP_FORCE_DEV_KERNARG=1
# Keep torch.cuda.is_available() fork-safe (see vllm/env_override.py).
ENV PYTORCH_NVML_BASED_CUDA_CHECK=1
# Workaround for ROCm profiler limits
RUN echo "ROCTRACER_MAX_EVENTS=10000000" > ${COMMON_WORKDIR}/libkineto.conf
ENV KINETO_CONFIG="${COMMON_WORKDIR}/libkineto.conf"
docker/Dockerfile.xpu
+2 -2
View File
@@ -28,8 +28,8 @@ COPY build_rust.sh build_rust.sh
# (rustc spawns enough concurrent processes to hit RLIMIT_NOFILE otherwise).
ENV CARGO_BUILD_JOBS=4
RUN --mount=type=cache,target=/root/.cargo/registry \
--mount=type=cache,target=/root/.cargo/git \
RUN --mount=type=cache,target=/root/.cargo/registry,sharing=locked \
--mount=type=cache,target=/root/.cargo/git,sharing=locked \
bash build_rust.sh
FROM ubuntu:24.04 AS vllm-base
docker/ci-rocm.hcl
+15 -6
View File
@@ -302,14 +302,22 @@ group "test-rocm-ci-with-wheel" {
}
# Image tags for the ci_base build. ci-bake-rocm.sh rewrites CI_BASE_IMAGE_TAG
# to the primary tag for this build. Non-nightly builds use a commit-scoped tag
# and also publish a content tag for reuse. NIGHTLY=1 builds on the stable branch
# can additionally set CI_BASE_IMAGE_TAG_STABLE to refresh rocm/vllm-dev:ci_base.
# to the primary tag for this build. Builds always publish a content-scoped tag
# when the ci_base content hash is available. Builds with BUILDKITE_COMMIT also
# publish a commit-scoped tag, either as the primary tag or an additional alias.
# NIGHTLY=1 builds on the stable branch can additionally set
# CI_BASE_IMAGE_TAG_STABLE to refresh rocm/vllm-dev:ci_base.
variable "CI_BASE_IMAGE_TAG" {
default = "rocm/vllm-dev:ci_base"
}
variable "CI_BASE_IMAGE_TAG_CONTENT" {
# Supplemental tags only. ci-bake-rocm.sh leaves these empty when the same ref
# is already the primary CI_BASE_IMAGE_TAG.
variable "CI_BASE_IMAGE_TAG_COMMIT_EXTRA" {
default = ""
}
variable "CI_BASE_IMAGE_TAG_CONTENT_EXTRA" {
default = ""
}
@@ -357,7 +365,8 @@ target "ci-base-rocm-ci" {
cache-from = concat(
compact([
CI_BASE_IMAGE_TAG != "" ? "type=registry,ref=${CI_BASE_IMAGE_TAG}" : "",
CI_BASE_IMAGE_TAG_CONTENT != "" ? "type=registry,ref=${CI_BASE_IMAGE_TAG_CONTENT}" : "",
CI_BASE_IMAGE_TAG_COMMIT_EXTRA != "" ? "type=registry,ref=${CI_BASE_IMAGE_TAG_COMMIT_EXTRA}" : "",
CI_BASE_IMAGE_TAG_CONTENT_EXTRA != "" ? "type=registry,ref=${CI_BASE_IMAGE_TAG_CONTENT_EXTRA}" : "",
CI_BASE_IMAGE_TAG_STABLE != "" ? "type=registry,ref=${CI_BASE_IMAGE_TAG_STABLE}" : "",
]),
# Import upstream dependency caches so RIXL/ROCShmem/DeepEP stages
@@ -365,7 +374,7 @@ target "ci-base-rocm-ci" {
get_cache_from_rocm_deps(),
)
cache-to = ["type=inline"]
tags = compact([CI_BASE_IMAGE_TAG, CI_BASE_IMAGE_TAG_CONTENT, CI_BASE_IMAGE_TAG_STABLE])
tags = compact([CI_BASE_IMAGE_TAG, CI_BASE_IMAGE_TAG_COMMIT_EXTRA, CI_BASE_IMAGE_TAG_CONTENT_EXTRA, CI_BASE_IMAGE_TAG_STABLE])
output = ["type=registry"]
}
docker/versions.json
+1 -1
View File
@@ -68,7 +68,7 @@
"default": "true"
},
"FLASHINFER_VERSION": {
"default": "0.6.12"
"default": "0.6.13"
},
"GDRCOPY_CUDA_VERSION": {
"default": "12.8"
docs/cli/README.md
+6
View File
@@ -80,6 +80,9 @@ vllm chat --url http://{vllm-serve-host}:{vllm-serve-port}/v1
# Quick chat with a single prompt
vllm chat --quick "hi"
# Print TTFT and throughput statistics after each response
vllm chat --stats
```
See [vllm chat](./chat.md) for the full reference of all available arguments.
@@ -97,6 +100,9 @@ vllm complete --url http://{vllm-serve-host}:{vllm-serve-port}/v1
# Quick complete with a single prompt
vllm complete --quick "The future of AI is"
# Print TTFT and throughput statistics after each response
vllm complete --stats
```
See [vllm complete](./complete.md) for the full reference of all available arguments.
docs/configuration/optimization.md
+1 -1
View File
@@ -297,7 +297,7 @@ The `fastokens` Python package (>= 0.2.0) must be installed; if it isn't,
vLLM raises a clear `ImportError` at tokenizer load. The override applies to
any `--tokenizer-mode` that ends up loading an HF fast tokenizer (`hf`,
`deepseek_v32`, `deepseek_v4`, …). Models that don't use the HF
fast tokenizer (`mistral`, `grok2`, `kimi_audio`) ignore the flag.
fast tokenizer (`mistral`, `kimi_audio`) ignore the flag.
Tokenizer-bound workloads — long shared prefixes, bursty short prompts,
batch detokenization — see the largest wins. If your bottleneck is GPU
docs/design/arch_overview.md
+1 -1
View File
@@ -178,7 +178,7 @@ incoming requests. The `AsyncLLMEngine` is designed for online serving, where it
can handle multiple concurrent requests and stream outputs to clients.
The OpenAI-compatible API server uses the `AsyncLLMEngine`. There is also a demo
API server that serves as a simpler example in [vllm/entrypoints/api_server.py](../../vllm/entrypoints/api_server.py).
API server that serves as a simpler example in [examples/applications/api_server/server.py](../../examples/applications/api_server/server.py).
The code for `AsyncLLMEngine` can be found in [vllm/engine/async_llm_engine.py](../../vllm/engine/async_llm_engine.py).
docs/design/attention_backends.md
+2 -2
View File
@@ -168,7 +168,7 @@ Priority is **1 = highest** (tried first).
| `FLASH_ATTN_DIFFKV` | | fp16, bf16 | `auto` | Any | Any | ❌ | ❌ | ❌ | ✅ | Decoder | Any |
| `FLEX_ATTENTION` | | fp16, bf16, fp32 | `auto`, `float16`, `bfloat16` | %16 | Any | ❌ | ✅ | ✅ | ❌ | Decoder, Encoder Only | Any |
| `ROCM_AITER_FA` | | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | 16, 32 | 64, 128, 256 | ✅ | ✅ | ❌ | ❌ | Decoder | N/A |
| `ROCM_AITER_UNIFIED_ATTN` | | fp16, bf16 | `auto` | %16 | Any | ✅ | ❌ | ✅ | ❌ | All | N/A |
| `ROCM_AITER_UNIFIED_ATTN` | | bf16 | `auto`, `bfloat16`, `fp8`, `fp8_e4m3` | %16 | Any | ✅ | ❌ | ✅ | ❌ | All | N/A |
| `ROCM_ATTN` | | fp16, bf16, fp32 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2` | %16 | 32, 64, 80, 96, 128, 160, 192, 224, 256 | ❌ | ✅ | ✅ | ❌ | Decoder, Encoder, Encoder Only | N/A |
| `TRITON_ATTN` | | fp16, bf16, fp32 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3`, `fp8_e5m2`, `int4_per_token_head`, `int8_per_token_head`, `fp8_per_token_head` | %16 | Any | ✅ | ✅ | ✅ | ❌ | All | Any |
| `TRITON_ATTN_DIFFKV` | | fp16, bf16 | `auto`, `bfloat16` | Any | Any | ❌ | ❌ | ❌ | ❌ | Decoder | Any |
@@ -220,7 +220,7 @@ MLA decode backends are selected using the standard
| Backend | Dtypes | KV Dtypes | Block Sizes | Head Sizes | Sink | Non-Causal | Sparse | MM Prefix | DCP | Attention Types | Compute Cap. |
| ------- | ------ | --------- | ----------- | ---------- | ---- | ---------- | ------ | --------- | --- | --------------- | ------------ |
| `CUTLASS_MLA` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | 128 | Any | ❌ | ❌ | ❌ | ❌ | ✅ | Decoder | 10.x |
| `FLASHINFER_MLA` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | 32, 64 | Any | ❌ | ❌ | ❌ | ❌ | ❌ | Decoder | 10.x |
| `FLASHINFER_MLA` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | 32, 64 | Any | ❌ | ❌ | ❌ | ❌ | ✅ | Decoder | 10.x |
| `FLASHINFER_MLA_SPARSE` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | 32, 64 | Any | ❌ | ❌ | ❌ | ❌ | ❌ | Decoder | 10.x |
| `FLASHINFER_MLA_SPARSE_SM120` | bf16 | `auto`, `fp8`, `fp8_e4m3`, `fp8_ds_mla` | 64, 256 | Any | ❌ | ❌ | ❌ | ❌ | ❌ | Decoder | 12.x |
| `FLASHMLA` | fp16, bf16 | `auto`, `float16`, `bfloat16`, `fp8`, `fp8_e4m3` | 64 | Any | ❌ | ❌ | ❌ | ❌ | ✅ | Decoder | 9.x-10.x |
docs/design/cuda_graphs_multimodal.md
+2 -1
View File
@@ -127,6 +127,7 @@ Models opt-in to encoder CUDA Graphs by implementing the [SupportsEncoderCudaGra
| Architecture | Models | CG for Image | CG for Video | Dual-Path Graph |
| ------------ | ------ | ------------ | ------------ | --------------- |
| `DeepseekOCRForCausalLM` | `DeepSeek-OCR` | ✅︎ | ❌︎ | ✅︎ |
| `Gemma3ForConditionalGeneration` | `Gemma3` | ✅︎ | ❌︎ | ❌︎ |
| `Glm4vForConditionalGeneration` | `GLM-4.1V, GLM-4.6V-Flash` | ✅︎ | ✅︎ | ❌︎ |
| `InternVLChatModel` | `InternVL3.5`, `InternVL3`, `InternVL2.5`, `InternVL2` | ✅︎ | ✅︎ | ❌︎ |
| `KimiVLForConditionalGeneration` | `Kimi-VL` | ✅︎ | ❌︎ | ❌︎ |
@@ -136,7 +137,7 @@ Models opt-in to encoder CUDA Graphs by implementing the [SupportsEncoderCudaGra
| `Qwen3VLForConditionalGeneration` | `Qwen3-VL` | ✅︎ | ✅︎ | ❌︎ |
| `Qwen3_5ForConditionalGeneration` | `Qwen3.5`, `Qwen3.6` | ✅︎ | ✅︎ | ❌︎ |
| `Qwen3_5MoeForConditionalGeneration` | `Qwen3.5-MoE`, `Qwen3.6-MoE` | ✅︎ | ✅︎ | ❌︎ |
| `Step3VLForConditionalGeneration` | `Step3-VL` | ✅︎ | ❌︎ | ❌︎ |
| `Step3VLForConditionalGeneration` | `Step3-VL` | ✅︎ | ❌︎ | ✅︎ |
!!! note
Encoder CUDA Graphs have currently been tested with `--mm-encoder-attn-backend=FLASH_ATTN` and `--mm-encoder-attn-backend=FLASHINFER` on Blackwell GPUs.
docs/design/moe_kernel_features.md
+1
View File
@@ -89,6 +89,7 @@ To be used with a particular `FusedMoEPrepareAndFinalizeModular` subclass, MoE k
| gpt oss triton | standard | N/A | N/A | <sup>5</sup> | Y | Y | [`triton_kernel_fused_experts`][vllm.model_executor.layers.fused_moe.experts.gpt_oss_triton_kernels_moe.triton_kernel_fused_experts],</br>[`OAITritonExperts`][vllm.model_executor.layers.fused_moe.experts.gpt_oss_triton_kernels_moe.OAITritonExperts] |
| marlin | standard,</br>batched | <sup>3</sup> / N/A | <sup>3</sup> / N/A | silu,</br>swigluoai | Y | Y | [`fused_marlin_moe`][vllm.model_executor.layers.fused_moe.experts.marlin_moe.fused_marlin_moe],</br>[`MarlinExperts`][vllm.model_executor.layers.fused_moe.experts.marlin_moe.MarlinExperts],</br>[`BatchedMarlinExperts`][vllm.model_executor.layers.fused_moe.experts.marlin_moe.BatchedMarlinExperts] |
| trtllm | standard | mxfp4,</br>nvfp4 | G(16),G(32) | <sup>5</sup> | N | Y | [`TrtLlmMxfp4ExpertsMonolithic`][vllm.model_executor.layers.fused_moe.experts.trtllm_mxfp4_moe.TrtLlmMxfp4ExpertsMonolithic],</br>[`TrtLlmMxfp4ExpertsModular`][vllm.model_executor.layers.fused_moe.experts.trtllm_mxfp4_moe.TrtLlmMxfp4ExpertsModular],</br>[`TrtLlmNvFp4ExpertsMonolithic`][vllm.model_executor.layers.fused_moe.experts.trtllm_nvfp4_moe.TrtLlmNvFp4ExpertsMonolithic],</br>[`TrtLlmNvfp4ExpertsModular`][vllm.model_executor.layers.fused_moe.experts.trtllm_nvfp4_moe.TrtLlmNvFp4ExpertsModular] |
| hpc | standard | fp8 | G(128),T | silu | Y | Y | [`HPCExperts`][vllm.model_executor.layers.fused_moe.hpc_moe.HPCExperts] |
| rocm aiter moe | standard | mxfp4,</br>fp8 | G(32),G(128),A,T | silu, gelu,</br>swigluoai | Y | N | `rocm_aiter_fused_experts`,</br>`AiterExperts` |
| cpu_fused_moe | standard | N/A | N/A | silu | N | N | [`CPUFusedMOE`][vllm.model_executor.layers.fused_moe.cpu_fused_moe.CPUFusedMOE] |
| naive batched<sup>4</sup> | batched | int8,</br>fp8 | G,A,T | silu, gelu | <sup>6</sup> | Y | [`NaiveBatchedExperts`][vllm.model_executor.layers.fused_moe.experts.fused_batched_moe.NaiveBatchedExperts] |
docs/examples/README.md
+3 -2
View File
@@ -9,9 +9,10 @@ vLLM's examples are organized into the following categories:
- **[`features/`](../../examples/features)** – Demonstrations of individual vLLM features: automatic prefix caching, speculative decoding, LoRA, structured outputs, prompt embedding, pause/resume, batch invariance, KV events, data parallelism, and more.
- **[`reasoning/`](../../examples/reasoning)** – Examples for reasoning with vLLM.
- **[`tool_calling/`](../../examples/tool_calling)** – Examples for function/tool calling with vLLM.
- **[`applications/`](../../examples/applications)** – Application examples such as chatbots and RAG (Retrieval-Augmented Generation).
- **[`applications/`](../../examples/applications)** – Application examples such as simpler api server, chatbots and RAG (Retrieval-Augmented Generation).
- **[`rl/`](../../examples/rl)** – Reinforcement learning examples.
- **[`deployment/`](../../examples/deployment)** – Examples for deploying vLLM in production.
- **[`ray_serving/`](../../examples/ray_serving)** – Scalable serving using Ray.
- **[`disaggregated/`](../../examples/disaggregated)** – Examples for disaggregated serving (separate prefill and decode), including various kv cache connectors (LMCache, Mooncake, FlexKV, P2P NCCL) and failure recovery.
- **[`disaggregated/`](../../examples/disaggregated)** – Examples for Disaggregated P/D (Prefill/Decoding) inference, including various kv cache connectors (LMCache, Mooncake, FlexKV, P2P NCCL) and failure recovery.
- **[`scale_out/`](../../examples/scale_out)** – Examples for Token In <> Token Out API Server.
- **[`observability/`](../../examples/observability)** – Metrics, logging, tracing (OpenTelemetry), and dashboards (Grafana, Perses).
docs/getting_started/installation/cpu.apple.inc.md
+4
View File
@@ -15,6 +15,10 @@ Currently the CPU implementation for macOS supports FP32 and FP16 datatypes.
- SDK: `XCode 15.4` or later with Command Line Tools
- Compiler: `Apple Clang >= 15.0.0`
!!! note
The macOS CPU build is smoke-tested in CI on the latest GA Apple Silicon
runner; other macOS or Apple Clang versions are best-effort.
--8<-- [end:requirements]
--8<-- [start:set-up-using-python]
docs/models/pooling_models/token_classify.md
+1
View File
@@ -45,6 +45,7 @@ The BAAI/bge-m3 model leverages token classification for sparse retrieval. For m
| ------------ | ------ | ----------------- | --------------------------- | --------------------------------------- |
| `BertForTokenClassification` | bert-based | `boltuix/NeuroBERT-NER` (see note), etc. | | |
| `ModernBertForTokenClassification` | ModernBERT-based | `disham993/electrical-ner-ModernBERT-base` | | |
| `OpenAIPrivacyFilterForTokenClassification` | gpt-oss-based encoder | `openai/privacy-filter` | | |
| `Qwen3ForTokenClassification`<sup>C</sup> | Qwen3-based | `bd2lcco/Qwen3-0.6B-finetuned` | | |
| `*Model`<sup>C</sup>, `*ForCausalLM`<sup>C</sup>, etc. | Generative models | N/A | \* | \* |
docs/models/supported_models.md
+1 -5
View File
@@ -414,8 +414,6 @@ th {
| `GraniteMoeHybridForCausalLM` | Granite 4.0 MoE Hybrid | `ibm-granite/granite-4.0-tiny-preview`, etc. | ✅︎ | ✅︎ |
| `GraniteMoeSharedForCausalLM` | Granite MoE Shared | `ibm-research/moe-7b-1b-active-shared-experts` (test model) | ✅︎ | ✅︎ |
| `GritLM` | GritLM | `parasail-ai/GritLM-7B-vllm`. | ✅︎ | ✅︎ |
| `Grok1ModelForCausalLM` | Grok1 | `hpcai-tech/grok-1`. | ✅︎ | ✅︎ |
| `Grok1ForCausalLM` | Grok2 | `xai-org/grok-2` | ✅︎ | ✅︎ |
| `HrmTextForCausalLM` | HRM-Text | `sapientinc/HRM-Text-1B`, etc. | | |
| `HunYuanDenseV1ForCausalLM` | Hunyuan Dense | `tencent/Hunyuan-7B-Instruct` | ✅︎ | ✅︎ |
| `HunYuanMoEV1ForCausalLM` | Hunyuan-A13B | `tencent/Hunyuan-A13B-Instruct`, `tencent/Hunyuan-A13B-Pretrain`, `tencent/Hunyuan-A13B-Instruct-FP8`, etc. | ✅︎ | ✅︎ |
@@ -488,9 +486,6 @@ th {
| `TeleFLMForCausalLM` | TeleFLM | `CofeAI/FLM-2-52B-Instruct-2407`, `CofeAI/Tele-FLM`, etc. | ✅︎ | ✅︎ |
| `Zamba2ForCausalLM` | Zamba2 | `Zyphra/Zamba2-7B-instruct`, `Zyphra/Zamba2-2.7B-instruct`, `Zyphra/Zamba2-1.2B-instruct`, etc. | | |
!!! note
Grok2 requires `tokenizer.tok.json` with `tiktoken` installed. You can optionally override MoE router renormalization with `moe_router_renormalize`.
Some models are supported only via the [Transformers modeling backend](#transformers). The purpose of the table below is to acknowledge models which we officially support in this way. The logs will say that the Transformers modeling backend is being used, and you will see no warning that this is fallback behaviour. This means that, if you have issues with any of the models listed below, please [make an issue](https://github.com/vllm-project/vllm/issues/new/choose) and we'll do our best to fix it!
| Architecture | Models | Example HF Models | [LoRA](../features/lora.md) | [PP](../serving/parallelism_scaling.md) |
@@ -634,6 +629,7 @@ These models primarily accept the [`LLM.generate`](./generative_models.md#llmgen
| `TarsierForConditionalGeneration` | Tarsier | T + I<sup>E+</sup> | `omni-search/Tarsier-7b`, `omni-search/Tarsier-34b` | | ✅︎ |
| `Tarsier2ForConditionalGeneration`<sup>^</sup> | Tarsier2 | T + I<sup>E+</sup> + V<sup>E+</sup> | `omni-research/Tarsier2-Recap-7b`, `omni-research/Tarsier2-7b-0115` | | ✅︎ |
| `UltravoxModel` | Ultravox | T + A<sup>E+</sup> | `fixie-ai/ultravox-v0_5-llama-3_2-1b` | ✅︎ | ✅︎ |
| `UnlimitedOCRForCausalLM` | Unlimited-OCR | T + I<sup>+</sup> | `baidu/Unlimited-OCR`, etc. | ✅︎ | ✅︎ |
Some models are supported only via the [Transformers modeling backend](#transformers). The purpose of the table below is to acknowledge models which we officially support in this way. The logs will say that the Transformers modeling backend is being used, and you will see no warning that this is fallback behaviour. This means that, if you have issues with any of the models listed below, please [make an issue](https://github.com/vllm-project/vllm/issues/new/choose) and we'll do our best to fix it!
docs/serving/online_serving/README.md
+2 -2
View File
@@ -119,9 +119,9 @@ For further details on profiling vLLM, please refer to [this page](../../contrib
- `/ping` - SageMaker health check
- `/invocations` - SageMaker-compatible endpoint (routes to the same inference functions as `/v1` endpoints)
## Disaggregated Everything
## Scale-Out APIs
### Tokens IN <> Tokens OUT
### Tokens IN <> Tokens OUT APIs
- `/inference/v1/generate` - Generate completions
- `/abort_requests` - Abort in-flight requests (only when `--tokens-only` is also set)
docs/usage/v1_guide.md
+1 -1
View File
@@ -125,7 +125,7 @@ We are working on enabling prefix caching and chunked prefill for more categorie
Models using selective state-space mechanisms instead of standard transformer attention are supported.
Models that use Mamba-2 and Mamba-1 layers (e.g., `Mamba2ForCausalLM`, `MambaForCausalLM`, `FalconMambaForCausalLM`) are supported.
Hybrid models that combine Mamba-2 and Mamba-1 layers with standard attention layers are also supported (e.g., `BambaForCausalLM`,
Hybrid models that combine Mamba-2 and Mamba-1 layers with standard attention layers are also supported (e.g.,
`Zamba2ForCausalLM`, `NemotronHForCausalLM`, `FalconH1ForCausalLM` and `GraniteMoeHybridForCausalLM`, `JambaForCausalLM`, `Plamo2ForCausalLM`).
Hybrid models with mechanisms different to Mamba are also supported (e.g, `Lfm2ForCausalLM`).
tests/entrypoints/serve/disagg/__init__.py → examples/__init__.py
View File
examples/applications/chatbot/api_client.py → examples/applications/api_server/client.py
+2 -2
View File
@@ -1,8 +1,8 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
"""Example Python client for `vllm.entrypoints.api_server`
"""Example Python client for `examples/applications/api_server/server.py`
Start the demo server:
python -m vllm.entrypoints.api_server --model <model_name>
python examples/applications/api_server/server.py --model <model_name>
NOTE: The API server is used only for demonstration and simple performance
benchmarks. It is not intended for production use.
vllm/entrypoints/api_server.py → examples/applications/api_server/server.py
+1 -1
View File
@@ -31,7 +31,7 @@ from vllm.utils.argparse_utils import FlexibleArgumentParser
from vllm.utils.system_utils import set_ulimit
from vllm.version import __version__ as VLLM_VERSION
logger = init_logger("vllm.entrypoints.api_server")
logger = init_logger("api_server")
app = FastAPI()
engine = None
examples/applications/chatbot/gradio_webserver.py
+1 -1
View File
@@ -2,7 +2,7 @@
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
"""Example for starting a Gradio Webserver
Start vLLM API server:
python -m vllm.entrypoints.api_server \
python examples/applications/api_server/server.py \
--model meta-llama/Llama-2-7b-chat-hf
Start Webserver:
examples/generate/multimodal/audio_language_offline.py
+2 -5
View File
@@ -463,16 +463,15 @@ def run_ultravox(question: str, audio_count: int) -> ModelRequestData:
# Voxtral
# Make sure to install mistral-common[audio].
def run_voxtral(question: str, audio_count: int) -> ModelRequestData:
from mistral_common.audio import Audio
from mistral_common.protocol.instruct.chunk import (
AudioChunk,
RawAudio,
TextChunk,
)
from mistral_common.protocol.instruct.messages import (
UserMessage,
)
from mistral_common.protocol.instruct.request import ChatCompletionRequest
from mistral_common.tokens.tokenizers.audio import Audio
from mistral_common.tokens.tokenizers.mistral import MistralTokenizer
model_name = "mistralai/Voxtral-Mini-3B-2507"
@@ -495,9 +494,7 @@ def run_voxtral(question: str, audio_count: int) -> ModelRequestData:
Audio.from_file(str(audio_assets[i].get_local_path()), strict=False)
for i in range(audio_count)
]
audio_chunks = [
AudioChunk(input_audio=RawAudio.from_audio(audio)) for audio in audios
]
audio_chunks = [AudioChunk.from_audio(audio) for audio in audios]
messages = [UserMessage(content=[*audio_chunks, text_chunk])]
examples/rl/rlhf_async_new_apis.py
+3 -4
View File
@@ -190,12 +190,11 @@ class TrainModel:
# Build platform-specific env vars for Ray
ray_env_vars = {
# Prevent Ray from setting CUDA_VISIBLE_DEVICES
"RAY_EXPERIMENTAL_NOSET_CUDA_ENV_VAR": "1",
}
ray_env_vars = {}
if current_platform.is_rocm():
# Workaround for RCCL bug. See https://github.com/ROCm/rocm-systems/issues/5756
ray_env_vars["RAY_EXPERIMENTAL_NOSET_HIP_VISIBLE_DEVICES"] = "1"
# For ROCm, BATCH_INVARIANT vllm is not supported
ray_env_vars["VLLM_ROCM_USE_SKINNY_GEMM"] = "0"
else:
examples/rl/rlhf_nccl.py
+16 -1
View File
@@ -29,6 +29,7 @@ causes unexpected behavior.
import os
import ray
import torch
from ray.util.placement_group import placement_group
from ray.util.scheduling_strategies import PlacementGroupSchedulingStrategy
from transformers import AutoModelForCausalLM
@@ -39,12 +40,24 @@ from vllm.distributed.weight_transfer.nccl_engine import (
NCCLTrainerSendWeightsArgs,
NCCLWeightTransferEngine,
)
from vllm.platforms import current_platform
from vllm.utils.network_utils import get_ip, get_open_port
MODEL_NAME = "facebook/opt-125m"
# MODEL_NAME = "inference-optimization/Qwen3-0.6B-W4A16-G128"
def get_assigned_gpu():
"""This is a temporary workaround for a runtime bug in RCCL on ROCm."""
if not current_platform.is_rocm():
return 0
assigned_gpu = int(ray.get_gpu_ids()[0])
os.environ.pop("CUDA_VISIBLE_DEVICES", None)
os.environ.pop("HIP_VISIBLE_DEVICES", None)
torch.accelerator.set_device_idx(assigned_gpu)
return assigned_gpu
class MyLLM(LLM):
"""Configure the vLLM worker for Ray placement group execution."""
@@ -58,9 +71,11 @@ class TrainModel:
"""Ray actor that wraps the training model on a dedicated GPU."""
def __init__(self, model_name: str):
assigned_gpu = get_assigned_gpu()
self.model = AutoModelForCausalLM.from_pretrained(
model_name,
).to("cuda:0")
).to(f"cuda:{assigned_gpu}")
self.port = get_open_port()
self.master_address = get_ip()
tests/entrypoints/serve/render/__init__.py → examples/scale_out/__init__.py
View File
examples/disaggregated/disaggregated_serving/example_mm_serve.py → examples/scale_out/example_mm_serve.py
View File
examples/generate/token_generation_client.py → examples/scale_out/token_generation_client.py
View File
pyproject.toml
+2 -1
View File
@@ -129,7 +129,8 @@ extend-exclude = ["tests/models/fixtures/*", "tests/prompts/*", "tests/tokenizer
"docs/governance/process.md", "docs/assets/contributing/vllm_bench_serve_timeline.html",
"tests/v1/engine/test_fast_incdec_prefix_err.py", ".git/*", "csrc/cpu/sgl-kernels/*",
"rust/src/chat/src/renderer/deepseek_v32/fixtures/*",
"rust/src/parser/src/tool/gemma4.rs", "rust/src/text/src/output/decoded.rs",
"rust/src/parser/src/tool/gemma4.rs", "rust/src/parser/src/unified/gemma4.rs",
"rust/src/text/src/output/decoded.rs",
"rust/src/tokenizer/src/incremental.rs", "rust/src/parser/src/reasoning/tests.rs"]
ignore-hidden = false
requirements/common.txt
+1 -1
View File
@@ -33,7 +33,7 @@ partial-json-parser # used for parsing partial JSON outputs
jsonschema >= 4.23.0 # required for MiniMax M3 tool schema validation
pyzmq >= 25.0.0
msgspec
mistral_common[image] >= 1.11.3
mistral_common[image] >= 1.11.5
opencv-python-headless >= 4.13.0 # required for video IO
pyyaml
six>=1.16.0; python_version > '3.11' # transitive dependency of pandas that needs to be the latest version for python 3.12
requirements/cuda.txt
+3 -2
View File
@@ -8,9 +8,10 @@ torch==2.11.0
torchaudio==2.11.0
# These must be updated alongside torch
torchvision==0.26.0 # Required for phi3v processor. See https://github.com/pytorch/vision?tab=readme-ov-file#installation for corresponding version
PyNvVideoCodec==2.0.4
# FlashInfer should be updated together with the Dockerfile
flashinfer-python==0.6.12
flashinfer-cubin==0.6.12
flashinfer-python==0.6.13
flashinfer-cubin==0.6.13
apache-tvm-ffi==0.1.9
tilelang==0.1.9
nvidia-cudnn-frontend>=1.19.1
requirements/kv_connectors.txt
+1 -1
View File
@@ -2,5 +2,5 @@ lmcache >= 0.3.9
# CuPy 14.1.0 imports pytest from cupy.testing._random. Use <14.1.0
# until a fixed newer release is verified for runtime images.
cupy-cuda13x < 14.1.0
nixl == 1.2.0 # Required for disaggregated prefill
nixl == 1.3.0
mooncake-transfer-engine >= 0.3.8
requirements/test/cuda.in
+4 -5
View File
@@ -1,3 +1,5 @@
-r ../common.txt
# testing
pytest
tensorizer==2.10.1
@@ -13,7 +15,6 @@ albumentations # required for Nemotron Parse in test_common.py
av # required for audio_in_video tests
backoff # required for phi4mm test
blobfile # required for kimi-vl test
einops # required for MPT, qwen-vl
httpx
librosa # required for audio tests
vector_quantize_pytorch # required for minicpmo_26 test
@@ -31,10 +32,9 @@ torchaudio==2.11.0
torchvision==0.26.0
transformers_stream_generator # required for qwen-vl test
matplotlib # required for qwen-vl test
mistral_common[image,audio] >= 1.11.3 # required for voxtral test
mistral_common[image,audio] >= 1.11.5 # required for voxtral test
num2words # required for smolvlm test
open_clip_torch==2.32.0 # Required for nemotron_vl test, Nemotron Parse in test_common.py
opencv-python-headless >= 4.13.0 # required for video test
datamodel_code_generator # required for minicpm3 test
lm-eval[api]>=0.4.12 # required for model evaluation test
mteb[bm25s]>=2, <3 # required for mteb test
@@ -55,11 +55,9 @@ grpcio-reflection==1.78.0
arctic-inference == 0.1.1; platform_machine == "x86_64" # Required for suffix decoding test
numba == 0.65.0 # Required for N-gram speculative decoding
numpy
runai-model-streamer[s3,gcs,azure]==0.15.7
fastsafetensors>=0.3.2
instanttensor>=0.1.5; platform_machine == "x86_64"
pydantic>=2.12 # 2.11 leads to error on python 3.13
decord==0.6.0; platform_machine == "x86_64"
# terratorch is temporarily disabled while PyPI has the `lightning` package
# in `quarantined` status (every published terratorch version transitively
@@ -73,6 +71,7 @@ gpt-oss >= 0.0.7; python_version > '3.11'
perceptron # required for isaac test
kaldi-native-fbank >= 1.18.7 # required for fireredasr2 test
cohere_melody>=0.9.0 # required for cohere command reasoning parser test
# Newer versions of datasets require torchcoded, that makes the tests fail in CI because of a missing library.
# Older versions are in conflict with teerratorch requirements.
requirements/test/cuda.txt
+291 -19
View File
@@ -9,6 +9,7 @@ aiohappyeyeballs==2.6.1
aiohttp==3.13.3
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# aiohttp-cors
# datasets
# fsspec
@@ -24,17 +25,34 @@ albumentations==1.4.6
alembic==1.16.4
# via optuna
annotated-doc==0.0.4
# via fastapi
# via
# fastapi
# typer
annotated-types==0.7.0
# via pydantic
anyio==4.6.2.post1
anthropic==0.112.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
anyio==4.14.1
# via
# anthropic
# httpx
# mcp
# openai
# sse-starlette
# starlette
# watchfiles
apache-tvm-ffi==0.1.9
# via
# -c requirements/cuda.txt
# xgrammar
arctic-inference==0.1.1
# via -r requirements/test/cuda.in
argcomplete==3.5.1
# via datamodel-code-generator
astor==0.8.1
# via depyf
attrs==24.2.0
# via
# aiohttp
@@ -59,6 +77,8 @@ bitsandbytes==0.49.2
# via -r requirements/test/cuda.in
black==24.10.0
# via datamodel-code-generator
blake3==1.0.9
# via -r requirements/test/../common.txt
blobfile==3.0.0
# via -r requirements/test/cuda.in
bm25s==0.2.13
@@ -76,12 +96,17 @@ bounded-pool-executor==0.0.3
buildkite-test-collector==0.1.9
# via -r requirements/test/cuda.in
cachetools==5.5.2
# via google-auth
# via
# -r requirements/test/../common.txt
# google-auth
cbor2==6.1.2
# via -r requirements/test/../common.txt
certifi==2024.8.30
# via
# httpcore
# httpx
# requests
# sentry-sdk
cffi==2.0.0
# via
# cryptography
@@ -98,9 +123,13 @@ click==8.1.7
# jiwer
# nltk
# ray
# rich-toolkit
# schemathesis
# typer
# uvicorn
cloudpickle==3.1.2
# via -r requirements/test/../common.txt
cohere-melody==0.9.0
# via -r requirements/test/cuda.in
colorama==0.4.6
# via
# perceptron
@@ -109,6 +138,10 @@ colorful==0.5.6
# via ray
colorlog==6.10.1
# via optuna
compressed-tensors==0.17.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
contourpy==1.3.0
# via matplotlib
coverage==7.10.6
@@ -147,30 +180,49 @@ decorator==5.1.1
# via librosa
decord==0.6.0
# via -r requirements/test/cuda.in
depyf==0.20.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
detect-installer==0.1.0
# via fastapi-cloud-cli
dill==0.3.8
# via
# datasets
# depyf
# evaluate
# lm-eval
# multiprocess
diskcache==5.6.3
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
distlib==0.3.9
# via virtualenv
distro==1.9.0
# via
# anthropic
# openai
dnspython==2.7.0
# via email-validator
docker==7.1.0
# via gpt-oss
docopt==0.6.2
# via num2words
docstring-parser==0.18.0
# via anthropic
einops==0.8.1
# via
# -r requirements/test/cuda.in
# -r requirements/test/../common.txt
# encodec
# vector-quantize-pytorch
# vocos
einx==0.3.0
# via vector-quantize-pytorch
email-validator==2.2.0
# via pydantic
# via
# fastapi
# pydantic
encodec==0.1.1
# via vocos
et-xmlfile==2.0.0
@@ -180,7 +232,17 @@ evaluate==0.4.3
fastapi==0.136.3
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# gpt-oss
# model-hosting-container-standards
fastapi-cli==0.0.27
# via fastapi
fastapi-cloud-cli==0.21.0
# via fastapi-cli
fastar==0.11.0
# via
# fastapi
# fastapi-cloud-cli
fastparquet==2024.11.0
# via genai-perf
fastrlock==0.8.2
@@ -192,6 +254,7 @@ fastsafetensors==0.3.2
filelock==3.16.1
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# blobfile
# datasets
# huggingface-hub
@@ -241,7 +304,10 @@ google-crc32c==1.7.1
google-resumable-media==2.7.2
# via google-cloud-storage
googleapis-common-protos==1.70.0
# via google-api-core
# via
# google-api-core
# opentelemetry-exporter-otlp-proto-grpc
# opentelemetry-exporter-otlp-proto-http
gpt-oss==0.0.8
# via -r requirements/test/cuda.in
graphql-core==3.2.6
@@ -252,6 +318,7 @@ grpcio==1.78.0
# via
# -r requirements/test/cuda.in
# grpcio-reflection
# opentelemetry-exporter-otlp-proto-grpc
# ray
grpcio-reflection==1.78.0
# via -r requirements/test/cuda.in
@@ -273,12 +340,22 @@ html2text==2025.4.15
# via gpt-oss
httpcore==1.0.6
# via httpx
httptools==0.8.0
# via uvicorn
httpx==0.27.2
# via
# -r requirements/test/cuda.in
# anthropic
# fastapi
# fastapi-cloud-cli
# huggingface-hub
# mcp
# model-hosting-container-standards
# openai
# perceptron
# schemathesis
httpx-sse==0.4.3
# via mcp
huggingface-hub==1.10.2
# via
# accelerate
@@ -312,6 +389,8 @@ idna==3.10
# httpx
# requests
# yarl
ijson==3.5.0
# via -r requirements/test/../common.txt
imagehash==4.3.2
# via -r requirements/test/cuda.in
imageio==2.37.0
@@ -324,6 +403,8 @@ iniconfig==2.0.0
# via pytest
instanttensor==0.1.5
# via -r requirements/test/cuda.in
interegular==0.3.3
# via lm-format-enforcer
isodate==0.7.2
# via azure-storage-blob
isort==5.13.2
@@ -331,15 +412,21 @@ isort==5.13.2
jinja2==3.1.6
# via
# datamodel-code-generator
# fastapi
# genai-perf
# lm-eval
# torch
jiter==0.15.0
# via
# anthropic
# openai
jiwer==3.0.5
# via -r requirements/test/cuda.in
jmespath==1.0.1
# via
# boto3
# botocore
# model-hosting-container-standards
joblib==1.4.2
# via
# librosa
@@ -348,7 +435,9 @@ joblib==1.4.2
jsonschema==4.23.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# hypothesis-jsonschema
# mcp
# mistral-common
# ray
jsonschema-rs==0.46.5
@@ -363,6 +452,10 @@ kaleido==0.2.1
# via genai-perf
kiwisolver==1.4.7
# via matplotlib
lark==1.2.2
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
lazy-loader==0.4
# via
# librosa
@@ -371,10 +464,20 @@ libnacl==2.1.0
# via tensorizer
librosa==0.10.2.post1
# via -r requirements/test/cuda.in
llguidance==1.7.6
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
llvmlite==0.47.0
# via numba
lm-eval==0.4.12
# via -r requirements/test/cuda.in
lm-format-enforcer==0.11.3
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
loguru==0.7.3
# via compressed-tensors
lxml==5.3.0
# via
# blobfile
@@ -396,12 +499,19 @@ mbstrdecoder==1.1.3
# dataproperty
# pytablewriter
# typepy
mcp==1.28.1
# via -r requirements/test/../common.txt
mdurl==0.1.2
# via markdown-it-py
mistral-common==1.11.3
mistral-common==1.11.5
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# -r requirements/test/cuda.in
model-hosting-container-standards==0.1.16
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
more-itertools==10.5.0
# via lm-eval
mpmath==1.3.0
@@ -416,6 +526,8 @@ msgpack==1.1.0
# via
# librosa
# ray
msgspec==0.21.1
# via -r requirements/test/../common.txt
mteb==2.8.3
# via -r requirements/test/cuda.in
multidict==6.1.0
@@ -432,6 +544,8 @@ networkx==3.2.1
# via
# scikit-image
# torch
ninja==1.13.0
# via -r requirements/test/../common.txt
nltk==3.9.1
# via rouge-score
num2words==0.5.14
@@ -443,7 +557,7 @@ numba==0.65.0
# librosa
numpy==2.2.6
# via
# -r requirements/test/cuda.in
# -r requirements/test/../common.txt
# accelerate
# albumentations
# bitsandbytes
@@ -487,6 +601,7 @@ numpy==2.2.6
# transformers
# tritonclient
# vocos
# xgrammar
nvidia-cublas==13.1.0.3
# via
# cuda-toolkit
@@ -528,9 +643,14 @@ nvidia-nvtx==13.0.85
# via cuda-toolkit
open-clip-torch==2.32.0
# via -r requirements/test/cuda.in
openai==2.44.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
openai-harmony==0.0.4
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# gpt-oss
opencensus==0.11.4
# via ray
@@ -539,7 +659,7 @@ opencensus-context==0.1.3
opencv-python-headless==4.13.0.90
# via
# -c requirements/common.txt
# -r requirements/test/cuda.in
# -r requirements/test/../common.txt
# albumentations
# mistral-common
openpyxl==3.1.5
@@ -547,24 +667,54 @@ openpyxl==3.1.5
opentelemetry-api==1.35.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# opentelemetry-exporter-otlp-proto-grpc
# opentelemetry-exporter-otlp-proto-http
# opentelemetry-exporter-prometheus
# opentelemetry-sdk
# opentelemetry-semantic-conventions
opentelemetry-exporter-otlp==1.35.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
opentelemetry-exporter-otlp-proto-common==1.35.0
# via
# opentelemetry-exporter-otlp-proto-grpc
# opentelemetry-exporter-otlp-proto-http
opentelemetry-exporter-otlp-proto-grpc==1.35.0
# via opentelemetry-exporter-otlp
opentelemetry-exporter-otlp-proto-http==1.35.0
# via opentelemetry-exporter-otlp
opentelemetry-exporter-prometheus==0.56b0
# via ray
opentelemetry-proto==1.35.0
# via ray
# via
# opentelemetry-exporter-otlp-proto-common
# opentelemetry-exporter-otlp-proto-grpc
# opentelemetry-exporter-otlp-proto-http
# ray
opentelemetry-sdk==1.35.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# opentelemetry-exporter-otlp-proto-grpc
# opentelemetry-exporter-otlp-proto-http
# opentelemetry-exporter-prometheus
# ray
opentelemetry-semantic-conventions==0.56b0
# via opentelemetry-sdk
opentelemetry-semantic-conventions-ai==0.4.13
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
optuna==3.6.1
# via genai-perf
orjson==3.11.5
# via genai-perf
outlines-core==0.2.14
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
packaging==24.2
# via
# accelerate
@@ -576,6 +726,7 @@ packaging==24.2
# fastparquet
# huggingface-hub
# lazy-loader
# lm-format-enforcer
# matplotlib
# optuna
# peft
@@ -595,6 +746,8 @@ pandas==2.2.3
# fastparquet
# genai-perf
# statsmodels
partial-json-parser==0.2.1.1.post7
# via -r requirements/test/../common.txt
pathspec==0.12.1
# via black
pathvalidate==3.2.1
@@ -609,6 +762,7 @@ perf-analyzer==0.1.0
# via genai-perf
pillow==10.4.0
# via
# -r requirements/test/../common.txt
# genai-perf
# imagehash
# imageio
@@ -642,8 +796,14 @@ pqdm==0.2.0
prometheus-client==0.22.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# opentelemetry-exporter-prometheus
# prometheus-fastapi-instrumentator
# ray
prometheus-fastapi-instrumentator==8.0.2
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
propcache==0.2.0
# via
# aiohttp
@@ -653,6 +813,7 @@ proto-plus==1.26.1
protobuf==6.33.6
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# google-api-core
# googleapis-common-protos
# grpcio-reflection
@@ -662,11 +823,14 @@ protobuf==6.33.6
# tensorizer
psutil==6.1.0
# via
# -r requirements/test/../common.txt
# accelerate
# peft
# tensorizer
py==1.11.0
# via pytest-forked
py-cpuinfo==9.0.0
# via -r requirements/test/../common.txt
py-spy==0.4.0
# via ray
pyarrow==23.0.0
@@ -679,6 +843,8 @@ pyasn1==0.6.1
# rsa
pyasn1-modules==0.4.2
# via google-auth
pybase64==1.4.3
# via -r requirements/test/../common.txt
pycountry==24.6.1
# via pydantic-extra-types
pycparser==2.22
@@ -688,26 +854,43 @@ pycryptodomex==3.22.0
pydantic==2.12.0
# via
# -c requirements/common.txt
# -r requirements/test/cuda.in
# -r requirements/test/../common.txt
# albumentations
# anthropic
# compressed-tensors
# datamodel-code-generator
# fastapi
# fastapi-cloud-cli
# gpt-oss
# lm-format-enforcer
# mcp
# mistral-common
# model-hosting-container-standards
# mteb
# openai
# openai-harmony
# pydantic-extra-types
# pydantic-settings
# ray
# xgrammar
pydantic-core==2.41.1
# via pydantic
pydantic-extra-types==2.10.5
# via mistral-common
# via
# fastapi
# mistral-common
pydantic-settings==2.14.2
# via
# fastapi
# mcp
pygments==2.18.0
# via
# pytest
# rich
pyjwt==2.11.0
# via msal
# via
# mcp
# msal
pyparsing==3.2.0
# via matplotlib
pyrate-limiter==4.4.0
@@ -749,6 +932,16 @@ python-dateutil==2.9.0.post0
# matplotlib
# pandas
# typepy
python-dotenv==1.2.2
# via
# pydantic-settings
# uvicorn
python-json-logger==4.1.0
# via -r requirements/test/../common.txt
python-multipart==0.0.32
# via
# fastapi
# mcp
python-rapidjson==1.20
# via tritonclient
pytrec-eval-terrier==0.5.7
@@ -761,12 +954,14 @@ pywavelets==1.9.0
# via imagehash
pyyaml==6.0.2
# via
# -r requirements/test/../common.txt
# accelerate
# albumentations
# datamodel-code-generator
# datasets
# genai-perf
# huggingface-hub
# lm-format-enforcer
# optuna
# peft
# ray
@@ -774,7 +969,12 @@ pyyaml==6.0.2
# schemathesis
# timm
# transformers
# uvicorn
# vocos
pyzmq==27.1.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
rapidfuzz==3.12.1
# via jiwer
ray==2.48.0
@@ -787,6 +987,7 @@ referencing==0.35.1
# jsonschema-specifications
regex==2026.2.28
# via
# -r requirements/test/../common.txt
# nltk
# open-clip-torch
# sacrebleu
@@ -795,6 +996,7 @@ regex==2026.2.28
requests==2.32.3
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# azure-core
# buildkite-test-collector
# datasets
@@ -807,6 +1009,7 @@ requests==2.32.3
# mistral-common
# msal
# mteb
# opentelemetry-exporter-otlp-proto-http
# pooch
# ray
# responses
@@ -820,8 +1023,15 @@ rich==13.9.4
# genai-perf
# mteb
# perceptron
# rich-toolkit
# schemathesis
# typer
rich-toolkit==0.20.1
# via
# fastapi-cli
# fastapi-cloud-cli
rignore==0.7.6
# via fastapi-cloud-cli
rouge-score==0.1.2
# via lm-eval
rpds-py==0.20.1
@@ -845,6 +1055,7 @@ sacrebleu==2.4.3
safetensors==0.7.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# accelerate
# open-clip-torch
# peft
@@ -880,9 +1091,17 @@ sentence-transformers==5.2.0
# via
# -r requirements/test/cuda.in
# mteb
sentencepiece==0.2.1
# via -r requirements/test/../common.txt
sentry-sdk==2.63.0
# via fastapi-cloud-cli
setproctitle==1.3.7
# via -r requirements/test/../common.txt
setuptools==77.0.3
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# model-hosting-container-standards
# pytablewriter
# torch
shellingham==1.5.4
@@ -892,6 +1111,7 @@ shellingham==1.5.4
six==1.16.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# junit-xml
# opencensus
# python-dateutil
@@ -900,8 +1120,9 @@ smart-open==7.1.0
# via ray
sniffio==1.3.1
# via
# anyio
# anthropic
# httpx
# openai
sortedcontainers==2.4.0
# via hypothesis
soundfile==0.12.1
@@ -920,10 +1141,17 @@ sqlalchemy==2.0.41
# optuna
sqlitedict==2.1.0
# via lm-eval
sse-starlette==3.4.5
# via mcp
starlette==1.3.1
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# fastapi
# mcp
# model-hosting-container-standards
# prometheus-fastapi-instrumentator
# sse-starlette
# starlette-testclient
starlette-testclient==0.4.1
# via schemathesis
@@ -931,6 +1159,8 @@ statsmodels==0.14.4
# via genai-perf
structlog==25.4.0
# via gpt-oss
supervisor==4.3.0
# via model-hosting-container-standards
sympy==1.13.3
# via
# einx
@@ -960,6 +1190,7 @@ tifffile==2025.3.30
tiktoken==0.12.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# gpt-oss
# lm-eval
# mistral-common
@@ -971,6 +1202,7 @@ timm==1.0.17
tokenizers==0.22.2
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# -r requirements/test/cuda.in
# transformers
torch==2.11.0+cu130
@@ -979,6 +1211,7 @@ torch==2.11.0+cu130
# -r requirements/test/cuda.in
# accelerate
# bitsandbytes
# compressed-tensors
# encodec
# instanttensor
# mteb
@@ -992,6 +1225,7 @@ torch==2.11.0+cu130
# torchvision
# vector-quantize-pytorch
# vocos
# xgrammar
torchaudio==2.11.0+cu130
# via
# -c requirements/cuda.txt
@@ -1007,6 +1241,7 @@ torchvision==0.26.0+cu130
# timm
tqdm==4.67.3
# via
# -r requirements/test/../common.txt
# datasets
# evaluate
# huggingface-hub
@@ -1014,6 +1249,7 @@ tqdm==4.67.3
# mteb
# nltk
# open-clip-torch
# openai
# optuna
# peft
# pqdm
@@ -1023,15 +1259,20 @@ tqdm==4.67.3
transformers==5.5.3
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# -r requirements/test/cuda.in
# compressed-tensors
# genai-perf
# peft
# sentence-transformers
# transformers-stream-generator
# xgrammar
transformers-stream-generator==0.0.5
# via -r requirements/test/cuda.in
triton==3.6.0
# via torch
# via
# torch
# xgrammar
tritonclient==2.64.0
# via -r requirements/test/cuda.in
typepy==1.3.2
@@ -1039,8 +1280,10 @@ typepy==1.3.2
# dataproperty
# pytablewriter
# tabledata
typer==0.15.2
typer==0.26.8
# via
# fastapi-cli
# fastapi-cloud-cli
# fastsafetensors
# huggingface-hub
# perceptron
@@ -1048,9 +1291,13 @@ typer==0.15.2
typing-extensions==4.15.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# aiosignal
# albumentations
# alembic
# anthropic
# anyio
# apache-tvm-ffi
# azure-core
# azure-identity
# azure-storage-blob
@@ -1060,9 +1307,13 @@ typing-extensions==4.15.0
# huggingface-hub
# librosa
# lm-eval
# mcp
# mistral-common
# mteb
# openai
# opentelemetry-api
# opentelemetry-exporter-otlp-proto-grpc
# opentelemetry-exporter-otlp-proto-http
# opentelemetry-sdk
# opentelemetry-semantic-conventions
# pqdm
@@ -1070,17 +1321,20 @@ typing-extensions==4.15.0
# pydantic-core
# pydantic-extra-types
# pytest-asyncio
# rich-toolkit
# schemathesis
# sentence-transformers
# sqlalchemy
# starlette
# torch
# typer
# typing-inspection
# xgrammar
typing-inspection==0.4.2
# via
# fastapi
# mcp
# pydantic
# pydantic-settings
tzdata==2024.2
# via pandas
urllib3==2.2.3
@@ -1090,23 +1344,41 @@ urllib3==2.2.3
# docker
# requests
# responses
# sentry-sdk
# tritonclient
uvicorn==0.35.0
# via gpt-oss
# via
# fastapi
# fastapi-cli
# fastapi-cloud-cli
# gpt-oss
# mcp
uvloop==0.22.1
# via uvicorn
vector-quantize-pytorch==1.21.2
# via -r requirements/test/cuda.in
virtualenv==20.31.2
# via ray
vocos==0.1.0
# via -r requirements/test/cuda.in
watchfiles==1.2.0
# via
# -r requirements/test/../common.txt
# uvicorn
wcwidth==0.2.13
# via ftfy
websockets==16.0
# via uvicorn
werkzeug==3.1.3
# via schemathesis
word2number==1.1
# via lm-eval
wrapt==1.17.2
# via smart-open
xgrammar==0.2.3
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
xxhash==3.5.0
# via
# datasets
requirements/test/nightly-torch.txt
+1 -1
View File
@@ -23,7 +23,7 @@ jiwer # required for audio tests
timm # required for internvl test
transformers_stream_generator # required for qwen-vl test
matplotlib # required for qwen-vl test
mistral_common[image,audio] >= 1.11.3 # required for voxtral test
mistral_common[image,audio] >= 1.11.5 # required for voxtral test
num2words # required for smolvlm test
opencv-python-headless >= 4.13.0 # required for video test
datamodel_code_generator # required for minicpm3 test
requirements/test/rocm.in
+2 -5
View File
@@ -15,7 +15,6 @@ albumentations # required for Nemotron Parse in test_common.py
av # required for audio_in_video tests
backoff # required for phi4mm test
blobfile # required for kimi-vl test
einops # required for MPT, qwen-vl
httpx
librosa # required for audio tests
vector_quantize_pytorch # required for minicpmo_26 test
@@ -30,10 +29,9 @@ tblib # for pickling test exceptions
timm>=1.0.17 # required for internvl and gemma3n-mm test
transformers_stream_generator # required for qwen-vl test
matplotlib # required for qwen-vl test
mistral_common[image,audio]>=1.11.3 # required for voxtral test
mistral_common[image,audio]>=1.11.5 # required for voxtral test
num2words # required for smolvlm test
open_clip_torch==2.32.0 # Required for nemotron_vl test, Nemotron Parse in test_common.py
opencv-python-headless>=4.13.0 # required for video test
datamodel_code_generator # required for minicpm3 test
lm-eval[api]>=0.4.12 # required for model evaluation test
mteb[bm25s]>=2, <3 # required for mteb test
@@ -54,11 +52,9 @@ grpcio-reflection==1.78.0
arctic-inference==0.1.1 # Required for suffix decoding test
numba==0.65.0 # Required for N-gram speculative decoding
numpy
runai-model-streamer[s3,gcs,azure]==0.15.7
fastsafetensors>=0.3.2
instanttensor>=0.1.5
pydantic>=2.12 # 2.11 leads to error on python 3.13
decord==0.6.0
# Prithvi tests
@@ -74,6 +70,7 @@ gpt-oss>=0.0.7; python_version > '3.11'
perceptron # required for isaac test
kaldi-native-fbank>=1.18.7 # required for fireredasr2 test
cohere_melody>=0.9.0 # required for cohere command reasoning parser test
# Newer versions of datasets require torchcoded, that makes the tests fail in CI because of a missing library.
# Older versions are in conflict with terratorch requirements.
requirements/test/rocm.txt
+3 -5
View File
@@ -130,6 +130,8 @@ cloudpickle==3.1.2
# via
# -r requirements/test/../common.txt
# tilelang
cohere-melody==0.9.0
# via -r requirements/test/rocm.in
colorama==0.4.6
# via
# perceptron
@@ -205,7 +207,6 @@ docstring-parser==0.17.0
einops==0.8.2
# via
# -r requirements/test/../common.txt
# -r requirements/test/rocm.in
# encodec
# vector-quantize-pytorch
# vocos
@@ -499,7 +500,7 @@ mcp==1.27.0
# via -r requirements/test/../common.txt
mdurl==0.1.2
# via markdown-it-py
mistral-common==1.11.3
mistral-common==1.11.5
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
@@ -561,7 +562,6 @@ numba==0.65.0
numpy==2.2.6
# via
# -r requirements/test/../common.txt
# -r requirements/test/rocm.in
# accelerate
# albumentations
# bitsandbytes
@@ -630,7 +630,6 @@ opencv-python-headless==4.13.0.92
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# -r requirements/test/rocm.in
# albumentations
# mistral-common
openpyxl==3.1.5
@@ -834,7 +833,6 @@ pydantic==2.12.5
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# -r requirements/test/rocm.in
# albumentations
# anthropic
# compressed-tensors
requirements/test/xpu.in
+2
View File
@@ -1,3 +1,5 @@
-r ../common.txt
# --- Test Infrastructure ---
tblib
pytest
requirements/test/xpu.txt
+317 -5
View File
@@ -11,6 +11,7 @@ aiohappyeyeballs==2.6.1
aiohttp==3.13.4
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# fsspec
# gpt-oss
# lm-eval
@@ -24,12 +25,25 @@ annotated-doc==0.0.4
# typer
annotated-types==0.7.0
# via pydantic
anthropic==0.112.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
anyio==4.13.0
# via
# anthropic
# httpx
# mcp
# openai
# sse-starlette
# starlette
# watchfiles
apache-tvm-ffi==0.1.12
# via xgrammar
arctic-inference==0.1.1
# via -r requirements/test/xpu.in
astor==0.8.1
# via depyf
attrs==26.1.0
# via
# aiohttp
@@ -39,6 +53,8 @@ audioread==3.0.1
# via
# -r requirements/test/xpu.in
# librosa
blake3==1.0.9
# via -r requirements/test/../common.txt
blobfile==3.0.0
# via -r requirements/test/xpu.in
bm25s==0.2.13
@@ -47,13 +63,20 @@ bm25s==0.2.13
# mteb
bounded-pool-executor==0.0.3
# via pqdm
cachetools==7.1.4
# via -r requirements/test/../common.txt
cbor2==6.1.2
# via -r requirements/test/../common.txt
certifi==2026.2.25
# via
# httpcore
# httpx
# requests
# sentry-sdk
cffi==2.0.0
# via soundfile
# via
# cryptography
# soundfile
chardet==5.2.0
# via mbstrdecoder
charset-normalizer==3.4.6
@@ -64,13 +87,22 @@ click==8.3.1
# via
# jiwer
# nltk
# rich-toolkit
# schemathesis
# typer
# uvicorn
cloudpickle==3.1.2
# via -r requirements/test/../common.txt
colorama==0.4.6
# via sacrebleu
compressed-tensors==0.17.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
coverage==7.13.5
# via pytest-cov
cryptography==49.0.0
# via pyjwt
dataproperty==1.1.0
# via
# pytablewriter
@@ -82,16 +114,35 @@ datasets==4.8.4
# mteb
decorator==5.2.1
# via librosa
depyf==0.20.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
detect-installer==0.1.0
# via fastapi-cloud-cli
dill==0.4.1
# via
# datasets
# depyf
# evaluate
# lm-eval
# multiprocess
diskcache==5.6.3
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
distro==1.9.0
# via
# anthropic
# openai
dnspython==2.8.0
# via email-validator
docker==7.1.0
# via gpt-oss
docopt==0.6.2
# via num2words
docstring-parser==0.18.0
# via anthropic
dpcpp-cpp-rt==2025.3.2
# via
# onemkl-sycl-blas
@@ -100,15 +151,30 @@ dpcpp-cpp-rt==2025.3.2
# onemkl-sycl-rng
# onemkl-sycl-sparse
# torch
einops==0.8.2
# via -r requirements/test/../common.txt
email-validator==2.3.0
# via
# fastapi
# pydantic
evaluate==0.4.6
# via lm-eval
fastapi==0.135.2
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# gpt-oss
# model-hosting-container-standards
fastapi-cli==0.0.27
# via fastapi
fastapi-cloud-cli==0.21.0
# via fastapi-cli
fastar==0.11.0
# via fastapi-cloud-cli
filelock==3.25.2
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# blobfile
# datasets
# huggingface-hub
@@ -124,10 +190,16 @@ fsspec==2026.2.0
# evaluate
# huggingface-hub
# torch
googleapis-common-protos==1.75.0
# via
# opentelemetry-exporter-otlp-proto-grpc
# opentelemetry-exporter-otlp-proto-http
gpt-oss==0.0.8
# via -r requirements/test/xpu.in
graphql-core==3.2.8
# via hypothesis-graphql
grpcio==1.81.1
# via opentelemetry-exporter-otlp-proto-grpc
h11==0.16.0
# via
# httpcore
@@ -140,11 +212,21 @@ html2text==2025.4.15
# via gpt-oss
httpcore==1.0.9
# via httpx
httptools==0.8.0
# via uvicorn
httpx==0.28.1
# via
# anthropic
# datasets
# fastapi
# fastapi-cloud-cli
# huggingface-hub
# mcp
# model-hosting-container-standards
# openai
# schemathesis
httpx-sse==0.4.3
# via mcp
huggingface-hub==1.10.2
# via
# accelerate
@@ -166,9 +248,12 @@ hypothesis-jsonschema==0.23.1
idna==3.11
# via
# anyio
# email-validator
# httpx
# requests
# yarl
ijson==3.5.0
# via -r requirements/test/../common.txt
imageio==2.37.3
# via scikit-image
impi-rt==2021.17.2
@@ -212,13 +297,22 @@ intel-sycl-rt==2025.3.2
# dpcpp-cpp-rt
# oneccl
# torch
interegular==0.3.3
# via lm-format-enforcer
jinja2==3.1.6
# via
# -c requirements/xpu.txt
# fastapi
# lm-eval
# torch
jiter==0.15.0
# via
# anthropic
# openai
jiwer==4.0.0
# via -r requirements/test/xpu.in
jmespath==1.1.0
# via model-hosting-container-standards
joblib==1.5.3
# via
# librosa
@@ -227,7 +321,9 @@ joblib==1.5.3
jsonschema==4.26.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# hypothesis-jsonschema
# mcp
# mistral-common
# schemathesis
jsonschema-rs==0.45.0
@@ -236,16 +332,30 @@ jsonschema-specifications==2025.9.1
# via jsonschema
junit-xml==1.9
# via schemathesis
lark==1.2.2
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
lazy-loader==0.5
# via
# librosa
# scikit-image
librosa==0.10.2.post1
# via -r requirements/test/xpu.in
llguidance==1.7.6
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
llvmlite==0.47.0
# via numba
lm-eval==0.4.12
# via -r requirements/test/xpu.in
lm-format-enforcer==0.11.3
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
loguru==0.7.3
# via compressed-tensors
lxml==6.0.2
# via
# blobfile
@@ -262,11 +372,14 @@ mbstrdecoder==1.1.4
# dataproperty
# pytablewriter
# typepy
mcp==1.28.1
# via -r requirements/test/../common.txt
mdurl==0.1.2
# via markdown-it-py
mistral-common==1.11.3
mistral-common==1.11.5
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# -r requirements/test/xpu.in
mkl==2025.3.1
# via
@@ -276,6 +389,10 @@ mkl==2025.3.1
# onemkl-sycl-rng
# onemkl-sycl-sparse
# torch
model-hosting-container-standards==0.1.16
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
modelscope==1.35.3
# via -r requirements/test/xpu.in
more-itertools==10.8.0
@@ -284,6 +401,8 @@ mpmath==1.3.0
# via sympy
msgpack==1.1.2
# via librosa
msgspec==0.21.1
# via -r requirements/test/../common.txt
mteb==2.12.7
# via -r requirements/test/xpu.in
multidict==6.7.1
@@ -298,6 +417,8 @@ networkx==3.6.1
# via
# scikit-image
# torch
ninja==1.13.0
# via -r requirements/test/../common.txt
nltk==3.9.4
# via rouge-score
num2words==0.5.14
@@ -308,6 +429,7 @@ numba==0.65.0
# librosa
numpy==2.2.6
# via
# -r requirements/test/../common.txt
# accelerate
# albumentations
# bm25s
@@ -333,6 +455,7 @@ numpy==2.2.6
# tifffile
# torchvision
# transformers
# xgrammar
oneccl==2021.17.2
# via
# oneccl-devel
@@ -356,15 +479,65 @@ onemkl-sycl-rng==2025.3.1
# via torch
onemkl-sycl-sparse==2025.3.1
# via torch
openai==2.44.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
openai-harmony==0.0.8
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# gpt-oss
opencv-python-headless==4.13.0.92
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# albumentations
# mistral-common
opentelemetry-api==1.43.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# opentelemetry-exporter-otlp-proto-grpc
# opentelemetry-exporter-otlp-proto-http
# opentelemetry-sdk
# opentelemetry-semantic-conventions
opentelemetry-exporter-otlp==1.43.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
opentelemetry-exporter-otlp-proto-common==1.43.0
# via
# opentelemetry-exporter-otlp-proto-grpc
# opentelemetry-exporter-otlp-proto-http
opentelemetry-exporter-otlp-proto-grpc==1.43.0
# via opentelemetry-exporter-otlp
opentelemetry-exporter-otlp-proto-http==1.43.0
# via opentelemetry-exporter-otlp
opentelemetry-proto==1.43.0
# via
# opentelemetry-exporter-otlp-proto-common
# opentelemetry-exporter-otlp-proto-grpc
# opentelemetry-exporter-otlp-proto-http
opentelemetry-sdk==1.43.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# opentelemetry-exporter-otlp-proto-grpc
# opentelemetry-exporter-otlp-proto-http
# opentelemetry-semantic-conventions-ai
opentelemetry-semantic-conventions==0.64b0
# via
# opentelemetry-sdk
# opentelemetry-semantic-conventions-ai
opentelemetry-semantic-conventions-ai==0.5.1
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
outlines-core==0.2.14
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
packaging==26.0
# via
# -c requirements/xpu.txt
@@ -373,6 +546,7 @@ packaging==26.0
# evaluate
# huggingface-hub
# lazy-loader
# lm-format-enforcer
# modelscope
# pooch
# pytest
@@ -384,10 +558,13 @@ pandas==3.0.1
# via
# datasets
# evaluate
partial-json-parser==0.2.1.1.post7
# via -r requirements/test/../common.txt
pathvalidate==3.3.1
# via pytablewriter
pillow==12.1.1
# via
# -r requirements/test/../common.txt
# imageio
# mistral-common
# scikit-image
@@ -410,16 +587,37 @@ portalocker==3.2.0
# via sacrebleu
pqdm==0.2.0
# via -r requirements/test/xpu.in
prometheus-client==0.25.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# prometheus-fastapi-instrumentator
prometheus-fastapi-instrumentator==8.0.2
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
propcache==0.4.1
# via
# aiohttp
# yarl
protobuf==7.35.1
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# googleapis-common-protos
# opentelemetry-proto
psutil==7.2.2
# via accelerate
# via
# -r requirements/test/../common.txt
# accelerate
py==1.11.0
# via pytest-forked
py-cpuinfo==9.0.0
# via -r requirements/test/../common.txt
pyarrow==23.0.1
# via datasets
pybase64==1.4.3
# via -r requirements/test/../common.txt
pycountry==26.2.16
# via pydantic-extra-types
pycparser==3.0
@@ -429,23 +627,41 @@ pycryptodomex==3.23.0
pydantic==2.12.5
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# albumentations
# anthropic
# compressed-tensors
# fastapi
# fastapi-cloud-cli
# gpt-oss
# lm-format-enforcer
# mcp
# mistral-common
# model-hosting-container-standards
# mteb
# openai
# openai-harmony
# pydantic-extra-types
# pydantic-settings
# xgrammar
pydantic-core==2.41.5
# via pydantic
pydantic-extra-types==2.11.1
# via mistral-common
# via
# fastapi
# mistral-common
pydantic-settings==2.14.2
# via
# fastapi
# mcp
pyelftools==0.32
# via triton-xpu
pygments==2.20.0
# via
# pytest
# rich
pyjwt==2.13.0
# via mcp
pyrate-limiter==4.1.0
# via schemathesis
pystemmer==3.0.0
@@ -480,19 +696,36 @@ python-dateutil==2.9.0.post0
# via
# pandas
# typepy
python-dotenv==1.2.2
# via
# pydantic-settings
# uvicorn
python-json-logger==4.1.0
# via -r requirements/test/../common.txt
python-multipart==0.0.32
# via
# fastapi
# mcp
pytrec-eval-terrier==0.5.10
# via mteb
pytz==2026.1.post1
# via typepy
pyyaml==6.0.3
# via
# -r requirements/test/../common.txt
# accelerate
# albumentations
# datasets
# huggingface-hub
# lm-format-enforcer
# schemathesis
# timm
# transformers
# uvicorn
pyzmq==27.1.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
rapidfuzz==3.12.1
# via
# -r requirements/test/xpu.in
@@ -503,6 +736,7 @@ referencing==0.37.0
# jsonschema-specifications
regex==2026.3.32
# via
# -r requirements/test/../common.txt
# nltk
# sacrebleu
# tiktoken
@@ -510,6 +744,7 @@ regex==2026.3.32
requests==2.33.1
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# datasets
# docker
# evaluate
@@ -518,6 +753,7 @@ requests==2.33.1
# mistral-common
# modelscope
# mteb
# opentelemetry-exporter-otlp-proto-http
# pooch
# schemathesis
# starlette-testclient
@@ -525,8 +761,15 @@ requests==2.33.1
rich==14.3.3
# via
# mteb
# rich-toolkit
# schemathesis
# typer
rich-toolkit==0.20.1
# via
# fastapi-cli
# fastapi-cloud-cli
rignore==0.7.6
# via fastapi-cloud-cli
rouge-score==0.1.2
# via lm-eval
rpds-py==0.30.0
@@ -538,6 +781,7 @@ sacrebleu==2.6.0
safetensors==0.7.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# accelerate
# timm
# transformers
@@ -564,10 +808,18 @@ scipy==1.17.1
# sentence-transformers
sentence-transformers==5.3.0
# via mteb
sentencepiece==0.2.1
# via -r requirements/test/../common.txt
sentry-sdk==2.63.0
# via fastapi-cloud-cli
setproctitle==1.3.7
# via -r requirements/test/../common.txt
setuptools==80.10.2
# via
# -c requirements/common.txt
# -c requirements/xpu.txt
# -r requirements/test/../common.txt
# model-hosting-container-standards
# modelscope
# pytablewriter
# torch
@@ -576,9 +828,14 @@ shellingham==1.5.4
six==1.17.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# junit-xml
# python-dateutil
# rouge-score
sniffio==1.3.1
# via
# anthropic
# openai
sortedcontainers==2.4.0
# via hypothesis
soundfile==0.13.1
@@ -593,15 +850,24 @@ soxr==0.5.0.post1
# mistral-common
sqlitedict==2.1.0
# via lm-eval
sse-starlette==3.4.5
# via mcp
starlette==1.3.1
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# fastapi
# mcp
# model-hosting-container-standards
# prometheus-fastapi-instrumentator
# sse-starlette
# starlette-testclient
starlette-testclient==0.4.1
# via schemathesis
structlog==25.5.0
# via gpt-oss
supervisor==4.3.0
# via model-hosting-container-standards
sympy==1.14.0
# via torch
tabledata==1.3.4
@@ -636,6 +902,7 @@ tifffile==2026.3.3
tiktoken==0.12.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# gpt-oss
# lm-eval
# mistral-common
@@ -644,19 +911,23 @@ timm==1.0.17
tokenizers==0.22.2
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# transformers
torch==2.12.0+xpu
# via
# -c requirements/xpu.txt
# accelerate
# compressed-tensors
# mteb
# sentence-transformers
# timm
# torchvision
# xgrammar
torchvision==0.27.0+xpu
# via timm
tqdm==4.67.3
# via
# -r requirements/test/../common.txt
# datasets
# evaluate
# huggingface-hub
@@ -664,13 +935,19 @@ tqdm==4.67.3
# modelscope
# mteb
# nltk
# openai
# pqdm
# sentence-transformers
# transformers
transformers==5.5.3
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# compressed-tensors
# sentence-transformers
# xgrammar
triton==3.7.1
# via xgrammar
triton-xpu==3.7.1
# via torch
typepy==1.3.4
@@ -680,36 +957,53 @@ typepy==1.3.4
# tabledata
typer==0.24.1
# via
# fastapi-cli
# fastapi-cloud-cli
# huggingface-hub
# transformers
typing-extensions==4.15.0
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
# aiosignal
# albumentations
# anthropic
# anyio
# apache-tvm-ffi
# chz
# fastapi
# grpcio
# huggingface-hub
# librosa
# lm-eval
# mcp
# mistral-common
# mteb
# openai
# opentelemetry-api
# opentelemetry-exporter-otlp-proto-grpc
# opentelemetry-exporter-otlp-proto-http
# opentelemetry-sdk
# opentelemetry-semantic-conventions
# pqdm
# pydantic
# pydantic-core
# pydantic-extra-types
# pytest-asyncio
# referencing
# rich-toolkit
# schemathesis
# sentence-transformers
# starlette
# torch
# typing-inspection
# xgrammar
typing-inspection==0.4.2
# via
# fastapi
# mcp
# pydantic
# pydantic-settings
umf==1.0.3
# via
# intel-cmplr-lib-ur
@@ -720,12 +1014,30 @@ urllib3==2.6.3
# docker
# modelscope
# requests
# sentry-sdk
uvicorn==0.42.0
# via gpt-oss
# via
# fastapi
# fastapi-cli
# fastapi-cloud-cli
# gpt-oss
# mcp
uvloop==0.22.1
# via uvicorn
watchfiles==1.2.0
# via
# -r requirements/test/../common.txt
# uvicorn
websockets==16.0
# via uvicorn
werkzeug==3.1.7
# via schemathesis
word2number==1.1
# via lm-eval
xgrammar==0.2.3
# via
# -c requirements/common.txt
# -r requirements/test/../common.txt
xxhash==3.6.0
# via
# datasets
rust/Cargo.lock
Generated
+74 -844
View File
File diff suppressed because it is too large Load Diff
rust/Cargo.toml
+8 -7
View File
@@ -23,7 +23,7 @@ license = "Apache-2.0"
[workspace.dependencies]
anyhow = "1.0.100"
arc-swap = "1.9.0"
async-openai = "0.33.1"
async-openai = { version = "0.33.1", default-features = false, features = ["native-tls"] }
async-trait = "0.1.89"
asynk-strim-attr = "0.1.0"
axum = "0.8.8"
@@ -37,22 +37,22 @@ easy-ext = "1.0.3"
educe = "0.6.0"
enum-as-inner = "0.7.0"
expect-test = "1.5.1"
fastokens = "0.2.0"
fastokens = { version = "0.2.1", default-features = false }
futures = "0.3.31"
half = { version = "2.7.1", features = ["bytemuck"] }
hex = "0.4.3"
hf-hub = { version = "0.5.0", features = ["tokio"] }
hf-hub = { version = "0.5.0", default-features = false, features = ["tokio"] }
http-body = "1.0.1"
indexmap = "2.13.0"
itertools = "0.14.0"
libc = "0.2.177"
llm-multimodal = { git = "https://github.com/vllm-project/llm-multimodal", rev = "5b558989844d1c7af3e43d0f604069ffd9c06320" }
llm-multimodal = { git = "https://github.com/vllm-project/llm-multimodal", rev = "046b669bd1c4faa2a7e05344d8cbf7b2befb37d5" }
mimalloc = "0.1.52"
minijinja = { version = "2.0", features = ["unstable_machinery", "json", "builtins", "loader", "loop_controls", "preserve_order"] }
minijinja-contrib = { version = "2.0", features = ["pycompat"] }
native-tls-vendored = { package = "native-tls", version = "0.2.18", features = ["vendored"] }
ndarray = { version = "0.16.1", features = ["serde"] }
openai-harmony = "0.0.8"
openai-harmony = { package = "oss-harmony", git = "https://github.com/oss-harmony/harmony", tag = "v0.0.11", default-features = false }
openai-protocol = "1.6.0"
parking_lot = "0.12.5"
paste = "1.0.15"
@@ -64,7 +64,7 @@ pyo3 = "0.28.3"
pythonize = "0.28.0"
rand = "0.9.2"
reasoning-parser = "1.2.2"
reqwest = { version = "0.12.8", default-features = false, features = ["rustls-tls"] }
reqwest = { version = "0.12.8", default-features = false, features = ["native-tls"] }
riptoken = { version = "0.3.0", default-features = false }
rmp-serde = "1.3.1"
rmpv = { version = "1.3.1", features = ["with-serde"] }
@@ -79,6 +79,7 @@ serde_with = "3.18.0"
serial_test = { version = "3.2.0", features = ["file_locks"] }
sha2 = "0.10.9"
socket2 = "0.6.3"
strum = { version = "0.27.2", features = ["derive"] }
subenum = "1.1.3"
subtle = "2.6"
task-local = "0.1.1"
@@ -120,7 +121,7 @@ vllm-parser = { path = "src/parser" }
vllm-server = { path = "src/server" }
vllm-text = { path = "src/text" }
vllm-tokenizer = { path = "src/tokenizer" }
winnow = "1.0.2"
winnow = { version = "1.0.2", features = ["simd"] }
xgrammar-structural-tag = "0.1.0"
zeromq = { version = "0.6.0", default-features = false, features = [
"tokio-runtime",
rust/deny.toml
+15
View File
@@ -0,0 +1,15 @@
[bans]
multiple-versions = "allow"
deny = [
# TLS / crypto provider
# We prefer the system's TLS (e.g. OpenSSL) over Rust implementations.
{ name = "rustls" },
{ name = "ring" },
{ name = "aws-lc-rs" },
{ name = "aws-lc-sys" },
{ name = "s2n-tls" },
{ name = "s2n-tls-sys" },
{ name = "boring" },
{ name = "boring-sys" },
]
rust/src/chat/Cargo.toml
+2
View File
@@ -21,9 +21,11 @@ serde.workspace = true
serde-json-fmt.workspace = true
serde_json.workspace = true
serde_with.workspace = true
strum.workspace = true
subenum.workspace = true
thiserror.workspace = true
thiserror-ext.workspace = true
time.workspace = true
tokio.workspace = true
tracing.workspace = true
trait-set.workspace = true
rust/src/chat/src/backend/hf.rs
+48 -8
View File
@@ -15,7 +15,9 @@ use crate::output::{
DefaultChatOutputProcessor, HarmonyChatOutputProcessor, validate_harmony_parser_overrides,
};
use crate::renderer::hf::{HfChatRenderer, MultimodalRenderInfo};
use crate::renderer::{DeepSeekV4ChatRenderer, DeepSeekV32ChatRenderer, DynChatRenderer};
use crate::renderer::{
DeepSeekV4ChatRenderer, DeepSeekV32ChatRenderer, DynChatRenderer, HarmonyChatRenderer,
};
use crate::request::ChatRequest;
use crate::{DynChatOutputProcessor, RendererSelection};
@@ -61,6 +63,7 @@ impl HfChatBackend {
)?),
RendererSelection::DeepSeekV32 => Arc::new(DeepSeekV32ChatRenderer::new()),
RendererSelection::DeepSeekV4 => Arc::new(DeepSeekV4ChatRenderer::new()),
RendererSelection::Harmony => Arc::new(HarmonyChatRenderer::new()?),
};
info!(
@@ -148,13 +151,15 @@ mod tests {
use std::sync::Arc;
use tempfile::tempdir;
use thiserror_ext::AsReport as _;
use vllm_text::Prompt;
use vllm_text::backend::hf::TokenizerSource;
use vllm_text::tokenizer::{DynTokenizer, Tokenizer};
use super::HfChatBackend;
use crate::RendererSelection;
use crate::backend::{ChatBackend, LoadModelBackendsOptions};
use crate::backend::{ChatBackend, LoadModelBackendsOptions, NewChatOutputProcessorOptions};
use crate::request::{ChatContent, ChatMessage, ChatRequest};
use crate::{ParserSelection, RendererSelection};
fn request_with_user_text(text: &str) -> ChatRequest {
ChatRequest {
@@ -219,12 +224,12 @@ mod tests {
Arc::new(TestTokenizer)
}
fn render_prompt(
fn backend_for_selection(
renderer: RendererSelection,
config_json: &str,
tokenizer_config_json: &str,
) -> String {
let backend = HfChatBackend::from_resolved_model_files(
) -> HfChatBackend {
HfChatBackend::from_resolved_model_files(
resolved_files(config_json, tokenizer_config_json),
"test-model".to_string(),
LoadModelBackendsOptions {
@@ -236,9 +241,15 @@ mod tests {
},
test_tokenizer(),
)
.unwrap();
.unwrap()
}
backend
fn render_prompt(
renderer: RendererSelection,
config_json: &str,
tokenizer_config_json: &str,
) -> String {
backend_for_selection(renderer, config_json, tokenizer_config_json)
.chat_renderer()
.render(&request_with_user_text("hello"))
.unwrap()
@@ -272,6 +283,35 @@ mod tests {
assert_eq!(prompt, "hello");
}
#[test]
fn auto_uses_harmony_renderer_and_output_processor_for_gpt_oss_model_type() {
let backend = backend_for_selection(
RendererSelection::Auto,
r#"{"model_type":"gpt_oss"}"#,
r#"{"chat_template":"{{ messages[0].content }}"}"#,
);
let prompt =
backend.chat_renderer().render(&request_with_user_text("hello")).unwrap().prompt;
assert!(matches!(prompt, Prompt::TokenIds(_)));
let mut request = request_with_user_text("hello");
let error = match backend.new_chat_output_processor(
&mut request,
NewChatOutputProcessorOptions {
tool_call_parser: &ParserSelection::Explicit("json".to_string()),
reasoning_parser: &ParserSelection::Auto,
},
) {
Ok(_) => panic!("gpt_oss should reject generic parser overrides"),
Err(error) => error,
};
assert_eq!(
error.to_report_string(),
"gpt_oss uses native Harmony output parsing; generic tool parser override `json` is not supported"
);
}
#[test]
fn language_model_only_skips_multimodal_preprocessor_config() {
let mut files = resolved_files(
rust/src/chat/src/error.rs
+11
View File
@@ -74,6 +74,17 @@ pub enum Error {
pub type Result<T> = std::result::Result<T, Error>;
impl Error {
/// Whether this error represents invalid user request parameters.
pub fn is_request_validation_error(&self) -> bool {
match self {
Self::PromptTooLong { .. } => true,
Self::Text(error) => error.is_request_validation_error(),
_ => false,
}
}
}
/// Format the available-parser suffix used in user-facing error messages.
fn available_parser_hint(available_names: &[String]) -> String {
if available_names.is_empty() {
rust/src/chat/src/lib.rs
+2 -2
View File
@@ -29,8 +29,8 @@ pub use parser::reasoning::{
pub use parser::tool::{ToolParser, ToolParserError, ToolParserFactory};
pub use renderer::hf::ChatTemplateContentFormatOption;
pub use renderer::{
ChatRenderer, DeepSeekV4ChatRenderer, DeepSeekV32ChatRenderer, DynChatRenderer, RenderedPrompt,
RendererSelection,
ChatRenderer, DeepSeekV4ChatRenderer, DeepSeekV32ChatRenderer, DynChatRenderer,
HarmonyChatRenderer, RenderedPrompt, RendererSelection,
};
pub use request::{
ChatContent, ChatContentPart, ChatMessage, ChatOptions, ChatRequest, ChatRole, ChatTool,
rust/src/chat/src/multimodal.rs
+7 -8
View File
@@ -245,16 +245,15 @@ pub(crate) async fn finalize_rendered_prompt(
return Ok((rendered.prompt, None));
}
let info = info.ok_or(Error::UnsupportedMultimodalRenderer)?;
let Prompt::Text(prompt) = rendered.prompt else {
bail_multimodal!("multimodal chat renderer must return a text prompt before expansion");
let mut prompt_token_ids = match rendered.prompt {
Prompt::Text(prompt) => info
.context
.tokenizer()
.encode(&prompt, request.add_special_tokens)
.map_err(|error| multimodal!("{error}"))?,
Prompt::TokenIds(token_ids) => token_ids,
};
let media_parts = extract_media_parts(request)?;
let mut prompt_token_ids = info
.context
.tokenizer()
.encode(&prompt, request.add_special_tokens)
.map_err(|error| multimodal!("{error}"))?;
let prepared = info.prepare_multimodal(media_parts, &mut prompt_token_ids, model_dtype).await?;
Ok((Prompt::TokenIds(prompt_token_ids), Some(prepared)))
rust/src/chat/src/output/default/mod.rs
+149 -15
View File
@@ -18,7 +18,8 @@ use crate::output::{ChatOutputProcessor, DynChatEventStream, DynDecodedTextEvent
use crate::parser::ParserSelection;
use crate::parser::reasoning::{ReasoningParser, ReasoningParserFactory};
use crate::parser::tool::{ToolParser, ToolParserFactory};
use crate::request::ChatRequest;
use crate::parser::unified::UnifiedParserFactory;
use crate::request::{ChatRequest, ChatTool};
use crate::{Error, Result as ChatResult};
/// Default request-scoped output processor used by Hugging Face style chat
@@ -46,20 +47,31 @@ impl DefaultChatOutputProcessor {
tool_call_parser: &ParserSelection,
reasoning_parser: &ParserSelection,
) -> ChatResult<Self> {
let tool_parsing_enabled = request.tool_parsing_enabled();
let tool_parser = if tool_parsing_enabled {
Some(Self::resolve_tool_parser(
request,
let parser = if tool_call_parser == reasoning_parser
&& let Some(parser) = Self::resolve_optional_unified_parser(
&request.tools,
model_id,
tokenizer.clone(),
tool_call_parser,
)?)
)? {
parser
} else {
None
let tool_parsing_enabled = request.tool_parsing_enabled();
let tool_parser = if tool_parsing_enabled {
Some(Self::resolve_tool_parser(
&request.tools,
model_id,
tool_call_parser,
)?)
} else {
None
};
let reasoning_parser =
Self::resolve_optional_reasoning_parser(model_id, tokenizer, reasoning_parser)?;
Box::new(CombinedParser::new(reasoning_parser, tool_parser)) as Box<dyn UnifiedParser>
};
let reasoning_parser =
Self::resolve_optional_reasoning_parser(model_id, tokenizer, reasoning_parser)?;
let parser: Box<dyn UnifiedParser> =
Box::new(CombinedParser::new(reasoning_parser, tool_parser));
apply_structural_tag_constraint(request, parser.structural_tag_model())?;
if parser.preserve_special_tokens() {
request.decode_options.skip_special_tokens = false;
@@ -84,7 +96,7 @@ impl DefaultChatOutputProcessor {
}
fn resolve_tool_parser(
request: &mut ChatRequest,
tools: &[ChatTool],
model_id: &str,
selection: &ParserSelection,
) -> ChatResult<Box<dyn ToolParser>> {
@@ -100,14 +112,36 @@ impl DefaultChatOutputProcessor {
ParserSelection::Explicit(name) => name.as_str(),
};
let parser = factory.create(parser_name, &request.tools)?;
apply_structural_tag_constraint(request, parser.as_ref())?;
let parser = factory.create(parser_name, tools)?;
TOOL_PARSER_LOG_ONCE.call_once(|| info!(parser_name, "using tool parser"));
Ok(parser)
}
fn resolve_optional_unified_parser(
tools: &[ChatTool],
model_id: &str,
tokenizer: DynTokenizer,
selection: &ParserSelection,
) -> ChatResult<Option<Box<dyn UnifiedParser>>> {
let factory = UnifiedParserFactory::global();
let parser_name = match selection {
ParserSelection::Auto => factory.resolve_name_for_model(model_id),
ParserSelection::None => None,
ParserSelection::Explicit(name) if factory.contains(name) => Some(name.as_str()),
ParserSelection::Explicit(_) => None,
};
let Some(parser_name) = parser_name else {
return Ok(None);
};
let parser = factory.create(parser_name, tools, tokenizer)?;
UNIFIED_PARSER_LOG_ONCE.call_once(|| info!(parser_name, "using unified parser"));
Ok(Some(parser))
}
fn resolve_optional_reasoning_parser(
model_id: &str,
tokenizer: DynTokenizer,
@@ -134,6 +168,7 @@ impl DefaultChatOutputProcessor {
static TOOL_PARSER_LOG_ONCE: Once = Once::new();
static REASONING_PARSER_LOG_ONCE: Once = Once::new();
static UNIFIED_PARSER_LOG_ONCE: Once = Once::new();
impl ChatOutputProcessor for DefaultChatOutputProcessor {
/// Transforms a raw generate-output token stream into structured chat
@@ -149,3 +184,102 @@ impl ChatOutputProcessor for DefaultChatOutputProcessor {
Ok(structured.boxed())
}
}
#[cfg(test)]
mod tests {
use std::sync::Arc;
use vllm_tokenizer::Tokenizer;
use super::DefaultChatOutputProcessor;
use crate::Error;
use crate::parser::ParserSelection;
use crate::request::ChatRequest;
struct FakeTokenizer;
impl Tokenizer for FakeTokenizer {
fn encode(
&self,
text: &str,
_add_special_tokens: bool,
) -> vllm_tokenizer::Result<Vec<u32>> {
Ok(text.chars().map(u32::from).collect())
}
fn decode(
&self,
token_ids: &[u32],
_skip_special_tokens: bool,
) -> vllm_tokenizer::Result<String> {
Ok(token_ids
.iter()
.map(|token_id| char::from_u32(*token_id).unwrap_or('\u{FFFD}'))
.collect())
}
fn token_to_id(&self, token: &str) -> Option<u32> {
match token {
"<|channel>" => Some(1),
"<channel|>" => Some(2),
_ => None,
}
}
}
fn tokenizer() -> Arc<FakeTokenizer> {
Arc::new(FakeTokenizer)
}
#[test]
fn equal_explicit_gemma4_uses_unified_parser() {
let mut request = ChatRequest::for_test();
let selection = ParserSelection::Explicit("gemma4".to_string());
DefaultChatOutputProcessor::new(
&mut request,
"other-model",
tokenizer(),
&selection,
&selection,
)
.unwrap();
}
#[test]
fn auto_auto_gemma4_model_uses_unified_parser() {
let mut request = ChatRequest::for_test();
DefaultChatOutputProcessor::new(
&mut request,
"google/gemma-4-27b-it",
tokenizer(),
&ParserSelection::Auto,
&ParserSelection::Auto,
)
.unwrap();
}
#[test]
fn mixed_gemma4_selection_uses_split_dummy_error() {
let mut request = ChatRequest::for_test();
let error = match DefaultChatOutputProcessor::new(
&mut request,
"other-model",
tokenizer(),
&ParserSelection::Auto,
&ParserSelection::Explicit("gemma4".to_string()),
) {
Ok(_) => panic!("expected mixed Gemma4 parser selection to fail"),
Err(error) => error,
};
let Error::ParserInitialization { error, .. } = error else {
panic!("expected parser initialization error");
};
assert_eq!(
error.to_string(),
"`gemma4` only provides a unified parser; the same reasoning parser and tool parser should be specified together"
);
}
}
rust/src/chat/src/output/default/structural_tag.rs
+12 -12
View File
@@ -2,12 +2,12 @@
use thiserror_ext::AsReport;
use vllm_engine_core_client::protocol::{StructuredOutputBackend, StructuredOutputsParams};
use vllm_parser::tool::StructuralTagModel;
use xgrammar_structural_tag::{
FunctionDefinition, FunctionToolParam, ToolChoice as StructuralTagToolChoice, ToolParam,
build_structural_tag,
};
use crate::parser::tool::ToolParser;
use crate::request::{ChatRequest, ChatToolChoice};
use crate::{Error, Result as ChatResult};
@@ -15,9 +15,9 @@ use crate::{Error, Result as ChatResult};
/// support and the request's tool choice.
pub(super) fn apply_structural_tag_constraint(
request: &mut ChatRequest,
parser: &dyn ToolParser,
model: Option<StructuralTagModel>,
) -> ChatResult<()> {
let Some(model) = parser.structural_tag_model() else {
let Some(model) = model else {
return Ok(());
};
let Some(tool_choice) = structural_tag_tool_choice(request) else {
@@ -77,7 +77,7 @@ fn structural_tag_tool_choice(request: &ChatRequest) -> Option<StructuralTagTool
mod tests {
use serde_json::{Value, json};
use vllm_engine_core_client::protocol::{StructuredOutputBackend, StructuredOutputsParams};
use vllm_parser::tool::{Qwen3CoderToolParser, Tool};
use vllm_parser::tool::{Qwen3CoderToolParser, Tool, ToolParser};
use super::*;
@@ -134,7 +134,7 @@ mod tests {
let mut request = request(ChatToolChoice::Auto, vec![chat_tool("search", Some(true))]);
let parser = qwen3_coder_parser(&request.tools);
apply_structural_tag_constraint(&mut request, parser.as_ref())
apply_structural_tag_constraint(&mut request, parser.structural_tag_model())
.expect("structural tag should build");
let tag = structural_tag_value(&request);
@@ -147,7 +147,7 @@ mod tests {
let mut request = request(ChatToolChoice::Auto, vec![chat_tool("search", None)]);
let parser = qwen3_coder_parser(&request.tools);
apply_structural_tag_constraint(&mut request, parser.as_ref())
apply_structural_tag_constraint(&mut request, parser.structural_tag_model())
.expect("structural tag decision should succeed");
assert!(request.sampling_params.structured_outputs.is_none());
@@ -163,7 +163,7 @@ mod tests {
});
let parser = qwen3_coder_parser(&request.tools);
apply_structural_tag_constraint(&mut request, parser.as_ref())
apply_structural_tag_constraint(&mut request, parser.structural_tag_model())
.expect("structural tag should build");
let params = structured_outputs(&request);
@@ -179,7 +179,7 @@ mod tests {
let mut request = request(ChatToolChoice::Required, vec![chat_tool("search", None)]);
let parser = qwen3_coder_parser(&request.tools);
apply_structural_tag_constraint(&mut request, parser.as_ref())
apply_structural_tag_constraint(&mut request, parser.structural_tag_model())
.expect("structural tag should build");
let tag = structural_tag_value(&request);
@@ -197,7 +197,7 @@ mod tests {
});
let parser = qwen3_coder_parser(&request.tools);
apply_structural_tag_constraint(&mut request, parser.as_ref())
apply_structural_tag_constraint(&mut request, parser.structural_tag_model())
.expect("structural tag should build");
let params = structured_outputs(&request);
@@ -218,7 +218,7 @@ mod tests {
);
let parser = qwen3_coder_parser(&request.tools);
apply_structural_tag_constraint(&mut request, parser.as_ref())
apply_structural_tag_constraint(&mut request, parser.structural_tag_model())
.expect("structural tag should build");
let tag = structural_tag_value(&request).to_string();
@@ -231,7 +231,7 @@ mod tests {
let mut request = request(ChatToolChoice::None, vec![chat_tool("search", Some(true))]);
let parser = qwen3_coder_parser(&request.tools);
apply_structural_tag_constraint(&mut request, parser.as_ref())
apply_structural_tag_constraint(&mut request, parser.structural_tag_model())
.expect("structural tag decision should succeed");
assert!(request.sampling_params.structured_outputs.is_none());
@@ -247,7 +247,7 @@ mod tests {
});
let parser = qwen3_coder_parser(&request.tools);
apply_structural_tag_constraint(&mut request, parser.as_ref())
apply_structural_tag_constraint(&mut request, parser.structural_tag_model())
.expect("structural tag decision should succeed");
let params = structured_outputs(&request);
rust/src/chat/src/output/default/unified.rs
+1 -1
View File
@@ -317,7 +317,7 @@ mod tests {
where
Self: Sized + 'static,
{
unreachable!("ScriptedParser is constructed directly in tests")
Ok(Box::new(Self::new([])))
}
fn parse_into(
rust/src/chat/src/output/harmony/mod.rs
+2 -19
View File
@@ -4,16 +4,10 @@
//! `DecodedTextEvent` token IDs directly and lets the official `openai-harmony`
//! parser recover the structured assistant message shape at token granularity.
use std::sync::LazyLock;
use anyhow::Context;
use asynk_strim_attr::{TryYielder, try_stream};
use futures::StreamExt as _;
use openai_harmony::chat::{Content as HarmonyContent, Message as HarmonyMessage, Role};
use openai_harmony::{
HarmonyEncoding, HarmonyEncodingName, StreamableParser, load_harmony_encoding,
};
use thiserror_ext::AsReport;
use openai_harmony::{HarmonyEncoding, StreamableParser};
use vllm_text::output::DecodedTextEvent;
use crate::Result as ChatResult;
@@ -24,6 +18,7 @@ use crate::output::{
generate_tool_call_id,
};
use crate::parser::ParserSelection;
use crate::renderer::harmony::encoding::harmony_encoding;
use crate::request::ChatRequest;
/// Request-scoped Harmony output processor used for `model_type == "gpt_oss"`.
@@ -384,18 +379,6 @@ async fn harmony_assistant_event_stream(
Ok(())
}
/// Lazily load the shared GPT-OSS Harmony encoding once per process.
fn harmony_encoding() -> Result<&'static HarmonyEncoding> {
static ENCODING: LazyLock<anyhow::Result<HarmonyEncoding>> = LazyLock::new(|| {
load_harmony_encoding(HarmonyEncodingName::HarmonyGptOss)
.context("failed to load harmony encoding for gpt-oss")
});
ENCODING.as_ref().map_err(|error| Error::HarmonyOutputParsing {
error: error.to_report_string().into(),
})
}
fn harmony_output_parsing_error(
error: impl Into<Box<dyn std::error::Error + Send + Sync>>,
) -> Error {
rust/src/chat/src/parser/mod.rs
+1
View File
@@ -1,5 +1,6 @@
pub mod reasoning;
pub mod tool;
pub mod unified;
use std::collections::HashMap;
use std::convert::Infallible;
rust/src/chat/src/parser/reasoning/mod.rs
+21 -10
View File
@@ -1,13 +1,13 @@
//! Reasoning parser registration and selection boundary for `vllm-chat`.
use std::sync::LazyLock;
use std::sync::{Arc, LazyLock};
pub use vllm_parser::reasoning::{
CohereCmdReasoningParser, DeepSeekR1ReasoningParser, DeepSeekV3ReasoningParser,
DeepSeekV4ReasoningParser, Gemma4ReasoningParser, Glm45ReasoningParser, KimiK2ReasoningParser,
KimiReasoningParser, MiniMaxM2ReasoningParser, MiniMaxM3ReasoningParser,
NemotronV3ReasoningParser, Qwen3ReasoningParser, ReasoningDelta, ReasoningError,
ReasoningParser, SeedOssReasoningParser, Step3ReasoningParser, Step3p5ReasoningParser,
DeepSeekV4ReasoningParser, Glm45ReasoningParser, KimiK2ReasoningParser, KimiReasoningParser,
MiniMaxM2ReasoningParser, MiniMaxM3ReasoningParser, NemotronV3ReasoningParser,
Qwen3ReasoningParser, ReasoningDelta, ReasoningError, ReasoningParser, SeedOssReasoningParser,
Step3ReasoningParser, Step3p5ReasoningParser,
};
use vllm_tokenizer::DynTokenizer;
@@ -33,8 +33,9 @@ pub mod names {
}
/// Constructor signature for one registered reasoning parser implementation.
type ReasoningParserCreator =
fn(DynTokenizer) -> vllm_parser::reasoning::Result<Box<dyn ReasoningParser>>;
type ReasoningParserCreator = Arc<
dyn Fn(DynTokenizer) -> vllm_parser::reasoning::Result<Box<dyn ReasoningParser>> + Send + Sync,
>;
/// Registry and model matcher for reasoning parsers.
pub type ReasoningParserFactory = ParserFactory<ReasoningParserCreator>;
@@ -58,7 +59,7 @@ impl ReasoningParserFactory {
.register_parser::<DeepSeekR1ReasoningParser>(names::DEEPSEEK_R1)
.register_parser::<DeepSeekV3ReasoningParser>(names::DEEPSEEK_V3)
.register_parser::<DeepSeekV4ReasoningParser>(names::DEEPSEEK_V4)
.register_parser::<Gemma4ReasoningParser>(names::GEMMA4)
.register_unified_dummy(names::GEMMA4)
.register_parser::<Glm45ReasoningParser>(names::GLM45)
.register_parser::<KimiReasoningParser>(names::KIMI)
.register_parser::<KimiK2ReasoningParser>(names::KIMI_K2)
@@ -109,7 +110,17 @@ impl ReasoningParserFactory {
where
T: ReasoningParser + 'static,
{
self.register_creator(name, T::create)
self.register_creator(name, Arc::new(T::create))
}
/// Register one unified-only parser name in the split reasoning registry.
pub fn register_unified_dummy(&mut self, name: &str) -> &mut Self {
let name = name.to_string();
let registered_name = name.clone();
self.register_creator(
&registered_name,
Arc::new(move |_| Err(ReasoningError::DummyUnifiedParser { name: name.clone() })),
)
}
/// Construct a parser from an exact name.
@@ -124,7 +135,7 @@ impl ReasoningParserFactory {
available_names: self.list(),
})?;
creator(tokenizer).map_err(|error| crate::Error::ParserInitialization {
creator.as_ref()(tokenizer).map_err(|error| crate::Error::ParserInitialization {
kind: "reasoning",
name: name.to_string(),
error: error.into(),
rust/src/chat/src/parser/reasoning/tests.rs
+2
View File
@@ -35,11 +35,13 @@ fn factory_contains_and_lists_registered_parsers() {
assert!(factory.contains(names::SEED_OSS));
assert!(factory.contains(names::STEP3P5));
assert!(factory.contains(names::MINIMAX_M3));
assert!(factory.contains(names::GEMMA4));
assert!(factory.list().contains(&names::QWEN3.to_string()));
assert!(factory.list().contains(&names::DEEPSEEK_V4.to_string()));
assert!(factory.list().contains(&names::SEED_OSS.to_string()));
assert!(factory.list().contains(&names::STEP3P5.to_string()));
assert!(factory.list().contains(&names::MINIMAX_M3.to_string()));
assert!(factory.list().contains(&names::GEMMA4.to_string()));
}
#[test]
rust/src/chat/src/parser/tool/mod.rs
+20 -9
View File
@@ -1,13 +1,13 @@
//! Tool parser registration and selection boundary for `vllm-chat`.
use std::sync::LazyLock;
use std::sync::{Arc, LazyLock};
pub use vllm_parser::tool::{
DeepSeekV3ToolParser, DeepSeekV4ToolParser, DeepSeekV31ToolParser, DeepSeekV32ToolParser,
Gemma4ToolParser, Glm45MoeToolParser, Glm47MoeToolParser, Granite4ToolParser, HermesToolParser,
HyV3ToolParser, Internlm2ToolParser, KimiK2ToolParser, Llama3JsonToolParser,
MinimaxM2ToolParser, MinimaxM3ToolParser, MistralToolParser, Phi4MiniJsonToolParser,
Qwen3CoderToolParser, Qwen3XmlToolParser, ToolParser, ToolParserError,
Glm45MoeToolParser, Glm47MoeToolParser, Granite4ToolParser, HermesToolParser, HyV3ToolParser,
Internlm2ToolParser, KimiK2ToolParser, Llama3JsonToolParser, MinimaxM2ToolParser,
MinimaxM3ToolParser, MistralToolParser, Phi4MiniJsonToolParser, Qwen3CoderToolParser,
Qwen3XmlToolParser, ToolParser, ToolParserError,
};
use crate::parser::ParserFactory;
@@ -40,7 +40,8 @@ pub mod names {
}
/// Constructor signature for one registered tool parser implementation.
type ToolParserCreator = fn(&[ChatTool]) -> vllm_parser::tool::Result<Box<dyn ToolParser>>;
type ToolParserCreator =
Arc<dyn Fn(&[ChatTool]) -> vllm_parser::tool::Result<Box<dyn ToolParser>> + Send + Sync>;
/// Registry and model matcher for tool parsers.
pub type ToolParserFactory = ParserFactory<ToolParserCreator>;
@@ -65,7 +66,7 @@ impl ToolParserFactory {
.register_parser::<DeepSeekV4ToolParser>(names::DEEPSEEK_V4)
.register_parser::<Glm45MoeToolParser>(names::GLM45)
.register_parser::<Glm47MoeToolParser>(names::GLM47)
.register_parser::<Gemma4ToolParser>(names::GEMMA4)
.register_unified_dummy(names::GEMMA4)
.register_parser::<Granite4ToolParser>(names::GRANITE4)
.register_parser::<HermesToolParser>(names::HERMES)
.register_parser::<HyV3ToolParser>(names::HY_V3)
@@ -126,7 +127,17 @@ impl ToolParserFactory {
where
T: ToolParser + 'static,
{
self.register_creator(name, T::create)
self.register_creator(name, Arc::new(T::create))
}
/// Register one unified-only parser name in the split tool registry.
pub fn register_unified_dummy(&mut self, name: &str) -> &mut Self {
let name = name.to_string();
let registered_name = name.clone();
self.register_creator(
&registered_name,
Arc::new(move |_| Err(ToolParserError::DummyUnifiedParser { name: name.clone() })),
)
}
/// Construct a parser from an exact name.
@@ -137,7 +148,7 @@ impl ToolParserFactory {
available_names: self.list(),
})?;
creator(tools).map_err(|error| crate::Error::ParserInitialization {
creator.as_ref()(tools).map_err(|error| crate::Error::ParserInitialization {
kind: "tool",
name: name.to_string(),
error: error.into(),
rust/src/chat/src/parser/unified.rs
+124
View File
@@ -0,0 +1,124 @@
//! Unified parser registration and selection boundary for `vllm-chat`.
use std::sync::LazyLock;
pub use vllm_parser::unified::{Gemma4UnifiedParser, UnifiedParser};
use vllm_tokenizer::DynTokenizer;
use crate::parser::ParserFactory;
use crate::request::ChatTool;
/// Canonical public names for registered unified parsers.
pub mod names {
pub const GEMMA4: &str = "gemma4";
}
/// Constructor signature for one registered unified parser implementation.
type UnifiedParserCreator =
fn(&[ChatTool], DynTokenizer) -> vllm_parser::unified::Result<Box<dyn UnifiedParser>>;
/// Registry and model matcher for unified parsers.
pub type UnifiedParserFactory = ParserFactory<UnifiedParserCreator>;
impl UnifiedParserFactory {
/// Get the global unified parser factory with built-in registrations and
/// model mappings.
pub fn global() -> &'static Self {
static INSTANCE: LazyLock<UnifiedParserFactory> = LazyLock::new(UnifiedParserFactory::new);
&INSTANCE
}
/// Create the default registry with built-in parser names and model
/// mappings.
pub fn new() -> Self {
let mut factory = Self::default();
factory.register_parser::<Gemma4UnifiedParser>(names::GEMMA4);
factory
.register_pattern("gemma-4", names::GEMMA4)
.register_pattern("gemma4", names::GEMMA4);
factory
}
/// Register one parser type that exposes a static `create()` constructor.
pub fn register_parser<T>(&mut self, name: &str) -> &mut Self
where
T: UnifiedParser + 'static,
{
self.register_creator(name, T::create)
}
/// Construct a parser from an exact name.
pub fn create(
&self,
name: &str,
tools: &[ChatTool],
tokenizer: DynTokenizer,
) -> crate::Result<Box<dyn UnifiedParser>> {
let creator = self.creator(name).ok_or_else(|| crate::Error::ParserUnavailableByName {
kind: "unified",
name: name.to_string(),
available_names: self.list(),
})?;
creator(tools, tokenizer).map_err(|error| crate::Error::ParserInitialization {
kind: "unified",
name: name.to_string(),
error: error.into(),
})
}
}
#[cfg(test)]
mod tests {
use std::sync::Arc;
use vllm_tokenizer::Tokenizer;
use super::{UnifiedParserFactory, names};
struct FakeTokenizer;
impl Tokenizer for FakeTokenizer {
fn encode(
&self,
text: &str,
_add_special_tokens: bool,
) -> vllm_tokenizer::Result<Vec<u32>> {
Ok(text.chars().map(u32::from).collect())
}
fn decode(
&self,
token_ids: &[u32],
_skip_special_tokens: bool,
) -> vllm_tokenizer::Result<String> {
Ok(token_ids
.iter()
.map(|token_id| char::from_u32(*token_id).unwrap_or('\u{FFFD}'))
.collect())
}
fn token_to_id(&self, token: &str) -> Option<u32> {
match token {
"<|channel>" => Some(1),
"<channel|>" => Some(2),
_ => None,
}
}
}
#[test]
fn factory_registers_gemma4() {
let factory = UnifiedParserFactory::new();
assert!(factory.contains(names::GEMMA4));
assert_eq!(
factory.resolve_name_for_model("google/gemma-4-27b-it"),
Some(names::GEMMA4)
);
factory.create(names::GEMMA4, &[], Arc::new(FakeTokenizer)).unwrap();
}
}
rust/src/chat/src/renderer/deepseek_v32/tests.rs
+7 -145
View File
@@ -1,82 +1,18 @@
use std::fs;
use std::path::PathBuf;
use expect_test::{ExpectFile, expect, expect_file};
use serde::Deserialize;
use serde_json::{Value, json};
use thiserror_ext::AsReport;
use super::DeepSeekV32ChatRenderer;
use crate::error::Error;
use crate::event::{AssistantContentBlock, AssistantToolCall};
use crate::renderer::test_utils::{FixtureRequestOptions, fixture_chat_request};
use crate::request::{
ChatContentPart, ChatMessage, ChatRequest, ChatTool, ChatToolChoice, GenerationPromptMode,
};
use crate::{ChatRenderer, ChatRole};
#[derive(Debug, Deserialize)]
struct FixtureRequest {
#[serde(default)]
tools: Vec<FixtureTool>,
messages: Vec<FixtureMessage>,
}
#[derive(Debug, Deserialize)]
struct FixtureTool {
function: FixtureToolFunction,
}
#[derive(Debug, Deserialize)]
struct FixtureToolFunction {
name: String,
description: Option<String>,
parameters: Value,
#[serde(default)]
strict: Option<bool>,
}
#[derive(Debug, Deserialize)]
#[serde(tag = "role", rename_all = "snake_case")]
enum FixtureMessage {
System {
content: String,
},
Developer {
content: String,
#[serde(default)]
tools: Vec<FixtureTool>,
},
User {
content: String,
},
Assistant {
#[serde(default)]
content: String,
#[serde(default)]
reasoning_content: String,
#[serde(default)]
tool_calls: Vec<FixtureToolCall>,
},
Tool {
content: String,
#[serde(default)]
tool_call_id: Option<String>,
},
}
#[derive(Debug, Deserialize)]
struct FixtureToolCall {
#[serde(default)]
id: Option<String>,
function: FixtureToolCallFunction,
}
#[derive(Debug, Deserialize)]
struct FixtureToolCallFunction {
name: String,
arguments: String,
}
fn render_request(request: &ChatRequest) -> String {
DeepSeekV32ChatRenderer::new()
.render(request)
@@ -115,88 +51,14 @@ fn thinking_request(messages: Vec<ChatMessage>) -> ChatRequest {
}
fn fixture_request(input_name: &str) -> ChatRequest {
let fixture = fs::read_to_string(fixture_path(input_name)).unwrap();
let fixture: FixtureRequest = serde_json::from_str(&fixture).unwrap();
let mut request = ChatRequest {
request_id: "deepseek-v32-fixture".to_string(),
messages: fixture
.messages
.into_iter()
.enumerate()
.map(|(index, message)| match message {
FixtureMessage::System { content } => ChatMessage::system(content),
FixtureMessage::Developer { content, tools } => ChatMessage::developer(
content,
(!tools.is_empty()).then(|| to_chat_tools(&tools)),
),
FixtureMessage::User { content } => ChatMessage::user(content),
FixtureMessage::Assistant {
content,
reasoning_content,
tool_calls,
} => {
let mut blocks = Vec::new();
if !reasoning_content.is_empty() {
blocks.push(AssistantContentBlock::Reasoning {
text: reasoning_content,
});
}
if !content.is_empty() {
blocks.push(AssistantContentBlock::Text { text: content });
}
blocks.extend(tool_calls.into_iter().enumerate().map(
|(tool_index, tool_call)| {
AssistantContentBlock::ToolCall(AssistantToolCall {
id: tool_call.id.unwrap_or_else(|| {
format!("fixture-tool-call-{index}-{tool_index}")
}),
name: tool_call.function.name,
arguments: tool_call.function.arguments,
})
},
));
ChatMessage::assistant_blocks(blocks)
}
FixtureMessage::Tool {
content,
tool_call_id,
} => ChatMessage::tool_response(
content,
tool_call_id.unwrap_or_else(|| format!("fixture-tool-response-{index}")),
),
})
.collect(),
tools: to_chat_tools(&fixture.tools),
tool_choice: if fixture.tools.is_empty() {
ChatToolChoice::None
} else {
ChatToolChoice::Auto
},
..ChatRequest::for_test()
};
if matches!(
request.messages.last().map(ChatMessage::role),
Some(ChatRole::Assistant)
) {
request.chat_options.generation_prompt_mode = GenerationPromptMode::NoGenerationPrompt;
}
request
.chat_options
.template_kwargs
.insert("thinking".to_string(), Value::Bool(true));
request
fixture_chat_request(&fixture_path(input_name), deepseek_fixture_options())
}
fn to_chat_tools(tools: &[FixtureTool]) -> Vec<ChatTool> {
tools
.iter()
.map(|tool| ChatTool {
name: tool.function.name.clone(),
description: tool.function.description.clone(),
parameters: tool.function.parameters.clone(),
strict: tool.function.strict,
})
.collect()
fn deepseek_fixture_options() -> FixtureRequestOptions {
FixtureRequestOptions {
enable_thinking: true,
no_generation_prompt_when_last_assistant: true,
}
}
fn fixture_path(name: &str) -> PathBuf {
rust/src/chat/src/renderer/deepseek_v4/tests.rs
+9 -165
View File
@@ -1,95 +1,13 @@
use std::fs;
use std::path::PathBuf;
use expect_test::{ExpectFile, expect, expect_file};
use serde::Deserialize;
use serde_json::Value;
use super::DeepSeekV4ChatRenderer;
use crate::ChatRenderer;
use crate::event::{AssistantContentBlock, AssistantToolCall};
use crate::request::{
ChatMessage, ChatRequest, ChatTool, ChatToolChoice, GenerationPromptMode, ReasoningEffort,
};
use crate::{ChatRenderer, ChatRole};
#[derive(Debug, Deserialize)]
#[serde(untagged)]
enum FixtureFile {
WithTools(FixtureRequest),
MessagesOnly(Vec<FixtureMessage>),
}
#[derive(Debug, Deserialize)]
struct FixtureRequest {
#[serde(default)]
tools: Vec<FixtureTool>,
messages: Vec<FixtureMessage>,
}
impl FixtureFile {
fn into_parts(self) -> (Vec<FixtureTool>, Vec<FixtureMessage>) {
match self {
Self::WithTools(req) => (req.tools, req.messages),
Self::MessagesOnly(messages) => (Vec::new(), messages),
}
}
}
#[derive(Debug, Deserialize)]
struct FixtureTool {
function: FixtureToolFunction,
}
#[derive(Debug, Deserialize)]
struct FixtureToolFunction {
name: String,
description: Option<String>,
parameters: Value,
#[serde(default)]
strict: Option<bool>,
}
#[derive(Debug, Deserialize)]
#[serde(tag = "role", rename_all = "snake_case")]
enum FixtureMessage {
System {
content: String,
},
Developer {
content: String,
#[serde(default)]
tools: Vec<FixtureTool>,
},
User {
content: String,
},
Assistant {
#[serde(default)]
content: String,
#[serde(default)]
reasoning_content: String,
#[serde(default)]
tool_calls: Vec<FixtureToolCall>,
},
Tool {
content: String,
#[serde(default)]
tool_call_id: Option<String>,
},
}
#[derive(Debug, Deserialize)]
struct FixtureToolCall {
#[serde(default)]
id: Option<String>,
function: FixtureToolCallFunction,
}
#[derive(Debug, Deserialize)]
struct FixtureToolCallFunction {
name: String,
arguments: String,
}
use crate::renderer::test_utils::{FixtureRequestOptions, fixture_chat_request};
use crate::request::{ChatMessage, ChatRequest, GenerationPromptMode, ReasoningEffort};
fn render_request(request: &ChatRequest) -> String {
DeepSeekV4ChatRenderer::new()
@@ -101,88 +19,14 @@ fn render_request(request: &ChatRequest) -> String {
}
fn fixture_request(input_name: &str) -> ChatRequest {
let fixture = fs::read_to_string(fixture_path(input_name)).unwrap();
let fixture: FixtureFile = serde_json::from_str(&fixture).unwrap();
let (fixture_tools, fixture_messages) = fixture.into_parts();
let mut request = ChatRequest {
request_id: "deepseek-v4-fixture".to_string(),
messages: fixture_messages
.into_iter()
.enumerate()
.map(|(index, message)| match message {
FixtureMessage::System { content } => ChatMessage::system(content),
FixtureMessage::Developer { content, tools } => ChatMessage::developer(
content,
(!tools.is_empty()).then(|| to_chat_tools(&tools)),
),
FixtureMessage::User { content } => ChatMessage::user(content),
FixtureMessage::Assistant {
content,
reasoning_content,
tool_calls,
} => {
let mut blocks = Vec::new();
if !reasoning_content.is_empty() {
blocks.push(AssistantContentBlock::Reasoning {
text: reasoning_content,
});
}
if !content.is_empty() {
blocks.push(AssistantContentBlock::Text { text: content });
}
blocks.extend(tool_calls.into_iter().enumerate().map(
|(tool_index, tool_call)| {
AssistantContentBlock::ToolCall(AssistantToolCall {
id: tool_call.id.unwrap_or_else(|| {
format!("fixture-tool-call-{index}-{tool_index}")
}),
name: tool_call.function.name,
arguments: tool_call.function.arguments,
})
},
));
ChatMessage::assistant_blocks(blocks)
}
FixtureMessage::Tool {
content,
tool_call_id,
} => ChatMessage::tool_response(
content,
tool_call_id.unwrap_or_else(|| format!("fixture-tool-response-{index}")),
),
})
.collect(),
tools: to_chat_tools(&fixture_tools),
tool_choice: if fixture_tools.is_empty() {
ChatToolChoice::None
} else {
ChatToolChoice::Auto
},
..ChatRequest::for_test()
};
if matches!(
request.messages.last().map(ChatMessage::role),
Some(ChatRole::Assistant)
) {
request.chat_options.generation_prompt_mode = GenerationPromptMode::NoGenerationPrompt;
}
request
.chat_options
.template_kwargs
.insert("thinking".to_string(), Value::Bool(true));
request
fixture_chat_request(&fixture_path(input_name), deepseek_fixture_options())
}
fn to_chat_tools(tools: &[FixtureTool]) -> Vec<ChatTool> {
tools
.iter()
.map(|tool| ChatTool {
name: tool.function.name.clone(),
description: tool.function.description.clone(),
parameters: tool.function.parameters.clone(),
strict: tool.function.strict,
})
.collect()
fn deepseek_fixture_options() -> FixtureRequestOptions {
FixtureRequestOptions {
enable_thinking: true,
no_generation_prompt_when_last_assistant: true,
}
}
fn fixture_path(name: &str) -> PathBuf {
rust/src/chat/src/renderer/harmony/encoding.rs
+21
View File
@@ -0,0 +1,21 @@
//! Shared Harmony encoding helper for the GPT-OSS renderer and output parser.
use std::sync::LazyLock;
use anyhow::Context as _;
use openai_harmony::{HarmonyEncoding, HarmonyEncodingName, load_harmony_encoding};
use thiserror_ext::AsReport as _;
use crate::error::{Error, Result};
/// Lazily load the shared GPT-OSS Harmony encoding once per process.
pub(crate) fn harmony_encoding() -> Result<&'static HarmonyEncoding> {
static ENCODING: LazyLock<anyhow::Result<HarmonyEncoding>> = LazyLock::new(|| {
load_harmony_encoding(HarmonyEncodingName::HarmonyGptOss)
.context("failed to load harmony encoding for gpt-oss")
});
ENCODING.as_ref().map_err(|error| Error::HarmonyOutputParsing {
error: error.to_report_string().into(),
})
}
rust/src/chat/src/renderer/harmony/fixtures/assistant_history.json
+14
View File
@@ -0,0 +1,14 @@
{
"add_generation_prompt": false,
"messages": [
{
"role": "user",
"content": "What is 2 + 2?"
},
{
"role": "assistant",
"reasoning_content": "Need simple arithmetic.",
"content": "4"
}
]
}
rust/src/chat/src/renderer/harmony/fixtures/assistant_history.txt
+7
View File
@@ -0,0 +1,7 @@
<|start|>system<|message|>You are ChatGPT, a large language model trained by OpenAI.
Knowledge cutoff: 2024-06
Current date: 2025-06-28
Reasoning: medium
# Valid channels: analysis, commentary, final. Channel must be included for every message.<|end|><|start|>user<|message|>What is 2 + 2?<|end|><|start|>assistant<|channel|>final<|message|>4<|end|>
rust/src/chat/src/renderer/harmony/fixtures/developer_tools.json
+27
View File
@@ -0,0 +1,27 @@
[
{
"role": "developer",
"content": "Use tools when needed.",
"tools": [
{
"function": {
"name": "lookup",
"description": "Lookup a record.",
"parameters": {
"type": "object",
"properties": {
"id": {
"type": "string"
}
},
"required": ["id"]
}
}
}
]
},
{
"role": "user",
"content": "Find record abc."
}
]
rust/src/chat/src/renderer/harmony/fixtures/developer_tools.txt
+23
View File
@@ -0,0 +1,23 @@
<|start|>system<|message|>You are ChatGPT, a large language model trained by OpenAI.
Knowledge cutoff: 2024-06
Current date: 2025-06-28
Reasoning: medium
# Valid channels: analysis, commentary, final. Channel must be included for every message.
Calls to these tools must go to the commentary channel: 'functions'.<|end|><|start|>developer<|message|># Instructions
Use tools when needed.
# Tools
## functions
namespace functions {
// Lookup a record.
type lookup = (_: {
id: string,
}) => any;
} // namespace functions<|end|><|start|>user<|message|>Find record abc.<|end|><|start|>assistant
rust/src/chat/src/renderer/harmony/fixtures/drop_analysis.json
+15
View File
@@ -0,0 +1,15 @@
[
{
"role": "user",
"content": "What is 2 + 2?"
},
{
"role": "assistant",
"reasoning_content": "This should be dropped.",
"content": "4"
},
{
"role": "user",
"content": "What is 3 + 5?"
}
]
rust/src/chat/src/renderer/harmony/fixtures/drop_analysis.txt
+7
View File
@@ -0,0 +1,7 @@
<|start|>system<|message|>You are ChatGPT, a large language model trained by OpenAI.
Knowledge cutoff: 2024-06
Current date: 2025-06-28
Reasoning: medium
# Valid channels: analysis, commentary, final. Channel must be included for every message.<|end|><|start|>user<|message|>What is 2 + 2?<|end|><|start|>assistant<|channel|>final<|message|>4<|end|><|start|>user<|message|>What is 3 + 5?<|end|><|start|>assistant

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