Winning the GPU MODE eigh competition

Batched symmetric eigendecomposition on a B200, from 67 ms to 6.8 ms in fifteen days, with a fleet of coding agents sharing one GPU.

7 September 2026

I won GPU MODE’s eigh competition. The task is batched real symmetric eigendecomposition on an NVIDIA B200: take a batch × n × n FP32 tensor, return eigenvectors and ascending eigenvalues in the torch.linalg.eigh convention, and be faster than everyone else on the geometric mean of 13 benchmark cases. The baseline, torch.linalg.eigh, runs the suite at a 67.2 ms geomean. My final submission runs it at 6.8 ms on the competition’s grader.

I did not write most of the code. Claude and Codex agents did, working in git worktrees against one shared GPU, with me steering. This post covers both halves: what the kernel ended up doing, and what it took to keep a fleet of agents productive on one B200 for two weeks. The worklog and the results ledger from the campaign are the source for every number here.

The problem

The input is symmetric up to FP32 roundoff. The checker never compares against a reference solver, because eigenvectors are not unique. Instead it checks three matrix identities in FP64, each relative to the matrix’s L1 norm and scaled by n·eps:

Plus ascending order and no NaNs. The gates are deliberately loose so that FP16, TF32, and FP8 are allowed internally. That one design decision shaped the whole solution. Almost every late win came from moving a boundary between two kernels from FP32 to FP16 and paying for it with one cheap repair step at the end.

The 13 ranked cases:

Nine of the thirteen cases are n512 at batch 640 or n1024 at batch 60, so those two shapes looked like most of the score. But the geomean is a product, so a 30 percent win on n32 is worth exactly as much as a 30 percent win on n2048. That arithmetic decided where the last week went.

What torch does, and why it is slow

An nsys profile of the baseline settled the first argument. On torch 2.12 with CUDA 13, torch.linalg.eigh at n512 is not batched Jacobi. It is cuSOLVER’s classic three-stage solver: Householder tridiagonalization (sytrd), a divide-and-conquer tridiagonal eigensolver, and a Householder back-transform. The tridiagonalization is BLAS-2. Each column needs a symmetric matrix-vector product with the trailing matrix, and the next column depends on it. That is n serial symv calls per matrix, and cuSOLVER’s version runs at about 15 percent of the memory roofline for that operation.

That chain was the wall for the whole competition. Every profile we took, on every day, said the same thing: the geomean is a tridiagonalization problem. The reflector chain cannot be parallelized across columns, tensor cores do not help a matrix-vector product, and the only levers are how fast one column goes and how much other work can overlap it.

The final pipeline

The shipped submission is a single 13k-line Python file with the CUDA sources inlined as strings, generated from csrc/kernels.cu, csrc/bindings.cpp, and a Python template. It has about 45 kernels and 51 custom ops. Every shape gets its own route. For the three big shapes the route is the same algorithm as cuSOLVER, rebuilt piece by piece:

1. Prescale. One fused pass computes each matrix’s max entry, divides by it, and writes an FP16 shadow copy of A next to the FP32 one.
2. One-stage blocked Householder tridiagonalization. A latrd kernel reduces a 32-column panel (16 columns at n2048) with the trailing matrix streamed into shared memory as FP16 tiles by 2D TMA through a ring of mbarrier-guarded buffers. The symv inner loop uses Blackwell’s fma.rn.f32.f16, which multiplies two FP16 values into an FP32 accumulator with an exact product in one instruction. At n512 one CTA owns one matrix, which fills the machine at batch 640. At n1024 and n2048 the panel is spread over a cluster of 2 or 8 CTAs, with the rows of the symv split across the cluster and a DSMEM st.async plus mbarrier handoff once per column. After each panel a fused mma.sync kernel applies the rank-2k update to the trailing matrix in 64×64 tiles and refreshes the FP16 shadow. The last 128 to 288 columns skip the blocked machinery and run an unblocked lane-per-row reduction from shared memory.
3. Tridiagonal divide and conquer. Leaves of 32 rows (22 on the small shapes) are solved by parallel Sturm bisection plus inverse iteration, one warp per matrix, with a residual check and a QL re-solve for the columns that fail it. Each merge level is two kernels: a prep kernel that sorts, deflates, and compacts, replacing a chain of about 130 torch launches, and a merge kernel that solves the secular equation with an FP32 packed-f32x2 seed and an FP64 polish and builds the eigenvector update. Merges too small to fill the machine run on 2- to 8-CTA clusters. Children pass between levels as FP16. The final level writes its eigenvalues sorted and its eigenvectors in sorted column order, so there is no separate sort and gather.
4. Back-transform. The Householder reflectors are aggregated into compact-WY blocks, 4 to 16 panels per block, and applied to the eigenvectors as FP16 tensor-core GEMMs with FP32 accumulation, carrying Z in FP16 between applies.
5. One Newton-Schulz step. Q ← Z (1.5 I − 0.5 ZᵀZ), also in FP16 with FP32 accumulation. This is the repair that makes all the FP16 above legal. Eigenvectors rounded to FP16 fail only the orthogonality gate, and one Newton-Schulz pass restores a 14× margin on every input family we could generate. Eigenvalues never go through FP16.
6. CUDA graph replay. Everything above is captured into a CUDA graph at import time, one per (n, batch) and per slot of a small ring, and replayed on each call after a fused copy-and-prescale of the input into the captured buffer. At n512 the batch is split into three chunks whose graphs are joined into one DAG so one chunk’s D&C overlaps another chunk’s reduction. At n2048 the DAG has seven nodes, and the left half of the D&C tree runs while the right half of the matrix is still being reduced. Graph replay was worth 16 percent at n512, 5 percent at n1024, 17 percent at n2048, and 26 percent at n176.

