LiquidCP: Decoupling Attention Tasks from Communication and Kernel Boundaries in Context Parallelism

LiquidCP: Decoupling Attention Tasks from Communication and Kernel Boundaries in Context Parallelism Hongyu Chen1 Yikang Ruan1 Yuhai Long2 Ziyu Huang1 Yiyu Luo1 Liukun Yu1 Yangjie Zhou3,1 Zihan Liu1 Weihao Cui1 Zijun Li1 Anbang Wu1 Yu Feng1 Jing Yang4,1 Jingwen Leng1 Bingsheng He3 Minyi Guo1 Shixuan Sun1,* 1 Shanghai Jiao Tong University

2 Sun Yat-sen University

arXiv:2611.04401v1 [cs.DC] 3 Nov 2026

3 National University of Singapore

4 Guizhou University

{chenhongyu2048,yikang.ruan,huang_ziyu,yiyu-luo,altair.liu}@sjtu.edu.cn {weihao,lzjzx1122,y-feng,leng-jw,sunshixuan}@sjtu.edu.cn {anbang,guo-my}@cs.sjtu.edu.cn {yj_zhou,dcsheb}@nus.edu.sg longyh27@mail2.sysu.edu.cn yu_liukun@outlook.com jyang23@gzu.edu.cn

histories. Decode context parallelism (DCP) instead shards history KV and exchanges queries and partial attention states [2, 37]. Efficient CP requires an end-to-end view of load balance and computation–communication coordination. Placement determines attention work and token loads across GPUs, while task partitioning determines attention parallelism, per-task overhead, and reduction work. Communication and reduction dependencies further determine when tasks become executable and when their outputs become usable. Consequently, balancing aggregate workload alone does not ensure efficient execution: task boundaries and the progression of dependent work also matter. Training and full prefill. Dynamic CP adapts placement and parallelism degrees to variable-length batches [4, 28]. FCP further combines computation- and memory-aware block placement with communication planning and stage coalescing [42]. Attention and tokenwise operations have different scaling characteristics, making complete-layer performance important when assessing placement: attention work grows approximately quadratically with sequence length, whereas many other operations scale with token count. On the GitHub dataset, Megatron Hybrid CP reduces coreattention latency by 46.6% relative to a ring CP baseline using FlashAttention-3 (FA3) [32], while increasing the latency of the remaining Transformer-layer operations by 38.3%. Execution introduces another source of overhead. The conventional ring baselines evaluated here invoke attention separately for each KV communication block, repeating setup and partial-result merging across steps (Section 2.2). Decoding. Evenly sharded history KV can still yield uneven task completion within each GPU: short queries in decode-only batches limit query parallelism, while heterogeneous histories produce tasks of unequal duration. In mixed batches that combine decoding and chunked prefill requests, long prefill queries can cause FA3’s batchwide heuristic to leave decoding histories unsplit. The FA3-based DCP paths of state-of-the-art inference frameworks [31, 37] execute communication, attention, and reduction as separate dependent kernel stages, so a few long tasks can delay subsequent work. Additional KV splits expose parallelism but also add task overhead and reduction work. An effective policy must therefore consider attention together with its dependent reductions, while allowing completed portions of the workload to advance (Section 2.2). In this paper, we present LiquidCP, which coordinates attentiontask formation and partitioning with fine-grained fused execution for CP. Building on composable partial attention states [25, 41], we

Abstract Context parallelism (CP) partitions long sequences across GPUs for LLM training and inference. For variable-length training and full prefill, attention-balanced placement can leave token loads uneven, while ring execution repeats overhead at each computation and communication step. In stage-serialized decode CP (DCP), long-running tasks delay dependent work and leave substantial GPU hardware idle. We propose LiquidCP, which unifies training and inference through a common model of Transfer, Compute, and Reduction tasks with tile-level dependencies, decoupling attention tasks from communication and kernel boundaries. For training and full prefill, MegaRing incorporates BuddyRing placement to balance attention and token loads, and coalesces ready KV regions during forward attention to amortize overhead. For decoding and mixed batches, MegaDCP selects history-KV splits by jointly modeling the completion of attention and its dependent reductions. A shared persistent GPU-resident engine executes both designs. Compared with the fastest baseline per configuration on eight H100 GPUs, MegaRing achieves geometric-mean core-attention speedups of 1.23× forward and 1.18× backward, and accelerates a Transformer layer by 1.08– 1.16×. MegaDCP achieves geometric-mean core-attention speedups of 1.20× for decode-only and 1.35× for mixed batches. In vLLM serving with KV prefix cache, it achieves an average time-to-first-token speedup of 1.22×.

Keywords Context Parallelism, Attention, Long-Context LLM, Machine Learning Systems

1

Introduction

Long-context training and inference impose substantial computation and memory demands on LLM systems. Context parallelism (CP) addresses these demands by partitioning sequences across GPUs along the token dimension, distributing attention computation and the storage of activations and KV caches [25, 28]. In training and full prefill, query and KV sequences are both long. Grouped-query attention (GQA) [1] uses fewer KV heads than query heads, favoring dataflows that keep queries local and exchange KV, typically around a ring [25, 45]. In decoding and chunked prefill over cached contexts, queries can be much shorter than their KV ∗ Corresponding author.

1

Hongyu Chen et al.

Normalized throughput

1.5 1.25

1.23/1.18×

1.20/1.35×

Seq. 0 1 2

1 0.75 0.5 0.25 0

Forward Backward

A ttn AG atron ing elin gF giA pp ma3 aR g Rin Ze Ma Meg Me Lla

(a) Training / full prefill

Decode-only Mixed

3 4 5 6 7

S A ng CP A2 +R La aD SG M AG vLLM Meg L L v

(b) Decoding

Rank 0

Rank 3

Q 0 7

Q 3 4

KV 0 7

O

…

KV 3 4

Attention

Attention

Red.

Red.

0 7

O 3 4

KV ring exchange

His. Append. 0 𝑁! 1 2 3 4 5 6 7

AG

Rank 1 𝑄: 𝐻"

His. Attn.

𝑆 𝐻# 𝐻" 𝐻$ 𝐻!

Rank 3 𝑄: 𝐻!

(a) Ring-based CP

Figure 1: Core-attention throughput on eight H100 GPUs. Throughput is normalized to the best evaluated baseline in each configuration (dashed line); labels show LiquidCP’s speedups.

Q 𝐻# 𝐻" 𝐻$ 𝐻! KV 0 1

𝑆 …

𝑆

Rank 0 𝑄: 𝐻#

A2A

𝑆 𝐻 𝐻 𝐻 𝐻 # # # #

New KV Chunk Attn.

part O

Final Red.

𝑂: 𝐻#

(b) DCP

Figure 2: CP dataflows across four ranks. Numbers identify sequence or history-KV partitions; colors indicate owning ranks or query-head groups. (a) Ring CP with zigzag placement. (b) DCP with TP = DCP = 4 and 𝐻𝑘𝑣 = 1, viewed from Rank 0. 𝐻𝑟 denotes rank 𝑟 ’s query-head group; hatched 𝑆 blocks denote partial attention states. AG, A2A, and Red. denote all-gather, all-to-all, and LSE-weighted reduction.

describe execution using Transfer, Compute, and Reduction tasks linked by tile-level dependencies (Section 3). This formulation exposes the trade-off between coalescing tasks to amortize overhead and splitting tasks to expose parallelism. LiquidCP develops policies for both regimes, accounting for input readiness and downstream reduction costs. For training and full prefill, MegaRing combines structured placement with readiness-driven task formation (Section 4). Its placement algorithm, BuddyRing, jointly selects sequence CP degrees and locations within aligned, nested GPU groups. It searches for placements satisfying attention-work and token-load tolerances while prioritizing reduced fragmentation of short sequences. MegaRing executes groups of different sizes together in one launch. During forward attention, it retains fine-grained KV transfers while dynamically forming Compute tasks from ready, unassigned KV regions across communication blocks and data owners. Computation granularity thus adapts to runtime input availability within prescribed bounds, enabling early computation while amortizing per-task overhead. For decoding and mixed batches, MegaDCP retains the existing history-KV placement and selects splits using a critical-wave model (Section 5). The model predicts the completion of history attention together with its dependent reductions, accounting for the additional work introduced by splitting and the overlap permitted by the executor. To realize these policies, we develop a shared GPU-resident engine that fuses communication, attention, and reductions within persistent kernels (Section 6). Building on prior persistent and fused execution techniques [6, 26, 33, 34], the engine coordinates KV arrival, partial-state accumulation, and output completion through device-side protocols. These mechanisms support readiness-driven forward task formation and allow dependent reductions and transfers to advance as their inputs become available, while unrelated attention tasks continue. We evaluate LiquidCP on eight NVIDIA H100 GPUs against training CP implementations [4, 16, 28, 35, 45] and the FA3-based DCP paths of vLLM and SGLang [31, 37], reporting speedups over the fastest evaluated baseline per configuration (Figure 1). Across five variable-length datasets, MegaRing achieves geometric-mean core-attention speedups of 1.23× forward and 1.18× backward, and accelerates a Qwen3-30B-A3B Transformer layer [39] by 1.08– 1.16×. On Mooncake trace-derived workloads, MegaDCP achieves geometric-mean core-attention speedups of 1.20× for decode-only

and 1.35× for mixed batches. In vLLM serving with preloaded history KV, MegaDCP achieves average speedups of 1.04× in time between tokens and 1.22× in time to first token. This paper makes three contributions: • Placement and task formation for training. BuddyRing searches structured CP-group placements under token- and attention-load tolerances, while MegaRing dynamically forms forward attention tasks from ready KV regions. • Reduction-aware scheduling for DCP. MegaDCP selects history-KV splits by modeling the completion of attention and its dependent reductions together. • Fine-grained fused execution for CP. A GPU-resident engine supports runtime task formation and coordinates attention, reductions, and transfers through CP-specific readiness and completion protocols.

