Beating hipBLASLt on a Gaming GPU

The setup

The RX 7900 XTX is a gaming card. It also happens to be one of the few ways to get 24 GB of VRAM without paying datacenter prices, which makes it an obvious thing to want to run inference on. The problem is that AMD's math libraries were not written with it in mind. hipBLASLt's tuning effort — the Tensile libraries that pick tile sizes, split-K factors, and workgroup mapping for a given GEMM shape — is concentrated on CDNA, the compute-only architecture in AMD's Instinct line. gfx1100 (RDNA3, the architecture under the 7900 XTX) ships a small fraction of that library coverage.

That asymmetry is a testable claim, not a vibe. If the vendor library is thin on this architecture, a hand-written kernel targeted specifically at gfx1100 should be able to close some of the gap, at least on some shapes. mojo-baro is an inference stack built to find out how much — GPU kernels written in Mojo, benchmarked against hipBLASLt through the same shim, on one card. This post covers the second of its two kernel results: a square fp16 GEMM.

The result

The kernel is a pipelined WMMA (wave matrix-multiply-accumulate) GEMM: 4x2 warps over a 128x128 output block, two LDS buffers so the next tile of A and B loads while the current one is being consumed, one barrier per K-step, and a two-deep global prefetch — tile k+2 sits in registers while tile k+1 waits in the second LDS buffer. A is stored XOR-swizzled to avoid bank conflicts; B is stored transposed. The whole thing compiles to 188 VGPR with zero register spills and fits in 32 KB of LDS, which matters because RDNA3's 64 KB LDS-per-WGP budget means a 32 KB kernel gets two blocks resident per workgroup processor instead of one.

Run across ten square sizes, 256³ through 4096³, against hipBLASLt fp16 through the same C++ shim, with a 10-second clock warm-up before every measurement:

size	256	512	768	1024	1536	2048	2560	3072	3584	4096
ours (GFLOP/s)	6372	30642	64204	74824	93632	91300	97974	99307	105786	90705
hipBLASLt (GFLOP/s)	6288	26324	54924	63623	69332	80203	87147	97224	85671	82437
ratio	1.01	1.16	1.17	1.18	1.35	1.14	1.12	1.02	1.24	1.10

The kernel is ahead of hipBLASLt at every size tested, from a rounding error at 256³ up to 1.35x at 1536³. That is the whole claim: ahead of hipBLASLt at these ten square fp16 sizes, on this card, reproducible with bench/fp16-templates.sh. Not the fastest GEMM in existence, not a general result about RDNA3 versus CDNA — one card, one kernel, one shape family, measured with a fixed protocol.

The part that actually mattered

The table above is the outcome. The interesting part is what it took to get there, because the first version of this kernel did not beat hipBLASLt — it lost, and lost badly, at every size at or below 1024³, running at 0.78-0.98x of the vendor. Same pipelining, same swizzling, same prefetch depth. The only thing that changed between that version and the one in the table is that tile geometry stopped being fixed.

A 128x128 output block is a lot of work per thread block. At small problem sizes that is a liability: the whole point of a big tile is to amortize LDS traffic and keep warps busy, but if the grid only has a handful of 128x128 blocks to hand out, most of the GPU's compute units sit idle waiting for work that was never divided finely enough to reach them. The kernel was correct and well-optimized for the tile it was given; the tile was wrong for the problem.

The fix was to make tile shape a compile-time kernel parameter — 128x128, 64x128, or 64x64, each with its own warp layout — and to have the launcher pick one based on how many 128x128-equivalent blocks the grid would actually produce for a given M and N:

blocks in grid	tile chosen	warps
≥ 96	128x128	4x2 (8 waves)
≥ 64	64x128	2x4 (8 waves)
< 64	64x64	2x2 (4 waves)

That single change — not a new optimization, just picking the right existing tool for the size in front of it — is what flipped the small sizes from behind hipBLASLt to ahead of it. It is the kind of result that is easy to miss if you only ever benchmark one size and call it done: a kernel that is a clear win at 4096³ can be a clear loss at 512³, and the two facts don't average into anything meaningful. One tile shape does not win everywhere.

Why these numbers should be trusted

Two things happened earlier in this project that are worth stating plainly, because a benchmark post that only shows wins invites the question of what got left out.

The first is a retracted claim. An earlier version of the project's baseline document recorded a different kernel — a register-tiled fp32 GEMM — as roughly 2x faster than hipBLASLt. That number was real in the sense that it was measured, and wrong in the sense that it was measuring an untuned vendor call. hipBLASLt was never given a fair shot: the shim allocated a fresh workspace on every call instead of caching one, it trusted the order the tuning heuristic returned candidates in rather than timing them, and it never touched splitK or wgm, two parameters that are only reachable through the hipblaslt_ext extension API, not the plain C API. Fixing those three defects took hipBLASLt from 2497 to 5201 GFLOP/s at 512³ on the same problem, and the 2x lead disappeared entirely. The corrected result — the two kernels trading places by size, roughly tied overall — is a smaller and less exciting claim, and it's the one that's true.

The rule that came out of that: a vendor baseline that looks easy to beat is a bug in your harness until proven otherwise. It is worth restating for the fp16 numbers above, because the same failure mode was available here too and the shim was built to avoid it from the start — cached workspace, all heuristic candidates timed, splitK/wgm set explicitly.

The second is about clock warm-up, which turned out to be worth more than any single kernel change at the high end. GPU clocks ramp from idle over roughly 400 milliseconds of sustained work and hold at their ceiling only as long as nothing interrupts them. At a 1-second warm-up before measurement, the identical fp16 kernel binary reported 66,000 GFLOP/s at 4096³. At a 10-second warm-up, the same binary, same build, reported 90,705. The clocks simply hadn't settled. Every bench in this project now warms for 10 seconds before it starts timing anything, and logs the warm-up duration (warmup_s) alongside the result so a future run can't silently regress this without it showing up in the receipt.

Method

Every number that goes into this project's baseline document is preregistered before the run that produces it: the question being asked, the instrument measuring it, the predicted range, and the condition that would falsify the prediction are all committed first. The result is recorded against that prediction whether it matched or not — a missed prediction stays in the file rather than getting quietly dropped. The benchmark harness (bench/run.py) also refuses to report a throughput number until it has checked the kernel's output against a correctness reference; a fast wrong answer doesn't get a GFLOP/s figure at all.

What this is not

This is one card. RDNA3 specifics — a 32-wide warp instead of CDNA's 64, a 64 KB LDS budget per workgroup processor, tile and occupancy parameters swept specifically for that memory system — are load-bearing throughout the kernel. None of the numbers here should be assumed to carry over to gfx942, gfx950, or an NVIDIA GPU without re-running the sweep on that hardware; the project treats "measured on this machine" as the actual scope of every claim, not a formality.

It's also worth flagging a comparison that doesn't exist and shouldn't be constructed: fp32 WMMA is not implementable on gfx1100 at all — it's an instruction set limitation, not a Mojo one, confirmed against llvm-mc's accepted mnemonics and LLVM's builtin table for the architecture. So the fp16 numbers in this post and any fp32 GEMM numbers from this same project are not comparable to each other. They run on different hardware paths inside the same chip.

The code, the benchmark harness, and the protocol files with the full history of what was tried and rejected are in mojo-baro (github.com/amarbaro/mojo-baro).