The small shapes are the same algorithm in a different body:

• n32. One kernel, one 128-thread CTA per matrix. Householder tridiagonalization with a block-parallel rank-2 update, Sturm multisection for the eigenvalues with four probe lanes per eigenvalue, inverse iteration for the eigenvectors, a Gram-Schmidt pass that runs only if an exact orthogonality residual says it has to, and the back-transform. Eager launch, because wrapping it in a graph costs 4.6 percent more than the kernel.
• n176 and n352. A 3-CTA cluster reduction, a D&C sized to the exact n, an FP16 back-transform, and a verify-then-repair wrapper: run the FP16 fast route as one replayed graph, compute the eigen-residual per matrix on device against a threshold tighter than the gate, and re-solve the few offenders on an FP32 route. On the scored families across 32 seeds, zero matrices needed repair. On adversarial test families, 2 to 7 out of 40 did.
• n512 clustered. This case has a two-cluster spectrum: 170 eigenvalues tightly packed around one value and 342 around another. A prefilter samples eight rows of each matrix, a moment classifier fused into the prescale pass confirms the structure from tr(A) and tr(A²), and the batch is dispatched to a different solver. It forms the projector onto the minority cluster from the two centers, sketches its range with a sparse signed embedding, checks the projector is idempotent on eight columns, and runs a blocked Householder QR on the 512×192 sketch for 160 columns with a pivoted-Cholesky reveal for the last ten. The reflectors that triangularize the minority panel also span its orthogonal complement, so the full Q = I − V T Vᵀ is formed in FP16 from one Gram, followed by one Newton-Schulz. The eigenvalues are the two centers, repeated. If any certificate fails, the batch falls through to the general route untouched. This route ran the case at 4.95 ms; the general route ran it at 19.4 ms.

Per-case timings from the final merged run on the competition grader, unlocked clocks:

The torch column was measured on my box with locked clocks and the final column on the grader with unlocked clocks, so the ratio flatters the final by about the clock difference. The last build measured on my own box was 8.42 ms locked, an 8.0× speedup in identical hardware state, and 7.95 ms at the grader’s clock.

Timeline

Fifteen days, 1,474 commits in the task directory, 836 ledger entries: 207 wins, 345 kills, 201 recorded facts, 46 neutrals. The days below are the ones where the frontier moved.

July 2 and 3: escape the library

The first day produced the dev loop (CUDA sources spliced into a single-file submission by a script) and the first custom kernels: a shared-memory Jacobi for n32 at 2.4× over torch and a 3-CTA cluster Jacobi for n176 at 1.9×. Day two added a block-Jacobi cluster kernel for n352 at 2.6×. Jacobi was the obvious first move for a saturated batch, and for n512 it also went in at 1.1×.

