[Chunked Prefill #2] Combined Walks by NSagan271 · Pull Request #362 · mstar-project/mstar

NSagan271 · 2026-10-07T02:26:11Z

What does this PR do?

Stack 2 of #266.

Allows for multiple different walks to run in one batch (e.g., allow decode rows to run along with a prefill batch).

Adds

def get_combined_graph_walks(self) -> dict[str, dict[str, set[str]]]:
        """``{node: {combined_walk: {constituent walks}}}``: walks the scheduler may batch together."""
        return {} # e.g., {"LLM": {"mixed_ar": {"prefill", "decode"}}}

to the base Model class, allowing a model to specify, per-node, which walks can be run together.

Scheduling automatically checks for all walks in the combined walk during the batch building process
Batch size caps: a step of one walk keeps that walk's cap. A step that mixes walks takes the combined walk's own max_batch_size when the model gives one, else the minimum cap among the walks it holds, so a low prefill cap never limits a decode-only step. max_rows_in_combined(graph_walk) also bounds how many rows of one walk a mixed step may hold (documented in adding_models.rst)
Speculation: by default a speculation chain stays on its own walk, merging only that walk's ready rows and yielding to the other walks of its combined walk; a step mixes walks when the scheduler builds one fresh. speculation_mixes_walks in the serving yaml (per node, or per walk) lets a chain merge the other walks' ready rows too. It is off here because a mixed step plans inline today; it is meant to be turned on by default with pre-planning later in the stack
Graph runtime semantics:
- The graph runtime has access to the combined graph walks; in scheduling passes, the resolution of constituent graph walks happens within the graph runtime (e.g., in Rust)
- Same semantics for postprocess: one roundtrip through the runtime where possible
Engine / models
- StepContext.force_prefill: a step leased on a packed capture plans prefill attention even when every row is one token long. Previously a packed bucket with bs == num_tokens got FlashInfer's decode wrapper and failed to capture.
- Qwen3.5 batches text prefill and decode under one combined walk. A mixed step runs its decode rows on the decode kernels and its prompt rows on the prefill/chunked kernels (CombinedSplit on the attention and GDN resources), which avoids planning and running every decode row as a packed prefill row. The prefill side is captured for as many rows as max_rows_in_combined lets in (4 for Qwen3.5)
- A row whose walk keeps its sampled token on the device (decode loop-back) keeps doing so in a mixed step
- BAGEL: the main -> cfg_text pre-fork of a mixed step is scoped to its prompt rows (KVStep.fork_rids), so a decode row sharing the step keeps its own cfg_text. Guidance labels stay per step, as on main
- Orpheus packed prefill captures at batch sizes [1, 4, 16, 64]. A packed graph's GPU work is set by its token count, so coarse batch buckets cost only minor CPU overhead and let a mixed step hold all the decode rows. BAGEL ([1..64]) and Qwen3-Omni ([1..32]) also widen their packed captures.
Init-time checks
- On each worker, a combined walk's walks must share one TP group, or all have world size 1.
- A node must not split its combined walk by capture key.
- Not checked: walks placed on different workers (e.g. PD disaggregation). Each worker then sees one walk of the combined walk, which is safe, so no config change is needed when PD is switched on.

How was it tested?

A/B against the previous PR in the stack, fixed output lengths, 2 restarts per side, H200.

Model	c=16	c=48
BAGEL i2t	level	level
Qwen3-Omni (Thinker TP2), audio→speech	+4.1% req/s, −3.9% p50	+8.4% req/s, −7.8% p50, −7.7% TTFT
Orpheus	+0.4% req/s	+0.1% req/s

CPU time in the scheduler (independent benchmark), compared to doing one call into the graph runtime per walk:

Scenario (p50 µs)	main	this PR, not combined	this PR, combined
decode step, 48 rows	118.5	111.4	116.6
decode step, 64 rows	151.9	142.4	147.8
63 decode + 1 prompt, until nothing is left, 64 rows	173.0 (2 steps)	169.3	183.1 (1 step)
speculation merge, 48 rows	18.4	23.3	24.7

Postprocess CPU time (independent benchmark), compared to doing one call into the graph runtime per walk:

Postprocess step (µs)	per-walk calls	one call	single-walk step of the same size
15 + 1 rows	301	276	278
63 + 1 rows	889	866	848
32 + 32 rows	1,092	865	848

BAGEL i2t was +6% in our runs but level with the base in a reviewer's runs on H100, so it is listed as level.

Qwen3.5-0.8B, paired A/B against main on one H100: both servers resident, alternating cells, variable output lengths (128 to 384 tokens) with staggered arrivals so prompts land beside decode rows.

Qwen3.5-0.8B, text tok/s vs main	c=32	c=64	c=128
median of 8 rounds	−1.9%	−2.1%	−1.1%
mean of 16 rounds (earlier run)	−3.0%	−3.6%	−1.7%

The node's CPU was shared with other jobs, so single rounds vary by several percent. The remaining gap is the inline plan of mixed steps (10 to 20% of steps), which the pre-planning PR later in the stack removes. With speculation_mixes_walks on, against off in the same process: −4.0% / −0.7% / −1.4% at c=32/64/128, from planning more mixed steps inline, hence off by default here. A new request's TTFT while 40 others decode (decode cap 32) is p50 35 ms, level with main.

Smoke tests at the head on H100: Orpheus with 24 TTS requests of staggered lengths at concurrency 8 (every clip transcribes back to its text with Whisper), and BAGEL with 8 chats and 2 images sent together so T2T decode rows share steps with T2I prompts (correct answers, images match their prompts); no failed requests or server errors on either.

Greedy Qwen3.5 outputs under mixed-step traffic differ from main only within main's own batch-to-batch variation. Suites: the modular suite passes except 15 tests that fail identically on main (stale test doubles, and misaki missing for Kokoro).

Checklist

ruff check . passes
Added or updated tests / docs where relevant

Review follow-ups on #362, mostly cutting per-step Python work under a combined walk: - speculate_node returns every walk's targets (each names its walk) and the caller picks; the async check is per walk in the runtime, so a batch may speculate when any of its walks is async (Python and Rust). - pop_walk_rids pops rows of several walks of one node in one runtime call, from columns (RequestWalks: rids, walks, walk_idx); all or none under check_ready. Both scheduler pop sites use it, and the TP head carries its walks as the same columns. A follower's batch now gets each walk's output signals from the pop rather than asking per walk at routing. - ScheduledBatch caches rids_by_walk; every row removal goes through discard_rid, which resets it. - The FIFO builder composes a combined step row by row: a row joins unless its walk (or the combined walk's cap, once the step mixes walks) would cap the step at or below the rows already in it. It takes raw caps plus the caller's own row count and walk, instead of a separate mixed cap. - The scheduler asks the engine for each (node, walk) cap once. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

merceod

Ran this on klone H200s (top of the stack c6da5d4 for the GPU runs, this head for the suites) against main 3c5f0dd: suites, greedy and seeded parity on Qwen3.5-0.8B, late-arrival latency under a full decode set, and BAGEL text and image requests sharing steps. Two things I think need changes, the rest are notes.

A new request cannot get its prefill in while a decode chain of the same node is at its cap. With combined walks, prefill_text and decode of LLM share the scheduling key (LLM, mixed). has_ready_excluding maps exclude_target through _key (micro_scheduler.py:979), so a ready prefill row never counts as other work and the speculation chain never yields for it, and the targeted merge returns None when the chain's own walk is at its cap (get_next_batch, "the caller is full"). On Qwen3.5-0.8B with 40 requests decoding (decode cap 32) and a short request sent every second, the newcomers' TTFT was p50 117-394 ms and p90 1.5-2.8 s over four runs. On the #354 head it is p50 30-64 ms / p90 214-824 ms over three runs, on main p50 31-37 / p90 276-534. At 16 concurrent every arm is at 25 ms, and Qwen3.5-9B TP2 (no speculation chain) does not show it at all. The chain should yield when a row of a different real walk is ready under the same key, or the merge should take fresh rows of the other walks even when the chain's own walk is full.
BAGEL mixed_text with guidance changes results. requires_cfg, the labels and the main to cfg_text pre-fork are decided for the whole step (bagel/submodules.py:909-979), so a decode row that shares a step with a guided prefill gets cfg_text/cfg_img segments and the fork too. Seeded T2I on the stack is deterministic (six prompts, same hashes across servers and against main), but with T2T requests decoding beside them one of the six images changed, and one of six greedy T2T outputs changed beside T2I traffic. On main both runs are identical to the alone runs. Scoping labels and forks per row (the fork_rids that #365 adds for KV, but for labels and the position manager too) would fix it.

Notes

The cap is the min over the walks present in the step, not over all constituents of the combined walk (batch_builder.py:160-183). Fine, but the body says min over the walks in the batch, so worth a line in the doc.
refuse_combined_walks_across_tp_groups filters by node, so a PD setup with TP on one side raises on both workers.
Rust pop_walk_rids returns None on bad lengths or indices where the Python runtime raises, which would show up as a follower stall rather than an error.

Numbers that held up: closed-loop 0.8B at c=32 is 42.8 req/s on the stack vs 29.8 on main, long-prompt TTFT p50 255 vs 360 ms, greedy outputs inside main's own batch-to-batch floor, suites clean (this head fixes the two tp_follow tests #354 breaks).

merceod · 2026-10-11T11:33:27Z

        under `exclude_target` is left out like any other — the speculative
        merge takes that one before its own rids, so it needs no yield.
        """
+        exclude_target = None if exclude_target is None else self._key(*exclude_target)


Mapping exclude_target through the combined key means a ready prefill_text row of the same node never counts as other work here, so the speculation chain never yields for a newcomer while its decode rows fill the cap. On Qwen3.5-0.8B with 40 requests decoding that is a p50 of 120-390 ms and a p90 of 1.5-2.8 s for a new request's first token, against about 30 / 300 ms on main and on #354.

One more thing from the tiny-pool runs, separate from the two above. On this head alone (its own worktree and venv, 48-page pool, 48 prompts of 1k to 3k tokens with 256-token outputs at c=8 and c=16), 7 and then 9 of the 48 requests fail with a 500 "KeyError: 'text_inputs'" from prepare_inputs on a decode walk, right after an OOM hold ("holding 2 of 5 requests so the retry is smaller"). The #364 head and main run the same workload with the same number of OOM holds and no failures, so the push-back and retry of a held decode row loses its text_inputs somewhere in this PR's scheduler changes. Log: 14 holds, 0 offloads, 19 failing requests.

Prefill starvation and KeyError fixed. An identifiable cause was bad cleanup of terminating loops that were still speculatively scheduled at time of OOM. Odd, because this should also happen in main.

@stephen-dwq Both fixes hold in my runs. With 40 requests decoding on Qwen3.5-0.8B, a new request's TTFT p50 is 2.4 s -> 35 ms, and on a 40-page Orpheus pool the KeyError went from 14/192 requests to none. main hits that KeyError too but rarely (1/192).

merceod · 2026-10-11T11:33:27Z

+            if "prefill_text" in row_walks:
                # cfg_text keeps the pre-text context, so it forks before
                # anything plans or writes.
                pre_forks = (("main", "cfg_text"),)


This fork and the labels above are decided for the whole step, so a decode row that shares a mixed_text step with a guided prefill gets cfg_text/cfg_img segments and the fork too. With T2T requests decoding beside seeded T2I prompts one of six images changed, and one of six greedy T2T outputs changed beside T2I traffic; on main both stay identical.

Correction on this one after more repeats. With 8 prompts and two runs per arm, main also changes outputs beside other traffic: seeded T2I next to T2T decode 8/8 identical in one run and 6/8 in the next, greedy T2T next to T2I traffic 6/8 (the 6-prompt main run above was 6/6 by luck). This head gives 8/8, 7/8 and 7/8 on the same checks. So the output differences are batch-composition numerics that main has too, not something this PR adds, and I am withdrawing that part. The batch-level labels and the main->cfg_text pre-fork for the decode rows of a mixed_text step are still worth a look for the pages they reserve, but I have no measurement of that cost.

For a think_mode image request, sharing a step with a guided prompt overwrites its cfg_text with main (19 tokens became 205 in my run), so cfg_text ends up holding the prompt it's meant to exclude. #365's fork_rids fixes this, so nothing to change here.

I cherry-picked the fork_rids change onto this branch: 11d96fc and a87fd4f

merceod · 2026-10-11T11:33:27Z

+
+
+def _min_cap(*caps: int | None) -> int | None:
+    return min((cap for cap in caps if cap is not None), default=None)


The cap is the min over the walks present in the step, not over every constituent of the combined walk. Fine as behaviour, but the body says min over the walks in the batch, worth saying in the doc too.

This was an intentional change; updated the body and the documentation to match

Gaurav-Shah05

Overall clean. the Thinker's wider capture set needs another look, since with it the default configs/qwen3omni_2gpu.yaml no longer boots on a H100 GPU, which I'll handle in #341.

I served Orpheus, BAGEL, Qwen3.5 and Qwen3-Omni from this branch on H100s, and mixed steps runs cleanly on all four, and throughput was level with the base. On BAGEL i2t at c=16 I got level throughput with base, so I couldn't reproduce the 6% gain, though my workload may not match yours and it may have been noise across boots.

Gaurav-Shah05 · 2026-10-12T14:55:19Z

    PREFILL_VISION_TOKEN_BUCKETS = [128, 256, 512, 1024, 2048, 4096, 8192, 16384]
-    PREFILL_VISION_CAPTURE_BATCH_SIZES = [1]
+    # coarse batch buckets, padded host-side; past 1 for mixed steps
+    PREFILL_VISION_CAPTURE_BATCH_SIZES = [1, 4, 16, 32]


With these captures the default configs/qwen3omni_2gpu.yaml runs out of memory at boot on an 80 GB H100 (29 Thinker buckets become 68), in the declare_inputs_for pre-seed that #341's budget doesn't count yet, so I'll fix that in #341.

I've applied the fix. It stopped pre-seeding template rows for packed buckets, since nothing reads them, which frees ~1 GB on the Thinker's GPU. This now boots and captures 64 of the 68 Thinker buckets.

Gaurav-Shah05 · 2026-10-12T15:01:53Z

+            if "prefill_text" in row_walks:
                # cfg_text keeps the pre-text context, so it forks before
                # anything plans or writes.
                pre_forks = (("main", "cfg_text"),)


For a think_mode image request, sharing a step with a guided prompt overwrites its cfg_text with main (19 tokens became 205 in my run), so cfg_text ends up holding the prompt it's meant to exclude. #365's fork_rids fixes this, so nothing to change here.

Gaurav-Shah05 · 2026-10-12T15:02:45Z

+def _tp_group_of(sharding: ShardingConfig, node: str, walk: str) -> ShardingGroup | None:
+    """The TP group ``node`` runs ``walk`` in, from the static config; None
+    when it has none or one of size 1."""
+    for group in sharding.groups:


_tp_group_of looks for each walk's group across the whole config, so a P/D config at TP2 on both sides, which the docs allow, fails to start on every rank even though no rank runs both walks. Could it compare only the walks this worker runs, so P/D needs no config change?

43b766f now keys on (node, walk) pairs that the worker runs

Gaurav-Shah05 · 2026-10-12T15:06:51Z

+        if key not in self._walks_of_key:
+            return
+        entries = [
+            entry for entry in self._scan_ready(request_state, key, None).get(key, [])


nit: this also runs check_ready on the chain's own waiting rows, which raised the speculate phase from 0.14 to 0.26 ms a step on Orpheus at c=128. Passing (key[0], own_walk) as the exclusion parks the same rows without those checks.

Done, 374b68c

Gaurav-Shah05 · 2026-10-12T15:08:48Z

        under `exclude_target` is left out like any other — the speculative
        merge takes that one before its own rids, so it needs no yield.
        """
+        exclude_target = None if exclude_target is None else self._key(*exclude_target)


@stephen-dwq Both fixes hold in my runs. With 40 requests decoding on Qwen3.5-0.8B, a new request's TTFT p50 is 2.4 s -> 35 ms, and on a 40-page Orpheus pool the KeyError went from 14/192 requests to none. main hits that KeyError too but rarely (1/192).

Gaurav-Shah05 · 2026-10-12T15:12:02Z

@merceod both of your findings are fixed by Stephen, and they hold in my runs. Could you take another look?

A targeted call now resolves the walk's cap through _room_left, which asks the engine, and the test's _FakeEngine has no get_max_batch_size -- so two tests error on this head while passing on the base. #362 happens to mask it by caching the lookup; the incomplete fake belongs here. None is no cap: these tests are about the TP-follow FIFO, not caps. Co-Authored-By: Claude <noreply@anthropic.com>

* Add batch builder abstraction * Bug fix: set backlog batches as in-flight to prevent double-scheduling * Backlog: a held or pending-remove row waits without a readiness check The fresh path already skipped held_until and pending_removes, but the backlog recheck asked the engine about every row. check_ready reloads an offloaded row, so a row held after an OOM went straight back to the OOM it was offloaded for. Such rows now count as blocked without the check. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * Speculation merge: a caller with no room left skips the scan and the backlog Speculation asks for a targeted merge every step. When its own rows fill the cap there is nothing to take, so return before the ready scan and before re-checking the backlog only to park it again. When its own rows and the live backlog fill the cap, no fresh row could join, so skip the ready scan. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * Give the tp-follow gate's fake engine a max batch size A targeted call now resolves the walk's cap through _room_left, which asks the engine, and the test's _FakeEngine has no get_max_batch_size -- so two tests error on this head while passing on the base. #362 happens to mask it by caching the lookup; the incomplete fake belongs here. None is no cap: these tests are about the TP-follow FIFO, not caps. Co-Authored-By: Claude <noreply@anthropic.com> * FIFO builder: hand a fresh row the cap pushed out back to its queue Everything past the cap parked, fresh rows included. A key holding a backlog is served before any other node, so a node whose cap is small enough to keep re-parking leftovers wins every pass: on wan22.yaml with dit capped at 1 and four requests, a fifth request's text_encoder waited 48.7 s for its first step against 0.2 s on main. A row that arrived fresh and was pushed out by the cap alone now goes back to its ready queue through `returned`, so the backlog drains every step as it did on main. A row held back for any other reason -- OOM backoff, pending remove, a foreign capture group -- still parks: it cannot run this step whoever holds it, and re-popping it every pass is what the backlog exists to avoid. The four tests that pinned the old "everything parks" contract now assert the new one; all four fail against the old builder. Co-Authored-By: Claude <noreply@anthropic.com> * Scheduling: refuse a builder that loses a row Nothing checked that a builder returns every row it was given. A row missing from scheduled, backlog and returned has already come off its ready queue and is marked in flight, so nothing would ever schedule it again -- the request hangs for good, with no evidence of where it went. One set per step, built from the request before the build (`_merged` folds fresh into backlog in place), and the failure names the rows. Co-Authored-By: Claude <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>

A model declares {node: {combined_walk: {walks}}} via Model.get_combined_graph_walks(); the worker builds a validated (node, real walk) -> combined walk map for the scheduler. - MicroScheduler keys scans, backlog, round-robin and caps on the combined walk, pops each real walk separately, and records per-rid walks in ScheduledBatch.request_walks. A batch whose rows share one real walk is labelled with that walk, so uniform steps keep their CUDA graphs. A combined key is capped at the smallest cap of its walks. - StepContext.request_walks / NodeInputs.graph_walk carry each row's real walk to prepare_inputs, the prefix cache, declare_step and KV publish. - The worker stores, routes and sends outputs once per real walk, and stops loops per walk. - TP: ScheduleTPNode.request_walks; followers pop per walk, all or none. - Speculation: speculate_node takes columnar (walks, sample_rids) and returns the first walk with a target; only that walk's rids continue. - EngineManager refuses split_batches_by_capture_key on a combined walk. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

LLM's prefill and decode rows share a "mixed" walk. Mixed steps replay the packed prefill CUDA graphs, where a decode row is a one-token segment. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

The smallest cap of the constituent walks throttled decode-only batches to the packed-prefill capture size (4 on BAGEL / Qwen3-Omni). The model now declares the combined walk's cap through its captures or max_batch_size. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

stephen-dwq

LGTM

vasilevklart · 2026-10-18T20:46:14Z

Just noticed a perf bug on main and 4512e23bf in #362 (sample under compiler.disable + orpheus/bagel row_outputs) fixes by every uncaptured step recompiling SamplerResource.sample on the new request ids, ~0.5 s with the gpu idle. On cuda thats packed prefill past the 8192 bucket and on rocm its every gdn prefill over 512 tok, with the fix it is +16% on mixed load and +24% on long prompts on h100 (qwen3.5-9b)

# Conflicts: # mstar/worker/micro_scheduler.py # test/modular/test_worker_drain.py

# Conflicts: # test/modular/test_worker_drain.py

# Conflicts: # mstar/worker/worker.py

… engine reads Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

NSagan271 · 2026-10-19T05:03:38Z

/review

github-actions · 2026-10-19T05:06:06Z

Automated review

The earlier automated inline comment (new CPU-only tests not listed in test/cpu-core.txt) is fixed: all seven new modules are now in the list.

This PR adds Model.get_combined_graph_walks, so the micro-scheduler and batch builder can put rows of several walks of one node into one step. Rows carry a per-row request_walks map, and the Python and Rust graph runtimes get the matching pop, route, set-in-flight and speculation entry points. Qwen3.5 uses it with a CombinedSplit on its attention and GDN resources. BAGEL, Orpheus and Qwen3-Omni widen their packed captures. The code matches the description.

Findings: none. The diff is about 8,600 lines and this light pass didn't cover all of it. Covered: the worker side (batch_builder.py, micro_scheduler.py, engine_manager.py, the scheduling, speculation, TP-follow and postprocess changes in worker.py, ScheduleTPNode), model/base.py, cuda_graph_config.py and the docs. Not reviewed: Python/Rust parity of the new graph-runtime methods (rust/src/graph/runtime.rs against mstar/graph/runtime/python.py), and the engine/resource side (engine.py beyond the loop-back handling, attn/wrappers.py, linear_attn/wrappers.py, kv/manager.py fork_rids, and the Qwen3.5, BAGEL and Qwen3-Omni submodule changes). Those need a thorough review.

The description lists results but not the commands behind them. A maintainer will want:

Test commands: the exact commands for the modular suite and python -m pytest @test/cpu-core.txt, and the names of the 15 tests that fail the same way on main.
Rust tests: whether cargo test and test/rust/ (including the extended test_runtime_parity.py) were run. Invariant 5 applies because both runtimes change.
Benchmark reproduction: the harness, command and flags for the H200 A/B table and the two scheduler/postprocess CPU-time tables. "Independent benchmark" doesn't say how to rerun them.
Python runtime: the follow-up live test says the Python graph runtime wasn't exercised on this head. Please say whether MSTAR_RUST_GRAPH=0 was run with mixed steps.

Questions

_build_speculation now excludes the speculation target's (node, walk) in has_ready_excluding and in the yield-away, where it used to exclude the in-flight step's (node, walk). For a same-node loop-back the two are the same. For a cross-node chain A→B, A's ready rows now count as other work, B's no longer do, and the yield-away schedules away from B instead of A. That behavior change applies to every model, not just ones with combined walks. Was it intended, and was a model with cross-node speculation checked for fairness?

_{Advisory, not a merge gate. Rules: AGENTS.md. Reply to discuss.}

…xt per window (#452) Orpheus streams SNAC audio per decode step with a small capture_group hook, instead of decoding the whole code sequence at the end. Reviewed on 53b0943 and verified on klone; the rebased head 9c6d531 carries the identical code (only the orpheus import list changed shape after #362), CI green.

Engine.check_stop_for_batch calls check_stop_batched(..., step_walk=...) since #362 (merged in from main). The ASR decoders' overrides predate it and do not take the keyword, so every call raised TypeError, was logged at debug, and the engine fell back to per-request check_stop: one .item() per row per step instead of one tolist. Both paths give the same stops, so this is overhead, not a correctness bug. The walk is still read per row, as before. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

* transcriptions: windows for bounded ASR models Long uploads are cut by the adapter's clip limit, uploads with no length in their header are measured, every decode stops at the route's cap, and a window without timestamps is one segment (Higgs and untimed Whisper captions). * fastconformer: shared encoder and log-mel NVIDIA's FastConformer encoder, offline or chunked for cache-aware checkpoints, NeMo's log-mel, and the encoder node, captured per mel bucket. * parakeet: Parakeet-TDT and Nemotron ASR Greedy TDT and RNN-T decoding in captured chunks of a compiled step; RNN-T loops over labels. * canary: Canary-1B-v2 Speech recognition in 25 languages and translation to or from English. Whisper and Canary share one KV-fit check. * canary_qwen: Canary-Qwen-2.5B Canary's encoder projected into a LoRA-merged Qwen3-1.7B, English only. * components: one encoder-decoder node for Whisper, Canary Whisper's and Canary's decoder nodes share the prefill/decode flow; Whisper keeps its token suppression as a process_logits hook. * asr: review fixes on the rebased stack - TokenStepNode.check_stop_batched takes step_walk, which the engine passes since #362; without it every call raised and fell back to the per-request stop check. - The FastConformer ASR adapters stream text without SentencePiece's leading space, and the route sends a window-joining space only before a window's first text, so the deltas add up to the done event's text. - Parakeet, Nemotron and Canary-Qwen refuse a task other than transcribe instead of returning a transcript for a translation. * perf(fastconformer): plan a batch's encoder replays by cost, not one series per bucket * perf(fastconformer): log-mel features for the whole batch in preprocess * perf(fastconformer): the position term as one GEMM per head, half-step residuals in the add * perf(fastconformer): Triton kernels for the conv module and the subsampling's bias, mask and ReLU * perf(fastconformer): each residual add in the kernel of the LayerNorm after it * perf(fastconformer): refit the replay cost model to the fused encoder * docs: Omnilingual ASR gets a segment and a caption cue per window A window without timestamps is one segment spanning it (merge_windows), so the CTC model's verbose_json and captions follow its windows, as Whisper's and Higgs-Audio's untimed windows do. * perf(transcriptions): skip the thread hop for an upload prepare() leaves alone * perf(media_io): skip the channel mean for mono audio in read_soundfile_capped * Canary-Qwen: take a custom instruction from model_kwargs, not /generate's text process_prompt used /generate's text prompt as the instruction, so the ASR harnesses' default prompt ("Transcribe the speech in this audio clip.") replaced the trained one, "Transcribe the following:", and Canary-Qwen was benchmarked under an instruction it was not trained on. The text prompt is now unused, as on the other FastConformer ASR models, and a custom instruction comes from model_kwargs' "instruction". Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * Docs: untimed transcription models get one segment per window The Omnilingual ASR adapter's docstring still said srt/vtt get one cue spanning the file and verbose_json an empty segment list. merge_windows now makes each window with text one segment spanning it, so each format gets one cue or segment per window, as docs/models.rst and docs/clients.rst already say. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * media_io: decode_audio is load_speech_file, so a missing FFmpeg stays a 500 The transcription route decoded uploads with its own copy of the libsndfile-then-FFmpeg chain, which turned every torchcodec RuntimeError into a 400, a missing FFmpeg included. load_speech_file already tells the two apart; the copy and its numpy resample fallback go (torchaudio is a dependency). * fastconformer: log-mel per replay group, not padded to the batch's longest clip A 600 s Parakeet upload batched with short clips padded every clip's STFT to 600 s. preprocess now plans the replay groups and computes each group's features together (a clip past every bucket alone), and hands the plan to the encoder so it isn't computed twice. * canary, canary_qwen: declare their KV labels and lay out each request Main refuses a paged pool with no labels at engine build (#832), so neither model started. Canary follows Whisper: the decoder's cache grows a slot per decode step and the encoder context is one rows span of the clip's frames. Canary-Qwen's prompt is a rows span: its audio placeholders are the same for every clip of a length, so it must never key a cached prefix. * parakeet, nemotron_asr: out of this PR, into a follow-up on top of it Review asked for the transducer to step through the graph as a loop, with its LSTM state and cursor in a resource, rather than run the whole greedy loop inside one forward. That rework lands with the two transducer models in their own PR, stacked on this one. This PR keeps Canary, Canary-Qwen, the shared FastConformer encoder, the encoder-decoder and token-step nodes, the Whisper refactor and the transcription route changes. Removed: mstar/model/parakeet/, components/transducer.py, their configs, registry, CLI and OpenAI adapter entries, benchmark entries, docs and tests. The streamed-deltas test of the transcription route now runs on Canary, the other windowed SentencePiece model. The encoder keeps its causal path, which its own tests cover. * audio_features: use the engine's fp32_precision_scope for the mel matmul Drops mstar/utils/precision.py, a second copy of what mstar/engine/fp32_precision.py already does, which restored only the legacy matmul flag and not the per-backend ones. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * canary, canary_qwen: cap the decode loop at the request's max_output_tokens Without get_loop_caps the loop cap fell back to the graph's max_iters, so speculation ran past a max-token stop and a smaller kv_cache override refused every request. Same override as Whisper's. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * transcriptions: no compression-ratio retry for the untimed ASR adapters The 2.4 threshold is Whisper's sequential-long-form fallback. Canary and Canary-Qwen decode greedily with no conditioning text, so a forced sequential retry could only replace a correct transcript with a sampled one. Omnilingual already opts out the same way. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * audio_features: only a failed decode is the upload's fault load_speech_file wrapped the resample too, so a server-side failure there (e.g. a CPU allocation failure, a RuntimeError) came back as a 400 'could not decode the audio upload'. Split the decode from the resample and wrap only the decode. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * canary: the prompt's KV is a rows span From the second decoder layer on, the prompt's self-attention KV depends on the clip through cross-attention, so its token ids don't key it. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * canary: refuse a checkpoint whose mel front end NemoLogMel doesn't compute Canary-Qwen already checks its preprocessor settings at load; Canary assumed them. Read processor_config.json's feature extractor and refuse a mismatch at startup instead of decoding wrong features. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * token_step: say decode is uncapped with accelerator graphs off Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: load_speech_file's blame split goes through _decode, and a resample failure stays the server's Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * canary, canary_qwen: one audio file per request, refused before decoding Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * asr extra installs soundfile; Higgs segments and a stale test comment documented Without soundfile every WAV/FLAC upload went through torchcodec and needed FFmpeg, which the docs said only M4A/WebM do. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> --------- Co-authored-by: Kirill Vasilev <290354918+vasilevklart@users.noreply.github.com> Co-authored-by: Joseph Tan <jdtan638@gmail.com> Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>

* whisper: read encoder dims, alignment heads and mel params from the checkpoint * whisper: add a native batched audio encoder * whisper: add timestamp decoding rules as a staged per-row state * whisper: capture the encoder per batch size, detect language, apply timestamp rules * whisper: native encoder, language detection walks, initial_prompt and timestamps * qwen3_asr: add the model package * qwen3_asr: read the AuT and Qwen3 configs from the checkpoint * qwen3_asr: packed AuT encoder and Qwen3 decoder submodules * qwen3_asr: add the Qwen3-ASR model graph and prompt * registry: add whisper_large_v3_turbo, qwen3_asr and qwen3_asr_realtime * cli: default configs and next steps for the ASR models * configs: add whisper_large_v3_turbo * configs: add qwen3_asr * configs: add qwen3_asr_realtime * benchmark: register the new ASR model types * benchmark: accept the new ASR model keys in asr_eval * benchmark: build the shared LibriSpeech ASR set from openslr * benchmark: protocol ASR benchmark over /v1/audio/transcriptions * benchmark: streaming partial-latency benchmark over /v1/realtime * benchmark: reference-implementation transcripts for WER parity * docs: Whisper turbo/v3 and Qwen3-ASR rows in the model table * tests: Whisper config, encoder parity, state machine and wiring * tests: pin Whisper timestamp rules to HF's logits processor * tests: Qwen3-ASR prompt, submodules, state machine and wiring * tests: GPU parity for Whisper on real weights * tests: GPU parity for Qwen3-ASR on real weights * api: add the Qwen3-ASR transcription adapter * tests: cover the Qwen3-ASR OpenAI adapter * whisper: carry the timestamp rule state as a tensor and advance it in the forward * whisper: route the rule state with the token instead of a host history * whisper: loop the timestamp rule state back through the decode walk * tests: advance reproduces rule_state token by token * tests: routed timestamp state in the decoder and the state machine * api: report Qwen3-ASR's detected language as an ISO code * tests: Qwen3-ASR adapter maps language names to codes * benchmark: realtime bench speaks vLLM's WebSocket dialect and reports delta latency * whisper: opt the submodules out of the engine's torch.compile * qwen3_asr: opt the submodules out of the engine's torch.compile * benchmark: realtime bench tolerates close frames * whisper: send the encoder samples and make the log-mel the encoder's first GPU step * whisper: encoder submodule pads the window and makes the log-mel batch on the GPU * qwen3_asr: send the encoder samples, the prompt only needs the frame count * qwen3_asr: encoder submodule makes the log-mel on the GPU * tests: Whisper encoder takes samples and the tiny config uses a 1 s window * tests: Qwen3-ASR encoder takes samples * tests: Whisper GPU parity checks the submodule's log-mel against the CPU transform * benchmark: vLLM realtime final text is the joined deltas * whisper: word timestamps from the alignment heads (median filter, DTW, word grouping) * whisper: name the align walk * whisper: teacher-forced cross-attention pass on the served decoder weights * whisper: align walk in the decoder and a float32 log-mel under the engine's cast * qwen3_asr: log-mel stays float32 under the engine's cast * whisper: run the align walk after decoding when word timestamps are requested * tests: word-timestamp alignment against HF attentions, DTW and word grouping * tests: align walk through the state machine and the decoder submodule * tests: Whisper GPU parity checks word starts against HF token timestamps * docs: speech recognition page (Whisper and Qwen3-ASR scaffold, walks, long-form, realtime, benchmarks) * docs: link the speech recognition page * whisper: select the prefill's last rows from the row lengths (a one-token language-detection prompt is planned as a decode step) * tests: decoder prefill selects each request's last row itself * benchmark: realtime bench keeps the final transcripts * tests: word-timestamp parity reads HF's output as a mapping * tests: word-timestamp parity compares starts with HF's previous-token ends * docs: serve the ASR API process with one OpenMP thread * docs: serve the ASR API process with a passive OpenMP wait policy * benchmark: keep the transcript text from vLLM's realtime chunk lines * whisper: deliver the decode loop's tokens in one message at the end * tests: the Whisper decode loop accumulates its tokens for the client * qwen3_asr: deliver the decode loop's tokens in one message at the end * docs: ASR tokens reach the client once per window * benchmark: reword a data split comment * docs: plainer wording on the ASR page * docs: plainer wording in the ASR model entries * qwen3_asr: reword config comments * qwen3_asr: reword comments and the audio count error * qwen3_asr: reword submodule comments * registry: reword the ASR entries * whisper: reword encoder comments * whisper: reword a timestamp rule comment * whisper: reword the walk comments * whisper: reword submodule comments * whisper: reword model comments and the audio count error * tests: reword parity test comments * tests: reword timestamp test docstrings * tests: reword Whisper test comments * tests: match the new audio count message * whisper: use PagedKVConfig after the engine rename * qwen3_asr: use PagedKVConfig after the engine rename * components: check the ASR request kwargs once * whisper: prompt length per walk, state keyed by the handle, greedy detection * whisper: check request kwargs and the concurrency cap in the config * qwen3_asr: check request kwargs and the concurrency cap in the config * api: Qwen3-ASR adapter streams through stream_text * cli: keep kokoro in the OpenAI example list * configs: cap whisper_large_v3_turbo at 64 concurrent requests * configs: cap whisper_large at 64 concurrent requests * configs: cap whisper_large_tp2 at 64 concurrent requests * configs: cap qwen3_asr at 64 requests and give it 2048 KV pages * configs: cap qwen3_asr_realtime at 64 concurrent requests * docs: the concurrency caps on the ASR page * tests: Whisper state by handle, prompt length per walk, greedy detection, cap * tests: Qwen3-ASR stream_text, cap and kwargs checks * tests: the request kwargs checker * registry: keep main's waypoint and Qwen3-TTS entries beside the ASR ones * move language and task computation into the process_prompt * engine: pre_plan as a serve-config key, off for whisper turbo Pre-planning hides host planning behind GPU work on a dedicated thread, and pays for it in GIL contention. Whether that trades well depends on how much GPU work a step has: measured on one H100, both arms on one build, whisper turbo (4 decoder layers) RTFx @1 109.3 -> 121.5 with it OFF RTFx @32 440.0 -> 467.5 whisper large-v3 (32 layers) RTFx @1 56.7 -> 52.4 with it OFF qwen3-asr (28 layers) RTFx @1 81.9 -> 75.3 with it OFF orpheus (28 layers) within noise either way so the shallow decoder wants it off and the deep ones want it on --- the sign flips inside one model family, which a global default cannot express. FlashInfer's `plan()` runs once per step, not per layer, so turbo pays the same planning cost as large-v3 with an eighth of the GPU work to hide it in. `pre_plan:` in the serve yaml, resolved once in `Worker.__init__` and handed to the engine, which needs the same answer for its slot count. MSTAR_PRE_PLAN_SPEC still overrides it, so an A/B needs no config edit. Only the turbo config sets it; large-v3 and qwen3_asr stay on the default, where the measurement says they belong. The c=32 deltas on the deep models are inside the +-9% restart variance of this benchmark; the c=1 deltas and the direction on all four models are not. Co-Authored-By: Claude <noreply@anthropic.com> * whisper: align with the detected language The prompt walk after detection has no language in its metadata, so the align walk fell back to <|en|>. Keep the sampled token instead. * whisper: count the sink page in the cross cache 768 pages minus the sink held 63 requests, so the 64th prefill offloaded another request. * asr: bound seed to int64 seed=2**63 passed the check, then failed its batch in the worker. * qwen3_asr: no compression-ratio retries Whisper's 2.4 bound is for 30 s windows. A 20 minute Qwen3-ASR window of ordinary speech can pass it and get re-decoded up to T=1.0. * api: stop long-form windows on disconnect * qwen3_asr: refuse <|audio_pad|> in the prompt A placeholder in the context took the audio's slot and failed the request with a 500 after the encoder ran. * whisper: keep the align walk within 448 positions A timestamped decode can leave 444 text tokens, one more than the align prompt has room for. * docs: Qwen3-ASR windowing, ASR route list * ci: run the ASR tests in CPU Core * shm: don't send from reused host buffers The stop check's host copies are views of pinned buffers the next step reuses, and ASR loops send their tokens when the loop ends (WER 95 %). The worker stops passing them; the store clones any it is given. * shm: drop host copies after the step's sends Not passing them cost Qwen3.5-0.8B 10 % at c64: per-token sends copied from the device again. Keep them for the step that made them; a later send, such as a loop's accumulated tokens, copies from the device. * whisper, qwen3_asr: batched decoder outputs and check_stop_batched Row outputs and one stop-buffer D2H per step instead of a clone and an .item() per request (the Qwen3.5 decoder's path). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * kv, cross attention: reuse a read-only step's plan; ASR split-worker configs A step whose KV segments are all span 0 (a decoder attending the encoder context it wrote at prefill) plans the same as the last one while the segments match and nothing touched the streams; the cross-attention manager skips rebuilding indptrs for the same context plan over the same rows. Whisper turbo on one H200, split+MPS: c=1 +16.5%, c=8..64 within noise. configs/*_split.yaml put the audio encoder and the decoder in separate workers on one GPU; under MPS, turbo c=8/32/64 +26/+42/+32% (a decode step no longer queues behind an encoder batch). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: capture the encoder at every batch size up to 8 At serving concurrency an encoder step is small (2-3 typical) and pads to the next captured size. Turbo, split+MPS, one H200, 4 rounds each: c=32 777 -> 826 RTFx (+6.3%), c=64 834 -> 956 (+14.6%). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: compile the encoder before capture At serving concurrency the encoder is most of the GPU time. Turbo, split+MPS, one H200, 2 rounds each: c=1 +3.6%, c=32 837 -> 914 RTFx (+9.1%), c=64 939 -> 1124 (+19.7%). The first start on a machine autotunes each capture bucket (over 20 min cold); with the Inductor cache ~1 min. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * kv: write K and V with one Triton kernel per layer KVCache.write_tokens scattered K and V with two advanced-index assignments per layer; at decode width their index math dominates the copy. One kernel copies each token's K and V rows into their (page, slot), reading strided k/v slices directly. The index-put path stays for non-CUDA tensors and unusual layouts. Qwen3-ASR 1.7B, H200, split encoder/decoder + MPS, 2 rounds: RTFx c1 95->103 (+8.4%), c8 472->516 (+9.4%), c32 1157->1219 (+5.4%); WER unchanged (1.41%). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_asr: fuse per-head q/k RMSNorm and RoPE into one kernel Qwen3Attention runs Qwen3's q/k RMSNorm and 1D RoPE as one Triton launch over every q and k head (mstar::qk_norm_rope_inplace), in place on the projected q and k, instead of two RMSNorm kernels and FlashInfer's RoPE per layer. Behind a custom op so inductor functionalizes the in-place write; the bare Triton kernel's mutation path read past a compacted buffer in the compiled decoder graph. Opt-in per model (Qwen3DenseLM fused_qk_norm_rope), on for Qwen3-ASR, and only when the position config is plain RoPE over the full head dim. More accurate in bf16 than the unfused path (no rounding between norm and rotation). Qwen3-ASR 1.7B, H200, split + MPS, same-run base, 2 rounds: RTFx c1 102->115, c8 515->569, c32 1232->1270; WER unchanged (1.41%). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * cross attention: plan one-token rows on the decode kernel A decoder's cross attention queries one token per request every step against the whole encoder context; it ran on the paged prefill kernel. Rows of one token each now plan a FlashInferDecodeWrapper (CUDA-core kernel for few query heads per KV head, as Whisper's MHA), prompt prefill keeps the prefill wrapper. FlashInferDecodeWrapper takes use_tensor_cores (default unchanged). Whisper turbo, H200, split + MPS, 2 rounds: RTFx c32 908->937 (+3.2%), c1/c8 within noise; WER unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * kv: no write addressing for a label that writes nothing A label whose segments are all span 0 (a decoder attending the encoder context it wrote at prefill) still uploaded its indptrs and ran the per-token page/offset unpacking, about a dozen small ops, every step, for zero tokens. It now gets an empty plan state; nothing reads the device indptrs otherwise. Whisper turbo, H200, split + MPS, same-run base, 2 rounds: RTFx c1 147->158, c8 495->543, c32 926->950; plan per decode step at c8 1.06->0.80 ms; WER unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * attention: CUDA-core decode kernel for MHA and narrow GQA FlashInfer's tensor-core decode pads each KV head's query rows to a tile and plans like a prefill; with few query heads per KV head the CUDA-core decode kernel and its lighter plan do better. decode_tensor_cores() picks tensor cores from a group size of 4 up, for self-attention decode wrappers and the cross-attention decode wrapper alike. Whisper turbo (MHA), H200, split + MPS, 2 rounds: RTFx c32 944->970 (+2.7%), c1/c8 within noise; WER unchanged. Qwen3-ASR (group 2) was within noise either way. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_asr: capture the AuT encoder stack The audio encoder's 24 layers and projector ran eagerly; at ASR clip lengths (13 tokens a second) its GEMMs are a few dozen rows, so the stack is launch-bound: a 7 s clip's stack is 6.3 ms eager and 1.9 ms replayed. It is now a piecewise packed region (AuTEncoder.encode) captured at batch sizes 1/2/4/8 and packed token buckets up to 4096; the conv frontend stays eager, and every window of every request is a segment of the replay's ragged attention. Batches past the largest bucket run eagerly. Padding to a bucket leaves the real rows bit-identical. Qwen3-ASR 1.7B, H200, split + MPS, 2 rounds: RTFx c1 115->124 (+8%, p50 68->63 ms), c8/c32/c64 within noise; WER within the unique-audio spread. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * submodule: cache get_device prepare_inputs calls get_device per request per step, and next(self.parameters()) walks the module tree each time (about a dozen walks per Whisper decode step at c=8). The device is cached on the submodule and dropped by _apply, so .to / .cuda / to_empty keep it right. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: capture the prefill; kv: KVStep.eager_write Whisper's prefill ran eagerly: four decoder layers of launches plus the logit rules, ~15% of turbo's decoder steps at c=8 and a quarter of their host time. It could not be captured as it was, because the step also writes each request's 1500-token encoder context, far more than the prompt tokens a packed bucket is sized by. KVStep.eager_write marks a cache whose planned tokens are written outside any captured graph: the KV manager keeps that label's addressing as an ordinary plan state instead of staging it into the slot's static buffers (a pre-plan records an event the writer waits on), and write_kv writes only the rows it is given, so a replay's padding rows (planned last, on the sink page) are skipped. Whisper now writes the context in preprocess, after the plan and before the replay, and captures the PREFILL walk as a packed graph over the prompt (batch sizes 1-32, buckets of 3-8 tokens a row, not compiled); the packed prefill picks each request's last row from the plan's offsets. Whisper turbo, H200, split + MPS, same-run base, 2 rounds: RTFx c1 157->165, c8 519->622 (+20%), c32 936->1093 (+17%), c64 1151->1231 (+7%); WER unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_asr: decode tokens loop back on the device and go to the client inline The decode loop routed two tensors per request per step (new_token into an accumulated cache, text_inputs back to the LLM), and at the end staged each request's tokens into shared memory one tensor apiece for the API server, which acknowledged each before it was freed. At high concurrency that per-row, per-step bookkeeping held the decoder's main thread. The loop now routes one signal, the sampled token, back to the LLM and to the client (Qwen3.5's shape): the node reads it off the sampler's slot masters (device_loopback_signals) and the worker sends the batch's tokens inline, one frame a step (inline_client_signals), so a decode step stores no tensor per request. The transcription handler joins the text chunks as before; prefill's first token is still persisted. MSTAR_DEVICE_LOOPBACK=0 keeps the tensor path. Qwen3-ASR 1.7B, H200, split + MPS, same-run base, 2 rounds: RTFx c64 1464->1667 (+14%), c32 1304->1347, c1 123->125, c8 flat; WER unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: decode token and rule state loop back on the device; tokens to the client inline Whisper's decode loop routed three tensors per request per step (new_token into an accumulated, persisted cache for the client and the align walk; text_inputs and ts_rules back to the decoder), and at the end staged each request's tokens into shared memory one tensor apiece. On turbo's short decode step that per-row bookkeeping was a third of the main thread. The loop now routes the sampled token back to the decoder and to the client inline (one frame a step), and the timestamp rule state back to the decoder, both kept on the device: the token off the sampler's slot masters, the rule state in the decoder's by-slot table (gathered in preprocess, written back in the captured step through the sampler's static slot index row; the new SamplerResource.loopback_slots). A row whose last step routed them (the first decode step, after the prefill) overrides its slot. The align walk teacher- forces the persisted first token plus the tokens the stop check recorded per request, stopping at the step that ended the loop. MSTAR_DEVICE_LOOPBACK=0 keeps the tensor path. Whisper turbo, H200, split + MPS, same-run base, 2 rounds: RTFx c1 165->178 (+7.6%), c8 616->661 (+7.3%), c32 1098->1189 (+8.3%), c64 1235->1250. Side by side on 79 requests (plain, segment and word timestamps, language detection, long-form): the same transcripts, timings and languages as base, up to the two near-ties base flips between its own runs. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_lm: fuse each residual add into the next norm Qwen3DenseLM ran DecoderLayer's residual adds as their own kernels before each RMSNorm. The adds are now deferred into the next norm, one in-place FlashInfer fused add+RMSNorm (RMSNorm.forward_residual, as Qwen3.5's language model does): 56 fewer launches a decode step on Qwen3-ASR's 28 layers. Qwen3-ASR 1.7B, H200, split + MPS, same-run base, 2 rounds: RTFx c1 125->129, c8 586->602, c32 1375->1394, c64 within noise; WER unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_asr: write K/V in the fused q/k norm + RoPE kernel Qwen3Attention now stores the normed, rotated k and the raw v straight into the paged cache from the q/k norm + RoPE kernel (mstar::qk_norm_rope_kv_write, a third group of programs copies v) and calls attention on the cache, instead of a separate KV-write launch per layer; KVManager.write_target hands it the layer and this step's (page, slot) addressing, padding rows still landing on the sink page. A backend that takes fresh K/V (dense) keeps write_kv. Qwen3-ASR 1.7B, H200, split + MPS, same-run base, 2 rounds: RTFx c1 129->132, c8 flat, c32 1300->1346, c64 1544->1728; WER unchanged. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * API server: wake non-streaming handlers on completion instead of a 5 ms poll collect_results slept in 5 ms steps between checks of the completion event, so every non-streaming response (transcriptions, chat, images, speech) sat for 0-5 ms after it was done. It now waits on the same loop-side event the streaming path uses, set by the message thread on completion or error; a 50 ms timeout remains for the disconnect and deadline checks. Its per-chunk wake is turned off, since only the completion matters to it. The message thread also prunes completed requests after the pass's chunks are drained rather than before, so a request whose completion and final chunk land in the same pass is released in it. Qwen3-ASR split+MPS, H200, 2 rounds: RTFx c1 133->140, c8 608->626, c32 1412->1446; turbo split+MPS c1 181->190, c32 1146->1190. Request profile 'last chunk -> finish' p50 3.8 ms -> 0.1 ms. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * Conductor and API message thread: wake on a message instead of a fixed 1 ms sleep Both loops slept 1 ms after every pass, so each hop through them (request ingest, the completion back to the API server, an inline token's round trip through the data worker) waited out the rest of that millisecond. They now block in the communicator's poll_for_messages for at most the same 1 ms, ended early by a message. The API server's thread also gets result chunks from the data worker thread, which now signals an eventfd the thread polls on whenever it queues one. H200, split+MPS, 2 rounds: Qwen3-ASR RTFx c1 138.5->144.5, c8 621->620, c32 1432->1453, c64 1811->1880; turbo c1 195->203, c8 692->693, c32 1210->1216, c64 1279->1326. Qwen3-ASR request profile 'conductor ingest -> first chunk' p50 12.0 -> 10.5 ms. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: stub communicator answers poll_for_messages; zip strict in the fused q/k test The conductor loop now waits in poll_for_messages, which the ingest test's stub communicator did not have. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * Deployment YAML: mps: true runs the workers under a conductor-owned MPS daemon Workers that share a GPU (rank_devices) time-slice it unless they are MPS clients. Until now that meant starting nvidia-cuda-mps-control by hand before mstar serve. With mps: true the conductor starts a control daemon in a private pipe directory before spawning workers, passes its client env to each worker before CUDA initializes, and quits it at shutdown. The Whisper and Qwen3-ASR split configs set it. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * API server: wake non-streaming handlers on completion instead of a 5 ms poll collect_results slept in 5 ms steps between checks of the completion event, so every non-streaming response (transcriptions, chat, images, speech) sat for 0-5 ms after it was done. It now waits on the same loop-side event the streaming path uses, set by the message thread on completion or error; a 50 ms timeout remains for the disconnect and deadline checks. Its per-chunk wake is turned off, since only the completion matters to it. The message thread also prunes completed requests after the pass's chunks are drained rather than before, so a request whose completion and final chunk land in the same pass is released in it. Qwen3-ASR split+MPS, H200, 2 rounds: RTFx c1 133->140, c8 608->626, c32 1412->1446; turbo split+MPS c1 181->190, c32 1146->1190. Request profile 'last chunk -> finish' p50 3.8 ms -> 0.1 ms. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> (cherry picked from commit 700286728026c4f9758636d81695059ad3888c1c) * Conductor and API message thread: wake on a message instead of a fixed 1 ms sleep Both loops slept 1 ms after every pass, so each hop through them (request ingest, the completion back to the API server, an inline token's round trip through the data worker) waited out the rest of that millisecond. They now block in the communicator's poll_for_messages for at most the same 1 ms, ended early by a message. The API server's thread also gets result chunks from the data worker thread, which now signals an eventfd the thread polls on whenever it queues one. H200, split+MPS, 2 rounds: Qwen3-ASR RTFx c1 138.5->144.5, c8 621->620, c32 1432->1453, c64 1811->1880; turbo c1 195->203, c8 692->693, c32 1210->1216, c64 1279->1326. Qwen3-ASR request profile 'conductor ingest -> first chunk' p50 12.0 -> 10.5 ms. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> (cherry picked from commit f3886895a00b273ed3c0feeaa06fc07b749a3403) * whisper: count the decode captures' padding pages in the startup check; tests after the main merge Main's #518 checks max_concurrent_requests against the decoder caches at startup, counting 4 self-attention pages per request. This PR captures the decode step, and its padding rows hold a page each, which is why the default self-attention cache has 2 more per request: the check counts them too (_cache_pages per resource), or it would accept a cap the cache can't serve. Tests: main's KV-sizing tests use this decoder's prefill batch cap and the padded sizing; this PR's own validation test goes (main's file covers it); the startup-config sweep builds Qwen3-ASR without weights like Whisper; the drain test's tensor-manager stub answers drop_host_copies. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: register the woken request under its submission key; stub communicator answers poll_for_messages Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_asr: import order after the accelerator-graph rename Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_asr: collapse repetition loops in the transcript, as the reference SDK does Review of #318: Qwen3-ASR had no repetition guard; the qwen-asr SDK runs detect_and_fix_repetitions (threshold 20) on every raw transcript in parse_asr_output. The adapter now does the same, before it splits off the language tag, as the SDK does. mstar/model/qwen3_asr/repetition.py is a port (Apache-2.0, qwen-asr 0.0.6): the same output as two regexes instead of a Python loop per character, which took 126 ms against 10 ms on a 19k-character transcript, on the API server's event loop. test_qwen3_asr_repetition checks it against the SDK function, copied verbatim, on known cases and 3,000 random strings. The decode itself still runs to max_output_tokens on a loop; this fixes the text the client gets, as the SDK does. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * transcriptions: parallel long form decodes the tail a timed window left unfinished Review of #318: long_form=parallel with verbose_json had 3.6-3.9% WER on the 605 s file, against 1.49%. Decoding with timestamps, Whisper ends a window by writing the next segment's start (and maybe a few words) and stopping, leaving that speech to the next window, as openai-whisper's seek loop expects. Sequential mode does re-hear it (the next window starts at the last closed segment's end); parallel mode's next window starts at the fixed cut, so it was lost. Plain json decodes without timestamps and is not affected. Now a parallel window that stopped inside a segment keeps its closed segments, as in sequential mode, and the audio from the last of them to the cut (WindowPlanner.cut) is decoded as one more window, all such tails at once after the first pass, and merged after the window they came from. A seek point under MIN_SEEK_SECONDS or a tail under half a second is left as it is. Streaming parallel (untimed text) is unchanged. Turbo on one H100, the 605 s file, before -> after: parallel verbose_json WER 2.57% -> 1.76%, the same as sequential verbose_json (1.76%); parallel json 1.49% either way. The segment gap the base left at 111.5-115.7 s (a window's dropped tail) is gone; the gaps left are the ones sequential mode has too. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: the tick-exception stub communicator answers poll_for_messages (the conductor's message wait) Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * ci: test/asr needs a GPU and real weights; not part of CPU Core Main's #581 requires every test module to be collected by CPU Core or excluded with a reason; this PR's test/asr parity tests are neither, which fails test/ci's selection check. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * transcriptions: no tail pass once the client is gone; streamed parallel windows decode untimed - The parallel tail pass ran even when the first pass stopped early because the client disconnected, submitting up to PARALLEL_INFLIGHT windows that were aborted at once. It now runs only once every window is in. - A streamed response carries text only, but a timed response_format still switched timestamps on for every parallel window, which can stop early and leave its tail to the next window (the tail pass cannot run for a stream). Streamed parallel windows now decode without timestamps; a single window is unchanged. - The tail list is a NamedTuple. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: back to main's decoder-cache sizing; padding rows hold no pages The two extra self-attention pages per request covered decode-capture padding rows, which on this engine hold none: KVManager skips padding rows at admit and commit, they read and write SINK_PAGE, and the graph runner frees them after every step and after capture. They cost ~2.5 GiB on large-v3 and made the startup check refuse a YAML sized as main sizes it. The turbo YAML's comment now matches large-v3's, and the shipped-YAML test covers turbo. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_asr: collapse repetition loops in the finished transcript only The fix ran in parse_transcript, which a realtime step also uses to hold back the unstable tail (the parse of the whole hypothesis minus the parse of its stable prefix). Collapsing a loop that straddles that boundary can make the whole shorter than its prefix and send unstable text as final. It now runs in the adapter's final_text, on the text of a finished transcription (the response, or the stream's transcript.text.done); partial text is left as the model wrote it. On the text after the language tag rather than the raw stream: a loop cannot span the tag. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * API server, conductor: guard the waits between passes; a testable wake backstop A poll can raise where a sleep could not (a closed wakeup fd, a terminated ZMQ context): the API message thread's wait is inside its loop's try, and the conductor logs and backs off. WAKE_BACKSTOP_S names the 50 ms a waiting handler sleeps at most, so the completion-wake test sets it to 5 s and fails (4.8 s) when the wake does not fire. Comments state what the waits do. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * docs: Qwen3-ASR long-file capacity, Whisper's cap check, the pre_plan serve key - A 20 minute Qwen3-ASR file holds ~122 pages for its prompt and up to 32 for its transcript, so 2048 pages decode about 13 at once, not 16; more concurrent long files run out of pages and wait on each other. - Whisper refuses a cap its caches cannot hold (an unset cap only warns). - pre_plan: in the serve config table; MSTAR_PRE_PLAN_SPEC overrides it. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: worker-metrics run-loop fake sets _enable_pre_plan, which this PR's Worker resolves in __init__ Main's new test builds the Worker without __init__ and turns pre-planning off with MSTAR_PRE_PLAN_SPEC=0; the run loop here reads the resolved _enable_pre_plan (the serve yaml's pre_plan: with the env var overriding). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * conductor: check mps: true against the device where the daemon starts Main's startup-config sweep builds every shipped YAML through Conductor.__init__ on a CPU host with worker launch stubbed; the split configs' mps: true failed it there. __init__ keeps the type check; the CUDA check and the daemon (and its temp directory) come with the worker launch. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: conductor launch stand-ins carry the mps attributes _launch_workers reads Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_asr: splice the audio embeddings in by slice, not by a device mask Prefill found the <|audio_pad|> placeholders on the device: int(slots.sum()) and a boolean-mask assignment each wait for the encoder's queued work, once per prefill row, in prepare_inputs on the GPU thread (AGENTS invariant 7). process_prompt already knows where it put the run, so it returns the start and the count as metadata (INTERNAL_KWARGS, so a client cannot set them), the prefill's step_metadata carries them, and prepare_inputs compares the count with the encoder output's shape on the host and assigns by slice. Tests (written, not run locally): the metadata matches the placeholder run with a forced language, free language and a realtime prefix; prefill's step_metadata carries it; prepare_inputs splices by it. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * request_kwargs: an integer field that overflows (inf, 1e400) is a ValueError naming the field int(float('inf')) raises OverflowError, which escaped the check's field-naming path as an unhandled error. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: numba is part of the asr extra; warn without it, compile at load Word timestamps' DTW was compiled by numba only when it happened to be installed (the qwen3_omni and all extras), and fell back to a plain-Python double loop on the GPU thread without a word. numba joins the asr extra, a missing numba is logged, and the decoder compiles the DTW when it loads instead of inside the first word-timestamp request; a compile failure falls back to Python with a warning. Test written, not run locally. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * asr benchmarks and parity tests: data under data/asr, not an undocumented MSTAR_WS_ROOT The four ASR benchmark scripts and the GPU parity tests defaulted to $MSTAR_WS_ROOT/commons/bench/data/asr, a workspace path behind a variable no doc mentions. They default to data/asr (the reference SDK to refs/Qwen3-ASR), and the parity tests' own variables have rows in environment_variables.rst. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * worker: resolve_enable_pre_plan, tested directly; a non-boolean pre_plan is refused The pre_plan tests re-implemented Worker.__init__'s rule and grepped the source to keep the copy honest. The rule is now a module function the worker calls and the tests call. It also refuses a pre_plan that is not a YAML boolean: bool("false") was true. Tests written, not run locally. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_asr: ASR_TEXT_TAG defined once, in a module the API server imports without the model Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: word timestamps on a TP-sharded decoder are a 400 at ingest The align walk re-runs the decoder unsharded, so with tp_size > 1 a word timestamp request failed with NotImplementedError after its whole decode. validate_config_yaml records the decoder's tp_size and get_initial_forward_pass_args refuses the request with a ValueError, which the conductor answers with a 400. Test written, not run locally. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * transcriptions: submit a parallel window's tail as soon as that window is in The tail pass waited for every first-pass window before submitting any tail, adding a whole serial window decode to parallel mode's latency. A window's tail now goes to the engine when that window's text arrives, beside the windows still decoding; a client that leaves releases it with the rest. Tests written, not run locally. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_asr: the per-request window cap follows max_audio_seconds instead of a constant MAX_WINDOWS_PER_REQUEST = 160 restated '20 min / 8 s' by hand; it is now computed from max_audio_seconds, the hop and n_window_infer (150 for 20 min), so the two cannot drift. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * docs: Whisper delivers a window's transcript at loop end; qwen3_asr_realtime in the clients table The Whisper model docstring still said the transcript is streamed token by token; the clients table left out qwen3_asr_realtime; models.rst called the 0.6B checkpoint 'trained for chunked streaming', which nothing here checks. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper, qwen3_asr: check_stop_batched takes the engine's step_walk Engine.check_stop_for_batch calls check_stop_batched(..., step_walk=...) since #362 (merged in from main). The ASR decoders' overrides predate it and do not take the keyword, so every call raised TypeError, was logged at debug, and the engine fell back to per-request check_stop: one .item() per row per step instead of one tolist. Both paths give the same stops, so this is overhead, not a correctness bug. The walk is still read per row, as before. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * check_stop_batched: one shared row matcher, and a test through the engine HostRows.first_values pairs each request with the first element of its row of a host buffer, with the forward-order fast path and the row map for a batch listed in another order. Qwen3.5, Qwen3-ASR and Whisper read their stop tokens through it instead of three copies of the same matching. Qwen3-ASR and Whisper also return early on a step of another walk, as check_stop does. New tests drive Engine.check_stop_for_batch with host rows for both ASR decoders, with the per-request check_stop replaced by one that fails, so a signature mismatch can no longer fall back silently. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * api: decode text chunks as one byte stream in /v1/realtime and chat Qwen3-ASR emits one byte-level BPE token per chunk, so a character's UTF-8 bytes can span chunks; decoding each chunk on its own turned it into replacement characters. A realtime step now joins the bytes and decodes once, as serving_transcriptions._text_of does. Chat does the same for a non-streamed reply and streams through an incremental decoder that holds a split character back until its last byte (a no-op for models whose chunks are already whole characters). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * attention: decode_tensor_cores is a declared flag, not a head-ratio heuristic The CUDA-core decode kernel was chosen for every model whose query-to-KV head ratio was under 4, which changed the kernel for models it was never measured on. AttentionConfig and CrossAttentionConfig now carry decode_tensor_cores (default True, main's kernel), settable under resources: in a deployment's YAML. FlashInferManager, the mixed wrapper's decode rows and FlashInferCrossManager read it; the heuristic is gone. Whisper (self and cross attention) and Qwen3-ASR's LLM declare False, the models it was measured on. A CPU test checks the flag reaches every decode wrapper the managers build. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * cross attention: decode kernel only when every row queries one token The kind was chosen by qo_indptr[-1] == rows, so spans [2, 0] took the decode kernel. It now follows FlashInferManager: not force_prefill, and every span is one. A captured key asserts its wrapper's kind matches the step, and the eager key is EagerSlotKey(is_decode=...) rather than a tuple. The two plan caches (_plan_inputs, _plan_keys) are one record per wrapper key: the same context plan object over the same rows skips building the arrays, and the same pages skip FlashInfer's plan. The docstring no longer quotes a model's measurement. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * sampler: declared loop-back state, and one routed-row override Whisper kept its timestamp rule state in a model-allocated table indexed by raw sampler slots (loopback_slots / slot_capacity), sized once while the sampler's masters grow, and never reset when a slot was reused. The state now lives in the sampler: SamplerSpec(loopback_state={name: LoopbackState(width, dtype, default)}). Its masters grow with the others and a (re)registered slot reads the default. A step reads it with loopback_state(name, rids) and the forward stores it with store_loopback_state(name, rids, values): under a lease into the slot's per-step row (captured with the step), persisted by commit when the SamplerStep lists it in loopback_writes; an eager step by request. loopback_slots and slot_capacity are gone. The three loop-back readers (Qwen3.5, Qwen3-ASR, Whisper) put a routed row's value over the one read by slot through SamplerResource.overwrite_rows, which sends the row index through a PinnedStager with a non-blocking copy instead of a synchronous torch.tensor(..., device=...). Same rows, same values. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * kv: read-only plan reuse survives what does not touch its rows The read-only plan was invalidated by every stream touch and every commit that wrote anything, so the graph runner's reset of a replay's padding rows dropped it after each padded step, as did any other request's ingest or prefill. The manager now keeps the cached step's real rows and bumps the read-only epoch only for a touch of one of them (or of every stream), or a commit that writes one of them, mirroring the decode plan cache's test. CPU tests cover a hit on an identical step, after a padding-row reset and after another request's write; a miss after a real-row write, reset, removal with the handle recycled, and a global touch (offload/reload on a GPU); staging into another slot; and an eager_write step's plan staying out of the static buffers, with write_target addressing it. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: the align walk must run where decode ran The align walk reads the tokens the decode walk recorded in the decoder's worker-local request state. validate_config_yaml now refuses a config whose decoder node serves decode and align on different ranks, and _prepare_alignment raises when the tokens are missing instead of aligning the persisted first token alone. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: drop test_kv_scatter_triton.py, which repeats main's test_kv_scatter.py Its two cases main's file lacks (int32 and empty indices, a captured replay reading its index buffers) move there. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: cross attention plan tests in the existing test_cross_attention_plan.py Moves the span, force_prefill, plan-skip and captured-kind tests into the file that already stubs the cross manager's wrappers. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * utils/attention: one @triton.jit helper for the per-head norm + RoPE _qk_norm_rope_kernel, _qk_pair_norm_rope_kernel and _qk_norm_rope_kv_kernel each restated the same RMSNorm and rotate-half RoPE; they now call _rms_norm_rope, which Triton inlines (same operations in the same order). The custom op's comment keeps the constraint and drops the debugging story. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_lm: the position resource says when RoPE may be fused Qwen3Attention read the position resource's private _config through getattr to decide whether its fused norm+RoPE kernel can stand in for apply_qk. PositionManager.fusable_rope_theta(head_dim) now answers that (plain rotate-half RoPE over the whole head dim, no scaling, interleave, RoPE dtype or Llama-3.1 factors) and returns the theta the kernel uses. The single-caller flag Qwen3DenseLM(fused_qk_norm_rope=...) is gone: the layer is always Qwen3Attention, which keeps the unfused path off CUDA or when the resource says no. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * mps: a hung control command or a dead conductor leaves no daemon behind - A control command that times out is a failed command, so start() gives up at its deadline and stop() goes straight to the kill; stop() kills a daemon that does not quit and removes the directory in a finally, and start() stops the daemon on any failure. - On Linux the daemon is exec'd with PR_SET_PDEATHSIG, so it gets SIGTERM if the conductor dies without stopping it (when started from the main thread, which the signal follows). - The conductor checks current_platform.is_cuda() rather than the device type, which is "cuda" on ROCm too. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * submodule_base: get_device's comment says what the cache does, not why it pays Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * docs: eager_write, write_target, device loop-back state, mps, ASR streaming - adding_models.rst: KVStep's eager_write, KVManager.write_target for a layer that writes its own K/V, SamplerStep.loopback_writes, and a Device loop-back section (loopback_tokens, overwrite_rows, declared loopback_state and its lifecycle). - serving.rst: the mps config key. - asr.rst: the decode loops stream per token; the *_split configs and mps. - environment_variables.rst: reads_device_loopback is not Qwen3.5-only. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: the MPS orphan test waits up to a minute for the stand-in conductor's imports Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: find the attention specs among specs that may have no config (the sampler's) Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * benchmark: the ASR harnesses behind #568's numbers ab.py: A/B arms of one tree (base/change via a patch, runtime switches, configs, MPS) and vLLM/sglang-omni baselines on the same cells, plus the phase-timing profile pass. regress568.py: other models (Qwen3-TTS, Higgs STT, Qwen3.5) across two or more source trees given as --arm LABEL=PATH, with TTS WER through Whisper. litmus_qwen35.py: the Qwen3.5-0.8B CPU-overhead litmus. benchmark/asr318_run/: run.py's shared helpers (servers, LibriSpeech/long-form data, WER, judging) and the correctness gate it drives (gate.py, oracle.py, gallery.py). assets/ (the language-detection clips) is not included: without asr_zh.wav the gate's detect check fails, everything else runs. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper, qwen3_asr: decode lengths on main's per-request loop caps Whisper: the conductor now owns the length budget. LENGTH_KWARGS = () so a requested max_output_tokens is cut to the room the prompt leaves in the 448-position table instead of refused; since main's KV layouts (#808), a long-form upload with a length failed at window 2, whose <|startofprev|> context pushed prompt + length past 448. check_stop's own positions_left budget and the prompt_len / detect_len request state go: they ended decode one iteration before the cap, and a second owner for the same limit. Qwen3-ASR: the decode loop runs to the context (max_position_embeddings), as every other decode loop does since #808; max_new_tokens stays the default cap, so a longer max_output_tokens is served rather than refused at 4096. Tests: Whisper's asked length is cut to the context; Qwen3-ASR's cap matches its stop check and fits its loop, and its prompt span matches the prefill. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: split zh/ja/th/lo/my/yue words per character, as openai-whisper Word timestamps grouped tokens only at spaces and punctuation, so a zh/ja sentence came back as one word. Split by character in the languages openai-whisper does, and end a piece at a replacement character the transcript itself holds instead of merging the rest. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_asr: decode tokens go to the client inline, not accumulated to the loop's end From #568: b51fdb8f5 and the qwen3_asr half of fa6fcb16a (batched decoder outputs and check_stop_batched), applied without conflicts on 246d285c9. The decode loop sent its tokens through Loop.accumulated_outputs: at the end of the loop each request's tokens went to the API server as one tensor apiece, ~3,450 for a 20-minute window. In a 16 x 20-minute run on #318, 8 of 16 requests failed with "result delivery timed out" after their compute had finished: the API process read and detokenized those tensors too slowly to meet its 15 s delivery deadline (27,626 late chunks logged afterwards). The loop now routes one signal, the sampled token, back to the LLM and to the client: the node reads it off the sampler's slot masters (device_loopback_signals) and the worker sends the batch's tokens inline, one frame a step (inline_client_signals), so tokens reach the API server while the window decodes. MSTAR_DEVICE_LOOPBACK=0 keeps the tensor path. docs/asr.rst now says Qwen3-ASR streams per step and Whisper still per window. Measured on #568 (Qwen3-ASR 1.7B, H200, split + MPS, same-run base, 2 rounds): RTFx c64 1464 -> 1667 (+14%), c32 1304 -> 1347, c1 123 -> 125, c8 flat; WER unchanged. Not yet re-run at 16 concurrent 20-minute files. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * docs/asr: serve ASR on the shared-memory arena Vasilev's Qwen3-ASR 1.7B run on #318: RTFx 81/218/350/412 on the file transport vs 87/252/494/558 on the arena at c=1/8/32/64, same 1.41% WER. MSTAR_SHM_ARENA is deployment-wide and env-only, so the ASR page says so. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * docs/asr: drop the unverified thread-pool hangs claim It was a correlation from a few runs, and main's own measurement points the other way: capping the API server's torch threads fixed Qwen3.5's preprocess stalls (PREPROCESS_TORCH_THREADS). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper, qwen3_asr: cap the API server's torch threads at 4 (PREPROCESS_TORCH_THREADS) Qwen3.5's declaration for the same problem: torch's one-thread-per-core pool spins between the API server's small CPU ops and starves its event loop. Pending an A/B against OMP_WAIT_POLICY=PASSIVE before the docs' manual launch step goes. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * docs/asr: Qwen3-ASR long-file capacity from a measured run The 13-files-on-2048-pages estimate didn't hold: Vasilev's 16 concurrent single-window 20 min files peaked at 1563/2048 pages and failed on result delivery, not pages. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * api server, conductor: drop the early wake between message-loop passes The message thread and the conductor go back to main's 1 ms sleep between passes. The data worker's per-pass results_wakeup signal cost Qwen3.5-0.8B 3-5% at c=64 (Vasilev's alternating-boot A/B on #318; no difference with it removed), against an unmeasured saving of under 1 ms per result. #837 on main already wakes collect_results on completion, the part that fixed the high-concurrency delivery timeouts. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: merge punctuation as openai-whisper's merge_punctuations Opening marks first, right to left, then closing marks, compared unstripped; each word keeps its own timing. Before, closing marks went first on stripped text, so [" Hello", ",", ' "', "world"] gave ' Hello, "' + 'world', and a merge moved the word's bounds onto the mark's. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test/asr: compare the served model's greedy text to the reference The Qwen3-ASR end-to-end test only ran the reference SDK and asserted its text was non-empty, and Whisper's "greedy tokens" test only ran HF generate the same way, so neither could fail on an M* regression. Both now start `mstar serve` with the shipped config and transcribe the benchmark's first 8 files (Whisper: those within one 30 s window), greedy with the language forced, against the reference run in the test process first (the qwen_asr SDK, HF generate), then freed so the server has the GPU. At most one file may differ, by at most 2 words: bf16 near-ties flip the odd homophone (9 of the 200 benchmark files differ), a regression breaks most files. test/asr/_served.py holds the shared pieces: the benchmark utterances, the server (own process group, killed on exit), the multipart request and the parity check. Whisper's unused _DenseKV stub goes, and its fixture now loads the checkpoint MSTAR_WHISPER_REPO names instead of always turbo. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: stage prefill's rule state and row lengths through a pinned buffer torch.tensor(<list>, device=) is a blocking copy that waits out the stream: Vasilev's sync audit on #318 counted 390 at prepare_inputs' ts_rules and 291 at preprocess' row_lens over 256 requests. Prefill rows now keep their rule state on the host and preprocess sends the batch's rules and row_lens in one non-blocking PinnedStager copy each; the rules-off decode state is a device constant built once. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_asr: check_stop_batched takes the engine's step_walk Engine.check_stop_for_batch passes step_walk=; the port in cf65a7525 took check_stop_batched from before #568 added the parameter, so every batched stop check raised TypeError and fell back to the per-request path (logged at debug). Same signature and early return as Qwen3.5 and #568. The test now calls it as the engine does. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper, qwen3_asr: fixes from the clean review of this round - split_words: a transcript cut inside a character (at the position table) read past the full decode (IndexError) or dropped its last tokens; guard the lookup and keep the tail as a piece. The align walk decodes without HF's clean-up, which rewrites text across token boundaries, so the pieces add up to the full decode as with openai-whisper's tiktoken. - Whisper preprocess stages only rule-state rows still on the host, so rows already on the device (#568's captured prefill template) pass through. - /v1/realtime decodes a step's text chunks as one byte stream (from #568's 1870aa921): Qwen3-ASR now emits one byte-level token per chunk, so a character split across chunks became U+FFFD and fed the next prefix. - Long-file sizing: docs, YAML and docstring agree again (~122 + 32 pages per 20 min file, ~13 at full size on 2048 pages), next to the measured peak. - Comments: a word in zh/ja/... is a run of tokens decoding to whole characters. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper tests, docs/asr: fixes from the GPU run; drop the OMP_WAIT_POLICY launch step - test_whisper_model's fake tokenizer takes clean_up_tokenization_spaces, which the align walk's word grouping now passes (dc5b7709a broke this CPU test). - test_whisper_parity reads process_prompt's ProcessPromptOutput by field; it indexed it as a dict and failed before reaching the comparison. - docs/asr: the thread cap replaces the OMP_WAIT_POLICY=PASSIVE launch step. A/B on a 192-core H200 host, 2 boots per arm, c=1..64: cap 4 matched PASSIVE on RTFx and latency for turbo and Qwen3-ASR; the default pool lost 4-9 % (29 % for Qwen3-ASR at c=64 with the server pinned to 16 cores). - docs/asr: 16 concurrent 20 min Qwen3-ASR files peaked at 2352 pages with a 4096-page pool and stalled on the shipped 2048 (measured on dc5b7709a). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_asr: cap the shipped config at 12 concurrent requests, and document the trade-off 16 concurrent 20 min files filled the 2048-page pool and the host pool and stalled on dc5b7709a. A 20 min file holds about 122 + 32 pages, so 12 fit any mix up to that length. docs/asr.rst gives the short-utterance throughput the cap costs (RTFx ~420 / 900 / 1280 at 8 / 32 / 64 concurrent on one H200) and when to raise it. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_asr_split: cap at 12 concurrent requests, as qwen3_asr 16 concurrent 20 min files stalled on 2048 pages; a 20 min file holds about 122 prompt pages plus up to 32 for its transcript. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * kv: restore @torch.compiler.disable on KVManager.layer_view Inserting write_target above layer_view took the decorator that was on layer_view (as on main). write_target keeps one too: like read_kv, write_slots and write_kv it reads this step's plan state, waits on the plan's event and returns the token count as a Python int, which a compiled caller would specialize on and recompile per step. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: word alignment off the GPU thread; DTW and words in the API server The align walk's forward now does only GPU work: the teacher-forced decoder rerun, head standardization, median filter and average, cropped to the rows that predict text and end-of-text and to the frames that hold audio. It emits one float32 tensor, that matrix plus a column of token ids (the text ids, and the language token on the end-of-text row); WhisperModel.postprocess tells it from token ids by dtype and runs the DTW, word grouping and punctuation merge in the API server's data worker, rendering the same <|startoflm|> + timestamp text the transcription route parses. The numba compile moves with it, to warmup_preprocess; the decoder no longer needs the tokenizer or numba. No host sync is left on the align walk: the forced prompt's language, task and audio length come on its step metadata (audio_frames from process_prompt's metadata, an INTERNAL_KWARG a client can't set, instead of a tensor edge read with int()), and the stop check, which already has each prompt walk's token on the host, keeps a word-timestamp request's detected language and first token, so the transcript edge and its .tolist()/.item() go. The prompt walks no longer write language/task into request state, and a request without word timestamps no longer records its decode tokens. The align step hands the stop check no buffers, so its matrix is not copied to the host there. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * benchmark: ASR harnesses launch without OMP_WAIT_POLICY=PASSIVE The API server's capped intra-op pool matches it (docs/asr.rst), so the harnesses serve as a deployment would. regress568's docstring no longer names the early wake, which is gone. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * benchmark: ASR harnesses serve mstar raised to their own concurrency configs/qwen3_asr.yaml (and the split config) now cap max_concurrent_requests at 12, while the 318 harness and the A/B harness drive up to 64 concurrent requests: past 12 the conductor queues them, so those cells would time the queue rather than the engine. A timed mstar server is now served from a benchmark copy of its config (the default one, or an arm's --config) with the cap raised to the most concurrent requests the run drives; a config already at or above it is served unchanged. The KV page count stays, since the benchmark utterances are short. Smoke, correctness and stress servers keep the shipped config, where queueing is the behaviour under test (on a 1-GPU box stress shares round 1's timed server, and so its raised cap). The cap each server ran with goes into the results (max_concurrent_requests on every row, and in the 318 harness's config column) and into both summaries. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * qwen3_asr: replay the encoder stack only on a leased step _replay_encode replayed whenever a bucket fit the batch. A fitting bucket is not a lease: can_lease can refuse the batch, or lease_slot find no slot, and then declare_step has declared and admitted the eager step, and a replay would run on a slot nobody holds. A leased step is the one that declares nothing (declare_step returns None exactly when piecewise_leases holds AUT_ENCODE), so the replay now also requires engine_inputs.step to be None, as Qwen3.5's vision block loop does. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: state which device each capture template's rule state lives on A replay's padding rows are stacked with the walk's real rows in preprocess, so each template's ts_rules has to sit where those rows do: on the device for decode, on the host for prefill (staged in one pinned copy). The prefill template already said so; the decode template now does too, and takes its row from _inactive_rules so the rules-off state on the device has one source. The Whisper template test pins the decode side as it already pinned the prefill's. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * kv: a step with sink rows is not reused as a read-only plan The read-only plan reuse keyed a span-0 step by its segments, slot bucket and epochs, not by KVStep.sink, so two span-0 steps with the same segments but different sink rows would share one plan, and a row the second step puts on SINK_PAGE would read its stream's pages (or the reverse). The steps that hit the reuse today (a decoder reading an encoder context written at prefill, e.g. Whisper's cross-attention) carry no sink, and the one model that declares sink rows (Bagel's unguided CFG rows) has no span-0 steps to repeat. So _is_readonly now requires an empty sink rather than the key growing a field no repeating step varies; it matches the decode plan cache (_remember_decode_plan), which also leaves sink steps alone. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: every decoder walk on the same ranks when word timestamps align Word timestamps record the transcript's first token and a detected language in the prompt walks' stop check, on whichever worker ran them, and the decode walk extends that transcript; the align walk reads all of it from the decoder's worker-local request state. validate_config_yaml only kept decode and align together, so a layout with the prompt walks on other ranks started and then failed each word-timestamp request at its align walk. It now requires prefill, detect_language, prefill_prompt, decode and align served by the same decoder ranks, naming each walk's ranks when they differ. The align walk's error now says which record is missing (the transcript or the language) and which prompt walk keeps it, and the stale comment that the decode loop's first step recorded the transcript is gone. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: record each decode step's transcript token once If check_stop_batched raises partway through a batch, Engine.check_stop_for_batch falls back to check_stop per request, and _record appended the token again for every row the batched pass had already recorded: the align walk then timed a transcript with repeated tokens. _record now appends only while the transcript holds exactly the tokens before this step (the prompt walk's, plus one per earlier decode loop iteration, from the step's own loop count), so a re-run of the same step is a no-op. No new state: the count is the step's dynamic_loop_iter_counts, which the engine keeps per executing batch. The batched-path test now starts the transcript with its prompt token, as a real request does. _record_prompt_token notes the transcript in request state is an interim, until the decode loop's outputs reach the align walk directly. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * whisper: word-timestamp cleanups - <|startoflm|> is one constant, WORDS_MARKER in whisper/config.py, which both WhisperModel's postprocess and the API server's adapter import (config.py loads no model modules). - alignment.unpack_output refuses a tensor without a frame column with a ValueError, instead of timing words off an empty matrix. - The DTW compiles with numba lazily, once, on its first use (or at warm_up, which calls the same function), so numba loads only in a process that runs the DTW; the "plain Python" warnings moved there with it. - Model.warmup_preprocess says it warms what process_prompt and postprocess need. - Tests: drop the decode capture template's device assert, which passed before that change too, and the INTERNAL_KWARGS membership assert, which restated a constant (test_model_internal_kwargs covers the data worker dropping such keys). The DTW test now pins the compile-once-on-first-use behaviour. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * benchmark: ASR harnesses raise only a cap the model accepts, read --config= Server.bench_…

… graph loop (#892) * transcriptions: windows for bounded ASR models Long uploads are cut by the adapter's clip limit, uploads with no length in their header are measured, every decode stops at the route's cap, and a window without timestamps is one segment (Higgs and untimed Whisper captions). * fastconformer: shared encoder and log-mel NVIDIA's FastConformer encoder, offline or chunked for cache-aware checkpoints, NeMo's log-mel, and the encoder node, captured per mel bucket. * parakeet: Parakeet-TDT and Nemotron ASR Greedy TDT and RNN-T decoding in captured chunks of a compiled step; RNN-T loops over labels. * canary: Canary-1B-v2 Speech recognition in 25 languages and translation to or from English. Whisper and Canary share one KV-fit check. * canary_qwen: Canary-Qwen-2.5B Canary's encoder projected into a LoRA-merged Qwen3-1.7B, English only. * components: one encoder-decoder node for Whisper, Canary Whisper's and Canary's decoder nodes share the prefill/decode flow; Whisper keeps its token suppression as a process_logits hook. * asr: review fixes on the rebased stack - TokenStepNode.check_stop_batched takes step_walk, which the engine passes since #362; without it every call raised and fell back to the per-request stop check. - The FastConformer ASR adapters stream text without SentencePiece's leading space, and the route sends a window-joining space only before a window's first text, so the deltas add up to the done event's text. - Parakeet, Nemotron and Canary-Qwen refuse a task other than transcribe instead of returning a transcript for a translation. * perf(fastconformer): plan a batch's encoder replays by cost, not one series per bucket * perf(fastconformer): log-mel features for the whole batch in preprocess * perf(fastconformer): the position term as one GEMM per head, half-step residuals in the add * perf(fastconformer): Triton kernels for the conv module and the subsampling's bias, mask and ReLU * perf(fastconformer): each residual add in the kernel of the LayerNorm after it * perf(fastconformer): refit the replay cost model to the fused encoder * docs: Omnilingual ASR gets a segment and a caption cue per window A window without timestamps is one segment spanning it (merge_windows), so the CTC model's verbose_json and captions follow its windows, as Whisper's and Higgs-Audio's untimed windows do. * perf(transcriptions): skip the thread hop for an upload prepare() leaves alone * perf(media_io): skip the channel mean for mono audio in read_soundfile_capped * Canary-Qwen: take a custom instruction from model_kwargs, not /generate's text process_prompt used /generate's text prompt as the instruction, so the ASR harnesses' default prompt ("Transcribe the speech in this audio clip.") replaced the trained one, "Transcribe the following:", and Canary-Qwen was benchmarked under an instruction it was not trained on. The text prompt is now unused, as on the other FastConformer ASR models, and a custom instruction comes from model_kwargs' "instruction". Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * Docs: untimed transcription models get one segment per window The Omnilingual ASR adapter's docstring still said srt/vtt get one cue spanning the file and verbose_json an empty segment list. merge_windows now makes each window with text one segment spanning it, so each format gets one cue or segment per window, as docs/models.rst and docs/clients.rst already say. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * media_io: decode_audio is load_speech_file, so a missing FFmpeg stays a 500 The transcription route decoded uploads with its own copy of the libsndfile-then-FFmpeg chain, which turned every torchcodec RuntimeError into a 400, a missing FFmpeg included. load_speech_file already tells the two apart; the copy and its numpy resample fallback go (torchaudio is a dependency). * fastconformer: log-mel per replay group, not padded to the batch's longest clip A 600 s Parakeet upload batched with short clips padded every clip's STFT to 600 s. preprocess now plans the replay groups and computes each group's features together (a clip past every bucket alone), and hands the plan to the encoder so it isn't computed twice. * canary, canary_qwen: declare their KV labels and lay out each request Main refuses a paged pool with no labels at engine build (#832), so neither model started. Canary follows Whisper: the decoder's cache grows a slot per decode step and the encoder context is one rows span of the clip's frames. Canary-Qwen's prompt is a rows span: its audio placeholders are the same for every clip of a length, so it must never key a cached prefix. * parakeet, nemotron_asr: out of this PR, into a follow-up on top of it Review asked for the transducer to step through the graph as a loop, with its LSTM state and cursor in a resource, rather than run the whole greedy loop inside one forward. That rework lands with the two transducer models in their own PR, stacked on this one. This PR keeps Canary, Canary-Qwen, the shared FastConformer encoder, the encoder-decoder and token-step nodes, the Whisper refactor and the transcription route changes. Removed: mstar/model/parakeet/, components/transducer.py, their configs, registry, CLI and OpenAI adapter entries, benchmark entries, docs and tests. The streamed-deltas test of the transcription route now runs on Canary, the other windowed SentencePiece model. The encoder keeps its causal path, which its own tests cover. * RecurrentStatePool: gather / scatter_ over the step's addressing A forward whose recurrent state has no backend kernel to plan against reads and writes its rows' slots through the pool instead of indexing pool.block() itself, so its eager and captured paths address the same slots and a replay's padding rows land on the sink. num_rows is the rows the step declared, padding rows included, which is what both slice by. From #383 (e9b665a, c1813dc), without its MiniCPM-o changes. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: gather / scatter_ round-trip, replay vs eager addressing, padding rows on the sink On CPU, admit / plan / commit driven by hand: a shorter eager step reads only its rows from the grown eager buffers; a replay padded to its bucket leaves the real slots as the eager forward does and writes nothing but the sink; the replay buffers are one per (bucket, capture slot) and the same storage every step; a freed slot reads as zeros. On CUDA, a graph captured over a bucket's addressing and replayed for four different batches matches the eager forward. * docs: RecurrentStatePool.gather / scatter_ are on main, not an open PR * Revert "parakeet, nemotron_asr: out of this PR, into a follow-up on top of it" This reverts commit 32c4258. Parakeet and Nemotron ASR come back as they are on #796's head (bf9f067), with the transducer's greedy loop still inside one forward. This is the follow-up's starting point; the graph loop with its state in the recurrent pool goes on top. * parakeet: the transducer as a decode loop, its state in a recurrent pool The transducer node no longer decodes a whole clip in one forward. prefill projects the encoder output once (with Nemotron's prompt) and persists each request's frames; decode is a Loop over the node that re-injects them every iteration, with a data-free clock loop-back so the node is a same-node speculation target. An iteration is 16 greedy steps: gather the rows' h, c, dec_out, cursor and symbols from a RecurrentStatePool slot, advance, scatter them back, emit the tokens (-1 where a step emitted none). check_stop ends a row's loop off the side-stream copy of its active flag, so the per-chunk host read is gone; the never-binding step bound is the loop cap and a stop-check guard. Capture moves to the engine's main runner: a BatchedAcceleratorGraphConfig per frame bucket on the decode walk, keyed by cg_key_info from the clip's frame count, which process_prompt computes and step_metadata carries. A fresh slot and a padding row start from the start token's state, computed once at load. The node asks the engine for full fp32 instead of setting it. The YAMLs size the pool at 129 slots for max_concurrent_requests 128, and startup refuses a cap the pool can't hold. * test: the transducer loop through the recurrent pool, against #796's decoder A CPU harness in the shape of test_gdn_cuda_graph.py: tiny TDT and RNN-T heads, the model's pool declaration under a YAML max_slots, a real step runner, and the node driven by hand from prepare_inputs to check_stop. Requests of 1 to 30 frames join and leave at different iterations, in both frame buckets, in eager steps and in steps padded to a capture bucket; a stopped request runs the speculated extra iteration and emits nothing. Each request's tokens equal the one-request reference loop's and a frozen copy of bf9f067's greedy_decode, for random heads, a head that never emits blank (NeMo's per-frame guard) and blank runs longer than a window. On CUDA the engine's runner captures and replays the decode iteration at 1, 3 and 64 rows, and must match eager. * docs: Parakeet and Nemotron keep their decoder state in a recurrent pool * RecurrentStatePool: gather / scatter_ refuse a pool without a sink slot With disable_sink_slot, rows with no slot address NO_SLOT (-1), which index_select and index_copy_ refuse: a bare IndexError on the CPU, and on CUDA a device-side assert that takes the worker down. Both now raise a ValueError naming the pool, as the bf16 GDN backend does for its kernel; the check is a bool per call, once at capture for a captured forward. * audio_features: use the engine's fp32_precision_scope for the mel matmul Drops mstar/utils/precision.py, a second copy of what mstar/engine/fp32_precision.py already does, which restored only the legacy matmul flag and not the per-backend ones. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: the engine's fp32_precision_scope in place of the removed helper Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * canary, canary_qwen: cap the decode loop at the request's max_output_tokens Without get_loop_caps the loop cap fell back to the graph's max_iters, so speculation ran past a max-token stop and a smaller kv_cache override refused every request. Same override as Whisper's. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * transcriptions: no compression-ratio retry for the untimed ASR adapters The 2.4 threshold is Whisper's sequential-long-form fallback. Canary and Canary-Qwen decode greedily with no conditioning text, so a forced sequential retry could only replace a correct transcript with a sampled one. Omnilingual already opts out the same way. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * parakeet: the decode tokens reach the client from the stop check's host copy The tokens join the active flag in the stop buffers. The stop check doesn't read them, but the worker then stages the batch's tokens to the host once, behind the step, and the shared-memory send takes each request's row from there instead of a device-to-host copy per request: register_outputs went from 0.57 to 0.25 ms per decode step at c=64 on 1x H100 (worker_phases). * audio_features: only a failed decode is the upload's fault load_speech_file wrapped the resample too, so a server-side failure there (e.g. a CPU allocation failure, a RuntimeError) came back as a 400 'could not decode the audio upload'. Split the decode from the resample and wrap only the decode. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * canary: the prompt's KV is a rows span From the second decoder layer on, the prompt's self-attention KV depends on the clip through cross-attention, so its token ids don't key it. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * canary: refuse a checkpoint whose mel front end NemoLogMel doesn't compute Canary-Qwen already checks its preprocessor settings at load; Canary assumed them. Read processor_config.json's feature extractor and refuse a mismatch at startup instead of decoding wrong features. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * token_step: say decode is uncapped with accelerator graphs off Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: load_speech_file's blame split goes through _decode, and a resample failure stays the server's Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * canary, canary_qwen: one audio file per request, refused before decoding Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * asr extra installs soundfile; Higgs segments and a stale test comment documented Without soundfile every WAV/FLAC upload went through torchcodec and needed FFmpeg, which the docs said only M4A/WebM do. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * parakeet: one walk, the encoder output decoded where it lands The transducer no longer projects a clip's frames in a step of its own and persists them for a second walk. One walk runs the encoder and then the decode loop, which re-injects the encoder output every iteration; each greedy step projects only the frames under the cursors (with Nemotron's prompt), in fp32. That drops a prefill step and a conductor round trip per request, about 0.5 ms of the worker's 3.5 ms a request at c=64 on 1x H100 (worker_phases: the step's GPU-thread and postprocess work, and the decode walk's message). The loop has no loop-back edge any more. The clock only made the node a same-node speculation target, and a speculated iteration past a clip's end is a whole chunk of steps for nothing; the worker still speculates a request's first iteration behind the encoder's step. Nemotron's prompt id travels in the step metadata, a host value, rather than as a tensor the loop would move to the device every iteration. * parakeet: 64 greedy steps a decode iteration, not 16 Every iteration costs each of its rows the engine's and the worker's per-step work: prepare_inputs, the store, route and send of its tokens, the stop check. At 16 steps a LibriSpeech test-clean clip ran 3.6 decode steps of the node at c=64 on 1x H100 (worker_phases; the two-walk loop, its speculated step past the end included), and that host work, not the GPU, bounded throughput. At 64 it runs 1.2, at the price of the idle steps of a row that ends early in an iteration. The loop cap and the stop check's guard count iterations of the node's chunk, which the loop test sets to 16 so that its short requests still span several iterations. * parakeet: a 2,048-frame decode bucket, 164 s of audio Past 41 s a clip decoded in the 8,192-frame bucket, whose captures hold at most 4 rows, so a batch of more than 4 minute-long clips ran its iterations eagerly. The middle bucket's captures hold 16. On the mixed workload (10% of the clips 60-120 s), 1x H100, 4 passes each: RTFx 1,862 -> 2,224 at c=32 and 3,090 -> 3,199 at c=64 (#796's transducer: 2,047 and 3,057). * TransducerAdapter: the compression-ratio opt-out now comes from _UntimedTranscriptionAdapter Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * parakeet, nemotron_asr: one audio file per request, refused before decoding Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * parakeet: one audio file per request, refused before decoding * RecurrentStatePool: gather / scatter_ read the step's plan, and gather takes a start state Every caller fetched the addressing and wrote its own torch.where over has_state for rows with no saved state (a request's first step, or a padding row on the sink, which held leftovers). gather now does that select itself, zeros by default or the given start, and scatter_ checks the row count with a message naming the pool. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * parakeet: read and write the transducer state through the pool's gather(start=) / scatter_ Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * parakeet: refuse a transducer state pool with fewer than 2 slots max_slots 1 leaves no usable slot (slot 0 is the sink), so every request hung waiting for one. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * parakeet: comments on the clock edge and slot lifetime match the code The clock carries each row's active flag, which check_stop reads, and a request's state slot frees at teardown, not at its last iteration. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * docs: the recurrent state pool's gather / scatter_ for models Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * RecurrentStatePool: say what a zero-span row addresses Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * parakeet: drop the stop check's copy of the loop cap get_loop_caps already has the runtime end each request's loop at that cap, so the stop check reads only the active flag. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * test: drop the frozen copy of the transducer step; the independent reference covers it Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> * RecurrentStatePool: cast the scatter index once per step gather passes the int32 slot indices to index_select as they are, and scatter_ casts them to the int64 index_copy_ needs once per step, cached on the step's addressing (keyed by whether a graph capture made it, so a capture records its own cast and each replay recomputes it from its own plan). Option (a) of the two: a staged int64 copy in plan would add an H2D for every pool user, GDN and Mamba-2 included. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> --------- Co-authored-by: Kirill Vasilev <290354918+vasilevklart@users.noreply.github.com> Co-authored-by: Joseph Tan <jdtan638@gmail.com> Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com> Co-authored-by: Naomi Sagan <nsagan@stanford.edu>

NSagan271 added this pull request to stack #355 October 7, 2026 02:26

NSagan271 marked this pull request as ready for review October 7, 2026 09:25

NSagan271 requested review from Gaurav-Shah05, garv901, merceod, stephen-dwq and vasilevklart October 7, 2026 09:26

NSagan271 force-pushed the cp/1-combined-walks branch from 0cb803d to 3926a8d Compare October 8, 2026 05:54

merceod mentioned this pull request Oct 8, 2026

Host-side overhead: per-row Python out of the decode step #368

Merged

NSagan271 force-pushed the cp/1-combined-walks branch from 99b37ef to fa43c91 Compare October 8, 2026 22:57

merceod mentioned this pull request Oct 11, 2026

[Chunked Prefill #1] Batch builder abstraction #354

Merged

2 tasks done

merceod requested changes Oct 11, 2026

View reviewed changes

merceod mentioned this pull request Oct 11, 2026

[Chunked Prefill #3] Declare input sequence length #364

Closed

2 tasks done

Gaurav-Shah05 reviewed Oct 12, 2026

View reviewed changes

Gaurav-Shah05 requested a review from merceod October 12, 2026 15:10

This was referenced Oct 13, 2026

[Chunked Prefill #4a] Chunked prefill core #365

Closed

[Chunked Prefill #6] Arrival-order batch builder #430

Closed

NSagan271 force-pushed the cp/1-combined-walks branch from 139d7e7 to d0a7eac Compare October 14, 2026 03:47

Base automatically changed from cp/3-batch-builder to main October 14, 2026 07:20

NSagan271 and others added 3 commits October 14, 2026 07:41

Orpheus: batch prefill and decode under one combined walk

034d631

LLM's prefill and decode rows share a "mixed" walk. Mixed steps replay the packed prefill CUDA graphs, where a decode row is a one-token segment. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

stephen-dwq approved these changes Oct 18, 2026

View reviewed changes

vasilevklart approved these changes Oct 18, 2026

View reviewed changes

This comment was marked as off-topic.

Sign in to view

merceod approved these changes Oct 19, 2026

View reviewed changes

NSagan271 and others added 4 commits October 19, 2026 03:17

Merge remote-tracking branch 'origin/main' into cp/1-combined-walks

e93b31a

# Conflicts: # mstar/worker/micro_scheduler.py # test/modular/test_worker_drain.py

Merge remote-tracking branch 'origin/main' into cp/1-combined-walks

cdfd56e

# Conflicts: # test/modular/test_worker_drain.py

Merge remote-tracking branch 'origin/main' into cp/1-combined-walks

8b48e88

# Conflicts: # mstar/worker/worker.py

Cosmos3 test: a real StepContext, which carries the per-row walks the…

cf01bda

… engine reads Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>

NSagan271 merged commit 73dc4b5 into main Oct 19, 2026
4 checks passed

NSagan271 deleted the cp/1-combined-walks branch October 19, 2026 05:14

arun-elr mentioned this pull request Oct 19, 2026

[Orpheus][Bug fix] Stream every SNAC frame, with left and right context per window #452

Merged

2 tasks done

vasilevklart mentioned this pull request Oct 20, 2026

Add Kimi Delta Attention on the recurrent pool #389

Open

Gaurav-Shah05 mentioned this pull request Oct 20, 2026

fix prefix caching under combined walks #768

Merged



		def _min_cap(*caps: int \| None) -> int \| None:
		return min((cap for cap in caps if cap is not None), default=None)

Conversation

NSagan271 commented Oct 7, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

What does this PR do?

How was it tested?

Checklist

Uh oh!

merceod left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Gaurav-Shah05 left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Gaurav-Shah05 commented Oct 12, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

stephen-dwq left a comment

Choose a reason for hiding this comment

Uh oh!

vasilevklart commented Oct 18, 2026

Uh oh!

This comment was marked as off-topic.

Uh oh!

NSagan271 commented Oct 19, 2026

Uh oh!

github-actions Bot commented Oct 19, 2026

Automated review

Questions

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

NSagan271 commented Oct 7, 2026 •

edited

Loading

Gaurav-Shah05 commented Oct 12, 2026 •

edited

Loading