Blogs›Together AI›FlashAttention-4 Algorithm and Kernel Co-Design
FlashAttention-4 introduces algorithm and kernel co-design for asymmetric hardware scaling, specifically targeting Blackwell GPUs. It features new pipelining for maximum overlap of tensor cores, softmax exponential, and memory operations. The forward pass includes software emulation of the exponential function and conditional online softmax rescaling. The backward pass utilizes tensor memory to store intermediate results, reducing shared-memory traffic, and supports a 2-CTA MMA mode. New tile sc. This post details the development and impact of the Kernels Lab, highlighting their work on FlashAttention and the ThunderKittens library for optimizing NVIDIA Blackwell GPUs. It also showcases the Together Megakernel implementation for real-time voice agent workloads, achieving significant latency reductions. The team's approach emphasizes academic-industry symbiosis and customer-facing collaboration for custom kernel optimization.
6/30/2026
This post details the Aurora paper on adaptive speculative decoding, which is shipped as the ATLAS speculator in production. It highlights how research across the full stack, from frontier agents to GPU kernels, contributes to the Together platform and production workloads.
4/1/2026
This post details the development and impact of the Kernels Lab, highlighting their work on FlashAttention and the ThunderKittens library for optimizing NVIDIA Blackwell GPUs. It also showcases the Together Megakernel implementation for real-time voice agent workloads, achieving significant latency reductions. The team's approach emphasizes academic-industry symbiosis and customer-facing collaboration for custom kernel optimization.
3/5/2026
This post details FlashAttention-4, an optimization for attention kernels on asymmetric hardware like Blackwell GPUs. Key contributions include: 1. New forward and backward software pipelines for maximum overlap of tensor cores, softmax exponential, and memory operations. 2. Forward pass: software emulation of exponential using FMA units and conditional online softmax rescaling. 3. Backward pass: storing intermediates in tensor memory, using 2-CTA MMA mode to reduce shared memory traffic and atomic operations, and supporting deterministic execution. 4. New tile scheduler for causal masks and variable sequence lengths.