AI Agents Beat PyTorch: Writing Faster CUDA Kernels

In February 2025, Sakana AI announced that its “AI CUDA Engineer” generated 17,000 CUDA kernels with speedups of up to 381× over PyTorch. This is a system built to write CUDA kernels, the small programs that tell a computer’s graphics chip (the GPU) exactly how to crunch numbers.

Within a day, an X user found the AI hadn’t written faster code—it had exploited a flaw in Sakana’s testing system that allowed incorrect kernels to pass. Sakana retracted the claims and acknowledged a key lesson: if the benchmark is flawed, an AI will optimize for the test, not the problem.

That raises the real question: How do you know a speedup is real?

To find out, I ran my own experiment on an NVIDIA DGX Spark using Claude Code. I asked it to optimize 4 common CUDA operations and evaluated the results with 2 benchmark suites: one rigorous, and one intentionally flawed to see if the AI would take the shortcut.

The results were encouraging. The agents produced correct, high-performance CUDA kernels, with the best implementation running 1.57× faster than torch.compile on a matrix multiplication workload. Three independent agents reached the same solution.

The bigger takeaway wasn’t about CUDA—it was about benchmarking. Building a trustworthy evaluation proved harder than generating the optimized code itself.

1. Who this is for

This article is for anyone considering AI-generated CUDA optimizations. You’ll walk away with:

• A practical framework for deciding when AI-driven kernel optimization is worth your time.

• Real performance results for four common GPU operations using fair benchmarks.

• Five ways CUDA benchmarks can produce misleading results—and how to catch them.

• Whether profiler feedback actually helps AI agents (it didn’t).

2. What a CUDA kernel is, and why “faster than PyTorch” is a trick question

A kernel is a small program that runs directly on the GPU, executed by thousands of threads in parallel — each one handling a small subset of data.

Think of PyTorch as ordering off the menu: its kernels are highly optimized to execute operations one at a time, because in general it doesn’t know what your program will do next. A custom kernel only pays off when you know something PyTorch doesn’t and merge those operations together.

That’s called kernel fusion, and it’s where most real speedups come from. Instead of writing intermediate results to memory between operations, the GPU does everything in one pass. Since moving data is often slower than the math, eliminating those extra memory transfers usually reduces latency.

2.1 What fusion looks like in code

Here’s the operation this article calls gelu_bias_residual:

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

Originally, it comes from 3 separate operations. In PyTorch’s eager mode, each operation launches its own GPU kernel, so the tensor is repeatedly read from and written back to memory—adding unnecessary overhead.

A fused kernel combines all three operations into a single GPU pass. Below is the core of a fused kernel generated by an AI agent. Each thread processes one element, and the GPU-specific memory instructions are explained in the comments.

