---
title: 'BurstAttention: Distributed Exact Long-Sequence Attention'
url: https://www.emergentmind.com/topics/burstattention
type: topic
---

# BurstAttention: Distributed Exact Long-Sequence Attention

Searching arXiv for BurstAttention and closely related attention-system papers.
BurstAttention is an efficient distributed attention framework for extremely long sequences. It is designed for exact full attention rather than approximate sparsification, and its central strategy is to partition attention along the sequence dimension across multiple GPUs while jointly optimizing global communication and local device-level memory access. In the original formulation, each device keeps its local query shard fixed, circulates key and value shards in a ring, aggregates partial results with online softmax accumulation, and executes local attention with FlashAttention-like tiling into SRAM; this combination is intended to reduce memory overhead, reduce communication, and overlap communication with computation for long-context training and inference [2403.09347].

## 1. Definition and problem setting

BurstAttention addresses the regime in which standard scaled dot-product attention becomes dominated by quadratic time and memory costs. With \(Q,K,V \in \mathbb{R}^{N \times d}\), attention is written as
\[
S = \frac{QK^\top}{\sqrt{d}}, \qquad P = \mathrm{softmax}(S), \qquad O = PV.
\]
The cost of computing \(QK^\top\) is \(O(N^2 d)\), and storing logits or probabilities incurs \(O(N^2)\) memory. In long-context settings, this makes attention the dominant bottleneck.

BurstAttention is explicitly a distributed solution to that bottleneck. Rather than modifying the mathematical definition of attention, it reorganizes the computation across a cluster. The sequence is split across \(G\) devices, so device \(i\) holds local shards
\[
Q_i, K_i, V_i \in \mathbb{R}^{\frac{N}{G}\times d}.
\]
Queries remain pinned on their home device, while key and value shards circulate through a logical ring so that each query shard eventually attends to the full sequence. The framework is presented as exact attention: it computes full softmax attention over all query-key pairs, and its reported perplexity on sampled C4 examples for LLaMA-7B is essentially unchanged across tensor parallelism, RingAttention, BurstAttention w/o LAO, and BurstAttention, with values \(9.901\), \(9.902\), \(9.904\), \(9.901\), and \(9.901\), respectively [2403.09347].

A common misconception is to treat BurstAttention as a new single-device attention operator or a sparse approximation. In its original usage, the term denotes a systems framework for distributed exact attention over very long sequences. That distinction matters because the framework’s main innovations concern communication pattern, activation storage, and local kernel organization rather than a new softmax rule or alternative similarity function.

## 2. Distributed decomposition and exact aggregation

At each ring step, device \(i\) combines its fixed local query shard \(Q_i\) with some current key-value shard \((K_j, V_j)\). The local computation is
\[
S_{i,j} = Q_i K_j^\top,\qquad
P_{i,j} = \mathrm{softmax}(S_{i,j}),\qquad
O_{i,j} = P_{i,j}V_j.
\]
These local softmaxes are not globally correct by themselves, because the true denominator must include logits against all key shards. BurstAttention resolves this by maintaining online rowwise statistics across ring steps.

For each query row, the framework keeps a running maximum \(m_i\), a running denominator \(l_i\), and a running numerator/output accumulator \(O_i\). For the current local block it computes
\[
m_{i,j} = \mathrm{rowmax}(S_{i,j}), \qquad
P_{i,j} = \exp(S_{i,j} - m_{i,j}), \qquad
l_{i,j} = \mathrm{rowsum}(P_{i,j}), \qquad
O_{i,j} = P_{i,j}V_j.
\]
These are merged with
\[
m_{\text{new}} = \max\{m_i,\; m_{i,j}\},
\]
\[
l_i \leftarrow e^{m_i - m_{\text{new}}} l_i + e^{m_{i,j} - m_{\text{new}}} l_{i,j},
\]
\[
O_i \leftarrow e^{m_i - m_{\text{new}}} O_i + e^{m_{i,j} - m_{\text{new}}} O_{i,j},
\]
followed by \(m_i \leftarrow m_{\text{new}}\). After all \(G\) shards have been processed,
\[
O_i = \mathrm{diag}(l_i)^{-1} O_i, \qquad
\mathrm{lse}_i = m_i + \log l_i.
\]

This formulation applies the usual online-softmax stability trick across distributed shards as well as within a single device. The significance is twofold. First, it avoids storing all local score or probability matrices \(S_{i,j}\) and \(P_{i,j}\). Second, it makes exact sequence-parallel attention feasible without reverting to global all-gather or all-reduce patterns that become increasingly expensive at long context lengths. In the original presentation, this inter-device mechanism is named **GAO**, or Global Attention Optimization [2403.09347].

