TechPrufer
Get started
Together AI
Together AI logoTogether AI · Feature trail

FlashAttention-4 Algorithm and Kernel Co-Design

3 posts · Mar 2026

What this trail covers

Together AI continues to push the boundaries of AI hardware utilization and software-hardware co-design, exemplified by their work on the NVIDIA Vera Rubin NVL72 platform. This post details the integration of ThunderKittens, a GEMM kernel, with the new ISA and features of Vera Rubin, achieving over 22 PFLOPS in NVFP4 GEMMs. The team has optimized the kernel by leveraging doubled K-step instructions, increased tensor memory, oversized shared memory, B-side collectors, and early A-release mechanisms. FlashAttention-4 pairs a new algorithm with a kernel co-design tuned for NVIDIA Blackwell GPUs, removing bottlenecks and keeping tensor cores busy. It is 2.7x faster than Triton and 1.3x faster than cuDNN 9.13, enabling more intelligent capabilities at a lower cost per token for long context workloads. Together Megakernel is a hand-optimized implementation that runs an entire model in a single kernel, targeting the HBM bandwidth ceiling, achieving significant performance improvements and better unit economics. together.compile automates kernel optimization, generating an optimized kernel stack at startup with a single function call, accelerating generation and improving startup times for image and video models. The Reinforcement Learning API applies the full Together stack to RL training, optimizing rollouts and accelerating weight distribution. ThunderAgent introduces a program-aware abstraction for agentic workflows, improving throughput, reducing memory imbalance, and reclaiming resources. ATLAS-2 introduces an online training flywheel for speculative decoding, continuously updating the speculator from live traffic for improved performance and adaptation. Cache-aware prefill–decode disaggregation (CPD) separates compute-heavy prefill from latency-sensitive decoding by classifying requests by cache hit rate and routing them to dedicated tiers, improving sustainable throughput for long-context inference.

All Together AI trails
  1. Inside the Together AI kernels teamThis 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.
  2. FlashAttention-4: Algorithm and Kernel Pipelining Co-Design for Asymmetric Hardware ScalingThis post introduces FlashAttention-4, an algorithm and kernel co-design specifically for asymmetric hardware scaling on Blackwell GPUs. It addresses bottlenecks in SFU units for softmax exponential (forward pass) and shared-memory traffic (backward pass) by implementing new pipelining strategies. Key innovations include a software emulation of the exponential function using FMA units, conditional online softmax rescaling, storing intermediate results in tensor memory for the backward pass, utilizing Blackwell's 2-CTA MMA mode to reduce shared memory traffic and atomic reductions, and a new tile scheduler to handle load imbalance. The post details hardware features of Blackwell (Tensor Memory, fully asynchronous 5th gen tensor cores, 2-CTA MMA) and provides a "feeds and speeds" analysis highlighting the bottlenecks. It achieves up to 1605 TFLOPs/s (71% utilization) on B200 with BF16, outperforming cuDNN and Triton.
  3. Key research and product announcements at the AI Native ConfThis post announces FlashAttention-4, which pairs a new algorithm with kernel co-design tuned for NVIDIA Blackwell GPUs, achieving 2.7x speedup over Triton and 1.3x over cuDNN 9.13. It also details Together Megakernel, a single-kernel implementation optimized for NVIDIA H100 HBM bandwidth, delivering a 3.6x performance improvement and 7.2x better unit economics for real-time voice agents. together.compile automates kernel optimization, generating optimized kernel stacks at startup, and has shown a 41% improvement in Flux Kontext benchmarks. The Reinforcement Learning API is introduced, applying the Together stack to RL training, with faster weight distribution (seconds within a datacenter, under a minute globally). ThunderAgent is open-sourced, improving agentic workflow throughput by 1.5–3.6x and RL rollout throughput by 1.8–3.9x. ATLAS-2 introduces an online training flywheel for speculative decoding, achieving a 1.2x performance improvement on established models. Cache-aware prefill–decode disaggregation (CPD) improves sustainable QPS by 35–40% for long-context inference on NVIDIA B200 GPUs.