How an Experiment Loop Made a GPU Kernel 232x Faster

A coding agent can write a fast-looking GPU kernel in minutes. The difficult part is proving that the kernel is correct across awkward inputs, measuring it without noise, and deciding what to try after the easy optimizations are gone.

That distinction explains a striking result from a recent GPU Mode contest. Sankalp used Codex in an automated research loop to optimize batched QR decomposition on NVIDIA B200 GPUs. Across two weeks and more than 1,500 submissions, the measured runtime fell from a roughly 419,000-microsecond baseline to 1,805 microseconds: about 232 times faster. The entry finished 12th among 183 participants.

The multiplier is attention-grabbing, but it is not the most reusable lesson. The real achievement was turning an open-ended performance problem into a controlled experimental system. The agent could edit code, run correctness checks, benchmark individual shapes, inspect profiles, record evidence, and either keep or reject a candidate. Human expertise then improved the questions, the search space, and the evaluator.

This is a useful blueprint for any work where a machine can run many experiments but cannot be trusted to judge success by appearance.

The task: compact Householder QR

QR decomposition factors a matrix A into:

A = Q R

Q is orthogonal: its columns are perpendicular unit vectors. R is upper triangular: everything below its diagonal is zero. QR is used in least-squares solvers, eigenvalue methods, and many numerical pipelines.

The contest did not ask for dense Q and R matrices. It required the compact format produced by torch.geqrf. The upper triangle of one output matrix stores R; the cells below the diagonal store Householder reflector vectors. A separate tau vector stores one scalar for each reflector. torch.linalg.householder_product can reconstruct Q from that compact representation.

For every submission, the checker rebuilt the factors and tested three relationships:

A ≈ Q R
Qᵀ Q ≈ I
Qᵀ A ≈ R

Only correct submissions were timed. The score used a geometric mean across matrix sizes, batch sizes, and input distributions. That matters: an optimization that wins on one friendly shape can lose overall, and lower precision that passes ordinary random inputs can fail badly on ill-conditioned or rank-deficient matrices.

Why ordinary Householder QR underuses a GPU

A Householder reflector zeroes the entries below one diagonal element. If v is the reflector vector and tau is its scale, applying it has the form:

H x = x - tau · v · (vᵀ x)

The algorithm advances one column at a time. Reflector j + 1 depends on the matrix left by reflector j, so the panel factorization contains a real serial dependency. A direct implementation repeatedly performs matrix-vector-style work. GPUs can execute that work, but their fastest units are designed for large matrix-matrix operations.

The classic solution is blocked Householder QR. Instead of applying every reflector immediately to the entire trailing matrix, it factors a narrow panel of b columns. Those reflectors are collected into a compact WY representation:

H₁ H₂ … Hᵦ = I - V T Vᵀ

The large trailing block can then be updated with matrix multiplications:

W = Vᵀ A_trail
Z = Tᵀ W
A_trail = A_trail - V Z

That reorganization does not remove the serial work. It confines it to a narrow panel and converts most of the remaining work into GEMMs, which can use tensor cores efficiently. LAPACK uses blocked reflector forms for the same reason: grouping reflectors turns much of a Level-2, matrix-vector workload into Level-3, matrix-matrix work.

Keep the unavoidable dependency chain narrow; send the large trailing update to the hardware’s matrix engines.

This architectural change supplied the largest class of gains, but it also opened a new optimization surface: panel width, precision, data layout, kernel fusion, dispatch overhead, small-matrix packing, and shape-specific paths.

The evaluator was the product

The kernel was only one component. The research harness made repeated improvement possible.

Each loop began with the current best candidate and a written hypothesis. The agent made a bounded code change, ran the cheapest useful correctness test, submitted valid candidates to the official evaluator, collected per-shape timings, and compared the result with the incumbent. A regression was reverted; an improvement was promoted and documented.

The agent proposes changes, but the verifier decides what survives.

The distinction between an attempted run and evidence is crucial. A timeout does not prove that a kernel is slow or wrong. A noisy one-off measurement does not prove an improvement. A successful compile says nothing about numerical accuracy. The loop should promote a candidate only after a completed correctness check and a comparable benchmark.

This makes the evaluator more important than the prose prompt. A detailed prompt can suggest good behavior; an executable gate determines what the system actually rewards.

Shape-by-shape evidence prevents fake wins

The contest covered sizes from very small matrices to 4096 × 4096, with different batch counts and conditioning. Those cases stress different parts of the machine.

For small matrices, launch overhead and occupancy can dominate. Packing many independent matrices into one launch may help more than improving arithmetic. At large sizes, there may be too few matrices in the batch to fill the device, while memory movement and tensor-core utilization become decisive. Mixed precision may accelerate the common path but violate tolerances on difficult distributions.

A single aggregate number hides these mechanisms. The harness therefore needs a result table that preserves:

• correctness by input family;
• latency by shape and batch size;
• compile and launch overhead;
• profiler observations;
• the exact candidate and parent revision;
• whether the result was complete, timed out, or invalid.

The aggregate score is still useful for ranking. The detailed table explains why the score moved and which bottleneck should be attacked next.

Profiling turns “make it faster” into a hypothesis

Without a profile, an agent often tweaks tile sizes, warps, and configuration values because those are easy edits. That can produce early gains, but it eventually becomes blind hill climbing.

