Auto-research with codex: How I achieved a 212x Faster Kernel over baseline with Codex in GPU Mode's qr_v2 problem

08 Jul, 2026

Table of Contents

• Intro
  • Contest in short
  • Problem intro
• Why this problem is auto-research-able
• Learning Enough to Ask Better Questions
• The math, in brief
• Make serial work small with the help of the blocked Householder algorithm
• Codex-maxxing
• Kernel progress breakthroughs
• Breakthrough ideas
• Introducing idea diversity to escape the local maxima
• Implementation Hints
• What I could have done better
• Conclusion
• References
• Acknowledgements

Intro

Contest in short

GPU Mode, in collab with Core Automation, recently hosted an auto-research themed contest. The problem statement was to implement batched square compact-Householder QR factorization aka QR decomposition. I placed 12th out of 183 participants ending up with a 212x speedup over the baseline solution. This post is about how I got there. I will go through my approach, learnings, and bottlenecks I ran into during the contest. It was my first serious attempt at auto-research. Some people will call this "loop engineering", and honestly that is fine too.

Note that you don't need to go through the mathematics or the problem itself in detail to follow most of this blog post. I have focused on my approach while keeping the math and the problem itself secondary as most people who will read this won't have participated in the contest.

You can check out the full contest page here: Problem Link and Leaderboard

This contest was part of GPU Mode's Linear Algebra Kernels in the Age of Research series.

Problem intro

We were given a batch of square FP32 CUDA matrices A with shape batch x n x n, and had to return the same compact Householder QR representation as torch.geqrf(A): an H matrix whose upper triangle is R and whose lower triangle stores Householder vectors, plus a tau vector of reflector coefficients. The checker rebuilt Q with torch.linalg.householder_product(H, tau), took R = triu(H), and verified:

$A\approx QR,Q^{\top }Q\approx I,Q^{\top }A\approx R$

Among correct submissions, the leaderboard ranked runtime by geometric mean across shapes and conditioning cases. The important sizes were batched square matrices like 512 x 512, with larger 1024, 2048, and 4096 cases too. Low-bit FP16, FP8, or NVFP4 was allowed internally, but returned factors still had to satisfy FP32-style QR checks.

A tiny 3 x 3 example is:

Here Q is orthogonal, which means its columns are unit-length and perpendicular to each other, and R is upper triangular, which means everything below the diagonal is zero. The contest was not asking us to print dense Q and R directly; it asked for the compact Householder version that lets the checker reconstruct Q and read R from the upper triangle.

For the 3×3 example above, the very first reflector maps the first column (12, 6, −4) straight onto (−14, 0, 0) in one shot — the −14 becomes $R_{11}$, the mirror's direction is what torch.geqrf stores below the diagonal, and its scale goes into tau.

Why this problem is auto-research-able

GPU Mode provides participants with the popcorn CLI making it agent friendly. Agents can use this to test, benchmark, and submit to the leaderboard directly. The checker also provided shape-wise feedback along with the overall geometric mean timing.

Astute observers will notice this is an apt setup for writing a loop. Agents yearn for tight feedback loops. They allow them to hillclimb to their heart's content.

GPU Mode contests usually give you some way to iterate on kernels. Either you submit directly, or a sponsor like Modal chips in credits. Here the organizers basically allowed unlimited submissions as long as you spaced them out. If you didn't, the queues got long and everybody's runs timed out. At one point the workspace even ran out of Modal credits because everyone had been hammering submissions. It's a nice way to make learning accessible.

Over the course of 14 days, I made over 1500 submissions.

Learning Enough to Ask Better Questions

I have known the basics of GPU kernel optimization (mainly in Triton with some understanding of CUDA) for a year, but haven't worked in this domain professionally. What I am trying to tell you is that I was an underdog among the people around me on the leaderboard. The person just above me on the leaderboard (CUDA Colonel) is a principal engineer at NVIDIA.

Anyway aura farming aside, since I know the basics and had recently read about GatedDeltaNet, I was fresh with the general GPU kernel lingo.

The better you know something, the better you can prompt the LLMs, because you convert unknown unknowns into known unknowns.

At the same time, it's worth noting that this contest was doable without domain knowledge - like you probably won't make it to the top 10, but you can get a respectable speedup over baseline by just relying on your harness/agent loop or whatever.

My first steps in the contest were to learn what QR decomposition is and how it can be done. There are a bunch of ways to do it - like Gram-Schmidt and Householder reflections. The contest mandated Householder reflections. I went back and forth with Claude and watched a few YouTube videos to build intuition. After my discussions with Claude, it was clear that we needed to use the blocked Householder algorithm as the main architecture with the trailing WY-update. As it turns out, GPT-5.5 also had a good idea about this. QR decomposition is a fairly well known problem.

