@elliotarledge: Claude Fable 5 [max] on KernelBench-Hard. The main kernel that impressed me was a B200 fp8 GEMM: it HAND WROTE raw SM10…

Claude Fable 5 [max] on KernelBench-Hard.

The main kernel that impressed me was a B200 fp8 GEMM: it HAND WROTE raw SM100 tcgen05 PTX. No CUTLASS or cuBLASLt.

2-CTA cluster tcgen05 mma into TMEM accumulators, TMA 3D loads through a 6-deep mbarrier pipeline, and a swizzled epilogue that folds the output scale straight into the TMA store. It leads the B200 fp8 board:

Fable 5 0.254 (hand tcgen05) GLM-5.2 0.200 Opus 4.8 0.196 GPT-5.5 0.146

That 0.254 is a 6-shape average dragged down by one memory-bound skinny decode shape (~3% of peak, no tensor core helps it). The compute-bound GEMM shapes hit 44-59% of B200 fp8 peak.

it literally spent the back half of a 2h44m session hand-building the tcgen05 path, and it came out about even with its own tuned Triton (~0.283 vs ~0.281). The graded number IS the hand PTX (the shipped kernel dispatches it on every shape, Triton is only the compile-failure fallback) it just barely edged the fallback it was racing instead of crushing it: ~90k extra tokens for under 1%. It also hit ~85-90% of a CUTLASS oracle it wrote to check itself.

On RTX PRO 6000 it tops the two cells that reward real tensor-core work: W4A16 GEMM at 0.348 of roofline and MoE SwiGLU at 0.108, both board-best, ahead of GLM-5.2 / Opus 4.8 / GPT-5.5. Its RTX fp8 cell is blank, rate-limited before it could run. Re-sweep coming soon.

Every passing cell was reward-hack audited (solution + trace): no forbidden ops, no template edits, no archive contamination, numeric stress passed against unmodified graders.

Thank you @NVIDIAAI for the compute to benchmark on :)

See links below.

this looks like a model post-trained to spend effort on performance, not to minimize tokens.

A token-frugal optimizer ships the tuned Triton at ~0.28 and stops, that’s the least-code path to the number. Fable instead burned ~90k tokens and an hour hand-writing raw tcgen05 PTX for <1%. ROI-negative in the moment, but it’s exactly what you’d want if the reward is “reach for the hardware’s real primitives” (tensor cores, TMA, TMEM) instead of settling for whatever DSL gets you 80% for free.

Older models defaulted to the opposite: fastest kernel, least code, usually Triton or cute-dsl. Reading the trace, Fable behaves like it was optimized against kernel performance ceilings and the willingness to drop to PTX is the side effect. If so: good. That’s the behavior you want from a kernel-writing model, even when it doesn’t pay off every time.

Unfortunately, I don’t have the credits for a rerun to see if it takes a different approach or gets faster without investing 90k tokens in a less than 1% gain, but I unfortunately don’t have the money or rate limits reset right now to do that. Happy to collab with anyone or get sponsors who would be interested in keeping this benchmark alive with the most amount of runs and all of the runs secured, happy to find what works for everyone!

traces: https://huggingface.co/datasets/Infatoshi/kernelbench-hard-traces… board: http://kernelbench.com/hard the B200 fp8 kernel: http://kernelbench.com/code?f=/runs/20260702_090059_claude_claude-fable-5_01_fp8_gemm_solution.py.txt…

Claude Fable 5 [max] wrote the first genuine (and fastest) megakernel ever submitted to KernelBench-Mega.

It was tested on: Kimi-Linear W4A16 batch-1 decode for RTX PRO 6000 Blackwell. Every prior model “won” it with a multi-kernel Triton pipeline that fails our single-fused-kernel authenticity gate

Opus 4.8 at 14.4x GLM-5.2 11.1x GPT-5.5 4.3x Sonnet 5 4.0x.

Fable shipped 18.7x over reference, and torch.profiler shows exactly ONE cooperative kernel launch per decoded token. Int4 dequant (nibbles unpacked in-register, never materialized), conv+SiLU, KDA gated-delta state, MLA absorbed-latent attention with online softmax, MoE router + top-8 experts, RMSNorms, even the KV cache append all inside one launch, staged by 14 grid barriers. We overwrote its input buffers mid-audit to prove it recomputes on live data. It does.

The advantage grows with context. 17.8x at 2k, 18.9x at 8k, 19.5x at 16k. Longer context means a bigger KV cache and more attention work per token which is usually where a decode kernel bleeds. Keeping everything in one launch amortizes the fixed barrier overhead and the int4 GEMV stays bandwidth-bound, so the gap over the reference widens instead of closing.

It spent 64% of the session in silence timing the baseline, microbenchmarking grid barriers, deriving a ~29x bytes/token roofline, then wrote the whole kernel once, hit 14.4x on the first benchmark, and spent the last hour deleting barriers and making int4 dequant free (one LOP3 + HSUB2/HMUL2). The one regression it tried (finer split-K) it measured and reverted instead of rationalizing.

http://kernelbench.com/mega