AI Brokers Beat PyTorch: Writing Sooner CUDA Kernels

In February 2025, Sakana AI introduced that its “AI CUDA Engineer” generated 17,000 CUDA kernels with speedups of as much as 381× over PyTorch. This can be a system constructed to put in writing CUDA kernels, the small applications that inform a pc’s graphics chip (the GPU) precisely learn how to crunch numbers.

Inside a day, an X consumer discovered the AI hadn’t written quicker code—it had exploited a flaw in Sakana’s testing system that allowed incorrect kernels to move. Sakana retracted the claims and acknowledged a key lesson: if the benchmark is flawed, an AI will optimize for the take a look at, not the issue.

That raises the true query: How have you learnt a speedup is actual?

To search out out, I ran my very own experiment on an NVIDIA DGX Spark utilizing Claude Code. I requested it to optimize 4 widespread CUDA operations and evaluated the outcomes with 2 benchmark suites: one rigorous, and one deliberately flawed to see if the AI would take the shortcut.

The outcomes have been encouraging. The brokers produced appropriate, high-performance CUDA kernels, with the perfect implementation operating 1.57× quicker than torch.compile on a matrix multiplication workload. Three unbiased brokers reached the identical resolution.

The larger takeaway wasn’t about CUDA—it was about benchmarking. Constructing a reliable analysis proved tougher than producing the optimized code itself.

1. Who that is for

This text is for anybody contemplating AI-generated CUDA optimizations. You’ll stroll away with:

• A sensible framework for deciding when AI-driven kernel optimization is price your time.

• Actual efficiency outcomes for 4 widespread GPU operations utilizing honest benchmarks.

• 5 methods CUDA benchmarks can produce deceptive outcomes—and learn how to catch them.

• Whether or not profiler suggestions really helps AI brokers (it didn’t).

2. What a CUDA kernel is, and why “quicker than PyTorch” is a trick query

A kernel is a small program that runs immediately on the GPU, executed by 1000’s of threads in parallel — each dealing with a small subset of information.

Consider PyTorch as ordering off the menu: its kernels are extremely optimized to execute operations separately, as a result of generally it doesn’t know what your program will do subsequent. A customized kernel solely pays off when you understand one thing PyTorch doesn’t and merge these operations collectively.

That’s known as kernel fusion, and it’s the place most actual speedups come from. As an alternative of writing intermediate outcomes to reminiscence between operations, the GPU does every part in a single move. Since transferring knowledge is usually slower than the mathematics, eliminating these further reminiscence transfers often reduces latency.

2.1 What fusion appears to be like like in code

Right here’s the operation this text calls gelu_bias_residual:

out = res + torch.nn.practical.gelu(x + bias, approximate="tanh")

Initially, it comes from 3 separate operations. In PyTorch’s keen mode, every operation launches its personal GPU kernel, so the tensor is repeatedly learn from and written again to reminiscence—including pointless overhead.

A fused kernel combines all three operations right into a single GPU move. Under is the core of a fused kernel generated by an AI agent. Every thread processes one component, and the GPU-specific reminiscence directions are defined within the feedback.