I found the concept interesting as matrix decompositions show up in several modern optimizer variants for LLM training, especially in methods that use matrix preconditioning, such as Shampoo-style optimizers and related approaches. Muon (used by Kimi) is another good example: instead of treating a weight update as one giant flattened vector, it keeps the matrix structure around and orthogonalizes the momentum update, usually through a few Newton-Schulz iterations that approximate the polar factor.

The math, in brief (optional)

Let's briefly go through the algebra behind Householder QR.

Householder builds Q as a product of reflectors, one per column:

$A=QR,Q=H_1{H}_2\cdots H_n,H_j=I-{\tau }_j{v}_j{v}_j^{\top }$

Each reflector is a rank-1 update, so applying one is just a subtraction:

$H_{j}x=x-{\tau }_j{v}_j(v_j^{\top }x)$

After all reflectors have been applied:

$R=A^{(n)}$

which is the same as stacking the reflectors into $Q$:

Make serial work small with the help of the blocked Householder algorithm

Householder QR zeroes out A below the diagonal one column at a time. Each step builds a reflector (basically a mirror) from the current column and applies it to everything on the right. The problem is reflector j+1 is built from the matrix after reflector j has already hit it. So you can't reorder the steps and you can't fuse them. It's inherently serial, and that serial matrix-vector work runs in the slow vector lanes of the SM while the tensor cores just sit there idle.

The classic fix is the blocked algorithm. You pick a narrow panel of b columns (say 32 or 64) and do all the serial work inside it. That's fine, because the panel is only b columns wide, so it stays cheap. Then, instead of applying the panel's b reflectors to the rest of the matrix one at a time, you compress them into a single rank-b update (the "WY representation") and hit the entire trailing block in one shot with three back-to-back matrix multiplies. The serial work stays confined to the panel, and everything else turns into GEMMs, which is exactly the shape the tensor cores want.

Concretely, the WY representation collapses a panel's $b$ reflectors into a single rank-$b$ update. Stack the panel's Householder vectors as columns of $V=[v_1,v_2,\dots ,v_b]$, build a small $b\times b$ upper-triangular $T$, then:

and the trailing-block update becomes three GEMM-shaped steps:

$W=V^{\top }{A}_{\text{trail}}$

$Z=T^{\top }W$

$A_{\text{trail}}\leftarrow A_{\text{trail}}-VZ$

The panel is where the serial matrix-vector work is stuck, but it is only $b$ columns wide. Everything to its right is the big trailing block, and that is pure GEMM. As the panel walks down the diagonal, the trailing block shrinks, and the finished columns split into R on and above the diagonal and the stored reflectors v below it, which is exactly the compact (H, tau) layout the checker wants.

If you want to understand / visualize the math properly, I highly recommend checking Mike's writeup. He placed 5th in the contest and shared his learnings in a problem focused way. He has provided peak visuals.

Other challenges

Two more things proved to be challenging: reliably using low precision internally (especially for ill-conditioned inputs), and coping with the wide spread of shapes (n = 32, 176, 352, 512, 1024, 2048, 4096) and batch sizes, where the largest matrices had too few batches to fill the tensor cores while n=32 was so small we had to pack many matrices into one kernel launch.

Codex-maxxing

Why I picked Codex

I participated with ChatGPT Pro (200 USD subscription) and Claude Pro (20 USD subscription). I also used Modal for profiling (they provide 30 USD worth of credits free every month btw).

I chose Codex mainly because:

1. I had the bigger subscription.
2. From prior experience, my intuition was that OpenAI models are better at Triton.
3. /compaction works very well in Codex.

Setting up the harness

After I got a basic understanding, I asked Codex to do the basic setup: add problem_statement.md, mention basic details in AGENTS.md on how to submit and use popcorn CLI, and maintain a log.md where we did bookkeeping of the submissions and their status (accept/reject along with shape-wise timings). If we have to contrast with Karpathy's auto-research, my AGENTS.md and problem_statement.md were initially my program.md.

Logs serve as the evidence of the ideas that worked and didn't work. Future agent sessions could read the logs and quickly check if an idea had been tried or not. I put more investment into logging after the 3000 µs mark as things started to get harder.

Just tell it what to do

The cool thing about Codex is you can just tell it to do stuff and it will actually do the stuff. You can make it work for hours if you give it a detailed enough prompt with targets. It knows all the math and code, and it has all the feedback it needs. I had this intuition, but I was still surprised by how much GPT-5.5 could push the performance beyond the baseline solution (which was the torch.geqrf/cuSolver function).

