Part of qwen3-tts — a pure C inference engine for Qwen3-TTS.
The engine is pure C and CPU by default. We added two opt-in GPU backends that leave the CPU path untouched:
make metal
) — 0.6B model make cuda
) — 1.7B model 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 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.
MTLCommandBuffer
per 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.
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).
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 | Batch speed-up |
|---|---|---|
| 1 | 19.5 s | — |
| 2 | 21.9 s | 1.78× |
| 4 | 27.7 s | 2.81× |
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.
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
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 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.
Full numbers, per-platform: docs/hardware-testing.md (Metal) and
docs/cuda-performance.md