PaperlayerView exact revision on arXiv

Source-linked AI summary

DeltaLog: Deferred Materialization of Recurrent States for Linear Attention Decoding

Junqing Lin, Jingwei Sun, Guangzhong Sun

arXiv:2608.15533v1cs.DCcs.LG

TL;DR

Linear-attention decoding still incurs substantial memory traffic when recurrent states are eagerly materialized after every token. DeltaLog defers dense write-backs using a bounded log of compact updates, accelerating recurrent-state update kernels by up to 1.86× and end-to-end serving by 1.05–1.20×.

  • Problem

    Eager recurrent-state materialization places full-state read–modify–write traffic on every generated token, despite each token introducing only compact update factors.

  • Method

    DeltaLog stores a dense base state with a bounded log of compact updates, appending most steps and periodically merging updates into the base.

  • Results

    Across GDN, KDA, and RWKV6, DeltaLog accelerates recurrent-state update kernels by up to 1.86× and achieves 1.05–1.20× end-to-end serving speedups over dense recurrent baselines.

  • Takeaways & Limitations

    The results identify recurrent-state materialization as a significant bottleneck and support optimizing the physical state-update schedule for linear-attention inference efficiency.

  • Takeaways & Limitations

    DeltaLog targets decoding rather than prefill, and end-to-end speedups are bounded by time spent outside the optimized update path.

Abstract

from arXiv · show

Linear attention models eliminate the quadratic prefix computation and context-growing KV cache of softmax attention by replacing pairwise token interactions with recurrent state updates. However, existing decoding implementations often materialize and write back the full recurrent state after every generated token, making state maintenance a major source of memory traffic, especially for models with large states and many heads. This paper presents DeltaLog, a recurrent-state decoding scheme that reduces this overhead without changing the model semantics. Specifically, DeltaLog represents the recurrent state as a dense base state together with a bounded log of recent compact updates. Most decode steps append only compact update factors to this log, while periodic merge steps fold the accumulated updates back into the dense base state. Thus, the model observes the same dense state as in eager decoding, but most full-state write-backs are replaced by lightweight append operations. We implement DeltaLog for GDN, KDA, and RWKV6 and integrate it into a prototype serving stack. Across these models, DeltaLog accelerates the recurrent-state update kernel by up to $1.86\times$, reduces profiled recurrent-state write traffic by up to $7.83\times$, and achieves $1.05$--$1.20\times$ end-to-end serving speedups over dense recurrent baselines.

1 Introduction

Recurrent linear-attention decoding avoids softmax attention’s prefix-scale computation but can incur excessive memory traffic by eagerly materializing dense recurrent states after every token. DeltaLog defers these write-backs by combining a dense base state with a bounded log of compact updates, preserving the logical recurrence while improving kernel and serving performance.

  • Motivation: Recurrent linear-attention models summarize the prefix in a fixed-size state, but eager implementations materialize and write the full dense state after each compact token update.The compact update contains O(BH(d_k+d_v)) values, whereas eager write-back touches O(BHd_kd_v) state elements per token.
  • Motivation: The resulting state-update tax comes from materialization policy rather than recurrence semantics, motivating deferred updates that subsequent reads incorporate exactly.Recent compact updates need not be stored as a freshly materialized dense tensor after every token.
  • DeltaLog: DeltaLog represents recurrent state as a dense base plus a bounded log of recent compact updates, so most decoding steps append factors instead of rewriting the full state.Periodic merging folds accumulated updates back into the dense base while preserving the state representation needed for decoding.
  • Evaluation: 1.19–1.69× lower latency on H200 and 1.30–1.86× on RTX 4090 are reported for one-token decoding versus the Fast Linear Attention dense baseline across GDN, KDA, and RWKV6.The comparison uses FLA, described as a highly optimized dense recurrent-state baseline.
  • Evaluation: 1.08–1.20× gains on Qwen3.6-35B-A3B and 1.05–1.12× on Kimi-Linear-48B-A3B-Instruc are reported in the serving prototype.The serving gains are smaller but positive than the decode-kernel improvements.

2 Background and Motivation

