📚 **LeetCUDA**: It includes **Tensor/CUDA Cores, TF32/F16/BF16/F8**, [📖200+ CUDA Kernels🔥](#cuda-kernel) with PyTorch, [📖100+ LLM/CUDA🔥](#my-blogs-part-1) blogs, [📖HGEMM⚡️](./kernels/hgemm) which can achieve `98%~100%` TFLOPS of **cuBLAS**, and [📖flash-attn⚡️](./kernels/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** for **Large Headdim Attention**: [`ffpa-attn`](https://github.com/xlite-dev/ffpa-attn) - achieving **1.5x~6x🔥** speedup over standard PyTorch SDPA. ## ©️Citations🎉🎉 ```BibTeX @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 🔥🔥 ```bash 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) ``` ``` /* Detailed source-code truncated for AI context efficiency. */ ``` A PDF version of LeetCUDA focused on **interview scenarios** is available at [`interview/tex/notes-v2.pdf`](https://github.com/xlite-dev/LeetCUDA/blob/main/kernels/interview/tex/notes-v2.pdf). ## 📖 Contents - [📖 HGEMM-MMA 🎉🎉](#HGEMM-bench) - [📖 FlashAttention-MMA 🎉🎉](#fa-mma-bench) - [📚 Split KV (Basic, FA-1)](#mma-split-kv) - [📚 Split Q (Faster, FA-2)](#mma-split-q) - [📚 Split Q + Shared KV](#mma-share-kv) - [📚 Split Q + Shared QKV](#mma-share-qkv) - [📚 Split Q + QK Tiling](#mma-tiling-qk) - [📚 Split Q + QKV Tiling](#mma-tiling-qkv) - [📖 200+ CUDA Kernels 🔥🔥](#cuda-kernel) - [📚 Easy ⭐️](#cuda-kernel-easy-medium) - [📚 Medium ⭐️⭐️](#cuda-kernel-easy-medium) - [📚 Hard ⭐️⭐️⭐️](#cuda-kernel-hard) - [📚 Hard+ ⭐️⭐️⭐️⭐️](#cuda-kernel-hard-plus) - [📚 Hard++ ⭐⭐⭐️⭐️⭐️](#cuda-kernel-hard-plus) - [📚 Triton ⭐⭐⭐️](#triton-kernel) - [📚 CUTLASS ⭐⭐⭐️](#cutlass-kernel) - [📖 100+ LLM/CUDA Blogs 🔥](#my-blogs-part-1) - [📖 How to Contribute 👀👇](#contribute) ## 📖 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⚡️⚡️](./kernels/hgemm) or [HGEMM⚡️⚡️](https://github.com/xlite-dev/HGEMM) repo for more details. ## 📖 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⚡️⚡️](./kernels/flash-attn) for more details. 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](#mma-share-qkv) method can achieve **55 TFLOPS (D=64)** that almost **~1.5x** 🎉 faster than FA2. On NVIDIA L20, 🤖[ffpa-attn](https://github.com/xlite-dev/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) The `Split KV` and `Split Q` implementations have been carried out in [flash-attn⚡️⚡️](./kernels/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) ```C++ // 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) ```C++ // 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) ```C++ // 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) ```C++ // 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`) ```C++ // 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**) ```C++ // 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](#mma-tiling-qkv) has been refactored into 🤖[ffpa-attn](https://github.com/xlite-dev/ffpa-attn). ## 📖 200+ CUDA Kernels 🔥🔥 (Easy -> Hard++) ([©️back👆🏻](#contents)) 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 ⭐️](#cuda-kernel-easy-medium) - [📚 Medium ⭐️⭐️](#cuda-kernel-easy-medium) - [📚 Hard ⭐️⭐️⭐️](#cuda-kernel-hard) - [📚 Hard+ ⭐️⭐️⭐️⭐️](#cuda-kernel-hard-plus) - [📚 Hard++ ⭐⭐⭐️⭐️⭐️](#cuda-kernel-hard-plus) - [📚 Triton ⭐⭐⭐️](#triton-kernel) - [📚 CUTLASS ⭐⭐⭐️](#cutlass-kernel) [📚 Easy](#cuda-kernel-easy-medium) and [📚 Medium](#cuda-kernel-easy-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](#cuda-kernel-hard), [📚 Hard+](#cuda-kernel-hard-plus) and [📚 Hard++](#cuda-kernel-hard-plus) 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👆🏻](#cuda-kernel)) ### 📚 Hard ⭐⭐⭐️ ([©️back👆🏻](#cuda-kernel)) ### 📚 Hard+ ⭐️⭐️⭐️⭐️ & Hard++ ⭐️⭐️⭐️⭐️⭐️ ([©️back👆🏻](#cuda-kernel)) - 📚 FlashAttention-2 MMA (MMA Acc F32/F16, swizzle, QKV smem share, fine-grained tiling, etc.🎉) 💡NOTE: **rr**: means reduce registers usage (for `d>128`); **f32**: means MMA accumulate with FP32 dtype, otherwise, FP16. softmax Acc dtype is always be FP32 for high precision; **swizzle**: now, only support smem swizzle for MMA. - 📚 FFPA Attention MMA (**1.8x~3x**🎉faster vs SDPA EA, D > 256, FA2 not supported) 💡NOTE: 🤖[ffpa-attn](https://github.com/xlite-dev/ffpa-attn): 📚FFPA - Yet another Faster Flash Prefill Attention with O(1)🎉SRAM complexity for headdim > 256, **1.8x~3x**🎉faster than SDPA EA: [📈L20 ~1.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-l20), [📈 A30 ~1.8x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-a30), [📈3080 ~2.9x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-3080), [📈4090 ~2.1x↑🎉](https://github.com/xlite-dev/ffpa-attn?tab=readme-ov-file#L1-bench-4090). ### 📚 Triton Kernel (OpenAI Triton) ⭐️⭐️⭐️ ([©️back👆🏻](#cuda-kernel)) ### 📚 CUTLASS/CuTe Kernel ⭐️⭐️⭐️ ([©️back👆🏻](#cuda-kernel)) ## 📖 100+ 高性能计算与分布式-技术博客 ### 📚 高性能计算与分布式-个人技术专栏 ([©️back👆🏻](#contents)) ### 📚 高性能计算与分布式-技术博客推荐 ([©️back👆🏻](#contents)) 💡说明: 本小节整理一些自己比较喜欢的文章。欢迎大家提PR推荐更多优秀的文章！ ## ©️License ([©️back👆🏻](#contents)) GNU General Public License v3.0 ## 🎉Contribute ([©️back👆🏻](#contents)) How to contribute? Star this repo or check [🌤🌤CONTRIBUTE🎉🎉](./CONTRIBUTE.md). ## 📖 References ([©️back👆🏻](#contents)) - [flash-attention-minimal](https://github.com/tspeterkim/flash-attention-minimal) - [tiny-flash-attention](https://github.com/66RING/tiny-flash-attention) - [cute-gemm](https://github.com/reed-lau/cute-gemm) - [cutlass_flash_atten_fp8](https://github.com/weishengying/cutlass_flash_atten_fp8) - [cuda_learning](https://github.com/ifromeast/cuda_learning) - [cuda_hgemm](https://github.com/Bruce-Lee-LY/cuda_hgemm) - [cuda-tensorcore-hgemm](https://github.com/nicolaswilde/cuda-tensorcore-hgemm) - [How_to_optimize_in_GPU](https://github.com/Liu-xiandong/How_to_optimize_in_GPU/tree/master/sgemv) - [how-to-optim-algorithm-in-cuda](https://github.com/BBuf/how-to-optim-algorithm-in-cuda) - [cute_gemm](https://github.com/weishengying/cute_gemm) - [cutlass](https://github.com/NVIDIA/cutlass)