Writing GPU kernels in Rust: how I made speakrs 60× faster than pyannote while dropping ONNX Runtime, cuDNN and cuBLAS

speakrs is my open-source speaker diarization library in Rust. It runs the pyannote community-1 pipeline: segmentation, speaker embeddings, then PLDA and VBx clustering. Three weeks ago its CUDA mode ran the models through ONNX Runtime. Today it runs on its own GPU kernels, written in Rust with cuda-oxide, and it can be built to need nothing but the NVIDIA driver. On the same RTX 4090 model, the full Earnings-21 set went from 2692 seconds to 75 seconds, a 35.8× speedup, and VoxConverse dev went from 1236 to 38 seconds, 32.4×. DER didn’t change. For comparison, pyannote in Python takes 2301 seconds, about 38 minutes, on the same VoxConverse dev set and GPU. That makes speakrs about 60× faster than pyannote, at 7.0% DER against pyannote’s 7.2%. On an RTX 5090, speakrs diarizes an hour of audio in about 1.2 seconds.

It didn’t start as a speed project. The first goal was to drop ONNX Runtime, a large native dependency, by running the same models on a small native CUDA backend that still used cuDNN and cuBLAS. I had already written GPU kernels in CUDA C++ with agents before this, and it went fine. But I love Rust, and I had just heard about cuda-oxide, a way to write GPU kernels in Rust. I wanted to try it, and moving off ONNX Runtime was a perfect real-world project to test it on. The first kernels beat cuDNN, so the experiment grew: could cuda-oxide replace cuDNN and cuBLAS completely, not just speed up a few layers?

I didn’t write those kernels by hand. AI coding agents wrote most of the code. That was far from all of it. I set the goals and decided what to speed up. I chose which GPUs to support, what “faster” had to mean and how to measure it, and what counted as done, with no accuracy loss allowed. I also made the calls on what to merge and what to drop.

Contents

• The result
• Where the speed came from
• Writing kernels in Rust with cuda-oxide
• The driver-only build
• Tuning per GPU
• How the agents did the work
• Per-GPU tuning in detail
• Lessons learned
• Try it
• What’s next

The result

Every run covers the full datasets: 216 VoxConverse dev files and 44 Earnings-21 files.

Dataset	Mode	pyannote	Before cuda-oxide	After cuda-oxide	RTFx, pyannote / before / after	vs before	vs pyannote	DER, speakrs / pyannote
VoxConverse dev	cuda	2301.3 s (38 min)	1236.3 s (21 min)	38.2 s	32× / 59× / 1912×	32.4×	60.2×	7.0% / 7.2%
VoxConverse dev	cuda-fast	2301.3 s (38 min)	604.4 s (10 min)	20.0 s	32× / 121× / 3646×	30.2×	115.1×	7.4% / 7.2%
Earnings-21	cuda	8036.8 s (2 h 14 min)	2692.1 s (45 min)	75.3 s	18× / 53× / 1876×	35.8×	106.7×	9.7% / 9.7%
Earnings-21	cuda-fast	8036.8 s (2 h 14 min)	1284.1 s (21 min)	39.1 s	18× / 110× / 3613×	32.8×	205.5×	9.2% / 9.7%

speakrs DER didn’t change from before cuda-oxide to after. RTFx is audio length divided by wall time. At 1912×, an hour of audio takes about 1.9 seconds with cuda and 1.0 second with cuda-fast, which uses a 2-second step instead of 1 second.

I also ran all eight benchmark datasets on that RTX 4090 and on an RTX 5090. DER is the same on both GPUs:

Dataset	Audio	cuda DER	cuda RTFx, 4090 / 5090	cuda-fast DER	cuda-fast RTFx, 4090 / 5090
VoxConverse dev	20.3 h	7.0%	1912× / 2922×	7.4%	3646× / 5633×
VoxConverse test	43.5 h	11.1%	1859× / 2888×	11.2%	3618× / 5678×
AMI IHM	18.7 h	17.0%	1909× / 2919×	17.4%	3563× / 5791×
AMI SDM	18.7 h	19.7%	1791× / 2927×	20.6%	3754× / 5838×
AISHELL-4	12.7 h	11.1%	1866× / 2844×	11.4%	3534× / 5621×
Earnings-21	39.3 h	9.7%	1876× / 2864×	9.2%	3613× / 5667×
ICSI	71.7 h	33.3%	1937× / 2888×	33.7%	3708× / 5725×
AVA-AVD	4.4 h	45.4%	2075× / 3095×	48.9%	3927× / 5971×