The important result of day two was a negative one. Three agents built the alternatives for the big shapes in parallel: spectral divide and conquer via the matrix sign function, a two-stage reduction (dense to band, band to tridiagonal), and a one-stage blocked Householder reduction. One-stage won at n512, 85.5 ms against 190 ms for two-stage, and spectral D&C passed every gate but ran at 213 ms against torch’s 122 ms at n1024. A custom tridiagonal D&C merge kernel went from a 1083 ms prototype to 155 ms to 52 ms in one day. The geomean stood at 53.9 ms.

July 4 and 5: the composed pipeline

Day three shipped the composed route (custom latrd reduction, custom D&C, WY back-transform) for n512, 180 ms to 138 ms, and for n1024, 118 ms to 91 ms. Jacobi was dead for the big shapes from that day. A row-unrolled latrd with an FP16 shadow copy of A took the reduction from 54 ms to 24 ms at n512. Glue fusion, a fused D&C prep kernel, and FP64-native secular solves each took a few more milliseconds. Geomean 24.9 ms.

Day four found the FP16 back-transform plus one Newton-Schulz step, which improved orthogonality margins while cutting time on all three big shapes, and gate-aware z-deflation in D&C, which cut the merge kernel by 25 percent. A 2D-TMA tile symv in the reduction took n2048 from 80.6 ms to 58.1 ms and n1024 from 36.7 to 30.1. The same day the grader’s toolchain turned out to be CUDA 12.9 rather than our 13.0, and a cluster size that was optimal locally was 30 percent slower under the grader’s compiler. Geomean 17.6 ms at the end of the day.

July 6: the diet

The best single day. Nsight showed the reduction kernel was issue-bound: its duration tracked dynamic instruction count almost one to one. So the day’s theme was deleting instructions without changing the math.

• fma.rn.f32.f16 in the symv inner loop replaced load, convert, FMA with one instruction. 16.2 to 15.0 ms.
• A packed f32x2 secular seed and branchless reciprocals in the merge kernel: 24 percent fewer instructions, 24.3 percent less time. 14.5 to 13.9 ms.
• The reduction kernel’s “wait” stall looked like physics until a per-SASS-line breakdown showed the FMA dependency chain was 1 to 2 percent of stall samples. The rest was glue: cluster barriers protecting a cluster of one CTA, a block barrier per tile, an integer division in address math. Deleting the barriers on the single-CTA route was worth 9.5 percent on every n512 case with bit-identical output. Replacing three cluster.sync per column with a DSMEM mbarrier handoff on the cluster routes gave 7 percent at n1024 and 11 percent at n2048.
• CUDA graph replay landed, after a root cause: every custom launch had been going to the legacy default queue, so capture recorded nothing, and one library call took a literal zero as its queue argument and produced all-zero outputs. A ring of graphs per route, ring depth equal to the harness’s own buffer count, made replay alias-safe.

17.6 to 12.2 ms in one day.

July 7: graphs on the grader, and a compiler regression

The grader kills any CUDA graph capture that begins inside a timed call. Probes showed capture at module import escapes the check and replay inside timed calls passes, so precapture at import became the standard, worth about 1 ms of official geomean. Splitting the n512 batch into two independent graph branches gave 7 percent on all five n512 cases. n32 got a direct solver that replaced Jacobi: 0.132 to 0.088 ms, a 2.7 percent geomean move from the smallest case.

The grader also moved to CUDA 13.3 that day. Our first 13.3 build regressed 26 percent, from 11.0 to 13.9 ms. The root cause was a stale #if __CUDACC_VER_MAJOR__ < 13 around a dispatch table. Fixed, 11.3. End of day 10.8 ms.

July 8 and 9: the small shapes and the tails

With the big shapes’ reduction near its wall, the geomean arithmetic pointed at the small cases. Verify-then-repair let n176 and n352 run an FP16 fast route with a per-matrix residual check and an FP32 re-solve for the failures: n176 1.81 to 1.51 ms, n352 3.41 to 3.05. Collapsing the last 128 columns of the reduction into a lane-per-row shared-memory tail took another 20 percent off both. A multisection eigenvalue solver in the n32 kernel: 0.085 to 0.067 ms. Together the smalls took the geomean from 10.8 to 9.5.

July 9 was the day with the most kills on the ledger, 55, and also nine shipped wins: an mma.sync fused trailing update, a 6-CTA-per-SM configuration for it, a three-chunk DAG at n512, pre-issued TMA for the next column, and a set of D&C prep fusions that were bit-identical and 0.5 to 0.7 percent each. 8.75 ms.

