[2410.20399] ThunderKittens: Simple, Fast, and Adorable AI Kernels
Abstract:The challenge of mapping AI architectures to GPU hardware is creating a critical bottleneck in AI progress. Despite substantial efforts, hand-written custom kernels fail to meet their theoretical performance thresholds, even on well-established operations like linear attention. The diverse hardware capabilities of GPUs might suggest that we need a wide variety of techniques to achieve high performance. However, our work explores whether a small number of key abstractions can drastically simplify the process. We present ThunderKittens (TK), a framework for writing performant AI kernels while remaining easy to use and maintain. Our abstractions map to the three levels of the GPU hierarchy: (1) at the warp-level, we provide 16x16 matrix tiles as basic data structures and PyTorch-like parallel compute operations over tiles, (2) at the thread-block level, we provide a template for overlapping asynchronous operations across parallel warps, and (3) at the grid-level, we provide support to help hide the block launch and tear-down, and memory costs. We show the value of TK by providing kernels that match or outperform prior kernels for a range of AI operations. We match CuBLAS and FlashAttention-3 on GEMM and attention inference performance and outperform the strongest baselines by $10-40\%$ on attention backwards, $8\times$ on state space models, and $14\times$ on linear attention.
# link_2e7zisd3ikk.pdf ## Metadata - PDFFormatVersion=1.5 - IsLinearized=false - IsAcroFormPresent=false - IsXFAPresent=false - IsCollectionPresent=false - IsSignaturesPresent=false - CreationDate=D:20241029004451Z - Creator=LaTeX with hyperref - ModDate=D:20241029004451Z - Custom.PTEX.Fullbanner=This is pdfTeX, Version 3.141592653-2.6-1.40.25 (TeX Live 2023) kpathsea version 6.3.5 - Producer=pdfTeX-1.40.25 - Trapped=False ## Contents ### Page 1 ThunderKittens: Simple, Fast, and Adorable AI KernelsBenjamin F. Spector, Simran Arora, Aaryan Singhal, Daniel Y. Fu, and Christopher R´e Stanford
Explore this link on the map →saved by
related reading
- Inside NVIDIA GPUs: Anatomy of high performance matmul kernels - Aleksa Gordićaleksagordic.com
- GPUs Go Brrr · Hazy Researchhazyresearch.stanford.edu
- How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance: a Worklogsiboehm.com
- AMD GPUs go brrr · Hazy Researchhazyresearch.stanford.edu
- How To Scale Your Modeljax-ml.github.io
- The Best GPUs for Deep Learning in 2023 — An In-depth Analysistimdettmers.com
- Transformer Inference Arithmetic | kipply's blogkipp.ly
- Look Ma, No Bubbles! Designing a Low-Latency Megakernel for Llama-1B · Hazy Researchhazyresearch.stanford.edu
- Flash Attention from Scratch Part 1: Introlubits.ch
- How to Think About GPUs | How To Scale Your Modeljax-ml.github.io
- CVPR2023_eff_tutorial_molchanov.pdfnvlabs.github.io
- Outperforming cuBLAS on H100: a Worklogcudaforfun.substack.com