Compare commits

...
Author SHA1 Message Date
mgoinandClaude Fable 5 62a0750892 [Model] Add LongCat-2.0 (sparse attention, MTP, n-gram fixes)
Support meituan-longcat/LongCat-2.0-FP8: the LongCat-Flash-Lite backbone
plus a DeepSeek-V3.2-style sparse attention indexer, reusing vLLM's
existing DSA infrastructure.

- Resolve the checkpoint's null model_type via a registered config class
  (arch LongcatCausalLM); alias the oe_* n-gram config fields.
- Wire the indexer into the dual-attention layer: only the first
  attention of each layer computes top-k indices (cli_factor), the
  second reuses them via the shared buffer. Adds a generic skip_topk
  override to DeepseekV2MLAAttention; the schedule lives in the model.
- Streaming-aware indexing: force the first index_init_tokens and last
  index_local_tokens into the top-k set (sparse_attn_indexer).
- MTP speculative decoding (one module iterated up to 3 draft steps,
  plain token embedding, FP8 indexer weights).
- Fix a latent LongCat weight-loading bug: the post-load
  mla_scale_q_lora/kv_lora fold breaks under incremental load_weights
  calls; fold at weight-load time instead.
- Fix two n-gram embedding bugs (also affect LongCat-Flash-Lite):
  rebuild the left-context from the accepted-token history so rejected
  draft tokens don't pollute it, and hash EOS-current positions with
  full look-back per the HF reference (unit test included).
