TechPrufer
Get started
Modal
Modal logoModal · Feature trail

Flash Attention Kernel Analysis

3 posts · Sep 2025 – Jun 2026

What this trail covers

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 Nvidia's Blackwell architecture. The platform also details the asynchronous pipeline architecture of FA4, including warp specialization and producer-consumer models for efficient execution.

All Modal trails
  1. Making FlashAttention-4 faster for inferenceThis post details contributions to the FlashAttention-4 kernel to improve LLM inference performance, specifically for decode-heavy workloads. Key changes include adding support for FP8 inputs to reduce memory and arithmetic bandwidth demands, and enabling arbitrary KV page sizes by implementing a path using `cp.async` instructions to handle irregular global memory accesses, which is more efficient for smaller page sizes than the original TMA-based approach. The post also discusses the benefits of adjusting parallelism strategies and the potential for tile-based programming models.
  2. Try GLM-5.1, the new frontier of open intelligence, on Modal | Modal BlogThis 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).
  3. Product updates: Updates to Volumes, JS and Go SDKs, and moreThis post details the reverse engineering of Nvidia's Flash Attention 4 kernel, explaining its performance improvements on Blackwell GPUs. It covers async pipelines and mathematical optimizations for exponentials and softmax, providing a technical breakdown of the kernel's workings.
  4. We reverse-engineered Flash Attention 4This post details the reverse-engineered architecture of Flash Attention 4 (FA4), focusing on its asynchronous pipeline of operations and optimizations for Nvidia's Blackwell architecture. It explains the 'life of a tile' through the GPU memory hierarchy, including the roles of shared memory, Tensor Memory, and the use of CUDA Cores for faster approximate exponentials and more efficient online softmax. The post highlights the dual nature of its asynchronous execution model compared to CPU-based async programming and the use of warp specialization for concurrency within program instances.