2 Background and Motivation 2.1 Preliminaries CP partitions attention along the token dimension, with dataflows determined by the relative query and KV lengths. A partial attention state contains a locally normalized output and the log-sum-exp (LSE) of scores over a KV subset. States over disjoint subsets can be merged into the complete output. Training and full prefill. For a sequence of length 𝑁 , query and KV lengths satisfy 𝑁𝑞 = 𝑁𝑘𝑣 = 𝑁 . Grouped-query attention (GQA) uses fewer KV heads than query heads (𝐻𝑘𝑣 < 𝐻𝑞 ), favoring KV exchange over exchanging queries and returning partial outputs. Ring CP retains 𝑁 /𝑃 queries on each of 𝑃 ranks and circulates KV shards, computing and merging partial results at each step. For causal attention, zigzag placement divides the sequence into 2𝑃 equal blocks and assigns blocks 𝑝 and 2𝑃 −1−𝑝 to rank 𝑝, balancing attention work (Figure 2(a)). In backward, 𝑑𝑄 accumulates locally, while 𝑑𝐾 and 𝑑𝑉 contributions return to their KV owners. For variable-length batches, dynamic CP [4, 15, 28, 35] adjusts persequence parallel degrees and placement to avoid over-partitioning short sequences. 2

Ring FA Megatron Hybrid CP

ProLong

ArXiv

MagiAttention Zeppelin

FreeLaw GitHub Pile-CC Dataset

(a) Performance under Workload Placement

1. Q allgather 2. Chunk attention

Physical SM ID

1.75 1.50 1.25 1.00 0.75 0.50 0.25 0 1.75 1.50 1.25 1.00 0.75 0.50 0.25 0

Normal

ProLong

ArXiv

Compute-only

One-step

FreeLaw GitHub Pile-CC Dataset

Physical SM ID

Forward (norm. time)

40 35 30 25 20 15 10 5 0

Backward (norm. time)

Attention (ms)

70 60 50 40 30 20 10 0

Other (ms)

LiquidCP: Decoupling Attention Tasks from Communication and Kernel Boundaries in Context Parallelism

(b) Execution Efficiency of Megatron

Figure 3: CP placement and execution efficiency. (a) Mean forward-plus-backward latency of core attention (top) and remaining layer operations (bottom). (b) Forward (top) and backward (bottom) latency under identical Megatron Hybrid CP assignments, normalized to One-step separately for each dataset and phase. Both panels use eight GPUs with a global 128K-token target.

3. History attention 4. History reduction

2

3

5. A2A pull 6. Final reduction 4

Tail idle 5

6

96 64 32 0 0 131 1

2

20

40

60

200

300

3

80

100

120

400

500

600

4

5

140 6

96 64 32 0 0

100

Timeline across serial stages (µs)

700

Figure 4: Streaming multiprocessor (SM) activity across serialized DCP stages. Timelines show a selected rank for decodeonly (top) and mixed (bottom) batches on eight GPUs with TP = 8 and DCP = 2. Colors indicate stage work; gray indicates tail idle. The horizontal axis concatenates each stage’s work span, from its earliest start to its latest completion.

Decoding. Under tensor parallelism (TP) of degree 𝑇 , each rank owns ℎ𝑞 = 𝐻𝑞 /𝑇 query heads. When 𝐻𝑘𝑣 < 𝑇 , head-wise partitioning replicates each KV head’s cache across 𝑇 /𝐻𝑘𝑣 ranks. Decode context parallelism (DCP), supported by vLLM [37] and SGLang [31], shards these replicated caches along the token dimension. We consider groups of 𝑃 = 𝑇 /𝐻𝑘𝑣 ranks sharing one KV head. For 𝑁𝑞 appended tokens and 𝑁𝑘𝑣 history tokens, exchanging queries and partial states moves substantially less data than exchanging history KV when 𝑁𝑞 𝑃ℎ𝑞 ≪ 𝑁𝑘𝑣 . In the all-to-all dataflow shown in Figure 2(b), each rank computes causal chunk attention using its local queries and replicated appended KV. Query all-gather enables history attention for all group query heads against each rank’s local history shard. KV splitting divides this shard into independent ranges, whose partial states are combined locally by history reduction. All-to-all returns these contributions to their query-head owners, where final reduction merges them with the local chunk result. Batches may contain only decoding requests (decode-only) or also prefill chunks appended to cached contexts (mixed).

2.2

131

1

timed region, preserving the same steps and local operations. Onestep further combines the local attention work of sequences with the same CP degree into a single call. In forward, removing transfers yields a 1.13–1.19× speedup, and merging steps yields a further 1.21–1.32× speedup (Figure 3(b)). The corresponding backward speedups are 1.15–1.25× and 1.05–1.10×. The first gap measures the aggregate impact of KV and gradients transfers under overlap. The second persists with KV already available and combines changes in invocation count, partial-result merging, and computation shapes. Insight. KV readiness constrains when computation can begin, but communication steps need not define its boundaries. This motivates exposing inputs through fine-grained transfers while coalescing ready KV regions across communication steps to amortize task and merge overhead. Execution tails in DCP.. The vLLM DCP path executes communication, attention, and reduction as sequential kernels [36]. Consequently, early-finishing SMs cannot advance to later stages until the current stage’s last task finishes. We examine this pattern using a stage-serialized variant of our executor with native FA3 and explicit inter-stage barriers. We define tail idle as an SM’s wait after its last task in a stage until that stage finishes. Its fraction is summed tail idle divided by total SM-time across the plotted stages: 40.1% for the decode-only batch and 48.1% for the mixed batch (Figure 4). The native FA3 heuristic [8, 9, 32] uses the batch’s maximum query length to bound KV splitting. In the plotted mixed batch, long prefill queries set this bound to one, leaving long-history decoding requests unsplit. Batch-wide query parallelism thus conceals longrunning decoding tasks that delay stage completion. In the decodeonly batch, splitting exposes more parallel work, but a few last wave tasks still keep other SMs waiting. Insight. KV splitting reduces work per attention task but adds partial results to merge. This motivates selecting splits by the completion of history attention and its dependent reductions, while allowing ready consumers to proceed once their own inputs are available.

Motivation

Attention and token load balance. Megatron Hybrid CP [28] initializes CP degrees from sequence lengths and places sequences by estimated attention work, without explicitly balancing per-rank token counts. Figure 3(a) separates layer latency into core attention (Attention) and remaining operations (Other). On the GitHub dataset, Megatron Hybrid CP reduces Attention latency by 46.6% relative to Ring FA, while increasing Other latency by 38.3%. Thus, attention gains can coexist with higher costs elsewhere in the layer. Attention work grows approximately quadratically with sequence length, whereas tokenwise work grows linearly with token count; balancing one does not guarantee balancing the other. This motivates a placement objective that accounts for both loads. Communication and ring-step overhead. To separate execution effects from placement, we compare three modes under identical Megatron Hybrid CP assignments. Normal retains overlapped ring communication and per-step computation. Compute-only makes KV available before timing and excludes KV/dKV transfers from the 3

Hongyu Chen et al. KV 0 1 2 3 new Q

G0 G1 G2 G3

0 1 2 3 4 5 6 7

0

O 0

7

R 7

T: Q

H0

R

H1

R

H2

R

G2

R

G3

H3

coalesce

τmin = 1, τmax = 4 tiles

O R H0 T

G1

From G1–G3

Q

T: KV KV

Qa 0

1

masked

Qb 0

1

2

Compute region

partial state

1

Transfer (T)

6

6

0

3

4

2

3

2

5

comp 0 a-1 a-0 comp 1 b-1,6

KV owner: G0

R − C < τmin: wait

X

5

4

b-4 b-2

G1 (local) G2

coalesced 2 ready tiles

G3

b-0,3

KV comm.

b-5

time 𝑡∗ (c) Execution timeline

Attn. comp.

Figure 6: Readiness-driven execution on GPU1 (G1) with fourGPU zigzag placement. Colors identify KV owners; numbers denote KV tile indices. (a) Valid KV ranges for 𝑄 𝑎 and 𝑄𝑏 ; gray regions are masked. (b) Logical KV order for 𝑄𝑏 before task assignment at 𝑡 ∗ . (c) Cooperative thread array (CTA) timelines with 𝜏min = 1 and 𝜏max = 4. Task labels identify the Q tile and processed KV tiles; arrows indicate dependencies.

Figure 5: Ring CP and DCP in the task model, viewed on GPU 0. Rows denote Q tiles (query-head groups 𝐻𝑟 in DCP), columns denote KV tiles, and colors identify data owners. Gray cells are masked; “new” denotes appended KV. Arrows, dashed regions, and R nodes denote Transfer, Compute, and Reduction, respectively.

3

5

[C, E) C=2 R=E=4 (b) Logical order of Qb at 𝑡 ∗

(b) DCP for decoding R Reduction

4

(a) KV ownership assigned ready pending

split

(a) Ring CP for training / full prefill

3

comm 0 0

ready KV regions across communication blocks in forward attention (Section 4). For short queries over evenly sharded history KV, MegaDCP chooses splits by jointly modeling history attention and its dependent reductions (Section 5).

A Common Task Model for Context-Parallel Attention

Attention Regions. An attention region 𝐴 = 𝐼 × 𝐽 consists of query tiles 𝐼 and KV tiles 𝐽 , evaluated under the attention mask. A region may contain noncontiguous tiles owned by different GPUs. For data reuse, our implementation fixes one query tile per forward/decoding region and one KV tile per backward region. For a query 𝑞𝑖 , let {𝑆𝑟 } partition its valid KV positions. Let 𝑜𝑖(𝑟 ) and lse𝑖(𝑟 ) denote the normalized attention output and LSE of scores over 𝑆𝑟 . Attention composition gives the complete output [8]:   Í (𝑟 ) 𝑜𝑖(𝑟 ) 𝑟 exp lse𝑖   . (1) 𝑜𝑖 = Í (𝑟 ) 𝑟 exp lse𝑖

4

MegaRing for Training and Full Prefill

MegaRing transfers KV in fine-grained tiles and forms forward Compute tasks by coalescing ready tiles across communication blocks and KV owners. Its placement algorithm, BuddyRing, balances token and attention loads while limiting sequence fragmentation.

4.1

Readiness-Driven Execution

Given a placement, MegaRing indexes each Q tile’s applicable KV tiles in a logical order, with local tiles first and remote tiles appended as their readiness is committed. Let 𝑅 denote the endpoint of the ready prefix and 𝐶 the endpoint of work already assigned to Compute tasks. Thus, [𝐶, 𝑅) contains ready, unassigned KV tiles. Communication CTAs extend the ready prefix, while compute CTAs claim intervals from [𝐶, 𝑅). After completing their transfers, communication CTAs join computation. When 𝑅 − 𝐶 ≥ 𝜏min , a compute CTA claims [𝐶, 𝐸), where