- Add --reasoning-parser longcat (<longcat_think> tags) and docs rows.
- Add longcat model types to the DeepGEMM Blackwell blocklist (non-ue8m0
  scales; measured GSM8K-neutral on vLLM but matches SGLang's guard).

Co-authored-by: Claude Fable 5 <noreply@anthropic.com>
Signed-off-by: mgoin <mgoin64@gmail.com>
2026-07-11 02:12:49 +00:00
c227aaa3f8 [ROCm] Enable DeepSeek-V4 DSpark speculative decoding on AMD (MI350X / MI355X, gfx950) (#47419)
Signed-off-by: larryli2-amd <larryli2@amd.com>
Signed-off-by: larryli2-amd <Larry.Li@amd.com>
Co-authored-by: Andreas Karatzas <akaratza@amd.com>
2026-07-10 23:22:19 +08:00
Michael GoinandGitHub 08dfd68610 [Model] Add LongCat-Flash-Lite (n-gram embedding) (#47857)
Signed-off-by: mgoin <mgoin64@gmail.com>
2026-07-10 07:17:50 -07:00
Tyler Michael SmithandGitHub 978a6dfa3f [Build/CI] Build arm64 PR and postmerge image builds for Blackwell SM10x and SM110 (#48041)
Signed-off-by: Tyler Michael Smith <tlrmchlsmth@gmail.com>
2026-07-10 10:12:50 -04:00
85c09e9885 fix: correct load_weights track logic and enable weight integrity for… (#41811)
Signed-off-by: Yipeng Hu <i26268@metax-tech.com>
Signed-off-by: HuYiPeng <144002351+MynameFelix@users.noreply.github.com>
Signed-off-by: Isotr0py <Isotr0py@outlook.com>
Co-authored-by: Yipeng Hu <i26268@metax-tech.com>
Co-authored-by: Isotr0py <Isotr0py@outlook.com>
2026-07-10 14:08:20 +00:00
b12cca6a23 [Bugfix] Fix turboquant FP8 cast failure for BF16 models on Ampere GPUs (#39988)
Signed-off-by: Xu Zhou <xuzhou9417@163.com>
Co-authored-by: Xu Zhou <xuzhou9417@163.com>
Co-authored-by: Hoseung Kim <ghyutjik123@gmail.com>
2026-07-10 06:55:28 -07:00
Wentao YeandGitHub e257faf87d [Refactor] Remove unused rocm kernel combine_topk_swa_indices_ragged (#48158)
Signed-off-by: yewentao256 <zhyanwentao@126.com>
2026-07-10 09:33:07 -04:00
FAN YUCHENandGitHub fabec87f63 [Model] Migrate MistralLarge3ForCausalLM to AutoWeightsLoader (#48153)
Signed-off-by: Yuchen Fan <functionhx@gmail.com>
2026-07-10 12:27:58 +00:00
7614b88ebd [Bugfix][Spec Decode] Fix DFlash draft/target layer-count mismatch (#48113)
Signed-off-by: Nick Hill <nickhill123@gmail.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
2026-07-10 04:42:53 -07:00
Isotr0pyandGitHub 68ea76e780 [Misc] Remove dead code in ViT functionality test (#48220)
Signed-off-by: Isotr0py <Isotr0py@outlook.com>
2026-07-10 11:17:42 +00:00
c241c7a2b0 [Rust Frontend] Add roundtrip fixtures for more chat parsers (#47883)
Co-authored-by: Bugen Zhao <i@bugenzhao.com>
Signed-off-by: reidliu41 <reid201711@gmail.com>
Signed-off-by: Bugen Zhao <i@bugenzhao.com>
2026-07-10 10:03:58 +00:00
e23b19309b Deepstream video backend (#42424)
Signed-off-by: Viranjan Pagar <vpagar@nvidia.com>
Signed-off-by: Isotr0py <Isotr0py@outlook.com>
Signed-off-by: Isotr0py <mozf@mail2.sysu.edu.cn>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Co-authored-by: Isotr0py <mozf@mail2.sysu.edu.cn>
Co-authored-by: Isotr0py <Isotr0py@outlook.com>
Co-authored-by: Roger Wang <hey@rogerw.io>
2026-07-10 02:23:30 -07:00
f36284a8d2 [CI] Add TORCH_NIGHTLY=1 build mode (run full suite on torch nightly) (#47180)
Co-authored-by: Andrey Talman <atalman@users.noreply.github.com>
Co-authored-by: Kevin H. Luu <khluu000@gmail.com>
2026-07-10 01:38:35 -07:00
Mingfei GuoandGitHub 424df4f65d [Model][CI/Build] Cosmos3: enable registry tests and register Cosmos3-Super (#48211)
Signed-off-by: Mingfei Guo <1800012773@pku.edu.cn>
2026-07-10 16:35:15 +08:00
Bugen ZhaoandGitHub 074bdd0d99 [Rust Frontend] Integrate MM video support (#47959)
Signed-off-by: Bugen Zhao <i@bugenzhao.com>
2026-07-10 08:15:33 +00:00
wenjun liuGitHubjun,du <jun.du@intel.com>Kunshang Ji
216ee58780 Add XPU nightly and release image publishing to DockerHub (#48126)
Signed-off-by: wenjun.liu <wenjun.liu@intel.com>
Signed-off-by: jun,du <jun.du@intel.com>
Co-authored-by: jun,du <jun.du@intel.com>
Co-authored-by: Kunshang Ji <kunshang.ji@intel.com>
2026-07-10 00:56:00 -07:00
433f291195 [CI] Right-size test-area timeouts from nightly durations (#48186)
Signed-off-by: khluu <khluu000@gmail.com>
Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
2026-07-10 00:53:16 -07:00
Chaojun ZhangandGitHub 28eaf05d56 [XPU] Enable v1/sample tests on XPU CI (#44472)
Signed-off-by: Chaojun Zhang <chaojun.zhang@intel.com>
2026-07-10 15:40:51 +08:00
Jiangyun ZhuandGitHub 300e33797f [Perf] fuse more rmsnorm and all-reduce in qwen3.5 (#46998)
Signed-off-by: zjy0516 <riverclouds.zhu@qq.com>
2026-07-10 15:37:51 +08:00
5715fde12c [Feature][Parser] Support include_reasoning param for non-Harmony models (#44301)
Signed-off-by: Alberto Perdomo <aperdomo@redhat.com>
Co-authored-by: Chauncey <chaunceyjiang@gmail.com>
2026-07-10 15:34:02 +08:00
114 changed files with 4905 additions and 1586 deletions
+12
View File
@@ -177,6 +177,18 @@ BRANCH=$4
IMAGE_TAG=$5
IMAGE_TAG_LATEST=${6:-} # only used for main branch, optional
# When TORCH_NIGHTLY=1, build the base CI image against PyTorch nightly so the
# entire existing pipeline runs on nightly torch (CUDA/GPU lane only). Delegate
# to the dedicated nightly build (PYTORCH_NIGHTLY=1, CUDA 13.0) and tag it at the
# normal IMAGE_TAG that every test step already pulls -- no separate image tag,
# no duplicate "vLLM Against PyTorch Nightly" pipeline section.
if [[ "${TORCH_NIGHTLY:-0}" == "1" ]]; then
echo "--- :warning: TORCH_NIGHTLY=1 -- building base image on PyTorch nightly"
SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
exec "${SCRIPT_DIR}/image_build_torch_nightly.sh" \
"${REGISTRY}" "${REPO}" "${BUILDKITE_COMMIT}" "${BRANCH}" "${IMAGE_TAG}"
fi
# build config
TARGET="test-ci"
VLLM_BAKE_FILE_PATH="${VLLM_BAKE_FILE_PATH:-docker/docker-bake.hcl}"
+3 -2
View File
@@ -19,13 +19,14 @@ if docker manifest inspect "$IMAGE" >/dev/null 2>&1; then
echo "Image found"
else
echo "Image not found, proceeding with build..."
# build for arm64 GPU targets: Grace/GH200 (sm_90) and DGX Spark/GB10
# build for arm64 GPU targets: Grace/GH200 (sm_90),
# Blackwell/Thor (sm_100/sm_103/sm_110), and DGX Spark/GB10
# (sm_121, family-covered by 12.0 under CUDA 13)
docker build --file docker/Dockerfile \
--platform linux/arm64 \
--build-arg max_jobs=16 \
--build-arg nvcc_threads=4 \
--build-arg torch_cuda_arch_list="9.0 12.0" \
--build-arg torch_cuda_arch_list="9.0 10.0 11.0 12.0" \
--build-arg USE_SCCACHE=1 \
--build-arg buildkite_commit="$BUILDKITE_COMMIT" \
--tag "$IMAGE" \
+1 -3
View File
@@ -72,9 +72,7 @@ steps:
pytest -v -s v1/test_oracle.py &&
pytest -v -s v1/test_request.py &&
pytest -v -s v1/test_outputs.py &&
pytest -v -s v1/sample/test_topk_topp_sampler.py &&
pytest -v -s v1/sample/test_logprobs.py &&
pytest -v -s v1/sample/test_logprobs_e2e.py'
pytest -v -s v1/sample'
- label: Basic Models Tests (Initialization)
timeout_in_minutes: 60
+18
View File
@@ -848,6 +848,23 @@ steps:
DOCKER_BUILDKIT: "1"
S3_BUCKET: "vllm-wheels"
- label: "Publish nightly XPU image to DockerHub"
depends_on:
- create-manifest-xpu
if: build.env("NIGHTLY") == "1"
agents:
queue: small_cpu_queue_release
commands:
- "bash .buildkite/scripts/xpu/push-nightly-builds-xpu.sh"
- "bash .buildkite/scripts/cleanup-nightly-builds.sh nightly- vllm/vllm-openai-xpu"
plugins:
- docker-login#v3.0.0:
username: vllmbot
password-env: DOCKERHUB_TOKEN
env:
DOCKER_BUILDKIT: "1"
DOCKERHUB_USERNAME: "vllmbot"
- label: "Publish nightly ROCm image to DockerHub"
depends_on:
- build-rocm-release-image
@@ -878,6 +895,7 @@ steps:
- create-multi-arch-manifest-cuda-12-9
- create-multi-arch-manifest-ubuntu2404
- create-multi-arch-manifest-cuda-12-9-ubuntu2404
- create-manifest-xpu
- build-rocm-release-image
- input-release-version
# Wait for CPU builds if their block steps were unblocked, so publish
@@ -29,7 +29,11 @@ if python3 -c "import torch; assert torch.version.hip" 2>/dev/null; then
TORCH_INDEX_URL=""
fi
else
TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
if [ "${TORCH_NIGHTLY:-0}" = "1" ]; then
TORCH_INDEX_URL="https://download.pytorch.org/whl/nightly/cu130"
else
TORCH_INDEX_URL="https://download.pytorch.org/whl/cu130"
fi
fi
echo ">>> Using PyTorch index: ${TORCH_INDEX_URL:-PyPI default}"
@@ -130,6 +130,22 @@ docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${ROCM_BASE_CACHE_KEY}-rocm
docker push vllm/vllm-openai-rocm:latest-base
docker push vllm/vllm-openai-rocm:v${RELEASE_VERSION}-base
# ---- XPU ----
docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu
docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu vllm/vllm-openai-xpu:latest-x86_64
docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:${COMMIT}-x86_64-xpu vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64
docker push vllm/vllm-openai-xpu:latest-x86_64
docker push vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64
docker manifest rm vllm/vllm-openai-xpu:latest || true
docker manifest rm vllm/vllm-openai-xpu:v${RELEASE_VERSION} || true
docker manifest create vllm/vllm-openai-xpu:latest vllm/vllm-openai-xpu:latest-x86_64 --amend
docker manifest create vllm/vllm-openai-xpu:v${RELEASE_VERSION} vllm/vllm-openai-xpu:v${RELEASE_VERSION}-x86_64 --amend
docker manifest push vllm/vllm-openai-xpu:latest
docker manifest push vllm/vllm-openai-xpu:v${RELEASE_VERSION}
# ---- CPU ----
# CPU images are behind separate block steps and may not have been built.
# All-or-nothing: inspect both arches first, then either publish everything
+32 -11
View File
@@ -23,6 +23,7 @@ NC='\033[0m' # No Color
# Default configuration
PIPELINE="ci"
DRY_RUN=true
TORCH_NIGHTLY=false
usage() {
cat <<EOF
@@ -34,12 +35,14 @@ Sets RUN_ALL=1 and NIGHTLY=1 environment variables.
SAFETY: Dry-run by default. Use --execute to actually trigger a build.
Options:
--execute Actually trigger the build (default: dry-run)
--pipeline Buildkite pipeline slug (default: ${PIPELINE})
--commit Override commit SHA (default: current HEAD)
--branch Override branch name (default: current branch)
--message Custom build message (default: auto-generated)
--help Show this help message
--execute Actually trigger the build (default: dry-run)
--pipeline Buildkite pipeline slug (default: ${PIPELINE})
--commit Override commit SHA (default: current HEAD)
--branch Override branch name (default: current branch)
--message Custom build message (default: auto-generated)
--torch-nightly Also build and run the full suite against torch nightly
(sets TORCH_NIGHTLY=1)
--help Show this help message
Prerequisites:
- bk CLI installed: brew tap buildkite/buildkite && brew install buildkite/buildkite/bk
@@ -49,6 +52,7 @@ Examples:
$(basename "$0") # Dry-run, show what would happen
$(basename "$0") --execute # Actually trigger the build
$(basename "$0") --pipeline ci-shadow # Dry-run with different pipeline
$(basename "$0") --torch-nightly # Dry-run a full torch-nightly run
EOF
exit 1
}
@@ -96,6 +100,10 @@ while [[ $# -gt 0 ]]; do
MESSAGE="$2"
shift 2
;;
--torch-nightly)
TORCH_NIGHTLY=true
shift
;;
--help|-h)
usage
;;
@@ -171,11 +179,17 @@ if [[ $(echo "$REMOTE_BRANCHES" | wc -l) -gt 5 ]]; then
fi
echo ""
# Environment variables passed to the build.
BUILD_ENV=("RUN_ALL=1" "NIGHTLY=1")
if [[ "$TORCH_NIGHTLY" == true ]]; then
BUILD_ENV+=("TORCH_NIGHTLY=1")
fi
log_info "Pipeline: ${PIPELINE}"
log_info "Branch: ${BRANCH}"
log_info "Commit: ${COMMIT}"
log_info "Message: ${MESSAGE}"
log_info "Environment: RUN_ALL=1, NIGHTLY=1"
log_info "Environment: ${BUILD_ENV[*]}"
echo ""
# Build the command
@@ -187,9 +201,10 @@ CMD=(bk build create
--commit "${COMMIT}"
--branch "${BRANCH}"
--message "${MESSAGE}"
--env "RUN_ALL=1"
--env "NIGHTLY=1"
)
for env_var in "${BUILD_ENV[@]}"; do
CMD+=(--env "${env_var}")
done
if [[ "$DRY_RUN" == true ]]; then
echo "=========================================="
@@ -210,8 +225,14 @@ if [[ "$DRY_RUN" == true ]]; then
echo " --commit '$(escape_for_shell "${COMMIT}")' \\"
echo " --branch '$(escape_for_shell "${BRANCH}")' \\"
echo " --message '$(escape_for_shell "${MESSAGE}")' \\"
echo " --env 'RUN_ALL=1' \\"
echo " --env 'NIGHTLY=1'"
last_idx=$(( ${#BUILD_ENV[@]} - 1 ))
for i in "${!BUILD_ENV[@]}"; do
if [[ $i -eq $last_idx ]]; then
echo " --env '$(escape_for_shell "${BUILD_ENV[$i]}")'"
else
echo " --env '$(escape_for_shell "${BUILD_ENV[$i]}")' \\"
fi
done
echo ""
echo "=========================================="
echo -e "${YELLOW}To actually trigger this build, run:${NC}"
@@ -0,0 +1,21 @@
#!/bin/bash
set -ex
ORIG_TAG_NAME="$BUILDKITE_COMMIT"
REPO="vllm/vllm-openai-xpu"
echo "Pushing original XPU tag ${ORIG_TAG_NAME}-xpu to nightly tags in ${REPO}"
aws ecr-public get-login-password --region us-east-1 | docker login --username AWS --password-stdin public.ecr.aws/q9t5s3a7
docker pull public.ecr.aws/q9t5s3a7/vllm-release-repo:"$ORIG_TAG_NAME"-x86_64-xpu
docker tag public.ecr.aws/q9t5s3a7/vllm-release-repo:"$ORIG_TAG_NAME"-x86_64-xpu ${REPO}:nightly-x86_64
docker push ${REPO}:nightly-x86_64
docker manifest rm ${REPO}:nightly || true
docker manifest rm ${REPO}:nightly-"$BUILDKITE_COMMIT" || true
docker manifest create ${REPO}:nightly ${REPO}:nightly-x86_64 --amend
docker manifest create ${REPO}:nightly-"$BUILDKITE_COMMIT" ${REPO}:nightly-x86_64 --amend
docker manifest push ${REPO}:nightly
docker manifest push ${REPO}:nightly-"$BUILDKITE_COMMIT"
+3 -3
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: V1 attention (H100-MI300)
key: v1-attention-h100-mi300
timeout_in_minutes: 30
timeout_in_minutes: 85
device: h100
source_file_dependencies:
- vllm/config/attention.py
@@ -16,7 +16,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 70
timeout_in_minutes: 95
depends_on:
- image-build-amd
source_file_dependencies:
@@ -30,7 +30,7 @@ steps:
- label: V1 attention (B200)
key: v1-attention-b200
timeout_in_minutes: 30
timeout_in_minutes: 80
device: b200-k8s
source_file_dependencies:
- vllm/config/attention.py
+2 -2
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Basic Correctness
key: basic-correctness
timeout_in_minutes: 30
timeout_in_minutes: 45
device: h200_18gb
source_file_dependencies:
- vllm/
@@ -19,6 +19,6 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 50
timeout_in_minutes: 70
depends_on:
- image-build-amd
+2 -2
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Benchmarks CLI Test
key: benchmarks-cli-test
timeout_in_minutes: 20
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/
@@ -23,7 +23,7 @@ steps:
num_gpus: 2
optional: true
working_dir: "/vllm-workspace/"
timeout_in_minutes: 10
timeout_in_minutes: 20
source_file_dependencies:
- benchmarks/attention_benchmarks/
- vllm/v1/attention/
+11 -11
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Sequence Parallel Correctness Tests (2 GPUs)
key: sequence-parallel-correctness-tests-2-gpus
timeout_in_minutes: 50
timeout_in_minutes: 80
working_dir: "/vllm-workspace/"
num_devices: 2
source_file_dependencies:
@@ -19,7 +19,7 @@ steps:
- label: Sequence Parallel Correctness Tests (2xH100)
key: sequence-parallel-correctness-tests-2xh100
timeout_in_minutes: 50
timeout_in_minutes: 75
working_dir: "/vllm-workspace/"
device: h100
optional: true
@@ -30,7 +30,7 @@ steps:
- label: AsyncTP Correctness Tests (2xH100)
key: asynctp-correctness-tests-2xh100
timeout_in_minutes: 50
timeout_in_minutes: 30
working_dir: "/vllm-workspace/"
device: h100
optional: true
@@ -41,7 +41,7 @@ steps:
- label: AsyncTP Correctness Tests (B200)
key: asynctp-correctness-tests-b200
timeout_in_minutes: 50
timeout_in_minutes: 30
working_dir: "/vllm-workspace/"
device: b200-k8s
optional: true
@@ -52,7 +52,7 @@ steps:
- label: Distributed Compile Unit Tests (2xH100)
key: distributed-compile-unit-tests-2xh100
timeout_in_minutes: 20
timeout_in_minutes: 45
working_dir: "/vllm-workspace/"
device: h100
num_devices: 2
@@ -66,7 +66,7 @@ steps:
- label: Fusion and Compile Unit Tests (2xB200)
key: fusion-and-compile-unit-tests-2xb200
timeout_in_minutes: 20
timeout_in_minutes: 30
working_dir: "/vllm-workspace/"
device: b200-k8s
source_file_dependencies:
@@ -96,7 +96,7 @@ steps:
- label: Fusion E2E Quick (H100)
key: fusion-e2e-quick-h100
timeout_in_minutes: 15
timeout_in_minutes: 25
working_dir: "/vllm-workspace/"
device: h100
num_devices: 1
@@ -115,7 +115,7 @@ steps:
- label: Fusion E2E Config Sweep (H100)
key: fusion-e2e-config-sweep-h100
timeout_in_minutes: 30
timeout_in_minutes: 25
working_dir: "/vllm-workspace/"
device: h100
num_devices: 1
@@ -149,7 +149,7 @@ steps:
- label: Fusion E2E TP2 Quick (H100)
key: fusion-e2e-tp2-quick-h100
timeout_in_minutes: 20
timeout_in_minutes: 35
working_dir: "/vllm-workspace/"
device: h100
num_devices: 2
@@ -167,7 +167,7 @@ steps:
- label: Fusion E2E TP2 AR-RMS Config Sweep (H100)
key: fusion-e2e-tp2-ar-rms-config-sweep-h100
timeout_in_minutes: 40
timeout_in_minutes: 30
working_dir: "/vllm-workspace/"
device: h100
num_devices: 2
@@ -207,7 +207,7 @@ steps:
- label: Fusion E2E TP2 (B200)
key: fusion-e2e-tp2-b200
timeout_in_minutes: 20
timeout_in_minutes: 45
working_dir: "/vllm-workspace/"
device: b200-k8s
num_devices: 2
+2 -2
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Platform Tests
key: platform-tests
timeout_in_minutes: 15
timeout_in_minutes: 20
device: h200_18gb
source_file_dependencies:
- vllm/envs.py
@@ -19,7 +19,7 @@ steps:
- label: Cudagraph
key: cudagraph
timeout_in_minutes: 20
timeout_in_minutes: 30
source_file_dependencies:
- tests/v1/cudagraph
- vllm/v1/cudagraph_dispatcher.py
+12 -12
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Distributed NixlConnector PD accuracy (4 GPUs)
key: distributed-nixlconnector-pd-accuracy-4-gpus
timeout_in_minutes: 30
timeout_in_minutes: 55
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -16,7 +16,7 @@ steps:
mirror:
amd:
device: mi300_4
timeout_in_minutes: 110
timeout_in_minutes: 85
depends_on:
- image-build-amd
source_file_dependencies:
@@ -29,7 +29,7 @@ steps:
- label: Distributed FlashInfer NixlConnector PD accuracy (4 GPUs)
key: distributed-flashinfer-nixlconnector-pd-accuracy-4-gpus
timeout_in_minutes: 30
timeout_in_minutes: 55
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -66,7 +66,7 @@ steps:
mirror:
amd:
device: mi300_4
timeout_in_minutes: 50
timeout_in_minutes: 60
depends_on:
- image-build-amd
source_file_dependencies:
@@ -79,7 +79,7 @@ steps:
- label: CrossLayer KV layout Distributed NixlConnector PD accuracy tests (4 GPUs)
key: crosslayer-kv-layout-distributed-nixlconnector-pd-accuracy-tests-4-gpus
timeout_in_minutes: 30
timeout_in_minutes: 55
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -91,7 +91,7 @@ steps:
mirror:
amd:
device: mi300_4
timeout_in_minutes: 110
timeout_in_minutes: 85
depends_on:
- image-build-amd
source_file_dependencies:
@@ -104,7 +104,7 @@ steps:
- label: Hybrid SSM NixlConnector PD accuracy tests (4 GPUs)
key: hybrid-ssm-nixlconnector-pd-accuracy-tests-4-gpus
timeout_in_minutes: 25
timeout_in_minutes: 60
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -116,7 +116,7 @@ steps:
mirror:
amd:
device: mi300_4
timeout_in_minutes: 60
timeout_in_minutes: 80
depends_on:
- image-build-amd
source_file_dependencies:
@@ -143,7 +143,7 @@ steps:
- label: MultiConnector (Nixl+Offloading) PD accuracy (2 GPUs)
key: multiconnector-nixl-offloading-pd-accuracy-2-gpus
timeout_in_minutes: 30
timeout_in_minutes: 40
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -158,7 +158,7 @@ steps:
- label: NixlConnector PD + Spec Decode acceptance (2 GPUs)
key: nixlconnector-pd-spec-decode-acceptance-2-gpus
timeout_in_minutes: 30
timeout_in_minutes: 45
device: a100
working_dir: "/vllm-workspace/tests"
num_devices: 2
@@ -172,7 +172,7 @@ steps:
mirror:
amd:
device: mi300_2
timeout_in_minutes: 60
timeout_in_minutes: 70
depends_on:
- image-build-amd
source_file_dependencies:
@@ -186,7 +186,7 @@ steps:
- label: MultiConnector (Nixl+Offloading) PD edge cases (2 GPUs)
key: multiconnector-nixl-offloading-pd-edge-cases-2-gpus
timeout_in_minutes: 30
timeout_in_minutes: 25
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
+10 -10
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Distributed Comm Ops
key: distributed-comm-ops
timeout_in_minutes: 20
timeout_in_minutes: 25
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -18,7 +18,7 @@ steps:
- label: Distributed DP Tests (2 GPUs)
key: distributed-dp-tests-2-gpus
timeout_in_minutes: 20
timeout_in_minutes: 35
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -55,7 +55,7 @@ steps:
- label: Distributed Compile + RPC Tests (2 GPUs)
key: distributed-compile-rpc-tests-2-gpus
timeout_in_minutes: 20
timeout_in_minutes: 65
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -78,7 +78,7 @@ steps:
- label: Distributed Torchrun + Shutdown Tests (2 GPUs)
key: distributed-torchrun-shutdown-tests-2-gpus
timeout_in_minutes: 20
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -133,7 +133,7 @@ steps:
- label: Distributed DP Tests (4 GPUs)
key: distributed-dp-tests-4-gpus
timeout_in_minutes: 30
timeout_in_minutes: 45
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -154,7 +154,7 @@ steps:
- label: Distributed Compile + Comm (4 GPUs)
key: distributed-compile-comm-4-gpus
timeout_in_minutes: 30
timeout_in_minutes: 70
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -176,7 +176,7 @@ steps:
- label: Distributed Tests (8xH100)
key: distributed-tests-8xh100
timeout_in_minutes: 10
timeout_in_minutes: 20
device: h100
num_devices: 8
working_dir: "/vllm-workspace/tests"
@@ -212,7 +212,7 @@ steps:
- label: Distributed Tests (2xH100-2xMI300)
key: distributed-tests-2xh100-2xmi300
timeout_in_minutes: 15
timeout_in_minutes: 30
device: h100
optional: true
working_dir: "/vllm-workspace/"
@@ -259,7 +259,7 @@ steps:
- label: Pipeline + Context Parallelism (4 GPUs)
key: pipeline-context-parallelism-4-gpus
timeout_in_minutes: 60
timeout_in_minutes: 55
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -274,7 +274,7 @@ steps:
- label: RayExecutorV2 (4 GPUs)
key: rayexecutorv2-4-gpus
timeout_in_minutes: 60
timeout_in_minutes: 45
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
+1 -1
View File
@@ -3,7 +3,7 @@ depends_on:
- image-build-cpu
steps:
- label: Docker Build Metadata
timeout_in_minutes: 10
timeout_in_minutes: 20
device: cpu-small
source_file_dependencies:
- .buildkite/release-pipeline.yaml
+5 -5
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: DeepSeek V2-Lite Sync EPLB Accuracy (4xH100)
key: deepseek-v2-lite-sync-eplb-accuracy-4xh100
timeout_in_minutes: 60
timeout_in_minutes: 25
device: h100
optional: true
num_devices: 4
@@ -14,7 +14,7 @@ steps:
- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (4xH100)
key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-4xh100
timeout_in_minutes: 60
timeout_in_minutes: 25
device: h100
optional: true
num_devices: 4
@@ -24,7 +24,7 @@ steps:
- label: Qwen3-30B-A3B-FP8-block Sync EPLB Accuracy (2xB200)
key: qwen3-30b-a3b-fp8-block-sync-eplb-accuracy-2xb200
timeout_in_minutes: 60
timeout_in_minutes: 20
device: b200-k8s
optional: true
num_devices: 2
@@ -34,7 +34,7 @@ steps:
- label: Qwen3-30B-A3B-FP8 DP4 Async EPLB Accuracy
key: qwen3-30b-a3b-fp8-dp4-async-eplb-accuracy
timeout_in_minutes: 60
timeout_in_minutes: 25
device: h100
optional: true
num_devices: 4
@@ -44,7 +44,7 @@ steps:
- label: DeepSeek V2-Lite Prefetch Offload Accuracy (H100)
key: deepseek-v2-lite-prefetch-offload-accuracy-h100
timeout_in_minutes: 60
timeout_in_minutes: 20
device: h100
optional: true
num_devices: 1
+11 -11
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Engine
key: engine
timeout_in_minutes: 15
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/compilation/
@@ -29,13 +29,13 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 60
timeout_in_minutes: 50
depends_on:
- image-build-amd
- label: Engine (1 GPU)
key: engine-1-gpu
timeout_in_minutes: 30
timeout_in_minutes: 45
source_file_dependencies:
- vllm/v1/engine/
- tests/v1/engine/
@@ -45,13 +45,13 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 40
timeout_in_minutes: 55
depends_on:
- image-build-amd
- label: e2e Scheduling (1 GPU)
key: e2e-scheduling-1-gpu
timeout_in_minutes: 30
timeout_in_minutes: 35
device: h200_18gb
source_file_dependencies:
- vllm/v1/
@@ -61,14 +61,14 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 60
timeout_in_minutes: 70
depends_on:
- image-build-amd
- label: e2e Core (1 GPU)
device: h200_35gb
key: e2e-core-1-gpu
timeout_in_minutes: 30
timeout_in_minutes: 40
source_file_dependencies:
- vllm/v1/
- tests/v1/e2e/general/
@@ -77,7 +77,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 35
timeout_in_minutes: 60
depends_on:
- image-build-amd
source_file_dependencies:
@@ -87,7 +87,7 @@ steps:
- label: V1 e2e (2 GPUs)
key: v1-e2e-2-gpus
timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
timeout_in_minutes: 25 # TODO: Fix timeout after we have more confidence in the test stability
optional: true
num_devices: 2
source_file_dependencies:
@@ -120,7 +120,7 @@ steps:
- label: V1 e2e (4 GPUs)
key: v1-e2e-4-gpus
timeout_in_minutes: 60 # TODO: Fix timeout after we have more confidence in the test stability
timeout_in_minutes: 20 # TODO: Fix timeout after we have more confidence in the test stability
optional: true
num_devices: 4
source_file_dependencies:
@@ -148,7 +148,7 @@ steps:
- label: V1 e2e (4xH100)
key: v1-e2e-4xh100
timeout_in_minutes: 60
timeout_in_minutes: 35
device: h100
num_devices: 4
optional: true
+10 -10
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Entrypoints Unit Tests
key: entrypoints-unit-tests
timeout_in_minutes: 10
timeout_in_minutes: 25
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/entrypoints
@@ -16,7 +16,7 @@ steps:
- label: Entrypoints Integration (LLM)
key: entrypoints-integration-llm
timeout_in_minutes: 40
timeout_in_minutes: 60
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -37,7 +37,7 @@ steps:
- label: Entrypoints Integration (API Server)
key: entrypoints-integration-api-server
device: h200_35gb
timeout_in_minutes: 130
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -56,7 +56,7 @@ steps:
- label: Entrypoints Integration (API Server OpenAI - Part 1)
key: entrypoints-integration-api-server-openai-part-1
timeout_in_minutes: 50
timeout_in_minutes: 45
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -68,13 +68,13 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 80
timeout_in_minutes: 65
depends_on:
- image-build-amd
- label: Entrypoints Integration (API Server OpenAI - Part 2)
key: entrypoints-integration-api-server-openai-part-2
timeout_in_minutes: 50
timeout_in_minutes: 45
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -109,7 +109,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 60
timeout_in_minutes: 65
depends_on:
- image-build-amd
@@ -126,7 +126,7 @@ steps:
- label: Entrypoints Integration (Speech to Text)
device: h200_35gb
key: entrypoints-integration-speech_to_text
timeout_in_minutes: 50
timeout_in_minutes: 45
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -138,7 +138,7 @@ steps:
- label: Entrypoints Integration (Multimodal)
device: h200_35gb
key: entrypoints-integration-multimodal
timeout_in_minutes: 50
timeout_in_minutes: 45
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- vllm/
@@ -160,7 +160,7 @@ steps:
- label: OpenAI API Correctness
key: openai-api-correctness
timeout_in_minutes: 30
timeout_in_minutes: 20
device: h200_18gb
source_file_dependencies:
- csrc/
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: EPLB Algorithm
key: eplb-algorithm
timeout_in_minutes: 15
timeout_in_minutes: 20
device: h200_18gb
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -27,7 +27,7 @@ steps:
- label: EPLB Execution # 17min
key: eplb-execution
timeout_in_minutes: 27
timeout_in_minutes: 25
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
@@ -39,7 +39,7 @@ steps:
- label: Elastic EP Scaling Test
key: elastic-ep-scaling-test
timeout_in_minutes: 20
timeout_in_minutes: 30
device: h100
working_dir: "/vllm-workspace/tests"
num_devices: 4
+19 -19
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: vLLM IR Tests
key: vllm-ir-tests
timeout_in_minutes: 10
timeout_in_minutes: 35
device: h200_18gb
working_dir: "/vllm-workspace/"
source_file_dependencies:
@@ -16,7 +16,7 @@ steps:
- label: Kernels Core Operation Test
key: kernels-core-operation-test
timeout_in_minutes: 75
timeout_in_minutes: 120
source_file_dependencies:
- csrc/
- tests/kernels/core
@@ -27,7 +27,7 @@ steps:
- label: Kernels MiniMax Reduce RMS Test (2 GPUs)
key: kernels-minimax-reduce-rms-test-2-gpus
timeout_in_minutes: 15
timeout_in_minutes: 20
num_devices: 2
device: h100
source_file_dependencies:
@@ -41,7 +41,7 @@ steps:
- label: Deepseek V4 Kernel Test (H100)
key: deepseek-v4-kernel-test-h100
timeout_in_minutes: 15
timeout_in_minutes: 30
device: h100
source_file_dependencies:
- csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
@@ -54,7 +54,7 @@ steps:
- label: Deepseek V4 Kernel Test (B200)
key: deepseek-v4-kernel-test-b200
timeout_in_minutes: 15
timeout_in_minutes: 20
device: b200-k8s
source_file_dependencies:
- csrc/fused_deepseek_v4_qnorm_rope_kv_insert_kernel.cu
@@ -65,7 +65,7 @@ steps:
- label: Kernels Attention Test %N
key: kernels-attention-test
timeout_in_minutes: 35
timeout_in_minutes: 65
source_file_dependencies:
- csrc/attention/
- vllm/v1/attention
@@ -79,7 +79,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 55
timeout_in_minutes: 90
depends_on:
- image-build-amd
source_file_dependencies:
@@ -106,7 +106,7 @@ steps:
- label: Kernels Quantization Test %N
key: kernels-quantization-test
timeout_in_minutes: 90
timeout_in_minutes: 60
source_file_dependencies:
- csrc/quantization/
- vllm/model_executor/layers/quantization
@@ -131,7 +131,7 @@ steps:
- label: Kernels MoE Test %N
key: kernels-moe-test
timeout_in_minutes: 25
timeout_in_minutes: 50
source_file_dependencies:
- csrc/quantization/cutlass_w8a8/moe/
- csrc/moe/
@@ -147,7 +147,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 50
timeout_in_minutes: 65
source_file_dependencies:
- csrc/quantization/cutlass_w8a8/moe/
- csrc/moe/
@@ -163,7 +163,7 @@ steps:
- label: Kernels Mamba Test
key: kernels-mamba-test
timeout_in_minutes: 45
timeout_in_minutes: 40
source_file_dependencies:
- csrc/mamba/
- tests/kernels/mamba
@@ -172,7 +172,7 @@ steps:
- pytest -v -s kernels/mamba
- label: Kernels KDA Test
timeout_in_minutes: 20
timeout_in_minutes: 25
device: h200_18gb
source_file_dependencies:
- vllm/model_executor/layers/fla/ops/kda.py
@@ -184,7 +184,7 @@ steps:
- label: Kernels DeepGEMM Test (H100)
key: kernels-deepgemm-test-h100
timeout_in_minutes: 45
timeout_in_minutes: 35
device: h100
num_devices: 1
source_file_dependencies:
@@ -211,7 +211,7 @@ steps:
- label: Kernels (B200)
key: kernels-b200
timeout_in_minutes: 30
timeout_in_minutes: 80
working_dir: "/vllm-workspace/"
device: b200-k8s
# optional: true
@@ -264,7 +264,7 @@ steps:
- label: Kernels Helion Test
key: kernels-helion-test
timeout_in_minutes: 30
timeout_in_minutes: 115
device: h100
source_file_dependencies:
- vllm/utils/import_utils.py
@@ -276,7 +276,7 @@ steps:
- label: Kernels FP8 MoE Test (1xH100)
key: kernels-fp8-moe-test-1xh100
timeout_in_minutes: 90
timeout_in_minutes: 40
device: h100
num_devices: 1
optional: true
@@ -293,7 +293,7 @@ steps:
- label: Kernels FP8 MoE Test (2xH100)
key: kernels-fp8-moe-test-2xh100
timeout_in_minutes: 90
timeout_in_minutes: 45
device: h100
num_devices: 2
optional: true
@@ -303,7 +303,7 @@ steps:
- label: Kernels Fp4 MoE Test (B200)
key: kernels-fp4-moe-test-b200
timeout_in_minutes: 60
timeout_in_minutes: 25
device: b200-k8s
num_devices: 1
optional: true
@@ -316,7 +316,7 @@ steps:
- label: Kernels FusedMoE Layer Test (2 H100s)
key: kernels-fusedmoe-layer-test-2-h100s
timeout_in_minutes: 90
timeout_in_minutes: 30
device: h100
num_devices: 2
source_file_dependencies:
+19 -19
View File
@@ -5,7 +5,7 @@ steps:
- label: LM Eval Small Models
device: h200_35gb
key: lm-eval-small-models
timeout_in_minutes: 75
timeout_in_minutes: 45
source_file_dependencies:
- csrc/
- vllm/model_executor/layers/quantization
@@ -56,7 +56,7 @@ steps:
- label: LM Eval Small Models (1xB200)
key: lm-eval-small-models-1xb200
timeout_in_minutes: 120
timeout_in_minutes: 50
device: b200-k8s
optional: true
source_file_dependencies:
@@ -80,7 +80,7 @@ steps:
- label: LM Eval Large Models EP (2xB200)
key: lm-eval-large-models-ep-2xb200
timeout_in_minutes: 120
timeout_in_minutes: 60
device: b200-k8s
optional: true
num_devices: 2
@@ -92,7 +92,7 @@ steps:
- label: LM Eval Qwen3.5 Models (2xB200)
key: lm-eval-qwen3-5-models-2xb200
timeout_in_minutes: 120
timeout_in_minutes: 45
device: b200-k8s
optional: true
num_devices: 2
@@ -109,7 +109,7 @@ steps:
- label: LM Eval Large Models (8xH200)
key: lm-eval-large-models-8xh200
timeout_in_minutes: 60
timeout_in_minutes: 50
device: h200
optional: true
num_devices: 8
@@ -118,7 +118,7 @@ steps:
mirror:
amd:
device: mi300_8
timeout_in_minutes: 180
timeout_in_minutes: 60
depends_on:
- image-build-amd
commands:
@@ -152,7 +152,7 @@ steps:
- label: LM Eval Humming f16 (A100 - TEMPORARY)
key: lm-eval-humming-f16-a100
timeout_in_minutes: 120
timeout_in_minutes: 75
device: a100
optional: true
num_devices: 1
@@ -167,7 +167,7 @@ steps:
- label: LM Eval Humming Act int8 (A100 - TEMPORARY)
key: lm-eval-humming-act-a100
timeout_in_minutes: 120
timeout_in_minutes: 45
device: a100
optional: true
num_devices: 1
@@ -182,7 +182,7 @@ steps:
- label: LM Eval Humming f16 (H100 - TEMPORARY)
key: lm-eval-humming-f16-h100
timeout_in_minutes: 120
timeout_in_minutes: 70
device: h100
optional: true
num_devices: 1
@@ -197,7 +197,7 @@ steps:
- label: LM Eval Humming Act fp8/int8 (H100 - TEMPORARY)
key: lm-eval-humming-act-h100
timeout_in_minutes: 120
timeout_in_minutes: 70
device: h100
optional: true
num_devices: 1
@@ -213,7 +213,7 @@ steps:
- label: LM Eval Humming f16 (B200 - TEMPORARY)
key: lm-eval-humming-f16-b200
timeout_in_minutes: 120
timeout_in_minutes: 50
device: b200-k8s
optional: true
num_devices: 1
@@ -228,7 +228,7 @@ steps:
- label: LM Eval Humming Act fp8/int8 (B200 - TEMPORARY)
key: lm-eval-humming-act-b200
timeout_in_minutes: 120
timeout_in_minutes: 50
device: b200-k8s
optional: true
num_devices: 1
@@ -244,7 +244,7 @@ steps:
- label: LM Eval TurboQuant KV Cache
key: lm-eval-turboquant-kv-cache
timeout_in_minutes: 75
timeout_in_minutes: 55
device: h200_18gb
source_file_dependencies:
- vllm/model_executor/layers/quantization/turboquant/
@@ -256,7 +256,7 @@ steps:
- label: GPQA Eval (GPT-OSS) (2xH100)
key: gpqa-eval-gpt-oss-2xh100
timeout_in_minutes: 120
timeout_in_minutes: 35
device: h100
optional: true
num_devices: 2
@@ -270,7 +270,7 @@ steps:
- label: GPQA Eval (GPT-OSS) (2xB200)
key: gpqa-eval-gpt-oss-2xb200
timeout_in_minutes: 120
timeout_in_minutes: 30
device: b200-k8s
optional: true
num_devices: 2
@@ -284,7 +284,7 @@ steps:
- label: GPQA Eval (GPT-OSS) (DGX Spark)
key: gpqa-eval-gpt-oss-spark
timeout_in_minutes: 120
timeout_in_minutes: 35
device: dgx-spark
optional: true
num_devices: 1
@@ -313,7 +313,7 @@ steps:
- label: LM Eval KV-Offload (2xH100)
key: kv-offload-medium
timeout_in_minutes: 60
timeout_in_minutes: 30
device: h100
num_devices: 2
source_file_dependencies:
@@ -327,7 +327,7 @@ steps:
- label: LM Eval KV-Offload (4xH100)
key: kv-offload-large
timeout_in_minutes: 60
timeout_in_minutes: 40
device: h100
num_devices: 4
source_file_dependencies:
@@ -341,7 +341,7 @@ steps:
- label: MRCR Eval Small Models
device: h200_35gb
timeout_in_minutes: 30
timeout_in_minutes: 25
source_file_dependencies:
- tests/evals/mrcr/
commands:
+3 -3
View File
@@ -5,7 +5,7 @@ steps:
- label: LoRA %N
device: h200_35gb
key: lora
timeout_in_minutes: 30
timeout_in_minutes: 40
source_file_dependencies:
- vllm/lora
- tests/lora
@@ -16,7 +16,7 @@ steps:
amd:
device: mi325_1
working_dir: "/vllm-workspace/tests"
timeout_in_minutes: 60
timeout_in_minutes: 65
source_file_dependencies:
- vllm/lora
- tests/lora
@@ -27,7 +27,7 @@ steps:
- label: LoRA TP (Distributed)
key: lora-tp-distributed
timeout_in_minutes: 30
timeout_in_minutes: 60
num_devices: 4
source_file_dependencies:
- vllm/lora
+15 -15
View File
@@ -5,7 +5,7 @@ steps:
- label: V1 Spec Decode
device: h200_35gb
key: v1-spec-decode
timeout_in_minutes: 30
timeout_in_minutes: 40
source_file_dependencies:
- vllm/config/
- vllm/distributed/
@@ -24,13 +24,13 @@ steps:
mirror:
amd:
device: mi300_1
timeout_in_minutes: 65
timeout_in_minutes: 75
depends_on:
- image-build-amd
- label: V1 Sample + Logits
key: v1-sample-logits
timeout_in_minutes: 30
timeout_in_minutes: 45
device: h200_18gb
source_file_dependencies:
- vllm/config/
@@ -64,7 +64,7 @@ steps:
- label: V1 Core + KV + Metrics
key: v1-core-kv-metrics
timeout_in_minutes: 30
timeout_in_minutes: 60
source_file_dependencies:
- vllm/config/
- vllm/distributed/
@@ -108,7 +108,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 60
timeout_in_minutes: 75
depends_on:
- image-build-amd
@@ -172,7 +172,7 @@ steps:
- label: Regression
key: regression
timeout_in_minutes: 20
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/config/
@@ -195,7 +195,7 @@ steps:
- label: Examples
device: h200_35gb
key: examples
timeout_in_minutes: 45
timeout_in_minutes: 40
working_dir: "/vllm-workspace/examples"
source_file_dependencies:
- vllm/entrypoints
@@ -237,7 +237,7 @@ steps:
- label: Metrics, Tracing (2 GPUs)
key: metrics-tracing-2-gpus
timeout_in_minutes: 20
timeout_in_minutes: 25
num_devices: 2
source_file_dependencies:
- vllm/config/
@@ -281,7 +281,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 20
timeout_in_minutes: 45
depends_on:
- image-build-amd
source_file_dependencies:
@@ -292,7 +292,7 @@ steps:
- label: Async Engine, Inputs, Utils, Worker
device: h200_35gb
key: async-engine-inputs-utils-worker
timeout_in_minutes: 50
timeout_in_minutes: 25
source_file_dependencies:
- vllm/assets/
- vllm/config/
@@ -319,7 +319,7 @@ steps:
key: async-engine-inputs-utils-worker-config-cpu
depends_on:
- image-build-cpu
timeout_in_minutes: 30
timeout_in_minutes: 65
source_file_dependencies:
- vllm/assets/
- vllm/config/
@@ -381,7 +381,7 @@ steps:
- label: Batch Invariance (A100)
key: batch-invariance-a100
timeout_in_minutes: 30
timeout_in_minutes: 40
device: a100
source_file_dependencies:
- vllm/v1/attention
@@ -395,7 +395,7 @@ steps:
- label: Batch Invariance (H100)
key: batch-invariance-h100
timeout_in_minutes: 30
timeout_in_minutes: 40
device: h100
source_file_dependencies:
- vllm/v1/attention
@@ -411,7 +411,7 @@ steps:
- label: Batch Invariance (B200)
key: batch-invariance-b200
timeout_in_minutes: 30
timeout_in_minutes: 35
device: b200-k8s
source_file_dependencies:
- vllm/v1/attention
@@ -430,7 +430,7 @@ steps:
- label: Acceptance Length Test (Large Models) # optional
device: h200_35gb
key: acceptance-length-test-large-models
timeout_in_minutes: 25
timeout_in_minutes: 20
gpu: h100
optional: true
num_gpus: 1
+1 -1
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Model Executor
key: model-executor
timeout_in_minutes: 35
timeout_in_minutes: 45
source_file_dependencies:
- vllm/engine/arg_utils.py
- vllm/config/model.py
+4 -4
View File
@@ -5,7 +5,7 @@ steps:
- label: Model Runner V2 Core Tests
device: h200_35gb
key: model-runner-v2-core-tests
timeout_in_minutes: 45
timeout_in_minutes: 35
source_file_dependencies:
- vllm/v1/worker/gpu/
- vllm/v1/worker/gpu_worker.py
@@ -27,7 +27,7 @@ steps:
- label: Model Runner V2 Examples
device: h200_35gb
key: model-runner-v2-examples
timeout_in_minutes: 45
timeout_in_minutes: 35
working_dir: "/vllm-workspace/examples"
source_file_dependencies:
- vllm/v1/worker/gpu/
@@ -63,7 +63,7 @@ steps:
- label: Model Runner V2 Distributed (2 GPUs)
key: model-runner-v2-distributed-2-gpus
timeout_in_minutes: 45
timeout_in_minutes: 30
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
@@ -84,7 +84,7 @@ steps:
- label: Model Runner V2 Pipeline Parallelism (4 GPUs)
key: model-runner-v2-pipeline-parallelism-4-gpus
timeout_in_minutes: 60
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
num_devices: 4
source_file_dependencies:
+4 -4
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Basic Models Tests (Initialization)
key: basic-models-tests-initialization
timeout_in_minutes: 45
timeout_in_minutes: 25
device: h200_18gb
source_file_dependencies:
- vllm/
@@ -17,7 +17,7 @@ steps:
- label: Basic Models Tests (Extra Initialization) %N
device: h200_35gb
key: basic-models-tests-extra-initialization
timeout_in_minutes: 45
timeout_in_minutes: 100
source_file_dependencies:
- vllm/model_executor/models/
- tests/models/test_initialization.py
@@ -32,7 +32,7 @@ steps:
- label: Basic Models Tests (Other)
device: h200_35gb
key: basic-models-tests-other
timeout_in_minutes: 45
timeout_in_minutes: 35
source_file_dependencies:
- vllm/
- tests/models/test_terratorch.py
@@ -50,7 +50,7 @@ steps:
key: basic-models-test-other-cpu
depends_on:
- image-build-cpu
timeout_in_minutes: 10
timeout_in_minutes: 20
source_file_dependencies:
- vllm/
- tests/models/test_utils.py
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Distributed Model Tests (2 GPUs)
key: distributed-model-tests-2-gpus
timeout_in_minutes: 50
timeout_in_minutes: 60
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
+8 -8
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Language Models Tests (Standard)
key: language-models-tests-standard
timeout_in_minutes: 25
timeout_in_minutes: 30
device: h200_18gb
source_file_dependencies:
- vllm/
@@ -21,7 +21,7 @@ steps:
- label: Language Models Tests (Extra Standard) %N
key: language-models-tests-extra-standard
timeout_in_minutes: 45
timeout_in_minutes: 40
source_file_dependencies:
- vllm/model_executor/models/
- tests/models/language/pooling/test_embedding.py
@@ -52,7 +52,7 @@ steps:
- label: Language Models Tests (Hybrid) %N
key: language-models-tests-hybrid
timeout_in_minutes: 75
timeout_in_minutes: 65
source_file_dependencies:
- vllm/
- tests/models/language/generation
@@ -67,7 +67,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 90
timeout_in_minutes: 70
depends_on:
- image-build-amd
commands:
@@ -78,7 +78,7 @@ steps:
- label: Language Models Test (Extended Generation) # 80min
device: h200_35gb
key: language-models-test-extended-generation
timeout_in_minutes: 110
timeout_in_minutes: 65
optional: true
source_file_dependencies:
- vllm/
@@ -92,7 +92,7 @@ steps:
- label: Language Models Test (PPL)
key: language-models-test-ppl
timeout_in_minutes: 110
timeout_in_minutes: 30
device: h200_18gb
optional: true
source_file_dependencies:
@@ -104,7 +104,7 @@ steps:
- label: Language Models Test (Extended Pooling) # 36min
device: h200_35gb
key: language-models-test-extended-pooling
timeout_in_minutes: 50
timeout_in_minutes: 70
optional: true
source_file_dependencies:
- vllm/
@@ -120,7 +120,7 @@ steps:
- label: Language Models Test (MTEB)
key: language-models-test-mteb
timeout_in_minutes: 110
timeout_in_minutes: 45
device: h200_18gb
optional: true
source_file_dependencies:
+7 -7
View File
@@ -20,7 +20,7 @@ steps:
- label: "Multi-Modal Models (Standard) 2: qwen3 + gemma"
key: multi-modal-models-standard-2-qwen3-gemma
timeout_in_minutes: 45
timeout_in_minutes: 50
device: h200_18gb
source_file_dependencies:
- vllm/
@@ -38,7 +38,7 @@ steps:
- label: "Multi-Modal Models (Standard) 3: llava + qwen2_vl"
device: h200_35gb
key: multi-modal-models-standard-3-llava-qwen2-vl
timeout_in_minutes: 45
timeout_in_minutes: 40
source_file_dependencies:
- vllm/
- tests/models/multimodal
@@ -54,7 +54,7 @@ steps:
- label: "Multi-Modal Models (Standard) 4: other + whisper"
device: h200_35gb
key: multi-modal-models-standard-4-other-whisper
timeout_in_minutes: 45
timeout_in_minutes: 50
source_file_dependencies:
- vllm/
- tests/models/multimodal
@@ -73,7 +73,7 @@ steps:
key: multi-modal-processor-cpu
depends_on:
- image-build-cpu
timeout_in_minutes: 60
timeout_in_minutes: 125
source_file_dependencies:
- vllm/
- tests/models/multimodal
@@ -84,7 +84,7 @@ steps:
- label: Multi-Modal Processor # 44min
key: multi-modal-processor
timeout_in_minutes: 60
timeout_in_minutes: 65
device: h200_18gb
source_file_dependencies:
- vllm/
@@ -96,7 +96,7 @@ steps:
- label: Multi-Modal Accuracy Eval (Small Models) # 50min
device: h200_35gb
key: multi-modal-accuracy-eval-small-models
timeout_in_minutes: 70
timeout_in_minutes: 30
working_dir: "/vllm-workspace/.buildkite/lm-eval-harness"
source_file_dependencies:
- vllm/multimodal/
@@ -164,7 +164,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 60
timeout_in_minutes: 75
depends_on:
- image-build-amd
source_file_dependencies:
+1 -1
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Plugin Tests (2 GPUs)
key: plugin-tests-2-gpus
timeout_in_minutes: 60
timeout_in_minutes: 35
working_dir: "/vllm-workspace/tests"
num_devices: 2
source_file_dependencies:
+5 -5
View File
@@ -5,7 +5,7 @@ steps:
- label: PyTorch Compilation Unit Tests
device: h200_35gb
key: pytorch-compilation-unit-tests
timeout_in_minutes: 10
timeout_in_minutes: 90
source_file_dependencies:
- vllm/__init__.py
- vllm/_aiter_ops.py
@@ -78,7 +78,7 @@ steps:
- label: PyTorch Compilation Passes Unit Tests
key: pytorch-compilation-passes-unit-tests
timeout_in_minutes: 20
timeout_in_minutes: 45
source_file_dependencies:
- vllm/__init__.py
- vllm/_aiter_ops.py
@@ -110,13 +110,13 @@ steps:
mirror:
amd:
device: mi300_1
timeout_in_minutes: 180
timeout_in_minutes: 65
depends_on:
- image-build-amd
- label: PyTorch Fullgraph Smoke Test
key: pytorch-fullgraph-smoke-test
timeout_in_minutes: 35
timeout_in_minutes: 60
source_file_dependencies:
- vllm/__init__.py
- vllm/_aiter_ops.py
@@ -152,7 +152,7 @@ steps:
- label: PyTorch Fullgraph
key: pytorch-fullgraph
timeout_in_minutes: 30
timeout_in_minutes: 40
device: h200_18gb
source_file_dependencies:
- vllm/__init__.py
+4 -4
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Quantization
key: quantization
timeout_in_minutes: 90
timeout_in_minutes: 60
source_file_dependencies:
- csrc/
- vllm/model_executor/layers/quantization
@@ -23,7 +23,7 @@ steps:
- label: Quantized Fusions
key: quantized-fusions
timeout_in_minutes: 30
timeout_in_minutes: 20
source_file_dependencies:
- tests/fusion
- vllm/model_executor/layers/fusion
@@ -35,7 +35,7 @@ steps:
- label: Quantized MoE Test (B200)
key: quantized-moe-test-b200
timeout_in_minutes: 60
timeout_in_minutes: 120
working_dir: "/vllm-workspace/"
device: b200-k8s
source_file_dependencies:
@@ -53,7 +53,7 @@ steps:
- label: Quantized Models Test
key: quantized-models-test
timeout_in_minutes: 60
timeout_in_minutes: 50
source_file_dependencies:
- vllm/model_executor/layers/quantization
- tests/models/quantization
+5 -5
View File
@@ -3,7 +3,7 @@ depends_on:
- image-build
steps:
- label: Rust Frontend OpenAI Coverage
timeout_in_minutes: 90
timeout_in_minutes: 30
device: h200_18gb
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -38,7 +38,7 @@ steps:
- pytest -v -s v1/sample/test_logprobs_e2e.py -k "test_prompt_logprobs_e2e_server"
- label: Rust Frontend Serve/Admin Coverage
timeout_in_minutes: 60
timeout_in_minutes: 25
device: h200_18gb
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -67,7 +67,7 @@ steps:
- pytest -v -s entrypoints/serve/tokenize/test_tokenization.py -k "not tokenizer_info"
- label: Rust Frontend Core Correctness
timeout_in_minutes: 30
timeout_in_minutes: 20
device: h200_18gb
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -81,7 +81,7 @@ steps:
- pytest -s entrypoints/openai/correctness/test_lmeval.py::test_lm_eval_accuracy_v1_engine
- label: Rust Frontend Tool Use
timeout_in_minutes: 60
timeout_in_minutes: 25
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
- rust/
@@ -95,7 +95,7 @@ steps:
- pytest -v -s tool_use --ignore=tool_use/mistral --models llama3.2 -k "not test_response_format_with_tool_choice_required and not test_parallel_tool_calls_false and not test_tool_call_and_choice"
- label: Rust Frontend Distributed
timeout_in_minutes: 30
timeout_in_minutes: 25
num_devices: 4
working_dir: "/vllm-workspace/tests"
source_file_dependencies:
@@ -4,7 +4,7 @@ steps:
- label: Rust Frontend Cargo Style + Clippy
key: rust-frontend-cargo-style-clippy
depends_on: []
timeout_in_minutes: 30
timeout_in_minutes: 20
device: cpu-medium
no_plugin: true
source_file_dependencies:
@@ -18,7 +18,7 @@ steps:
- label: Rust Frontend Cargo Tests
key: rust-frontend-cargo-tests
depends_on: []
timeout_in_minutes: 30
timeout_in_minutes: 20
device: cpu-medium
no_plugin: true
source_file_dependencies:
+1 -1
View File
@@ -5,7 +5,7 @@ steps:
- label: Samplers Test
device: h200_35gb
key: samplers-test
timeout_in_minutes: 75
timeout_in_minutes: 40
source_file_dependencies:
- vllm/model_executor/layers
- vllm/sampling_metadata.py
+11 -11
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Spec Decode Eagle
key: spec-decode-eagle
timeout_in_minutes: 30
timeout_in_minutes: 25
device: h200_18gb
source_file_dependencies:
- vllm/v1/spec_decode/
@@ -15,7 +15,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 45
timeout_in_minutes: 60
depends_on:
- image-build-amd
source_file_dependencies:
@@ -29,7 +29,7 @@ steps:
- label: Spec Decode Eagle Nightly B200
key: spec-decode-eagle-nightly-b200
timeout_in_minutes: 30
timeout_in_minutes: 25
device: b200-k8s
optional: true
source_file_dependencies:
@@ -41,7 +41,7 @@ steps:
- label: Spec Decode Speculators + MTP
key: spec-decode-speculators-mtp
timeout_in_minutes: 30
timeout_in_minutes: 20
device: h200_18gb
source_file_dependencies:
- vllm/v1/spec_decode/
@@ -82,7 +82,7 @@ steps:
- label: Spec Decode Ngram + Suffix
key: spec-decode-ngram-suffix
timeout_in_minutes: 30
timeout_in_minutes: 20
device: h200_18gb
source_file_dependencies:
- vllm/v1/spec_decode/
@@ -93,7 +93,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 65
timeout_in_minutes: 55
# TODO(akaratza): Test after Torch >= 2.12 bump
soft_fail: true
depends_on:
@@ -109,7 +109,7 @@ steps:
- label: Spec Decode Draft Model
key: spec-decode-draft-model
timeout_in_minutes: 30
timeout_in_minutes: 45
device: h200_18gb
source_file_dependencies:
- vllm/v1/spec_decode/
@@ -120,7 +120,7 @@ steps:
mirror:
amd:
device: mi325_1
timeout_in_minutes: 50
timeout_in_minutes: 55
depends_on:
- image-build-amd
source_file_dependencies:
@@ -134,7 +134,7 @@ steps:
- label: Spec Decode Draft Model Nightly B200
key: spec-decode-draft-model-nightly-b200
timeout_in_minutes: 30
timeout_in_minutes: 40
device: b200-k8s
optional: true
source_file_dependencies:
@@ -146,7 +146,7 @@ steps:
- label: Speculators Correctness
key: speculators-correctness
timeout_in_minutes: 60
timeout_in_minutes: 30
device: h100
optional: true
num_devices: 1
@@ -159,7 +159,7 @@ steps:
- pytest -v -s v1/spec_decode/test_speculators_correctness.py -m slow_test
- label: Spec Decode MTP hybrid (B200)
timeout_in_minutes: 30
timeout_in_minutes: 20
device: b200-k8s
optional: true
source_file_dependencies:
+1 -1
View File
@@ -4,7 +4,7 @@ depends_on:
steps:
- label: Weight Loading Multiple GPU # 33min
key: weight-loading-multiple-gpu
timeout_in_minutes: 45
timeout_in_minutes: 50
working_dir: "/vllm-workspace/tests"
num_devices: 2
optional: true
+12 -2
View File
@@ -70,6 +70,15 @@ endif()
#
set(TORCH_SUPPORTED_VERSION_CUDA "2.11.0")
set(TORCH_SUPPORTED_VERSION_ROCM "2.11.0")
# TORCH_NIGHTLY=1 builds run against unpinned nightly wheels, so the supported-
# version check would always warn. Only treat it as a nightly build when the
# value is exactly "1" (the bootstrap exports TORCH_NIGHTLY=0 by default, which
# must NOT suppress the warning for normal builds).
if (DEFINED ENV{TORCH_NIGHTLY} AND "$ENV{TORCH_NIGHTLY}" STREQUAL "1")
set(TORCH_NIGHTLY_BUILD TRUE)
else()
set(TORCH_NIGHTLY_BUILD FALSE)
endif()
#
# Try to find python package with an executable that exactly matches
@@ -175,7 +184,7 @@ endif()
if (NOT HIP_FOUND AND NOT PYTORCH_FOUND_HIP AND CUDA_FOUND)
set(VLLM_GPU_LANG "CUDA")
if (NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
if (NOT TORCH_NIGHTLY_BUILD AND NOT Torch_VERSION VERSION_EQUAL ${TORCH_SUPPORTED_VERSION_CUDA})
message(WARNING "Pytorch version ${TORCH_SUPPORTED_VERSION_CUDA} "
"expected for CUDA build, saw ${Torch_VERSION} instead.")
endif()
@@ -188,7 +197,7 @@ elseif(HIP_FOUND OR PYTORCH_FOUND_HIP)
enable_language(HIP)
# ROCm 5.X and 6.X
if (ROCM_VERSION_DEV_MAJOR GREATER_EQUAL 5 AND
if (NOT TORCH_NIGHTLY_BUILD AND ROCM_VERSION_DEV_MAJOR GREATER_EQUAL 5 AND
Torch_VERSION VERSION_LESS ${TORCH_SUPPORTED_VERSION_ROCM})
message(WARNING "Pytorch version >= ${TORCH_SUPPORTED_VERSION_ROCM} "
"expected for ROCm build, saw ${Torch_VERSION} instead.")
@@ -381,6 +390,7 @@ if(VLLM_GPU_LANG STREQUAL "CUDA" OR VLLM_GPU_LANG STREQUAL "HIP")
"csrc/libtorch_stable/cuda_view.cu"
"csrc/libtorch_stable/cuda_utils_kernels.cu"
"csrc/libtorch_stable/activation_kernels.cu"
"csrc/libtorch_stable/ngram_embedding_kernels.cu"
"csrc/libtorch_stable/quantization/activation_kernels.cu"
"csrc/libtorch_stable/quantization/w8a8/int8/scaled_quant.cu"
"csrc/libtorch_stable/quantization/w8a8/fp8/common.cu"
@@ -0,0 +1,96 @@
// N-gram embedding index kernel for LongCat-Flash (n-gram embedding variant).
//
// Adapted from SGLang:
// https://github.com/sgl-project/sglang/blob/main/python/sglang/jit_kernel/csrc/ngram_embedding.cuh
//
// For each position, computes the hashed n-gram embedding ids that index the
// concatenated embedder table. Integer tensors are int32 except ``row_indices``
// (int64); the token table is ``[max_running_reqs, max_context_len]`` int32,
// where a negative entry marks an ignored token (e.g. an EOS boundary).
#include "torch_utils.h"
#include "ops.h"
#include <cstdint>
namespace vllm::ngram_embedding {
constexpr int kBlockThreads = 256;
__global__ void ComputeNGramIdsKernel(
int batch_size, int ne_n, int ne_k,
int* ne_weights, // [ne_n-1, ne_k, ne_n]
int* ne_mods, // [ne_n-1, ne_k]
int* exclusive_ne_embedder_size_sums, // [(ne_n-1)*ne_k + 1]
int* exclusive_req_len_sums, // [batch_size + 1]
int* ne_token_table, // [max_running_reqs, max_context_len]
int max_context_len,
const int64_t* __restrict__ row_indices, // [batch_size]
int* column_starts, // [batch_size]
int* n_gram_ids // [token_num, (ne_n-1)*ne_k]
) {
const int req_id = blockIdx.x % batch_size;
const int config_id = (blockIdx.x - req_id) / batch_size;
// n and k are offset from their physical meaning: n = real_n - 2, k = real_k
// - 1 (they index into ne_weights / ne_mods).
const int k = config_id % ne_k;
const int n = (config_id - config_id % ne_k) / ne_k;
const int ne_weight_base_idx = n * ne_k * ne_n + k * ne_n;
const int ne_mod = ne_mods[n * ne_k + k];
for (int i = exclusive_req_len_sums[req_id] + threadIdx.x;
i < exclusive_req_len_sums[req_id + 1]; i += blockDim.x) {
uint64_t n_gram_id = 0;
const int64_t current_token_offset = i - exclusive_req_len_sums[req_id];
const int64_t req_token_table_index =
row_indices[req_id] * static_cast<int64_t>(max_context_len);
const int64_t current_token_table_index =
req_token_table_index + column_starts[req_id] + current_token_offset;
for (int j = 0; j < n + 2; j++) {
if (current_token_table_index - j < req_token_table_index) {
break; // outside this request's range
}
if (ne_token_table[current_token_table_index - j] < 0) {
break; // ignored token
}
const uint64_t term =
(uint64_t)ne_token_table[current_token_table_index - j] *
(uint64_t)ne_weights[ne_weight_base_idx + j];
n_gram_id += term % ne_mod;
}
n_gram_id %= ne_mod;
n_gram_id += exclusive_ne_embedder_size_sums[n * ne_k + k];
n_gram_ids[i * (ne_n - 1) * ne_k + n * ne_k + k] = (int)(n_gram_id);
}
}
} // namespace vllm::ngram_embedding
void ngram_compute_n_gram_ids(
int64_t ne_n, int64_t ne_k, torch::stable::Tensor& ne_weights,
torch::stable::Tensor& ne_mods,
torch::stable::Tensor& exclusive_ne_embedder_size_sums,
torch::stable::Tensor& exclusive_req_len_sums,
torch::stable::Tensor& ne_token_table, torch::stable::Tensor& row_indices,
torch::stable::Tensor& column_starts, torch::stable::Tensor& n_gram_ids) {
const int batch_size = static_cast<int>(exclusive_req_len_sums.size(0) - 1);
const int max_context_len = static_cast<int>(ne_token_table.size(1));
const int num_configs = (static_cast<int>(ne_n) - 1) * static_cast<int>(ne_k);
const int grid_size = num_configs * batch_size;
if (grid_size <= 0) return;
const torch::stable::accelerator::DeviceGuard device_guard(
ne_weights.get_device_index());
const cudaStream_t stream = get_current_cuda_stream();
vllm::ngram_embedding::ComputeNGramIdsKernel<<<
grid_size, vllm::ngram_embedding::kBlockThreads, 0, stream>>>(
batch_size, static_cast<int>(ne_n), static_cast<int>(ne_k),
ne_weights.mutable_data_ptr<int32_t>(),
ne_mods.mutable_data_ptr<int32_t>(),
exclusive_ne_embedder_size_sums.mutable_data_ptr<int32_t>(),
exclusive_req_len_sums.mutable_data_ptr<int32_t>(),
ne_token_table.mutable_data_ptr<int32_t>(), max_context_len,
row_indices.const_data_ptr<int64_t>(),
column_starts.mutable_data_ptr<int32_t>(),
n_gram_ids.mutable_data_ptr<int32_t>());
}
+9
View File
@@ -554,3 +554,12 @@ void cp_gather_indexer_k_quant_cache(
// quant_block_size * 4]
const torch::stable::Tensor& block_table, // [batch_size, num_blocks]
const torch::stable::Tensor& cu_seq_lens); // [batch_size + 1]
// LongCat n-gram embedding index kernel (see ngram_embedding_kernels.cu).
void ngram_compute_n_gram_ids(
int64_t ne_n, int64_t ne_k, torch::stable::Tensor& ne_weights,
torch::stable::Tensor& ne_mods,
torch::stable::Tensor& exclusive_ne_embedder_size_sums,
torch::stable::Tensor& exclusive_req_len_sums,
torch::stable::Tensor& ne_token_table, torch::stable::Tensor& row_indices,
torch::stable::Tensor& column_starts, torch::stable::Tensor& n_gram_ids);
+13
View File
@@ -598,9 +598,22 @@ STABLE_TORCH_LIBRARY_FRAGMENT(_C, ops) {
"Tensor? initial_state_idx,"
"Tensor? cu_chunk_seqlen,"
"Tensor? last_chunk_indices) -> ()");
// LongCat n-gram embedding index kernel. All tensor args are marked mutable
// to match the (non-const) stable-Tensor& C++ signature; only ne_token_table
// and n_gram_ids are actually written in place.
ops.def(
"ngram_compute_n_gram_ids(int ne_n, int ne_k, Tensor(a!) ne_weights, "
"Tensor(b!) ne_mods, Tensor(c!) exclusive_ne_embedder_size_sums, "
"Tensor(d!) exclusive_req_len_sums, Tensor(e!) ne_token_table, "
"Tensor(f!) row_indices, Tensor(g!) column_starts, "
"Tensor(h!) n_gram_ids) -> ()");
}
STABLE_TORCH_LIBRARY_IMPL(_C, CUDA, ops) {
// LongCat n-gram embedding index kernel.
ops.impl("ngram_compute_n_gram_ids", TORCH_BOX(&ngram_compute_n_gram_ids));
// Per-token group quantization
ops.impl("per_token_group_fp8_quant", TORCH_BOX(&per_token_group_quant_fp8));
ops.impl("per_token_group_fp8_quant_packed",
@@ -93,6 +93,28 @@ To address this, manually trigger a build on Buildkite to accomplish two objecti
<img width="60%" alt="Buildkite new build popup" src="https://github.com/user-attachments/assets/3b07f71b-bb18-4ca3-aeaf-da0fe79d315f" />
</p>
You can also trigger this build from the command line with
[`.buildkite/scripts/trigger-ci-build.sh`](../../../.buildkite/scripts/trigger-ci-build.sh)
(dry-run by default; pass `--execute` to actually trigger it).
## Test against PyTorch nightly
The steps above test against a specific PyTorch RC/stable wheel pinned in the
requirements files. To instead build and run the CI suite against the latest
PyTorch **nightly** wheels, set the `TORCH_NIGHTLY=1` environment variable on
the build (or apply the `ready-torch-nightly` label to the PR).
When `TORCH_NIGHTLY=1`, the base CI image is built against PyTorch nightly
(`image_build_torch_nightly.sh`, `PYTORCH_NIGHTLY=1`, CUDA 13.0) and tagged at
the normal image tag, so the entire existing pipeline runs on nightly torch --
there is no separate pipeline section to trigger. Combine it with `RUN_ALL=1`
to run the full suite (the `ready-torch-nightly` label and
`trigger-ci-build.sh --torch-nightly` both set this for you). This is the
configuration to use for a scheduled "vLLM vs PyTorch nightly" run.
Use `.buildkite/scripts/trigger-ci-build.sh --torch-nightly` to trigger it from
the command line.
## Update all the different vLLM platforms
Rather than attempting to update all vLLM platforms in a single pull request, it's more manageable
+49
View File
@@ -879,6 +879,55 @@ vllm serve Qwen/Qwen3-VL-30B-A3B-Instruct \
Works with common video formats like MP4 when using OpenCV backends.
#### GPU Video Decoding with DeepStream (NVDEC)
By default vLLM decodes video on the CPU. On NVIDIA GPUs you can instead decode
directly on the hardware video engine (NVDEC) with the DeepStream backend, which
keeps decoding off the CPU and can significantly increase video throughput.
Install the backend (Linux x86-64 only):
```bash
pip install vllm[deepstream]
```
The pip wheel bundles the DeepStream libraries but still relies on a few system
packages that pip cannot install. On Ubuntu:
```bash
apt-get install -y \
gstreamer1.0-tools gstreamer1.0-plugins-base gstreamer1.0-plugins-good \
gstreamer1.0-plugins-bad gstreamer1.0-libav \
python3-gi python3-gst-1.0 libv4l-0 cuda-libraries-13-0
```
Select the backend either with an environment variable:
```bash
export VLLM_VIDEO_LOADER_BACKEND=deepstream
vllm serve Qwen/Qwen3-VL-30B-A3B-Instruct
```
or per request via `--media-io-kwargs`:
```bash
vllm serve Qwen/Qwen3-VL-30B-A3B-Instruct \
--media-io-kwargs '{"video": {"backend": "deepstream"}}'
```
**Parameters:**
- `pool_size`: Number of GPU decode workers in the process-wide decode pool
(clamped to `[1, 16]`). When unset it defaults to
`VLLM_MEDIA_LOADING_THREAD_COUNT` (default `8`). The pool is a singleton, so
the first request's value wins.
```bash
# Example: 12 decode workers
vllm serve Qwen/Qwen3-VL-30B-A3B-Instruct \
--media-io-kwargs '{"video": {"backend": "deepstream", "pool_size": 12}}'
```
#### Pre-extracted Frame Sequences with `media_io_kwargs`
When you extract video frames on the client side and send them as `video/jpeg` (base64-concatenated JPEG frames), you can preserve the original video metadata by using `media_io_kwargs` in your request. This enables more accurate video understanding by preserving temporal information that would otherwise be lost during client-side frame extraction.
+63
View File
@@ -351,6 +351,69 @@ print(response.choices[0].message.reasoning)
print(response.choices[0].message.content)
```
## Suppressing Reasoning Output
You can suppress reasoning content from API responses using the `include_reasoning` parameter. When set to `false`, reasoning tokens are still generated (so model quality is unaffected) but excluded from the response. This reduces network traffic without changing inference behavior.
The parameter is supported in both the Chat Completions API and the Responses API, for streaming and non-streaming requests.
When `include_reasoning=false`, vLLM also suppresses per-token metadata (logprobs and token IDs) to prevent leaking reasoning content through decoded token text in logprob entries or raw token IDs.
### Chat Completions API
```python
from openai import OpenAI
client = OpenAI(base_url="http://localhost:8000/v1", api_key="EMPTY")
model = client.models.list().data[0].id
# Reasoning is included by default (include_reasoning=True)
response = client.chat.completions.create(
model=model,
messages=[{"role": "user", "content": "What is 15 * 37?"}],
extra_body={"include_reasoning": False},
)
msg = response.choices[0].message
assert msg.content # Content is still present
assert not getattr(msg, "reasoning", None) # Reasoning is suppressed
```
Streaming works the same way, reasoning deltas are omitted from chunks:
```python
stream = client.chat.completions.create(
model=model,
messages=[{"role": "user", "content": "What is 15 * 37?"}],
stream=True,
extra_body={"include_reasoning": False},
)
for chunk in stream:
delta = chunk.choices[0].delta
# delta.reasoning will always be None
if delta.content:
print(delta.content, end="", flush=True)
```
### Responses API
```python
from openai import OpenAI
client = OpenAI(base_url="http://localhost:8000/v1", api_key="EMPTY")
response = client.responses.create(
model=client.models.list().data[0].id,
input="What is 15 * 37?",
include_reasoning=False,
)
# No "reasoning" items in output
types = [item.type for item in response.output]
assert "reasoning" not in types
```
## Limitations
- The reasoning content is only available for online serving's chat completion endpoint (`/v1/chat/completions`), Anthropic Messages API (`/v1/messages`) and the Responses API (`/v1/responses`).
+3 -1
View File
@@ -405,6 +405,8 @@ th {
| `Lfm2MoeForCausalLM` | LFM2MoE | `LiquidAI/LFM2-8B-A1B-preview`, etc. | ✅︎ | ✅︎ |
| `LlamaForCausalLM` | Llama 3.1, Llama 3, Llama 2, LLaMA, Yi | `meta-llama/Meta-Llama-3.1-405B-Instruct`, `meta-llama/Meta-Llama-3.1-70B`, `meta-llama/Meta-Llama-3-70B-Instruct`, `meta-llama/Llama-2-70b-hf`, `01-ai/Yi-34B`, etc. | ✅︎ | ✅︎ |
| `LongcatFlashForCausalLM` | LongCat-Flash | `meituan-longcat/LongCat-Flash-Chat`, `meituan-longcat/LongCat-Flash-Chat-FP8` | ✅︎ | ✅︎ |
| `LongcatFlashNgramForCausalLM` | LongCat-Flash-Lite | `meituan-longcat/LongCat-Flash-Lite` | ✅︎ | ✅︎ |
| `LongcatCausalLM` | LongCat-2.0 | `meituan-longcat/LongCat-2.0-FP8` | ✅︎ | ✅︎ |
| `MambaForCausalLM` | Mamba | `state-spaces/mamba-130m-hf`, `state-spaces/mamba-790m-hf`, `state-spaces/mamba-2.8b-hf`, etc. | | ✅︎ |
| `Mamba2ForCausalLM` | Mamba2 | `mistralai/Mamba-Codestral-7B-v0.1`, etc. | | ✅︎ |
| `MellumForCausalLM` | Mellum 2 | `JetBrains/Mellum2-12B-A2.5B-Base`, etc. | | ✅︎ |
@@ -512,7 +514,7 @@ These models primarily accept the [`LLM.generate`](./generative_models.md#llmgen
| `ChameleonForConditionalGeneration` | Chameleon | T + I | `facebook/chameleon-7b`, etc. | | ✅︎ |
| `CheersForConditionalGeneration` | Cheers | T + I | `ai9stars/Cheers` | | ✅︎ |
| `Cohere2VisionForConditionalGeneration` | Command A Vision, Command-A+ | T + I<sup>+</sup> | `CohereLabs/command-a-vision-07-2025`, `CohereLabs/command-a-plus-05-2026`, etc. | | ✅︎ |
| `Cosmos3ForConditionalGeneration` | Cosmos3 (understanding tower) | T + I<sup>E+</sup> + V<sup>E+</sup> | `nvidia/Cosmos3-Nano` | | ✅︎ |
| `Cosmos3ForConditionalGeneration` | Cosmos3 (understanding tower) | T + I<sup>E+</sup> + V<sup>E+</sup> | `nvidia/Cosmos3-Nano`, `nvidia/Cosmos3-Super` | | ✅︎ |
| `DeepseekVLV2ForCausalLM` | DeepSeek-VL2 | T + I<sup>+</sup> | `deepseek-ai/deepseek-vl2-tiny`, `deepseek-ai/deepseek-vl2-small`, `deepseek-ai/deepseek-vl2`, etc. | | ✅︎ |
| `DeepseekOCRForCausalLM` | DeepSeek-OCR | T + I<sup>+</sup> | `deepseek-ai/DeepSeek-OCR`, etc. | ✅︎ | ✅︎ |
| `DeepseekOCR2ForCausalLM` | DeepSeek-OCR-2 | T + I<sup>+</sup> | `deepseek-ai/DeepSeek-OCR-2`, etc. | ✅︎ | ✅︎ |
+3 -1
View File
@@ -2167,7 +2167,7 @@ checksum = "11d3d7f243d5c5a8b9bb5d6dd2b1602c0cb0b9db1621bafc7ed66e35ff9fe092"
[[package]]
name = "llm-multimodal"
version = "1.7.1"
source = "git+https://github.com/smg-project/llm-multimodal?rev=7d74582aeaf0e4086a44964382655d22f1af0686#7d74582aeaf0e4086a44964382655d22f1af0686"
source = "git+https://github.com/smg-project/llm-multimodal?rev=c8a29dcc755139fdc26185f400ea48c6d6d48273#c8a29dcc755139fdc26185f400ea48c6d6d48273"
dependencies = [
"anyhow",
"base64 0.22.1",
@@ -2473,6 +2473,7 @@ dependencies = [
"portable-atomic",
"portable-atomic-util",
"rawpointer",
"serde",
]
[[package]]
@@ -5093,6 +5094,7 @@ dependencies = [
"llm-multimodal",
"minijinja",
"minijinja-contrib",
"ndarray 0.17.2",
"oss-harmony",
"paste",
"reqwest",
+2 -2
View File
@@ -53,12 +53,12 @@ hyper-util = { version = "0.1.20", features = [
indexmap = "2.13.0"
itertools = "0.14.0"
libc = "0.2.177"
llm-multimodal = { git = "https://github.com/smg-project/llm-multimodal", rev = "7d74582aeaf0e4086a44964382655d22f1af0686" }
llm-multimodal = { git = "https://github.com/smg-project/llm-multimodal", rev = "c8a29dcc755139fdc26185f400ea48c6d6d48273" }
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"] }
ndarray = { version = "0.17", features = ["serde"] }
openai-harmony = { package = "oss-harmony", git = "https://github.com/oss-harmony/harmony", tag = "v0.0.11", default-features = false }
openai-protocol = "1.6.0"
openssl = "0.10"
+1
View File
@@ -42,6 +42,7 @@ anyhow.workspace = true
bytes.workspace = true
clap.workspace = true
expect-test.workspace = true
ndarray.workspace = true
paste.workspace = true
rmp-serde.workspace = true
serial_test.workspace = true
+13 -4
View File
@@ -10,7 +10,7 @@ use crate::backend::{
NewChatOutputProcessorOptions,
};
use crate::error::Result;
use crate::multimodal::MultimodalModelInfo;
use crate::multimodal::{MultimodalConfigFiles, MultimodalModelInfo};
use crate::output::{
DefaultChatOutputProcessor, HarmonyChatOutputProcessor, validate_harmony_parser_overrides,
};
@@ -46,8 +46,12 @@ impl HfChatBackend {
MultimodalModelInfo::from_paths(
model_id.clone(),
(!model_type.is_empty()).then_some(model_type.to_string()),
files.config_path.as_deref(),
files.preprocessor_config_path.as_deref(),
MultimodalConfigFiles {
config: files.config_path.as_deref(),
preprocessor_config: files.preprocessor_config_path.as_deref(),
video_preprocessor_config: files.video_preprocessor_config_path.as_deref(),
processor_config: files.processor_config_path.as_deref(),
},
tokenizer.clone(),
)?
};
@@ -139,8 +143,11 @@ pub(super) async fn load_model_backends(
fn resolve_multimodal_render_info(
info: Option<&MultimodalModelInfo>,
) -> Option<MultimodalRenderInfo> {
use llm_multimodal::Modality;
info.map(|info| MultimodalRenderInfo {
placeholder_token: info.placeholder_token().to_string(),
image_token: info.placeholder_token(Modality::Image).map(str::to_string),
video_token: info.placeholder_token(Modality::Video).map(str::to_string),
})
}
@@ -192,6 +199,8 @@ mod tests {
tokenizer_config_path: Some(tokenizer_config_path),
generation_config_path: None,
preprocessor_config_path: None,
video_preprocessor_config_path: None,
processor_config_path: None,
chat_template_path: None,
config_path: Some(config_path),
}
+31 -1
View File
@@ -1,5 +1,5 @@
use thiserror::Error;
use thiserror_ext::Macro;
use thiserror_ext::{AsReport as _, Macro};
type BoxedError = Box<dyn std::error::Error + Send + Sync>;
@@ -18,6 +18,8 @@ pub enum Error {
UnsupportedMultimodalRenderer,
#[error("unsupported multimodal content: {0}")]
UnsupportedMultimodalContent(&'static str),
#[error("`{modality}` input is not supported by this model")]
UnsupportedModality { modality: String },
#[error("multimodal preprocessing error: {0}")]
Multimodal(#[message] String),
#[error("{kind} parsing is not available for model `{model_id}`")]
@@ -80,11 +82,39 @@ impl Error {
match self {
Self::PromptTooLong { .. } => true,
Self::Text(error) => error.is_request_validation_error(),
Self::UnsupportedMultimodalRenderer
| Self::UnsupportedMultimodalContent(_)
| Self::UnsupportedModality { .. } => true,
_ => false,
}
}
}
impl From<llm_multimodal::MediaConnectorError> for Error {
fn from(error: llm_multimodal::MediaConnectorError) -> Self {
Self::Multimodal(error.to_report_string())
}
}
impl From<llm_multimodal::MultiModalError> for Error {
fn from(error: llm_multimodal::MultiModalError) -> Self {
Self::Multimodal(error.to_report_string())
}
}
impl From<llm_multimodal::TransformError> for Error {
fn from(error: llm_multimodal::TransformError) -> Self {
Self::Multimodal(error.to_report_string())
}
}
impl From<llm_multimodal::registry::ModelRegistryError> for Error {
fn from(error: llm_multimodal::registry::ModelRegistryError) -> Self {
Self::Multimodal(error.to_report_string())
}
}
/// Format the available-parser suffix used in user-facing error messages.
fn available_parser_hint(available_names: &[String]) -> String {
if available_names.is_empty() {
File diff suppressed because it is too large Load Diff
+446
View File
@@ -0,0 +1,446 @@
//! Prompt placeholder expansion shared across modalities.
use std::collections::{HashMap, VecDeque};
use llm_multimodal::{Modality, PromptReplacement};
use vllm_engine_core_client::protocol::multimodal::PlaceholderRange;
use vllm_engine_core_client::protocol::tensor::WireTensor;
use super::PreparedMedia;
use crate::error::{Error, Result, bail_multimodal};
/// One modality's queue of pending placeholder replacements for prompt
/// expansion.
struct ExpansionLane<'a> {
modality: Modality,
marker_token_id: u32,
embed_token_id: u32,
placeholder_token: String,
replacements: VecDeque<&'a PromptReplacement>,
}
impl<'a> ExpansionLane<'a> {
fn from_prepared(media: &'a PreparedMedia) -> Option<Self> {
if media.replacements.is_empty() {
return None;
}
Some(Self {
modality: media.modality,
marker_token_id: media.placeholder.marker_token_id,
embed_token_id: media.placeholder.embed_token_id,
placeholder_token: media.placeholder.token.clone(),
replacements: media.replacements.iter().collect(),
})
}
}
/// Replace rendered placeholder markers with model-specific replacement
/// tokens across all modalities in one left-to-right pass.
///
/// Each prepared modality consumes its own marker occurrences in order,
/// matching the original media-part order within that modality; markers of
/// different modalities may interleave freely.
///
/// The returned ranges point into the already-expanded prompt, grouped per
/// modality in item order.
pub(super) fn expand_prompt_token_ids(
prompt_token_ids: &mut Vec<u32>,
prepared: &[PreparedMedia],
) -> Result<HashMap<Modality, Vec<PlaceholderRange>>> {
let mut lanes = prepared.iter().filter_map(ExpansionLane::from_prepared).collect::<Vec<_>>();
if lanes.is_empty() {
return Ok(HashMap::new());
}
let replacement_growth = lanes
.iter()
.flat_map(|lane| lane.replacements.iter())
.fold(0usize, |total, replacement| {
total.saturating_add(replacement.tokens.len().saturating_sub(1))
});
let expanded_len = prompt_token_ids.len().saturating_add(replacement_growth);
let mut expanded = Vec::with_capacity(expanded_len);
let mut ranges = HashMap::<Modality, Vec<PlaceholderRange>>::new();
for &token in prompt_token_ids.iter() {
let lane = lanes
.iter_mut()
.find(|lane| lane.marker_token_id == token && !lane.replacements.is_empty());
let Some(lane) = lane else {
expanded.push(token);
continue;
};
let replacement = lane.replacements.pop_front().expect("lane queue is non-empty");
debug_assert_eq!(replacement.modality, lane.modality);
if replacement.tokens.is_empty() {
bail_multimodal!(
"placeholder token `{}` expanded to no tokens",
lane.placeholder_token
);
}
let replacement_len = replacement.tokens.len();
let is_embed = {
let mask = replacement
.tokens
.iter()
.map(|&token| token as u32 == lane.embed_token_id)
.collect::<Vec<_>>();
WireTensor::from_bool(vec![replacement_len], mask).map_err(Error::Multimodal)?
};
let expanded_offset = expanded.len();
expanded.extend(replacement.tokens.iter().map(|&token| token as u32));
ranges.entry(lane.modality).or_default().push(PlaceholderRange {
offset: expanded_offset,
length: replacement_len,
is_embed: Some(is_embed),
});
}
for lane in &lanes {
if !lane.replacements.is_empty() {
bail_multimodal!(
"placeholder token `{}` was not found in tokenized prompt for {} remaining `{}` item(s)",
lane.placeholder_token,
lane.replacements.len(),
lane.modality
);
}
}
*prompt_token_ids = expanded;
Ok(ranges)
}
#[cfg(test)]
mod tests {
use llm_multimodal::TokenId;
use vllm_engine_core_client::protocol::tensor::WireArrayData;
use super::super::tests::{
LLAMA4_IMAGE_END_ID, LLAMA4_IMAGE_ID, LLAMA4_IMAGE_START_ID, LLAMA4_PATCH_ID,
LLAMA4_TILE_X_SEPARATOR_ID, LLAMA4_TILE_Y_SEPARATOR_ID, QWEN3_IMAGE_PAD_ID,
QWEN3_VIDEO_PAD_ID,
};
use super::super::{PreparedMedia, ResolvedPlaceholder};
use super::*;
/// Build prepared media directly from placeholder token IDs.
fn prepared_media(
modality: Modality,
placeholder_token: &str,
marker_token_id: u32,
embed_token_id: u32,
replacements: Vec<PromptReplacement>,
) -> PreparedMedia {
PreparedMedia {
modality,
placeholder: ResolvedPlaceholder {
token: placeholder_token.to_string(),
marker_token_id,
embed_token_id,
},
replacements,
items: Vec::new(),
}
}
/// Llama4 image prepared media: the `<|image|>` marker expands to
/// sequences whose embed positions are the `<|patch|>` tokens.
fn llama4_prepared(replacements: Vec<PromptReplacement>) -> PreparedMedia {
prepared_media(
Modality::Image,
"<|image|>",
LLAMA4_IMAGE_ID,
LLAMA4_PATCH_ID,
replacements,
)
}
fn qwen3_image_prepared(replacements: Vec<PromptReplacement>) -> PreparedMedia {
prepared_media(
Modality::Image,
"<|image_pad|>",
QWEN3_IMAGE_PAD_ID,
QWEN3_IMAGE_PAD_ID,
replacements,
)
}
fn qwen3_video_prepared(replacements: Vec<PromptReplacement>) -> PreparedMedia {
prepared_media(
Modality::Video,
"<|video_pad|>",
QWEN3_VIDEO_PAD_ID,
QWEN3_VIDEO_PAD_ID,
replacements,
)
}
fn llama4_single_tile_replacement() -> PromptReplacement {
PromptReplacement::sequence(
Modality::Image,
"<|image|>",
vec![
LLAMA4_IMAGE_START_ID as TokenId,
LLAMA4_IMAGE_ID as TokenId,
LLAMA4_PATCH_ID as TokenId,
LLAMA4_PATCH_ID as TokenId,
LLAMA4_IMAGE_END_ID as TokenId,
],
)
}
fn llama4_multi_tile_replacement() -> PromptReplacement {
PromptReplacement::sequence(
Modality::Image,
"<|image|>",
vec![
LLAMA4_IMAGE_START_ID as TokenId,
LLAMA4_PATCH_ID as TokenId,
LLAMA4_TILE_X_SEPARATOR_ID as TokenId,
LLAMA4_PATCH_ID as TokenId,
LLAMA4_TILE_Y_SEPARATOR_ID as TokenId,
LLAMA4_IMAGE_ID as TokenId,
LLAMA4_PATCH_ID as TokenId,
LLAMA4_IMAGE_END_ID as TokenId,
],
)
}
fn assert_bool_mask(range: &PlaceholderRange, expected: &[bool]) {
let tensor = range.is_embed.as_ref().expect("is_embed mask");
assert_eq!(tensor.dtype, "bool");
assert_eq!(tensor.shape, vec![expected.len()]);
assert_eq!(
tensor.data,
WireArrayData::RawView(expected.iter().map(|value| u8::from(*value)).collect())
);
}
#[test]
fn expand_prompt_tokens_marks_only_llama4_patch_tokens_as_embed() {
let mut prompt_token_ids = vec![1, LLAMA4_IMAGE_ID, 2];
let prepared = vec![llama4_prepared(vec![llama4_multi_tile_replacement()])];
let ranges = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap();
let ranges = &ranges[&Modality::Image];
assert_eq!(
prompt_token_ids,
vec![
1,
LLAMA4_IMAGE_START_ID,
LLAMA4_PATCH_ID,
LLAMA4_TILE_X_SEPARATOR_ID,
LLAMA4_PATCH_ID,
LLAMA4_TILE_Y_SEPARATOR_ID,
LLAMA4_IMAGE_ID,
LLAMA4_PATCH_ID,
LLAMA4_IMAGE_END_ID,
2,
]
);
assert_eq!(ranges[0].offset, 1);
assert_eq!(ranges[0].length, 8);
assert_bool_mask(
&ranges[0],
&[false, true, false, true, false, false, true, false],
);
}
#[test]
fn expand_prompt_tokens_errors_when_placeholder_missing() {
let mut prompt_token_ids = vec![1, 2, 3];
let prepared = vec![llama4_prepared(vec![llama4_single_tile_replacement()])];
let error = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap_err();
assert!(matches!(error, Error::Multimodal(message) if message.contains("not found")));
}
#[test]
fn expand_prompt_tokens_ignores_empty_replacements() {
let mut prompt_token_ids = vec![1, LLAMA4_IMAGE_ID, 2];
let original_prompt_token_ids = prompt_token_ids.clone();
let prepared = vec![llama4_prepared(Vec::new())];
let ranges = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap();
assert!(ranges.is_empty());
assert_eq!(prompt_token_ids, original_prompt_token_ids);
}
#[test]
fn expand_prompt_tokens_leaves_prompt_unchanged_when_later_placeholder_missing() {
let mut prompt_token_ids = vec![1, LLAMA4_IMAGE_ID, 2];
let original_prompt_token_ids = prompt_token_ids.clone();
let prepared = vec![llama4_prepared(vec![
llama4_single_tile_replacement(),
llama4_single_tile_replacement(),
])];
let error = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap_err();
assert!(matches!(error, Error::Multimodal(message) if message.contains("not found")));
assert_eq!(prompt_token_ids, original_prompt_token_ids);
}
#[test]
fn expand_prompt_tokens_errors_when_replacement_is_empty() {
let mut prompt_token_ids = vec![1, LLAMA4_IMAGE_ID, 2];
let original_prompt_token_ids = prompt_token_ids.clone();
let prepared = vec![llama4_prepared(vec![PromptReplacement::sequence(
Modality::Image,
"<|image|>",
Vec::new(),
)])];
let error = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap_err();
assert!(
matches!(error, Error::Multimodal(message) if message.contains("expanded to no tokens"))
);
assert_eq!(prompt_token_ids, original_prompt_token_ids);
}
#[test]
fn expand_prompt_tokens_skips_llama4_image_marker_inside_replacement() {
let mut prompt_token_ids = vec![1, LLAMA4_IMAGE_ID, 2, LLAMA4_IMAGE_ID, 3];
let prepared = vec![llama4_prepared(vec![
llama4_single_tile_replacement(),
llama4_single_tile_replacement(),
])];
let ranges = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap();
let ranges = &ranges[&Modality::Image];
assert_eq!(
prompt_token_ids,
vec![
1,
LLAMA4_IMAGE_START_ID,
LLAMA4_IMAGE_ID,
LLAMA4_PATCH_ID,
LLAMA4_PATCH_ID,
LLAMA4_IMAGE_END_ID,
2,
LLAMA4_IMAGE_START_ID,
LLAMA4_IMAGE_ID,
LLAMA4_PATCH_ID,
LLAMA4_PATCH_ID,
LLAMA4_IMAGE_END_ID,
3,
]
);
assert_eq!(ranges[0].offset, 1);
assert_eq!(ranges[0].length, 5);
assert_bool_mask(&ranges[0], &[false, false, true, true, false]);
assert_eq!(ranges[1].offset, 7);
assert_eq!(ranges[1].length, 5);
assert_bool_mask(&ranges[1], &[false, false, true, true, false]);
}
#[test]
fn expand_prompt_tokens_interleaves_image_and_video_prepared_media() {
let mut prompt_token_ids = vec![
1,
QWEN3_IMAGE_PAD_ID,
2,
QWEN3_VIDEO_PAD_ID,
3,
QWEN3_IMAGE_PAD_ID,
4,
];
let prepared = vec![
qwen3_image_prepared(vec![
PromptReplacement::repeated(
Modality::Image,
"<|image_pad|>",
QWEN3_IMAGE_PAD_ID as TokenId,
2,
),
PromptReplacement::repeated(
Modality::Image,
"<|image_pad|>",
QWEN3_IMAGE_PAD_ID as TokenId,
3,
),
]),
qwen3_video_prepared(vec![PromptReplacement::repeated(
Modality::Video,
"<|video_pad|>",
QWEN3_VIDEO_PAD_ID as TokenId,
4,
)]),
];
let ranges = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap();
assert_eq!(
prompt_token_ids,
vec![
1,
QWEN3_IMAGE_PAD_ID,
QWEN3_IMAGE_PAD_ID,
2,
QWEN3_VIDEO_PAD_ID,
QWEN3_VIDEO_PAD_ID,
QWEN3_VIDEO_PAD_ID,
QWEN3_VIDEO_PAD_ID,
3,
QWEN3_IMAGE_PAD_ID,
QWEN3_IMAGE_PAD_ID,
QWEN3_IMAGE_PAD_ID,
4,
]
);
let image_ranges = &ranges[&Modality::Image];
assert_eq!(image_ranges[0].offset, 1);
assert_eq!(image_ranges[0].length, 2);
assert_bool_mask(&image_ranges[0], &[true, true]);
assert_eq!(image_ranges[1].offset, 9);
assert_eq!(image_ranges[1].length, 3);
assert_bool_mask(&image_ranges[1], &[true, true, true]);
let video_ranges = &ranges[&Modality::Video];
assert_eq!(video_ranges[0].offset, 4);
assert_eq!(video_ranges[0].length, 4);
assert_bool_mask(&video_ranges[0], &[true, true, true, true]);
}
#[test]
fn expand_prompt_tokens_error_names_modality_with_leftover_replacements() {
let mut prompt_token_ids = vec![1, QWEN3_IMAGE_PAD_ID, 2];
let original_prompt_token_ids = prompt_token_ids.clone();
let prepared = vec![
qwen3_image_prepared(vec![PromptReplacement::repeated(
Modality::Image,
"<|image_pad|>",
QWEN3_IMAGE_PAD_ID as TokenId,
2,
)]),
qwen3_video_prepared(vec![PromptReplacement::repeated(
Modality::Video,
"<|video_pad|>",
QWEN3_VIDEO_PAD_ID as TokenId,
4,
)]),
];
let error = expand_prompt_token_ids(&mut prompt_token_ids, &prepared).unwrap_err();
assert!(matches!(
error,
Error::Multimodal(message)
if message.contains("<|video_pad|>") && message.contains("`video`")
));
assert_eq!(prompt_token_ids, original_prompt_token_ids);
}
}
+141
View File
@@ -0,0 +1,141 @@
//! Image-modality preparation: batch preprocessing and per-item feature
//! build.
use std::sync::Arc;
use itertools::izip;
use llm_multimodal::{FieldLayout, ImageFrame, Modality, PreprocessedEncoderInputs};
use vllm_engine_core_client::protocol::dtype::ModelDtype;
use vllm_engine_core_client::protocol::multimodal::{
MmBatchedField, MmField, MmFieldElem, MmFlatField, MmKwargsItem, MmSharedField, MmSlice,
SliceSpec,
};
use super::{ModalitySupport, MultimodalModelInfo, PreparedItem, PreparedMedia, tensor};
use crate::error::{Error, Result, bail_multimodal, multimodal};
impl MultimodalModelInfo {
/// Preprocess all fetched image frames as one batch and build per-item
/// features.
pub(super) async fn prepare_images(
&self,
frames: Vec<Arc<ImageFrame>>,
uuids: Vec<Option<String>>,
model_dtype: ModelDtype,
) -> Result<PreparedMedia> {
let support = self.image.as_ref().ok_or_else(|| Error::UnsupportedModality {
modality: Modality::Image.to_string(),
})?;
let preprocessed = self.preprocess_images(support, &frames).await?;
let replacements =
self.spec
.prompt_replacements_for(&self.context, &preprocessed, Modality::Image)?;
if replacements.len() != frames.len() {
bail_multimodal!(
"number of image prompt replacements {} does not match number of images {}",
replacements.len(),
frames.len()
);
}
let items = self.build_image_items(preprocessed, &frames, uuids, model_dtype)?;
Ok(PreparedMedia {
modality: Modality::Image,
placeholder: support.placeholder.clone(),
replacements,
items,
})
}
/// Preprocess fetched image frames with the model's resolved vision
/// processor.
///
/// The processor work is CPU-heavy relative to request wiring, so it runs
/// in a blocking task and returns owned tensors ready for wire
/// conversion.
async fn preprocess_images(
&self,
support: &ModalitySupport,
image_frames: &[Arc<ImageFrame>],
) -> Result<PreprocessedEncoderInputs> {
let config = support.config.clone();
let processor = support.processor;
let images = image_frames.iter().map(|frame| frame.data().clone()).collect::<Vec<_>>();
// TODO: is it still necessary given that we've already in a dedicated runtime?
tokio::task::spawn_blocking(move || Ok(processor.preprocess(&images, &config)?))
.await
.map_err(|error| multimodal!("image preprocessing task failed: {error}"))?
}
/// Convert one batch of preprocessed image tensors into per-item engine
/// kwargs.
///
/// Tensor fields are sliced per item according to the model spec's field
/// layout declarations.
fn build_image_items(
&self,
preprocessed: PreprocessedEncoderInputs,
frames: &[Arc<ImageFrame>],
uuids: Vec<Option<String>>,
model_dtype: ModelDtype,
) -> Result<Vec<PreparedItem>> {
let len = frames.len();
let tensors = tensor::collect_tensors(preprocessed, "pixel_values", model_dtype)?;
let mut items = Vec::with_capacity(len);
for (index, (frame, uuid)) in izip!(frames, uuids).enumerate() {
let mut data = MmKwargsItem::new();
for (key, tensor) in &tensors {
let keep_on_cpu = self.spec.keep_on_cpu_keys.contains(key);
let (value, field) = match self.spec.field_layouts.get(key) {
Some(FieldLayout::Batched) => (
tensor.batched_value_at(index)?,
MmField::Batched(MmBatchedField { keep_on_cpu }),
),
Some(FieldLayout::Flat { sizes_key }) => {
let sizes = tensors.get(sizes_key).ok_or_else(|| {
multimodal!("flat tensor sizes key `{sizes_key}` is missing")
})?;
let (start, end) = tensor::flat_range_for_index(sizes, sizes_key, index)?;
(
tensor.flat_value_range(start, end)?,
MmField::Flat(MmFlatField {
slices: vec![MmSlice::Slice(SliceSpec {
start: Some(0),
stop: Some((end - start) as isize),
step: None,
})],
dim: 0,
keep_on_cpu,
}),
)
}
None => (
tensor.clone(),
MmField::Shared(MmSharedField {
batch_size: len,
keep_on_cpu,
}),
),
};
data.insert(
key.clone(),
MmFieldElem {
data: Some(value.try_into()?),
field,
},
);
}
items.push(PreparedItem {
data,
hash: frame.hash.clone(),
uuid,
});
}
Ok(items)
}
}
+32 -14
View File
@@ -1,7 +1,7 @@
use std::collections::HashMap;
use half::{bf16, f16};
use llm_multimodal::{ModelSpecificValue, PreprocessedEncoderInputs as PreprocessedImages};
use llm_multimodal::{ModelSpecificValue, PreprocessedEncoderInputs};
use vllm_engine_core_client::protocol::dtype::ModelDtype;
use vllm_engine_core_client::protocol::multimodal::MmKwargValue as ProtocolKwargValue;
use vllm_engine_core_client::protocol::tensor::{ShapeExt as _, WireTensor};
@@ -25,25 +25,31 @@ pub(super) enum KwargValue {
Passthrough(ProtocolKwargValue),
}
/// Collect `pixel_values` and model-specific outputs into one tensor map.
/// Collect the primary encoder input and model-specific outputs into one
/// tensor map.
///
/// `primary_key` names the encoder-input tensor as the model's forward kwargs
/// expect it (e.g. `pixel_values` for images, `pixel_values_videos` for
/// videos).
pub(super) fn collect_tensors(
preprocessed: PreprocessedImages,
preprocessed: PreprocessedEncoderInputs,
primary_key: &str,
float_dtype: ModelDtype,
) -> Result<HashMap<String, KwargValue>> {
let PreprocessedImages {
let PreprocessedEncoderInputs {
encoder_input,
model_specific,
..
} = preprocessed;
let pixel_values = {
let primary_value = {
let shape = encoder_input.shape().to_vec();
let data = encoder_input.into_iter().collect();
KwargValue::from_f32_tensor(data, shape, float_dtype)?
};
let mut tensors = HashMap::new();
tensors.insert("pixel_values".to_string(), pixel_values);
tensors.insert(primary_key.to_string(), primary_value);
for (key, value) in model_specific {
tensors.insert(key, KwargValue::from_model_specific(value, float_dtype)?);
}
@@ -124,10 +130,22 @@ impl TryFrom<KwargValue> for ProtocolKwargValue {
}
impl KwargValue {
/// Extract one image from a batched tensor field.
/// First-axis length for tensor values; `None` for passthrough kwargs.
pub(super) fn first_dim(&self) -> Option<usize> {
match self {
Self::F32Tensor { shape, .. }
| Self::F16Tensor { shape, .. }
| Self::Bf16Tensor { shape, .. }
| Self::I64Tensor { shape, .. }
| Self::U32Tensor { shape, .. } => shape.first().copied(),
Self::Passthrough(_) => None,
}
}
/// Extract one media item from a batched tensor field.
///
/// Batched fields use their first axis as image index and drop that axis in
/// the per-feature value, matching vLLM's batched-field semantics.
/// Batched fields use their first axis as media-item index and drop that
/// axis in the per-feature value, matching vLLM's batched-field semantics.
pub(super) fn batched_value_at(&self, index: usize) -> Result<Self> {
match self {
Self::F32Tensor { data, shape } => {
@@ -154,9 +172,9 @@ impl KwargValue {
}
}
/// Extract one image's variable-length range from a flat tensor field.
/// Extract one media item's variable-length range from a flat tensor field.
///
/// Flat fields keep the first axis as the sliced length for this image.
/// Flat fields keep the first axis as the sliced length for this item.
pub(super) fn flat_value_range(&self, start: usize, end: usize) -> Result<Self> {
match self {
Self::F32Tensor { data, shape } => {
@@ -184,10 +202,10 @@ impl KwargValue {
}
}
/// Compute the first-axis range for one image in a flat tensor.
/// Compute the first-axis range for one media item in a flat tensor.
///
/// `sizes_key` names a companion tensor whose entries are cumulative slice
/// sizes per image.
/// sizes per media item.
pub(super) fn flat_range_for_index(
sizes: &KwargValue,
sizes_key: &str,
@@ -195,7 +213,7 @@ pub(super) fn flat_range_for_index(
) -> Result<(usize, usize)> {
let sizes = tensor_as_usize_vec(sizes)?;
let size = *sizes.get(index).ok_or_else(|| {
multimodal!("flat tensor sizes key `{sizes_key}` has no entry for image {index}")
multimodal!("flat tensor sizes key `{sizes_key}` has no entry for media item {index}")
})?;
let start = sizes[..index].iter().sum::<usize>();
Ok((start, start + size))
+316
View File
@@ -0,0 +1,316 @@
//! Video-modality preparation: per-clip preprocessing, config resolution,
//! and per-item feature build.
use std::sync::Arc;
use itertools::izip;
use llm_multimodal::{FieldLayout, Modality, PreprocessedEncoderInputs, VideoClip};
use thiserror_ext::AsReport as _;
use tracing::warn;
use vllm_engine_core_client::protocol::dtype::ModelDtype;
use vllm_engine_core_client::protocol::multimodal::{
MmBatchedField, MmField, MmFieldElem, MmFlatField, MmKwargsItem, MmSharedField, MmSlice,
SliceSpec,
};
use super::{ModalitySupport, MultimodalModelInfo, PreparedItem, PreparedMedia, tensor};
use crate::error::{Error, Result, bail_multimodal, multimodal};
/// Forward-kwargs name of the primary video encoder input.
///
/// Video-capable vLLM models read `pixel_values_videos` alongside
/// `video_grid_thw`, mirroring the HF processor output naming.
const VIDEO_PRIMARY_KEY: &str = "pixel_values_videos";
impl MultimodalModelInfo {
/// Preprocess fetched video clips one at a time and build per-item
/// features.
///
/// Unlike images, each clip runs through the preprocessor independently
/// (a batch of one), so its tensors are complete per item and need no
/// cross-item slicing.
pub(super) async fn prepare_videos(
&self,
clips: Vec<Arc<VideoClip>>,
uuids: Vec<Option<String>>,
model_dtype: ModelDtype,
) -> Result<PreparedMedia> {
let support = self.video.as_ref().ok_or_else(|| Error::UnsupportedModality {
modality: Modality::Video.to_string(),
})?;
let mut replacements = Vec::with_capacity(clips.len());
let mut items = Vec::with_capacity(clips.len());
for (clip, uuid) in izip!(&clips, uuids) {
let preprocessed = self.preprocess_video_clip(support, Arc::clone(clip)).await?;
let mut clip_replacements =
self.spec
.prompt_replacements_for(&self.context, &preprocessed, Modality::Video)?;
if clip_replacements.len() != 1 {
bail_multimodal!(
"expected exactly one prompt replacement per video clip, got {}",
clip_replacements.len()
);
}
replacements.push(clip_replacements.pop().unwrap());
items.push(self.build_video_item(
preprocessed,
clip.hash.clone(),
uuid,
model_dtype,
)?);
}
Ok(PreparedMedia {
modality: Modality::Video,
placeholder: support.placeholder.clone(),
replacements,
items,
})
}
/// Preprocess one decoded video clip with the model's resolved vision
/// processor.
async fn preprocess_video_clip(
&self,
support: &ModalitySupport,
clip: Arc<VideoClip>,
) -> Result<PreprocessedEncoderInputs> {
let config = support.config.clone();
let processor = support.processor;
tokio::task::spawn_blocking(move || {
// Prefer the borrowed-RGB fast path, which avoids materializing a
// `DynamicImage` per sampled frame after media decode.
if let Some(rgb_video) = clip.rgb_video() {
match rgb_video.frame_refs() {
Ok(frame_refs) => match processor.preprocess_video_rgb(&frame_refs, &config) {
Ok(preprocessed) => return Ok(preprocessed),
Err(error) => warn!(
error = %error.as_report(),
"RGB video preprocessing fast path failed; falling back to materialized frames"
),
},
Err(error) => warn!(
error,
"RGB video frame refs are invalid; falling back to materialized frames"
),
}
}
let frames = clip.materialized_frames().map_err(|error| multimodal!("{error}"))?;
Ok(processor.preprocess_video(&frames, &config)?)
})
.await
.map_err(|error| multimodal!("video preprocessing task failed: {error}"))?
}
/// Convert one preprocessed video clip into engine kwargs.
///
/// The clip is a batch of one, so no per-item slicing is required: the
/// primary tensor ships as a full-range flat field (the engine re-batches
/// flat fields by concatenating along the declared dim, matching vLLM's
/// `flat_from_sizes` treatment of video patches), and batched metadata
/// tensors drop their singleton batch axis.
fn build_video_item(
&self,
preprocessed: PreprocessedEncoderInputs,
hash: String,
uuid: Option<String>,
model_dtype: ModelDtype,
) -> Result<PreparedItem> {
let tensors = tensor::collect_tensors(preprocessed, VIDEO_PRIMARY_KEY, model_dtype)?;
let mut data = MmKwargsItem::new();
for (key, tensor) in tensors {
let keep_on_cpu = self.spec.keep_on_cpu_keys.contains(&key);
let (value, field) = if key == VIDEO_PRIMARY_KEY {
let len = tensor
.first_dim()
.ok_or_else(|| multimodal!("video encoder input `{key}` is not a tensor"))?;
(
tensor,
MmField::Flat(MmFlatField {
slices: vec![MmSlice::Slice(SliceSpec {
start: Some(0),
stop: Some(len as isize),
step: None,
})],
dim: 0,
keep_on_cpu,
}),
)
} else if matches!(
self.spec.field_layouts.get(&key),
Some(FieldLayout::Batched)
) {
(
tensor.batched_value_at(0)?,
MmField::Batched(MmBatchedField { keep_on_cpu }),
)
} else {
(
tensor,
MmField::Shared(MmSharedField {
batch_size: 1,
keep_on_cpu,
}),
)
};
data.insert(
key,
MmFieldElem {
data: Some(value.try_into()?),
field,
},
);
}
Ok(PreparedItem { data, hash, uuid })
}
}
#[cfg(test)]
mod tests {
use std::collections::HashMap;
use std::sync::Arc;
use llm_multimodal::ModelSpecificValue;
use ndarray::ArrayD;
use vllm_engine_core_client::protocol::multimodal::MmKwargValue;
use super::super::tests::{
QWEN3_IMAGE_PAD_ID, QWEN3_VIDEO_PAD_ID, qwen3_vl_info, qwen3_vl_tokenizer,
};
use super::super::{MultimodalConfigFiles, MultimodalModelInfo};
use super::*;
#[test]
fn from_paths_resolves_video_config_from_dedicated_file_or_processor_config() {
let dir = tempfile::tempdir().unwrap();
let config_path = dir.path().join("config.json");
std::fs::write(
&config_path,
serde_json::json!({
"model_type": "qwen3_vl",
"image_token_id": QWEN3_IMAGE_PAD_ID,
"video_token_id": QWEN3_VIDEO_PAD_ID,
})
.to_string(),
)
.unwrap();
let info_for = |files: MultimodalConfigFiles<'_>| {
MultimodalModelInfo::from_paths(
"qwen3-vl-test".to_string(),
Some("qwen3_vl".to_string()),
files,
Arc::new(qwen3_vl_tokenizer()),
)
};
// Dedicated video preprocessor config file.
let video_config_path = dir.path().join("video_preprocessor_config.json");
std::fs::write(&video_config_path, r#"{"size":{"shortest_edge":128}}"#).unwrap();
let info = info_for(MultimodalConfigFiles {
config: Some(&config_path),
video_preprocessor_config: Some(&video_config_path),
..Default::default()
})
.unwrap()
.unwrap();
assert!(info.video.is_some());
// `video_processor` section of the combined processor config.
let processor_config_path = dir.path().join("processor_config.json");
std::fs::write(
&processor_config_path,
r#"{"video_processor":{"size":{"shortest_edge":128}}}"#,
)
.unwrap();
let info = info_for(MultimodalConfigFiles {
config: Some(&config_path),
processor_config: Some(&processor_config_path),
..Default::default()
})
.unwrap()
.unwrap();
assert!(info.video.is_some());
// Neither source: video support still resolves on the image config.
let info = info_for(MultimodalConfigFiles {
config: Some(&config_path),
..Default::default()
})
.unwrap()
.unwrap();
assert!(info.video.is_some());
// Malformed dedicated file is a real error, not a silent fallback.
std::fs::write(&video_config_path, r#"{"size""#).unwrap();
let error = match info_for(MultimodalConfigFiles {
config: Some(&config_path),
video_preprocessor_config: Some(&video_config_path),
..Default::default()
}) {
Err(error) => error,
Ok(_) => panic!("malformed video preprocessor config should fail"),
};
assert!(matches!(
error,
Error::Multimodal(message)
if message.contains("failed to parse video_preprocessor_config.json")
));
}
#[test]
fn build_video_item_names_primary_tensor_and_layouts() {
let info = qwen3_vl_info();
// One clip flattened to 6 patches with 4 features each.
let preprocessed = PreprocessedEncoderInputs {
encoder_input: ArrayD::zeros(vec![6, 4]),
feature_token_counts: vec![6],
item_sizes: vec![(32, 32)],
model_specific: HashMap::from([
(
"video_grid_thw".to_string(),
ModelSpecificValue::int_2d(vec![1, 2, 3], 1, 3),
),
(
"patches_per_video".to_string(),
ModelSpecificValue::int_1d(vec![6]),
),
]),
};
let item = info
.build_video_item(
preprocessed,
"<hash>".to_string(),
None,
ModelDtype::Float32,
)
.unwrap();
let primary = &item.data[VIDEO_PRIMARY_KEY];
assert!(matches!(
&primary.field,
MmField::Flat(MmFlatField { slices, dim: 0, .. })
if matches!(
slices.as_slice(),
[MmSlice::Slice(SliceSpec { start: Some(0), stop: Some(6), step: None })]
)
));
// Batched metadata drops its singleton batch axis per item.
let grid = &item.data["video_grid_thw"];
assert!(matches!(&grid.field, MmField::Batched(_)));
let MmKwargValue::Tensor(grid_tensor) = grid.data.as_ref().unwrap() else {
panic!("expected tensor value for video_grid_thw");
};
assert_eq!(grid_tensor.shape, vec![3]);
assert_eq!(item.hash, "<hash>");
}
}
+83 -8
View File
@@ -31,9 +31,14 @@ pub use template::{load_chat_template, resolve_chat_template};
pub use self::format::ChatTemplateContentFormatOption;
#[derive(Debug, Clone)]
/// Template-visible placeholder tokens per supported modality.
///
/// A `None` token means the loaded model does not support that modality, and
/// content parts of that modality are rejected during rendering.
#[derive(Debug, Clone, Default)]
pub struct MultimodalRenderInfo {
pub placeholder_token: String,
pub image_token: Option<String>,
pub video_token: Option<String>,
}
/// Hugging Face chat-template renderer backed by the local Jinja chat-template
@@ -254,6 +259,7 @@ enum TemplateContent {
enum TemplateContentPart {
Text { text: String },
Image,
Video,
}
#[derive(Debug, Serialize)]
@@ -417,9 +423,17 @@ fn to_template_openai_content(
}
// All multimodal contents are normalized to `{ "type": <modality> }`.
ChatContentPart::ImageUrl { .. } => {
multimodal.ok_or(Error::UnsupportedMultimodalContent("image_url"))?;
multimodal
.and_then(|multimodal| multimodal.image_token.as_ref())
.ok_or(Error::UnsupportedMultimodalContent("image_url"))?;
Ok(TemplateContentPart::Image)
}
ChatContentPart::VideoUrl { .. } => {
multimodal
.and_then(|multimodal| multimodal.video_token.as_ref())
.ok_or(Error::UnsupportedMultimodalContent("video_url"))?;
Ok(TemplateContentPart::Video)
}
})
.collect(),
}
@@ -437,9 +451,16 @@ fn to_template_string_content(
match part {
ChatContentPart::Text { text } => out.push_str(text),
ChatContentPart::ImageUrl { .. } => {
let multimodal =
multimodal.ok_or(Error::UnsupportedMultimodalContent("image_url"))?;
out.push_str(&multimodal.placeholder_token);
let image_token = multimodal
.and_then(|multimodal| multimodal.image_token.as_ref())
.ok_or(Error::UnsupportedMultimodalContent("image_url"))?;
out.push_str(image_token);
}
ChatContentPart::VideoUrl { .. } => {
let video_token = multimodal
.and_then(|multimodal| multimodal.video_token.as_ref())
.ok_or(Error::UnsupportedMultimodalContent("video_url"))?;
out.push_str(video_token);
}
}
}
@@ -468,7 +489,7 @@ fn append_continue_final_message_tag(message: &mut TemplateMessage) -> Result<St
// Pick the last text part in the message.
TemplateContent::OpenAi(parts) => parts.iter_mut().rev().find_map(|part| match part {
TemplateContentPart::Text { text } => Some(text),
TemplateContentPart::Image => None,
TemplateContentPart::Image | TemplateContentPart::Video => None,
}),
};
let text = text.ok_or_else(|| {
@@ -577,7 +598,8 @@ mod tests {
) -> Result<crate::RenderedPrompt> {
HfChatRenderer::new(Some(template.to_string()), HashMap::new(), content_format)?
.with_multimodal(Some(MultimodalRenderInfo {
placeholder_token: "<image>".to_string(),
image_token: Some("<image>".to_string()),
video_token: Some("<video>".to_string()),
}))
.render(request)
}
@@ -590,6 +612,14 @@ mod tests {
])])
}
fn video_request() -> ChatRequest {
sample_request(vec![ChatMessage::user(vec![
ChatContentPart::text("a"),
ChatContentPart::video_url("https://example.com/demo.mp4"),
ChatContentPart::text("b"),
])])
}
#[test]
fn string_content_format_replaces_image_with_placeholder_text() {
let rendered = render_mm(
@@ -602,6 +632,18 @@ mod tests {
assert_eq!(rendered.prompt, Prompt::Text("a<image>b".to_string()));
}
#[test]
fn string_content_format_replaces_video_with_placeholder_text() {
let rendered = render_mm(
"{{ messages[0].content }}",
&video_request(),
ChatTemplateContentFormatOption::String,
)
.unwrap();
assert_eq!(rendered.prompt, Prompt::Text("a<video>b".to_string()));
}
#[test]
fn openai_content_format_normalizes_image_url_for_template() {
let rendered = render_mm(
@@ -614,6 +656,39 @@ mod tests {
assert_eq!(rendered.prompt, Prompt::Text("a<|image_pad|>b".to_string()));
}
#[test]
fn openai_content_format_normalizes_video_url_for_template() {
let rendered = render_mm(
"{% for item in messages[0].content %}{% if item.type == 'video' %}<|video_pad|>{% else %}{{ item.text }}{% endif %}{% endfor %}",
&video_request(),
ChatTemplateContentFormatOption::OpenAi,
)
.unwrap();
assert_eq!(rendered.prompt, Prompt::Text("a<|video_pad|>b".to_string()));
}
#[test]
fn video_parts_are_rejected_when_model_lacks_video_support() {
let error = HfChatRenderer::new(
Some("{{ messages[0].content }}".to_string()),
HashMap::new(),
ChatTemplateContentFormatOption::String,
)
.unwrap()
.with_multimodal(Some(MultimodalRenderInfo {
image_token: Some("<image>".to_string()),
video_token: None,
}))
.render(&video_request())
.unwrap_err();
assert!(matches!(
error,
Error::UnsupportedMultimodalContent("video_url")
));
}
#[test]
fn chat_template_supports_pycompat_templates() {
let request = sample_request(vec![ChatMessage::text(ChatRole::User, "<think>hello")]);
+36 -1
View File
@@ -36,7 +36,13 @@ pub enum ChatContentPart {
detail: Option<ImageDetail>,
uuid: Option<String>,
},
/// One video URL/data URL content block.
VideoUrl {
video_url: String,
uuid: Option<String>,
},
// ImageData...
// VideoData...
// ImageEmbeds...
}
@@ -55,12 +61,21 @@ impl ChatContentPart {
}
}
/// Construct one video URL content part with the given URL string.
pub fn video_url(video_url: impl Into<String>) -> Self {
Self::VideoUrl {
video_url: video_url.into(),
uuid: None,
}
}
/// Return the text content of this part when it's a text block, or an
/// "unsupported multimodal content" error otherwise.
pub(crate) fn as_text(&self) -> Result<&str> {
match self {
Self::Text { text } => Ok(text),
Self::ImageUrl { .. } => Err(Error::UnsupportedMultimodalContent("image_url")),
Self::VideoUrl { .. } => Err(Error::UnsupportedMultimodalContent("video_url")),
}
}
@@ -73,7 +88,7 @@ impl ChatContentPart {
pub(crate) fn is_multimodal(&self) -> bool {
match self {
Self::Text { .. } => false,
Self::ImageUrl { .. } => true,
Self::ImageUrl { .. } | Self::VideoUrl { .. } => true,
}
}
}
@@ -555,6 +570,26 @@ mod tests {
assert_eq!(content, ChatContent::Text("hello".to_string()));
}
#[test]
fn chat_content_video_url_part_round_trips_through_serde() {
let content = ChatContent::Parts(vec![ChatContentPart::VideoUrl {
video_url: "https://example.com/demo.mp4".to_string(),
uuid: Some("video-1".to_string()),
}]);
let value = to_value(&content).unwrap();
assert_eq!(
value,
json!([{
"type": "video_url",
"video_url": "https://example.com/demo.mp4",
"uuid": "video-1",
}])
);
let decoded: ChatContent = serde_json::from_value(value).unwrap();
assert_eq!(decoded, content);
}
#[test]
fn chat_content_deserializes_from_openai_text_blocks() {
let content: ChatContent =
+42
View File
@@ -117,6 +117,19 @@ impl RoundtripCase {
}
}
/// MiniMax M3 invoke format with `<mm:think>` reasoning tags.
fn minimax_m3() -> Self {
Self {
model_id: "MiniMaxAI/MiniMax-M3",
assistant_stop_suffix: "[e~[\n",
tool_call_parser: ParserSelection::Auto,
reasoning_parser: ParserSelection::Auto,
thinking_behavior: ThinkingBehavior::Always { value: true },
json_fmt: compact_json_fmt(),
sort_json_keys: false,
}
}
/// DeepSeek V4 DSML tool-call format.
fn deepseek_v4() -> Self {
Self {
@@ -143,6 +156,19 @@ impl RoundtripCase {
}
}
/// GLM-4.5 XML-like argument format with `<think>` reasoning tags.
fn glm45() -> Self {
Self {
model_id: "zai-org/GLM-4.5",
assistant_stop_suffix: "",
tool_call_parser: ParserSelection::Auto,
reasoning_parser: ParserSelection::Auto,
thinking_behavior: ThinkingBehavior::Toggleable { default: true },
json_fmt: compact_json_fmt(),
sort_json_keys: false,
}
}
/// GLM-4.7 XML-like argument format with `<think>` reasoning tags.
fn glm47() -> Self {
Self {
@@ -209,6 +235,19 @@ impl RoundtripCase {
}
}
/// Nemotron V3 with `<think>` / `</think>` reasoning tags.
fn nemotron_v3() -> Self {
Self {
model_id: "nvidia/NVIDIA-Nemotron-3-Nano-30B-A3B-BF16",
assistant_stop_suffix: "<|im_end|>\n",
tool_call_parser: ParserSelection::Auto,
reasoning_parser: ParserSelection::Auto,
thinking_behavior: ThinkingBehavior::Always { value: true },
json_fmt: compact_json_fmt(),
sort_json_keys: false,
}
}
/// GPT-OSS Harmony token-id renderer and native Harmony output processor.
fn gpt_oss() -> Self {
Self {
@@ -247,11 +286,14 @@ roundtrip_tests! {
qwen3 => [reasoning_and_content, tool_call_mix],
qwen35 => [reasoning_and_content, tool_call_mix],
minimax_m25 => [reasoning_and_content, tool_call_mix],
minimax_m3 => [reasoning_and_content, tool_call_mix],
deepseek_v4 => [reasoning_and_content, tool_call_mix],
deepseek_v32 => [tool_call_mix],
glm45 => [reasoning_and_content, tool_call_mix],
glm47 => [reasoning_and_content, tool_call_mix],
seed_oss => [reasoning_and_content],
step3p5 => [reasoning_and_content],
nemotron_v3 => [reasoning_and_content],
gemma4 => [tool_call_mix], // Gemma4 strips reasoning in history if there's no tool call
kimi_k25 => [tool_call_mix], // Kimi K2.5 strips reasoning in history
gpt_oss => [tool_call_mix], // Harmony strips reasoning in history if there's no tool call
@@ -290,7 +290,10 @@ fn convert_content(content: MessageContent) -> Result<ChatContent, ApiError> {
detail: image_url.detail,
uuid,
}),
_ => bail_invalid_request!("Only text and image_url content parts are supported."),
ContentPart::VideoUrl { video_url, uuid } => Ok(ChatContentPart::VideoUrl {
video_url: video_url.url,
uuid,
}),
})
.try_collect()
.map(ChatContent::Parts),
@@ -772,28 +775,34 @@ mod tests {
}
#[test]
fn prepare_chat_request_rejects_video_content_parts() {
fn prepare_chat_request_accepts_video_content_parts() {
let request = ChatCompletionRequest {
messages: vec![ChatMessage::User {
content: MessageContent::Parts(vec![ContentPart::VideoUrl {
video_url: VideoUrl {
url: "https://example.com/video.mp4".to_string(),
},
uuid: Some("video-uuid".to_string()),
}]),
name: None,
}],
..base_request()
};
let error = prepare_chat_request(
let prepared = prepare_chat_request(
request,
&served(&["Qwen/Qwen1.5-0.5B-Chat"]),
ResolvedRequestContext::default(),
)
.unwrap_err();
.expect("request is valid");
expect!["Only text and image_url content parts are supported."]
.assert_eq(&error.to_error_response().error.message);
assert_eq!(
prepared.chat_request.messages,
vec![VllmChatMessage::user(vec![ChatContentPart::VideoUrl {
video_url: "https://example.com/video.mp4".to_string(),
uuid: Some("video-uuid".to_string()),
}])]
);
}
#[test]
@@ -91,7 +91,11 @@ pub enum ContentPart {
uuid: Option<String>,
},
#[serde(rename = "video_url")]
VideoUrl { video_url: VideoUrl },
VideoUrl {
video_url: VideoUrl,
#[serde(skip_serializing_if = "Option::is_none")]
uuid: Option<String>,
},
}
#[serde_with::skip_serializing_none]
+8 -4
View File
@@ -504,7 +504,7 @@ impl ChatRenderer for FakeChatBackend {
let placeholder = self
.multimodal_model_info
.as_ref()
.map(|info| info.placeholder_token())
.and_then(|info| info.placeholder_token(llm_multimodal::Modality::Image))
.unwrap_or("<image>");
let mut prompt = String::new();
for message in &request.messages {
@@ -544,7 +544,9 @@ fn render_fake_content(content: &ChatContent, placeholder: &str) -> vllm_chat::R
for part in parts {
match part {
ChatContentPart::Text { text } => out.push_str(text),
ChatContentPart::ImageUrl { .. } => out.push_str(placeholder),
ChatContentPart::ImageUrl { .. } | ChatContentPart::VideoUrl { .. } => {
out.push_str(placeholder)
}
}
}
out
@@ -565,8 +567,10 @@ fn qwen_multimodal_model_info() -> vllm_chat::multimodal::MultimodalModelInfo {
let info = vllm_chat::multimodal::MultimodalModelInfo::from_paths(
"qwen2-vl-test".to_string(),
Some("qwen2_vl".to_string()),
Some(&config_path),
None,
vllm_chat::multimodal::MultimodalConfigFiles {
config: Some(&config_path),
..Default::default()
},
Arc::new(fake_chat_tokenizer()),
)
.expect("load multimodal info")
@@ -41,6 +41,10 @@ pub struct ResolvedModelFiles {
pub tokenizer_config_path: Option<PathBuf>,
pub generation_config_path: Option<PathBuf>,
pub preprocessor_config_path: Option<PathBuf>,
/// Video-specific preprocessor config, when provided by the model repo.
pub video_preprocessor_config_path: Option<PathBuf>,
/// Combined processor config, which may embed a `video_processor` section.
pub processor_config_path: Option<PathBuf>,
pub chat_template_path: Option<PathBuf>,
pub config_path: Option<PathBuf>,
}
@@ -70,6 +74,11 @@ fn resolve_local_model_files(model_dir: &Path) -> Result<ResolvedModelFiles> {
tokenizer_config_path,
generation_config_path: local_file_if_exists(model_dir, "generation_config.json"),
preprocessor_config_path: local_file_if_exists(model_dir, "preprocessor_config.json"),
video_preprocessor_config_path: local_file_if_exists(
model_dir,
"video_preprocessor_config.json",
),
processor_config_path: local_file_if_exists(model_dir, "processor_config.json"),
chat_template_path: discover_chat_template_in_dir(model_dir),
config_path: local_file_if_exists(model_dir, "config.json"),
})
@@ -107,6 +116,10 @@ async fn resolve_remote_model_files(model_id: &str) -> Result<ResolvedModelFiles
download_if_present(&repo, model_id, &siblings, "generation_config.json").await?;
let preprocessor_config_path =
download_if_present(&repo, model_id, &siblings, "preprocessor_config.json").await?;
let video_preprocessor_config_path =
download_if_present(&repo, model_id, &siblings, "video_preprocessor_config.json").await?;
let processor_config_path =
download_if_present(&repo, model_id, &siblings, "processor_config.json").await?;
let chat_template_name = siblings
.contains("chat_template.json")
.then_some("chat_template.json")
@@ -123,6 +136,8 @@ async fn resolve_remote_model_files(model_id: &str) -> Result<ResolvedModelFiles
tokenizer_config_path,
generation_config_path,
preprocessor_config_path,
video_preprocessor_config_path,
processor_config_path,
chat_template_path,
config_path,
})
@@ -143,6 +158,8 @@ fn resolve_cached_model_files(model_id: &str) -> Result<Option<ResolvedModelFile
})?;
let generation_config_path = cache_repo.get("generation_config.json");
let preprocessor_config_path = cache_repo.get("preprocessor_config.json");
let video_preprocessor_config_path = cache_repo.get("video_preprocessor_config.json");
let processor_config_path = cache_repo.get("processor_config.json");
let chat_template_path = discover_chat_template_in_dir(model_dir);
let config_path = cache_repo.get("config.json");
@@ -151,6 +168,8 @@ fn resolve_cached_model_files(model_id: &str) -> Result<Option<ResolvedModelFile
tokenizer_config_path,
generation_config_path,
preprocessor_config_path,
video_preprocessor_config_path,
processor_config_path,
chat_template_path,
config_path,
}))
+3
View File
@@ -1257,6 +1257,9 @@ setup(
"mistral_common[audio]",
], # Required for audio processing
"video": [], # Kept for backwards compatibility
# NVIDIA DeepStream (NVDEC) GPU video-decode backend. Linux x86-64
# only; also needs system GStreamer + libv4l (see docs).
"deepstream": ["nvidia-deepstream-videodecode-cu13>=9.0.2"],
"flashinfer": [], # Kept for backwards compatibility
# Optional deps for Helion kernel development
# NOTE: When updating helion version, also update CI files:
@@ -0,0 +1,158 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
"""E2E tests for ``include_reasoning`` with non-Harmony reasoning models.
Verifies that reasoning content is included by default and suppressed
when ``include_reasoning=False``, for both streaming and non-streaming
Chat Completions.
"""
import openai
import pytest
import pytest_asyncio
from tests.utils import RemoteOpenAIServer
MODEL_NAME = "Qwen/Qwen3-0.6B"
MESSAGES = [{"role": "user", "content": "What is 1+1? Be concise."}]
@pytest.fixture(scope="module")
def server():
args = [
"--reasoning-parser",
"qwen3",
"--max-model-len",
"2048",
"--enforce-eager",
"--gpu-memory-utilization",
"0.4",
]
with RemoteOpenAIServer(MODEL_NAME, args) as remote_server:
yield remote_server
@pytest_asyncio.fixture
async def client(server):
async with server.get_async_client() as async_client:
yield async_client
@pytest.mark.asyncio
async def test_include_reasoning_true_non_streaming(client: openai.AsyncOpenAI):
"""Default: reasoning content appears in non-streaming response."""
response = await client.chat.completions.create(
model=MODEL_NAME,
messages=MESSAGES,
max_tokens=200,
extra_body={"include_reasoning": True},
)
msg = response.choices[0].message
reasoning = getattr(msg, "reasoning", None) or getattr(
msg, "reasoning_content", None
)
assert reasoning, "Expected reasoning content when include_reasoning=True"
assert msg.content, "Expected content in response"
@pytest.mark.asyncio
async def test_include_reasoning_false_non_streaming(client: openai.AsyncOpenAI):
"""Reasoning content is suppressed when include_reasoning=False."""
response = await client.chat.completions.create(
model=MODEL_NAME,
messages=MESSAGES,
max_tokens=200,
extra_body={"include_reasoning": False},
)
msg = response.choices[0].message
reasoning = getattr(msg, "reasoning", None) or getattr(
msg, "reasoning_content", None
)
assert not reasoning, (
f"Expected no reasoning when include_reasoning=False, got: {reasoning}"
)
assert msg.content, "Expected content in response even without reasoning"
@pytest.mark.asyncio
async def test_include_reasoning_true_streaming(client: openai.AsyncOpenAI):
"""Default: reasoning deltas appear in streaming response."""
stream = await client.chat.completions.create(
model=MODEL_NAME,
messages=MESSAGES,
max_tokens=200,
stream=True,
extra_body={"include_reasoning": True},
)
reasoning_parts = []
content_parts = []
async for chunk in stream:
delta = chunk.choices[0].delta if chunk.choices else None
if delta:
r = getattr(delta, "reasoning", None) or getattr(
delta, "reasoning_content", None
)
if r:
reasoning_parts.append(r)
if delta.content:
content_parts.append(delta.content)
reasoning_text = "".join(reasoning_parts)
content_text = "".join(content_parts)
assert reasoning_text, "Expected reasoning deltas when include_reasoning=True"
assert content_text, "Expected content deltas in streaming response"
@pytest.mark.asyncio
async def test_include_reasoning_false_streaming(client: openai.AsyncOpenAI):
"""Reasoning deltas are suppressed in streaming when include_reasoning=False."""
stream = await client.chat.completions.create(
model=MODEL_NAME,
messages=MESSAGES,
max_tokens=200,
stream=True,
extra_body={"include_reasoning": False},
)
reasoning_parts = []
content_parts = []
async for chunk in stream:
delta = chunk.choices[0].delta if chunk.choices else None
if delta:
r = getattr(delta, "reasoning", None) or getattr(
delta, "reasoning_content", None
)
if r:
reasoning_parts.append(r)
if delta.content:
content_parts.append(delta.content)
reasoning_text = "".join(reasoning_parts)
content_text = "".join(content_parts)
assert not reasoning_text, (
f"Expected no reasoning deltas when include_reasoning=False, "
f"got: {reasoning_text[:100]}"
)
assert content_text, "Expected content deltas even without reasoning"
@pytest.mark.asyncio
async def test_default_includes_reasoning(client: openai.AsyncOpenAI):
"""Without specifying include_reasoning, reasoning appears (default=True)."""
response = await client.chat.completions.create(
model=MODEL_NAME,
messages=MESSAGES,
max_tokens=200,
)
msg = response.choices[0].message
reasoning = getattr(msg, "reasoning", None) or getattr(
msg, "reasoning_content", None
)
assert reasoning, "Expected reasoning content by default"
@@ -197,46 +197,6 @@ def _ragged_from_rows(
)
def _ref_combine_topk_swa_ragged(
device: torch.device,
) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]:
expected_ragged = torch.tensor(
[
100,
101,
7,
8,
9,
110,
111,
8,
9,
10,
120,
121,
122,
9,
10,
11,
150,
27,
28,
29,
160,
161,
28,
29,
30,
],
dtype=torch.int32,
device=device,
)
expected_lens = torch.tensor([5, 5, 6, 4, 5], dtype=torch.int32, device=device)
expected_indptr = torch.zeros(6, dtype=torch.int32, device=device)
torch.cumsum(expected_lens, dim=0, out=expected_indptr[1:])
return expected_ragged, expected_indptr, expected_lens
@torch.inference_mode()
def test_compute_global_topk_ragged_indices_and_indptr() -> None:
from vllm.models.deepseek_v4.amd.rocm import (
@@ -369,55 +329,6 @@ def test_sparse_attn_decode_ragged_kernel() -> None:
torch.testing.assert_close(actual, expected, atol=2e-2, rtol=2e-2)
@torch.inference_mode()
def test_combine_topk_swa_indices_ragged() -> None:
from vllm.models.deepseek_v4.amd.rocm import (
combine_topk_swa_indices_ragged,
)
device = torch.device("cuda")
topk_indices = torch.tensor(
[
[100, 101, 102, 103],
[110, 111, 112, 113],
[120, 121, 122, 123],
[130, 131, 132, 133],
[140, 141, 142, 143],
],
dtype=torch.int32,
device=device,
)
query_start_loc = torch.tensor([0, 3, 5], dtype=torch.int32, device=device)
seq_lens = torch.tensor([6, 4], dtype=torch.int32, device=device)
gather_lens = torch.tensor([4, 3], dtype=torch.int32, device=device)
window_size = 3
compress_ratio = 2
topk = 4
M = 20
N = 8
actual_ragged, actual_indptr, actual_lens = combine_topk_swa_indices_ragged(
topk_indices,
query_start_loc,
seq_lens,
gather_lens,
window_size,
compress_ratio,
topk,
M,
N,
)
expected_ragged, expected_indptr, expected_lens = _ref_combine_topk_swa_ragged(
device
)
torch.testing.assert_close(
actual_ragged[: expected_ragged.numel()], expected_ragged
)
torch.testing.assert_close(actual_indptr, expected_indptr)
torch.testing.assert_close(actual_lens, expected_lens)
@requires_gfx950
@torch.inference_mode()
def test_decode_num_splits_heuristic(monkeypatch) -> None:
+117
View File
@@ -0,0 +1,117 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
"""LongCat n-gram embedding id computation vs a pure-Python reference.
Guards the hash-id semantics of ``ngram_compute_n_gram_ids`` plus the
EOS-position fixup (``compute_eos_position_ngram_ids``): an EOS *current*
token hashes with its full look-back, while later positions' look-back stops
at the EOS boundary (LongCat reference behavior).
"""
from types import SimpleNamespace
import pytest
import torch
from vllm import _custom_ops as ops
from vllm.model_executor.models.longcat_flash_ngram import (
compute_eos_position_ngram_ids,
)
from vllm.platforms import current_platform
VOCAB = 163840
N, K = 5, 4 # oe_neighbor_num, oe_split_num (LongCat-2.0)
M = int(100.567 * VOCAB)
EOS = 2
NUM_EMB = K * (N - 1)
SIZES = [M + c * 2 + 1 for c in range(NUM_EMB)]
OFFSETS = [0]
for s in SIZES:
OFFSETS.append(OFFSETS[-1] + s)
def _ref_ids(tokens: list[int]) -> list[list[int]]:
"""Reference n-gram ids on raw tokens (HF LongCat semantics)."""
out = []
for pos in range(len(tokens)):
row = []
for i in range(N - 1):
n_real = i + 2
for j in range(K):
cfg = i * K + j
mod = SIZES[cfg]
h = 0
for delta in range(n_real):
p = pos - delta
if p < 0:
break
t = tokens[p]
if delta > 0 and t == EOS:
break # look-back stops at an EOS boundary
h += (t * pow(VOCAB, delta, mod)) % mod
row.append(h % mod + OFFSETS[cfg])
out.append(row)
return out
@pytest.mark.skipif(not current_platform.is_cuda(), reason="CUDA kernel")
def test_ngram_ids_match_reference_with_eos():
device = "cuda"
torch.manual_seed(0)
tokens = torch.randint(3, VOCAB, (14,), dtype=torch.int32).tolist()
tokens[5] = EOS
tokens[6] = EOS # double EOS: second one's look-back stops at the first
tokens[13] = EOS # trailing EOS (chat-template turn boundary)
ctx_len = N - 1
toks_neg = [-t if t == EOS else t for t in tokens]
width = ctx_len + len(tokens)
table = torch.full((1, width), -1, dtype=torch.int32, device=device)
table[0, ctx_len:] = torch.tensor(toks_neg, dtype=torch.int32, device=device)
ne_weights = torch.zeros(N - 1, K, N, dtype=torch.int32)
ne_mods = torch.zeros(N - 1, K, dtype=torch.int32)
for i in range(N - 1):
for j in range(K):
mod = SIZES[i * K + j]
ne_mods[i, j] = mod
for delta in range(N):
ne_weights[i, j, delta] = pow(VOCAB, delta, mod)
ngram = SimpleNamespace(
n=N,
k=K,
num_embedders=NUM_EMB,
ne_weights=ne_weights.to(device),
ne_mods=ne_mods.to(device),
exclusive_sizes=torch.tensor(OFFSETS, dtype=torch.int32, device=device),
)
T = len(tokens)
qsl = torch.tensor([0, T], dtype=torch.int32, device=device)
row_indices = torch.zeros(1, dtype=torch.int64, device=device)
column_starts = torch.full((1,), ctx_len, dtype=torch.int32, device=device)
got = torch.empty(T, NUM_EMB, dtype=torch.int32, device=device)
ops.ngram_compute_n_gram_ids(
N,
K,
ngram.ne_weights,
ngram.ne_mods,
ngram.exclusive_sizes,
qsl,
table,
row_indices,
column_starts,
got,
)
cur = torch.tensor(tokens, dtype=torch.int32, device=device)
tok_req = torch.zeros(T, dtype=torch.int64, device=device)
col = ctx_len + torch.arange(T, device=device)
eos_tok = (cur == EOS).nonzero(as_tuple=True)[0]
got[eos_tok] = compute_eos_position_ngram_ids(
ngram, EOS, table, tok_req, col, eos_tok
)
want = torch.tensor(_ref_ids(tokens), dtype=torch.int32)
torch.testing.assert_close(got.cpu(), want, rtol=0, atol=0)
@@ -0,0 +1,79 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
import json
import os
import tempfile
import pytest
from vllm.model_executor.model_loader.weight_utils import (
filter_duplicate_safetensors_files,
)
def test_filter_duplicate_safetensors_files_missing_weight():
with tempfile.TemporaryDirectory() as tmpdir:
existing_file = os.path.join(tmpdir, "model-00001-of-00002.safetensors")
with open(existing_file, "wb") as f:
f.write(b"")
existing_file2 = os.path.join(tmpdir, "model-00002-of-00002.safetensors")
with open(existing_file2, "wb") as f:
f.write(b"")
index_file = os.path.join(tmpdir, "model.safetensors.index.json")
index_content = {
"weight_map": {
"layer.0.weight": "model-00001-of-00002.safetensors",
"layer.1.weight": "model-00002-of-00002.safetensors",
"layer.2.weight": "model-00003-of-00002.safetensors",
}
}
with open(index_file, "w") as f:
json.dump(index_content, f)
hf_weights_files = [
os.path.join(tmpdir, "model-00001-of-00002.safetensors"),
os.path.join(tmpdir, "model-00002-of-00002.safetensors"),
]
with pytest.raises(FileNotFoundError) as exc_info:
filter_duplicate_safetensors_files(
hf_weights_files=hf_weights_files,
hf_folder=tmpdir,
index_file="model.safetensors.index.json",
)
assert "model-00003-of-00002.safetensors" in str(exc_info.value)
def test_filter_duplicate_safetensors_files_all_exist():
with tempfile.TemporaryDirectory() as tmpdir:
existing_files = []
for i in range(1, 3):
file_path = os.path.join(tmpdir, f"model-0000{i}-of-00002.safetensors")
with open(file_path, "wb") as f:
f.write(b"")
existing_files.append(file_path)
index_file = os.path.join(tmpdir, "model.safetensors.index.json")
index_content = {
"weight_map": {
"layer.0.weight": "model-00001-of-00002.safetensors",
"layer.1.weight": "model-00002-of-00002.safetensors",
}
}
with open(index_file, "w") as f:
json.dump(index_content, f)
filter_duplicate_safetensors_files(
hf_weights_files=existing_files,
hf_folder=tmpdir,
index_file="model.safetensors.index.json",
)
if __name__ == "__main__":
test_filter_duplicate_safetensors_files_missing_weight()
test_filter_duplicate_safetensors_files_all_exist()
@@ -369,6 +369,12 @@ VLM_TEST_SETTINGS = {
vllm_output_post_proc=model_utils.qwen2_vllm_to_hf_output,
patch_hf_runner=model_utils.qwen3_vl_patch_hf_runner,
image_size_factors=[(0.25,), (0.25, 0.25, 0.25), (0.25, 0.2, 0.15)],
marks=[
pytest.mark.skip(
reason="Transformers has no cosmos3_omni mapping, so the HF "
"reference runner cannot load the checkpoint."
)
],
),
"deepseek_vl_v2": VLMTestInfo(
models=["Isotr0py/deepseek-vl2-tiny"], # model repo using dynamic module
@@ -1,438 +0,0 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
"""
Consolidated test for ViT attention backend functionality across multiple models.
This test validates that each multimodal model can successfully generate outputs
using different ViT attention backends. Tests are parametrized by model and backend.
"""
from typing import Any
import pytest
from transformers import AutoProcessor
from vllm import LLM, SamplingParams
from vllm.multimodal.utils import encode_image_url
from vllm.multimodal.video import sample_frames_from_video
from vllm.platforms import current_platform
from vllm.v1.attention.backends.registry import AttentionBackendEnum
from ....utils import create_new_process_for_each_test
from ...utils import dummy_hf_overrides
# Dots.OCR prompt from official repository
# https://github.com/rednote-hilab/dots.ocr/blob/d72d1d8c5bdd0362eb264f714cdbd1e5daa7cdff/dots_ocr/utils/prompts.py#L3
# ruff: noqa: E501
DOTS_OCR_PROMPT = """Please output the layout information from the PDF image, including each layout element's bbox, its category, and the corresponding text content within the bbox.
1. Bbox format: [x1, y1, x2, y2]
2. Layout Categories: The possible categories are ['Caption', 'Footnote', 'Formula', 'List-item', 'Page-footer', 'Page-header', 'Picture', 'Section-header', 'Table', 'Text', 'Title'].
3. Text Extraction & Formatting Rules:
- Picture: For the 'Picture' category, the text field should be omitted.
- Formula: Format its text as LaTeX.
- Table: Format its text as HTML.
- All Others (Text, Title, etc.): Format their text as Markdown.
4. Constraints:
- The output text must be the original text from the image, with no translation.
- All layout elements must be sorted according to human reading order.
5. Final Output: The entire output must be a single JSON object.
"""
VIDEO_PLACEHOLDER = "<|vision_start|><|video_pad|><|vision_end|>"
# Model configurations
MODEL_CONFIGS: dict[str, dict[str, Any]] = {
"dots_ocr": {
"model_name": "rednote-hilab/dots.ocr",
"interface": "llm_chat",
"max_model_len": 32768,
"max_num_seqs": 1,
"limit_mm_per_prompt": {"image": 1},
"sampling_params": {
"temperature": 0.1,
"max_tokens": 16384,
"top_p": 0.9,
"stop_token_ids": None,
},
"use_specific_image": "stop_sign",
"prompt_builder": "build_dots_ocr_prompt",
"output_validator": lambda x: len(x) > 10 and "stop" in x.lower(),
},
"ernie45_vl": {
"model_name": "baidu/ERNIE-4.5-VL-28B-A3B-PT",
"interface": "llm_generate",
"max_model_len": 16384,
"max_num_seqs": 2,
"sampling_params": {
"temperature": 0.0,
"max_tokens": 256,
"stop_token_ids": None,
},
"use_processor": True,
"question": "What is the content of each image?",
},
"glm4_1v": {
"model_name": "zai-org/GLM-4.1V-9B-Thinking",
"interface": "llm_generate",
"max_model_len": 32768,
"max_num_seqs": 2,
"sampling_params": {
"temperature": 0.0,
"max_tokens": 256,
"stop_token_ids": None,
},
"use_processor": True,
"question": "What is the content of each image?",
},
"glm_ocr": {
"model_name": "zai-org/GLM-OCR",
"interface": "llm_generate",
"max_model_len": 131072,
"max_num_seqs": 2,
"sampling_params": {
"temperature": 0.0,
"max_tokens": 256,
"stop_token_ids": None,
},
"use_processor": True,
"question": "Text Recognition:",
},
"keye_vl": {
"model_name": "Kwai-Keye/Keye-VL-8B-Preview",
"interface": "llm_generate",
"max_model_len": 8192,
"max_num_seqs": 5,
"sampling_params": {
"temperature": 0.0,
"max_tokens": 256,
"stop_token_ids": None,
},
"supported_backends": {
AttentionBackendEnum.FLASH_ATTN,
AttentionBackendEnum.ROCM_AITER_FA,
},
"use_processor": True,
"question": "What is the content of each image?",
},
"ovis2_5": {
"model_name": "AIDC-AI/Ovis2.5-2B",
"interface": "llm_generate",
"max_model_len": 8192,
"max_num_seqs": 2,
"sampling_params": {
"temperature": 0.0,
"max_tokens": 256,
"stop_token_ids": None,
},
"prompt_builder": "build_ovis_prompt",
"question": "What is the content of each image?",
},
"qwen2_5_vl": {
"model_name": "Qwen/Qwen2.5-VL-3B-Instruct",
"interface": "vllm_runner",
"media_type": "video",
"max_model_len": 4000,
"max_num_seqs": 1,
"limit_mm_per_prompt": {"video": 1},
"sampling_params": {
"max_tokens": 128,
},
"runner_kwargs": {
"runner": "generate",
"dtype": "bfloat16",
},
"video_params": {
"num_frames": 16,
"pruning_rates": [0.0, 0.75],
},
},
"qwen2_5_omni": {
"model_name": "Qwen/Qwen2.5-Omni-3B",
"interface": "llm_generate",
"max_model_len": 32768,
"max_num_seqs": 2,
"limit_mm_per_prompt": {"image": 3, "video": 3, "audio": 3},
"sampling_params": {
"temperature": 0.6,
"top_p": 0.95,
"top_k": 20,
"max_tokens": 16384,
},
"use_processor": True,
"question": "What is the content of each image?",
},
"qwen3_omni": {
"model_name": "Qwen/Qwen3-Omni-30B-A3B-Instruct",
"interface": "llm_generate",
"max_model_len": 32768,
"max_num_seqs": 2,
"limit_mm_per_prompt": {"image": 3, "video": 3, "audio": 3},
"sampling_params": {
"temperature": 0.6,
"top_p": 0.95,
"top_k": 20,
"max_tokens": 16384,
},
"use_processor": True,
"question": "What is the content of each image?",
},
}
# Prompt builder functions
def build_dots_ocr_prompt(images, config):
"""Build Dots.OCR specific prompt with OCR instructions."""
# Use only stop_sign image for Dots.OCR
image = images[0] # Already filtered to stop_sign
image_url = encode_image_url(image)
placeholders = [{"type": "image_url", "image_url": {"url": image_url}}]
messages = [
{
"role": "user",
"content": [
*placeholders,
{
"type": "text",
"text": f"<|img|><|imgpad|><|endofimg|>{DOTS_OCR_PROMPT}",
},
],
},
]
return messages
def build_processor_prompt(images, config):
"""Build prompt using AutoProcessor.apply_chat_template()."""
processor = AutoProcessor.from_pretrained(
config["model_name"], trust_remote_code=True
)
image_urls = [encode_image_url(img) for img in images]
placeholders = [{"type": "image", "image": url} for url in image_urls]
messages = [
{
"role": "user",
"content": [
*placeholders,
{"type": "text", "text": config["question"]},
],
},
]
return processor.apply_chat_template(
messages, tokenize=False, add_generation_prompt=True
)
def build_ovis_prompt(images, config):
"""Build Ovis2.5 specific prompt with custom format."""
image_urls = [encode_image_url(img) for img in images]
placeholders = "\n".join(
f"Image-{i}: <image>\n" for i, _ in enumerate(image_urls, start=1)
)
return (
f"<|im_start|>user\n\n{placeholders}\n{config['question']}<|im_end|>\n"
"<|im_start|>assistant\n"
)
def build_qwen2_5_video_prompt():
"""Build Qwen2.5-VL video prompt with EVS placeholder."""
return (
f"<|im_start|>system\nYou are a helpful assistant.<|im_end|>\n"
f"<|im_start|>user\n{VIDEO_PLACEHOLDER}"
"Describe this video with a short sentence (no more than 20 words)"
"<|im_end|><|im_start|>assistant\n"
)
# Handler functions
def run_llm_generate_test(config, mm_encoder_attn_backend, image_assets):
"""Standard LLM.generate() interface handler."""
images = [asset.pil_image for asset in image_assets]
# Build prompt
if config.get("use_processor"):
prompt = build_processor_prompt(images, config)
else:
prompt_builder_name = config.get("prompt_builder", "build_ovis_prompt")
prompt_builder = globals()[prompt_builder_name]
prompt = prompt_builder(images, config)
# Determine limit_mm_per_prompt
limit_mm_per_prompt = config.get("limit_mm_per_prompt", {"image": len(images)})
# Create engine
llm = LLM(
model=config["model_name"],
trust_remote_code=True,
max_model_len=config["max_model_len"],
max_num_seqs=config["max_num_seqs"],
limit_mm_per_prompt=limit_mm_per_prompt,
mm_encoder_attn_backend=mm_encoder_attn_backend,
hf_overrides=dummy_hf_overrides,
load_format="dummy",
seed=42,
)
# Generate
sampling_params = SamplingParams(**config["sampling_params"])
outputs = llm.generate(
{
"prompt": prompt,
"multi_modal_data": {"image": images},
},
sampling_params=sampling_params,
)
# Validate
for o in outputs:
generated_text = o.outputs[0].text
validator = config.get("output_validator", lambda x: len(x) > 10)
assert validator(generated_text), (
f"Validation failed for {config['model_name']}: {generated_text}"
)
def run_llm_chat_test(config, mm_encoder_attn_backend, image_assets):
"""LLM.chat() interface handler for Dots.OCR."""
# Filter to stop_sign image only
stop_sign_image = [
asset.pil_image for asset in image_assets if asset.name == "stop_sign"
][0]
# Build messages
messages = build_dots_ocr_prompt([stop_sign_image], config)
# Create engine
llm = LLM(
model=config["model_name"],
trust_remote_code=True,
max_model_len=config["max_model_len"],
max_num_seqs=config["max_num_seqs"],
limit_mm_per_prompt=config["limit_mm_per_prompt"],
mm_encoder_attn_backend=mm_encoder_attn_backend,
hf_overrides=dummy_hf_overrides,
load_format="dummy",
seed=42,
)
# Generate using chat
sampling_params = SamplingParams(**config["sampling_params"])
outputs = llm.chat(messages=messages, sampling_params=sampling_params)
# Validate
for o in outputs:
generated_text = o.outputs[0].text
validator = config.get("output_validator", lambda x: len(x) > 10)
assert validator(generated_text), (
f"Validation failed for {config['model_name']}: {generated_text}"
)
def run_video_test(config, mm_encoder_attn_backend, video_assets, vllm_runner):
"""Video test with EVS (Efficient Video Sampling) handler."""
for pruning_rate in config["video_params"]["pruning_rates"]:
num_frames = config["video_params"]["num_frames"]
# Sample frames from video
sampled_vids = [
sample_frames_from_video(asset.np_ndarrays, num_frames)
for asset in video_assets
]
# Build prompt and prepare video
prompt = build_qwen2_5_video_prompt()
prompts = [prompt]
videos = [sampled_vids[0]]
# Run with vllm_runner context manager
with vllm_runner(
config["model_name"],
max_model_len=config["max_model_len"],
max_num_seqs=config["max_num_seqs"],
limit_mm_per_prompt=config["limit_mm_per_prompt"],
tensor_parallel_size=1,
video_pruning_rate=pruning_rate,
mm_encoder_attn_backend=mm_encoder_attn_backend,
hf_overrides=dummy_hf_overrides,
load_format="dummy",
**config["runner_kwargs"],
) as vllm_model:
outputs = vllm_model.generate_greedy(
prompts,
config["sampling_params"]["max_tokens"],
videos=videos,
)
# Validate output
assert len(outputs) == 1, f"Expected 1 output, got {len(outputs)}"
output_ids, output_text = outputs[0]
assert len(output_ids) > 0, "Generated no output IDs"
assert len(output_text) > 0, "Generated empty text"
assert isinstance(output_text, str), (
f"Output is not string: {type(output_text)}"
)
# Main test function
@pytest.mark.parametrize("model_key", list(MODEL_CONFIGS.keys()))
@pytest.mark.parametrize(
"mm_encoder_attn_backend",
[None] + current_platform.get_supported_vit_attn_backends(),
)
@pytest.mark.skip(reason="Broken test due to memory segmentation fault")
@create_new_process_for_each_test()
def test_vit_backend_functionality(
model_key: str,
mm_encoder_attn_backend: AttentionBackendEnum | None,
image_assets,
video_assets,
vllm_runner,
request,
):
"""Test ViT attention backend functionality for multimodal models.
This test validates that each model can successfully generate outputs
using different ViT attention backends. The test:
1. Filters unsupported backends per model
2. Applies appropriate GPU marks
3. Routes to the correct test handler based on interface
4. Validates output meets minimum requirements
"""
config = MODEL_CONFIGS[model_key]
# Step 1: Backend filtering
if (
"supported_backends" in config
and mm_encoder_attn_backend is not None
and mm_encoder_attn_backend not in config["supported_backends"]
):
pytest.skip(
f"{model_key} does not support {mm_encoder_attn_backend} backend now."
)
# Step 2: Apply GPU marks dynamically
if "gpu_marks" in config:
for mark in config["gpu_marks"]:
request.applymarker(mark)
# Step 3: Route to appropriate handler
if config.get("media_type") == "video":
run_video_test(config, mm_encoder_attn_backend, video_assets, vllm_runner)
elif config["interface"] == "llm_chat":
run_llm_chat_test(config, mm_encoder_attn_backend, image_assets)
elif config["interface"] == "llm_generate":
run_llm_generate_test(config, mm_encoder_attn_backend, image_assets)
else:
raise ValueError(f"Unknown interface: {config['interface']}")
+15 -2
View File
@@ -382,6 +382,19 @@ _TEXT_GENERATION_EXAMPLE_MODELS = {
"LongcatFlashForCausalLM": _HfExamplesInfo(
"meituan-longcat/LongCat-Flash-Chat", trust_remote_code=True
),
"LongcatFlashNgramForCausalLM": _HfExamplesInfo(
"meituan-longcat/LongCat-Flash-Lite",
trust_remote_code=True,
# Shrink the ~62GB n-gram tables (ngram_vocab_size_ratio * vocab_size)
# so the dummy-weight init test fits in CI memory.
hf_overrides={"ngram_vocab_size_ratio": 1},
),
"LongcatCausalLM": _HfExamplesInfo(
"meituan-longcat/LongCat-2.0-FP8",
# Shrink the huge n-gram tables (~264M rows at the checkpoint's
# oe_vocab_size_ratio=100.567) so dummy-weight init fits in CI memory.
hf_overrides={"ngram_vocab_size_ratio": 1},
),
"MambaForCausalLM": _HfExamplesInfo("state-spaces/mamba-130m-hf"),
"Mamba2ForCausalLM": _HfExamplesInfo(
"mistralai/Mamba-Codestral-7B-v0.1",
@@ -799,9 +812,9 @@ _MULTIMODAL_EXAMPLE_MODELS = {
),
"Cosmos3ForConditionalGeneration": _HfExamplesInfo(
"nvidia/Cosmos3-Nano",
extras={"super": "nvidia/Cosmos3-Super"},
max_model_len=4096,
min_transformers_version="4.57",
is_available_online=False,
),
"DeepseekVLV2ForCausalLM": _HfExamplesInfo(
"deepseek-ai/deepseek-vl2-tiny",
@@ -1403,7 +1416,7 @@ _SPECULATIVE_DECODING_EXAMPLE_MODELS = {
# ),
# [DFlash]
"DFlashDraftModel": _HfExamplesInfo(
"Qwen/Qwen3.5-4B",
"Qwen/Qwen3-4B",
speculative_model="z-lab/Qwen3-4B-DFlash-b16",
use_original_num_layers=True, # Need all layers since DFlash has >1 layer,
max_model_len=8192, # Reduce max len to ensure test runs in low-VRAM CI env
+5 -3
View File
@@ -48,9 +48,11 @@ def test_registry_imports(model_arch):
"(see #41376)"
)
# DSpark draft model is NVIDIA-only; class is stubbed to None on ROCm/XPU.
if model_arch == "DSparkDraftModel" and not current_platform.is_cuda():
pytest.skip("DSparkDraftModel is only supported on CUDA")
# DSpark draft model is supported on CUDA and ROCm; stubbed to None on XPU.
if model_arch == "DSparkDraftModel" and not (
current_platform.is_cuda() or current_platform.is_rocm()
):
pytest.skip("DSparkDraftModel is only supported on CUDA and ROCm")
# Ensure all model classes can be imported successfully
model_cls = ModelRegistry._try_load_model_cls(model_arch)
+7 -2
View File
@@ -529,8 +529,13 @@ def dummy_hf_overrides(
}
)
# Update num_hidden_layers for non-Longcat architectures
if model_arch != "LongcatFlashForCausalLM" and model_arch != "LongCatFlashMTPModel":
# Update num_hidden_layers for non-Longcat architectures (Longcat derives it
# from num_layers for its dual-attention layers).
if model_arch not in (
"LongcatFlashForCausalLM",
"LongCatFlashMTPModel",
"LongcatFlashNgramForCausalLM",
):
update_dict["num_hidden_layers"] = num_hidden_layers
text_config.update(update_dict)
+1
View File
@@ -49,4 +49,5 @@ def mock_request():
req = MagicMock(spec=ChatCompletionRequest)
req.tools = []
req.tool_choice = "auto"
req.include_reasoning = True
return req
+1
View File
@@ -43,6 +43,7 @@ def _make_request(**chat_template_kwargs):
request = MagicMock(spec=ChatCompletionRequest)
request.tools = []
request.tool_choice = "auto"
request.include_reasoning = True
request.chat_template_kwargs = chat_template_kwargs or None
return request
@@ -896,6 +896,7 @@ def _make_delegating_request():
req = MagicMock(spec=ChatCompletionRequest)
req.tools = []
req.tool_choice = "auto"
req.include_reasoning = True
return req
+456
View File
@@ -0,0 +1,456 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
"""Tests for include_reasoning suppression in the unified Parser interface.
Covers non-streaming (parser.parse() + build_response_output_items),
streaming (parse_delta), and ParsableContext.append_output() paths.
"""
import json
import os
import pytest
_STRICT_TOOL_CALLING_ENV = "VLLM_ENFORCE_STRICT_TOOL_CALLING"
_STRICT_TOOL_CALLING_ENV_VALUE = os.environ.get(_STRICT_TOOL_CALLING_ENV)
os.environ[_STRICT_TOOL_CALLING_ENV] = "0"
from vllm.entrypoints.openai.chat_completion.protocol import ( # noqa: E402
ChatCompletionRequest,
)
from vllm.entrypoints.openai.engine.protocol import DeltaMessage # noqa: E402
from vllm.entrypoints.openai.responses.protocol import ResponsesRequest # noqa: E402
from vllm.parser.abstract_parser import DelegatingParser # noqa: E402
from vllm.reasoning.basic_parsers import BaseThinkingReasoningParser # noqa: E402
from vllm.tool_parsers.hermes_tool_parser import Hermes2ProToolParser # noqa: E402
@pytest.fixture(scope="module", autouse=True)
def restore_strict_tool_calling_env():
yield
if _STRICT_TOOL_CALLING_ENV_VALUE is None:
os.environ.pop(_STRICT_TOOL_CALLING_ENV, None)
else:
os.environ[_STRICT_TOOL_CALLING_ENV] = _STRICT_TOOL_CALLING_ENV_VALUE
class ThinkReasoningParser(BaseThinkingReasoningParser):
@property
def start_token(self) -> str:
return "<think>"
@property
def end_token(self) -> str:
return "</think>"
MODEL_OUTPUT_REASONING_AND_CONTENT = (
"<think>let me think about this</think>The answer is 42."
)
MODEL_OUTPUT_REASONING_AND_TOOL = (
"<think>I need to call a tool</think>"
'<tool_call>\n{"name": "get_weather", '
'"arguments": {"city": "Dallas"}}\n</tool_call>'
)
MODEL_OUTPUT_CONTENT_ONLY = "The answer is 42."
@pytest.fixture(scope="module")
def tokenizer():
from vllm.tokenizers import get_tokenizer
return get_tokenizer("Qwen/Qwen3-32B")
def make_responses_request(**kwargs) -> ResponsesRequest:
defaults = dict(model="test-model", input="test input")
defaults.update(kwargs)
return ResponsesRequest(**defaults)
def make_chat_request(**kwargs) -> ChatCompletionRequest:
defaults = dict(
model="test-model",
messages=[{"role": "user", "content": "hi"}],
)
defaults.update(kwargs)
return ChatCompletionRequest(**defaults)
def make_parser(tokenizer, reasoning=False, tool=False):
class TestParser(DelegatingParser):
reasoning_parser_cls = ThinkReasoningParser if reasoning else None
tool_parser_cls = Hermes2ProToolParser if tool else None
return TestParser(tokenizer)
# ── Non-streaming: parser.parse() + build_response_output_items ──────
def parse_and_build(parser, request, model_output, enable_auto_tools=False):
"""Mirrors the non-streaming path in _make_response_output_items /
ParsableContext.append_output(): parse suppress reasoning build items.
"""
from vllm.entrypoints.openai.responses.utils import (
build_response_output_items,
)
reasoning, content, tool_calls = parser.parse(
model_output, request, enable_auto_tools=enable_auto_tools
)
if not request.include_reasoning:
reasoning = None
return build_response_output_items(
reasoning=reasoning,
content=content,
tool_calls=tool_calls,
)
class TestNonStreamingIncludeReasoning:
def test_include_reasoning_true_has_reasoning_item(self, tokenizer):
"""Default: reasoning items appear in output."""
parser = make_parser(tokenizer, reasoning=True)
request = make_responses_request(include_reasoning=True)
outputs = parse_and_build(parser, request, MODEL_OUTPUT_REASONING_AND_CONTENT)
types = [o.type for o in outputs]
assert "reasoning" in types
assert "message" in types
def test_include_reasoning_false_no_reasoning_item(self, tokenizer):
"""Reasoning item is suppressed when include_reasoning=False."""
parser = make_parser(tokenizer, reasoning=True)
request = make_responses_request(include_reasoning=False)
outputs = parse_and_build(parser, request, MODEL_OUTPUT_REASONING_AND_CONTENT)
types = [o.type for o in outputs]
assert "reasoning" not in types
assert "message" in types
assert outputs[0].content[0].text == "The answer is 42."
def test_include_reasoning_false_content_preserved(self, tokenizer):
"""Content is extracted correctly even when reasoning is suppressed."""
parser = make_parser(tokenizer, reasoning=True)
request = make_responses_request(include_reasoning=False)
outputs = parse_and_build(parser, request, MODEL_OUTPUT_REASONING_AND_CONTENT)
message = next(o for o in outputs if o.type == "message")
assert message.content[0].text == "The answer is 42."
def test_include_reasoning_false_tool_calls_preserved(self, tokenizer):
"""Tool calls still work when reasoning is suppressed."""
parser = make_parser(tokenizer, reasoning=True, tool=True)
request = make_responses_request(
include_reasoning=False,
tools=[
{
"type": "function",
"name": "get_weather",
"parameters": {
"type": "object",
"properties": {"city": {"type": "string"}},
},
}
],
)
outputs = parse_and_build(
parser,
request,
MODEL_OUTPUT_REASONING_AND_TOOL,
enable_auto_tools=True,
)
types = [o.type for o in outputs]
assert "reasoning" not in types
assert "function_call" in types
fc = next(o for o in outputs if o.type == "function_call")
assert fc.name == "get_weather"
assert json.loads(fc.arguments) == {"city": "Dallas"}
def test_no_reasoning_parser_include_false_is_noop(self, tokenizer):
"""include_reasoning=False is harmless when no reasoning parser."""
parser = make_parser(tokenizer, reasoning=False)
request = make_responses_request(include_reasoning=False)
outputs = parse_and_build(parser, request, MODEL_OUTPUT_CONTENT_ONLY)
assert len(outputs) == 1
assert outputs[0].type == "message"
assert outputs[0].content[0].text == MODEL_OUTPUT_CONTENT_ONLY
def test_default_include_reasoning_is_true(self, tokenizer):
"""ResponsesRequest defaults to include_reasoning=True."""
request = make_responses_request()
assert request.include_reasoning is True
def test_include_reasoning_false_suppresses_all_reasoning(self, tokenizer):
"""Reasoning is suppressed regardless of request type."""
parser = make_parser(tokenizer, reasoning=True)
request = make_responses_request(include_reasoning=False)
outputs = parse_and_build(parser, request, MODEL_OUTPUT_REASONING_AND_CONTENT)
assert all(o.type != "reasoning" for o in outputs)
# ── Streaming: parse_delta ───────────────────────────────────────────
def stream_text(parser, tokenizer, text, request, prompt_token_ids=None):
token_ids = tokenizer.encode(text, add_special_tokens=False)
results: list[DeltaMessage | None] = []
for i, tid in enumerate(token_ids):
delta_text = tokenizer.decode([tid])
is_last = i == len(token_ids) - 1
result = parser.parse_delta(
delta_text,
[tid],
request,
prompt_token_ids=prompt_token_ids,
finished=is_last,
)
prompt_token_ids = None
results.append(result)
return results
def collect_fields(results):
all_reasoning = "".join(r.reasoning for r in results if r and r.reasoning)
all_content = "".join(r.content for r in results if r and r.content)
all_tool_calls = [tc for r in results if r and r.tool_calls for tc in r.tool_calls]
return all_reasoning, all_content, all_tool_calls
class TestParseDeltaIncludeReasoning:
def test_streaming_include_true_emits_reasoning(self, tokenizer):
"""With include_reasoning=True, reasoning deltas are emitted."""
parser = make_parser(tokenizer, reasoning=True)
request = make_responses_request(include_reasoning=True)
results = stream_text(
parser,
tokenizer,
MODEL_OUTPUT_REASONING_AND_CONTENT,
request,
prompt_token_ids=[],
)
reasoning, content, _ = collect_fields(results)
assert "let me think about this" in reasoning
assert "42" in content
def test_streaming_include_false_suppresses_reasoning(self, tokenizer):
"""With include_reasoning=False, no reasoning deltas are emitted."""
parser = make_parser(tokenizer, reasoning=True)
request = make_responses_request(include_reasoning=False)
results = stream_text(
parser,
tokenizer,
MODEL_OUTPUT_REASONING_AND_CONTENT,
request,
prompt_token_ids=[],
)
reasoning, content, _ = collect_fields(results)
assert reasoning == ""
assert "42" in content
def test_streaming_include_false_content_still_works(self, tokenizer):
"""Content is correctly extracted in streaming even with suppression."""
parser = make_parser(tokenizer, reasoning=True)
request = make_responses_request(include_reasoning=False)
results = stream_text(
parser,
tokenizer,
MODEL_OUTPUT_REASONING_AND_CONTENT,
request,
prompt_token_ids=[],
)
_, content, _ = collect_fields(results)
assert "The answer is 42" in content
def test_streaming_include_false_tool_calls_preserved(self, tokenizer):
"""Tool calls stream correctly when reasoning is suppressed."""
parser = make_parser(tokenizer, reasoning=True, tool=True)
request = make_responses_request(
include_reasoning=False,
tools=[
{
"type": "function",
"name": "get_weather",
"parameters": {
"type": "object",
"properties": {"city": {"type": "string"}},
},
}
],
)
results = stream_text(
parser,
tokenizer,
MODEL_OUTPUT_REASONING_AND_TOOL,
request,
prompt_token_ids=[],
)
reasoning, content, tool_calls = collect_fields(results)
assert reasoning == ""
assert len(tool_calls) > 0
assert tool_calls[0].function.name == "get_weather"
tool_args = "".join(
tc.function.arguments for tc in tool_calls if tc.function.arguments
)
assert json.loads(tool_args) == {"city": "Dallas"}
def test_streaming_no_reasoning_parser_include_false(self, tokenizer):
"""No crash when reasoning parser absent and include_reasoning=False."""
parser = make_parser(tokenizer, reasoning=False)
request = make_responses_request(include_reasoning=False)
results = stream_text(
parser,
tokenizer,
MODEL_OUTPUT_CONTENT_ONLY,
request,
prompt_token_ids=[],
)
reasoning, content, _ = collect_fields(results)
assert reasoning == ""
assert "42" in content
def test_streaming_chat_completion_include_false(self, tokenizer):
"""parse_delta also respects ChatCompletionRequest.include_reasoning."""
parser = make_parser(tokenizer, reasoning=True)
request = make_chat_request(include_reasoning=False)
results = stream_text(
parser,
tokenizer,
MODEL_OUTPUT_REASONING_AND_CONTENT,
request,
prompt_token_ids=[],
)
reasoning, content, _ = collect_fields(results)
assert reasoning == ""
assert "42" in content
def test_streaming_reasoning_only_deltas_become_none(self, tokenizer):
"""Deltas that carry only reasoning become None (not empty)."""
parser = make_parser(tokenizer, reasoning=True)
request = make_responses_request(include_reasoning=False)
results = stream_text(
parser,
tokenizer,
MODEL_OUTPUT_REASONING_AND_CONTENT,
request,
prompt_token_ids=[],
)
for r in results:
if r is not None:
assert r.reasoning is None
# ── ParsableContext.append_output() ───────────────────────────────────
class TestParsableContextIncludeReasoning:
def _make_context(self, tokenizer, request):
from vllm.entrypoints.openai.responses.context import ParsableContext
class TestParser(DelegatingParser):
reasoning_parser_cls = ThinkReasoningParser
tool_parser_cls = None
return ParsableContext(
tokenizer=tokenizer,
parser_cls=TestParser,
response_messages=[],
request=request,
available_tools=None,
chat_template=None,
chat_template_content_format="auto",
)
def test_process_include_false_suppresses_reasoning(self, tokenizer):
"""ParsableContext.process() suppresses reasoning items."""
from vllm.outputs import CompletionOutput, RequestOutput
request = make_responses_request(include_reasoning=False)
ctx = self._make_context(tokenizer, request)
output = RequestOutput(
request_id="test",
prompt=None,
prompt_token_ids=[],
prompt_logprobs=None,
outputs=[
CompletionOutput(
index=0,
text=MODEL_OUTPUT_REASONING_AND_CONTENT,
token_ids=tokenizer.encode(
MODEL_OUTPUT_REASONING_AND_CONTENT,
add_special_tokens=False,
),
cumulative_logprob=None,
logprobs=None,
finish_reason="stop",
)
],
finished=True,
)
ctx.append_output(output)
types = [getattr(m, "type", None) for m in ctx.response_messages]
assert "reasoning" not in types
assert "message" in types
def test_process_include_true_has_reasoning(self, tokenizer):
"""ParsableContext.process() includes reasoning by default."""
from vllm.outputs import CompletionOutput, RequestOutput
request = make_responses_request(include_reasoning=True)
ctx = self._make_context(tokenizer, request)
output = RequestOutput(
request_id="test",
prompt=None,
prompt_token_ids=[],
prompt_logprobs=None,
outputs=[
CompletionOutput(
index=0,
text=MODEL_OUTPUT_REASONING_AND_CONTENT,
token_ids=tokenizer.encode(
MODEL_OUTPUT_REASONING_AND_CONTENT,
add_special_tokens=False,
),
cumulative_logprob=None,
logprobs=None,
finish_reason="stop",
)
],
finished=True,
)
ctx.append_output(output)
types = [getattr(m, "type", None) for m in ctx.response_messages]
assert "reasoning" in types
assert "message" in types
+31
View File
@@ -235,6 +235,37 @@ def rms_norm(
torch.ops._C.rms_norm(out, input, weight, epsilon)
# LongCat n-gram embedding index kernel (see csrc/.../ngram_embedding_kernels.cu).
def ngram_compute_n_gram_ids(
ne_n: int,
ne_k: int,
ne_weights: torch.Tensor,
ne_mods: torch.Tensor,
exclusive_ne_embedder_size_sums: torch.Tensor,
exclusive_req_len_sums: torch.Tensor,
ne_token_table: torch.Tensor,
row_indices: torch.Tensor,
column_starts: torch.Tensor,
n_gram_ids: torch.Tensor,
) -> None:
"""Compute concatenated (offset) n-gram ids for a ragged prefill batch.
Writes ``n_gram_ids`` of shape ``[token_num, (ne_n-1)*ne_k]``.
"""
torch.ops._C.ngram_compute_n_gram_ids(
ne_n,
ne_k,
ne_weights,
ne_mods,
exclusive_ne_embedder_size_sums,
exclusive_req_len_sums,
ne_token_table,
row_indices,
column_starts,
n_gram_ids,
)
def fused_add_rms_norm(
input: torch.Tensor,
residual: torch.Tensor,
+33 -2
View File
@@ -518,9 +518,15 @@ class SpeculativeConfig:
"architectures": ["Qwen3_5MoeMTP" if is_moe else "Qwen3_5MTP"],
}
)
if hf_config.model_type == "longcat_flash":
if hf_config.model_type in ("longcat_flash", "longcat_flash_ngram"):
hf_config.model_type = "longcat_flash_mtp"
n_predict = getattr(hf_config, "num_nextn_predict_layers", 1)
# LongCat-2.0 ships one MTP module applied for up to
# mtp_num_layers draft steps (mtp_replicate_modules).
n_predict = (
getattr(hf_config, "num_nextn_predict_layers", None)
or getattr(hf_config, "mtp_num_layers", None)
or 1
)
hf_config.update(
{"n_predict": n_predict, "architectures": ["LongCatFlashMTPModel"]}
)
@@ -942,6 +948,31 @@ class SpeculativeConfig:
"`num_speculative_tokens` was not provided"
)
if self.method == "dspark":
# DSpark is a semi-autoregressive *block* drafter. A
# speculative length smaller than the checkpoint's block
# feeds the block / Markov-head machinery an unsupported
# layout and yields incorrect (garbled) output rather than
# merely lower acceptance. Require num_speculative_tokens to
# be at least the block size (e.g. 5 or 7 for DeepSeek-V4).
dspark_block_size = getattr(
self.draft_model_config.hf_config,
"dspark_block_size",
None,
)
if (
dspark_block_size is not None
and self.num_speculative_tokens < dspark_block_size
):
raise ValueError(
"DSpark requires num_speculative_tokens >= "
f"dspark_block_size ({dspark_block_size}); got "
f"{self.num_speculative_tokens}. Smaller values "
"produce incorrect output. Use "
f"num_speculative_tokens={dspark_block_size} or "
"larger (e.g. 7)."
)
self.draft_tensor_parallel_size = (
SpeculativeConfig._verify_and_get_draft_tp(
self.target_parallel_config,
+2
View File
@@ -70,6 +70,8 @@ DEFAULT_V2_MODEL_RUNNER_ARCHITECTURES = frozenset(
"DeepseekV2ForCausalLM",
"Qwen2MoeForCausalLM",
"GraniteMoeForCausalLM",
"LongcatFlashNgramForCausalLM",
"LongcatCausalLM",
}
)
@@ -600,19 +600,8 @@ class OpenAIServingChat(GenerateBaseServing):
prompt_token_ids=res.prompt_token_ids,
finished=output.finish_reason is not None,
)
if delta_message is not None:
if delta_message.tool_calls:
tools_streamed[i] = True
if (
delta_message.reasoning
and not request.include_reasoning
):
delta_message.reasoning = None
if not (
delta_message.content or delta_message.tool_calls
):
delta_message = None
if delta_message is not None and delta_message.tool_calls:
tools_streamed[i] = True
# handle streaming just a content delta (no parsers)
else:
@@ -627,13 +616,22 @@ class OpenAIServingChat(GenerateBaseServing):
# "control token" for tool calls or the parser otherwise
# wasn't ready to send a token, then
# get the next token without streaming a chunk
# When reasoning is hidden, suppress per-token
# metadata (logprobs, token_ids) on every chunk to
# prevent leaking reasoning tokens through decoded
# token text in logprob entries or raw token IDs.
hide_stream_metadata = (
not request.include_reasoning and parser is not None
)
if hide_stream_metadata:
logprobs = None
if delta_message is None:
# NOTE: If return_token_ids is enabled, we still need to
# send a chunk with token_ids even if delta_message is None
# to ensure all tokens are included in the response
if (
output.finish_reason is None
and not request.return_token_ids
if output.finish_reason is None and (
not request.return_token_ids or hide_stream_metadata
):
continue
delta_message = DeltaMessage()
@@ -666,6 +664,10 @@ class OpenAIServingChat(GenerateBaseServing):
delta=True,
)
include_token_ids = (
request.return_token_ids and not hide_stream_metadata
)
if output.finish_reason is None:
# Send token-by-token response for each request.n
choice_data = ChatCompletionResponseStreamChoice(
@@ -674,9 +676,7 @@ class OpenAIServingChat(GenerateBaseServing):
logprobs=logprobs,
finish_reason=None,
token_ids=(
as_list(output.token_ids)
if request.return_token_ids
else None
as_list(output.token_ids) if include_token_ids else None
),
)
@@ -704,9 +704,7 @@ class OpenAIServingChat(GenerateBaseServing):
finish_reason=finish_reason_,
stop_reason=output.stop_reason,
token_ids=(
as_list(output.token_ids)
if request.return_token_ids
else None
as_list(output.token_ids) if include_token_ids else None
),
)
@@ -883,12 +881,16 @@ class OpenAIServingChat(GenerateBaseServing):
enable_auto_tools=self.enable_auto_tools,
model_output_token_ids=token_ids,
)
suppress_metadata = not request.include_reasoning and parser is not None
if not request.include_reasoning:
reasoning = None
if suppress_metadata:
logprobs = None
else:
reasoning = None
content = output.text
tool_calls = []
suppress_metadata = False
auto_tools_called = False
is_named_tool_choice = (
@@ -982,7 +984,9 @@ class OpenAIServingChat(GenerateBaseServing):
else "stop",
stop_reason=output.stop_reason,
token_ids=(
as_list(output.token_ids) if request.return_token_ids else None
as_list(output.token_ids)
if request.return_token_ids and not suppress_metadata
else None
),
routed_experts=routed_experts_b64,
)
+5 -1
View File
@@ -310,7 +310,9 @@ class ParsableContext(ConversationContext):
self.finish_reason: str | None = None
self.enable_auto_tools = enable_auto_tools
self.response_parser = response_parser
self.response_parser = response_parser or (
parser_cls(tokenizer, request.tools) if parser_cls is not None else None
)
self.parser_cls = parser_cls
self.request = request
@@ -344,6 +346,8 @@ class ParsableContext(ConversationContext):
enable_auto_tools=self.enable_auto_tools,
model_output_token_ids=completion.token_ids,
)
if not self.request.include_reasoning:
reasoning = None
self.response_messages.extend(
build_response_output_items(
reasoning=reasoning,
@@ -161,6 +161,15 @@ class ResponsesRequest(OpenAIBaseModel):
previous_response_id: str | None = None
prompt: ResponsePrompt | None = None
reasoning: Reasoning | None = None
include_reasoning: bool = Field(
default=True,
description=(
"Whether to include reasoning content in the response. "
"When false, reasoning tokens are still generated but "
"excluded from the output. This reduces network traffic "
"without affecting model inference."
),
)
service_tier: Literal["auto", "default", "flex", "scale", "priority"] = "auto"
store: bool | None = True
stream: bool | None = False
@@ -1068,6 +1068,9 @@ class OpenAIServingResponses(GenerateBaseServing):
enable_auto_tools=self.enable_auto_tools,
model_output_token_ids=final_output.token_ids,
)
if not request.include_reasoning:
reasoning = None
logprobs = None
return build_response_output_items(
reasoning=reasoning,
content=content,
@@ -1342,11 +1345,15 @@ class OpenAIServingResponses(GenerateBaseServing):
) -> AsyncGenerator[StreamingResponsesResponse, None]:
processor = SimpleStreamingEventProcessor(tools=request.tools)
hide_stream_metadata = not request.include_reasoning and self.parser is not None
def _get_logprobs(
output: CompletionOutput,
) -> list[response_text_delta_event.Logprob]:
if not request.is_include_output_logprobs():
return []
if hide_stream_metadata:
return []
return self._create_stream_response_logprobs(
token_ids=output.token_ids,
logprobs=output.logprobs,
@@ -844,17 +844,14 @@ class QwenGatedDeltaNetAttention(GatedDeltaNetAttention):
def forward(
self,
hidden_states: torch.Tensor,
output: torch.Tensor,
):
self._forward_method(hidden_states, output)
) -> torch.Tensor:
return self._forward_method(hidden_states)
def _output_projection(
self,
core_attn_out: torch.Tensor,
z: torch.Tensor,
output: torch.Tensor,
num_tokens: int,
):
) -> torch.Tensor:
"""Part 3: RMSNormGated + output linear projection.
The RMSNormGated + quant sequence is eligible for fusion
@@ -866,13 +863,13 @@ class QwenGatedDeltaNetAttention(GatedDeltaNetAttention):
core_attn_out = self.norm(core_attn_out, z)
core_attn_out = core_attn_out.reshape(z_shape_og)
core_attn_out = core_attn_out.flatten(-2) # ... h d -> ... (h d)
output[:num_tokens], _ = self.out_proj(core_attn_out)
output, _ = self.out_proj(core_attn_out)
return output
def forward_hip(
self,
hidden_states: torch.Tensor,
output: torch.Tensor,
):
) -> torch.Tensor:
"""ROCm forward using AITER Triton fused projection+attention when
available, otherwise falling back to the generic CUDA path."""
if GDN_AITER_TRITON_AVAILABLE:
@@ -901,15 +898,14 @@ class QwenGatedDeltaNetAttention(GatedDeltaNetAttention):
use_aiter=True,
)
self._output_projection(core_attn_out, z, output, num_tokens)
return self._output_projection(core_attn_out, z)
else:
self.forward_cuda(hidden_states, output)
return self.forward_cuda(hidden_states)
def forward_cuda(
self,
hidden_states: torch.Tensor,
output: torch.Tensor,
):
) -> torch.Tensor:
"""
Forward pass with three parts:
1. Input projection
@@ -964,13 +960,12 @@ class QwenGatedDeltaNetAttention(GatedDeltaNetAttention):
# ============================================================
# Part 3: Output Projection
# ============================================================
self._output_projection(core_attn_out, z, output, num_tokens)
return self._output_projection(core_attn_out, z)
def forward_xpu(
self,
hidden_states: torch.Tensor,
output: torch.Tensor,
):
) -> torch.Tensor:
"""
Forward pass with three parts:
1. Input projection
@@ -1013,13 +1008,13 @@ class QwenGatedDeltaNetAttention(GatedDeltaNetAttention):
core_attn_out = self.norm(core_attn_out, z)
core_attn_out = core_attn_out.reshape(z_shape_og)
core_attn_out = core_attn_out.flatten(-2) # ... h d -> ... (h d)
output[:num_tokens], _ = self.out_proj(core_attn_out)
out, _ = self.out_proj(core_attn_out)
return out
def forward_cpu(
self,
hidden_states: torch.Tensor,
output: torch.Tensor,
):
) -> torch.Tensor:
assert not hasattr(self, "in_proj_qkv"), "lora isn't supported on CPU."
mixed_qkvz, _ = self.in_proj_qkvz(hidden_states)
@@ -1063,7 +1058,8 @@ class QwenGatedDeltaNetAttention(GatedDeltaNetAttention):
core_attn_out = self.norm(core_attn_out, z)
core_attn_out = core_attn_out.reshape(z_shape_og)
core_attn_out = core_attn_out.flatten(-2) # ... h d -> ... (h d)
output[:num_tokens], _ = self.out_proj(core_attn_out)
out, _ = self.out_proj(core_attn_out)
return out
def _warmup_prefill_kernels(self, qkv_or_qkvz: torch.Tensor, v_dim: int) -> None:
"""Warm up GDN prefill kernels during V1 profiling.
@@ -250,6 +250,40 @@ def fused_indexer_q_rope_quant(
return q_fp8, weights_out
def _mask_init_and_local_tokens(
logits: torch.Tensor,
row_starts: torch.Tensor | None,
row_ends: torch.Tensor,
num_init_tokens: int,
num_local_tokens: int,
) -> None:
"""Force streaming tokens into the top-k set (streaming-aware indexing):
scatter +inf into the first ``num_init_tokens`` and last
``num_local_tokens`` columns of each row's valid range
``[row_starts, row_ends)``. Out-of-range writes are clamped into
``[row_starts, row_ends)`` so they never leak into other rows' ranges.
"""
device = logits.device
ends = row_ends.to(device=device, dtype=torch.int64).reshape(-1)
if row_starts is None:
starts = torch.zeros_like(ends)
else:
starts = row_starts.to(device=device, dtype=torch.int64).reshape(-1)
last = (ends - 1).clamp_max_(logits.shape[1] - 1).clamp_min_(starts)
if num_init_tokens > 0:
init_idx = starts[:, None] + torch.arange(
num_init_tokens, dtype=torch.int64, device=device
)
init_idx = torch.minimum(init_idx, last[:, None])
logits.scatter_(1, init_idx, float("inf"))
if num_local_tokens > 0:
local_idx = last[:, None] - torch.arange(
num_local_tokens, dtype=torch.int64, device=device
)
local_idx = torch.maximum(local_idx, starts[:, None])
logits.scatter_(1, local_idx, float("inf"))
def _gather_workspace_shapes(
total_seq_lens: int,
head_dim: int,
@@ -313,6 +347,8 @@ def sparse_attn_indexer(
dcp_world_size: int = 1,
cp_kv_cache_interleave_size: int = 1,
skip_topk_buffer_clear: bool = False,
num_init_tokens: int = 0,
num_local_tokens: int = 0,
) -> torch.Tensor:
# careful! this will be None in dummy run
attn_metadata = get_forward_context().attn_metadata
@@ -471,6 +507,14 @@ def sparse_attn_indexer(
cu_seqlen_ke,
clean_logits=False,
)
if num_init_tokens > 0 or num_local_tokens > 0:
_mask_init_and_local_tokens(
logits,
cu_seqlen_ks,
cu_seqlen_ke,
num_init_tokens,
num_local_tokens,
)
num_rows = logits.shape[0]
ops.top_k_per_row_prefill(
logits,
@@ -572,6 +616,17 @@ def sparse_attn_indexer(
num_rows = logits.shape[0]
topk_indices = topk_indices_buffer[:num_padded_tokens, :topk_tokens]
if num_init_tokens > 0 or num_local_tokens > 0:
# seq_lens is (B, next_n) whenever next_n > 1, so flattening
# yields the per-row context length.
_mask_init_and_local_tokens(
logits,
None,
seq_lens.reshape(-1)[:num_rows],
num_init_tokens,
num_local_tokens,
)
use_cooperative_topk = (
current_platform.is_cuda()
and topk_tokens in (512, 1024, 2048)
@@ -668,6 +723,8 @@ def sparse_attn_indexer_fake(
dcp_world_size: int = 1,
cp_kv_cache_interleave_size: int = 1,
skip_topk_buffer_clear: bool = False,
num_init_tokens: int = 0,
num_local_tokens: int = 0,
) -> torch.Tensor:
return topk_indices_buffer
@@ -706,6 +763,8 @@ class SparseAttnIndexer(CustomOp):
topk_indices_buffer: torch.Tensor,
skip_k_cache_insert: bool = False,
use_fp4_cache: bool = False,
num_init_tokens: int = 0,
num_local_tokens: int = 0,
):
super().__init__()
self.k_cache = k_cache
@@ -718,6 +777,8 @@ class SparseAttnIndexer(CustomOp):
self.topk_indices_buffer = topk_indices_buffer
self.skip_k_cache_insert = skip_k_cache_insert
self.use_fp4_cache = use_fp4_cache
self.num_init_tokens = num_init_tokens
self.num_local_tokens = num_local_tokens
# DCP scalars are constant for the run; resolve them here (config is set
# during model construction) and pass them into the custom op, rather
# than threading them through per-step metadata.
@@ -781,6 +842,8 @@ class SparseAttnIndexer(CustomOp):
self.dcp_rank,
self.dcp_world_size,
self.cp_kv_cache_interleave_size,
num_init_tokens=self.num_init_tokens,
num_local_tokens=self.num_local_tokens,
)
def forward_xpu(
@@ -803,6 +866,10 @@ class SparseAttnIndexer(CustomOp):
assert isinstance(q_quant, torch.Tensor), (
"AMD sparse_attn_indexer expects a single FP8 q_quant tensor"
)
assert self.num_init_tokens == 0 and self.num_local_tokens == 0, (
"Streaming-aware indexing (index_init_tokens/index_local_tokens) is "
"not supported on the ROCm sparse_attn_indexer path yet"
)
if rocm_aiter_ops.is_enabled():
return torch.ops.vllm.rocm_aiter_sparse_attn_indexer(
hidden_states,
@@ -595,6 +595,14 @@ def filter_duplicate_safetensors_files(
weight_files_in_index = set()
for weight_name in weight_map:
weight_files_in_index.add(os.path.join(hf_folder, weight_map[weight_name]))
# Check if files referenced in model.safetensors.index.json actually exist.
# Raise error if any file is missing.
hf_weights_files_set = set(hf_weights_files)
missing_files = weight_files_in_index - hf_weights_files_set
if missing_files:
raise FileNotFoundError(
f"Weight files referenced in index but missing: {missing_files}"
)
# Filter out any fields that are not found in the index file.
hf_weights_files = [f for f in hf_weights_files if f in weight_files_in_index]
return hf_weights_files
+29 -1
View File
@@ -499,7 +499,7 @@ class LlamaBidirectionalConfig(VerifyAndUpdateConfig):
"last": "LAST",
}
pooling_type = pooling_type_map.get(hf_config.pooling, None)
pooling_type = pooling_type_map.get(hf_config.pooling)
if pooling_type is None:
raise ValueError(f"pool_type {hf_config.pooling!r} not supported")
@@ -809,6 +809,32 @@ class VoyageQwen3BidirectionalEmbedModelConfig(VerifyAndUpdateConfig):
model_config.hf_config.embedding_size = model_config.hf_config.num_labels
class LongcatFlashNgramForCausalLMConfig(VerifyAndUpdateConfig):
@staticmethod
def verify_and_update_config(vllm_config: "VllmConfig") -> None:
# LongCat-Flash-Lite's zero-expert MoE trips a data-dependent assert
# under torch.compile, and its n-gram inputs_embeds are only wired for
# FULL cudagraph capture (PIECEWISE prefill drops them). Default to
# no-compile + FULL cudagraph (prefill runs eager) unless the user
# configured compilation explicitly.
from vllm.config.compilation import CompilationMode, CUDAGraphMode
compilation_config = vllm_config.compilation_config
if compilation_config.mode is None:
compilation_config.mode = CompilationMode.NONE
if compilation_config.cudagraph_mode is None:
compilation_config.cudagraph_mode = CUDAGraphMode.FULL
# LongCat-2.0 sparse attention (DSA indexer) requires the same
# kv-cache dtype normalization as DeepSeek-V3.2.
hf_config = vllm_config.model_config.hf_config
if hasattr(hf_config, "index_topk"):
cache_config = vllm_config.cache_config
if cache_config.cache_dtype == "bfloat16":
cache_config.cache_dtype = "auto"
logger.info("Using bfloat16 kv-cache for LongCat sparse attention")
MODELS_CONFIG_MAP: dict[str, type[VerifyAndUpdateConfig]] = {
"ColBERTJinaRobertaModel": JinaRobertaModelConfig,
"ColQwen3_5": ColQwen3_5Config,
@@ -822,6 +848,8 @@ MODELS_CONFIG_MAP: dict[str, type[VerifyAndUpdateConfig]] = {
"Gemma4ForConditionalGeneration": Gemma4Config,
"Gemma4UnifiedForConditionalGeneration": Gemma4Config,
"GptOssForCausalLM": GptOssForCausalLMConfig,
"LongcatFlashNgramForCausalLM": LongcatFlashNgramForCausalLMConfig,
"LongcatCausalLM": LongcatFlashNgramForCausalLMConfig,
"GteModel": SnowflakeGteNewModelConfig,
"GteNewForSequenceClassification": GteNewModelConfig,
"GteNewModel": GteNewModelConfig,
+34 -17
View File
@@ -710,6 +710,8 @@ class Indexer(nn.Module):
self.max_model_len,
self.max_total_seq_len,
self.topk_indices_buffer,
num_init_tokens=getattr(config, "index_init_tokens", 0),
num_local_tokens=getattr(config, "index_local_tokens", 0),
)
self.is_inplace_rope = is_inplace_rope
@@ -820,45 +822,56 @@ def _try_load_fp8_indexer_wk(
name, tensor, buf, params_dict, loaded_params, pp_missing_layer_names
):
"""
We fuse the WK and weights_proj projections, but in some checkpoints WK is stored
in FP8 with a separate weight_scale_inv, while weights_proj is stored in BF16.
Upcasting to BF16 during loading enables the fusion. This function loads the FP8 WK
weights and scale, and when both are available, dequantizes to BF16 and stores into
the fused wk_weights_proj.weight parameter.
We fuse the WK and weights_proj projections, but in some checkpoints one
or both are stored in FP8 with a separate weight_scale_inv while the fused
parameter is BF16. Upcasting to BF16 during loading enables the fusion.
This function buffers the FP8 weight and scale, and when both are
available, dequantizes to BF16 and stores into the corresponding shard of
the fused wk_weights_proj.weight.
"""
if "indexer.wk." not in name or "wk_weights" in name:
return False # Weight is not an isolated WK weight for the indexer, ignore.
if "wk_weights" in name:
return False # Already-fused parameter name, ignore.
if "indexer.wk." in name:
sub_name, shard_id = ".wk.", 0
elif "indexer.weights_proj." in name:
sub_name, shard_id = ".weights_proj.", 1
else:
return False # Not an isolated indexer projection weight, ignore.
is_weight = name.endswith(".weight") and tensor.dtype == torch.float8_e4m3fn
is_scale = "weight_scale" in name
if not is_weight and not is_scale:
return False # WK is not in FP8 format, ignore.
return False # Projection is not in FP8 format, ignore.
# Buffer this tensor (weight or scale) until both have arrived.
layer_prefix = name.rsplit(".wk.", 1)[0] # e.g. "model.layers.0.self_attn.indexer"
# layer_prefix is e.g. "model.layers.0.self_attn.indexer"
layer_prefix = name.rsplit(sub_name, 1)[0]
fused_name = f"{layer_prefix}.wk_weights_proj.weight"
if any(
name.startswith(missing_layer_name)
for missing_layer_name in pp_missing_layer_names
):
return True
entry = buf.setdefault(layer_prefix, {})
entry = buf.setdefault((layer_prefix, shard_id), {})
entry["weight" if is_weight else "scale"] = tensor
if "weight" not in entry or "scale" not in entry:
return True # still waiting for the other param
# We have both weight and scale: dequantize FP8 to BF16.
# We have both weight and scale: dequantize FP8 to BF16. Derive the block
# shape per axis: narrow projections (e.g. a 32-row weights_proj under
# 128x128 quantization) span a partial row block.
weight_fp8, scale_inv = entry["weight"], entry["scale"]
del buf[layer_prefix]
block_size = weight_fp8.shape[1] // scale_inv.shape[1]
del buf[(layer_prefix, shard_id)]
row_block = weight_fp8.shape[0] // scale_inv.shape[0]
col_block = weight_fp8.shape[1] // scale_inv.shape[1]
weight_bf16 = scaled_dequantize(
weight_fp8,
scale_inv,
group_shape=GroupShape(block_size, block_size),
group_shape=GroupShape(row_block, col_block),
out_dtype=torch.bfloat16,
)
# Load the dequantized weight into shard 0 of the fused buffer.
# Load the dequantized weight into its shard of the fused buffer.
param = params_dict[fused_name]
param.weight_loader(param, weight_bf16, 0)
param.weight_loader(param, weight_bf16, shard_id)
loaded_params.add(fused_name)
return True
@@ -975,6 +988,7 @@ class DeepseekV2MLAAttention(nn.Module):
topk_indices_buffer: torch.Tensor | None = None,
input_size: int | None = None,
reduce_results: bool = True,
skip_topk: bool | None = None,
) -> None:
super().__init__()
self.hidden_size = hidden_size
@@ -1077,7 +1091,10 @@ class DeepseekV2MLAAttention(nn.Module):
# Refer: https://arxiv.org/abs/2603.12201 for more details.
_skip_topk = False
is_mtp_layer = False
if self.is_v32:
if self.is_v32 and skip_topk is not None:
# The caller decides the skip schedule directly.
_skip_topk = skip_topk
elif self.is_v32:
_index_topk_freq = getattr(config, "index_topk_freq", 1)
_index_topk_pattern = getattr(config, "index_topk_pattern", None)
_index_skip_topk_offset = getattr(config, "index_skip_topk_offset", 2)
+78 -10
View File
@@ -64,13 +64,18 @@ from vllm.model_executor.layers.vocab_parallel_embedding import (
VocabParallelEmbedding,
)
from vllm.model_executor.model_loader.weight_utils import default_weight_loader
from vllm.model_executor.models.deepseek_v2 import DeepseekV2MLAAttention
from vllm.model_executor.models.deepseek_v2 import (
DeepseekV2MLAAttention,
_try_load_fp8_indexer_wk,
)
from vllm.platforms import current_platform
from vllm.sequence import IntermediateTensors
from .interfaces import SupportsLoRA, SupportsPP
from .utils import (
AutoWeightsLoader,
PPMissingLayer,
get_pp_missing_layer_names,
is_pp_missing_parameter,
make_empty_intermediate_tensors_factory,
make_layers,
@@ -318,7 +323,7 @@ class LongcatMoe(nn.Module):
hidden_states = hidden_states.view(-1, hidden_dim)
# Align to FusedMoE padded hidden size to avoid dim mismatch
padded_hidden = self.experts.hidden_size
padded_hidden = self.experts.moe_config.hidden_dim
if hidden_dim < padded_hidden:
hidden_states_padded = torch.nn.functional.pad(
hidden_states,
@@ -348,6 +353,18 @@ class LongcatMoe(nn.Module):
return final_hidden_states.view(num_tokens, hidden_dim)
def maybe_replace_indexer_k_norm(
attn: DeepseekV2MLAAttention, config: PretrainedConfig
) -> None:
"""LongCat's DSA indexer normalizes K with RMSNorm (index_k_norm_type),
where the shared Indexer defaults to DeepSeek-V3.2's LayerNorm."""
if (
getattr(attn, "indexer", None) is not None
and getattr(config, "index_k_norm_type", None) == "rms"
):
attn.indexer.k_norm = RMSNorm(config.index_head_dim, eps=1e-6)
class FlashDecoderLayer(nn.Module):
"""Flash decoder layer with dual attention and MLP structure."""
@@ -359,12 +376,18 @@ class FlashDecoderLayer(nn.Module):
quant_config: QuantizationConfig | None = None,
prefix: str = "",
enable_eplb: bool = False,
topk_indices_buffer: torch.Tensor | None = None,
) -> None:
super().__init__()
self.layer_idx = int(prefix.split(sep=".")[-1])
self.hidden_size = config.hidden_size
max_position_embeddings = getattr(config, "max_position_embeddings", 8192)
# Cross-Layer Indexing: with cli_factor=2 the first attention of each
# layer computes DSA top-k indices and the second reuses them via the
# shared buffer (attention sublayer id = 2 * layer_idx + i).
cli_factor = getattr(config, "cli_factor", 1) or 1
# Dual attention structure
self.self_attn = nn.ModuleList(
[
@@ -386,10 +409,14 @@ class FlashDecoderLayer(nn.Module):
if "self_attn" in getattr(config, "disable_quant_module", [])
else quant_config,
prefix=f"{prefix}.self_attn.{i}",
topk_indices_buffer=topk_indices_buffer,
skip_topk=(2 * self.layer_idx + i) % cli_factor != 0,
)
for i in range(2)
]
)
for attn in self.self_attn:
maybe_replace_indexer_k_norm(attn, config)
self.input_layernorm = nn.ModuleList(
[RMSNorm(config.hidden_size, eps=config.rms_norm_eps) for i in range(2)]
)
@@ -490,6 +517,17 @@ class FlashModel(nn.Module):
self.vocab_size = config.vocab_size
self.is_v32 = hasattr(config, "index_topk")
if self.is_v32:
topk_indices_buffer = torch.empty(
vllm_config.scheduler_config.max_num_batched_tokens,
config.index_topk,
dtype=torch.int32,
device=current_platform.device_type,
)
else:
topk_indices_buffer = None
if get_pp_group().is_first_rank:
self.embed_tokens = VocabParallelEmbedding(
config.vocab_size,
@@ -506,6 +544,7 @@ class FlashModel(nn.Module):
cache_config=cache_config,
quant_config=quant_config,
prefix=prefix,
topk_indices_buffer=topk_indices_buffer,
),
prefix=f"{prefix}.layers",
)
@@ -572,15 +611,37 @@ class FlashModel(nn.Module):
("fused_qkv_a_proj", "kv_a_proj_with_mqa", 1),
(".gate_up_proj", ".gate_proj", 0),
(".gate_up_proj", ".up_proj", 1),
# Fused indexer wk + weights_proj (shard 0 = wk, 1 = weights_proj)
("wk_weights_proj", "wk", 0),
("wk_weights_proj", "weights_proj", 1),
]
expert_params_mapping = self.get_expert_mapping()
loaded_params: set[str] = set()
pp_missing_layer_names = get_pp_missing_layer_names(self)
params_dict = dict(self.named_parameters())
_pending_wk_fp8: dict = {}
# Drop checkpoint indexer weights for sublayers without an indexer.
indexer_present_prefixes = {
n.rsplit(".indexer.", 1)[0] for n in params_dict if ".indexer." in n
}
for name, loaded_weight in weights:
if "rotary_emb.inv_freq" in name:
continue
if ".indexer." in name and (
name.rsplit(".indexer.", 1)[0] not in indexer_present_prefixes
):
continue
if _try_load_fp8_indexer_wk(
name,
loaded_weight,
_pending_wk_fp8,
params_dict,
loaded_params,
pp_missing_layer_names,
):
continue
for param_name, weight_name, shard_id in stacked_params_mapping:
if weight_name not in name:
continue
@@ -653,6 +714,21 @@ class FlashModel(nn.Module):
continue
if is_pp_missing_parameter(name, self):
continue
# Fold the MLA LoRA scaling at load time: load_weights may
# run incrementally, so a post-load fold can miss weights
# arriving in a later call.
if name.endswith(".q_a_layernorm.weight") and getattr(
self.config, "mla_scale_q_lora", False
):
loaded_weight = loaded_weight.float() * (
(self.config.hidden_size / self.config.q_lora_rank) ** 0.5
)
elif name.endswith(".kv_a_layernorm.weight") and getattr(
self.config, "mla_scale_kv_lora", False
):
loaded_weight = loaded_weight.float() * (
(self.config.hidden_size / self.config.kv_lora_rank) ** 0.5
)
param = params_dict[name]
weight_loader = getattr(
param, "weight_loader", default_weight_loader
@@ -687,14 +763,6 @@ class FlashModel(nn.Module):
).split([self_attn.qk_nope_head_dim, self_attn.v_head_dim], dim=1)
self_attn.w_kc = w_kc.transpose(1, 2).contiguous().transpose(1, 2)
self_attn.w_vc = w_vc.contiguous().transpose(1, 2)
if self.config.mla_scale_q_lora:
self_attn.q_a_layernorm.weight.data *= (
self.config.hidden_size / self.config.q_lora_rank
) ** 0.5
if self.config.mla_scale_kv_lora:
self_attn.kv_a_layernorm.weight.data *= (
self.config.hidden_size / self.config.kv_lora_rank
) ** 0.5
return loaded_params
+64 -13
View File
@@ -20,10 +20,14 @@ from vllm.model_executor.layers.vocab_parallel_embedding import (
VocabParallelEmbedding,
)
from vllm.model_executor.model_loader.weight_utils import default_weight_loader
from vllm.model_executor.models.longcat_flash import FlashConfig
from vllm.model_executor.models.longcat_flash import (
FlashConfig,
maybe_replace_indexer_k_norm,
)
from vllm.platforms import current_platform
from vllm.sequence import IntermediateTensors
from .deepseek_v2 import DeepseekV2DecoderLayer
from .deepseek_v2 import DeepseekV2DecoderLayer, _try_load_fp8_indexer_wk
from .utils import maybe_prefix
@@ -34,6 +38,7 @@ class LongCatMultiTokenPredictorLayer(nn.Module):
prefix: str,
vllm_config: VllmConfig,
quant_config: QuantizationConfig | None = None,
topk_indices_buffer: torch.Tensor | None = None,
) -> None:
super().__init__()
self.enorm = RMSNorm(config.hidden_size, eps=config.rms_norm_eps)
@@ -45,7 +50,10 @@ class LongCatMultiTokenPredictorLayer(nn.Module):
quant_config=quant_config,
prefix="eh_proj",
)
self.mtp_block = DeepseekV2DecoderLayer(vllm_config, prefix)
self.mtp_block = DeepseekV2DecoderLayer(
vllm_config, prefix, topk_indices_buffer=topk_indices_buffer
)
maybe_replace_indexer_k_norm(self.mtp_block.self_attn, config)
self.final_layernorm = RMSNorm(config.hidden_size, eps=config.rms_norm_eps)
def forward(
@@ -84,6 +92,15 @@ class LongCatMultiTokenPredictor(nn.Module):
vllm_config.model_config.hf_config.intermediate_size = config.intermediate_size
self.mtp_start_layer_idx = config.num_hidden_layers * 2
self.num_mtp_layers = 1
if hasattr(config, "index_topk"):
topk_indices_buffer = torch.empty(
vllm_config.scheduler_config.max_num_batched_tokens,
config.index_topk,
dtype=torch.int32,
device=current_platform.device_type,
)
else:
topk_indices_buffer = None
self.layers = torch.nn.ModuleDict(
{
str(idx): LongCatMultiTokenPredictorLayer(
@@ -91,6 +108,7 @@ class LongCatMultiTokenPredictor(nn.Module):
prefix=f"{prefix}.layers.{idx}",
vllm_config=vllm_config,
quant_config=quant_config,
topk_indices_buffer=topk_indices_buffer,
)
for idx in range(
self.mtp_start_layer_idx,
@@ -126,8 +144,10 @@ class LongCatMultiTokenPredictor(nn.Module):
class LongCatFlashMTP(nn.Module):
def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""):
super().__init__()
# LongCat MTP without MoE layers
vllm_config.model_config.hf_config.n_routed_experts = None
# LongCat MTP has no MoE layers: clear n_routed_experts so the predictor
# builds a dense MLP. object.__setattr__ bypasses the ngram remote
# config's strict validation (it rejects setting the int field to None).
object.__setattr__(vllm_config.model_config.hf_config, "n_routed_experts", None)
self.config = FlashConfig(**vllm_config.model_config.hf_config.__dict__)
self.quant_config = (
None
@@ -176,6 +196,9 @@ class LongCatFlashMTP(nn.Module):
("gate_up_proj", "up_proj", 1),
("fused_qkv_a_proj", "q_a_proj", 0),
("fused_qkv_a_proj", "kv_a_proj_with_mqa", 1),
# Fused indexer wk + weights_proj (shard 0 = wk, 1 = weights_proj)
("wk_weights_proj", "wk", 0),
("wk_weights_proj", "weights_proj", 1),
]
new_to_old_names_mapping = {
@@ -186,6 +209,13 @@ class LongCatFlashMTP(nn.Module):
"model.mtp.layers.0.hnorm.m.weight": "hnorm.weight",
"model.mtp.layers.0.input_layernorm.weight": "model.layers.0.input_layernorm.weight", # noqa: E501
"model.mtp.layers.0.post_attention_layernorm.weight": "model.layers.0.post_attention_layernorm.weight", # noqa: E501
"model.mtp.layers.0.self_attn.indexer.k_norm.weight": "model.layers.0.self_attn.indexer.k_norm.weight", # noqa: E501
"model.mtp.layers.0.self_attn.indexer.wq_b.weight": "model.layers.0.self_attn.indexer.wq_b.weight", # noqa: E501
"model.mtp.layers.0.self_attn.indexer.wq_b.weight_scale_inv": "model.layers.0.self_attn.indexer.wq_b.weight_scale_inv", # noqa: E501
"model.mtp.layers.0.self_attn.indexer.wk.weight": "model.layers.0.self_attn.indexer.wk.weight", # noqa: E501
"model.mtp.layers.0.self_attn.indexer.wk.weight_scale_inv": "model.layers.0.self_attn.indexer.wk.weight_scale_inv", # noqa: E501
"model.mtp.layers.0.self_attn.indexer.weights_proj.weight": "model.layers.0.self_attn.indexer.weights_proj.weight", # noqa: E501
"model.mtp.layers.0.self_attn.indexer.weights_proj.weight_scale_inv": "model.layers.0.self_attn.indexer.weights_proj.weight_scale_inv", # noqa: E501
"model.mtp.layers.0.self_attn.kv_a_layernorm.weight": "model.layers.0.self_attn.kv_a_layernorm.weight", # noqa: E501
"model.mtp.layers.0.self_attn.kv_a_proj_with_mqa.weight": "model.layers.0.self_attn.kv_a_proj_with_mqa.weight", # noqa: E501
"model.mtp.layers.0.self_attn.kv_a_proj_with_mqa.weight_scale_inv": "model.layers.0.self_attn.kv_a_proj_with_mqa.weight_scale_inv", # noqa: E501
@@ -209,15 +239,29 @@ class LongCatFlashMTP(nn.Module):
params_dict = dict(self.named_parameters())
loaded_params: set[str] = set()
_pending_wk_fp8: dict = {}
for name, loaded_weight in weights:
if "rotary_emb.inv_freq" in name:
continue
# MTP embeds plain tokens (mtp_disable_over_tokenizer); its
# checkpoint n-gram tables are unused.
if "ngram_embeddings" in name:
continue
spec_layer = self.get_spec_layer_idx_from_weight_name(self.config, name)
if spec_layer is None:
continue
name = self._rewrite_spec_layer_name(
spec_layer, name, new_to_old_names_mapping
)
if _try_load_fp8_indexer_wk(
name,
loaded_weight,
_pending_wk_fp8,
params_dict,
loaded_params,
[],
):
continue
for param_name, weight_name, shard_id in stacked_params_mapping:
# Skip non-stacked layers and experts (experts handled below).
if weight_name not in name:
@@ -258,6 +302,21 @@ class LongCatFlashMTP(nn.Module):
):
continue
# Fold the MLA LoRA scaling at load time (see
# FlashModel.load_weights).
if name.endswith(".q_a_layernorm.weight") and getattr(
self.config, "mla_scale_q_lora", False
):
loaded_weight = loaded_weight.float() * (
(self.config.hidden_size / self.config.q_lora_rank) ** 0.5
)
elif name.endswith(".kv_a_layernorm.weight") and getattr(
self.config, "mla_scale_kv_lora", False
):
loaded_weight = loaded_weight.float() * (
(self.config.hidden_size / self.config.kv_lora_rank) ** 0.5
)
param = params_dict[name]
weight_loader = getattr(param, "weight_loader", default_weight_loader)
weight_loader(param, loaded_weight)
@@ -287,14 +346,6 @@ class LongCatFlashMTP(nn.Module):
).split([self_attn.qk_nope_head_dim, self_attn.v_head_dim], dim=1)
self_attn.w_kc = w_kc.transpose(1, 2).contiguous().transpose(1, 2)
self_attn.w_vc = w_vc.contiguous().transpose(1, 2)
if self.config.mla_scale_q_lora:
self_attn.q_a_layernorm.weight.data *= (
self.config.hidden_size / self.config.q_lora_rank
) ** 0.5
if self.config.mla_scale_kv_lora:
self_attn.kv_a_layernorm.weight.data *= (
self.config.hidden_size / self.config.kv_lora_rank
) ** 0.5
return loaded_params
def _rewrite_spec_layer_name(
@@ -0,0 +1,435 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
"""Inference-only LongCat-Flash-Lite (n-gram embedding) model.
``LongcatFlashNgramForCausalLM`` is LongCat-Flash (MLA dual-attention +
zero-expert MoE + YaRN) plus an n-gram embedding input layer: each position's
embedding fuses the token embedding with hashed embeddings of the preceding
``n`` tokens. That per-request token history is isolated in a Model-Runner-V2
:class:`LongcatNgramModelState` (mirroring ``DiffusionGemmaModelState``), so
``get_model_state_cls`` makes the model MRV2-only.
"""
from collections.abc import Iterable
from typing import Any
import torch
from torch import nn
from vllm import _custom_ops as ops
from vllm.config import VllmConfig
from vllm.distributed import get_pp_group
from vllm.model_executor.layers.logits_processor import LogitsProcessor
from vllm.model_executor.layers.vocab_parallel_embedding import (
ParallelLMHead,
VocabParallelEmbedding,
)
from vllm.v1.worker.gpu.input_batch import InputBatch
from vllm.v1.worker.gpu.model_states.default import DefaultModelState
from vllm.v1.worker.gpu.states import RequestState
from .interfaces import SupportsLoRA, SupportsPP
from .longcat_flash import FlashConfig, FlashModel
from .utils import AutoWeightsLoader, PPMissingLayer, maybe_prefix
def uses_ngram_embedding(config: FlashConfig) -> bool:
return getattr(config, "ngram_vocab_size_ratio", None) is not None
def compute_eos_position_ngram_ids(
ngram: "NgramEmbedding",
eos_id: int,
table: torch.Tensor,
tok_req: torch.Tensor,
col: torch.Tensor,
eos_tok: torch.Tensor,
) -> torch.Tensor:
"""Hash ids for positions whose current token is EOS.
The ``ngram_compute_n_gram_ids`` kernel stops the n-gram walk at a negated
(EOS) table entry even at delta 0, but per the reference semantics an EOS
*current* token hashes normally with its full look-back; only later
positions' look-back stops at it. This recomputes those (rare) ids,
mirroring the kernel's per-config walk with delta 0 forced to the
un-negated EOS id.
"""
device = table.device
n, k = ngram.n, ngram.k
num_emb = ngram.num_embedders
deltas = torch.arange(n, device=device)
back = table[tok_req[eos_tok, None], col[eos_tok, None] - deltas] # [E, n]
back[:, 0] = eos_id
valid = (back >= 0).cumprod(dim=1).bool()
toks = torch.where(valid, back, 0).to(torch.int64) # [E, n]
w = ngram.ne_weights.view(num_emb, n).to(torch.int64) # [cfg, n]
m = ngram.ne_mods.view(num_emb).to(torch.int64) # [cfg]
# Config i*k+j is an (i+2)-gram: only deltas < i+2 participate.
cfg_n = torch.arange(num_emb, device=device) // k + 2
dmask = deltas[None, :] < cfg_n[:, None] # [cfg, n]
terms = (toks[:, None, :] * w[None]) % m[None, :, None] * dmask[None]
h = terms.sum(-1) % m[None]
return (h + ngram.exclusive_sizes[:-1].to(torch.int64)).to(torch.int32)
def _config_dtype(config: FlashConfig) -> torch.dtype:
dt = getattr(config, "torch_dtype", None) or getattr(config, "dtype", None)
if isinstance(dt, torch.dtype):
return dt
return getattr(torch, str(dt), None) or torch.bfloat16
class NgramEmbedding(nn.Module):
"""Token embedding fused with hashed n-gram embeddings.
TP-sharded: the ``k*(n-1)`` per-embedder tables are concatenated into one
:class:`VocabParallelEmbedding` (``oe_embedder``) with per-embedder offsets,
and the projections are stacked into one ``oe_projection`` applied with a
single ``bmm``. Hashing math is ported from the HF reference.
"""
def __init__(self, config: FlashConfig, base_embeddings: nn.Module) -> None:
super().__init__()
self.config = config
self.word_embeddings = base_embeddings
self.m = config.ngram_vocab_size_ratio * config.vocab_size
self.k = config.emb_split_num
self.n = config.emb_neighbor_num
self.pad_id = config.pad_token_id
self.eos_token_id = config.eos_token_id
self._dtype = _config_dtype(config)
self._init_ngram_embeddings()
def _init_ngram_embeddings(self) -> None:
self.num_embedders = self.k * (self.n - 1)
oe_dim = self.config.hidden_size // self.num_embedders
self.oe_dim = oe_dim
# Exclusive prefix sums of per-embedder table sizes; each embedder's
# local id is offset into the single concatenated table.
sizes = [int(self.m + i * 2 + 1) for i in range(self.num_embedders)]
offsets = [0]
for s in sizes:
offsets.append(offsets[-1] + s)
self._offsets = offsets # len num_embedders + 1
self._sizes = sizes
self.oe_embedder = VocabParallelEmbedding(
offsets[-1], oe_dim, params_dtype=self._dtype
)
# Stacked projections: oe_projection[i] = post_projs[i].weight.T
self.oe_projection = nn.Parameter(
torch.empty(
self.num_embedders, oe_dim, self.config.hidden_size, dtype=self._dtype
),
requires_grad=False,
)
# Precomputed tables for the CUDA n-gram id kernel (ngram_embedding
# _kernels.cu): ne_weights[i][j][delta] = vocab^delta mod ne_mods[i][j],
# ne_mods[i][j] = m + 2*(i*k+j) + 1. Registered as non-persistent buffers
# so they follow the module to the device (not part of the checkpoint).
vocab = self.config.vocab_size
ne_weights = torch.zeros(self.n - 1, self.k, self.n, dtype=torch.int32)
ne_mods = torch.zeros(self.n - 1, self.k, dtype=torch.int32)
for i in range(self.n - 1):
for j in range(self.k):
mod = int(self.m + 2 * (i * self.k + j) + 1)
ne_mods[i, j] = mod
for delta in range(self.n):
ne_weights[i, j, delta] = pow(vocab, delta, mod)
self.register_buffer("ne_weights", ne_weights, persistent=False)
self.register_buffer("ne_mods", ne_mods, persistent=False)
self.register_buffer(
"exclusive_sizes",
torch.tensor(offsets, dtype=torch.int32),
persistent=False,
)
def load_weight(self, weight_name: str, loaded_weight: torch.Tensor) -> str:
"""Split a per-embedder checkpoint weight into the sharded layout.
Returns the destination parameter's qualified name (relative to the
enclosing model) so the caller can mark it loaded for completeness
checks.
"""
if "ngram_embeddings.embedders." in weight_name:
index = int(
weight_name.split("ngram_embeddings.embedders.")[1].split(".")[0]
)
lo, hi = self._offsets[index], self._offsets[index + 1]
assert hi - lo == loaded_weight.shape[0], (
f"{hi - lo=} {loaded_weight.shape[0]=}"
)
shard = self.oe_embedder.shard_indices
tp_start, tp_end = shard.org_vocab_start_index, shard.org_vocab_end_index
load_start, load_end = max(lo, tp_start), min(hi, tp_end)
if load_start < load_end:
self.oe_embedder.weight.data[
load_start - tp_start : load_end - tp_start
] = loaded_weight[load_start - lo : load_end - lo]
return "ngram_embeddings.oe_embedder.weight"
elif "ngram_embeddings.post_projs." in weight_name:
index = int(
weight_name.split("ngram_embeddings.post_projs.")[1].split(".")[0]
)
self.oe_projection.data[index].copy_(loaded_weight.t())
return "ngram_embeddings.oe_projection"
else:
raise AssertionError(f"Unexpected ngram weight: {weight_name}")
def embed_batched(
self, input_ids: torch.Tensor, oe_ids: torch.Tensor
) -> torch.Tensor:
"""Fused n-gram embedding for a flat batch given precomputed ids.
Args:
input_ids: ``[num_tokens]`` current token per position.
oe_ids: ``[num_tokens, num_embedders]`` global (offset) n-gram ids,
as produced by the ``ngram_compute_n_gram_ids`` kernel.
Returns: ``[num_tokens, hidden]``.
"""
word = self.word_embeddings(input_ids) # [N, H]
flat = oe_ids.permute(1, 0).contiguous() # [num_embedders, N]
oe = self.oe_embedder(flat) # [num_embedders, N, oe_dim]
proj = torch.bmm(oe, self.oe_projection) # [num_embedders, N, H]
all_h = torch.cat([word.unsqueeze(0), proj], dim=0) # [ne+1, N, H]
return all_h.mean(dim=0) # [N, H]
class FlashNgramModel(FlashModel):
"""FlashModel whose input embedding is an :class:`NgramEmbedding`."""
def __init__(self, *, vllm_config: VllmConfig, prefix: str = "") -> None:
# Each FlashDecoderLayer is a *dual* layer (2 attentions), so the number
# of decoder layers is ``num_layers``. The ngram HF config sets
# ``num_hidden_layers`` to a multiple of that (attention-module count),
# which FlashModel would otherwise build as too many (dead) layers.
hf = vllm_config.model_config.hf_config
num_layers = getattr(hf, "num_layers", None)
if num_layers is not None and hf.num_hidden_layers != num_layers:
hf.num_hidden_layers = num_layers
super().__init__(vllm_config=vllm_config, prefix=prefix)
if get_pp_group().is_first_rank and uses_ngram_embedding(self.config):
self.ngram_embeddings = NgramEmbedding(self.config, self.embed_tokens)
else:
self.ngram_embeddings = None
def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]:
# Names arrive with the ``model.`` prefix already stripped (routed here
# by AutoWeightsLoader). Split the concatenated/sharded ngram tables and
# stacked projections; delegate everything else to FlashModel.
loaded: set[str] = set()
rest: list[tuple[str, torch.Tensor]] = []
for name, w in weights:
if "ngram_embeddings." in name:
# Drop the checkpoint's n-gram tables when the module is
# disabled (e.g. ngram_vocab_size_ratio overridden to None).
if self.ngram_embeddings is not None:
loaded.add(self.ngram_embeddings.load_weight(name, w))
else:
rest.append((name, w))
loaded |= super().load_weights(rest)
return loaded
class LongcatFlashNgramForCausalLM(nn.Module, SupportsLoRA, SupportsPP):
"""LongCat-Flash-Lite for causal LM (MRV2-only, n-gram embedding)."""
packed_modules_mapping = {
"qkv_proj": ["q_proj", "k_proj", "v_proj"],
"gate_up_proj": ["gate_proj", "up_proj"],
}
def __init__(self, *, vllm_config: VllmConfig, prefix: str = "") -> None:
super().__init__()
if not vllm_config.use_v2_model_runner:
raise NotImplementedError(
"LongcatFlashNgramForCausalLM (LongCat-Flash-Lite) requires the "
"V2 model runner for its n-gram embedding state; it is selected "
"automatically unless VLLM_USE_V2_MODEL_RUNNER=0 is set."
)
config = FlashConfig(**vllm_config.model_config.hf_config.__dict__)
config.intermediate_size = getattr(
config, "ffn_hidden_size", config.intermediate_size
)
self.config = config
self.quant_config = vllm_config.quant_config
self.model = FlashNgramModel(
vllm_config=vllm_config, prefix=maybe_prefix(prefix, "model")
)
if get_pp_group().is_last_rank:
self.lm_head = ParallelLMHead(
config.vocab_size,
config.hidden_size,
quant_config=self.quant_config,
prefix=maybe_prefix(prefix, "lm_head"),
)
else:
self.lm_head = PPMissingLayer()
self.logits_processor = LogitsProcessor(config.vocab_size)
self.make_empty_intermediate_tensors = (
self.model.make_empty_intermediate_tensors
)
@staticmethod
def get_model_state_cls() -> type["LongcatNgramModelState"]:
return LongcatNgramModelState
def embed_input_ids(self, input_ids: torch.Tensor) -> torch.Tensor:
return self.model.embed_input_ids(input_ids)
def forward(
self,
input_ids: torch.Tensor | None,
positions: torch.Tensor,
intermediate_tensors=None,
inputs_embeds: torch.Tensor | None = None,
):
# inputs_embeds is produced by LongcatNgramModelState.prepare_inputs.
return self.model(input_ids, positions, intermediate_tensors, inputs_embeds)
def compute_logits(self, hidden_states: torch.Tensor) -> torch.Tensor | None:
return self.logits_processor(self.lm_head, hidden_states)
def get_expert_mapping(self):
return self.model.get_expert_mapping()
def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]:
# AutoWeightsLoader routes ``model.*`` to FlashNgramModel.load_weights
# (which handles the ngram split) and ``lm_head.*`` to the head. MTP
# weights are not part of this model.
loader = AutoWeightsLoader(self, skip_prefixes=["model.mtp."])
return loader.load_weights(weights)
class LongcatNgramModelState(DefaultModelState):
"""n-gram input embedding state for LongCat-Flash models.
``prepare_inputs`` computes the fused n-gram embedding for the batch into
a persistent ``inputs_embeds`` buffer handed to the model forward. Each
position's left-context is gathered from the runner's authoritative
``all_token_ids`` history, which keeps it correct under chunked prefill,
request resumption, and speculative decoding (rejected draft tokens never
enter the history).
"""
def __init__(self, vllm_config, model, encoder_cache, device) -> None:
super().__init__(vllm_config, model, encoder_cache, device)
config = model.config
self.ngram = model.model.ngram_embeddings
self.n = int(config.emb_neighbor_num)
self.ctx_len = self.n - 1
self.eos_id = int(config.eos_token_id)
self.device = device
self._inputs_embeds_buf = torch.zeros(
self.max_num_tokens,
config.hidden_size,
dtype=self.dtype,
device=device,
)
def prepare_inputs(
self, input_batch: InputBatch, req_states: RequestState
) -> dict[str, Any]:
model_inputs = super().prepare_inputs(input_batch, req_states) # positions
num_tokens = input_batch.num_tokens
num_padded = input_batch.num_tokens_after_padding
input_ids = input_batch.input_ids[:num_tokens]
embeds = self._inputs_embeds_buf[:num_padded]
oe_ids = self._compute_oe_ids(input_batch, req_states)
embeds[:num_tokens].copy_(self.ngram.embed_batched(input_ids, oe_ids))
model_inputs["inputs_embeds"] = embeds
return model_inputs
def prepare_dummy_inputs(self, num_reqs: int, num_tokens: int) -> dict[str, Any]:
# FULL cudagraph replay reads only the captured buffers, so capture must
# reference the same persistent ``inputs_embeds`` buffer prepare_inputs
# re-fills (the base class wires this for multimodal models only).
model_inputs = super().prepare_dummy_inputs(num_reqs, num_tokens) # positions
model_inputs["inputs_embeds"] = self._inputs_embeds_buf[:num_tokens]
return model_inputs
def _compute_oe_ids(
self, input_batch: InputBatch, req_states: RequestState
) -> torch.Tensor:
"""Batched global n-gram ids ``[num_tokens, num_embedders]``.
Assembles an ephemeral per-request token table (``[n-1] context ++
current tokens``, EOS-negated) and runs the ``ngram_compute_n_gram_ids``
CUDA kernel for the whole batch. The left-context is gathered from the
authoritative ``all_token_ids`` history each step (rather than rolled
incrementally) so rejected speculative-draft tokens never pollute it.
"""
device = self.device
num_tokens = input_batch.num_tokens
num_reqs = input_batch.num_reqs
ctx_len = self.ctx_len
idx_mapping = input_batch.idx_mapping[:num_reqs].long()
qsl = input_batch.query_start_loc[: num_reqs + 1].to(torch.int32)
cur = input_batch.input_ids[:num_tokens].to(torch.int32)
cur_neg = torch.where(cur == self.eos_id, -cur, cur)
req_lens = qsl[1:] - qsl[:-1]
max_len = int(req_lens.max().item())
width = ctx_len + max_len
# Left-context: the ctx_len accepted tokens preceding this batch's
# first position, EOS-negated; -1 marks the sequence start.
p0 = req_states.num_computed_tokens.gpu[idx_mapping].long() # [R]
ctx_pos = p0[:, None] + torch.arange(-ctx_len, 0, device=device) # [R, C]
in_range = ctx_pos >= 0
ctx = req_states.all_token_ids.gpu[
idx_mapping[:, None], ctx_pos.clamp_min(0)
].to(torch.int32)
ctx = torch.where(ctx == self.eos_id, -ctx, ctx)
ctx = torch.where(in_range, ctx, ctx.new_full((), -1))
# table[r] = [context(n-1) | current tokens | pad(-1)]
table = torch.full((num_reqs, width), -1, dtype=torch.int32, device=device)
table[:, :ctx_len] = ctx
tok_req = torch.repeat_interleave(
torch.arange(num_reqs, device=device), req_lens.long()
)
col = ctx_len + (
torch.arange(num_tokens, device=device) - qsl[:-1].long()[tok_req]
)
table[tok_req, col] = cur_neg
column_starts = torch.full(
(num_reqs,), ctx_len, dtype=torch.int32, device=device
)
row_indices = torch.arange(num_reqs, dtype=torch.int64, device=device)
n_gram_ids = torch.empty(
num_tokens, self.ngram.num_embedders, dtype=torch.int32, device=device
)
ops.ngram_compute_n_gram_ids(
self.n,
self.ngram.k,
self.ngram.ne_weights,
self.ngram.ne_mods,
self.ngram.exclusive_sizes,
qsl,
table,
row_indices,
column_starts,
n_gram_ids,
)
# The kernel stops the n-gram walk at a negated (EOS) table entry even
# at delta 0, but per the reference semantics an EOS *current* token
# hashes normally with its full look-back; only later positions'
# look-back stops at it. Recompute the (rare) EOS-position ids here.
eos_tok = (cur == self.eos_id).nonzero(as_tuple=True)[0]
if eos_tok.numel():
n_gram_ids[eos_tok] = compute_eos_position_ngram_ids(
self.ngram, self.eos_id, table, tok_req, col, eos_tok
)
return n_gram_ids.long()
+80 -51
View File
@@ -2,62 +2,91 @@
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
from collections.abc import Iterable
import regex as re
import regex
import torch
from vllm.model_executor.models.deepseek_v2 import DeepseekV3ForCausalLM
from vllm.model_executor.models.utils import AutoWeightsLoader, WeightsMapper
class MistralLarge3ForCausalLM(DeepseekV3ForCausalLM):
# fmt: off
remapping = {
r"layers\.(\d+)\.attention_norm\.weight": r"model.layers.\1.input_layernorm.weight", # noqa: E501
r"layers\.(\d+)\.attention\.wq_a\.(\w+)": r"model.layers.\1.self_attn.q_a_proj.\2", # noqa: E501
r"layers\.(\d+)\.attention\.q_a_norm\.weight": r"model.layers.\1.self_attn.q_a_layernorm.weight", # noqa: E501
r"layers\.(\d+)\.attention\.wq_b\.(\w+)": r"model.layers.\1.self_attn.q_b_proj.\2", # noqa: E501
r"layers\.(\d+)\.attention\.wkv_a_with_mqa\.(\w+)": r"model.layers.\1.self_attn.kv_a_proj_with_mqa.\2", # noqa: E501
r"layers\.(\d+)\.attention\.kv_a_norm\.weight": r"model.layers.\1.self_attn.kv_a_layernorm.weight", # noqa: E501
r"layers\.(\d+)\.attention\.wkv_b\.(\w+)": r"model.layers.\1.self_attn.kv_b_proj.\2", # noqa: E501
r"layers\.(\d+)\.attention\.wo\.(\w+)": r"model.layers.\1.self_attn.o_proj.\2", # noqa: E501
r"layers\.(\d+)\.ffn_norm\.weight": r"model.layers.\1.post_attention_layernorm.weight", # noqa: E501
r"layers\.(\d+)\.feed_forward\.w1\.(\w+)": r"model.layers.\1.mlp.gate_proj.\2", # noqa: E501
r"layers\.(\d+)\.feed_forward\.w2\.(\w+)": r"model.layers.\1.mlp.down_proj.\2", # noqa: E501
r"layers\.(\d+)\.feed_forward\.w3\.(\w+)": r"model.layers.\1.mlp.up_proj.\2", # noqa: E501
r"layers\.(\d+)\.gate\.weight": r"model.layers.\1.mlp.gate.weight", # noqa: E501
r"layers\.(\d+)\.shared_experts\.w1\.(\w+)": r"model.layers.\1.mlp.shared_experts.gate_proj.\2", # noqa: E501
r"layers\.(\d+)\.shared_experts\.w2\.(\w+)": r"model.layers.\1.mlp.shared_experts.down_proj.\2", # noqa: E501
r"layers\.(\d+)\.shared_experts\.w3\.(\w+)": r"model.layers.\1.mlp.shared_experts.up_proj.\2", # noqa: E501
r"layers\.(\d+)\.experts\.(\d+)\.w1\.(\w+)": r"model.layers.\1.mlp.experts.\2.gate_proj.\3", # noqa: E501
r"layers\.(\d+)\.experts\.(\d+)\.w2\.(\w+)": r"model.layers.\1.mlp.experts.\2.down_proj.\3", # noqa: E501
r"layers\.(\d+)\.experts\.(\d+)\.w3\.(\w+)": r"model.layers.\1.mlp.experts.\2.up_proj.\3", # noqa: E501
r"norm\.weight": "model.norm.weight", # noqa: E501
r"tok_embeddings\.weight": "model.embed_tokens.weight", # noqa: E501
r"output\.weight": "lm_head.weight", # noqa: E501
}
# fmt: on
# WeightsMapper applies all matching patterns sequentially (no break on first
# match). This is safe here because every pattern is anchored at both ends
# (\A...\Z) and after substitution the resulting key always starts with
# "model." or "lm_head.", so no later pattern can accidentally match again.
hf_to_vllm_mapper = WeightsMapper(
orig_to_new_regex={ # noqa: B950
regex.compile(
r"\Alayers\.(\d+)\.attention_norm\.weight\Z"
): r"model.layers.\1.input_layernorm.weight",
regex.compile(
r"\Alayers\.(\d+)\.attention\.wq_a\.(\w+)\Z"
): r"model.layers.\1.self_attn.q_a_proj.\2",
regex.compile(
r"\Alayers\.(\d+)\.attention\.q_a_norm\.weight\Z"
): r"model.layers.\1.self_attn.q_a_layernorm.weight",
regex.compile(
r"\Alayers\.(\d+)\.attention\.wq_b\.(\w+)\Z"
): r"model.layers.\1.self_attn.q_b_proj.\2",
regex.compile(
r"\Alayers\.(\d+)\.attention\.wkv_a_with_mqa\.(\w+)\Z"
): r"model.layers.\1.self_attn.kv_a_proj_with_mqa.\2",
regex.compile(
r"\Alayers\.(\d+)\.attention\.kv_a_norm\.weight\Z"
): r"model.layers.\1.self_attn.kv_a_layernorm.weight",
regex.compile(
r"\Alayers\.(\d+)\.attention\.wkv_b\.(\w+)\Z"
): r"model.layers.\1.self_attn.kv_b_proj.\2",
regex.compile(
r"\Alayers\.(\d+)\.attention\.wo\.(\w+)\Z"
): r"model.layers.\1.self_attn.o_proj.\2",
regex.compile(
r"\Alayers\.(\d+)\.ffn_norm\.weight\Z"
): r"model.layers.\1.post_attention_layernorm.weight",
regex.compile(
r"\Alayers\.(\d+)\.feed_forward\.w1\.(\w+)\Z"
): r"model.layers.\1.mlp.gate_proj.\2",
regex.compile(
r"\Alayers\.(\d+)\.feed_forward\.w2\.(\w+)\Z"
): r"model.layers.\1.mlp.down_proj.\2",
regex.compile(
r"\Alayers\.(\d+)\.feed_forward\.w3\.(\w+)\Z"
): r"model.layers.\1.mlp.up_proj.\2",
regex.compile(
r"\Alayers\.(\d+)\.gate\.weight\Z"
): r"model.layers.\1.mlp.gate.weight",
regex.compile(
r"\Alayers\.(\d+)\.shared_experts\.w1\.(\w+)\Z"
): r"model.layers.\1.mlp.shared_experts.gate_proj.\2",
regex.compile(
r"\Alayers\.(\d+)\.shared_experts\.w2\.(\w+)\Z"
): r"model.layers.\1.mlp.shared_experts.down_proj.\2",
regex.compile(
r"\Alayers\.(\d+)\.shared_experts\.w3\.(\w+)\Z"
): r"model.layers.\1.mlp.shared_experts.up_proj.\2",
regex.compile(
r"\Alayers\.(\d+)\.experts\.(\d+)\.w1\.(\w+)\Z"
): r"model.layers.\1.mlp.experts.\2.gate_proj.\3",
regex.compile(
r"\Alayers\.(\d+)\.experts\.(\d+)\.w2\.(\w+)\Z"
): r"model.layers.\1.mlp.experts.\2.down_proj.\3",
regex.compile(
r"\Alayers\.(\d+)\.experts\.(\d+)\.w3\.(\w+)\Z"
): r"model.layers.\1.mlp.experts.\2.up_proj.\3",
regex.compile(r"\Anorm\.weight\Z"): "model.norm.weight",
regex.compile(r"\Atok_embeddings\.weight\Z"): "model.embed_tokens.weight",
regex.compile(r"\Aoutput\.weight\Z"): "lm_head.weight",
},
orig_to_new_suffix={
".qscale_act": ".input_scale",
".qscale_weight": ".weight_scale",
},
)
# Bypass super().load_weights() and construct AutoWeightsLoader(self)
# directly (same pattern as Qwen2ForCausalLM). Any logic in the parent
# class's load_weights is a thin wrapper around AutoWeightsLoader, and
# we must apply hf_to_vllm_mapper before the loader walks the tree.
def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]:
return super().load_weights(map(self._remap_mistral_to_ds, weights))
def _remap_mistral_to_ds(
self, weight: tuple[str, torch.Tensor]
) -> tuple[str, torch.Tensor]:
"""Remap Mistral parameters to DeepseekV2 parameters."""
name, loaded_weight = weight
for k, v in self.remapping.items():
match = re.fullmatch(k, name)
if match:
name = re.sub(k, v, name)
break
else:
raise ValueError(f"Cannot remap {name}")
# Remapping scale names. We could do this in the regex above but it
# would triple the number of lines for most layers.
if name.endswith(".qscale_act"):
name = re.sub(r"\.qscale_act$", ".input_scale", name)
elif name.endswith(".qscale_weight"):
name = re.sub(r"\.qscale_weight$", ".weight_scale", name)
return name, loaded_weight
loader = AutoWeightsLoader(self)
return loader.load_weights(weights, mapper=self.hf_to_vllm_mapper)
@@ -5,6 +5,7 @@ import copy
from collections.abc import Iterable
from functools import partial
import regex
import torch
import torch.nn as nn
@@ -22,7 +23,7 @@ from vllm.model_executor.models.deepseek_v2 import (
from vllm.model_executor.models.mistral_large_3 import MistralLarge3ForCausalLM
from .interfaces import SupportsMultiModal
from .utils import make_empty_intermediate_tensors_factory, maybe_prefix
from .utils import WeightsMapper, make_empty_intermediate_tensors_factory, maybe_prefix
logger = init_logger(__name__)
@@ -107,11 +108,13 @@ class EagleMistralLarge3Model(DeepseekV2Model):
class EagleMistralLarge3ForCausalLM(MistralLarge3ForCausalLM):
remapping = MistralLarge3ForCausalLM.remapping | {
r"eagle_linear\.weight": r"model.fc.weight",
r"eagle_linear\.qscale_act": r"model.fc.input_scale",
r"eagle_linear\.qscale_weight": r"model.fc.weight_scale",
}
hf_to_vllm_mapper = MistralLarge3ForCausalLM.hf_to_vllm_mapper | WeightsMapper(
orig_to_new_regex={
regex.compile(r"\Aeagle_linear\.weight\Z"): r"model.fc.weight",
regex.compile(r"\Aeagle_linear\.qscale_act\Z"): r"model.fc.input_scale",
regex.compile(r"\Aeagle_linear\.qscale_weight\Z"): r"model.fc.weight_scale",
},
)
def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""):
target_layer_num = vllm_config.model_config.get_num_layers(
@@ -754,6 +754,14 @@ class DFlashQwen3ForCausalLM(Qwen3ForCausalLM):
needs_squeeze = hidden_states.dim() == 1
if needs_squeeze:
hidden_states = hidden_states.unsqueeze(0)
expected = self.model.fc.input_size
if hidden_states.shape[-1] != expected:
raise ValueError(
f"DFlash drafter expects {expected} concatenated aux hidden "
f"features but received {hidden_states.shape[-1]}. This usually "
"means the draft model's target_layer_ids reference layers that "
"do not exist in the target model (incompatible draft/target pair)."
)
result = self.model.fc(hidden_states)
if needs_squeeze:
result = result.squeeze(0)
+5 -11
View File
@@ -381,15 +381,15 @@ class Qwen3NextAttention(nn.Module):
def forward(
self,
positions: torch.Tensor,
output: torch.Tensor,
hidden_states: torch.Tensor,
):
) -> torch.Tensor:
qkv, _ = self.qkv_proj(hidden_states)
q, k, v, gate = self._project_qkv_gate(qkv, positions)
attn_output = self.attn(q, k, v)
if gate is not None:
attn_output = attn_output * torch.sigmoid(gate)
output[:], _ = self.o_proj(attn_output)
output, _ = self.o_proj(attn_output)
return output
class Qwen3NextDecoderLayer(nn.Module):
@@ -484,21 +484,15 @@ class Qwen3NextDecoderLayer(nn.Module):
else:
hidden_states, residual = self.input_layernorm(hidden_states, residual)
self_attention_output = torch.empty_like(hidden_states)
if self.layer_type == "linear_attention":
self.linear_attn(
hidden_states=hidden_states,
output=self_attention_output,
)
hidden_states = self.linear_attn(hidden_states=hidden_states)
elif self.layer_type == "full_attention":
self.self_attn(
hidden_states = self.self_attn(
hidden_states=hidden_states,
output=self_attention_output,
positions=positions,
)
else:
raise ValueError("Invalid layer_type")
hidden_states = self_attention_output
if self.layer_scale:
if len(hidden_states.shape) == 2:
+9
View File
@@ -145,6 +145,15 @@ _TEXT_GENERATION_MODELS = {
# For decapoda-research/llama-*
"LLaMAForCausalLM": ("llama", "LlamaForCausalLM"),
"LongcatFlashForCausalLM": ("longcat_flash", "LongcatFlashForCausalLM"),
"LongcatFlashNgramForCausalLM": (
"longcat_flash_ngram",
"LongcatFlashNgramForCausalLM",
),
# LongCat-2.0 (LongCat-Flash + n-gram embedding + LongCat Sparse Attention)
"LongcatCausalLM": (
"longcat_flash_ngram",
"LongcatFlashNgramForCausalLM",
),
"MambaForCausalLM": ("mamba", "MambaForCausalLM"),
"Mamba2ForCausalLM": ("mamba2", "Mamba2ForCausalLM"),
"MellumForCausalLM": ("mellum", "MellumForCausalLM"),
+4 -4
View File
@@ -15,16 +15,16 @@ from .quant_config import DeepseekV4FP8Config
# default that mypy sees; the ROCm/XPU branches override at runtime and are
# kept type-compatible via ``# type: ignore[assignment]``.
if current_platform.is_rocm():
from .amd.dspark import ( # type: ignore[assignment]
DSparkDeepseekV4ForCausalLM,
)
from .amd.model import DeepseekV4ForCausalLM
from .amd.mtp import DeepSeekV4MTP
# DSpark is NVIDIA-only for now.
DSparkDeepseekV4ForCausalLM = None # type: ignore[assignment]
elif current_platform.is_xpu():
from .xpu.model import DeepseekV4ForCausalLM # type: ignore[assignment]
from .xpu.mtp import DeepSeekV4MTP # type: ignore[assignment]
DSparkDeepseekV4ForCausalLM = None # type: ignore[assignment]
DSparkDeepseekV4ForCausalLM = None # type: ignore[assignment, misc]
else:
from .nvidia.dspark import ( # type: ignore[assignment]
DSparkDeepseekV4ForCausalLM,
+499
View File
@@ -0,0 +1,499 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
"""DSpark draft model for DeepSeek-V4 on ROCm/AMD (gfx950).
ROCm port of ``nvidia/dspark.py``. Follows the same nvidia->amd recipe used for
``amd/mtp.py``:
* import ``DeepseekV4DecoderLayer`` from the AMD ``.model`` (aiter/triton
attention + MHC CustomOp path) instead of the nvidia one;
* route the MHC head through the ``HCHeadOp`` CustomOp dispatcher (aiter /
tilelang / triton / torch) instead of calling the tilelang kernels directly,
and gate the trailing ``mhc_post`` on ``use_fused_mhc`` (False on the aiter
path, where the decoder layer already applies hc_post in-layer);
* drop the mega-MoE weight path (``make_deepseek_v4_expert_params_mapping`` /
``use_mega_moe`` / ``finalize_mega_moe_weights`` do not exist in amd/model.py).
Everything else the semi-autoregressive drafting hooks, the Markov head, the
sliding-window context-KV insert, and the checkpoint ``mtp.*`` weight remap is
pure torch / Triton and shared with the nvidia implementation unchanged.
"""
from collections.abc import Iterable
import regex as re
import torch
import torch.nn as nn
from vllm.config import VllmConfig, get_current_vllm_config
from vllm.distributed import (
get_tensor_model_parallel_rank,
get_tensor_model_parallel_world_size,
)
from vllm.logger import init_logger
from vllm.model_executor.layers.fused_moe import (
fused_moe_make_expert_params_mapping,
)
from vllm.model_executor.layers.layernorm import RMSNorm
from vllm.model_executor.layers.linear import ReplicatedLinear
from vllm.model_executor.layers.logits_processor import LogitsProcessor
from vllm.model_executor.layers.mhc import HCHeadOp
from vllm.model_executor.layers.vocab_parallel_embedding import (
ParallelLMHead,
VocabParallelEmbedding,
)
from vllm.model_executor.model_loader.weight_utils import default_weight_loader
from vllm.model_executor.models.qwen3_dspark import (
DSparkMarkovHead,
)
from vllm.model_executor.models.utils import maybe_prefix
from .model import (
DeepseekV4DecoderLayer,
)
logger = init_logger(__name__)
# MoE expert scale suffix differs by expert dtype (mirrors deepseek_v4 loaders):
# fp4 experts register ``.weight_scale``; block-fp8 experts ``.weight_scale_inv``.
_EXPERT_SCALE_RE = re.compile(r"\.experts\.\d+\.w[123]\.scale$")
class DSparkDeepseekV4Model(nn.Module):
def __init__(self, *, vllm_config: VllmConfig, prefix: str = "") -> None:
super().__init__()
assert vllm_config.speculative_config is not None
config = vllm_config.speculative_config.draft_model_config.hf_config
self.config = config
self.hidden_size = config.hidden_size
self.hc_mult = config.hc_mult
self.hc_eps = config.hc_eps
self.rms_norm_eps = config.rms_norm_eps
self.num_hidden_layers = config.num_hidden_layers
self.target_layer_ids = tuple(config.dspark_target_layer_ids)
self.num_dspark_layers = getattr(config, "n_mtp_layers", None) or 3
# Shared with the target (aliased by the speculator's loading utility).
self.embed_tokens = VocabParallelEmbedding(
config.vocab_size,
config.hidden_size,
prefix=maybe_prefix(prefix, "embed_tokens"),
)
self.main_proj = ReplicatedLinear(
config.hidden_size * len(self.target_layer_ids),
config.hidden_size,
bias=False,
return_bias=False,
quant_config=vllm_config.quant_config,
prefix=maybe_prefix(prefix, "main_proj"),
)
self.main_norm = RMSNorm(config.hidden_size, eps=config.rms_norm_eps)
current_vllm_config = get_current_vllm_config()
self.layers = nn.ModuleList(
[
DeepseekV4DecoderLayer(
current_vllm_config,
prefix=maybe_prefix(prefix, f"layers.{self.num_hidden_layers + i}"),
)
for i in range(self.num_dspark_layers)
]
)
# Heads: final norm + hc_head, and the Markov head
# Loaded from the "final" MTP layer weights (mtp.*) in the target checkpoint
self.norm = RMSNorm(config.hidden_size, eps=config.rms_norm_eps)
hc_dim = self.hc_mult * config.hidden_size
self.hc_head_fn = nn.Parameter(
torch.empty(self.hc_mult, hc_dim, dtype=torch.float32),
requires_grad=False,
)
self.hc_head_base = nn.Parameter(
torch.empty(self.hc_mult, dtype=torch.float32), requires_grad=False
)
self.hc_head_scale = nn.Parameter(
torch.empty(1, dtype=torch.float32), requires_grad=False
)
draft_vocab_size = (
getattr(config, "draft_vocab_size", None) or config.vocab_size
)
self.markov_head = DSparkMarkovHead(
config.vocab_size,
draft_vocab_size,
config.dspark_markov_rank,
prefix=maybe_prefix(prefix, "markov_head"),
)
# MHC head CustomOp dispatcher (aiter / tilelang / triton / torch),
# replacing the direct nvidia tilelang kernel call.
self.hc_head_op = HCHeadOp()
def embed_input_ids(self, input_ids: torch.Tensor) -> torch.Tensor:
return self.embed_tokens(input_ids)
def combine_hidden_states(self, aux_hidden_states: torch.Tensor) -> torch.Tensor:
"""main_x = main_norm(main_proj(concat of target aux hidden states)).
``aux_hidden_states`` is [T, hidden_size * len(target_layer_ids)].
"""
return self.main_norm(self.main_proj(aux_hidden_states))
@torch.inference_mode()
def precompute_and_store_context_kv(
self,
main_x: torch.Tensor,
context_positions: torch.Tensor,
context_slot_mappings: list[torch.Tensor | None] | None = None,
) -> None:
"""Insert the sliding-window context KV for every draft layer.
Mirrors the reference DSparkAttention: each layer derives its context KV
from the SAME projected target hidden ``main_x``, via that layer's own
``wkv`` + ``kv_norm`` + RoPE + quant, then writes it at the
layer's context slots.
``context_slot_mappings`` is a per-layer list (each entry is the context
slot mapping for that layer's kv-cache group, since the hybrid manager may
place draft layers in different groups). ``None`` (or a ``None`` entry)
runs the projection to reserve workspace but writes nothing (profiling).
"""
for i, layer in enumerate(self.layers):
slot_mapping = (
None if context_slot_mappings is None else context_slot_mappings[i]
)
attn = layer.attn
# Optimized DSV4 MLA path: wkv part of the fused wq_a|wkv projection
# (q_lora part discarded), then RoPE/quant/insert via the fused op.
qr_kv, _ = attn.fused_wqa_wkv(main_x)
kv = qr_kv[..., attn.q_lora_rank :]
kv = attn.kv_norm(kv)
if slot_mapping is None:
continue
_insert_context_kv(attn, kv, context_positions, slot_mapping)
def forward(
self,
input_ids: torch.Tensor,
positions: torch.Tensor,
inputs_embeds: torch.Tensor | None = None,
) -> torch.Tensor:
if inputs_embeds is None:
inputs_embeds = self.embed_input_ids(input_ids)
# Expand to hc_mult copies for hyper-connections ([T, H] -> [T, hc, H]).
hidden_states = inputs_embeds.unsqueeze(-2).repeat(1, self.hc_mult, 1)
residual = post_mix = res_mix = None
for layer in self.layers:
hidden_states, residual, post_mix, res_mix = layer(
hidden_states,
positions,
input_ids,
post_mix,
res_mix,
residual,
)
# On the fused-MHC path the trailing hc_post must be applied here; on the
# aiter unfused path (ROCm default) the decoder layer already applied
# hc_post in-layer and returned None mixes, so this is skipped. Mirrors
# amd/mtp.py.
last_layer = self.layers[-1]
if last_layer.use_fused_mhc:
hidden_states = last_layer.hc_post(
hidden_states, residual, post_mix, res_mix
)
# hc_head reduces the hc copies; return the PRE-norm head hidden.
hidden_states = self.hc_head_op(
hidden_states,
self.hc_head_fn,
self.hc_head_scale,
self.hc_head_base,
self.rms_norm_eps,
self.hc_eps,
)
return hidden_states
def _insert_context_kv(
attn: nn.Module,
kv: torch.Tensor,
positions: torch.Tensor,
slot_mapping: torch.Tensor,
) -> None:
"""RoPE + quant + paged-cache insert of (already kv_norm'd) context KV.
Reuses the DSV4 fused insert ops (which also process a query; we pass a dummy
query and discard it, since context tokens have no query). Mirrors
``DeepseekV4Attention._fused_qnorm_rope_kv_insert``.
"""
swa_cache = attn.swa_cache_layer.kv_cache
block_size = attn.swa_cache_layer.block_size
cos_sin_cache = attn.rotary_emb.cos_sin_cache
cache_dtype = swa_cache.dtype
n_ctx = kv.shape[0]
dummy_q = torch.zeros(
(n_ctx, attn.n_local_heads, attn.head_dim),
dtype=kv.dtype,
device=kv.device,
)
if cache_dtype == torch.uint8:
# fp8_ds_mla UE8M0 paged layout (the gfx950 aiter SWA cache path).
swa_2d = swa_cache.view(swa_cache.shape[0], -1)
torch.ops._C.fused_deepseek_v4_qnorm_rope_kv_rope_quant_insert(
dummy_q,
kv,
swa_2d,
slot_mapping,
positions,
cos_sin_cache,
attn.padded_heads,
attn.eps,
block_size,
)
elif cache_dtype == torch.bfloat16:
swa_3d = swa_cache.view(-1, block_size, attn.head_dim)
torch.ops._C.fused_deepseek_v4_qnorm_rope_kv_rope_full_cache_bf16_insert(
dummy_q,
kv,
swa_3d,
slot_mapping,
positions,
cos_sin_cache,
attn.eps,
block_size,
)
else: # per-tensor fp8 (torch.float8_e4m3fn)
# NOTE(rocm): unreachable on ROCm/aiter, where the SWA cache dtype is
# uint8 (fp8_ds_mla) or bfloat16. This branch relies on FlashInfer-only
# attributes (``_flashinfer_fp8_*``) that the aiter attention layer does
# not define; kept for parity with the nvidia path.
swa_3d = swa_cache.view(-1, block_size, attn.head_dim)
dummy_q_fp8 = torch.zeros_like(dummy_q, dtype=torch.float8_e4m3fn)
torch.ops._C.fused_deepseek_v4_qnorm_rope_kv_rope_full_cache_fp8_insert(
dummy_q,
kv,
dummy_q_fp8,
swa_3d,
slot_mapping,
positions,
cos_sin_cache,
attn._flashinfer_fp8_kv_scale,
attn._flashinfer_fp8_q_scale_inv,
attn.eps,
block_size,
)
class DSparkDeepseekV4ForCausalLM(nn.Module):
# Draft weights ship in the target checkpoint (mtp.*) without embed/head, so
# load_dspark_model always aliases the target's.
has_own_embed_tokens = False
has_own_lm_head = False
def __init__(self, *, vllm_config: VllmConfig, prefix: str = "") -> None:
super().__init__()
assert vllm_config.speculative_config is not None
self.draft_model_config = vllm_config.speculative_config.draft_model_config
self.config = self.draft_model_config.hf_config
self.model = DSparkDeepseekV4Model(
vllm_config=vllm_config, prefix=maybe_prefix(prefix, "model")
)
# Shared with the target (aliased by the speculator's load utility).
self.lm_head = ParallelLMHead(
self.config.vocab_size,
self.config.hidden_size,
prefix=maybe_prefix(prefix, "lm_head"),
)
self.logits_processor = LogitsProcessor(self.config.vocab_size)
# --- Hooks used by the speculator -------------------------------------
def embed_input_ids(self, input_ids: torch.Tensor) -> torch.Tensor:
return self.model.embed_input_ids(input_ids)
def combine_hidden_states(self, aux_hidden_states: torch.Tensor) -> torch.Tensor:
return self.model.combine_hidden_states(aux_hidden_states)
def get_draft_kv_cache_layer_names(self) -> list[str]:
# DSV4 MLA path: each draft layer's sliding-window cache is a separate
# layer, named by its prefix.
return [layer.attn.swa_cache_layer.prefix for layer in self.model.layers]
def precompute_and_store_context_kv(
self,
context_states: torch.Tensor,
context_positions: torch.Tensor,
context_slot_mappings: list[torch.Tensor | None] | None = None,
) -> None:
self.model.precompute_and_store_context_kv(
context_states, context_positions, context_slot_mappings
)
def forward(
self,
input_ids: torch.Tensor,
positions: torch.Tensor,
inputs_embeds: torch.Tensor | None = None,
) -> torch.Tensor:
# Returns the pre-norm hc_head hidden ([T, hidden_size]).
return self.model(input_ids, positions, inputs_embeds)
def compute_logits(self, hidden_states: torch.Tensor) -> torch.Tensor:
"""Base logits U_k = lm_head(norm(head_hidden))."""
return self.logits_processor(self.lm_head, self.model.norm(hidden_states))
def compute_draft_logits(self, hidden_states: torch.Tensor) -> torch.Tensor:
# Full-vocab draft: base logits, no d2t scatter.
return self.compute_logits(hidden_states)
def map_draft_to_target(self, draft_ids: torch.Tensor) -> torch.Tensor:
return draft_ids # full-vocab: draft ids are target ids
def markov_embed(self, token_ids: torch.Tensor) -> torch.Tensor:
return self.model.markov_head.embed(token_ids)
def markov_bias(self, markov_embed: torch.Tensor) -> torch.Tensor:
return self.model.markov_head.bias(markov_embed, self.logits_processor)
# --- Weight loading ----------------------------------------------------
def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]:
"""Load the ``mtp.{0,1,2}.*`` draft weights from the target checkpoint.
Non-mtp weights (embed/head/main layers) belong to the target model and
are skipped here. ``embed_tokens``/``lm_head`` are aliased from the target.
"""
# AMD DeepseekV4MoE has no mega-MoE path; always use the standard
# per-expert fused-MoE mapping (mirrors amd/mtp.py).
expert_mapping = fused_moe_make_expert_params_mapping(
self,
ckpt_gate_proj_name="w1",
ckpt_down_proj_name="w2",
ckpt_up_proj_name="w3",
num_experts=self.config.n_routed_experts,
)
expert_scale_suffix = (
".weight_scale"
if getattr(self.config, "expert_dtype", "fp4") == "fp4"
else ".weight_scale_inv"
)
# (param_name, ckpt_shard_name, shard_id) for non-expert stacked params.
stacked_params_mapping = [
("gate_up_proj", "w1", 0),
("gate_up_proj", "w3", 1),
("attn.fused_wqa_wkv", "attn.wq_a", 0),
("attn.fused_wqa_wkv", "attn.wkv", 1),
]
params_dict = dict(self.named_parameters())
loaded_params: set[str] = set()
tp_size = get_tensor_model_parallel_world_size()
tp_rank = get_tensor_model_parallel_rank()
n_local_head = self.config.num_attention_heads // tp_size
head_start = n_local_head * tp_rank
head_end = n_local_head * (tp_rank + 1)
for name, loaded_weight in weights:
mapped = self._remap_dspark_name(name)
if mapped is None:
continue
name = mapped
# ``.scale`` -> per-method scale suffix.
if name.endswith(".scale"):
suffix = (
expert_scale_suffix
if _EXPERT_SCALE_RE.search(name)
else ".weight_scale_inv"
)
name = name.removesuffix(".scale") + suffix
# E8M0 expert scales: keep raw exponent bytes.
if ".experts." in name:
if (
"weight_scale" in name
and loaded_weight.dtype == torch.float8_e8m0fnu
):
loaded_weight = loaded_weight.view(torch.uint8)
for param_name, weight_name, expert_id, shard_id in expert_mapping:
if weight_name not in name:
continue
name_mapped = name.replace(weight_name, param_name)
param = params_dict[name_mapped]
success = param.weight_loader(
param,
loaded_weight,
name_mapped,
shard_id=shard_id,
expert_id=expert_id,
return_success=True,
)
if success:
loaded_params.add(name_mapped)
break
continue
# Stacked rules only apply to decoder-layer weights. Head-stack params
# (main_proj/norm/hc_head/markov_head) load directly — otherwise e.g.
# "markov_w1" would collide with the "w1" shard rule.
is_layer_param = name.startswith("model.layers.")
for param_name, weight_name, stacked_shard_id in stacked_params_mapping:
if not is_layer_param or weight_name not in name:
continue
name = name.replace(weight_name, param_name)
param = params_dict[name]
param.weight_loader(param, loaded_weight, stacked_shard_id)
loaded_params.add(name)
break
else:
if "attn_sink" in name:
narrow = loaded_weight[head_start:head_end]
params_dict[name][: narrow.shape[0]].copy_(narrow)
loaded_params.add(name)
continue
if ".shared_experts.w2" in name:
name = name.replace(
".shared_experts.w2", ".shared_experts.down_proj"
)
if name.endswith(".ffn.gate.bias"):
name = name.replace(
".ffn.gate.bias", ".ffn.gate.e_score_correction_bias"
)
param = params_dict[name]
weight_loader = getattr(param, "weight_loader", default_weight_loader)
weight_loader(param, loaded_weight)
loaded_params.add(name)
logger.info_once("DSpark draft model loaded: %d params", len(loaded_params))
return loaded_params
def _remap_dspark_name(self, name: str) -> str | None:
"""Map a checkpoint ``mtp.{i}.*`` name to this model's parameter path.
Returns None for non-mtp weights (owned by the target model).
"""
m = re.match(r"mtp\.(\d+)\.(.*)", name)
if m is None:
return None
stage = int(m.group(1))
rest = m.group(2)
# The confidence head is not wired into inference yet; drop its weights.
if rest.startswith("confidence_head."):
return None
# Head-stack params live at model level (mtp.last), context combiner at
# model level (mtp.0); everything else is a per-layer decoder block.
head_prefixes = (
"norm.",
"hc_head_fn",
"hc_head_base",
"hc_head_scale",
"markov_head.",
)
if rest.startswith(("main_proj.", "main_norm.")) or rest.startswith(
head_prefixes
):
return f"model.{rest}"
return f"model.layers.{stage}.{rest}"
+45 -5
View File
@@ -40,7 +40,11 @@ from vllm.model_executor.layers.vocab_parallel_embedding import (
VocabParallelEmbedding,
)
from vllm.model_executor.model_loader.weight_utils import default_weight_loader
from vllm.model_executor.models.interfaces import SupportsPP
from vllm.model_executor.models.interfaces import (
EagleModelMixin,
SupportsEagle3,
SupportsPP,
)
from vllm.model_executor.models.utils import (
AutoWeightsLoader,
PPMissingLayer,
@@ -437,7 +441,7 @@ class DeepseekV4DecoderLayer(nn.Module):
)
class DeepseekV4Model(nn.Module):
class DeepseekV4Model(nn.Module, EagleModelMixin):
def __init__(self, *, vllm_config: VllmConfig, prefix: str = ""):
super().__init__()
@@ -573,7 +577,19 @@ class DeepseekV4Model(nn.Module):
hidden_states = intermediate_tensors["hidden_states"]
residual, post_mix, res_mix = None, None, None
for layer in islice(self.layers, self.start_layer, self.end_layer):
# EAGLE3 / DSpark / DFlash aux hidden states: reconstructed (post-mhc)
# hidden state at the configured target layers, averaged over the
# hc_mult streams to [T, hidden_size]. Empty unless a draft model set
# aux_hidden_state_layers.
aux_hidden_states: list[torch.Tensor] = []
# On the fused path the final layer's hc_post output is reused below
# (avoids computing hc_post twice when the last layer is also an aux
# layer).
final_aux_recon: torch.Tensor | None = None
for idx, layer in enumerate(
islice(self.layers, self.start_layer, self.end_layer),
start=self.start_layer,
):
hidden_states, residual, post_mix, res_mix = layer(
hidden_states,
positions,
@@ -582,8 +598,30 @@ class DeepseekV4Model(nn.Module):
res_mix,
residual,
)
if (idx + 1) in self.aux_hidden_state_layers:
# On the unfused (aiter) path the layer already applied hc_post,
# so hidden_states is the reconstructed stream; on the fused
# path reconstruct it via hc_post before averaging.
if layer.use_fused_mhc:
aux_recon = layer.hc_post(
hidden_states, residual, post_mix, res_mix
)
final_aux_recon = aux_recon
else:
aux_recon = hidden_states
aux_hidden_states.append(aux_recon.mean(dim=1))
if layer is not None and layer.use_fused_mhc:
hidden_states = layer.hc_post(hidden_states, residual, post_mix, res_mix)
# Reuse the last layer's hc_post output if it was already computed
# for the aux hidden state above; otherwise compute it now.
if (
final_aux_recon is not None
and self.end_layer in self.aux_hidden_state_layers
):
hidden_states = final_aux_recon
else:
hidden_states = layer.hc_post(
hidden_states, residual, post_mix, res_mix
)
if not get_pp_group().is_last_rank:
return IntermediateTensors({"hidden_states": hidden_states})
@@ -601,6 +639,8 @@ class DeepseekV4Model(nn.Module):
self.hc_eps,
)
hidden_states = self.norm(hidden_states)
if len(aux_hidden_states) > 0:
return hidden_states, aux_hidden_states
return hidden_states
def load_weights(self, weights: Iterable[tuple[str, torch.Tensor]]) -> set[str]:
@@ -751,7 +791,7 @@ def _make_deepseek_v4_weights_mapper(expert_dtype: str) -> WeightsMapper:
)
class DeepseekV4ForCausalLM(nn.Module, SupportsPP):
class DeepseekV4ForCausalLM(nn.Module, SupportsPP, SupportsEagle3):
model_cls = DeepseekV4Model
# Default mapper assumes the original FP4-expert checkpoint layout.

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