Research · Compiler
Compiler
Building a principled ML compiler stack and pushing GPU kernels — matmul, cuBLAS, and FlashAttention from first principles.
AUG 2026 · BENCHMARKS · LLM
Outperforming vLLM and Llama.cpp on Gemma4-12B
Emmy, an AI-driven compiler, generates CUDA kernels automatically instead of by hand. On Gemma 4 12B in FP16 its kernels reach up to 1.6× over cuBLAS on the RTX 5090, achieved via TMA transport, hybrid FP16/FP32 accumulation, and FlashAttention-2 parity, shipped as a drop-in vLLM plugin.
JUL 2026 · TUTORIALS · GPU
Learning FlashAttention the Hard Way — Part 2
Part 2 takes Part 1's twisted monoid down to the metal: a generic compiler moveset walks one attention kernel from 25959µs to 206µs on an RTX 5090 — on par with FlashAttention-2 — with a measured number for every rung, the optimizations measurement refused, and the reasons FA-3/4 cannot land on consumer silicon.
JUL 2026 · TUTORIALS · ALGEBRA
Learning FlashAttention the Hard Way — Part 1
FlashAttention looks like a bespoke GPU kernel, but it is not much different from a regular parallel reduction. Safe softmax becomes an associative reduction via denominator trick, Welford's variance, and flash attention is the same concept. This part is the complete theoretical foundation, with the numerical-stability analysis and a test for when any loop is secretly associative.
MAY 2026 · TUTORIALS · GPU
Modern GPU Matmul Optimization
How to optimize a matmul kernel on a modern GPU, one optimization at a time: register tiling, vectorized loads, shared-memory staging, cp.async, TMA, warp specialization, split-K, and tensor cores. Each is introduced, then demonstrated by toggling it on a real kernel and watching the generated code diff and the latency change on an RTX 5090.
MAY 2026 · TUTORIALS · GPU
A Principled ML Compiler Stack in 5,000 Lines of Python — Part 3
Part 3 of the from-scratch ML compiler walkthrough: replacing the hand-coded heuristics from part 2 with an SP-MCTS search loop over Tile-IR rewrite parameters. The same six-IR pipeline, the same sixteen rules — just a tree walk on top picking the parameters that bench fastest.
MAY 2026 · TUTORIALS · GPU
A Principled ML Compiler Stack in 5,000 Lines of Python — Part 2
Part 2 of the from-scratch ML compiler walkthrough: how Loop IR is lowered to a GPU schedule. An overview of the Tile IR through three running examples (pointwise, reduction, matmul) where each picks up the rules they need from a stack of small rewrites.
APR 2026 · TUTORIALS · GPU
A Principled ML Compiler Stack in 5,000 Lines of Python — Part 1
ML compilers look like black boxes. They're not. I built one from scratch (tracing, fusion, scheduling, CUDA codegen) in 5,000 lines of Python. I walk a transformer's RMSNorm layer from PyTorch through decomposition, fusion, and tile-level scheduling, ending with the emitted CUDA kernel.
APR 2026 · TUTORIALS · GPU
Surfacing a 60% performance bug in cuBLAS
While benchmarking an FP32 SGEMM kernel on the RTX 5090, I found cuBLAS dispatching a tiny kernel for huge batched workloads — stuck at ~40% FMA utilization across the entire size range. The same library binary correctly escalates to 73% on the RTX PRO 6000 and 82% on the H200.