From 8cf77351fd970364ba3629a3b275120e01400987 Mon Sep 17 00:00:00 2001 From: tzhouam Date: Tue, 3 Nov 2026 11:29:12 +0800 Subject: [PATCH 1/8] [Core] Align vLLM-Omni with vLLM 0.33.0 Adapt the final upstream APIs and runtime dependencies, retain compatibility regressions, and update installation documentation and Dockerfiles. Signed-off-by: tzhouam --- .buildkite/amd/scripts/aiter_mha_cache.py | 5 + .buildkite/amd/scripts/build-ci-image.sh | 2 + .buildkite/amd/test-amd-ready.yml | 4 +- .buildkite/cuda/scripts/build-ci-image.sh | 7 +- .buildkite/cuda/scripts/check_vllm_runtime.py | 27 +++- .buildkite/cuda/test-ready.yml | 3 +- .buildkite/intel/pipeline-intel.yml | 2 +- docker/Dockerfile.ci | 59 ++++++- docker/Dockerfile.cuda | 47 +++++- docker/Dockerfile.npu | 4 + docker/Dockerfile.npu.ci | 4 + docker/Dockerfile.rocm | 29 ++-- docker/Dockerfile.xpu | 12 +- docs/configuration/README.md | 2 +- docs/configuration/environment_variables.md | 4 +- docs/contributing/README.md | 2 +- docs/contributing/ci/test_execution_guide.md | 15 +- docs/contributing/profiling.md | 2 +- .../feature/attention_backend_selection.md | 8 + docs/design/fullduplex.md | 2 +- docs/design/minicpm_o45_mrv2_performance.md | 2 +- docs/design/module/ar_runtime.md | 35 +++++ docs/design/module/cache_management.md | 14 ++ docs/design/module/error_contracts.md | 4 + .../module/input_output_modality_contracts.md | 2 +- docs/design/module/profiling.md | 7 + docs/getting_started/installation/README.md | 33 +++- .../installation/gpu/cuda.inc.md | 105 ++++++++++--- .../installation/gpu/musa.inc.md | 16 +- .../installation/gpu/rocm.inc.md | 51 +++--- .../installation/gpu/xpu.inc.md | 26 +++- .../installation/npu/npu.inc.md | 18 ++- docs/getting_started/quickstart.md | 17 +- docs/serving/realtime_duplex_api.md | 9 +- .../attention_backends/dense_backends.md | 14 ++ .../attention_backends/huggingface_hub.md | 5 + .../offline_inference/text_to_image.md | 18 ++- .../offline_inference/text_to_speech.md | 17 +- .../examples/online_serving/image_to_image.md | 31 ++-- .../examples/online_serving/qwen2_5_omni.md | 27 ++-- .../examples/online_serving/qwen3_omni.md | 5 +- .../examples/online_serving/text_to_image.md | 9 +- .../text_to_speech/README.md | 17 +- .../streaming_video_generation/README.md | 6 + .../online_serving/text_to_speech/README.md | 5 + recipes/MOSS-Image/MOSS-Image-Base.md | 2 +- recipes/MiniMaxAI/MiniMax-H3.md | 14 +- recipes/README.md | 8 + recipes/inclusionAI/Ming-3.0-audio-flash.md | 2 +- recipes/models/MiniMaxAI/MiniMax-H3.yaml | 9 +- recipes/models/Qwen/Qwen-Image-2.1.yaml | 14 +- recipes/models/Qwen/Qwen-Image.yaml | 6 +- recipes/models/fishaudio/drama-3.yaml | 2 +- .../meituan-longcat/LongCat-Image-Edit.yaml | 24 +-- recipes/models/nvidia/Cosmos3-Nano.yaml | 5 +- .../nvidia/Cosmos3-Super-Image2Video.yaml | 5 +- .../nvidia/Cosmos3-Super-Text2Image.yaml | 5 +- recipes/models/nvidia/Cosmos3-Super.yaml | 5 +- .../stabilityai/stable-audio-open-1.0.yaml | 4 +- .../stable-diffusion-3.5-medium.yaml | 19 +-- requirements/common.txt | 2 +- requirements/cuda-bootstrap.txt | 12 ++ requirements/cudnn-runtime.txt | 3 + tests/buildkite/test_aiter_mha_cache.py | 13 ++ tests/buildkite/test_amd_pipeline.py | 4 +- tests/buildkite/test_check_vllm_runtime.py | 44 ++++-- tests/buildkite/test_rocm_dockerfile.py | 26 ++-- tests/config/test_config_import_cycle.py | 4 +- .../test_generation_scheduler_restore.py | 6 +- .../sched/test_omni_ar_materialized_cache.py | 145 ++++++++++++++++++ ...t_omni_ar_scheduler_payloadless_outputs.py | 8 +- .../sched/test_omni_ar_scheduler_streaming.py | 36 +++++ .../test_omni_sched_deferred_free_fence.py | 26 ++++ ...test_omni_scheduler_async_chunk_prewarm.py | 15 +- ...mni_scheduler_input_coordinator_cleanup.py | 2 + tests/diffusion/attention/test_kernels_hub.py | 24 +++ .../test_wan_vae_fastpath_verdicts.py | 10 +- .../layers/test_adalayernorm_cuda_path.py | 20 +++ .../model_loader/test_diffusers_loader.py | 2 + .../minimax_h3/test_minimax_h3_regenerate.py | 3 + .../minimax_h3/test_minimax_h3_vae_ops.py | 6 +- tests/diffusion/models/pi0/test_pi0_units.py | 3 +- .../test_qwen_image_scheduler_device.py | 60 ++++++++ .../test_transformer_sana_video_sp.py | 18 ++- .../test_distributed_layerwise_backend.py | 6 +- .../quantization/test_mxfp4_config.py | 4 +- .../online_serving/test_duplex_client_live.py | 4 +- tests/engine/test_duplex_import_boundary.py | 4 +- tests/engine/test_orchestrator.py | 88 ++++++++++- tests/engine/test_output_processor.py | 39 ++++- .../engine/test_output_processor_mrv2_text.py | 2 + tests/engine/test_tracing.py | 51 +++++- .../entrypoints/duplex/test_duplex_serving.py | 35 ++++- .../openai_api/test_api_server_guards.py | 1 + .../openai_api/test_duplex_api_server.py | 5 + .../openai_api/test_serving_audio_generate.py | 17 +- .../openai_api/test_serving_speech.py | 3 +- tests/helpers/cudnn.py | 41 +++++ tests/helpers/omni_scheduler.py | 24 +++ tests/metrics/test_stat_logger.py | 5 +- .../higgs_audio_v3/test_codec_snake_kernel.py | 10 +- .../minicpmo_4_5/test_cuda_graph_wrapper.py | 27 +++- .../models/moss_tts/test_codec_cudagraph.py | 13 +- .../moss_tts/test_shared_reference_encoder.py | 14 +- .../test_qwen3_omni_deepstack_device.py | 3 + .../models/qwen3_omni/test_talker_sampler.py | 81 +++++++++- .../qwen3_tts/test_cuda_graph_decoder.py | 11 +- tests/platforms/npu/test_aux_output.py | 12 +- tests/platforms/test_rocm_ir_op_priority.py | 63 ++++++++ tests/profile/test_omni_torch_profiler.py | 26 +++- .../test_profile_iteration_controls.py | 57 ++++++- tests/worker/test_img2img_mm_prefix.py | 125 +++++++++++++++ tests/worker/test_moe_prefill_warmup.py | 4 +- tests/worker/test_omni_v2_profile_dispatch.py | 2 + tests/worker/test_process_gpu_memory.py | 21 +-- tests/worker/test_process_gpu_memory_cuda.py | 25 +++ tests/worker_v2/test_omni_ar_model_runner.py | 2 + .../test_omni_generation_model_runner.py | 1 + tests/worker_v2/test_omni_gpu_model_runner.py | 53 +++++++ .../worker_v2/test_shared_stage_contracts.py | 2 + tools/build_gradio_pillow13.py | 96 ++++++++++++ tools/gradio_pillow13/NOTICE.md | 16 ++ tools/gradio_pillow13/backport.patch | 21 +++ tools/gradio_pillow13/build-constraints.txt | 23 +++ tools/gradio_pillow13/manifest.json | 26 ++++ tools/prepare_ci_telemetry.py | 36 +++++ vllm_omni/config/omni_config.py | 2 +- vllm_omni/core/sched/omni_ar_scheduler.py | 59 +++++-- .../core/sched/omni_generation_scheduler.py | 9 +- vllm_omni/core/sched/omni_scheduler_mixin.py | 19 ++- .../attention/backends/flash_attn_hub.py | 4 +- vllm_omni/diffusion/diffusion_engine.py | 11 +- .../diffusion/diffusion_kv/initialization.py | 16 +- .../autoencoders/wan_vae_fastpath/forwards.py | 54 ++++--- vllm_omni/diffusion/layers/adalayernorm.py | 7 +- .../diffusion/models/magi2/modeling_magi2.py | 22 ++- .../models/minimax_h3/ops/vae/qk_norm_rope.py | 63 +++++++- .../models/qwen_image/pipeline_qwen_image.py | 1 + .../qwen_image/pipeline_qwen_image_edit.py | 3 +- .../pipeline_qwen_image_edit_plus.py | 3 +- .../qwen_image/qwen_image_transformer.py | 125 ++++++++++++--- .../models/z_image/pipeline_z_image.py | 6 +- vllm_omni/diffusion/vllm_config.py | 3 + .../diffusion/worker/diffusion_worker.py | 3 + vllm_omni/engine/orchestrator.py | 111 ++++++++++---- vllm_omni/engine/stage_pool.py | 1 + vllm_omni/entrypoints/async_omni.py | 6 +- vllm_omni/entrypoints/cli/benchmark/main.py | 5 +- vllm_omni/entrypoints/cli/main.py | 2 +- vllm_omni/entrypoints/cli/serve.py | 4 +- vllm_omni/entrypoints/duplex/serving.py | 4 + vllm_omni/entrypoints/duplex/warmup.py | 2 +- vllm_omni/entrypoints/omni_base.py | 32 +++- vllm_omni/entrypoints/openai/api_server.py | 72 +++++++-- .../entrypoints/openai/realtime/connection.py | 10 +- vllm_omni/entrypoints/openai/serving_chat.py | 4 + .../entrypoints/serve/profile/protocol.py | 12 +- .../models/cosyvoice3/v1_sampling.py | 6 +- .../models/fish_omni/model_state.py | 2 +- .../models/moss_tts/local_kernels.py | 7 +- .../models/moss_tts/local_model_state.py | 2 +- .../models/moss_tts/packed_codec.py | 2 +- .../models/qwen3_omni/talker_sampler.py | 18 ++- vllm_omni/outputs/output_processor.py | 34 +++- vllm_omni/patch.py | 4 +- vllm_omni/platforms/cuda/platform.py | 16 +- vllm_omni/platforms/npu/profiler.py | 6 +- vllm_omni/platforms/npu/worker/base.py | 5 +- vllm_omni/platforms/rocm/platform.py | 10 +- vllm_omni/profiler/omni_torch_profiler.py | 32 ++-- vllm_omni/utils/mm_prefix.py | 64 ++++++++ vllm_omni/worker/base.py | 5 + vllm_omni/worker/gpu_ar_model_runner.py | 19 ++- .../worker/gpu_generation_model_runner.py | 12 +- vllm_omni/worker/sampler_warmup.py | 2 +- vllm_omni/worker_v2/native_output_worker.py | 26 +++- vllm_omni/worker_v2/omni_ar_model_runner.py | 5 +- .../worker_v2/omni_generation_model_runner.py | 14 +- vllm_omni/worker_v2/omni_model_runner.py | 59 ++++++- vllm_omni/worker_v2/staged_writes.py | 18 ++- 180 files changed, 2994 insertions(+), 497 deletions(-) create mode 100644 requirements/cuda-bootstrap.txt create mode 100644 requirements/cudnn-runtime.txt create mode 100644 tests/core/sched/test_omni_ar_materialized_cache.py create mode 100644 tests/diffusion/models/qwen_image/test_qwen_image_scheduler_device.py create mode 100644 tests/helpers/cudnn.py create mode 100644 tests/platforms/test_rocm_ir_op_priority.py create mode 100644 tests/worker/test_img2img_mm_prefix.py create mode 100644 tests/worker/test_process_gpu_memory_cuda.py create mode 100644 tools/build_gradio_pillow13.py create mode 100644 tools/gradio_pillow13/NOTICE.md create mode 100644 tools/gradio_pillow13/backport.patch create mode 100644 tools/gradio_pillow13/build-constraints.txt create mode 100644 tools/gradio_pillow13/manifest.json create mode 100644 tools/prepare_ci_telemetry.py create mode 100644 vllm_omni/utils/mm_prefix.py diff --git a/.buildkite/amd/scripts/aiter_mha_cache.py b/.buildkite/amd/scripts/aiter_mha_cache.py index acd268ee480..23ff7327bdd 100644 --- a/.buildkite/amd/scripts/aiter_mha_cache.py +++ b/.buildkite/amd/scripts/aiter_mha_cache.py @@ -105,6 +105,7 @@ "pyproject.toml", "setup.py", "tools/install_torchcodec_rocm.sh", + "tools/build_gradio_pillow13.py", ) _READ_SIZE = 1024 * 1024 _MAX_HTTP_BYTES = 2 * 1024 * 1024 @@ -389,7 +390,11 @@ def dependency_identity(root: Path) -> str: if not requirement_dir.is_dir(): return "" requirement_files = sorted(path for path in requirement_dir.glob("*.txt") if path.is_file()) + backport_files = sorted((root / "tools/gradio_pillow13").glob("*")) + if not backport_files: + return "" files = [root / relative for relative in _IMAGE_DEPENDENCY_PATHS] + files.extend(backport_files) files.extend(requirement_files) if any(not path.is_file() for path in files): return "" diff --git a/.buildkite/amd/scripts/build-ci-image.sh b/.buildkite/amd/scripts/build-ci-image.sh index 3954cd567f9..96025944a17 100755 --- a/.buildkite/amd/scripts/build-ci-image.sh +++ b/.buildkite/amd/scripts/build-ci-image.sh @@ -25,6 +25,8 @@ dependency_files=( pyproject.toml setup.py tools/install_torchcodec_rocm.sh + tools/build_gradio_pillow13.py + tools/gradio_pillow13/* ) while IFS= read -r dependency_file; do dependency_files+=("$dependency_file") diff --git a/.buildkite/amd/test-amd-ready.yml b/.buildkite/amd/test-amd-ready.yml index 6a7da834bc1..20e4cd2ae91 100644 --- a/.buildkite/amd/test-amd-ready.yml +++ b/.buildkite/amd/test-amd-ready.yml @@ -1223,7 +1223,7 @@ steps: marker="core_model and omni and rocm and MI325 and cards_1" pytest --collect-only -q "$${nodes[@]}" \ --shard-id=0 --num-shards=1 \ - -m "$$marker" --run-level core_model \ + -m "$$marker" --run-level advanced_model \ > "$$ROCM_MINICPMO_DUPLEX_ARTIFACT_DIR/collection.txt" count=$$(awk '/::/ { n++ } END { print n + 0 }' \ "$$ROCM_MINICPMO_DUPLEX_ARTIFACT_DIR/collection.txt") @@ -1247,7 +1247,7 @@ steps: timeout --signal=TERM --kill-after=1m 35m \ pytest -s -v "$${nodes[@]}" \ --shard-id=0 --num-shards=1 \ - -m "core_model and omni and rocm and MI325 and cards_1" --run-level core_model \ + -m "core_model and omni and rocm and MI325 and cards_1" --run-level advanced_model \ --junitxml="$$ROCM_MINICPMO_DUPLEX_ARTIFACT_DIR/pytest.xml" || status=$$? if [[ "$$status" -eq 0 ]]; then if python -c 'import sys, xml.etree.ElementTree as ET; root = ET.parse(sys.argv[1]).getroot(); suites = [root] if root.tag == "testsuite" else list(root.findall("testsuite")); tests = sum(int(suite.attrib.get("tests", "0")) for suite in suites); skipped = sum(int(suite.attrib.get("skipped", "0")) for suite in suites); ok = tests == 2 and skipped == 0; print(f"Expected exactly two executed MiniCPM-o duplex tests; reported tests={tests}, skipped={skipped}.", file=sys.stderr) if not ok else None; raise SystemExit(0 if ok else 5)' "$$ROCM_MINICPMO_DUPLEX_ARTIFACT_DIR/pytest.xml" diff --git a/.buildkite/cuda/scripts/build-ci-image.sh b/.buildkite/cuda/scripts/build-ci-image.sh index 9da69bc8be3..0e0282f8ce0 100755 --- a/.buildkite/cuda/scripts/build-ci-image.sh +++ b/.buildkite/cuda/scripts/build-ci-image.sh @@ -1,4 +1,6 @@ #!/bin/bash +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project set -euo pipefail # Build and push the CI Docker image with registry-based layer caching. @@ -28,8 +30,9 @@ aws ecr-public get-login-password --region "$REGION" \ # This is because dependencies don't change often, so most PRs will have the same hash. # # The Dockerfile is also included to avoid churn in the cached image based on the vLLM version. -DEP_FILES="pyproject.toml setup.py requirements/*.txt docker/Dockerfile.ci" -CACHE_KEY=$(cat ${DEP_FILES} | sha256sum | cut -c1-16) +DEP_FILES=(pyproject.toml setup.py requirements/*.txt docker/Dockerfile.ci + tools/build_gradio_pillow13.py tools/prepare_ci_telemetry.py tools/gradio_pillow13/*) +CACHE_KEY=$(cat "${DEP_FILES[@]}" | sha256sum | cut -c1-16) CACHE_TAG="deps-cache-${CACHE_KEY}" echo "Cache key: ${CACHE_TAG}" diff --git a/.buildkite/cuda/scripts/check_vllm_runtime.py b/.buildkite/cuda/scripts/check_vllm_runtime.py index f79f3676d08..616fde12af9 100644 --- a/.buildkite/cuda/scripts/check_vllm_runtime.py +++ b/.buildkite/cuda/scripts/check_vllm_runtime.py @@ -10,15 +10,27 @@ def main() -> None: - dockerfile = Path(__file__).resolve().parents[3] / "docker" / "Dockerfile.ci" + root = Path(__file__).resolve().parents[3] + dockerfile = root / "docker" / "Dockerfile.ci" + text = dockerfile.read_text() match = re.search( r"^ARG VLLM_PRECOMPILED_WHEEL_COMMIT=([0-9a-f]{40})$", - dockerfile.read_text(), + text, re.MULTILINE, ) if match is None: raise RuntimeError(f"Missing vLLM commit pin in {dockerfile}") expected_commit = match.group(1) + version_match = re.search(r"^ARG VLLM_VERSION=(\S+)$", text, re.MULTILINE) + if version_match is None: + raise RuntimeError(f"Missing vLLM version pin in {dockerfile}") + expected_version = version_match.group(1) + cudnn_requirements = root / "requirements/cudnn-runtime.txt" + cudnn_match = re.match(r"nvidia-cudnn-cu13==(\d+)\.(\d+)\.(\d+)\.\d+(?:\s|$)", cudnn_requirements.read_text()) + if cudnn_match is None: + raise RuntimeError(f"Missing cuDNN runtime pin in {cudnn_requirements}") + major, minor, patch = map(int, cudnn_match.groups()) + expected_cudnn = major * 10000 + minor * 100 + patch import torch import vllm @@ -31,13 +43,20 @@ def main() -> None: assert len(installed_commit) >= 8 and expected_commit.startswith(installed_commit), ( f"Expected vLLM {expected_commit}, got {__commit_id__} ({vllm.__version__})" ) + # A final wheel can share its source SHA with an RC while containing + # different native binaries. Both identities must match this CI cycle. + assert vllm.__version__ == expected_version, f"Expected vLLM version {expected_version}, got {vllm.__version__}" assert torch.cuda.is_available(), "vLLM runtime check requires a GPU worker" - for module in ("flashinfer", "vllm._C_stable_libtorch"): + actual_cudnn = torch.backends.cudnn.version() + assert actual_cudnn == expected_cudnn, f"Expected cuDNN {expected_cudnn}, got {actual_cudnn}" + # Python dependency metadata cannot detect a retained connector wheel's + # Torch ABI mismatch. Import it on the GPU worker before running models. + for module in ("flashinfer", "vllm._C_stable_libtorch", "lmcache.cuda_ops"): importlib.import_module(module) print( f"Verified vLLM CUDA runtime: vllm={vllm.__version__} " f"commit={expected_commit} torch={torch.__version__} " - f"cuda={torch.version.cuda} device={torch.cuda.get_device_name()}" + f"cuda={torch.version.cuda} cudnn={actual_cudnn} device={torch.cuda.get_device_name()}" ) diff --git a/.buildkite/cuda/test-ready.yml b/.buildkite/cuda/test-ready.yml index c7893416fba..adc3f140920 100644 --- a/.buildkite/cuda/test-ready.yml +++ b/.buildkite/cuda/test-ready.yml @@ -204,12 +204,13 @@ steps: " mirror_hardwares: h100_1 + # The live client waits for a learned turn terminator; use real weights. - label: "Omni · MiniCPM-o 4.5 Duplex Test" source_file_dependencies: omni_minicpmo_4_5_duplex_function commands: - | timeout 35m bash -c " - pytest -s -v tests/e2e/online_serving/test_minicpmo_4_5_duplex.py tests/e2e/online_serving/test_duplex_client_live.py -m 'core_model and cuda' --run-level 'core_model' + pytest -s -v tests/e2e/online_serving/test_minicpmo_4_5_duplex.py tests/e2e/online_serving/test_duplex_client_live.py -m 'core_model and cuda' --run-level 'advanced_model' " mirror_hardwares: h100_1 diff --git a/.buildkite/intel/pipeline-intel.yml b/.buildkite/intel/pipeline-intel.yml index 46485ededea..4e6983ec737 100644 --- a/.buildkite/intel/pipeline-intel.yml +++ b/.buildkite/intel/pipeline-intel.yml @@ -10,7 +10,7 @@ steps: DOCKER_BUILDKIT: "1" # Buildkite will automatically replace this with the actual commit hash VLLM_IMAGE_TAG: "${BUILDKITE_COMMIT}" - VLLM_VERSION: "v0.32.0" + VLLM_VERSION: "v0.33.0" priority: 100 timeout_in_minutes: 60 soft_fail: false diff --git a/docker/Dockerfile.ci b/docker/Dockerfile.ci index 20278bcca9d..4aed0652563 100644 --- a/docker/Dockerfile.ci +++ b/docker/Dockerfile.ci @@ -1,10 +1,10 @@ -# Align the CUDA base and commit-index wheel with the final v0.32.0 release. -# Pin the published multi-platform image digest so CI uses the same release -# stack even if the upstream tag is updated. -ARG VLLM_PRECOMPILED_WHEEL_COMMIT=5df63a7203cdb9328d824a27b682e77c7b9928be -ARG VLLM_VERSION=0.32.0 +# Install the final v0.33.0 package from its frozen commit index. +# Use the matching published v0.33.0 image and verify the selected runtime +# from its exact frozen wheel below. +ARG VLLM_PRECOMPILED_WHEEL_COMMIT=4ba787ec9959e566d3b587a796dff07673324da5 +ARG VLLM_VERSION=0.33.0 ARG VLLM_BASE_IMAGE=vllm/vllm-openai -ARG VLLM_BASE_TAG=v${VLLM_VERSION}@sha256:befccf0e7e4089aa23926bb6f887a85c0d46c2a99e6de5c5ffdf5cb7957c5b34 +ARG VLLM_BASE_TAG=v0.33.0@sha256:d825baf9ce4b1b6bfc5ecc7b79880c7e6748c65a72d697ce3d47607aafa84c6e # We use registry-based layer caching for the dependencies, # where the tag is determined based on the hash of the contents @@ -12,6 +12,11 @@ ARG VLLM_BASE_TAG=v${VLLM_VERSION}@sha256:befccf0e7e4089aa23926bb6f887a85c0d46c2 # everything except the actual project dir, since that should not # change much across branches. FROM ${VLLM_BASE_IMAGE}:${VLLM_BASE_TAG} AS deps +ARG VLLM_BASE_IMAGE +ARG VLLM_BASE_TAG +# The base overrides Torch's NCCL dependency for its optional EFA build. +# Resolve the official wheel's declared dependencies for the Omni runtime. +ENV UV_OVERRIDE="" ARG VLLM_PRECOMPILED_WHEEL_COMMIT ARG VLLM_VERSION # Keep APP_DIR outside /workspace so k8s Buildkite workspace mounts cannot @@ -38,19 +43,57 @@ RUN apt-get update && \ # for CI COPY pyproject.toml setup.py ./ COPY requirements/ requirements/ +COPY tools/build_gradio_pillow13.py tools/build_gradio_pillow13.py +COPY tools/prepare_ci_telemetry.py tools/prepare_ci_telemetry.py +COPY tools/gradio_pillow13/ tools/gradio_pillow13/ # Scope the override to Omni; dependency source builds need their own versions. ENV VLLM_OMNI_VERSION_OVERRIDE=0.0.0 # Editable install: deps + project metadata are cached in this layer, while # the package itself is linked to APP_DIR (populated later by COPY --link). # Do NOT use a non-editable install here — that copies an empty stub into # site-packages and breaks ``import vllm_omni.entrypoints``. -RUN uv pip install --system \ +RUN python3 tools/build_gradio_pillow13.py --output-dir /tmp/gradio-pillow13 +# This exact base contains an unused exporter requiring a newer SDK than +# LMCache supports. Validate its known metadata before removing that orphan. +RUN python3 tools/prepare_ci_telemetry.py "${VLLM_BASE_IMAGE}:${VLLM_BASE_TAG}" && \ + uv pip uninstall --system opentelemetry-exporter-otlp-common +RUN uv pip install --system -r requirements/cuda-bootstrap.txt \ + "nvidia-nccl-cu13==2.30.7" \ + /tmp/gradio-pillow13/gradio-6.31.0+omni.pillow13-py3-none-any.whl \ "vllm @ https://wheels.vllm.ai/${VLLM_PRECOMPILED_WHEEL_COMMIT}/vllm-${VLLM_VERSION}-cp38-abi3-manylinux_2_28_$(uname -m).whl" \ -e ".[dev]" +# The final base's LMCache wheel was compiled against Torch 2.13. Its CUDA +# extension references a c10 symbol removed in Torch 2.14. Follow LMCache's +# source workflow so the unchanged release is compiled against this runtime. +RUN uv pip install --system -r requirements/cuda-bootstrap.txt \ + ninja "packaging>=24.2" "setuptools>=77.0.3,<81.0.0" \ + "setuptools_scm>=8" "grpcio==1.78.0" "grpcio-tools==1.78.0" wheel +# CUDA builder images can lack the math-library headers supplied by the +# installed NVIDIA packages. Keep all published LMCache CUDA architectures. +# Disable the installer cache so an older Torch-specific source wheel cannot +# satisfy this rebuild merely because its Python package version is unchanged. +RUN PYTHON_SITE=$(python3 -c 'import sysconfig; print(sysconfig.get_path("purelib"))') && \ + test -f "${PYTHON_SITE}/nvidia/cu13/include/cusparse.h" && \ + test -f "${PYTHON_SITE}/nvidia/cusparselt/include/cusparseLt.h" && \ + CUDA_VISIBLE_DEVICES="" CUDA_HOME=/usr/local/cuda BUILD_WITH_CUDA=1 MAX_JOBS=3 NVCC_THREADS=1 \ + TORCH_CUDA_ARCH_LIST="7.5;8.0;8.6;8.9;9.0;10.0;11.0;12.0" \ + CPATH="${PYTHON_SITE}/nvidia/cu13/include:${PYTHON_SITE}/nvidia/cusparselt/include${CPATH:+:${CPATH}}" \ + uv pip install --system --no-cache --no-binary lmcache \ + --no-build-isolation-package lmcache --reinstall-package lmcache \ + -r requirements/cuda-bootstrap.txt "lmcache==0.5.6" +RUN uv pip check --system && rm -rf /tmp/gradio-pillow13 +# cuDNN 9.24 causes the original Qwen-Image 2.1 pixel gate to fail under +# Torch 2.14. Keep Torch's declared Python dependencies and load the verified +# NVIDIA 9.25.1 wheel in a private target outside Python's import path. +RUN uv pip install --target /opt/vllm-omni-cudnn --no-deps --only-binary=:all: \ + --require-hashes -r requirements/cudnn-runtime.txt +ENV LD_PRELOAD="/opt/vllm-omni-cudnn/nvidia/cudnn/lib/libcudnn.so.9 ${LD_PRELOAD:-}" +ENV LD_LIBRARY_PATH="/opt/vllm-omni-cudnn/nvidia/cudnn/lib:${LD_LIBRARY_PATH:-/usr/local/nvidia/lib:/usr/local/nvidia/lib64}" + # Fail the image build if dependency resolution replaced the selected vLLM. # The CPU builder has no NVIDIA driver; CUDA imports are checked at the start # of each GPU test job by .buildkite/cuda/scripts/check_vllm_runtime.py. -RUN python3 -c "from vllm._version import __commit_id__; commit = (__commit_id__ or '').removeprefix('g'); assert len(commit) >= 8 and '${VLLM_PRECOMPILED_WHEEL_COMMIT}'.startswith(commit), __commit_id__" +RUN python3 -c "from vllm._version import __commit_id__, __version__; commit = (__commit_id__ or '').removeprefix('g'); assert len(commit) >= 8 and '${VLLM_PRECOMPILED_WHEEL_COMMIT}'.startswith(commit), __commit_id__; assert __version__ == '${VLLM_VERSION}', __version__" # The final stage just copies the working dir on top of the # dependencies image. This uses --link so that we can skip diff --git a/docker/Dockerfile.cuda b/docker/Dockerfile.cuda index f6fe19b9edf..75f880303f9 100644 --- a/docker/Dockerfile.cuda +++ b/docker/Dockerfile.cuda @@ -1,6 +1,14 @@ -ARG BASE_IMAGE=vllm/vllm-openai:v0.32.0 +ARG BASE_IMAGE=vllm/vllm-openai:v0.33.0@sha256:d825baf9ce4b1b6bfc5ecc7b79880c7e6748c65a72d697ce3d47607aafa84c6e FROM ${BASE_IMAGE} +ARG BASE_IMAGE +# The base's optional EFA NCCL override conflicts with Torch 2.14 metadata. +# Use the official wheel's declared NCCL dependency for the Omni runtime. +ENV UV_OVERRIDE="" +# Verify the final runtime using the same frozen wheel as CUDA CI, on the +# matching published base image. Omni runtime qualification is pending. +ARG VLLM_PRECOMPILED_WHEEL_COMMIT=4ba787ec9959e566d3b587a796dff07673324da5 +ARG VLLM_VERSION=0.33.0 ARG COMMON_WORKDIR=/app WORKDIR ${COMMON_WORKDIR} @@ -15,7 +23,42 @@ RUN mkdir -p ${COMMON_WORKDIR}/vllm-omni # Step 2: Copy vllm-omni code and install COPY . ${COMMON_WORKDIR}/vllm-omni -RUN cd ${COMMON_WORKDIR}/vllm-omni && uv pip install --python "$(python3 -c 'import sys; print(sys.executable)')" --no-cache-dir "." +# Apply the same guarded bootstrap repair and dependency pins as CUDA CI. +RUN python3 vllm-omni/tools/prepare_ci_telemetry.py "${BASE_IMAGE}" && \ + uv pip uninstall --python "$(python3 -c 'import sys; print(sys.executable)')" opentelemetry-exporter-otlp-common +RUN cd ${COMMON_WORKDIR}/vllm-omni && uv pip install --python "$(python3 -c 'import sys; print(sys.executable)')" --no-cache-dir \ + -r requirements/cuda-bootstrap.txt \ + "nvidia-nccl-cu13==2.30.7" \ + "vllm @ https://wheels.vllm.ai/${VLLM_PRECOMPILED_WHEEL_COMMIT}/vllm-${VLLM_VERSION}-cp38-abi3-manylinux_2_28_$(uname -m).whl" \ + "." + +# The published final base still ships an LMCache wheel with the Torch 2.13 +# CUDA ABI. Rebuild the same release against Torch 2.14, retaining every +# published CUDA architecture. +RUN uv pip install --python "$(python3 -c 'import sys; print(sys.executable)')" \ + -r ${COMMON_WORKDIR}/vllm-omni/requirements/cuda-bootstrap.txt \ + ninja "packaging>=24.2" "setuptools>=77.0.3,<81.0.0" \ + "setuptools_scm>=8" "grpcio==1.78.0" "grpcio-tools==1.78.0" wheel +RUN PYTHON_SITE=$(python3 -c 'import sysconfig; print(sysconfig.get_path("purelib"))') && \ + test -f "${PYTHON_SITE}/nvidia/cu13/include/cusparse.h" && \ + test -f "${PYTHON_SITE}/nvidia/cusparselt/include/cusparseLt.h" && \ + CUDA_VISIBLE_DEVICES="" CUDA_HOME=/usr/local/cuda BUILD_WITH_CUDA=1 MAX_JOBS=3 NVCC_THREADS=1 \ + TORCH_CUDA_ARCH_LIST="7.5;8.0;8.6;8.9;9.0;10.0;11.0;12.0" \ + CPATH="${PYTHON_SITE}/nvidia/cu13/include:${PYTHON_SITE}/nvidia/cusparselt/include${CPATH:+:${CPATH}}" \ + uv pip install --python "$(python3 -c 'import sys; print(sys.executable)')" \ + --no-cache --no-binary lmcache --no-build-isolation-package lmcache --reinstall-package lmcache \ + -r ${COMMON_WORKDIR}/vllm-omni/requirements/cuda-bootstrap.txt "lmcache==0.5.6" +RUN uv pip check --python "$(python3 -c 'import sys; print(sys.executable)')" + +# cuDNN 9.24 causes the original Qwen-Image 2.1 pixel gate to fail under +# Torch 2.14. Keep Torch's declared Python dependencies and load the verified +# NVIDIA 9.25.1 wheel in a private target outside Python's import path. +RUN uv pip install --target /opt/vllm-omni-cudnn --no-deps --only-binary=:all: \ + --require-hashes -r ${COMMON_WORKDIR}/vllm-omni/requirements/cudnn-runtime.txt +ENV LD_PRELOAD="/opt/vllm-omni-cudnn/nvidia/cudnn/lib/libcudnn.so.9 ${LD_PRELOAD:-}" +ENV LD_LIBRARY_PATH="/opt/vllm-omni-cudnn/nvidia/cudnn/lib:${LD_LIBRARY_PATH:-/usr/local/nvidia/lib:/usr/local/nvidia/lib64}" + +RUN python3 -c "from vllm._version import __commit_id__, __version__; commit = (__commit_id__ or '').removeprefix('g'); assert len(commit) >= 8 and '${VLLM_PRECOMPILED_WHEEL_COMMIT}'.startswith(commit), __commit_id__; assert __version__ == '${VLLM_VERSION}', __version__" RUN ln -sf /usr/bin/python3 /usr/bin/python diff --git a/docker/Dockerfile.npu b/docker/Dockerfile.npu index f8a6567aff7..632f89154c0 100644 --- a/docker/Dockerfile.npu +++ b/docker/Dockerfile.npu @@ -1,3 +1,7 @@ +# Legacy Ascend 0.32 runtime; this default does not qualify the vLLM 0.33 +# alignment checkout. A matching 0.33 vLLM-Ascend image and NPU validation +# are required before using this checkout. Override VLLM_ASCEND_IMAGE and +# VLLM_ASCEND_VERSION only with that verified runtime; do not relabel 0.32. ARG VLLM_ASCEND_IMAGE=quay.io/atlas-ci/vllm-ascend ARG VLLM_ASCEND_VERSION=v0.32.0 ARG VLLM_ASCEND_DEVICE_SUFFIX= diff --git a/docker/Dockerfile.npu.ci b/docker/Dockerfile.npu.ci index cca7f6a2494..cc43f6a408a 100644 --- a/docker/Dockerfile.npu.ci +++ b/docker/Dockerfile.npu.ci @@ -1,3 +1,7 @@ +# Legacy Ascend 0.32 runtime; this default does not qualify the vLLM 0.33 +# alignment checkout. A matching 0.33 vLLM-Ascend image and NPU validation +# are required before using this checkout. Override VLLM_ASCEND_IMAGE and +# VLLM_ASCEND_VERSION only with that verified runtime; do not relabel 0.32. ARG VLLM_ASCEND_IMAGE=quay.io/atlas-ci/vllm-ascend ARG VLLM_ASCEND_VERSION=v0.32.0 ARG VLLM_ASCEND_DEVICE_SUFFIX= diff --git a/docker/Dockerfile.rocm b/docker/Dockerfile.rocm index 19e635bfeb7..6ceddc81e95 100644 --- a/docker/Dockerfile.rocm +++ b/docker/Dockerfile.rocm @@ -1,7 +1,6 @@ -# Keep the ROCm runtime aligned with the vLLM release targeted by -# docker/Dockerfile.ci. A stale base image can build successfully while failing -# at runtime when vLLM-Omni imports APIs added by the newer vLLM release. -ARG BASE_IMAGE=vllm/vllm-openai-rocm:v0.32.0 +# Use the published v0.33.0 base, then build the frozen matching source +# below. This combination still requires Omni runtime qualification. +ARG BASE_IMAGE=vllm/vllm-openai-rocm:v0.33.0 # Keep the dependency-heavy portion independent from the source tree. The # registry cache is keyed on these inputs, while the final source copy can be @@ -9,14 +8,12 @@ ARG BASE_IMAGE=vllm/vllm-openai-rocm:v0.32.0 FROM ${BASE_IMAGE} AS deps # NOTE: REMINDER to vLLM-Omni rebase maintainer -# BASE_IMAGE pins the default prebuilt runtime when USE_NIGHTLY_BUILD=0. -# VLLM_VERSION_OR_COMMIT_HASH is consumed only by the opt-in source rebuild when -# USE_NIGHTLY_BUILD=1. Keep both aligned with docker/Dockerfile.ci's -# release version. This is enforced by tests/buildkite/test_rocm_dockerfile.py. An -# intentional custom source pin must update that regression in the same change, -# and should return to the stable release as soon as one is available. -ARG USE_NIGHTLY_BUILD=0 -ARG VLLM_VERSION_OR_COMMIT_HASH=v0.32.0 +# Build from the exact release-branch source until a matching prebuilt ROCm image is available. +# Set USE_NIGHTLY_BUILD=0 only with a BASE_IMAGE that already contains the +# target runtime. CUDA CI and the source commit are checked together in +# tests/buildkite/test_rocm_dockerfile.py. +ARG USE_NIGHTLY_BUILD=1 +ARG VLLM_VERSION_OR_COMMIT_HASH=4ba787ec9959e566d3b587a796dff07673324da5 ARG ARG_PYTORCH_ROCM_ARCH ENV PYTORCH_ROCM_ARCH=${ARG_PYTORCH_ROCM_ARCH:-${PYTORCH_ROCM_ARCH}} @@ -75,11 +72,15 @@ WORKDIR ${COMMON_WORKDIR}/vllm-omni RUN mkdir -p vllm_omni && touch vllm_omni/__init__.py COPY pyproject.toml setup.py ./ COPY requirements/ requirements/ +COPY tools/build_gradio_pillow13.py tools/build_gradio_pillow13.py +COPY tools/gradio_pillow13/ tools/gradio_pillow13/ # Scope the override to Omni; dependency source builds need their own versions. ENV VLLM_OMNI_VERSION_OVERRIDE=0.0.0 -RUN uv pip install --python "$(python3 -c 'import sys; print(sys.executable)')" \ - --no-cache-dir -e ".[dev]" --no-build-isolation && \ +RUN python3 tools/build_gradio_pillow13.py --output-dir /tmp/gradio-pillow13 && \ + uv pip install --python "$(python3 -c 'import sys; print(sys.executable)')" \ + --no-cache-dir /tmp/gradio-pillow13/gradio-6.31.0+omni.pillow13-py3-none-any.whl \ + -e ".[dev]" --no-build-isolation && \ python3 -m pip uninstall -y onnxruntime && \ python3 -m pip install --no-cache-dir --force-reinstall --no-deps \ "onnxruntime-rocm>=1.22.2" && \ diff --git a/docker/Dockerfile.xpu b/docker/Dockerfile.xpu index 6105513c3ee..1d96f74c20e 100644 --- a/docker/Dockerfile.xpu +++ b/docker/Dockerfile.xpu @@ -1,4 +1,4 @@ -# Target vLLM version: v0.32.0. Keep VLLM_VERSION aligned with the target vLLM +# Target vLLM version: v0.33.0. Keep VLLM_VERSION aligned with the target vLLM # release on every rebase. # # This file layers vLLM-Omni on top of the published upstream XPU image, which @@ -8,8 +8,8 @@ # pulls that base automatically, so there is nothing to build here beyond the # Omni layers. Mirrors how docker/Dockerfile.ci consumes vllm/vllm-openai. # -# If the tag does not exist yet (targeting an unreleased vLLM commit), build the -# base from upstream's own Dockerfile and point VLLM_BASE at the result: +# If the matching image cannot be fetched, build the base from upstream's +# Dockerfile at the selected tag and point VLLM_BASE at the result: # # git clone --depth 1 --branch $VLLM_VERSION https://github.com/vllm-project/vllm /tmp/vllm # docker build -t vllm-openai-xpu:local --target vllm-openai \ @@ -18,7 +18,7 @@ # # .buildkite/intel/scripts/run-xpu-test.sh automates that pull-then-build fallback. ARG VLLM_BASE_IMAGE=vllm/vllm-openai-xpu -ARG VLLM_VERSION=v0.32.0 +ARG VLLM_VERSION=v0.33.0 ARG VLLM_BASE=${VLLM_BASE_IMAGE}:${VLLM_VERSION} FROM ${VLLM_BASE} AS vllm-omni @@ -41,8 +41,10 @@ ENV VLLM_OMNI_TARGET_DEVICE=xpu # vLLM XPU wheel requires triton==3.8.0+xpu; installing the old triton-xpu # overlay would replace that runtime with an incompatible version. RUN --mount=type=cache,target=/root/.cache/uv \ + python3 tools/build_gradio_pillow13.py --output-dir /tmp/gradio-pillow13 && \ uv pip install --no-build-isolation --excludes <(printf "torch\ntriton-xpu\n") \ - ".[dev]" + /tmp/gradio-pillow13/gradio-6.31.0+omni.pillow13-py3-none-any.whl ".[dev]" && \ + uv pip check FROM vllm-omni AS vllm-omni-openai diff --git a/docs/configuration/README.md b/docs/configuration/README.md index 9749cb20d0c..6c33a996ac4 100644 --- a/docs/configuration/README.md +++ b/docs/configuration/README.md @@ -2,7 +2,7 @@ This section lists the most common options for running vLLM-Omni. -For options within a vLLM Engine, please refer to the [vLLM 0.32 configuration guide](https://docs.vllm.ai/en/v0.32.0/configuration/index.html). +For options within a vLLM Engine, please refer to the [vLLM 0.33 configuration guide](https://docs.vllm.ai/en/v0.33.0/configuration/index.html). Each model defines fixed topology in a registered `PipelineConfig` and runtime overrides in a deploy YAML. diff --git a/docs/configuration/environment_variables.md b/docs/configuration/environment_variables.md index 83d15b757e5..e31ad4796d9 100644 --- a/docs/configuration/environment_variables.md +++ b/docs/configuration/environment_variables.md @@ -99,7 +99,7 @@ environment variables. | `VLLM_OMNI_ENABLE_RUN_ENDPOINT` | `1` to enable; default off | Read when the server initializes its endpoints | Indicates that `POST /v1/run` should be enabled for independent stage execution; this should only be used by trusted callers. | Experimental | | `VLLM_OMNI_EVENT_DRIVEN_ORCH` | Defaults to on; `0`, `false`, `no` or `off` selects the legacy poll loop | Orchestrator stage-output loop; resolved at `Orchestrator` construction | Unset selects event-driven orchestration for every pipeline. Values are stripped and case-normalized; `1`, `true`, `yes` or `on` keep it enabled and any other value selects the legacy poll loop. The serving-side final-output loop always wakes on orchestrator output, independent of this setting. Set before server startup. | Experimental | | `VLLM_OMNI_INPUT_WAIT_TIMEOUT_S` | Float seconds; default `600`; `<=0` disables | Full-payload input coordinator, not async-chunk transfer; read when the scheduler module imports in each worker | Environment-only setting. A non-float logs a warning and uses `600`. | Stable operational control | -| `VLLM_OMNI_MOE_PREFILL_WARMUP` | Defaults to on; `0`, `false`, `no` or `off` disables | AR stages with fused-MoE layers (CUDA/Triton, v1 model runner); read in each worker during startup warmup | Environment-only. Unset or any other value keeps it on; values are stripped and case-normalized. When on, startup runs one dummy forward per fused-MoE tuning batch size and its odd `+1` neighbour, for sizes above the largest CUDA graph capture size up to `max_num_batched_tokens`, so `fused_moe_kernel` configs do not JIT-compile during the first requests. Skipped on the v2 model runner. This adds startup time (tens of seconds on a cold cache for Qwen3-Omni). | Experimental performance control | +| `VLLM_OMNI_MOE_PREFILL_WARMUP` | Defaults to on; `0`, `false`, `no` or `off` disables | AR stages with fused-MoE layers (CUDA/Triton, v1 and v2 model runners); read in each worker during startup warmup | Environment-only. Unset or any other value keeps it on; values are stripped and case-normalized. When on, startup runs one dummy forward per fused-MoE tuning batch size and its odd `+1` neighbour, for sizes above the largest CUDA graph capture size up to `max_num_batched_tokens`, so `fused_moe_kernel` configs do not JIT-compile during the first requests. This adds startup time (tens of seconds on a cold cache for Qwen3-Omni). | Experimental performance control | | `VLLM_OMNI_ORCH_MONITOR_PATH` | Filesystem path; default `/vllm_omni_orch_monitor_.json` | Orchestrator monitor enabled by `--enable-orch-monitor`; read when the monitor is created | Environment-only path override. Parent directories are created; write errors are logged. | Diagnostic | | `VLLM_OMNI_SPEAKER_REGISTRATION_POLICY` | `overwrite` or `immutable`; default `overwrite` | Speech server; read when speaker storage initializes | Environment-only setting. `immutable` rejects re-registering an existing uploaded voice name until it is deleted; any other value raises `ValueError` at startup. | Experimental | | `VLLM_OMNI_VIDEO_SYNC_TIMEOUT` | Float seconds; default `600` | Synchronous Videos API; read when the API server module imports | Environment-only setting. A non-float raises `ValueError` during import. | Experimental | @@ -191,7 +191,7 @@ their keys only. ## Inherited vLLM variables vLLM-Omni also reads variables through its aligned vLLM dependency. Refer to -the [vLLM 0.32 environment-variable reference](https://docs.vllm.ai/en/v0.32.0/configuration/env_vars.html) +the [vLLM 0.33 environment-variable reference](https://docs.vllm.ai/en/v0.33.0/configuration/env_vars/) for their definitions. This includes vLLM launch, cache, logging, plugin, ROCm, XPU, ModelScope, and FlashInfer workspace settings. diff --git a/docs/contributing/README.md b/docs/contributing/README.md index 10d87377fb5..79b89ae5bab 100644 --- a/docs/contributing/README.md +++ b/docs/contributing/README.md @@ -19,7 +19,7 @@ source .venv/bin/activate vLLM-Omni is quickly evolving, please see the [installation guide](../getting_started/installation/README.md) for details. It's recommended to build from source to provide the latest development environment. !!! tip - This source checkout requires Python 3.11 to 3.13, matching the Python floor of vLLM 0.32. We recommend developing with Python 3.12 to match our CI environment. + This source checkout requires Python 3.11 to 3.13, matching the Python floor of vLLM 0.33. We recommend developing with Python 3.12 to match our CI environment. ### Adding a new model to vLLM-Omni diff --git a/docs/contributing/ci/test_execution_guide.md b/docs/contributing/ci/test_execution_guide.md index e2c88fb435c..ebbc41f02d8 100644 --- a/docs/contributing/ci/test_execution_guide.md +++ b/docs/contributing/ci/test_execution_guide.md @@ -15,13 +15,24 @@ vLLM-Omni is built based on vLLM. You can follow [install guide](../../getting_s ### Test Case Dependencies -When running test cases, you may need to install the following dependencies: +After installing the matching platform runtime, run from the repository root. +The `[dev]` extra includes VoxCPM 2.0.3 and its Gradio dependency; use the +[temporary source backport](https://docs.vllm.ai/projects/vllm-omni/en/latest/getting_started/installation/#source-development-and-demo-dependencies) +in the same dependency install. These setup commands do not establish CI or +native model acceptance. ```bash -uv pip install ".[dev]" +python tools/build_gradio_pillow13.py --output-dir /tmp/gradio-pillow13 +uv pip install ".[dev]" \ + /tmp/gradio-pillow13/gradio-6.31.0+omni.pillow13-py3-none-any.whl +uv pip check apt-get install -y espeak-ng jq ``` +The CUDA CI image verifies the vLLM 0.33 runtime and imports the rebuilt +`lmcache.cuda_ops` before model tests. See the installation guide for the +base-image dependency repairs; metadata checks alone do not verify native ABI. + ## Running Tests Our test scripts use the pytest framework. First, please use `git clone https://github.com/vllm-project/vllm-omni.git` to download the vllm-omni source code. Then, in the root directory of vllm-omni, you can run the following commands in your local test environment to execute the corresponding test cases. diff --git a/docs/contributing/profiling.md b/docs/contributing/profiling.md index b86f3c4f29f..350a4305d1a 100644 --- a/docs/contributing/profiling.md +++ b/docs/contributing/profiling.md @@ -87,7 +87,7 @@ Supported fields: | `warmup_iterations` | Torch-profiler warmup iterations. | | `active_iterations` | Torch-profiler active iterations. | -For detailed explanations of the fields, please refer to upstream vLLM implementation [vllm/config/profiler.py](https://github.com/vllm-project/vllm/blob/v0.20.1/vllm/config/profiler.py) +For detailed explanations of the fields, please refer to upstream vLLM implementation [vllm/config/profiler.py](https://github.com/vllm-project/vllm/blob/v0.33.0/vllm/config/profiler.py) ### Minimal configurations by output diff --git a/docs/design/feature/attention_backend_selection.md b/docs/design/feature/attention_backend_selection.md index d1d1315f11b..6add83bbc7c 100644 --- a/docs/design/feature/attention_backend_selection.md +++ b/docs/design/feature/attention_backend_selection.md @@ -53,6 +53,14 @@ The active `OmniPlatform` owns compatibility policy through - account for head size and other platform-visible constraints; and - return a qualified class path implementing `AttentionBackend`. +Hub preflight and execution share the same version resolver. FlashAttention 2 +tries versions 1, 2, then 3 so Torch stable-ABI builds remain available when +older per-Torch builds do not match. An explicit selection still fails when +no compatible published build resolves. + +Explicit backend validation must avoid initializing unrelated runtimes. The +cuDNN version probe belongs to automatic Blackwell fallback selection. + The selector must not duplicate device capability or package-availability policy that belongs to the platform. diff --git a/docs/design/fullduplex.md b/docs/design/fullduplex.md index 0fbea375622..6a26cc8753e 100644 --- a/docs/design/fullduplex.md +++ b/docs/design/fullduplex.md @@ -245,7 +245,7 @@ vllm_omni/ │ │ ├── session_attachment.py DuplexSessionAttachmentRegistry (resume tokens, replay journal) │ │ ├── audio_encoding.py encode_audio, injected into DuplexOmniEngine for the plugin's data plane │ │ ├── openai.py dispatch_realtime_websocket (turn-based /v1/realtime, when duplex mode is not selected) -│ │ ├── warmup.py startup warmup probe; wait_for_duplex_warmup holds clients until it finishes +│ │ ├── warmup.py startup warmup probe; holds native session admission after WebSocket acceptance │ │ └── websocket.py websocket send/close/receive helpers │ └── openai/api_server.py builds DuplexOmni for duplex models; session-backed app state ├── protocol/ SHARED WIRE CODEC (no engine / no model / no transport) diff --git a/docs/design/minicpm_o45_mrv2_performance.md b/docs/design/minicpm_o45_mrv2_performance.md index 00283ac8a19..ffbb3a13ee7 100644 --- a/docs/design/minicpm_o45_mrv2_performance.md +++ b/docs/design/minicpm_o45_mrv2_performance.md @@ -106,7 +106,7 @@ stages: num_speculative_tokens: 7 ``` -vLLM 0.32 resolves a draft config with `architectures: ["DSparkDraftModel"]` +vLLM 0.33 resolves a draft config with `architectures: ["DSparkDraftModel"]` and `model_type: qwen3` to `Qwen3DSparkModel`. Drafters that read Thinker layers require the Thinker on MRv2; the V1 Thinker rejects them at load. Validation covers CUDA turn sessions with greedy text and speech output. diff --git a/docs/design/module/ar_runtime.md b/docs/design/module/ar_runtime.md index 44785f4b1d8..aba6faf0234 100644 --- a/docs/design/module/ar_runtime.md +++ b/docs/design/module/ar_runtime.md @@ -32,6 +32,41 @@ last_reviewed: 2026-07-16 The AR runtime extends vLLM scheduling and worker execution for omni-stage inputs and outputs while preserving vLLM scheduling and cache semantics. +After applying sampled output, the AR scheduler publishes only completed KV +blocks to upstream prefix caching, including on its plain-decode path. Async +placeholders and speculative steps use `num_computed_tokens - num_in_flight_tokens` +as the materialized frontier. Stale output and requests that were not running +before the update do not publish blocks. Publication precedes stop cleanup so +the last completed blocks remain reusable after a request finishes. + +For Bagel image editing, the vLLM 0.33 prefix-range helper receives shallow +`image` views of the custom `img2img` features so both VAE and ViT embedding +spans retain bidirectional attention. The original features keep their modality, +encoder data, cache identifiers and placeholder masks. Upstream still computes +the ranges and applies sliding-window and padding rules. This adapter is installed +from `vllm_omni.patch` for both upstream runner implementations; it is unnecessary +with vLLM versions that build the ranges inline. + +MOSS Local's fused and separate depth-attention paths share an explicitly +rounded initial weighted value and the same FMA accumulation order. This +preserves bitwise parity when the compiler inlines attention into sampling, +including SM89 code generation affected by LLVM SLP reduction reordering. + +## Qwen3-Omni Talker sampling + +The MRv2 Talker adapter follows the selected vLLM sampler's dispatch contract, +including its `has_topp_only_rows` predicate. Seeded rows pass that predicate to +top-k/top-p masking before drawing from their request-local V1 generators. +Unseeded and greedy rows keep upstream sampling. When mixed rows have already +been masked, their upstream call disables repeated masking and fused sampling. + +Custom samplers register their selected JIT kernels while the runner's startup +registry is active. The Talker's seeded V1 stream uses native top-k/top-p +masking even when unseeded sampling selects FlashInfer, so it registers those +native kernels explicitly. AR workers also warm the existing fused-MoE prefill +batch sizes on both runner implementations before CUDA graph capture and JIT +monitoring. These startup passes leave EPLB load recording disabled. + ## Full-payload input scheduling For a downstream, non-async stage with `requires_full_payload_input`, diff --git a/docs/design/module/cache_management.md b/docs/design/module/cache_management.md index 22245f89568..c9bd47d2256 100644 --- a/docs/design/module/cache_management.md +++ b/docs/design/module/cache_management.md @@ -26,6 +26,20 @@ last_reviewed: 2026-07-16 Cache management defines reusable state, cache identity, validity, lifecycle, and eviction across AR and diffusion execution. +When upstream reports a multimodal encoder-cache size mismatch, both AR and +generation schedulers drain that request into a terminal error output and +discard its pending prewarm payload. A rejected cache identifier must not +leave a request waiting indefinitely or replay the same error on later steps. + +When HiSparse defers asynchronous KV loads, a completed nonempty scheduler +step advances the upstream fence and drains both deferred frees and loads. +Empty steps leave the fence unchanged so a reused block cannot be loaded +while an earlier GPU step still reads it. + +Diffusion supplies worker attention specs to the native KV planner and uses +its default grouping. Its model-config adapter declares that it does not use +a registered model's custom `SupportsKVCacheGrouping` hook. + ## Candidate invariants ### CACHE-INV-001: Cache identity is complete diff --git a/docs/design/module/error_contracts.md b/docs/design/module/error_contracts.md index cc47d60de7c..cf5fa8b54df 100644 --- a/docs/design/module/error_contracts.md +++ b/docs/design/module/error_contracts.md @@ -75,6 +75,10 @@ with code 503 in HTTP and SSE responses. Cache feedback for cancelled requests is applied without failing an unrelated request; other engine errors keep their existing classification. +When tracing is enabled, pipeline and stage error spans use that same resolved +classification, including upstream `RequestError.code` and cache-miss +`ServiceUnavailableError`. Unclassified stage failures use `stage_error`. + Offline generation discards non-fatal request-scoped errors whose requests are no longer registered, including queued errors from an aborted batch. Global errors and fatal errors still propagate, even when a fatal error diff --git a/docs/design/module/input_output_modality_contracts.md b/docs/design/module/input_output_modality_contracts.md index 703832cab9a..28031afaadf 100644 --- a/docs/design/module/input_output_modality_contracts.md +++ b/docs/design/module/input_output_modality_contracts.md @@ -92,7 +92,7 @@ public rendering of its error fields belong to `error_contracts.md`. ## Current audio decoding compatibility -Omni retains soundfile-first automatic audio decoding when using vLLM 0.32. +Omni retains soundfile-first automatic audio decoding when using vLLM 0.33. Decoder rounding can affect BF16 model inputs and generated responses. Missing soundfile dependencies and unsupported formats fall back to upstream decoders; explicit backend selections keep upstream behavior. Encoded-size, diff --git a/docs/design/module/profiling.md b/docs/design/module/profiling.md index 0da52e53939..db65e003403 100644 --- a/docs/design/module/profiling.md +++ b/docs/design/module/profiling.md @@ -32,6 +32,13 @@ validation to the engine configuration. See the [offline profiling guide](../../contributing/profiling.md#offline-qwen25-omni-and-cosyvoice3-examples) for commands and the shared helper API. +AR, diffusion and NPU worker profiling forward vLLM's per-session +`profiler_kwargs` to the +Torch profiler constructor. Overrides apply only while profiling is inactive; +omitting them on a later session restores the configured defaults. Stage/rank +trace naming and the default trace export handler remain in place unless the +caller explicitly overrides the handler. + ## Candidate invariants ### PROFILE-INV-001: Profiling is opt-in diff --git a/docs/getting_started/installation/README.md b/docs/getting_started/installation/README.md index 2fade2c013b..a109727c89e 100644 --- a/docs/getting_started/installation/README.md +++ b/docs/getting_started/installation/README.md @@ -1,7 +1,11 @@ # Installation !!! important - [vLLM-Omni 0.32.0](https://github.com/vllm-project/vllm-omni/releases/tag/v0.32.0) is published with a Python wheel and CUDA Docker images. It requires the matching vLLM 0.32.0 release and Python `>=3.11,<3.14`; the CUDA installation uses CUDA 13.0 and PyTorch 2.13.0. This source checkout also targets vLLM 0.32.0 at `5df63a7203cdb9328d824a27b682e77c7b9928be`, but `main` can contain changes beyond the tagged Omni release. Choose the installation path for your platform: CUDA release artifacts do not establish ROCm, XPU, MUSA, or NPU runtime compatibility. The platform guides retain their separate version and validation requirements. + This source checkout targets vLLM 0.33.0 at `4ba787ec9959e566d3b587a796dff07673324da5`, with CUDA 13.0 and PyTorch 2.14.0. Full CI and ROCm/XPU runtime validation remain pending. CUDA Dockerfiles use the published 0.33.0 image and verify the frozen final wheel; ROCm builds the same commit from source on its matching 0.33.0 base. The published Omni 0.32.0 wheel instructions below retain their matching vLLM 0.32.0 dependency. + +The CUDA Dockerfiles rebuild LMCache 0.5.6 against Torch 2.14 because the base's +Torch 2.13 extension fails to import. They retain its CUDA architecture coverage +and check dependency consistency. Release-container validation remains pending. vLLM-Omni supports the following hardware platforms: @@ -11,3 +15,30 @@ vLLM-Omni supports the following hardware platforms: - [Intel XPU](gpu.md) - [MThreads MUSA](gpu.md) - [NPU](npu.md) + +## Source development and demo dependencies + +For this 0.33 source checkout, the `[dev]` and `[demo]` extras need a temporary +Gradio source backport to satisfy vLLM's Pillow 13 requirement. From the repository +root, after installing the matching vLLM runtime for your platform. The builder +requires Git and uv 0.11.8 or newer: + +```bash +python tools/build_gradio_pillow13.py --output-dir /tmp/gradio-pillow13 +uv pip install -e '.[dev,demo]' \ + /tmp/gradio-pillow13/gradio-6.31.0+omni.pillow13-py3-none-any.whl +uv pip check +``` + +Select just `.[dev]` or `.[demo]` when appropriate, keeping the explicit wheel +in the same install command. Preserve the platform's target-device and build +options shown in the GPU installation instructions. The isolated build verifies +pinned source and wheel hashes and retains the original licenses. + +The local distribution `gradio==6.31.0+omni.pillow13` derives from the exact +official Gradio 6.31.0 source archive. It backports the Pillow `<14` bound and +I-mode PNG error guidance from upstream [Gradio PR #14005](https://github.com/gradio-app/gradio/pull/14005), +merged to `7.0-dev`; it is not an official Gradio 6 release. Hash-pinned source, +patch and build-tool versions are recorded under `tools/gradio_pillow13/`; +the builder preserves the original licenses. I-mode PNG output requires an explicit `I;16` +conversion for in-range values or a format such as TIFF to preserve 32-bit values. diff --git a/docs/getting_started/installation/gpu/cuda.inc.md b/docs/getting_started/installation/gpu/cuda.inc.md index 72fd61a4c52..a0f9c8a0715 100644 --- a/docs/getting_started/installation/gpu/cuda.inc.md +++ b/docs/getting_started/installation/gpu/cuda.inc.md @@ -1,16 +1,19 @@ + + # --8<-- [start:requirements] - GPU: compute capability 7.5 or higher (e.g., T4, RTX20xx, A100, L4, H100, etc.) # --8<-- [end:requirements] + # --8<-- [start:set-up-using-python] vLLM-Omni depends on the matching major/minor release of vLLM. This source -checkout targets vLLM 0.32.0 at commit -`5df63a7203cdb9328d824a27b682e77c7b9928be` at the `v0.32.0` tag on `releases/v0.32.0`. +checkout targets vLLM 0.33.0 at commit +`4ba787ec9959e566d3b587a796dff07673324da5` on `releases/v0.33.0` (tagged `v0.33.0`). The combined Python metadata range for this checkout and vLLM is -`>=3.11,<3.14`. The tested CUDA runtime uses Python 3.12, PyTorch 2.13.0 and -CUDA 13.0 (`cu130`). +`>=3.11,<3.14`. The final CUDA wheel uses PyTorch 2.14.0 and CUDA 13.0 +(`cu130`); the current validation environment uses Python 3.12. The pre-built Omni wheel instructions below install the published 0.32.0 release with its matching vLLM dependency. @@ -28,6 +31,7 @@ Therefore, it is recommended to install vLLM and vLLM-Omni with a **fresh new** These pre-built wheel instructions install the published vLLM-Omni 0.32.0 release. vLLM-Omni is built based on vLLM. Please install it with command below. + ```bash uv pip install vllm==0.32.0 --torch-backend=cu130 ``` @@ -39,6 +43,7 @@ uv pip install vllm-omni==0.32.0 ``` To run Gradio demos, also install the optional extras: + ```bash uv pip install 'vllm-omni[demo]==0.32.0' ``` @@ -48,60 +53,110 @@ uv pip install 'vllm-omni[demo]==0.32.0' # --8<-- [start:build-wheel-from-source] #### Installation of vLLM -If you do not need to modify source code of vLLM, install the release-branch -0.32.0 wheel directly from the tested commit: + +If you do not need to modify source code of vLLM, install the final +0.33.0 wheel directly from the frozen commit: ```bash uv pip install --torch-backend=cu130 \ - "vllm @ https://wheels.vllm.ai/5df63a7203cdb9328d824a27b682e77c7b9928be/vllm-0.32.0-cp38-abi3-manylinux_2_28_x86_64.whl" + "vllm @ https://wheels.vllm.ai/4ba787ec9959e566d3b587a796dff07673324da5/vllm-0.33.0-cp38-abi3-manylinux_2_28_x86_64.whl" ``` This command selects the Linux x86_64 CUDA 13.0 wheel and matching PyTorch backend. It requires a CUDA 13 compatible driver. For another architecture, CUDA variant or an existing incompatible PyTorch installation, build vLLM -from source instead. The [vLLM 0.32.0 release](https://github.com/vllm-project/vllm/releases/tag/v0.32.0) -and its commit-index and PyPI wheels are available. The PyPI x86_64 wheel -embeds the same release commit and CUDA/PyTorch requirements. The command -above retains the exact tested artifact. The published vLLM-Omni 0.32.0 -wheel is available through the pre-built wheel instructions above; use the -source checkout for development. +from source instead. The final commit-index wheel is available; native model +and full CI validation remain pending. Install +Omni from this source checkout. #### Installation of vLLM-Omni + Since vllm-omni is rapidly evolving, it's recommended to install it from source + ```bash git clone https://github.com/vllm-project/vllm-omni.git cd vllm-omni uv pip install -e . ``` -To run Gradio demos, install with optional extras: +The Torch 2.14 CUDA environment declares cuDNN 9.24 as an exact Python +package dependency. This checkout's CUDA images load a separate, pinned +cuDNN 9.25.1 runtime for correct `CUDNN_ATTN` image outputs. The normal wheel +installer verifies the pinned NVIDIA wheel hash and retains its license +outside Python site-packages. +The original Qwen-Image 2.1 IT2I test passes on L20X with this runtime; +native H100 and full CI qualification remain pending. + +For a source installation, configure the same runtime before starting Python: + +```bash +OMNI_CUDNN_RUNTIME_DIR="$(python -c 'import sys; print(sys.prefix)')/omni-cudnn-runtime" +uv pip install --target "$OMNI_CUDNN_RUNTIME_DIR" --no-deps --only-binary=:all: \ + --require-hashes -r requirements/cudnn-runtime.txt +export LD_PRELOAD="$OMNI_CUDNN_RUNTIME_DIR/nvidia/cudnn/lib/libcudnn.so.9${LD_PRELOAD:+:$LD_PRELOAD}" +export LD_LIBRARY_PATH="$OMNI_CUDNN_RUNTIME_DIR/nvidia/cudnn/lib${LD_LIBRARY_PATH:+:$LD_LIBRARY_PATH}" +python -c 'import torch; assert torch.backends.cudnn.version() == 92501' +uv pip check +``` + +The output directory must be fresh. Reuse its exported library paths for +subsequent launches. See NVIDIA's [cuDNN compatibility policy](https://docs.nvidia.com/deeplearning/cudnn/backend/v9.25.1/developer/forward-compatibility.html) +and [support matrix](https://docs.nvidia.com/deeplearning/cudnn/backend/v9.25.1/reference/support-matrix.html) +for binary, driver and hardware requirements. Python packages and the +selected vLLM/Torch binaries remain installed with their original metadata. + +To run Gradio demos from this source checkout, build the temporary Gradio +backport and pass its wheel in the same normal dependency install: + ```bash -uv pip install -e '.[demo]' +python tools/build_gradio_pillow13.py --output-dir /tmp/gradio-pillow13 +uv pip install -e '.[demo]' \ + /tmp/gradio-pillow13/gradio-6.31.0+omni.pillow13-py3-none-any.whl ``` +Development installs use the same built wheel: + +```bash +uv pip install -e '.[dev]' \ + /tmp/gradio-pillow13/gradio-6.31.0+omni.pillow13-py3-none-any.whl +uv pip check +``` + +See the [source dependency instructions](https://docs.vllm.ai/projects/vllm-omni/en/latest/getting_started/installation/#source-development-and-demo-dependencies) +for exact source provenance and qualification limits. This local Gradio 6.31 +backport supplies the upstream Pillow 13 compatibility changes; browser, +native model and full CI acceptance remain pending. +
(Optional) Installation of vLLM from source If you want to check, modify or debug with source code of vLLM, install the library from source with the following instructions: ```bash git clone https://github.com/vllm-project/vllm.git cd vllm -git checkout 5df63a7203cdb9328d824a27b682e77c7b9928be +git checkout 4ba787ec9959e566d3b587a796dff07673324da5 ``` + Set up environment variables to get pre-built wheels. If there are internet problems, just download the whl file manually. And set `VLLM_PRECOMPILED_WHEEL_LOCATION` as your local absolute path of whl file. + ```bash -# CUDA 13.0 release-branch wheel for the exact source commit on x86_64 -export VLLM_PRECOMPILED_WHEEL_LOCATION=https://wheels.vllm.ai/5df63a7203cdb9328d824a27b682e77c7b9928be/vllm-0.32.0-cp38-abi3-manylinux_2_28_x86_64.whl +# CUDA 13.0 final wheel for the exact source commit on x86_64 +export VLLM_PRECOMPILED_WHEEL_LOCATION=https://wheels.vllm.ai/4ba787ec9959e566d3b587a796dff07673324da5/vllm-0.33.0-cp38-abi3-manylinux_2_28_x86_64.whl ``` + Install vllm with command below (If you have no existing PyTorch). + ```bash VLLM_USE_PRECOMPILED=1 uv pip install --editable . --torch-backend=cu130 ``` + Install vllm with command below (If you already have PyTorch). + ```bash python use_existing_torch.py uv pip install -r requirements/build/cuda.txt uv pip install --no-build-isolation --editable . ``` +
# --8<-- [end:build-wheel-from-source] @@ -127,6 +182,7 @@ tag for a reproducible release environment rather than the rolling `latest` tag. Example deployment on two GPUs (adjust hardware and memory capacity for your model): + ```bash docker run --runtime nvidia --gpus 2 \ -v ~/.cache/huggingface:/root/.cache/huggingface \ @@ -146,17 +202,18 @@ docker run --runtime nvidia --gpus 2 \ #### Build docker image -The production Dockerfile defaults to the published -`vllm/vllm-openai:v0.32.0` CUDA 13 image, which identifies release commit -`5df63a7203cdb9328d824a27b682e77c7b9928be`: +The production Dockerfile bootstraps from the published +`vllm/vllm-openai:v0.33.0` CUDA 13 image and verifies the frozen +vLLM 0.33.0 wheel with its matching dependencies: ```bash DOCKER_BUILDKIT=1 docker build -f docker/Dockerfile.cuda -t vllm-omni-cuda . ``` -The CI Dockerfile pins the final CUDA base by digest and installs the exact -0.32.0 commit-index wheel used by the tested local runtime. Full CI validation -of this final image stack remains pending: +Both CUDA Dockerfiles pin the same bootstrap image by digest and install the +final wheel. The telemetry repair validates that exact base before removing +its unused incompatible exporter; retained exporters and LMCache remain pinned +in the normal dependency solve. Full CI validation remains pending: ```bash DOCKER_BUILDKIT=1 docker build -f docker/Dockerfile.ci -t vllm-omni-cuda-ci . diff --git a/docs/getting_started/installation/gpu/musa.inc.md b/docs/getting_started/installation/gpu/musa.inc.md index a7cbc848f58..15cffc26548 100644 --- a/docs/getting_started/installation/gpu/musa.inc.md +++ b/docs/getting_started/installation/gpu/musa.inc.md @@ -1,8 +1,12 @@ + + + # --8<-- [start:requirements] - GPU: Moore Threads GPU with MUSA SDK installed (validated on MTT S5000) # --8<-- [end:requirements] + # --8<-- [start:set-up-using-python] vLLM-Omni for MUSA requires building from source. Pre-built wheels are not currently available. @@ -40,12 +44,20 @@ cd vllm-omni VLLM_OMNI_TARGET_DEVICE=musa pip install -e . --no-build-isolation ``` -For Gradio demos: +For Gradio demos from this source checkout, retain the MUSA target and +include the temporary Gradio backport in the normal dependency install: ```bash -pip install -e '.[demo]' --no-build-isolation +python tools/build_gradio_pillow13.py --output-dir /tmp/gradio-pillow13 +VLLM_OMNI_TARGET_DEVICE=musa pip install -e '.[demo]' --no-build-isolation \ + /tmp/gradio-pillow13/gradio-6.31.0+omni.pillow13-py3-none-any.whl +python -m pip check ``` +See the [source dependency instructions](https://docs.vllm.ai/projects/vllm-omni/en/latest/getting_started/installation/#source-development-and-demo-dependencies) +for provenance and qualification limits. The CPU serialization checks do not +establish MUSA runtime or browser acceptance. + #### Environment Variables ```bash diff --git a/docs/getting_started/installation/gpu/rocm.inc.md b/docs/getting_started/installation/gpu/rocm.inc.md index 524f1202d32..e4f42765df0 100644 --- a/docs/getting_started/installation/gpu/rocm.inc.md +++ b/docs/getting_started/installation/gpu/rocm.inc.md @@ -11,10 +11,11 @@ For ROCm, vLLM-Omni currently recommends the setup steps through Docker Images. vLLM-Omni depends on the matching major/minor release of vLLM. This source -checkout targets vLLM 0.32.0. The matching ROCm wheel and final 0.32.0 -ROCm base image are published; ROCm compatibility for this source checkout -has not been validated. The pre-built Omni wheel instructions below install -the published 0.30.0 release with its matching vLLM dependency. +checkout targets vLLM 0.33.0. The Dockerfile retains the frozen source build +on the published 0.33.0 base image. The published upstream 0.33.0 ROCm +image and this source-build path still require Omni runtime validation. +The pre-built Omni wheel instructions below retain the published 0.30.0 +release with its matching vLLM dependency. The Dockerfile's `BASE_IMAGE` pin applies only to Docker builds. The `vllm-omni` package does not install vLLM as a dependency, so non-Docker source @@ -52,14 +53,10 @@ uv pip install onnxruntime-rocm #### Installation of vLLM -The published `0.32.0+rocm723` wheel is available from -`https://wheels.vllm.ai/rocm/0.32.0/rocm723` for Linux x86_64 and Python 3.12. -It requires the matching ROCm 7.2.3 PyTorch/Triton builds, rather than the -CUDA packages. Its version file omits the source commit ID, so it has not -been verified as the exact source pin below. The older 0.31.0 ROCm wheel is -not a validated substitute. An experimental source build must use that -exact release commit and upstream's matching ROCm/PyTorch requirements; -both paths still need ROCm runtime validation with this checkout. +Build vLLM from the frozen 0.33.0 source commit using upstream's matching +ROCm/PyTorch requirements, or use the Docker source build below. The older +0.32.0 ROCm wheel does not provide the 0.33 APIs used by this checkout. +ROCm runtime validation remains pending. #### Installation of vLLM-Omni @@ -73,13 +70,26 @@ VLLM_OMNI_TARGET_DEVICE=rocm uv pip install -e . uv pip install -e . --no-build-isolation ``` +For source development or Gradio demos, include the temporary Gradio backport +in the normal install while retaining the ROCm target: + +```bash +python tools/build_gradio_pillow13.py --output-dir /tmp/gradio-pillow13 +VLLM_OMNI_TARGET_DEVICE=rocm uv pip install --no-build-isolation -e '.[dev,demo]' \ + /tmp/gradio-pillow13/gradio-6.31.0+omni.pillow13-py3-none-any.whl +uv pip check +``` + +See the [source dependency instructions](https://docs.vllm.ai/projects/vllm-omni/en/latest/getting_started/installation/#source-development-and-demo-dependencies) +for provenance and qualification limits. ROCm runtime acceptance remains pending. +
(Optional) Installation of vLLM from source If you want to check, modify or debug with source code of vLLM, install the library from source with the following instructions: ```bash git clone https://github.com/vllm-project/vllm.git cd vllm -git checkout 5df63a7203cdb9328d824a27b682e77c7b9928be +git checkout 4ba787ec9959e566d3b587a796dff07673324da5 python3 -m pip install -r requirements/rocm.txt python3 setup.py develop ``` @@ -90,14 +100,11 @@ python3 setup.py develop #### Build docker image -The source build defaults to the published upstream base image -`vllm/vllm-openai-rocm:v0.32.0`, matching this checkout's target release. -The optional vLLM source rebuild also defaults to `v0.32.0`. ROCm runtime -validation with this checkout remains pending. -Older or custom bases must provide the vLLM APIs checked by the Dockerfile's -image-build canary. This upstream base is distinct from the prebuilt -`vllm/vllm-omni-rocm` images discussed below; their published-tag availability -is tracked in [#7405](https://github.com/vllm-project/vllm-omni/issues/7405). +The Dockerfile starts from `vllm/vllm-openai-rocm:v0.33.0` and defaults to +`USE_NIGHTLY_BUILD=1` with source commit `4ba787ec9959e566d3b587a796dff07673324da5`. Its development dependency +install builds the hash-pinned temporary Gradio source backport and installs +the resulting local wheel with the normal resolver. ROCm runtime validation +remains pending. ```bash DOCKER_BUILDKIT=1 docker build -f docker/Dockerfile.rocm -t vllm-omni-rocm . @@ -108,7 +115,7 @@ To select the upstream ROCm base explicitly, pass `BASE_IMAGE`: ```bash DOCKER_BUILDKIT=1 docker build \ -f docker/Dockerfile.rocm \ - --build-arg BASE_IMAGE=vllm/vllm-openai-rocm:v0.32.0 \ + --build-arg BASE_IMAGE=vllm/vllm-openai-rocm:v0.33.0 \ -t vllm-omni-rocm . ``` diff --git a/docs/getting_started/installation/gpu/xpu.inc.md b/docs/getting_started/installation/gpu/xpu.inc.md index 18201f81adf..335760c42e5 100644 --- a/docs/getting_started/installation/gpu/xpu.inc.md +++ b/docs/getting_started/installation/gpu/xpu.inc.md @@ -16,6 +16,22 @@ vLLM-Omni currently recommends using the Docker image setup steps below. # --8<-- [start:build-wheel-from-source] +For development or Gradio demos in a source environment with the matching +upstream XPU runtime already installed, run from the Omni repository root: + +```bash +python tools/build_gradio_pillow13.py --output-dir /tmp/gradio-pillow13 +VLLM_OMNI_TARGET_DEVICE=xpu uv pip install --no-build-isolation \ + --excludes <(printf 'torch\ntriton-xpu\n') -e '.[dev,demo]' \ + /tmp/gradio-pillow13/gradio-6.31.0+omni.pillow13-py3-none-any.whl +uv pip check +``` + +Keep the upstream PyTorch and Triton packages paired. See the +[source dependency instructions](https://docs.vllm.ai/projects/vllm-omni/en/latest/getting_started/installation/#source-development-and-demo-dependencies) +for the temporary Gradio backport's provenance and qualification limits. +XPU runtime acceptance remains pending. + # --8<-- [end:build-wheel-from-source] # --8<-- [start:build-docker] @@ -23,15 +39,15 @@ vLLM-Omni currently recommends using the Docker image setup steps below. #### Build docker image ```bash -export VLLM_VERSION=v0.32.0 +export VLLM_VERSION=v0.33.0 DOCKER_BUILDKIT=1 docker build -f docker/Dockerfile.xpu -t vllm-omni-xpu --shm-size=4g \ --build-arg VLLM_VERSION="$VLLM_VERSION" . ``` -This layers vLLM-Omni on top of the published `vllm/vllm-openai-xpu:` -base image, which Docker pulls automatically. The Dockerfile defaults to -v0.32.0. Its matching vLLM wheel requires PyTorch 2.14.0 and -`triton==3.8.0+xpu`. The Dockerfile retains the upstream runtime instead of +This layers vLLM-Omni on top of the selected `vllm/vllm-openai-xpu:` +base image. The Dockerfile targets the published v0.33.0 tag; if the image +cannot be fetched, build the upstream source image at that tag as shown below. Keep its +PyTorch and Triton packages paired. The Dockerfile retains the upstream runtime instead of replacing Triton with the older `triton-xpu` package, and checks installed dependency requirements after adding Omni. XPU runtime validation with this checkout remains pending. To target a different vLLM release, pass diff --git a/docs/getting_started/installation/npu/npu.inc.md b/docs/getting_started/installation/npu/npu.inc.md index 5c3b18b71eb..7d090c65a27 100644 --- a/docs/getting_started/installation/npu/npu.inc.md +++ b/docs/getting_started/installation/npu/npu.inc.md @@ -1,8 +1,11 @@ + + # --8<-- [start:requirements] For detailed hardware and software requirements, please refer to the [vLLM-Ascend installation documentation](https://docs.vllm.ai/projects/ascend/en/latest/installation.html). # --8<-- [end:requirements] + # --8<-- [start:installation-release] The recommended way to use vLLM-Omni on NPU is through the vLLM-Ascend pre-built Docker images: @@ -72,9 +75,9 @@ We are keeping [issue #886](https://github.com/vllm-project/vllm-omni/issues/886 # --8<-- [start:installation-main] You can also install vLLM-Omni from source if you want to use the latest features -or bug fixes. This alignment checkout targets vLLM 0.32.0 and requires a +or bug fixes. This alignment checkout targets vLLM 0.33.0 and requires a matching vLLM-Ascend runtime validated on your NPU. The v0.30.0 Ascend images -above do not establish compatibility with this vLLM 0.32.0 checkout. Check +above do not establish compatibility with this vLLM 0.33.0 checkout. Check [issue #886](https://github.com/vllm-project/vllm-omni/issues/886) for aligned versions and NPU validation before selecting a development image. @@ -97,7 +100,7 @@ vLLM-Omni offers Docker images for Ascend NPU deployment. You can just pull the Supported images as following. | image name | Hardware | OS | -|-|-|-| +| - | - | - | | image-tag | A2 | Ubuntu | | image-tag-a3 | A3 | Ubuntu | | image-tag-a5 | A5 | Ubuntu | @@ -211,6 +214,15 @@ git clone https://github.com/vllm-project/vllm-omni.git cd vllm-omni ``` +!!! warning "vLLM 0.33 alignment is not yet qualified on NPU" + Both `docker/Dockerfile.npu` and `docker/Dockerfile.npu.ci` retain the + legacy `v0.32.0` Ascend runtime default. The commands below illustrate + hardware suffix selection; they do not produce a validated vLLM 0.33 + image. Before building this checkout, obtain a matching vLLM 0.33.0 and + vLLM-Ascend runtime, pass its verified image repository and tag through + `VLLM_ASCEND_IMAGE` and `VLLM_ASCEND_VERSION`, and validate on your NPU. + A final vLLM package alone does not establish vLLM-Ascend compatibility. + All hardware variants use `docker/Dockerfile.npu`. Choose the device suffix for your hardware: === "A2" diff --git a/docs/getting_started/quickstart.md b/docs/getting_started/quickstart.md index 218c401fb5c..1991a1c15f5 100644 --- a/docs/getting_started/quickstart.md +++ b/docs/getting_started/quickstart.md @@ -20,22 +20,23 @@ source .venv/bin/activate # On CUDA (Linux x86_64, CUDA 13.0) uv pip install --torch-backend=cu130 \ - "vllm @ https://wheels.vllm.ai/5df63a7203cdb9328d824a27b682e77c7b9928be/vllm-0.32.0-cp38-abi3-manylinux_2_28_x86_64.whl" + "vllm @ https://wheels.vllm.ai/4ba787ec9959e566d3b587a796dff07673324da5/vllm-0.33.0-cp38-abi3-manylinux_2_28_x86_64.whl" git clone https://github.com/vllm-project/vllm-omni.git cd vllm-omni uv pip install -e . ``` -These commands install the vLLM 0.32.0 release-branch wheel at the exact -tested tag commit, with PyTorch 2.13.0 and CUDA 13.0. The -[vLLM 0.32.0 release](https://github.com/vllm-project/vllm/releases/tag/v0.32.0), -PyPI wheel and matching ROCm wheel are published. ROCm runtime validation -with this checkout remains pending. A vLLM-Omni 0.32 wheel is not yet -published; install Omni from source as shown above. For published -vLLM-Omni wheels and additional installation methods, see the +These commands select vLLM 0.33.0 at the frozen release-branch commit, +with PyTorch 2.14.0 and CUDA 13.0. The final package is available; native model, +full CI and ROCm/XPU runtime validation remain pending. Install Omni from this source +checkout. For published Omni wheels and other installation methods, see the [installation guide](installation/README.md). +For source development or Gradio demos, also follow the +[source dependency instructions](installation/README.md#source-development-and-demo-dependencies) +to install the temporary Gradio wheel together with the requested extras. + !!! note It is important to install the same major & minor version of vLLM and vLLM Omni, otherwise things may not work as expected. If the versions are misaligned, you will see a warning when you import vLLM Omni. diff --git a/docs/serving/realtime_duplex_api.md b/docs/serving/realtime_duplex_api.md index eb8c29c0a4f..b4cc059614b 100644 --- a/docs/serving/realtime_duplex_api.md +++ b/docs/serving/realtime_duplex_api.md @@ -518,9 +518,12 @@ record rejection and retry counts when measuring service capacity. PersonaPlex startup warmup is opt-in: set `duplex_session.warmup_frames` to a positive frame count to enable it. Frontends with nonzero admission budgets -then warm their owned groups locally and gate their own duplex connections -until warmup completes; clients wait at most `DUPLEX_WARMUP_CLIENT_WAIT_S` (180 s) -before being admitted, as in single-frontend mode. +then warm their owned groups locally. Native duplex endpoints accept the +WebSocket transport immediately, then hold session configuration and engine +admission until warmup completes; clients wait at most +`DUPLEX_WARMUP_CLIENT_WAIT_S` (180 s) before session admission, as in +single-frontend mode. The warmup wait does not consume the WebSocket opening +handshake deadline. `/health` retains the engine-health check; it is not an all-frontend warmup barrier, and a 200 response does not mean that the accepting frontend has free admission slots. diff --git a/docs/user_guide/diffusion/attention_backends/dense_backends.md b/docs/user_guide/diffusion/attention_backends/dense_backends.md index 2cb7019bedb..a637ea33a42 100644 --- a/docs/user_guide/diffusion/attention_backends/dense_backends.md +++ b/docs/user_guide/diffusion/attention_backends/dense_backends.md @@ -18,6 +18,16 @@ conservative reference when validating another backend. vllm-omni serve --diffusion-attention-backend TORCH_SDPA ``` +For Qwen-Image transformers, CUDA BF16 with `--enforce-eager` and the resolved +`TORCH_SDPA` backend also uses the Diffusers normalization and FP32 complex +rotary rounding order. Its dense text QKV projections use three GEMMs while +retaining packed checkpoint weights, tensor-parallel sharding, and LoRA routing. +Quantized projection methods and an explicitly configured batch-invariant GEMM +policy keep their existing dispatch. This reference path uses more kernel +launches than the fused path; benchmark it separately when choosing a backend. +Compiled execution, other attention backends, CPU, FP16, and ROCm retain the +existing fused or activation-dtype rotary paths. + ## `FLASH_ATTN` `FLASH_ATTN` uses the installed FlashAttention implementation. On Blackwell it @@ -67,6 +77,10 @@ for mask-heavy DiTs and is automatically preferred on Blackwell when cuDNN 9.5 or newer is available and the higher-priority TRTLLM route is not compatible. +For this checkout's vLLM 0.33 / Torch 2.14 CUDA environment, use the pinned +cuDNN 9.25.1 runtime in the [CUDA source installation instructions](../../../getting_started/installation/gpu.md). +The CUDA Docker and CI images configure that runtime at process startup. + ```bash vllm-omni serve --diffusion-attention-backend CUDNN_ATTN ``` diff --git a/docs/user_guide/diffusion/attention_backends/huggingface_hub.md b/docs/user_guide/diffusion/attention_backends/huggingface_hub.md index cfd220e829c..b97c2b592fc 100644 --- a/docs/user_guide/diffusion/attention_backends/huggingface_hub.md +++ b/docs/user_guide/diffusion/attention_backends/huggingface_hub.md @@ -19,6 +19,11 @@ If `kernels` is unavailable, an explicit `FLASH_ATTN_HUB` or `FLASH_ATTN_3_HUB` selection raises. Automatic / `FLASH_ATTN` selection in the Diffusers adapter may still walk hub then local FlashAttention backends. +For FlashAttention 2, the loader tries Hub API versions 1, 2, then 3. +Version 3 provides Torch stable-ABI CUDA builds compatible with the Torch 2.14 +runtime used by vLLM 0.33. An explicit Hub selection raises if none of these +versions has a compatible build. + ## Usage ```bash diff --git a/docs/user_guide/examples/offline_inference/text_to_image.md b/docs/user_guide/examples/offline_inference/text_to_image.md index 24cab3a3c3d..056bd95efd0 100644 --- a/docs/user_guide/examples/offline_inference/text_to_image.md +++ b/docs/user_guide/examples/offline_inference/text_to_image.md @@ -2,7 +2,6 @@ Source . - Generate images from text prompts using vLLM-Omni's diffusion pipeline entrypoints. - `text_to_image.py`: command-line script for single image generation with advanced options. @@ -22,15 +21,15 @@ This folder provides several entrypoints for experimenting with text-to-image di ### Supported Models -| Model | Image Shape | Peak VRAM (GiB) * | Model Weights (GiB) | +| Model | Image Shape | Peak VRAM (GiB) * | Model Weights (GiB) | | ----- | ----------- | ----------- | ----------------- | | `Qwen/Qwen-Image` | 1024 x 1024 | 60.0 | 53.7 | -| `Qwen/Qwen-Image-2512` |1024 x 1024 | 60.0 | 53.7 | +| `Qwen/Qwen-Image-2512` | 1024 x 1024 | 60.0 | 53.7 | | `Tongyi-MAI/Z-Image-Turbo` | 1024 x 1024 | 24.8 | 19.2 | | `stepfun-ai/NextStep-1.1` | 512 x 512 | 71.8 | 28.1 | | `meituan-longcat/LongCat-Image` | 1024 x 1024 | 71.2 | 27.3 | | `AIDC-AI/Ovis-Image-7B` | 1024 x 1024 | 71.8 | 17.1 | -| `OmniGen2/OmniGen2` | 1024 x 1024 | 20.1 | 14.7 | +| `OmniGen2/OmniGen2` | 1024 x 1024 | 20.1 | 14.7 | | `stabilityai/stable-diffusion-3.5-medium` | 1024 x 1024 | 20.1 | 15.6 | | `black-forest-labs/FLUX.1-dev` | 1024 x 1024 | 77.6 | 31.4 | | `black-forest-labs/FLUX.2-klein-4B` | 1024 x 1024 | 72.7 | 14.9 | @@ -90,7 +89,7 @@ python text_to_image.py \ | `--enable-cpu-offload` | flag | off | Enable CPU offloading for diffusion models | | `--lora-path` | str | — | Path to PEFT LoRA adapter folder or checkpoint file | | `--lora-scale` | float | `1.0` | Scale factor for LoRA weights | -| `--lora-backend` | str |`"peft"`| LoRA backend for loading LoRA adapters. Default: peft. Available options: peft, distill | +| `--lora-backend` | str | `"peft"` | LoRA backend for loading LoRA adapters. Default: peft. Available options: peft, distill | **NextStep-1.1 specific arguments:** @@ -102,7 +101,7 @@ python text_to_image.py \ | `--use-norm` | flag | off | Apply layer normalization to sampled tokens | > If you encounter OOM errors, try using `--vae-use-slicing` and `--vae-use-tiling` to reduce memory usage. - +> > Qwen-Image currently publishes best-effort presets at `1328x1328`, `1664x928`, `928x1664`, `1472x1140`, `1140x1472`, `1584x1056`, and `1056x1584`. Adjust `--height/--width` accordingly for the most reliable outcomes. ## More CLI Examples @@ -242,7 +241,7 @@ python text_to_image.py \ LoRA adapters must be in PEFT format. A typical adapter directory structure: -``` +```text lora_adapter/ ├── adapter_config.json └── adapter_model.safetensors @@ -257,7 +256,10 @@ lora_adapter/ pip install 'vllm-omni[demo]' ``` - Or, if installing from source: `pip install -e '.[demo]'` + For this 0.33 source checkout, follow the + [source demo dependency instructions](https://docs.vllm.ai/projects/vllm-omni/en/latest/getting_started/installation/#source-development-and-demo-dependencies) + from the repository root. Build the temporary Gradio wheel and include it + with `.[demo]` in the same normal install command. Launch the Gradio demo: diff --git a/docs/user_guide/examples/offline_inference/text_to_speech.md b/docs/user_guide/examples/offline_inference/text_to_speech.md index 940a6b05ca4..f748fca472c 100644 --- a/docs/user_guide/examples/offline_inference/text_to_speech.md +++ b/docs/user_guide/examples/offline_inference/text_to_speech.md @@ -46,12 +46,23 @@ Single-stage native AR TTS at 48 kHz. Pipeline: `feat_encoder → MiniCPM4 → F ### Prerequisites +The native model requires `voxcpm==2.0.3`, including its Gradio dependency. +After the matching vLLM and Omni source installation, run from the repository +root using the [source dependency instructions](https://docs.vllm.ai/projects/vllm-omni/en/latest/getting_started/installation/#source-development-and-demo-dependencies): + ```bash -pip install voxcpm -# or, for a local source checkout: -pip install -e /path/to/voxcpm +python tools/build_gradio_pillow13.py --output-dir /tmp/gradio-pillow13 +uv pip install 'voxcpm==2.0.3' \ + /tmp/gradio-pillow13/gradio-6.31.0+omni.pillow13-py3-none-any.whl +# Alternatively, add an editable VoxCPM checkout with matching 2.0.3 metadata: +uv pip install 'voxcpm==2.0.3' -e /path/to/voxcpm \ + /tmp/gradio-pillow13/gradio-6.31.0+omni.pillow13-py3-none-any.whl +uv pip check ``` +The selected CPU dependency/serialization checks do not establish native +VoxCPM inference or output acceptance. + ### Quick start ```bash diff --git a/docs/user_guide/examples/online_serving/image_to_image.md b/docs/user_guide/examples/online_serving/image_to_image.md index cc235754744..43ab5ceb8fc 100644 --- a/docs/user_guide/examples/online_serving/image_to_image.md +++ b/docs/user_guide/examples/online_serving/image_to_image.md @@ -2,7 +2,6 @@ Source . - This example demonstrates how to deploy image-to-image models for online image editing service using vLLM-Omni. Supported models include Qwen-Image-Edit, BAGEL, JoyAI-Image-Edit, and other image-to-image pipelines. @@ -36,7 +35,6 @@ vllm serve jdopensource/JoyAI-Image-Edit-Diffusers --omni --port 8092 ### Start with Parameters - Or use the startup script: ```bash @@ -144,7 +142,10 @@ python openai_chat_client.py --input input1.png input2.png --prompt "Combine the pip install 'vllm-omni[demo]' ``` - Or, if installing from source: `pip install -e '.[demo]'` + For this 0.33 source checkout, follow the + [source demo dependency instructions](https://docs.vllm.ai/projects/vllm-omni/en/latest/getting_started/installation/#source-development-and-demo-dependencies) + from the repository root. Build the temporary Gradio wheel and include it + with `.[demo]` in the same normal install command. ```bash python gradio_demo.py @@ -346,7 +347,7 @@ The response contains multiple images in `choices[0].message.content` — one pe #### Qwen-Image-Layered Parameters | Parameter | Type | Default | Description | -|-----------|------|---------|-------------| +| ----------- | ------ | --------- | ------------- | | `layers` | int | 4 | Number of layers to decompose | | `resolution` | int | 640 | Resolution for dimension calculation (640 or 1024) | | `cfg_scale` | float | 4.0 | Classifier-free guidance scale (alias for `true_cfg_scale`) | @@ -382,17 +383,17 @@ endpoint, pass the supported generation controls as top-level form fields directly. For image dimensions and count, use `size` and `n` rather than `height`, `width`, or `num_outputs_per_prompt`. -| Parameter | Type | Default | Description | -| ------------------------ | ----- | ------- | ------------------------------------- | -| `height` | int | None | Output image height in pixels | -| `width` | int | None | Output image width in pixels | -| `size` | str | None | Output image size (e.g., "1024x1024") | -| `num_inference_steps` | int | 50 | Number of denoising steps | -| `guidance_scale` | float | 1.0 | CFG guidance scale | -| `seed` | int | None | Random seed (reproducible) | -| `negative_prompt` | str | None | Negative prompt | -| `num_outputs_per_prompt` | int | 1 | Number of images to generate | -| `layers` | int | 4 | Number of layers (Qwen-Image-Layered) | +| Parameter | Type | Default | Description | +| ------------------------ | ----- | ------- | -------------------------------------------- | +| `height` | int | None | Output image height in pixels | +| `width` | int | None | Output image width in pixels | +| `size` | str | None | Output image size (e.g., "1024x1024") | +| `num_inference_steps` | int | 50 | Number of denoising steps | +| `guidance_scale` | float | 1.0 | CFG guidance scale | +| `seed` | int | None | Random seed (reproducible) | +| `negative_prompt` | str | None | Negative prompt | +| `num_outputs_per_prompt` | int | 1 | Number of images to generate | +| `layers` | int | 4 | Number of layers (Qwen-Image-Layered) | | `resolution` | int | 640 | Resolution, 640 or 1024 (Qwen-Image-Layered) | ## Response Format diff --git a/docs/user_guide/examples/online_serving/qwen2_5_omni.md b/docs/user_guide/examples/online_serving/qwen2_5_omni.md index e6d3f252f92..222497fac2c 100644 --- a/docs/user_guide/examples/online_serving/qwen2_5_omni.md +++ b/docs/user_guide/examples/online_serving/qwen2_5_omni.md @@ -2,7 +2,6 @@ Source . - ## 🛠️ Installation Please refer to [README.md](https://github.com/vllm-project/vllm-omni/tree/main/README.md) @@ -16,12 +15,13 @@ vllm serve Qwen/Qwen2.5-Omni-7B --omni --port 8091 ``` To use a custom deploy config, launch the server with the command below: + ```bash vllm serve Qwen/Qwen2.5-Omni-7B --omni --port 8091 --deploy-config /path/to/deploy_config.yaml ``` Omni keeps soundfile as the first automatic audio decoder to preserve input -waveforms when upgrading to vLLM 0.32. Switching decoders can change BF16 audio +waveforms when upgrading to vLLM 0.33. Switching decoders can change BF16 audio features and mixed-media answers. Formats unsupported by soundfile continue through upstream's torchcodec/PyAV fallback. To explicitly select a decoder, use `--media-io-kwargs '{"audio":{"audio_backend":"torchcodec"}}'` (or @@ -30,6 +30,7 @@ use `--media-io-kwargs '{"audio":{"audio_backend":"torchcodec"}}'` (or ### Send Multi-modal Request Get into the example folder + ```bash cd examples/online_serving/qwen2_5_omni ``` @@ -48,7 +49,6 @@ The Python client supports the following command-line arguments: - `--audio-path` (or `-a`): Path to local audio file or URL. If not provided and query-type uses audio, uses default audio URL. Supports local file paths (automatically encoded to base64) or HTTP/HTTPS URLs and common audio formats: MP3, WAV, OGG, FLAC, M4A. Example: `--audio-path /path/to/audio.wav` or `--audio-path https://example.com/audio.mp3` - `--prompt` (or `-p`): Custom text prompt/question. If not provided, uses default prompt for the selected query type. Example: `--prompt "What are the main activities shown in this video?"` - For example, to use mixed modalities with all local files: ```bash @@ -60,19 +60,20 @@ python openai_chat_completion_client_for_multimodal_generation.py \ --prompt "Analyze all the media content and provide a comprehensive summary." ``` -#### Send request via curl +#### Send request via curl ```bash bash run_curl_multimodal_generation.sh mixed_modalities ``` ## Modality control + You can control output modalities to specify which types of output the model should generate. This is useful when you only need text output and want to skip audio generation stages for better performance. ### Supported modalities | Modalities | Output | -|------------|--------| +| ------------ | -------- | | `["text"]` | Text only | | `["audio"]` | Text + Audio | | `["text", "audio"]` | Text + Audio | @@ -147,7 +148,9 @@ print(response.choices[1].message.audio) # Audio response ``` ## Streaming Output + If you want to enable streaming output, please set the argument as below. The final output will be obtained just after generated by corresponding stage. Now we only support text streaming output. Other modalities can output normally. + ```bash python openai_chat_completion_client_for_multimodal_generation.py \ --query-type mixed_modalities \ @@ -163,7 +166,10 @@ python openai_chat_completion_client_for_multimodal_generation.py \ pip install 'vllm-omni[demo]' ``` - Or, if installing from source: `pip install -e '.[demo]'` + For this 0.33 source checkout, follow the + [source demo dependency instructions](https://docs.vllm.ai/projects/vllm-omni/en/latest/getting_started/installation/#source-development-and-demo-dependencies) + from the repository root. Build the temporary Gradio wheel and include it + with `.[demo]` in the same normal install command. This Web UI demo allows users to interact with the model through a web browser. @@ -180,12 +186,14 @@ The convenience script launches both the vLLM server and Gradio demo together: ``` This script will: + 1. Start the vLLM server in the background 2. Wait for the server to be ready 3. Launch the Gradio demo 4. Handle cleanup when you press Ctrl+C The script supports the following arguments: + - `--model`: Model name/path (default: Qwen/Qwen2.5-Omni-7B) - `--server-port`: Port for vLLM server (default: 8091) - `--gradio-port`: Port for Gradio demo (default: 7861) @@ -196,18 +204,19 @@ The script supports the following arguments: #### Option 2: Manual Launch (Two-Step Process) -**Step 1: Launch the vLLM API server** +##### Step 1: Launch the vLLM API server ```bash vllm serve Qwen/Qwen2.5-Omni-7B --omni --port 8091 ``` To use a custom deploy config: + ```bash vllm serve Qwen/Qwen2.5-Omni-7B --omni --port 8091 --deploy-config /path/to/deploy_config.yaml ``` -**Step 2: Run the Gradio demo** +##### Step 2: Run the Gradio demo In a separate terminal: @@ -220,7 +229,7 @@ Then open `http://localhost:7861/` on your local browser to interact with the we The gradio script supports the following arguments: - `--model`: Model name/path (should match the server model) -- `--api-base`: Base URL for the vLLM API server (default: http://localhost:8091/v1) +- `--api-base`: Base URL for the vLLM API server (default: ) - `--ip`: Host/IP for Gradio server (default: 127.0.0.1) - `--port`: Port for Gradio server (default: 7861) - `--share`: Share the Gradio demo publicly (creates a public link) diff --git a/docs/user_guide/examples/online_serving/qwen3_omni.md b/docs/user_guide/examples/online_serving/qwen3_omni.md index 3c56422ecd9..1bdaecab2c5 100644 --- a/docs/user_guide/examples/online_serving/qwen3_omni.md +++ b/docs/user_guide/examples/online_serving/qwen3_omni.md @@ -311,7 +311,10 @@ python openai_chat_completion_client_for_multimodal_generation.py \ pip install 'vllm-omni[demo]' ``` - Or, if installing from source: `pip install -e '.[demo]'` + For this 0.33 source checkout, follow the + [source demo dependency instructions](https://docs.vllm.ai/projects/vllm-omni/en/latest/getting_started/installation/#source-development-and-demo-dependencies) + from the repository root. Build the temporary Gradio wheel and include it + with `.[demo]` in the same normal install command. This Web UI demo allows users to interact with the model through a web browser. diff --git a/docs/user_guide/examples/online_serving/text_to_image.md b/docs/user_guide/examples/online_serving/text_to_image.md index 43fb965948e..5286467adc8 100644 --- a/docs/user_guide/examples/online_serving/text_to_image.md +++ b/docs/user_guide/examples/online_serving/text_to_image.md @@ -2,7 +2,6 @@ Source . - This example demonstrates how to deploy Qwen-Image model for online image generation service using vLLM-Omni. ## Start Server @@ -12,6 +11,7 @@ This example demonstrates how to deploy Qwen-Image model for online image genera ```bash vllm serve Qwen/Qwen-Image --omni --port 8091 ``` + !!! note If you encounter Out-of-Memory (OOM) issues or have limited GPU memory, you can enable VAE slicing and tiling to reduce memory usage, --vae-use-slicing --vae-use-tiling @@ -112,7 +112,10 @@ python openai_chat_client.py --prompt "A beautiful landscape painting" --output pip install 'vllm-omni[demo]' ``` - Or, if installing from source: `pip install -e '.[demo]'` + For this 0.33 source checkout, follow the + [source demo dependency instructions](https://docs.vllm.ai/projects/vllm-omni/en/latest/getting_started/installation/#source-development-and-demo-dependencies) + from the repository root. Build the temporary Gradio wheel and include it + with `.[demo]` in the same normal install command. ```bash python gradio_demo.py @@ -166,7 +169,7 @@ curl -X POST http://localhost:8091/v1/images/generations \ LoRA adapters must be in PEFT (Parameter-Efficient Fine-Tuning) format. A typical LoRA adapter directory structure: -``` +```text lora_adapter/ ├── adapter_config.json └── adapter_model.safetensors diff --git a/examples/offline_inference/text_to_speech/README.md b/examples/offline_inference/text_to_speech/README.md index d0162f11533..4f43c199879 100644 --- a/examples/offline_inference/text_to_speech/README.md +++ b/examples/offline_inference/text_to_speech/README.md @@ -638,12 +638,23 @@ Single-stage native AR TTS at 48 kHz. Pipeline: `feat_encoder → MiniCPM4 → F ### Prerequisites +The native model requires `voxcpm==2.0.3`, including its Gradio dependency. +After the matching vLLM and Omni source installation, run from the repository +root using the [source dependency instructions](https://docs.vllm.ai/projects/vllm-omni/en/latest/getting_started/installation/#source-development-and-demo-dependencies): + ```bash -pip install voxcpm -# or, for a local source checkout: -pip install -e /path/to/voxcpm +python tools/build_gradio_pillow13.py --output-dir /tmp/gradio-pillow13 +uv pip install 'voxcpm==2.0.3' \ + /tmp/gradio-pillow13/gradio-6.31.0+omni.pillow13-py3-none-any.whl +# Alternatively, add an editable VoxCPM checkout with matching 2.0.3 metadata: +uv pip install 'voxcpm==2.0.3' -e /path/to/voxcpm \ + /tmp/gradio-pillow13/gradio-6.31.0+omni.pillow13-py3-none-any.whl +uv pip check ``` +The selected CPU dependency/serialization checks do not establish native +VoxCPM inference or output acceptance. + ### Quick start ```bash diff --git a/examples/online_serving/streaming_video_generation/README.md b/examples/online_serving/streaming_video_generation/README.md index 16021023069..d1ec941cfc9 100644 --- a/examples/online_serving/streaming_video_generation/README.md +++ b/examples/online_serving/streaming_video_generation/README.md @@ -63,6 +63,12 @@ pip install av websockets pip install 'vllm-omni[demo]' websockets ``` +For this 0.33 source checkout, install `.[demo]` with the temporary Gradio wheel +using the [source dependency instructions](https://docs.vllm.ai/projects/vllm-omni/en/latest/getting_started/installation/#source-development-and-demo-dependencies) +from the repository root, while retaining the `av` and `websockets` client +dependencies above. The published-package command above applies to released +versions and does not replace the source setup. + ## Run The Client ```bash diff --git a/examples/online_serving/text_to_speech/README.md b/examples/online_serving/text_to_speech/README.md index ee56976ead8..87e92b588a4 100644 --- a/examples/online_serving/text_to_speech/README.md +++ b/examples/online_serving/text_to_speech/README.md @@ -1034,6 +1034,11 @@ reports the last-spoken word, useful for the voice-agent barge-in case. Single-stage native AR TTS at 48 kHz. +Install the real `voxcpm==2.0.3` runtime and its temporary Gradio backport +using the [offline VoxCPM2 prerequisites](../../offline_inference/text_to_speech/README.md#voxcpm2) +before launching the server. Those dependency checks do not establish native +model or browser acceptance. + ### Launch ```bash diff --git a/recipes/MOSS-Image/MOSS-Image-Base.md b/recipes/MOSS-Image/MOSS-Image-Base.md index 79e98d66334..87ae6964f81 100644 --- a/recipes/MOSS-Image/MOSS-Image-Base.md +++ b/recipes/MOSS-Image/MOSS-Image-Base.md @@ -31,7 +31,7 @@ The eager CUDA path fuses RMSNorm/AdaLN, Q/K RMSNorm with RoPE, SiTU and gated r ## Environment -Linux, Python 3.12, PyTorch 2.13.0+cu130, vLLM 0.31.0, Transformers 5.14.1 and Diffusers 0.40.0. See the [installation guide](../../docs/getting_started/installation/README.md). +Historical reference environment: Linux, Python 3.12, PyTorch 2.13.0+cu130, vLLM 0.31.0, Transformers 5.14.1 and Diffusers 0.40.0. For this source checkout, use final vLLM 0.33.0 and its matching dependencies in the [installation guide](../../docs/getting_started/installation/README.md). The historical kernel and model results do not establish qualification for the new runtime. ## Deploy diff --git a/recipes/MiniMaxAI/MiniMax-H3.md b/recipes/MiniMaxAI/MiniMax-H3.md index e38899cc9ab..72d4f372ccb 100644 --- a/recipes/MiniMaxAI/MiniMax-H3.md +++ b/recipes/MiniMaxAI/MiniMax-H3.md @@ -622,16 +622,18 @@ Select the device with `HIP_VISIBLE_DEVICES` (not `CUDA_VISIBLE_DEVICES`), drop CUDA-only `FLASHINFER_DISABLE_VERSION_CHECK`, and install without the `[fa4]` extra (FA4 is CUDA-only). The VAE uses AITER GroupNorm on ROCm. -Install (ROCm wheel + source vLLM-Omni): +For this source checkout, build the frozen vLLM 0.33.0 ROCm runtime from +source using the repository Dockerfile, from the Omni repository root: ```bash -pip install "vllm==0.26.0+rocm723" \ - --extra-index-url https://wheels.vllm.ai/rocm/0.26.0/rocm723 -VLLM_OMNI_TARGET_DEVICE=rocm pip install -e . --no-build-isolation +docker build -f docker/Dockerfile.rocm -t vllm-omni-rocm:v033-source . ``` -Prebuilt image: `vllm/vllm-omni-rocm:minimax-h3`. All tasks work out of the box: -the image bundles TorchCodec (for image+audio Ref2VA) and `ffmpeg` (for +See the [ROCm installation guide](../../docs/getting_started/installation/gpu/rocm.inc.md) +for source/runtime details. Container and native H3 validation on ROCm remain +pending for this checkout; an older ROCm wheel cannot provide its 0.33 APIs. + +Historical prebuilt image: `vllm/vllm-omni-rocm:minimax-h3`. That image bundles TorchCodec (for image+audio Ref2VA) and `ffmpeg` (for video-reference Ref2VA). ### ROCm single GPU diff --git a/recipes/README.md b/recipes/README.md index 05b7d243554..eb6615e5a2e 100644 --- a/recipes/README.md +++ b/recipes/README.md @@ -25,6 +25,14 @@ published recipe when available and otherwise to the corresponding recipe in this directory. The recipe files remain the source of truth for commands and validation details. +This source checkout targets final vLLM 0.33.0. Use the matching runtime in +[the installation guide](../docs/getting_started/installation/README.md). +Recipes explicitly pinned to an older Omni release, commit, external PR or +container retain their historical runtime. Their measured hardware and quality +results do not establish acceptance for this checkout. Source installs must not +combine current Omni with an older vLLM wheel; ROCm, XPU and NPU qualification +limits are recorded in the installation guide. + Example layout: ```text diff --git a/recipes/inclusionAI/Ming-3.0-audio-flash.md b/recipes/inclusionAI/Ming-3.0-audio-flash.md index e8df8c4fb1e..bb75864b766 100644 --- a/recipes/inclusionAI/Ming-3.0-audio-flash.md +++ b/recipes/inclusionAI/Ming-3.0-audio-flash.md @@ -25,7 +25,7 @@ ASR and Speech QA share the Ming3 audio thinker; requests for audio output addit ### Environment -- Linux, Python 3.11–3.13, and vLLM-Omni matching vLLM 0.32.0. +- Linux, Python 3.11–3.13, and vLLM-Omni matching final vLLM 0.33.0. - Client examples use Bash, `curl`, `jq`, and GNU `base64`. - Run commands from the repository root. Speech output also requires the [text normalization resources](#text-normalization-resources). diff --git a/recipes/models/MiniMaxAI/MiniMax-H3.yaml b/recipes/models/MiniMaxAI/MiniMax-H3.yaml index 4cc68d878a7..f7f7ecd8ae7 100644 --- a/recipes/models/MiniMaxAI/MiniMax-H3.yaml +++ b/recipes/models/MiniMaxAI/MiniMax-H3.yaml @@ -42,7 +42,7 @@ model: command: | uv venv source .venv/bin/activate - uv pip install vllm==0.26.0 + uv pip install --torch-backend=cu130 vllm==0.33.0 note: "H3 also needs current vLLM-Omni main for the modular FL2VA/Ref2VA pipeline — see the extra install block below." architecture: dense parameter_count: "64B" @@ -900,7 +900,7 @@ guide: | ```bash uv venv source .venv/bin/activate - uv pip install vllm==0.26.0 + uv pip install --torch-backend=cu130 vllm==0.33.0 git clone --depth 1 https://github.com/vllm-project/vllm-omni.git /path/to/vllm-omni cd /path/to/vllm-omni uv pip install -e . @@ -1094,6 +1094,11 @@ guide: | models currently need the source branch from [vLLM-Omni #9167](https://github.com/vllm-project/vllm-omni/pull/9167): + This external fork's instructions retain its vLLM 0.32.0 runtime. They do + not install or qualify this checkout's final vLLM 0.33.0 alignment; follow + the [source installation guide](https://docs.vllm.ai/projects/vllm-omni/en/latest/getting_started/installation/gpu/#build-wheel-from-source) + for this checkout. FastH3 on the 0.33.0 runtime remains unverified. + ```bash sudo apt-get update sudo apt-get install -y git ffmpeg diff --git a/recipes/models/Qwen/Qwen-Image-2.1.yaml b/recipes/models/Qwen/Qwen-Image-2.1.yaml index 9bfd0af689d..a5ee49e6adb 100644 --- a/recipes/models/Qwen/Qwen-Image-2.1.yaml +++ b/recipes/models/Qwen/Qwen-Image-2.1.yaml @@ -156,18 +156,18 @@ guide: | ## Prerequisites - Support is not in a tagged release and, as of this writing, not yet merged - (vllm-project/vllm-omni#7759). + This source checkout includes the pipeline and targets final vLLM 0.33.0. + Full image-quality qualification remains pending; successful generation alone + does not establish reference parity. - The pipeline itself is known-good: brought up from PR #7759 on one NVIDIA - GB300, it generated a 1024x1024 PNG in **4.5 s** over 40 steps at **34.0 GB** - peak. + Historical bring-up from PR #7759 on one NVIDIA GB300 generated a 1024x1024 + PNG in **4.5 s** over 40 steps at **34.0 GB** peak. Those measurements do not + qualify the current source/runtime. ```bash git clone https://github.com/vllm-project/vllm-omni.git && cd vllm-omni - git fetch origin pull/7759/head:qwen-image-2.1 && git checkout qwen-image-2.1 uv venv --python 3.12 --seed && source .venv/bin/activate - uv pip install vllm==0.29.0 --torch-backend=auto && uv pip install -e . + uv pip install vllm==0.33.0 --torch-backend=cu130 && uv pip install -e . ``` ## Offline inference diff --git a/recipes/models/Qwen/Qwen-Image.yaml b/recipes/models/Qwen/Qwen-Image.yaml index 404081d437d..268139fa2d1 100644 --- a/recipes/models/Qwen/Qwen-Image.yaml +++ b/recipes/models/Qwen/Qwen-Image.yaml @@ -25,8 +25,8 @@ omni: - t2i dependencies: - - note: "vLLM-Omni must be installed from source and pins vllm==0.18.0 for diffusion support" - command: "git clone https://github.com/vllm-project/vllm-omni.git && cd vllm-omni && uv pip install -e . vllm==0.18.0" + - note: "vLLM-Omni must be installed from source with final vllm==0.33.0" + command: "git clone https://github.com/vllm-project/vllm-omni.git && cd vllm-omni && uv pip install --torch-backend=cu130 -e . vllm==0.33.0" features: {} @@ -84,7 +84,7 @@ guide: | cd vllm-omni uv venv source .venv/bin/activate - uv pip install -e . vllm==0.18.0 + uv pip install --torch-backend=cu130 -e . vllm==0.33.0 ``` ## Usage diff --git a/recipes/models/fishaudio/drama-3.yaml b/recipes/models/fishaudio/drama-3.yaml index 04ef4a2d9f0..133e4afe181 100644 --- a/recipes/models/fishaudio/drama-3.yaml +++ b/recipes/models/fishaudio/drama-3.yaml @@ -88,7 +88,7 @@ guide: | ```bash uv venv && source .venv/bin/activate - uv pip install vllm==0.31.0 --torch-backend=auto + uv pip install vllm==0.33.0 --torch-backend=cu130 uv pip install git+https://github.com/vllm-project/vllm-omni.git ``` diff --git a/recipes/models/meituan-longcat/LongCat-Image-Edit.yaml b/recipes/models/meituan-longcat/LongCat-Image-Edit.yaml index 4b14bcd1710..9bc29a634a0 100644 --- a/recipes/models/meituan-longcat/LongCat-Image-Edit.yaml +++ b/recipes/models/meituan-longcat/LongCat-Image-Edit.yaml @@ -24,12 +24,8 @@ omni: - i2i dependencies: - - note: "vLLM-Omni must be installed from source with vllm==0.12.0 for LongCat-Image-Edit" - command: "git clone https://github.com/vllm-project/vllm-omni.git && cd vllm-omni && uv pip install -e . vllm==0.12.0" - - note: "xformers CUDA 12.8 build required for the diffusion attention kernels" - command: "uv pip install -U xformers --index-url https://download.pytorch.org/whl/cu128" - - note: "diffusers from source (needed by the image-edit pipeline)" - command: "git clone https://github.com/huggingface/diffusers.git && cd diffusers && uv pip install -e ." + - note: "vLLM-Omni must be installed from source with vllm==0.33.0 for LongCat-Image-Edit" + command: "git clone https://github.com/vllm-project/vllm-omni.git && cd vllm-omni && uv pip install --torch-backend=cu130 -e . vllm==0.33.0" features: {} @@ -56,9 +52,8 @@ guide: | ## Prerequisites - Hardware: 1x GPU with >=40 GB VRAM - - vLLM-Omni (runs on top of vLLM 0.12.0) - - diffusers (latest from source) - - xformers (latest) + - vLLM-Omni with final vLLM 0.33.0 and its CUDA 13 / PyTorch 2.14 runtime + - diffusers 0.40.0 (installed by the source checkout) ## Installation @@ -68,15 +63,8 @@ guide: | cd vllm-omni uv venv source .venv/bin/activate - uv pip install -e . vllm==0.12.0 - - # Update xformers to the latest version - uv pip install -U xformers --index-url https://download.pytorch.org/whl/cu128 - - # Update diffusers to the latest version - git clone https://github.com/huggingface/diffusers.git - cd diffusers - uv pip install -e . + uv pip install --torch-backend=cu130 -e . vllm==0.33.0 + uv pip check ``` ## Usage diff --git a/recipes/models/nvidia/Cosmos3-Nano.yaml b/recipes/models/nvidia/Cosmos3-Nano.yaml index 78bc08bd598..e07d7191218 100644 --- a/recipes/models/nvidia/Cosmos3-Nano.yaml +++ b/recipes/models/nvidia/Cosmos3-Nano.yaml @@ -113,7 +113,10 @@ guide: | - 1× H200 (or H100 / A100) for the documented single-GPU profile - The release-tested `vllm/vllm-omni:cosmos3` container, or vLLM-Omni installed - on top of `vllm==0.21.0` + on top of final `vllm==0.33.0` for this source checkout. The published + Cosmos3 container and its historical results use an older runtime and do + not qualify the current source. The separate Reasoner environment below, + where present, retains its own plugin/runtime requirements. - Prompts should be JSON-upsampled for best quality (see [cosmos-framework prompt upsampling](https://github.com/NVIDIA/cosmos-framework/blob/main/docs/prompt_upsampling.md)) diff --git a/recipes/models/nvidia/Cosmos3-Super-Image2Video.yaml b/recipes/models/nvidia/Cosmos3-Super-Image2Video.yaml index 7ff0151aa02..0b6ab55442d 100644 --- a/recipes/models/nvidia/Cosmos3-Super-Image2Video.yaml +++ b/recipes/models/nvidia/Cosmos3-Super-Image2Video.yaml @@ -98,7 +98,10 @@ guide: | - 8× H200, H100, or A100 for the documented full-node profile - The release-tested `vllm/vllm-omni:cosmos3` container, or vLLM-Omni installed - on top of `vllm==0.21.0` + on top of final `vllm==0.33.0` for this source checkout. The published + Cosmos3 container and its historical results use an older runtime and do + not qualify the current source. The separate Reasoner environment below, + where present, retains its own plugin/runtime requirements. ## Launch command (vLLM-Omni) diff --git a/recipes/models/nvidia/Cosmos3-Super-Text2Image.yaml b/recipes/models/nvidia/Cosmos3-Super-Text2Image.yaml index a9a3e5ecc96..46f187b1dd8 100644 --- a/recipes/models/nvidia/Cosmos3-Super-Text2Image.yaml +++ b/recipes/models/nvidia/Cosmos3-Super-Text2Image.yaml @@ -85,7 +85,10 @@ guide: | - 8× H100 for the documented full-node profile (4×H200 / 4×GB200 alternative below) - The release-tested `vllm/vllm-omni:cosmos3` container, or vLLM-Omni installed - on top of `vllm==0.21.0` + on top of final `vllm==0.33.0` for this source checkout. The published + Cosmos3 container and its historical results use an older runtime and do + not qualify the current source. The separate Reasoner environment below, + where present, retains its own plugin/runtime requirements. ## Launch command (vLLM-Omni) diff --git a/recipes/models/nvidia/Cosmos3-Super.yaml b/recipes/models/nvidia/Cosmos3-Super.yaml index db81e918544..8dd3d68cfb9 100644 --- a/recipes/models/nvidia/Cosmos3-Super.yaml +++ b/recipes/models/nvidia/Cosmos3-Super.yaml @@ -120,7 +120,10 @@ guide: | - 8× H200, H100, or A100 for the documented full-node profile - The release-tested `vllm/vllm-omni:cosmos3` container, or vLLM-Omni installed - on top of `vllm==0.21.0` + on top of final `vllm==0.33.0` for this source checkout. The published + Cosmos3 container and its historical results use an older runtime and do + not qualify the current source. The separate Reasoner environment below, + where present, retains its own plugin/runtime requirements. - Prompts should be JSON-upsampled for best quality (see [cosmos-framework prompt upsampling](https://github.com/NVIDIA/cosmos-framework/blob/main/docs/prompt_upsampling.md)) diff --git a/recipes/models/stabilityai/stable-audio-open-1.0.yaml b/recipes/models/stabilityai/stable-audio-open-1.0.yaml index 7ea67214878..f637dc3a97a 100644 --- a/recipes/models/stabilityai/stable-audio-open-1.0.yaml +++ b/recipes/models/stabilityai/stable-audio-open-1.0.yaml @@ -29,8 +29,8 @@ omni: - t2a dependencies: - - note: "vLLM 0.22.0 or newer" - command: "uv pip install 'vllm>=0.22.0'" + - note: "Final vLLM 0.33.0 for this source checkout" + command: "uv pip install --torch-backend=cu130 vllm==0.33.0" - note: "vllm-omni provides the audio generation backend" command: "uv pip install git+https://github.com/vllm-project/vllm-omni.git" - note: "soundfile (recommended) or scipy for WAV output" diff --git a/recipes/models/stabilityai/stable-diffusion-3.5-medium.yaml b/recipes/models/stabilityai/stable-diffusion-3.5-medium.yaml index aa05c675c8a..8518613a574 100644 --- a/recipes/models/stabilityai/stable-diffusion-3.5-medium.yaml +++ b/recipes/models/stabilityai/stable-diffusion-3.5-medium.yaml @@ -29,14 +29,14 @@ omni: - t2i dependencies: - - note: "Pin vllm==0.12.0 for Stable Diffusion 3.5" - command: "uv pip install vllm==0.12.0" + - note: "Pin vllm==0.33.0 for Stable Diffusion 3.5" + command: "uv pip install --torch-backend=cu130 vllm==0.33.0" brand: NVIDIA - note: "vllm-omni provides the image generation backend" command: "uv pip install git+https://github.com/vllm-project/vllm-omni.git" brand: NVIDIA - - note: "AMD ROCm: install vLLM for ROCm and vLLM-Omni from source" - command: "VLLM_OMNI_TARGET_DEVICE=rocm uv pip install git+https://github.com/vllm-project/vllm-omni.git vllm==0.20.0+rocm721 --extra-index-url https://wheels.vllm.ai/rocm/0.20.0/rocm721" + - note: "AMD ROCm: build the frozen vLLM 0.33.0 source runtime; native validation pending" + command: "git clone https://github.com/vllm-project/vllm-omni.git && cd vllm-omni && docker build -f docker/Dockerfile.rocm -t vllm-omni-rocm:v033-source ." brand: AMD features: {} @@ -83,7 +83,7 @@ guide: | ## Prerequisites - - vLLM-Omni on top of vLLM 0.12.0 + - vLLM-Omni on top of vLLM 0.33.0 - `diffusers` library ## Installation @@ -91,19 +91,20 @@ guide: | ```bash uv venv source .venv/bin/activate - uv pip install vllm==0.12.0 + uv pip install --torch-backend=cu130 vllm==0.33.0 uv pip install git+https://github.com/vllm-project/vllm-omni.git ``` ### AMD ROCm (MI300X, MI325X, MI355X) - > **Note:** The vLLM wheel for ROCm requires Python 3.12, ROCm 7.2.1, and glibc >= 2.35. - > If your environment does not meet these requirements, please use the Docker-based setup. + > The Dockerfile builds the frozen vLLM 0.33.0 source over its published + > bootstrap image. ROCm build and native model validation remain pending; + > this is not a claim that a final ROCm wheel is available. ```bash uv venv --python 3.12 source .venv/bin/activate - VLLM_OMNI_TARGET_DEVICE=rocm uv pip install git+https://github.com/vllm-project/vllm-omni.git vllm==0.20.0+rocm721 --extra-index-url https://wheels.vllm.ai/rocm/0.20.0/rocm721 + git clone https://github.com/vllm-project/vllm-omni.git && cd vllm-omni && docker build -f docker/Dockerfile.rocm -t vllm-omni-rocm:v033-source . ``` ## Python Usage diff --git a/requirements/common.txt b/requirements/common.txt index d2df485f576..805f4ece121 100644 --- a/requirements/common.txt +++ b/requirements/common.txt @@ -1,6 +1,6 @@ # Common dependencies for all platforms opentelemetry-api>=1.28.0 -# Match vLLM v0.32.0's supported Transformers range. The release runtime +# Use the 5.16+ Transformers APIs shared with vLLM v0.33.0. The release runtime # uses the 5.16+ model and tokenizer APIs; the former <5.15 cap conflicts. # MOSS-TTS-Nano's list/set weight-loading compatibility is handled by # transformers_keys_to_ignore_compat in model_executor/models/utils.py. diff --git a/requirements/cuda-bootstrap.txt b/requirements/cuda-bootstrap.txt new file mode 100644 index 00000000000..d328c0aac32 --- /dev/null +++ b/requirements/cuda-bootstrap.txt @@ -0,0 +1,12 @@ +lmcache==0.5.6 +opentelemetry-api==1.40.0 +opentelemetry-exporter-http-transport==0.66b2 +opentelemetry-exporter-otlp==1.40.0 +opentelemetry-exporter-otlp-proto-common==1.40.0 +opentelemetry-exporter-otlp-proto-grpc==1.40.0 +opentelemetry-exporter-otlp-proto-http==1.40.0 +opentelemetry-exporter-prometheus==0.61b0 +opentelemetry-proto==1.40.0 +opentelemetry-sdk==1.40.0 +opentelemetry-semantic-conventions==0.61b0 +opentelemetry-semantic-conventions-ai==0.5.1 diff --git a/requirements/cudnn-runtime.txt b/requirements/cudnn-runtime.txt new file mode 100644 index 00000000000..41d86b4735a --- /dev/null +++ b/requirements/cudnn-runtime.txt @@ -0,0 +1,3 @@ +nvidia-cudnn-cu13==9.25.1.1 \ + --hash=sha256:195199e8c7f8bca0c366ea675dd2441a406646cf95a98bcda66c999d57dd1f26 \ + --hash=sha256:6ed974d46b1b5fd123c0d9ff031c857142aa94257c6df41c40f6b5fd11a01d50 diff --git a/tests/buildkite/test_aiter_mha_cache.py b/tests/buildkite/test_aiter_mha_cache.py index e854ec5dd9a..b1144ceb84b 100644 --- a/tests/buildkite/test_aiter_mha_cache.py +++ b/tests/buildkite/test_aiter_mha_cache.py @@ -155,12 +155,25 @@ def test_image_dependency_identity_tracks_the_rocm_image_inputs(tmp_path: Path) requirements = root / "requirements" requirements.mkdir() (requirements / "rocm.txt").write_text("torch\n", encoding="utf-8") + backport = root / "tools/gradio_pillow13" + backport.mkdir() + for filename in ("manifest.json", "backport.patch", "build-constraints.txt", "NOTICE.md"): + (backport / filename).write_text(filename, encoding="utf-8") original = cache.dependency_identity(root) assert len(original) == 64 assert cache.dependency_identity(root) == original (requirements / "rocm.txt").write_text("torch==2\n", encoding="utf-8") assert cache.dependency_identity(root) != original + (requirements / "rocm.txt").write_text("torch\n", encoding="utf-8") + for path in (root / "tools/build_gradio_pillow13.py", *backport.iterdir()): + content = path.read_bytes() + path.write_bytes(content + b"changed\n") + assert cache.dependency_identity(root) != original, path + path.write_bytes(content) + assert cache.dependency_identity(root) == original + shutil.rmtree(backport) + assert cache.dependency_identity(root) == "" def test_stage_and_restore_verify_module_identity_and_provenance(tmp_path: Path) -> None: diff --git a/tests/buildkite/test_amd_pipeline.py b/tests/buildkite/test_amd_pipeline.py index 1cbd8c4ef98..4f10954d349 100644 --- a/tests/buildkite/test_amd_pipeline.py +++ b/tests/buildkite/test_amd_pipeline.py @@ -3182,7 +3182,9 @@ def test_minicpmo_ready_jobs_select_marked_core_model_nodes() -> None: assert step["artifact_paths"] == [f"{artifact}/**/*"] assert f"--kill-after=1m {command_timeout}" in command assert _MINICPMO_READY_MARKER in command - assert "--run-level core_model" in command + # The public client waits for a model-generated turn terminator. + run_level = "advanced_model" if label == "ROCm · MiniCPM-o 4.5 Duplex Coverage" else "core_model" + assert f"--run-level {run_level}" in command assert command.count("--shard-id=0 --num-shards=1") == 2 assert "test_minicpmo_4_5_duplex_mrv2.py" not in command assert "test_minicpmo_4_5_ready.json" not in command diff --git a/tests/buildkite/test_check_vllm_runtime.py b/tests/buildkite/test_check_vllm_runtime.py index 8e4801e07de..515601a147e 100644 --- a/tests/buildkite/test_check_vllm_runtime.py +++ b/tests/buildkite/test_check_vllm_runtime.py @@ -3,6 +3,7 @@ """The CUDA runtime gate must verify release wheel build commits.""" +import importlib import runpy import sys from pathlib import Path @@ -15,12 +16,17 @@ SCRIPT = Path(__file__).resolve().parents[2] / ".buildkite/cuda/scripts/check_vllm_runtime.py" -@pytest.mark.parametrize("commit", ["g5df63a720", "g82eb574d9", "gdeadbeef0", None, "g5df6"]) -def test_release_wheel_commit_check(monkeypatch, commit): - # All candidates advertise the same release version. Only the wheel built - # from the pinned source can pass; absent/truncated build metadata fails. +@pytest.mark.parametrize( + "commit", ["g4ba787ec9", "g61cbf5429", "gcbafc4690", "g4135590ff", "g82eb574d9", "gdeadbeef0", None, "g5df6"] +) +@pytest.mark.parametrize("lmcache_abi_error", [False, True]) +@pytest.mark.parametrize("version", ["0.33.0", "0.33.0rc3.dev3+g4ba787ec9", "0.33.0.post1"]) +@pytest.mark.parametrize("cudnn", [92501, 92400]) +def test_release_wheel_commit_check(monkeypatch, commit, lmcache_abi_error, version, cudnn): + # Source and package identities are independent: RC and final wheels can + # share a commit, and a different build can advertise the final version. fake_vllm = ModuleType("vllm") - vars(fake_vllm)["__version__"] = "0.32.0" + vars(fake_vllm)["__version__"] = version fake_version = ModuleType("vllm._version") vars(fake_version)["__commit_id__"] = commit monkeypatch.setitem(sys.modules, "vllm", fake_vllm) @@ -29,16 +35,36 @@ def test_release_wheel_commit_check(monkeypatch, commit): sys.modules, "torch", SimpleNamespace( - __version__="2.13.0", + __version__="2.14.0", version=SimpleNamespace(cuda="13.0"), + backends=SimpleNamespace(cudnn=SimpleNamespace(version=lambda: cudnn)), cuda=SimpleNamespace(is_available=lambda: True, get_device_name=lambda: "test GPU"), ), ) - for name in ("flashinfer", "vllm._C_stable_libtorch"): + for name in ("flashinfer", "vllm._C_stable_libtorch", "lmcache.cuda_ops"): monkeypatch.setitem(sys.modules, name, ModuleType(name)) + if lmcache_abi_error: + original_import = importlib.import_module + + def import_module(name: str) -> ModuleType: + if name == "lmcache.cuda_ops": + raise ImportError("undefined symbol: _ZN3c104cuda29c10_cuda_check_implementationEiPKcS2_jb") + return original_import(name) + + monkeypatch.setattr(importlib, "import_module", import_module) main = runpy.run_path(str(SCRIPT))["main"] - if commit == "g5df63a720": - main() + if commit == "g4ba787ec9": + if version != "0.33.0": + with pytest.raises(AssertionError, match="Expected vLLM version"): + main() + elif cudnn != 92501: + with pytest.raises(AssertionError, match="Expected cuDNN"): + main() + elif lmcache_abi_error: + with pytest.raises(ImportError, match="undefined symbol:.*c10_cuda_check_implementation"): + main() + else: + main() else: with pytest.raises(AssertionError, match="Expected vLLM"): main() diff --git a/tests/buildkite/test_rocm_dockerfile.py b/tests/buildkite/test_rocm_dockerfile.py index 294fb7cb86f..4bc306ba149 100644 --- a/tests/buildkite/test_rocm_dockerfile.py +++ b/tests/buildkite/test_rocm_dockerfile.py @@ -31,9 +31,9 @@ def _line_index(lines: list[str], prefix: str) -> int: return matches[0] -def test_rocm_base_tracks_cuda_vllm_release() -> None: - # CI installs the exact final wheel on its digest-pinned release image; - # production CUDA/ROCm images track the same stable release. +def test_rocm_base_tracks_cuda_bootstrap_release() -> None: + # Bootstrap images must exist; the candidate runtime is installed later + # from the pinned CUDA wheel or the matching ROCm source commit. assert _docker_arg(CUDA_CI_DOCKERFILE, "VLLM_BASE_IMAGE") == "vllm/vllm-openai" assert "@sha256:" in _docker_arg(CUDA_CI_DOCKERFILE, "VLLM_BASE_TAG") assert len(_docker_arg(CUDA_CI_DOCKERFILE, "VLLM_PRECOMPILED_WHEEL_COMMIT")) == 40 @@ -45,11 +45,11 @@ def test_rocm_base_tracks_cuda_vllm_release() -> None: assert separator, f"expected a tagged ROCm base image, got {rocm_base}" assert cuda_image == "vllm/vllm-openai" assert image_ref == "vllm/vllm-openai-rocm" - assert image_tag == cuda_release_tag == f"v{_docker_arg(CUDA_CI_DOCKERFILE, 'VLLM_VERSION')}" + assert image_tag == cuda_release_tag == _docker_arg(CUDA_CI_DOCKERFILE, "VLLM_BASE_TAG").split("@")[0] -def test_rocm_defaults_to_prebuilt_base_image() -> None: - assert _docker_arg(ROCM_DOCKERFILE, "USE_NIGHTLY_BUILD") == "0" +def test_rocm_defaults_to_building_the_candidate_runtime() -> None: + assert _docker_arg(ROCM_DOCKERFILE, "USE_NIGHTLY_BUILD") == "1" def test_amd_build_uses_registry_cache() -> None: @@ -74,6 +74,8 @@ def test_amd_cache_key_tracks_dependency_inputs_and_architecture() -> None: "setup.py", "requirements", "tools/install_torchcodec_rocm.sh", + "tools/build_gradio_pillow13.py", + "tools/gradio_pillow13/*", "ROCM_ARCH", ): assert cache_input in script @@ -112,12 +114,16 @@ def test_rocm_build_context_excludes_git_history() -> None: def test_rocm_source_ref_tracks_cuda_vllm_release() -> None: - # The opt-in ROCm source rebuild defaults to the released CUDA vLLM tag; - # CI itself pins the release wheel via VLLM_PRECOMPILED_WHEEL_COMMIT. - cuda_release_tag = _docker_arg(CUDA_RELEASE_DOCKERFILE, "BASE_IMAGE").rpartition(":")[2] + # Both runtime installations must track the candidate, independently of + # the older published images used to bootstrap their dependencies. + cuda_version = _docker_arg(CUDA_RELEASE_DOCKERFILE, "VLLM_VERSION") rocm_source_ref = _docker_arg(ROCM_DOCKERFILE, "VLLM_VERSION_OR_COMMIT_HASH") - assert rocm_source_ref == cuda_release_tag == f"v{_docker_arg(CUDA_CI_DOCKERFILE, 'VLLM_VERSION')}" + assert cuda_version == _docker_arg(CUDA_CI_DOCKERFILE, "VLLM_VERSION") + assert rocm_source_ref == _docker_arg(CUDA_CI_DOCKERFILE, "VLLM_PRECOMPILED_WHEEL_COMMIT") + assert _docker_arg(CUDA_RELEASE_DOCKERFILE, "VLLM_PRECOMPILED_WHEEL_COMMIT") == _docker_arg( + CUDA_CI_DOCKERFILE, "VLLM_PRECOMPILED_WHEEL_COMMIT" + ) def test_rocm_dockerfile_contains_vllm_api_canary() -> None: diff --git a/tests/config/test_config_import_cycle.py b/tests/config/test_config_import_cycle.py index e149b63ee44..5711c964fbb 100644 --- a/tests/config/test_config_import_cycle.py +++ b/tests/config/test_config_import_cycle.py @@ -29,8 +29,8 @@ def _top_level_import_modules(path: Path) -> set[str]: def _assert_isolated_import_succeeds(script: str) -> None: env = os.environ.copy() - env.pop("CUDA_VISIBLE_DEVICES", None) - env.pop("HIP_VISIBLE_DEVICES", None) + env["CUDA_VISIBLE_DEVICES"] = "" + env["HIP_VISIBLE_DEVICES"] = "" result = subprocess.run( [sys.executable, "-c", script], cwd=REPO_ROOT, diff --git a/tests/core/sched/test_generation_scheduler_restore.py b/tests/core/sched/test_generation_scheduler_restore.py index c4debba5bf7..889cdb14d8a 100644 --- a/tests/core/sched/test_generation_scheduler_restore.py +++ b/tests/core/sched/test_generation_scheduler_restore.py @@ -17,6 +17,7 @@ from vllm.v1.core.sched.request_queue import SchedulingPolicy, create_request_queue from vllm.v1.request import Request, RequestStatus +from tests.helpers.omni_scheduler import make_kv_cache_manager from vllm_omni.core.sched.omni_generation_scheduler import OmniGenerationScheduler, _has_async_chunk_payload_to_run from vllm_omni.core.sched.omni_scheduling_coordinator import OmniSchedulingCoordinator from vllm_omni.engine.serialization import serialize_additional_information @@ -127,8 +128,11 @@ def _make_generation_scheduler(waiting_request, *, use_v2_model_runner=False): def test_generation_scheduler_preserves_auxiliary_output_metadata() -> None: request = Request("aux-contract", [1, 2], SamplingParams(max_tokens=2), pooling_params=None) request.block_hashes.append(b"block-hash") + # Hashes are emitted only for a completed hash block. One scheduled token + # does not cover the default 16-token block, so count a full block computed. + request.num_computed_tokens = 16 scheduler = _make_generation_scheduler(request, use_v2_model_runner=True) - scheduler.aux_output_connector = AuxOutputSchedulerConnector() + scheduler.aux_output_connector = AuxOutputSchedulerConnector(make_kv_cache_manager(block_size=16)) output = scheduler._build_generation_scheduler_output( new_reqs_data=[], diff --git a/tests/core/sched/test_omni_ar_materialized_cache.py b/tests/core/sched/test_omni_ar_materialized_cache.py new file mode 100644 index 00000000000..2524d6bcd12 --- /dev/null +++ b/tests/core/sched/test_omni_ar_materialized_cache.py @@ -0,0 +1,145 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project +"""Materialized AR outputs must publish their completed prefix-cache blocks.""" + +import pytest + +# Import Omni before Request so the request type is the runtime's patched type. +# isort: off +import vllm_omni # noqa: F401 +from vllm.sampling_params import SamplingParams +from vllm.utils.hashing import sha256 +from vllm.v1.core.kv_cache_utils import get_request_block_hasher, init_none_hash +from vllm.v1.core.sched.output import SchedulerOutput +from vllm.v1.request import Request, RequestStatus +# isort: on + +from tests.core.sched.test_fast_allocate import BLOCK, _manager +from tests.core.sched.test_omni_ar_scheduler_logprobs import _make_scheduler_stub +from vllm_omni.core.sched.omni_ar_scheduler import OmniARAsyncScheduler, OmniARScheduler +from vllm_omni.outputs import OmniModelRunnerOutput + +pytestmark = [pytest.mark.core_model, pytest.mark.cpu] + + +def _request(request_id, tokens): + return Request( + request_id, + list(tokens), + SamplingParams(max_tokens=200, stop_token_ids=[99]), + None, + block_hasher=get_request_block_hasher(BLOCK, sha256), + ) + + +def _prepared(*, trailing_inflight=0, sync_spec=False): + init_none_hash(sha256) + manager = _manager(enable_caching=True) + request = _request("source", range(31)) + request.status = RequestStatus.RUNNING + # Async allocation can run before the previous sampled token is delivered: + # reserve two blocks, but the known 31-token sequence hashes only one. + scheduled = 3 if sync_spec else 1 + computed = 34 if sync_spec else 32 + trailing_inflight + assert manager.allocate_slots(request, computed) is not None + request.append_output_token_ids(7) + request.num_computed_tokens = computed + request.num_in_flight_tokens = scheduled + trailing_inflight + request.num_output_placeholders = 0 if sync_spec else 1 + trailing_inflight + request.prefill_stats = None + cls = OmniARScheduler if sync_spec else OmniARAsyncScheduler + stub = _make_scheduler_stub([request]) + scheduler = cls.__new__(cls) + scheduler.__dict__.update(stub.__dict__) + scheduler.kv_cache_manager = manager + scheduler.max_model_len = 256 + scheduler.num_spec_tokens = 2 if sync_spec else 0 + scheduler._process_kv_transfer_trigger = lambda *args: False + scheduler._handle_stopped_request = lambda *args: True + before_free = [] + + def free_request(req): + before_free.append(_hit(manager, list(req.all_token_ids))) + manager.free(req) + scheduler.requests.pop(req.request_id, None) + scheduler.finished_req_ids.add(req.request_id) + scheduler.finished_req_ids_dict[req.client_index].add(req.request_id) + return None, None + + scheduler._free_request = free_request + return scheduler, request, manager, before_free + + +def _hit(manager, tokens): + query = _request("lookup", tokens) + return manager.get_computed_blocks(query)[1] + + +def _frame(*, general=False, token=8, sync_spec=False): + scheduler_output = SchedulerOutput.make_empty() + scheduler_output.num_scheduled_tokens = {"source": 3 if sync_spec else 1} + scheduler_output.total_num_scheduled_tokens = 3 if sync_spec else 1 + scheduler_output.scheduled_spec_decode_tokens = {"source": [9, 10]} if sync_spec else {} + runner_output = OmniModelRunnerOutput( + req_ids=["source"], + req_id_to_index={"source": 0}, + sampled_token_ids=[[token]], + logprobs=None, + prompt_logprobs_dict={}, + pooler_output=[None] if general else None, + ) + return scheduler_output, runner_output + + +@pytest.mark.parametrize("general", [False, True], ids=["fast", "general"]) +@pytest.mark.parametrize("trailing_inflight", [0, 1], ids=["settled", "one_step_inflight"]) +def test_materialized_async_output_exposes_completed_prefix(mocker, general, trailing_inflight): + scheduler, request, manager, _ = _prepared(trailing_inflight=trailing_inflight) + lookup = [*request.all_token_ids, 8] + assert _hit(manager, lookup) == BLOCK + + publication = mocker.spy(manager, "cache_blocks") + scheduler.update_from_output(*_frame(general=general)) + + publication.assert_called_once_with(request, 2 * BLOCK) + assert request.num_computed_tokens - request.num_in_flight_tokens == 2 * BLOCK + assert _hit(manager, lookup) == 2 * BLOCK + + +@pytest.mark.parametrize("general", [False, True], ids=["fast", "general"]) +def test_stop_publishes_completed_prefix_before_free(general): + scheduler, request, manager, before_free = _prepared() + lookup = [*request.all_token_ids, 99] + assert _hit(manager, lookup) == BLOCK + + scheduler.update_from_output(*_frame(general=general, token=99)) + + assert before_free == [2 * BLOCK] + assert _hit(manager, lookup) == 2 * BLOCK + + +@pytest.mark.parametrize("preempted", [False, True], ids=["discarded", "preempted"]) +def test_stale_output_does_not_publish_unconfirmed_prefix(preempted): + scheduler, request, manager, _ = _prepared() + request.num_output_placeholders = 0 + request.num_stale_output_tokens = 1 + request.drop_stale_output = not preempted + if preempted: + request.status = RequestStatus.PREEMPTED + lookup = [*request.all_token_ids, 8] + assert _hit(manager, lookup) == BLOCK + + scheduler.update_from_output(*_frame()) + + assert _hit(manager, lookup) == BLOCK + + +def test_sync_spec_output_exposes_only_verified_prefix(): + scheduler, request, manager, _ = _prepared(sync_spec=True) + lookup = [*request.all_token_ids, 8] + assert _hit(manager, lookup) == BLOCK + + scheduler.update_from_output(*_frame(sync_spec=True)) + + assert request.num_computed_tokens == 2 * BLOCK + assert _hit(manager, lookup) == 2 * BLOCK diff --git a/tests/core/sched/test_omni_ar_scheduler_payloadless_outputs.py b/tests/core/sched/test_omni_ar_scheduler_payloadless_outputs.py index 947799f0534..01982a74d6b 100644 --- a/tests/core/sched/test_omni_ar_scheduler_payloadless_outputs.py +++ b/tests/core/sched/test_omni_ar_scheduler_payloadless_outputs.py @@ -8,6 +8,7 @@ # isort: off import vllm_omni # noqa: F401 - import for side effects (patch vLLM) +from vllm.outputs import RequestError from vllm.v1.core.sched.output import SchedulerOutput from vllm.v1.engine import FinishReason from vllm.v1.metrics.stats import IterationStats, LoRARequestStates, RequestStateStats @@ -164,7 +165,12 @@ def finish_kv_failure(failed_ids, outputs): sched._handle_failed_kv_load_outputs = finish_kv_failure elif termination == "grammar_error": - sched.grammar_compile_error_reqs = {request.request_id} + sched.grammar_compile_error_reqs = { + request.request_id: RequestError( + code="validation_error", + message="Invalid structured outputs: fixture", + ) + } elif termination == "input_timeout": OmniSchedulerMixin._finish_input_timeout_requests(sched, {request.request_id}) elif termination == "context_overflow": diff --git a/tests/core/sched/test_omni_ar_scheduler_streaming.py b/tests/core/sched/test_omni_ar_scheduler_streaming.py index 227a77ae8e8..dbfc9f3215b 100644 --- a/tests/core/sched/test_omni_ar_scheduler_streaming.py +++ b/tests/core/sched/test_omni_ar_scheduler_streaming.py @@ -586,6 +586,40 @@ def test_stage0_streaming_update_keeps_all_computed_tokens_without_placeholder() assert sched._new_prompt_len_snapshot[session.request_id] == 2 +def test_stage0_streaming_update_keeps_full_hash_block() -> None: + """A fully computed length that is a multiple of hash_block_size must keep that block. + + The pin deletes ``block_hashes[keep_end // hash_block_size:]`` with + ``keep_end = num_tokens - 1``. Clamping the computed frontier to + ``num_tokens - 1`` before that delete drops hash 0 when the length is 16. + """ + sched = _make_scheduler(stage_id=0) + # Request stores this list. Compare against a copy so the extension + # inside _update_request_as_session does not rewrite the expected prefix. + prompt = list(range(16)) + expected_prefix = list(prompt) + session = Request( + request_id="req-ar-hash-block", + prompt_token_ids=prompt, + sampling_params=SamplingParams(max_tokens=8), + pooling_params=None, + arrival_time=100.0, + block_hasher=None, + ) + session.status = RequestStatus.WAITING_FOR_STREAMING_REQ + session.num_computed_tokens = 16 + session.block_hashes.append("h0") + + sched._update_request_as_session(session, _make_update([10, 20])) + + assert session.prompt_token_ids[:16] == expected_prefix + assert list(session._all_token_ids[:16]) == expected_prefix + assert session.prompt_token_ids[-2:] == [10, 20] + assert list(session._all_token_ids[-2:]) == [10, 20] + assert session.block_hashes[0] == "h0" + assert session.num_computed_tokens >= 16 + + def test_stage0_basic_window_rebuilds_below_low_watermark() -> None: sched = _make_scheduler(stage_id=0) session = _make_request() @@ -1405,6 +1439,8 @@ def test_overflow_guard_survives_eager_abort_until_fifo_cleanup(mocker, fifo_pub other.status = RequestStatus.RUNNING sched.running.append(other) engine = StageEngineCoreProc.__new__(StageEngineCoreProc) + # EngineCore.__init__ sets this; __new__ skips it. False is the unpaused default. + engine.reject_new_requests = False engine.scheduler = sched engine.engines_running = False engine.batch_queue = None diff --git a/tests/core/sched/test_omni_sched_deferred_free_fence.py b/tests/core/sched/test_omni_sched_deferred_free_fence.py index db1f9fb1c6b..a4242349a18 100644 --- a/tests/core/sched/test_omni_sched_deferred_free_fence.py +++ b/tests/core/sched/test_omni_sched_deferred_free_fence.py @@ -53,6 +53,7 @@ def _make_update_sched(scheduler_cls) -> MagicMock: sched.perf_metrics = None sched.chunk_transfer_adapter = None sched.defer_block_free = True + sched.defer_async_loads = False sched.processed_step_seq = 0 return sched @@ -83,6 +84,31 @@ def test_update_from_output_advances_fence_and_drains(scheduler_cls) -> None: sched._drain_deferred_frees.assert_called_once_with() +@pytest.mark.parametrize("scheduler_cls", [OmniARScheduler, OmniGenerationScheduler]) +def test_completed_step_starts_fenced_async_kv_load(scheduler_cls) -> None: + """A queued HiSparse load must become runnable when its GPU fence completes.""" + sched = _make_update_sched(scheduler_cls) + sched.defer_async_loads = True + pending_loads = {"remote-request": 1} + started_loads = [] + + def drain_loads(): + for request_id, fence in list(pending_loads.items()): + if fence <= sched.processed_step_seq: + started_loads.append(request_id) + del pending_loads[request_id] + + sched._drain_deferred_loads.side_effect = drain_loads + empty_step, model_output = _empty_frame(total_num_scheduled_tokens=0) + scheduler_cls.update_from_output(sched, empty_step, model_output) + assert not started_loads + + completed_step, model_output = _empty_frame(total_num_scheduled_tokens=3) + scheduler_cls.update_from_output(sched, completed_step, model_output) + assert started_loads == ["remote-request"] + assert not pending_loads + + @pytest.mark.parametrize("scheduler_cls", [OmniARScheduler, OmniGenerationScheduler]) def test_update_from_output_zero_token_step_leaves_fence(scheduler_cls) -> None: """0-token steps do not advance the fence upstream (their seq was never diff --git a/tests/core/sched/test_omni_scheduler_async_chunk_prewarm.py b/tests/core/sched/test_omni_scheduler_async_chunk_prewarm.py index 99c381eff66..8393d5d3c1f 100644 --- a/tests/core/sched/test_omni_scheduler_async_chunk_prewarm.py +++ b/tests/core/sched/test_omni_scheduler_async_chunk_prewarm.py @@ -14,6 +14,7 @@ import msgspec import pytest import torch +from vllm.outputs import RequestError from vllm.v1.core.sched.output import SchedulerOutput from vllm.v1.request import RequestStatus @@ -253,10 +254,21 @@ def test_scheduler_output_defaults_to_no_prewarms() -> None: def _finish_via_error_outputs(scheduler, request_id: str) -> None: - scheduler.grammar_compile_error_reqs = {request_id} + scheduler.grammar_compile_error_reqs = { + request_id: RequestError( + code="validation_error", + message="Invalid structured outputs: fixture", + ) + } scheduler._finish_error_requests(defaultdict(list)) +def _finish_via_encoder_cache_mismatch(scheduler, request_id: str) -> None: + scheduler.encoder_cache_mismatch_reqs = {request_id} + scheduler._finish_error_requests(defaultdict(list)) + assert scheduler.encoder_cache_mismatch_reqs == set() + + @_SCHEDULERS @pytest.mark.parametrize( "finish", @@ -265,6 +277,7 @@ def _finish_via_error_outputs(scheduler, request_id: str) -> None: pytest.param(lambda s, rid: s.finish_requests(rid, RequestStatus.FINISHED_ABORTED), id="abort-single-id"), pytest.param(lambda s, rid: s._finish_input_timeout_requests({rid}), id="input-timeout"), pytest.param(_finish_via_error_outputs, id="error-outputs"), + pytest.param(_finish_via_encoder_cache_mismatch, id="encoder-cache-mismatch"), ], ) def test_finish_paths_discard_undrained_prewarm(scheduler_cls, base_scheduler, finish) -> None: diff --git a/tests/core/sched/test_omni_scheduler_input_coordinator_cleanup.py b/tests/core/sched/test_omni_scheduler_input_coordinator_cleanup.py index e90834d4acf..7393c3267ea 100644 --- a/tests/core/sched/test_omni_scheduler_input_coordinator_cleanup.py +++ b/tests/core/sched/test_omni_scheduler_input_coordinator_cleanup.py @@ -190,6 +190,8 @@ def test_native_oversized_chunk_uses_real_finish_cleanup_and_emits_one_error( scheduler.log_stats = False # Match vLLM 0.32 initialization with KV-fetch tracking disabled. scheduler._kv_fetch_stages = None + # Scheduler.__init__ creates this map; finish pops it. Empty is the real default. + scheduler._pending_encoder_input_ids = {} scheduler.perf_metrics = None scheduler._last_stats_time = time.monotonic() diff --git a/tests/diffusion/attention/test_kernels_hub.py b/tests/diffusion/attention/test_kernels_hub.py index c85fb58f3a8..c1e81fe34e9 100644 --- a/tests/diffusion/attention/test_kernels_hub.py +++ b/tests/diffusion/attention/test_kernels_hub.py @@ -19,9 +19,33 @@ def _clear_hub_module_cache(monkeypatch: pytest.MonkeyPatch) -> None: monkeypatch.setattr(flash_attn_hub, "_hub_modules", {}) +@pytest.mark.core_model +@pytest.mark.cpu +def test_flash_attention_hub_uses_torch_stable_abi_build(monkeypatch): + """Torch 2.14 has no FA2 v1/v2 binaries; v3 supplies a stable-ABI build.""" + from vllm_omni.diffusion.attention.backends.flash_attn_hub import _load_hub_module + + kernel = types.SimpleNamespace(flash_attn_func=lambda *args, **kwargs: None) + attempted = [] + + def get_kernel(repo_id, *, version): + attempted.append(version) + if version in (1, 2): + raise FileNotFoundError("Torch version 2.13 does not match environment Torch version 2.14") + assert version == 3 + return kernel + + monkeypatch.setitem(sys.modules, "kernels", types.SimpleNamespace(get_kernel=get_kernel)) + assert _load_hub_module("kernels-community/flash-attn2") is kernel + assert 3 in attempted + + @pytest.mark.core_model @pytest.mark.cpu def test_explicit_kernels_hub_selection_does_not_fallback(monkeypatch: pytest.MonkeyPatch): + # Explicit FA/Sage selections must not initialize the unrelated cuDNN runtime. + monkeypatch.setattr(torch.backends.cudnn, "version", lambda: pytest.fail("Unexpected cuDNN initialization")) + from vllm_omni.platforms.cuda.platform import CudaOmniPlatform # Temporarily hide/remove kernels module if present diff --git a/tests/diffusion/distributed/test_wan_vae_fastpath_verdicts.py b/tests/diffusion/distributed/test_wan_vae_fastpath_verdicts.py index 50c5d89d503..2d9c67a0b7f 100644 --- a/tests/diffusion/distributed/test_wan_vae_fastpath_verdicts.py +++ b/tests/diffusion/distributed/test_wan_vae_fastpath_verdicts.py @@ -19,7 +19,7 @@ @pytest.fixture(params=["_SPATIAL_PAD_VERDICTS", "_CONV_OUT_LAYOUT_VERDICTS"]) def verdict_cache(request, monkeypatch): - cache = weakref.WeakKeyDictionary() + cache: weakref.WeakKeyDictionary[nn.Module, dict[tuple, bool]] = weakref.WeakKeyDictionary() monkeypatch.setattr(fp, request.param, cache) return request.param, cache @@ -133,8 +133,8 @@ def run(width): assert torch.equal(output, expected) assert len(cache[conv]) <= 2 - # One comparison on random operands decides each new shape. - check = [matches] + # Acceptance requires all eight probes; rejection stops at the first mismatch. + check = [True] * 8 if matches else [False] run(4) run(5) assert comparisons == check * 2 @@ -206,7 +206,7 @@ def run(x): @pytest.mark.parametrize("history", [None, 1, 2]) def test_causal_conv_forward_shares_spatial_padding_verdicts(cpu_assembly, monkeypatch, matches: bool, history) -> None: """A normally called (e.g. hooked) convolution keeps the verified cuDNN padding path.""" - cache = weakref.WeakKeyDictionary() + cache: weakref.WeakKeyDictionary[nn.Module, dict[tuple, bool]] = weakref.WeakKeyDictionary() monkeypatch.setattr(fp, "_SPATIAL_PAD_VERDICTS", cache) generator = torch.Generator().manual_seed(0) conv = fp.WanCausalConv3d(2, 2, 3, padding=1) @@ -234,7 +234,7 @@ def compare(a, b): monkeypatch.setattr(fp, "_bitwise_equal", compare) for _ in range(2): assert torch.equal(fp.causal_conv_forward(conv, x, cache_x), expected) - check = [matches] + check = [True] * 8 if matches else [False] assert comparisons == check # Verified once, then reused. assert list(cache[conv].values()) == [matches] # The cached call site shares the verdict instead of probing again. diff --git a/tests/diffusion/layers/test_adalayernorm_cuda_path.py b/tests/diffusion/layers/test_adalayernorm_cuda_path.py index 6e6c8e13f07..57de4fcf6b5 100644 --- a/tests/diffusion/layers/test_adalayernorm_cuda_path.py +++ b/tests/diffusion/layers/test_adalayernorm_cuda_path.py @@ -69,6 +69,26 @@ def test_fused_fast_path_matches_native(dtype, affine, bs, seq, hidden): assert_close(out, fp32_reference(x, scale, shift, 1e-6), dtype, loose=True) +@pytest.mark.parametrize("dtype", [torch.bfloat16, torch.float16]) +@torch.no_grad() +def test_fused_modulation_preserves_intermediate_rounding(dtype): + # Mean zero and variance five eliminate reduction-order differences, so + # this checks the eager rounding between multiplication and addition. + x = ( + torch.tensor([-3.0, -1.0, 1.0, 3.0], device="cuda", dtype=dtype) + .repeat(128) + .reshape(1, 1, 512) + .expand(1, 4, 512) + .contiguous() + ) + scale = torch.linspace(-0.8, 0.9, 512, device="cuda").to(dtype).reshape(1, 1, 512) + shift = torch.linspace(0.4, -0.3, 512, device="cuda").to(dtype).reshape(1, 1, 512) + module = make_module(512, False, 0.0, "cuda", dtype) + actual = _adaln_fused_forward(module, x, scale, shift) + assert actual is not None + torch.testing.assert_close(actual, module.forward_native(x, scale, shift), rtol=0, atol=0) + + @pytest.mark.parametrize("dtype", DTYPES) def test_fused_noncontiguous_x(dtype): # Non-contiguous x routes to the forward_native fallback (the fused kernel diff --git a/tests/diffusion/model_loader/test_diffusers_loader.py b/tests/diffusion/model_loader/test_diffusers_loader.py index 7a7acb7b468..ab135179c11 100644 --- a/tests/diffusion/model_loader/test_diffusers_loader.py +++ b/tests/diffusion/model_loader/test_diffusers_loader.py @@ -2181,6 +2181,8 @@ def load_weights(self, weights): if model.load_called: with open(os.path.join(temp_dir, f"rank_{rank}_success.flag"), "w") as f: f.write("ok") + # Keep successful ranks alive until every peer has finished initialization/loading. + dist.barrier() finally: dist.destroy_process_group() diff --git a/tests/diffusion/models/minimax_h3/test_minimax_h3_regenerate.py b/tests/diffusion/models/minimax_h3/test_minimax_h3_regenerate.py index 4daca07a9d7..ccabb7e7784 100644 --- a/tests/diffusion/models/minimax_h3/test_minimax_h3_regenerate.py +++ b/tests/diffusion/models/minimax_h3/test_minimax_h3_regenerate.py @@ -80,6 +80,9 @@ def test_adapter_fuses_every_lora_into_the_native_layout(tmp_path): lora = _write_bundle(tmp_path) adapter = _load(tmp_path) assert adapter is not None + # This CPU test reconstructs real deltas on the host. + # Production fusion selects the serving accelerator for checkpoint loading. + adapter.fusion._device = torch.device("cpu") def delta(prefix, target): return lora[f"{prefix}{target}.lora_B.weight"] @ lora[f"{prefix}{target}.lora_A.weight"] diff --git a/tests/diffusion/models/minimax_h3/test_minimax_h3_vae_ops.py b/tests/diffusion/models/minimax_h3/test_minimax_h3_vae_ops.py index c46ca0a4d88..1cd407e66ed 100644 --- a/tests/diffusion/models/minimax_h3/test_minimax_h3_vae_ops.py +++ b/tests/diffusion/models/minimax_h3/test_minimax_h3_vae_ops.py @@ -11,7 +11,11 @@ from tests.helpers.mark import hardware_marks from vllm_omni.platforms import current_omni_platform -pytestmark = [pytest.mark.core_model, pytest.mark.diffusion, *hardware_marks(res={"cuda": "L4"}, num_cards=1)] +pytestmark = [ + pytest.mark.core_model, + pytest.mark.diffusion, + *hardware_marks(res={"cuda": ["L4", "H100", "B200"]}, num_cards=1), +] def _selected_operators(): diff --git a/tests/diffusion/models/pi0/test_pi0_units.py b/tests/diffusion/models/pi0/test_pi0_units.py index cae81cd2855..4bcbdc14c51 100644 --- a/tests/diffusion/models/pi0/test_pi0_units.py +++ b/tests/diffusion/models/pi0/test_pi0_units.py @@ -1,5 +1,5 @@ # SPDX-License-Identifier: Apache-2.0 -# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project """CPU unit tests for the π0 (Pi-Zero) VLA model. Deliberately lightweight: pure utility code (config plumbing, mask building, @@ -546,6 +546,7 @@ def _selfconsist_model(device): model = Pi0ForActionPrediction( Pi0Config(max_action_dim=32, max_state_dim=32, chunk_size=50, num_inference_steps=10, dtype="float32") ) + assert _SELFCONSIST_CKPT is not None state = safetensors.torch.load_file(os.path.join(_SELFCONSIST_CKPT, "model.safetensors")) filled = model.load_weights(list(state.items())) expected = {n for n, _ in model.named_parameters() if "rotary_emb" not in n and not n.endswith(".inv_freq")} diff --git a/tests/diffusion/models/qwen_image/test_qwen_image_scheduler_device.py b/tests/diffusion/models/qwen_image/test_qwen_image_scheduler_device.py new file mode 100644 index 00000000000..b76dbc4243c --- /dev/null +++ b/tests/diffusion/models/qwen_image/test_qwen_image_scheduler_device.py @@ -0,0 +1,60 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project + +"""Weight-free consumer check for Qwen-Image's device-local Euler schedule.""" + +import numpy as np +import pytest +import torch +from diffusers.schedulers.scheduling_flow_match_euler_discrete import FlowMatchEulerDiscreteScheduler +from torch import nn +from vllm.platforms import current_platform + +from tests.helpers.mark import hardware_marks, hardware_test +from vllm_omni.diffusion.models.qwen_image.pipeline_qwen_image import QwenImagePipeline +from vllm_omni.diffusion.models.qwen_image.pipeline_qwen_image_edit import QwenImageEditPipeline +from vllm_omni.diffusion.models.qwen_image.pipeline_qwen_image_edit_plus import QwenImageEditPlusPipeline + +pytestmark = [pytest.mark.core_model, pytest.mark.diffusion, *hardware_marks(res={"cuda": "L4"}, num_cards=1)] + + +@hardware_test(res={"cuda": "L4"}, num_cards=1) +@pytest.mark.parametrize("pipeline_cls", [QwenImagePipeline, QwenImageEditPipeline, QwenImageEditPlusPipeline]) +@pytest.mark.parametrize("dtype", [torch.bfloat16, torch.float16]) +@torch.inference_mode() +def test_qwen_image_euler_steps_match_device_local_diffusers(pipeline_cls, dtype): + # Qwen/Qwen-Image scheduler settings; no checkpoint or model weights. + config = { + "num_train_timesteps": 1000, + "shift": 1.0, + "use_dynamic_shifting": True, + "base_shift": 0.5, + "max_shift": 0.9, + "base_image_seq_len": 256, + "max_image_seq_len": 8192, + "shift_terminal": 0.02, + } + device = torch.device(current_platform.device_type) + pipeline = pipeline_cls.__new__(pipeline_cls) + nn.Module.__init__(pipeline) + pipeline.device = device + pipeline.scheduler = FlowMatchEulerDiscreteScheduler(**config) + timesteps, steps = pipeline.prepare_timesteps(20, sigmas=None, image_seq_len=1024) + + # Build the Diffusers pipeline's device-local schedule independently, then + # consume both schedules with the same BF16 predictions and latent state. + reference = FlowMatchEulerDiscreteScheduler(**config) + slope = (config["max_shift"] - config["base_shift"]) / (config["max_image_seq_len"] - config["base_image_seq_len"]) + mu = config["base_shift"] + (1024 - config["base_image_seq_len"]) * slope + reference.set_timesteps(20, sigmas=np.linspace(1.0, 1 / 20, 20), mu=mu, device=device) + pipeline.scheduler.set_begin_index(0) + reference.set_begin_index(0) + prediction = (torch.arange(1, 65, device=device, dtype=torch.float32) / 64).to(dtype) + prediction = prediction.view(1, 1, 64).expand(1, 1024, 64).contiguous() + actual = torch.zeros_like(prediction) + expected = actual.clone() + assert steps == 20 + for actual_t, reference_t in zip(timesteps, reference.timesteps): + actual = pipeline.scheduler.step(prediction, actual_t, actual, return_dict=False)[0] + expected = reference.step(prediction, reference_t, expected, return_dict=False)[0] + torch.testing.assert_close(actual, expected, rtol=0, atol=0) diff --git a/tests/diffusion/models/sana_video/test_transformer_sana_video_sp.py b/tests/diffusion/models/sana_video/test_transformer_sana_video_sp.py index f41b66fe1b9..4c1af3f8d69 100644 --- a/tests/diffusion/models/sana_video/test_transformer_sana_video_sp.py +++ b/tests/diffusion/models/sana_video/test_transformer_sana_video_sp.py @@ -422,11 +422,15 @@ def __init__(self, world_size: int, tls: threading.local) -> None: def all_reduce(self, tensor: torch.Tensor) -> torch.Tensor: self.slots[self.tls.rank] = tensor self.barrier.wait(timeout=60) - total = self.slots[0].clone() - for part in self.slots[1:]: + parts: list[torch.Tensor] = [] + for part in self.slots: + assert part is not None + parts.append(part) + total = parts[0].clone() + for part in parts[1:]: total = total + part if self.tls.rank == 0: - self.reduce_inputs.append([part.clone() for part in self.slots]) + self.reduce_inputs.append([part.clone() for part in parts]) self.barrier.wait(timeout=60) return total @@ -434,8 +438,11 @@ def all_gather(self, tensor: torch.Tensor, dim: int = 0, separate_tensors: bool assert separate_tensors self.slots[self.tls.rank] = tensor self.barrier.wait(timeout=60) - assert all(part.shape == tensor.shape for part in self.slots), "all_gather requires equal shards" - parts = [part.clone() for part in self.slots] + parts: list[torch.Tensor] = [] + for part in self.slots: + assert part is not None + assert part.shape == tensor.shape, "all_gather requires equal shards" + parts.append(part.clone()) if self.tls.rank == 0: with self.lock: self.gather_calls += 1 @@ -535,6 +542,7 @@ def run_rank(rank: int) -> torch.Tensor: outs = list(pool.map(run_rank, range(world_size))) num_layers = _TINY_SP_CONFIG["num_layers"] + assert isinstance(num_layers, int) # One attn1 state reduce per block, its inputs distinct per rank (the ranks # really processed different frames), one neighbour halo exchange per block # per rank, and only the final frame gather is a full collective. diff --git a/tests/diffusion/offloader/test_distributed_layerwise_backend.py b/tests/diffusion/offloader/test_distributed_layerwise_backend.py index b3ac19baefe..4362395bc80 100644 --- a/tests/diffusion/offloader/test_distributed_layerwise_backend.py +++ b/tests/diffusion/offloader/test_distributed_layerwise_backend.py @@ -93,7 +93,7 @@ def test_buffer_only_block_reports_offloaded_state(self, patched_offload_runtime hook.offload_layer() assert not hook.is_materialized - def test_initialize_failure_keeps_next_block_materialized(self, monkeypatch): + def test_initialize_failure_keeps_next_block_materialized(self, monkeypatch, patched_offload_runtime): current_block = nn.Linear(2, 2) next_block = nn.Linear(2, 2) expected = {name: tensor.detach().clone() for name, tensor in next_block.state_dict().items()} @@ -2931,7 +2931,9 @@ def fail_second_hook(*args, **kwargs): assert registry is None or registry.get_hook("distributed_layerwise_offload") is None torch.testing.assert_close(block.weight, expected) - def test_multirank_enable_failure_cleanup_skips_restore_collective(self, monkeypatch, mocker): + def test_multirank_enable_failure_cleanup_skips_restore_collective( + self, monkeypatch, mocker, patched_offload_runtime + ): backend = DistributedLayerwiseOffloadBackend( OffloadConfig( strategy=OffloadStrategy.DISTRIBUTED_LAYER_WISE, diff --git a/tests/diffusion/quantization/test_mxfp4_config.py b/tests/diffusion/quantization/test_mxfp4_config.py index 0193b23cb45..deefd7a3d6d 100644 --- a/tests/diffusion/quantization/test_mxfp4_config.py +++ b/tests/diffusion/quantization/test_mxfp4_config.py @@ -1112,7 +1112,9 @@ def pack(x, dtype): ) cls = mxfp4_config.NPUMxfp4OnlineLinearMethod if online else mxfp4_config.NPUMxfp4LinearMethod method = cls(config, prefix="blocks.0.attn1.to_qkv") - layer = torch.nn.Linear(64, 2, bias=False, dtype=torch.bfloat16) + # Match create_weights' inference parameters: native replacement preserves + # the existing parameter's requires_grad flag when packing integer weights. + layer = torch.nn.Linear(64, 2, bias=False, dtype=torch.bfloat16).requires_grad_(False) layer.orig_dtype = torch.bfloat16 if not online: layer.register_parameter("weight_scale", torch.nn.Parameter(torch.full((2, 2), 127, dtype=torch.uint8), False)) diff --git a/tests/e2e/online_serving/test_duplex_client_live.py b/tests/e2e/online_serving/test_duplex_client_live.py index abf83727a76..083b541ebdd 100644 --- a/tests/e2e/online_serving/test_duplex_client_live.py +++ b/tests/e2e/online_serving/test_duplex_client_live.py @@ -6,7 +6,9 @@ :class:`DuplexClient`: session handshake with reference audio, paced PCM streaming, a speak response consumed via :class:`ResponseHandle` with incremental playback acks, a forced transport drop recovered by automatic -``session.resume``, and a clean close. +``session.resume``, and a clean close. The live server needs real weights (``--run-level +advanced_model``): random weights may never emit the turn terminator that +completes the response. The assertions remain basic client lifecycle checks. """ from __future__ import annotations diff --git a/tests/engine/test_duplex_import_boundary.py b/tests/engine/test_duplex_import_boundary.py index 7b44fc61cf4..77a1bcb4ef1 100644 --- a/tests/engine/test_duplex_import_boundary.py +++ b/tests/engine/test_duplex_import_boundary.py @@ -21,8 +21,8 @@ def _assert_isolated_import_succeeds(script: str) -> None: # unrelated, inconsistent device-selection environment prevent the child # interpreter from reaching the boundary assertion. env = os.environ.copy() - env.pop("CUDA_VISIBLE_DEVICES", None) - env.pop("HIP_VISIBLE_DEVICES", None) + env["CUDA_VISIBLE_DEVICES"] = "" + env["HIP_VISIBLE_DEVICES"] = "" result = subprocess.run( [sys.executable, "-c", script], cwd=REPO_ROOT, diff --git a/tests/engine/test_orchestrator.py b/tests/engine/test_orchestrator.py index 4d5a85dfce6..91988e49a62 100644 --- a/tests/engine/test_orchestrator.py +++ b/tests/engine/test_orchestrator.py @@ -19,7 +19,7 @@ import pytest import torch from vllm.lora.request import LoRARequest -from vllm.outputs import CompletionOutput, RequestOutput +from vllm.outputs import CompletionOutput, RequestError, RequestOutput from vllm.sampling_params import SamplingParams from vllm.v1.engine import EngineCoreOutput, EngineCoreOutputs, EngineCoreRequest, FinishReason from vllm.v1.engine.exceptions import EngineDeadError @@ -258,6 +258,7 @@ async def collective_rpc_async( class FakeOutputProcessor: def __init__(self, *, request_outputs: list[Any] | None = None) -> None: self.request_outputs = list(request_outputs or []) + self.request_states: dict[str, object] = {} self.add_request_calls: list[tuple[tuple[Any, ...], dict[str, Any]]] = [] self.abort_calls: list[list[str]] = [] @@ -1285,7 +1286,7 @@ async def test_request_watermarking(stage_type: str, watermarking: bool | None) else: await pool.process_llm_raw_outputs( 0, - SimpleNamespace(outputs=[output], timestamp=1.0, scheduler_stats=None), + OmniEngineCoreOutputs(outputs=[output], timestamp=1.0, scheduler_stats=None, weight_version=None), ) # Then ensure the watermarking behavior is as expected @@ -1350,7 +1351,7 @@ async def test_strict_watermark_failure_replaces_llm_output_with_error() -> None ] outputs = await pool.process_llm_raw_outputs( - 0, SimpleNamespace(outputs=raw_outputs, timestamp=1.0, scheduler_stats=None) + 0, OmniEngineCoreOutputs(outputs=raw_outputs, timestamp=1.0, scheduler_stats=None, weight_version=None) ) good_output, bad_output = outputs @@ -3344,6 +3345,7 @@ async def test_orchestrator_records_scheduler_stats_without_outputs(orchestrator timestamp=1.0, scheduler_stats=scheduler_stats, finished_requests=None, + weight_version=None, ) ) @@ -3397,6 +3399,7 @@ async def test_orchestrator_does_not_build_iteration_stats_for_finished_only_bat timestamp=1.0, scheduler_stats=None, finished_requests={"req-finished"}, + weight_version=None, ) ) @@ -3688,6 +3691,85 @@ async def test_duplex_session_request_error_finish_is_delivered_as_request_error assert output_queue.empty() +_PAUSE_MESSAGE = "Generation is paused. Please retry after it resumes." + + +@pytest.mark.asyncio +async def test_raw_stage_error_prefers_request_error_message(mocker): + pools = _build_stage_pools([[FakeStageClient()], [FakeStageClient()]]) + for pool in pools: + mocker.patch.object(pool, "process_llm_raw_outputs", new_callable=mocker.AsyncMock, return_value=[]) + output_queue: asyncio.Queue[ErrorMessage] = asyncio.Queue() + orchestrator = Orchestrator( + request_async_queue=asyncio.Queue(), + output_async_queue=output_queue, + rpc_async_queue=asyncio.Queue(), + stage_pools=[], + ) + orchestrator.stage_pools = pools + orchestrator.request_states["stuck"] = OrchestratorRequestState(request_id="stuck", final_stage_id=1) + mocker.patch.object(orchestrator, "_handle_kv_ready_raw_outputs", new_callable=mocker.AsyncMock) + mocker.patch.object(orchestrator, "_cleanup_request_ids", new_callable=mocker.AsyncMock) + raw = EngineCoreOutputs( + outputs=[ + OmniEngineCoreOutput( + request_id="stuck", + new_token_ids=[], + finish_reason=FinishReason.ERROR, + stop_reason=None, + error=RequestError(code="engine_paused", message=_PAUSE_MESSAGE, retryable=True), + ) + ] + ) + + await orchestrator._process_llm_stage_outputs(0, 0, raw, set()) + + error = output_queue.get_nowait() + assert error.error == _PAUSE_MESSAGE + assert error.error_type == "engine_paused" + assert error.request_id == "stuck" + assert output_queue.empty() + + +@pytest.mark.asyncio +async def test_duplex_session_error_is_not_emitted_twice(mocker): + pools = _build_stage_pools([[FakeStageClient()]]) + pause = OmniRequestOutput.from_error("duplex-req", _PAUSE_MESSAGE, error_type="engine_paused") + mocker.patch.object(pools[0], "process_llm_raw_outputs", new_callable=mocker.AsyncMock, return_value=[pause]) + output_queue: asyncio.Queue[ErrorMessage] = asyncio.Queue() + orchestrator = Orchestrator( + request_async_queue=asyncio.Queue(), + output_async_queue=output_queue, + rpc_async_queue=asyncio.Queue(), + stage_pools=[], + ) + orchestrator.stage_pools = pools + orchestrator.request_states["duplex-req"] = OrchestratorRequestState( + request_id="duplex-req", + session_owned=True, + ) + mocker.patch.object(orchestrator, "_handle_kv_ready_raw_outputs", new_callable=mocker.AsyncMock) + raw = EngineCoreOutputs( + outputs=[ + OmniEngineCoreOutput( + request_id="duplex-req", + new_token_ids=[], + finish_reason=FinishReason.ERROR, + error=RequestError(code="engine_paused", message=_PAUSE_MESSAGE, retryable=True), + ) + ] + ) + + processed = await orchestrator._process_llm_stage_outputs(0, 0, raw, set()) + + assert processed == [] + error = output_queue.get_nowait() + assert error.error == _PAUSE_MESSAGE + assert error.error_type == "engine_paused" + assert error.request_id == "duplex-req" + assert output_queue.empty() + + @pytest.mark.asyncio async def test_interaction_wait_does_not_block_playback_or_abort(monkeypatch: pytest.MonkeyPatch) -> None: requests: asyncio.Queue = asyncio.Queue() diff --git a/tests/engine/test_output_processor.py b/tests/engine/test_output_processor.py index 231240e23f0..221797c3ff9 100644 --- a/tests/engine/test_output_processor.py +++ b/tests/engine/test_output_processor.py @@ -10,7 +10,7 @@ import numpy as np import pytest import torch -from vllm.outputs import PoolingRequestOutput, RequestOutput +from vllm.outputs import PoolingRequestOutput, RequestError, RequestOutput from vllm.pooling_params import PoolingParams from vllm.sampling_params import RequestOutputKind from vllm.v1.engine import EngineCoreEvent, EngineCoreEventType, FinishReason @@ -890,6 +890,41 @@ def test_pooling_completion_uses_upstream_output_and_finishes_request(cache_miss assert request.external_req_id not in processor.external_req_ids +def test_admission_error_returns_external_pause_message(): + """A queue=None generation request must surface RequestError, not a chunk.""" + kwargs = dict(_NO_DETOK_STATE_KWARGS) + kwargs["request_id"] = "internal" + kwargs["external_req_id"] = "ext" + state = OmniRequestState(**kwargs, output_kind=RequestOutputKind.CUMULATIVE, stream_input=True) + processor = MultimodalOutputProcessor(tokenizer=None, log_stats=False, engine_core_output_type="audio") + processor.request_states[state.request_id] = state + processor.external_req_ids[state.external_req_id].append(state.request_id) + pause = "Generation is paused. Please retry after it resumes." + + processed = processor.process_outputs( + [ + OmniEngineCoreOutput( + request_id="internal", + new_token_ids=[], + finish_reason=FinishReason.ERROR, + error=RequestError(code="engine_paused", message=pause, retryable=True), + multimodal_output={"audio": torch.tensor([1.0, 2.0])}, + ) + ] + ) + + assert processed.reqs_to_abort == ["internal"] + assert len(processed.request_outputs) == 1 + output = processed.request_outputs[0] + assert output.request_id == "ext" + assert output.error == pause + assert output.error_type == "engine_paused" + assert output.finished is True + assert output.outputs == [] + assert state.mm_accumulated.is_empty + assert "internal" not in processor.request_states + + def test_no_detokenizer_generation_cache_miss_keeps_multimodal_completion(): """A generation-stage error must retain its completion type and payload.""" processor = MultimodalOutputProcessor(tokenizer=None, log_stats=False, output_modality=OutputModality.AUDIO) @@ -950,6 +985,7 @@ def test_no_detokenizer_process_outputs_returns_nonterminal_audio_chunk(monkeypa ec_transfer_params=None, routed_experts=None, num_cached_tokens=0, + error=None, ) result = processor.process_outputs([engine_output]) @@ -997,6 +1033,7 @@ def _audio_engine_output(*, is_segment_finished: bool, is_last_chunk: bool): routed_experts=None, num_cached_tokens=0, is_segment_finished=is_segment_finished, + error=None, ) diff --git a/tests/engine/test_output_processor_mrv2_text.py b/tests/engine/test_output_processor_mrv2_text.py index e3f072fbd35..3a3b397c47e 100644 --- a/tests/engine/test_output_processor_mrv2_text.py +++ b/tests/engine/test_output_processor_mrv2_text.py @@ -90,6 +90,8 @@ def _rehydrated_output(*, pooling_output=None, finish_reason=FinishReason.STOP): output_type=None, num_generation_tokens=None, finished=True, + # Upstream EngineCoreOutput.error is read before any other field. + error=None, ) diff --git a/tests/engine/test_tracing.py b/tests/engine/test_tracing.py index aa016823390..4658ee1000d 100644 --- a/tests/engine/test_tracing.py +++ b/tests/engine/test_tracing.py @@ -20,6 +20,7 @@ from opentelemetry.sdk.trace.export import SimpleSpanProcessor from opentelemetry.sdk.trace.export.in_memory_span_exporter import InMemorySpanExporter from opentelemetry.trace.propagation.tracecontext import TraceContextTextMapPropagator +from vllm.outputs import RequestError from vllm.v1.engine import EngineCoreOutputs, FinishReason from vllm.v1.engine.exceptions import EngineDeadError @@ -33,7 +34,7 @@ ) from vllm_omni.engine import OmniEngineCoreOutput from vllm_omni.engine.async_omni_engine import AsyncOmniEngine -from vllm_omni.engine.messages import StageSubmissionMessage +from vllm_omni.engine.messages import ErrorMessage, StageSubmissionMessage from vllm_omni.engine.orchestrator import Orchestrator, OrchestratorRequestState from vllm_omni.engine.stage_pool import StagePool from vllm_omni.engine.stage_runtime import StageRuntimeInfo @@ -174,6 +175,54 @@ async def test_orchestrator_closes_diffusion_request_and_stage(recording, outcom assert all("private" not in str(s.attributes) for s in spans) +@pytest.mark.asyncio +@pytest.mark.parametrize("cache_miss", [False, True], ids=["engine-paused", "mm-cache-miss"]) +async def test_raw_stage_error_trace_matches_frontend_classification(recording, mocker, cache_miss): + tracer, exporter = recording + pools = _build_stage_pools([[FakeStageClient(final_output=True)]]) + mocker.patch.object(pools[0], "process_llm_raw_outputs", new_callable=mocker.AsyncMock, return_value=[]) + output_queue: asyncio.Queue[ErrorMessage] = asyncio.Queue() + orch = Orchestrator( + request_async_queue=asyncio.Queue(), + output_async_queue=output_queue, + rpc_async_queue=asyncio.Queue(), + stage_pools=pools, + tracer=tracer, + ) + request_trace = RequestTrace(tracer, "failed", HEADERS) + request_trace.start_stage(0, 0, "llm") + orch.request_states["failed"] = OrchestratorRequestState( + request_id="failed", final_stage_id=0, sampling_params_list=[_sampling_params()], trace=request_trace + ) + mocker.patch.object(orch, "_handle_kv_ready_raw_outputs", new_callable=mocker.AsyncMock) + raw = EngineCoreOutputs( + outputs=[ + OmniEngineCoreOutput( + request_id="failed", + new_token_ids=[], + finish_reason=FinishReason.ERROR, + error=RequestError(code="engine_paused", message="Generation is paused.", retryable=True), + mm_cache_miss_hashes=["missing"] if cache_miss else None, + ) + ] + ) + + await orch._process_llm_stage_outputs(0, 0, raw, set()) + + error = output_queue.get_nowait() + expected_type = "ServiceUnavailableError" if cache_miss else "engine_paused" + assert error.error_type == expected_type + assert error.status_code == (503 if cache_miss else None) + assert error.mm_cache_miss_hashes == (["missing"] if cache_miss else None) + assert output_queue.empty() + assert not orch.request_states + spans = exporter.get_finished_spans() + assert {span.name for span in spans} == {"omni.pipeline", "omni.stage"} + assert len(spans) == 2 + assert all(span.attributes["error.type"] == expected_type for span in spans) + assert all(span.status.status_code == trace.StatusCode.ERROR for span in spans) + + @pytest.mark.asyncio @pytest.mark.parametrize("error_type", ["invalid_request_error", None]) async def test_stage_error_preserves_trace_error_type(recording, error_type): diff --git a/tests/entrypoints/duplex/test_duplex_serving.py b/tests/entrypoints/duplex/test_duplex_serving.py index c3753afa10a..bb3715c90a1 100644 --- a/tests/entrypoints/duplex/test_duplex_serving.py +++ b/tests/entrypoints/duplex/test_duplex_serving.py @@ -13,7 +13,7 @@ from typing import Any import pytest -from fastapi import WebSocketDisconnect +from fastapi import FastAPI, WebSocketDisconnect from vllm_omni.config.stage_config import DuplexSessionRuntimeConfig from vllm_omni.engine.duplex import commands @@ -39,6 +39,7 @@ class FakeWebSocket: def __init__(self, query: dict[str, str] | None = None) -> None: + self.app = FastAPI() self.query_params = dict(query or {}) self.sent: list[dict[str, Any]] = [] self.accepted = False @@ -265,6 +266,38 @@ async def _open( # --------------------------------------------------------------------------- # +@pytest.mark.asyncio +@pytest.mark.parametrize("route_name", ["realtime_websocket", "duplex_websocket"]) +async def test_transport_accepts_before_warmup_but_session_open_waits(route_name: str) -> None: + from vllm_omni.entrypoints.openai import api_server + + accepted = asyncio.Event() + + class HandshakeSocket(FakeWebSocket): + async def accept(self) -> None: + await super().accept() + accepted.set() + + omni = FakeOmni() + ws = HandshakeSocket({"duplex": "1", "autostart": "0"}) + warmup_done = asyncio.Event() + ws.app.state.duplex_warmup_done = warmup_done + ws.app.state.openai_serving_duplex = _handler(omni) + ws.feed(_session_update()) + task = asyncio.create_task(getattr(api_server, route_name)(ws)) + try: + await asyncio.wait_for(accepted.wait(), timeout=1.0) + assert not omni.handles + assert not ws.sent + warmup_done.set() + await ws.wait_for("session.created") + assert len(omni.handles) == 1 + finally: + warmup_done.set() + ws.disconnect() + await asyncio.wait_for(task, timeout=2.0) + + @pytest.mark.asyncio async def test_session_update_opens_a_session_and_announces_the_server_id() -> None: omni = FakeOmni() diff --git a/tests/entrypoints/openai_api/test_api_server_guards.py b/tests/entrypoints/openai_api/test_api_server_guards.py index f7cf8e56a7a..5a4dee6cbad 100644 --- a/tests/entrypoints/openai_api/test_api_server_guards.py +++ b/tests/entrypoints/openai_api/test_api_server_guards.py @@ -1504,6 +1504,7 @@ async def warmup(self): monkeypatch.setattr(api_server, "process_lora_modules", lambda modules, _defaults: modules or []) monkeypatch.setattr(api_server, "OpenAIServingModels", _FakeModels) monkeypatch.setattr(api_server, "OnlineRenderer", _FakeCtor) + monkeypatch.setattr(api_server, "OnlineDerenderer", _FakeCtor) monkeypatch.setattr(api_server, "OpenAIServingResponses", _FakeCtor) monkeypatch.setattr(api_server, "OmniOpenAIServingChat", _FakeCtor) monkeypatch.setattr(api_server, "OmniOpenAIServingChatBatch", _FakeCtor) diff --git a/tests/entrypoints/openai_api/test_duplex_api_server.py b/tests/entrypoints/openai_api/test_duplex_api_server.py index 4962b403b2f..f532096cf6f 100644 --- a/tests/entrypoints/openai_api/test_duplex_api_server.py +++ b/tests/entrypoints/openai_api/test_duplex_api_server.py @@ -386,6 +386,7 @@ async def test_duplex_app_state_wires_only_the_session_surfaces(monkeypatch, moc """ monkeypatch.setattr(api_server, "OpenAIServingModels", _FakeModels) renderer_ctor = mocker.patch.object(api_server, "OnlineRenderer") + derenderer_ctor = mocker.patch.object(api_server, "OnlineDerenderer") monkeypatch.setattr(api_server, "OmniOpenAIServingChat", lambda **kwargs: SimpleNamespace(**kwargs)) engine = _FakeDuplexOmni() state = State() @@ -401,7 +402,9 @@ async def test_duplex_app_state_wires_only_the_session_surfaces(monkeypatch, moc assert isinstance(state.openai_serving_duplex, OmniDuplexSessionHandler) assert state.openai_serving_chat is not None renderer_ctor.assert_called_once() + derenderer_ctor.assert_called_once() assert renderer_ctor.call_args.kwargs["log_error_stack"] is log_error_stack + assert derenderer_ctor.call_args.kwargs["log_error_stack"] is log_error_stack assert state.openai_serving_chat.online_renderer is state.online_renderer is renderer_ctor.return_value state.online_renderer.warmup.assert_called_once_with() assert state.engine_client is engine @@ -583,6 +586,7 @@ async def test_duplex_chat_initialization_respects_model_gates( """Neither gate may construct or warm up a chat renderer.""" monkeypatch.setattr(api_server, "OpenAIServingModels", _FakeModels) renderer_ctor = mocker.patch.object(api_server, "OnlineRenderer") + derenderer_ctor = mocker.patch.object(api_server, "OnlineDerenderer") chat_ctor = mocker.patch.object(api_server, "OmniOpenAIServingChat") engine = _FakeDuplexOmni() engine._capabilities = DuplexCapabilities(supports_chat_completions=supports_chat) @@ -594,6 +598,7 @@ async def test_duplex_chat_initialization_respects_model_gates( assert state.openai_serving_chat is None assert state.openai_serving_chat_batch is None renderer_ctor.assert_not_called() + derenderer_ctor.assert_not_called() chat_ctor.assert_not_called() assert isinstance(state.openai_serving_duplex, OmniDuplexSessionHandler) diff --git a/tests/entrypoints/openai_api/test_serving_audio_generate.py b/tests/entrypoints/openai_api/test_serving_audio_generate.py index 6462a3a4210..ea30b3008c7 100644 --- a/tests/entrypoints/openai_api/test_serving_audio_generate.py +++ b/tests/entrypoints/openai_api/test_serving_audio_generate.py @@ -84,7 +84,12 @@ def _make_server(engine_client=None): @pytest.fixture -def test_app(): +def test_app(monkeypatch): + # Exercise seeded request wiring with a CPU generator in this CPU API test. + monkeypatch.setattr( + "vllm_omni.platforms.current_omni_platform", + SimpleNamespace(device_type="cpu"), + ) server = _make_server() original_fn = server.create_audio_generate @@ -95,7 +100,7 @@ def test_app(): async def patched_create_audio_generate(*args, **kwargs): return await original_fn(*args, **kwargs) - patched_create_audio_generate.__signature__ = new_sig + setattr(patched_create_audio_generate, "__signature__", new_sig) server.create_audio_generate = patched_create_audio_generate app = FastAPI() @@ -641,7 +646,11 @@ async def gen_runtime_error(*args, **kwargs): @pytest.mark.parametrize( ("exc", "expected_message", "expected_stage_id"), [ - (OmniEngineDeadError("engine dead", error_stage_id=2), "engine dead", 2), + ( + OmniEngineDeadError("engine dead", error_stage_id=2), + OmniEngineDeadError.CLIENT_MESSAGE, + 2, + ), (EngineGenerateError("engine generate failed"), "engine generate failed", None), ], ) @@ -693,7 +702,7 @@ def test_with_all_params(self, client): assert response.headers["content-type"] == "audio/wav" def test_missing_input_rejected(self, client): - payload = {} + payload: dict[str, object] = {} response = client.post("/v1/audio/generate", json=payload) assert response.status_code == 422 diff --git a/tests/entrypoints/openai_api/test_serving_speech.py b/tests/entrypoints/openai_api/test_serving_speech.py index 336069e74f1..7c357a257c0 100644 --- a/tests/entrypoints/openai_api/test_serving_speech.py +++ b/tests/entrypoints/openai_api/test_serving_speech.py @@ -5184,7 +5184,8 @@ def test_api_server_create_speech_engine_error_response_includes_request_and_sta assert isinstance(response, JSONResponse) assert response.status_code == 500 assert response.body.decode("utf-8") == ( - '{"error":{"message":"engine dead","type":"InternalServerError","param":null,' + '{"error":{"message":"The server had an error while processing your request.",' + '"type":"InternalServerError","param":null,' '"code":500,"request_id":"speech-req-1","error_stage_id":1}}' ) terminate_mock.assert_called_once() diff --git a/tests/helpers/cudnn.py b/tests/helpers/cudnn.py new file mode 100644 index 00000000000..9e1defd5312 --- /dev/null +++ b/tests/helpers/cudnn.py @@ -0,0 +1,41 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project +"""cuDNN numeric policy for CUDA-graph parity tests. + +Graph replay and eager convolution can disagree by about one ulp when TF32 +or cuDNN benchmark autotune is left on. These helpers set the conv-level +flag (not the backend-wide one) and restore the previous values. +""" + +from __future__ import annotations + +from collections.abc import Iterator +from contextlib import contextmanager + +import torch + + +@contextmanager +def ieee_cudnn_fp32() -> Iterator[None]: + """Force IEEE FP32 conv math, then restore the previous conv precision.""" + conv = torch.backends.cudnn.conv + previous = conv.fp32_precision + conv.fp32_precision = "ieee" + try: + yield + finally: + conv.fp32_precision = previous + + +@contextmanager +def deterministic_cudnn() -> Iterator[None]: + """Turn benchmark off and deterministic algorithms on, then restore both.""" + benchmark = torch.backends.cudnn.benchmark + deterministic = torch.backends.cudnn.deterministic + torch.backends.cudnn.benchmark = False + torch.backends.cudnn.deterministic = True + try: + yield + finally: + torch.backends.cudnn.benchmark = benchmark + torch.backends.cudnn.deterministic = deterministic diff --git a/tests/helpers/omni_scheduler.py b/tests/helpers/omni_scheduler.py index e5a4c58a550..018da1356bd 100644 --- a/tests/helpers/omni_scheduler.py +++ b/tests/helpers/omni_scheduler.py @@ -7,9 +7,33 @@ from types import MethodType +import torch +import vllm.v1.core.single_type_kv_cache_manager as native_kv_managers +from vllm.v1.core.kv_cache_manager import KVCacheManager +from vllm.v1.kv_cache_interface import FullAttentionSpec, KVCacheConfig, KVCacheGroupSpec + from vllm_omni.core.sched.omni_ar_scheduler import OmniARScheduler +def make_kv_cache_manager(block_size: int) -> KVCacheManager: + """Build a small real manager for scheduler-side CPU connector tests.""" + # A foreign spec can populate the registry before native specs are loaded. + native_kv_managers.register_all_kvcache_specs(None) + spec = FullAttentionSpec(block_size=block_size, num_kv_heads=1, head_size=8, dtype=torch.float16) + config = KVCacheConfig( + num_blocks=8, + kv_cache_tensors=[], + kv_cache_groups=[KVCacheGroupSpec(["layer0"], spec)], + ) + return KVCacheManager( + config, + max_model_len=32, + scheduler_block_size=block_size, + hash_block_size=block_size, + enable_caching=False, + ) + + def bind_omits_transfer_helpers(sched) -> None: """Bind the real omit helpers so a MagicMock scheduler does not skip save_async. diff --git a/tests/metrics/test_stat_logger.py b/tests/metrics/test_stat_logger.py index 627f5aafc2a..b98c15743b6 100644 --- a/tests/metrics/test_stat_logger.py +++ b/tests/metrics/test_stat_logger.py @@ -18,8 +18,6 @@ from vllm.v1.metrics.stats import IterationStats, SchedulerStats from vllm.v1.spec_decode.metrics import SpecDecodingStats -from vllm_omni.engine.orchestrator import Orchestrator -from vllm_omni.engine.stage_pool import StagePool from vllm_omni.metrics import definitions as defs from vllm_omni.metrics.stat_logger import ( _ENGINE_INDEX_MAP, @@ -733,6 +731,9 @@ def test_register_replica_preserves_existing_metrics(stats_config, initial_map): @pytest.mark.asyncio @pytest.mark.parametrize("initial_replicas", [0, 1]) async def test_orchestrator_records_stats_from_late_replica(stats_config, initial_replicas): + from vllm_omni.engine.orchestrator import Orchestrator + from vllm_omni.engine.stage_pool import StagePool + pool = Mock(spec=StagePool, num_replicas=initial_replicas, stage_vllm_config=stats_config) pool.available_replica_ids.return_value = list(range(initial_replicas)) pool.process_llm_raw_outputs = AsyncMock(return_value=[]) diff --git a/tests/model_executor/models/higgs_audio_v3/test_codec_snake_kernel.py b/tests/model_executor/models/higgs_audio_v3/test_codec_snake_kernel.py index 208c841aa98..50fbb2d20b8 100644 --- a/tests/model_executor/models/higgs_audio_v3/test_codec_snake_kernel.py +++ b/tests/model_executor/models/higgs_audio_v3/test_codec_snake_kernel.py @@ -8,11 +8,19 @@ import torch import torch.nn as nn +from tests.helpers.cudnn import deterministic_cudnn from tests.helpers.mark import hardware_test from vllm_omni.platforms import current_omni_platform pytestmark = [pytest.mark.core_model] + +@pytest.fixture +def with_deterministic_cudnn(): + with deterministic_cudnn(): + yield + + requires_cuda = pytest.mark.skipif(not current_omni_platform.is_cuda(), reason="requires CUDA") _DAC_CONFIG = { @@ -124,7 +132,7 @@ def test_fuse_dac_decoder_keeps_state_dict_and_falls_back_on_cpu(): @hardware_test(res={"cuda": "L4"}, num_cards=1) @pytest.mark.parametrize("dtype", [torch.bfloat16, torch.float32]) @torch.inference_mode() -def test_fused_decoder_graph_replays_changed_inputs(dtype): +def test_fused_decoder_graph_replays_changed_inputs(dtype, with_deterministic_cudnn): reference, fused = _dac_decoder_pair("cuda", dtype) x = torch.randn(2, _DAC_CONFIG["hidden_size"], 8, device="cuda").to(dtype) stream = torch.cuda.Stream() diff --git a/tests/model_executor/models/minicpmo_4_5/test_cuda_graph_wrapper.py b/tests/model_executor/models/minicpmo_4_5/test_cuda_graph_wrapper.py index cc3889135d7..6bcd05de373 100644 --- a/tests/model_executor/models/minicpmo_4_5/test_cuda_graph_wrapper.py +++ b/tests/model_executor/models/minicpmo_4_5/test_cuda_graph_wrapper.py @@ -12,6 +12,7 @@ from vllm.platforms import current_platform import vllm_omni.model_executor.models.minicpmo_4_5.cuda_graph_wrapper as wrapper_module +from tests.helpers.cudnn import ieee_cudnn_fp32 from tests.helpers.mark import hardware_marks from vllm_omni.model_executor.models.cosyvoice3.code2wav_core.hifigan import ( HiFTGenerator, @@ -2082,8 +2083,16 @@ def forward_chunk(self, x, mask, cnn_cache, att_cache, cnn_buf, att_buf): return x +@pytest.fixture +def with_ieee_cudnn_fp32(): + with ieee_cudnn_fp32(): + yield + + @pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required for graph capture") -def test_whole_euler_query_bucket_keeps_current_first_cache_layout(monkeypatch: pytest.MonkeyPatch) -> None: +def test_whole_euler_query_bucket_keeps_current_first_cache_layout( + monkeypatch: pytest.MonkeyPatch, with_ieee_cudnn_fp32 +) -> None: """A narrow padded chunk replayed on a wider capture keeps the cache aligned. The capture mask has to put the cache block after the full capture width, @@ -2287,7 +2296,9 @@ def _ragged_mask(lengths: list[int], width: int, offset: int) -> torch.Tensor: @pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required for graph capture") -def test_whole_euler_ragged_rows_match_per_row_exact_solves(monkeypatch: pytest.MonkeyPatch) -> None: +def test_whole_euler_ragged_rows_match_per_row_exact_solves( + monkeypatch: pytest.MonkeyPatch, with_ieee_cudnn_fp32 +) -> None: """One graph solves rows of different lengths exactly as each row's own unpadded solve. Chunk 1 runs through a 16-frame query bucket at width 8: its CNN cache must @@ -2358,7 +2369,9 @@ def test_whole_euler_ragged_rows_match_per_row_exact_solves(monkeypatch: pytest. @pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required for graph capture") -def test_whole_euler_mixed_histories_match_independent_streams(monkeypatch: pytest.MonkeyPatch) -> None: +def test_whole_euler_mixed_histories_match_independent_streams( + monkeypatch: pytest.MonkeyPatch, with_ieee_cudnn_fp32 +) -> None: """Different reference lengths share a graph without exposing padded history.""" pool = torch.cuda.graph_pool_handle() monkeypatch.setattr(current_platform, "get_global_graph_pool", lambda: pool) @@ -2462,7 +2475,9 @@ def test_whole_euler_mixed_histories_do_not_enable_npu_arena(monkeypatch: pytest @pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required for graph capture") -def test_whole_euler_mixed_prompts_reuse_caches_after_streaming_trim(monkeypatch: pytest.MonkeyPatch) -> None: +def test_whole_euler_mixed_prompts_reuse_caches_after_streaming_trim( + monkeypatch: pytest.MonkeyPatch, with_ieee_cudnn_fp32 +) -> None: """Mixed prompt lengths retain exact per-row caches through steady-state trimming.""" pool = torch.cuda.graph_pool_handle() monkeypatch.setattr(current_platform, "get_global_graph_pool", lambda: pool) @@ -2547,7 +2562,9 @@ def test_whole_euler_mixed_prompts_reuse_caches_after_streaming_trim(monkeypatch @pytest.mark.skipif(not torch.cuda.is_available(), reason="CUDA required for graph capture") -def test_whole_euler_padded_ragged_replay_after_a_wider_replay(monkeypatch: pytest.MonkeyPatch) -> None: +def test_whole_euler_padded_ragged_replay_after_a_wider_replay( + monkeypatch: pytest.MonkeyPatch, with_ieee_cudnn_fp32 +) -> None: """A padded ragged replay is exact after a wider replay of the same graph batch. Every graph of one batch size shares the lengths buffer. The 16-frame diff --git a/tests/model_executor/models/moss_tts/test_codec_cudagraph.py b/tests/model_executor/models/moss_tts/test_codec_cudagraph.py index 5bbf738f5c1..65a6451461f 100644 --- a/tests/model_executor/models/moss_tts/test_codec_cudagraph.py +++ b/tests/model_executor/models/moss_tts/test_codec_cudagraph.py @@ -12,6 +12,7 @@ import torch import torch.nn as nn +from tests.helpers.cudnn import ieee_cudnn_fp32 from tests.helpers.mark import hardware_marks from vllm_omni.model_executor.models.moss_tts.audio_tokenizer import ( MossAudioTokenizerDecoderOutput, @@ -24,6 +25,14 @@ pytest.mark.core_model, *hardware_marks(res={"cuda": "L4", "rocm": "MI325"}, num_cards=1), ] + + +@pytest.fixture(scope="module", autouse=True) +def _ieee_cudnn_fp32(): + with ieee_cudnn_fp32(): + yield + + DEVICE = torch.device("cuda:0") NUM_QUANTIZERS = 8 DOWNSAMPLE_RATE = 4 # synthetic; real checkpoint uses 1920 @@ -95,13 +104,13 @@ def batch_decode( @pytest.fixture(scope="module") -def model(): +def model(_ieee_cudnn_fp32): torch.manual_seed(0) return SyntheticCodecModel().to(DEVICE).eval() @pytest.fixture(scope="module") -def wrapper(model): +def wrapper(model, _ieee_cudnn_fp32): w = MossTTSCUDAGraphCodecWrapper( model=model, capture_sizes=[25, 50, 100], diff --git a/tests/model_executor/models/moss_tts/test_shared_reference_encoder.py b/tests/model_executor/models/moss_tts/test_shared_reference_encoder.py index 2ccddd1a5b3..a699a3ad630 100644 --- a/tests/model_executor/models/moss_tts/test_shared_reference_encoder.py +++ b/tests/model_executor/models/moss_tts/test_shared_reference_encoder.py @@ -113,7 +113,13 @@ def test_client_round_trip_and_per_clip_errors(host_dir): def test_host_merges_clips_from_concurrent_clients(host_dir): gate = threading.Event() recorder = _Recorder(gate) - host = shared.SharedReferenceEncoderHost(host_dir, [(recorder, None)]) + entered = threading.Event() + + def encode(wavs): + entered.set() + return recorder(wavs) + + host = shared.SharedReferenceEncoderHost(host_dir, [(encode, None)]) try: results = {} @@ -123,16 +129,20 @@ def send(name, value): # Encoding is held until both clients' clips are queued. warm = threading.Thread(target=lambda: host.encode([torch.full((1, 10), 1.0)])) warm.start() + assert entered.wait(2), "warm singleton must enter the encoder before clients queue" threads = [threading.Thread(target=send, args=(n, v)) for n, v in (("a", 2.0), ("b", 4.0))] for thread in threads: thread.start() - threading.Timer(0.5, gate.set).start() + with host._ready: + assert host._ready.wait_for(lambda: len(host._pending) == 4, timeout=2) + gate.set() for thread in (warm, *threads): thread.join(10) # Both clients' clips were encoded together, not one request at a time. assert sum(recorder.batches) == 5 and max(recorder.batches) >= 4 assert all(int(c[0, 0]) == 2 for c in results["a"]) and all(int(c[0, 0]) == 4 for c in results["b"]) finally: + gate.set() host.close() diff --git a/tests/model_executor/models/qwen3_omni/test_qwen3_omni_deepstack_device.py b/tests/model_executor/models/qwen3_omni/test_qwen3_omni_deepstack_device.py index 3ce8f447962..e3a7646ad2a 100644 --- a/tests/model_executor/models/qwen3_omni/test_qwen3_omni_deepstack_device.py +++ b/tests/model_executor/models/qwen3_omni/test_qwen3_omni_deepstack_device.py @@ -46,6 +46,9 @@ def test_deepstack_buffers_use_execution_device(mocker, vision_limit, num_tokens device_config=DeviceConfig(device="cpu"), scheduler_config=mocker.Mock(spec=SchedulerConfig, max_num_batched_tokens=4), quant_config=None, + # Disabled-feature default. A spec Mock otherwise returns a truthy + # child mock, and _mark_tower_model then requires a PP group. + use_v2_model_runner=False, ) vllm_config.with_hf_config.return_value = vllm_config diff --git a/tests/model_executor/models/qwen3_omni/test_talker_sampler.py b/tests/model_executor/models/qwen3_omni/test_talker_sampler.py index f5ce2111bea..dd3bf601191 100644 --- a/tests/model_executor/models/qwen3_omni/test_talker_sampler.py +++ b/tests/model_executor/models/qwen3_omni/test_talker_sampler.py @@ -2,11 +2,13 @@ # SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project from types import SimpleNamespace +from unittest.mock import Mock, create_autospec, patch import numpy as np import pytest import torch from vllm.sampling_params import SamplingParams +from vllm.v1.worker.gpu.sample.sampler import Sampler from vllm_omni.model_executor.models.qwen3_omni.talker_sampler import Qwen3OmniTalkerSampler @@ -23,7 +25,9 @@ def __init__(self) -> None: def add_request(self, req_idx: int, sampling_params: SamplingParams) -> None: pass - def _sample_random(self, logits, expanded_idx_mapping, idx_mapping_np, pos, top_k, top_p, use_fused_sampler): + def _sample_random( + self, logits, expanded_idx_mapping, idx_mapping_np, pos, top_k, top_p, has_topp_only_rows, use_fused_sampler + ): self.upstream_rows.append(int(logits.shape[0])) return logits.argmax(dim=-1), logits @@ -38,10 +42,34 @@ def __call__(self, logits, input_batch): None, None, False, + False, ) return sampled +def test_custom_sampler_registers_native_masking_kernels(monkeypatch): + """Seeded Talker sampling uses native masking even with FlashInfer available.""" + from vllm.model_executor.warmup.jit_warmup import JitWarmupRegistry + from vllm.v1.sample.ops import topk_topp_sampler + + monkeypatch.setattr( + topk_topp_sampler, + "current_platform", + SimpleNamespace(is_cpu=lambda: False, is_cuda_alike=lambda: True), + ) + registry = JitWarmupRegistry(SimpleNamespace()) + with registry.activate(): + Qwen3OmniTalkerSampler(_BaseSampler()) + + assert set(registry._registrations) == { + topk_topp_sampler._topk_topp, + topk_topp_sampler._topp_split_stats, + topk_topp_sampler._topp_split_step, + topk_topp_sampler._topp_split_mask, + } + assert len(registry) == 4 + + def _sampler(requests: dict[int, SamplingParams]) -> tuple[Qwen3OmniTalkerSampler, _BaseSampler]: base = _BaseSampler() sampler = Qwen3OmniTalkerSampler(base) @@ -53,7 +81,7 @@ def _sampler(requests: dict[int, SamplingParams]) -> tuple[Qwen3OmniTalkerSample def _step(sampler: Qwen3OmniTalkerSampler, slots: list[int], logits: torch.Tensor) -> torch.Tensor: mapping = torch.tensor(slots, dtype=torch.int64) sampled, _ = sampler.base_sampler._sample_random( - logits.clone(), mapping, np.array(slots), torch.zeros(len(slots), dtype=torch.int64), None, None, False + logits.clone(), mapping, np.array(slots), torch.zeros(len(slots), dtype=torch.int64), None, None, False, False ) return sampled @@ -128,3 +156,52 @@ def test_unseeded_step_stays_on_upstream_sampler(): sampled = _step(sampler, [0, 1], logits) assert torch.equal(sampled, logits.argmax(dim=-1)) assert base.upstream_rows == [2] + + +@pytest.mark.parametrize("has_topp_only_rows", [False, True]) +@pytest.mark.parametrize("seeded_rows", [0, 1, 2]) +def test_target_sampler_dispatch_preserves_top_p_only_rows(has_topp_only_rows, seeded_rows): + """Exercise the target runtime dispatcher, rather than an older hook stub.""" + base = Mock(spec=Sampler) + logits = torch.zeros(2, _VOCAB) + mapping = torch.arange(2, dtype=torch.int64) + slots = np.arange(2) + pos = torch.zeros(2, dtype=torch.int64) + base._sample_random = create_autospec( + Sampler._sample_random.__get__(base, Sampler), + side_effect=lambda rows, *args: (rows.argmax(dim=-1), rows), + ) + original = base._sample_random + base.apply_sampling_params.side_effect = lambda logits, *args, **kwargs: logits + base.sampling_states = Mock() + base.sampling_states.get_top_k_top_p.return_value = (None, None, has_topp_only_rows) + base.sampling_states.any_greedy.return_value = False + base.sampling_states.any_explicit_seed.return_value = False + base.use_flashinfer = False + base.use_xpu_sampler = False + base.penalties_state = Mock(device=torch.device("cpu")) + sampler = Qwen3OmniTalkerSampler(base) + for row in range(seeded_rows): + sampler.add_request(row, SamplingParams(temperature=0.9, seed=row)) + with patch( + "vllm_omni.model_executor.models.qwen3_omni.talker_sampler.apply_top_k_top_p", return_value=logits + ) as mask: + sampled, processed = Sampler.sample(base, logits, mapping, mapping, slots, pos, mapping, mapping, slots) + + assert sampled.shape == (2,) + assert processed is logits + if seeded_rows: + mask.assert_called_once_with(logits, None, None, has_topp_only_rows=has_topp_only_rows) + else: + mask.assert_not_called() + if seeded_rows == 0: + original.assert_called_once_with(logits, mapping, slots, pos, None, None, has_topp_only_rows, False) + elif seeded_rows == 1: + original.assert_called_once() + # The other row was already masked, so its upstream call must not mask again. + assert original.call_args.args[-4:] == (None, None, False, False) + assert original.call_args.args[0].shape == (1, _VOCAB) + assert original.call_args.args[1].tolist() == [1] + assert sampled[1] == 0 + else: + original.assert_not_called() diff --git a/tests/model_executor/models/qwen3_tts/test_cuda_graph_decoder.py b/tests/model_executor/models/qwen3_tts/test_cuda_graph_decoder.py index 454cacf3a14..4c093745eda 100644 --- a/tests/model_executor/models/qwen3_tts/test_cuda_graph_decoder.py +++ b/tests/model_executor/models/qwen3_tts/test_cuda_graph_decoder.py @@ -20,6 +20,7 @@ import torch import torch.nn as nn +from tests.helpers.cudnn import ieee_cudnn_fp32 from tests.helpers.mark import hardware_test pytestmark = [pytest.mark.core_model] @@ -51,6 +52,12 @@ CUDAGraphDecoderWrapper = _mod.CUDAGraphDecoderWrapper +@pytest.fixture(scope="module", autouse=True) +def _ieee_cudnn_fp32(): + with ieee_cudnn_fp32(): + yield + + class SyntheticDecoder(nn.Module): """A small decoder mimicking Qwen3TTSTokenizerV2Decoder's interface. @@ -86,14 +93,14 @@ def forward(self, codes): @pytest.fixture(scope="module") -def decoder(): +def decoder(_ieee_cudnn_fp32): """Create a synthetic decoder on CUDA with fixed weights.""" torch.manual_seed(42) return SyntheticDecoder().to(DEVICE).eval() @pytest.fixture(scope="module") -def wrapper(decoder): +def wrapper(decoder, _ieee_cudnn_fp32): """Create a warmed-up CUDAGraphDecoderWrapper.""" w = CUDAGraphDecoderWrapper( decoder=decoder, diff --git a/tests/platforms/npu/test_aux_output.py b/tests/platforms/npu/test_aux_output.py index 8a4362d6dcf..89a51539b16 100644 --- a/tests/platforms/npu/test_aux_output.py +++ b/tests/platforms/npu/test_aux_output.py @@ -14,9 +14,11 @@ AuxOutputSchedulerConnector, PackedBlockHashes, ) +from vllm.forward_context import ForwardContext from vllm.model_executor.layers.fused_moe import routed_experts_capturer from vllm.v1.kv_cache_interface import KVCacheConfig +from tests.helpers.omni_scheduler import make_kv_cache_manager from vllm_omni.platforms.npu.worker.aux_output import NPUAuxOutputMixin pytestmark = [pytest.mark.core_model, pytest.mark.cpu] @@ -46,7 +48,13 @@ def runner(monkeypatch): # Exercise the real capturer, block store, worker and scheduler connectors # with CPU tensors; importing vllm-ascend or owning NPU hardware is unnecessary. monkeypatch.setattr(routed_experts_capturer, "current_platform", SimpleNamespace(device_type="cpu")) - monkeypatch.setattr(routed_experts_capturer, "get_forward_context", lambda: SimpleNamespace(dp_metadata=None)) + # Upstream RoutedExpertsCapturer reads batch_descriptor when dp_metadata + # is absent. ForwardContext.batch_descriptor defaults to None. + monkeypatch.setattr( + routed_experts_capturer, + "get_forward_context", + lambda: ForwardContext(no_compile_layers={}, attn_metadata={}, slot_mapping={}), + ) monkeypatch.setattr(worker, "get_tp_group", lambda: SimpleNamespace(is_first_rank=True)) runner = _Runner() runner.model = torch.nn.Sequential(_AscendMoE()) @@ -105,7 +113,7 @@ def test_routing_reaches_scheduler_and_owns_legacy_batch_snapshot(runner): _capture(runner, [[7, 8]] * 4) output = runner._finish_omni_aux_output(pending, step, torch.tensor([[5]])) request = SimpleNamespace(request_id="r", num_tokens=5) - rows = AuxOutputSchedulerConnector().take_output(request, output) + rows = AuxOutputSchedulerConnector(make_kv_cache_manager(block_size=4)).take_output(request, output) np.testing.assert_array_equal(rows, expected) assert output["r"].token_start == 0 assert runner.synchronized == 1 diff --git a/tests/platforms/test_rocm_ir_op_priority.py b/tests/platforms/test_rocm_ir_op_priority.py new file mode 100644 index 00000000000..1589303c7cd --- /dev/null +++ b/tests/platforms/test_rocm_ir_op_priority.py @@ -0,0 +1,63 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project + +"""ROCm IR-priority contract for the Omni custom-op AITER gate.""" + +from types import SimpleNamespace +from unittest.mock import patch + +import pytest +from vllm.config import CompilationConfig, CompilationMode, DeviceConfig, VllmConfig +from vllm.config.compilation import CUDAGraphMode + +pytestmark = [pytest.mark.core_model, pytest.mark.cpu] + + +@pytest.mark.parametrize( + "gfx90a,rdna4,aiter,mode,custom_ops,expected", + [ + (True, False, True, CompilationMode.NONE, ["all"], ["vllm_c", "native"]), + (False, True, True, CompilationMode.NONE, ["all"], ["aiter", "vllm_c", "native"]), + (False, False, True, CompilationMode.VLLM_COMPILE, ["all"], ["aiter", "native"]), + (False, False, False, CompilationMode.NONE, ["all"], ["vllm_c", "native"]), + (False, False, True, CompilationMode.NONE, ["none"], ["vllm_c", "native"]), + ], + ids=["gfx90a", "rdna4-no-cudagraph", "inductor", "aiter-disabled", "custom-op-disabled"], +) +def test_rocm_ir_op_priority( + monkeypatch: pytest.MonkeyPatch, + gfx90a: bool, + rdna4: bool, + aiter: bool, + mode: CompilationMode, + custom_ops: list[str], + expected: list[str], +) -> None: + # ROCm import reads the architecture; these policy tests need no device. + with patch("torch.cuda.get_device_properties", return_value=SimpleNamespace(gcnArchName="gfx942")): + import vllm.platforms.rocm as rocm + + from vllm_omni.platforms.rocm import platform as rocm_platform_mod + from vllm_omni.platforms.rocm.platform import RocmOmniPlatform + + monkeypatch.setattr(rocm, "on_gfx90a", lambda: gfx90a) + # A copied upstream RDNA4 gate must not replace Omni's custom-op policy. + monkeypatch.setattr(rocm, "on_rdna4", lambda: rdna4) + monkeypatch.setattr(rocm_platform_mod.envs, "VLLM_ROCM_USE_AITER", aiter) + monkeypatch.setattr(rocm_platform_mod.envs, "VLLM_ROCM_USE_AITER_RMSNORM", aiter) + config = VllmConfig( + device_config=DeviceConfig(device="cpu"), + compilation_config=CompilationConfig( + backend="inductor", + mode=mode, + custom_ops=custom_ops, + # The Omni AITER gate must work without CUDA graph capture. + cudagraph_mode=CUDAGraphMode.NONE, + ), + ) + + priority = RocmOmniPlatform.get_default_ir_op_priority(config) + + assert priority.rms_norm == expected + assert priority.fused_add_rms_norm == expected + assert priority.gelu_and_mul_sparse == ["native"] diff --git a/tests/profile/test_omni_torch_profiler.py b/tests/profile/test_omni_torch_profiler.py index 7bfc8f8ccb7..7b594a4f75e 100644 --- a/tests/profile/test_omni_torch_profiler.py +++ b/tests/profile/test_omni_torch_profiler.py @@ -1,3 +1,6 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project + # tests/test_omni_torch_profiler.py from __future__ import annotations @@ -151,7 +154,7 @@ def fake_config(tmp_path): @pytest.fixture def fake_profiler_factory(monkeypatch): - created = {} + created: dict[str, object] = {} def fake_profile(*args, **kwargs): profiler = FakeTorchProfiler(on_trace_ready=kwargs.get("on_trace_ready")) @@ -174,6 +177,27 @@ def wrapper(fake_config, fake_profiler_factory): ) +def test_session_profiler_kwargs_override_and_reset_defaults(wrapper, fake_profiler_factory): + wrapper._active = False + overrides = {"record_shapes": False, "with_stack": False} + wrapper.set_profiler_kwargs(overrides) + assert fake_profiler_factory["kwargs"]["record_shapes"] is False + assert fake_profiler_factory["kwargs"]["with_stack"] is False + assert fake_profiler_factory["kwargs"]["on_trace_ready"] == wrapper._on_trace_ready + assert overrides == {"record_shapes": False, "with_stack": False} + + wrapper.set_profiler_kwargs(None) + assert fake_profiler_factory["kwargs"]["record_shapes"] is True + assert fake_profiler_factory["kwargs"]["with_stack"] is True + + +def test_session_profiler_kwargs_do_not_replace_an_active_profiler(wrapper): + wrapper._active = True + original = wrapper.profiler + wrapper.set_profiler_kwargs({"record_shapes": False}) + assert wrapper.profiler is original + + def test_set_trace_filename_creates_timestamped_session_dir(wrapper, monkeypatch, tmp_path): class FixedDatetime: @classmethod diff --git a/tests/profile/test_profile_iteration_controls.py b/tests/profile/test_profile_iteration_controls.py index 91aab1d5b8a..7148a41fa33 100644 --- a/tests/profile/test_profile_iteration_controls.py +++ b/tests/profile/test_profile_iteration_controls.py @@ -11,6 +11,7 @@ from types import ModuleType, SimpleNamespace import pytest +from vllm.config import ProfilerConfig from vllm.profiler.wrapper import WorkerProfiler from vllm_omni.diffusion.diffusion_engine import DiffusionEngine @@ -27,7 +28,7 @@ class RecordingProfiler(WorkerProfiler): """Use the real vLLM scheduling state machine without a device profiler.""" def __init__(self): - super().__init__(SimpleNamespace(delay_iterations=0, max_iterations=0)) + super().__init__(ProfilerConfig()) self.events = [] def _start(self): @@ -37,6 +38,7 @@ def _stop(self): self.events.append("stop") +@pytest.mark.parametrize("profiler_overrides", [None, {"record_shapes": False}]) @pytest.mark.parametrize("route", ["inline", "subprocess"]) @pytest.mark.parametrize( "controls", @@ -48,9 +50,10 @@ def _stop(self): {"delay_iterations": 0, "max_iterations": 0}, ], ) -async def test_frontend_profile_controls_reach_diffusion_profiler(route, controls, mocker): +async def test_frontend_profile_controls_reach_diffusion_profiler(route, controls, profiler_overrides, mocker): profiler = RecordingProfiler() profiler.start = mocker.Mock(wraps=profiler.start) + profiler.set_profiler_kwargs = mocker.Mock(wraps=profiler.set_profiler_kwargs) worker = SimpleNamespace(_get_profiler=lambda: profiler) def worker_rpc(*, method, args, kwargs=None): @@ -69,7 +72,10 @@ async def dispatch(*, method, args, kwargs=None, stage_ids=None): return await StageDiffusionProc._handle_collective_rpc(client, method, None, args, kwargs or {}) frontend = SimpleNamespace(collective_rpc=dispatch) - await AsyncOmni.start_profile(frontend, "test-profile", stages=[1], **controls) + await AsyncOmni.start_profile( + frontend, "test-profile", stages=[1], profiler_kwargs=profiler_overrides, **controls + ) + profiler.set_profiler_kwargs.assert_called_once_with(profiler_overrides) profiler.start.assert_called_once_with(**controls) assert profiler._delay_iters == controls.get("delay_iterations", 0) assert profiler._max_iters == controls.get("max_iterations", 0) @@ -99,16 +105,23 @@ def _npu_worker_class(monkeypatch): return module.OmniNPUWorkerBase +@pytest.mark.parametrize("profiler_overrides", [None, {"record_shapes": False}]) @pytest.mark.parametrize("kind", ["ar", "diffusion", "npu"]) @pytest.mark.parametrize("controls", [{}, {"delay_iterations": 2, "max_iterations": 3}]) -def test_worker_profile_forwards_controls_and_preserves_default_call(monkeypatch, mocker, kind, controls): - profiler = SimpleNamespace(start=mocker.Mock(), stop=mocker.Mock()) +def test_worker_profile_forwards_controls_and_preserves_default_call( + monkeypatch, mocker, kind, controls, profiler_overrides +): + profiler = RecordingProfiler() + profiler.start = mocker.Mock(wraps=profiler.start) + profiler.stop = mocker.Mock(wraps=profiler.stop) + profiler.set_profiler_kwargs = mocker.Mock(wraps=profiler.set_profiler_kwargs) worker = SimpleNamespace(profiler=profiler, _get_profiler=lambda: profiler) cls = {"ar": OmniGPUWorkerBase, "diffusion": DiffusionWorker}.get(kind) if cls is None: cls = _npu_worker_class(monkeypatch) - cls.profile(worker, True, "test-profile", **controls) + cls.profile(worker, True, "test-profile", profiler_kwargs=profiler_overrides, **controls) profiler.start.assert_called_once_with(**controls) + profiler.set_profiler_kwargs.assert_called_once_with(profiler_overrides) cls.profile(worker, False) profiler.stop.assert_called_once_with() @@ -147,3 +160,35 @@ def test_diffusion_execution_advances_native_iteration_schedule(execution, mocke assert profiler.is_running worker.profile(False) assert profiler.events == ["start", "stop", "start", "stop"] + + +def test_npu_profiler_constructor_overrides_and_defaults(monkeypatch, mocker, tmp_path): + from vllm_omni.platforms.npu.profiler import NPUTorchProfilerWrapper + + torch_npu = ModuleType("torch_npu") + backend = mocker.Mock() + monkeypatch.setattr(torch_npu, "profiler", backend, raising=False) + monkeypatch.setitem(sys.modules, "torch_npu", torch_npu) + config = ProfilerConfig( + profiler="torch", + torch_profiler_dir=str(tmp_path), + torch_profiler_record_shapes=True, + torch_profiler_with_stack=True, + ) + profiler = NPUTorchProfilerWrapper(config, "stage1_rank0", 0) + defaults = backend.profile.call_args.kwargs + assert defaults["record_shapes"] is True + assert defaults["with_modules"] is True + assert defaults["activities"] == [backend.ProfilerActivity.CPU, backend.ProfilerActivity.NPU] + handler = mocker.Mock() + profiler.set_profiler_kwargs({"record_shapes": False, "with_modules": False, "on_trace_ready": handler}) + overrides = backend.profile.call_args.kwargs + assert overrides["record_shapes"] is False + assert overrides["with_modules"] is False + assert overrides["on_trace_ready"] is handler + assert overrides["experimental_config"] is not None + profiler.set_profiler_kwargs(None) + restored = backend.profile.call_args.kwargs + assert restored["record_shapes"] is True + assert restored["with_modules"] is True + assert restored["on_trace_ready"] is defaults["on_trace_ready"] diff --git a/tests/worker/test_img2img_mm_prefix.py b/tests/worker/test_img2img_mm_prefix.py new file mode 100644 index 00000000000..248e1ca7394 --- /dev/null +++ b/tests/worker/test_img2img_mm_prefix.py @@ -0,0 +1,125 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project +"""Bagel's real prompt mask retains both image-prefix attention spans.""" + +import sys +from dataclasses import replace +from types import ModuleType + +import pytest +from PIL import Image +from tokenizers import Tokenizer +from tokenizers.models import WordLevel +from transformers import PreTrainedTokenizerFast +from vllm.multimodal.inputs import MultiModalFeatureSpec, MultiModalKwargsItem, PlaceholderRange +from vllm.multimodal.parse import MultiModalDataItems +from vllm.transformers_utils.configs.bagel import BagelConfig +from vllm.v1.attention.backends import utils + +from vllm_omni.model_executor.models.bagel.bagel import ( + Img2ImgProcessorItems, + OmniBagelMultiModalProcessor, + OmniBagelProcessingInfo, +) +from vllm_omni.utils.mm_prefix import apply_img2img_mm_prefix_patch + +pytestmark = [pytest.mark.core_model, pytest.mark.cpu] + + +@pytest.fixture +def bagel_feature(mocker): + processor = object.__new__(OmniBagelMultiModalProcessor) + info = object.__new__(OmniBagelProcessingInfo) + config = BagelConfig( + vit_config={"image_size": 980, "patch_size": 14}, + latent_patch_size=2, + vae_config={"downsample": 8}, + max_latent_size=64, + ) + tokenizer = PreTrainedTokenizerFast( + tokenizer_object=Tokenizer(WordLevel({"": 0, "<|fim_middle|>": 1}, unk_token="")), + unk_token="", + ) + mocker.patch.object(info, "get_hf_config", return_value=config) + mocker.patch.object(info, "get_tokenizer", return_value=tokenizer) + processor.info = info + mm_items = MultiModalDataItems({"img2img": Img2ImgProcessorItems([Image.new("RGB", (1024, 672))])}) + update = processor._get_prompt_updates(mm_items, {}, {})[0].resolve(0) + assert update.modality == "img2img" + mask = update.content.is_embed(update.content.full) + assert mask.device.type == "cpu" + position = PlaceholderRange(offset=0, length=len(update.content.full), is_embed=mask) + # Includes the VAE and ViT markers, with one non-embedding separator. + assert position.extract_embeds_range() == [(0, 2689), (2691, 7592)] + return MultiModalFeatureSpec( + data=MultiModalKwargsItem({}), + modality="img2img", + identifier="encoder-id", + mm_position=position, + mm_hash="processor-id", + ) + + +def test_real_bagel_img2img_ranges(bagel_feature): + features = [[bagel_feature]] + assert utils.compute_mm_prefix_ranges(features) == {0: [(0, 2689), (2691, 7592)]} + assert features[0][0] is bagel_feature + assert bagel_feature.modality == "img2img" + + +@pytest.mark.parametrize( + "sliding_window,clamp_sliding_window,span_pad", + [(None, False, 0), (1024, False, 0), (1024, True, 0), (None, False, 16)], +) +def test_delegates_native_range_options(bagel_feature, sliding_window, clamp_sliding_window, span_pad): + original = utils.compute_mm_prefix_ranges._omni_mm_prefix_original + expected = original([[replace(bagel_feature, modality="image")]], sliding_window, clamp_sliding_window, span_pad) + assert utils.compute_mm_prefix_ranges([[bagel_feature]], sliding_window, clamp_sliding_window, span_pad) == expected + + +def test_other_modalities_and_metadata_unchanged(bagel_feature): + original = utils.compute_mm_prefix_ranges._omni_mm_prefix_original + others = [replace(bagel_feature, modality=modality) for modality in ("image", "video", "audio", "prompt_embeds")] + assert utils.compute_mm_prefix_ranges([others, []]) == original([others, []]) + identity = (bagel_feature.data, bagel_feature.mm_position, bagel_feature.identifier, bagel_feature.mm_hash) + assert utils.compute_mm_prefix_ranges([[bagel_feature], others, []]) == { + 0: [(0, 2689), (2691, 7592)], + 1: original([others])[0], + 2: [], + } + assert bagel_feature.data is identity[0] + assert bagel_feature.mm_position is identity[1] + assert bagel_feature.identifier is identity[2] + assert bagel_feature.mm_hash is identity[3] + assert bagel_feature.modality == "img2img" + + +def test_installation_handles_both_import_orders_and_is_idempotent(monkeypatch): + original = utils.compute_mm_prefix_ranges._omni_mm_prefix_original + monkeypatch.setattr(utils, "compute_mm_prefix_ranges", original) + runners = [] + for name in ("vllm.v1.worker.gpu_model_runner", "vllm.v1.worker.gpu.model_states.default"): + module = ModuleType(name) + setattr(module, "compute_mm_prefix_ranges", original) + monkeypatch.setitem(sys.modules, name, module) + runners.append(module) + apply_img2img_mm_prefix_patch() + wrapped = utils.compute_mm_prefix_ranges + assert wrapped is not original + assert all(module.compute_mm_prefix_ranges is wrapped for module in runners) + namespace: dict[str, object] = {} + exec("from vllm.v1.attention.backends.utils import compute_mm_prefix_ranges", namespace) + assert namespace["compute_mm_prefix_ranges"] is wrapped + apply_img2img_mm_prefix_patch() + assert utils.compute_mm_prefix_ranges is wrapped + + +def test_conflicting_runner_alias_fails_before_any_rebinding(monkeypatch): + original = utils.compute_mm_prefix_ranges._omni_mm_prefix_original + monkeypatch.setattr(utils, "compute_mm_prefix_ranges", original) + module = ModuleType("vllm.v1.worker.gpu_model_runner") + setattr(module, "compute_mm_prefix_ranges", lambda: None) + monkeypatch.setitem(sys.modules, module.__name__, module) + with pytest.raises(RuntimeError, match="Conflicting"): + apply_img2img_mm_prefix_patch() + assert utils.compute_mm_prefix_ranges is original diff --git a/tests/worker/test_moe_prefill_warmup.py b/tests/worker/test_moe_prefill_warmup.py index 3c0853eeb89..a79819843b0 100644 --- a/tests/worker/test_moe_prefill_warmup.py +++ b/tests/worker/test_moe_prefill_warmup.py @@ -113,7 +113,7 @@ def test_warmup_runs_on_v2_model_runner(moe_detection): sizes = warm_up_moe_prefill_kernels(worker) assert sizes == moe_prefill_warmup_sizes(8, 32768) - assert [num_tokens for num_tokens, _ in worker.model_runner.calls] == sizes + assert worker.model_runner.calls == [(size, {"skip_eplb": True, "remove_lora": False}) for size in sizes] class _FakeMoERunner: @@ -123,7 +123,7 @@ class _FakeMoERunner: @pytest.fixture def fake_fused_moe_module(monkeypatch): module = ModuleType("vllm.model_executor.layers.fused_moe") - module.MoERunner = _FakeMoERunner # type: ignore[attr-defined] + setattr(module, "MoERunner", _FakeMoERunner) monkeypatch.setitem(sys.modules, "vllm.model_executor.layers.fused_moe", module) diff --git a/tests/worker/test_omni_v2_profile_dispatch.py b/tests/worker/test_omni_v2_profile_dispatch.py index 20a4d88fda8..63109ab173e 100644 --- a/tests/worker/test_omni_v2_profile_dispatch.py +++ b/tests/worker/test_omni_v2_profile_dispatch.py @@ -18,6 +18,8 @@ def test_dispatch_eager_and_manager_paths(monkeypatch, dp_size): # dp_size/dp_rank are read-only upstream properties resolved from # parallel_config, so the stub only needs the config namespace. runner.parallel_config = SimpleNamespace(data_parallel_size=dp_size, data_parallel_rank=1) + # Non-spec default: one newly sampled token and no draft steps. + runner.decode_query_len = 1 import vllm.v1.worker.gpu.dp_utils as dp_utils diff --git a/tests/worker/test_process_gpu_memory.py b/tests/worker/test_process_gpu_memory.py index d5cc2eda83c..6262b984f01 100644 --- a/tests/worker/test_process_gpu_memory.py +++ b/tests/worker/test_process_gpu_memory.py @@ -1,3 +1,6 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project + """Tests for process-scoped GPU memory accounting.""" import os @@ -52,24 +55,6 @@ def test_spaces(self, mocker: MockerFixture): class TestGetProcessGpuMemory: - @pytest.mark.skipif(not os.path.exists("/dev/nvidia0"), reason="No GPU") - def test_returns_memory_for_current_process(self): - import torch - - from vllm_omni.worker.gpu_memory_utils import get_process_gpu_memory - - if not torch.cuda.is_available(): - pytest.skip("CUDA not available") - - device = torch.device("cuda:0") - tensor = torch.zeros(1000, 1000, device=device) - - memory = get_process_gpu_memory(0) - assert memory >= 0 - - del tensor - torch.accelerator.empty_cache() - def test_raises_on_invalid_device(self, mocker: MockerFixture): from vllm_omni.worker.gpu_memory_utils import get_process_gpu_memory diff --git a/tests/worker/test_process_gpu_memory_cuda.py b/tests/worker/test_process_gpu_memory_cuda.py new file mode 100644 index 00000000000..8173395c59b --- /dev/null +++ b/tests/worker/test_process_gpu_memory_cuda.py @@ -0,0 +1,25 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project +"""GPU allocation coverage for process-scoped memory accounting.""" + +import pytest + +from tests.helpers.mark import hardware_test + +pytestmark = pytest.mark.core_model + + +@hardware_test(res={"cuda": "L4"}, num_cards=1) +def test_returns_memory_for_current_process(): + import torch + + from vllm_omni.worker.gpu_memory_utils import get_process_gpu_memory + + if not torch.accelerator.is_available(): + pytest.skip("CUDA not available") + + tensor = torch.zeros(1000, 1000, device="cuda:0") + memory = get_process_gpu_memory(0) + assert memory >= 0 + del tensor + torch.accelerator.empty_cache() diff --git a/tests/worker_v2/test_omni_ar_model_runner.py b/tests/worker_v2/test_omni_ar_model_runner.py index fb7d89284ef..e043f901496 100644 --- a/tests/worker_v2/test_omni_ar_model_runner.py +++ b/tests/worker_v2/test_omni_ar_model_runner.py @@ -156,6 +156,8 @@ def test_voxcpm2_streaming_output_owns_channels_and_finalizes_once(monkeypatch, def test_last_pp_rank_orchestration_and_kv_resolver( monkeypatch, needs_history, include_hidden, engine_output_type ) -> None: + # Native prompt-logprob preparation otherwise pins memory using CUDA. + monkeypatch.setattr("vllm.utils.torch_utils.PIN_MEMORY", False) runner = OmniARModelRunner.__new__(OmniARModelRunner) input_batch = SimpleNamespace(req_ids=["req"], num_reqs=1, seq_lens=torch.tensor([3])) input_batch.idx_mapping, input_batch.query_start_loc = torch.tensor([0]), torch.tensor([0, 1]) diff --git a/tests/worker_v2/test_omni_generation_model_runner.py b/tests/worker_v2/test_omni_generation_model_runner.py index e6334a8911c..d9062be2374 100644 --- a/tests/worker_v2/test_omni_generation_model_runner.py +++ b/tests/worker_v2/test_omni_generation_model_runner.py @@ -436,6 +436,7 @@ def test_prewarm_runs_after_finish_and_only_on_consecutive_idle_steps(cached): from vllm_omni.core.sched.output import OmniRequestPrewarm, OmniSchedulerOutput runner = object.__new__(OmniGenerationModelRunner) + runner.encoder_cache = None runner._prev_step_idle = False runner._failed_optional_model_hooks = set() events = [] diff --git a/tests/worker_v2/test_omni_gpu_model_runner.py b/tests/worker_v2/test_omni_gpu_model_runner.py index 511c281a1ee..f67ec8b1a54 100644 --- a/tests/worker_v2/test_omni_gpu_model_runner.py +++ b/tests/worker_v2/test_omni_gpu_model_runner.py @@ -71,6 +71,55 @@ def _make_runner(): return runner +@pytest.mark.parametrize("hook_fails", [False, True]) +def test_load_model_captures_custom_sampler_warmups_without_leaking_context(monkeypatch, hook_fails): + from vllm.model_executor.warmup.jit_warmup import JitWarmupRegistry + from vllm.v1.worker.gpu import model_runner as upstream_model_runner + + from vllm_omni.worker_v2 import omni_model_runner + + runner = _make_runner() + runner.jit_warmup_registry = JitWarmupRegistry(SimpleNamespace()) + runner.use_aux_hidden_state_outputs = False + runner.model = SimpleNamespace(has_preprocess=False) + state = object.__new__(OmniModelState) + state.model = runner.model + sampler_kernel = object() + outer_kernel = object() + outer_registry = JitWarmupRegistry(SimpleNamespace()) + + def custom_sampler(sampler): + JitWarmupRegistry.register(sampler_kernel) + if hook_fails: + raise RuntimeError("custom sampler initialization failed") + return sampler, None + + runner.model.mrv2_custom_sampler = custom_sampler + + def upstream_load(self): + # Match the target's ordering: model state and its sampler hook are + # initialized after the model-construction registry context ends. + self.model_state = upstream_model_runner.init_model_state(None, self.model, None, None) + self.model_state.custom_sampler(object()) + + monkeypatch.setattr(GPUModelRunner, "load_model", upstream_load) + monkeypatch.setattr(omni_model_runner, "init_omni_model_state", lambda *args: state) + monkeypatch.setattr(runner, "_configure_cudagraph_output_contract", lambda: None) + + original_init_model_state = upstream_model_runner.init_model_state + with outer_registry.activate(): + if hook_fails: + with pytest.raises(RuntimeError, match="custom sampler initialization failed"): + runner.load_model() + else: + runner.load_model() + JitWarmupRegistry.register(outer_kernel) + + assert set(runner.jit_warmup_registry._registrations) == {sampler_kernel} + assert set(outer_registry._registrations) == {outer_kernel} + assert upstream_model_runner.init_model_state is original_init_model_state + + @pytest.mark.parametrize("moss", [False, True]) @pytest.mark.parametrize("sender_fails", [False, True]) def test_shutdown_closes_model_owned_first_audio_and_other_resources(moss, sender_fails): @@ -225,6 +274,8 @@ def test_duplex_kv_reanchor_runs_after_block_table_writes(): apply_duplex_kv_reanchor=lambda r, scheduler_output: order.append(("reanchor", r, scheduler_output)) ) runner.aux_output_connector = None + # Encoder payload updates are disabled in this fixture. + runner.encoder_cache = None output = object() runner.kv_connector = SimpleNamespace(no_forward=lambda _s: output) runner._merge_ec_connector_no_forward = lambda _s, value: value @@ -577,6 +628,7 @@ def test_dummy_forward_uses_upstream_execution_state(runner_kind, randomize_inpu runner.is_encoder_decoder = False runner.eplb = MagicMock() runner.kv_connector = MagicMock() + runner.pcp_manager = None runner.input_buffers = object() runner.kv_cache_config = object() runner.attn_groups = [] @@ -646,6 +698,7 @@ def test_drafter_target_layers_reach_execute_state(cg_mode, monkeypatch): runner.cudagraph_manager = MagicMock() runner.cudagraph_manager.run_fullgraph.return_value = (hidden, aux) runner.use_aux_hidden_state_outputs = True + runner.pcp_manager = None runner.model_config = SimpleNamespace() runner.vllm_config = SimpleNamespace() runner.req_states = SimpleNamespace() diff --git a/tests/worker_v2/test_shared_stage_contracts.py b/tests/worker_v2/test_shared_stage_contracts.py index 682614c0a6c..1b92d746561 100644 --- a/tests/worker_v2/test_shared_stage_contracts.py +++ b/tests/worker_v2/test_shared_stage_contracts.py @@ -63,10 +63,12 @@ def test_model_sampler_hook_returning_none_keeps_runner_samplers(): def test_logits_vocab_is_adopted_before_parent_sampler_initialization(monkeypatch): import vllm.v1.worker.gpu.model_runner as upstream + from vllm.model_executor.warmup.jit_warmup import JitWarmupRegistry from vllm_omni.worker_v2 import omni_model_runner as module runner = object.__new__(module.OmniGPUModelRunner) + runner.jit_warmup_registry = JitWarmupRegistry(None) runner.vocab_size = 100 runner.req_states = SimpleNamespace(vocab_size=100) runner.model_config = SimpleNamespace(model_arch_config=SimpleNamespace(vocab_size=100)) diff --git a/tools/build_gradio_pillow13.py b/tools/build_gradio_pillow13.py new file mode 100644 index 00000000000..b133ecca003 --- /dev/null +++ b/tools/build_gradio_pillow13.py @@ -0,0 +1,96 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project +"""Build the pinned Gradio 6.31 Pillow 13 source backport with Python 3.11-3.13.""" + +import argparse +import hashlib +import json +import os +import shutil +import subprocess +import sys +import tarfile +import tempfile +import urllib.request +from pathlib import Path + +INPUTS = Path(__file__).resolve().parent / "gradio_pillow13" + + +def checked(path: Path, expected: str) -> Path: + with path.open("rb") as stream: + if hashlib.file_digest(stream, "sha256").hexdigest() != expected: + raise RuntimeError(f"SHA256 mismatch: {path.name}") + return path + + +def main() -> None: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--output-dir", type=Path, required=True) + output = parser.parse_args().output_dir.resolve() + if not (3, 11) <= sys.version_info[:2] <= (3, 13): + raise RuntimeError("Python 3.11-3.13 required") + manifest = json.loads((INPUTS / "manifest.json").read_text()) + for name, expected in manifest["local_inputs"].items(): + checked(INPUTS / name, expected) + output.mkdir(parents=True, exist_ok=True) + cache = output / ".gradio-pillow13-inputs" + cache.mkdir(exist_ok=True) + for name, pin in manifest["inputs"].items(): + path = cache / name + if not path.exists(): + with urllib.request.urlopen(pin["url"], timeout=60) as source: + path.write_bytes(source.read()) + checked(path, pin["sha256"]) + with tempfile.TemporaryDirectory(prefix="gradio-pillow13-", dir=output) as scratch_name: + scratch = Path(scratch_name) + with tarfile.open(cache / "gradio-6.31.0.tar.gz") as archive: + # The archive hash is pinned above; use the standard safe filter where available. + if hasattr(tarfile, "data_filter"): + archive.extractall(scratch, filter="data") + else: + archive.extractall(scratch) + source = scratch / "gradio-6.31.0" + subprocess.run( + ["git", "apply", "--unidiff-zero", str(INPUTS / "backport.patch")], + cwd=source, + env=dict(os.environ, GIT_CEILING_DIRECTORIES=str(scratch)), + check=True, + ) + # Regenerate dynamic metadata after the source patch; the official sdist omits MCP requirements. + (source / "PKG-INFO").unlink() + shutil.copyfile(cache / "requirements-mcp.txt", source / "requirements-mcp.txt") + shutil.copyfile(INPUTS / "NOTICE.md", source / "gradio/PILLOW13_BACKPORT_NOTICE.md") + with (output / "gradio-pillow13-build.log").open("w") as log: + subprocess.run( + [ + "uv", + "build", + "--no-config", + "--wheel", + "--python", + sys.executable, + "--default-index", + "https://pypi.org/simple", + "--build-constraints", + str(INPUTS / "build-constraints.txt"), + "--out-dir", + str(scratch / "dist"), + str(source), + ], + check=True, + stdout=log, + stderr=subprocess.STDOUT, + timeout=600, + ) + built = checked(scratch / "dist" / manifest["wheel"], manifest["expected_wheel_sha256"]) + destination = output / manifest["wheel"] + shutil.copyfile(built, destination) + receipt = {"wheel": destination.name, "sha256": manifest["expected_wheel_sha256"], "inputs": manifest["inputs"]} + (output / "gradio-pillow13-provenance.json").write_text(json.dumps(receipt, indent=2) + "\n") + print(destination) + + +if __name__ == "__main__": + main() diff --git a/tools/gradio_pillow13/NOTICE.md b/tools/gradio_pillow13/NOTICE.md new file mode 100644 index 00000000000..1294d6bf07b --- /dev/null +++ b/tools/gradio_pillow13/NOTICE.md @@ -0,0 +1,16 @@ +This is an unpublished experimental Gradio6.31 Pillow13 backport. + +Local distribution version: 6.31.0+omni.pillow13 +Base: official PyPI Gradio6.31.0 source archive, SHA256 +40cf094fe8ea1b8ada0fc4c10b685c4c5f09c4061d34c073be25bab342e14e22 +Base upstream source:07341609c68bb98701e7db9ba9e2ad6c0fb2acb4 + +Pillow compatibility changes are backported from the Gradio maintainers' +Apache-2.0 code merged in PR14005 at d6972ed9d1185634cea083e0263839543d5ceb8d: +https://github.com/gradio-app/gradio/pull/14005 + +Changes: permit Pillow>=8,<14; reject I-mode PNG encoding with the exact +upstream conversion/TIFF guidance; identify this artifact with a local version. +The original Apache-2.0 license and bundled third-party licenses remain intact. +This artifact is not an official Gradio release and has no upstream support +claim. Small CPU probes do not qualify the full Gradio/demos/model feature set. diff --git a/tools/gradio_pillow13/backport.patch b/tools/gradio_pillow13/backport.patch new file mode 100644 index 00000000000..517d966fcd9 --- /dev/null +++ b/tools/gradio_pillow13/backport.patch @@ -0,0 +1,21 @@ +--- official6.31/requirements.txt ++++ 6.31.0+omni.pillow13/requirements.txt +@@ -16 +16 @@ +-pillow>=8.0,<13.0 ++pillow>=8.0,<14.0 +--- official6.31/gradio/processing_utils.py ++++ 6.31.0+omni.pillow13/gradio/processing_utils.py +@@ -134 +134,8 @@ + def encode_pil_to_bytes(pil_image, format="png"): ++ if format.lower() == "png" and pil_image.mode == "I": ++ raise OSError( ++ "Cannot save a 32-bit integer (I-mode) image as PNG without losing data. " ++ "If all pixel values are between 0 and 65535, explicitly convert the " ++ "image with image.convert('I;16') before returning it. Otherwise, use " ++ "a format such as TIFF to preserve the 32-bit integer values." ++ ) +--- official6.31/gradio/package.json ++++ 6.31.0+omni.pillow13/gradio/package.json +@@ -3 +3 @@ +- "version": "6.31.0", ++ "version": "6.31.0+omni.pillow13", diff --git a/tools/gradio_pillow13/build-constraints.txt b/tools/gradio_pillow13/build-constraints.txt new file mode 100644 index 00000000000..eedac8a9caf --- /dev/null +++ b/tools/gradio_pillow13/build-constraints.txt @@ -0,0 +1,23 @@ +# This file was autogenerated by uv via the following command: +# uv pip compile --python /usr/bin/python3.12 build-requirements.in --output-file build-constraints.txt +hatch-fancy-pypi-readme==25.1.0 + # via -r build-requirements.in +hatch-requirements-txt==0.4.1 + # via -r build-requirements.in +hatchling==1.33.0 + # via + # -r build-requirements.in + # hatch-fancy-pypi-readme + # hatch-requirements-txt +packaging==26.3 + # via + # hatch-requirements-txt + # hatchling +pathspec==1.1.1 + # via hatchling +pluggy==1.7.0 + # via hatchling +tomlkit==0.15.2 + # via hatchling +trove-classifiers==2026.9.21.13 + # via hatchling diff --git a/tools/gradio_pillow13/manifest.json b/tools/gradio_pillow13/manifest.json new file mode 100644 index 00000000000..6d47af45b4f --- /dev/null +++ b/tools/gradio_pillow13/manifest.json @@ -0,0 +1,26 @@ +{ + "schema": 1, + "base_version": "6.31.0", + "version": "6.31.0+omni.pillow13", + "wheel": "gradio-6.31.0+omni.pillow13-py3-none-any.whl", + "expected_wheel_sha256": "1c80d33042642f5ad61d4f82b6634893a4830e5c4b047025380b97d61da25778", + "base_commit": "07341609c68bb98701e7db9ba9e2ad6c0fb2acb4", + "upstream_patch_commit": "d6972ed9d1185634cea083e0263839543d5ceb8d", + "upstream_patch_url": "https://github.com/gradio-app/gradio/pull/14005", + "license": "Apache-2.0; preserve bundled third-party notices", + "inputs": { + "gradio-6.31.0.tar.gz": { + "url": "https://files.pythonhosted.org/packages/68/98/a76721a95bebc678b8a89c174d72292b3e99b2e4b387b77965d2797a1b52/gradio-6.31.0.tar.gz", + "sha256": "40cf094fe8ea1b8ada0fc4c10b685c4c5f09c4061d34c073be25bab342e14e22" + }, + "requirements-mcp.txt": { + "url": "https://raw.githubusercontent.com/gradio-app/gradio/07341609c68bb98701e7db9ba9e2ad6c0fb2acb4/requirements-mcp.txt", + "sha256": "21951cbe1c5714b99b47a81a968d84c0ed342cfd0a7528ac368318decab49e88" + } + }, + "local_inputs": { + "backport.patch": "3d6691ad0ddc9b4ba95e9ddf677661be1525baf18f2628255273813eb28fab40", + "build-constraints.txt": "35498847b65225c90c83f4ace3f2676f7155f1d33b5e20652a6a941a1237f77d", + "NOTICE.md": "4c450fa5b01628255717a87f8fdcd362d6a0b56cf8a6766e8fb97ad2469d5f1f" + } +} diff --git a/tools/prepare_ci_telemetry.py b/tools/prepare_ci_telemetry.py new file mode 100644 index 00000000000..6fa26646a88 --- /dev/null +++ b/tools/prepare_ci_telemetry.py @@ -0,0 +1,36 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project + +"""Guard the orphan-exporter repair for the pinned final vLLM CUDA image.""" + +import importlib.metadata as metadata +import sys +from pathlib import Path + +from packaging.requirements import Requirement +from packaging.utils import canonicalize_name + +BASE_IMAGE = "vllm/vllm-openai:v0.33.0@sha256:d825baf9ce4b1b6bfc5ecc7b79880c7e6748c65a72d697ce3d47607aafa84c6e" +COMMON = "opentelemetry-exporter-otlp-common" + + +def main() -> None: + # This immutable image's orphan has no shared or OS-owned payload files. + # Refuse other bootstraps rather than inferring that the same repair is safe. + if sys.argv[1:] != [BASE_IMAGE]: + raise RuntimeError("Telemetry repair requires the pinned vLLM CUDA base image") + pins = Path(__file__).resolve().parents[1] / "requirements/cuda-bootstrap.txt" + expected = dict(line.split("==") for line in pins.read_text().splitlines()) + actual = {name: metadata.version(name) for name in expected} + if actual != expected or metadata.version(COMMON) != "0.66b2": + raise RuntimeError("Telemetry bootstrap package versions differ from the pinned base") + for distribution in metadata.distributions(): + if not distribution.metadata.get("Name"): + raise RuntimeError("Unnamed installed distribution; cannot establish orphan status") + if any(canonicalize_name(Requirement(r).name) == COMMON for r in distribution.requires or ()): + raise RuntimeError(f"The common exporter is required by {distribution.metadata['Name']}") + print("Verified the pinned CUDA bootstrap and its unused common exporter") + + +if __name__ == "__main__": + main() diff --git a/vllm_omni/config/omni_config.py b/vllm_omni/config/omni_config.py index 0a803cf4be7..b1a522670b1 100644 --- a/vllm_omni/config/omni_config.py +++ b/vllm_omni/config/omni_config.py @@ -1430,7 +1430,7 @@ def _validate_stage_engine_override_ownership( def _frontend_cli_fields() -> frozenset[str]: """Return vLLM server fields that are consumed before stage startup.""" - from vllm.entrypoints.launchers.cli_args import FrontendArgs + from vllm.entrypoints.common.cli_args import FrontendArgs return frozenset(config_field.name for config_field in fields(FrontendArgs)) diff --git a/vllm_omni/core/sched/omni_ar_scheduler.py b/vllm_omni/core/sched/omni_ar_scheduler.py index e9973cc213e..e5d78fe4069 100644 --- a/vllm_omni/core/sched/omni_ar_scheduler.py +++ b/vllm_omni/core/sched/omni_ar_scheduler.py @@ -8,7 +8,7 @@ from collections import defaultdict from collections.abc import Iterable, Iterator from dataclasses import replace -from typing import Any +from typing import Any, TypedDict import numpy as np from vllm.compilation.cuda_graph import CUDAGraphStat @@ -36,6 +36,13 @@ logger = init_logger(__name__) +class _ChunkSaveKwargs(TypedDict, total=False): + new_token_ids: list[int] + confirmed_num_computed_tokens: int | None + duplex_perf_chunk: tuple[int, int] + segment_generation: int | None + + def _should_emit_engine_output( model_config: Any, *, @@ -124,6 +131,23 @@ class OmniARScheduler(OmniSchedulerMixin, VLLMScheduler): max_num_active_reqs: int + def _update_request_with_output( + self, request: Request, new_token_ids: list[int], is_stale: bool = False + ) -> tuple[list[int], bool]: + status_before_update = request.status + has_output_placeholders = request.num_output_placeholders > 0 + new_token_ids, stopped = super()._update_request_with_output(request, new_token_ids, is_stale=is_stale) + # vLLM 0.33 moved this publication into update_from_output, which + # Omni overrides. Keep both the plain-decode and general paths on + # the upstream frontier: materialized KV only, never in-flight KV. + if ( + not is_stale + and status_before_update == RequestStatus.RUNNING + and (has_output_placeholders or self.num_spec_tokens > 0) + ): + self.kv_cache_manager.cache_blocks(request, request.num_computed_tokens - request.num_in_flight_tokens) + return new_token_ids, stopped + def reset_prefix_cache(self, reset_running_requests: bool = False, reset_connector: bool = False) -> bool: model_config = self.vllm_config.model_config if ( @@ -848,6 +872,9 @@ def update_from_output( ): self.processed_step_seq += 1 self._drain_deferred_frees() + # HiSparse async loads must wait for the same completed-step fence. + if getattr(self, "defer_async_loads", False): + self._drain_deferred_loads() perf_stats: PerfStats | None = None if self.perf_metrics and self.perf_metrics.is_enabled(): @@ -1394,7 +1421,7 @@ def update_from_output( or finished ) ): - save_kwargs = { + save_kwargs: _ChunkSaveKwargs = { "new_token_ids": new_token_ids, "confirmed_num_computed_tokens": confirmed_num_computed_tokens, } @@ -1689,6 +1716,15 @@ def _update_request_as_session(self, session: Request, update: StreamingUpdate) if self._streaming_update_overflows(session, update): return session._omni_segment_generation = int(getattr(session, "_omni_segment_generation", 0) or 0) + 1 + # The pin truncates at ``num_tokens - 1`` and asserts that + # ``num_computed_tokens`` does not pass that index. Omni still keeps + # every computed token: when the computed span already includes the + # final token, append a sentinel so the truncation drops the copy + # and the ``hash_block_size`` slice stays in range. + if session.num_computed_tokens >= session.num_tokens and session._all_token_ids: + session._all_token_ids.append(session._all_token_ids[-1]) + if session.num_tokens > 0 and session.num_computed_tokens > session.num_tokens - 1: + session.num_computed_tokens = session.num_tokens - 1 super()._update_request_as_session(session, update) if hasattr(update, "model_intermediate_buffer"): session.model_intermediate_buffer = update.model_intermediate_buffer @@ -2008,13 +2044,18 @@ def _handle_stopped_request(self, request: Request) -> bool: if finished: return True overflow = getattr(self, "_streaming_context_overflow", None) - if not overflow or request.request_id not in overflow: - return False - # Whether the overflow was recorded by this call's queued update or - # earlier makes no difference: a session in the overflow map is - # terminal, and upstream has just put it back into admission. - self._finish_overflowed_streaming_session(request) - return True + if overflow and request.request_id in overflow: + # Whether the overflow was recorded by this call's queued update or + # earlier makes no difference: a session in the overflow map is + # terminal. Baseline upstream enqueued before returning; this pin + # does not, but a prior admission must still be taken back out. + self._finish_overflowed_streaming_session(request) + return True + # Baseline ``Scheduler._handle_stopped_request`` enqueued here. The pin + # moved that enqueue into ``update_from_output``, which omni replaces. + # Direct callers still need the request in a waiting queue. + self._enqueue_waiting_request(request) + return False def _finish_overflowed_streaming_session(self, request: Request) -> None: """Take a terminal session back out of admission. diff --git a/vllm_omni/core/sched/omni_generation_scheduler.py b/vllm_omni/core/sched/omni_generation_scheduler.py index 7912eb300f3..569a887977d 100644 --- a/vllm_omni/core/sched/omni_generation_scheduler.py +++ b/vllm_omni/core/sched/omni_generation_scheduler.py @@ -405,7 +405,11 @@ def _handle_stopped_request(self, request: Request) -> bool: request.status = RequestStatus.WAITING self._enqueue_waiting_request(request) return False - return super()._handle_stopped_request(request) + finished = super()._handle_stopped_request(request) + if not finished: + # The pin no longer enqueues inside the base stop handler. + self._enqueue_waiting_request(request) + return finished def _requeue_completed_native_chunks(self) -> None: """Return completed native generation chunks to the admission queue. @@ -877,6 +881,9 @@ def update_from_output( ): self.processed_step_seq += 1 self._drain_deferred_frees() + # HiSparse async loads must wait for the same completed-step fence. + if getattr(self, "defer_async_loads", False): + self._drain_deferred_loads() perf_stats: PerfStats | None = None if self.perf_metrics and self.perf_metrics.is_enabled(): diff --git a/vllm_omni/core/sched/omni_scheduler_mixin.py b/vllm_omni/core/sched/omni_scheduler_mixin.py index bdd1c422572..0ab1170b422 100644 --- a/vllm_omni/core/sched/omni_scheduler_mixin.py +++ b/vllm_omni/core/sched/omni_scheduler_mixin.py @@ -865,6 +865,7 @@ def _make_omni_engine_output( is_segment_finished: bool | None = False, new_prompt_len_snapshot: int | None = None, num_generation_tokens: int | None = None, + error: Any = None, ) -> OmniEngineCoreOutput: """Build the common request-output envelope used by LLM schedulers.""" pooling_output_payload = None @@ -893,6 +894,7 @@ def _make_omni_engine_output( is_segment_finished=is_segment_finished, new_prompt_len_snapshot=new_prompt_len_snapshot, num_generation_tokens=num_generation_tokens, + error=error, ) def _append_request_output( @@ -1011,11 +1013,19 @@ def _reject_invalid_grammar_tokens(self, request: Request, new_token_ids: list[i return True def _finish_error_requests(self, outputs: dict[int, list[EngineCoreOutput]]) -> None: - """Drain grammar-compilation and unavailable-encoder errors into terminal outputs.""" + """Drain grammar and encoder-cache errors into terminal outputs.""" + # Pin types this as dict[str, RequestError] and copies the payload onto + # EngineCoreOutput.error. Snapshot first: replacing the attribute drops + # the dict the caller still needs for the error field. grammar_error_reqs = getattr(self, "grammar_compile_error_reqs", None) - error_req_ids = set(grammar_error_reqs or ()) - if grammar_error_reqs: - grammar_error_reqs.clear() + request_errors: dict[str, Any] = dict(grammar_error_reqs) if isinstance(grammar_error_reqs, dict) else {} + if isinstance(grammar_error_reqs, dict): + self.grammar_compile_error_reqs: dict[str, Any] = {} + error_req_ids = set(request_errors) + encoder_cache_mismatch_reqs = getattr(self, "encoder_cache_mismatch_reqs", None) + if encoder_cache_mismatch_reqs: + error_req_ids.update(encoder_cache_mismatch_reqs) + encoder_cache_mismatch_reqs.clear() ec_connector = getattr(self, "ec_connector", None) if ec_connector is not None: error_req_ids.update(ec_connector.take_unavailable_requests()) @@ -1027,6 +1037,7 @@ def _finish_error_requests(self, outputs: dict[int, list[EngineCoreOutput]]) -> request, new_token_ids=[], finish_reason=request.get_finished_reason(), + error=request_errors.get(request.request_id), ) def _remove_stopped_requests_from_queues( diff --git a/vllm_omni/diffusion/attention/backends/flash_attn_hub.py b/vllm_omni/diffusion/attention/backends/flash_attn_hub.py index b24c9969492..527444d824f 100644 --- a/vllm_omni/diffusion/attention/backends/flash_attn_hub.py +++ b/vllm_omni/diffusion/attention/backends/flash_attn_hub.py @@ -28,7 +28,9 @@ def _load_hub_module(repo_id: str): logger.info("Loading %s kernel from HuggingFace Hub...", repo_id) last_error = None - for version in (1, 2): + # FA2 v3 provides Torch stable-ABI builds for the vLLM 0.33 runtime. + versions = (1, 2, 3) if repo_id == "kernels-community/flash-attn2" else (1, 2) + for version in versions: try: return get_kernel(repo_id, version=version) except Exception as exc: diff --git a/vllm_omni/diffusion/diffusion_engine.py b/vllm_omni/diffusion/diffusion_engine.py index 32e8b129ada..832d5abb6dd 100644 --- a/vllm_omni/diffusion/diffusion_engine.py +++ b/vllm_omni/diffusion/diffusion_engine.py @@ -1005,11 +1005,12 @@ def _run_rpc_task(self, task: _RpcTask) -> None: fut.set_exception(TimeoutError(f"RPC call to {task.method} timed out before execution.")) return + result: object = None try: if task.method == "pause_scheduler": - result = self._run_pause_barrier(remaining) + self._run_pause_barrier(remaining) elif task.method == "resume_scheduler": - result = self._open_scheduling_gate() + self._open_scheduling_gate() else: result = self.executor.collective_rpc( method=task.method, @@ -1668,6 +1669,7 @@ def profile( *, delay_iterations: int | None = None, max_iterations: int | None = None, + profiler_kwargs: dict | None = None, ) -> None: """Start or stop profiling on all diffusion workers. @@ -1676,6 +1678,7 @@ def profile( profile_prefix: Optional prefix for trace filename. delay_iterations: Per-session worker steps before profiling starts. max_iterations: Per-session limit on profiled worker steps. + profiler_kwargs: Per-session profiler constructor overrides. """ if is_start: if profile_prefix is None: @@ -1685,11 +1688,13 @@ def profile( logger.info("Stopping diffusion profiling...") try: - kwargs: dict[str, int] = {} + kwargs: dict[str, Any] = {} if delay_iterations is not None: kwargs["delay_iterations"] = delay_iterations if max_iterations is not None: kwargs["max_iterations"] = max_iterations + if profiler_kwargs is not None: + kwargs["profiler_kwargs"] = profiler_kwargs if kwargs: self.collective_rpc(method="profile", args=(is_start, profile_prefix), kwargs=kwargs) else: diff --git a/vllm_omni/diffusion/diffusion_kv/initialization.py b/vllm_omni/diffusion/diffusion_kv/initialization.py index bac0a4d963d..97b9b42e739 100644 --- a/vllm_omni/diffusion/diffusion_kv/initialization.py +++ b/vllm_omni/diffusion/diffusion_kv/initialization.py @@ -7,9 +7,12 @@ import torch from vllm.config import VllmConfig +from vllm.v1.core.kv_cache_planning import ( + get_global_kv_cache_groups, + get_kv_cache_configs, +) from vllm.v1.core.kv_cache_utils import ( generate_scheduler_kv_cache_config, - get_kv_cache_configs, resolve_kv_cache_block_sizes, ) from vllm.v1.kv_cache_interface import KVCacheConfig, KVCacheSpec @@ -55,7 +58,16 @@ def build_native_kv_cache_configs( required_layout=get_connector_required_kv_cache_layout(vllm_config), ) single_type_kv_cache_manager.register_all_kvcache_specs(vllm_config) - worker_configs = get_kv_cache_configs(vllm_config, worker_specs, available_memory) + # Planning now takes the merged groups, and NIXL DCP interleave is finalized + # from those groups before the configs are built. + global_kv_cache_groups = get_global_kv_cache_groups(vllm_config, worker_specs) + vllm_config.adjust_dcp_kv_cache_interleave_size(global_kv_cache_groups) + worker_configs = get_kv_cache_configs( + vllm_config, + worker_specs, + available_memory, + global_kv_cache_groups, + ) scheduler_config = generate_scheduler_kv_cache_config(worker_configs) # Workers only receive a KVCacheConfig, so the resolved layout travels on it # and is adopted into each worker's own CacheConfig before allocation. diff --git a/vllm_omni/diffusion/distributed/autoencoders/wan_vae_fastpath/forwards.py b/vllm_omni/diffusion/distributed/autoencoders/wan_vae_fastpath/forwards.py index 3fd566c6bda..788ea7f15f4 100644 --- a/vllm_omni/diffusion/distributed/autoencoders/wan_vae_fastpath/forwards.py +++ b/vllm_omni/diffusion/distributed/autoencoders/wan_vae_fastpath/forwards.py @@ -272,6 +272,15 @@ def _bitwise_equal(a: torch.Tensor, b: torch.Tensor) -> bool: _PROBE_SEED = 0 +# One Gaussian draw is not a bitwise proof. On sm_90 with cuDNN 9.24, implicit +# spatial padding and an explicit pad are each deterministic, and fp32 matches, +# but bf16/fp16 still differ by one ulp on a minority of inputs. Seed 0 can +# match every element (fp16 2x2 with 32 channels, bf16 6x10 with a cached +# frame) while the next seed does not; that false accept is then reused for +# every later activation of the shape. Eight fixed seeds caught every such +# shape exercised by the lossless decoder tests. Trials stay sequential so a +# probe still holds one input, the reference copy, and two outputs. +_PROBE_TRIALS = 8 def _random_like(t: torch.Tensor, generator: torch.Generator, std: float = 1.0) -> torch.Tensor: @@ -286,28 +295,33 @@ def _equal_on_random_operands( reference: Callable[[torch.Tensor, torch.Tensor], torch.Tensor], reference_input: Callable[[torch.Tensor], torch.Tensor], ) -> bool: - """Compare two convolution formulations on seeded random stand-ins for ``x`` and ``weight``. + """Compare two convolution formulations on several seeded stand-ins for ``x`` and ``weight``. Agreement on the real operands proves little: all-zero activations or - weights give the same bytes under any kernel. Which kernel cuDNN runs - depends only on the verdict key's metadata, never on values, so a match on - random data is what makes a verdict reusable for every later input and - weight update; it is also the only comparison needed. A private generator - leaves the global RNG streams untouched. + weights give the same bytes under any kernel. cuDNN picks the kernel from + the verdict key's metadata, but two deterministic low-precision kernels + can still disagree by an ulp on some inputs, so every one of + ``_PROBE_TRIALS`` probes has to match before the verdict is reused for a + later input or weight update. A private generator leaves the global RNG + streams untouched. ``reference`` convolves ``reference_input(x)`` (the pre-padded or - channels-first copy). The probe input is drawn twice from the same + channels-first copy). Each probe input is drawn twice from the same generator state, so each formulation's input is freed before the other's - is created; callers probe before running the real convolution. The first - call for a shape then holds at most one probe input, the reference copy - and the two probe outputs, no more than comparing real operands did. + is created; callers probe before running the real convolution. Trials are + sequential, so the first call for a shape holds at most one probe input, + the reference copy and the two probe outputs. """ - generator = torch.Generator(device=x.device).manual_seed(_PROBE_SEED) - probe_weight = _random_like(weight, generator, std=weight[0].numel() ** -0.5) - state = generator.get_state() - reference_out = reference(reference_input(_random_like(x, generator)), probe_weight) - generator.set_state(state) - return _bitwise_equal(candidate(_random_like(x, generator), probe_weight), reference_out) + fan_in_std = weight[0].numel() ** -0.5 + for seed in range(_PROBE_SEED, _PROBE_SEED + _PROBE_TRIALS): + generator = torch.Generator(device=x.device).manual_seed(seed) + probe_weight = _random_like(weight, generator, std=fan_in_std) + state = generator.get_state() + reference_out = reference(reference_input(_random_like(x, generator)), probe_weight) + generator.set_state(state) + if not _bitwise_equal(candidate(_random_like(x, generator), probe_weight), reference_out): + return False + return True def _conv_with_spatial_padding( @@ -321,10 +335,10 @@ def _conv_with_spatial_padding( Upstream pads the input tensor explicitly and convolves with ``padding=0``. Handing the spatial padding to cuDNN avoids materializing the padded copy, - but is only bit-identical if cuDNN selects the same kernel for both - formulations. Exact paths compare both on seeded random operands for each - (conv, shape) before running the selected formulation; later calls reuse - the verdict until it is evicted (see :func:`_equal_on_random_operands`). + but is only bit-identical when the two formulations agree on every probe in + :func:`_equal_on_random_operands`. The first call for each (conv, shape) + compares both before running the real convolution once with the selected + formulation; later calls reuse the verdict until it is evicted. Internal channels-last decoder convolutions skip this bitwise check under their tolerance-based contract. """ diff --git a/vllm_omni/diffusion/layers/adalayernorm.py b/vllm_omni/diffusion/layers/adalayernorm.py index a250382721f..7dc2a54eb52 100644 --- a/vllm_omni/diffusion/layers/adalayernorm.py +++ b/vllm_omni/diffusion/layers/adalayernorm.py @@ -279,7 +279,12 @@ def _adaln_fused_forward_impl(module: "AdaLayerNorm", x: torch.Tensor, scale: to if runtime_key in _FAILED_ADALN_KEYS: return None try: - _adaln_scale_shift_layernorm_kernel[(rows,)](*args, num_warps=cfg[1]) + # Keep the half-precision rounding between modulation multiply + # and add; contraction otherwise produces a single BF16/FP16 FMA. + if is_half: + _adaln_scale_shift_layernorm_kernel[(rows,)](*args, num_warps=cfg[1], enable_fp_fusion=False) + else: + _adaln_scale_shift_layernorm_kernel[(rows,)](*args, num_warps=cfg[1]) except Exception as exc: # noqa: BLE001 - fail closed after optimized-path failure _FAILED_ADALN_KEYS.add(runtime_key) logger.warning( diff --git a/vllm_omni/diffusion/models/magi2/modeling_magi2.py b/vllm_omni/diffusion/models/magi2/modeling_magi2.py index 0cabe3409d7..e612211b0b6 100644 --- a/vllm_omni/diffusion/models/magi2/modeling_magi2.py +++ b/vllm_omni/diffusion/models/magi2/modeling_magi2.py @@ -50,6 +50,18 @@ class Modality(IntEnum): TIME = 3 +def _add_stored_gemm_outputs(left: torch.Tensor, right: torch.Tensor) -> torch.Tensor: + """Add GEMM outputs after each result has been stored. + + Inductor rewrites ``gemm + gemm`` as ``addmm`` and adds in the accumulator, + skipping the bf16 store eager inserts between the GEMM and the add. + Promoting the compiled add restores that store-then-add rounding. + """ + if torch.compiler.is_compiling(): + return (left.float() + right.float()).to(dtype=left.dtype) + return left + right + + class Magi2Attention(nn.Module): def __init__( self, @@ -315,7 +327,10 @@ def _shared_experts( # separately keeps both GEMM outputs and the down-projection inputs dense. shared = swiglu7(self.shared_expert_fc1(normalized)) modality = swiglu7(self.modality_specific_shared_expert_fc1(normalized, dispatcher)) - return self.shared_expert_fc2(shared) + self.modality_specific_shared_expert_fc2(modality, dispatcher) + return _add_stored_gemm_outputs( + self.shared_expert_fc2(shared), + self.modality_specific_shared_expert_fc2(modality, dispatcher), + ) def route_input( self, @@ -331,7 +346,10 @@ def combine( routed: torch.Tensor, dispatcher: ModalityDispatcher, ) -> torch.Tensor: - return self.merge_linear(routed) + self._shared_experts(normalized, dispatcher) + return _add_stored_gemm_outputs( + self.merge_linear(routed), + self._shared_experts(normalized, dispatcher), + ) class Magi2PreAdapter(nn.Module): diff --git a/vllm_omni/diffusion/models/minimax_h3/ops/vae/qk_norm_rope.py b/vllm_omni/diffusion/models/minimax_h3/ops/vae/qk_norm_rope.py index f964c15772b..b63dddec995 100644 --- a/vllm_omni/diffusion/models/minimax_h3/ops/vae/qk_norm_rope.py +++ b/vllm_omni/diffusion/models/minimax_h3/ops/vae/qk_norm_rope.py @@ -9,6 +9,61 @@ if HAS_TRITON: + @triton.jit + def _f32_add(left, right): + return tl.inline_asm_elementwise( + "add.rn.f32 $0, $1, $2;", + "=f,f,f", + [left, right], + dtype=tl.float32, + is_pure=True, + pack=1, + ) + + @triton.jit + def _f32_fma(left, right, acc): + return tl.inline_asm_elementwise( + "fma.rn.f32 $0, $1, $2, $3;", + "=f,f,f,f", + [left, right, acc], + dtype=tl.float32, + is_pure=True, + pack=1, + ) + + @triton.jit + def _fold_high_plus_low(values, half: tl.constexpr, heads_per_program: tl.constexpr): + # Shuffle-down by half the active width adds the high lane into the low lane. + view = tl.reshape(values, [2, half, heads_per_program]) + lane = tl.arange(0, 2)[:, None, None] + high = tl.sum(tl.where(lane == 1, view, 0.0), axis=0) + low = tl.sum(tl.where(lane == 0, view, 0.0), axis=0) + return _f32_add(high, low) + + @triton.jit + def _rms_sum_squares(values, heads_per_program: tl.constexpr): + """Sum squares in PyTorch's N=64 vectorized RMSNorm order. + + Each group of four lanes is an ``fma.rn`` chain, matching one float4 + thread. Groups then combine like a warp shuffle-down. ``tl.sum`` does + not keep that parenthesization. + """ + + dims = tl.arange(0, 64)[None, :] + chunk_ids = tl.arange(0, 16)[:, None] + chunks = tl.zeros([16, heads_per_program], dtype=tl.float32) + for group in tl.static_range(16): + acc = tl.zeros([heads_per_program], dtype=tl.float32) + for offset in tl.static_range(4): + val = tl.sum(tl.where(dims == group * 4 + offset, values, 0.0), axis=1) + acc = _f32_fma(val, val, acc) + chunks = tl.where(chunk_ids == group, acc[None, :], chunks) + stage = _fold_high_plus_low(chunks, 8, heads_per_program) + stage = _fold_high_plus_low(stage, 4, heads_per_program) + stage = _fold_high_plus_low(stage, 2, heads_per_program) + stage = _fold_high_plus_low(stage, 1, heads_per_program) + return tl.reshape(stage, [heads_per_program]) + @triton.jit def _qk_rms_norm_rope_exact_kernel( q_ptr, @@ -48,7 +103,7 @@ def _qk_rms_norm_rope_exact_kernel( k_offsets = token * k_stride_token + heads[:, None] * k_stride_head + dims[None, :] * k_stride_dim x_ptrs = tl.where(qk_domain == 0, q_ptr + q_offsets, k_ptr + k_offsets) x = tl.load(x_ptrs, mask=mask, other=0.0).to(tl.float32) - inv_rms = tl.rsqrt(tl.sum(x * x, axis=1) / head_dim + eps) + inv_rms = tl.rsqrt(_rms_sum_squares(x, heads_per_program) / head_dim + eps) normalized = (x * inv_rms[:, None]).to(tl.float16) rotary_half = rotary_dim // 2 @@ -164,9 +219,9 @@ def try_qk_norm_rope_exact( sin = sin.reshape(tokens, sin.shape[-1]) q_output = torch.empty_like(q) k_output = torch.empty_like(k) - # The first two launch dimensions are part of the exactness contract on - # the validated SM90 and SM103 targets: changing the warp-to-row mapping - # changes the FP32 RMS reduction order. + # The RMS sum is an explicit PyTorch vectorized-RMS parenthesization, so + # these launch dimensions no longer change its rounding. They stay fixed + # at the shape validated on SM90 and SM103. heads_per_program = 8 grid = (tokens, triton.cdiv(heads, heads_per_program), 2) launch_args = { diff --git a/vllm_omni/diffusion/models/qwen_image/pipeline_qwen_image.py b/vllm_omni/diffusion/models/qwen_image/pipeline_qwen_image.py index 2511797aee2..208402b0fff 100644 --- a/vllm_omni/diffusion/models/qwen_image/pipeline_qwen_image.py +++ b/vllm_omni/diffusion/models/qwen_image/pipeline_qwen_image.py @@ -622,6 +622,7 @@ def prepare_timesteps(self, num_inference_steps, sigmas, image_seq_len): timesteps, num_inference_steps = retrieve_timesteps( self.scheduler, num_inference_steps, + device=self.device, sigmas=sigmas, mu=mu, ) diff --git a/vllm_omni/diffusion/models/qwen_image/pipeline_qwen_image_edit.py b/vllm_omni/diffusion/models/qwen_image/pipeline_qwen_image_edit.py index a9b7ef95a48..39993082dab 100644 --- a/vllm_omni/diffusion/models/qwen_image/pipeline_qwen_image_edit.py +++ b/vllm_omni/diffusion/models/qwen_image/pipeline_qwen_image_edit.py @@ -1,5 +1,5 @@ # SPDX-License-Identifier: Apache-2.0 -# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project import inspect import json @@ -607,6 +607,7 @@ def prepare_timesteps(self, num_inference_steps, sigmas, image_seq_len): timesteps, num_inference_steps = retrieve_timesteps( self.scheduler, num_inference_steps, + device=self.device, sigmas=sigmas, mu=mu, ) diff --git a/vllm_omni/diffusion/models/qwen_image/pipeline_qwen_image_edit_plus.py b/vllm_omni/diffusion/models/qwen_image/pipeline_qwen_image_edit_plus.py index 8b5c4e3b141..af9f919e35a 100644 --- a/vllm_omni/diffusion/models/qwen_image/pipeline_qwen_image_edit_plus.py +++ b/vllm_omni/diffusion/models/qwen_image/pipeline_qwen_image_edit_plus.py @@ -1,5 +1,5 @@ # SPDX-License-Identifier: Apache-2.0 -# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project import json import logging @@ -597,6 +597,7 @@ def prepare_timesteps(self, num_inference_steps, sigmas, image_seq_len): timesteps, num_inference_steps = retrieve_timesteps( self.scheduler, num_inference_steps, + device=self.device, sigmas=sigmas, mu=mu, ) diff --git a/vllm_omni/diffusion/models/qwen_image/qwen_image_transformer.py b/vllm_omni/diffusion/models/qwen_image/qwen_image_transformer.py index 131fdc71ed1..5d5953a040f 100644 --- a/vllm_omni/diffusion/models/qwen_image/qwen_image_transformer.py +++ b/vllm_omni/diffusion/models/qwen_image/qwen_image_transformer.py @@ -5,6 +5,7 @@ from collections import OrderedDict from collections.abc import Callable, Iterable +from contextlib import nullcontext from math import prod from typing import TYPE_CHECKING, Any @@ -16,12 +17,15 @@ from diffusers.models.embeddings import TimestepEmbedding, Timesteps from diffusers.models.modeling_outputs import Transformer2DModelOutput from diffusers.models.normalization import AdaLayerNormContinuous +from diffusers.models.normalization import RMSNorm as DiffusersRMSNorm +from vllm import envs from vllm.logger import init_logger from vllm.model_executor.layers.linear import ( ColumnParallelLinear, QKVParallelLinear, ReplicatedLinear, RowParallelLinear, + UnquantizedLinearMethod, ) from vllm.model_executor.model_loader.weight_utils import default_weight_loader @@ -60,6 +64,7 @@ select01_modulation_native, ) from vllm_omni.diffusion.layers.rope import RotaryEmbedding +from vllm_omni.platforms import current_omni_platform logger = init_logger(__name__) @@ -655,6 +660,45 @@ def forward(self, hidden_states: torch.Tensor) -> torch.Tensor: return hidden_states +class _ReferenceAdaLayerNorm(AdaLayerNorm): + """Keep PyTorch reduction and activation rounding in the eager reference.""" + + def forward_cuda(self, x: torch.Tensor, scale: torch.Tensor, shift: torch.Tensor) -> torch.Tensor: + return self.forward_native(x, scale, shift) + + +class _ReferenceTextQKVMethod(UnquantizedLinearMethod): + """Separate text GEMMs while retaining packed weights and LoRA dispatch.""" + + def __init__(self, original: UnquantizedLinearMethod): + # Preserve the configured GEMM provider instead of selecting a new one. + self._gemm_impl = original._gemm_impl + + def apply(self, layer: nn.Module, x: torch.Tensor, bias: torch.Tensor | None = None) -> torch.Tensor: + if envs.VLLM_BATCH_INVARIANT or x.device.type != "cuda" or x.dtype != torch.bfloat16: + return super().apply(layer, x, bias) + sizes = layer.output_partition_sizes + weights = layer.weight.split(sizes, dim=0) + biases = bias.split(sizes, dim=0) if bias is not None else [None] * len(sizes) + return torch.cat([self._gemm_impl(layer, x, w, b) for w, b in zip(weights, biases)], dim=-1) + + +def _reference_qk_norm_rope( + q: torch.Tensor, + k: torch.Tensor, + norm_q: nn.Module, + norm_k: nn.Module, + freqs: torch.Tensor, +) -> tuple[torch.Tensor, torch.Tensor]: + from diffusers.models.transformers.transformer_qwenimage import apply_rotary_emb_qwen + + # FP32 complex rotation belongs outside compiled/FA paths (#7494). + return ( + apply_rotary_emb_qwen(norm_q(q), freqs, use_real=False), + apply_rotary_emb_qwen(norm_k(k), freqs, use_real=False), + ) + + class QwenImageCrossAttention(nn.Module): def __init__( self, @@ -671,6 +715,7 @@ def __init__( out_dim: int | None = None, quant_config: QuantizationConfig | None = None, prefix: str = "", + reference_eager: bool = False, ) -> None: super().__init__() assert dim % num_heads == 0 @@ -692,9 +737,6 @@ def __init__( self.query_num_heads = self.to_qkv.num_heads self.kv_num_heads = self.to_qkv.num_kv_heads - self.norm_q = nn.RMSNorm(head_dim, eps=eps) if qk_norm else nn.Identity() - self.norm_k = nn.RMSNorm(head_dim, eps=eps) if qk_norm else nn.Identity() - self.inner_dim = out_dim if out_dim is not None else head_dim * self.total_num_heads assert context_pre_only is not None @@ -730,9 +772,6 @@ def __init__( prefix=f"{prefix}.to_out.0", ) - self.norm_added_q = nn.RMSNorm(head_dim, eps=eps) - self.norm_added_k = nn.RMSNorm(head_dim, eps=eps) - self.attn = Attention( num_heads=self.query_num_heads, head_size=self.head_dim, @@ -741,6 +780,24 @@ def __init__( num_kv_heads=self.kv_num_heads, ) self.rope = RotaryEmbedding(is_neox_style=False) + # TORCH_SDPA is the existing conservative correctness reference. Do not + # change the selected attention backend or compiled/FA arithmetic. + self.reference_precision = bool( + reference_eager + and current_omni_platform.is_cuda() + and self.attn.attn_backend is not None + and self.attn.attn_backend.get_name() == "SDPA" + ) + norm_cls = DiffusersRMSNorm if self.reference_precision else nn.RMSNorm + self.norm_q = norm_cls(head_dim, eps=eps) if qk_norm else nn.Identity() + self.norm_k = norm_cls(head_dim, eps=eps) if qk_norm else nn.Identity() + self.norm_added_q = norm_cls(head_dim, eps=eps) + self.norm_added_k = norm_cls(head_dim, eps=eps) + # Keep QKVParallelLinear and its loaders, TP shards, quantization and + # outer LoRA wrappers. Only the concrete dense method changes its GEMM + # partitioning; explicit batch-invariant execution still delegates. + if self.reference_precision and type(self.add_kv_proj.quant_method) is UnquantizedLinearMethod: + self.add_kv_proj.quant_method = _ReferenceTextQKVMethod(self.add_kv_proj.quant_method) try: config = get_forward_context().omni_diffusion_config @@ -785,7 +842,8 @@ def forward( # sharded by `_sp_plan` while `txt_freqs` stays replicated, so a joint # table built from the two would not line up with the local tokens. use_fused_joint = ( - self.qk_norm + not self.reference_precision + and self.qk_norm and qk_norm_rope_table is not None and not sp_active and _fused_cuda_supported(img_query, img_key, self.head_dim, qk_norm_rope_table.shape[-1], interleaved=True) @@ -805,6 +863,14 @@ def forward( qk_norm_rope_table, self.eps, ) + elif self.reference_precision: + img_query, img_key = _reference_qk_norm_rope(img_query, img_key, self.norm_q, self.norm_k, vid_freqs) + txt_query, txt_key = _reference_qk_norm_rope( + txt_query, txt_key, self.norm_added_q, self.norm_added_k, txt_freqs + ) + joint_query = torch.cat([txt_query, img_query], dim=1) + joint_key = torch.cat([txt_key, img_key], dim=1) + joint_value = torch.cat([txt_value, img_value], dim=1) else: img_query, img_key = _qwen_image_qk_norm_rope( img_query, @@ -882,6 +948,8 @@ def forward( class QwenImageTransformerBlock(nn.Module): + reference_precision = False + def __init__( self, dim: int, @@ -892,6 +960,7 @@ def __init__( zero_cond_t: bool = False, quant_config: QuantizationConfig | None = None, prefix: str = "", + reference_eager: bool = False, ): super().__init__() @@ -913,7 +982,6 @@ def __init__( prefix=f"{prefix}.img_mod.1", ), ) - self.img_norm1 = AdaLayerNorm(dim, elementwise_affine=False, eps=eps) self.attn = QwenImageCrossAttention( dim=dim, num_heads=num_attention_heads, @@ -922,8 +990,12 @@ def __init__( head_dim=attention_head_dim, quant_config=quant_config, prefix=f"{prefix}.attn", + reference_eager=reference_eager, ) - self.img_norm2 = AdaLayerNorm(dim, elementwise_affine=False, eps=eps) + self.reference_precision = self.attn.reference_precision + norm_cls = _ReferenceAdaLayerNorm if self.reference_precision else AdaLayerNorm + self.img_norm1 = norm_cls(dim, elementwise_affine=False, eps=eps) + self.img_norm2 = norm_cls(dim, elementwise_affine=False, eps=eps) self.img_mlp = FeedForward( dim=dim, dim_out=dim, @@ -944,9 +1016,9 @@ def __init__( prefix=f"{prefix}.txt_mod.1", ), ) - self.txt_norm1 = AdaLayerNorm(dim, elementwise_affine=False, eps=eps) + self.txt_norm1 = norm_cls(dim, elementwise_affine=False, eps=eps) # Text doesn't need separate attention - it's handled by img_attn joint computation - self.txt_norm2 = AdaLayerNorm(dim, elementwise_affine=False, eps=eps) + self.txt_norm2 = norm_cls(dim, elementwise_affine=False, eps=eps) self.txt_mlp = FeedForward( dim=dim, dim_out=dim, @@ -986,7 +1058,9 @@ def forward( txt_mod1, txt_mod2 = txt_mod_params.chunk(2, dim=-1) # Each [B, 3*dim] # Process image stream - norm1 + modulation - use_fused_select01 = modulate_index is not None and can_use_qwen_select01_triton(hidden_states) + use_fused_select01 = ( + not self.reference_precision and modulate_index is not None and can_use_qwen_select01_triton(hidden_states) + ) if use_fused_select01: img_modulated, img_gate1 = fused_layernorm_select01( hidden_states, @@ -1165,19 +1239,12 @@ def __init__( self.guidance_embeds = guidance_embeds self.quant_config = quant_config - if not use_layer3d_rope: - self.pos_embed = QwenEmbedRope(theta=10000, axes_dim=list(axes_dims_rope), scale_rope=True) - else: - self.pos_embed = QwenEmbedLayer3DRope(theta=10000, axes_dim=list(axes_dims_rope), scale_rope=True) - self.time_text_embed = QwenTimestepProjEmbeddings( embedding_dim=self.inner_dim, use_additional_t_cond=use_additional_t_cond, quant_config=quant_config, ) - self.txt_norm = nn.RMSNorm(joint_attention_dim, eps=1e-6) - # Entry projections (image/text) are kept full precision — # small sensitive layers at the network boundary (see #2728). self.img_in = ReplicatedLinear( @@ -1206,11 +1273,23 @@ def __init__( zero_cond_t=zero_cond_t, quant_config=quant_config, prefix=f"transformer_blocks.{i}", + reference_eager=od_config.enforce_eager and od_config.dtype == torch.bfloat16, ) for i in range(num_layers) ] ) + self.reference_precision = any(block.reference_precision for block in self.transformer_blocks) + norm_cls = DiffusersRMSNorm if self.reference_precision else nn.RMSNorm + self.txt_norm = norm_cls(joint_attention_dim, eps=1e-6) + # Diffusers constructs these complex tables on CPU before moving them; + # do the same computation rather than copying any reference operand. + with torch.device("cpu") if self.reference_precision else nullcontext(): + if not use_layer3d_rope: + self.pos_embed = QwenEmbedRope(theta=10000, axes_dim=list(axes_dims_rope), scale_rope=True) + else: + self.pos_embed = QwenEmbedLayer3DRope(theta=10000, axes_dim=list(axes_dims_rope), scale_rope=True) + # Final modulation and output projection are kept full precision — # they produce the output latent and are precision-sensitive # (see #2728). @@ -1365,9 +1444,11 @@ def forward( # One packed table per forward for the joint QK RMSNorm + RoPE + cat of # every block. Never mutate the caller-owned dict: rebind to a copy. - qk_norm_rope_table = _packed_qk_norm_rope_table( - image_rotary_emb[0], image_rotary_emb[1], hidden_states.shape[0], hidden_states.dtype - ) + qk_norm_rope_table = None + if not self.reference_precision: + qk_norm_rope_table = _packed_qk_norm_rope_table( + image_rotary_emb[0], image_rotary_emb[1], hidden_states.shape[0], hidden_states.dtype + ) if qk_norm_rope_table is not None: attention_kwargs = {**(attention_kwargs or {}), QK_NORM_ROPE_TABLE_KEY: qk_norm_rope_table} elif attention_kwargs is not None: diff --git a/vllm_omni/diffusion/models/z_image/pipeline_z_image.py b/vllm_omni/diffusion/models/z_image/pipeline_z_image.py index b40fd57bfa8..f123b844b8f 100644 --- a/vllm_omni/diffusion/models/z_image/pipeline_z_image.py +++ b/vllm_omni/diffusion/models/z_image/pipeline_z_image.py @@ -220,7 +220,11 @@ def __init__( od_config, hf_config=text_encoder_config, ).to(self._execution_device) - if text_encoder_config.tie_word_embeddings: + # Online FP8 rewrites lm_head.weight in place (transpose + fp8). Sharing + # that storage with the token embedding turns vocab lookup into an OOB + # gather. Leave the tied head detached whenever a quant config is set; + # prompt encoding reads hidden states, not logits. + if text_encoder_config.tie_word_embeddings and not od_config.quantization_config: self.text_encoder.lm_head.weight = self.text_encoder.get_input_embeddings().weight vae_config = DistributedAutoencoderKL.load_config(model, subfolder="vae", local_files_only=local_files_only) diff --git a/vllm_omni/diffusion/vllm_config.py b/vllm_omni/diffusion/vllm_config.py index 514ba671484..c5a9c495eae 100644 --- a/vllm_omni/diffusion/vllm_config.py +++ b/vllm_omni/diffusion/vllm_config.py @@ -93,6 +93,9 @@ class _DiffusionVllmModelConfig: # narrow compatibility surface here instead of manufacturing a second # model configuration object in the Worker backend. is_mm_prefix_lm: bool = False + # Diffusion supplies attention specs directly and uses native default + # grouping rather than a registered model's SupportsKVCacheGrouping hook. + supports_kv_cache_grouping: bool = False rswa_window: int | None = None _attention_num_heads: int | None = field(default=None, init=False, repr=False) _attention_num_kv_heads: int | None = field(default=None, init=False, repr=False) diff --git a/vllm_omni/diffusion/worker/diffusion_worker.py b/vllm_omni/diffusion/worker/diffusion_worker.py index 9037d4fa14e..760c95c1b6b 100644 --- a/vllm_omni/diffusion/worker/diffusion_worker.py +++ b/vllm_omni/diffusion/worker/diffusion_worker.py @@ -685,6 +685,7 @@ def profile( *, delay_iterations: int | None = None, max_iterations: int | None = None, + profiler_kwargs: dict | None = None, ) -> None: """Start or stop profiling for this GPU worker. @@ -693,6 +694,7 @@ def profile( profile_prefix: Optional prefix for trace filename. delay_iterations: Per-session worker steps before profiling starts. max_iterations: Per-session limit on profiled worker steps. + profiler_kwargs: Per-session profiler constructor overrides. """ profiler = self._get_profiler() if profiler is None: @@ -704,6 +706,7 @@ def profile( filename = profile_prefix or f"diffusion_rank{self.rank}_{int(time.time())}" profiler.set_trace_filename(filename) + profiler.set_profiler_kwargs(profiler_kwargs) kwargs: dict[str, int] = {} if delay_iterations is not None: kwargs["delay_iterations"] = delay_iterations diff --git a/vllm_omni/engine/orchestrator.py b/vllm_omni/engine/orchestrator.py index a7da4df3d47..48b179ca77d 100644 --- a/vllm_omni/engine/orchestrator.py +++ b/vllm_omni/engine/orchestrator.py @@ -125,6 +125,23 @@ def _event_driven_orch_enabled() -> bool: return value.strip().lower() in ("1", "true", "yes", "on") +def _stage_request_error_text(eco: Any, fallback: str) -> tuple[str, str | None]: + """Prefer RequestError.message over stop_reason for a stage failure. + + A paused engine sets ``error`` and leaves ``stop_reason`` unset. Bare + ``FinishReason.ERROR`` outputs keep the historical fallback. + """ + request_error = getattr(eco, "error", None) + message = getattr(request_error, "message", None) + code = getattr(request_error, "code", None) + if isinstance(message, str) and message: + return message, code if isinstance(code, str) else None + stop_reason = getattr(eco, "stop_reason", None) + if isinstance(stop_reason, str) and stop_reason: + return stop_reason, None + return fallback, None + + def _build_terminal_empty_output( request_id: str, *, @@ -976,6 +993,7 @@ async def _process_llm_stage_outputs( pool = self.stage_pools[stage_id] await self._route_upstream_first_audio(stage_id, replica_id, raw_outputs) await self._handle_kv_ready_raw_outputs(stage_id, raw_outputs) + reported_session_errors: set[str] = set() for eco in raw_outputs.outputs: mm_cache_miss_hashes = getattr(eco, "mm_cache_miss_hashes", None) @@ -1014,21 +1032,27 @@ async def _process_llm_stage_outputs( # A non-session ERROR is fatal to the whole request. Handle it # before normal terminal routing so it reaches the client once. + # Cleanup pops the output-processor entry, so process_outputs will + # not see this output. Prefer RequestError.message: a paused engine + # sets error and leaves stop_reason unset. if ( getattr(eco, "finish_reason", None) == FinishReason.ERROR and not segment_finished and not req_state.session_owned ): - reason = ( - "Multi-modal processor cache miss. Please retry the request." - if mm_cache_miss_hashes - else getattr(eco, "stop_reason", None) - ) - await self._handle_stage_error( - stage_id, - eco, - error=reason if isinstance(reason, str) and reason else "Stage request failed", - ) + reason, error_type = _stage_request_error_text(eco, "Stage request failed") + if mm_cache_miss_hashes: + reason = "Multi-modal processor cache miss. Please retry the request." + error_type = None + if error_type is None: + await self._handle_stage_error(stage_id, eco, error=reason) + else: + await self._handle_stage_error( + stage_id, + eco, + error=reason, + error_type=error_type, + ) continue # Only streaming requests maintain the per-stage segment snapshot. @@ -1052,12 +1076,18 @@ async def _process_llm_stage_outputs( ): raw_terminal_request_ids.add(req_state.request_id) - await self._report_duplex_session_request_error( - stage_id, - replica_id, - eco, - req_state, - ) + if ( + await self._report_duplex_session_request_error( + stage_id, + replica_id, + eco, + req_state, + ) + is True + ): + reported_session_errors.add(req_state.request_id) + if request_id is not None: + reported_session_errors.add(request_id) if isinstance(getattr(eco, "stage_work", None), list): self._collect_duplex_stage_work(stage_id, eco, req_state) @@ -1067,6 +1097,12 @@ async def _process_llm_stage_outputs( raw_outputs, iteration_stats=iteration_stats, ) + if reported_session_errors: + # Duplex already enqueued the client error. Drop a processor + # from_error for the same id so it is not delivered again. + processed = [ + output for output in processed if getattr(output, "request_id", None) not in reported_session_errors + ] if self._stat_logger is not None and (raw_outputs.scheduler_stats is not None or iteration_stats is not None): self._stat_logger.record( @@ -1571,12 +1607,26 @@ async def _handle_processed_outputs(self, stage_id: int, replica_id: int, output await self._route_output(stage_id, replica_id, output, req_state, stage_metrics) - async def _handle_stage_error(self, stage_id: int, output: Any, *, error: str | None = None) -> None: + async def _handle_stage_error( + self, + stage_id: int, + output: Any, + *, + error: str | None = None, + error_type: str | None = None, + ) -> None: """Emit a frontend-visible error and clean up request state.""" if self._cfg_tracker.is_companion(output.request_id): parent_id = self._cfg_tracker.get_parent_id(output.request_id) or output.request_id else: parent_id = output.request_id + resolved_error_type: str | None + if getattr(output, "mm_cache_miss_hashes", None): + resolved_error_type = "ServiceUnavailableError" + elif error_type is not None: + resolved_error_type = error_type + else: + resolved_error_type = getattr(output, "error_type", None) await self.output_async_queue.put( ErrorMessage( request_id=parent_id, @@ -1585,9 +1635,7 @@ async def _handle_stage_error(self, stage_id: int, output: Any, *, error: str | status_code=503 if getattr(output, "mm_cache_miss_hashes", None) else getattr(output, "error_status_code", None), - error_type="ServiceUnavailableError" - if getattr(output, "mm_cache_miss_hashes", None) - else getattr(output, "error_type", None), + error_type=resolved_error_type, mm_cache_miss_hashes=getattr(output, "mm_cache_miss_hashes", None), ) ) @@ -1595,7 +1643,7 @@ async def _handle_stage_error(self, stage_id: int, output: Any, *, error: str | [parent_id, *self._cfg_tracker.cleanup_parent(parent_id)], abort=True, release_owners=True, - trace_error=getattr(output, "error_type", None) or "stage_error", + trace_error=resolved_error_type or "stage_error", ) def _update_stage_replica_waiting(self, stage_id: int, replica_id: int, n_waiting: int) -> None: @@ -2025,7 +2073,7 @@ async def _report_duplex_session_request_error( replica_id: int | None, eco: Any, req_state: OrchestratorRequestState, - ) -> None: + ) -> bool: """Turn a scheduler-side error finish of a duplex session request into a request error. A duplex session request is resumable, so its terminal outputs are @@ -2039,13 +2087,12 @@ async def _report_duplex_session_request_error( the serving layer can fail and close the session. """ if not req_state.session_owned: - return + return False if getattr(eco, "finish_reason", None) != FinishReason.ERROR: - return + return False if getattr(eco, "is_segment_finished", False): - return - stop_reason = getattr(eco, "stop_reason", None) - error = stop_reason if isinstance(stop_reason, str) and stop_reason else "duplex session request failed" + return False + error, error_type = _stage_request_error_text(eco, "duplex session request failed") logger.error( "[Orchestrator] Duplex session request %s failed at stage-%s replica-%s: %s", req_state.request_id, @@ -2056,12 +2103,14 @@ async def _report_duplex_session_request_error( await self.output_async_queue.put( ErrorMessage( error=error, + error_type=error_type, fatal=False, request_id=req_state.request_id, stage_id=stage_id, mm_cache_miss_hashes=getattr(eco, "mm_cache_miss_hashes", None), ) ) + return True async def _finish_raw_terminal_requests( self, @@ -3666,13 +3715,13 @@ async def _dispatch_message(self, msg: EngineQueueMessage) -> bool: if isinstance(msg, NextStageInputMessage): await self._handle_next_stage_input(msg) return True - msg_type = msg.type + msg_type = getattr(msg, "type") if msg_type == "add_request": - await self._handle_add_request(msg) + await self._handle_add_request(cast(StageSubmissionMessage, msg)) elif msg_type == "streaming_update": - await self._handle_streaming_update(msg) + await self._handle_streaming_update(cast(StageSubmissionMessage, msg)) elif msg_type == "add_companion_request": - await self._handle_add_companion(msg) + await self._handle_add_companion(cast(AddCompanionRequestMessage, msg)) elif msg_type == "interaction": interaction = cast(InteractionMessage, msg) try: diff --git a/vllm_omni/engine/stage_pool.py b/vllm_omni/engine/stage_pool.py index bcb3f6f2817..687849956a6 100644 --- a/vllm_omni/engine/stage_pool.py +++ b/vllm_omni/engine/stage_pool.py @@ -1362,6 +1362,7 @@ async def process_llm_raw_outputs( raw_outputs.outputs, raw_outputs.timestamp, iteration_stats, + weight_version=raw_outputs.weight_version, ) if self._watermarkers: self._discard_completed_llm_watermark_state(raw_outputs.outputs, processor.request_states) diff --git a/vllm_omni/entrypoints/async_omni.py b/vllm_omni/entrypoints/async_omni.py index 0045cb2b32a..cf932488892 100644 --- a/vllm_omni/entrypoints/async_omni.py +++ b/vllm_omni/entrypoints/async_omni.py @@ -1201,6 +1201,7 @@ async def start_profile( # type: ignore[override] *, delay_iterations: int | None = None, max_iterations: int | None = None, + profiler_kwargs: dict[str, Any] | None = None, ) -> list[Any]: """Start profiling specified stages. @@ -1211,16 +1212,19 @@ async def start_profile( # type: ignore[override] stages: List of stage IDs to profile. If None, profiles all stages. delay_iterations: Optional per-session delayed start (vLLM compat). max_iterations: Optional per-session iteration cap (vLLM compat). + profiler_kwargs: Optional per-session profiler constructor overrides. """ validate_profile_iteration_bounds(delay_iterations, max_iterations) # Mirror upstream ExecutorBase.profile: forward the per-session # controls as kwargs only when set, so workers without them keep # working with plain (is_start, profile_prefix) args. - kwargs: dict[str, int] = {} + kwargs: dict[str, Any] = {} if delay_iterations is not None: kwargs["delay_iterations"] = delay_iterations if max_iterations is not None: kwargs["max_iterations"] = max_iterations + if profiler_kwargs is not None: + kwargs["profiler_kwargs"] = profiler_kwargs return await self.collective_rpc( method="profile", args=(True, profile_prefix), diff --git a/vllm_omni/entrypoints/cli/benchmark/main.py b/vllm_omni/entrypoints/cli/benchmark/main.py index 342b4fb4f37..54f870b16ca 100644 --- a/vllm_omni/entrypoints/cli/benchmark/main.py +++ b/vllm_omni/entrypoints/cli/benchmark/main.py @@ -1,3 +1,6 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project + """Top-level ``vllm bench`` command integration for vLLM-Omni. ``OmniBenchmarkSubcommand`` creates the benchmark command group, discovers @@ -12,7 +15,7 @@ import typing from vllm.entrypoints.cli.types import CLISubcommand -from vllm.entrypoints.serve.utils.api_utils import VLLM_SUBCMD_PARSER_EPILOG +from vllm.entrypoints.common.startup import VLLM_SUBCMD_PARSER_EPILOG from vllm_omni.entrypoints.cli.benchmark.base import OmniBenchmarkSubcommandBase diff --git a/vllm_omni/entrypoints/cli/main.py b/vllm_omni/entrypoints/cli/main.py index 17be480b27a..3c2a694f44f 100644 --- a/vllm_omni/entrypoints/cli/main.py +++ b/vllm_omni/entrypoints/cli/main.py @@ -26,7 +26,7 @@ def main(): if "VLLM_LOGGING_COLOR" not in os.environ: os.environ["VLLM_LOGGING_COLOR"] = "1" - from vllm.entrypoints.serve.utils.api_utils import VLLM_SUBCMD_PARSER_EPILOG, cli_env_setup + from vllm.entrypoints.common.startup import VLLM_SUBCMD_PARSER_EPILOG, cli_env_setup from vllm.logger import configure_logging_from_args import vllm_omni.entrypoints.cli.benchmark.main diff --git a/vllm_omni/entrypoints/cli/serve.py b/vllm_omni/entrypoints/cli/serve.py index f113cd69a6d..a36e8912c4b 100644 --- a/vllm_omni/entrypoints/cli/serve.py +++ b/vllm_omni/entrypoints/cli/serve.py @@ -21,8 +21,8 @@ import uvloop from vllm.entrypoints.cli.types import CLISubcommand -from vllm.entrypoints.launchers.cli_args import make_arg_parser, validate_parsed_serve_args -from vllm.entrypoints.serve.utils.api_utils import VLLM_SUBCMD_PARSER_EPILOG +from vllm.entrypoints.common.cli_args import make_arg_parser, validate_parsed_serve_args +from vllm.entrypoints.common.startup import VLLM_SUBCMD_PARSER_EPILOG from vllm.logger import init_logger from vllm.utils.argparse_utils import FlexibleArgumentParser diff --git a/vllm_omni/entrypoints/duplex/serving.py b/vllm_omni/entrypoints/duplex/serving.py index 3084376c62c..5933deca58d 100644 --- a/vllm_omni/entrypoints/duplex/serving.py +++ b/vllm_omni/entrypoints/duplex/serving.py @@ -51,6 +51,7 @@ InvalidResumeTokenError, ResumeToken, ) +from vllm_omni.entrypoints.duplex.warmup import wait_for_duplex_warmup from vllm_omni.entrypoints.duplex.websocket import ( MAX_EVENT_BYTES, SendJson, @@ -177,6 +178,9 @@ def __init__( async def handle_realtime_session(self, websocket: WebSocket) -> None: await websocket.accept() + # Complete the transport handshake before holding session admission; + # cold warmup can exceed the client's WebSocket opening timeout. + await wait_for_duplex_warmup(websocket) envelope = RealtimeEnvelope.from_query_params(websocket.query_params) async def send_json(payload: Mapping[str, object]) -> None: diff --git a/vllm_omni/entrypoints/duplex/warmup.py b/vllm_omni/entrypoints/duplex/warmup.py index 0bf2d59fe29..560a475ffc9 100644 --- a/vllm_omni/entrypoints/duplex/warmup.py +++ b/vllm_omni/entrypoints/duplex/warmup.py @@ -7,7 +7,7 @@ the hold-clients preamble is ``wait_for_duplex_warmup`` here. Video-required models default to one short audio-plus-frame turn so the four -pipeline stages compile real shapes before a client is accepted. That is +pipeline stages compile real shapes before a client session is admitted. That is separate from the per-stage JIT kernel registry, which can finish in ~0.01 s without compiling anything. """ diff --git a/vllm_omni/entrypoints/omni_base.py b/vllm_omni/entrypoints/omni_base.py index 7e2bb4b9a5d..ecd1c8ee226 100644 --- a/vllm_omni/entrypoints/omni_base.py +++ b/vllm_omni/entrypoints/omni_base.py @@ -13,6 +13,7 @@ import huggingface_hub import vllm.envs as envs from omegaconf import OmegaConf +from vllm.config.profiler import validate_profile_iteration_bounds from vllm.exceptions import RetryableRequestError from vllm.logger import init_logger from vllm.transformers_utils.repo_utils import file_or_path_exists @@ -521,7 +522,7 @@ def _handle_output_message( self._raise_nonfatal_error_message(msg) if not isinstance(msg, OutputMessage): - logger.warning("[%s] got unexpected msg type: %s", self.__class__.__name__, msg.type) + logger.warning("[%s] got unexpected msg type: %s", self.__class__.__name__, getattr(msg, "type")) return True, None, None, None req_id = msg.request_id @@ -794,6 +795,10 @@ def start_profile( self, profile_prefix: str | None = None, stages: list[int] | None = None, + *, + delay_iterations: int | None = None, + max_iterations: int | None = None, + profiler_kwargs: dict[str, Any] | None = None, ) -> list[Any]: """Start profiling specified stages. @@ -802,11 +807,34 @@ def start_profile( Args: profile_prefix: Optional prefix for the trace file names. stages: List of stage IDs to profile. If None, profiles all stages. + delay_iterations: Optional per-session delayed start (vLLM compat). + max_iterations: Optional per-session iteration cap (vLLM compat). + profiler_kwargs: Optional per-session profiler constructor overrides. Returns: List of results from each stage. """ - return self.engine.collective_rpc(method="profile", args=(True, profile_prefix), stage_ids=stages) + validate_profile_iteration_bounds(delay_iterations, max_iterations) + kwargs: dict[str, Any] = {} + if delay_iterations is not None: + kwargs["delay_iterations"] = delay_iterations + if max_iterations is not None: + kwargs["max_iterations"] = max_iterations + if profiler_kwargs is not None: + kwargs["profiler_kwargs"] = profiler_kwargs + # Omit kwargs when unset so existing profile RPC call shapes stay stable. + if kwargs: + return self.engine.collective_rpc( + method="profile", + args=(True, profile_prefix), + kwargs=kwargs, + stage_ids=stages, + ) + return self.engine.collective_rpc( + method="profile", + args=(True, profile_prefix), + stage_ids=stages, + ) def stop_profile(self, stages: list[int] | None = None) -> list[Any]: """Stop profiling specified stages. diff --git a/vllm_omni/entrypoints/openai/api_server.py b/vllm_omni/entrypoints/openai/api_server.py index 2d20737cbb5..904754308d0 100644 --- a/vllm_omni/entrypoints/openai/api_server.py +++ b/vllm_omni/entrypoints/openai/api_server.py @@ -35,8 +35,8 @@ from starlette.types import ASGIApp, Receive, Scope, Send from vllm.engine.protocol import EngineClient from vllm.entrypoints.anthropic.serving import AnthropicServingMessages +from vllm.entrypoints.common.cli_args import make_arg_parser, resolve_default_chat_template_kwargs from vllm.entrypoints.generate.base.protocol import RequestResponseMetadata -from vllm.entrypoints.launchers.cli_args import make_arg_parser from vllm.entrypoints.launchers.launcher import serve_http from vllm.entrypoints.launchers.utils.server_utils import get_uvicorn_log_config from vllm.entrypoints.mcp.tool_server import DemoToolServer, MCPToolServer, ToolServer @@ -71,6 +71,7 @@ validate_json_request, with_cancellation, ) +from vllm.entrypoints.serve.utils.fingerprint import set_default_fingerprint_mode from vllm.entrypoints.serve.utils.orca_metrics import metrics_header from vllm.entrypoints.serve.utils.request_logger import RequestLogger from vllm.entrypoints.speech_to_text.transcription.serving import ( @@ -81,6 +82,7 @@ ) from vllm.logger import configure_logging_from_args, init_logger from vllm.renderers.chat_utils import ChatTemplateConfig, load_chat_template +from vllm.renderers.online_derenderer import OnlineDerenderer from vllm.renderers.online_renderer import OnlineRenderer from vllm.tasks import POOLING_TASKS from vllm.tool_parsers import ToolParserManager @@ -105,7 +107,6 @@ _warmup_duplex_realtime, lookup_duplex_plugin, startup_warmup_kind, - wait_for_duplex_warmup, ) from vllm_omni.entrypoints.duplex_omni import DuplexOmni from vllm_omni.entrypoints.openai import app_state as openai_app_state @@ -786,6 +787,7 @@ async def _init_duplex_app_state( state.serving_tokens = None # Replaced by the chat init below when the model serves chat. state.online_renderer = None + state.online_derenderer = None for attribute in ( "openai_serving_chat_batch", "openai_serving_completion", @@ -880,7 +882,20 @@ def _init_chat_services( enable_batch: bool, ) -> None: """Share chat setup while callers retain capability and endpoint gating.""" - state.online_renderer = OnlineRenderer( + # Serving classes cache system_fingerprint at init. Match upstream + # init_generate_state: configure the mode before those constructors run. + set_default_fingerprint_mode( + getattr(args, "fingerprint_mode", "full"), + getattr(args, "fingerprint_value", None), + ) + default_chat_template_kwargs = resolve_default_chat_template_kwargs(args) + # Pin CLI fields. getattr keeps test namespaces that predate these flags + # on the same disabled defaults the upstream constructors use. + trust_request_mm_kwargs = getattr(args, "trust_request_mm_kwargs", False) + tool_strict_level = getattr(args, "tool_strict_level", None) + reasoning_effort_map = getattr(args, "reasoning_effort_map", None) + reasoning_parser = args.structured_outputs_config.reasoning_parser + renderer_kwargs = dict( model_config=engine_client.model_config, renderer=engine_client.renderer, request_logger=request_logger, @@ -890,10 +905,17 @@ def _init_chat_services( enable_auto_tools=args.enable_auto_tool_choice, exclude_tools_when_tool_choice_none=args.exclude_tools_when_tool_choice_none, tool_parser=args.tool_call_parser, - reasoning_parser=args.structured_outputs_config.reasoning_parser, - default_chat_template_kwargs=args.default_chat_template_kwargs, + tool_strict_level=tool_strict_level, + reasoning_parser=reasoning_parser, + default_chat_template_kwargs=default_chat_template_kwargs, + reasoning_effort_map=reasoning_effort_map, log_error_stack=args.log_error_stack, ) + state.online_renderer = OnlineRenderer( + **renderer_kwargs, + trust_request_mm_kwargs=trust_request_mm_kwargs, + ) + state.online_derenderer = OnlineDerenderer(**renderer_kwargs) chat_kwargs = dict( engine_client=engine_client, models=state.openai_serving_models, @@ -902,17 +924,19 @@ def _init_chat_services( request_logger=request_logger, chat_template=resolved_chat_template, chat_template_content_format=args.chat_template_content_format, - default_chat_template_kwargs=args.default_chat_template_kwargs, + default_chat_template_kwargs=default_chat_template_kwargs, trust_request_chat_template=args.trust_request_chat_template, return_tokens_as_token_ids=args.return_tokens_as_token_ids, enable_auto_tools=args.enable_auto_tool_choice, exclude_tools_when_tool_choice_none=args.exclude_tools_when_tool_choice_none, tool_parser=args.tool_call_parser, - reasoning_parser=args.structured_outputs_config.reasoning_parser, + tool_strict_level=tool_strict_level, + reasoning_parser=reasoning_parser, enable_prompt_tokens_details=args.enable_prompt_tokens_details, enable_force_include_usage=args.enable_force_include_usage, enable_log_outputs=args.enable_log_outputs, enable_log_deltas=args.enable_log_deltas, + enable_per_request_metrics=getattr(args, "enable_per_request_metrics", False), ) state.openai_serving_chat = OmniOpenAIServingChat(**chat_kwargs) if enable_chat else None state.openai_serving_chat_batch = ( @@ -1070,11 +1094,14 @@ async def omni_init_app_state( logger.info("Supported tasks: %s", supported_tasks) resolved_chat_template = await _resolve_chat_template(engine_client, args) + default_chat_template_kwargs = resolve_default_chat_template_kwargs(args) + tool_strict_level = getattr(args, "tool_strict_level", None) chat_template_config = ChatTemplateConfig( chat_template=resolved_chat_template, chat_template_content_format=args.chat_template_content_format, trust_request_chat_template=args.trust_request_chat_template, + trust_request_mm_kwargs=getattr(args, "trust_request_mm_kwargs", False), ) if args.tool_server == "demo": @@ -1162,11 +1189,14 @@ async def omni_init_app_state( return_tokens_as_token_ids=args.return_tokens_as_token_ids, enable_auto_tools=args.enable_auto_tool_choice, tool_parser=args.tool_call_parser, + tool_strict_level=tool_strict_level, tool_server=tool_server, reasoning_parser=args.structured_outputs_config.reasoning_parser, enable_prompt_tokens_details=args.enable_prompt_tokens_details, enable_force_include_usage=args.enable_force_include_usage, + enable_per_request_metrics=getattr(args, "enable_per_request_metrics", False), enable_log_outputs=args.enable_log_outputs, + default_chat_template_kwargs=default_chat_template_kwargs, ) if "generate" in supported_tasks else None @@ -1181,6 +1211,9 @@ async def omni_init_app_state( return_tokens_as_token_ids=args.return_tokens_as_token_ids, enable_prompt_tokens_details=args.enable_prompt_tokens_details, enable_force_include_usage=args.enable_force_include_usage, + enable_log_outputs=args.enable_log_outputs, + enable_log_deltas=getattr(args, "enable_log_deltas", True), + enable_per_request_metrics=getattr(args, "enable_per_request_metrics", False), ) if "generate" in supported_tasks else None @@ -1234,7 +1267,7 @@ async def omni_init_app_state( request_logger=request_logger, chat_template=resolved_chat_template, chat_template_content_format=args.chat_template_content_format, - default_chat_template_kwargs=args.default_chat_template_kwargs, + default_chat_template_kwargs=default_chat_template_kwargs, trust_request_chat_template=args.trust_request_chat_template, ) state.openai_serving_transcription = ( @@ -1269,10 +1302,14 @@ async def omni_init_app_state( return_tokens_as_token_ids=args.return_tokens_as_token_ids, enable_auto_tools=args.enable_auto_tool_choice, tool_parser=args.tool_call_parser, + tool_strict_level=tool_strict_level, reasoning_parser=args.structured_outputs_config.reasoning_parser, enable_prompt_tokens_details=args.enable_prompt_tokens_details, enable_force_include_usage=args.enable_force_include_usage, - default_chat_template_kwargs=args.default_chat_template_kwargs, + enable_log_outputs=getattr(args, "enable_log_outputs", False), + enable_log_deltas=getattr(args, "enable_log_deltas", True), + default_chat_template_kwargs=default_chat_template_kwargs, + disabled_thinking_effort=getattr(args, "anthropic_disabled_thinking_effort", "auto"), ) if "generate" in supported_tasks else None @@ -1282,6 +1319,7 @@ async def omni_init_app_state( engine_client, state.openai_serving_models, state.online_renderer, + online_derenderer=state.online_derenderer, request_logger=request_logger, return_tokens_as_token_ids=args.return_tokens_as_token_ids, enable_prompt_tokens_details=args.enable_prompt_tokens_details, @@ -1914,7 +1952,6 @@ async def realtime_websocket(websocket: WebSocket): if use_duplex_realtime and duplex_handler is not None: if await _reject_multi_api_duplex(websocket): return - await wait_for_duplex_warmup(websocket) await duplex_handler.handle_realtime_session(websocket) return @@ -1944,7 +1981,6 @@ async def duplex_websocket(websocket: WebSocket): """Alias of ``/v1/realtime?duplex=1``: the same Realtime duplex session protocol.""" if await _reject_multi_api_duplex(websocket): return - await wait_for_duplex_warmup(websocket) handler = getattr(websocket.app.state, "openai_serving_duplex", None) if handler is None: await websocket.accept() @@ -3419,7 +3455,19 @@ async def start_profile(raw_request: Request, request: ProfileRequest | None = N stages = request.stages if request else None logger.info("Starting profiler for stages: %s", stages if stages else "all") engine_client = raw_request.app.state.engine_client - await engine_client.start_profile(stages=stages) + # Forward only supplied per-session overrides. AsyncOmni.start_profile + # already accepts these; a stages-only body stays a stages-only call. + profile_kwargs: dict[str, Any] = {} + if request is not None: + if request.profile_prefix is not None: + profile_kwargs["profile_prefix"] = request.profile_prefix + if request.delay_iterations is not None: + profile_kwargs["delay_iterations"] = request.delay_iterations + if request.max_iterations is not None: + profile_kwargs["max_iterations"] = request.max_iterations + if request.profiler_kwargs is not None: + profile_kwargs["profiler_kwargs"] = request.profiler_kwargs + await engine_client.start_profile(stages=stages, **profile_kwargs) logger.info("Profiler started.") return JSONResponse(content={"status": "SUCCESS"}) except Exception as e: diff --git a/vllm_omni/entrypoints/openai/realtime/connection.py b/vllm_omni/entrypoints/openai/realtime/connection.py index 419261aeefe..9b0bc9906c1 100644 --- a/vllm_omni/entrypoints/openai/realtime/connection.py +++ b/vllm_omni/entrypoints/openai/realtime/connection.py @@ -1911,7 +1911,15 @@ async def _build_full_prompt( if text: messages.append({"role": "system", "content": text}) - request = ChatCompletionRequest(model=self.model_name, messages=messages) + # Prompt-length probes reserve output tokens themselves. Pin + # ChatCompletionRequest.build_tok_params leaves max_output_tokens + # unset when max_tokens is omitted, so pass 0 to measure the prompt + # against the full context window. + request = ChatCompletionRequest( + model=self.model_name, + messages=messages, + max_completion_tokens=0, + ) _, (engine_input,) = await chat_handler._preprocess_chat( request, messages, diff --git a/vllm_omni/entrypoints/openai/serving_chat.py b/vllm_omni/entrypoints/openai/serving_chat.py index a802d091a09..4ce4e64b7ed 100644 --- a/vllm_omni/entrypoints/openai/serving_chat.py +++ b/vllm_omni/entrypoints/openai/serving_chat.py @@ -2783,6 +2783,9 @@ def _create_text_choice( else (output.finish_reason if output.finish_reason else "stop") ), stop_reason=output.stop_reason, + weight_versions=( + getattr(output, "weight_versions", None) if request.return_weight_versions else None + ), ) choices.append(choice_data) continue @@ -2920,6 +2923,7 @@ def _create_text_choice( ), stop_reason=output.stop_reason, token_ids=(as_list(output.token_ids) if request.return_token_ids else None), + weight_versions=(getattr(output, "weight_versions", None) if request.return_weight_versions else None), ) choices.append(choice_data) diff --git a/vllm_omni/entrypoints/serve/profile/protocol.py b/vllm_omni/entrypoints/serve/profile/protocol.py index ba424e924a7..8730458dd65 100644 --- a/vllm_omni/entrypoints/serve/profile/protocol.py +++ b/vllm_omni/entrypoints/serve/profile/protocol.py @@ -2,13 +2,23 @@ # SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project """Request models for Omni profile control routes.""" +from typing import Any + from pydantic import BaseModel, Field class ProfileRequest(BaseModel): - """Request model for profiling endpoints.""" + """Request model for profiling endpoints. + + ``stages`` is Omni-specific. The other fields match upstream + ``StartProfileRequest`` and are forwarded to ``start_profile``. + """ stages: list[int] | None = Field( default=None, description="List of stage IDs to profile. If None, profiles all stages.", ) + profile_prefix: str | None = None + delay_iterations: int | None = None + max_iterations: int | None = None + profiler_kwargs: dict[str, Any] | None = None diff --git a/vllm_omni/model_executor/models/cosyvoice3/v1_sampling.py b/vllm_omni/model_executor/models/cosyvoice3/v1_sampling.py index aa429c2f479..3ba00364dcd 100644 --- a/vllm_omni/model_executor/models/cosyvoice3/v1_sampling.py +++ b/vllm_omni/model_executor/models/cosyvoice3/v1_sampling.py @@ -12,7 +12,7 @@ import torch from vllm.triton_utils import tl, triton -from vllm.v1.sample.ops import topk_topp_sampler +from vllm.v1.sample.ops.topk_topp_sampler import random_sample from vllm_omni.platforms import current_omni_platform from vllm_omni.utils.seeded_exponential import torch_exponential_noise, torch_exponential_policy @@ -165,7 +165,7 @@ def _use_triton(probs: torch.Tensor) -> bool: def categorical_sample(probs: torch.Tensor, generators: dict[int, torch.Generator]) -> torch.Tensor: """Return an int64 token vector; replacements supply only rejected rows.""" if not _use_triton(probs): - return topk_topp_sampler.random_sample(probs, generators).reshape(-1).long() + return random_sample(probs, generators).reshape(-1).long() return _cuda_sample(probs, generators)[0] @@ -191,7 +191,7 @@ def nucleus_sample( keep = (cumulative - probs) < top_p.unsqueeze(1) ranks = torch.arange(probs.shape[1], device=probs.device) keep &= (top_k.unsqueeze(1) <= 0) | (ranks.unsqueeze(0) < top_k.unsqueeze(1)) - draws = topk_topp_sampler.random_sample(probs * keep.to(probs.dtype), generators).reshape(-1, 1).long() + draws = random_sample(probs * keep.to(probs.dtype), generators).reshape(-1, 1).long() sampled = indices.gather(1, draws).squeeze(1) repeated = None if history is not None: diff --git a/vllm_omni/model_executor/models/fish_omni/model_state.py b/vllm_omni/model_executor/models/fish_omni/model_state.py index 55db893bbae..f0f695ae585 100644 --- a/vllm_omni/model_executor/models/fish_omni/model_state.py +++ b/vllm_omni/model_executor/models/fish_omni/model_state.py @@ -24,9 +24,9 @@ import numpy as np import torch +import vllm.v1.worker.gpu.buffer_utils as buffer_utils from vllm.logger import init_logger from vllm.utils.platform_utils import is_uva_available -from vllm.v1.worker.gpu import buffer_utils from vllm.v1.worker.gpu.input_batch import get_num_sampled_and_rejected from vllm.v1.worker.gpu.model_states.mamba_hybrid import MambaHybridModelState from vllm.v1.worker.gpu.sample.output import SamplerOutput diff --git a/vllm_omni/model_executor/models/moss_tts/local_kernels.py b/vllm_omni/model_executor/models/moss_tts/local_kernels.py index 6eb9f75052c..1e485bc8279 100644 --- a/vllm_omni/model_executor/models/moss_tts/local_kernels.py +++ b/vllm_omni/model_executor/models/moss_tts/local_kernels.py @@ -79,10 +79,13 @@ def _attend( scores = tl.where(s[None, :] == POSITION, sp[:, None], scores) probs = tl.exp(scores - tl.max(scores, axis=1)[:, None]) probs = probs / tl.sum(probs, axis=1)[:, None] - result = v * tl.sum(tl.where(s[None, :] == POSITION, probs, 0.0), axis=1)[:, None] + # Pin the accumulator's rounding. LLVM SLP can otherwise contract the + # initial product with the first add differently when this body is + # inlined into sampling on SM89, breaking fused/unfused bitwise parity. + result = libdevice.mul_rn(v, tl.sum(tl.where(s[None, :] == POSITION, probs, 0.0), axis=1)[:, None]) for p in tl.static_range(POSITION): vp = tl.load(V + base + p * D, dmask, 0, eviction_policy=CACHE_LOAD).to(tl.float32) - result += vp * tl.sum(tl.where(s[None, :] == p, probs, 0.0), axis=1)[:, None] + result = tl.fma(vp, tl.sum(tl.where(s[None, :] == p, probs, 0.0), axis=1)[:, None], result) if LATE and not LAST: tl.store(K + base + POSITION * D, k, dmask, eviction_policy="evict_last") tl.store(V + base + POSITION * D, v, dmask, eviction_policy="evict_last") diff --git a/vllm_omni/model_executor/models/moss_tts/local_model_state.py b/vllm_omni/model_executor/models/moss_tts/local_model_state.py index 5cbac9e5fd8..8638d967ec7 100644 --- a/vllm_omni/model_executor/models/moss_tts/local_model_state.py +++ b/vllm_omni/model_executor/models/moss_tts/local_model_state.py @@ -9,12 +9,12 @@ import numpy as np import torch +import vllm.v1.worker.gpu.buffer_utils as buffer_utils from vllm.config import CUDAGraphMode from vllm.forward_context import set_forward_context from vllm.logger import init_logger from vllm.utils.platform_utils import is_uva_available from vllm.utils.torch_utils import async_tensor_h2d -from vllm.v1.worker.gpu import buffer_utils from vllm.v1.worker.gpu.input_batch import get_num_sampled_and_rejected from vllm.v1.worker.gpu.sample.bad_words import BadWordsState from vllm.v1.worker.gpu.sample.logit_bias import LogitBiasState diff --git a/vllm_omni/model_executor/models/moss_tts/packed_codec.py b/vllm_omni/model_executor/models/moss_tts/packed_codec.py index 4cbf17bcd76..49cc56eae74 100644 --- a/vllm_omni/model_executor/models/moss_tts/packed_codec.py +++ b/vllm_omni/model_executor/models/moss_tts/packed_codec.py @@ -29,12 +29,12 @@ import numpy as np import torch import torch.nn as nn +import vllm.v1.worker.gpu.buffer_utils as buffer_utils from torch.cuda import CUDAGraph from vllm.config import VllmConfig from vllm.logger import init_logger from vllm.platforms import current_platform from vllm.triton_utils import tl, triton -from vllm.v1.worker.gpu import buffer_utils logger = init_logger(__name__) diff --git a/vllm_omni/model_executor/models/qwen3_omni/talker_sampler.py b/vllm_omni/model_executor/models/qwen3_omni/talker_sampler.py index ea0148c0215..65088c67cf2 100644 --- a/vllm_omni/model_executor/models/qwen3_omni/talker_sampler.py +++ b/vllm_omni/model_executor/models/qwen3_omni/talker_sampler.py @@ -20,7 +20,7 @@ import torch from vllm.logger import init_logger from vllm.sampling_params import SamplingParams -from vllm.v1.sample.ops.topk_topp_sampler import apply_top_k_top_p, random_sample +from vllm.v1.sample.ops.topk_topp_sampler import apply_top_k_top_p, random_sample, register_top_k_top_p_warmups from vllm.v1.worker.gpu.input_batch import InputBatch from vllm.v1.worker.gpu.sample.output import SamplerOutput from vllm.v1.worker.gpu.sample.sampler import Sampler @@ -35,6 +35,9 @@ class Qwen3OmniTalkerSampler(OmniSampler): def __init__(self, base_sampler: Sampler): super().__init__(base_sampler) + # Seeded rows use native masking even when upstream selects FlashInfer. + # Declare its kernels while the runner collects custom-sampler warmups. + register_top_k_top_p_warmups() self._generators: dict[int, torch.Generator] = {} self._sample_random = base_sampler._sample_random base_sampler._sample_random = self._sample_like_v1 @@ -81,6 +84,7 @@ def _sample_like_v1( pos: torch.Tensor, top_k: torch.Tensor | None, top_p: torch.Tensor | None, + has_topp_only_rows: bool, use_fused_sampler: bool, ) -> tuple[torch.Tensor, torch.Tensor]: original: Callable[..., tuple[torch.Tensor, torch.Tensor]] = self._sample_random @@ -107,9 +111,16 @@ def _sample_like_v1( ) self._logged_fallback = True return original( - processed_logits, expanded_idx_mapping, idx_mapping_np, pos, top_k, top_p, use_fused_sampler + processed_logits, + expanded_idx_mapping, + idx_mapping_np, + pos, + top_k, + top_p, + has_topp_only_rows, + use_fused_sampler, ) - masked = apply_top_k_top_p(processed_logits, top_k, top_p) + masked = apply_top_k_top_p(processed_logits, top_k, top_p, has_topp_only_rows=has_topp_only_rows) if len(generators) == n_rows: if not self._logged_v1: logger.info("Qwen3-Omni Talker sampled %d rows with the V1 generator", n_rows) @@ -141,6 +152,7 @@ def _sample_like_v1( None, None, False, + False, ) sampled[other_idx] = other_sampled.to(torch.int64) return sampled, masked diff --git a/vllm_omni/outputs/output_processor.py b/vllm_omni/outputs/output_processor.py index 5e7c6bfb4c8..7ad0f50d25c 100644 --- a/vllm_omni/outputs/output_processor.py +++ b/vllm_omni/outputs/output_processor.py @@ -710,6 +710,7 @@ def process_outputs( engine_core_outputs: list[EngineCoreOutput], engine_core_timestamp: float | None = None, iteration_stats: IterationStats | None = None, + weight_version: str | None = None, ) -> OutputProcessorOutput: default_mm_type = _modality_to_type_string(self.output_modality) @@ -718,12 +719,40 @@ def process_outputs( # that would trigger upstream's `assert detokenizer is not None`. upstream_outputs: list[EngineCoreOutput] = [] mm_only_outputs: list[EngineCoreOutput] = [] + admission_outputs: list[OmniRequestOutput] = [] + admission_aborts: list[str] = [] for eco in engine_core_outputs: req_state = self.request_states.get(eco.request_id) if req_state is None: continue + # Paused-engine and other pre-execution failures ship RequestError + # and nothing was run. StagePool registers queue=None, so upstream's + # queue.put short-circuit does not run here or in super(). Key on + # error, not finish_reason: pooling cache misses use ERROR with + # error is None and must stay on the normal path. weight_version is + # forwarded only into super(), which is the detokenizer path. + error = getattr(eco, "error", None) + if error is not None: + if req_state.queue is not None: + req_state.queue.put(error.to_exception()) + else: + message = getattr(error, "message", None) + code = getattr(error, "code", None) + external_id = getattr(req_state, "external_req_id", None) or eco.request_id + admission_outputs.append( + OmniRequestOutput.from_error( + external_id, + message if isinstance(message, str) and message else str(error), + error_type=code if isinstance(code, str) else None, + ) + ) + if req_state.streaming_input: + admission_aborts.append(eco.request_id) + self._finish_request(req_state) + continue + # Accumulate multimodal tensors regardless of path. if isinstance(req_state, OmniRequestState): mm_type = getattr(eco, "output_type", None) or default_mm_type @@ -770,10 +799,11 @@ def process_outputs( upstream_outputs, engine_core_timestamp=engine_core_timestamp, iteration_stats=iteration_stats, + weight_version=weight_version, ) return OutputProcessorOutput( - request_outputs=mm_request_outputs + upstream_processed.request_outputs, - reqs_to_abort=mm_reqs_to_abort + upstream_processed.reqs_to_abort, + request_outputs=admission_outputs + mm_request_outputs + upstream_processed.request_outputs, + reqs_to_abort=admission_aborts + mm_reqs_to_abort + upstream_processed.reqs_to_abort, ) def _process_mm_only_outputs( diff --git a/vllm_omni/patch.py b/vllm_omni/patch.py index 7f34a750518..a46d3852739 100644 --- a/vllm_omni/patch.py +++ b/vllm_omni/patch.py @@ -29,12 +29,14 @@ from vllm_omni.inputs.data import OmniTokensPrompt from vllm_omni.model_executor.layers.rotary_embedding import OmniMRotaryEmbedding from vllm_omni.request import OmniRequest, OmniStreamingUpdate +from vllm_omni.utils.mm_prefix import apply_img2img_mm_prefix_patch _PATCH_LOGGER = logging.getLogger("vllm_omni.patch") # Decoder rounding is visible to audio encoders and can change mixed-media # answers. Preserve Omni's existing default while retaining explicit choices. apply_soundfile_first_audio_decode_patch() +apply_img2img_mm_prefix_patch() # Omni's deploy config selects a runner per stage. vLLM 0.31's config property # otherwise defaults to v2 independently of the runner selected by our workers, @@ -730,7 +732,7 @@ def _patched_free_callback(self, ptr: int) -> tuple: data.cpu_backup_tensor = None if data.is_asleep: device, size, d_mem, _ = data.handle - result = (device, size, d_mem, []) + result: tuple = (device, size, d_mem, []) _PATCH_LOGGER.debug( "[xpumem] Free callback: asleep entry %s -> empty handle", ptr, diff --git a/vllm_omni/platforms/cuda/platform.py b/vllm_omni/platforms/cuda/platform.py index 7c6cf55c60f..1d34a1a057d 100644 --- a/vllm_omni/platforms/cuda/platform.py +++ b/vllm_omni/platforms/cuda/platform.py @@ -116,12 +116,6 @@ def get_diffusion_attn_backend_cls( else: flash_attn_supported = compute_supported and packages_available - # cuDNN 9.5+ ships Blackwell FMHA kernels. If the runtime is older, - # skip the CUDNN_ATTN default rather than selecting a backend whose - # kernel selector may reject the runtime shape. - cudnn_version = torch.backends.cudnn.version() or 0 - cudnn_blackwell_ready = cudnn_version >= 90500 - # FlashInfer edges cuDNN by ~4% at the kernel level on sm_120 but # regresses ~2x at e2e on HV-1.5 because its dense-prefill path cannot # represent every diffusion attention mask without changing semantics. @@ -232,7 +226,7 @@ def get_diffusion_attn_backend_cls( ) if backend_upper in ("FLASH_ATTN_HUB", "FLASH_ATTN_3_HUB"): - # Same loader/cache as execution (Hub versions 1 then 2). A missing + # Same loader/cache as execution, including FA2 stable-ABI builds. A missing # torch/CUDA variant used to degrade silently to native/SDPA (#6971). hub_repo = ( "kernels-community/flash-attn3" @@ -329,7 +323,11 @@ def get_diffusion_attn_backend_cls( ) return DiffusionAttentionBackendEnum.TRTLLM_ATTN.get_path() - if is_blackwell and cudnn_blackwell_ready: + # Torch 2.14 can initialize CUDA while querying the cuDNN version. + # Only the automatic Blackwell fallback needs this probe; explicit + # selections and other architectures must not initialize cuDNN here. + cudnn_version = (torch.backends.cudnn.version() or 0) if is_blackwell else 0 + if is_blackwell and cudnn_version >= 90500: from vllm_omni.diffusion.attention.backends.cudnn_attn import CuDNNAttentionBackend # Unknown head_size is a capability probe (SP auto-pad), not a @@ -360,7 +358,7 @@ def get_diffusion_attn_backend_cls( ) return DiffusionAttentionBackendEnum.FLASHINFER_ATTN.get_path() - if is_blackwell and not cudnn_blackwell_ready: + if is_blackwell and cudnn_version < 90500: logger.warning( "Detected Blackwell %s but cuDNN %d < 9.5 — no tuned Blackwell FMHA. " "Automatic backend selection will use FlashInfer when available, otherwise SDPA.", diff --git a/vllm_omni/platforms/npu/profiler.py b/vllm_omni/platforms/npu/profiler.py index 40bf5f91ead..52c1ae0e47a 100644 --- a/vllm_omni/platforms/npu/profiler.py +++ b/vllm_omni/platforms/npu/profiler.py @@ -1,5 +1,5 @@ # SPDX-License-Identifier: Apache-2.0 -# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project from __future__ import annotations @@ -67,7 +67,7 @@ def _create_profiler( os.makedirs(npu_trace_dir, exist_ok=True) self._trace_path = npu_trace_dir - return torch_npu.profiler.profile( + kwargs = dict( activities=npu_activities, with_stack=False, profile_memory=profiler_config.torch_profiler_with_memory, @@ -83,6 +83,8 @@ def _create_profiler( # this specific handler format on_trace_ready=torch_npu.profiler.tensorboard_trace_handler(npu_trace_dir), ) + kwargs.update(self._profiler_kwargs) + return torch_npu.profiler.profile(**kwargs) @override def _on_stop_hook(self) -> None: diff --git a/vllm_omni/platforms/npu/worker/base.py b/vllm_omni/platforms/npu/worker/base.py index 97f9f84c6e1..e7d6561627b 100644 --- a/vllm_omni/platforms/npu/worker/base.py +++ b/vllm_omni/platforms/npu/worker/base.py @@ -71,12 +71,14 @@ def profile( *, delay_iterations: int | None = None, max_iterations: int | None = None, + profiler_kwargs: dict | None = None, ): """Override to set trace filename before starting the profiler. NPUWorker's profile() accepts profile_prefix, so we use it to generate a descriptive trace filename for OmniProfiler. - Per-session iteration controls are handled by the profiler wrapper. + Per-session iteration controls and constructor overrides are handled + by the profiler wrapper. """ if self.profiler is None: raise RuntimeError( @@ -92,6 +94,7 @@ def profile( stage_id = getattr(self.vllm_config.model_config, "stage_id", 0) filename = profile_prefix or f"stage{stage_id}_rank{self.rank}_{int(time.time())}" self.profiler.set_trace_filename(filename) + self.profiler.set_profiler_kwargs(profiler_kwargs) kwargs: dict[str, int] = {} if delay_iterations is not None: kwargs["delay_iterations"] = delay_iterations diff --git a/vllm_omni/platforms/rocm/platform.py b/vllm_omni/platforms/rocm/platform.py index d29e9043742..3d8377d36e6 100644 --- a/vllm_omni/platforms/rocm/platform.py +++ b/vllm_omni/platforms/rocm/platform.py @@ -208,6 +208,7 @@ def get_default_ir_op_priority(cls, vllm_config: VllmConfig) -> IrOpPriorityConf Preserves omni-specific is_custom_op_enabled('rms_norm') check. """ from vllm.config.compilation import CompilationMode + from vllm.platforms.rocm import on_gfx90a # TODO(luka/TJ) use aiter, vllm_c, native by default on ROCm cc = vllm_config.compilation_config @@ -218,7 +219,14 @@ def get_default_ir_op_priority(cls, vllm_config: VllmConfig) -> IrOpPriorityConf # Necessary on ROCm because it's common that users # enable rms_norm to use the aiter kernel. # TODO(luka/TJ) remove env vars completely - if cc.is_custom_op_enabled("rms_norm") and envs.VLLM_ROCM_USE_AITER and envs.VLLM_ROCM_USE_AITER_RMSNORM: + # AITER rmsnorm is not built for gfx90a and segfaults (upstream 2622ee21aa). + # Keep the Omni custom-op gate; do not switch to cudagraph_mode / on_rdna4. + if ( + cc.is_custom_op_enabled("rms_norm") + and envs.VLLM_ROCM_USE_AITER + and envs.VLLM_ROCM_USE_AITER_RMSNORM + and not on_gfx90a() + ): rms_norm = ["aiter"] + default else: rms_norm = default diff --git a/vllm_omni/profiler/omni_torch_profiler.py b/vllm_omni/profiler/omni_torch_profiler.py index 69f893b05fa..42245b1866e 100644 --- a/vllm_omni/profiler/omni_torch_profiler.py +++ b/vllm_omni/profiler/omni_torch_profiler.py @@ -1,5 +1,5 @@ # SPDX-License-Identifier: Apache-2.0 -# SPDX-FileCopyrightText: Copyright contributors to the vLLM project +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project from __future__ import annotations @@ -9,7 +9,6 @@ from typing import Any, Literal import torch -from typing_extensions import override from vllm.config import ProfilerConfig from vllm.config.profiler import _is_uri_path from vllm.logger import init_logger @@ -63,6 +62,7 @@ def __init__( self._table_path: str | None = None self._activities = activities + self._profiler_kwargs: dict[str, Any] = {} self._session_dir: str | None = None self._artifact_paths: dict[str, str | None] = {} self._memory_history_enabled = False @@ -98,14 +98,23 @@ def _create_profiler( Override in subclasses for platform-specific profiler creation. """ - return torch.profiler.profile( - activities=[TorchProfilerActivityMap[a] for a in activities], - record_shapes=profiler_config.torch_profiler_record_shapes, - profile_memory=profiler_config.torch_profiler_with_memory, - with_stack=profiler_config.torch_profiler_with_stack, - with_flops=profiler_config.torch_profiler_with_flops, - on_trace_ready=self._on_trace_ready, - ) + kwargs = { + "activities": [TorchProfilerActivityMap[a] for a in activities], + "record_shapes": profiler_config.torch_profiler_record_shapes, + "profile_memory": profiler_config.torch_profiler_with_memory, + "with_stack": profiler_config.torch_profiler_with_stack, + "with_flops": profiler_config.torch_profiler_with_flops, + "on_trace_ready": self._on_trace_ready, + } + kwargs.update(self._profiler_kwargs) + return torch.profiler.profile(**kwargs) + + def set_profiler_kwargs(self, profiler_kwargs: dict | None) -> None: + """Apply constructor overrides to the next inactive profiling session.""" + if self._active: + return + self._profiler_kwargs = dict(profiler_kwargs or {}) + self.profiler = self._create_profiler(self.profiler_config, self._activities) def set_trace_filename(self, filename: str) -> None: """Set the trace filename before starting profiling. @@ -383,13 +392,11 @@ def _write_excel_artifact(self, name: str, sheets: dict[str, list[dict]]) -> str self._artifact_paths[name] = path return path - @override def _start(self) -> None: self._ensure_session_dir() self._try_enable_memory_history() self.profiler.start() - @override def _stop(self) -> None: """Stop profiler, export trace via on_trace_ready, and dump table.""" self.profiler.stop() @@ -488,7 +495,6 @@ def _on_stop_hook(self) -> None: if self.dump_cpu_time_total and rank == 0: logger.info(self.profiler.key_averages().table(sort_by="self_cpu_time_total", row_limit=50)) - @override def annotate_context_manager(self, name: str): return torch.profiler.record_function(name) diff --git a/vllm_omni/utils/mm_prefix.py b/vllm_omni/utils/mm_prefix.py new file mode 100644 index 00000000000..14be63da7b1 --- /dev/null +++ b/vllm_omni/utils/mm_prefix.py @@ -0,0 +1,64 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project +"""Keep Bagel's custom image modality visible to vLLM's prefix-range builder.""" + +from __future__ import annotations + +import sys +from collections.abc import Sequence +from dataclasses import replace +from functools import wraps +from typing import TYPE_CHECKING + +if TYPE_CHECKING: + from vllm.multimodal.inputs import MultiModalFeatureSpec + +_RUNNER_MODULES = ( + "vllm.v1.worker.gpu_model_runner", + "vllm.v1.worker.gpu.model_states.default", +) + + +def apply_img2img_mm_prefix_patch() -> None: + from vllm.v1.attention.backends import utils + + original = utils.compute_mm_prefix_ranges + if getattr(original, "_omni_img2img_mm_prefix", False): + return + + @wraps(original) + def compute_mm_prefix_ranges( + mm_features_per_req: Sequence[Sequence[MultiModalFeatureSpec]], + sliding_window: int | None = None, + clamp_sliding_window: bool = False, + span_pad: int = 0, + ) -> dict[int, list[tuple[int, int]]]: + # Alias only the prefix builder's host-side views. Encoder batching, + # kwargs, caches, and embedding modality tags keep their native keys. + if any(feature.modality == "img2img" for features in mm_features_per_req for feature in features): + mm_features_per_req = [ + [ + replace(feature, modality="image") if feature.modality == "img2img" else feature + for feature in features + ] + for features in mm_features_per_req + ] + return original(mm_features_per_req, sliding_window, clamp_sliding_window, span_pad) + + # Validate previously imported aliases before changing any module binding. + loaded_runners = [] + for name in _RUNNER_MODULES: + module = sys.modules.get(name) + if module is None: + continue + alias = getattr(module, "compute_mm_prefix_ranges", None) + if alias is not None and alias is not original: + raise RuntimeError(f"Conflicting compute_mm_prefix_ranges alias in {name}") + if alias is original: + loaded_runners.append(module) + + setattr(compute_mm_prefix_ranges, "_omni_img2img_mm_prefix", True) + setattr(compute_mm_prefix_ranges, "_omni_mm_prefix_original", original) + utils.compute_mm_prefix_ranges = compute_mm_prefix_ranges + for module in loaded_runners: + setattr(module, "compute_mm_prefix_ranges", compute_mm_prefix_ranges) diff --git a/vllm_omni/worker/base.py b/vllm_omni/worker/base.py index cc301325f5a..82fded6ea76 100644 --- a/vllm_omni/worker/base.py +++ b/vllm_omni/worker/base.py @@ -199,6 +199,7 @@ def profile( *, delay_iterations: int | None = None, max_iterations: int | None = None, + profiler_kwargs: dict | None = None, ): """Override to set trace filename before starting the profiler. @@ -207,6 +208,7 @@ def profile( profile_prefix: Optional prefix for trace filename (vLLM compat). delay_iterations: Optional per-session delayed start (vLLM compat). max_iterations: Optional per-session iteration cap (vLLM compat). + profiler_kwargs: Optional per-session profiler constructor overrides. Generate a stage-specific trace filename before delegating the lifecycle and iteration controls to the profiler. @@ -225,6 +227,9 @@ def profile( stage_id = getattr(self.vllm_config.model_config, "stage_id", 0) filename = profile_prefix or f"stage{stage_id}_rank{self.rank}_{int(time.time())}" self.profiler.set_trace_filename(filename) + # WorkerProfiler.start accepts only the iteration bounds. + # Session kwargs go through set_profiler_kwargs, matching Worker.profile. + self.profiler.set_profiler_kwargs(profiler_kwargs) kwargs: dict[str, int] = {} if delay_iterations is not None: kwargs["delay_iterations"] = delay_iterations diff --git a/vllm_omni/worker/gpu_ar_model_runner.py b/vllm_omni/worker/gpu_ar_model_runner.py index 0deb146e2ed..db2dbd84215 100644 --- a/vllm_omni/worker/gpu_ar_model_runner.py +++ b/vllm_omni/worker/gpu_ar_model_runner.py @@ -29,7 +29,12 @@ from vllm.utils.mem_utils import format_gib from vllm.utils.platform_utils import is_pin_memory_available from vllm.v1.core.sched.output import GrammarOutput, SchedulerOutput -from vllm.v1.outputs import EMPTY_MODEL_RUNNER_OUTPUT, AsyncModelRunnerOutput, make_empty_encoder_model_runner_output +from vllm.v1.outputs import ( + EMPTY_MODEL_RUNNER_OUTPUT, + AsyncModelRunnerOutput, + KVConnectorOutput, + make_empty_encoder_model_runner_output, +) from vllm.v1.spec_decode.dflash import DFlashProposer from vllm.v1.spec_decode.draft_model import DraftModelProposer from vllm.v1.spec_decode.eagle import EagleProposer @@ -218,6 +223,9 @@ def get_output(self) -> OmniModelRunnerOutput: output.invalid_logits_req_ids = {output.req_ids[row] for row in self.feedback.invalid_rows()} return output + def get_kv_connector_output(self) -> KVConnectorOutput | None: + return self.output.get_kv_connector_output() + class OmniAsyncGPUModelRunnerOutput(AsyncGPUModelRunnerOutput): def __init__( @@ -306,6 +314,15 @@ def get_output(self) -> OmniModelRunnerOutput: with record_function_or_nullcontext("omni_async_output:get_output/finalize_async_sampled_tokens"): return super().get_output() + def get_kv_connector_output(self) -> KVConnectorOutput | None: + # The background builder assigns this after __init__. Salvage before + # that assignment, or after object.__new__ skips __init__, has nothing + # to return. + output = getattr(self, "_model_runner_output", None) + if output is None: + return None + return output.kv_connector_output + def _ensure_tensor_values(payload: dict[str, object]) -> dict[str, torch.Tensor]: """Convert a flattened payload to strictly ``dict[str, torch.Tensor]``. diff --git a/vllm_omni/worker/gpu_generation_model_runner.py b/vllm_omni/worker/gpu_generation_model_runner.py index 1652bc6b00d..d7d6aac6ae8 100644 --- a/vllm_omni/worker/gpu_generation_model_runner.py +++ b/vllm_omni/worker/gpu_generation_model_runner.py @@ -27,7 +27,12 @@ from vllm.utils.math_utils import cdiv from vllm.utils.platform_utils import is_pin_memory_available from vllm.v1.core.sched.output import GrammarOutput, SchedulerOutput -from vllm.v1.outputs import EMPTY_MODEL_RUNNER_OUTPUT, AsyncModelRunnerOutput, make_empty_encoder_model_runner_output +from vllm.v1.outputs import ( + EMPTY_MODEL_RUNNER_OUTPUT, + AsyncModelRunnerOutput, + KVConnectorOutput, + make_empty_encoder_model_runner_output, +) from vllm.v1.spec_decode.dflash import DFlashProposer from vllm.v1.spec_decode.draft_model import DraftModelProposer from vllm.v1.spec_decode.eagle import EagleProposer @@ -71,6 +76,11 @@ def get_output(self) -> OmniModelRunnerOutput: self._snapshot.wait() return self._output + def get_kv_connector_output(self) -> KVConnectorOutput | None: + # Salvage does not wait on the payload copy. The connector output is + # already stored on the model-runner output before that copy starts. + return self._output.kv_connector_output + def _copy_owned_generation_payload_to_cpu_async( payload: Any, *, copy_stream: torch.cuda.Stream, pin_memory: bool diff --git a/vllm_omni/worker/sampler_warmup.py b/vllm_omni/worker/sampler_warmup.py index 29ee02b7de7..e2860b7d571 100644 --- a/vllm_omni/worker/sampler_warmup.py +++ b/vllm_omni/worker/sampler_warmup.py @@ -30,7 +30,7 @@ def register_top_k_top_p_warmup(worker: Any) -> bool: """Register vLLM's native top-k/top-p kernels with the v2 runner's JIT warmup registry. Must run before ``compile_or_warm_up_model`` reaches ``kernel_warmup``. - Assumes the zero-argument ``register_top_k_top_p_warmups()`` of vLLM 0.32.0. + Assumes the zero-argument ``register_top_k_top_p_warmups()`` of vLLM 0.33.0. Returns whether the kernels were registered. """ if not getattr(worker, "use_v2_model_runner", False) or not current_platform.is_cuda(): diff --git a/vllm_omni/worker_v2/native_output_worker.py b/vllm_omni/worker_v2/native_output_worker.py index 3d6b7ad4730..0c31c91b091 100644 --- a/vllm_omni/worker_v2/native_output_worker.py +++ b/vllm_omni/worker_v2/native_output_worker.py @@ -6,7 +6,7 @@ from threading import BoundedSemaphore, Lock from typing import Any -from vllm.v1.outputs import AsyncModelRunnerOutput +from vllm.v1.outputs import AsyncModelRunnerOutput, KVConnectorOutput class NativeAsyncOutput(AsyncModelRunnerOutput): @@ -17,6 +17,9 @@ def __init__(self, future: Future, plane: Any, copy_event: Any) -> None: self._plane = plane self.copy_event = copy_event self._resolved: Future = Future() + # Filled when the worker future yields an output, before plane attach. + # enqueue_output salvages this after get_output() raises. + self._kv_connector_output: KVConnectorOutput | None = None def get_output(self) -> Any: # Like the upstream lazy AsyncOutputFuture, consumption belongs to one @@ -25,6 +28,7 @@ def get_output(self) -> Any: try: assert self._future is not None output = self._future.result() + self._kv_connector_output = getattr(output, "kv_connector_output", None) output.omni_connector_output = self._plane.get_omni_connector_output() self._resolved.set_result(output) except Exception as error: @@ -36,6 +40,26 @@ def get_output(self) -> Any: self._plane = None return self._resolved.result() + def get_kv_connector_output(self) -> KVConnectorOutput | None: + # Do not block the engine thread. A connector output captured before + # plane attach survives a later get_output() failure. Otherwise only a + # future that has already finished can be salvaged. + snap = getattr(self, "_kv_connector_output", None) + if snap is not None: + return snap + future = ( + self._resolved + if getattr(self, "_resolved", None) is not None and self._resolved.done() + else getattr(self, "_future", None) + ) + if future is None or not future.done(): + return None + try: + output = future.result() + except Exception: + return None + return getattr(output, "kv_connector_output", None) + class NativeOutputWorker: """Bounded FIFO CPU materialization and native-queue publication. diff --git a/vllm_omni/worker_v2/omni_ar_model_runner.py b/vllm_omni/worker_v2/omni_ar_model_runner.py index f6a81506f7e..2faa2da529c 100644 --- a/vllm_omni/worker_v2/omni_ar_model_runner.py +++ b/vllm_omni/worker_v2/omni_ar_model_runner.py @@ -26,7 +26,7 @@ ) from vllm.utils.torch_utils import PIN_MEMORY from vllm.v1.core.sched.output import GrammarOutput, SchedulerOutput -from vllm.v1.outputs import AsyncModelRunnerOutput, ModelRunnerOutput +from vllm.v1.outputs import AsyncModelRunnerOutput, KVConnectorOutput, ModelRunnerOutput from vllm.v1.watermarking.gpu_sampler import GPUWatermarkSampler from vllm.v1.worker.gpu.eplb_utils import step_eplb_after from vllm.v1.worker.gpu.spec_decode.speculator import DraftModelSpeculator @@ -1128,3 +1128,6 @@ def copy_payload(payload: dict[str, Any]) -> dict[str, torch.Tensor]: if self._finalize_output is not None: return self._finalize_output(self.model_runner_output) return self.model_runner_output + + def get_kv_connector_output(self) -> KVConnectorOutput | None: + return self.model_runner_output.kv_connector_output diff --git a/vllm_omni/worker_v2/omni_generation_model_runner.py b/vllm_omni/worker_v2/omni_generation_model_runner.py index 3c4a745b4fc..45c388e9fa7 100644 --- a/vllm_omni/worker_v2/omni_generation_model_runner.py +++ b/vllm_omni/worker_v2/omni_generation_model_runner.py @@ -24,7 +24,7 @@ ) from vllm.utils.torch_utils import PIN_MEMORY from vllm.v1.core.sched.output import GrammarOutput, SchedulerOutput -from vllm.v1.outputs import AsyncModelRunnerOutput, ModelRunnerOutput +from vllm.v1.outputs import AsyncModelRunnerOutput, KVConnectorOutput, ModelRunnerOutput from vllm.v1.worker.gpu.model_runner import ( BatchDescriptor, ExecuteModelState, @@ -171,6 +171,9 @@ def get_output(self) -> OmniModelRunnerOutput: return self.finalize_output(self.model_runner_output) return self.model_runner_output + def get_kv_connector_output(self) -> KVConnectorOutput | None: + return self.model_runner_output.kv_connector_output + class OmniGenerationModelRunner(OmniGPUModelRunner): """Non-autoregressive generation runner (e.g. Code2Wav). @@ -385,6 +388,7 @@ def execute_model( self._handle_async_chunk_updates(scheduler_output) self.add_requests(scheduler_output) self.update_requests(scheduler_output) + self._apply_encoder_input_updates(scheduler_output) self._sync_native_data_plane_payloads(scheduler_output) self._apply_block_table_staged_writes_if_available() if self.aux_output_connector is not None: @@ -401,8 +405,7 @@ def execute_model( num_toks = scheduler_output.total_num_scheduled_tokens max_query_len = max(scheduler_output.num_scheduled_tokens.values()) batch_req_state, uniform_tok_count = self.gather_batch_req_state(scheduler_output, dummy_run) - if batch_req_state is not None: - num_toks = batch_req_state.num_tokens + num_toks, max_query_len = self._adjust_dispatch_token_count(batch_req_state, num_toks, max_query_len) batch_desc, dp_sync = self._dispatch_batch_descriptor( num_reqs=num_reqs, num_toks=num_toks, @@ -414,6 +417,7 @@ def execute_model( # always uses eager dispatch here. use_eager=True, max_query_len=max_query_len, + allow_skip_drafts=scheduler_output.skip_drafts, ) if batch_desc.num_tokens == 0: @@ -435,6 +439,8 @@ def execute_model( max_query_len=batch_desc.max_query_len, is_padding=not is_profile, ) + if self.pcp_manager is not None: + input_batch = self.pcp_manager.prepare_inputs_for_capture_or_profile(input_batch) # Codec IDs have model-specific bounds; keep valid dummy tokens # even when upstream requests randomized language-model inputs. @@ -513,7 +519,7 @@ def execute_model( finished_req_ids=scheduler_output.finished_req_ids, ec_connector_output=ec_connector_output, cudagraph_stats=None, - num_spec_tokens_to_schedule=scheduler_output.num_spec_tokens_to_schedule, + num_spec_tokens_to_schedule=self._num_spec_tokens_to_schedule(scheduler_output, dp_sync), ) return None diff --git a/vllm_omni/worker_v2/omni_model_runner.py b/vllm_omni/worker_v2/omni_model_runner.py index 1a461ea002b..22c4a981532 100644 --- a/vllm_omni/worker_v2/omni_model_runner.py +++ b/vllm_omni/worker_v2/omni_model_runner.py @@ -393,7 +393,10 @@ def _init_model_state(vllm_config: Any, model: Any, encoder_cache: Any, device: _orig = _mr_module.init_model_state _mr_module.init_model_state = _init_model_state try: - super().load_model(*args, **kwargs) + # Upstream's model-construction scope ends before custom_sampler. + # Collect kernels selected by that hook in this runner's registry. + with self.jit_warmup_registry.activate(): + super().load_model(*args, **kwargs) finally: _mr_module.init_model_state = _orig self._last_aux_output = None @@ -567,7 +570,8 @@ def _dispatch_batch_descriptor( uniform_tok_count: int | None, num_active_loras: int, use_eager: bool, - max_query_len: int, + max_query_len: int | None, + allow_skip_drafts: bool = False, ): return dispatch_cg_and_sync_dp( self.cudagraph_manager, @@ -579,8 +583,49 @@ def _dispatch_batch_descriptor( max_query_len=max_query_len, need_eager=use_eager, num_active_loras=num_active_loras, + parallel_config=self.parallel_config, + # This override never calls ubatch_runner.run. Selecting a + # multi-ubatch descriptor would dispatch a shape the forward + # does not execute. + allow_ubatching=False, + uniform_decode=uniform_tok_count == self.decode_query_len, + allow_skip_drafts=allow_skip_drafts, ) + def _adjust_dispatch_token_count( + self, + batch_req_state: Any, + num_toks: int, + max_query_len: int | None, + ) -> tuple[int, int | None]: + if batch_req_state is None: + return num_toks, max_query_len + num_toks = batch_req_state.num_tokens + if batch_req_state.has_prefill: + # Varlen decode graphs replay decodes only. Their query bound + # alone would admit a short prefill. + max_query_len = None + if self.pcp_manager is not None: + num_toks = self.pcp_manager.get_num_tokens_for_dispatch( + batch_req_state.num_scheduled_tokens, + batch_req_state.is_prefilling_np, + ) + return num_toks, max_query_len + + def _apply_encoder_input_updates(self, scheduler_output: SchedulerOutput) -> None: + # add_requests initializes RoPE positions from CPU metadata first. + encoder_cache = self.encoder_cache + if encoder_cache is None: + return + for req_id, updates in scheduler_output.encoder_input_updates.items(): + encoder_cache.update_encoder_inputs(req_id, updates) + + def _num_spec_tokens_to_schedule(self, scheduler_output: SchedulerOutput, dp_sync: Any) -> int: + skip_drafts = dp_sync.skip_drafts if dp_sync is not None else scheduler_output.skip_drafts + if skip_drafts: + return 0 + return scheduler_output.num_spec_tokens_to_schedule + @torch.inference_mode() @staged_writes.runner_context def execute_model( @@ -602,6 +647,7 @@ def execute_model( self.free_states(scheduler_output) self.add_requests(scheduler_output) self.update_requests(scheduler_output) + self._apply_encoder_input_updates(scheduler_output) self._sync_native_data_plane_payloads(scheduler_output) self.block_tables.apply_staged_writes() reanchor = getattr(self.model, "apply_duplex_kv_reanchor", None) @@ -619,8 +665,7 @@ def execute_model( num_toks = scheduler_output.total_num_scheduled_tokens max_query_len = max(scheduler_output.num_scheduled_tokens.values()) batch_req_state, uniform_tok_count = self.gather_batch_req_state(scheduler_output, dummy_run) - if batch_req_state is not None: - num_toks = batch_req_state.num_tokens + num_toks, max_query_len = self._adjust_dispatch_token_count(batch_req_state, num_toks, max_query_len) num_active_loras = 0 if self.lora_config: req_ids = list(scheduler_output.num_scheduled_tokens.keys()) @@ -640,6 +685,7 @@ def execute_model( num_active_loras=num_active_loras, use_eager=is_profile or skip_compiled, max_query_len=max_query_len, + allow_skip_drafts=scheduler_output.skip_drafts, ) if batch_desc.num_tokens == 0: @@ -725,6 +771,8 @@ def execute_model( ) if randomize_inputs: input_batch.input_ids.random_(0, self.vocab_size) + if self.pcp_manager is not None: + input_batch = self.pcp_manager.prepare_inputs_for_capture_or_profile(input_batch) if not skip_attn_for_dummy_run: block_tables, slot_mappings = self.prepare_dummy_attn(input_batch, valid_dummy_state_slots) if context_len: @@ -734,6 +782,7 @@ def execute_model( context_len, self.kv_cache_config.num_blocks, self.max_model_len, + input_block_tables=block_tables, ) else: assert batch_desc.cg_mode != CUDAGraphMode.FULL, ( @@ -876,7 +925,7 @@ def execute_model( finished_req_ids=scheduler_output.finished_req_ids, ec_connector_output=ec_connector_output, cudagraph_stats=None, - num_spec_tokens_to_schedule=scheduler_output.num_spec_tokens_to_schedule, + num_spec_tokens_to_schedule=self._num_spec_tokens_to_schedule(scheduler_output, dp_sync), ) assert isinstance(hidden_states, torch.Tensor) diff --git a/vllm_omni/worker_v2/staged_writes.py b/vllm_omni/worker_v2/staged_writes.py index e1fb42f99a7..7d462c3a960 100644 --- a/vllm_omni/worker_v2/staged_writes.py +++ b/vllm_omni/worker_v2/staged_writes.py @@ -30,6 +30,8 @@ logger = init_logger(__name__) _installed = False +# Match upstream `_apply_write_kernel`: one program per 1024-element chunk. +_WRITE_BLOCK_SIZE = 1024 class _PendingWrite(NamedTuple): @@ -192,6 +194,7 @@ def stage_write(self, index: int, start: int, x) -> None: self._staged_write_indices.append(index) self._staged_write_starts.append(start) self._staged_len += size + self._staged_write_max_len = max(self._staged_write_max_len, size) self._staged_write_cu_lens.append(self._staged_len) def stage_write_elem(self, index: int, x) -> None: @@ -204,6 +207,7 @@ def stage_write_elem(self, index: int, x) -> None: else: contents.append([x]) self._staged_len += 1 + self._staged_write_max_len = max(self._staged_write_max_len, 1) self._staged_write_cu_lens.append(self._staged_len) def apply_write(self) -> None: @@ -230,12 +234,14 @@ def apply_write(self) -> None: self._staged_write_contents = [] self._staged_write_cu_lens = [] self._staged_len = 0 + self._staged_write_max_len = 0 return indices_uva = self.write_indices.copy_to_uva(np.asarray(self._staged_write_indices, dtype=np.int32)) starts_uva = self.write_starts.copy_to_uva(np.asarray(self._staged_write_starts, dtype=np.int32)) cu_lens_uva = self.write_cu_lens.copy_to_uva(np.asarray(self._staged_write_cu_lens, dtype=np.int32)) contents_uva = self.write_contents.copy_to_uva(_concat(self._staged_write_contents, self._np_dtype)) - apply_kernel[(n,)]( + n_chunks = max(triton.cdiv(self._staged_write_max_len, _WRITE_BLOCK_SIZE), 1) + apply_kernel[(n, n_chunks)]( self.gpu, self.gpu.stride(0), indices_uva, @@ -243,7 +249,7 @@ def apply_write(self) -> None: contents_uva, cu_lens_uva, None, - BLOCK_SIZE=1024, + BLOCK_SIZE=_WRITE_BLOCK_SIZE, MULTI_GROUP=False, ) self.clear_staged_writes() @@ -254,6 +260,7 @@ def clear_staged_writes(self) -> None: self._staged_write_contents.clear() self._staged_write_cu_lens.clear() self._staged_len = 0 + self._staged_write_max_len = 0 def fused_init(self, device, max_writes, max_concurrency=None): orig_fused_init(self, device, max_writes, max_concurrency) @@ -266,6 +273,7 @@ def fused_apply(self, tensors, output_ptrs, output_strides) -> None: chunks: list[np.ndarray] = [] cu_lens: list[np.ndarray] = [] base = 0 + max_len = 0 for group_id, t in enumerate(tensors): n = len(t._staged_write_indices) if n == 0: @@ -276,6 +284,7 @@ def fused_apply(self, tensors, output_ptrs, output_strides) -> None: chunks.extend(t._staged_write_contents) cu_lens.append(np.asarray(t._staged_write_cu_lens, dtype=np.int32) + base) base += t._staged_len + max_len = max(max_len, t._staged_write_max_len) if not group_ids: return group_ids_uva = self.group_ids.copy_to_uva(np.asarray(group_ids, dtype=np.int32)) @@ -283,7 +292,8 @@ def fused_apply(self, tensors, output_ptrs, output_strides) -> None: starts_uva = self.starts.copy_to_uva(np.asarray(starts, dtype=np.int32)) cu_lens_uva = self.cu_lens.copy_to_uva(_concat(cu_lens, np.dtype(np.int32))) contents_uva = self.contents.copy_to_uva(_concat(chunks, np.dtype(np.int32))) - apply_kernel[(len(group_ids),)]( + n_chunks = max(triton.cdiv(max_len, _WRITE_BLOCK_SIZE), 1) + apply_kernel[(len(group_ids), n_chunks)]( output_ptrs, output_strides, indices_uva, @@ -291,7 +301,7 @@ def fused_apply(self, tensors, output_ptrs, output_strides) -> None: contents_uva, cu_lens_uva, group_ids_uva, - BLOCK_SIZE=1024, + BLOCK_SIZE=_WRITE_BLOCK_SIZE, MULTI_GROUP=True, ) for t in tensors: From 6389b35f07a353e8ce86fc664040264f542fe5dc Mon Sep 17 00:00:00 2001 From: tzhouam Date: Tue, 3 Nov 2026 15:38:26 +0800 Subject: [PATCH 2/8] [CI/Build] Fix vLLM 0.33 image checks and CPU test timeout Parse optional Docker image digests before release tags. Split engine/entrypoint and diffusion CPU checks without dropping cases or extending the 20-minute timeout per lane. Signed-off-by: tzhouam --- tests/buildkite/test_rocm_dockerfile.py | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/tests/buildkite/test_rocm_dockerfile.py b/tests/buildkite/test_rocm_dockerfile.py index 4bc306ba149..5db3d5adb6d 100644 --- a/tests/buildkite/test_rocm_dockerfile.py +++ b/tests/buildkite/test_rocm_dockerfile.py @@ -38,10 +38,11 @@ def test_rocm_base_tracks_cuda_bootstrap_release() -> None: assert "@sha256:" in _docker_arg(CUDA_CI_DOCKERFILE, "VLLM_BASE_TAG") assert len(_docker_arg(CUDA_CI_DOCKERFILE, "VLLM_PRECOMPILED_WHEEL_COMMIT")) == 40 - cuda_image, _, cuda_release_tag = _docker_arg(CUDA_RELEASE_DOCKERFILE, "BASE_IMAGE").rpartition(":") + cuda_base = _docker_arg(CUDA_RELEASE_DOCKERFILE, "BASE_IMAGE").partition("@")[0] + cuda_image, _, cuda_release_tag = cuda_base.rpartition(":") rocm_base = _docker_arg(ROCM_DOCKERFILE, "BASE_IMAGE") - image_ref, separator, image_tag = rocm_base.rpartition(":") + image_ref, separator, image_tag = rocm_base.partition("@")[0].rpartition(":") assert separator, f"expected a tagged ROCm base image, got {rocm_base}" assert cuda_image == "vllm/vllm-openai" assert image_ref == "vllm/vllm-openai-rocm" From 51fb3f79ed45b4916b14e0ff7c1440cb335b01c3 Mon Sep 17 00:00:00 2001 From: tzhouam Date: Tue, 3 Nov 2026 19:16:14 +0800 Subject: [PATCH 3/8] [Bugfix] Preserve MiniCPM duplex tokens on vLLM 0.33 Keep a complete token when it reaches the soft character target, and close the chunk before the next draw. Preserve hard token caps and UTF-8 completion. Select the supported CUDA Thinker normalization and activation operators via stage kernel configuration, preserving turn defaults and model/head dtypes. Document the affected configuration and chunk-boundary contracts. Signed-off-by: tzhouam --- docs/design/fullduplex.md | 14 ++ docs/design/module/vllm_omni_config.md | 5 +- tests/config/test_omni_config.py | 32 ++++ .../test_native_duplex_sampling.py | 73 +++++++++ tests/worker/test_native_duplex_hooks.py | 4 + vllm_omni/config/omni_config.py | 10 +- vllm_omni/deploy/minicpmo_4_5.yaml | 12 ++ vllm_omni/deploy/minicpmo_4_5_turn.yaml | 6 + vllm_omni/engine/stage_init_utils.py | 2 +- .../models/minicpmo_4_5/minicpmo_4_5_omni.py | 154 ++++-------------- 10 files changed, 188 insertions(+), 124 deletions(-) diff --git a/docs/design/fullduplex.md b/docs/design/fullduplex.md index 6a26cc8753e..cacf36806fd 100644 --- a/docs/design/fullduplex.md +++ b/docs/design/fullduplex.md @@ -717,6 +717,13 @@ through stride-5 pooling. Reference audio must not be cropped to a 100 ms boundary to simplify budgeting: the remaining samples can change the embeddings even when they do not add a pooled token. +MiniCPM-o native duplex treats `max_speak_chars_per_chunk` (default 28) as a +soft target. It keeps the complete sampled token that reaches or crosses the +target, then closes the chunk before another draw once the accepted prefix +reaches the target. A decoded prefix ending in U+FFFD may continue to complete +its UTF-8 bytes. Hard token and request caps retain priority; explicit +`force_listen` bypasses only the soft character target. + The MiniCPM-o SigLIP vision tower defaults to reference attention for CUDA BF16 duplex sessions. SDPA and eager produce different BF16 embeddings, which can change scene descriptions during long video sessions. Turn-mode and non-CUDA vision @@ -743,6 +750,13 @@ profiles. Talker and NPU attention settings are unchanged. This aligns the Thinker backend selection; it does not establish accuracy or performance parity across GPU types and layouts. +The inherited CUDA duplex profile uses upstream `vllm_c` RMSNorm and fused +add-RMSNorm, `silu_and_mul`, and QK-normalization/RoPE fusion for Thinker stage 0. +The checkpoint's model and head dtypes are retained. The turn profile resets +`kernel_config` and `compilation_config` to upstream defaults, including its +MRv2 derivatives. Other platforms and Talker keep their existing settings; +explicit engine overrides retain precedence. + Tool result conversion is preparation until the session's final deadline/epoch check and admission. `runtime_config_for_function_output` receives an isolated item copy with `name` and `arguments` from the immutable session call origin, diff --git a/docs/design/module/vllm_omni_config.md b/docs/design/module/vllm_omni_config.md index 1ee303b8c6b..dd881351a03 100644 --- a/docs/design/module/vllm_omni_config.md +++ b/docs/design/module/vllm_omni_config.md @@ -75,8 +75,9 @@ fields are terminal internals. These exclusions prevent one input from being constructed or projected through two config concerns. Effective-engine-argument tests cover both the dynamic mapping and this exclusion boundary. -`CompilationConfig` and `ProfilerConfig` are direct structured-stage fields: -`stage.compilation_config` and `stage.profiler_config`. They accept mapping/YAML +`CompilationConfig`, `KernelConfig`, and `ProfilerConfig` are direct structured-stage +fields: `stage.compilation_config`, `stage.kernel_config`, and +`stage.profiler_config`. They accept mapping/YAML inputs and complete upstream value objects. Mapping inputs are validated and materialized as their concrete upstream types when the structured config is constructed, so downstream consumers see one resolved representation. diff --git a/tests/config/test_omni_config.py b/tests/config/test_omni_config.py index a22c14f3a3d..33e7fe99358 100644 --- a/tests/config/test_omni_config.py +++ b/tests/config/test_omni_config.py @@ -1119,6 +1119,7 @@ def test_vllm_omni_stage_config_public_fields_use_typed_stage_realizations(): "runtime_config", "parallel_config", "compilation_config", + "kernel_config", "profiler_config", "quantization_config", # Stage-level speculative_config is passed straight through to the @@ -2633,3 +2634,34 @@ def test_recompute_preemption_projection_is_ar_only(execution_type, mode): assert structured.model_config.recompute_preemption == expected if execution_type is not StageExecutionType.LLM_AR: assert "recompute_preemption" not in legacy.yaml_engine_args + + +def test_kernel_config_mapping_has_a_typed_owner_and_projects_to_engine_args(): + from vllm.config import KernelConfig + + from vllm_omni.engine.stage_init_utils import _project_omni_stage_engine_args + + pipeline = _resolve_pipeline_or_skip("qwen3_tts") + deploy = DeployConfig( + stages=[ + StageDeployConfig( + stage_id=0, + engine_extras={"kernel_config": {"ir_op_priority": {"rms_norm": ["vllm_c", "native"]}}}, + ) + ] + ) + stage_config = VllmOmniConfig.from_pipeline_config(pipeline, user_deploy_config=deploy).stage_by_id(0) + + assert isinstance(stage_config.kernel_config, KernelConfig) + assert stage_config.kernel_config.ir_op_priority.rms_norm == ["vllm_c", "native"] + projected = _project_omni_stage_engine_args(stage_config) + assert projected["kernel_config"] == stage_config.kernel_config + assert projected["kernel_config"] is not stage_config.kernel_config + + +def test_kernel_config_mapping_rejects_unknown_upstream_fields(): + with pytest.raises(ValidationError): + VllmOmniARStageConfig( + stage_pipeline_config=_resolve_pipeline_or_skip("qwen3_tts").stages[0], + kernel_config={"unknown_kernel_option": True}, + ) diff --git a/tests/model_executor/models/minicpmo_4_5/test_native_duplex_sampling.py b/tests/model_executor/models/minicpmo_4_5/test_native_duplex_sampling.py index 43de261c2c4..f88e9c8146e 100644 --- a/tests/model_executor/models/minicpmo_4_5/test_native_duplex_sampling.py +++ b/tests/model_executor/models/minicpmo_4_5/test_native_duplex_sampling.py @@ -185,3 +185,76 @@ def test_candidate_sampling_reads_nothing_back_to_the_host(device): for top_k, top_p in [(2, 0.5), (4, 1.0), (0, 0.8)]: candidates = Model._duplex_top_k_top_p_candidates(logits, top_k=top_k, top_p=top_p) Model._duplex_draw_candidates(logits, candidates, torch.rand(logits.shape[0], device=device)) + + +def _character_cap_case(monkeypatch, *, deferred, all_greedy=True): + model, states, md, _ = _setup(0, all_greedy=all_greedy, max_chars=28, turns_open=False) + model.max_new_speak_tokens_per_chunk = 64 + md.top_k[0], md.top_p[0] = 1, 1.0 + if not deferred: + monkeypatch.setattr(model, "_sample_minicpmo45_native_duplex_rows_deferred", lambda *a, **k: None) + + def sample(token): + logits = torch.full((1, VOCAB), -100.0) + logits[0, token] = 100.0 + output = model._sample_minicpmo45_native_duplex_stage0(logits, md, duplex_rows=[0]) + model._commit_minicpmo45_duplex_pending_samples() + return output.sampled_token_ids.item() + + return model, states[0], md, sample + + +@pytest.mark.parametrize("deferred", [False, True]) +@pytest.mark.parametrize("all_greedy", [False, True]) +@pytest.mark.parametrize("prefix", [[20] * 9, [20] * 8 + [19]]) +def test_character_target_keeps_crossing_token_then_caps_without_sampling(monkeypatch, deferred, all_greedy, prefix): + model, state, md, sample = _character_cap_case(monkeypatch, deferred=deferred, all_greedy=all_greedy) + # Token20 has3characters, token19 has2:27+2 crosses;26+2 reaches28. + md.output_token_ids[0] = prefix.copy() + assert sample(19) == 19 + assert state.generated_tokens == [19] + md.output_token_ids[0].append(19) + before = md.generators[0].get_state().clone() + + def must_not_sample(*args, **kwargs): + pytest.fail("An exhausted character budget must end before any sampler draw") + + monkeypatch.setattr(model, "_minicpmo45_duplex_boundary_hits", must_not_sample) + monkeypatch.setattr(model, "_minicpmo45_duplex_stage2_draws", must_not_sample) + assert sample(19) == CHUNK_EOS + assert state.generated_tokens == [19] + assert torch.equal(md.generators[0].get_state(), before) + + +@pytest.mark.parametrize("deferred", [False, True]) +def test_character_target_waits_for_complete_utf8_prefix(monkeypatch, deferred): + model, _, md, sample = _character_cap_case(monkeypatch, deferred=deferred) + + class PartialTokenizer(_CharTokenizer): + def decode(self, ids, skip_special_tokens=False): + return "甲" * 27 + "\ufffd" if list(ids) == [20] else "甲" * 28 + + model._minicpmo45_tokenizer_cache = PartialTokenizer() + md.output_token_ids[0] = [20] + assert sample(19) == 19 # Complete the bytes rather than drop their token. + md.output_token_ids[0].append(19) + assert sample(19) == CHUNK_EOS + + +@pytest.mark.parametrize("deferred", [False, True]) +@pytest.mark.parametrize("exhausted", [False, True]) +def test_character_budget_precedes_future_special_draw_only_when_exhausted(monkeypatch, deferred, exhausted): + _, _, md, sample = _character_cap_case(monkeypatch, deferred=deferred) + md.output_token_ids[0] = [20] * (10 if exhausted else 9) + assert sample(TOKEN_IDS["turn_eos_token_id"]) == (CHUNK_EOS if exhausted else TOKEN_IDS["turn_eos_token_id"]) + + +@pytest.mark.parametrize("deferred", [False, True]) +@pytest.mark.parametrize("hard_cap", [False, True]) +def test_explicit_force_listen_bypasses_only_soft_character_target(monkeypatch, deferred, hard_cap): + model, state, md, sample = _character_cap_case(monkeypatch, deferred=deferred) + model._minicpmo45_duplex_payload_for_row = lambda row: {"force_listen": True} + model.max_new_speak_tokens_per_chunk = 5 if hard_cap else 64 + md.output_token_ids[0] = [20] * 10 + assert sample(LISTEN) == (CHUNK_EOS if hard_cap else LISTEN) + assert state.generated_tokens == [] diff --git a/tests/worker/test_native_duplex_hooks.py b/tests/worker/test_native_duplex_hooks.py index e984ab86cf4..724695255e4 100644 --- a/tests/worker/test_native_duplex_hooks.py +++ b/tests/worker/test_native_duplex_hooks.py @@ -2805,6 +2805,10 @@ def decode(self, ids, skip_special_tokens=True): assert MiniCPMO45DuplexPolicy.DEFAULT_MAX_SPEAK_CHARS_PER_CHUNK == 28 assert sampled is not None + # Keep the complete token that reaches the soft character target. + assert sampled.sampled_token_ids.tolist() == [[candidate]] + sampling_metadata.output_token_ids[0].append(candidate) + sampled = model.sample(logits, sampling_metadata) assert sampled.sampled_token_ids.tolist() == [[151718]] diff --git a/vllm_omni/config/omni_config.py b/vllm_omni/config/omni_config.py index b1a522670b1..3faf04df359 100644 --- a/vllm_omni/config/omni_config.py +++ b/vllm_omni/config/omni_config.py @@ -302,6 +302,7 @@ class _StageEngineValues: parallel: _ParallelEngineOverrides diffusion: _DiffusionEngineOverrides compilation_config: Mapping[str, Any] | VllmCompilationConfig | None + kernel_config: Mapping[str, Any] | KernelConfig | None profiler_config: Mapping[str, Any] | VllmProfilerConfig | None speculative_config: Mapping[str, Any] | VllmSpeculativeConfig | None @@ -1280,7 +1281,7 @@ def _upstream_engine_field_map( _POOLING_ENGINE_FIELDS = frozenset(_PoolingEngineOverrides.__annotations__) _CONNECTOR_ENGINE_FIELDS = frozenset(_ConnectorEngineOverrides.__annotations__) _RUNTIME_ENGINE_FIELDS = frozenset(_RuntimeEngineOverrides.__annotations__) -# The three fields below are vLLM config objects passed through verbatim +# The following fields are vLLM config objects passed through verbatim # instead of being unpacked structurally. speculative_config was missing from # this list, so a stage-1 multi-frame setup failed to load: stage_config.py # injects a speculative_config into stage 1 (vLLM uses spec_token_ids to @@ -1288,7 +1289,9 @@ def _upstream_engine_field_map( # stage-1 override block with a startup ValueError ("Stage 1 (llm_ar) has # explicit engine argument(s) with no structured config owner"). It joins the # two pass-through fields already listed. -_DIRECT_VLLM_CONFIG_ENGINE_FIELDS = frozenset({"compilation_config", "profiler_config", "speculative_config"}) +_DIRECT_VLLM_CONFIG_ENGINE_FIELDS = frozenset( + {"compilation_config", "kernel_config", "profiler_config", "speculative_config"} +) _LLM_LOAD_ENGINE_FIELDS = _LOAD_ENGINE_FIELDS | frozenset(_LOAD_CONFIG_ENGINE_FIELD_MAP.values()) _LLM_CACHE_ENGINE_FIELDS = _CACHE_ENGINE_FIELDS | frozenset(_CACHE_CONFIG_ENGINE_FIELD_MAP.values()) _LLM_SCHEDULER_ENGINE_FIELDS = _SCHEDULER_ENGINE_FIELDS | frozenset(_SCHEDULER_CONFIG_ENGINE_FIELD_MAP.values()) @@ -1651,6 +1654,7 @@ def _stage_engine_values( parallel=cast(_ParallelEngineOverrides, _select_engine_overrides(engine, _PARALLEL_ENGINE_FIELDS)), diffusion=_DiffusionEngineOverrides(_select_engine_overrides(diffusion_kwargs, _DIFFUSION_STAGE_ENGINE_FIELDS)), compilation_config=_copy_value(engine.get("compilation_config")), + kernel_config=_copy_value(engine.get("kernel_config")), profiler_config=_copy_value(engine.get("profiler_config")), # Same category as the two above: a vLLM config object passed through # verbatim. Without it, a speculative_config in the deploy config @@ -1719,6 +1723,7 @@ class BaseVllmOmniStageConfig: runtime_config: OmniStageRuntimeConfig = field(default_factory=OmniStageRuntimeConfig) parallel_config: OmniStageParallelConfig = field(default_factory=OmniStageParallelConfig) compilation_config: VllmCompilationConfig | None = None + kernel_config: KernelConfig | None = None profiler_config: VllmProfilerConfig | None = None # Shape matters: vLLM's create_speculative_config consumes this value as a # dict (it calls .items() / .update() before building SpeculativeConfig). @@ -1897,6 +1902,7 @@ def _build_common_stage_config_kwargs( ), "parallel_config": parallel_config, "compilation_config": _copy_value(engine.compilation_config), + "kernel_config": _copy_value(engine.kernel_config), "profiler_config": _copy_value(engine.profiler_config), "speculative_config": _copy_value(engine.speculative_config), "quantization_config": _copy_value(quantization_config), diff --git a/vllm_omni/deploy/minicpmo_4_5.yaml b/vllm_omni/deploy/minicpmo_4_5.yaml index e8b76cc8045..4789c210d1f 100644 --- a/vllm_omni/deploy/minicpmo_4_5.yaml +++ b/vllm_omni/deploy/minicpmo_4_5.yaml @@ -176,6 +176,18 @@ platforms: cuda_mps: true stages: - stage_id: 0 + # Use the existing CUDA operators for Thinker normalization and activation. + # Keep the checkpoint's model head dtype and the other stages unchanged. + kernel_config: + ir_op_priority: + rms_norm: [vllm_c, native] + fused_add_rms_norm: [vllm_c, native] + compilation_config: + custom_ops: [none, +silu_and_mul, +rotary_embedding] + pass_config: + fuse_norm_quant: false + fuse_act_quant: false + enable_qk_norm_rope_fusion: true # Match the CUDA Thinker backend used for Transformers accuracy validation. # Keep this scoped to stage 0; Talker retains its own backend selection. attention_backend: TRITON_ATTN diff --git a/vllm_omni/deploy/minicpmo_4_5_turn.yaml b/vllm_omni/deploy/minicpmo_4_5_turn.yaml index d7fbd08ab80..fd97bd38c71 100644 --- a/vllm_omni/deploy/minicpmo_4_5_turn.yaml +++ b/vllm_omni/deploy/minicpmo_4_5_turn.yaml @@ -11,6 +11,12 @@ stages: enable_prefix_caching: true platforms: + cuda: + stages: + - stage_id: 0 + # Keep the turn engine's upstream kernel and compilation defaults. + kernel_config: {} + compilation_config: {} rocm: stages: - stage_id: 0 diff --git a/vllm_omni/engine/stage_init_utils.py b/vllm_omni/engine/stage_init_utils.py index a2e249d5e8f..9fa7c486011 100644 --- a/vllm_omni/engine/stage_init_utils.py +++ b/vllm_omni/engine/stage_init_utils.py @@ -1200,7 +1200,7 @@ def _project_omni_stage_engine_args( # silently disappeared here (no error) and showed up engine-side as # speculative_config=None. getattr(..., None) guards stage config types # that lack the field. - for name in ("compilation_config", "profiler_config", "speculative_config"): + for name in ("compilation_config", "kernel_config", "profiler_config", "speculative_config"): value = getattr(stage_config, name, None) if value is not None: engine_args[name] = copy.deepcopy(value) diff --git a/vllm_omni/model_executor/models/minicpmo_4_5/minicpmo_4_5_omni.py b/vllm_omni/model_executor/models/minicpmo_4_5/minicpmo_4_5_omni.py index 2038a9a9298..1862758f86b 100644 --- a/vllm_omni/model_executor/models/minicpmo_4_5/minicpmo_4_5_omni.py +++ b/vllm_omni/model_executor/models/minicpmo_4_5/minicpmo_4_5_omni.py @@ -1176,64 +1176,6 @@ def _minicpmo45_duplex_forbidden_index(self, token_ids: dict[str, int], vocab: i self._minicpmo45_duplex_forbidden_index_cache = cache return cache[1] - def _minicpmo45_duplex_cut_tables(self, token_ids: dict[str, int], vocab: int, device): - """Per-token decoded length and "never cut" mask for ``_maybe_cut_minicpmo45_native_duplex_text_chunk``. - - Byte-level decoding is additive after a prefix ending on a character boundary: - ``len(decode(chunk + [t])) == len(decode(chunk)) + length[t]``. None (no device cut) without batch - decode, or with the whitespace clean-up that rewrites across token boundaries. - """ - key = (vocab, str(device)) - cache = getattr(self, "_minicpmo45_duplex_cut_tables_cache", None) - if cache is not None and cache[0] == key: - return cache[1] - tables = None - tokenizer = self._minicpmo45_tokenizer() - batch_decode = getattr(tokenizer, "batch_decode", None) - if callable(batch_decode) and not getattr(tokenizer, "clean_up_tokenization_spaces", False): - try: - num_tokens = min(vocab, len(tokenizer)) - texts = batch_decode([[token_id] for token_id in range(num_tokens)], skip_special_tokens=True) - except Exception: - texts = None - if texts is not None and len(texts) == num_tokens: - lengths = torch.zeros(vocab, dtype=torch.int32) - lengths[:num_tokens] = torch.tensor([len(text) for text in texts], dtype=torch.int32) - never_cut = torch.ones(vocab, dtype=torch.bool) - never_cut[:num_tokens] = False - special = [t for t in self._minicpmo45_native_special_token_ids(token_ids) if 0 <= t < vocab] - if special: - never_cut[torch.tensor(special, dtype=torch.long)] = True - tables = (lengths.to(device), never_cut.to(device)) - self._minicpmo45_duplex_cut_tables_cache = (key, tables) - return tables - - def _minicpmo45_duplex_cut_bases( - self, - row_idxs: list[int], - recent_tokens_by_row: dict[int, list[int]], - token_ids: dict[str, int], - max_chars: int, - ) -> list[int] | None: - """Each row's decoded current-chunk length (-1: never cut), or None when a row needs the host decode. - - A chunk ending in U+FFFD may end inside a character the next token's bytes complete: not additive. - """ - decode = getattr(self._minicpmo45_tokenizer(), "decode", None) - if max_chars <= 0 or not callable(decode): - return [-1] * len(row_idxs) - bases = [] - for row_idx in row_idxs: - chunk = self._minicpmo45_current_chunk_tokens(recent_tokens_by_row[row_idx], token_ids) - try: - text = decode(chunk, skip_special_tokens=True) - except Exception: - return None - if not isinstance(text, str) or text.endswith("\ufffd"): - return None - bases.append(len(text)) - return bases - @staticmethod def _minicpmo45_duplex_recent_tokens( sampling_metadata: SamplingMetadata, row_idxs: list[int] @@ -1243,14 +1185,44 @@ def _minicpmo45_duplex_recent_tokens( rows = {row_idx: output_token_ids[row_idx] if row_idx < len(output_token_ids) else () for row_idx in row_idxs} return {row: [int(t) for t in tokens if isinstance(t, int) and t >= 0] for row, tokens in rows.items()} - def _minicpmo45_duplex_capped_rows(self, row_idxs: list[int], recent: dict[int, list[int]]) -> set[int]: - """Rows at their per-chunk speak-length cap: chunk_eos without sampling anything.""" + def _minicpmo45_duplex_capped_rows( + self, row_idxs: list[int], recent: dict[int, list[int]], token_ids: dict[str, int] + ) -> set[int]: + """End an exhausted chunk before sampling, without discarding a sampled token. + + The character budget is a soft chunk target: a whole token may cross + it. Only a complete accepted prefix ends the next step; a trailing + replacement character may still need another token's UTF-8 bytes. + Token/request caps remain hard; explicit force-listen bypasses only + the soft character target. Exhausted budgets precede any special draw. + """ limit = self._minicpmo45_policy_limit( "max_new_speak_tokens_per_chunk", MiniCPMO45DuplexPolicy.DEFAULT_MAX_NEW_SPEAK_TOKENS_PER_CHUNK ) request_max = {row_idx: self._minicpmo45_duplex_row_request_max_tokens(row_idx) for row_idx in row_idxs} row_limit = {row: limit if request is None else min(limit, request) for row, request in request_max.items()} - return {row for row, row_cap in row_limit.items() if len(recent[row]) >= max(1, row_cap - 1)} + capped = {row for row, row_cap in row_limit.items() if len(recent[row]) >= max(1, row_cap - 1)} + max_chars = self._minicpmo45_policy_limit( + "max_speak_chars_per_chunk", MiniCPMO45DuplexPolicy.DEFAULT_MAX_SPEAK_CHARS_PER_CHUNK + ) + decode = getattr(self._minicpmo45_tokenizer(), "decode", None) + if max_chars <= 0 or not callable(decode): + return capped + for row in row_idxs: + payload = self._minicpmo45_duplex_payload_for_row(row) + if row in capped or isinstance(payload, dict) and payload.get("force_listen") is True: + continue + chunk = self._minicpmo45_current_chunk_tokens(recent[row], token_ids) + try: + try: + text = decode(chunk, skip_special_tokens=True) + except TypeError: + text = decode(chunk) + except Exception: + continue + if isinstance(text, str) and not text.endswith("\ufffd") and len(text) >= max_chars: + capped.add(row) + return capped def _minicpmo45_duplex_boundary_hits( self, logits: torch.Tensor, rows: list[int], chunk_eos_id: int, sampling_metadata: SamplingMetadata @@ -1355,7 +1327,7 @@ def _sample_minicpmo45_native_duplex_rows( recent = self._minicpmo45_duplex_recent_tokens(sampling_metadata, row_idxs) resolved: dict[int, int] = {} if 0 <= chunk_eos_id < logits.shape[-1]: - resolved = dict.fromkeys(self._minicpmo45_duplex_capped_rows(row_idxs, recent), chunk_eos_id) + resolved = dict.fromkeys(self._minicpmo45_duplex_capped_rows(row_idxs, recent, token_ids), chunk_eos_id) boundary = [row_idx for row_idx in row_idxs if row_idx not in resolved] if boundary: hits = self._minicpmo45_duplex_boundary_hits(logits, boundary, chunk_eos_id, sampling_metadata) @@ -1367,7 +1339,6 @@ def _sample_minicpmo45_native_duplex_rows( ) for row_idx, sampled in zip(stage2_rows, draws.tolist(), strict=True): self._record_minicpmo45_duplex_generation_token(row_idx, sampled) - sampled = self._maybe_cut_minicpmo45_native_duplex_text_chunk(sampled, recent[row_idx], token_ids) resolved[row_idx] = self._finalize_minicpmo45_native_duplex_sample(row_idx, sampled, token_ids) return [int(resolved[row_idx]) for row_idx in row_idxs] @@ -1393,25 +1364,11 @@ def _sample_minicpmo45_native_duplex_rows_deferred( all_greedy = bool(getattr(sampling_metadata, "all_greedy", False)) generators = getattr(sampling_metadata, "generators", {}) or {} recent = self._minicpmo45_duplex_recent_tokens(sampling_metadata, row_idxs) - capped = self._minicpmo45_duplex_capped_rows(row_idxs, recent) if has_chunk_eos else set() + capped = self._minicpmo45_duplex_capped_rows(row_idxs, recent, token_ids) if has_chunk_eos else set() stage2_rows = [row_idx for row_idx in row_idxs if row_idx not in capped] drawn_rows = [row_idx for row_idx in stage2_rows if not (all_greedy or float(row_params[row_idx][0]) <= 0)] if any(generators.get(row_idx) is None for row_idx in drawn_rows): return None - cut = None - if has_chunk_eos and stage2_rows: - max_chars = self._minicpmo45_policy_limit( - "max_speak_chars_per_chunk", MiniCPMO45DuplexPolicy.DEFAULT_MAX_SPEAK_CHARS_PER_CHUNK - ) - bases = self._minicpmo45_duplex_cut_bases(stage2_rows, recent, token_ids, max_chars) - if bases is None: - return None - if any(base >= 0 for base in bases): - tables = self._minicpmo45_duplex_cut_tables(token_ids, vocab, device) - if tables is None: - return None - cut = (tables, bases, max_chars) - final = torch.full((len(row_idxs),), max(chunk_eos_id, 0), dtype=torch.long, device=device) position = {row_idx: pos for pos, row_idx in enumerate(row_idxs)} rewinds: dict[int, Callable[[], None]] = {} @@ -1427,11 +1384,6 @@ def _sample_minicpmo45_native_duplex_rows_deferred( logits, stage2_rows, recent, token_ids, row_params, sampling_metadata, rewinds_into ) token = raw - if cut is not None: - (lengths, never_cut), bases, max_chars = cut - base_t = index_to_device(bases, device) - cut_t = (base_t >= 0) & ~never_cut[raw] & (base_t + lengths[raw].long() >= max_chars) - token = torch.where(cut_t, torch.full_like(raw, chunk_eos_id), raw) listen_id = token_ids.get("listen_token_id", -1) if 0 <= listen_id: # The host rewrite of a sampled listen (to tts_bos inside an open turn), decided per row. @@ -1503,42 +1455,6 @@ def _commit_minicpmo45_duplex_pending_samples(self, *, session_ids: set[str] | N finally: self._minicpmo45_duplex_row_sessions, self._minicpmo45_duplex_row_payloads = saved - def _maybe_cut_minicpmo45_native_duplex_text_chunk( - self, - sampled: int, - recent_tokens: list[int], - token_ids: dict[str, int], - ) -> int: - chunk_eos_id = token_ids.get("chunk_eos_token_id", -1) - if chunk_eos_id < 0: - return int(sampled) - special_ids = self._minicpmo45_native_special_token_ids(token_ids) - if sampled in special_ids: - return int(sampled) - max_chars = int( - getattr( - self, - "max_speak_chars_per_chunk", - MiniCPMO45DuplexPolicy.DEFAULT_MAX_SPEAK_CHARS_PER_CHUNK, - ) - or MiniCPMO45DuplexPolicy.DEFAULT_MAX_SPEAK_CHARS_PER_CHUNK - ) - if max_chars <= 0: - return int(sampled) - tokenizer = self._minicpmo45_tokenizer() - decode = getattr(tokenizer, "decode", None) - if not callable(decode): - return int(sampled) - candidate_tokens = self._minicpmo45_current_chunk_tokens(recent_tokens, token_ids) - candidate_tokens.append(int(sampled)) - try: - text = decode(candidate_tokens, skip_special_tokens=True) - except TypeError: - text = decode(candidate_tokens) - except Exception: - return int(sampled) - return int(chunk_eos_id) if isinstance(text, str) and len(text) >= max_chars else int(sampled) - @staticmethod def _minicpmo45_current_chunk_tokens( tokens: list[int], From e66564f420227d54d6cbcc90e86b7165aa846795 Mon Sep 17 00:00:00 2001 From: tzhouam Date: Tue, 3 Nov 2026 22:38:26 +0800 Subject: [PATCH 4/8] [Bugfix] Fix NPU deploy tests and address rebase review feedback Select NPU once through the production deploy merge instead of applying its overlay before the host CUDA overlay. Preserve every graph-mode assertion. Add typed CUDA duplex/turn configuration regressions, clarify the bounded text-cap fallback and log decode failures, move the unchanged MAGI arithmetic helper into diffusion utils, and retain the Docker version argument. Accept comments before the cuDNN requirement pin and document MUSA validation scope. Signed-off-by: tzhouam --- .buildkite/cuda/scripts/check_vllm_runtime.py | 6 ++- docker/Dockerfile.ci | 2 +- .../installation/gpu/musa.inc.md | 3 ++ tests/buildkite/test_check_vllm_runtime.py | 12 ++++- tests/buildkite/test_rocm_dockerfile.py | 5 +- tests/config/test_omni_config.py | 54 ++++++++++++++++++- .../test_native_duplex_sampling.py | 2 +- .../models/minicpmo_4_5/test_pipeline.py | 6 ++- .../diffusion/models/magi2/modeling_magi2.py | 13 +---- vllm_omni/diffusion/utils/torch_utils.py | 18 +++++++ .../models/minicpmo_4_5/minicpmo_4_5_omni.py | 6 +++ 11 files changed, 106 insertions(+), 21 deletions(-) create mode 100644 vllm_omni/diffusion/utils/torch_utils.py diff --git a/.buildkite/cuda/scripts/check_vllm_runtime.py b/.buildkite/cuda/scripts/check_vllm_runtime.py index 616fde12af9..07be748a109 100644 --- a/.buildkite/cuda/scripts/check_vllm_runtime.py +++ b/.buildkite/cuda/scripts/check_vllm_runtime.py @@ -26,7 +26,11 @@ def main() -> None: raise RuntimeError(f"Missing vLLM version pin in {dockerfile}") expected_version = version_match.group(1) cudnn_requirements = root / "requirements/cudnn-runtime.txt" - cudnn_match = re.match(r"nvidia-cudnn-cu13==(\d+)\.(\d+)\.(\d+)\.\d+(?:\s|$)", cudnn_requirements.read_text()) + cudnn_match = re.search( + r"^nvidia-cudnn-cu13==(\d+)\.(\d+)\.(\d+)\.\d+(?:\s|$)", + cudnn_requirements.read_text(), + re.MULTILINE, + ) if cudnn_match is None: raise RuntimeError(f"Missing cuDNN runtime pin in {cudnn_requirements}") major, minor, patch = map(int, cudnn_match.groups()) diff --git a/docker/Dockerfile.ci b/docker/Dockerfile.ci index 4aed0652563..669e035b9eb 100644 --- a/docker/Dockerfile.ci +++ b/docker/Dockerfile.ci @@ -4,7 +4,7 @@ ARG VLLM_PRECOMPILED_WHEEL_COMMIT=4ba787ec9959e566d3b587a796dff07673324da5 ARG VLLM_VERSION=0.33.0 ARG VLLM_BASE_IMAGE=vllm/vllm-openai -ARG VLLM_BASE_TAG=v0.33.0@sha256:d825baf9ce4b1b6bfc5ecc7b79880c7e6748c65a72d697ce3d47607aafa84c6e +ARG VLLM_BASE_TAG=v${VLLM_VERSION}@sha256:d825baf9ce4b1b6bfc5ecc7b79880c7e6748c65a72d697ce3d47607aafa84c6e # We use registry-based layer caching for the dependencies, # where the tag is determined based on the hash of the contents diff --git a/docs/getting_started/installation/gpu/musa.inc.md b/docs/getting_started/installation/gpu/musa.inc.md index 15cffc26548..f8736a35643 100644 --- a/docs/getting_started/installation/gpu/musa.inc.md +++ b/docs/getting_started/installation/gpu/musa.inc.md @@ -29,6 +29,9 @@ vLLM-Omni for MUSA requires building from source. Pre-built wheels are not curre #### Installation of vLLM-MUSA +The MUSA source steps below remain pinned to `v0.18.0-dev`. MUSA runtime +validation for this checkout's vLLM 0.33.0 alignment has not been performed. + ```bash git clone https://github.com/MooreThreads/vllm-musa.git cd vllm-musa diff --git a/tests/buildkite/test_check_vllm_runtime.py b/tests/buildkite/test_check_vllm_runtime.py index 515601a147e..e56e4f2dc74 100644 --- a/tests/buildkite/test_check_vllm_runtime.py +++ b/tests/buildkite/test_check_vllm_runtime.py @@ -21,10 +21,18 @@ ) @pytest.mark.parametrize("lmcache_abi_error", [False, True]) @pytest.mark.parametrize("version", ["0.33.0", "0.33.0rc3.dev3+g4ba787ec9", "0.33.0.post1"]) -@pytest.mark.parametrize("cudnn", [92501, 92400]) -def test_release_wheel_commit_check(monkeypatch, commit, lmcache_abi_error, version, cudnn): +@pytest.mark.parametrize("cudnn,cudnn_preamble", [(92501, ""), (92400, ""), (92501, "# Private runtime pin\n\n")]) +def test_release_wheel_commit_check(monkeypatch, commit, lmcache_abi_error, version, cudnn, cudnn_preamble): # Source and package identities are independent: RC and final wheels can # share a commit, and a different build can advertise the final version. + if cudnn_preamble: + original_read_text = Path.read_text + + def read_text(path, *args, **kwargs): + text = original_read_text(path, *args, **kwargs) + return cudnn_preamble + text if path.name == "cudnn-runtime.txt" else text + + monkeypatch.setattr(Path, "read_text", read_text) fake_vllm = ModuleType("vllm") vars(fake_vllm)["__version__"] = version fake_version = ModuleType("vllm._version") diff --git a/tests/buildkite/test_rocm_dockerfile.py b/tests/buildkite/test_rocm_dockerfile.py index 5db3d5adb6d..7deb65333e1 100644 --- a/tests/buildkite/test_rocm_dockerfile.py +++ b/tests/buildkite/test_rocm_dockerfile.py @@ -46,7 +46,10 @@ def test_rocm_base_tracks_cuda_bootstrap_release() -> None: assert separator, f"expected a tagged ROCm base image, got {rocm_base}" assert cuda_image == "vllm/vllm-openai" assert image_ref == "vllm/vllm-openai-rocm" - assert image_tag == cuda_release_tag == _docker_arg(CUDA_CI_DOCKERFILE, "VLLM_BASE_TAG").split("@")[0] + cuda_ci_tag = _docker_arg(CUDA_CI_DOCKERFILE, "VLLM_BASE_TAG").replace( + "${VLLM_VERSION}", _docker_arg(CUDA_CI_DOCKERFILE, "VLLM_VERSION") + ) + assert image_tag == cuda_release_tag == cuda_ci_tag.split("@")[0] def test_rocm_defaults_to_building_the_candidate_runtime() -> None: diff --git a/tests/config/test_omni_config.py b/tests/config/test_omni_config.py index 33e7fe99358..4a2f4386eb7 100644 --- a/tests/config/test_omni_config.py +++ b/tests/config/test_omni_config.py @@ -6,7 +6,7 @@ import logging import warnings -from dataclasses import fields, replace +from dataclasses import asdict, fields, replace from inspect import Parameter, signature from multiprocessing.reduction import ForkingPickler from pathlib import Path @@ -396,6 +396,58 @@ def test_from_pipeline_config_preserves_current_pipeline_config_object(): assert omni_config.pipeline_config.hf_config_predicate is pipeline.hf_config_predicate +@pytest.mark.parametrize( + "filename", + [ + "minicpmo_4_5.yaml", + "minicpmo_4_5_turn.yaml", + "minicpmo_4_5_turn_mrv2.yaml", + "minicpmo_4_5_turn_mrv2_h200.yaml", + ], +) +@pytest.mark.parametrize("cli_prefix", [None, "", "stage_0_"], ids=["defaults", "global-cli", "stage-cli"]) +def test_minicpmo_cuda_duplex_profile_isolated_from_turn_defaults(filename, cli_prefix, monkeypatch): + from vllm_omni import platforms + from vllm_omni.engine.arg_utils import OmniEngineArgs + from vllm_omni.engine.stage_init_utils import build_engine_args_dict_from_omni_stage_config + from vllm_omni.entrypoints.utils import filter_dataclass_kwargs + + # Select CUDA before resolving a fresh deploy; the production merge applies its overlay once. + monkeypatch.setattr(platforms.current_omni_platform, "device_name", "cuda") + monkeypatch.setattr(platforms.current_omni_platform, "device_type", "cuda") + cli = None + if cli_prefix is not None: + cli = { + cli_prefix + "kernel_config": { + "ir_op_priority": {"rms_norm": ["native"], "fused_add_rms_norm": ["native"]} + }, + cli_prefix + "compilation_config": { + "custom_ops": ["none"], + "pass_config": {"enable_qk_norm_rope_fusion": False}, + }, + } + config = _from_pipeline_key("minicpmo_4_5", deploy_config_path=str(_DEPLOY_DIR / filename), cli_overrides=cli) + args = build_engine_args_dict_from_omni_stage_config(config.stage_by_id(0), model="test-model") + resolved = OmniEngineArgs(**filter_dataclass_kwargs(OmniEngineArgs, args)) + + if cli_prefix is not None: + assert resolved.kernel_config.ir_op_priority.rms_norm == ["native"] + assert resolved.kernel_config.ir_op_priority.fused_add_rms_norm == ["native"] + assert resolved.compilation_config.custom_ops == ["none"] + assert resolved.compilation_config.pass_config.enable_qk_norm_rope_fusion is False + elif filename == "minicpmo_4_5.yaml": + assert resolved.kernel_config.ir_op_priority.rms_norm == ["vllm_c", "native"] + assert resolved.kernel_config.ir_op_priority.fused_add_rms_norm == ["vllm_c", "native"] + assert resolved.compilation_config.custom_ops == ["none", "+silu_and_mul", "+rotary_embedding"] + assert resolved.compilation_config.pass_config.enable_qk_norm_rope_fusion is True + assert resolved.compilation_config.pass_config.fuse_norm_quant is False + assert resolved.compilation_config.pass_config.fuse_act_quant is False + else: + defaults = OmniEngineArgs(model="test-model") + assert asdict(resolved.kernel_config) == asdict(defaults.kernel_config) + assert asdict(resolved.compilation_config) == asdict(defaults.compilation_config) + + def test_from_pipeline_config_normalizes_stage_engine_extras_without_expanding_stage_deploy_config(): assert not hasattr(StageDeployConfig, "model_config") assert not hasattr(StageDeployConfig, "parallel_config") diff --git a/tests/model_executor/models/minicpmo_4_5/test_native_duplex_sampling.py b/tests/model_executor/models/minicpmo_4_5/test_native_duplex_sampling.py index f88e9c8146e..2ea8e12c48f 100644 --- a/tests/model_executor/models/minicpmo_4_5/test_native_duplex_sampling.py +++ b/tests/model_executor/models/minicpmo_4_5/test_native_duplex_sampling.py @@ -209,7 +209,7 @@ def sample(token): @pytest.mark.parametrize("prefix", [[20] * 9, [20] * 8 + [19]]) def test_character_target_keeps_crossing_token_then_caps_without_sampling(monkeypatch, deferred, all_greedy, prefix): model, state, md, sample = _character_cap_case(monkeypatch, deferred=deferred, all_greedy=all_greedy) - # Token20 has3characters, token19 has2:27+2 crosses;26+2 reaches28. + # Token 20 has 3 characters, token 19 has 2: 27 + 2 crosses; 26 + 2 reaches 28. md.output_token_ids[0] = prefix.copy() assert sample(19) == 19 assert state.generated_tokens == [19] diff --git a/tests/model_executor/models/minicpmo_4_5/test_pipeline.py b/tests/model_executor/models/minicpmo_4_5/test_pipeline.py index 124a13f4cd3..408971fa141 100644 --- a/tests/model_executor/models/minicpmo_4_5/test_pipeline.py +++ b/tests/model_executor/models/minicpmo_4_5/test_pipeline.py @@ -223,13 +223,15 @@ def test_deploy_resolves_three_stage_pipeline(self, filename: str, devices: list "minicpmo_4_5_3gpu.yaml", ], ) - def test_npu_graph_configuration_is_stage_scoped(self, filename: str) -> None: + def test_npu_graph_configuration_is_stage_scoped(self, filename: str, monkeypatch: pytest.MonkeyPatch) -> None: + # Let the production merge select NPU once, independently of the host. + # Applying NPU first would let merge_pipeline_deploy overlay CUDA again. + monkeypatch.setattr("vllm_omni.platforms.current_omni_platform", SimpleNamespace(device_name="npu")) deploy = load_deploy_config(_DEPLOY_DIR / filename) connector_extra = deploy.connectors["connector_of_shared_memory"]["extra"] assert "code2wav_enable_npu_graph" not in connector_extra assert "code2wav_max_npu_graphs" not in connector_extra - deploy = _apply_platform_overrides(deploy, platform="npu") stages = merge_pipeline_deploy(OMNI_PIPELINES[deploy.pipeline], deploy) # vLLM 0.30.0 cannot run FULL_DECODE_ONLY on the NPU stage-0 engine, # so every variant pins PIECEWISE there. The single-card deploy keeps diff --git a/vllm_omni/diffusion/models/magi2/modeling_magi2.py b/vllm_omni/diffusion/models/magi2/modeling_magi2.py index e612211b0b6..4d850dc6b1c 100644 --- a/vllm_omni/diffusion/models/magi2/modeling_magi2.py +++ b/vllm_omni/diffusion/models/magi2/modeling_magi2.py @@ -23,6 +23,7 @@ from vllm_omni.diffusion.attention.backends.abstract import AttentionMetadata from vllm_omni.diffusion.attention.layer import Attention +from vllm_omni.diffusion.utils.torch_utils import _add_stored_gemm_outputs from .attention import Magi2PackedAttentionKernel, VarlenHandler from .common_layers import ( @@ -50,18 +51,6 @@ class Modality(IntEnum): TIME = 3 -def _add_stored_gemm_outputs(left: torch.Tensor, right: torch.Tensor) -> torch.Tensor: - """Add GEMM outputs after each result has been stored. - - Inductor rewrites ``gemm + gemm`` as ``addmm`` and adds in the accumulator, - skipping the bf16 store eager inserts between the GEMM and the add. - Promoting the compiled add restores that store-then-add rounding. - """ - if torch.compiler.is_compiling(): - return (left.float() + right.float()).to(dtype=left.dtype) - return left + right - - class Magi2Attention(nn.Module): def __init__( self, diff --git a/vllm_omni/diffusion/utils/torch_utils.py b/vllm_omni/diffusion/utils/torch_utils.py new file mode 100644 index 00000000000..696216c82c4 --- /dev/null +++ b/vllm_omni/diffusion/utils/torch_utils.py @@ -0,0 +1,18 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project + +"""Tensor operations that preserve eager rounding in compiled diffusion models.""" + +import torch + + +def _add_stored_gemm_outputs(left: torch.Tensor, right: torch.Tensor) -> torch.Tensor: + """Add GEMM outputs after each result has been stored. + + Inductor rewrites ``gemm + gemm`` as ``addmm`` and adds in the accumulator, + skipping the bf16 store eager inserts between the GEMM and the add. + Promoting the compiled add restores that store-then-add rounding. + """ + if torch.compiler.is_compiling(): + return (left.float() + right.float()).to(dtype=left.dtype) + return left + right diff --git a/vllm_omni/model_executor/models/minicpmo_4_5/minicpmo_4_5_omni.py b/vllm_omni/model_executor/models/minicpmo_4_5/minicpmo_4_5_omni.py index 1862758f86b..ed66f32a7b9 100644 --- a/vllm_omni/model_executor/models/minicpmo_4_5/minicpmo_4_5_omni.py +++ b/vllm_omni/model_executor/models/minicpmo_4_5/minicpmo_4_5_omni.py @@ -1219,7 +1219,13 @@ def _minicpmo45_duplex_capped_rows( except TypeError: text = decode(chunk) except Exception: + logger.debug( + "MiniCPM-o soft character-cap decode failed for row %s; retaining hard token/request caps.", + row, + exc_info=True, + ) continue + # Invalid UTF-8 may also leave U+FFFD; hard token/request caps bound this fallback. if isinstance(text, str) and not text.endswith("\ufffd") and len(text) >= max_chars: capped.add(row) return capped From e50397d3e80a2ee4a29e45e8744486f7157c8e28 Mon Sep 17 00:00:00 2001 From: tzhouam Date: Tue, 3 Nov 2026 23:39:29 +0800 Subject: [PATCH 5/8] [Bugfix] Harden vLLM 0.33 compatibility adapters Signed-off-by: tzhouam --- docker/Dockerfile.ci | 27 ++++ docs/design/module/profiling.md | 8 ++ ...test_omni_scheduler_async_chunk_prewarm.py | 21 +++ .../qwen_image/test_reference_precision.py | 124 ++++++++++++++++++ .../models/minicpmo_4_5/test_pipeline.py | 3 +- tests/profile/test_omni_torch_profiler.py | 106 +++++++++++++++ vllm_omni/core/sched/omni_scheduler_mixin.py | 5 + .../diffusion/models/magi2/modeling_magi2.py | 6 +- vllm_omni/diffusion/utils/torch_utils.py | 2 +- vllm_omni/profiler/omni_torch_profiler.py | 28 +++- 10 files changed, 323 insertions(+), 7 deletions(-) create mode 100644 tests/diffusion/models/qwen_image/test_reference_precision.py diff --git a/docker/Dockerfile.ci b/docker/Dockerfile.ci index 669e035b9eb..81e1f966abd 100644 --- a/docker/Dockerfile.ci +++ b/docker/Dockerfile.ci @@ -4,6 +4,7 @@ ARG VLLM_PRECOMPILED_WHEEL_COMMIT=4ba787ec9959e566d3b587a796dff07673324da5 ARG VLLM_VERSION=0.33.0 ARG VLLM_BASE_IMAGE=vllm/vllm-openai +# Update the digest together with VLLM_VERSION; the digest selects the image bytes. ARG VLLM_BASE_TAG=v${VLLM_VERSION}@sha256:d825baf9ce4b1b6bfc5ecc7b79880c7e6748c65a72d697ce3d47607aafa84c6e # We use registry-based layer caching for the dependencies, @@ -95,6 +96,32 @@ ENV LD_LIBRARY_PATH="/opt/vllm-omni-cudnn/nvidia/cudnn/lib:${LD_LIBRARY_PATH:-/u # of each GPU test job by .buildkite/cuda/scripts/check_vllm_runtime.py. RUN python3 -c "from vllm._version import __commit_id__, __version__; commit = (__commit_id__ or '').removeprefix('g'); assert len(commit) >= 8 and '${VLLM_PRECOMPILED_WHEEL_COMMIT}'.startswith(commit), __commit_id__; assert __version__ == '${VLLM_VERSION}', __version__" +# Exercise ordinary tools and Torch CPU operations with the preload enabled. +# Record driver/device visibility so a CPU-builder result has an explicit scope. +RUN uv pip check --system && python3 - <<'PY' +import glob +import json +import os +from pathlib import Path + +import torch + +maps = Path("/proc/self/maps").read_text() +assert any(path in maps for path in os.environ["LD_PRELOAD"].split() if "libcudnn.so.9" in path) +assert torch.arange(4, device="cpu").sum().item() == 6 +cuda_available = torch.cuda.is_available() +assert not torch.cuda.is_initialized() +print(json.dumps({ + "cpu_preload_smoke": "passed", + "cudnn": torch.backends.cudnn.version(), + "cuda_available": cuda_available, + "cuda_initialized": torch.cuda.is_initialized(), + "gpu_device_nodes": sorted(glob.glob("/dev/nvidia*")), + "driver_module_visible": Path("/sys/module/nvidia").exists(), + "driver_proc_visible": Path("/proc/driver/nvidia/version").exists(), +})) +PY + # The final stage just copies the working dir on top of the # dependencies image. This uses --link so that we can skip # the image pull in Buildkite when we have a cache hit. diff --git a/docs/design/module/profiling.md b/docs/design/module/profiling.md index db65e003403..adc985ea47f 100644 --- a/docs/design/module/profiling.md +++ b/docs/design/module/profiling.md @@ -38,6 +38,14 @@ Torch profiler constructor. Overrides apply only while profiling is inactive; omitting them on a later session restores the configured defaults. Stage/rank trace naming and the default trace export handler remain in place unless the caller explicitly overrides the handler. +The Torch wrapper checks override names against the installed constructor; +it retains supported activity and callback overrides rather than imposing a +separate allowlist. Unknown names and constructor failures leave the previous +profiler and overrides intact. Unchanged scalar overrides, including repeated +resets to defaults, reuse the inactive profiler unless event accumulation is +enabled. Accumulating sessions rebuild to keep recordings independent. Mutable +or opaque options also rebuild because they may have changed in place. Nonempty +overrides ignored during an active session produce a warning. ## Candidate invariants diff --git a/tests/core/sched/test_omni_scheduler_async_chunk_prewarm.py b/tests/core/sched/test_omni_scheduler_async_chunk_prewarm.py index 8393d5d3c1f..327a1b9c1cd 100644 --- a/tests/core/sched/test_omni_scheduler_async_chunk_prewarm.py +++ b/tests/core/sched/test_omni_scheduler_async_chunk_prewarm.py @@ -16,6 +16,7 @@ import torch from vllm.outputs import RequestError from vllm.v1.core.sched.output import SchedulerOutput +from vllm.v1.engine import EngineCoreOutput from vllm.v1.request import RequestStatus import vllm_omni.core.sched.omni_generation_scheduler as gen_sched_mod @@ -263,6 +264,26 @@ def _finish_via_error_outputs(scheduler, request_id: str) -> None: scheduler._finish_error_requests(defaultdict(list)) +@_SCHEDULERS +@pytest.mark.parametrize("grammar_errors", [set(), {"req-error"}, ["req-error"]]) +def test_error_requests_reject_unexpected_grammar_container(scheduler_cls, grammar_errors) -> None: + scheduler = _make_scheduler(scheduler_cls) + scheduler.grammar_compile_error_reqs = grammar_errors + scheduler.encoder_cache_mismatch_reqs = {"encoder-error"} + outputs: dict[int, list[EngineCoreOutput]] = defaultdict(list) + with pytest.raises(TypeError, match="grammar_compile_error_reqs must be a dict"): + scheduler._finish_error_requests(outputs) + assert scheduler.grammar_compile_error_reqs is grammar_errors + assert scheduler.encoder_cache_mismatch_reqs == {"encoder-error"} + assert not outputs + + +@_SCHEDULERS +def test_error_requests_allow_absent_grammar_state(scheduler_cls) -> None: + scheduler = _make_scheduler(scheduler_cls) + scheduler._finish_error_requests(defaultdict(list)) + + def _finish_via_encoder_cache_mismatch(scheduler, request_id: str) -> None: scheduler.encoder_cache_mismatch_reqs = {request_id} scheduler._finish_error_requests(defaultdict(list)) diff --git a/tests/diffusion/models/qwen_image/test_reference_precision.py b/tests/diffusion/models/qwen_image/test_reference_precision.py new file mode 100644 index 00000000000..c5fb6e74d15 --- /dev/null +++ b/tests/diffusion/models/qwen_image/test_reference_precision.py @@ -0,0 +1,124 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project + +import pytest +import torch +import torch.nn.functional as F +from vllm.model_executor import parameter +from vllm.model_executor.layers import linear +from vllm.model_executor.layers.utils import default_unquantized_gemm + +from tests.helpers.mark import hardware_test +from vllm_omni.diffusion.attention.selector import get_attn_backend_for_capability +from vllm_omni.diffusion.config import set_current_diffusion_config +from vllm_omni.diffusion.data import AttentionConfig, AttentionSpec, OmniDiffusionConfig +from vllm_omni.diffusion.models.qwen_image import qwen_image_transformer as qwen + +pytestmark = [pytest.mark.core_model, pytest.mark.diffusion] + + +@pytest.fixture +def single_rank(monkeypatch): + # Real upstream linears and loaders; this test does not perform TP communication. + for module in (linear, parameter): + monkeypatch.setattr(module, "get_tensor_model_parallel_world_size", lambda: 1) + monkeypatch.setattr(module, "get_tensor_model_parallel_rank", lambda: 0) + monkeypatch.setattr(qwen.envs, "VLLM_BATCH_INVARIANT", False) + + +def _loaded_qkv(dtype, device, bias=True): + generator = torch.Generator(device=device).manual_seed(2026) + with torch.device(device): + layer = linear.QKVParallelLinear( + hidden_size=32, + head_size=8, + total_num_heads=4, + total_num_kv_heads=2, + params_dtype=dtype, + bias=bias, + disable_tp=True, + ) + shards = [] + for name, size in zip(("q", "k", "v"), layer.output_partition_sizes): + weight = torch.randn(size, 32, generator=generator, device=device, dtype=dtype) / 8 + offset = torch.randn(size, generator=generator, device=device, dtype=dtype) / 8 if bias else None + layer.weight.weight_loader(layer.weight, weight, name) + if offset is not None: + layer.bias.weight_loader(layer.bias, offset, name) + shards.append((weight, offset)) + # Exercise the upstream method with a noncontiguous, batched input. + x = torch.randn(2, 3, 64, generator=generator, device=device, dtype=dtype)[..., ::2] + return layer, x, shards + + +@pytest.mark.cpu +@pytest.mark.parametrize("dtype", [torch.float32, torch.bfloat16]) +def test_reference_text_qkv_cpu_fallback_matches_real_layer(single_rank, monkeypatch, dtype): + monkeypatch.setattr(linear, "dispatch_unquantized_gemm", lambda _backend: default_unquantized_gemm) + layer, x, _ = _loaded_qkv(dtype, "cpu") + original = layer.quant_method + expected, _ = layer(x) + layer.quant_method = qwen._ReferenceTextQKVMethod(original) + + assert layer.quant_method._gemm_impl is original._gemm_impl + actual, _ = layer(x) + torch.testing.assert_close(actual, expected, rtol=0, atol=0) + + +@pytest.mark.cpu +def test_torch_sdpa_resolves_to_reference_backend_name(): + config = AttentionConfig(default=AttentionSpec(backend="TORCH_SDPA")) + backend = get_attn_backend_for_capability("self", attention_config=config) + assert backend.get_name() == "SDPA" + + +@hardware_test(res={"cuda": "L4"}, num_cards=1) +@pytest.mark.parametrize("dtype", [torch.bfloat16, torch.float32]) +@pytest.mark.parametrize("bias", [False, True]) +@torch.inference_mode() +def test_reference_text_qkv_cuda_preserves_provider_and_loaded_partitions(single_rank, mocker, dtype, bias): + layer, x, shards = _loaded_qkv(dtype, "cuda", bias=bias) + original = layer.quant_method + packed, _ = layer(x) + replacement = qwen._ReferenceTextQKVMethod(original) + assert replacement._gemm_impl is original._gemm_impl + # Observe calls to the actual configured provider; do not substitute its computation. + calls = mocker.Mock(wraps=replacement._gemm_impl) + replacement._gemm_impl = calls + layer.quant_method = replacement + actual, _ = layer(x) + + assert all(call.args[0] is layer and call.args[1] is x for call in calls.call_args_list) + if dtype == torch.bfloat16: + assert [call.args[2].shape[0] for call in calls.call_args_list] == layer.output_partition_sizes + independent = torch.cat([F.linear(x, weight, offset) for weight, offset in shards], dim=-1) + # Separate and packed BF16 GEMMs can round differently; neither is an FP32 reference. + torch.testing.assert_close(actual, independent, rtol=1e-2, atol=1e-2) + torch.testing.assert_close(actual, packed, rtol=1e-2, atol=1e-2) + else: + assert calls.call_count == 1 + torch.testing.assert_close(actual, packed, rtol=0, atol=0) + + +@hardware_test(res={"cuda": "L4"}, num_cards=1) +@pytest.mark.parametrize("reference_eager", [False, True]) +def test_qwen_reference_gate_uses_resolved_sdpa(single_rank, reference_eager): + config = OmniDiffusionConfig( + dtype=torch.bfloat16, + enforce_eager=reference_eager, + diffusion_attention_config=AttentionConfig(default=AttentionSpec(backend="TORCH_SDPA")), + ) + with set_current_diffusion_config(config): + attention = qwen.QwenImageCrossAttention( + dim=32, + num_heads=4, + head_dim=8, + added_kv_proj_dim=32, + reference_eager=reference_eager, + ) + + assert attention.attn.attn_backend.get_name() == "SDPA" + assert attention.reference_precision is reference_eager + expected_method = qwen._ReferenceTextQKVMethod if reference_eager else linear.UnquantizedLinearMethod + assert type(attention.add_kv_proj.quant_method) is expected_method + assert type(attention.to_qkv.quant_method) is linear.UnquantizedLinearMethod diff --git a/tests/model_executor/models/minicpmo_4_5/test_pipeline.py b/tests/model_executor/models/minicpmo_4_5/test_pipeline.py index 408971fa141..efc80e2fb77 100644 --- a/tests/model_executor/models/minicpmo_4_5/test_pipeline.py +++ b/tests/model_executor/models/minicpmo_4_5/test_pipeline.py @@ -226,7 +226,8 @@ def test_deploy_resolves_three_stage_pipeline(self, filename: str, devices: list def test_npu_graph_configuration_is_stage_scoped(self, filename: str, monkeypatch: pytest.MonkeyPatch) -> None: # Let the production merge select NPU once, independently of the host. # Applying NPU first would let merge_pipeline_deploy overlay CUDA again. - monkeypatch.setattr("vllm_omni.platforms.current_omni_platform", SimpleNamespace(device_name="npu")) + monkeypatch.setattr(current_omni_platform, "device_name", "npu") + monkeypatch.setattr(current_omni_platform, "device_type", "npu") deploy = load_deploy_config(_DEPLOY_DIR / filename) connector_extra = deploy.connectors["connector_of_shared_memory"]["extra"] assert "code2wav_enable_npu_graph" not in connector_extra diff --git a/tests/profile/test_omni_torch_profiler.py b/tests/profile/test_omni_torch_profiler.py index 7b594a4f75e..8bafe375e70 100644 --- a/tests/profile/test_omni_torch_profiler.py +++ b/tests/profile/test_omni_torch_profiler.py @@ -6,6 +6,7 @@ import gzip from dataclasses import dataclass +from functools import wraps from pathlib import Path from types import SimpleNamespace @@ -156,6 +157,7 @@ def fake_config(tmp_path): def fake_profiler_factory(monkeypatch): created: dict[str, object] = {} + @wraps(profiler_mod.torch.profiler.profile) def fake_profile(*args, **kwargs): profiler = FakeTorchProfiler(on_trace_ready=kwargs.get("on_trace_ready")) created["profiler"] = profiler @@ -198,6 +200,110 @@ def test_session_profiler_kwargs_do_not_replace_an_active_profiler(wrapper): assert wrapper.profiler is original +@pytest.mark.parametrize("overrides", [[], "record_shapes"]) +def test_session_profiler_kwargs_reject_non_mapping_without_changing_state(wrapper, overrides): + wrapper._active = False + original = wrapper.profiler + with pytest.raises(TypeError, match="profiler_kwargs must be a mapping"): + wrapper.set_profiler_kwargs(overrides) + assert wrapper.profiler is original + assert wrapper._profiler_kwargs == {} + + +def test_session_profiler_kwargs_reject_unknown_keyword_before_construction(wrapper, fake_profiler_factory): + wrapper._active = False + original = wrapper.profiler + with pytest.raises(TypeError, match="Invalid profiler_kwargs.*record_shape"): + wrapper.set_profiler_kwargs({"record_shape": True}) + assert wrapper.profiler is original + assert fake_profiler_factory["profiler"] is original + assert wrapper._profiler_kwargs == {} + + +def test_session_profiler_kwargs_reuse_unchanged_scalar_options_and_defaults(wrapper): + wrapper._active = False + original = wrapper.profiler + wrapper.set_profiler_kwargs(None) + assert wrapper.profiler is original + + wrapper.set_profiler_kwargs({"record_shapes": False}) + changed = wrapper.profiler + assert changed is not original + wrapper.set_profiler_kwargs({"record_shapes": False}) + assert wrapper.profiler is changed + + wrapper.set_profiler_kwargs(None) + restored = wrapper.profiler + assert restored is not changed + wrapper.set_profiler_kwargs({}) + assert wrapper.profiler is restored + + +@pytest.mark.parametrize("acc_events", [False, True]) +def test_session_profiler_kwargs_reused_cpu_profiler_keeps_sessions_separate(fake_config, monkeypatch, acc_events): + traces = [] + + def capture(_self, profiler): + traces.append({event.key for event in profiler.key_averages()}) + + monkeypatch.setattr(OmniTorchProfilerWrapper, "_on_trace_ready", capture) + wrapper = OmniTorchProfilerWrapper(fake_config, "worker0", 0, activities=["CPU"]) + wrapper._active = False + profilers = [] + for factory in (profiler_mod.torch.ones, profiler_mod.torch.zeros): + wrapper.set_profiler_kwargs({"acc_events": True} if acc_events else None) + profilers.append(wrapper.profiler) + with wrapper.profiler: + factory(2) + assert "aten::ones" in traces[0] and "aten::zeros" not in traces[0] + assert "aten::zeros" in traces[1] and "aten::ones" not in traces[1] + assert (profilers[0] is profilers[1]) is not acc_events + + +def test_session_profiler_kwargs_preserve_callbacks_activities_and_mutable_options(wrapper, fake_profiler_factory): + wrapper._active = False + activities = [profiler_mod.torch.profiler.ProfilerActivity.CPU] + + def callback(_profiler): + pass + + options = {"activities": activities, "on_trace_ready": callback, "acc_events": True} + wrapper.set_profiler_kwargs(options) + assert fake_profiler_factory["kwargs"]["activities"] is activities + assert fake_profiler_factory["kwargs"]["on_trace_ready"] is callback + assert fake_profiler_factory["kwargs"]["acc_events"] is True + original = wrapper.profiler + activities.append(profiler_mod.torch.profiler.ProfilerActivity.CUDA) + wrapper.set_profiler_kwargs(options) + assert wrapper.profiler is not original + assert fake_profiler_factory["kwargs"]["activities"] == activities + + +def test_session_profiler_kwargs_restore_state_when_backend_construction_fails(wrapper, monkeypatch): + wrapper._active = False + wrapper.set_profiler_kwargs({"record_shapes": False}) + original = wrapper.profiler + + def fail(*_args): + raise RuntimeError("profiler backend unavailable") + + monkeypatch.setattr(wrapper, "_create_profiler", fail) + with pytest.raises(RuntimeError, match="profiler backend unavailable"): + wrapper.set_profiler_kwargs({"record_shapes": True}) + assert wrapper.profiler is original + assert wrapper._profiler_kwargs == {"record_shapes": False} + + +def test_session_profiler_kwargs_log_active_overrides_without_changing_state(wrapper, caplog): + wrapper._active = True + original = wrapper.profiler + with caplog.at_level("WARNING"): + wrapper.set_profiler_kwargs({"record_shapes": False}) + assert "Ignoring profiler_kwargs" in caplog.text + assert wrapper.profiler is original + assert wrapper._profiler_kwargs == {} + + def test_set_trace_filename_creates_timestamped_session_dir(wrapper, monkeypatch, tmp_path): class FixedDatetime: @classmethod diff --git a/vllm_omni/core/sched/omni_scheduler_mixin.py b/vllm_omni/core/sched/omni_scheduler_mixin.py index 0ab1170b422..71ef24ee1d5 100644 --- a/vllm_omni/core/sched/omni_scheduler_mixin.py +++ b/vllm_omni/core/sched/omni_scheduler_mixin.py @@ -1018,6 +1018,11 @@ def _finish_error_requests(self, outputs: dict[int, list[EngineCoreOutput]]) -> # EngineCoreOutput.error. Snapshot first: replacing the attribute drops # the dict the caller still needs for the error field. grammar_error_reqs = getattr(self, "grammar_compile_error_reqs", None) + if grammar_error_reqs is not None and not isinstance(grammar_error_reqs, dict): + raise TypeError( + "grammar_compile_error_reqs must be a dict of RequestError values, " + f"got {type(grammar_error_reqs).__name__}" + ) request_errors: dict[str, Any] = dict(grammar_error_reqs) if isinstance(grammar_error_reqs, dict) else {} if isinstance(grammar_error_reqs, dict): self.grammar_compile_error_reqs: dict[str, Any] = {} diff --git a/vllm_omni/diffusion/models/magi2/modeling_magi2.py b/vllm_omni/diffusion/models/magi2/modeling_magi2.py index 4d850dc6b1c..ffcab28f069 100644 --- a/vllm_omni/diffusion/models/magi2/modeling_magi2.py +++ b/vllm_omni/diffusion/models/magi2/modeling_magi2.py @@ -23,7 +23,7 @@ from vllm_omni.diffusion.attention.backends.abstract import AttentionMetadata from vllm_omni.diffusion.attention.layer import Attention -from vllm_omni.diffusion.utils.torch_utils import _add_stored_gemm_outputs +from vllm_omni.diffusion.utils.torch_utils import add_stored_gemm_outputs from .attention import Magi2PackedAttentionKernel, VarlenHandler from .common_layers import ( @@ -316,7 +316,7 @@ def _shared_experts( # separately keeps both GEMM outputs and the down-projection inputs dense. shared = swiglu7(self.shared_expert_fc1(normalized)) modality = swiglu7(self.modality_specific_shared_expert_fc1(normalized, dispatcher)) - return _add_stored_gemm_outputs( + return add_stored_gemm_outputs( self.shared_expert_fc2(shared), self.modality_specific_shared_expert_fc2(modality, dispatcher), ) @@ -335,7 +335,7 @@ def combine( routed: torch.Tensor, dispatcher: ModalityDispatcher, ) -> torch.Tensor: - return _add_stored_gemm_outputs( + return add_stored_gemm_outputs( self.merge_linear(routed), self._shared_experts(normalized, dispatcher), ) diff --git a/vllm_omni/diffusion/utils/torch_utils.py b/vllm_omni/diffusion/utils/torch_utils.py index 696216c82c4..d593cbab78b 100644 --- a/vllm_omni/diffusion/utils/torch_utils.py +++ b/vllm_omni/diffusion/utils/torch_utils.py @@ -6,7 +6,7 @@ import torch -def _add_stored_gemm_outputs(left: torch.Tensor, right: torch.Tensor) -> torch.Tensor: +def add_stored_gemm_outputs(left: torch.Tensor, right: torch.Tensor) -> torch.Tensor: """Add GEMM outputs after each result has been stored. Inductor rewrites ``gemm + gemm`` as ``addmm`` and adds in the accumulator, diff --git a/vllm_omni/profiler/omni_torch_profiler.py b/vllm_omni/profiler/omni_torch_profiler.py index 42245b1866e..53b4d783636 100644 --- a/vllm_omni/profiler/omni_torch_profiler.py +++ b/vllm_omni/profiler/omni_torch_profiler.py @@ -5,7 +5,9 @@ import os import subprocess +from collections.abc import Mapping from datetime import datetime +from inspect import signature from typing import Any, Literal import torch @@ -107,14 +109,36 @@ def _create_profiler( "on_trace_ready": self._on_trace_ready, } kwargs.update(self._profiler_kwargs) + try: + signature(torch.profiler.profile).bind(**kwargs) + except TypeError as exc: + raise TypeError(f"Invalid profiler_kwargs: {exc}") from exc return torch.profiler.profile(**kwargs) def set_profiler_kwargs(self, profiler_kwargs: dict | None) -> None: """Apply constructor overrides to the next inactive profiling session.""" + if profiler_kwargs is not None and not isinstance(profiler_kwargs, Mapping): + raise TypeError("profiler_kwargs must be a mapping or None") + kwargs = dict(profiler_kwargs or {}) if self._active: + if kwargs: + logger.warning("Ignoring profiler_kwargs because profiling is active") + return + # Mutable or opaque options may have changed in place since construction. + if ( + all(value is None or type(value) in (bool, int, float, str, bytes) for value in kwargs.values()) + and kwargs == self._profiler_kwargs + and not kwargs.get("acc_events", False) + ): return - self._profiler_kwargs = dict(profiler_kwargs or {}) - self.profiler = self._create_profiler(self.profiler_config, self._activities) + previous_kwargs = self._profiler_kwargs + self._profiler_kwargs = kwargs + try: + profiler = self._create_profiler(self.profiler_config, self._activities) + except Exception: + self._profiler_kwargs = previous_kwargs + raise + self.profiler = profiler def set_trace_filename(self, filename: str) -> None: """Set the trace filename before starting profiling. From 67ace00a087dacb8b61fdbe2134a860261a8db73 Mon Sep 17 00:00:00 2001 From: tzhouam Date: Tue, 3 Nov 2026 23:53:46 +0800 Subject: [PATCH 6/8] [CI] Resolve cuDNN preload paths in CPU smoke check Signed-off-by: tzhouam --- docker/Dockerfile.ci | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docker/Dockerfile.ci b/docker/Dockerfile.ci index 81e1f966abd..d9c92fc8568 100644 --- a/docker/Dockerfile.ci +++ b/docker/Dockerfile.ci @@ -107,7 +107,7 @@ from pathlib import Path import torch maps = Path("/proc/self/maps").read_text() -assert any(path in maps for path in os.environ["LD_PRELOAD"].split() if "libcudnn.so.9" in path) +assert any(os.path.realpath(path) in maps for path in os.environ["LD_PRELOAD"].split() if "libcudnn.so.9" in path) assert torch.arange(4, device="cpu").sum().item() == 6 cuda_available = torch.cuda.is_available() assert not torch.cuda.is_initialized() From 3a803b6c4d7db44857f9b041f22b290d559a931b Mon Sep 17 00:00:00 2001 From: tzhouam Date: Wed, 4 Nov 2026 08:54:36 +0800 Subject: [PATCH 7/8] [Bugfix] Fix Simple test fixtures and parallel isolation Signed-off-by: tzhouam --- .../test_omni_ar_scheduler_kv_transfer.py | 1 + .../test_omni_ar_scheduler_stale_drain.py | 1 + .../sched/test_omni_ar_scheduler_streaming.py | 3 ++ ...mni_generation_scheduler_update_session.py | 2 + .../test_omni_sched_deferred_free_fence.py | 1 + .../test_omni_sched_ec_request_finish.py | 1 + .../test_omni_sched_prefill_stats_finalize.py | 1 + .../test_multiproc_engine_concurrency.py | 13 +++-- .../test_chunk_transfer_adapter.py | 3 ++ .../omni_connectors/test_shm_connector.py | 47 ++++++++++++------- .../openai_api/test_serving_speech.py | 10 +++- 11 files changed, 61 insertions(+), 22 deletions(-) diff --git a/tests/core/sched/test_omni_ar_scheduler_kv_transfer.py b/tests/core/sched/test_omni_ar_scheduler_kv_transfer.py index 5241cd3bcd5..7491f521349 100644 --- a/tests/core/sched/test_omni_ar_scheduler_kv_transfer.py +++ b/tests/core/sched/test_omni_ar_scheduler_kv_transfer.py @@ -113,6 +113,7 @@ def _run_finished_save_step(mocker, request, *, inter_stage_outputs=None): adapter._confirmed_num_computed_tokens.return_value = 1 sched = mocker.MagicMock() + sched.grammar_compile_error_reqs = {} sched.requests = {request.request_id: request} sched.perf_metrics = None sched.defer_block_free = False diff --git a/tests/core/sched/test_omni_ar_scheduler_stale_drain.py b/tests/core/sched/test_omni_ar_scheduler_stale_drain.py index e3c28cdf08c..6f9c56d9f13 100644 --- a/tests/core/sched/test_omni_ar_scheduler_stale_drain.py +++ b/tests/core/sched/test_omni_ar_scheduler_stale_drain.py @@ -93,6 +93,7 @@ def _replace_streaming_session(session: Request) -> None: def _make_drain_sched(session: Request) -> MagicMock: sched = MagicMock() + sched.grammar_compile_error_reqs = {} sched.requests = {session.request_id: session} sched.perf_metrics = None sched.structured_output_manager.accept_tokens.return_value = True diff --git a/tests/core/sched/test_omni_ar_scheduler_streaming.py b/tests/core/sched/test_omni_ar_scheduler_streaming.py index dbfc9f3215b..df983542f6b 100644 --- a/tests/core/sched/test_omni_ar_scheduler_streaming.py +++ b/tests/core/sched/test_omni_ar_scheduler_streaming.py @@ -165,6 +165,7 @@ def _run_resumable_segment_stop( inter_stage_output=None, ): sched = MagicMock() + sched.grammar_compile_error_reqs = {} sched._hold_payloadless_outputs = False sched.requests = {session.request_id: session} sched.perf_metrics = None @@ -325,6 +326,7 @@ def test_running_decode_step_without_inter_stage_payload_does_not_raise( session.additional_information = {"omni_final_stage_id": 0} if omit_chunk_transfer else None sched = mocker.MagicMock() + sched.grammar_compile_error_reqs = {} sched._hold_payloadless_outputs = False sched.requests = {session.request_id: session} sched.perf_metrics = None @@ -419,6 +421,7 @@ def test_stale_async_frame_is_dropped_before_output_processing() -> None: num_output_placeholders = session.num_output_placeholders sched = MagicMock() + sched.grammar_compile_error_reqs = {} sched._hold_payloadless_outputs = False sched.requests = {session.request_id: session} sched.perf_metrics = None diff --git a/tests/core/sched/test_omni_generation_scheduler_update_session.py b/tests/core/sched/test_omni_generation_scheduler_update_session.py index 884f043acbb..c07097bf56a 100644 --- a/tests/core/sched/test_omni_generation_scheduler_update_session.py +++ b/tests/core/sched/test_omni_generation_scheduler_update_session.py @@ -119,6 +119,7 @@ def test_resumable_generation_stop_marks_segment_boundary() -> None: session.num_in_flight_tokens = 1 sched = MagicMock() + sched.grammar_compile_error_reqs = {} sched._first_chunk_express = False sched._express_min_slack_s = 0.0 sched.requests = {session.request_id: session} @@ -190,6 +191,7 @@ def test_rejected_grammar_finishes_and_frees_generation_request(prompt_complete: session.num_in_flight_tokens = 1 sched = MagicMock() + sched.grammar_compile_error_reqs = {} sched._first_chunk_express = False sched._express_min_slack_s = 0.0 sched.requests = {session.request_id: session} diff --git a/tests/core/sched/test_omni_sched_deferred_free_fence.py b/tests/core/sched/test_omni_sched_deferred_free_fence.py index a4242349a18..43deea7fe8a 100644 --- a/tests/core/sched/test_omni_sched_deferred_free_fence.py +++ b/tests/core/sched/test_omni_sched_deferred_free_fence.py @@ -49,6 +49,7 @@ def _make_update_sched(scheduler_cls) -> MagicMock: """MagicMock scheduler driving the real ``update_from_output`` with an empty batch; only the fence state carries real values.""" sched = MagicMock() + sched.grammar_compile_error_reqs = {} sched.requests = {} sched.perf_metrics = None sched.chunk_transfer_adapter = None diff --git a/tests/core/sched/test_omni_sched_ec_request_finish.py b/tests/core/sched/test_omni_sched_ec_request_finish.py index f1b747adc04..96219247062 100644 --- a/tests/core/sched/test_omni_sched_ec_request_finish.py +++ b/tests/core/sched/test_omni_sched_ec_request_finish.py @@ -120,6 +120,7 @@ def _make_finish_sched(request: Request) -> MagicMock: """MagicMock scheduler driving the real ``update_from_output`` through a finishing request; ``_free_request`` reports EC transfer params.""" sched = MagicMock() + sched.grammar_compile_error_reqs = {} sched._first_chunk_express = False sched._express_min_slack_s = 0.0 sched.requests = {request.request_id: request} diff --git a/tests/core/sched/test_omni_sched_prefill_stats_finalize.py b/tests/core/sched/test_omni_sched_prefill_stats_finalize.py index 3438466b7a0..87aae356c25 100644 --- a/tests/core/sched/test_omni_sched_prefill_stats_finalize.py +++ b/tests/core/sched/test_omni_sched_prefill_stats_finalize.py @@ -53,6 +53,7 @@ def _make_request() -> Request: def _make_sched(request: Request) -> MagicMock: sched = MagicMock() + sched.grammar_compile_error_reqs = {} sched._first_chunk_express = False sched._express_min_slack_s = 0.0 sched.requests = {request.request_id: request} diff --git a/tests/diffusion/test_multiproc_engine_concurrency.py b/tests/diffusion/test_multiproc_engine_concurrency.py index 520e0c2e37e..ee5eef7f9e3 100644 --- a/tests/diffusion/test_multiproc_engine_concurrency.py +++ b/tests/diffusion/test_multiproc_engine_concurrency.py @@ -3,12 +3,14 @@ import asyncio import multiprocessing as mp +import os import queue import signal import threading import time import weakref from collections import OrderedDict +from multiprocessing.process import BaseProcess from types import SimpleNamespace from unittest.mock import MagicMock, Mock @@ -709,7 +711,7 @@ def test_results_are_correctly_routed(self): a_enqueued, b_complete = _inject_interleave(executor) wt = _start_worker(req_q, res_q, count=2) - results: dict[str, object] = {} + results: dict[str, DiffusionOutput] = {} def _a(): results["A"] = engine.collective_rpc( @@ -750,7 +752,7 @@ def test_results_are_correctly_routed(self): a_enqueued, b_complete = _inject_interleave(executor) wt = _start_worker(req_q, res_q, count=2) - results: dict[str, object] = {} + results: dict[str, DiffusionOutput] = {} def _a(): # request execution path results["A"] = engine.add_req_and_wait_for_response(_mock_request("A")) @@ -1544,8 +1546,11 @@ def run_worker(): signal.signal(signal.SIGTERM, signal.SIG_IGN) ready.set() stop.wait() + # This event-only worker must not spend its 0.2s shutdown budget + # flushing the test runner's inherited coverage data. + os._exit(0) - dead = ctx.Process(target=lambda: None) + dead = ctx.Process(target=os._exit, args=(0,)) worker = ctx.Process(target=run_worker) monkeypatch.setattr(executor_module, "_WORKER_SHUTDOWN_GRACE_S", 0.2) monkeypatch.setattr(executor_module, "_WORKER_TERMINATE_GRACE_S", 0.1) @@ -1673,7 +1678,7 @@ def _poll_flag(get_flag, *, timeout=5.0, interval=0.05) -> bool: return False -def _make_short_lived_process() -> mp.Process: +def _make_short_lived_process() -> BaseProcess: """Spawn a real subprocess that exits immediately. The process must be started with ``"fork"`` (or the platform default) diff --git a/tests/distributed/omni_connectors/test_chunk_transfer_adapter.py b/tests/distributed/omni_connectors/test_chunk_transfer_adapter.py index 5acec75eda1..c2788803255 100644 --- a/tests/distributed/omni_connectors/test_chunk_transfer_adapter.py +++ b/tests/distributed/omni_connectors/test_chunk_transfer_adapter.py @@ -2859,6 +2859,7 @@ def test_generation_scheduler_calls_cleanup_on_finished(monkeypatch, mocker: Moc from vllm_omni.core.sched.omni_generation_scheduler import OmniGenerationScheduler scheduler = mocker.MagicMock() + scheduler.grammar_compile_error_reqs = {} scheduler._first_chunk_express = False scheduler.aux_output_connector = None scheduler._express_min_slack_s = 0.0 @@ -2943,6 +2944,7 @@ def test_ar_scheduler_defers_cleanup_and_queues_save_on_finished(mocker: MockerF adapter_mock.save_async = lambda *a, **kw: save_calls.append((a, kw)) scheduler = mocker.MagicMock() + scheduler.grammar_compile_error_reqs = {} scheduler.aux_output_connector = None scheduler.chunk_transfer_adapter = adapter_mock bind_omits_transfer_helpers(scheduler) @@ -3457,6 +3459,7 @@ def _build_deferred_finish_scheduler(mocker, *, running, pending_finish_reqs): adapter_mock.cleanup = lambda *a, **kw: cleanup_calls.append((a, kw)) scheduler = mocker.MagicMock() + scheduler.grammar_compile_error_reqs = {} scheduler.chunk_transfer_adapter = adapter_mock scheduler.connector = None scheduler.ec_connector = None diff --git a/tests/distributed/omni_connectors/test_shm_connector.py b/tests/distributed/omni_connectors/test_shm_connector.py index 6abd3a55176..e017ae572f5 100644 --- a/tests/distributed/omni_connectors/test_shm_connector.py +++ b/tests/distributed/omni_connectors/test_shm_connector.py @@ -458,18 +458,21 @@ def test_wakeups_can_be_disabled_and_close_unlinks_fifo(monkeypatch): receiver.close() -def _sibling_receiver(stage_id, ready, result): +def _sibling_receiver(stage_id, ready, start, result): receiver = _stage_connector(stage_id) generation = receiver.get_wakeup_generation() ready.set() - start = time.monotonic() + assert start.wait(5) + started = time.monotonic() woke = receiver.wait_for_change(generation, timeout=5) - result.put((woke, time.monotonic() - start)) + result.put((woke, time.monotonic() - started)) receiver.close() -def _sibling_sender(stage_id, key): +def _sibling_sender(stage_id, key, ready, start): sender = _stage_connector(stage_id - 1) + ready.set() + assert start.wait(5) sender.put(str(stage_id - 1), str(stage_id), key, {"value": 3}) @@ -478,20 +481,30 @@ def test_sibling_stage_processes_share_wakeups(): import multiprocessing ctx = multiprocessing.get_context("fork") - ready, result = ctx.Event(), ctx.Queue() + ready, sender_ready, start, result = ctx.Event(), ctx.Event(), ctx.Event(), ctx.Queue() key = f"wake_{uuid.uuid4().hex}" - receiver = ctx.Process(target=_sibling_receiver, args=(78, ready, result)) - receiver.start() - assert ready.wait(5) - time.sleep(0.05) - sender = ctx.Process(target=_sibling_sender, args=(78, key)) - sender.start() - sender.join(5) - woke, waited = result.get(timeout=5) - receiver.join(5) - assert woke and waited < 1.0 - cleanup = make_connector({}) - cleanup.cleanup(key) + receiver = ctx.Process(target=_sibling_receiver, args=(78, ready, start, result)) + sender = ctx.Process(target=_sibling_sender, args=(78, key, sender_ready, start)) + try: + receiver.start() + assert ready.wait(5) + sender.start() + assert sender_ready.wait(5) + start.set() + sender.join(5) + woke, waited = result.get(timeout=5) + receiver.join(5) + assert woke and waited < 1.0 + finally: + start.set() + for process in (receiver, sender): + if process.pid is not None: + if process.is_alive(): + process.kill() + process.join(5) + process.close() + cleanup = make_connector({}) + cleanup.cleanup(key) def test_same_stage_receivers_are_independent_across_close_and_restart(): diff --git a/tests/entrypoints/openai_api/test_serving_speech.py b/tests/entrypoints/openai_api/test_serving_speech.py index 7c357a257c0..93f059545a9 100644 --- a/tests/entrypoints/openai_api/test_serving_speech.py +++ b/tests/entrypoints/openai_api/test_serving_speech.py @@ -431,7 +431,15 @@ def client(test_app): @pytest.mark.asyncio @pytest.mark.parametrize("enable_scale_out", [False, True], ids=["default", "scale-out"]) -async def test_server_worker_applies_speech_only_endpoint_restrictions(mocker: MockerFixture, enable_scale_out): +async def test_server_worker_applies_speech_only_endpoint_restrictions( + mocker: MockerFixture, enable_scale_out, request: pytest.FixtureRequest +): + # SageMaker route registration imports a dependency that changes root.level. + # Restore it through setLevel so later tests also clear cached log filters. + root_logger = logging.getLogger() + previous_level = root_logger.level + request.addfinalizer(lambda: root_logger.setLevel(previous_level)) + class FakeCompletionHandler: def __init__(self): self.create_completion = mocker.AsyncMock(return_value={"id": "cmpl", "choices": []}) From d8397156120955c4e24141519356bbc8977243c4 Mon Sep 17 00:00:00 2001 From: tzhouam Date: Wed, 4 Nov 2026 11:43:19 +0800 Subject: [PATCH 8/8] [Bugfix] Publish synthetic image cache atomically --- tests/helpers/media.py | 39 ++++--- tests/helpers/tests/test_synthetic_image.py | 110 ++++++++++++++++++++ 2 files changed, 137 insertions(+), 12 deletions(-) create mode 100644 tests/helpers/tests/test_synthetic_image.py diff --git a/tests/helpers/media.py b/tests/helpers/media.py index 25e41434803..2470f4804ec 100644 --- a/tests/helpers/media.py +++ b/tests/helpers/media.py @@ -522,7 +522,8 @@ def generate_synthetic_image( """ Random colored squares on white background. Caches JPEG by ``width`` / ``height`` unless ``force_regenerate`` is true. Cache root: ``cache_dir`` - if given, else the default temp subdirectory. + if given, else the default temp subdirectory. Incomplete cached images are + regenerated, and complete JPEGs are published atomically for parallel readers. """ if seed is not None: random.seed(seed) @@ -532,16 +533,21 @@ def generate_synthetic_image( cache_path = root / f"synth_image_w{width}_h{height}.jpg" if not force_regenerate and cache_path.is_file(): - from PIL import Image as PILImage - - image = PILImage.open(cache_path) - image.load() + # Decode the same snapshot that is returned, even if another worker + # replaces the cache between decoding and payload construction. image_bytes = cache_path.read_bytes() - return { - "np_array": np.array(image).copy(), - "base64": base64.b64encode(image_bytes).decode("utf-8"), - "file_path": str(cache_path.resolve()), - } + try: + with Image.open(io.BytesIO(image_bytes)) as image: + image.load() + cached_pixels = np.array(image).copy() + except OSError: + logger.warning("Cached synthetic image %s is incomplete; regenerating.", cache_path) + else: + return { + "np_array": cached_pixels, + "base64": base64.b64encode(image_bytes).decode("utf-8"), + "file_path": str(cache_path.resolve()), + } from PIL import ImageDraw @@ -556,8 +562,17 @@ def generate_synthetic_image( border_width = random.randint(1, 5) draw.rectangle([x, y, x + square_size, y + square_size], fill=color, outline=(0, 0, 0), width=border_width) - image.save(str(cache_path), format="JPEG", quality=85, optimize=True) - image_bytes = cache_path.read_bytes() + buffer = io.BytesIO() + image.save(buffer, format="JPEG", quality=85, optimize=True) + image_bytes = buffer.getvalue() + temp_cache_path = root / f".{cache_path.name}.{uuid.uuid4().hex}.tmp" + temp_cache = temp_cache_path.open("xb") + try: + with temp_cache: + temp_cache.write(image_bytes) + os.replace(temp_cache_path, cache_path) + finally: + temp_cache_path.unlink(missing_ok=True) return { "np_array": np.array(image).copy(), diff --git a/tests/helpers/tests/test_synthetic_image.py b/tests/helpers/tests/test_synthetic_image.py new file mode 100644 index 00000000000..92261944e21 --- /dev/null +++ b/tests/helpers/tests/test_synthetic_image.py @@ -0,0 +1,110 @@ +# SPDX-License-Identifier: Apache-2.0 +# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project + +import base64 +import io +import threading +from concurrent.futures import ThreadPoolExecutor +from pathlib import Path + +import numpy as np +import pytest +from PIL import Image + +from tests.helpers import media + +pytestmark = [pytest.mark.core_model, pytest.mark.cpu] + + +@pytest.mark.parametrize("existing_cache", [False, True]) +def test_readers_do_not_see_an_incomplete_jpeg(tmp_path, monkeypatch, existing_cache): + previous = media.generate_synthetic_image(32, 24, cache_dir=tmp_path) if existing_cache else None + encoding = threading.Event() + release = threading.Event() + save = Image.Image.save + + def pause_writer(image, fp, *args, **kwargs): + if threading.current_thread() is not threading.main_thread(): + # Reproduce Pillow's final-path truncation before encoding starts. + # Encoding into memory must leave the published cache untouched. + if isinstance(fp, (str, Path)): + with open(fp, "w+b") as stream: + encoding.set() + assert release.wait(10) + return save(image, stream, *args, **kwargs) + encoding.set() + assert release.wait(10) + return save(image, fp, *args, **kwargs) + + monkeypatch.setattr(Image.Image, "save", pause_writer) + with ThreadPoolExecutor(max_workers=1) as pool: + writer = pool.submit(media.generate_synthetic_image, 32, 24, cache_dir=tmp_path, force_regenerate=True) + try: + assert encoding.wait(10) + reader = media.generate_synthetic_image(32, 24, cache_dir=tmp_path) + with Image.open(io.BytesIO(base64.b64decode(reader["base64"]))) as image: + image.load() + assert image.size == (32, 24) + if previous is not None: + assert reader["base64"] == previous["base64"] + finally: + release.set() + published = writer.result(timeout=10) + + path = Path(published["file_path"]) + assert path.read_bytes() == base64.b64decode(published["base64"]) + assert list(tmp_path.iterdir()) == [path] + + +@pytest.mark.parametrize("failure_at", ["encode", "publish"]) +def test_failed_image_regeneration_preserves_cache(tmp_path, monkeypatch, failure_at): + previous = media.generate_synthetic_image(32, 24, cache_dir=tmp_path) + path = Path(previous["file_path"]) + + def fail(*args, **kwargs): + raise OSError("image publication failed") + + if failure_at == "encode": + monkeypatch.setattr(Image.Image, "save", fail) + else: + monkeypatch.setattr(media.os, "replace", fail) + with pytest.raises(OSError, match="image publication failed"): + media.generate_synthetic_image(32, 24, cache_dir=tmp_path, force_regenerate=True) + + assert path.read_bytes() == base64.b64decode(previous["base64"]) + assert list(tmp_path.iterdir()) == [path] + + +@pytest.mark.parametrize("broken_bytes", [b"", b"not a JPEG", b"\xff\xd8\xff\xe0"]) +def test_incomplete_image_cache_is_regenerated(tmp_path, broken_bytes): + path = tmp_path / "synth_image_w32_h24.jpg" + path.write_bytes(broken_bytes) + + result = media.generate_synthetic_image(32, 24, cache_dir=tmp_path) + + assert path.read_bytes() == base64.b64decode(result["base64"]) + with Image.open(path) as image: + image.load() + assert image.size == (32, 24) + assert list(tmp_path.iterdir()) == [path] + + +def test_cached_pixels_and_payload_use_the_same_generation(tmp_path, monkeypatch): + media.generate_synthetic_image(32, 24, cache_dir=tmp_path) + cache = tmp_path / "synth_image_w32_h24.jpg" + replacement = tmp_path / "replacement.jpg" + Image.new("RGB", (32, 24), "black").save(replacement) + open_image = Image.open + + def replace_after_decode(fp, *args, **kwargs): + image = open_image(fp, *args, **kwargs) + image.load() + replacement.replace(cache) + return image + + with monkeypatch.context() as patch: + patch.setattr(Image, "open", replace_after_decode) + result = media.generate_synthetic_image(32, 24, cache_dir=tmp_path) + + with Image.open(io.BytesIO(base64.b64decode(result["base64"]))) as image: + np.testing.assert_array_equal(result["np_array"], np.array(image))