July 10 and 11: calibration and the last local wins

On July 10 I checked the live board. The leader was at 7.2 ms official. Our own official score was 8.16 and our local locked number 8.75, so locked-clock numbers were about 7.5 percent pessimistic. That killed a set of “we are near the floor” conclusions that had been calibrated against a stale 8.4 figure and reopened every shape. I also lowered the held-out accuracy bar from 2× to 1.2× margin. The last local wins were small and many: a zero-fill removal in the reduction, an FP16 D&C carrier, fused verify norms on the smalls. The last local measurement was 8.42 ms locked.

July 12 to 14: remote only

After July 11 there was no local GPU. Every measurement from then on was a submission to the competition’s remote B200 through popcorn-cli: a cold nvcc build plus the benchmark, seven to nine minutes, on a shared and sometimes overloaded service with unlocked clocks. The same code measured 7.81 ms there. The whole workflow changed, and that is its own section below.

The wins in this phase came from the clustered n512 case, which had been running the general route at 19.4 ms and had the largest slack of any case. A dedicated route had been developed on a side branch for days and written off as a wash at 18.5 ms, and my own notes from July 13 said a geomean under 7.5 was not reachable. A Codex session restructured it that day, 19.4 to 15.0 ms, and then took it through about two dozen shipped iterations on July 14: shared-row padding, a sparse sketch shrunk from 16 nonzeros per column to 2, a packed-upper larft worth 10 percent on its own, a half-carried Gram, live-row compact-WY updates. 4.95 ms at the end of the day. The same day shipped a class of direct-FP16 wins across all big shapes: when a consumer is going to round a boundary to FP16 anyway, have the producer write FP16 directly. Each was worth 0.3 to 2.7 percent per shape. Remote geomean 7.81 to 6.78 ms in three days, 72 ledger wins and 274 commits on July 14 alone.

July 15 and 16: production port

The last two days were a Codex-driven rewrite of the shipped kernels into a standalone package without the graph machinery, benchmark-case routes, or static output buffers: a solver you can import and call on any batch size and any stream. It measures 8.0 ms warmed on a B200, and the code went from a 13k-line generated file to a 3.1k-line module. Graph replay was tried in that package too and measured 0.35 percent slower than eager, so the graph-free version is the production default.

What did not work

The ledger has 345 kills. The decision-grade ones, each with the mechanism that killed it:

• Two-stage reduction. Dense to band is BLAS-3 and fast. The band-to-tridiagonal bulge chase has a serial depth that does not depend on bandwidth: about 1850 to 2000 rounds whether the band is 4 or 32 wide. Built four separate times, including the ELPA-style diamond schedule. Always lost.
• Tensor cores for the symv. A standalone probe put the tcgen05 MMA chain at 5.5k cycles per column against 24.4k scalar. Wired into the production kernel it was 3.7 percent slower at n512 and 15 percent slower at n1024. A matrix-vector product has no N dimension, so the MMA ran at 1.78 percent pipe utilization, and the scalar path was already hiding the same DRAM traffic behind sixteen warps of memory parallelism while the MMA design funnelled everything through one issuing warp. Three agent generations, one clean answer.
• FP8 anywhere. The panel’s shared-memory footprint caps the reduction at one CTA per SM, and halving it would allow two. But FP8’s error of about 0.016 times the norm exceeds the eigenvalue gaps at n512, so Newton refinement of eigenvectors diverges on close pairs, and one refinement step costs more than the whole prize. Killed in four forms.
• Jacobi for the big shapes. One-sided block Jacobi at tensor-core rates needs FP32-grade rotation applies on mixed spectra, and B200 has no FP32 tensor core. The kill is hardware-shaped.
• Spectral divide and conquer. Passed the gates at n2048 with a Polar-Express sign iteration and a 21 to 28 ms floor against a 29 ms bar. Died on the orthonormalization: revealing the rank of an oversampled fuzzy basis at batch 8 costs more than the tridiagonalization it replaces.
• MRRR and twisted factorizations. Numerically fine, and the tree fixes cluster rank collapse. Economically dead: the root values alone cost as much as the whole D&C tail.
• Triton and TLX rewrites of the smalls. The fusion thesis was refuted because the smalls already replay as one graph.
• Adaptive secular tolerances at n512. The entire 1 percent win came from ill-conditioned roots exiting early, which is exactly the set where accuracy is decided. Per-root speed and accuracy were the same physical quantity.

