# A Case Study in CUDA Kernel Fusion: Implementing FlashAttention-2 on NVIDIA Hopper Architecture using the CUTLASS Library

arXiv: https://arxiv.org/abs/2312.11918 · HTML: https://mlsanity.com/paper/2312.11918/

## 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 (most cited first, 5 of 5)

- [Attention Is All You Need](https://mlsanity.com/paper/1706.03762/) (2017)
- [Language Models are Few-Shot Learners](https://mlsanity.com/paper/2005.14165/) (2020)
- [FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness](https://mlsanity.com/paper/2205.14135/) (2022)
- [FlashAttention-2: Faster Attention with Better Parallelism and Work Partitioning](https://mlsanity.com/paper/2307.08691/) (2023)
- [Online normalizer calculation for softmax](https://mlsanity.com/paper/1805.02867/) (2018)

## Built on by (most cited first, 5 of 5)

- [SageAttention: Accurate 8-Bit Attention for Plug-and-play Inference Acceleration](https://mlsanity.com/paper/2410.02367/) (2024)
- [vAttention: Dynamic Memory Management for Serving LLMs without PagedAttention](https://mlsanity.com/paper/2405.04437/) (2024)
- [ThunderKittens: Simple, Fast, and Adorable AI Kernels](https://mlsanity.com/paper/2410.20399/) (2024)
- [FlashAttention on a Napkin: A Diagrammatic Approach to Deep Learning IO-Awareness](https://mlsanity.com/paper/2412.03317/) (2024)
- [A Scalable Distributed Framework for Multimodal GigaVoxel Image Registration](https://mlsanity.com/paper/2509.25044/) (2025)

## 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.