That’s about 229 hours of audio in under 8 minutes with cuda on the RTX 4090, or under 4 minutes with cuda-fast. The RTX 5090 takes under 5 minutes with cuda and under 2.5 minutes with cuda-fast. The old ONNX Runtime rows for VoxConverse test and AMI IHM ran on an L40S (50× and 49×), so I haven’t put them in the speedup table.

Where the speed came from

The work landed as a stack of pull requests over about three weeks. Each was measured with alternating A/B rounds on a fixed file set, with DER or byte-identical RTTMs as the accuracy check. The GPU and file set depended on what I could rent that day, so these factors don’t multiply out to the headline. The headline is the direct before/after above.

Step	What changed	Measured gain	Where
#36	Native CUDA backend instead of ONNX Runtime, FP32	1.81× (cuda), 1.62× (cuda-fast)	RTX 5070 Ti, 216 files
#36	TF32 embedding, persistent LSTM algorithm, CUDA graphs	1.33×, 1.37×	same
#36	Three cuda-oxide kernels: ResNet 3×3 conv, BiLSTM, SincNet	1.23×, 1.21×	same
#37	A speakrs kernel for every layer; driver-only build	1.54–1.55× over cuDNN/cuBLAS (driver-only)	RTX 5060 Ti, 4060 Ti, 30 files
#38	Winograd for wide layers, LSTM tiles, exact batch sizes	+15.5–16.6% (driver-only)	same
#39	Measured per-GPU recipes, speakrs cuda tune	+11% (cuda)	T4, 10 files
#40	FP16 tensor cores for the ResNet trunk	1.69× on T4, +7.9% on RTX 4060 Ti	10 and 30 files
#43	speakrs kernels by default on every GPU	+26.8% (cuda)	RTX 4090, 10 files
#44	Decode and clustering overlap GPU inference	+20.4% (cuda)	RTX 4090, 10 files
#51, #61	A100: FP16 tiles for the early trunk, a resident LSTM recurrence	3.5–4.3% and 2.5% less time (cuda)	A100 40 GB, 10 files
#54	Upload overlapping audio windows once per batch	7.9% less time (cuda), 9.2% (cuda-fast)	RTX 4090, 10 files
#56	Residual block hidden activations in FP16	5.9% less time on RTX 4090, 7.8% on T4	10 files
#57	Embedding batches overlap the next decode	2.5% less time (cuda)	RTX 4090, 10 files
#58	Residual stream between blocks in FP16	8.2% less time on RTX 4090, 10.0% on T4	10 files
#59, #60	FP16 through the 128- and 256-channel tiles	3.5% less time on RTX 4090, 4.9% on A100, 8.9% on T4	10 files
#62	Coalesced FP16 epilogue stores	5.3% less time on T4	10 files
#75	RTX 5090 recipe and sm_120 FP16 Winograd kernels	1.57× (cuda), 1.59× (cuda-fast)	RTX 5090, 216 files

The first step was about the dependency, not speed. #36 loads the same model weights, exported from the original ONNX models, as safetensors and runs them on cudarc with cuBLAS and cuDNN. That removed ONNX Runtime from the CUDA build, and it also turned out 1.81× faster than ONNX Runtime in FP32, more than any single kernel since. The total for #36 was about 2.96× for cuda and 2.69× for cuda-fast, the product of three measured steps.

With the native backend in place, I tried cuda-oxide. The three kernels in #36 were the test of how good and how ready it was, and they replaced cuDNN only where they were faster. The ResNet 3×3 convolutions were 1.19–3.1× faster per conv, the BiLSTM recurrence 1.10–2.95× per stage, and the fused SincNet conv, abs and max-pool 2.1× faster with bit-identical output. Everything else still went through the libraries.