In backward, each region uses the complete forward outputs and LSEs, together with the output gradient 𝑑𝑂, to produce additive contributions to 𝑑𝑄, 𝑑𝐾, and 𝑑𝑉 . If the regions cover every valid query–key interaction exactly once, summing their gradient contributions yields the complete attention gradients. Primitive Operations. The model uses three logical primitives (Figure 5). Transfer makes operand tiles or partial results available on their consuming GPU. Compute evaluates one region and produces its partial results. Reduction merges output/LSE states or sums gradient contributions for the same target, either incrementally or as a separate task. These primitives may be fused within a single kernel. Dependencies track the tiles and partial states consumed by each task. A task may execute once its own inputs are available, independently of unrelated tasks. A target becomes complete only after all required contributions have been combined.

𝐸 = 𝐶 + min(𝑅 − 𝐶, 𝜏max ).

(2)

Assignment advances 𝐶 to 𝐸 and fixes the task’s interval; later arrivals are handled by subsequent tasks. Once all required KV tiles are ready, a final nonempty remainder smaller than 𝜏min may also be claimed. Each task merges its partial output and LSE into the Q tile’s running state using Equation 1. The output is complete only after all required KV contributions have been merged. Section 6.1 details readiness tracking, task assignment, and result accumulation. At 𝑡 ∗ in Figure 6, 𝑄𝑏 has 𝐶 = 2 and 𝑅 = 4. Compute CTA 1 claims [2, 4), combining KV tiles 0 and 3 from G0 and G3 into one task. Both thresholds are measured in KV tiles. 𝜏min trades earlier dispatch against amortization of Q loading and partial-state updates; 𝜏max limits per-task work to preserve parallelism when many tiles are ready. For zigzag-partitioned sequences, we use the empirical defaults   𝐻 𝜏min = , 𝜏max = 4𝐻, (3) 2 where 𝐻 = 𝐿/(2𝑔𝐵 KV ) is the number of KV tiles in a zigzag halfshard. Here, 𝐿 is the sequence length, 𝑔 its CP degree, and 𝐵 KV

Design Implications. Placement determines per-rank attention and token loads. For a fixed placement, region size trades task and merge overhead against input readiness and task parallelism. Coalescing amortizes costs, but waiting to form larger regions can delay dispatch and leave fewer tasks available for parallel execution. Splitting exposes more tasks while creating more partial states and reduction work. For long queries, MegaRing uses BuddyRing to balance attention and token loads through sequence placement, and coalesces 4

LiquidCP: Decoupling Attention Tasks from Communication and Kernel Boundaries in Context Parallelism

the number of tokens per KV tile. The thresholds are therefore sequence-specific. Backward retains one KV tile per task, accumulating its 𝑑𝐾 and 𝑑𝑉 in registers over the applicable local Q tiles, which are already available. Coalescing multiple KV tiles would require additional gradient accumulators and increase register pressure. We therefore leave backward coalescing to future work.

4.2

G0 G1 G2 G3 Part.

Seq.

4 x 6K

S0 24K

S1 24K

4 x 6K

S1 24K

2 x 12K

S2 14K

2 x 7K

S2 14K

4 x 3.5K

S3 2K

2 x 1K

S3 2K

1 x 2K

𝐷! = 18.75%, 𝐷" = 14.20% (a) Megatron Hybrid CP

Workload Placement with BuddyRing

G0 G1 G2 G3 Part. 2 x 12K

𝐷! = 9.38%, 𝐷" = 0.89% (b) BuddyRing

Figure 7: Placement of a 64K-token batch on four GPUs. Colors identify sequences; gray indicates GPUs outside each sequence’s group. “Part.” denotes CP degree × tokens per GPU, with K = 1024. 𝐷𝑇 and 𝐷𝑊 report token and attention imbalance, respectively.

Placement space. Given sequence lengths {𝐿𝑖 } and a power-oftwo GPU count 𝑁 , BuddyRing assigns sequence 𝑖 to an aligned group (𝑠𝑖 , 𝑔𝑖 ): 𝑔𝑖 ∈ {1, 2, 4, . . . , 𝑁 } consecutive GPUs starting at 𝑠𝑖 , with 𝑠𝑖 mod 𝑔𝑖 = 0. These disjoint or nested groups form a buddy hierarchy. Admissible group sizes and tile-alignment requirements are specified in Appendix A.1. Sequences remain local when 𝑔𝑖 = 1 and use zigzag placement otherwise. The groups of different sizes are executed in one MegaRing kernel launch.

with twelve for Megatron Hybrid CP, limiting fragmentation to preserve computational efficiency.

Placement objective. Each group member receives 𝐿𝑖 /𝑔𝑖 tokens and estimated attention work 𝑎𝑖 /𝑔𝑖 , where 𝑎𝑖 = 𝐿𝑖 (𝐿𝑖 + 1)/2 for causal attention. Let 𝐷𝑇 and 𝐷𝑊 denote the maximum absolute relative deviations of per-rank token and attention loads from their respective means 𝑇¯ and 𝑊¯ . BuddyRing permits deviations up to 𝜖𝑇 and 𝜖𝑊 (10% and 5% by default), so small imbalances do not force additional splitting. A placement is feasible when its violation 𝑉 = max (0, 𝐷𝑇 − 𝜖𝑇 , 𝐷𝑊 − 𝜖𝑊 )

Seq. S0 24K

5

MegaDCP for Decoding

MegaDCP targets decode-only and mixed batches containing chunked prefill. It keeps history KV evenly sharded and combines dependency-driven execution with KV splitting guided by the joint completion of attention and history reduction.

5.1

Execution Model

On each GPU, MegaDCP executes the batch’s DCP stages as a graph of Transfer, Compute, and Reduction tasks within one persistent kernel. The shared attention queue places local chunk tasks before history tasks, allowing chunk attention to overlap query transfers. A history-attention task becomes eligible when all its queries are available; a history-reduction task becomes eligible when all contributing partial attention results are ready. Completed history contributions are published for transfer to their query owners at output-block granularity. Final reduction requires the block’s local chunk and history results and all remote history contributions. Each compute CTA switches to reduction once all attention tasks are claimed, allowing ready reductions to overlap with attention on other CTAs. Section 6 describes task dispatch and completion signaling.

(4)

is zero. BuddyRing ranks placements lexicographically by 𝑉 , a splitting Í profile, the communication proxy 𝑖 𝐿𝑖 (𝑔𝑖 − 1)/𝑔𝑖 , the number of distinct groups, and finally 𝐷𝑊 and 𝐷𝑇 . The splitting profile compares, for each length class from shortest to longest, the number of split sequences followed by their summed split depths log2 𝑔𝑖 . Once both load tolerances are satisfied, this ordering prioritizes avoiding fragmentation of short sequences. Search. BuddyRing uses large sequences to establish a placement and smaller sequences to fill residual capacity. Sequence 𝑖 is structural if 𝜅𝑖 = max(𝐿𝑖 /𝑇¯, 𝑎𝑖 /𝑊¯ ) ≥ 0.5; the remaining sequences are fillers. A Pareto beam search places structural sequences in descending 𝜅𝑖 order, exploring their legal groups while retaining partial placements with different load, fragmentation, and communication trade-offs. Fillers are initially placed greedily on single GPUs, and the best completed placements undergo local refinement. If no feasible placement is found, BuddyRing progressively allows fillers to split, processing longer length classes first and trying smaller groups before larger ones. The search stops at the first feasible relaxation stage; otherwise, it returns the best-ranked placement across all stages. Appendix A.1 specifies candidate pruning, filler placement, local refinement, and relaxation. The planner runs once per batch and returns group assignments ordered by decreasing group size to MegaRing’s executor. In Figure 7, BuddyRing keeps the 2K sequence local. Megatron Hybrid CP splits it across two GPUs when expanding groups initialized with an 8K-token per-GPU target. BuddyRing’s placement meets both load tolerances, reducing 𝐷𝑇 from 18.75% to 9.38% and 𝐷𝑊 from 14.20% to 0.89%. It uses nine sequence shards, compared

5.2

Critical-Wave KV Splitting

Critical-wave model. Splitting changes both when partial results become ready and when attention CTAs release warps for reduction. MegaDCP’s critical-wave model jointly simulates attention and history reduction on a GPU. One time unit is the per-CTA time to process a history KV tile with all 𝐾 compute CTAs executing attention. With 𝑘 active attention CTAs, the modeled per-tile time is 𝜏 (𝑘) = max (𝜌, 𝑘/𝐾) .

(5)

The 𝑘/𝐾 term models shared HBM bandwidth, while 𝜌 imposes a calibrated per-tile time floor. An attention task 𝑡 carries 𝑐 attn (𝑡) = 𝑛𝑡 +𝛿 attn tile-equivalents of work, comprising 𝑛𝑡 KV tiles and a fixed per-task work overhead 𝛿 attn . The simulator advances each running attention task at rate 1/𝜏 (𝑘), updating 𝑘 whenever an attention task starts or finishes. 5

Hongyu Chen et al.

History-reduction duration is approximated by the load-independent cost 𝑐 red (𝑢) = 𝛿 red + 𝛾𝑣𝑢 𝑠𝑢 . (6)

CTA0 CTA1 CTA2 CTA3

Here, 𝑣𝑢 counts output vectors (query token–head pairs), 𝑠𝑢 counts contributing splits per vector, 𝛿 red is the per-task overhead, and 𝛾 is the cost per partial vector, in normalized time units. The simulator dispatches each candidate plan’s attention tasks in order to the first available compute CTA. Here, 𝑘 counts CTAs executing chunk or history attention. A CTA releases its warps to history reduction after completing its last attention task. Each reduction uses one warp and starts in queue order when its warp and inputs are available. The score is the latest simulated completion time across attention and history-reduction tasks; lower is better. The model omits query transfers, result exchange, and final reduction, so its score serves as a proxy for kernel latency.

S1: 176

0

S2: 2

History red.

121.5

58 2

(b) FA3-derived (3, 3, 1)

59 96 96 88 88 50

Idle 152

192 176 (a) NoSplit (1, 1, 1) 64 64 64 59

CTA0 CTA1 CTA2 CTA3

104 2

100 Time (model time units)

(c) Critical-wave (2, 2, 1) 150