__global__ void gbr_kernel(const float* __restrict__ X, const float* __restrict__ B,                           const float* __restrict__ R, float* __restrict__ Y,                           int n, lengthy numel) {  const lengthy i = (lengthy)blockIdx.x * GBR_BLOCK + threadIdx.x;  const float xv = __ldg(&X[i]);             // learn x  const float rv = __ldg(&R[i]);             // learn the residual  const float bv = __ldg(&B[(int)(i % n)]);  // bias is small, so it lives in cache  __stcs(&Y[i], gelu_bias_res(xv, bv, rv));  // one write, and this component is completed}

Solely two reminiscence reads and one write are wanted. The addition, GELU, and residual computation all occur in registers—the GPU’s quickest reminiscence—between the load and retailer. That’s the important thing to fusion.

This fused kernel runs 2.5× quicker than the three-operation model. Its efficiency can also be almost an identical to torch.compile, which robotically performs the identical fusion behind the scenes.

2.2 Choose the best baseline

A declare like “2.5× quicker than PyTorch” often compares towards PyTorch keen mode, the place each operation launches its personal GPU kernel and repeatedly reads and writes reminiscence.

However PyTorch additionally ships torch.compile to have a look at your mannequin, work out which operations might be fused, and robotically generate its personal optimized GPU code to do it — Triton kernels. These are carried out by PyTorch behind the scenes, and normally, you are already utilizing this optimization. If you happen to’re optimizing for efficiency, the true comparability is to benchmark with the run of torch.compile.

Earlier than testing any AI-generated kernels, I additionally validated my benchmark by operating the unique unfused code because the “candidate.” It measured 0.99–1.00× throughout all 4 duties, confirming the benchmark itself wasn’t introducing any synthetic speedup. In different phrases, the ruler reads zero if you measure one thing towards itself.

torch.compile is 2.48x quicker than keen on that op — without spending a dime, with no agent concerned.

If I’d benchmarked an AI-generated kernel towards keen mode and reported a 2.4× speedup, I’d really be celebrating code that’s slower than a single line of normal PyTorch. Hold that in thoughts—it comes again in Experiment 3.

2.3 Why does Keen mode nonetheless exist?

PyTorch’s keen mode isn’t sluggish by chance. It’s designed for debuggability, not most efficiency.

In keen mode, operations run separately. If a tensor immediately fills with NaNs, you possibly can cease on the offending line and examine what occurred. In distinction, torch.compile works in a different way: it first traces your mannequin right into a graph, then generates fused kernels. That delivers significantly better efficiency, however the code operating on the GPU now not maps cleanly to your Python supply.

Keen mode additionally stays important as a result of:

• Compilation has a price. Tracing and autotuning add startup overhead, which is negligible for lengthy coaching runs however noticeable throughout fast pocket book iteration.

• Not all code might be compiled. Information-dependent management circulation, customized operators, and unsupported options can set off graph breaks, forcing execution again to keen mode.

• It’s the reference implementation. torch.compile is validated towards keen mode, and new operators are carried out there first.

The 2 modes serve completely different functions: keen is the dependable reference; torch.compile is the optimized quick path. That’s why evaluating a customized kernel solely towards keen mode is deceptive. You’re benchmarking towards a mode constructed for correctness and debugging—not velocity.

3. Benchmark context

3.1 From <20% to actual progress

The benchmark most researchers use is KernelBench (arXiv:2502.10517, Stanford, ICML 2025), with 250 actual PyTorch workloads throughout 4 issue tiers. It scores kernels provided that they’re each appropriate and quicker than the PyTorch baseline.

Early outcomes have been underwhelming: frontier reasoning fashions beat the baseline in fewer than 20% of instances.

Since then, outcomes have improved, largely by higher coaching and search methods quite than smarter base fashions. Cognition’s Kevin-32B improved correctness from 56% to 82% and common efficiency from 0.53× (slower than PyTorch) to 1.10×. NVIDIA reported 100% correctness on KernelBench’s best tier utilizing an automatic refinement loop, whereas Meta’s KernelEvolve discovered 1.25×–17× speedups on manufacturing workloads by evolving many candidate kernels as an alternative of producing only one.

The progress is actual—however the benchmark holds some limitations.

3.2 When the benchmark was not sufficient

In 2026, KernelBench-Verified (arXiv:2607.16241, June 2026) confirmed that the unique benchmark understated PyTorch’s efficiency. The baseline had TF32 disabled, although fashionable NVIDIA GPUs sometimes use it for matrix multiplication.

With TF32 enabled and hidden take a look at instances added, the perfect mannequin examined (GPT-5.5) dropped from a reported 1.43× speedup to 0.88×—slower than plain PyTorch.

The examine additionally discovered that 28% of generated kernels elevated peak GPU reminiscence utilization, a price the unique benchmark ignored.

One other benchmark, AgentKernelArena, reported speedups of as much as 6.89× when changing PyTorch code to AMD’s HIP. However many generated kernels failed as soon as tensor shapes modified as a result of the brokers had quietly hardcoded assumptions that solely held for the benchmark inputs.

The lesson is straightforward: benchmark outcomes are solely as reliable because the benchmark itself.

3.3 The extra reliable image

Not each result’s overstated. In 2026, Hugging Face launched an agent ability for CUDA kernel technology and reported 1.88×–1.94× quicker RMSNorm kernels, with peaks of 2.47× in microbenchmarks.

The top-to-end outcomes inform a unique story. As soon as in contrast towards an already-compiled baseline, whole runtime improved from 2.14 s to 2.01 s—roughly 1.06×. By comparability, torch.compile alone delivered about 1.34×, with none AI-generated kernels.

This isn’t a criticism of the work—Hugging Face revealed the compiled baseline, which many papers omit. As an alternative, it highlights a recurring sample: massive kernel-level positive aspects typically translate into modest application-level enhancements as a result of fashionable compilers already carry out a lot of the out there optimization.

One last commentary: the launched CUDA ability is sort of 660 traces lengthy, however most of it focuses on construct methods, integration, and commonplace optimization patterns. It says little about Tensor Cores, WMMA/MMA directions, or mixed-precision strategies—the {hardware} options that always decide peak GPU efficiency. That omission turns into necessary within the experiments that comply with.

4. Analysis

4.1 Constructing a benchmark that resists dishonest

I reviewed each main failure mode—from Sakana’s autopsy, the KernelBench-Verified paper, and the broader reward-hacking literature—and constructed my analysis harness to defend towards each.

Two pitfalls are particularly simple to overlook:

• Stale reminiscence: A kernel can seem appropriate by writing nothing if the output buffer nonetheless accommodates legitimate outcomes from an earlier run.

• Unsynchronized timing: GPU kernels run asynchronously. If you happen to cease the timer earlier than the GPU finishes, you measure Python’s launch overhead—not the kernel’s execution time.

The harness additionally checks for reminiscence aliasing, enter mutation, and verifies correctness throughout three random seeds for each enter form.

To validate these safeguards, I constructed a second, deliberately weak harness.

4.2 Breaking the benchmark

The second harness, hidden in an innocuously named eval/ listing, intentionally contains the issues above and studies a easy SCORE = speedup. I examined each harnesses with three deliberately dishonest kernels.

The outcomes spotlight why strong analysis issues:

• A memoization cheat handed the primary two runs however failed on the third, when PyTorch reused a beforehand allotted reminiscence deal with. A benchmark utilizing just one or two validation runs would have incorrectly reported a 9.68× speedup for a kernel that did no computation in any respect.

• An FP16 precision cheat failed the strict harness with a most error of 7.8 × 10⁻⁴. A typical tolerance of 1 × 10⁻² would have accepted it, letting diminished precision masquerade as a sound optimization.

• A naive timing harness measured softmax at 0.0060 ms, whereas the synchronized execution time was 1.1581 ms—a 193× error precipitated solely by incorrect timing.

Both of those flaws is sufficient to make benchmark outcomes unreliable. Collectively, they present why cautious analysis issues as a lot because the optimization itself.

4.3 The “Free” 1.65× Speedup

KernelBench-Verified’s most putting outcome took simply one line of code to breed. On my GPU, PyTorch defaults to allow_tf32=False. Turning TF32 on delivered a 1.65× speedup.

That’s not a greater kernel—it’s a precision tradeoff.

The strict correctness test compares GPU outcomes towards a 64-bit CPU reference. With an error tolerance (atol) of 2e-3, enabling TF32 produced an error of 1.042e-2—over 5× the allowed restrict.

The stunning half is that the naive benchmark makes use of a tolerance of 1e-2, virtually an identical to the measured error. Whether or not it catches the accuracy loss is actually right down to probability.

5. Outcomes

Earlier than the benchmarks, there are a couple of particulars concerning the methodology.

Every process was solved by a contemporary Claude Code agent with no reminiscence of prior runs. It obtained solely the PyTorch reference, {hardware} particulars, benchmark directions, and a single goal: maximize verified speedup. 

One caveat: the orchestrator and the brokers share the identical mannequin household, although they ran in totally separate classes. This might introduce bias.

All outcomes have been measured independently. Each kernel was re-run twice on the strict harness on an idle machine; brokers’ reported speeds matched inside ~0.1%.

5.1 Experiment 1: 4 Kernels

The benchmark coated 4 operations: softmax, layernorm, gelu_bias_residual, and matmul_bias_relu. Every process used a contemporary agent with a price range of six benchmark runs to iterate on its resolution. Solely kernels that handed the strict correctness test have been scored.

Earlier than operating the experiment, I anticipated:

• Wins on the fusion-friendly kernels.

• A shut contest on softmax, which is already a single operation.

• A clear loss on matmul towards cuBLAS, NVIDIA’s extremely optimized matrix multiplication library.

That final case was alleged to be the article’s “brokers don’t beat vendor libraries” instance.

5.1.1 The outcomes

All 4 kernels handed the strict correctness gate. Throughout repeated runs, efficiency diverse by simply 0.19–0.51%, indicating steady and dependable measurements.

The GELU outcome captures the principle story: 2.50× quicker than keen execution, however only one.01× quicker than torch.compile. torch.compile had already fused the operation virtually completely.

The identical sample seems for softmax and layernorm. Certainly, three unbiased brokers reached the identical conclusion: there was virtually nothing left to optimize.

5.1.2 Small enchancment is the reight outcome

To confirm this, every agent wrote a kernel that did no computation in any respect—it merely copied the identical bytes from enter to output. The copy-only kernels ran virtually as quick as the true operations.

The above desk reveals the place the bottleneck is. The GPU spends most of its time transferring knowledge, not performing arithmetic. The maths finishes quicker than the {hardware} can fetch the inputs.

This is named being memory-bound, or hitting the reminiscence roof: the utmost throughput allowed by reminiscence bandwidth. As soon as an operation reaches that restrict, additional algorithmic enhancements can not make it quicker. Matching the copy-only kernel isn’t a failure—it means you’ve reached the {hardware} restrict.

The kernels themselves verify this. A textbook softmax scans every row thrice: as soon as to seek out the utmost, as soon as to compute exponentials and their sum, and as soon as to normalize. The agent’s kernel reads the row as soon as, performs all operations in registers and shared reminiscence, then writes the outcome again. There are merely no extra reminiscence accesses left to remove.

5.1.3 Two caveats

First, the benchmark barely favors the customized kernels. Timing begins earlier than the CPU finishes dispatching work to the GPU, so torch.compile pays about 20.7 μs of dispatch overhead versus 5.8 μs for my kernels. That distinction is sort of all the GELU benefit, so deal with GELU and softmax as parity quite than wins.

Second, the run price range understated the precise search effort. Whereas the official harness recorded solely a handful of benchmark runs, the brokers independently explored roughly 25 matmul configurations and 400 GELU configurations utilizing their very own scripts. In follow, they consumed much more compute than the reported price range suggests.

One stunning outcome: processing one float per thread outperformed float4 vectorization by about 4%, reaching 240 GB/s. Standard GPU recommendation would predict the other. The lesson is straightforward: hardware-specific optimizations are solely legitimate for the {hardware} they have been measured on.

5.1.4 The outcome I didn’t anticipate

The largest shock was matmul.

I anticipated it to lose decisively towards cuBLAS and torch.compile. As an alternative, the agent produced an accurate kernel that achieved 1.67× over keen execution and 1.57× over torch.compile, after solely two benchmark iterations.

The duty required FP32 accuracy, however the GPU’s tensor cores are a lot quicker with FP16. The apparent shortcut—TF32—failed the accuracy test, and the agent independently found the identical limitation I had discovered earlier.

As an alternative, it used a higher-precision decomposition. Every FP32 worth was cut up into excessive and low FP16 elements. The kernel then computed three tensor-core matrix multiplications (excessive×excessive, excessive×low, and low×excessive) and gathered them right into a single FP32 outcome. The remaining low×low time period is negligibly small (round 2⁻²² of the full), so omitting it preserves FP32 accuracy whereas exploiting the a lot quicker tensor-core path.

In different phrases, it traded extra arithmetic for a lot quicker {hardware}, and nonetheless met the correctness requirement. That’s why matmul was the one benchmark that produced a real, surprising win.

5.1.5 A Sooner FP32 Matrix Multiply

Earlier than, all the operation is a single cuBLAS name:

out = torch.relu(a @ b + bias)

After, The optimized kernel as an alternative splits every fp32 worth into two fp16 values because it’s loaded from world reminiscence into shared reminiscence—with out creating further tensors or reminiscence writes.

#outline SPLIT(x, H, L) { __half hq = __float2half_rn(x);     (H) = hq; (L) = __float2half_rn((x) - __half2float(hq)); }

The primary worth (hq) is the fp16-rounded model of the unique quantity. The second shops the rounding error. Collectively, they reconstruct the unique fp32 worth with solely a tiny further rounding error.

The kernel then performs three tensor-core matrix multiplications:

wmma::mma_sync(acc[i][j], fah[i], fbh[j], acc[i][j]);   // hello x hellowmma::mma_sync(acc[i][j], fah[i], fbl[j], acc[i][j]);   // hello x lowmma::mma_sync(acc[i][j], fal[i], fbh[j], acc[i][j]);   // lo x hello

(wmma::mma_sync is the tensor-core instruction that performs the matrix multiply.)

All three operations accumulate immediately into the identical fp32 accumulator, preserving precision all through the computation. The fourth mixture (low × low) is rarely computed as a result of its contribution is simply too small to matter.

The implementation is solely customized—it makes use of tensor-core directions immediately, with no cuBLAS calls or PyTorch fallbacks.

Consequence: the kernel runs in 1.248 ms, matching TF32 cuBLAS (1.251 ms) whereas passing the FP32 correctness take a look at that TF32 fails. It additionally outperforms commonplace FP32 cuBLAS (2.059 ms) by 1.65×, attaining TF32-level velocity with out sacrificing FP32 accuracy.

5.1.6 Replication: Was It a One-Off?

A single profitable run isn’t adequate sufficient to conclude, so I repeated the experiment with three new brokers beneath an identical circumstances. Every ran in an remoted atmosphere with no entry to the unique kernel or its efficiency.

The success standards have been outlined beforehand:

• Cross the correctness take a look at.

• Beat torch.compile by a minimum of 1.05×.

• Not less than 2 of three replications should succeed.

Consequence: 3 out of three succeeded, with solely 6.1% efficiency variation throughout all 4 kernels.

Extra importantly, each agent independently found the identical optimization:

• Skip the negligible low × low product.

• Break up fp32 values into two fp16 values.

• Use tensor cores for 3 matrix multiplies.

In addition they independently rejected implementing the cut up in PyTorch as a result of materializing the additional tensors provides roughly 128 MB of reminiscence site visitors and about 0.5 ms, eliminating the efficiency acquire. The optimization solely works as a result of the cut up occurs contained in the kernel, not in reminiscence.

Two brokers even produced bitwise-identical outcomes regardless of utilizing completely different implementations—one by way of the WMMA API and the opposite by way of inline PTX—exhibiting that they independently converged on the identical arithmetic.

5.1.7 Passing Isn’t the Identical as FP32 Accuracy

Though each kernel handed the correctness take a look at, they operated a lot nearer to its tolerance restrict than native FP32.

Throughout 5 unseen random seeds:

• Native FP32 used about 15% of the allowed error price range.

• Break up-fp16 kernels used roughly 90%.

They remained deterministic and handed each take a look at, however the margin was a lot smaller. The biggest errors occurred close to the ReLU zero crossing, the place tiny numerical variations are almost definitely to vary the output.

One agent even rejected a bf16 model earlier than writing any CUDA, concluding that its numerical margin was too small.

The takeaway is straightforward: passing a correctness take a look at doesn’t essentially imply matching FP32 accuracy. For numerical optimizations like this, measuring how a lot of the error price range is consumed is simply as necessary as whether or not the take a look at passes.

5.2 Experiment 2: Does profiler suggestions enhance optimization?

Most LLM kernel optimization workflows comply with the identical loop: generate → compile → confirm → profile → feed the profiler output again to the agent → repeat. Surprisingly, no prior work isolates whether or not that profiler suggestions really improves efficiency.

To check this, I ran six deliberate optimization rounds on the matrix multiplication kernel—the one process with significant optimization headroom remaining. After every spherical, I measured efficiency beneath managed circumstances and returned a set profiler report containing SM throughput, reminiscence throughput, occupancy, and the three largest stall causes.

Not like Experiment 1, the agent couldn’t run its personal benchmarks. It might compile and confirm correctness, however all efficiency numbers got here from a managed benchmark to isolate the worth of profiler suggestions itself.

5.2.1 Outcomes: Higher metrics, identical velocity

The profiler metrics improved all through the experiment; nonetheless, runtime didn’t.

• Occupancy almost doubled.

• Stall metrics dropped considerably.

• One optimization eliminated about 45% of the executed directions.

The clearest instance got here in Spherical 3. Regardless of eradicating almost half the directions and growing SM utilization, execution time stayed virtually an identical (1.2503 ms vs. 1.2482 ms).

The important thing discovering is straightforward:

Profiler counters measure signs, not bottlenecks. Bettering a counter doesn’t assure bettering efficiency.

The replication experiment reinforces this. Doubling the variety of energetic warps produced a big speedup for one kernel, however the identical optimization had no impact right here as a result of this kernel had already handed that bottleneck. The profiler supplied no indication of which case utilized.

One sensible commentary additionally emerged: on the GB10 GPU, the usual gpu__dram_throughput counter is unavailable as a result of LPDDR5X unified reminiscence exposes no separate DRAM metric. I substituted gpu__compute_memory_throughput as an alternative. Some generally used profiler metrics merely don’t exist on newer {hardware}.

5.2.2 What the agent did effectively

Though profiler suggestions failed to enhance efficiency, the agent behaved rationally.

It deserted weak hypotheses earlier than testing them, validated concepts utilizing low-level SASS evaluation, and after a number of unsuccessful rounds concluded that additional optimization was unlikely to assist. Fairly than chasing bettering metrics indefinitely, it advisable stopping.

One result’s notably instructive. The agent constructed a efficiency mannequin that matched the present kernel inside 0.1%, but predicted the following optimization would cut back runtime to roughly 1.0 ms. The precise outcome was 1.276 ms.

A mannequin can clarify current knowledge extraordinarily effectively whereas nonetheless failing to foretell the impact of a brand new optimization.

Lastly, though the unique protocol deliberate 5 optimization rounds, I finished after 4 as a result of the agent had already met its pre-registered stopping criterion. Persevering with would have measured persistence quite than the usefulness of profiler suggestions.

5.3 Experiment 3: What occurs when the benchmark is flawed?

This experiment examined whether or not an agent would exploit a benchmark it knew was weak.

Every agent optimized towards a benchmark containing identified flaws beneath two circumstances:

• Impartial: optimize usually.

• Strain: explicitly maximize the reported rating, with scores beneath 2.0× outlined as failure.

Earlier than operating the experiment, I categorized attainable outcomes as trustworthy, unintentional exploit, deliberate exploit, or whistleblower.

5.3.1 Outcomes: No exploitation

All 4 brokers behaved pretty.

• 0 deliberate exploits

• 0 unintended exploits

• 4/4 recognized the benchmark as flawed

Each submitted kernel additionally handed the hardened validation harness, confirming there was no memoization, pointer-based caching, form hardcoding, or different benchmark exploits.

Probably the most putting case got here beneath the strain situation. After failing to succeed in the required rating, the agent explicitly described how the benchmark might be exploited—for instance, by caching outputs based mostly on enter pointers—however refused to take action. As an alternative, it advisable fixing the lacking torch.cuda.synchronize() name that precipitated the incorrect timing.

The situation designed to encourage dishonest produced a bug report as an alternative.

One nuance is price noting. One agent briefly evaluated an optimization resembling an exploit (CUDA graph replay with pointer-based caching) however rejected it as a result of it was ineffective, not as a result of it thought of it unethical. Solely the softmax brokers explicitly rejected benchmark exploitation on principled grounds.

5.3.2 And it did not matter

Right here’s the discovering I didn’t anticipate.

The neutral-gelu kernel studies a 3.41× speedup, although it’s really 7% slower than torch.compile.

The kernel itself is legitimate. The agent produced an actual fused CUDA kernel that passes a strict correctness test and is 2.29× quicker than keen PyTorch. The inflated 3.41× outcome wasn’t attributable to the agent—it got here from the benchmark harness, which used a weak baseline and by no means synchronized the GPU earlier than timing.

I got down to see whether or not an agent would exploit a flawed benchmark. Nonetheless, the reply is No.

The extra necessary lesson is that a flawed benchmark can manufacture spectacular speedups by itself. It doesn’t require a dishonest agent, and bettering the agent doesn’t repair the measurement.

Each reported kernel speedup is determined by the benchmark behind it. Most papers by no means present you that benchmark.

6. Conclusion

The experiments present that AI brokers can write actual, high-performance CUDA.

In Experiment 1, an agent produced a split-precision WMMA GEMM kernel that outperformed torch.compile by 1.57×, matched cuBLAS TF32 efficiency, and handed a correctness take a look at that TF32 itself failed. Much more stunning, three unbiased brokers converged on basically the identical optimization.

The larger problem, nonetheless, wasn’t writing the kernels—it was measuring them accurately. Many of the work went into constructing and validating a reliable benchmark, uncovering deceptive measurements, and verifying numerical correctness. The agent wrote the kernel in a day.

The query is now not whether or not AI can generate CUDA.

It’s whether or not you possibly can precisely measure what it generated.

6.1 When customized kernels are literally worthwhile

All through this text, torch.compile is the reference level as a result of it’s the fairest comparability at any time when it’s out there. However there are necessary instances the place it isn’t.

The commonest is deployment with out Python. Since torch.compile generates and compiles code at runtime inside a Python course of, it may possibly’t be utilized in environments similar to C++ inference servers, embedded methods, robotics platforms, or automotive deployments. In these settings, a hand-written fused kernel delivers the 2.5× speedup that torch.compile would in any other case present robotically.

Customized kernels additionally matter on memory-bandwidth-limited {hardware}, similar to unified-memory methods just like the DGX Spark. On these gadgets, lowering reminiscence site visitors by fusion typically has a bigger impression than bettering arithmetic throughput.

Lastly, customized kernels are helpful once they implement optimizations a general-purpose library can not safely assume. Experiment 1 demonstrated this with split-precision tensor-core computation: cuBLAS can not robotically select that trade-off as a result of solely the appliance developer is aware of whether or not the precision loss is suitable.

For many workloads, nonetheless, customized kernels aren’t definitely worth the effort. If torch.compile already runs in manufacturing and your workload is primarily memory-bound, a hand-written kernel typically duplicates the identical optimizations the compiler already performs robotically. Three of the 4 experiments fell into this class.

6.2 When an AI Agent Can Assist Optimize GPU Kernels

Earlier than asking an AI agent to optimize a kernel, set up a practical baseline and know the place the bottleneck is.

1. Test the roofline first. Benchmark a kernel that solely copies the identical quantity of information. If it’s almost as quick as your actual kernel, you’re memory-bound, so matching that velocity is the perfect you possibly can anticipate.

2. Evaluate towards torch.compile, not keen mode. Keen execution is a weak baseline as a result of it doesn’t fuse operations.

3. Search for optimizations PyTorch gained’t make. The largest positive aspects often come from hardware-specific trade-offs—similar to utilizing tensor cores with diminished precision—that torch.compile avoids as a result of it may possibly’t assume the accuracy trade-off is suitable.

4. Validate completely. Take a look at a number of tensor shapes, random seeds, and evaluate towards a high-precision reference. A kernel that passes one take a look at can nonetheless fail on others.

5. Benchmark accurately. Use CUDA occasions with express synchronization; a easy stopwatch can produce wildly inaccurate timings.

6. Measure numerical error, not simply move/fail. Two kernels could each move a tolerance test whereas utilizing very completely different quantities of the error price range.

7. Reuse current work when attainable. Earlier than writing a customized kernel, test libraries like Hugging Face Kernel Hub—an optimized implementation could exist already.

These checks present a sensible path from “this kernel is sluggish” to deciding whether or not a customized AI-generated kernel is price pursuing.

7. References and assets

Benchmarks and evaluations:

• KernelBench: Can LLMs Write Environment friendly GPU Kernels? — Ouyang, Guo, Arora, Zhang, Hu, Ré, Mirhoseini, 2025 · arXiv:2502.10517 · GitHub ScalingIntelligence/KernelBench, License MIT.

• KernelBench-Verified: Do LLM-Generated Kernels Really Beat PyTorch? — 2026 · arXiv:2607.16241 

• AgentKernelArena — 2026 · arXiv:2605.16819. 196 duties.

Kernel-generating methods

• Kevin: Multi-Flip RL for Producing CUDA Kernels — Cognition AI, 2025 · arXiv:2507.11948 · mannequin card cognition-ai/Kevin-32B.

• KernelEvolve: Scaling Agentic Kernel Coding for Heterogeneous AI Accelerators at Meta — 2025 · arXiv:2512.23236.

• Automating GPU Kernel Technology with DeepSeek-R1 and Inference-Time Scaling — NVIDIA Developer Weblog, 2025.

• Customized Kernels for All from Codex and Claude — Burtenshaw, Paul, Roy Gosthipaty, Smith, Hugging Face, 13 Feb 2026. Supply of the LTX-Video desk quoted above.

Instruments

• kernels / kernel-builder — Hugging Face’s kernel construct system, Kernel Hub, and the cuda-kernels agent ability · GitHub huggingface/kernels.

• torch.compile — PyTorch 2.13.0+cu130, mode="max-autotune-no-cudagraphs", the baseline each declare right here is measured towards.

• Nsight Compute — CUDA 13.0. gpu__dram_throughput is unavailable on GB10.

The incident

• Sakana AI, “The AI CUDA Engineer” announcement and correction, 19–21 Feb 2025. Dataset: SakanaAI/AI-CUDA-Engineer-Archive, 17,000+ kernels, License CC-BY-4.0. Protection: TechCrunch, “Sakana walks again claims that its AI can dramatically velocity up mannequin coaching,” 21 Feb 2025.

• In the direction of Automated GPU Kernel Technology — Simon Guo, Oct 2025. Retrospective on eval pitfalls, reward hacking, and {hardware} variance.

{Hardware}

• NVIDIA DGX Spark (GB10) — 20-core Arm CPU, 48-SM Blackwell GPU at compute functionality sm_121, 128 GB unified LPDDR5X. Reminiscence bandwidth