Linear attention replaces token-level KV history with a fixed-shape recurrent state, but eager decoding materializes that dense state on every token. Because each token contributes only compact update factors, this creates a bandwidth-heavy cost that grows with serving concurrency and can dominate decode latency.

  • Recurrent linear attention: Linear attention aggregates prefix key–value interactions into a matrix state whose shape is independent of context length.Unlike softmax attention’s token-level KV history, the recurrent state carries the prefix in a fixed-size matrix.
  • Recurrent linear attention: Modern recurrent variants share a common step structure: transform the previous state, then add an outer product of token-dependent vectors.Decay, gating, and delta-rule corrections instantiate this common formulation with model-specific update semantics.
  • Eager state materialization: Eager decoding reads and writes the full dense state for every generated token, despite each token contributing only O(d_k+d_v) update factors.This replaces compact logical updates with a full-state read–modify–write operation on the critical path.
  • Eager state materialization: With B requests and H heads, dense state traffic scales with [B, H, d_k, d_v], making fixed per-request states an aggregate bandwidth cost at higher concurrency.The update factors occupy only O(BH(d_k+d_v)) elements, while the dense recurrent-state tensor is maintained and accessed each iteration.
  • Serving impact: 44.6% of decode latency is consumed by recurrent-state materialization at B=256 for a 1k-token context, versus 25.2% at B=64.The corresponding shares at 2k and 4k contexts also rise monotonically with batch size in Qwen3.6-35B-A3B.
  • Bandwidth bottleneck: 1 GiB is occupied by the recurrent state at B=512, H=32, d_k=d_v=128, and a 32-bit state; one eager step moves roughly 2 GiB for state read and write alone.The resulting arithmetic intensity is far below modern GPU machine balance, placing eager updates in the memory-bound regime.

3 Method

DeltaLog replaces per-token dense recurrent-state write-back with a dense base plus a bounded log of compact updates. Periodic merges bound replay while preserving exact-arithmetic equivalence to eager decoding.

  • State representation and cycle: DeltaLog stores a dense base with a bounded log of recent compact updates, appending during M−1 steps and materializing at the M-th step.The live-log capacity is C = M−1, so each cycle contains C append steps followed by one merge step.
  • Append–merge decoding cycle: Append steps record compact updates while leaving the dense base unchanged; merge steps materialize the represented logical state, clear the log, and reset accumulated decay.This decouples frequent compact updates from infrequent dense-state materialization while keeping log growth bounded.
  • Correctness invariant: In exact arithmetic, the base-plus-log representation remains logically equivalent to the state produced by eager dense decoding after both append and merge steps.The recurrence is closed under appending a decayed compact contribution and under merging the represented state into the dense base.
  • Traffic tradeoff: Increasing M lowers amortized dense-state write cost but raises replay traffic, making the optimal merge interval dependent on the operator and kernel.DeltaLog writes M−1 log entries per M-step cycle, while replay work grows with the accumulated log.

4 End-to-End System Implementation

DeltaLog is integrated into a vLLM-based, graph-captured serving runtime without changing model weights or prefill kernels. The implementation uses phase-specific CUDA graphs, base-plus-log state allocation, and separate append/merge scheduling to preserve efficient continuous batching.

  • System integration: The prototype replaces eager dense decoding in a vLLM-based serving stack while leaving model weights and prefill kernels unchanged.The integration must replay phase-specific decode graphs, manage base-plus-log recurrent state, and preserve correctness with continuous batching.
  • Graph selection and execution: Because append and merge phases require distinct schedules and state semantics, the runtime captures separate CUDA-graph variants for each uniform decode-batch descriptor.Both graph variants expose the same model-level inputs and outputs but commit recurrent state differently.
  • State-slot lifecycle: Each recurrent-state slot stores a dense prefill base plus bounded key/value logs, log-decay metadata, snapshots, and live-log metadata.A newly assigned slot initializes the dense base from the prefill output and creates the auxiliary representation.
  • Phase-specialized decode scheduling: The runtime maintains separate append and merge queues, routing requests to merge execution at boundaries while admitting new requests directly into the append queue.New requests begin with empty auxiliary state and are not forced into merge executions without deferred updates.

5 Experiments