Figure 8: Illustrative split plans under the critical-wave model. Tuples list per-sequence KV partition counts; (1, 1, 1) denotes NoSplit. Numbers inside attention blocks label KVtile counts; hatched blocks are history reductions, drawn on the CTAs whose warps execute them. Block lengths reflect concurrency-dependent processing rates. Gray marks idle time, dashed lines mark modeled completion, and arrows show reduction dependencies. Illustrative parameters: 𝛿 attn = 𝛿 red = 4, 𝛾 = 1, 𝜌 = 3/4, and 𝑣𝑢 = 1. computation and coordinate through readiness flags and completion counters. Appendix A.3 details the publication protocols and CTA configurations.

Task ordering. Decode-only plans with splitting use releaseaware longest-processing-time-first (LPT) ordering. Query transfers are prioritized by the attention work they unlock. History tasks are ordered by the transfer epoch of their last required query block, then by decreasing 𝑐 attn . Unsplit and mixed-batch plans retain the original FIFO order of attention tasks in the shared chunk/history queue.

6.1

Plan selection. The planner searches over per-sequence split counts, starting from NoSplit: one Compute task per query tile over its full local history. MegaDCP uses the native FA3 heuristic’s split counts as the FA3-derived reference plan and initial persequence search upper bounds. For decode requests with long histories, MegaDCP raises these bounds to at least four, preserving splitting options in mixed batches. At each iteration, the planner evaluates increasing the split count of each sequence below its limit by one. It also evaluates joint increments for sequences whose tasks tie at the attention tail, since changing one sequence alone may leave the score unchanged. The best-scoring candidate is accepted only if it improves the current score. Search stops when no candidate improves the score or no further splitting is allowed. MegaDCP selects the best-scoring plan among NoSplit, the search result, and the FA3-derived reference. A split plan is adopted only if its modeled completion time is at least 10% lower than NoSplit. Appendix A.2 specifies parameter calibration, split limits, tie-breaking, and the construction of transfer epochs. Figure 8 compares plans for local histories of 192, 176, and 2 KV tiles on four compute CTAs. Tuples give their partition counts in this order. NoSplit (1, 1, 1) leaves the long histories on two CTAs and finishes at 152 model time units. The FA3-derived plan (3, 3, 1) exposes more parallelism but creates a second wave of long tasks and more reduction work, finishing at 121.5. The selected plan (2, 2, 1) distributes the long histories across four concurrent tasks while limiting task and reduction overhead, finishing at 104.

6

S0: 192 CTA0 CTA1 CTA2 2 CTA3

MegaRing Specialization

Communication CTAs pull remote KV tiles directly from their owners over NVLink into a local buffer, so each tile is transferred once and reused by all local Q tiles and query heads. MegaRing exposes out-of-order KV arrivals through a per-sequence committed-prefix counter and a logical-to-physical KV map. Only fully written tiles enter this contiguous, monotonic prefix. For causal zigzag placement, a barrier ensures that all KV needed by front-half queries is committed before publishing KV needed only by back-half queries. Compute CTAs combine the applicable remote prefix with local KV to claim eligible intervals (Section 4.1). Forward epilogues merge partial outputs and LSE under a per-Q-tile lock; a separate mergedKV counter tracks completion. Communication CTAs join the compute scheduler after completing their transfers. Backward accumulates 𝑑𝑄 locally and stages 𝑑𝐾 and 𝑑𝑉 for their KV owners. After local completion, communication CTAs apply Tensor Memory Accelerator (TMA) reduce-add operations to the owners’ gradient accumulators and publish system-scope completion signals once these operations finish. Owners wait for all expected signals before final gradient conversion.

6.2

MegaDCP Specialization

MegaDCP dispatches host-prepared descriptors from a unified chunk/history attention queue and separate history- and finalreduction queues. After its last attention task, each compute CTA synchronizes, reuses its shared memory, and repartitions registers for reduction. Communication CTAs pull query blocks and update local readiness counters. History attention waits for its required query blocks, and attention tasks publish completion flags after writing partial outputs and LSE to local GPU memory. History reductions wait for their inputs and track completion per destination and output block. The last contributor to a remote output block notifies the query

Execution Engine Implementation

MegaRing and MegaDCP share a GPU-resident engine that implements the task primitives of Section 3 using the persistent execution structure of prior megakernels [6, 33]. The host supplies placement and task metadata; CTAs initially specialize in communication or 6

LiquidCP: Decoupling Attention Tasks from Communication and Kernel Boundaries in Context Parallelism

AllGather Llama3 AllGather

Ring FA Megatron Hybrid CP

64K

7 Evaluation 7.1 Experimental Setup Platform and workloads. All experiments run on one node with eight NVIDIA H100 SXM GPUs, each with 80 GB HBM and 450 GB/s NVLink bandwidth per direction. MegaRing uses causal attention with 32 query heads, eight KV heads, and head dimension 128 [16, 39]. We sample 20 batches per dataset and token budget (64K, 128K, and 256K) from ProLong [14], ArXiv, FreeLaw, GitHub, and Pile-CC [13], covering distributions dominated by long or short sequences (Figure 15(a)). For MegaDCP, offline scheduling with prefix caching generates 300 snapshots from the Mooncake trace [27, 30] at arrival-rate scales of 1×, 2×, and 4×. Replaying each snapshot at DCP degrees 2, 4, and 8 yields 630 decode-only and 270 mixed cases. Each decoding request contributes 16 query tokens per iteration, one target and 15 draft tokens as in speculative decoding [11, 22]; chunked prefill processes uncached suffixes in chunks of up to 4K tokens. Appendix B provides additional setup details and workload statistics.

128K 1.04

1.40 1.44

MegaRing All CP MegaRing Hybrid

256K

1.10 1.31 1.25

1.63

1.08 1.42

1.08 1.17 1.17 1.26

300 0 600 1.18

1.31

1.22 1.37 1.34

1.20

1.43

1.15 1.13

1.18 1.15

1.11

1.08 1.04

0.91

300 0

ng

Lo

Pro

Backward

Weighted TFLOPS / GPU

600 1.03 1.17

MagiAttention Zeppelin

Forward

owner, which pulls the results and records receive completion after the transfer finishes. Final reduction waits for local chunk-attention flags, the local history-reduction counter, and all remote receivecompletion flags to indicate that its enclosing output block’s inputs are complete.

C C C iv aw iv aw iv aw ub ub ub ng ng ArX reeL GitH ile-C roLo ArX reeL GitH ile-C roLo ArX reeL GitH ile-C P P P P P F F F

Figure 9: Workload-weighted forward and backward attention throughput per GPU. Hybrid annotations show speedup over the best baseline per dataset and token budget. vLLM AG+RS Eager vLLM AG+RS Graph

vLLM A2A Eager vLLM A2A Graph

DCP = 2 1.19

1.31

1.13

1.30

1.21

1.11

1.29

0 500

1.25

1.45

1.56

1.24

1.37

1.43

1.25

1.32

250

Baselines. MegaRing baselines include zigzag NCCL Ring Attention [25, 29, 45]; overlapped KV all-gather per sequence or across packed batches (Llama-3 style [16]); the placement algorithms of Megatron Hybrid CP [28] and Zeppelin [4], reimplemented with FA3/NCCL; and MagiAttention’s public variable-length API [35]. MegaRing All CP zigzag-partitions every sequence across all GPUs; MegaRing Hybrid uses BuddyRing placement. MegaDCP baselines reproduce the DCP paths of vLLM and SGLang on a shared FA3 [32] attention kernel with Q all-gather. vLLM-AG-RS [21, 36], vLLM-A2A [2, 37], and SGLang-AG-AR [31, 43] use output reduce-scatter, packed output/LSE all-to-all, and output all-reduce, respectively. SGLang communicates in FP32, whereas vLLM and MegaDCP use BF16. We evaluate eager execution and CUDA Graph replay, reporting Graph results unless noted otherwise.

0

1

2

4

1

2

4

1

2

1.34

Mixed

Weighted effective TFLOPS / GPU

50

DCP = 8 1.17

1.15

100

MegaDCP Eager

Decode-only

150

SGLang Eager SGLang Graph

DCP = 4

4

Arrival-rate scale

Figure 10: Workload-weighted attention throughput per GPU on decode-only (top) and mixed (bottom) batches. MegaDCP uses eager execution; baselines are evaluated with both eager execution and CUDA Graph replay. Annotations compare MegaDCP with the fastest CUDA Graph baseline.

7.2

Core Attention Performance

MegaRing. MegaRing Hybrid outperforms the best baseline in all 15 forward configurations and 14 of 15 backward configurations, with geometric-mean speedups of 1.23× and 1.18×, respectively (Figure 9). On ProLong and ArXiv, All CP already matches or exceeds the best baseline, and BuddyRing adds at most 7%. Placement matters more on FreeLaw, GitHub, and Pile-CC, where splitting short sequences across all eight GPUs leaves too little work per task. All CP falls to 0.17× the best baseline on Pile-CC, whereas Hybrid reaches 1.63× in forward at 64K. Gains diminish at larger token budgets as baselines better amortize per-step costs. Backward benefits less on these short-sequence-heavy datasets because fixed KV tiles prevent coalescing of small ready regions. On Pile-CC backward at 256K, Hybrid’s latency is 8.8% higher than MagiAttention’s.

Measurement. Core-attention timings include CP communication and result reduction and use the slowest rank’s execution time. Throughput is effective attention FLOPs per second per GPU; workload-weighted throughput divides total FLOPs by total time and GPU count. Case latency is the mean over 40 timed iterations for MegaRing and the median over 60 for MegaDCP. Appendix B details the warm-up and measurement procedures. For variablelength workloads, we select the best communication-CTA count from {4, 8, 12, 16} per configuration. Core-attention timings exclude host-side planning, which runs once per batch on the CPU and is reused by every layer. MegaDCP’s planner, including descriptor packing, takes under 3 ms at p95 per decoding step and is included in the end-to-end serving results; BuddyRing depends only on sequence lengths, so it can run in the data loader ahead of each training step.

MegaDCP.. MegaDCP achieves geometric-mean speedups of 1.20× on decode-only batches and 1.35× on mixed batches over the best CUDA Graph baseline (Figure 10). Decode-only speedups increase from 1.11–1.19× at DCP degree 2 to 1.29–1.31× at degree 8: shorter local histories increase the relative cost of transfers and reductions, making inter-task overlap more valuable. Mixed-batch speedups grow with arrival rate, reaching 1.56× at degree 2 and 7

