Squeezing 4.2× throughput out of Yoruba TTS with a hand-tuned Zig kernel.

The problem

Yoruba is a tonal language. Three pitches that carry semantic meaning, plus contour variation that no English-pretrained vocoder gets right. The model we host — afrispeech-tts, an open-weights checkpoint we like a lot — produces 80-band mel-spectrograms at 22kHz that then need to become audio.

The reference inference path is mel → HiFi-GAN-style vocoder → 16-bit PCM. Two stages, two device-to-device copies, a healthy amount of layer-norm in the middle. On a single L40S, the reference path got us about 12 real-time-factor (RTF) — meaning we could synthesize 12 seconds of audio per second of GPU time.

That is fine. But our launch customer wants live narration on long-form radio archives, and 12× RTF was eating GPUs we did not want to eat. So: where is the time going?

The baseline

Profiling the reference graph showed something boring: 41% of wall-clock was spent inside two ops that ran back-to-back on identical tensors — the last upsampling block of the vocoder and the prebatch normalization before the de-emphasis filter. Both are memory-bound. Both run in a kernel each. The intermediate tensor lives just long enough to be flushed out of L2 before the next kernel reads it back in.

This is the classic "fuse it" situation, except cuBLAS won't fuse it for you and TensorRT's plugin path costs more in operational complexity than the savings buy us. So we wrote a kernel.

The fused kernel

Tomoul's engine is Zig. The op layer compiles to PTX via a small intrinsics shim. For this kernel we wanted three things:

• Branchless decode. The mel-bin indexing is data-dependent but the upper bound is fixed; predicate the load, don't branch.
• One read, one write. Load the intermediate tile once into shared memory, do both ops in-register, write the final audio tile out.
• No allocator traffic. The activation buffer comes from the per-stream arena; the kernel never touches the device allocator.

The Zig glue is small. The interesting half:

1// fuse upsample + denorm + de-emphasis, one pass over the tile
2pub fn fusedTailKernel(
3  mel: [*]const f16,
4  gain: [*]const f16,
5  out: [*]i16,
6  n_frames: u32,
7) callconv(.PtxKernel) void {
8  const tid = ptx.threadIdx.x;
9  const tile = sharedTile(256);   // shared mem, one-shot
10  tile[tid] = upsample(mel, tid);
11  ptx.barrier();
12  const y = deEmphasize(denorm(tile[tid], gain[tid]));
13  out[tid] = saturatingI16(y);    // no allocator, no copy
14}

The full kernel is 84 lines of Zig including the boilerplate to register it. The autotuner picks a tile size of 192 or 256 depending on n_frames; everything else is fixed.

Numbers

Same model, same prompts (200 hours of afrispeech test split), same hardware:

Quality: PESQ and STOI within noise on the eval set. We expected this — the fused kernel is bitwise-identical to the reference path modulo accumulation order in fp16.

Pricing: afrispeech-tts went from $0.018 / 1k char to $0.0042 / 1k char on the menu, effective last week.

Takeaways

1. Profile first, fuse second. The 41% was hiding in plain sight. We caught it in two hours of NSight.
2. Boring code is fine. 84 lines of Zig. No fancy IR, no autotuning DSL. The whole point of running our own engine is that we can drop down a layer when it pays.
3. Open weights deserve hand-tuning too. Hyperscalers won't write you a custom kernel for a 2B-parameter African-language TTS model. We will. That is the wedge.

Code is open at github.com/tomoul/tomoul — look in src/ops/tts_tail.zig. Reach out if you have a kernel you'd like us to look at.