## 3. Two-level optimization: GAO and LAO

BurstAttention is organized around two coupled optimizations. GAO governs communication and aggregation across devices. **LAO**, or Local Attention Optimization, governs memory access and compute inside each GPU. This two-level design is the defining feature of the framework.

Within a device, local \(Q_i, K_j, V_j\) blocks are further tiled along the sequence dimension so that attention kernels operate in SRAM rather than repeatedly spilling large intermediates to HBM. Each thread block reads tiles of \(Q_i, K_j, V_j\) from HBM into SRAM, computes score and probability blocks in SRAM, accumulates outputs with online softmax, and writes back only the accumulated outputs that must persist. The paper states that when BurstAttention runs on a single device, there is no need for GAO and LAO plays the same role as FlashAttention. LAO is therefore not a separate attention algorithm; it is the single-device, FlashAttention-like half of the broader distributed framework [2403.09347].

This architecture is best understood as a synthesis of two ideas that earlier systems often optimized separately. FlashAttention-style methods reduce local HBM traffic but remain single-device in their basic form. Sequence-parallel distributed methods reduce per-device sequence burden but may still store too many intermediate states or communicate too much. BurstAttention combines sequence partitioning, online softmax accumulation, and SRAM tiling so that the cluster-level and device-level optimizations reinforce one another.

A concise comparison of the communication formulas reported for major baselines is useful:

| Method | Forward communication | Backward or total communication |
|---|---:|---:|
| TP (Megatron V3) | \(\Theta(4BZNd)\) | \(\Theta(8BZNd)\) total |
| RingAttention | \(\Theta(2BZNd)\) | \(\Theta(6BZNd)\) |
| BurstAttention | \(\Theta(2BZNd)\) | \(\Theta(3BZNd + 2BZN)\) |

These expressions summarize the framework’s stated motivation: lower communication than tensor parallelism, and lower backward communication than RingAttention.

## 4. Backward pass, memory behavior, and overlap

The backward pass is where BurstAttention most clearly differentiates itself from RingAttention. Rather than storing all intermediate \(S_{i,j}\) and \(P_{i,j}\), BurstAttention stores only the final output \(O_i\) and the rowwise \(\mathrm{lse}_i\), then recomputes local logits and probabilities during backpropagation. It first forms
\[
D_i = \mathrm{rowsum}(dO_i \circ O_i),
\]
and for each ring step uses
\[
S_{j,i} = Q_j K_i^\top, \qquad
P_{j,i} = \exp(S_{j,i} - \mathrm{lse}_j),
\]
\[
dV_i = dV_i + P_{j,i}^\top dO_j,
\]
\[
dP_{j,i} = dO_j V_i^\top,
\]
\[
dS_{j,i} = P_{j,i} \circ (dP_{j,i} - D_j),
\]
\[
dK_i = dK_i + dS_{j,i}^\top Q_j, \qquad
dQ_j = dQ_j + dS_{j,i} K_i.
\]

The memory implication is direct: the framework stores compact normalization summaries rather than the full \(N \times N\) attention matrix. The communication implication is equally important. RingAttention is reported to require \(\Theta(6BZNd)\) backward communication, whereas BurstAttention reduces this to \(\Theta(3BZNd + 2BZN)\) by using online-softmax-based aggregation and recomputation [2403.09347].

The system implementation also relies on double buffering and asynchronous communication. One buffer can be consumed by local attention compute while another simultaneously receives the next shard or sends the current one onward. This is summarized in the appendix runtime model:
\[
T_{\mathrm{total}} =
\max(T_{\mathrm{attn\_f}}, t_{\mathrm{comm\_attn\_f}})
+ \max(T_{\mathrm{attn\_b}}, t_{\mathrm{comm\_attn\_b}})
+ T_{\mathrm{ffn}} + t_{\mathrm{comm\_weights}}.
\]
The use of \(\max(\text{compute}, \text{communication})\) rather than a simple sum reflects the framework’s goal of hiding communication behind computation.

These design choices explain why BurstAttention is primarily a systems method. Its benefits depend on peer-to-peer communication, asynchronous overlap, and suitable cluster bandwidth. They also explain its limitations: the method is most compelling when sequence length is so large that distributed activation memory and communication dominate runtime.

## 5. Empirical behavior and scaling

The original evaluation uses two hardware settings: a single node with **8× NVIDIA A100 GPUs** connected via **PCIe**, and a distributed multi-node setting with **4 nodes × 8 A100**, for **32× A100**, connected by **600 Gb/s RoCE**. The models are **LLaMA-2 7B** and **LLaMA-2 13B**. Inference tables cover sequence lengths from **4,096** through **262,144**, and training studies include settings up to **128K** and beyond [2403.09347].

