BlogsModalFlash Attention Kernel Analysis

Flash Attention Kernel Analysis

Flash Attention Kernel Analysis

3
posts
2025–2026

Modal's platform provides deep integration with FlashAttention kernels, enabling significant performance optimizations for LLM inference. Recent contributions focus on adapting FlashAttention-4 for decode-heavy workloads by adjusting parallelism strategies (e.g., switching to KV parallelism) and supporting irregular global memory accesses via `cp.async` and Tensor Memory Accelerator (TMA) optimizations. This includes adding support for FP8 inputs, enabling arbitrary KV page sizes, and optimizing for LLM inference, including support for models like GLM-5. The platform also offers a free endpoint for GLM-5 and sample deployment code.

2026

Making FlashAttention-4 faster for inference

6/11/2026

This post details contributions to the FlashAttention-4 kernel to improve LLM inference performance, specifically for decode-heavy workloads. Key technical contributions include: 1. Adding support for FP8 inputs, reducing memory and arithmetic bandwidth demands and enabling smaller KV caches. 2. Implementing support for arbitrary KV page sizes and optimizing performance for smaller pages by introducing a `cpasync`-based path to handle irregular memory accesses, repurposing an idle warp for this task. These changes leverage techniques like split KV parallelism and optimized memory access patterns to boost throughput.

Try GLM-5.1, the new frontier of open intelligence, on Modal | Modal Blog

2/11/2026

This post details the integration of GLM-5, a frontier open weights foundation language model, with Modal's infrastructure. It discusses the technical challenges of running large models like GLM-5, including its 700GB weight size and the need for multi-GPU deployment. The post outlines the specific parallelism strategies (tensor parallelism in MoE layers with DeepGemm kernels, data parallelism in attention layers with FlashMLA kernels) used for GLM-5 inference. It also mentions ongoing work to integrate inference-optimized paths for Flash Attention 4 and provides reproducible deployment code for running GLM-5 with SGLang on Modal. Additionally, it announces a free endpoint for GLM-5 and provides integration examples for various AI frontend frameworks (OpenCode, OpenClaw, Claude Code via LiteLLM, Vercel AI SDK).

2025

We reverse-engineered Flash Attention 4

9/26/2025

Reverse-engineered and explained the internal architecture of Flash Attention 4, focusing on its asynchronous pipeline, warp specialization, and the 'life of a tile' through the GPU memory hierarchy. Detailed the implementation of faster approximate exponentials and an efficient online softmax using CUDA Cores and a cubic polynomial approximation, contrasting it with traditional SFU usage. Explained the producer/consumer model with barriers and the mapping of pipeline steps onto warps.