feat(qwen35): overlap Unified prefill with active decode

[CAICAIIs]

Summary

Closes #715.

Qwen3.5 previously executed Unified prefill chunks serially before active decode. This PR adds an opt-in Shared-SM overlap path so one prefill chunk can run on a dedicated CUDA stream while active decode continues on the model stream.

The default remains unchanged: --decode-overlap off.

Changes

• Add Qwen35DecodeOverlap::{Off, SharedSm}.
• Add --decode-overlap off|stream to the server and serving benchmark CLI.
• Run at most one in-flight prefill on an owned CUDA stream.
• Poll completion events while decode is active and wait explicitly when the last decoder retires.
• Drain the prefill stream on error, cancellation, drop, and shutdown.
• Route prefill and decode through separate cuBLAS handles.
• Keep stream-specific routing limited to Qwen3.5 async prefill.
• Add mode-aware ITL_STEP actions:
  overlap_launch, overlap_decode, overlap_complete, and overlap_wait.
• Reject unsupported combinations before model loading:
  • tensor parallelism with overlap;
  • scheduler-policy=auto with decode-overlap=stream;
  • Qwen3.5 Green Context overlap.

Results

Verified on 1x NVIDIA RTX 5090 32 GB, CUDA 12.8, sm_120, with Qwen3.5-4B.

Workload	Off	Stream	Change
1024/256, c1 TPOT	6.99 ms	7.07 ms	+1.20%
1024/256, c16 TPOT	11.48 ms	11.13 ms	-3.04%
1024/128, QPS 8 TPOT	12.12 ms	11.31 ms	-6.68%
1024/128, QPS 12 TPOT	14.85 ms	13.70 ms	-7.74%
1024/128, QPS 16 TPOT	16.88 ms	15.52 ms	-8.05%
Mixed-load ITL p99	58.23 ms	26.15 ms	-55.1%

All HTTP runs completed with zero failures, timeouts, error strings, or zero-token outputs. The c1 tradeoff and higher loaded TTFT keep overlap opt-in.

Nsight Systems confirmed concurrent prefill/decode kernel execution across the two CUDA streams. The trace contained 28,082 cross-stream kernel-overlap pairs.

Verification

• Release workspace library tests: 73 passed, 6 ignored.
• Qwen3.5 default-off scheduler e2e: passed.
• Shared-SM lifecycle e2e: passed; observed two overlap_wait actions.
• Qwen3.5 HF golden gate: 2 passed, 2 ignored.
• Qwen3 and Qwen3.5 release checks: passed.
• Stream-routing regression test: passed.
• cargo fmt --all -- --check: passed.
• git diff --check: passed.

The evidence covers single-GPU Qwen3.5-4B on RTX 5090. It does not claim tensor-parallel overlap, Green Context support, prefix-cache coverage, broad hardware coverage, vLLM parity, SOTA, or production readiness.

[FeathBow]

Ran this on sm_89, CUDA 12.9, Qwen3.5-4B.

#	Check	Result
1	qwen35 lib tests / kernels lib tests	74 + 15 passed
2	Qwen3.5 hf_golden_gate	2 passed, deltas match your table
3	Qwen3 hf_golden_gate (cross-model probe of the shared kernels change)	1 passed
4	Shared-SM lifecycle e2e	pass
5	Boundary parity: bg = max_batch - 1 decoders so bucket 4/32 actually replays during an 8192-token prefill, overlap witnessed structurally, all 16 injection steps asserted bitwise (token id, selected logprob bits, top-5 ids and value bits)	pass at both boundaries
6	Gate-lifted probe: max_batch=64, 40 decoders overlapping successive prefills, oracles on both sides, off-vs-off / stream-vs-stream controls	no corruption; all divergence reproduces in off-vs-off

Not covered: TP2, sm_120, HTTP layer, performance (cross-arch numbers aren't comparable).

[FeathBow]

should be pegainfer. As-is, cargo build --features qwen35 fails with E0433.

btw consider unifying this enum with the same-named one in pegainfer-server/src/config.rs (different variants and fallibility).

[FeathBow]

Could you add the exact server command (--max-batch, --max-prefill-tokens, CUDA Graph, features)? These binaries predate the --max-batch <= 32 gate and the server default (64) no longer starts — if the runs used the default, the table needs a re-run at <= 32.

[FeathBow]

Two suggestions here

1. Add const _: () = assert!(MAX_SHARED_SM_DECODE_BATCH <= ops::GEMM_LT_MAX_N); and a back-reference at the tune filter (weights.rs:691) — the equality is load-bearing. (Non-bucket max_batch is fine: the loader rounds at weights.rs:203; verified at mb=5.)

2. The "workspace-backed" rationale in this comment needs an update: cuBLAS 12.9 says cublasSetStream() "unconditionally resets the cuBLAS library workspace back to the default workspace pool", and gemm_cuda calls it every invocation — so the 32 MiB bound in cublas_init (linear.cu:234) detaches after the first GEMM and is a dead allocation today (remove it, or re-bind per stream). The accurate rationale for the gate: both paths share one handle's default pool, replay-vs-concurrent-call isolation is undocumented, and the docs flag multi-stream single-handle as nondeterministic — per-stream workspaces/handles are the documented fix, and the clean path to bucket-64 later. The gate itself is right; my sm_89 probes found no corruption, which bounds nothing on other archs.

[CAICAIIs]

Ran this on sm_89, CUDA 12.9, Qwen3.5-4B.

Check Result

1 qwen35 lib tests / kernels lib tests 74 + 15 passed
2 Qwen3.5 hf_golden_gate 2 passed, deltas match your table
3 Qwen3 hf_golden_gate (cross-model probe of the shared kernels change) 1 passed
4 Shared-SM lifecycle e2e pass
5 Boundary parity: bg = max_batch - 1 decoders so bucket 4/32 actually replays during an 8192-token prefill, overlap witnessed structurally, all 16 injection steps asserted bitwise (token id, selected logprob bits, top-5 ids and value bits) pass at both boundaries
6 Gate-lifted probe: max_batch=64, 40 decoders overlapping successive prefills, oracles on both sides, off-vs-off / stream-vs-stream controls no corruption; all divergence reproduces in off-vs-off
Not covered: TP2, sm_120, HTTP layer, performance (cross-arch numbers aren't comparable).

Thanks for the validation! I pushed a follow-up commit (7b32330d) addressing the review notes: the bench_serving Qwen3.5 crate rename, the Shared-SM max-batch boundary, the stale cuBLAS workspace rationale, and the pre-cap HTTP evidence wording.

I also validated the follow-up on RTX 5090 / sm_120 / CUDA 12.8:

• bench_serving qwen35 compile: pass
• Qwen3.5 e2e_scheduler: pass (TP2 ignored on single GPU)
• Qwen3.5 hf_golden_gate: pass (TP2 ignored on single GPU)

[FeathBow]

LGTM on the impl :)

Two non-blocking suggestions: retain decoder-side parity in the Shared-SM E2E, and rerun the HTTP A/B with explicit --max-batch 32.