Profiles expose the limiting resource. Is the kernel spending time in PyTorch dispatch between Triton kernels? Are tensor cores idle while vector lanes perform panel work? Is conversion between FP32 and FP16 repeated? Are memory loads poorly coalesced? Does one shape suffer from low occupancy? Each answer suggests a different experiment.

Sankalp used both the remote evaluator and NVIDIA profiling. The strongest algorithmic direction was to make more of the computation matrix-shaped. Later improvements came from details such as reducing framework transitions, retaining useful intermediate representations, specializing awkward shapes, and studying instructions available on Blackwell.

NVIDIA’s current tcgen05 programming guide shows why hardware knowledge matters. Blackwell introduces fifth-generation tensor-core matrix instructions, tensor memory for accumulators, and CTA cooperation. An agent can search and implement these mechanisms, but a useful profile and a person who understands the missing capability can focus that search dramatically.

One incumbent creates a local maximum

The simplest optimization loop keeps one best candidate. Every new idea must beat it immediately or disappear. That policy is efficient for small parameter changes, but hostile to structural changes.

Suppose a new data layout is initially slower because the surrounding kernels are still designed for the old layout. The layout may enable a later fusion or remove a conversion, but it cannot survive long enough to receive that second change. The loop mistakes “unfinished” for “bad.”

A small beam of candidates fixes this. Keep several idea families alive:

• an exploit candidate near the current best;
• a near-miss that wins important shapes;
• a structural candidate with greater upside;
• optionally, a simplification candidate that removes overhead or dead machinery.

Each beam entry should record its parent, hypothesis, affected functions, best evidence, and next decision. Candidates should not live forever; retire them when repeated tests reject the mechanism or the profile shows that the targeted cost is no longer important. The point is to give promising structures enough runway to mature.

Idea diversity is not the same as generating many random variants. Five agents changing the same block size are one idea with five parameter values. Useful diversity means different explanations of the bottleneck.

Domain knowledge improves the loop in three places

The result does not show that expertise is obsolete. It shows where expertise has the highest leverage.

First, experts design the verifier. Numerical code needs tolerances, adversarial distributions, and independent invariants. A weak checker rewards shortcuts that do not generalize.

Second, experts choose representations. Recognizing that blocked Householder QR converts trailing updates into GEMMs is a higher-value move than tuning an unblocked kernel indefinitely.

Third, experts notice missing experiments. The postmortem found several: distribution-aware fast paths, more aggressive removal of library calls, keeping the trailing matrix in lower precision rather than repeatedly converting it, earlier use of a candidate beam, and deeper use of Blackwell-specific tensor instructions.

An agent expands implementation throughput. Expertise improves the map it explores.

A practical autoresearch workspace

A durable setup needs less orchestration than many teams expect. Start with a small set of files whose responsibilities are obvious:

optimization/
├── AGENTS.md # standing experiment rules
├── problem.md # interface, constraints, evaluator
├── candidate.py # current promoted implementation
├── experiments.md # hypotheses and decisions
├── beams.md # active idea families
├── results/ # raw correctness and timing output
├── profiles/ # profiler captures and notes
└── archive/ # rejected or superseded variants

The standing rules should be operational. Require a sanity check before expensive evaluation. Save raw output. Treat timeouts as inconclusive. Promote only from completed evidence. Record why a candidate was rejected so a fresh context does not repeat it. Profile after meaningful architecture changes or unexplained gains.

Then make the first loop deliberately narrow:

1. Define one immutable correctness gate.
2. Choose representative fast and full benchmark suites.
3. Establish a reproducible baseline.
4. Run one hypothesis per experiment.
5. Keep an append-only result ledger.
6. Add a beam only when structural ideas begin losing to the incumbent too early.
7. Review failures and evaluator blind spots regularly.

Automation should arrive after the experiment is trustworthy. A fast loop around an unstable benchmark merely produces wrong conclusions faster.

Where this approach works—and where it does not

GPU kernels are unusually suitable because they offer a hard correctness checker, a numerical score, short experiments, and a reversible artifact. The same shape appears in compiler optimization, database query planning, compression, model training, scheduling, and code-size reduction.

The method is weaker when evaluation is subjective, delayed, easy to game, or poorly correlated with the real goal. Product design, maintainability, and security cannot usually be reduced to one latency number. They can still use agent-driven experiments, but promotion needs multiple gates and human judgment.

There is also a compute budget. More than 1,500 submissions over 14 days consumed real accelerator time and shared queue capacity. A responsible loop separates cheap local checks from expensive remote evaluation, spaces submissions, caches comparable results, and stops exploring dead families.

The lasting lesson

The 232x speedup did not come from one magical prompt. It came from a chain of improvements: a formal checker, per-shape measurements, persistent logs, profiles, better numerical understanding, blocked QR, mixed-precision experiments, shape specialization, and a search policy that eventually preserved multiple ideas.

That is what productive autonomy looks like. Give the agent a bounded artifact it can change, a reliable way to observe reality, memory of previous attempts, and permission to reject its own work. Keep humans at the points where taste, representation, and evaluator design matter most.

The code generator is useful. The evidence loop is the system.

Sources: Sankalp’s experiment report, Hacker News discussion, GPU Mode QR leaderboard, Mike Lazo’s QR optimization write-up, PyTorch geqrf documentation, LAPACK QR factorization guide, and NVIDIA tcgen05 guide.