Hongyu Chen et al. MegaDCP naive

64K

128K

(a) Execution ablation

256K

1.02×

1.11× 1.05×

1.12× 1.08×

1.02×

300 1.20×

1.00×

1.03×

ArXiv ArXiv FreeLaw FreeLaw Forward Backward Forward Backward

1.26×

1.29×

150 2

4

8

2

4

8

2

Arrival 2×

4

0

8

2

Arrival 4×

4

8

2

Arrival 1×

4

8

2

Arrival 2×

4

8

Arrival 4×

DCP size

Figure 12: Mean per-case median kernel latency on decodeonly (left) and mixed (right) batches. Annotations give the ratio of mean latencies, MegaDCP naive over MegaDCP.

(b) Placement comparison

Mean layer latency (ms)

AllGather Llama3 AllGather

a 4× rate. Section 7.4 examines the contributions of planning and cross-stage execution. Uniform-length MegaRing and uniform-history MegaDCP results appear in Appendices B.4 and B.5.

Ring FA Megatron Hybrid CP OtherBwd

MagiAttention Zeppelin

OtherFwd

MegaRing All CP MegaRing Hybrid

CoreBwd

CoreFwd

100 1.16 1.08

1.12

1.16

50

0

L Pro

on

g

iv ArX

e Fre

La

w

Git

b Hu

1.16

C e- C Pil

Figure 13: Transformer-layer latency averaged over 20 cases per dataset, decomposed into forward/backward and Core Attention/Other operations. Each MegaRing variant uses its best communication-CTA count per dataset. Annotations above the fastest non-MegaRing baseline show Hybrid’s speedup. Zeppelin denotes our placement-derived baseline without linear-module remapping.

Sources of MegaRing Gains

Execution mechanisms. Starting from ring CP with external reductions (L1), L2–L5 cumulatively add fused reductions, persistent execution with a fixed queue, CTA role switching, and readinessdriven coalescing, while retaining All-CP placement (Figure 11(a)). Their incremental gains are 1.25–1.31×, 6.1–12.9%, 1.2–2.8%, and 4.0–7.5%, respectively. Fusing reductions yields a substantial improvement; coalescing further amortizes Q loads and partial-result updates across larger tasks. L6 adds BuddyRing placement, improving mean throughput by another 3–8%.

p50 per-request mean TBT (ms)

MegaDCP MegaDCP naive vLLM A2A vLLM AG+RS p50 request-mean TBT (left axis) p50 TTFT (right axis) Batch size = 64 Batch size = 128

Placement. We compare BuddyRing with Megatron-derived and Zeppelin-derived placements using the same MegaRing kernel and sampled batches at 128K tokens. Across all five datasets, BuddyRing achieves geometric-mean forward/backward speedups of 1.11/1.13× over Megatron-derived placement and 1.05/1.07× over Zeppelinderived placement, with gains on every dataset. Figure 11(b) shows ArXiv and FreeLaw; the larger gains on FreeLaw support balancing token and attention loads while limiting short-sequence partitioning. BuddyRing meets the 10% token-load and 5% attention-load tolerances in all 100 batches. On FreeLaw, mean per-batch (𝐷𝑇 , 𝐷𝑊 ) values are (6.4%, 4.3%) for BuddyRing, (44.8%, 33.3%) for Megatronderived, and (18.7%, 6.8%) for Zeppelin-derived placement. Planner details are in Appendix B.

7.4

450

1.01×

1.01×

30

Arrival 1×

Figure 11: MegaRing ablations. (a) Mean forward throughput on ArXiv at 64K, 128K, and 256K. (b) Workload-weighted forward/backward throughput on ArXiv and FreeLaw at 128K, comparing placements on the same MegaRing kernel with the baselines’ FA3/NCCL execution. MR denotes MegaRing.

7.3

1.02× 1.02×

60

0 0

1.11×

1.00×

90

20

10 5

15

2 1 0.5

10

1

2 Arrival scale

4

0.2 0.1

20

10 5

15

2 1 0.5

10

1

2 Arrival scale

4

0.2 0.1

p50 TTFT (s, log scale)

0

300

600

1.02×

p50 per-request mean TBT (ms)

300

Mixed batch

120

600

TFLOPS / GPU

TFLOPS / GPU

600

MegaDCP

Decode-only

Zeppelin MR + Zeppelin MR Hybrid

p50 TTFT (s, log scale)

Megatron MR + Megatron MR All CP

L4: + role switching L5: + coalescing L6: + BuddyRing

P50 latency (µs)

L1: Ring CP L2: + fused reduce L3: + persistent

Figure 14: vLLM serving performance across arrival-rate scales and maximum batch sizes. Solid curves show median per-request mean TBT; dashed curves show median TTFT. All backends use CUDA Graphs. prefill chunks. Shortening these tasks gives the largest gain at DCP degree 2 and a 4× arrival rate, where local histories are longest. Cross-stage execution. At DCP degree 4 and a 4× arrival rate, MegaDCP reduces mean attention-pipeline latency from 128 to 107 𝜇s on decode-only batches and from 570 to 394 𝜇s on mixed batches relative to vLLM-A2A Graph. Stage profiling shows earlier query-transfer and result-exchange completion, consistent with processing query and output blocks as their dependencies become ready. Appendix B.3 provides the timelines and profiling methodology.

Sources of MegaDCP Gains

Split planning and ordering. We compare MegaDCP with FA3derived splits and FIFO ordering on the same engine (MegaDCP naive). Critical-wave planning and release-aware ordering provide 1.00–1.03× speedups on decode-only batches. For mixed batches, both variants retain FIFO ordering, and critical-wave split planning provides 1.02–1.29× speedups (Figure 12). For decode-only batches, both variants already split histories and overlap attention with reduction, leaving little additional benefit. For mixed batches, the heuristic leaves long-history decoding requests unsplit alongside

7.5

End-to-End Performance

We integrate both designs into workloads using Qwen3-30BA3B [39]; Appendix B.6 gives the setup and further analysis. For 8

LiquidCP: Decoupling Attention Tasks from Communication and Kernel Boundaries in Context Parallelism

Transformer-layer training at a 128K token budget, MegaRing Hybrid achieves 1.08–1.16× forward-plus-backward speedups across the five datasets (Figure 13). Faster attention and balanced token loads preserve the gains across the remaining layer operations. We evaluate vLLM serving with preloaded history KV and singletoken decoding on a Mooncake trace segment. Relative to the best vLLM baseline per configuration and metric, MegaDCP achieves 0.99–1.10× speedups in median per-request mean time between tokens (TBT) and 1.02–1.37× in median time to first token (TTFT) (Figure 14). TBT gains are limited by unchanged non-attention computation; TTFT benefits more from the attention gains on mixed batches.

8

MegaRing improves core-attention throughput on variable-length training batches by 1.23× (forward) and 1.18× (backward), and MegaDCP by 1.20× on decode-only and 1.35× on mixed batches, in geometric mean over the best baseline per configuration.

References [1] Joshua Ainslie, James Lee-Thorp, Michiel de Jong, Yury Zemlyanskiy, Federico Lebrón, and Sumit Sanghai. 2023. GQA: Training Generalized Multi-Query Transformer Models from Multi-Head Checkpoints. In Proceedings of the 2023 Conference on Empirical Methods in Natural Language Processing (EMNLP). https: //arxiv.org/abs/2305.13245 [2] Nidhi Bhatia, Ankit More, Ritika Borkar, Tiyasa Mitra, Ramon Matas, Ritchie Zhao, Maximilian Golub, Dheevatsa Mudigere, Brian Pharris, and Bita Darvish Rouhani. 2025. Helix Parallelism: Rethinking Sharding Strategies for Interactive Multi-Million-Token LLM Decoding. arXiv:2507.07120 https://arxiv.org/abs/25 07.07120 arXiv:2507.07120. [3] William Brandon, Aniruddha Nrusimha, Kevin Qian, Zachary Ankner, Tian Jin, Zhiye Song, and Jonathan Ragan-Kelley. 2023. Striped Attention: Faster Ring Attention for Causal Transformers. arXiv:2311.09431 https://arxiv.org/abs/2311 .09431 arXiv:2311.09431. [4] Chang Chen, Tiancheng Chen, Jiangfei Duan, Qianchao Zhu, Zerui Wang, Qinghao Hu, Peng Sun, Xiuhong Li, Chao Yang, and Torsten Hoefler. 2026. Zeppelin: Balancing Variable-length Workloads in Data Parallel Large Model Training. In Proceedings of the 21st European Conference on Computer Systems. ACM, 1879– 1893. doi:10.1145/3767295.3769369 [5] Jiefei Chen, Binbin Lin, Jinming Ma, et al. 2026. NanoCP: Request-Level Dynamic Context Parallelism for Data-Expert Parallel Decoding. arXiv:2605.21100 https: //arxiv.org/abs/2605.21100 arXiv:2605.21100. [6] Xinhao Cheng, Zhihao Zhang, Yu Zhou, et al. 2026. MPK: A Compiler and Runtime for Mega-Kernelizing Tensor Programs. In 20th USENIX Symposium on Operating Systems Design and Implementation (OSDI 26). USENIX Association, 1909–1926. https://www.usenix.org/conference/osdi26/presentation/cheng [7] Tri Dao, Daniel Y. Fu, Stefano Ermon, Atri Rudra, and Christopher Ré. 2022. FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness. arXiv:2205.14135 https://arxiv.org/abs/2205.14135 arXiv:2205.14135. [8] Tri Dao, Daniel Haziza, Francisco Massa, and Grigory Sizov. 2023. Flash-Decoding for long-context inference. Stanford CRFM technical blog. https://crfm.stanfor d.edu/2023/10/12/flashdecoding.html [9] Dao AI Lab. 2025. FlashAttention Hopper Split Heuristics. Software documentation. https://github.com/Dao-AILab/flash-attention/blob/v2.8.3/hopper/heurist ics.h Version 2.8.3; num_splits_heuristic; accessed 2026-09-25. [10] Kalyanmoy Deb, Amrit Pratap, Sameer Agarwal, and T. Meyarivan. 2002. A Fast and Elitist Multiobjective Genetic Algorithm: NSGA-II. IEEE Transactions on Evolutionary Computation 6, 2 (2002), 182–197. doi:10.1109/4235.996017 [11] DeepSeek-AI. 2024. DeepSeek-V3 Technical Report. arXiv:2412.19437 https: //arxiv.org/abs/2412.19437 arXiv:2412.19437. [12] Jiarui Fang and Shangchun Zhao. 2024. USP: A Unified Sequence Parallelism Approach for Long Context Generative AI. arXiv:2405.07719 https://arxiv.org/ abs/2405.07719 arXiv:2405.07719. [13] Leo Gao, Stella Biderman, Sid Black, et al. 2020. The Pile: An 800GB Dataset of Diverse Text for Language Modeling. arXiv:2101.00027 https://arxiv.org/abs/21 01.00027 arXiv:2101.00027. [14] Tianyu Gao, Alexander Wettig, Howard Yen, and Danqi Chen. 2025. How to Train Long-Context Language Models (Effectively). In Proceedings of the 63rd Annual Meeting of the Association for Computational Linguistics (Volume 1: Long Papers). Association for Computational Linguistics, 7376–7399. doi:10.18653/v1/2025.acllong.366 [15] Hao Ge, Junda Feng, Qi Huang, Fangcheng Fu, Xiaonan Nie, Lei Zuo, Haibin Lin, Bin Cui, and Xin Liu. 2025. ByteScale: Communication-Efficient Scaling of LLM Training with a 2048K Context Length on 16384 GPUs. In Proceedings of the ACM SIGCOMM 2025 Conference. ACM, 963–978. doi:10.1145/3718958.3754352 [16] Aaron Grattafiori, Abhimanyu Dubey, Abhinav Jauhri, Abhinav Pandey, Abhishek Kadian, Ahmad Al-Dahle, Aiesha Letman, Akhil Mathur, Alan Schelten, Alex Vaughan, et al. 2024. The Llama 3 Herd of Models. arXiv:2407.21783 https://arxiv.org/abs/2407.21783 arXiv:2407.21783. [17] Diandian Gu, Peng Sun, Qinghao Hu, et al. 2024. LoongTrain: Efficient Training of Long-Sequence LLMs with Head-Context Parallelism. arXiv:2406.18485 https: //arxiv.org/abs/2406.18485 arXiv:2406.18485. [18] Sam Ade Jacobs, Masahiro Tanaka, Chengming Zhang, Minjia Zhang, Shuaiwen Leon Song, Samyam Rajbhandari, and Yuxiong He. 2023. DeepSpeed Ulysses: System Optimizations for Enabling Training of Extreme Long Sequence Transformer Models. arXiv:2309.14509 https://arxiv.org/abs/2309.14509 arXiv:2309.14509.

