Compare commits

...
Author SHA1 Message Date
Sy03 dedb045225 [CI] Incorporate upstream Server VAD test quarantine
Merge upstream #7314 for the known Qwen3-Omni Server VAD input_backpressure failure tracked in #7279. Preserve all MOSS-TTS runtime changes and existing validation coverage.

Signed-off-by: Sy03 <1370724210@qq.com>
2026-09-09 19:53:28 +08:00
Sy03 28dd49a541 [Maintenance] Refresh MOSS-TTS optimization on main
Signed-off-by: Sy03 <1370724210@qq.com>
2026-09-09 12:46:57 +08:00
Sy03 92cecf0296 [Bugfix] Score complete long audio in Seed-TTS WER
Retain audio beyond the default Whisper 30-second window and enable long-form timestamp decoding. Extend the existing transcription regression with a 61-second case while preserving the short-audio path. Clarify the existing test loader and per-turn request copy for static checking.

Signed-off-by: Sy03 <1370724210@qq.com>
2026-09-08 17:40:59 +08:00
Sy03 8231f5c978 [Bugfix] Use reference transcripts for MOSS-TTS Local continuation
Build the official assistant audio prefix when Local cloning requests include a reference transcript. Preserve the audio-only fallback, conditioning cache identity, seed and output budget. Cover the prompt contract and fallback in the existing adapter tests.

Signed-off-by: Sy03 <1370724210@qq.com>
2026-09-08 16:00:43 +08:00
Sy03 d530938b5b [Model] Reuse frame-local KV in MOSS-TTS Local
Cache each position once while generating the audio codebooks, preserving frame-local RoPE positions and the full-prefix reference path. Share compiled batch and position shapes to avoid specialization-limit fallback. Add one compact parity regression covering frame reuse.

Signed-off-by: Sy03 <1370724210@qq.com>
2026-09-08 06:37:00 +08:00
Sy03 30d28a9a04 [Test] Consolidate MOSS-TTS batching regression coverage
Combine overlapping row/snapshot checks, cover masked and valid attention rows in one batch, reuse the real codec session in grouping tests, and move output-policy coverage into the existing entrypoint module. Preserve both CUDA synchronization regressions and reduce the added suite from 29 to 15 cases.

Signed-off-by: Sy03 <1370724210@qq.com>
2026-09-08 03:03:32 +08:00
Sy03 55855fee9c [Model] Reduce MOSS-TTS Local intermediate output overhead
Keep internal talker control outputs FINAL_ONLY while codec audio remains incremental. Reapply pipeline sampling constraints after default streaming coercion, log routine coercions at debug level, and cover explicit/default parameters with regression tests. Document matched high-concurrency statistics settings.

Signed-off-by: Sy03 <1370724210@qq.com>
2026-09-08 01:15:46 +08:00
Sy03 dcd6588a36 [Perf][MOSS-TTS] Avoid synchronization in batched index transfers
Construct MTP row offsets and mixed prefill/decode stop indices on the CPU, then enqueue nonblocking device transfers before batched writes. Add CUDA regressions for synchronization, row placement, stop flags, and snapshot ownership.

Signed-off-by: Sy03 <1370724210@qq.com>
2026-09-08 00:39:50 +08:00
Sy03 251ef4f48f [Model] Optimize MOSS-TTS Local batched execution and streaming codec
Signed-off-by: Sy03 <1370724210@qq.com>
2026-09-07 16:20:49 +08:00
25 changed files with 843 additions and 114 deletions

No files matched your search