My initial few sessions were manual prompts to implement a solution in Triton and optimize for the n=512 and n=1024 shapes, as they had the most weightage in the final geometric mean.

Steering with /goal

However, if you want to make the model loop until a certain objective is achieved, use /goal. You can give specific, achievable, quantitative goals. I found that giving a good numeric goal followed by specific criteria worked well.

Example: "Use only Triton or CUDA and beat our active best's n=512 timings. Try several ideas either by submitting directly to the leaderboard or using Modal profiling. Remove cuSolver altogether in the new set of experiments. We will only use it as a fallback." At one point, this goal ran for over a day.

I gave inputs every 2-3 hours to move the model in the direction I wanted. I also let it run without supervision overnight on some days. In the initial few days, my instructions were mostly around steering the model to try out different ideas (more on this later) for different shapes, shouting at it to use more Triton, less PyTorch, and fewer fallbacks! A lesson I learned here: I often had the itch to check on my agents frequently, but you gotta trust it and let it do its work. Get out of the agent's way you must; steer it back only when it gets stuck.

Checking in without pausing the loop

When using /goal, you can ask the model questions without pausing the loop by using /btw or /side. This creates a temporary thread with context from your main conversation. I liked this way of checking on the agent and providing oversight.

I would ask questions like: Are you winning son? What algorithm/changes are there in the current active submission? Explain this concept to me. What are the ideas you are currently working on? What's the progress? What are some recent breakthroughs? Can we transfer it to another shape? What are your next best ideas? Based on the responses, I would consult Claude to improve my understanding of the concepts and bottlenecks involved and then provide my thoughts to Codex. After questioning, I would just go back to the main thread and dump thoughts to steer the model.

