Using agents to replace ONNX Runtime, cuDNN and cuBLAS with cuda-oxide kernels for a 32× speedup

speakrs is our 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. On VoxConverse dev that’s about 60× faster than pyannote on the same GPU, at 7.0% DER against pyannote’s 7.2%. On an RTX 5090, speakrs diarizes an hour of audio in about 1.2 seconds.

I didn’t write those kernels by hand. AI coding agents wrote most of the code. I set the goals, decided what counted as done, and made the calls on what to keep. That split matters here, because the most expensive mistake in the project was in the acceptance checks, not in the kernels.

The result

Both columns use the full datasets and the same timer: WAV loading, inference, clustering and RTTM output count; model loading and scoring don’t. The before rows are speakrs before cuda-oxide, running ONNX Runtime’s CUDA execution provider on an RTX 4090 with an AMD EPYC 7B13. The new rows ran on an RTX 4090 on Beam with 16 CPU cores and 8 GiB of RAM. It’s the same GPU model on a different host. The VoxConverse dev rows are the median of 3 runs; the Earnings-21 rows are one run.

Dataset	Mode	Before cuda-oxide	After cuda-oxide	Speedup	DER (unchanged)
VoxConverse dev, 216 files	`cuda`	1236.3 s (59×)	38.2 s (1912×)	32.4×	7.0%
VoxConverse dev, 216 files	`cuda-fast`	604.4 s (121×)	20.0 s (3646×)	30.2×	7.4%
Earnings-21, 44 files	`cuda`	2692.1 s (53×)	75.3 s (1876×)	35.8×	9.7%
Earnings-21, 44 files	`cuda-fast`	1284.1 s (110×)	39.1 s (3613×)	32.8×	9.2%

The numbers in brackets are RTFx: 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.

We 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.

60× faster than pyannote

We benchmarked pyannote community-1 on CUDA on the full VoxConverse dev set, in a container with the same GPU, CPU count, RAM and driver as the speakrs rows: one RTX 4090, 16 CPU cores and 8 GiB of RAM.

Implementation	DER	Time	RTFx	vs pyannote
speakrs `cuda`	7.0%	38.2 s	1912×	60× faster
speakrs `cuda-fast`	7.4%	20.0 s	3646×	115× faster
pyannote `community-1` (CUDA)	7.2%	2301.3 s	32×

speakrs cuda is 60× faster than pyannote, with 0.2 points better DER. The pyannote row is from an earlier run, so this compares separate runs. When we first ran both in one container, speakrs cuda was 36× faster. speakrs has gotten faster since, and pyannote can’t be rerun on that host.

On VoxConverse test, speakrs cuda is about 75× faster than pyannote at the same 11.1% DER. On Earnings-21, against an earlier pyannote run on the same GPU model, it’s about 107× faster at the same 9.7% DER. On AMI IHM, where the pyannote run used an L40S, speakrs matches its 17.0% DER.

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 we 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 mattered more than any single kernel. PR #36 loads the same model weights, exported from the original ONNX models, as safetensors and runs them on cudarc with cuBLAS and cuDNN. That alone was 1.81× faster than ONNX Runtime in FP32, and it removed a large native dependency. The total for #36 was about 2.96× for cuda and 2.69× for cuda-fast, the product of three measured steps. We couldn’t run one direct A/B because the old ONNX Runtime build ran out of memory on that box in a single-process full-set run.

The three kernels in #36 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.

From #37 on, the goal changed: 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. We 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 we 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. We 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 we 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. We 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!. We 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. We built our 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 our 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. We 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 we 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, we’ve since switched to it (#65).

Two problems in our first report to the cuda-oxide team didn’t hold up. We retested them on current cuda-oxide main and couldn’t reproduce either one: an add rewritten as an or in shared addresses, and a ptxas address-lowering bug. We withdrew both.

Some of the bugs were ours. An early kernel cast a generic shared pointer to a u32. It passed Compute Sanitizer 21 times on an sm_120 card and did invalid shared writes on sm_89. We now lint our PTX for that truncation. In another, the loader preferred an architecture-specific cubin we had never measured, so on RTX 50 cards every #36 layer would have quietly fallen back to cuDNN. We caught that one before it shipped.

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`

We 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 we 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 we’ve tested, so the libraries don’t load there either.

Tuning per GPU

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 our FP16 kernels were slower than our 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.

So kernel choice is per device. Each layer’s kernel is picked in this order:

a tune file written by speakrs cuda tune for this GPU, driver and build
a built-in recipe measured on that GPU, matched by exact device name and SM count
a default for the device class

speakrs cuda tune times each layer’s candidates on your GPU, which took 18–26 seconds on the RTX cards we tried. It only picks kernels that are allowed in the selected math mode and have end-to-end accuracy evidence. Results on the first 10 VoxConverse test files, with process start-up and model loading included in the time:

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×

Every run had 16 CPU cores, and each value is the median of 3 runs. The GPU-specific builds run within about 1% of cuda, so going driver-only gives up almost no speed. In an earlier build, on the same subset, 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.

When we 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

How the agents did the work

My job was the goals and the acceptance checks: which GPUs to support, what “faster” had to mean, and what accuracy loss was acceptable: none, meaning byte-identical RTTMs or DER that didn’t get worse. I also decided what to merge and what to drop. 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.

The biggest lesson is about 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”. Proving 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 we merged one PR, #36, whose three kernels were built for sm_75 and ran unchanged on every GPU. The end-to-end benchmark for replacing the libraries was scheduled for the last phase.

Then we 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. That included the A100 fix, the driver-only proof on four GPUs, the per-GPU recipes and the tune command, Winograd for the wide layers, and the T4 beating cuDNN.

Most of the kernels existed in the first days. The harness was the bottleneck, and nothing capped it. It did catch real bugs, including the unmeasured cubin and the truncated shared pointer above, but a cheap end-to-end check would have caught most of them, or made them harmless, 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.
Put a time limit on harness work. 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 for Winograd and the A100 confirmation cost $0.21.

We also asked several AI advisors to forecast what each idea would gain, then scored them against what we measured. The common mistake was turning a peak-rate ratio, like “FP16 is 2× TF32”, into an end-to-end gain without checking how much of the runtime that layer takes. Every advisor got the RTX 5060 Ti FP16 result wrong. The most useful suggestion was a cheap experiment instead of a forecast: round the operands through FP16 inside the existing TF32 kernels and check whether the output changes. It didn’t change on the RTX 4060 Ti, and that led to the T4’s FP16 kernels.

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 we haven’t tested, I’d like to see your RTFx and tune output in an issue.