+1 -1
View File
@@ -93,7 +93,7 @@ steps:
commands: commands:
- | - |
timeout 15m bash -c " timeout 15m bash -c "
pytest -sv tests/engine/ tests/model_executor/ -m 'core_model and cuda' --run-level 'core_model' pytest -sv tests/engine/ tests/model_executor/ tests/worker/test_batched_omni_output.py -m 'core_model and cuda' --run-level 'core_model'
" "
mirror_hardwares: l4_1 mirror_hardwares: l4_1
@@ -259,21 +259,25 @@ python examples/online_serving/text_to_speech/indextts2/speech_client.py \
0.6B DualAR TTS at 44.1 kHz, 11 languages, zero-shot voice cloning. 0.6B DualAR TTS at 44.1 kHz, 11 languages, zero-shot voice cloning.
### Prerequisites ### Prerequisites
None beyond the base install: the neural audio codec is implemented in tree and None beyond the base install: the neural audio codec is implemented in tree and
its weights (`codec.pth`) ship with the checkpoint. its weights (`codec.pth`) ship with the checkpoint.
### Launch ### Launch
```bash ```bash
vllm serve Audio8/Audio8-TTS-Preview-0.6b --omni --port 8092 vllm serve Audio8/Audio8-TTS-Preview-0.6b --omni --port 8092
# or: # or:
./audio8_tts/run_server.sh ./audio8_tts/run_server.sh
``` ```
The deploy config auto-loads from `vllm_omni/deploy/audio8_tts.yaml` (HF The deploy config auto-loads from `vllm_omni/deploy/audio8_tts.yaml` (HF
`model_type` is `arktts`). Do **not** pass `--trust-remote-code`: vllm-omni `model_type` is `arktts`). Do **not** pass `--trust-remote-code`: vllm-omni
registers its own `arktts` config, and transformers would otherwise prefer the registers its own `arktts` config, and transformers would otherwise prefer the
checkpoint's remote code and bypass it. checkpoint's remote code and bypass it.
### Text-only synthesis ### Text-only synthesis
```bash ```bash
curl -X POST http://localhost:8092/v1/audio/speech \ curl -X POST http://localhost:8092/v1/audio/speech \
-H "Content-Type: application/json" \ -H "Content-Type: application/json" \
@@ -284,6 +288,7 @@ curl -X POST http://localhost:8092/v1/audio/speech \
``` ```
### Voice cloning ### Voice cloning
```bash ```bash
curl -X POST http://localhost:8092/v1/audio/speech \ curl -X POST http://localhost:8092/v1/audio/speech \
-H "Content-Type: application/json" \ -H "Content-Type: application/json" \
@@ -293,18 +298,21 @@ curl -X POST http://localhost:8092/v1/audio/speech \
"ref_text": "The exact transcript of the reference recording." "ref_text": "The exact transcript of the reference recording."
}' --output cloned.wav }' --output cloned.wav
``` ```
`ref_audio` also accepts a `data:audio/wav;base64,...` URL. `ref_text` is `ref_audio` also accepts a `data:audio/wav;base64,...` URL. `ref_text` is
mandatory whenever `ref_audio` is present and must match the recording. mandatory whenever `ref_audio` is present and must match the recording.
Uploading a voice via `POST /v1/audio/voices` lets the server reuse the encoded Uploading a voice via `POST /v1/audio/voices` lets the server reuse the encoded
reference codes across requests (`voice: "<name>"`). reference codes across requests (`voice: "<name>"`).
### Streaming ### Streaming
```bash ```bash
python audio8_tts/speech_client.py --text "Welcome to Audio8 TTS." --stream --output out.pcm python audio8_tts/speech_client.py --text "Welcome to Audio8 TTS." --stream --output out.pcm
ffplay -f s16le -ar 44100 -ac 1 out.pcm ffplay -f s16le -ar 44100 -ac 1 out.pcm
``` ```
### Gradio demo ### Gradio demo
```bash ```bash
./audio8_tts/run_gradio_demo.sh # server + demo ./audio8_tts/run_gradio_demo.sh # server + demo
python audio8_tts/gradio_demo.py --api-base http://localhost:8092 # demo only python audio8_tts/gradio_demo.py --api-base http://localhost:8092 # demo only
@@ -328,7 +336,7 @@ these as a lower bound. 20 prompts from `seed_tts_smoke/en`, mean audio 4.5 s,
deploy defaults: deploy defaults:
| | HF reference (batch=1) | vllm-omni c=1 | c=4 | c=8 | | | HF reference (batch=1) | vllm-omni c=1 | c=4 | c=8 |
|---|---|---|---|---| | --- | --- | --- | --- | --- |
| RTF (mean) | 1.008 | **0.19** | 0.30 | 0.54 | | RTF (mean) | 1.008 | **0.19** | 0.30 | 0.54 |
| Throughput (req/s) | 0.23 | 1.18 | 2.84 | **2.93** | | Throughput (req/s) | 0.23 | 1.18 | 2.84 | **2.93** |
| E2E latency (mean, ms) | 4375 | 850 | 1349 | 2361 | | E2E latency (mean, ms) | 4375 | 850 | 1349 | 2361 |
@@ -347,7 +355,7 @@ job holding ~80% SM, no repeats), so read them as directional, not as measured
speedups. Same setup, varying stage 0's `max_num_seqs` (stage 1 stays at 1): speedups. Same setup, varying stage 0's `max_num_seqs` (stage 1 stays at 1):
| stage-0 `max_num_seqs` | client concurrency | req/s | RTF | TTFP mean / p99 (ms) | underrun p99 (s) | | stage-0 `max_num_seqs` | client concurrency | req/s | RTF | TTFP mean / p99 (ms) | underrun p99 (s) |
|---|---|---|---|---|---| | --- | --- | --- | --- | --- | --- |
| 4 (default) | 1 | 1.18 | 0.19 | 63 / 67 | 0.00 | | 4 (default) | 1 | 1.18 | 0.19 | 63 / 67 | 0.00 |
| 4 | 4 | 2.84 | 0.30 | 105 / 148 | 0.00 | | 4 | 4 | 2.84 | 0.30 | 105 / 148 | 0.00 |
| 4 | 8 | 2.93 | 0.54 | 1124 / 1642 | 0.00 | | 4 | 8 | 2.93 | 0.54 | 1124 / 1642 | 0.00 |
@@ -376,6 +384,7 @@ underrun. Raise it to 8 if first-packet latency under load matters more than a
rare sub-300 ms gap. rare sub-300 ms gap.
### Notes ### Notes
- Output: 44.1 kHz mono; ~21.5 codec frames per second. - Output: 44.1 kHz mono; ~21.5 codec frames per second.
- No built-in speaker presets. Omit `voice` for a random timbre, or clone one. - No built-in speaker presets. Omit `voice` for a random timbre, or clone one.
- `max_new_tokens` caps generated codec frames (1 frame ~= 46 ms). - `max_new_tokens` caps generated codec frames (1 frame ~= 46 ms).
@@ -553,6 +562,46 @@ python examples/online_serving/text_to_speech/ming_flash_omni_tts/speech_client.
--- ---
## MOSS-TTS Local Transformer v1.5
For a single H200, the optional
[`moss_tts_local_h200.yaml`](../../../vllm_omni/deploy/moss_tts_local_h200.yaml)
deployment places both the talker and codec on logical GPU 0:
```bash
CUDA_VISIBLE_DEVICES=0 VLLM_OMNI_EVENT_DRIVEN_ORCH=1 \
vllm serve OpenMOSS-Team/MOSS-TTS-Local-Transformer-v1.5 \
--omni --trust-remote-code \
--deploy-config vllm_omni/deploy/moss_tts_local_h200.yaml \
--disable-log-stats
```
Run from the repository root and select an available physical GPU with
`CUDA_VISIBLE_DEVICES`. The command enables event-driven orchestration and
disables detailed per-request statistics logging to reduce CPU overhead at
high concurrency. Remove `--disable-log-stats` when those statistics are needed;
keep these settings identical when comparing performance.
This preset configures 256 request slots per stage,
a 32 GiB talker KV cache, talker CUDA Graph buckets through 512 scheduled tokens,
and codec batch buckets through 256. It requires more than 80 GiB of GPU memory;
request capacity also depends on input and generated lengths. The default
[`moss_tts_local.yaml`](../../../vllm_omni/deploy/moss_tts_local.yaml) remains
available for smaller deployments.
The preset enables the optional codec backend with
`hf_overrides.codec_attention_backend: triton` on stage 1. It preserves the
streaming ring-cache mask and uses BF16 attention with 64-dimensional heads;
other attention shapes use PyTorch SDPA. Set the backend to `sdpa` to use the
default attention implementation. Codec terminal tails share execution only
when they already map to the same padded CUDA Graph; returned audio retains
each request's actual length.
Voice cloning requests use `ref_audio` and `ref_text`. For streaming output,
set `stream: true`, `stream_format: "audio"`, and `response_format: "pcm"`;
the native PCM format is 48 kHz, stereo, signed 16-bit little-endian. Compute
audio throughput as `PCM bytes / (48000 * 2 * 2) / elapsed seconds`, and compare
the same requests, concurrency, and reference-cache state.
## MOSS-TTS-Nano ## MOSS-TTS-Nano
Single-stage 0.1B AR LM + MOSS-Audio-Tokenizer-Nano codec at 48 kHz mono. Every request must include `ref_audio`; there are no built-in speaker presets. Single-stage 0.1B AR LM + MOSS-Audio-Tokenizer-Nano codec at 48 kHz mono. Every request must include `ref_audio`; there are no built-in speaker presets.
+10 -2
View File
@@ -16,7 +16,7 @@ Use this recipe for 24 kHz multilingual TTS with voice cloning (20 languages inc
Chinese and English). Choose a variant based on your latency and quality requirements: Chinese and English). Choose a variant based on your latency and quality requirements:
| Model | Params | Use case | | Model | Params | Use case |
|---|---|---| | --- | --- | --- |
| MOSS-TTS | 8B | General TTS, highest quality | | MOSS-TTS | 8B | General TTS, highest quality |
| MOSS-TTS-v1.5 | 8B | General TTS upgrade of 1.0: 31 languages, steadier cloning, `[pause Xs]` markers (set `language` for best results); same `MossTTSDelay` API | | MOSS-TTS-v1.5 | 8B | General TTS upgrade of 1.0: 31 languages, steadier cloning, `[pause Xs]` markers (set `language` for best results); same `MossTTSDelay` API |
| MOSS-TTS-Realtime | 1.7B | Lowest latency (TTFB ~180 ms), streaming-first | | MOSS-TTS-Realtime | 1.7B | Lowest latency (TTFB ~180 ms), streaming-first |
@@ -24,9 +24,17 @@ Chinese and English). Choose a variant based on your latency and quality require
| MOSS-SoundEffect | 8B | Sound effect synthesis from text description | | MOSS-SoundEffect | 8B | Sound effect synthesis from text description |
| MOSS-VoiceGenerator | 1.7B | Zero-shot voice design | | MOSS-VoiceGenerator | 1.7B | Zero-shot voice design |
All variants share the same codec (`OpenMOSS-Team/MOSS-Audio-Tokenizer`, ~7 GB) and The variants above share the same codec (`OpenMOSS-Team/MOSS-Audio-Tokenizer`, ~7 GB) and
output 24 kHz mono audio. output 24 kHz mono audio.
MOSS-TTS-Local-Transformer-v1.5 uses MOSS-Audio-Tokenizer-v2 and outputs 48 kHz
stereo audio. For Local voice cloning through `/v1/audio/speech`, provide an
accurate `ref_text` transcript alongside `ref_audio`. The adapter uses the
reference transcript and audio as a continuation prefix, then generates only
the requested `input` speech. Without a nonblank `ref_text`, Local uses
audio-reference generation. Reference transcripts must match the reference
audio; they are not style instructions.
## References ## References
- Offline inference example: [`examples/offline_inference/text_to_speech/moss_tts/`](../../examples/offline_inference/text_to_speech/moss_tts/) - Offline inference example: [`examples/offline_inference/text_to_speech/moss_tts/`](../../examples/offline_inference/text_to_speech/moss_tts/)
@@ -26,6 +26,7 @@ _MODULE_NAME = "vllm_omni.benchmarks.data_modules.seed_tts_dataset"
if _MODULE_NAME not in sys.modules: if _MODULE_NAME not in sys.modules:
_spec = importlib.util.spec_from_file_location(_MODULE_NAME, _MODULE_PATH) _spec = importlib.util.spec_from_file_location(_MODULE_NAME, _MODULE_PATH)
assert _spec is not None and _spec.loader is not None
_mod = importlib.util.module_from_spec(_spec) _mod = importlib.util.module_from_spec(_spec)
sys.modules[_MODULE_NAME] = _mod sys.modules[_MODULE_NAME] = _mod
_spec.loader.exec_module(_mod) _spec.loader.exec_module(_mod)
@@ -237,7 +238,8 @@ def test_seed_tts_design_dataset_rejects_missing_description(seed_tts_design_roo
assert req.seed_tts_utterance_id == "ok" assert req.seed_tts_utterance_id == "ok"
def test_seed_tts_whisper_transcribe_passes_attention_mask(monkeypatch): @pytest.mark.parametrize("audio_seconds", [0.1, 61])
def test_seed_tts_whisper_transcribe_passes_attention_mask(monkeypatch, audio_seconds):
from vllm_omni.benchmarks.data_modules import seed_tts_eval from vllm_omni.benchmarks.data_modules import seed_tts_eval
calls = {} calls = {}
@@ -252,8 +254,9 @@ def test_seed_tts_whisper_transcribe_passes_attention_mask(monkeypatch):
return self return self
class FakeProcessor: class FakeProcessor:
def __call__(self, wav, *, sampling_rate, return_tensors, return_attention_mask=False): def __call__(self, wav, *, sampling_rate, return_tensors, return_attention_mask=False, **kwargs):
calls["return_attention_mask"] = return_attention_mask calls["return_attention_mask"] = return_attention_mask
calls["processor_kwargs"] = kwargs
assert sampling_rate == 16000 assert sampling_rate == 16000
assert return_tensors == "pt" assert return_tensors == "pt"
assert len(wav) > 0 assert len(wav) > 0
@@ -262,9 +265,10 @@ def test_seed_tts_whisper_transcribe_passes_attention_mask(monkeypatch):
attention_mask=FakeTensor("mask") if return_attention_mask else None, attention_mask=FakeTensor("mask") if return_attention_mask else None,
) )
def get_decoder_prompt_ids(self, *, language, task): def get_decoder_prompt_ids(self, *, language, task, no_timestamps=True):
assert language == "english" assert language == "english"
assert task == "transcribe" assert task == "transcribe"
calls["no_timestamps"] = no_timestamps
return [(1, 2)] return [(1, 2)]
def batch_decode(self, predicted_ids, *, skip_special_tokens): def batch_decode(self, predicted_ids, *, skip_special_tokens):
@@ -283,10 +287,18 @@ def test_seed_tts_whisper_transcribe_passes_attention_mask(monkeypatch):
monkeypatch.setattr(seed_tts_eval, "_en_model", FakeModel()) monkeypatch.setattr(seed_tts_eval, "_en_model", FakeModel())
monkeypatch.setattr(seed_tts_eval, "_device", "cuda:1") monkeypatch.setattr(seed_tts_eval, "_device", "cuda:1")
text = seed_tts_eval._transcribe_en_f32_16k(np.ones(1600, dtype=np.float32)) text = seed_tts_eval._transcribe_en_f32_16k(np.ones(int(audio_seconds * 16000), dtype=np.float32))
assert text == "hello" assert text == "hello"
assert calls["return_attention_mask"] is True assert calls["return_attention_mask"] is True
assert calls["input_features"].device == "cuda:1" assert calls["input_features"].device == "cuda:1"
assert calls["generate_kwargs"]["attention_mask"].device == "cuda:1" assert calls["generate_kwargs"]["attention_mask"].device == "cuda:1"
assert calls["generate_kwargs"]["forced_decoder_ids"] == [(1, 2)] assert calls["generate_kwargs"]["forced_decoder_ids"] == [(1, 2)]
if audio_seconds > 30:
assert calls["processor_kwargs"] == {"truncation": False, "padding": "longest"}
assert calls["generate_kwargs"]["return_timestamps"] is True
assert calls["no_timestamps"] is False
else:
assert calls["processor_kwargs"] == {}
assert "return_timestamps" not in calls["generate_kwargs"]
assert calls["no_timestamps"] is True
@@ -166,6 +166,56 @@ def test_moss_tts_seed_falls_back_to_stage_default(adapter_cls, mocker):
assert prepared.tts_params["seed"] == [42] assert prepared.tts_params["seed"] == [42]
@pytest.mark.parametrize(
"variant,ref_text,mode",
[
("local", " Reference. ", "continuation"),
("local", None, "generation"),
("local", " ", "generation"),
("tts", "Reference.", "generation"),
],
)
def test_moss_reference_transcript_mode(variant, ref_text, mode, mocker):
import torch
reference = [torch.ones((3, 12), dtype=torch.int64)]
unified = torch.arange(52, dtype=torch.int64).reshape(1, 4, 13)
processor = mocker.Mock(return_value={"input_ids": unified})
server = mocker.Mock(_moss_variant="local", uploaded_speakers={"speaker": {}})
server._voice_created_at.return_value = 123
request = OpenAICreateSpeechRequest(
input="Target.", ref_text=ref_text, language="English", voice="Speaker", seed=0, max_new_tokens=2048
)
adapter = MossTTSAdapter(SpeechServingContext(server=server))
adapter._moss_variant = variant
mocker.patch.object(adapter, "_get_moss_processor", return_value=processor)
mocker.patch.object(
adapter, "_encode_moss_references", new=mocker.AsyncMock(return_value=(reference, {0: "reference-key"}))
)
prepared = asyncio.run(adapter.build(request, [], has_inline_ref_audio=False))
if mode == "continuation":
processor.build_user_message.assert_called_once_with(text="Reference. Target.", language="English")
processor.build_assistant_message.assert_called_once_with(audio_codes_list=reference)
conversation = [processor.build_user_message.return_value, processor.build_assistant_message.return_value]
else:
processor.build_user_message.assert_called_once_with(text="Target.", language="English", reference=reference)
processor.build_assistant_message.assert_not_called()
conversation = [processor.build_user_message.return_value]
processor.assert_called_once_with(conversations=[conversation], mode=mode)
assert prepared.prompt["prompt_token_ids"] == unified[0, :, 0].tolist()
assert torch.equal(prepared.tts_params["codes"]["ref"], unified[0, :, 1:])
assert prepared.tts_params["max_new_frames"] == [2048]
assert prepared.tts_params["ref_audio_cache_key"] == "reference-key"
assert prepared.tts_params["voice_name"] == ["speaker"]
assert prepared.tts_params["voice_created_at"] == [123]
assert prepared.tts_params["seed"] == [0]
assert prepared.prompt["cache_salt"]
assert request.input == "Target."
def test_qwen3_tts_metadata(): def test_qwen3_tts_metadata():
assert Qwen3TTSAdapter.backend == "ar" assert Qwen3TTSAdapter.backend == "ar"
assert issubclass(Qwen3TTSAdapter, ARTTSAdapter) assert issubclass(Qwen3TTSAdapter, ARTTSAdapter)
+25 -1
View File
@@ -1,3 +1,6 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
from __future__ import annotations from __future__ import annotations
import queue import queue
@@ -97,9 +100,11 @@ class FakeAsyncOmniEngine:
) -> None: ) -> None:
self.model = model self.model = model
self.config_path = None self.config_path = None
self.stage_configs: list[Any] = []
self.stage_metadata = stage_metadata or [THREE_STAGE_META[-1]] self.stage_metadata = stage_metadata or [THREE_STAGE_META[-1]]
self.num_stages = len(self.stage_metadata) self.num_stages = len(self.stage_metadata)
self.stage_configs = [
StageConfig(stage_id=i, model_stage="dummy-model").to_omegaconf() for i in range(self.num_stages)
]
self.default_sampling_params_list = default_sampling_params_list or [ self.default_sampling_params_list = default_sampling_params_list or [
SamplingParams(max_tokens=8) for _ in range(self.num_stages) SamplingParams(max_tokens=8) for _ in range(self.num_stages)
] ]
@@ -222,6 +227,25 @@ def test_resolve_sampling_params_list_preserves_stage_constraints():
assert caller_params.stop_token_ids == [7] assert caller_params.stop_token_ids == [7]
@pytest.mark.parametrize("use_defaults", [False, True])
def test_moss_local_output_policy_preserves_codec_streaming(use_defaults):
from vllm_omni.model_executor.models.moss_tts.pipeline import MOSS_TTS_LOCAL_PIPELINE
base = _make_base()
base.engine.num_stages = 2
base.sampling_constraints_list = [stage.sampling_constraints for stage in MOSS_TTS_LOCAL_PIPELINE.stages]
base.default_sampling_params_list = [
base._apply_sampling_constraints(SamplingParams(), constraints)
for constraints in base.sampling_constraints_list
]
caller = [SamplingParams(output_kind=RequestOutputKind.DELTA) for _ in range(2)]
result = base.resolve_sampling_params_list(None if use_defaults else caller, allow_delta_coercion=True)
assert [params.output_kind for params in result] == [RequestOutputKind.FINAL_ONLY, RequestOutputKind.DELTA]
assert all(params.output_kind == RequestOutputKind.DELTA for params in caller)
assert base.default_sampling_params_list[0].output_kind == RequestOutputKind.FINAL_ONLY
def _stage_spec( def _stage_spec(
stage_id: int, stage_id: int,
*, *,
@@ -1,5 +1,5 @@
# SPDX-License-Identifier: Apache-2.0 # SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM project # SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
"""CPU parity tests for MOSS-TTS Local Depth RoPE caching.""" """CPU parity tests for MOSS-TTS Local Depth RoPE caching."""
import pytest import pytest
@@ -42,3 +42,16 @@ def test_rope_cache_matches_reference_and_reuses_storage() -> None:
# Smaller slices retain the fixed allocation used by CUDA Graph. # Smaller slices retain the fixed allocation used by CUDA Graph.
attn.prepare_rope_cache(6, device, torch.float32) attn.prepare_rope_cache(6, device, torch.float32)
assert attn._rope_cos_cache.data_ptr() == cache_ptr assert attn._rope_cos_cache.data_ptr() == cache_ptr
def test_frame_local_kv_matches_full_prefix_and_reuses_slots() -> None:
model = MossTTSLocalDepthTransformer(_test_config()).eval()
model.h[0].attn.prepare_rope_cache(12, torch.device("cpu"), torch.float32)
cache = (torch.empty(2, 4, 12, 8), torch.empty(2, 4, 12, 8))
with torch.inference_mode():
for _ in range(2):
inputs = torch.randn(2, 12, 32)
for position in range(12):
expected = model._forward_prefix(inputs[:, : position + 1])[:, -1:]
actual = model._forward_prefix(inputs[:, position : position + 1], cache, position)
torch.testing.assert_close(actual, expected, atol=1e-6, rtol=1e-5)
@@ -0,0 +1,58 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
import copy
import pytest
import torch
import torch.nn.functional as F
from vllm_omni.model_executor.models.moss_tts.audio_tokenizer_v2 import (
MossAudioTokenizerMultiheadAttention,
StreamingExecutionContext,
)
from vllm_omni.model_executor.models.moss_tts.streaming_attention import masked_attention
pytestmark = [pytest.mark.core_model, pytest.mark.cuda]
@pytest.mark.parametrize(
"batch,heads,queries,keys", [(2, 20, 1, 125), (4, 20, 15, 125), (1, 12, 480, 400), (4, 12, 32, 250)]
)
def test_masked_attention_matches_sdpa_with_strides_and_invalid_rows(batch, heads, queries, keys):
torch.manual_seed(7)
q = torch.randn(batch, queries, 3, heads, 64, device="cuda", dtype=torch.bfloat16)[:, :, 0].permute(0, 2, 1, 3)
k = torch.randn(batch, heads, keys, 64, device="cuda", dtype=torch.bfloat16)
v = torch.randn_like(k)
positions = torch.arange(keys, device="cuda").roll(17)
query_positions = torch.arange(queries, device="cuda") + keys - queries
mask = (positions[None, :] <= query_positions[:, None])[None, None].expand(batch, 1, -1, -1).clone()
mask[0, :, 0] = False # Cover an all-masked row alongside valid rows.
actual = masked_attention(q, k, v, mask)
expected = F.scaled_dot_product_attention(q, k, v, mask)
assert bool(actual.isfinite().all())
torch.testing.assert_close(actual, expected, atol=0.008, rtol=0.01)
assert actual[0, :, 0].count_nonzero() == 0
def test_streaming_attention_preserves_ring_wrap_reordering_and_slot_reset():
torch.manual_seed(9)
reference = MossAudioTokenizerMultiheadAttention(
128, 2, causal=True, context=9, device="cuda", dtype=torch.bfloat16
)
candidate = copy.deepcopy(reference)
candidate._streaming_attention = masked_attention
for model in (reference, candidate):
model._streaming_state = model._init_streaming_state(4)
for step in range(8):
slots = torch.tensor([1, 0, 3] if step % 2 else [0, 1, 3], device="cuda")
valid = torch.tensor([True, True, False], device="cuda")
context = StreamingExecutionContext(state_slot_ids=slots, valid_rows=valid)
x = torch.randn(3, 3, 128, device="cuda", dtype=torch.bfloat16)
with torch.inference_mode():
expected = reference(x, x, x, execution_context=context)
actual = candidate(x, x, x, execution_context=context)
torch.testing.assert_close(actual, expected, atol=0.01, rtol=0.02)
torch.testing.assert_close(candidate._streaming_state.offset, reference._streaming_state.offset, rtol=0, atol=0)
if step == 4:
for model in (reference, candidate):
model._streaming_state.reset_slots(torch.tensor([0], device="cuda"))
@@ -0,0 +1,89 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
from types import SimpleNamespace
import pytest
import torch
from torch import nn
from vllm_omni.model_executor.models.moss_tts.modeling_moss_tts_codec import (
MossTTSCodecDecoder,
_MossCodecStreamSession,
)
pytestmark = [pytest.mark.core_model, pytest.mark.cpu]
class CausalCodec(nn.Module):
downsample_rate = 1
def __init__(self):
super().__init__()
self.weight = nn.Parameter(torch.zeros(()))
def initialize_decoder_state_pool(self, capacity, scratch):
self.state = torch.zeros(capacity + scratch)
def reset_decoder_state_slots(self, slots):
self.state[slots] = 0
def decode_streaming_batch(self, codes, lengths, slots, valid_rows):
audio = codes.sum(0).float().cumsum(-1) + self.state[slots, None]
self.state[slots] = audio[:, -1]
return SimpleNamespace(audio=audio[:, None], audio_lengths=lengths)
def session():
return _MossCodecStreamSession(CausalCodec(), state_capacity=8, n_vq=2, vllm_config=None)
def test_terminal_padding_preserves_audio_live_state_and_reused_slots():
s = session()
assert [s.acquire(), s.acquire()] == [0, 1]
plan = {0: torch.ones(2, 3), 1: torch.ones(2, 15)}
with pytest.raises(ValueError, match="Only terminal"):
s.step(plan)
assert s._codec.state.count_nonzero() == 0
s.step({0: torch.ones(2, 1), 1: torch.full((2, 1), 2)})
actual = s.step(plan, terminal_slots={0})
torch.testing.assert_close(actual[0], torch.tensor([[4.0, 6.0, 8.0]]))
torch.testing.assert_close(actual[1], torch.arange(6.0, 36.0, 2.0)[None])
s.release(0, state_already_reset=True)
assert s.acquire() == 0
resumed = s.step({0: torch.ones(2, 1), 1: torch.ones(2, 1)})
assert [resumed[i].item() for i in (0, 1)] == [2, 36]
@pytest.mark.parametrize("graph_capacity, expected_calls", [(None, 5), (4, 5), (8, 3)])
def test_only_existing_compatible_graph_buckets_are_coalesced(mocker, graph_capacity, expected_calls):
s = session()
s._cudagraph_wrapper = (
None
if graph_capacity is None
else SimpleNamespace(
batch_sizes=[1, graph_capacity],
_select_frame_size=lambda size, allow_padding: next((b for b in [1, 8, 15] if b >= size), None),
decode=lambda *args, **kwargs: None,
)
)
step = mocker.spy(s, "step")
decoder = object.__new__(MossTTSCodecDecoder)
nn.Module.__init__(decoder)
decoder._stream_max_step_frames = 15
decoder._stream_state_capacity = 8
decoder._stream_req_slots = {}
decoder._ensure_stream_session = lambda: s
lengths = [3, 7, 8, 1, 15]
finished = [True, True, False, True, False]
result = decoder._decode_streaming_batch(
[
(i, str(i), torch.ones(2, length), done)
for i, (length, done) in enumerate(zip(lengths, finished, strict=True))
]
)
assert step.call_count == expected_calls
assert [result[i].shape[-1] for i in range(5)] == lengths
if graph_capacity == 8:
assert list(step.call_args_list[0].args[0]) == [0, 1, 2]
assert step.call_args_list[0].kwargs["terminal_slots"] == {0, 1}
assert list(step.call_args_list[1].args[0]) == [3] # Preserve the exact one-frame graph.
+76
View File
@@ -0,0 +1,76 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
import pytest
import torch
from vllm_omni.model_executor.models.moss_tts.modeling_moss_tts_talker import MossTTSLocalTalkerForGeneration
from vllm_omni.worker.gpu_ar_model_runner import _snapshot_tensor_payload_to_cpu_async
pytestmark = [pytest.mark.core_model, pytest.mark.cuda]
def test_async_packed_snapshot_survives_source_reuse():
values = [
torch.arange(12, device="cuda").view(1, 12),
torch.empty((0, 12), device="cuda", dtype=torch.long),
torch.arange(24, device="cuda").view(2, 12),
]
expected = [v.cpu() for v in values]
snapshot = _snapshot_tensor_payload_to_cpu_async(
{"list": values, "tuple": tuple(values)}, copy_stream=torch.cuda.Stream(), pin_memory=True
)
for v in values:
v.fill_(-99)
snapshot.wait()
torch.testing.assert_close(snapshot.payload, {"list": expected, "tuple": tuple(expected)}, rtol=0, atol=0)
@pytest.mark.parametrize("offsets", [[0, 1], [1, 4]])
def test_mtp_preserves_rows_and_snapshot_without_cuda_sync(monkeypatch, offsets):
import vllm_omni.worker.gpu_model_runner as mod
from tests.worker.test_omni_gpu_model_runner import _make_runner, _noop_forward_context
monkeypatch.setattr(mod.current_omni_platform, "set_forward_context", _noop_forward_context)
runner = _make_runner(req_ids=("r1", "r2"), hidden_size=4)
runner.model.gpu_resident_buffer_keys = {("codes", "audio")}
for name in ("talker_mtp_input_ids", "talker_mtp_inputs_embeds", "last_talker_hidden", "text_step"):
buffer = getattr(runner, name)
buffer.gpu = buffer.gpu.to("cuda")
embeds = torch.ones((2, 4), device="cuda")
codes = torch.arange(2, device="cuda").view(2, 1)
runner.talker_mtp = lambda *args, **kwargs: (embeds, codes)
output = torch.full((5, 4), -10.0, device="cuda")
previous = torch.cuda.get_sync_debug_mode()
try:
torch.cuda.set_sync_debug_mode("error")
runner._talker_mtp_forward(["r1", "r2"], output, offsets)
finally:
torch.cuda.set_sync_debug_mode(previous)
expected = torch.full((5, 4), -10.0)
expected[offsets] = 1.0
torch.testing.assert_close(output.cpu(), expected, rtol=0, atol=0)
codes.fill_(999)
assert runner.model_intermediate_buffer["r1"]["codes"]["audio"].item() == 0
assert runner.model_intermediate_buffer["r2"]["codes"]["audio"].item() == 1
def test_mixed_stop_flags_do_not_synchronize_cuda():
model = object.__new__(MossTTSLocalTalkerForGeneration)
torch.nn.Module.__init__(model)
model.n_vq = 12
model.audio_pad_token_id = 1024
valid = torch.arange(12, device="cuda").view(1, 12)
stopped = torch.full_like(valid, 1024)
infos = [{"audio_codes": {"current": valid}}, {}, {"audio_codes": {"current": stopped}}]
hidden = torch.ones((3, 4), device="cuda")
previous = torch.cuda.get_sync_debug_mode()
try:
torch.cuda.set_sync_debug_mode("error")
output = model.make_omni_output(hidden, runtime_additional_information=infos)
finally:
torch.cuda.set_sync_debug_mode(previous)
assert model._batch_should_continue.tolist() == [True, True, False]
torch.testing.assert_close(output.multimodal_outputs["codes"]["audio"], [valid, valid[:0], stopped])
assert all("current" not in info.get("audio_codes", {}) for info in infos)
prefill = model.make_omni_output(hidden[:1], runtime_additional_information=[{}])
assert model._batch_should_continue is None and prefill.multimodal_outputs == {}
@@ -464,25 +464,33 @@ def _transcribe_en_f32_16k(wav_f32: np.ndarray) -> str:
return "" return ""
with _lock: with _lock:
assert _en_processor is not None and _en_model is not None and _device is not None assert _en_processor is not None and _en_model is not None and _device is not None
# Whisper's default feature extraction truncates at 30 seconds. Keep
# abnormal tails so WER evaluates the complete generated response.
long_audio = len(wav_f32) > 30 * 16000
processor_kwargs = {"truncation": False, "padding": "longest"} if long_audio else {}
try: try:
inputs = _en_processor( inputs = _en_processor(
wav_f32, wav_f32,
sampling_rate=16000, sampling_rate=16000,
return_tensors="pt", return_tensors="pt",
return_attention_mask=True, return_attention_mask=True,
**processor_kwargs,
) )
except TypeError: except TypeError:
inputs = _en_processor(wav_f32, sampling_rate=16000, return_tensors="pt") inputs = _en_processor(wav_f32, sampling_rate=16000, return_tensors="pt", **processor_kwargs)
input_features = inputs.input_features.to(_device) input_features = inputs.input_features.to(_device)
attention_mask = getattr(inputs, "attention_mask", None) attention_mask = getattr(inputs, "attention_mask", None)
if attention_mask is None and isinstance(inputs, dict): if attention_mask is None and isinstance(inputs, dict):
attention_mask = inputs.get("attention_mask") attention_mask = inputs.get("attention_mask")
generate_kwargs: dict[str, Any] = {} generate_kwargs: dict[str, Any] = {}
if long_audio:
generate_kwargs["return_timestamps"] = True
if attention_mask is not None: if attention_mask is not None:
generate_kwargs["attention_mask"] = attention_mask.to(_device) generate_kwargs["attention_mask"] = attention_mask.to(_device)
with torch.no_grad(): with torch.no_grad():
try: try:
forced = _en_processor.get_decoder_prompt_ids(language="english", task="transcribe") prompt_kwargs = {"no_timestamps": False} if long_audio else {}
forced = _en_processor.get_decoder_prompt_ids(language="english", task="transcribe", **prompt_kwargs)
predicted_ids = _en_model.generate(input_features, forced_decoder_ids=forced, **generate_kwargs) predicted_ids = _en_model.generate(input_features, forced_decoder_ids=forced, **generate_kwargs)
except Exception: except Exception:
predicted_ids = _en_model.generate( predicted_ids = _en_model.generate(
@@ -548,14 +556,13 @@ def _expand_seed_tts_turn_outputs(
turn_pcm = getattr(output, "tts_turn_pcm_bytes", None) turn_pcm = getattr(output, "tts_turn_pcm_bytes", None)
session_pcm = getattr(output, "tts_output_pcm_bytes", None) session_pcm = getattr(output, "tts_output_pcm_bytes", None)
for turn_index, turn in enumerate(turns): for turn_index, turn in enumerate(turns):
expanded_requests.append( turn_request = replace(
replace( request,
request, seed_tts_utterance_id=turn.utterance_id,
prompt=turn.target_text, seed_tts_turns=(),
seed_tts_utterance_id=turn.utterance_id,
seed_tts_turns=(),
)
) )
turn_request.prompt = turn.target_text
expanded_requests.append(turn_request)
turn_output = copy(output) turn_output = copy(output)
if isinstance(turn_pcm, list) and turn_index < len(turn_pcm): if isinstance(turn_pcm, list) and turn_index < len(turn_pcm):
turn_output.tts_output_pcm_bytes = turn_pcm[turn_index] turn_output.tts_output_pcm_bytes = turn_pcm[turn_index]
+64
View File
@@ -0,0 +1,64 @@
# MOSS-TTS Local Transformer v1.5: both stages on one H200.
# Select one free physical GPU with CUDA_VISIBLE_DEVICES; devices below are logical.
# Capacity and graph capture require more than 80 GiB of GPU memory.
async_chunk: true
connectors:
shm:
name: SharedMemoryConnector
extra:
connector_get_sleep_s: 0.005
connector_get_max_wait_first_chunk: 1000
connector_get_max_wait: 300
initial_codec_chunk_frames: 1
codec_chunk_frames: 15
codec_left_context_frames: 0
stages:
- stage_id: 0
max_num_seqs: 256
gpu_memory_utilization: 0.60
kv_cache_memory_bytes: 34359738368 # 32 GiB; leave room for codec state and graphs.
enable_prefix_caching: false
async_scheduling: true
max_num_batched_tokens: 65536
max_model_len: 4096
devices: "0"
compilation_config:
# Capture mixed prefill/decode batches by total scheduled tokens.
cudagraph_capture_sizes: [1, 2, 4, 8, 16, 24, 32, 40, 48, 56, 64, 72, 80, 88, 96, 104, 112, 120, 128, 160, 192, 224, 256, 320, 384, 448, 512]
max_cudagraph_capture_size: 512
output_connectors:
to_stage_1: shm
default_sampling_params:
temperature: 1.7
top_p: 0.8
top_k: 25
max_tokens: 4096
seed: 42
repetition_penalty: 1.0
- stage_id: 1
max_num_seqs: 256
gpu_memory_utilization: 0.30
enforce_eager: false
enable_prefix_caching: false
async_scheduling: true
compilation_config:
# Codec execution-batch buckets; frame counts come from the connector.
cudagraph_capture_sizes: [1, 2, 4, 8, 16, 32, 64, 128, 160, 192, 256]
cudagraph_num_of_warmups: 1
max_num_batched_tokens: 65536
max_model_len: 65536
devices: "0"
hf_overrides:
codec_attention_backend: triton
input_connectors:
from_stage_0: shm
default_sampling_params:
temperature: 0.0
top_p: 1.0
top_k: -1
max_tokens: 65536
seed: 42
repetition_penalty: 1.0
+6 -3
View File
@@ -339,6 +339,7 @@ class OmniBase(PDDisaggregationMixin):
allow_delta_coercion: bool = False, allow_delta_coercion: bool = False,
) -> Sequence[Any]: ) -> Sequence[Any]:
"""Resolve request parameters; pipeline sampling constraints override caller values.""" """Resolve request parameters; pipeline sampling constraints override caller values."""
normalized: Sequence[Any]
if sampling_params_list is None: if sampling_params_list is None:
normalized = self.default_sampling_params_list normalized = self.default_sampling_params_list
# Set the output kind to delta since no params were specified # Set the output kind to delta since no params were specified
@@ -354,7 +355,9 @@ class OmniBase(PDDisaggregationMixin):
if len(normalized) != self.num_stages: if len(normalized) != self.num_stages:
raise ValueError(f"Expected {self.num_stages} sampling params, got {len(normalized)}") raise ValueError(f"Expected {self.num_stages} sampling params, got {len(normalized)}")
if sampling_params_list is not None: # Streaming coercion may also change a constrained output kind in
# the defaults (for example, an internal TTS stage's FINAL_ONLY).
if sampling_params_list is not None or allow_delta_coercion:
normalized = [ normalized = [
self._apply_sampling_constraints(params, constraints) self._apply_sampling_constraints(params, constraints)
for params, constraints in zip(normalized, self.sampling_constraints_list, strict=True) for params, constraints in zip(normalized, self.sampling_constraints_list, strict=True)
@@ -364,7 +367,7 @@ class OmniBase(PDDisaggregationMixin):
@staticmethod @staticmethod
def _get_sampling_constraints_list(stage_configs: Sequence[Any]) -> list[dict[str, Any]]: def _get_sampling_constraints_list(stage_configs: Sequence[Any]) -> list[dict[str, Any]]:
"""Extract each stage's required sampling settings from runtime configs.""" """Extract each stage's required sampling settings from runtime configs."""
constraints_list = [] constraints_list: list[dict[str, Any]] = []
for stage_config in stage_configs: for stage_config in stage_configs:
constraints = getattr(stage_config, "sampling_constraints", {}) constraints = getattr(stage_config, "sampling_constraints", {})
if not isinstance(constraints, Mapping): if not isinstance(constraints, Mapping):
@@ -398,7 +401,7 @@ class OmniBase(PDDisaggregationMixin):
req_state = self.request_states.get(request_id) req_state = self.request_states.get(request_id)
prom = getattr(self, "prom_metrics", None) prom = getattr(self, "prom_metrics", None)
metrics = getattr(req_state, "metrics", None) metrics = getattr(req_state, "metrics", None)
if metrics is None or prom is None: if req_state is None or metrics is None or prom is None:
return return
if str(request_id) in metrics.e2e_done or getattr(req_state, "failure_recorded", False): if str(request_id) in metrics.e2e_done or getattr(req_state, "failure_recorded", False):
return return
@@ -6,7 +6,7 @@ Both variants share the same build/validate flow (``_build_moss_tts_params``
handles each); they are registered under distinct model-type names. handles each); they are registered under distinct model-type names.
""" """
from typing import TYPE_CHECKING, Any from typing import TYPE_CHECKING, Any, cast
from vllm.inputs import tokens_input from vllm.inputs import tokens_input
@@ -61,7 +61,7 @@ class _MossTTSAdapterBase(ARTTSAdapter):
# the process-wide speaker cache. # the process-wide speaker cache.
encoder = build_reference_encoder( encoder = build_reference_encoder(
self._get_moss_processor(), self._get_moss_processor(),
variant=self._moss_variant, variant=cast(str, self._moss_variant),
speaker_cache=self._speaker_cache, speaker_cache=self._speaker_cache,
) )
self._moss_ref_encoder = encoder self._moss_ref_encoder = encoder
@@ -101,7 +101,7 @@ class _MossTTSAdapterBase(ARTTSAdapter):
return await encode_request_references( return await encode_request_references(
self._get_moss_ref_encoder(), self._get_moss_ref_encoder(),
request.ref_audio, cast(str, request.ref_audio),
request.ref_audio_2 if two_speaker else None, request.ref_audio_2 if two_speaker else None,
resolve_ref_audio=self._resolve_ref_audio, resolve_ref_audio=self._resolve_ref_audio,
get_artifact_key=self._get_resolved_ref_audio_artifact_key, get_artifact_key=self._get_resolved_ref_audio_artifact_key,
@@ -179,15 +179,16 @@ class _MossTTSAdapterBase(ARTTSAdapter):
v = self._moss_variant v = self._moss_variant
params: dict[str, Any]
# ---- Legacy nano path (unchanged) ---- # ---- Legacy nano path (unchanged) ----
if v is None: # moss_tts_nano if v is None: # moss_tts_nano
params: dict[str, Any] = { params = {
"text": [request.input or ""], "text": [request.input or ""],
"mode": ["voice_clone"], "mode": ["voice_clone"],
} }
if request.max_new_tokens is not None: if request.max_new_tokens is not None:
params["max_new_frames"] = [request.max_new_tokens] params["max_new_frames"] = [request.max_new_tokens]
wav_list, sr, cache_key = await self._resolve_ref_audio(request.ref_audio) wav_list, sr, cache_key = await self._resolve_ref_audio(cast(str, request.ref_audio))
params["prompt_audio_array"] = [[wav_list, sr]] params["prompt_audio_array"] = [[wav_list, sr]]
params["ref_audio_cache_key"] = cache_key params["ref_audio_cache_key"] = cache_key
return params return params
@@ -201,13 +202,13 @@ class _MossTTSAdapterBase(ARTTSAdapter):
# talker reads for short prompts; full Realtime support needs a # talker reads for short prompts; full Realtime support needs a
# separate processor.from_module path which we don't wire here. # separate processor.from_module path which we don't wire here.
if v == "realtime": if v == "realtime":
params: dict[str, Any] = { params = {
"text": [request.input or ""], "text": [request.input or ""],
"mode": ["voice_clone"], "mode": ["voice_clone"],
} }
if request.max_new_tokens is not None: if request.max_new_tokens is not None:
params["max_new_frames"] = [request.max_new_tokens] params["max_new_frames"] = [request.max_new_tokens]
wav_list, sr, cache_key = await self._resolve_ref_audio(request.ref_audio) wav_list, sr, cache_key = await self._resolve_ref_audio(cast(str, request.ref_audio))
params["prompt_audio_array"] = [[wav_list, sr]] params["prompt_audio_array"] = [[wav_list, sr]]
params["ref_audio_cache_key"] = cache_key params["ref_audio_cache_key"] = cache_key
return params return params
@@ -254,13 +255,20 @@ class _MossTTSAdapterBase(ARTTSAdapter):
# Build the unified-codes prompt: (L, 1+n_vq) where col 0 is text/special # Build the unified-codes prompt: (L, 1+n_vq) where col 0 is text/special
# tokens and cols 1..n_vq are the delay-pattern audio code grid (mostly # tokens and cols 1..n_vq are the delay-pattern audio code grid (mostly
# audio_pad_code outside the reference block). # audio_pad_code outside the reference block).
user_msg = proc.build_user_message(**user_kwargs) if v == "local" and request.ref_text and request.ref_text.strip():
batch = proc(conversations=[[user_msg]], mode="generation") reference = user_kwargs.pop("reference")
user_kwargs["text"] = request.ref_text.strip() + " " + (request.input or "")
user_msg = proc.build_user_message(**user_kwargs)
assistant_msg = proc.build_assistant_message(audio_codes_list=reference)
batch = proc(conversations=[[user_msg, assistant_msg]], mode="continuation")
else:
user_msg = proc.build_user_message(**user_kwargs)
batch = proc(conversations=[[user_msg]], mode="generation")
unified = batch["input_ids"][0] # torch.LongTensor (L, 1+n_vq) unified = batch["input_ids"][0] # torch.LongTensor (L, 1+n_vq)
text_ids: list[int] = unified[:, 0].tolist() text_ids: list[int] = unified[:, 0].tolist()
audio_codes: torch.Tensor = unified[:, 1:].contiguous().to(torch.int64) audio_codes: torch.Tensor = unified[:, 1:].contiguous().to(torch.int64)
params: dict[str, Any] = { params = {
"prompt_token_ids": text_ids, "prompt_token_ids": text_ids,
"codes": {"ref": audio_codes}, "codes": {"ref": audio_codes},
} }
+2 -2
View File
@@ -278,9 +278,9 @@ def maybe_coerce_to_message_type(params: SamplingParams, is_streaming: bool):
if params.output_kind == target_type: if params.output_kind == target_type:
return params return params
elif is_streaming and params.output_kind == RequestOutputKind.FINAL_ONLY: elif is_streaming and params.output_kind == RequestOutputKind.FINAL_ONLY:
logger.warning("Request appears to be streaming, but got request type final only!") logger.debug("Coercing FINAL_ONLY output to DELTA for streaming")
elif not is_streaming and params.output_kind == RequestOutputKind.DELTA: elif not is_streaming and params.output_kind == RequestOutputKind.DELTA:
logger.warning("Request appears to not be streaming, but got request type delta!") logger.debug("Coercing DELTA output to FINAL_ONLY for non-streaming")
if not params.skip_clone: if not params.skip_clone:
params = params.clone() params = params.clone()
+7 -8
View File
@@ -48,7 +48,7 @@ class StageRequestStats:
finish_reason: str | None = None finish_reason: str | None = None
request_id: str | None = None request_id: str | None = None
postprocess_time_ms: float = 0.0 postprocess_time_ms: float = 0.0
diffusion_metrics: dict[str, float] = None diffusion_metrics: dict[str, float] | None = None
audio_generated_frames: int = 0 audio_generated_frames: int = 0
audio_sample_rate: int = 0 audio_sample_rate: int = 0
audio_duration_s: float = 0.0 audio_duration_s: float = 0.0
@@ -181,11 +181,11 @@ class OrchestratorAggregator:
self.e2e_total_ms = 0.0 self.e2e_total_ms = 0.0
self.e2e_total_tokens = 0 self.e2e_total_tokens = 0
self.e2e_count = 0 self.e2e_count = 0
self.e2e_done = set() self.e2e_done: set[str] = set()
self.wall_start_ts = float(wall_start_ts) self.wall_start_ts = float(wall_start_ts)
self.last_finish_ts = float(wall_start_ts) self.last_finish_ts = float(wall_start_ts)
self.stage_first_ts = [None for _ in range(self.num_stages)] self.stage_first_ts: list[float | None] = [None for _ in range(self.num_stages)]
self.stage_last_ts = [None for _ in range(self.num_stages)] self.stage_last_ts: list[float | None] = [None for _ in range(self.num_stages)]
self.accumulated_gen_time_ms: defaultdict[str, defaultdict[int, float]] = defaultdict( self.accumulated_gen_time_ms: defaultdict[str, defaultdict[int, float]] = defaultdict(
lambda: defaultdict(float) lambda: defaultdict(float)
) # {request_id: {stage_id:accumulated_gen_time_ms}} ) # {request_id: {stage_id:accumulated_gen_time_ms}}
@@ -618,6 +618,7 @@ class OrchestratorAggregator:
final_output_type: str | None = None, final_output_type: str | None = None,
) -> None: ) -> None:
stats = self._as_stage_request_stats(stage_id, req_id, metrics, final_output_type) stats = self._as_stage_request_stats(stage_id, req_id, metrics, final_output_type)
assert stats.stage_id is not None
self.stage_total_tokens[stats.stage_id] += int(stats.num_tokens_out) self.stage_total_tokens[stats.stage_id] += int(stats.num_tokens_out)
if stats.stage_id == 0: if stats.stage_id == 0:
self.stage_total_tokens[stats.stage_id] += int(stats.num_tokens_in) self.stage_total_tokens[stats.stage_id] += int(stats.num_tokens_in)
@@ -774,10 +775,8 @@ class OrchestratorAggregator:
final_stage_id_map = self.final_stage_id_for_e2e final_stage_id_map = self.final_stage_id_for_e2e
stage_wall_time_ms = [ stage_wall_time_ms = [
((self.stage_last_ts[i] - self.stage_first_ts[i]) * 1000.0) ((last - first) * 1000.0) if first is not None and last is not None else 0.0
if (self.stage_first_ts[i] is not None and self.stage_last_ts[i] is not None) for first, last in zip(self.stage_first_ts, self.stage_last_ts, strict=True)
else 0.0
for i in range(self.num_stages)
] ]
overall_summary = { overall_summary = {
@@ -1,3 +1,6 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
# Copyright 2026 OpenMOSS and the HuggingFace Inc. team. All rights reserved. # Copyright 2026 OpenMOSS and the HuggingFace Inc. team. All rights reserved.
# #
# Licensed under the Apache License, Version 2.0 (the "License"); # Licensed under the Apache License, Version 2.0 (the "License");
@@ -801,7 +804,16 @@ class MossAudioTokenizerMultiheadAttention(StreamingModule):
else: else:
attn_bias = None attn_bias = None
x = F.scaled_dot_product_attention(q, k, v, attn_bias, dropout_p=0.0) streaming_attention = getattr(self, "_streaming_attention", None)
if (
streaming_attention is not None
and attn_bias is not None
and q.dtype == torch.bfloat16
and q.shape[-1] == 64
):
x = streaming_attention(q, k, v, attn_bias)
else:
x = F.scaled_dot_product_attention(q, k, v, attn_bias, dropout_p=0.0)
x = x.transpose(1, 2).reshape(B, T, self.embed_dim) x = x.transpose(1, 2).reshape(B, T, self.embed_dim)
x = apply_weights_per_step(self.out_projs, self.weights_per_step_schedule, x, offset_cpu) x = apply_weights_per_step(self.out_projs, self.weights_per_step_schedule, x, offset_cpu)
@@ -145,15 +145,22 @@ class _MossCodecStreamSession:
terminal_slots = terminal_slots or set() terminal_slots = terminal_slots or set()
if not terminal_slots.issubset(slot_codes): if not terminal_slots.issubset(slot_codes):
raise ValueError("terminal_slots must be a subset of the decode batch slots.") raise ValueError("terminal_slots must be a subset of the decode batch slots.")
step_lengths = {int(codes.shape[1]) for codes in slot_codes.values()} step_t = max(int(codes.shape[1]) for codes in slot_codes.values())
if len(step_lengths) != 1: if any(int(codes.shape[1]) != step_t and slot not in terminal_slots for slot, codes in slot_codes.items()):
raise ValueError(f"MOSS codec streaming step needs uniform T, got {sorted(step_lengths)}") raise ValueError("Only terminal codec rows may be padded to a larger step length")
(step_t,) = step_lengths
slots = list(slot_codes) slots = list(slot_codes)
if any(slot not in self._leased_slots for slot in slots): if any(slot not in self._leased_slots for slot in slots):
raise RuntimeError(f"Streaming decode references an unleased state slot: {slots}") raise RuntimeError(f"Streaming decode references an unleased state slot: {slots}")
codes_step = torch.stack( codes_step = torch.stack(
[slot_codes[slot].to(device=self._device, dtype=torch.long) for slot in slots], [
torch.nn.functional.pad(
slot_codes[slot].to(device=self._device, dtype=torch.long),
(0, step_t - int(slot_codes[slot].shape[1])),
)
if int(slot_codes[slot].shape[1]) != step_t
else slot_codes[slot].to(device=self._device, dtype=torch.long)
for slot in slots
],
dim=1, dim=1,
) )
state_slot_ids = torch.tensor(slots, device=self._device, dtype=torch.long) state_slot_ids = torch.tensor(slots, device=self._device, dtype=torch.long)
@@ -197,9 +204,9 @@ class _MossCodecStreamSession:
self._reset_slot_ids(terminal_slot_ids) self._reset_slot_ids(terminal_slot_ids)
audio = audio_tensor.detach().to("cpu", torch.float32) audio = audio_tensor.detach().to("cpu", torch.float32)
audio_length = step_t * self._samples_per_frame
out: dict[int, torch.Tensor] = {} out: dict[int, torch.Tensor] = {}
for row, slot in enumerate(slots): for row, slot in enumerate(slots):
audio_length = int(slot_codes[slot].shape[1]) * self._samples_per_frame
out[slot] = audio[row, ..., :audio_length].contiguous() out[slot] = audio[row, ..., :audio_length].contiguous()
return out return out
@@ -518,6 +525,8 @@ class MossTTSCodecDecoder(nn.Module):
outputs: dict[int, torch.Tensor] = {} outputs: dict[int, torch.Tensor] = {}
grouped: dict[int, list[tuple[int, str, int, torch.Tensor, bool]]] = {} grouped: dict[int, list[tuple[int, str, int, torch.Tensor, bool]]] = {}
max_step_frames = max(1, int(self._stream_max_step_frames)) max_step_frames = max(1, int(self._stream_max_step_frames))
graph = session._cudagraph_wrapper
coalesce_tails = graph is not None and max(graph.batch_sizes, default=0) >= self._stream_state_capacity
for output_index, request_id, codes_nq_t, finished in items: for output_index, request_id, codes_nq_t, finished in items:
slot = self._stream_req_slots.get(request_id) slot = self._stream_req_slots.get(request_id)
@@ -548,9 +557,13 @@ class MossTTSCodecDecoder(nn.Module):
) )
continue continue
grouped.setdefault(int(codes_nq_t.shape[1]), []).append( frame_count = int(codes_nq_t.shape[1])
(output_index, request_id, slot, codes_nq_t, finished) # Only combine tails that already used the same padded graph.
) # Preserve exact-size graphs (including T=1) and the eager path.
group_frames = frame_count
if coalesce_tails and finished:
group_frames = graph._select_frame_size(frame_count, allow_padding=True) or frame_count
grouped.setdefault(group_frames, []).append((output_index, request_id, slot, codes_nq_t, finished))
for group in grouped.values(): for group in grouped.values():
plan = {slot: codes_nq_t for _, _, slot, codes_nq_t, _ in group} plan = {slot: codes_nq_t for _, _, slot, codes_nq_t, _ in group}
@@ -756,6 +769,17 @@ class MossTTSCodecDecoder(nn.Module):
codec.eval() codec.eval()
if device.type != "cpu": if device.type != "cpu":
codec.decoder.to(dtype=torch.bfloat16) codec.decoder.to(dtype=torch.bfloat16)
attention_backend = getattr(self.vllm_config.model_config.hf_config, "codec_attention_backend", "sdpa")
if attention_backend != "sdpa":
if attention_backend != "triton" or device.type != "cuda":
raise ValueError(f"Unsupported codec attention backend/device: {attention_backend}/{device.type}")
from vllm_omni.model_executor.models.moss_tts.audio_tokenizer_v2 import MossAudioTokenizerMultiheadAttention
from vllm_omni.model_executor.models.moss_tts.streaming_attention import masked_attention
for module in codec.decoder.modules():
if isinstance(module, MossAudioTokenizerMultiheadAttention):
module._streaming_attention = masked_attention
logger.info("Enabled Triton masked attention for the streaming codec decoder")
build_decode_lut = getattr(codec.quantizer, "build_decode_lut", None) build_decode_lut = getattr(codec.quantizer, "build_decode_lut", None)
if callable(build_decode_lut): if callable(build_decode_lut):
lut_dtype = torch.bfloat16 if device.type == "cuda" else torch.float32 lut_dtype = torch.bfloat16 if device.type == "cuda" else torch.float32
@@ -1,3 +1,5 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
# Copyright 2026 OpenMOSS and the vLLM-Omni team. All rights reserved. # Copyright 2026 OpenMOSS and the vLLM-Omni team. All rights reserved.
# #
# Licensed under the Apache License, Version 2.0 (the "License"). # Licensed under the Apache License, Version 2.0 (the "License").
@@ -101,8 +103,13 @@ class _MossTTSLocalAttention(nn.Module):
odd = x[..., 1::2] odd = x[..., 1::2]
return torch.stack((-odd, even), dim=-1).reshape_as(x) return torch.stack((-odd, even), dim=-1).reshape_as(x)
def forward(self, hidden_states: torch.Tensor) -> torch.Tensor: def forward(
"""``hidden_states``: ``(B, S, H)``. Re-prefills with a fresh causal mask over ``[0, S)`` every call.""" self,
hidden_states: torch.Tensor,
kv_cache: tuple[torch.Tensor, torch.Tensor] | None = None,
position: int = 0,
) -> torch.Tensor:
"""Run a causal prefix, or one new token using frame-local K/V."""
batch_size, seq_len, _ = hidden_states.shape batch_size, seq_len, _ = hidden_states.shape
qkv = self.c_attn(hidden_states) qkv = self.c_attn(hidden_states)
query, key, value = qkv.split(self.embed_dim, dim=-1) query, key, value = qkv.split(self.embed_dim, dim=-1)
@@ -110,14 +117,22 @@ class _MossTTSLocalAttention(nn.Module):
key = key.view(batch_size, seq_len, self.n_head, self.head_dim) key = key.view(batch_size, seq_len, self.n_head, self.head_dim)
value = value.view(batch_size, seq_len, self.n_head, self.head_dim) value = value.view(batch_size, seq_len, self.n_head, self.head_dim)
cos, sin = self._rope_cos_sin(seq_len, hidden_states.device, hidden_states.dtype) cos, sin = self._rope_cos_sin(position + seq_len, hidden_states.device, hidden_states.dtype)
cos, sin = cos[:, position:], sin[:, position:]
query = query * cos + self._rotate_half(query) * sin query = query * cos + self._rotate_half(query) * sin
key = key * cos + self._rotate_half(key) * sin key = key * cos + self._rotate_half(key) * sin
query = query.transpose(1, 2) query = query.transpose(1, 2)
key = key.transpose(1, 2) key = key.transpose(1, 2)
value = value.transpose(1, 2) value = value.transpose(1, 2)
attn_output = F.scaled_dot_product_attention(query, key, value, is_causal=True) if kv_cache is not None:
assert seq_len == 1, "Cached Local Depth execution consumes one token at a time"
k_cache, v_cache = kv_cache
k_cache[:, :, position : position + 1].copy_(key)
v_cache[:, :, position : position + 1].copy_(value)
key = k_cache[:, :, : position + 1]
value = v_cache[:, :, : position + 1]
attn_output = F.scaled_dot_product_attention(query, key, value, is_causal=kv_cache is None)
attn_output = attn_output.transpose(1, 2).reshape(batch_size, seq_len, self.embed_dim) attn_output = attn_output.transpose(1, 2).reshape(batch_size, seq_len, self.embed_dim)
return self.c_proj(attn_output) return self.c_proj(attn_output)
@@ -141,8 +156,13 @@ class _MossTTSLocalBlock(nn.Module):
self.ln_2 = nn.LayerNorm(hidden_size, eps=eps) self.ln_2 = nn.LayerNorm(hidden_size, eps=eps)
self.mlp = _MossTTSLocalMLP(hidden_size, inner_size) self.mlp = _MossTTSLocalMLP(hidden_size, inner_size)
def forward(self, hidden_states: torch.Tensor) -> torch.Tensor: def forward(
hidden_states = hidden_states + self.attn(self.ln_1(hidden_states)) self,
hidden_states: torch.Tensor,
kv_cache: tuple[torch.Tensor, torch.Tensor] | None = None,
position: int = 0,
) -> torch.Tensor:
hidden_states = hidden_states + self.attn(self.ln_1(hidden_states), kv_cache, position)
hidden_states = hidden_states + self.mlp(self.ln_2(hidden_states)) hidden_states = hidden_states + self.mlp(self.ln_2(hidden_states))
return hidden_states return hidden_states
@@ -156,11 +176,9 @@ class MossTTSLocalDepthTransformer(nn.Module):
and codebook-0's head (``audio_lm_heads[0]``) simultaneously. and codebook-0's head (``audio_lm_heads[0]``) simultaneously.
- codebooks 1..n_vq-1 are sampled sequentially: each sampled code is - codebooks 1..n_vq-1 are sampled sequentially: each sampled code is
re-embedded (``audio_embeddings[c]``) and appended as the next re-embedded (``audio_embeddings[c]``) and appended as the next
position, re-prefilling the block over the growing (<=n_vq) sequence position. K/V for earlier positions are reused within the frame, so
with a fresh causal mask each call -- mathematically identical to each position's projections and MLP run once. A new frame overwrites
incremental KV-cache decoding since attention is strictly causal and every cache position before reading it, including CUDA Graph replay.
only the last position is ever read, but avoids any cache plumbing
given the trivially short sequence length.
""" """
def __init__(self, gpt2_config, hidden_size: int | None = None) -> None: def __init__(self, gpt2_config, hidden_size: int | None = None) -> None:
@@ -174,11 +192,13 @@ class MossTTSLocalDepthTransformer(nn.Module):
self.ln_f = nn.LayerNorm(self.hidden_size, eps=eps) self.ln_f = nn.LayerNorm(self.hidden_size, eps=eps)
self._compiled_forward_prefix = None self._compiled_forward_prefix = None
def _forward_prefix(self, seq_embeds: torch.Tensor) -> torch.Tensor: def _forward_prefix(
hidden_states = seq_embeds self,
for block in self.h: seq_embeds: torch.Tensor,
hidden_states = block(hidden_states) kv_cache: tuple[torch.Tensor, torch.Tensor] | None = None,
return self.ln_f(hidden_states) position: int = 0,
) -> torch.Tensor:
return self.ln_f(self.h[0](seq_embeds, kv_cache, position))
def setup_compile(self) -> None: def setup_compile(self) -> None:
if self._compiled_forward_prefix is not None: if self._compiled_forward_prefix is not None:
@@ -189,14 +209,18 @@ class MossTTSLocalDepthTransformer(nn.Module):
return return
self._compiled_forward_prefix = torch.compile( self._compiled_forward_prefix = torch.compile(
self._forward_prefix, self._forward_prefix,
dynamic=False, # Share batch/position shapes instead of exhausting Dynamo's
# recompilation budget while warming request-size graph buckets.
dynamic=True,
options={"epilogue_fusion": False}, options={"epilogue_fusion": False},
) )
logger.info("MOSS-TTS local depth prefix enabled with torch.compile") logger.info("MOSS-TTS local depth frame-local KV execution enabled with torch.compile")
def _run_prefix(self, seq_embeds: torch.Tensor) -> torch.Tensor: def _run_prefix(
self, seq_embeds: torch.Tensor, kv_cache: tuple[torch.Tensor, torch.Tensor], position: int
) -> torch.Tensor:
forward_prefix = self._compiled_forward_prefix or self._forward_prefix forward_prefix = self._compiled_forward_prefix or self._forward_prefix
return forward_prefix(seq_embeds) return forward_prefix(seq_embeds, kv_cache, position)
@torch.no_grad() @torch.no_grad()
def generate_frame( def generate_frame(
@@ -240,9 +264,13 @@ class MossTTSLocalDepthTransformer(nn.Module):
for block in self.h: for block in self.h:
block.attn.prepare_rope_cache(n_vq, backbone_last_hidden.device, dtype) block.attn.prepare_rope_cache(n_vq, backbone_last_hidden.device, dtype)
embeds = backbone_last_hidden.new_zeros((batch_size, n_vq, self.hidden_size), dtype=dtype) attn = self.h[0].attn
embeds[:, 0, :] = backbone_last_hidden.to(dtype) cache_shape = (batch_size, attn.n_head, n_vq, attn.head_dim)
hidden = self._run_prefix(embeds[:, :1, :]) kv_cache = (
backbone_last_hidden.new_empty(cache_shape, dtype=dtype),
backbone_last_hidden.new_empty(cache_shape, dtype=dtype),
)
hidden = self._run_prefix(backbone_last_hidden[:, None, :].to(dtype), kv_cache, 0)
local_hidden = hidden[:, 0, :] local_hidden = hidden[:, 0, :]
binary_logits = local_text_lm_head(local_hidden).float() binary_logits = local_text_lm_head(local_hidden).float()
@@ -292,9 +320,9 @@ class MossTTSLocalDepthTransformer(nn.Module):
codes[:, channel_index] = channel_token codes[:, channel_index] = channel_token
if channel_index + 1 < n_vq: if channel_index + 1 < n_vq:
embeds[:, channel_index + 1, :] = audio_embeddings[channel_index](channel_token).to(dtype) embeds = audio_embeddings[channel_index](channel_token).to(dtype)[:, None, :]
hidden = self._run_prefix(embeds[:, : channel_index + 2, :]) hidden = self._run_prefix(embeds, kv_cache, channel_index + 1)
local_hidden = hidden[:, channel_index + 1, :] local_hidden = hidden[:, 0, :]
return should_continue, codes return should_continue, codes
@@ -1,3 +1,6 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
# Copyright 2026 OpenMOSS and the vLLM-Omni team. All rights reserved. # Copyright 2026 OpenMOSS and the vLLM-Omni team. All rights reserved.
# #
# Licensed under the Apache License, Version 2.0 (the "License"). # Licensed under the Apache License, Version 2.0 (the "License").
@@ -1671,20 +1674,19 @@ class MossTTSLocalTalkerForGeneration(nn.Module):
self._batch_state = [(info["audio_state"] if isinstance(info, dict) else {}) for info in info_dicts] self._batch_state = [(info["audio_state"] if isinstance(info, dict) else {}) for info in info_dicts]
self._batch_state_spans = kwargs.get("request_token_spans") self._batch_state_spans = kwargs.get("request_token_spans")
should_values: list[torch.Tensor] = [] should_codes: list[torch.Tensor] = []
have_mtp_output = False should_indices: list[int] = []
per_req_codes: list[torch.Tensor] = [] per_req_codes: list[torch.Tensor] = []
have_codes = False have_codes = False
for info in info_dicts: for index, info in enumerate(info_dicts):
audio_codes = (info.get("audio_codes", {}) or {}) if isinstance(info, dict) else {} audio_codes = (info.get("audio_codes", {}) or {}) if isinstance(info, dict) else {}
current = audio_codes.pop("current", None) current = audio_codes.pop("current", None)
if isinstance(current, torch.Tensor) and current.numel() > 0: if isinstance(current, torch.Tensor) and current.numel() > 0:
codes = current.to(device=hidden.device, dtype=torch.long) codes = current.to(device=hidden.device, dtype=torch.long)
if codes.dim() == 1: if codes.dim() == 1:
codes = codes.unsqueeze(0) codes = codes.unsqueeze(0)
should_continue = codes.ne(self.audio_pad_token_id).any(dim=-1) should_codes.append(codes[:1])
should_values.append(should_continue.reshape(-1)[:1]) should_indices.append(index)
have_mtp_output = True
# Do not boolean-index a CUDA tensor here. Even for the # Do not boolean-index a CUDA tensor here. Even for the
# common single-row case, ``codes[should_continue]`` must # common single-row case, ``codes[should_continue]`` must
# discover a dynamic output size on the host and therefore # discover a dynamic output size on the host and therefore
@@ -1694,9 +1696,19 @@ class MossTTSLocalTalkerForGeneration(nn.Module):
per_req_codes.append(codes) per_req_codes.append(codes)
have_codes = True have_codes = True
else: else:
should_values.append(torch.ones((1,), device=hidden.device, dtype=torch.bool))
per_req_codes.append(hidden.new_empty((0, self.n_vq), dtype=torch.long)) per_req_codes.append(hidden.new_empty((0, self.n_vq), dtype=torch.long))
self._batch_should_continue = torch.cat(should_values, dim=0) if have_mtp_output and should_values else None self._batch_should_continue = None
if should_codes:
active = torch.cat(should_codes, dim=0).ne(self.audio_pad_token_id).any(dim=-1)
if len(should_codes) == len(info_dicts):
self._batch_should_continue = active
else:
self._batch_should_continue = torch.ones(len(info_dicts), device=hidden.device, dtype=torch.bool)
# Keep mixed prefill/decode output packing asynchronous.
indices = torch.tensor(should_indices, device="cpu", dtype=torch.long).to(
hidden.device, non_blocking=True
)
self._batch_should_continue.index_copy_(0, indices, active)
if not have_codes: if not have_codes:
return OmniOutput(text_hidden_states=hidden, multimodal_outputs={}) return OmniOutput(text_hidden_states=hidden, multimodal_outputs={})
@@ -1,8 +1,13 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
# Copyright 2026 OpenMOSS and the vLLM-Omni team. All rights reserved. # Copyright 2026 OpenMOSS and the vLLM-Omni team. All rights reserved.
# #
# Licensed under the Apache License, Version 2.0 (the "License"). # Licensed under the Apache License, Version 2.0 (the "License").
"""Pipeline topology for all MOSS-TTS variants (2-stage: talker → codec).""" """Pipeline topology for all MOSS-TTS variants (2-stage: talker → codec)."""
from vllm.sampling_params import RequestOutputKind
from vllm_omni.config.stage_config import ( from vllm_omni.config.stage_config import (
PipelineConfig, PipelineConfig,
StageExecutionType, StageExecutionType,
@@ -100,6 +105,10 @@ MOSS_TTS_LOCAL_PIPELINE = PipelineConfig(
sampling_constraints={ sampling_constraints={
"detokenize": False, "detokenize": False,
"stop_token_ids": [151645], "stop_token_ids": [151645],
# The worker connector streams codes directly to the codec.
# This internal stage only needs to publish its terminal
# result; codec audio output remains incremental.
"output_kind": RequestOutputKind.FINAL_ONLY,
}, },
), ),
StagePipelineConfig( StagePipelineConfig(
@@ -0,0 +1,119 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
"""Masked streaming attention preserving ring-cache position semantics."""
import torch
from vllm.triton_utils import tl, triton
@triton.jit
def _attention(
q_ptr,
k_ptr,
v_ptr,
mask_ptr,
out_ptr,
qs0: tl.constexpr,
qs1: tl.constexpr,
qs2: tl.constexpr,
qs3: tl.constexpr,
ks0: tl.constexpr,
ks1: tl.constexpr,
ks2: tl.constexpr,
ks3: tl.constexpr,
vs0: tl.constexpr,
vs1: tl.constexpr,
vs2: tl.constexpr,
vs3: tl.constexpr,
ms0: tl.constexpr,
ms2: tl.constexpr,
ms3: tl.constexpr,
num_heads: tl.constexpr,
q_len: tl.constexpr,
kv_len: tl.constexpr,
head_dim: tl.constexpr,
block_m: tl.constexpr,
block_n: tl.constexpr,
):
rows = tl.program_id(0) * block_m + tl.arange(0, block_m)
bh = tl.program_id(1)
batch, head = bh // num_heads, bh % num_heads
dims = tl.arange(0, head_dim)
q = tl.load(
q_ptr + batch * qs0 + head * qs1 + rows[:, None] * qs2 + dims[None, :] * qs3,
mask=rows[:, None] < q_len,
other=0,
)
acc = tl.full((block_m, head_dim), 0.0, tl.float32)
maximum = tl.full((block_m,), -float("inf"), tl.float32)
denominator = tl.full((block_m,), 0.0, tl.float32)
for first in range(tl.cdiv(kv_len, block_n)):
cols = first * block_n + tl.arange(0, block_n)
k = tl.load(
k_ptr + batch * ks0 + head * ks1 + cols[None, :] * ks2 + dims[:, None] * ks3,
mask=cols[None, :] < kv_len,
other=0,
)
allowed = tl.load(
mask_ptr + batch * ms0 + rows[:, None] * ms2 + cols[None, :] * ms3,
mask=(rows[:, None] < q_len) & (cols[None, :] < kv_len),
other=0,
)
score = tl.dot(q, k).to(tl.float32) * (head_dim**-0.5)
score = tl.where(allowed & (cols[None, :] < kv_len), score, -float("inf"))
next_max = tl.maximum(maximum, tl.max(score, axis=1))
safe_max = tl.where(next_max == -float("inf"), 0.0, next_max)
rescale = tl.exp(maximum - safe_max)
p = tl.exp(score - safe_max[:, None])
denominator = denominator * rescale + tl.sum(p, axis=1)
v = tl.load(
v_ptr + batch * vs0 + head * vs1 + cols[:, None] * vs2 + dims[None, :] * vs3,
mask=cols[:, None] < kv_len,
other=0,
)
acc = acc * rescale[:, None] + tl.dot(p.to(v.dtype), v)
maximum = next_max
out = acc / tl.where(denominator > 0, denominator, 1.0)[:, None]
tl.store(
out_ptr + ((batch * num_heads + head) * q_len + rows[:, None]) * head_dim + dims[None, :],
out,
mask=rows[:, None] < q_len,
)
@torch.library.custom_op("moss_streaming::masked_attention", mutates_args=())
def masked_attention(q: torch.Tensor, k: torch.Tensor, v: torch.Tensor, mask: torch.Tensor) -> torch.Tensor:
b, h, t, d = q.shape
out = torch.empty(q.shape, device=q.device, dtype=q.dtype)
# Small query tiles expose parallelism for short streaming chunks. Larger
# tiles reuse K/V for steady chunks. Keep 64-key softmax blocks to retain
# the same accumulation granularity across both paths.
bm = 16 if t <= 32 else 64
num_warps = 8 if k.shape[2] <= 256 else 4
_attention[(triton.cdiv(t, bm), b * h)](
q,
k,
v,
mask,
out,
*q.stride(),
*k.stride(),
*v.stride(),
mask.stride(0),
mask.stride(2),
mask.stride(3),
h,
t,
k.shape[2],
d,
bm,
64,
num_warps=num_warps,
num_stages=2,
)
return out
@masked_attention.register_fake
def _(q: torch.Tensor, k: torch.Tensor, v: torch.Tensor, mask: torch.Tensor) -> torch.Tensor:
return torch.empty(q.shape, device=q.device, dtype=q.dtype)
+4 -1
View File
@@ -1,3 +1,6 @@
# SPDX-License-Identifier: Apache-2.0
# SPDX-FileCopyrightText: Copyright contributors to the vLLM-Omni project
from dataclasses import dataclass, field from dataclasses import dataclass, field
from typing import Any from typing import Any
@@ -52,7 +55,7 @@ class OmniModelRunnerOutput(ModelRunnerOutput):
""" """
multimodal_outputs: list[dict[str, object]] | None = None multimodal_outputs: list[dict[str, object]] | None = None
inter_stage_outputs: list[dict[str, object]] | None = None inter_stage_outputs: list[dict[str, Any] | None] | None = None
# IDs of requests whose KV cache has been extracted from GPU/NPU to CPU. # IDs of requests whose KV cache has been extracted from GPU/NPU to CPU.
# The Scheduler can safely free the block tables for these requests. # The Scheduler can safely free the block tables for these requests.
kv_extracted_req_ids: list[str] | None = None kv_extracted_req_ids: list[str] | None = None
+50 -11
View File
@@ -15,7 +15,7 @@ from collections.abc import Callable, Sequence
from contextlib import nullcontext from contextlib import nullcontext
from copy import copy from copy import copy
from dataclasses import replace from dataclasses import replace
from typing import Any, NamedTuple from typing import Any, NamedTuple, cast
import numpy as np import numpy as np
import torch import torch
@@ -87,14 +87,39 @@ def _clone_cuda_tensor_payload(value: Any, sources: list[torch.Tensor]) -> Any:
return value.detach().clone() return value.detach().clone()
if isinstance(value, dict): if isinstance(value, dict):
return {k: _clone_cuda_tensor_payload(v, sources) for k, v in value.items()} return {k: _clone_cuda_tensor_payload(v, sources) for k, v in value.items()}
if isinstance(value, list): if isinstance(value, (list, tuple)):
return [_clone_cuda_tensor_payload(v, sources) for v in value] if (
if isinstance(value, tuple): value
return tuple(_clone_cuda_tensor_payload(v, sources) for v in value) and all(
isinstance(v, torch.Tensor)
and v.is_cuda
and v.dtype == value[0].dtype
and v.device == value[0].device
and v.layout == torch.strided
and v.is_contiguous()
for v in value
)
# Keep frame/state packing bounded; large payloads use individual snapshots.
and sum(v.numel() * v.element_size() for v in value) <= 1024 * 1024
):
packed = torch.cat([v.detach().reshape(-1) for v in value])
sources.append(packed)
return _PackedTensorPayload(packed, [v.shape for v in value], isinstance(value, tuple))
items = [_clone_cuda_tensor_payload(v, sources) for v in value]
return tuple(items) if isinstance(value, tuple) else items
return value return value
def _copy_tensor_payload_to_cpu(value: Any, pin_memory: bool) -> Any: def _copy_tensor_payload_to_cpu(value: Any, pin_memory: bool) -> Any:
if isinstance(value, _PackedTensorPayload):
flat = _copy_tensor_payload_to_cpu(value.tensor, pin_memory)
items = []
offset = 0
for shape in value.shapes:
count = shape.numel()
items.append(flat[offset : offset + count].view(shape))
offset += count
return tuple(items) if value.is_tuple else items
if isinstance(value, torch.Tensor): if isinstance(value, torch.Tensor):
if value.device.type != "cuda": if value.device.type != "cuda":
return value return value
@@ -110,6 +135,12 @@ def _copy_tensor_payload_to_cpu(value: Any, pin_memory: bool) -> Any:
return value return value
class _PackedTensorPayload(NamedTuple):
tensor: torch.Tensor
shapes: list[torch.Size]
is_tuple: bool
class _AsyncCPUPayloadSnapshot: class _AsyncCPUPayloadSnapshot:
def __init__( def __init__(
self, self,
@@ -181,7 +212,7 @@ class OmniAsyncGPUModelRunnerOutput(AsyncGPUModelRunnerOutput):
if kwargs: if kwargs:
raise TypeError(f"Unexpected OmniAsyncGPUModelRunnerOutput kwargs: {sorted(kwargs)}") raise TypeError(f"Unexpected OmniAsyncGPUModelRunnerOutput kwargs: {sorted(kwargs)}")
self._model_runner_output = None self._model_runner_output: OmniModelRunnerOutput | None = None
self._invalid_req_indices = invalid_req_indices self._invalid_req_indices = invalid_req_indices
self.async_copy_ready_event = torch.Event() self.async_copy_ready_event = torch.Event()
@@ -211,7 +242,7 @@ class OmniAsyncGPUModelRunnerOutput(AsyncGPUModelRunnerOutput):
self._num_nans_cpu = self._num_nans.to("cpu", non_blocking=True) if self._num_nans is not None else None self._num_nans_cpu = self._num_nans.to("cpu", non_blocking=True) if self._num_nans is not None else None
self.async_copy_ready_event.record() self.async_copy_ready_event.record()
self._model_runner_output_builder = model_runner_output_builder self._model_runner_output_builder: Callable[[], OmniModelRunnerOutput] | None = model_runner_output_builder
self._background_exception: BaseException | None = None self._background_exception: BaseException | None = None
self._background_thread: threading.Thread | None = None self._background_thread: threading.Thread | None = None
self._cuda_device = cuda_device self._cuda_device = cuda_device
@@ -226,6 +257,7 @@ class OmniAsyncGPUModelRunnerOutput(AsyncGPUModelRunnerOutput):
if self._model_runner_output is not None: if self._model_runner_output is not None:
return return
with record_function_or_nullcontext("omni_async_output:get_output/build_model_runner_output"): with record_function_or_nullcontext("omni_async_output:get_output/build_model_runner_output"):
assert self._model_runner_output_builder is not None
self._model_runner_output = self._model_runner_output_builder() self._model_runner_output = self._model_runner_output_builder()
self._model_runner_output_builder = None self._model_runner_output_builder = None
@@ -321,6 +353,9 @@ class GPUARModelRunner(OmniGPUModelRunner, OmniConnectorModelRunnerMixin, Duplex
outputs per request while keeping Async output semantics. outputs per request while keeping Async output semantics.
""" """
execute_model_state: ExecuteModelState | None
kv_extracted_req_ids: list[str] | None
def __init__(self, *args, **kwargs): def __init__(self, *args, **kwargs):
super().__init__(*args, **kwargs) super().__init__(*args, **kwargs)
self.input_ids = self._make_buffer(self.max_num_tokens, dtype=torch.int32) self.input_ids = self._make_buffer(self.max_num_tokens, dtype=torch.int32)
@@ -436,6 +471,8 @@ class GPUARModelRunner(OmniGPUModelRunner, OmniConnectorModelRunnerMixin, Duplex
if not isinstance(info, dict): if not isinstance(info, dict):
return None return None
val = info.get("omni_final_stage_id") val = info.get("omni_final_stage_id")
if val is None:
return None
try: try:
return int(val) return int(val)
except (TypeError, ValueError): except (TypeError, ValueError):
@@ -1485,7 +1522,7 @@ class GPUARModelRunner(OmniGPUModelRunner, OmniConnectorModelRunnerMixin, Duplex
) )
selected_indices.append(idx) selected_indices.append(idx)
selected_req_ids.append(req_id) selected_req_ids.append(req_id)
selected_req_infos.append(req_info) selected_req_infos.append(cast(dict[str, Any], req_info))
if not selected_indices: if not selected_indices:
return multimodal_outputs return multimodal_outputs
@@ -1552,7 +1589,7 @@ class GPUARModelRunner(OmniGPUModelRunner, OmniConnectorModelRunnerMixin, Duplex
def _build_multimodal_outputs( def _build_multimodal_outputs(
self, self,
per_req_payloads: list[dict[str, object] | None] | None, per_req_payloads: Sequence[dict[str, object] | None] | None,
) -> list[dict[str, torch.Tensor] | None] | None: ) -> list[dict[str, torch.Tensor] | None] | None:
"""Build per-request multimodal output payloads (dedicated channel). """Build per-request multimodal output payloads (dedicated channel).
@@ -1648,8 +1685,8 @@ class GPUARModelRunner(OmniGPUModelRunner, OmniConnectorModelRunnerMixin, Duplex
def _build_omni_step_outputs( def _build_omni_step_outputs(
self, self,
pooler_inter: list[dict[str, object]], pooler_inter: Sequence[dict[str, object] | None] | None,
pooler_client: list[dict[str, object]], pooler_client: Sequence[dict[str, object] | None] | None,
*, *,
defer_full_payload_d2h: bool, defer_full_payload_d2h: bool,
) -> tuple[ ) -> tuple[
@@ -1935,6 +1972,8 @@ class GPUARModelRunner(OmniGPUModelRunner, OmniConnectorModelRunnerMixin, Duplex
) )
pooler_output.append(flatten_payload(payload)) pooler_output.append(flatten_payload(payload))
pooler_inter: Sequence[dict[str, object] | None] | None
pooler_client: Sequence[dict[str, object] | None] | None
pooler_output = pooler_output or [] pooler_output = pooler_output or []
if self._async_chunk and stage_sends_async_output(self.model_config): if self._async_chunk and stage_sends_async_output(self.model_config):
pooler_inter, pooler_client = partition_payload_list(pooler_output) pooler_inter, pooler_client = partition_payload_list(pooler_output)
+38 -15
View File
@@ -78,6 +78,8 @@ def _filter_mrope_kwargs_for_model(model: object, kwargs: dict[str, Any]) -> dic
class OmniGPUModelRunner(GPUModelRunner): class OmniGPUModelRunner(GPUModelRunner):
intermediate_tensors: IntermediateTensors | None
def __init__(self, *args, **kwargs): def __init__(self, *args, **kwargs):
super().__init__(*args, **kwargs) super().__init__(*args, **kwargs)
self.model_intermediate_buffer: dict[str, dict[str, Any]] = {} self.model_intermediate_buffer: dict[str, dict[str, Any]] = {}
@@ -842,7 +844,9 @@ class OmniGPUModelRunner(GPUModelRunner):
return None return None
@torch.inference_mode() @torch.inference_mode()
def extract_multimodal_outputs(self, hidden_states: torch.Tensor | list[torch.Tensor] | OmniOutput) -> dict: def extract_multimodal_outputs(
self, hidden_states: torch.Tensor | list[torch.Tensor] | OmniOutput
) -> tuple[Any, Any]:
if ( if (
hasattr(self.model, "have_multimodal_outputs") hasattr(self.model, "have_multimodal_outputs")
and self.model.have_multimodal_outputs and self.model.have_multimodal_outputs
@@ -1407,14 +1411,14 @@ class OmniGPUModelRunner(GPUModelRunner):
def _process_additional_information_updates( def _process_additional_information_updates(
self, self,
hidden_states: torch.Tensor, hidden_states: torch.Tensor,
multimodal_outputs: object, multimodal_outputs: Any,
num_scheduled_tokens_np: np.ndarray, num_scheduled_tokens_np: np.ndarray,
scheduler_output: "SchedulerOutput", scheduler_output: "SchedulerOutput",
combined_hidden_states: dict[str, torch.Tensor] | None = None, combined_hidden_states: dict[str, torch.Tensor] | None = None,
combined_multimodal_outputs: dict[str, object] | None = None, combined_multimodal_outputs: dict[str, Any] | None = None,
req_ids_filter: set[str] | None = None, req_ids_filter: set[str] | None = None,
req_ids: list[str] | None = None, req_ids: list[str] | None = None,
query_start_loc_cpu: object | None = None, query_start_loc_cpu: Any = None,
) -> None: ) -> None:
"""Process model-provided per-request updates and merge into model_intermediate_buffer.""" """Process model-provided per-request updates and merge into model_intermediate_buffer."""
req_ids = req_ids if req_ids is not None else self.input_batch.req_ids req_ids = req_ids if req_ids is not None else self.input_batch.req_ids
@@ -1476,9 +1480,9 @@ class OmniGPUModelRunner(GPUModelRunner):
def _collect_additional_information_for_prefill( def _collect_additional_information_for_prefill(
self, self,
num_scheduled_tokens_np: np.ndarray, num_scheduled_tokens_np: np.ndarray,
) -> dict[str, dict]: ) -> None:
"""Overlay per-request prompt_embeds for the prefill portion and collect """Overlay per-request prompt_embeds for the prefill portion and collect
additional_information slices for this step. Returns a map req_id -> dict.""" additional_information slices for this step."""
for req_index, req_id in enumerate(self.input_batch.req_ids): for req_index, req_id in enumerate(self.input_batch.req_ids):
req_state = self.requests[req_id] req_state = self.requests[req_id]
pe_cpu = getattr(req_state, "prompt_embeds_cpu", None) pe_cpu = getattr(req_state, "prompt_embeds_cpu", None)
@@ -1707,9 +1711,9 @@ class OmniGPUModelRunner(GPUModelRunner):
# Overlay custom prompt_embeds per request for the prompt portion; # Overlay custom prompt_embeds per request for the prompt portion;
# collect additional_information (tensor/list) for prefill portion only # collect additional_information (tensor/list) for prefill portion only
decode_req_ids = [] decode_req_ids: list[str] = []
decode_start_offsets = [] decode_start_offsets: list[int] = []
decode_batch_items = [] decode_batch_items: list[tuple[str, int, dict[str, Any]]] = []
batch_decode_preprocess = getattr(self.model, "preprocess_decode_batch", None) batch_decode_preprocess = getattr(self.model, "preprocess_decode_batch", None)
def flush_decode_batch() -> None: def flush_decode_batch() -> None:
@@ -1717,6 +1721,7 @@ class OmniGPUModelRunner(GPUModelRunner):
if not decode_batch_items: if not decode_batch_items:
return return
assert callable(batch_decode_preprocess)
req_ids_b = [item[0] for item in decode_batch_items] req_ids_b = [item[0] for item in decode_batch_items]
start_offsets_b = [item[1] for item in decode_batch_items] start_offsets_b = [item[1] for item in decode_batch_items]
req_infos_b = [item[2] for item in decode_batch_items] req_infos_b = [item[2] for item in decode_batch_items]
@@ -1761,7 +1766,7 @@ class OmniGPUModelRunner(GPUModelRunner):
# mimo-audio check # mimo-audio check
req_state = self.requests.get(req_id) req_state = self.requests.get(req_id)
req_infos = self._maybe_attach_mimo_audio_req_infos(req_state, req_infos, req_id) req_infos = cast(dict[str, Any], self._maybe_attach_mimo_audio_req_infos(req_state, req_infos, req_id))
start_offset = int(self.query_start_loc.cpu[req_index]) start_offset = int(self.query_start_loc.cpu[req_index])
sched_tokens = int(num_scheduled_tokens_np[req_index]) sched_tokens = int(num_scheduled_tokens_np[req_index])
@@ -1968,11 +1973,29 @@ class OmniGPUModelRunner(GPUModelRunner):
if start_offsets is None: if start_offsets is None:
id_to_index = self.input_batch.req_id_to_index id_to_index = self.input_batch.req_id_to_index
start_offsets = [int(self.query_start_loc.cpu[id_to_index[req_id]]) for req_id in decode_req_ids] start_offsets = [int(self.query_start_loc.cpu[id_to_index[req_id]]) for req_id in decode_req_ids]
for idx, (req_id, start_offset) in enumerate(zip(decode_req_ids, start_offsets, strict=True)): if start_offsets == list(range(decode_batch_size)):
inputs_embeds[start_offset : start_offset + 1] = req_embeds[idx : idx + 1] inputs_embeds[:decode_batch_size].copy_(req_embeds[:decode_batch_size])
if code_predictor_codes is not None: else:
update_dict = {out_key[0]: {out_key[1]: code_predictor_codes[idx : idx + 1]}} # A device tensor constructor waits for the preceding MTP graph.
self._update_intermediate_buffer(req_id, update_dict) # Enqueue the small host index transfer without synchronizing it.
offsets = torch.tensor(start_offsets, device="cpu", dtype=torch.long).to(
inputs_embeds.device, non_blocking=True
)
inputs_embeds.index_copy_(0, offsets, req_embeds[:decode_batch_size])
if code_predictor_codes is not None:
if out_key in getattr(self.model, "gpu_resident_buffer_keys", set()):
# One owned batch snapshot protects all rows from graph reuse.
owned_codes = code_predictor_codes[:decode_batch_size].detach().clone()
for req_id, row in zip(decode_req_ids, owned_codes.split(1), strict=True):
req_state = self.requests.get(req_id)
if req_state is not None:
existing = self.model_intermediate_buffer.setdefault(req_id, {})
existing.setdefault(out_key[0], {})[out_key[1]] = row
req_state.additional_information_cpu = existing
else:
for idx, req_id in enumerate(decode_req_ids):
update_dict = {out_key[0]: {out_key[1]: code_predictor_codes[idx : idx + 1]}}
self._update_intermediate_buffer(req_id, update_dict)
def _model_forward( def _model_forward(
self, self,