Part of qwen3-tts — a pure C inference engine for Qwen3-TTS.
TL;DR
The engine is pure C and CPU by default. We added two opt-in GPU backends that leave the CPU path untouched:
-
Apple Metal (
make metal) — 0.6B model RTF 1.5 → 0.60 on an M1, 0.36 on an M2 Pro; streaming first-audio in 314 ms. -
NVIDIA CUDA (
make cuda) — 1.7B model RTF 0.44 on a ~270 GB/s GPU (RTX 4060-class), scaling to ~0.12 on a 4090. - Server request-batching on both (plus CPU): ~2.8× throughput on Metal, ~3.35× on CUDA — serving N concurrent users in roughly the time of one.
And two optimizations that looked obvious and that we measured and threw away (ICB and MMA). More on that at the end — it's the most useful part.
RTF = processing time ÷ audio duration. < 1.0 = faster than real time.
The one architectural idea: keep everything resident
The naive way to "use the GPU" is to offload one operation at a time — send the activation over, run a matmul, copy the result back. For a model that runs 16 sequential Code-Predictor passes per audio frame plus a 28-layer Talker step per token, that's death by a thousand PCIe/round-trip cuts. We measured it: per-op offload was slower than the CPU.
The fix is the same on both backends: make the whole step resident. Weights, KV cache, and activations live on the device. A full decode step is encoded into one command buffer / one kernel graph, committed once, waited on once. The CPU orchestrates; it never babysits individual matmuls.
- On CUDA that's resident weights + CUDA Graphs + cuBLAS for the pointwise convs.
- On Metal that's MTLBuffers cached by pointer + one
MTLCommandBufferper step + simdgroup matvec kernels.
The biggest single win on Metal came from the Code Predictor. It was sync-round-trip-bound — 16 waits per frame. Moving the whole 16-pass RVQ loop (embed → 5 transformer layers → argmax, ×16) onto the device as one command buffer with one sync/frame took 0.6B from RTF 1.36 → 0.89. Direct kernel fusions (per-head RMSNorm+RoPE, residual-add+norm) took it to 0.72, and int4 weights to 0.60 — the practical floor on an M1, where the step is dispatch-bound.
The rented-Mac trick (and why it works)
We develop on an M1. To see what an M2 does, we rented a Scaleway Mac mini M2 Pro by the hour. Two things made this painless:
1. Metal shaders compile at runtime. The MSL kernels are embedded as a string and compiled by the local Metal driver at startup (newLibraryWithSource). Consequence: a binary built on an M1 drives an M2's GPU at full speed — no rebuild needed, because the shaders are recompiled for whatever GPU is present. Only the CPU SIMD paths are baked in at build time.
2. One curl to bootstrap a bare box. A fresh macOS box has curl but not much else. A single script installs the Command Line Tools headlessly, clones the repo, pulls the models from HuggingFace, builds natively, and runs the benchmark:
curl -fsSL https://raw.githubusercontent.com/.../bootstrap_m2.sh | bash
That gave us the first cross-backend RTF matrix on real silicon — CPU vs Metal (M1 & M2) vs CUDA (a real NVIDIA GPU, not a datacenter part). A gotcha worth writing down: macOS has no timeout and no setsid (both GNU-only). If your bench scripts wrap curl in timeout, they silently no-op on a Mac.
Measured on the M2 Pro (0.6B, Metal, int8):
| Mode | RTF | First audio (TTFA) |
|---|---|---|
| CLI / warm server (single) | 0.36–0.39 | — |
| Streaming (single client) | 0.36 | 314 ms |
int8 is the sweet spot on Apple Silicon: it's bandwidth-rich, so int4's nibble-unpacking doesn't pay off (that's the x86 lever).
Batching is a throughput lever, not a speed-up
This one trips people up, so it's worth being blunt: batching does not make a single request faster. In batch mode each step does the work of B slots, so per-request RTF rises. What you buy is throughput — you serve B concurrent users in roughly the wall-clock of one, because each weight row is read once (from DRAM) and reused for all B (matvec → matmat).
We wired continuous request-batching (--serve --batch-size N) into all three backends. On the M2 Pro, 0.6B:
| Concurrent requests | Wall time | Per-request RTF | Throughput speed-up |
|---|---|---|---|
| 1 | 19.5 s | 1.01 | — |
| 2 | 21.9 s | 1.13 | 1.78× |
| 4 | 27.7 s | 1.44 | 2.81× |
So the honest per-request cost: RTF climbs from 1.01 (one request) to 1.44 (four) as the batch fills — the price of the B-wide work per step. What stays cheap is first
audio: streaming TTFA is 314 ms on the 0.6B and 517 ms on the 1.7B even under a full batch, so it still feels responsive. You trade a lone request's 0.36
single-stream RTF for ~2.8× aggregate throughput.
4 requests served in 27.7 s instead of ~78 s serial. The 1.7B model scales identically (2.82× at B=4). CUDA does 3.35× at B=8. And crucially, the batch output is bit-identical to single-stream — batching never changes what a user hears.
The bit-identity bug that taught us something
Our first batched Metal matvec accumulated scalar (element by element). The single-stream kernel accumulated with float4 dot products. Same math, different floating-point order → a ~1e-2 difference per step. In isolation this is benign (the argmax still matches). But TTS has a feedback loop: this step's hidden state → sampled token → next step's embedding. That tiny difference compounded, flipped a few token choices, and the batched audio diverged completely from single-stream (mel-corr 0.44).
The fix was three lines: vectorize the batched matvec to float4 so it matches the reference kernel's FP order exactly. Result: bit-identical (RMS-rel 0.000), and the batch server produced mel-corr 1.00000 vs single-stream.
Lesson: a numerically "close enough" kernel is not close enough inside a feedback loop. Match the reference's accumulation order.
The two optimizations we killed with data
The most valuable engineering this session was not shipping two things that looked obviously good.
ICB (Indirect Command Buffers) — the "CUDA-graphs for Metal" idea. Pre-encode the step once, replay it. Before writing it, we profiled: the encode was only 12% of the Talker step; the other 88% was GPU execution. ICB removes the 12% at best (~6–8% overall) for a large, invasive rewrite. And the Code Predictor — not the Talker — dominates the frame anyway. Not built.
MMA matmat (simdgroup_float8x8) for batched matvecs. The matrix units are ~4.6× faster than a simdgroup matvec on a big B=32 tile. But at our batch sizes (B≤8) the 8×8 tile is half-empty and the kernel underutilizes the GPU. Measured on the M2: better scaling ratio (3.27× vs 2.81×) but a 2× worse baseline → net loss. It only wins at large B. Kept opt-in, off by default.
Both were killed by a five-minute measurement instead of a two-day build. That's the whole point.
Takeaways
- Resident-everything beats per-op offload for autoregressive TTS — the round-trips, not the matmuls, are the enemy.
- Runtime-compiled Metal shaders make cross-device testing trivial: build once, run on any Apple GPU.
- Batching = throughput, not latency. Say so in your docs before someone benchmarks a single request and files a bug.
- Inside a feedback loop, match the reference kernel's FP order or watch a "correct" kernel diverge.
- Measure before you build. ICB and MMA both looked like free wins; the profiler said otherwise.
Full numbers, per-platform: docs/hardware-testing.md (Metal) and docs/cuda-performance.md (CUDA). It's all opt-in — the CPU path still ships as the default, pure C, zero heavy deps.
Top comments (4)
The residency insight generalizes way past TTS — "per-op offload was slower than CPU" is the exact wall you hit doing small-batch LLM decode too. When compute per kernel is tiny, launch/round-trip overhead is the whole budget, and the fix is always the same shape: stop babysitting individual matmuls, fold the whole step into one graph, commit once, sync once. CUDA Graphs and one command buffer per frame are doing the same job your Code-Predictor fusion did — collapsing 16 syncs/frame into 1. Two things I'd love more on. The RTF-vs-throughput story: single-stream RTF 0.60 and ~2.8× batched throughput optimize different things, and for a server the batched number is the real one — how do you handle the latency tax on the first request that opens a batch window? And huge respect for the "measured and threw away ICB and MMA" section — negative results feel like failures, so most posts cut them, but that's the most useful part here. The runtime-compiled MSL trick (build on M1, drive an M2 at full speed) is a genuinely clever portability property. What made ICB lose — dispatch overhead not amortizing at your batch sizes?
Really fair point — and the one I'd have asked too. The post has the throughput number (~2.8×) but not the per-request RTF / TTFA in batch mode, which is exactly what matters if you're running the server. I did record it, just didn't put the table in: per-request RTF sits around ~1.0 on the 0.6B (~1.27 on the 1.7B), with TTFA ~314 ms / ~517 ms. So that's the real trade — you give up the 0.39 single-stream RTF to buy ~2.8× aggregate across users. You're right that the two numbers optimize completely different things, and the batched one is the honest figure for a server.
On the first-request latency: we don't use a fixed batch window at all. It's continuous batching — a new request gets prefilled on a side thread while the active slots keep
generating, then joins in flight, so nobody waits for a window to fill. The only tax is that the batched kernels always do B-width work, so a lone request in batch mode runs at ~1.0 instead of the 0.39 a dedicated single-stream path gives. The clean fix (haven't built it yet) is adaptive: single-stream kernel when one slot is hot, flip to batched the moment a second request lands. And thanks — glad the "measured it and threw it away" bits landed, that's usually the part I most want to read in someone else's post too.
I will update this post adding TTFA and RTF for the server batching section, it's what is missing and really the most interesting part! Thanks a lot for your comment :)
The bit-identity bug is the real lesson here - a numerically "close enough" kernel is a ticking bomb the moment there's a feedback loop, and that's exactly the kind of bug that passes every unit test and only shows up as "the audio sounds a bit off" three requests deep. Do you have a standing regression test for FP accumulation order now, or is it still a manual mel-corr check per kernel change?
Interesting take on adding GPU backends to a C-based TTS engine—good to see Metal and CUDA being compared in a real-world deployment scenario. I’ve seen similar latency trade-offs when using VoltageGPU for containerized inference workloads; it’s a good reminder that not all optimizations pan out as expected without profiling. The rented-Mac trick is clever for quick benchmarking.