cuda-oxide was clearly ready for real kernels, so I asked a bigger question: could it replace the libraries altogether? From #37 on, the goal was a speakrs kernel for every layer, and a build that loads neither cuDNN nor cuBLAS. That covers the full ResNet34 trunk, the segmentation convolutions and dense layers, a cuBLAS-free bidirectional LSTM, and one filterbank kernel that runs each 512-point real DFT as a packed 256-point complex FFT. On RTX cards the driver-only build ended up about 1.8× faster than the same backend on cuDNN and cuBLAS: 190× to 344× on an RTX 5060 Ti, and 208× to 369× on an RTX 4060 Ti, after #38.

Some gains weren’t kernels at all. #38 skips embedding windows with no active speaker, adds exact batch sizes so partial batches stop running one window at a time, and caches kernel hashes once per process, which cut T4 startup from 2.3 s to 0.6 s. #44 decodes upcoming files on CPU workers and clusters one file while the GPU works on the next. Once the GPU was this fast, the CPU stages were in the way, and overlapping them was worth another 20% on the RTX 4090. #54 stopped copying every 10-second window to the GPU on its own, about 20 MB per 32-window batch that covers only about 41 seconds of audio, and uploads the stretch of audio once instead. #57 prepares the next embedding batch while the GPU runs the current one.

After the FP16 convolutions, the activations between them moved to FP16 too. #56 keeps each residual block’s hidden activation in FP16, #58 keeps the residual stream between blocks in FP16, and #59 and #60 pass halves straight through the 128- and 256-channel tiles. Each step was small, but together they took roughly a sixth off the RTX 4090’s time and a quarter off the T4’s.

Writing kernels in Rust with cuda-oxide

cuda-oxide is the SIMT kernel model in NVIDIA’s cuda-rust project. You write #[kernel] functions in Rust and it compiles them to PTX. I used it on a pinned nightly.

What worked well:

• Kernels live in Rust next to the host code, with the same types and layouts. That made fused kernels easy to change and re-measure.
• #[unroll] and ptx_asm! covered nearly everything I needed.
• The generated PTX was readable, so reviews could reason about it line by line.
• Exact numerics were possible. The SincNet kernel matches cuDNN bit for bit in FP32.
• Shipping works without making users install anything. I build PTX and cubins offline, commit them, and embed them with include_str!. Users stay on stable Rust with no CUDA toolkit, and a CI check fails if the committed artifacts are stale.

The rough edges, most of which I filed upstream:

• Scalar intrinsics such as fma_rn_f32, mul_rn_f32 and ex2_approx_f32 are gated to sm_80, though the PTX instructions exist on sm_75. I used ptx_asm! instead (#1462). The FP16 mma.sync m16n8k8 intrinsics had the same problem on Turing; that issue is now closed (#1467).
• The docs said you need a barrier before reading another thread’s shared-memory writes, but not what the compiler may move across sync_threads() for raw pointers or ptx_asm!. I made every shared access a side-effecting ptx_asm! to be safe. That docs issue is now closed (#1463).
• There’s no documented way for a library crate to ship prebuilt kernels to users on stable Rust. I built my own artifact manifest and staleness check (#1464).
• Plain shared loads from a u32 address can only be written as ptx_asm!. LLVM then sees the address as an opaque integer, hoists part of it out of the loop, and ptxas can’t fold the offset into the load. In one kernel, removing my workaround cost 24 SASS instructions. That request was closed as not planned (#1465).
• ptxas fully unrolls small constant-trip loops even when the PTX keeps them rolled, then issues every iteration’s loads at once. I emit .pragma "nounroll" by hand and asked for a #[nounroll] attribute (#1468).
• #[unroll] handles counted while loops but not for loops or helper functions. One large warp tile exceeded its clone budget and spilled accumulators to local memory, so I shrank the tiles.
• Tensor-core code is still mostly raw PTX. Every mma.sync, ldmatrix and cp.async group is a ptx_asm! site. Where cuda-oxide has a native form, such as typed vector stores and warp::sync_mask, I’ve since switched to it (#65).

The driver-only build

From #37, a build that enables only a GPU target feature doesn’t compile cuDNN or cuBLAS support at all:

speakrs = { version = "0.7", features = ["cuda-rtx40"] }

That binary opens only libcuda and the driver’s PTX JIT, and needs driver 570 or newer. In every driver-only run, ldd showed no cuDNN or cuBLAS link and strace showed neither library opened. The NVIDIA container toolkit mounts the driver from the host, so a GPU image can use a slim base instead of a CUDA and cuDNN runtime image and leave out about 2 GB: cuDNN 9 is about 1.2 GB and cuBLAS 12 about 0.8 GB as NVIDIA ships them.

GPU	Feature
RTX 20 series, T4	cuda-rtx20
RTX 30 series, A10, other Ampere	cuda-rtx30
RTX 40 series, L4, L40S, other Ada	cuda-rtx40
A100	cuda-a100
H100	cuda-sm90
RTX 50 series	cuda-rtx50

I started by shipping only sm_75 PTX and letting the driver JIT it forward. That worked, but a cold JIT of all six kernel areas took 1.9 seconds, against 5 ms warm. So I now ship exact-architecture cubins for sm_75, 80, 86, 89, 90 and 120 next to the PTX.

With every kernel area and architecture compiled in, the packaged crate reached 13.9 MiB compressed, over crates.io’s 10 MiB limit. The kernels now ship in their own crates, speakrs-ptx-sm75, speakrs-ptx-sm80 and speakrs-ptx-sm120, which the CUDA features pull in (#66). The speakrs crate itself is 0.58 MB, and a CPU or CoreML build never downloads the kernel crates.

The plain cuda feature keeps cuDNN and cuBLAS as a fallback for a layer without a speakrs kernel. Since #43, every layer has a kernel on the GPUs I’ve tested, so the libraries don’t load there either.

Tuning per GPU

No kernel is fastest on every GPU. The A100 has fast TF32 tensor cores, the T4 has no TF32 and relies on FP16, and the RTX 5090 is fastest with nearly the whole model in FP16. So each GPU picks its own kernel for each layer, from the first of these that exists:

1. a tune file from speakrs cuda tune, written for this GPU, driver and build
2. a built-in recipe measured on this exact GPU model
3. a default for its GPU generation

speakrs cuda tune times each layer’s candidates on your GPU in under 30 seconds on the RTX cards I tried. It only picks kernels with end-to-end accuracy evidence.

RTFx on the first 10 VoxConverse test files, including start-up and model loading, median of 3 runs:

GPU	cuda	GPU-specific build
T4	294×	294×
A10	432×	436×
L4	310×	312×
A100 40 GB	1296×	1296×
RTX 4090	1453×	1438×
RTX 5090	2500×	2501×

The driver-only, GPU-specific builds are within about 1% of cuda. In an earlier build, speakrs’s kernels were 1.4× to 2.0× faster than cuDNN and cuBLAS on the RTX 4090, L4, A10 and H100. The H100 hasn’t been tuned yet. The details for each GPU are near the end.

How the agents did the work

AI coding agents wrote most of the kernel code, ran the benchmarks on rented GPUs and wrote up what they found. Other agents reviewed their work. I set the bar for accepting a change: byte-identical RTTMs or DER that didn’t get worse, plus a measured speedup.

The biggest lesson was where the time went. From October 2 to midday on October 7, about five days went into a proof harness. The bar was “provably no slower on any layer, batch size, math mode, PTX tier or GPU”. That needed a separate proof for every combination, plus noise rules, paired timing guards, f64 reference outputs and mutation tests to show each gate could catch a cheat. Every review found a gap, every gap became a new harness feature, and each feature needed its own proof. In those five days I merged one PR, #36.

Then I dropped the harness for a rough check: byte-identical RTTMs or unchanged DER, and a measured RTFx gain on a fixed file set. More shipped in the next eight hours than in the five days before, including the A100 fix, the driver-only proof on four GPUs, the per-GPU recipes, Winograd for the wide layers, and the T4 beating cuDNN. Most of the kernels had existed since the first days. The harness did catch real bugs, such as an unmeasured cubin and a truncated shared-memory pointer, but a cheap end-to-end check would have caught most of them in hours.

What I’d tell anyone doing this with agents:

• Measure the payoff first. Profile end to end and work out how much of the runtime a stage can move before anyone builds proof machinery for it.
• Keep checks rough until something ships. A CPU or f64 reference for accuracy, a timed A/B against the library, and an end-to-end DER and speed check on fixed files. If the harness takes longer than the kernels, stop and reassess.
• Profile on every new GPU. Compare which execution unit, layout and algorithm each hot kernel uses with what the vendor library chose. A win on a consumer card says nothing about a datacenter card.
• Let workers fix their own scratch bugs. Early on, agents stopped and waited on every failure, even typos in their own scripts. GPU time was cheap next to that: one round of three rentals cost $0.21.

Per-GPU tuning in detail

The same kernel can win on one GPU and lose on another. On the A100 the driver-only build was 1.4× slower than cuDNN, even after several rounds of tuning the wide C128 and C256 layers. One profile of the whole pipeline showed the cause: the 14 early C32 and C64 convolutions ran scalar FP32 FMA, while cuDNN ran them on TF32 tensor cores, about 8× faster than FP32 on an A100. Consumer cards have roughly equal dense TF32 and FP32 throughput, so the same kernels beat cuDNN there and the problem never showed. With TF32 tensor-core kernels for those layers, the A100 build went from 1.4× slower than cuDNN to 1.17× faster. Those layers later moved to FP16 tiles on the A100 too (#51).

The T4 needed the opposite. Turing has no TF32, so #38 added a Winograd convolution for the 64-channel layers in plain FP32 FMA, which put the T4 ahead of cuDNN. #40 then moved the trunk’s 3×3 convolutions to FP16 tensor cores. The kernels round each operand to FP16 with a fixed ×1024 scale and accumulate in FP32. On the T4 the trunk runs in 136 ms per 32-window batch against cuDNN’s 272 ms, and end to end the T4 got 1.69× faster. FP16 overflows above ±65504, so a layer whose weights would overflow skips FP16 at model load, and the kernel sets a flag when a scaled activation overflows or isn’t finite. The host then recomputes that batch with the non-FP16 kernels. On the full VoxConverse dev set, the T4’s DER went from 7.0125% to 7.0118%.

FP16 didn’t help everywhere. On the RTX 5060 Ti my FP16 kernels were slower than my TF32 Winograd on the C128 and C256 layers, 0.64–0.71× the speed, so that card stays on TF32.

The RTX 5090 went the other way. Its recipe, from #75, runs the whole trunk in FP16 at large batches. It pairs FP16 taps into m16n8k16 products, which issue at the same rate as m16n8k8 on that card, and adds fused FP16 Winograd kernels for the 64-, 128- and 256-channel layers. Weights whose Winograd transform could leave the FP16 range stay on the direct FP16 tiles. Each step was measured over all 216 VoxConverse dev files:

Step	cuda	cuda-fast
FP16 early trunk recipe	39.0 → 35.0 s	20.4 → 18.3 s
FP16 taps paired into m16n8k16	34.0 → 30.9 s	17.6 → 16.1 s
Whole trunk in FP16 at large batches	31.0 → 27.3 s	16.0 → 14.2 s
FP16 Winograd trunk kernels	27.0 → 25.7 s	14.0 → 13.2 s
64-channel layers on FP16 Winograd	25.3 → 24.9 s	13.1 → 12.8 s

Two host-side changes in between, which apply to every GPU, account for the gaps between rows. Over the 216 files, cuda DER went from 7.0123% to 7.0119%, with 2 RTTMs changed. The recipe matches only an RTX 5090 with 170 SMs, so other sm_120 cards keep the earlier kernels until someone measures them.

When I tuned these GPUs, the libraries still won a few individual layers, each by 2–8 µs, which doesn’t move end-to-end speed:

GPU	Layer	Batch	speakrs	cuDNN/cuBLAS
T4	layer4.0.shortcut	4	149 µs	141 µs
T4	linear0 / linear1 (FP32)	1	18 / 13 µs	15 / 11 µs
A100	layer4.0.conv1	4	59 µs	57 µs
L4, A10, RTX 4090	seg_1	32	31 / 42 / 15 µs	28 / 37 / 13 µs

Lessons learned

cuda-oxide is ready. Every speakrs kernel is now Rust, from the FFT filterbank to the FP16 tensor-core Winograd tiles. The rough edges above cost me workarounds, not rewrites, and none of them stopped a kernel from shipping. If you already write Rust, you don’t need a separate CUDA C++ codebase for your kernels.

Replacing ONNX Runtime with agents was easy, and worth it. It was the first step and the simplest one: export the weights, write a small native backend for a fixed model, and check the output against the old path. Agents are good at exactly that kind of work. It removed a large native dependency, and the native backend alone was 1.81× faster than ONNX Runtime. If you run a fixed model through ONNX Runtime on CUDA, I’d do this first.

Replacing cuDNN and cuBLAS mostly wasn’t worth it. It worked: the driver-only build is faster than the libraries on every GPU I tuned, and a GPU image is about 2 GB smaller. But a vendor library is tuned for every GPU NVIDIA ships, and my kernels are tuned only for the GPUs I’ve rented and measured. The A100 was 1.4× slower than cuDNN until I profiled it on an A100. The T4 needed FP16 tensor cores, and the RTX 5090 needed its own recipe. Each new GPU means renting it, profiling it and tuning it again. GPU rentals for this project, on Vast.ai, Modal and Beam, came to about $150. The money was small, but it only bought tuning for the GPUs I tried, and every GPU I add costs more rentals and more of my time. The H100 still hasn’t been tuned. Unless your users run a small, known set of GPUs, or the 2 GB really matters, keep cuDNN and cuBLAS and write kernels only for the layers they’re slow at.

Try it

If you know which GPU you’ll run on, use its feature. You only need the NVIDIA driver:

speakrs = { version = "0.7", features = ["cuda-rtx40"] }

If you don’t, use cuda, which includes kernels for every supported generation and can fall back to cuDNN 9 and cuBLAS 12:

speakrs = { version = "0.7", features = ["cuda"] }

On a GPU without a built-in recipe, run the tuner once with the same feature as your application, since a tune file only applies to the build that wrote it:

cargo build --release --no-default-features --features cuda-rtx40 --bin speakrs
./target/release/speakrs cuda tune --models-dir /path/to/models

speakrs 0.7 also ships Python bindings. The Linux x86_64 wheel includes CPU and CUDA, and the macOS arm64 wheel includes CPU and CoreML:

pip install speakrs

import speakrs

pipeline = speakrs.Pipeline.from_pretrained("cuda")
print(pipeline.diarize_file("meeting.wav").rttm())

The CUDA wheel runs the same kernels, and its RTTMs match the Rust pipeline exactly.

Every GPU from the RTX 20 series and T4 onward is supported. The benchmarks have every table, the hardware and the timing method, and the PRs linked above have the per-step numbers. If you run speakrs on a GPU I haven’t tested, I’d like to see your RTFx and tune output in an issue.

What’s next

Next I’m doing the same for speech-to-text. scriptrs is my Rust transcription library, and its new CUDA backend runs NVIDIA’s Parakeet TDT 0.6B v2 and v3 without ONNX Runtime.

	RTFx	vs ONNX Runtime	WER	Wall time
scriptrs	948×	3.7×	1.93%	20.5 s
NVIDIA NeMo (PyTorch)	477×	1.9×	1.93%	40.8 s
parakeet-rs (ONNX Runtime CUDA)	253×	1×	1.92%	76.9 s

scriptrs is already about twice as fast as NeMo and 3.7 times as fast as ONNX Runtime, with the same WER. It holds up on long recordings too: across the 39 hours of Earnings-21 earnings calls, it transcribes at 2500× real time with 8.24% WER, where NeMo gets 9.58%.

After that comes dialogrs, which combines scriptrs and speakrs into a faster alternative to WhisperX: a transcript with speaker labels. On the same RTX 5090, over all 44 Earnings-21 calls:

	WER	cpWER	DER	RTFx	vs WhisperX	Wall time
dialogrs, Parakeet v3 + speakrs cuda	8.24%	16.66%	9.69%	1401×	17.7×	1 min 41 s
WhisperX, large-v3 + pyannote community-1	9.74%	17.82%	9.70%	79×	1×	29 min 44 s

dialogrs is 17.7× faster than WhisperX, with lower WER and cpWER and the same DER.

If you want the full write-ups of these and to know when they’re ready, subscribe to the email list below. I’m also working on an API that serves all of this at a very low cost, which opens soon. Here’s how it compares with pyannoteAI.