Related Work

Context parallelism and placement. CP uses ring KV exchange, KV all-gather, or head redistribution [16, 18, 25], with hybrid dataflows [12, 17] and improvements in overlap and causal balance [3, 23]. Serving systems adapt parallelism or shard history KV [2, 5, 31, 37, 38]. Variable-length training uses adaptive partitioning, placement, and block mapping [15, 19, 24, 28]; Zeppelin [4] remaps between attention and linear-module layouts. FCP [42] balances computation and memory through fixed-size block placement and flexible peer-to-peer communication. BuddyRing jointly selects sequence CP degrees and aligned power-of-two groups, seeking a single layout that meets token- and attention-load tolerances while limiting short-sequence fragmentation. Attention task granularity. Single-GPU attention uses tiling, KV splitting, and composable partial states [7–9, 32, 41]. Across GPUs, MagiAttention [35] combines balanced dispatch with multi-stage overlap, while UltraAttn [40] schedules hierarchically tiled attention. FCP [42] coalesces planned communication and computation sub-stages, separating scheduling granularity from stage-execution granularity. MegaRing retains fine-grained KV transfers and dynamically forms forward Compute tasks from ready, unassigned KV regions across communication blocks and owners. MegaDCP selects history-KV splits by jointly modeling attention and dependent reductions, and orders tasks to advance downstream work. GPU-resident execution. Megakernels provide persistent execution and dependency management [6, 20, 33]. ParallelKittens [34] and mKernel [26] fuse communication and computation, including Ring Attention; UniEP [44] targets expert-parallel mixture-ofexperts (MoE) training. LiquidCP specializes these mechanisms for out-of-order KV arrival, partial-state accumulation, and outputblock completion, supporting its task policies and fine-grained progression through attention, reductions, and transfers.

9

Conclusion

This paper presented LiquidCP, which decouples attention tasks from communication and kernel boundaries in context parallelism. Its common task model of Transfer, Compute, and Reduction primitives places training/full prefill and decoding at opposite ends of one design space: for long queries, MegaRing coalesces ready KV regions and BuddyRing balances attention and token loads, while for short queries, MegaDCP splits history KV by jointly modeling attention and its dependent reductions. On eight H100 GPUs, 9

Hongyu Chen et al.

[19] Chenyu Jiang, Zhenkun Cai, Ye Tian, Zhen Jia, Yida Wang, and Chuan Wu. 2025. DCP: Addressing Input Dynamism In Long-Context Training via Dynamic Context Parallelism. In Proceedings of the ACM SIGOPS 31st Symposium on Operating Systems Principles. ACM, 221–236. doi:10.1145/3731569.3764849 [20] Hongyi Jin, Bohan Hou, Guanjie Wang, et al. 2026. Event Tensor: A Unified Abstraction for Compiling Dynamic Megakernel. In Proceedings of Machine Learning and Systems, Vol. 8. MLSys, 1917–1933. https://proceedings.mlsys.org/ paper_files/paper/2026/hash/53d3f45797970d323bd8a0d379c525aa-AbstractConference.html [21] Woosuk Kwon, Zhuohan Li, Siyuan Zhuang, Ying Sheng, Lianmin Zheng, Cody Hao Yu, Joseph Gonzalez, Hao Zhang, and Ion Stoica. 2023. Efficient Memory Management for Large Language Model Serving with PagedAttention. In Proceedings of the 29th Symposium on Operating Systems Principles. ACM, 611–626. doi:10.1145/3600006.3613165 [22] Yaniv Leviathan, Matan Kalman, and Yossi Matias. 2023. Fast Inference from Transformers via Speculative Decoding. In Proceedings of the 40th International Conference on Machine Learning (Proceedings of Machine Learning Research, Vol. 202). PMLR, 19274–19286. https://proceedings.mlr.press/v202/levia than23a.html [23] Dacheng Li, Rulin Shao, Anze Xie, Eric P. Xing, Xuezhe Ma, Ion Stoica, Joseph E. Gonzalez, and Hao Zhang. 2023. DISTFLASHATTN: Distributed Memoryefficient Attention for Long-context LLMs Training. arXiv:2310.03294 https: //arxiv.org/abs/2310.03294 arXiv:2310.03294. [24] Haoyang Li, Fangcheng Fu, Sheng Lin, et al. 2025. Hydraulis: Balancing Large Transformer Model Training via Co-designing Parallel Strategies and Data Assignment. Proceedings of the ACM on Management of Data 3, 6 (2025), 1–30. doi:10.1145/3769802 [25] Hao Liu, Matei Zaharia, and Pieter Abbeel. 2023. Ring Attention with Blockwise Transformers for Near-Infinite Context. arXiv:2310.01889 https://arxiv.org/abs/ 2310.01889 arXiv:2310.01889. [26] Ziming Mao, Yihan Zhang, Shawn Wei Chew, Shuang Ma, Costin Raiciu, Yang Zhou, Scott Shenker, and Ion Stoica. 2026. mKernel: Fast Multi-GPU, MultiNode Fused Kernels. arXiv:2609.13585 https://arxiv.org/abs/2609.13585 arXiv:2609.13585. [27] Mooncake. 2026. Mooncake FAST25 Trace Release. Software documentation. https://github.com/kvcache-ai/Mooncake/tree/main/FAST25-release Official trace release and schema documentation; accessed 2026-09-25. [28] NVIDIA. 2026. Megatron Core: Hybrid Context Parallel Scheduling and BalancedCPScheduler. Software documentation. https://docs.nvidia.com/megatroncore/developer-guide/latest/apidocs/core/core.pipeline_parallel.hybrid_cp_sch edule.html Official API documentation; accessed 2026-09-25. [29] NVIDIA. 2026. NVIDIA Collective Communication Library (NCCL) Documentation. Software documentation. https://docs.nvidia.com/deeplearning/nccl/userguide/docs/index.html Official documentation; accessed 2026-09-25. [30] Ruoyu Qin, Zheming Li, Weiran He, Jialei Cui, Feng Ren, Mingxing Zhang, Yongwei Wu, Weimin Zheng, and Xinran Xu. 2025. Mooncake: Trading More Storage for Less Computation — A KVCache-centric Architecture for Serving LLM Chatbot. In 23rd USENIX Conference on File and Storage Technologies (FAST 25). USENIX Association, 155–170. https://www.usenix.org/conference/fast25/p resentation/qin [31] SGLang. 2026. SGLang: Serving Framework Source Code. Software documentation. https://github.com/sgl-project/sglang Source repository; accessed 2026-09-25. [32] Jay Shah, Ganesh Bikshandi, Ying Zhang, Vijay Thakkar, Pradeep Ramani, and Tri Dao. 2024. FlashAttention-3: Fast and Accurate Attention with Asynchrony and Low-precision. In Advances in Neural Information Processing Systems, Vol. 37. Neural Information Processing Systems Foundation, Inc. (NeurIPS), 68658–68685. doi:10.52202/079017-2193 [33] Benjamin Spector, Jordan Juravsky, Stuart Sul, Owen Dugan, Dylan Lim, Dan Fu, Simran Arora, and Chris Ré. 2025. Look Ma, No Bubbles! Designing a Low-Latency Megakernel for Llama-1B. Hazy Research technical blog. https: //hazyresearch.stanford.edu/blog/2025-05-27-no-bubbles [34] Stuart H. Sul, Simran Arora, Benjamin F. Spector, and Christopher Ré. 2025. ParallelKittens: Systematic and Practical Simplification of Multi-GPU AI Kernels. arXiv:2511.13940 https://arxiv.org/abs/2511.13940 arXiv:2511.13940. [35] Zewei Tao and Yunpeng Huang. 2025. MagiAttention: A Distributed Attention Towards Linear Scalability for Ultra-Long Context, Heterogeneous Mask Training. Software and technical documentation. https://github.com/SandAIorg/MagiAttention/ [36] vLLM. 2025. vLLM Attention Operations: Context-Parallel LSE All-Gather and Output Reduce-Scatter. Software documentation. https://docs.vllm.ai/en/v0.11. 1/api/vllm/attention/ops/common/ Version 0.11.1; cp_lse_ag_out_rs; accessed 2026-09-25. [37] vLLM. 2026. vLLM: Context Parallel Deployment. Software documentation. https://docs.vllm.ai/en/latest/serving/context_parallel_deployment/ Official documentation; accessed 2026-09-25. [38] Bingyang Wu, Shengyu Liu, Yinmin Zhong, Peng Sun, Xuanzhe Liu, and Xin Jin. 2024. LoongServe: Efficiently Serving Long-Context Large Language Models with

