Claude Fable 5 writes first-ever verified megakernel on KernelBench-Mega for an 18.7x GPU speedup

VIEWS104.3KBOOKMARKS322LIKES1KRETWEETS52REPLIES36

it took Claude Fable 2.5 hours to write a fused megakernel which delivers a >18x speed-up over a PyTorch baseline

now please recall that: - Fable is not the full Mythos model - Anthropic can spend much more than just 2.5h and ~550k tokens on this - they probably have better harnesses

Anthropic is definitely doing some sweet autoresearch internally. Especially architecture research bros are probably so happy at Anthropic. Imagine vibe-testing a new arch / tweak some arch and wanting to test it in a semi-optimized way. Just let 10T Mythos cook for a day.

Elliot Arledge@elliotarledge

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

10h104.3K1K322

Zephyr@zephyr_z9

BRUH Just imagine how hard Anthropic is optimizing its inference by using the big boi Mythos (Fable 5 is distilled from it) The gap between Anthropic and the rest of the industry is increasing (both the capability & economic level)

Elliot Arledge@elliotarledge

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

3h50.8K34987

Lisan al Gaib@scaling01

if you understand this touch some fucking grass

Elliot Arledge@elliotarledge

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.

8h12.6K10129

Elliot Arledge@elliotarledge

DUDE!!! https://kernelbench.com/runs/20260701_172615_claude_claude-fable-5_02_kimi_linear_decode_solution.py.txt

11h673123

Elliot Arledge@elliotarledge

wasnt able to finish h100 and b200 sweeps due to rate limits (used the entirety of two 20x max coding plans (400 USD))

im open to sponsors to help keep this benchmark alive for such pricey models. hit me with a dm and we can figure something out so everyone can benefit :)

11h68814

Rafa Schwinger 🇻🇦@Rafa_Schwinger

@elliotarledge Had a similar experience but on metal

11h3843

Rafa Schwinger 🇻🇦@Rafa_Schwinger

@elliotarledge Metal retardkernels

Basically a single run of fable I barely have to do anything and it is better than like one week of Opus surpervised

For flux klein and qwen 3.6

11h801

Elliot Arledge@elliotarledge

@Rafa_Schwinger are you writing metal megakernels?

11h3652

Elliot Arledge@elliotarledge

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

11h234K760360

Archimedes@AlphaTrader00

@elliotarledge Great work Elliot!

Also, I don't think most people are mentally prepared for this

11h2052

Elliot Arledge@elliotarledge

@scaling01 i dont understand it XD

8h1042

Elliot Arledge@elliotarledge

@AlphaTrader00 all good. i post about this because it interests me and i like sharing :)

11h1681

Lisan al Gaib@scaling01

I still like the benchmarks tho

Lisan al Gaib@scaling01

if you understand this touch some fucking grass

8h1.6K30

Youssof Al Toukhi@Youssofal_

@elliotarledge Told you it’s insane

11h1465

jp@jeevanpillay

@elliotarledge find funding and drop autoresearch on it :)

6h312

No One@NoOnemklz

@scaling01 We are falling into event horizon. We are in singularity. It’s over.

10h5853

cemkaya@CEMKAYA2704

@elliotarledge @gpusteve is this what you guys make ai's do in wafer ?

10h471

Elliot Arledge@elliotarledge

@teortaxesTex very shocking to me as a kernel engineer

11h1394

AuspexLabs@AuspexLabs

Really good application of this But fused kernels opportunities are sometimes fairly clear depending on what you’re working on. Start breaking down step time - some of this falls out if you’re looking for it. If the prompt was “make an epic fused kernel” it’s cool but not amazing

6h119

steve@gpusteve

@CEMKAYA2704 @elliotarledge kernel gen w ai ye

10h82