AM
Aarav Mehta
last week
A few sections moved fast for me, but re-reading with the figures open made everything land. Highly recommend.
Free preview available. Sign in and subscribe to unlock the full book.
Vizuara AI Labs · advanced
Kernel Engineering
From silicon to speculative decoding — GPU kernels that run modern LLMs.
Write GPU kernels from first principles — the memory hierarchy, a GEMM from naive to near-cuBLAS, FlashAttention, and frontier inference kernels.
72 capsules518 figures~30 hoursby Dr. Raj Dandekar
Read on your Kindle or e-reader
Download the EPUB and read offline — perfect for the train. Works on Kindle, Kobo, Apple Books & more.
00Start Here5 capsules
The mental models. Why kernels decide who wins, how to think in speed-of-light terms, and the exact skill map a GPU kernel engineer is hired for.
01Why kernels run the worldfree28 min02The three regimes: compute, memory, overheadBRRRfree21 min03Speed-of-light thinking & the rooflineROOFLINE🔒27 min04The kernel engineer's skill mapCAREER🔒31 min05How to use this site🔒23 min01The GPU, From Silicon Up14 capsules
A tour of the H100/B200 from the die down to the wire. Every component, what it costs, and why it exists — the vocabulary the rest of the site assumes.
06The Streaming MultiprocessorSM🔒21 min07CUDA cores & the FP32/INT pipes🔒20 min08Tensor coresTC🔒22 min09The warp scheduler & latency hiding🔒21 min10The register fileRMEM🔒26 min11Shared memory & L1SMEM🔒28 min12The L2 cache & partitionsL2🔒26 min13HBM, global memory & the packageHBM🔒28 min14GPCs, TPCs & the chip floorplanGPC🔒24 min15A100 → H100 → B200: what changed🔒26 min16The roofline model in practiceROOFLINE🔒24 min17Occupancy: the balancing act🔒23 min18Arithmetic intensityAI🔒26 min19Memory coalescing🔒22 min02The CUDA Programming Model12 capsules
From the abstract launch to the metal. Threads to grids, the compilation story, and the primitives you build every kernel out of.
20Threads, warps, blocks, grids🔒26 min21SIMT & warp divergence🔒23 min22Anatomy of a kernel launch🔒22 min23The memory spaces🔒25 min24PTX vs SASS: the compilation storyPTX🔒28 min25Compute capability & targeting🔒25 min26Shared-memory bank conflictsSMEM🔒26 min27Atomics & reductions🔒24 min28Streams, events & async🔒25 min29Your first kernel, end to end🔒24 min30GPU Puzzles, walkthrough IPRACTICE🔒24 min31GPU Puzzles, walkthrough IIPRACTICE🔒27 min03The GEMM Worklog15 capsules
The heart of the course. We rebuild matrix multiply from a 1.3%-of-cuBLAS naive kernel to a 94% warptiled monster — one optimization, one measurement, one figure at a time. Then we do it again on tensor cores.
32Kernel 1: Naive1.3%🔒25 min33Kernel 2: Global memory coalescing8.5%🔒22 min34Kernel 3: Shared-memory tiling12.8%🔒24 min35Kernel 4: 1D block-tiling36.5%🔒23 min36Kernel 5: 2D block-tiling68.7%🔒18 min37Kernel 6: Vectorized memory access78.4%🔒20 min38Kernel 7: Autotuning the tiles84.8%🔒23 min39Kernel 8: Warptiling93.7%🔒25 min40Double buffering & cp.async🔒22 min41What cuBLAS is actually doing🔒25 min42Benchmarking without lying to yourself🔒26 min43The ladder, end to endRECAP🔒24 min44Tensor cores I: the WMMA GEMMTC🔒22 min45Tensor cores II: fragments & swizzlingTC🔒26 min46Tensor cores III: to cuBLAS speedTC🔒27 min04Kernels for Inference12 capsules
Where the GEMM skills meet real LLMs. Fusion, softmax, attention, FlashAttention, the KV cache and the quantized kernels that serve tokens at scale.
47Prefill vs decode: two different machines🔒25 min48Operator fusion🔒23 min49Softmax from scratch (and online)🔒26 min50RMSNorm & LayerNorm kernels🔒26 min51Attention, the naive way🔒24 min52FlashAttention I: tiling attentionFA🔒22 min53FlashAttention II: better work partitioningFA2🔒23 min54FlashAttention III: Hopper & asyncFA3🔒23 min55The KV cache & PagedAttention🔒26 min56Quantization kernels: FP8, INT4, W4A16FP8🔒28 min57The SwiGLU kernel🔒22 min58Batched decode: the GEMV problem🔒24 min05The Frontier10 capsules
The cutting edge, as of now. Hopper's async engine, Blackwell's tensor memory and NVFP4, DeepSeek's open kernels, CUTLASS the hard way, and how to debug when it all breaks.
59Hopper's TMA: async bulk copyTMA🔒25 min60WGMMA & warp specializationWGMMA🔒27 min61Beating cuBLAS on an H100WORKLOG🔒28 min62Blackwell: tcgen05 & tensor memoryB200🔒24 min63NVFP4 & microscaling formatsNVFP4🔒21 min64DeepSeek's open kernels: FlashMLA & DeepGEMM🔒25 min65DSpark: speculative decoding as a kernel problemDSPARK🔒25 min66CUTLASS the hard wayCUTLASS🔒27 min67CuTe, the DSL landscape & Triton🔒26 min68Debugging kernels: the vLLM workflowDEBUG🔒29 min06AI × Kernels4 capsules
The newest frontier: can models write the kernels? KernelBench, test-time search, RL, and the honest picture of where AI-generated kernels win and where they still fail.
69KernelBench & measuring AI kernels🔒28 min70Monkeys & search: test-time scaling🔒27 min71The CRFM experiments🔒25 min72Kevin, RL & what's still human🔒24 minRatings & reviews
1,204 readersRate this book
DO
Daniel Okafor
2 months ago
This is the resource I wish I had when I started. Clear mental models, zero fluff.
LC
Liam Chen
3 months ago
Concise capsules I can finish over coffee, yet each one taught me something I actually use.
MB
Marcus Bell
5 months ago
Rare mix of depth and readability. The worked examples are the clearest I have seen on this subject.
JB
Jonas Bergström
6 months ago
Loved the from-first-principles approach. It rebuilt my intuition rather than just handing me formulas.