Elastic Sequence Parallelism. In Proceedings of the ACM SIGOPS 30th Symposium on Operating Systems Principles. ACM, 640–654. doi:10.1145/3694715.3695948 [39] An Yang, Anfeng Li, Baosong Yang, et al. 2025. Qwen3 Technical Report. arXiv:2505.09388 https://arxiv.org/abs/2505.09388 arXiv:2505.09388. [40] Haoyu Yang, Zan Zong, Yuyang Jin, Kinman Lei, Jiaao He, Qigang Yang, and Jidong Zhai. 2025. UltraAttn: Efficiently Parallelizing Attention through Hierarchical Context-Tiling. In Proceedings of the International Conference for High Performance Computing, Networking, Storage and Analysis. ACM, 1769–1784. doi:10.1145/3712285.3759894 [41] Zihao Ye, Lequn Chen, Ruihang Lai, Wuwei Lin, Yineng Zhang, Stephanie Wang, Tianqi Chen, Baris Kasikci, Vinod Grover, Arvind Krishnamurthy, et al. 2025. FlashInfer: Efficient and Customizable Attention Engine for LLM Inference Serving. Proceedings of Machine Learning and Systems 7 (2025). [42] Yilong Zhao, Xiaonan Nie, Kan Zhu, Shuang Ma, Zhichao Lai, Hongxiang Hao, Yang Zhou, Baris Kasikci, and Ion Stoica. 2026. Unleashing Scalable Context Parallelism for Foundation Models Pre-Training via FCP. Proceedings of Machine Learning and Systems 8 (2026), 2129–2142. [43] Lianmin Zheng, Liangsheng Yin, Zhiqiang Xie, et al. 2024. SGLang: Efficient Execution of Structured Language Model Programs. In Advances in Neural Information Processing Systems, Vol. 37. Neural Information Processing Systems Foundation, Inc. (NeurIPS), 62557–62583. doi:10.52202/079017-2000 [44] Size Zheng, Xuegui Zheng, Li-wen Chang, and Jidong Zhai. 2026. UniEP: Unified Expert-Parallel MoE MegaKernel for LLM Training. arXiv:2604.19241 https: //arxiv.org/abs/2604.19241 arXiv:2604.19241. [45] Zilin Zhu. 2026. Ring Flash Attention: Zigzag Ring Attention Implementation. Software documentation. https://github.com/zhuzilin/ring-flash-attention/blob /main/ring_flash_attn/zigzag_ring_flash_attn.py Source repository; accessed 2026-09-25.

10

LiquidCP: Decoupling Attention Tasks from Communication and Kernel Boundaries in Context Parallelism

A A.1

grid-search the time unit, 𝜌, and 𝛿 attn to minimize root-mean-square error (RMSE) in per-CTA busy time. For reduction, we schedule tasks on warps in queue order from observed CTA start times and fit 𝛿 red and 𝛾 to sorted CTA busy times by least squares. We use the fitted parameters (Table 2) across all evaluated DCP degrees and workloads. Table 2: Critical-wave model parameters measured on H100. Overheads and costs are in model time units.

Technical Details BuddyRing Search Details

Candidate groups. BuddyRing tests power-of-two group sizes up to min(𝑁 , 8). Causal splits require 256𝑔𝑖 to divide 𝐿𝑖 , keeping all 2𝑔𝑖 zigzag chunks aligned to 128 tokens. For a structural sequence, the minimum group size is the smallest legal 𝑔 satisfying 𝐿𝑖 /𝑔 ≤ (1 + 𝜖𝑇 )𝑇¯ and 𝑎𝑖 /𝑔 ≤ (1 + 𝜖𝑊 )𝑊¯ , or the largest legal size if none qualifies. The splitting profile in Section 4.2 uses length cutoffs of 2K, 4K, 8K, and 16K tokens. Table 1: Default BuddyRing parameters. Parameter Load tolerances 𝜖𝑇 / 𝜖𝑊 Structural threshold on 𝜅𝑖 Beam width State-merging bin (of mean load) Filler smooth-max 𝜆 Top-ranked placements refined Maximum repair iterations GPUs per load extreme / max. sequences per GPU

Value 10% / 5% 0.5 64 2% 8 8 32 2/4

Value

Time per KV tile with all compute CTAs active Per-tile time floor 𝜌 History-attention overhead 𝛿 attn History-reduction overhead 𝛿 red Cost per merged partial vector 𝛾

1.08 𝜇s 0.43 32 2.5 0.47

Fit quality. On the calibration batch, per-CTA RMSE is 1.2 𝜇s for attention (measured range 42–67 𝜇s) and 0.8 𝜇s for reduction (16–35 𝜇s). The three CTAs with two attention tasks finish in 64.9– 67.3 𝜇s, compared with the predicted 65.5–66.4 𝜇s. The fixed cost of 32 tile-equivalents per attention task exceeds the 5–14 KV tiles each task processes in this batch.

Pareto beam search. For each structural sequence, the beam tries every allowed group for each retained placement and discards dominated candidates. For load 𝑋 ∈ {𝑇 ,𝑊 }, overload is max(0, max𝑟 𝑋𝑟 /𝑋¯ − 1 − 𝜖𝑋 ) and spread is (max𝑟 𝑋𝑟 − min𝑟 𝑋𝑟 )/𝑋¯ . For matching per-GPU load bins, BuddyRing keeps the best placement by splitting profile, communication proxy, then distinct group count. The beam keeps the best candidate for each of these three criteria, each spread, and the larger overload. It fills remaining slots by decreasing NSGA-II crowding distance over all criteria [10].

Split planning. With 𝐾 compute CTAs, decode requests with at most 16 query tokens and at least 𝐾 local history KV tiles receive a split limit of at least four. Score ties favor earlier attention completion, then fewer attention tasks, then fewer total splits. The planner also uses 𝑛𝑡 + 𝛿 attn for chunk-attention tasks. For unsplit sequences, reduction tasks copy runs of output vectors. With ℎ𝑞 local query heads, MegaDCP chooses the largest run in {1, ℎ𝑞 , 2ℎ𝑞 , 4ℎ𝑞 , min(8ℎ𝑞 , 32)} that leaves at least 0.8 × 12𝐾 tasks for the 12𝐾 warps.

Filler placement. In decreasing 𝜅𝑖 order, each filler goes to the GPU that minimizes the larger post-placement overload. Ties favor lower smax𝑟 (𝑇𝑟 /𝑇¯ ) + smax𝑟 (𝑊𝑟 /𝑊¯ ), then lower GPU index, with Í smax𝑟 (𝑥𝑟 ) = 𝜆 −1 log 𝑟 𝑒 𝜆𝑥𝑟 .

Release-aware ordering. In split decode-only plans, each history task divides its modeled cost equally among the 16-token query blocks it needs. Blocks transfer in decreasing order of their summed costs, in epochs of max(1, ⌊𝑀/𝑃⌋) transfers for 𝑀 communication CTAs and DCP degree 𝑃. History tasks run in order of their last required block’s epoch, then by decreasing 𝑐 attn .

Local repair. Repair considers the highest-𝜅𝑖 sequences on the most and least loaded GPUs for each measure (Table 1). Local sequences can move from a high-load GPU to a low-load GPU or swap between the two. A sequence on a high-load GPU can move to its parent group or, if split, to a same-size or child group containing a low-load GPU. A group containing a marked GPU can move one sequence to each child. Repair takes the best legal improvement until none remains or the iteration limit is reached.

A.3

Execution Engine Details

MegaRing KV publication. After storing a KV tile for all heads, a communication CTA probes consecutive slots from the committed prefix using acquire–release compare-and-swap (CAS). A failed CAS makes the occupied slot’s KV data visible. After inserting the physical index at 𝑝, the CTA raises the prefix count to at least 𝑝 + 1 with a release atomicMax. Producer warps acquire-load the count before claiming work. Once transfers finish, communication CTAs synchronize locally and reuse shared memory for compute.

Relaxation. If placement is infeasible, fillers enter the beam search by length class, longest first, trying group sizes 2, 4, then 8. BuddyRing skips stages with unchanged groups. The final stage permits every legal group for every sequence.

A.2

Parameter

MegaDCP CTA configuration. Each 384-thread compute CTA has three warp groups. Per-thread register quotas change from 24/240/240 for attention to 168/168/168 for reduction. The leader claims 12 history-reduction descriptors at once, one per warp. For final reduction, eight warps execute one descriptor at a time per CTA.

Critical-Wave Model Details

Calibration. We calibrate the model in Section 5.2 using per-SM stage timings from Figure 4 (H100, decode-only, DCP degree 2). Of 135 history-attention tasks, the first 132 map to CTAs in index order. The first three CTAs to finish each take one more task. We 11

Hongyu Chen et al.

MegaDCP block transfers. Each 128-row history-attention tile waits for its 16-token query blocks, which communication CTAs pull in host order. After writing the send buffer, history-reduction warps increment the counter for the destination rank and 16-token output block. For a remote block, the last contributing warp sets a ready flag on the query owner’s GPU with system-scope release ordering. The owner’s communication CTA acquire-reads the flag, pulls the output and LSE, then signals local receive completion.

B

Additional Evaluation Details

B.1

Workloads and Baselines

CDF (%)

100

ProLong ArXiv FreeLaw GitHub Pile-CC

75 50