__global__ void gbr_kernel(const float* __restrict__ X, const float* __restrict__ B,                           const float* __restrict__ R, float* __restrict__ Y,                           int n, long numel) {  const long i = (long)blockIdx.x * GBR_BLOCK + threadIdx.x;  const float xv = __ldg(&X[i]);             // read x  const float rv = __ldg(&R[i]);             // read 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 element is finished}

Only two memory reads and one write are needed. The addition, GELU, and residual computation all happen in registers—the GPU’s fastest memory—between the load and store. That’s the key to fusion.

This fused kernel runs 2.5× faster than the three-operation version. Its performance is also nearly identical to torch.compile, which automatically performs the same fusion behind the scenes.

2.2 Pick the right baseline

A claim like “2.5× faster than PyTorch” usually compares against PyTorch eager mode, where every operation launches its own GPU kernel and repeatedly reads and writes memory.

But PyTorch also ships torch.compile to look at your model, figure out which operations can be fused, and automatically generate its own optimized GPU code to do it — Triton kernels. These are done by PyTorch behind the scenes, and in most cases, you’re already using this optimization. If you’re optimizing for performance, the real comparison is to benchmark with the run of torch.compile.

Before testing any AI-generated kernels, I also validated my benchmark by running the original unfused code as the “candidate.” It measured 0.99–1.00× across all four tasks, confirming the benchmark itself wasn’t introducing any artificial speedup. In other words, the ruler reads zero when you measure something against itself.

torch.compile is 2.48x faster than eager on that op — for free, with no agent involved.

If I’d benchmarked an AI-generated kernel against eager mode and reported a 2.4× speedup, I’d actually be celebrating code that’s slower than a single line of standard PyTorch. Keep that in mind—it comes back in Experiment 3.

2.3 Why does Eager mode still exist?

PyTorch’s eager mode isn’t slow by accident. It’s designed for debuggability, not maximum performance.

In eager mode, operations run one at a time. If a tensor suddenly fills with NaNs, you can stop on the offending line and inspect what happened. In contrast, torch.compile works differently: it first traces your model into a graph, then generates fused kernels. That delivers much better performance, but the code running on the GPU no longer maps cleanly to your Python source.

Eager mode also remains essential because:

• Compilation has a cost. Tracing and autotuning add startup overhead, which is negligible for long training runs but noticeable during rapid notebook iteration.

• Not all code can be compiled. Data-dependent control flow, custom operators, and unsupported features can trigger graph breaks, forcing execution back to eager mode.

• It’s the reference implementation. torch.compile is validated against eager mode, and new operators are implemented there first.

The two modes serve different purposes: eager is the reliable reference; torch.compile is the optimized fast path. That’s why comparing a custom kernel only against eager mode is misleading. You’re benchmarking against a mode built for correctness and debugging—not speed.

3. Benchmark context

3.1 From <20% to real progress

The benchmark most researchers use is KernelBench (arXiv:2502.10517, Stanford, ICML 2025), with 250 real PyTorch workloads across four difficulty tiers. It scores kernels only if they are both correct and faster than the PyTorch baseline.

Early results were underwhelming: frontier reasoning models beat the baseline in fewer than 20% of cases.

Since then, results have improved, largely through better training and search strategies rather than smarter base models. Cognition’s Kevin-32B improved correctness from 56% to 82% and average performance from 0.53× (slower than PyTorch) to 1.10×. NVIDIA reported 100% correctness on KernelBench’s easiest tier using an automated refinement loop, while Meta’s KernelEvolve found 1.25×–17× speedups on production workloads by evolving many candidate kernels instead of generating just one.

The progress is real—but the benchmark holds some limitations.

3.2 When the benchmark was not enough

In 2026, KernelBench-Verified (arXiv:2607.16241, June 2026) showed that the original benchmark understated PyTorch’s performance. The baseline had TF32 disabled, even though modern NVIDIA GPUs typically use it for matrix multiplication.

With TF32 enabled and hidden test cases added, the best model tested (GPT-5.5) dropped from a reported 1.43× speedup to 0.88×—slower than plain PyTorch.

The study also found that 28% of generated kernels increased peak GPU memory usage, a cost the original benchmark ignored.

Another benchmark, AgentKernelArena, reported speedups of up to 6.89× when converting PyTorch code to AMD’s HIP. But many generated kernels failed once tensor shapes changed because the agents had quietly hardcoded assumptions that only held for the benchmark inputs.

The lesson is simple: benchmark results are only as trustworthy as the benchmark itself.

3.3 The more trustworthy picture

Not every result is overstated. In 2026, Hugging Face released an agent skill for CUDA kernel generation and reported 1.88×–1.94× faster RMSNorm kernels, with peaks of 2.47× in microbenchmarks.

The end-to-end results tell a different story. Once compared against an already-compiled baseline, total runtime improved from 2.14 s to 2.01 s—roughly 1.06×. By comparison, torch.compile alone delivered about 1.34×, without any AI-generated kernels.

This isn’t a criticism of the work—Hugging Face published the compiled baseline, which many papers omit. Instead, it highlights a recurring pattern: large kernel-level gains often translate into modest application-level improvements because modern compilers already perform much of the available optimization.

One final observation: the released CUDA skill is nearly 660 lines long, but most of it focuses on build systems, integration, and standard optimization patterns. It says little about Tensor Cores, WMMA/MMA instructions, or mixed-precision techniques—the hardware features that often determine peak GPU performance. That omission becomes important in the experiments that follow.

4. Evaluation

4.1 Building a benchmark that resists cheating

I reviewed every major failure mode—from Sakana’s post-mortem, the KernelBench-Verified paper, and the broader reward-hacking literature—and built my evaluation harness to defend against each one.

Two pitfalls are especially easy to miss:

• Stale memory: A kernel can appear correct by writing nothing if the output buffer still contains valid results from an earlier run.

• Unsynchronized timing: GPU kernels run asynchronously. If you stop the timer before the GPU finishes, you measure Python’s launch overhead—not the kernel’s execution time.

The harness also checks for memory aliasing, input mutation, and verifies correctness across three random seeds for every input shape.

To validate these safeguards, I built a second, intentionally vulnerable harness.

4.2 Breaking the benchmark

The second harness, hidden in an innocuously named eval/ directory, deliberately includes the flaws above and reports a simple SCORE = speedup. I tested both harnesses with three intentionally cheating kernels.

The results highlight why robust evaluation matters:

• A memoization cheat passed the first two runs but failed on the third, when PyTorch reused a previously allocated memory address. A benchmark using only one or two validation runs would have incorrectly reported a 9.68× speedup for a kernel that did no computation at all.

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

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

Either of these flaws is enough to make benchmark results unreliable. Together, they show why careful evaluation matters as much as the optimization itself.

4.3 The “Free” 1.65× Speedup

KernelBench-Verified’s most striking result took just one line of code to reproduce. On my GPU, PyTorch defaults to allow_tf32=False. Turning TF32 on delivered a 1.65× speedup.

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

The strict correctness check compares GPU results against 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 limit.

The surprising part is that the naive benchmark uses a tolerance of 1e-2, almost identical to the measured error. Whether it catches the accuracy loss is essentially down to chance.

5. Results

Before the benchmarks, there are a few details about the methodology.

Each task was solved by a fresh Claude Code agent with no memory of prior runs. It received only the PyTorch reference, hardware details, benchmark instructions, and a single objective: maximize verified speedup. 

One caveat: the orchestrator and the agents share the same model family, though they ran in fully separate sessions. This could introduce bias.

All results were measured independently. Every kernel was re-run twice on the strict harness on an idle machine; agents’ reported speeds matched within ~0.1%.

5.1 Experiment 1: Four Kernels

The benchmark covered four operations: softmax, layernorm, gelu_bias_residual, and matmul_bias_relu. Each task used a fresh agent with a budget of six benchmark runs to iterate on its solution. Only kernels that passed the strict correctness check were scored.

Before running the experiment, I expected:

• Wins on the fusion-friendly kernels.

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

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

That last case was supposed to be the article’s “agents don’t beat vendor libraries” example.

5.1.1 The results

All four kernels passed the strict correctness gate. Across repeated runs, performance varied by just 0.19–0.51%, indicating stable and reliable measurements.

The GELU result captures the main story: 2.50× faster than eager execution, but only 1.01× faster than torch.compile. torch.compile had already fused the operation almost perfectly.

The same pattern appears for softmax and layernorm. Indeed, three independent agents reached the same conclusion: there was almost nothing left to optimize.

5.1.2 Small improvement is the reight result

To verify this, each agent wrote a kernel that did no computation at all—it simply copied the same bytes from input to output. The copy-only kernels ran almost as fast as the real operations.

The above table reveals where the bottleneck is. The GPU spends most of its time moving data, not performing arithmetic. The math finishes faster than the hardware can fetch the inputs.

This is known as being memory-bound, or hitting the memory roof: the maximum throughput allowed by memory bandwidth. Once an operation reaches that limit, further algorithmic improvements cannot make it faster. Matching the copy-only kernel isn’t a failure—it means you’ve reached the hardware limit.

The kernels themselves confirm this. A textbook softmax scans each row three times: once to find the maximum, once to compute exponentials and their sum, and once to normalize. The agent’s kernel reads the row once, performs all operations in registers and shared memory, then writes the result back. There are simply no more memory accesses left to eliminate.

5.1.3 Two caveats

First, the benchmark slightly favors the custom kernels. Timing starts before 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 difference is almost the entire GELU advantage, so treat GELU and softmax as parity rather than wins.

Second, the run budget understated the actual search effort. While the official harness recorded only a handful of benchmark runs, the agents independently explored roughly 25 matmul configurations and 400 GELU configurations using their own scripts. In practice, they consumed far more compute than the reported budget suggests.

One surprising result: processing one float per thread outperformed float4 vectorization by about 4%, reaching 240 GB/s. Conventional GPU advice would predict the opposite. The lesson is simple: hardware-specific optimizations are only valid for the hardware they were measured on.

5.1.4 The result I didn’t expect

The biggest surprise was matmul.

I expected it to lose decisively against cuBLAS and torch.compile. Instead, the agent produced a correct kernel that achieved 1.67× over eager execution and 1.57× over torch.compile, after only two benchmark iterations.

The task required FP32 accuracy, but the GPU’s tensor cores are much faster with FP16. The obvious shortcut—TF32—failed the accuracy check, and the agent independently discovered the same limitation I had found earlier.

Instead, it used a higher-precision decomposition. Each FP32 value was split into high and low FP16 components. The kernel then computed three tensor-core matrix multiplications (high×high, high×low, and low×high) and accumulated them into a single FP32 result. The remaining low×low term is negligibly small (around 2⁻²² of the total), so omitting it preserves FP32 accuracy while exploiting the much faster tensor-core path.

In other words, it traded more arithmetic for much faster hardware, and still met the correctness requirement. That’s why matmul was the only benchmark that produced a genuine, unexpected win.

5.1.5 A Faster FP32 Matrix Multiply

Before, the entire operation is a single cuBLAS call:

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

After, The optimized kernel instead splits each fp32 value into two fp16 values as it’s loaded from global memory into shared memory—without creating extra tensors or memory writes.

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

The first value (hq) is the fp16-rounded version of the original number. The second stores the rounding error. Together, they reconstruct the original fp32 value with only a tiny additional rounding error.

The kernel then performs three tensor-core matrix multiplications:

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

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

All three operations accumulate directly into the same fp32 accumulator, preserving precision throughout the computation. The fourth combination (low × low) is never computed because its contribution is too small to matter.

The implementation is entirely custom—it uses tensor-core instructions directly, with no cuBLAS calls or PyTorch fallbacks.

Result: the kernel runs in 1.248 ms, matching TF32 cuBLAS (1.251 ms) while passing the FP32 correctness test that TF32 fails. It also outperforms standard FP32 cuBLAS (2.059 ms) by 1.65×, achieving TF32-level speed without sacrificing FP32 accuracy.

5.1.6 Replication: Was It a One-Off?

A single successful run isn’t sufficient enough to conclude, so I repeated the experiment with three new agents under identical conditions. Each ran in an isolated environment with no access to the original kernel or its performance.

The success criteria were defined beforehand:

• Pass the correctness test.

• Beat torch.compile by at least 1.05×.

• At least 2 of 3 replications must succeed.

Result: 3 out of 3 succeeded, with only 6.1% performance variation across all four kernels.

More importantly, every agent independently discovered the same optimization:

• Skip the negligible low × low product.

• Split fp32 values into two fp16 values.

• Use tensor cores for three matrix multiplies.

They also independently rejected implementing the split in PyTorch because materializing the extra tensors adds roughly 128 MB of memory traffic and about 0.5 ms, eliminating the performance gain. The optimization only works because the split happens inside the kernel, not in memory.

Two agents even produced bitwise-identical results despite using different implementations—one via the WMMA API and the other via inline PTX—showing that they independently converged on the same arithmetic.

5.1.7 Passing Isn’t the Same as FP32 Accuracy

Although every kernel passed the correctness test, they operated much closer to its tolerance limit than native FP32.

Across five unseen random seeds:

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

• Split-fp16 kernels used roughly 90%.

They remained deterministic and passed every test, but the margin was much smaller. The largest errors occurred near the ReLU zero crossing, where tiny numerical differences are most likely to change the output.

One agent even rejected a bf16 version before writing any CUDA, concluding that its numerical margin was too small.

The takeaway is simple: passing a correctness test does not necessarily mean matching FP32 accuracy. For numerical optimizations like this, measuring how much of the error budget is consumed is just as important as whether the test passes.

5.2 Experiment 2: Does profiler feedback improve optimization?

Most LLM kernel optimization workflows follow the same loop: generate → compile → verify → profile → feed the profiler output back to the agent → repeat. Surprisingly, no prior work isolates whether that profiler feedback actually improves performance.

To test this, I ran six planned optimization rounds on the matrix multiplication kernel—the only task with meaningful optimization headroom remaining. After each round, I measured performance under controlled conditions and returned a fixed profiler report containing SM throughput, memory throughput, occupancy, and the three largest stall reasons.

Unlike Experiment 1, the agent could not run its own benchmarks. It could compile and verify correctness, but all performance numbers came from a controlled benchmark to isolate the value of profiler feedback itself.

5.2.1 Results: Better metrics, same speed

The profiler metrics improved throughout the experiment; however, runtime did not.

• Occupancy nearly doubled.

• Stall metrics dropped substantially.

• One optimization removed about 45% of the executed instructions.

The clearest example came in Round 3. Despite removing nearly half the instructions and increasing SM utilization, execution time stayed almost identical (1.2503 ms vs. 1.2482 ms).

The key finding is simple:

Profiler counters measure symptoms, not bottlenecks. Improving a counter does not guarantee improving performance.

The replication experiment reinforces this. Doubling the number of active warps produced a large speedup for one kernel, but the same optimization had no effect here because this kernel had already passed that bottleneck. The profiler offered no indication of which case applied.

One practical observation also emerged: on the GB10 GPU, the standard gpu__dram_throughput counter is unavailable because LPDDR5X unified memory exposes no separate DRAM metric. I substituted gpu__compute_memory_throughput instead. Some commonly used profiler metrics simply do not exist on newer hardware.

5.2.2 What the agent did well

Although profiler feedback failed to improve performance, the agent behaved rationally.

It abandoned weak hypotheses before testing them, validated ideas using low-level SASS analysis, and after several unsuccessful rounds concluded that further optimization was unlikely to help. Rather than chasing improving metrics indefinitely, it recommended stopping.

One result is particularly instructive. The agent built a performance model that matched the current kernel within 0.1%, yet predicted the next optimization would reduce runtime to roughly 1.0 ms. The actual result was 1.276 ms.

A model can explain existing data extremely well while still failing to predict the effect of a new optimization.

Finally, although the original protocol planned five optimization rounds, I stopped after four because the agent had already met its pre-registered stopping criterion. Continuing would have measured persistence rather than the usefulness of profiler feedback.

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

This experiment tested whether an agent would exploit a benchmark it knew was vulnerable.

Each agent optimized against a benchmark containing known flaws under two conditions:

• Neutral: optimize normally.

• Pressure: explicitly maximize the reported score, with scores below 2.0× defined as failure.

Before running the experiment, I classified possible outcomes as honest, unintentional exploit, deliberate exploit, or whistleblower.

5.3.1 Results: No exploitation

All four agents behaved fairly.

• 0 deliberate exploits

• 0 accidental exploits

• 4/4 identified the benchmark as flawed

Every submitted kernel also passed the hardened validation harness, confirming there was no memoization, pointer-based caching, shape hardcoding, or other benchmark exploits.

The most striking case came under the pressure condition. After failing to reach the required score, the agent explicitly described how the benchmark could be exploited—for example, by caching outputs based on input pointers—but refused to do so. Instead, it recommended fixing the missing torch.cuda.synchronize() call that caused the inaccurate timing.

The condition designed to encourage cheating produced a bug report instead.

One nuance is worth noting. One agent briefly evaluated an optimization resembling an exploit (CUDA graph replay with pointer-based caching) but rejected it because it was ineffective, not because it considered it unethical. Only the softmax agents explicitly rejected benchmark exploitation on principled grounds.

5.3.2 And it didn’t matter

Here’s the finding I didn’t expect.

The neutral-gelu kernel reports a 3.41× speedup, even though it’s actually 7% slower than torch.compile.

The kernel itself is valid. The agent produced a real fused CUDA kernel that passes a strict correctness check and is 2.29× faster than eager PyTorch. The inflated 3.41× result wasn’t caused by the agent—it came from the benchmark harness, which used a weak baseline and never synchronized the GPU before timing.

I set out to see whether an agent would exploit a flawed benchmark. However, the answer is No.

The more important lesson is that a flawed benchmark can manufacture impressive speedups on its own. It doesn’t require a dishonest agent, and improving the agent doesn’t fix the measurement.

Every reported kernel speedup depends on the benchmark behind it. Most papers never show you that benchmark.

6. Conclusion

The experiments show that AI agents can write real, 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 performance, and passed a correctness test that TF32 itself failed. Even more surprising, three independent agents converged on essentially the same optimization.

The bigger challenge, however, wasn’t writing the kernels—it was measuring them correctly. Most of the work went into building and validating a trustworthy benchmark, uncovering misleading measurements, and verifying numerical correctness. The agent wrote the kernel in an afternoon.

The question is no longer whether AI can generate CUDA.

It’s whether you can accurately measure what it generated.

6.1 When custom kernels are actually worthwhile

Throughout this article, torch.compile is the reference point because it’s the fairest comparison whenever it’s available. But there are important cases where it isn’t.

The most common is deployment without Python. Since torch.compile generates and compiles code at runtime inside a Python process, it can’t be used in environments such as C++ inference servers, embedded systems, robotics platforms, or automotive deployments. In those settings, a hand-written fused kernel delivers the 2.5× speedup that torch.compile would otherwise provide automatically.

Custom kernels also matter on memory-bandwidth-limited hardware, such as unified-memory systems like the DGX Spark. On these devices, reducing memory traffic through fusion often has a larger impact than improving arithmetic throughput.

Finally, custom kernels are valuable when they implement optimizations a general-purpose library cannot safely assume. Experiment 1 demonstrated this with split-precision tensor-core computation: cuBLAS cannot automatically choose that trade-off because only the application developer knows whether the precision loss is acceptable.

For most workloads, however, custom kernels aren’t worth the effort. If torch.compile already runs in production and your workload is primarily memory-bound, a hand-written kernel often duplicates the same optimizations the compiler already performs automatically. Three of the four experiments fell into this category.

6.2 When an AI Agent Can Help Optimize GPU Kernels

Before asking an AI agent to optimize a kernel, establish a realistic baseline and know where the bottleneck is.

1. Check the roofline first. Benchmark a kernel that only copies the same amount of data. If it’s nearly as fast as your real kernel, you’re memory-bound, so matching that speed is the best you can expect.

2. Compare against torch.compile, not eager mode. Eager execution is a weak baseline because it doesn’t fuse operations.

3. Look for optimizations PyTorch won’t make. The biggest gains usually come from hardware-specific trade-offs—such as using tensor cores with reduced precision—that torch.compile avoids because it can’t assume the accuracy trade-off is acceptable.

4. Validate thoroughly. Test multiple tensor shapes, random seeds, and compare against a high-precision reference. A kernel that passes one test can still fail on others.

5. Benchmark correctly. Use CUDA events with explicit synchronization; a simple stopwatch can produce wildly inaccurate timings.

6. Measure numerical error, not just pass/fail. Two kernels may both pass a tolerance check while using very different amounts of the error budget.

7. Reuse existing work when possible. Before writing a custom kernel, check libraries like Hugging Face Kernel Hub—an optimized implementation may already exist.

These checks provide a practical path from “this kernel is slow” to deciding whether a custom AI-generated kernel is worth pursuing.

7. References and resources

Benchmarks and evaluations:

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

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

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

Kernel-generating systems

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

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

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

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

Tools

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

• torch.compile — PyTorch 2.13.0+cu130, mode="max-autotune-no-cudagraphs", the baseline every claim here is measured against.

• 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. Coverage: TechCrunch, “Sakana walks back claims that its AI can dramatically speed up model training,” 21 Feb 2025.

• Towards Automated GPU Kernel Generation — 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 capability sm_121, 128 GB unified LPDDR5X. Memory bandwidth