A Case Study in CUDA Kernel Fusion: Implementing FlashAttention-2 on NVIDIA Hopper Architecture using the CUTLASS Library
Not on the map. Why?
Abstract (arXiv)
We provide an optimized implementation of the forward pass of FlashAttention-2, a popular memory-aware scaled dot-product attention algorithm, as a custom fused CUDA kernel targeting NVIDIA Hopper architecture and written using the open-source CUTLASS library. In doing so, we explain the challenges and techniques involved in fusing online-softmax with back-to-back GEMM kernels, utilizing the Hopper-specific Tensor Memory Accelerator (TMA) and Warpgroup Matrix-Multiply-Accumulate (WGMMA) instructions, defining and transforming CUTLASS Layouts and Tensors, overlapping copy and GEMM operations, and choosing optimal tile sizes for the Q, K and V attention matrices while balancing the register pressure and shared memory utilization. In head-to-head benchmarks on a single H100 PCIe GPU for some common choices of hyperparameters, we observe 20-50% higher FLOPs/s over a version of FlashAttention-2 optimized for last-generation NVIDIA Ampere architecture.
Key facts. Cited by 5 papers in the map (data through 14 Oct 2026). Builds on 5 of its 19 references that are in the map. First posted on arXiv 19 Dec 2023.
Builds on
5 of 19 references.
References · 5
- 2017Attention Is All You Needcited by 60,361
- 2020Language Models are Few-Shot Learnerscited by 32,419
- 2022FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awarenesscited by 3,673
- 2023FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioningcited by 2,598
- 2018Online normalizer calculation for softmaxcited by 162
Built on by
Citing papers per year: 2024: 4; 2025: 1.
5 papers.
- 2024SageAttention: Accurate 8-Bit Attention for Plug-and-play Inference Accelerationcited by 160
- 2024vAttention: Dynamic Memory Management for Serving LLMs without PagedAttentioncited by 98
- 2024ThunderKittens: Simple, Fast, and Adorable AI Kernelscited by 36
- 2024FlashAttention on a Napkin: A Diagrammatic Approach to Deep Learning IO-Awarenesscited by 7
- 2025A Scalable Distributed Framework for Multimodal GigaVoxel Image Registrationcited by 1
- 2025A Scalable Distributed Framework for Multimodal GigaVoxel Image Registration29 Sep 2025
- 2024FlashAttention on a Napkin: A Diagrammatic Approach to Deep Learning IO-Awareness4 Dec 2024
- 2024ThunderKittens: Simple, Fast, and Adorable AI Kernels27 Oct 2024
- 2024SageAttention: Accurate 8-Bit Attention for Plug-and-play Inference Acceleration3 Oct 2024
- 2024vAttention: Dynamic Memory Management for Serving LLMs without PagedAttention7 May 2024