📚 LeetCUDA: Modern CUDA Learn Notes with PyTorch for Beginners 🐑
📚 LeetCUDA: It includes Tensor/CUDA Cores, TF32/F16/BF16/F8, 📖200+ CUDA Kernels🔥 with PyTorch, 📖100+ LLM/CUDA🔥 blogs, 📖HGEMM⚡️ which can achieve 98%~100% TFLOPS of cuBLAS, and 📖flash-attn⚡️ using Tensor Cores with pure MMA PTX. ♥️ Please consider to leave a ⭐️ Star to support me, my bro ~ ♥️
Please also check out our production-ready Kernel Library: ffpa-attn - Fast and Memory-Efficient Exact Attention (BF16/FP16/FP8/FP4) for Large Headdim, 1.5x~15x🔥🔥 speedup over standard PyTorch SDPA.
©️Citations🎉🎉
@misc{LeetCUDA@2025,
title={LeetCUDA: A Modern CUDA Learn Notes with PyTorch for Beginners},
url={https://github.com/xlite-dev/LeetCUDA.git},
note={Open-source software available at https://github.com/xlite-dev/LeetCUDA.git},
author={DefTruth and Many Others},
year={2025}
}
📖 Quick Start 🔥🔥
git clone https://github.com/xlite-dev/LeetCUDA.git && cd LeetCUDA
git submodule update --init --recursive --force && cd kernels/interview
# Install the latest CUDNN library for benchmarks (remove the old version first)
apt remove -y libcudnn9-cuda-13 libcudnn9-dev-cuda-13 libcudnn9-headers-cuda-13
apt install -y cudnn9-cuda-13 ccache # Also install ccache for faster rebuilds
# Build for target architecture (ccache accelerated when available):
./build.sh --arch sm_89 # Ada Lovelace (L20, RTX 40 series, CUDA Toolkit >= 13.2)
./build.sh --arch sm_90a # Hopper (H100/H200, CUDA Toolkit >= 13.2)
./build.sh --arch sm_120a # Blackwell (RTX 5090 / PRO 5000/6000, CUDA Toolkit >= 13.2)
./build.sh --arch all # All three architectures (sm_89, sm_90a, sm_120a)
./build.sh --clean # Remove build artifacts (*.o, *.bin, *.ptx)
# Then, run the notes_v2_sm120a.bin with bench mode (e.g., NVIDIA RTX 5090, Blackwell SM_120a)
# Baseline: cuBLAS v13.3.0.5-1 (290T); cuDNN v9.25.0.15 SDPA (222T), PyTorch v2.11 SDPA (210T)
# Speedup: Flash-Attention 2/3 -> ~1.37x (F16 Acc vs cuDNN), ~1.01x (F32 Acc vs cuDNN), ~1.07x
# (F32 Acc vs PyTorch SDPA); HGEMM w/ Pipe & SMEM & Block Swizzle -> 1.07x (F16 Acc vs cuBLAS)
./notes_v2_sm120a.bin --bench --mnk 4096,4096,4096 --bhnd 1,32,16384,128 # MMA ACC F16/F32 Acc
| Kernel | Max Err | TFLOPS/cu{BLAS,DNN} |
|----------------------------------------------------------|-----------|---------------------|
| HGEMM CuTe Swizzle (S=2, BLK_SW=0) | 0.000e+00 | 307.7/295.4 (1.04x) |
| HGEMM CuTe Swizzle (S=2, BLK_SW=1) | 0.000e+00 | 307.3/295.4 (1.04x) |
| HGEMM CuTe Swizzle (S=3, BLK_SW=0) | 0.000e+00 | 315.3/295.4 (1.07x) |
| HGEMM CuTe Swizzle (S=3, BLK_SW=1) | 0.000e+00 | 317.4/295.4 (1.07x) |
| FA2 MMA Stages (Sk=1, Pad, F16Acc) | 1.831e-04 | 220.2/222.9 (0.99x) |
| FA2 MMA Stages (Sk=2, Pad, F16Acc) | 1.831e-04 | 254.7/222.9 (1.14x) |
| FA2 MMA Stages (Sk=1, Pad, F32Acc) | 1.526e-05 | 165.6/221.6 (0.75x) |
| FA2 MMA Stages (Sk=2, Pad, F32Acc) | 1.526e-05 | 180.1/221.6 (0.81x) |
| FA2 CuTe MMA Stages (Sk=1, F32Acc) | 1.526e-05 | 199.0/221.6 (0.90x) |
| FA2 CuTe MMA Stages (Sk=2, F32Acc) | 1.526e-05 | 201.2/221.6 (0.91x) |
| FA2 TMA MMA WS (1 Consumer WG) (Sk=1, Sv=1, F16Acc) | 1.831e-04 | 262.5/222.9 (1.18x) |
| FA2 TMA MMA WS (1 Consumer WG) (Sk=2, Sv=1, F16Acc) | 1.831e-04 | 292.0/222.9 (1.31x) |
| FA2 TMA MMA WS (1 Consumer WG) (Sk=3, Sv=1, F16Acc) | 1.831e-04 | 296.9/222.9 (1.33x) |
| FA2 TMA MMA WS (1 Consumer WG) (Sk=2, Sv=1, F32Acc) | 1.526e-05 | 201.9/221.6 (0.91x) |
| FA2 TMA MMA WS (1 Consumer WG) (Sk=3, Sv=1, F32Acc) | 1.526e-05 | 202.2/221.6 (0.91x) |
| FA3 TMA MMA WS (2 Consumer WG) (Sk=1, Sv=1, F16Acc) | 9.155e-05 | 305.2/222.9 (1.37x) |
| FA3 TMA MMA WS (2 Consumer WG) (Sk=1, Sv=1, F32Acc) | 1.526e-05 | 212.0/221.6 (0.96x) |
| FA2 CuTe TMA MMA WS (1 Consumer WG) (Sk=2, Sv=1, F32Acc) | 1.526e-05 | 220.8/221.6 (1.00x) |
| FA2 CuTe TMA MMA WS (1 Consumer WG) (Sk=3, Sv=1, F32Acc) | 1.526e-05 | 222.8/221.6 (1.01x) |
# Speedup: Split-D for large headdim (e.g, D=320) ~2.20x faster than cuDNN SDPA (with F32 Acc)
./notes_v2_sm120a.bin --bench --bhnd 1,32,16384,320 # Split-D for large headdims (e.g, D=320)
| Kernel | Max Err | TFLOPS/cu{BLAS,DNN} |
|----------------------------------------------------------|-----------|---------------------|
| FA Split-D CuTe TMA MMA WS (D=320, Sk=1, Sv=1) | 1.526e-05 | 127.2/83.1 (1.53x) |
| FA Split-D CuTe TMA MMA WS (D=320, Sk=2, Sv=2) | 1.526e-05 | 182.6/83.1 (2.20x) |
A PDF version of LeetCUDA focused on interview scenarios is available at interview/tex/notes-v2.pdf.
📖 Contents
📖 HGEMM Benchmark 🎉🎉
Currently, on NVIDIA L20, RTX 4090 and RTX 3080 Laptop, compared with cuBLAS's default Tensor Cores algorithm, the HGEMM (WMMA/MMA/CuTe) in this repo (blue🔵) can achieve 98%~100% of its (orange🟠) performance. Please check toy-hgemm library⚡️⚡️ or HGEMM⚡️⚡️ repo for more details.
| 📚Feature |
📚Feature |
📚Feature |
📚Feature |
| ✔️CUDA/Tensor Cores |
✔️Loop over K |
✔️Tile Block(BMxBK) |
✔️Tile Threads(T 8x8) |
| ✔️WMMA(m16n16k16) |
✔️MMA(m16n8k16) |
✔️Pack LDST(128 bits) |
✔️SMEM Padding |
| ✔️Copy Async |
✔️Tile MMAs |
✔️Tile Warps |
✔️Multi Stages(2~4) |
| ✔️Register Double Buffers |
✔️Block Swizzle |
✔️Warp Swizzle |
✔️SMEM Swizzle(CuTe/MMA) |
| ✔️Collective Store(Shfl) |
✔️Layout NN |
✔️Layout TN |
✔️SGEMM FP32/TF32 |
📖 FA2-MMA Benchmark 🎉🎉
I have also implemented FlashAttention-2 using pure MMA PTX instructions, which supports features such as Multi-Stages, Tile MMA, Tile Warp, Shared KV SMEM, Fully Shared QKV SMEM, Prefetch Q s2r, Prefetch K/V g2s, QKV Fine-grained Tiling, Collective Store, etc. Please refer to flash-attn⚡️⚡️ for more details.
| 📚Feature |
📚Feature |
📚Feature |
📚Feature |
| ✔️Tensor Cores |
✔️Loop over N/D |
✔️Tile Block(Br, Bc) |
✔️MMA(m16n8k16) |
| ✔️Pack LDST(128 bits) |
✔️SMEM Swizzle/Padding |
✔️Copy Async |
✔️Tile MMAs |
| ✔️Tile Warps |
✔️Multi Stages(1/2) |
✔️Collective Store(Shfl) |
✔️Split KV/Q |
| ✔️Shared QKV SMEM |
✔️Prefetch Q s2r |
✔️Prefetch KV g2s |
✔️QKV Fine-grained Tiling |
Currently, for small-scale attention (B<=4, H <=48, SeqLen <= 8192, D <= 64) it can run faster than FA2/SDPA on some Devices. For example, on NVIDIA RTX 3080 Laptop, 📚 Split Q + Fully Shared QKV SMEM method can achieve 55 TFLOPS (D=64) that almost ~1.5x 🎉 faster than FA2. On NVIDIA L20, 🤖ffpa-attn method can achieve 104 TFLOPS (D=512) that almost ~1.8x 🎉 faster than SDPA (EFFICIENT ATTENTION). However, for large-scale attention, there remains a performance gap. Stay tuned for updates ~ (MMA Acc F16/F32, softmax Acc F32 vs FA2 MMA/softmax Acc F32, 👇Benchmark)
| Algorithm |
(B,H,N,D) |
RTX 3080 Laptop |
L20 |
RTX 4090 |
| FlashAttention-2 |
(1,8,8192,64) |
37 TFLOPS |
100 TFLOPS |
145 TFLOPS |
| share-qkv+stage2 |
(1,8,8192,64) |
55 TFLOPS |
99 TFLOPS |
221 TFLOPS |
| FlashAttention-2 |
(1,48,8192,64) |
37 TFLOPS |
109 TFLOPS |
163 TFLOPS |
| share-qkv+stage2 |
(1,48,8192,64) |
48 TFLOPS |
107 TFLOPS |
224 TFLOPS |
| SDPA(EFFICIENT ATTENTION) |
(1,48,8192,512) |
16 TFLOPS |
58 TFLOPS |
85 TFLOPS |
| 🤖ffpa-attn |
(1,48,8192,512) |
39 TFLOPS |
104 TFLOPS |
200 TFLOPS |
| Precision Errors vs FA2/SDPA |
/ |
max: < ~1e-3 |
min: ~0.0 |
mean: < ~1e-5 |
The Split KV and Split Q implementations have been carried out in flash-attn⚡️⚡️ for performance comparison. The Split KV method, which involves splitting all QKV across MMA (Warps), is slower than Split Q method, which splitting Q across MMA(Warps) and keep access KV for all MMA(Warps).
- 📚 Split KV (Basic, FlashAttention-1)
// Split QKV across MMA(Warps) using naive matmul MMA&Warp tiling policy.
// case: The layout of 8 MMA(2x4) [after] kWarpTileSeqLenQxkWarpTileSeqLenK(2x2) -> 32x2,32x2=64x64:
// | [64,64] | warp_KV 0 | warp_KV 1 | warp_KV 2 | warp_KV 3 |
// | warp_QP 0 |-- MMA 0,MMA 0 --|-- MMA 2,MMA 2 --|-- MMA 4,MMA 4 --|-- MMA 6,MMA 6 --|
// | warp_QP 0 |-- MMA 0,MMA 0 --|-- MMA 2,MMA 2 --|-- MMA 4,MMA 4 --|-- MMA 6,MMA 6 --|
// | warp_QP 1 |-- MMA 1,MMA 1 --|-- MMA 3,MMA 2 --|-- MMA 5,MMA 5 --|-- MMA 7,MMA 7 --|
// | warp_QP 1 |-- MMA 1,MMA 1 --|-- MMA 3,MMA 2 --|-- MMA 5,MMA 5 --|-- MMA 7,MMA 7 --|
__global__ void // Q, K, V, O -> [B, H, N, D]
flash_attn_mma_stages_split_kv_kernel(half* Q, half* K, half* V, half* O, ...);
- 📚 Split Q (Faster, FlashAttention-2)
// Split Q across MMA(Warps) and keep access KV for all MMA(Warps),
// in order to reduce the comm between warps via smem and warp shuffle.
// case: MMA = m16n8k16, Br=16x4=64, Bc=8x8=64, layout: 4 warps
// | 64x64 | warp_KV 0 |
// | warp_QP 0 | MMA 0 ... MMA 0 (x8) |
// | warp_QP 1 | MMA 1 ... MMA 1 (x8) |
// | warp_QP 2 | MMA 2 ... MMA 2 (x8) |
// | warp_QP 3 | MMA 3 ... MMA 3 (x8) |
__global__ void // Q, K, V, O -> [B, H, N, D]
flash_attn_mma_stages_split_q_kernel(half* Q, half* K, half* V, half* O, ...);
- 📚 Split Q + Shared KV SMEM (1/2 SRAM vs FA2)
// K, V shared the same shared memory, improve block occupancy.
__global__ void // Q, K, V, O -> [B, H, N, D]
flash_attn_mma_stages_split_q_shared_kv_kernel(half* Q, half* K, half* V, half* O, ...);
- 📚 Split Q + Fully Shared QKV SMEM (1/4 SRAM vs FA2)
// Q, K, V fully shared the same shared memory and prefetch Q s2r, improve block occupancy
// and reduce Q SMEM IO-Access.
__global__ void // Q, K, V, O -> [B, H, N, D]
flash_attn_mma_stages_split_q_shared_qkv_kernel(half* Q, half* K, half* V, half* O, ...);
- 📚 Split Q + QK Fine-grained Tiling (O(16xd) SRAM vs FA2 O(4xBrxd) SRAM,
Headdim -> 1024)
// Fine-grained tiling at the MMA level for Q@K^T results in a constant SRAM usage of
// 64 * kMmaAtomK for Q and K. For V, the SRAM complexity is O(kMmaAtomK * d), leading to
// an overall SRAM complexity of O(kMmaAtomK * d). Consequently, this approach allows us to
// extend D (head dimension) up to 1024.
__global__ void // Q, K, V, O -> [B, H, N, D]
flash_attn_mma_stages_split_q_tiling_qk_kernel(half* Q, half* K, half* V, half* O, ...);
- 📚 Split Q + Fully QKV Fine-grained Tiling (O(2xBrx16)~O(1) SRAM vs FA2 O(4xBrxd) SRAM)
// Fine-grained tiling at the MMA level for all Q@K^T and P@V results in a constant SRAM usage of
// Br * 16 or Bc * 16 for Q, K, V, leading to an overall SRAM complexity of O(Br * 16). Consequently,
// this approach allows us to run faster than SDPA w or w/o MMA Acc F32.
__global__ void // Q, K, V, O -> [B, H, N, D]
flash_attn_mma_stages_split_q_tiling_qkv_kernel(half* Q, half* K, half* V, half* O, ...);
💡NOTE: 📚Split Q + Fully QKV Fine-grained Tiling has been refactored into 🤖ffpa-attn.
📖 200+ CUDA Kernels 🔥🔥 (Easy -> Hard++) (©️back👆🏻)
The kernels listed here will guide you through a step-by-step progression, ranging from easy to very challenging topics. The workflow for each topic will be as follows: custom CUDA kernel implementation -> PyTorch Python bindings -> Run tests. 👉TIPS: * = Tensor Cores (WMMA, MMA, CuTe), otherwise, CUDA Cores; / = not supported; ✔️ = supported; ❔ = TODO. Contents are listed as follows:
📚 Easy and 📚 Medium sections cover operations such as element-wise, mat_trans, warp/block reduce, nms, relu, gelu, swish, layer-norm, rms-norm, online-softmax, dot-prod, embedding and basic usage for FP32, FP16, BF16 and FP8 . 📚 Hard, 📚 Hard+ and 📚 Hard++ sections delve deeper into advanced topics, primarily focusing on operations like sgemv, sgemm, hgemv, hgemm and flash-attention. These sections also provide numerous kernels implemented using Tensor Cores with pure MMA PTX.
📚 Easy ⭐️ & Medium ⭐️⭐️ (©️back👆🏻)
LeetCUDA 是一本面向初学者的开源 CUDA 学习笔记，并附带一个包含 200+ 个手写 CUDA 核函数的库（包括逐元素运算、HGEMM、Flash-Attention 等），以及 100 多篇 CUDA 和大语言模型相关的博客文章。