The pattern across the list: most of these were killed on paper first, then re-litigated when the premise changed, and a few came back. The n32 direct solver had been killed on July 6 with the argument that its Gram-Schmidt was unconditional. It was not, and the same solver shipped on July 7 at 33 percent faster than the Jacobi it replaced. The clustered route was written off twice before it became the single largest win of the campaign.

Running agents against one GPU

One human, one orchestrator session, and subagents in their own git worktrees off main. The subagent templates were pinned to Claude Opus, with a few Fable agents on the hardest missions, and Codex sessions joined on July 13 and did the clustered route and the production port. The notes record 114 distinct Claude agent ids and 28 Codex ones. Five agent roles:

• kernel-agent owns a set of shapes and is free to rewrite anything on its path.
• architecture-explorer is told to escape the local minimum: megakernels, warp-specialized pipelines, different algorithms.
• precision-optimizer widens gate margins at zero performance cost.
• simplifier makes the submission shorter and more readable at zero performance cost.
• research-scout is read-only, never touches the GPU, and returns a ranked, sourced list of ideas. 37 scout documents were written over the campaign.

The tooling that made this work was small and mostly boring:

• gpu_run.sh. One flock on one global lock file. Every GPU command from every worktree runs through it, one job at a time, with VRAM freed between jobs and every transition logged. Without it, parallel agents corrupt each other’s timings and OOM the box. It sources the grader’s exact CUDA toolchain, exports a per-worktree scratch directory for profiler outputs, and re-locks the clocks to 1800 MHz after any job, because Nsight drops the lock and the next benchmark would otherwise measure at 1095 MHz. Two incidents shaped it: a nested lock froze the queue for six minutes until a re-entrancy guard went in, and an agent wrapped it in an outer timeout that killed jobs still waiting in the queue.
• ab_bench.sh. Locked clocks, alternating A/B against a baseline git ref, three rounds, and a win requires B faster than A in every round. Locked clocks made timings portable between the B300 dev box and the B200 grader. Later the small shapes got a stricter rule: they must be measured in the grader’s cold-L2 regime, because one n176 change read 4.6 percent faster with a hot L2 and 4.8 percent slower with a cold one.
• ship_bar.sh. The full shippability check as one command: all 39 correctness tests, the benchmark geomean measured with the grader’s own timing function, a benchmark-mode check that runs the solver on reused inputs and rechecks every iteration, an assertion that graph replay is actually on, and a content scan. The word “stream” is banned anywhere in the submission text, and an uninitialized buffer that only showed up in the grader’s benchmark mode is the reason the reused-input check exists.
• integrate.sh. Merges an agent branch into main only if ship_bar.sh passes on the merged tree, then blesses the ratchet and regenerates the priors. Agent-claimed numbers never landed without a harness measurement.
• Per-case ratchet. After a per-call environment scan costing 21 µs passed a green bar, because it only showed up as 16 percent on n32 and the aggregate moved by less than noise, the bar got teeth: any case more than 2 percent slower than its blessed best (4 percent for sub-millisecond cases), or the geomean more than 0.5 percent slower, fails integration.
• ledger.jsonl and make_brief.py. Every measured candidate, win or kill or neutral, goes into a structured ledger with the shapes it applies to, the delta, the mechanism, and the evidence pointer. The brief compiler emits the base contract plus the ledger’s prior results for the agent’s shapes plus the mission text. Agents launch with compiled briefs, not hand-written boilerplate, so dead ends die once.
• PRIORS.md with a premise register. Each kill is a measurement under premises: the compiler version, the kernel generation, the hardware, the accuracy budget. A small premise file records when each of those changed, and the priors document is regenerated after every integration with the premise table at the top. On July 9 I wrote into the map that the priors record is extensive but not hard truth, and that a mission into closed territory starts with a premise check, not deference. Most of July 6 was re-litigation of earlier kills whose premises had expired.
• agent_init.sh. A CLI bug kept cutting agent worktrees from a stale base ref, which cost about ten interventions on one day and produced a phantom regression that survived two A/B rounds. The init script merges local main, verifies sentinels, and prints the current frontier before the agent does anything.
• A 20-minute health check. Completion notifications only fire at the end of an agent’s run, so a stalled agent burns idle GPU hours silently. A session cron looked for new queue-log entries and new commits on each agent’s branch, and flagged agents with neither, a queue job running longer than ten minutes, or clocks off 1800 with no live profiler.
• A dashboard. One HTML page regenerated from the performance log after every run: the frontier, and each case’s first-versus-newest bar. Most of my useful interventions started as a bar that had not moved in days.

