Sankalp uses Codex loop to place 12th in GPU Mode's B200 QR contest

Why it matters

Sankalp's result shows how a strict benchmark can turn a coding agent into an experimental system. Engineers still have to choose the architecture, interpret profiles and change the search policy when the agent stalls.

The independently verifiable result is Sankalp's 12th-place finish among 183 entrants at 1,804.779 microseconds on an NVIDIA B200. He reached that result after more than 1,500 submissions over 14 days using OpenAI Codex. The project was part of GPU Mode's Linear Algebra Kernels in the Age of Research series and its qr_v2 competition.

In a July 8 account of the project, Sankalp wrote: "Over the course of 14 days, I made over 1500 submissions." Sankalp's blog reports a 232x improvement over an approximate PyTorch baseline of 419,000 microseconds; that comparison is specific to this benchmark workload and was not independently reproduced in the supplied materials.

The competition ended on June 29, about six and a half weeks before this article. Sankalp's final result was roughly 48% slower than the winning 1,220.774-microsecond submission. His finish remains notable because he entered with about a year of experience learning GPU-kernel optimization, mostly in Triton, and said he had never worked in the field professionally.

Sankalp described an optimization loop that combined code generation, benchmarking, profiling and leaderboard submissions. He chose the architecture and interpreted the results. The workflow details below come from his account; the independently measurable outcome is his leaderboard time and rank.

A benchmark designed for agents

GPU Mode, a platform for GPU programming competitions and kernel benchmarks, and Core Automation asked competitors to implement batched, square compact-Householder QR factorization. Each submission had to accept FP32 CUDA matrices and return the same compact representation used by PyTorch's torch.geqrf: an H matrix containing the upper-triangular R result and stored Householder vectors, plus reflector coefficients in a tau vector.

The checker reconstructed Q, tested orthogonality and residual error, and ranked correct submissions by geometric-mean runtime across matrix shapes and input conditions. The workload covered matrices through 4,096 by 4,096, including difficult conditioning cases. Competitors could use lower precision internally, but their output still had to pass FP32-style checks.

That combination gave Codex a fast, quantitative definition of progress. GPU Mode's popcorn command-line interface let the agent test, benchmark and submit candidates. Shape-level timings showed where each change helped or hurt, while profiling supplied another layer of evidence.

According to Sankalp's account, he maintained an AGENTS.md file containing operating instructions, a problem statement, experiment records and timestamped submission logs. Later Codex sessions could read which approaches had failed instead of rediscovering them. He also gave the agent numerical targets and let some optimization runs continue overnight, checking in every two or three hours to ask what it had changed and which bottleneck it was pursuing.

The resulting directory eventually contained 560 named submission variants, 119 Modal B200 probe and comparison scripts, and 68 per-experiment documents. Modal, a cloud infrastructure provider, supplied GPU credits for the competition, according to Sankalp. The records preserved failed approaches, profiling results and submission history so later Codex sessions could build on earlier work.

Turning serial work into matrix multiplication

Sankalp first used Claude and educational material to understand Householder QR, according to his account of the technical work. He settled on a blocked Householder design with a trailing WY update, then used profiling, correctness checks and repeated leaderboard submissions to guide further optimization.

The core performance problem was sequential dependency. A conventional Householder implementation processes columns in order, with each reflector depending on the matrix produced by the previous step. That leaves substantial work in slower matrix-vector operations while the B200's tensor cores wait.

The blocked approach confines the serial work to a narrow panel and converts the larger trailing update into matrix multiplication, where the GPU has much more parallel work available. Sankalp reported reaching about 5,000 microseconds on the heavily weighted 512 by 512 case within his first day of using that architecture.

Further gains required changes across the stack. The submission history in his post moved through custom Triton panels, grouped WY updates, CUDA graph replay, fused layout assembly, fixed-shape kernel specialization and a custom Cholesky path for the largest matrices. Sankalp reported that the tracked full-table result dropped from 108,803 microseconds to about 1,805 microseconds. This progression was separate from the approximate 419,000-microsecond PyTorch baseline used in his 232x calculation.

Sankalp's account of the later profiles says launch overhead and panel processing dominated, while the kernel was rarely limited by raw memory bandwidth or compute capacity. His optimization table records work to reduce launches, fuse reductions, specialize fixed shapes, combine V/T assembly and remove copies, concatenations and temporary representations.

Sankalp wrote that he gave Codex access to Modal profiling and used Torch profiling, NVIDIA Nsight Systems and, later, NCU. These are self-reported methodology details. In his account, Codex implemented and measured changes while he inspected the results, set targets and redirected the search.

The human changed the search process

According to Sankalp's account, optimization became harder after the result fell below 3,000 microseconds. He described Codex spending more time on parameter tuning and small variants of ideas it had already tried. Sankalp responded by changing how experiments competed for attention.

He instructed Codex to maintain a beam of three to five candidate families instead of preserving a single incumbent and rejecting every slower attempt. The candidates included changes near the current best result, promising near-misses and higher-risk structural ideas.

In the same account, Sankalp described using headless claude -p calls as advisers, assigning sub-agents to seek optimization ideas and clearing accumulated context before fresh runs. His AGENTS.md instructions treated timeouts as inconclusive, preserved timestamped logs, required completed output as evidence and recorded why candidates were promoted or rejected. The file also warned against repeating rejected ideas without a material change.

Codex supplied persistence, code generation and a large search budget. Sankalp selected the architecture, revised the feedback loop, recognized repetitive behavior and changed the experiment policy.

In his review of the top 10 entries, Sankalp said faster competitors used input-distribution detectors, removed more library calls, implemented custom triangular inverses and handled lower-precision data more aggressively. These were Sankalp's observations rather than independently audited findings. In the conclusion of the same post, he wrote that he "wasn't able to use tcgen05 instructions" on the B200 to exploit its tensor cores further.

GPU Mode turns contests into an agent test bed

In a Core Automation essay on AI-written systems code, GPU Mode co-founder Mark Saroufim wrote that he and co-founder Andreas Kopf started a GPU programming reading group in late 2023. The group later expanded into a YouTube channel, kernel competitions and hackathons.

Saroufim used the phrase "data starvation" for the shortage of GPU-kernel material available online to train coding models. GPU Mode's contests generate measurable examples of agents and humans working on that sparse class of systems problems.

Sankalp's 12th-place finish offers a grounded view of that activity. A newcomer used an agent to approach the top of a competitive leaderboard on a narrow, measurable task. The work still required him to learn the problem, structure the experiments, challenge repetitive behavior and interpret profiling results. His result shows the value of an evaluator that can produce trustworthy evidence after every run.