SGLang. Against the best eager baseline, MegaDCP’s geometricmean speedups are 3.72× for decode-only and 1.52× for mixed batches, consistent with greater launch sensitivity for short queries. Measurement protocol. MegaRing and MegaDCP run 10 and 40 warm-ups, respectively, before timing. CUDA Graph baselines add three eager calls before capture, then use graph replay. Uniform batches use the per-configuration communication-CTA sweep in Section 7.1.

B.2

The L1 baseline in Figure 11(a) uses eight attention launches, seven NCCL transfers, and eight external reductions.

Cache length (tok.) Cache ratio (%)

Our Megatron baseline uses the hybrid-CP scheduler from Megatron-LM [28] (368fa88) to select power-of-two CP degrees from token targets, balance attention work, and expand groups onto idle ranks. Our padded implementation of Zeppelin’s Algorithm 2 [4] allocates GPUs to long sequences in proportion to squared length and packs short sequences locally. Its linear modules retain the attention layout.

25 0

1

128

16K

1M 0 512

Sequence tokens

0

(a) Training datasets

8K

128K

50

100

(b) Inference trace

With the MegaRing kernel fixed, BuddyRing gains up to 1.16/1.18× in forward/backward throughput over Megatron-derived placement on FreeLaw. With placement fixed, MegaRing outperforms FA3/NCCL by 1.11–1.34× for Megatron and 1.62–2.13× for Zeppelin (Figure 11(b)).

Figure 15: Training and inference workload distributions. (a) Sequence-length CDFs for five training datasets. (b) CDFs of cache length and cache ratio in the Mooncake trace. Table 3: Mooncake snapshots by arrival-rate scale. D/M denotes decode-only/mixed batches. Request counts are means within each batch type. Lengths are means over request occurrences.

Snapshots (D/M)

1×

2×

4×

B.3

MegaDCP Stage Profiling Decode-only

All

P50 completion time (μs)

Arrival-rate scale

MegaRing Ablation Details

87/13 69/31 54/46 210/90

Mean requests per batch Decode-only: decode 8.28 17.12 34.54 17.93 Mixed: decode 8.38 17.26 34.26 24.67 Mixed: chunk prefill 1.15 1.13 1.48 1.31

Mixed batch

175

1600

150

1400

MegaDCP Q transfer Attention History red. / publish Receive Final reduction

1200

125

1000

100

800 75

vLLM-A2A Graph Q all-gather History attention A2A pack All-to-all Unpack + reduction Chunk attention Final reduction

600

50

400

25

200

0 6 × 10

1

102

0 2 × 10

2

102

103

Global effective workload (GFLOP)

Mean length (tokens) Decode KV history 12,156 12,479 12,174 12,259 Prefill KV history 20,992 16,574 12,040 14,523 Prefill chunk 2,420 3,486 2,772 2,939

Batch

Query transfer

Attention

Result exchange

Final reduction

Decode-only Mixed

20/29 67/104

73/71 297/376

102/109 376/506

107/128 394/570

Figure 16: Stage completion at DCP degree 4 and 4× arrivals. Entries give mean times (𝜇s) for MegaDCP/vLLM-A2A Graph. Exchange includes history reduction. Both use a common host-side start. MegaDCP (solid) timestamps the last task’s completion in each stage in-kernel and adds launch latency. vLLM-A2A Graph (dashed) accumulates CUDA-event durations scaled per case to latency without profiling events.

We simulate Mooncake requests in 20 ms steps using trace arrival times and input/output lengths. For each scale in Table 3, we divide interarrival times by that factor and sample 100 snapshots between 300 and 3,600 s. Each decoding step accepts 0–4 of 15 drafts with probabilities (0.20, 0.25, 0.25, 0.20, 0.10), committing 2.75 tokens on average including the target. MegaDCP microbenchmarks use 32 query heads, head dimension 128, and TP degree 8. DCP degrees 2, 4, and 8 use four, two, and one KV heads, respectively, so each group shares one KV head.

Figure 16 orders cases by total effective attention FLOPs across GPUs and smooths per-case medians with a 1 : 2 : 1 moving The public mKernel [26] (68644c9) and ParallelKittens/ThunderKittens [34] average. These cross-system timings reflect overlap alongside other (34b15f7) ring kernels lack the causal mask required for the MegaRimplementation differences. MegaDCP is faster in 53 of 54 decodeing comparison. vLLM uses more Triton-fused operators than only cases and all 46 mixed cases. Mean mixed-batch speedups 12

LiquidCP: Decoupling Attention Tasks from Communication and Kernel Boundaries in Context Parallelism

of 1.81× in the lowest FLOP quartile and 1.20× in the highest are consistent with less scope for overlap as attention dominates.

B.4

With four requests, history splitting exposes parallelism across all lengths (1.22–1.37×). Across all 72 configurations, geometric-mean speedups over vLLM-AG-RS Graph and SGLang-AG-AR Graph are 1.30× and 1.70×, respectively.

Uniform-Length MegaRing Results

BuddyRing and Zeppelin assign each 𝑇 /𝐵-token sequence to the same 8/𝐵 GPUs (Table 4). Their comparison isolates execution differences, including fused reductions and readiness-driven coalescing. Table 4: MegaRing Hybrid forward/backward speedups (×) for 𝐵 equal-length sequences totaling 𝑇 tokens. Baselines, in column order, are all-gather, Llama-3 style, Ring FA, Megatron, MagiAttention, and Zeppelin. 𝑇

AG

𝐵

L3

Ring

Meg.

Magi

B.6

Both experiments use Qwen3-30B-A3B [39] (32 query heads, four KV heads, head dimension 128). Transformer-layer training. Attention uses eight-way CP and the MoE block uses eight-way expert parallelism (EP). All weights, including the expert router, are randomly initialized. In Figure 13, Core Attention includes CP communication. Other covers projections, EP communication, and remaining operations. On ProLong, Other latency for Megatron, Zeppelin, and MagiAttention is close to uniform partitioning. It rises with token imbalance on shortsequence-heavy workloads. Both MegaRing variants stay close to uniform partitioning.

Zep.

64K 64K 64K

1 1.05/1.22 1.03/1.15 1.26/1.27 1.29/1.28 1.06/1.31 1.29/1.28 2 1.20/1.49 1.14/1.29 1.85/1.84 1.13/1.11 1.16/1.29 1.13/1.11 4 1.53/2.03 1.38/1.51 3.06/2.47 1.07/1.03 1.34/1.31 1.09/1.03

128K 128K 128K

1 1.00/1.10 1.02/1.06 1.09/1.13 1.11/1.14 1.00/1.19 1.12/1.13 2 1.04/1.26 1.06/1.16 1.30/1.38 1.31/1.32 1.03/1.17 1.04/1.05 4 1.23/1.50 1.17/1.24 1.80/1.84 1.18/1.15 1.12/1.14 1.10/0.99

256K 256K 256K

1 1.02/1.02 1.00/1.02 1.02/0.99 1.02/1.00 1.01/1.10 1.02/1.00 2 1.02/1.11 1.02/1.08 1.11/1.10 1.13/1.15 0.99/1.10 0.99/1.00 4 1.10/1.25 1.10/1.13 1.36/1.26 1.37/1.34 1.21/1.10 1.00/0.97

vLLM serving. We replay a contiguous 303 s Mooncake segment [30] (100 warm-up and 1,000 measured requests). Only new prompt tokens need chunked prefill on the decode side of disaggregated serving. With TP and EP degrees of 8 and DCP degree 2, each rank pair shares one KV head. Figure 14 varies arrival-rate scale (1, 2, 4) and maximum batch size (64, 128). Speedups relative to the best vLLM baseline per configuration and metric are averaged arithmetically. At 4× arrivals and batch size 128, attention takes a larger share of each iteration and TBT speedup reaches 1.10×. At batch size 64, faster iterations shorten the wait for new prompts. With the shared engine, MegaDCP naive reaches up to 1.07× for TBT and 1.29× for TTFT before critical-wave planning and release-aware ordering.

Against the fastest baseline per configuration, Hybrid’s geometricmean speedup is 1.04× in both directions (forward 0.99–1.13×, backward 0.97–1.15×). At 64K (8K tokens per GPU), per-step transfers and merges cost proportionally more. For 𝐵 = 4, Ring FA’s eight-way split of each 16K sequence also makes placement a factor. For a single 128K or 256K sequence, all methods use eight GPUs and the fastest baselines, mostly all-gather variants, hide communication behind attention (Hybrid speedup 0.99–1.06×).

B.5

Uniform-History MegaDCP Results

Against vLLM-A2A Graph, the strongest trace baseline, MegaDCP achieves 1.24× on 24 decode-only configurations and 1.20× on 48 mixed configurations (geometric means). The respective ranges are 1.03–1.37× and 1.05–1.53× (Table 5). Table 5: MegaDCP speedups (×) over vLLM-A2A Graph at DCP degrees 2/4/8 for history length 𝐿KV per request. Each batch has 𝐵𝑑 decode requests with 16 query tokens each and 𝐵𝑝 prefill requests with one 2K-token chunk each. Workload 𝐵𝑝

𝐵𝑑

𝐿KV 16K

32K

64K

128K

0

4 1.25/1.35/1.35 1.22/1.32/1.36 1.26/1.37/1.32 1.24/1.28/1.30 32 1.34/1.29/1.37 1.24/1.16/1.15 1.15/1.14/1.03 1.10/1.11/1.09

1

4 1.05/1.40/1.41 1.08/1.14/1.30 1.07/1.15/1.22 1.17/1.18/1.07 32 1.42/1.53/1.38 1.27/1.38/1.29 1.15/1.24/1.23 1.12/1.11/1.07

2

4 1.15/1.33/1.35 1.17/1.33/1.24 1.07/1.12/1.13 1.08/1.06/1.06 32 1.42/1.36/1.35 1.27/1.25/1.23 1.16/1.16/1.12 1.10/1.08/1.07

End-to-End Evaluation Details

In decode-only batches of 32 requests, speedups fall from 1.29–1.37× at 16K to 1.09–1.11× at 128K as attention dominates. 13

Related documents

• Language Models Can Control Their Own Attention#924873
• Language Models Can Control Their Own Attention#920639
• Splitting Prompt Prefill from Response Replay for Context-Parallel Long-Context LLM Post-Training#1117845
• Splitting Prompt Prefill from Response Replay for Context-Parallel Long-Context LLM Post-Training#1108735
• SPLASH: Switching Parallel Layouts of Attention with Seamless Handoff for LLM Serving#1141940