Before we dive into the project log, some context: I worked on this during my internship at Modular (@clattner_llvm’s startup) under Abdul Dakkak, Hengjie Wang, and Jiexiang Liu. Because of it, I went from knowing ZERO kernel programming to writing code that outperformed

why target the matrix-multiplication kernel? arguably the most important kernel. Matmul makes up more than 80% of modern LLM (GPT, Claude…) execution our 6% improvement translates to ~5% end-to-end speedup. Millions of dollars shaved off inference

Kernel 1) Naive 🐣 pure, simple matrix multiplication. exactly like middle school performance? a BRUTAL 0.3% of SOTA this is THE baseline

Kernel 2) Tiling 🧩 repetitive global loads = bottleneck tile 64x64 into shared memory compute all ops per-tile TMA streams tiles insanely fast stash intermediates in Tensor Memory status: 8.7% SOTA

Kernel 3) Swizzling 🍭 shared memory bank conflicts were DESTROYING performance 4 banks bottlenecked us BAD but XOR operations (swizzling) just SCRAMBLED data across all 32 banks now it's parallel we hit 16.4% SOTA

Kernel 4) Packing Output and using TMA 📦 stopped writing 4B chunks like an amateur. now we pack output in shared memory with stmatrix, then TMA store for 32B bulk transfers ONE tile to global memory. FAST. 16.4% SOTA

Kernel 5) 2SM MMA & Multicasting 🤝 profiling -> global memory bound we are still bound by global memory accesses. fix 1) multicasting: SMs share tiles instead of duplicate loads fix 2) 2xSM MMA: two SMs collaborate on one 256x256x16 op doubled arithmetic intensity for FREE 20%

Kernel 6) Pipelining & Warp Specialization ⏩ still memory bound, so we rewrote the flow this is THE kernel; we stop hardware idling now a 5-stage circular buffer: while MMA computes tile N, TMA loads N+1 loader warps on one side, compute warps on the other load→load→load

Kernel 7) Double-Buffered Output 🔄 we killed idle hardware, but stores were still serial now we pipeline the writeback split output into 8 tiles of 32 cols. while writing tile N, compute N+1. NO TAIL STALLS. SMs stay busy status: 85% SOTA

Kernel 8) CLC Persistent Kernels ♾️ removed CTA launch overhead & major write bottlenecks blackwell on-chip scheduler dynamically assigns tiles to free SMs COMPLETELY eliminate launch overhead and this is the culmination status: 100.6% SOTA

Kernel 9) Thread Block Swizzling 🔀 L2 cache trick: zig-zag across N then M MAXIMIZE reuse of loaded tiles across units let auto-tuner pick optimal swizzle status: up to 106% vs cuBLAS

final results 🚀 we just blew past cuBLAS by up to 6% on Gemma 3-27B the future of AI compute? it's hardware-agnostic... it's mojo 🔥

@AliesTaha The plot here seems misleading? IIRC it’s not possible to go over 2000TFLOPs of bf16 on B200 due to power limits (on non-zero data).

@apaszke @apaszke you’re absolutely right about the 2000 TFLOPS limit. The original plot uses a simpler timing approach- just a loop with a timer, calculating performance as (2MNK/(time/n_runs)). This was a rough back-of-napkin comparison between our implementation and cuBLAS (both

@AliesTaha 🫡

@ollama a reply from ollama is craaaazy

@AliesTaha You have insane leverage if you can outperform CUDA.. you need to partner up with AMD ASAP..

@ndimasTech @LisaSu adopt mojo as AMD's official programming language 🔥

@AliesTaha Great results! But it’s not hardware agnostic. All the perf squeeze is from NVIDIA specific mojo abstractions.

@1duooo good point, and while max perf extraction does require hardware-specific optimizations, Mojo itself remains a hardware-agnostic language ; if you call a matmul function in mojo on an AMD gpu, perf will parallel/surpass ROCm

@AliesTaha Congrats on the win, but what's 'Mojo' in this case? C++ wizardry or a hint at a certain magical framework

@Reddy2399 elevator pitch: > same or better performance compared to CUDA/ROCm > pythonic syntax so ML community can actually write their own custom kernels > hardware agnostic so that your codebase isn't locked to any one company’s GPUs > open-source

@AliesTaha Impressive leap! Why choose pure Mojo over CUDA?

@VibeCodeTeddy Mojo is hardware-agnostic, so you're not stuck with NVIDIA, and can port the same code over to say an AMD GPU instantly. plus it's pythonic- if you could write python code that matched c++ perf, why would you use c++?

@AliesTaha Ultra solid

@AliesTaha this is a really cool write up!

@AliesTaha goat

@AliesTaha Awesome results and great writeup! Laconic and in-depth, thank you!

@AliesTaha Beating cuBLAS is impressive, but doing it in 170 LOC of pure Mojo without any CUDA is the real headline here

@AliesTaha Thanks for sharing, Ali! Can you link to the kernel and benchmarking code, please?

@AliesTaha @__tinygrad__ can we get a response to this?

@AliesTaha Great work. Are you using any new tricks to beat SOTA or is it purely a good foundation on Mojo?

@AliesTaha Impressive work. 106% of cuBLAS performance with 170 lines is exactly the kind of efficiency we need more of.

@AliesTaha the alpha of smarter algorithms, bravo ;)

@AliesTaha r/nextfuckinglevel

@AliesTaha insane