How I placed 5th in the GPU MODE qr_v2 competition with common sense engineering, linear algebra tricks, an agent swarm, and a dream.
I recently placed 5th in GPU MODE's qr_v2 competition, where a bunch of kernel writers and I raced to make batched QR decomposition lightning fast on NVIDIA's B200. If you missed the recent Twitter drama, QR has been getting attention because it can show up inside various rivaling deep learning optimizer algorithms. The competition made one thing very clear: the default PyTorch path is really slow for this workload.
My final submission reached a 1499µs geometric mean on the validation set of 12 sizes, roughly 280× faster than torch.geqrf. Along the way I learned the B200 architecture from the ground up and found a few almost-heretical math tricks that make Householder QR much faster than I initially assumed possible.
If you have never heard of Householder QR decomposition, then you're like me two weeks ago. I will build up the algorithm from scratch, draw the B200 GPU architecture in enough detail to make the bottlenecks visible, and then walk through the experiments that moved my score.
This is also a story about agentic research orchestration. I'd like to share the orchestration workflow that helped me learn the math, launch nearly 2,000 focused experiments, and generate custom HTML reports that made the next research direction clear. Join me on this journey to learn about fun linear algebra tricks and glimpse what kernel autoresearch looks like in mid 2026!
Before we get into the details, here's what the SOTA frontier looked like over the course of roughly 1,900 experiments.
The math starts simply. QR decomposition takes a matrix and splits it into two pieces: an orthonormal matrix \(Q\), and an upper-triangular matrix \(R\).
\[ A \;=\; \underbrace{\vphantom{R}Q}_{\substack{\text{orthonormal:}\\ Q^{\top}Q\,=\,I}}\;\underbrace{\vphantom{Q}R}_{\substack{\text{upper}\\ \text{triangular}}} \]QR matters because it is a stable way to solve least-squares problems and orthogonalize vectors. A triangular \(R\) makes systems solvable by cheap back-substitution, and an orthonormal \(Q\) inverts for free (\(Q^{-1} = Q^{\top}\)) without amplifying numerical error. A least-squares fit, for example, becomes \(x = R^{-1}Q^{\top}b\): no normal equations, no squared condition number. QR also shows up in ML systems more often than I appreciated at the start: orthogonalized optimizer updates (the idea behind Muon and its descendants, now training frontier models like Kimi K2), GaLore-style gradient projections built on periodic SVDs, and the randomized sketches inside low-rank methods all involve solving batched QR problems.
With actual numbers, the structure is easier to see. \(R\) keeps the upper triangle and zeroes everything below the diagonal. The columns of \(Q\) have unit length and are mutually perpendicular, so multiplying by \(Q\) rotates or reflects space without stretching it. For example, \(\lVert q_1\rVert^2 = \tfrac{4}{9}+\tfrac{1}{9}+\tfrac{4}{9} = 1\), and \(q_1 \cdot q_2 = -\tfrac{4}{9}+\tfrac{2}{9}+\tfrac{2}{9} = 0\).
\[ \underbrace{\begin{pmatrix} 2 & 0 & 3 \\ 1 & 3 & 6 \\ 2 & 3 & 3 \end{pmatrix}}_{A} \;=\; \underbrace{\begin{pmatrix} \tfrac{2}{3} & -\tfrac{2}{3} & \tfrac{1}{3} \\[2pt] \tfrac{1}{3} & \tfrac{2}{3} & \tfrac{2}{3} \\[2pt] \tfrac{2}{3} & \tfrac{1}{3} & -\tfrac{2}{3} \end{pmatrix}}_{Q,\;\;Q^{\top}Q\,=\,I} \; \underbrace{\begin{pmatrix} 3 & 3 & 6 \\ 0 & 3 & 3 \\ 0 & 0 & 3 \end{pmatrix}}_{R} \]Householder QR is a QR algorithm built from one simple operation: reflect one vector from the starting matrix across a high-dimensional line, then repeat that move to knock out the lower-triangular terms.
Drag the slider to rotate the mirror. The yellow vector \(x\) stays fixed. The blue vector is its reflection. The purple segment is always perpendicular to the axis of reflection, with length 2× the distance from \(x\) to that axis. From this geometric observation, it is easy to figure out what the reflected vector is.
One special angle makes the reflection land exactly on the first axis. That mirror is the Householder reflector for this vector.
The idea: walk across \(A\) left to right. Reflect column 1 onto the axis, so everything under the diagonal becomes zero. Move to column 2 and do the same using only the part below the diagonal. Finished rows are never touched again. After \(n-1\) columns the matrix is upper triangular. That is \(R\). And conveniently, reflections are orthonormal and stack multiplicatively, so the stack itself is the other factor: \(A = (H_1 H_2 \cdots H_{n-1})\,R = QR\).
The stepper above is correct, but it works one skinny column at a time: dot products and small vector updates, the exact work a GPU is worst at. The fix used by every fast QR is blocking: factor a few columns first, then hit the rest of the matrix with all of their mirrors at once, as big matrix multiplies. One guarantee makes this legal: the green cells from the previous demo are finished. No later mirror ever touches a finished row or column, so batching mirrors up cannot disturb them.
First, notation. Stack the saved normals side by side as the columns of one matrix, which everyone calls \(Y\). Thanks to the pinned 1s, \(Y\) is exactly the lower staircase from the packing step: each column has a 1 in its diagonal slot and its tail below.
Now try to apply two mirrors in one shot by multiplying them out:
\[ H_1 H_2 = (I - \tau_1 v_1 v_1^{\top})(I - \tau_2 v_2 v_2^{\top}) \] \[ = I - \tau_1 v_1 v_1^{\top} - \tau_2 v_2 v_2^{\top} + \tau_1\tau_2\, v_1 (v_1^{\top} v_2)\, v_2^{\top} \]The first two terms are each mirror acting alone. The last term is the catch: the mirrors overlap, and the overlap is the single number \(v_1^{\top}v_2\). Drop it and the answer is simply wrong. So a block of mirrors is not the sum of its parts. Something has to carry the cross terms.
That something is a small matrix \(T\). Collect the pieces:
\[ H_1 H_2 = I - Y\, T\, Y^{\top}, \qquad T = \begin{pmatrix} \tau_1 & -\tau_1\tau_2\, v_1^{\top}v_2 \\ 0 & \tau_2 \end{pmatrix} \]Multiply it back out to check: the diagonal of \(T\) reproduces the two solo mirrors, and the corner entry is exactly the cross term. With \(b\) reflectors the same thing happens: \(T\) is a \(b \times b\) upper triangle whose entries soak up every pairwise overlap \(v_i^{\top}v_j\). The block apply becomes two big multiplies with a tiny triangle in the middle,
\[ C \;\leftarrow\; C - Y\, T\, (Y^{\top} C), \]and \(Y^{\top}C\) plus the outer multiply are exactly the dense matrix-matrix work tensor cores want.
One more trick, the one my fastest kernels used. Building \(T\) is a serial chore: each new column of \(T\) depends on all the previous ones, so it is a chain of \(b\) small dependent steps, exactly the kind of work this whole section is trying to escape. But look at what \(T\) is made of: every entry is some overlap \(v_i^{\top}v_j\). All of those overlaps live in one object, the Gram matrix \(G = Y^{\top}Y\), and computing \(G\) is a single matrix multiply, tensor-core food like everything else here. Better still, \(T\) never needs to be built at all, because its inverse can be read straight off \(G\): the strict upper triangle of \(G\), with \(1/\tau\) on the diagonal.
\[ T^{-1} = \begin{pmatrix} 1/\tau_1 & v_1^{\top}v_2 \\ 0 & 1/\tau_2 \end{pmatrix} \]Multiply this against the \(T\) above and everything cancels to \(I\), so it really is the inverse. Why does holding \(T^{-1}\) help? Because the update only ever needs the product \(W = T\,(Y^{\top}C)\), and when you know the inverse, a product becomes a solve. Mechanically the block update is now three steps:
\[ W_0 = Y^{\top} C \qquad \text{(big multiply)} \] \[ T^{-1} W = W_0 \;\Rightarrow\; W \qquad \text{(tiny triangular solve, } T^{-1} \text{ read off } G \text{ and } \tau\text{)} \] \[ C \leftarrow C - Y\,W \qquad \text{(big multiply)} \]Nothing serial survives except the tiny solve, and the solve's matrix costs nothing to form. That is what solve_gram_tau32 does in the final kernels, and on late panels the whole sandwich, RHS, solve, and update, runs fused in one launch. Not every experiment took this route: for weeks the campaign ran explicit \(T\) against Gram-plus-solve, and Gram is the side that survived into the final submission.
Panel 1 is factored with the simple reflector loop. The wide unfinished block waits.
Two variations became important later. First, even the panel factorization can be broken into smaller panels: zero the first slice, apply its update to the rest, then continue. Second, for tall panels, you can sometimes form the Gram matrix \(A^{\top}A\), Cholesky-factor it, and recover the reflectors from a triangular solve. That second route is faster when it works, but it spends numerical margin, so it needs guards.
Now the competition contract is easier to state. The leaderboard did not ask for an explicit \(Q\) and \(R\). It asked competitors to match the compact LAPACK form returned by torch.geqrf. The upper triangle stores \(R\), the entries below the diagonal store the Householder vectors, and the separate \(\tau\) buffer has length \(n\), matching torch.geqrf. Only the first \(n-1\) entries are nontrivial reflectors. From those two buffers, the checker rebuilds \(Q = H_1 H_2 \cdots H_{n-1}\).
Scoring used a validation set of twelve cases. Each case was a batch of square fp32 matrices on the B200. The score was the geometric mean of the twelve runtimes, so every shape mattered multiplicatively:
\[ \mathrm{score} = \left(\prod_{i=1}^{12} t_i\right)^{1/12} \]This is why the experiments below are so shape-specific. A real win on \(n=1024\) could vanish if it made \(n=512\) worse, and a tiny-looking launch-overhead fix could matter a lot on \(n=32\).
| Case | Batch | Why it mattered |
|---|---|---|
| \(n=32\) | 20 | Too small for much math. Launch overhead and routing dominate. |
| \(n=176\) | 40 | Small enough that wasted panel codegen shows up immediately. |
| \(n=352\) | 40 | The bridge between tiny custom paths and the larger blocked kernels. |
| \(n=512\), dense | 640 | The high-batch workhorse. Regressions here were expensive. |
| \(n=1024\), dense | 60 | Large enough for tensor-core updates, still frequent enough to drive the score. |
| \(n=2048\), dense | 8 | Low batch count. The GPU can starve unless the kernel exposes enough work inside each matrix. |
| \(n=4096\), dense | 2 | The giant case. It rewarded different math more than another round of panel tuning. |
| \(n=512\), mixed | 640 | Looked like the dense case, but precision shortcuts could fail here first. |
| \(n=1024\), mixed | 60 | Forced the routing logic to know when fast math was still safe. |
| \(n=512\), rank-deficient | 640 | Had exploitable structure. The right answer was to stop doing unnecessary columns. |
| \(n=512\), clustered | 640 | Another structured family where cheap detectors could unlock a specialized path. |
| \(n=1024\), near-rank | 60 | The tail could be synthesized after factoring the live prefix. |
Correctness was not a suggestion. The checker rebuilt \(Q\) from \((H,\tau)\) in fp64, then checked both \(QR \approx A\) and \(Q^{\top}Q \approx I\). The tolerances were purely relative: factorization residual \(\le 20\,n\,\varepsilon\), orthogonality \(\le 100\,n\,\varepsilon\). The returned factors were fp32, but at the boundary they still had to satisfy a strict fp64 verifier.
The B200 matters because it strongly prefers matrix-shaped work. It is two reticle-sized dies presented as one logical GPU. Each SM has vector lanes for scalar and skinny work, plus fifth-generation tensor cores for dense matrix tiles. If the algorithm cannot be expressed as matrix-matrix work, most of the machine sits outside the hot path.
Figures are from NVIDIA's Blackwell architecture materials and CUDA device queries on the competition B200 runners (compute capability 10.0). This is a schematic, not a floorplan. The practical rule of thumb: keep ~148 independent chunks of work live, and enough warps inside each chunk to hide memory latency.
Each tile is one SM. Click any tile to zoom into the parts the kernels fought over: warp schedulers, registers, shared memory, vector lanes for mat-vec work, and tensor cores for mat-mat work.
Click any SM tile to open it up.
click a component inside the SM
A kernel launches thread blocks. Each block lands on an SM and executes as warps of 32 threads. In the mat-vec lane, bytes stream in, get used once, and leave. The lane looks busy while the math meter crawls. In the mat-mat lane, a tile parks in shared memory and the tensor core reuses the same bytes over and over. Similar traffic, very different work done. Householder panel factorization is the top lane. The blocked trailing update is the bottom lane.
Three constraints shaped almost every experiment.
Each reflector waits for the last: norm, projection, update. Once trailing updates became fast GEMMs, panel time became the floor. I once tested the extreme (w322): one raw CUDA kernel per matrix, zero launch overhead, full unblocked QR. It ran 459× slower. Launch count was never the real cost. Giving up tensor-core-shaped updates was.
n=2048 and n=4096 have tiny batches. During panel work, there are far fewer independent jobs than 148 SMs.
At n=32, the math ends quickly. Launch overhead dominates, so CUDA graph capture mattered more than another tiny kernel tweak.
Here is blocked Householder placed onto the B200. It alternates between two modes. During panel factor, each matrix runs roughly one resident panel kernel, about one CTA and one SM, because panel columns form a dependency chain. During trailing update, the remaining \(m \times m\) block becomes independent 128×128 GEMM tiles. Pick a shape and flip phases. Small shapes fill the machine through batch count. n=4096 at batch 2 is the hard case: the factor phase runs on 2 of 148 SMs, and the update work shrinks as the factorization progresses.
Over fifteen days I ran roughly 1,900 numbered experiments, with 652 logged in enough detail to compare scores. The chart below shows the changes that improved the frontier and the negative results that redirected the search.
The torch.geqrf baseline scored a 419 ms geometric mean on the hosted B200. A first custom CUDA compact-Householder kernel, then a panel-blocked variant, then per-shape dispatch cut it 5× in a day. The lesson was basic: the library path was not built for this exact mix of small batched QR problems. I kept checking whether cuSOLVER or a related library had a batched compact-QR primitive I had missed. It did not.
The first stable architecture was blocked-WY Householder: 32-column panels for n=512 (w33), tiled trailing updates (w47), the n=1024 port (w52), multi-CTA panels (w76), and Triton tl.dot updates (w100/w101). This is where the diagrams above became code.
The old path launched a small kernel chain for every panel column: norm, tau, project, update. Between launches, panel state bounced through global memory. w194 rewrote that path from scratch. One Triton kernel holds the whole 32-column panel in registers and shared memory, generates all 32 reflectors in a single launch, and hands the wide trailing update to two library GEMMs. Same math, far fewer launches. The score nearly halved in one step.
With the panel path cleaned up, smaller improvements started landing: compile-time-constant panel offsets (w198), lazily built helper kernels (w210), a custom 32-reflector solve (w222), TF32 for the n=1024 update GEMMs (w231). Individually minor, and mostly invisible in the final submission: the T-matrix machinery got restructured away, and TF32 was later overshadowed by fp16-resident state. Still, the stretch halved the score again in three days. This is the part of the project that only works when iteration is cheap.
The n=512 and n=1024 cases were not one distribution. They came as dense, clustered, rank-deficient, mixed, and near-rank families. Cheap detectors routed them to specialized paths (w259, w317, w537). In parallel, the trailing state moved to fp16 residency (w389–w499): halve memory traffic where tolerance allowed, keep reflector generation and tau in fp32 where it did not.
The blocked scheme still leaves one stubborn serial lump: the 32-column panel factor. The trick that kept paying was to apply the blocking idea to the panel itself. Factor 16 columns, hit the other 16 with their compound update, a small matrix multiply instead of a scalar recurrence, then factor those. Why is two stages faster than one straight factor-32? Registers. A 32-wide factor keeps twice the live state per thread, and that costs occupancy exactly where latency-hiding matters most. Halving the width kept more warps resident, and the cross-half work became tensor-core shaped. n512 panel emission moved about 3.3× in isolation (w290–w292). Repeating the move on n1024 took captured factor time from 2.64 to 1.88 ms with a 16+16 split (w529), and the shipped kernel ends at an 8+8+16 staircase (w642, w664): the exact three factor kernels visible in the NCU strip below. The idea is convergent: the fourth-place finisher's published repo uses the same recursion as split trees like 96→48+48.
Six hundred experiments bought 11%. Most of the work was small and local: panel widths, register caps, packed copies, n4096 projected tails (w703), and factor-emitted Gram/Y kernels (w1010). The long sweep around w850–w896 is a useful example of what agents are good for: forty-six cutoff tests in one direction. The hard part was deciding whether the direction deserved that many tests in the first place.
One thread ran the whole campaign: the library triangular solve was never the right tool, at any scale. w222/w229 swapped it for a hand-written 32-reflector recurrence (5.40 to 4.82 ms hosted). w343 extended that to n176/n352 (both cases about 37% faster). w297 settled the Gram-vs-T fork: delete the explicit T matrix entirely and solve straight from the Gram matrix \(Y^{\top}Y\) plus tau in one fused kernel (ranked 4.213 to 3.813 ms). The final submission runs this solve_gram_tau32 path on every mid shape, fused together with the RHS build and trailing update on late panels. And w1431 did it one last time on the large shapes: a one-CTA triangular inverse plus GEMM against cuSOLVER's serialized solve, 202µs vs 619µs at batch 8. Only n32 escapes the pattern, by never solving anything at all. Independent confirmation came after the deadline: the second-place finisher's Discord writeup calls out the same library solve and the same fix, custom blocked inverse kernels.
Some frontier steps mattered because they made future experiments usable. Around w343–w379, a fast small-shape route kept passing hosted evals and failing ranked submissions. I pulled it and accepted the slower score, because reliable leaderboard attempts were worth more than the microseconds. Two hard-won corollaries: the eval ledger itself can lie, since a polling timeout records a failure even when the submission actually landed and scored, so ranked outcomes got confirmed against the public board directly. And a passing submission is not a safe submission: one route (w466) passed 6/6 runs with a public score of 2.59 ms and a secret score of 9.81 ms, which is why a slower stability reference stayed on the shelf next to every fast artifact.
The same logic applied at w1406. The ranked evaluator had a 240-second test phase, and my constexpr-specialized kernels needed 267 Triton compiles, 354 seconds cold. Making panel offsets runtime arguments where specialization was not paying cut that to 100 compiles and 96 seconds. It barely moved the score, but every later ranked submission depended on it.
Two ideas that had failed repeatedly landed within 48 hours. CholeskyQR-HR, the Householder-reconstruction variant of the Gram-matrix idea, worked for underfilled n=2048/n=4096 once it had conditioning guards (w1420). CUDA graph replay (w1422) removed the launch overhead of roughly 270–560 kernels per factorization, moving 1921→1827µs and giving later changes a faster base to build on.
Once graphs made launches cheap, the profiler pointed at library calls inside the large-shape route. A one-CTA-per-matrix triangular inverse beat cuSOLVER trsm 3× at batch 8 (w1431). Fused unpivoted LU beat torch.linalg.lu 1.9×. TF32 routing (w1664) and fused sign/pack (w1864) closed the ledger at 1537.85µs. The final board run landed at 1499.09µs, 5th place, 280× from baseline.
Where the final submission landed, case by case. The bars are hosted B200 times for the last promoted lineage (w1864), on a log axis. The dashed line is the geometric mean, which is the score itself: 1535.8µs hosted, 1499.09µs on the final board.
One level deeper than the bars: every kernel type in the n=512 and n=1024 dense routes, one NCU-measured firing each, laid end to end. Hover any segment (or its legend entry) to see exactly what you are looking at, including its register count. Orange is custom panel-factor work, yellow is the Gram/tau solve, blue is library GEMMs, grey is casts and copies, purple is the route detector. Honesty note: this capture is from the w1359 generation, whose kernels match the final source name-for-name. The shipped kernels repeat these types many times per case. The n2048 route was never profiled end to end, so it gets no strip, only its two surviving leg numbers (LU 432µs, tri-inverse 202µs).
The main infrastructure lesson is boring and important: the repo had to make the next experiment obvious. Every agent needed to know the current best version, what had already failed, and what evidence would count as progress.
Every experiment lived in submissions/research/<id>/ with its evidence. Only the ranked leaderboard could promote into submissions/sota/. Local A/B runs were fast enough to drive iteration. Hosted B200 runs sanity-checked the direction. The ranked board was the final gate, not the steering wheel. 377 commits in 15 days recorded it all, and the standing rule was to never block the next experiment. PRIORS.md, a 2,400-line belief file, kept each new agent from repeating old mistakes.
Every eval appended to performance.jsonl, and dashboards regenerated per-case bars and the best-so-far frontier after every run. The point was not decoration. With twelve cases and a thousand experiments, the only way either I or an agent could cut to what actually mattered was evidence reduced to one glance: which bar moved, which did not, and when it changed.
nsys/ncu profiles sat next to the experiment that produced them. The loop that worked was simple: read the profile, form one directed bet, and measure exactly that axis. The register cap at w563 is a clean example. NCU showed 183 registers and 12.5% occupancy, the bet was maxnreg=128 on that one launch, and the measurement was that kernel alone.
One more move paid for itself many times: a dedicated rewrite for readability. No private methods, short docstrings, and an outline of the algorithm and math at the top of the file. When the code read top-down, agent patches got sharper and my reviews got faster. That rewrite also exposed the w194 jump.
None of what follows ever moved the frontier chart, which is exactly why PRIORS.md existed: a conclusion that stops future work from being repeated is worth as much as a win. A sample of what it held:
I tried a visual ideas DAG, and it was useful while the tree was small. Once there were hundreds of submissions, the graph became too unwieldy to be the main interface. What actually worked was a frontier of priors: short, stable conclusions in PRIORS.md that told the next agent which directions were alive, which were dead, and why. An auto-generated prior frontier would probably be better. For this project, the hand-maintained file was enough.
The repo also had a rule for killing ideas: no experiment could be closed without a code audit, raw logs, and a hardware-level explanation. That rule was not bureaucracy. CholeskyQR was rejected too early and only came back after a forced re-read of the logs.
This is the view I needed before promotion: candidate versus incumbent, case by case. The first toggle shows w537, a clean one-shape win. The second shows w132, the more important reject: rankdef got faster and the aggregate improved, but mixed slowed, so the route was not isolated yet.
The best prompts were not "make it faster." They looked more like contracts: here is the per-case table you must not regress, here is the profile showing the kernel is register-bound, and here is the code path you are allowed to touch. A specific objective gave the agent room to iterate. A scalar score gave it room to cheat the intent.
1. The single-orchestrator pattern fails quietly. Asking one orchestrator agent to "manage subagents toward a faster kernel" looked good on paper and failed in practice. This is not inherent to orchestration. Current models just were not good enough at the orchestration taste I needed: keeping the target narrow, forcing a plausible idea through ugly implementation cycles, and rejecting it only after the hardware argument died. The subagents drifted off-objective, evaluated along convenient metrics, introduced bugs, and stopped too early on ideas that still had plausible hardware arguments.
2. One strong model in a tight loop beat elaborate orchestration. I ran both Claude Opus 4.8 and GPT 5.5 Codex. Codex was marginally faster to iterate with, and it won most sessions here. Multi-agent workflows gave me breadth, but they also produced shallow experiments and not enough careful follow-up. The setup that worked was one Codex session at a time, pointed at one direction I had chosen.
3. Separate worktrees beat vague delegation. What worked was sharper: pick a few concrete directions, give each one a separate worktree, then merge the survivors back into the clean timeline. "Make the n=1024 panel fit in fewer registers" is a direction. "Get CholeskyQR past conditioning failures" is a direction. Cholesky panels and graph capture each took double-digit failed iterations before landing. My job was to keep those directions alive when the recent diffs were bad but the hardware argument still made sense.
4. "Improve the score" is a bad objective. The geometric mean hides regressions. An agent chasing the aggregate can win 8% on n=1024 while losing 5% on n=512, and you have burned a day learning very little. The useful objectives were narrower: improve n=1024 dense with zero regression elsewhere. Under that constraint, agents did something genuinely useful. They separated code paths until the change could not touch unrelated shapes. The complementary move is doing the geometric-mean arithmetic before choosing a lane at all: at w324 I computed exactly how much cumulative improvement the podium required and proved single-bucket wins could not reach it, no matter how good. That arithmetic, not any profile, is what forced the CholeskyQR and graph-capture pivot.
5. Presentation to me is part of the system. Per-shape charts, frontier plots, and linked profiler traces turned "read 40 logs" into "look at one dashboard for 90 seconds." Most of my useful interventions began as something visible: a shape whose bar had not moved in days, or a profile where the gaps between kernels were larger than the kernels.
6. Force the code to be skimmable. It is easy to stare at outputs and profiles while avoiding the actual submission file. That is a mistake. I had to force agents to produce code in patterns I would want to read: clear files, real modules, consistent function names, and routes that could be skimmed top-down. The 30-minute refactor that makes the current SOTA reference point readable usually pays for itself. Review gets sharper, the next agent stops patching around ghosts, and every fork starts from code that is still readable.
After a leaderboard ends, GPU MODE exposes submission code to logged-in users, so I could read what beat me. Two things stood out. First, everyone at the top specialized to the benchmark shapes. The organizers accepted that and policed correctness instead. Second, the podium split around the core technical question in this post: how much Householder QR do you replace with Gram-matrix Cholesky methods, and how do you survive the precision cliff?
| Rank | Score (geomean) | Approach | Vs. mine |
|---|---|---|---|
| 1 · 10billiontokens | 1220.8 µs | Pure Triton, 12.6k lines. A CholeskyQR–Householder hybrid on nearly every shape: register-resident 16×16 block Cholesky fused with compact-Householder recovery, guarded by precision-classifier kernels that route ill-conditioned panels to Householder. Survives fp16 precision with compensated arithmetic: three-fp16-MMA emulated fp32 dots, TF32×2 splits, MXFP8 two-term dots. | I reached the same destination, guarded CholeskyQR, but only for n≥2048. The winner made it the backbone, spending the precision budget with emulated-precision MMAs where I spent it with fp32 fallbacks and selective TF32. |
| 2 · gau.nernst | 1227.7 µs | Proof the other fork also works: pure blocked compact-WY Householder, no Cholesky anywhere, and no input-structure detection at all. Per his post-competition summary in the GPU MODE Discord: three warp-specialized panel kernel families, all producer-consumer (one warp makes reflectors while others apply them): 1 SM per matrix (n32/512), 2 SMs per matrix via thread-block clusters with a shared-to-cluster TMA broadcast (n176/352/1024 and large-shape tails), and multi-SM with reflectors through global memory and a polled flag (n2048/4096). Custom blocked triangular-inverse kernels instead of torch's solve (16×16 diagonal inverses, then doubled 32→64 by block substitution), cuBLAS/cuBLASLt GEMMs with per-shape TF32/fp16 precision, and hand-tuned panel-width schedules like n=512's (96,96,64,32,32,192). | Structurally the closest to my mid-shape routes, and the most instructive contrasts. His 2-SM cluster panels worked where my cluster/DSM attempts stopped at parity. The difference is warp-specialized broadcast of reflectors rather than trying to widen the recurrence itself. We independently made the same discovery about torch's triangular solve and both replaced it with custom inverse kernels. And where I leaned on structure detection, he used none, tuning precision per shape instead. His note that the n512 tests are the stress-tested ones matches my n512 precision cliff exactly. |
| 3 · dhu.randhar | 1279.3 µs | A 14.5k-line machine-generated file inlining ~170 CUDA kernels using Blackwell tcgen05 tensor-core MMA. Exact-fp32 blocked-WY for all n≤1024. CholeskyQR pipeline only for the scored dense n=2048/4096. Structure detectors with fused column-norm probes and multi-seed validation. | The same large-shape-only Cholesky conclusion I reached, plus a lesson I didn't act on: first-class use of tcgen05/TMEM, the Blackwell-native MMA path my Triton kernels never touched. |
| 4 · zhongmingee | 1438.9 µs | 25.8k lines of NVRTC source-JIT: CUDA compiled lazily on the target GPU, sm_100 tcgen05/TMEM kernels, 3×TF32 emulated-fp32 Gram feeding CholeskyQR, everything wrapped in CUDA graph capture with per-route memory pools. His post-competition repo (fishmingyu/qrv2-gpu-mode, adapting gau.nernst's wide-panel kernel design with credit) documents the rest: recursive panel split trees like 96→48+48, per-matrix classify-permute-restore for mixed batches instead of sampling representatives, PDL enabled only where it pays, and a guarded CholeskyQR path whose ALGORITHM.md insists the conditioning guard "is part of the algorithm's numerical domain, not merely a performance heuristic." | Reached the same endgame ideas I did, but pushed both harder: graph replay everywhere and Cholesky with explicit precision emulation. Two mirrors worth noting: his split trees are the same recursion as my 8+8+16 staged panels, arrived at independently, and his guard-is-correctness framing is exactly the argument behind my w1420 conditioning guards. One real difference: his mixed-batch route classifies and permutes every matrix individually, where mine detected at batch level. |
| 5 · michaelmelons (me) | 1499.1 µs | Blocked-WY Householder with fp16-resident state and value-based structure routing for small and mid shapes. Graph-replayed, conditioning-guarded CholeskyQR-HR with custom LU and triangular-inverse kernels for n≥2048. Roughly 1,900 experiments of accumulated micro-decisions. | (this row is me) |
The lesson is not that one approach was obviously right. Second place proved a pure Householder route could be excellent. First and fourth place proved that Cholesky could carry much more of the workload if you engineered around precision hard enough. My mistake was staying between those theses until the final weekend. The system was good at local search and honest evaluation, but I was late to force the bigger architectural bet.
The surprising part was not that agents magically wrote a fast kernel. They did not. The surprising part was how much faster I could learn when the loop was tight: profile, form a hypothesis, fork the code, measure the result, and write down what changed.
Householder QR is old math. The B200 is new hardware. The research lived in the gap between them, where a thousand small bets slowly turned into a kernel that was 280× faster than the PyTorch baseline.
That is the version of autoresearch I believe in right now. Not a magic button, and not a replacement for taste. A way to learn fast, stay honest, and make failed experiments compound instead of disappear.