Experiments evaluate DeltaLog from standalone recurrent decode kernels through integrated serving, across GDN, KDA, and RWKV6 on two GPUs. DeltaLog improves kernel and end-to-end performance while preserving tolerance-equivalent outputs, with benefits shaped by operator structure and recurrent-state size.

  • Experimental Design: Experiments measure one-token recurrent decode latency, write-back mechanisms, and serving impact using kernel- and serving-level evidence.Kernel tests span operators, batch sizes, and GPUs; serving tests assess integrated decode performance.
  • Kernel Performance: 1.69–1.86× GDN, 1.30–1.55× KDA, and 1.41–1.63× RWKV6 FLA-relative speedups occur on RTX 4090.On H200, corresponding speedups are 1.34–1.60×, 1.19–1.60×, and 1.52–1.69×.
  • Kernel Performance: At B=128, best FLA-relative settings are GDN M=8 with 0.103 ms and 1.57×, KDA M=4 with 0.116 ms and 1.34×, and RWKV6 M=4 with 0.110 ms and 1.60×.The measured interval curves are U-shaped because short intervals merge too frequently, while long intervals increase append latency.
  • Serving and Numerical Equivalence: 1.08–1.20× Qwen/GDN and 1.05–1.12× Kimi/KDA TPOT speedups improve all six recorded serving configurations.Maximum output relative errors are 0.00303 for GDN, 0.00450 for KDA, and 0.00376 for RWKV6, supporting tolerance-equivalence to dense references.

6 Discussion and Limitations

DeltaLog targets large-batch recurrent decoding, where recurrent-state write-back can account for substantial memory traffic. Its write-traffic benefits are corroborated by H200 GDN profiling, but end-to-end speedups remain limited by components outside the optimized update path.

  • DeltaLog is designed for large-batch recurrent decoding, where recurrent-state write-back can constitute substantial memory traffic.This focus is reflected in the state-size sweep and end-to-end serving results.
  • H200 GDN counter profiling under a representative configuration confirms the expected reduction in recurrent-state write traffic.
  • At the serving level, realized speedups are limited by components outside the optimized update path.

7 Related Work

DeltaLog is positioned among recurrent linear-attention architectures, efficient recurrent algorithms and kernels, and IO-aware inference systems. Unlike prior approaches that modify architectures, optimize parallel execution, or target Transformer KV-cache bottlenecks, it reduces decode-time recurrent-state write-back while preserving the trained model.

  • Positioning: DeltaLog reduces decode-time recurrent-state write-back without modifying the architecture or trained model semantics.The paper contrasts this approach with prior work that modifies architectures or optimizes parallel execution.
  • Recurrent linear attention: Linear attention replaces token-level histories with constant-size recurrent statistics, while later models add decay, gating, delta-rule corrections, or finer-grained updates.Examples include RetNet and GLA for decay or gating, DeltaNet and Gated DeltaNet for delta rules, and RWKV6 and Kimi Linear for dynamic or finer-grained updates.
  • Efficient recurrent algorithms and kernels: Prior recurrent algorithms use parallel or chunkwise formulations to address sequential dependencies that complicate parallel training and prefill.These formulations are applied to recurrent linear attention, while systems such as FLA, ThunderKittens, FlashRNN, and MetaAttention optimize kernels, abstractions, or code generation.
  • IO-aware serving: IO-aware attention and serving systems reduce softmax-attention or Transformer KV-cache costs through tiling, scheduling, batching, paging, offloading, and disaggregation.The passage distinguishes these methods from recurrent linear attention, whose decode-time bottleneck is dominated by fixed-state maintenance.

8 Conclusion

DeltaLog defers recurrent-state materialization by combining a dense base with a bounded log of compact updates, avoiding most full-state write-backs while preserving recurrence semantics. Across GDN, KDA, and RWKV6, it delivers 1.05–1.20× end-to-end serving speedups over dense recurrent baselines.

  • Conclusion: DeltaLog represents recurrent state as a dense base plus a bounded log of compact updates, preserving the original recurrence semantics while avoiding most full dense-state write-backs.Most decode steps append compact updates, while periodic merges fold them into the dense base state.
  • Conclusion: Recurrent-state materialization is a significant bottleneck in linear-attention decoding, making physical state-update scheduling an effective direction for improving inference efficiency.The conclusion identifies state-update scheduling as the optimization target enabled by DeltaLog.
Loading 2608.15533v1…