if a /goal or some loop type thing is running, then either you can queue up the instruction in codex to ask stuff or you can do /btw. (img me checking on codex to ask what's it doing, what next ideas are etc.). if you do esc to interrupt the loop, then the /goal pauses https://t.co/Yx6EpH5PKo pic.twitter.com/imqZc6RPkz

— sankalp (@dejavucoder) June 27, 2026

The baseline torch.geqrf ran at around 419 ms (419,000 µs) for this problem. I was able to reach 5000 µs within a day for the n = 512 shape (the one weighted most heavily in the geomean). This mainly came from switching to the blocked Householder algorithm described above.

When optimising kernels, making the work more matrix-shaped so the tensor cores stop being idle is your life's purpose.

Optimisations were much harder after the 3000 µs point. I had to get more involved in the loop in terms of learning the concepts and steering the model. I gave Modal profiling access to Codex and let it run torch profiling / nsys profiling to test out different ideas, compare implementations, and sweep parameters faster. (Later on, the organizers also provided a way to do NCU profiling.)

Kernel progress breakthroughs

Breakthrough ideas

The rough structural evolution of the QR kernel, from baseline/library-heavy paths toward custom panel work, fused assembly, and GEMM-shaped trailing updates:

There was a lot of back and forth on profiling the entire shape and then identifying the bottlenecks. For this problem, launch overhead and panel overhead dominated. We were rarely memory or compute bound. Most of the effort was spent finding ways to reduce panel overhead and make the WY-update more GEMM-able.

Questions I found myself repeatedly asking:

• What is panel overhead/launch overhead? How do we solve it?

• What info are we missing and how can we get it?

• What's the latest profile that you have done? What's your take on it?

• Look for fusion candidates. Are there Triton fusion-level candidates?

• Look for possible compiler-based optimizations. What can we convert from a runtime signal to a static signal? The idea was to give compiler hints. I once had a 200 µs jump because of this

• Asked it a few times to look at Triton-generated compiled artifacts

• What are some numerical tricks to exploit the precision slack?

• Can you deploy sub-agents to do some math and find possible optimizations?

• Deploy agents to search for bugs that may be bogging us down

• Are there reduction fusions possible? I saw Codex discover one and then I started hammering it often. Reductions are operations like doing a sum or finding max. Since work needs to be done to iterate through the entire sequence, we can perform multiple of these in a single go.

Introducing idea diversity to escape the local maxima

A major challenge I started facing in the 3000 -> 1800 µs range was the model getting stuck in local maxima. This looked like endless hand-tuning of parameters and small variants of the same idea.

A couple of years back, agents got stuck in (doom) loops because they were not smart enough or just didn't have the knowledge (or the verification loop was not robust enough). People used to experiment with sampling, temperature variation and different decoding strategies for this.

Then over time, models got smart enough, pretrained on newer knowledge, and we got tons of RL scaling and inference-time compute (training the model to think for a larger number of tokens).

Now the models struggle with finding new ideas and "research taste" - what next best idea/experiment should we do given the verifier feedback, prior evidence, and our evals (profiler feedback in our case). Good idea generation is the next great adventure. I recently wrote an article too on this theme.

I used the following strategies to help the model get out of local maxima:

• Beam of candidates: For a long time, I did a dumb thing. I kept a single best candidate against which new candidates were tested. If the agent tries a new structural idea or a significantly big change, then it's highly likely that it would score less than our current best submission. However, after a few iterations, that change may outperform our best candidate. With this observation, I introduced some instructions to maintain a beam of 3-5 candidates.

• Human in the loop: I would act as the secret sauce to steer the model when it was stuck for long times without improvement

• Encourage the model to take more risks and try ambitious ideas. You will not believe it but this worked. Also, Claude would often give up after a few rounds, with excuses like "we have exhausted all optimizations to reach x geomean"; Codex, on the other hand, is more persistent.

• Use a stronger advisor model that produces more varied ideas. In my harness, this would look like an instruction in AGENTS.md to encourage the model to use headless calls claude -p to get ideas, provide profiling data, etc.

• Instruct the model to frequently use sub-agents to try out ambitious ideas, search the web for blogs and papers, go through the list I mentioned above, and find micro-optimizations

• Profile using NCU and Modal, compare the results, and work on bottlenecks that both Modal and NCU profiling pointed out

• Cleanup

  • Clear context, start fresh
  • Clean up the environment, move older submission files to archives
  • Occasional iterations to simplify code, remove dead code, renaming/refactoring

• A multi-agent swarm approach is something I thought of but didn't try.

I think the strong advisor strategy (like GPT-5.6 Sol or Fable) is going to be a standard strategy in auto-research flows. Think very big models with state-of-the-art training. Claude Code provides an /advisor command for this too.

We're bringing the advisor strategy to the Claude Platform.

Pair Opus as an advisor with Sonnet or Haiku as an executor, and get near Opus-level intelligence in your agents at a fraction of the cost. pic.twitter.com/fRkegyMs5t

— Claude (@claudeai) April 9, 2026

Implementation Hints

My directory structure looked like this by the end:

qr/
├── submission.py                    # the live entry
├── submission_*.py            (560) # named submission variants (crystal_rain, blue_reply, …)
├── modal_b200_*.py            (119) # Modal B200 probe / compare scripts
│
├── AGENTS.md                        # top-level notes / logs
├── attempts_log.md
├── claude_ideas.md
├── leaps.md
├── problem_statement.md
│
├── docs/                      ( 68) # per-experiment writeups & status docs
├── code/                            # the actual QR kernel source tree
├── scripts/                         # summarizers, timing, submit helpers
├── archive/                         # cleaned-out old submissions & probes
│   ├── submissions_20260616_17/
│   ├── submissions_20260618_20/
│   ├── probe_scripts_cleanup_20260627/
│   └── …
├── submit_logs/                     # logs provided by the evaluator for each submission
└── profile.*/                       # captured NCU profile runs

Here's what my AGENTS.md looked like by the end:

Leaderboard Agent Notes

This workspace is for leaderboard-style optimization work. Follow these standing instructions unless the user explicitly overrides them.

Submission Discipline

• Don't hesitate to use sub-agents. Give them relevant instructions so they can do their task.
• Profile after major changes or major score gains.
• Use an advisor model when you are stuck or need fresh ideas. It can help break tunnel vision.
• Don't be afraid to implement hard tasks.
• Don't hesitate to take risks.
• Be open to new ideas and search the internet at times for new idea exposure.
• Before submitting, run the cheapest available sanity check for the candidate file.
• Keep submission logs under submit_logs/.
• Always save submit output into a timestamped log.
• Space leaderboard submissions out. Do not stack rapid back-to-back submissions unless explicitly asked.
• Treat a timeout as inconclusive, not as a correctness/performance rejection.
• Treat a completed residual failure or timing regression as real evidence.
• Only completed pass/fail/timing output is evidence.

Beam Search Discipline

Do not optimize as a single-incumbent hill climb. Maintain a small beam of active idea families so that local negative results do not prematurely kill useful ingredients.

• Keep a live beam note in docs/.
• Each beam entry should record the parent candidate, hypothesis, exact changed functions or gates, current best candidate/log, and next singleton, combination, or kill decision.
• Keep at least 3 active beams when there is enough work:
  • one exploit beam near the current best,
  • one near-miss beam,
  • one structural/high-risk beam from profiling or external ideas,
  • one cleanup/compile-time beam only if it has not consumed the whole search.
• Use sub-agents to work different beams, not many variants of the same tiny parameter unless explicitly asked for a parallel sweep.
• Do not declare a family dead after isolated singletons fail.
• If two ideas are individually neutral or slightly slower but touch independent costs, try combining them before retiring the family unless correctness risk is high.
• Preserve near-misses as beam material when they show a repeatable isolated win, improve one important case, remove overhead, or change an algorithmic surface that can combine with another beam.
• Kill a beam only with a clear reason: inherent correctness failure, repeated meaningful regression after a reasonable retune, singleton and plausible combinations both lose, profile evidence shows the targeted cost is no longer material, or implementation cost is blocking higher-value beams.
• After every 3-5 submissions, update the beam note with the current beam ranking and next combination candidates.

Promotion / Git Hygiene

• When a new candidate is promoted into the active file, stage the promoted candidate, active file, evidence docs/logs, and any archive moves.
• Commit after promotion unless the user says not to.
• Use a lower-case commit subject and describe both what changed and what made the improvement in the commit body.

Profiling / Evidence Habits

• Prefer current active evidence over older sidecar timings.
• Keep raw profile logs and JSON under submit_logs/.
• When a candidate is rejected or promoted, update the relevant doc in docs/.
• Use rg for searching when possible.
• Archive older candidates, profile logs, and probe scripts so the root stays navigable.

Known Tried Ideas

• Do not mirror every attempt summary here. Check docs/ before repeating an idea, and update the relevant doc when an idea is rejected or promoted.
• Do not repeat rejected ideas as-is. If revisiting one, make the algorithmic difference explicit in the doc/log.
• If a platform or evaluator rule rejects a class of ideas, record it clearly so future agents do not rediscover the same invalid path.

What I could have done better

While there is scope for improvement in my simple harness, most of my shortcomings were problem-specific. I had these realizations after scanning the top 10 submissions and reading Mike's writeup.

• The n=512 and n=1024 cases had multiple input distributions such as dense, clustered, rank-deficient, mixed, and near-rank families. The faster kernels had written data detectors and exploited distributions, e.g. low-rank cases have lots of zeros, so how do we exploit this? I could have pushed Codex more here or at least inquired more in this regard.

• Top 10 solutions more aggressively removed library functions. For example, the 2nd and 5th solutions used custom triangular inverse instead of using PyTorch triangular solve. My solution had lots of back and forth between PyTorch and Triton.

• Could have kept the trailing matrix resident in fp16 instead of repeatedly moving between representations. This was an unknown unknown for me and was purely a domain expertise miss.

• Have the beam of candidates from the start

• Write a more robust profiler to test the mixed precision cases

• Wasn't able to use tcgen05 instructions, NVIDIA's fifth-generation tensor-core instructions on Blackwell, to exploit B200 tensor cores more

Conclusion

I placed 12th out of 183, with a 212x speedup over baseline — and more importantly, I learned a lot about GPU kernel optimization, auto-research (or loop engineering?) basics, and some B200-specific details.

I also concluded that domain expertise accelerates one both in terms of harness design and human-in-the-loop improvement. I also had a couple of observations. There is a spectrum for how domain-specific engineers want to make things. On the left, people want to make a general harness that can self-improve. On the other end, people are interested in making very problem/environment-specific harnesses. I personally lean more towards problem specific preference.

I hoped you enjoy reading this blog and learnt something new. If you liked this blog, please like/upvote/share!

By the way, the second contest in the series, eigen decomposition, is currently going on. See you on the leaderboard.

References

I used the following to revise or learn new concepts.

Modal GPU glossary

How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance: a Worklog

Outperforming cuBLAS on NVIDIA B200

Outperforming cuBLAS on H100: a Worklog

Fast QR decomposition on NVIDIA B200 GPUs - Mike's writeup helped me revise the problem itself lol.

Additional references that I wish to go through:

Simon's blog - CuteDSL and B200-specific things.

gau-nernst blog - gau-nernst placed 2nd.

Acknowledgements

Mark Saroufim, Rohan Anil, for hosting the contest. Sinatras for detecting reward hacks and encouraging me initially.

Mike placed 5th and wrote an amazing writeup that decodes the problem and has great visuals.

Levidiamode for valuable posts around GPU kernels, Tokenbender for nudging me to write this, and CUDA colonel for providing tips on the GPU Mode server.

Claude Opus 4.8 and GPT-5.5 (Codex) helped with editing and math-related writing.