The main headline result is that, compared with **tensor parallelism (Megatron-V3) + FlashAttention**, BurstAttention **reduces communication overhead by 40%** and achieves **1.37× speedup during training at 128K sequence length on 32× A100**. The introduction also reports about **2× speedup during training 128K sequence length on 8× A100**.

First-token latency results illustrate the long-context regime in which the method is most effective. For **LLaMA-7B**, BurstAttention reports **6.49** at **65,536**, **16.01** at **131,072**, and **49.32** at **262,144**, while **TP V3 + Flash** reports **12.25**, **28.73**, and **75.52**, and RingAttention is already **OOM** at those larger lengths. For **LLaMA-13B**, BurstAttention reports **9.92** at **65,536**, **25.91** at **131,072**, and **78.80** at **262,144**, while **TP V3 + Flash** reports **19.06**, **45.46**, and **119.03**, and RingAttention again becomes **OOM** earlier [2403.09347].

The empirical pattern is consistent across the paper’s analyses. BurstAttention scales better than tensor parallelism because it partitions along the sequence dimension directly and relies on ring communication rather than repeated collectives. It scales better than RingAttention because it avoids storing large intermediate local attention states and lowers backward communication. The ablation against **BurstAttention w/o LAO** shows that GAO alone is insufficient: LAO materially reduces latency and extends the maximum supported sequence length. The framework is therefore most accurately characterized as a joint cluster-level and memory-hierarchy-level optimization.

## 6. Extensions, related meanings, and limitations

Later work embeds BurstAttention inside a broader training system for million-token contexts. **BurstEngine** introduces BurstAttention as an optimized distributed attention mechanism with topology-aware ring communication, fine-grained communication-computation overlap, and a cheaper backward formulation than RingAttention. In that later treatment, backward communication per GPU is written as \(4Nd \to 3Nd+2N\), and the attention-only benchmark at **1M** sequence length reports **1.05× speedup over LoongTrain-USP** and **1.33× speedup over LoongTrain’s DoubleRingAttention**; the full BurstEngine system reports roughly **1.2×** speedup and supports **7B at 4M sequence on 64 GPUs** and **14B at 2M sequence on 64 GPUs** [2509.19836]. This suggests continuity rather than redefinition: BurstAttention remains an exact sequence-parallel distributed attention method, while the surrounding system adds checkpointing, fused LM head/loss, and workload balancing for masked attention.

The name should also be distinguished from several adjacent uses of “burst” terminology. **“Attention to Burstiness: Low-Rank Bilinear Prompt Tuning”** studies heavy-tailed statistics in ViT self-attention and proposes a burstiness-aware prompt reparameterization; it is not a new attention mechanism called BurstAttention [2506.22908]. **BLASST** is a FlashAttention-compatible runtime block-pruning method based on online-softmax thresholding for long-context LLM inference, with reported speedups of **1.62x** for prefill at **74.7% sparsity** and **1.48x** for decode at **73.2% sparsity**; it is a dynamic sparse execution method rather than a distributed exact framework [2512.12087]. **Block-Sparse FlashAttention** likewise performs exact-score block pruning after \(QK^\top\) score computation and reports up to **1.10x** on real-world reasoning benchmarks and up to **1.24x** for needle-in-a-haystack retrieval, again as a FlashAttention-compatible inference kernel rather than a cluster-level exact-attention framework [2512.07011]. **SystolicAttention** is a hardware/software co-design for executing FlashAttention inside a single systolic array, with **1.77x** and **4.83x** higher attention FLOPs/s utilization than **AWS NeuronCore-v2** and **Google TPUv5e**, respectively, and is relevant as accelerator-level related work rather than as BurstAttention itself [2507.11331].

The original BurstAttention also has clear scope limits. Its strongest benefits appear in the extremely long-sequence regime; for shorter contexts, the extra distributed machinery may offer less benefit. It lowers activation memory more effectively than tensor parallelism, but tensor parallelism has lower parameter memory, which is why the original work discusses combining BurstAttention with **ZeRO**. It assumes a multi-GPU cluster with peer-to-peer communication and asynchronous overlap. The original paper does not deeply explore topology-aware routing beyond the ring, and it does not present an extensive treatment of masking variants. These constraints do not diminish the method’s central contribution: BurstAttention established a concrete template for exact long-context attention in which online-softmax aggregation is lifted from a single-kernel optimization to a distributed systems principle [2403.09347].

Source: https://www.emergentmind.com/topics/burstattention