Agent slots were scored by benchmark weight times current milliseconds times estimated slack from a named floor, until the geomean arithmetic on July 10 replaced that with “percent per case, weighted equally.” A mechanism that had just won on one shape was transferred to its siblings before any new direction was opened. Long-horizon ideas that were not ready to ship lived on bet/* branches with a resumable status file, so an agent could pick one up cold. The clustered route was one of those bets.

When the GPU went away

From July 12 the only B200 was the competition’s own, reachable by popcorn-cli. The constraints inverted. Compute and tokens were abundant. Measurements were serialized, took seven to nine minutes each including the cold build, failed intermittently, and returned nothing useful on a build error through the CLI. The notes cite 574 distinct submission ids over three days.

The workflow became a shared runner script that submits once, captures the submission id early, kills the CLI’s fragile live poll, and polls the results endpoint directly with a five-minute request timeout, because the endpoint took over two minutes to answer and a short timeout had made every retrieval look like a failure. Submission id capture is serialized across agents with a directory lock so ids from concurrent agents cannot cross. A usage log records who submitted what and when.

Three rules replaced the local A/B:

1. De-risk offline first. Lock the algorithm in numpy, and where possible prove a kernel change with a host-side schedule proof before spending a submission. The n352 panel change on July 14 came with a script that exhaustively matched the old and new owner schedules for all five panels and 512 threads, and only then went to the remote.
2. Measure as a candidate-parent-candidate bracket. Unlocked clocks drift, so a single candidate run is not evidence. Every keep in the notes from this phase quotes three runs and divides the affected cases’ ratios by the untouched cases’ ratio to cancel fleet drift.
3. An empty result is inconclusive, not a kill. The service sometimes returned a terminal status with no timings at exactly the 600-second mark. Several ideas got a retry on that basis and a few of them shipped.

What I learned about orchestrating agents

• The scarce resource decides the workflow. With a local GPU the bottleneck was the queue, and the tooling was about serializing access and making every measurement trustworthy. Without one, the bottleneck was eight-minute measurements, and the tooling was about spending as few of them as possible. The same agents did much better work once the harness matched the constraint.
• Structured kills are worth as much as wins. The ledger held 345 of them by the end. The value was not in the count. It was that each one named its premise, so an agent could tell the difference between “this mechanism is dead” and “this mechanism lost under conditions that no longer hold.” Most of the biggest day’s wins were the second kind.
• Narrow missions beat “make it faster”. An agent asked to improve the geomean will trade a 5 percent loss on one shape for an 8 percent win on another and call it a day. An agent asked to improve n1024 mixed with zero regression anywhere else separates the code paths until its change cannot touch the other shapes. The per-case ratchet enforces the second framing mechanically.
• Agents claim wins that harnesses have to verify. Every integration ran the full bar on the merged tree. The one regression that got through did so because a check lived in the integration script but not in the ship bar, so an agent’s own green bar and the integration’s green bar were checking different things. After that, the same check ran in both.
• Aggregate metrics are not diagnoses. The “wait” stall on the reduction kernel and the “structural wall” on the spectral solver were both wrong at the level of the summary and right only after being decomposed: per SASS line for the first, per rank per level for the second. The two most expensive mistakes of the campaign were believing a summary number.
• Calibrate against the real board. Ten days of local numbers accumulated a set of floor models that were all off by the same 7.5 percent, because nobody had submitted a calibration run. One submission fixed that and reopened several shapes. My own notes then declared under 7.5 unreachable three days before the fleet measured 6.78.
• Per-case bars beat prose. The dashboards that mattered were a table of which case moved, by how much, and when. The 18k-line worklog was necessary for the agents and nearly useless for me. The ledger and the performance log were the opposite.

Epilogue

The production port is the same numerical kernels without the competition scaffolding: no graph, no static outputs, any batch size, any stream, 3.1k lines, 8.0 ms warmed on a B200 for the same 13 cases. A graph-replay variant exists as a separate package for services that can prepare shapes at startup. One adversarial arrowhead family from the LAPACK test set still fails the FP32 fallback at n352 in the production build, and that is the open bug.