RedFuser: Fused Reduction Compiler

• RedFuser is a compiler framework that fuses cascaded reductions into single-loop kernels for efficient online computation.
• It employs formal algebraic conditions such as decomposability, commutativity, and distributivity to legally transform serial reduction chains.
• The framework supports diverse AI applications including safe softmax, attention, MoE routing, and FP8 quantization while reducing memory traffic.

Searching arXiv for the RedFuser paper and closely related prior work on fused attention and online softmax. Reduction Fuser, usually abbreviated RedFuser, is a compiler framework and formal methodology for automatically fusing cascaded reductions—sequences of reduction loops with inter-loop data dependencies—into single-loop kernels with incremental, or streaming, computation on AI accelerators (Tang et al., 24 Feb 2026). The framework addresses a class of workloads in which the root value of one reduction is required to compute the next, a structure that is common in safe softmax, attention, mixture-of-experts routing, and FP8 quantization pipelines. Its central claim is that, under specific algebraic conditions, chains of dependent reduction trees can be transformed into one fused reduction tree with same-level dependencies, numerically stable online update rules, and $O(1)$ on-chip state per stream, thereby eliminating redundant memory traffic and synchronization while matching hand-written kernels such as FlashAttention on supported patterns (Tang et al., 24 Feb 2026).

1. Problem setting and computational motivation

RedFuser is formulated around cascaded reductions, which are sequences of reductions with strong data dependencies. A reduction computes a single value over an axis using an associative operator, and on GPUs it is ordinarily implemented as a hierarchy of reduction trees in which partial results are computed per-thread, per-warp, and per-block before a final root is produced. In a cascaded structure, the output root of reduction $i$ is needed before reduction $i+1$ can proceed, so the trees form a serial chain rather than an independently parallel set (Tang et al., 24 Feb 2026).

The motivating examples are drawn from AI workloads. In safe softmax, a max-reduction determines a running maximum, a sum-reduction accumulates exponentials, and normalization produces the final probabilities. Within attention, softmax is positioned between two GEMMs, $QK^T$ and $PV$, and both GEMMs are themselves reductions along $K$. The same structural pattern appears in MoE routing, where a GEMM for expert scores is followed by softmax and then TopK, and in FP8 per-token quantization + GEMM, where an abs-max reduction determines a scale before quantization and matrix multiplication (Tang et al., 24 Feb 2026).

The difficulty for existing AI compilers is not elementwise fusion per se, but the lack of a formal legal transformation for multi-reduction chains. TVM, XLA, and PyTorch Dynamo/Inductor can fuse elementwise operations with single reductions, but inter-loop dependencies across reductions prevent naïve loop fusion. The practical consequences are redundant re-reads from global memory, kernel launches that must serialize the chain, and the inability to overlap memory access and compute across reductions. The paper frames RedFuser as a response to this gap: a way to load inputs once, maintain constant-size on-chip state, and derive the reweighting or rescaling rules needed when upstream reduction results change (Tang et al., 24 Feb 2026).

2. Formal model of cascaded reduction fusion

The formal core begins with the standard definition of a reduction over an index set $I$: $R = \bigoplus_{i \in I} f(i),$ where $\oplus$ is associative and often commutative. Cascaded reductions are then written over inputs $\bm{X} \in \mathbb{S}^{M \times L_0}$, with the $i$0-th reduction result $i$1 depending on previous outputs $i$2: $i$3 Because $i$4 depends on prior reduction outputs, the corresponding reduction trees execute as a chain, and this chain imposes serial dependencies that block straightforward fusion (Tang et al., 24 Feb 2026).

RedFuser’s fusion theory states that such cascaded reductions can be legally fused into a single tree when three conditions hold. First, decomposability requires

$i$5

Second, $i$6 must form a commutative monoid, meaning $i$7 is associative, commutative, and has identity $i$8. Third, $i$9 must distribute over $i+1$0: $i+1$1 Under these conditions, the fused level-$i+1$2 local result depends only on level-$i+1$3 results of prior reductions rather than on final roots. The significance of this result is that the dependency structure is shifted from “root-to-root” to same-level dependency, which makes fusion compatible with hierarchical GPU reduction schedules (Tang et al., 24 Feb 2026).

A further technical point is the handling of non-invertible cases. When $i+1$4 is invertible in $i+1$5, the derivation factors it out directly. When invertibility fails, RedFuser replaces $i+1$6 with $i+1$7, which equals $i+1$8 if invertible and otherwise the identity element $i+1$9. The paper states that substituting $QK^T$0 yields a robust fused formula that coincides with the original when invertibility holds (Tang et al., 24 Feb 2026). This suggests that the framework is designed not merely around an ideal algebraic case, but around a repair mechanism that preserves legality across a broader class of practical reductions.

3. Fixed-point test and incremental computation

Automatic discovery of $QK^T$1 and $QK^T$2 is based on a fixed-point identity. RedFuser chooses $QK^T$3 such that $QK^T$4 is invertible in $QK^T$5, and then checks whether

$QK^T$6

If the identity holds, the decomposition follows constructively; if it fails, the reduction is marked NotFusable (Tang et al., 24 Feb 2026). This test is the basis of the ACRF legality analysis and gives the framework a formal criterion rather than a pattern-only heuristic.

The second key step is the derivation of an incremental (streaming) form. Instead of caching all prior-level outputs, RedFuser updates only the running state and the current input, reducing storage from $QK^T$7 to $QK^T$8 while remaining semantically exact (Tang et al., 24 Feb 2026). The canonical example is numerically stable streaming softmax. For a sequence $QK^T$9, define

$PV$0

On a new element $PV$1,

$PV$2

For a downstream GEMM-like weighted sum, RedFuser maintains

$PV$3

and updates

$PV$4

The paper explicitly states that these are the online softmax rules used in FlashAttention, generalized by RedFuser’s fusion theory to cascaded reductions (Tang et al., 24 Feb 2026, Dao et al., 2022).

The same rescale-then-accumulate template is applied beyond softmax. For FP8 Quant+GEMM, the upstream abs-max reduction determines a scale; when the running maximum changes, previously accumulated sums are rescaled before the new contribution is added. The paper also gives incremental formulas for MoE softmax+TopK and for a Sum+Sum pattern with square-and-normalize dependency (Tang et al., 24 Feb 2026). A common misconception is that RedFuser is simply an attention-specific implementation of online softmax. The supported examples indicate otherwise: attention is one instance of a broader class of cascaded reductions that satisfy the same algebraic template.

4. Compiler architecture and GPU mapping

RedFuser is implemented as a framework integrated with TVM, using Relax IR and TensorIR (TIR) for front-end lowering and analysis. It traverses the TIR AST, inlines functions, reorders loops, and extracts a mathematical representation of reduction chains, including operators, axes, and dependencies. The legality stage then applies ACRF (Automatic Cascaded Reductions Fusion), checks decomposability through the fixed-point identity, verifies the commutative-monoid and distributivity conditions, and repairs non-invertible cases by identity substitution (Tang et al., 24 Feb 2026).

Once a pattern is accepted, RedFuser instantiates fused and streaming formulas as scalar-level IR templates. The generated update logic follows three operations: save previous state, correct current state by rescaling with the inverse, and perform the reduction update. This is then mapped onto GPU execution hierarchy through two strategies. In Single-segment, one CTA streams the whole sequence with $PV$5 state and avoids inter-block synchronization. In Multi-segment, the sequence is split into segments processed by multiple CTAs in parallel, and partial results are merged at the next level with the fused formulas. The paper describes this as analogous to FlashDecoding’s chunked combine (Tang et al., 24 Feb 2026).

The backend is hardware-conscious. Tensorization converts the scalar IR to tile-level IR through blockization, block-level buffer management, and conversion to standardized TileOps. Parallelization then binds tiles to blocks and generates TileLang code. The optimization stack includes thread-level mapping, software pipelining, warp specialization, cp.async on Ampere, TMA on Hopper, MMA/WGMMA for GEMM cores, bank-conflict avoidance, and auto-vectorization. Auto-tuning chooses tile sizes, threads per block, pipeline depth, and, for multi-segment execution, the number of segments (Tang et al., 24 Feb 2026). In architectural terms, RedFuser is not only a legality transformation but also a schedule-and-code-generation pipeline that targets the GPU memory hierarchy—registers, shared memory, and global memory—with the specific goal of preserving intermediate values on-chip.

5. Supported patterns and empirical performance

The supported pattern set extends beyond softmax followed by GEMM. The paper lists attention (MHA, MLA), MoE routing, FP8 per-token Quant → GEMM, Sum+Sum chains with dependency, and non-ML reductions such as variance and inertia (Tang et al., 24 Feb 2026). This breadth is central to RedFuser’s positioning: it generalizes the logic of hand-optimized attention kernels to other cascaded reductions that satisfy the same formal conditions.

The evaluation was conducted on NVIDIA A10 (24GB) and H800 (80GB) with CUDA 12.8, with additional appendix results on A100 and AMD MI308X. Workloads include MHA for BERT and ViT configurations, MLA for DeepSeek-R1-like decode, MoE routing for Switch Transformer, ERNIE, DeepSeek-V2-Lite, and Qwen3 variants, FP8 PerToken Quant+GEMM for ERNIE, DeepSeek-R1, and Qwen3 configurations, and non-ML kernels for variance and moment of inertia. Baselines are PyTorch Eager, PyTorch Dynamo/Inductor (Triton), TVM, and hand-optimized kernels FlashAttention2 and FlashMLA (Tang et al., 24 Feb 2026, Dao, 2023).

Workload	Reported comparison	Result
MHA	vs. FlashAttention2	averages 1.09× FlashAttention2 performance
MHA on LLaMA-65B configs	vs. Dynamo / TVM	2.8× over Dynamo; 2.6× over TVM
MLA	vs. FlashMLA / Dynamo / TVM	102% of FlashMLA; 2.4× over Dynamo; 8.7× over TVM
MoE routing	vs. Dynamo / TVM	1.7× over Dynamo; 6.6× over TVM
FP8 Quant+GEMM	vs. Dynamo / TVM	3.4× over Dynamo; 12.1× over TVM
Variance	average speedups	4.8× (A10), 4.0× (A100), 2.9× (H800), 4.6× (MI308X)
Moment of inertia	average speedups	6.4×, 5.5×, 6.2×, 11.6×

The ablation results distinguish between fusion levels and between incremental and non-incremental execution. Intra-block fusion is reported as best, while inter-block and intra-warp are similar and intra-thread is worst. Non-incremental execution requires full prior-level results to be cached on-chip and, in the reported experiments, is only feasible for sequence length $PV$6. Incremental execution avoids these constraints and enables longer segments and more flexible parallelism, including configurations that peak at Waves per SM = 3 (Tang et al., 24 Feb 2026). A plausible implication is that the practical value of RedFuser comes as much from enlarging the scheduling search space as from reducing memory traffic directly.

6. Correctness, limitations, and relation to prior work

RedFuser’s correctness claim rests on the direct derivation of fused expressions from associativity, distributivity, and decomposability, together with the fixed-point legality test. For softmax-like patterns, numerical stability is maintained through the standard safe-softmax strategy: subtract the running max, rescale prior sums when the maximum increases, and normalize with the current sum. The paper states that the final result is identical to the unfused computation, while also noting that floating-point reordering may affect least significant bits; it recommends FP32 accumulators for sums and exponentials, with BF16 or FP16 inputs paired with FP32 accumulation as a common regime (Tang et al., 24 Feb 2026).

The limitations are explicit. Fusion requires decomposability into $PV$7 and distributivity of $PV$8 over $PV$9, so non-associative reductions or dependencies not expressible in this form are not currently fusible. The domain assumptions emphasize real numbers and common reductions such as sum, product, and max/min. The system lacks a global cost model, so fusion can increase register or shared-memory usage and correction compute in ways that lead to suboptimal choices. The implementation and evaluation focus primarily on GPUs; support for other accelerators is future work. The paper also describes the evaluation as inference-centric and notes that training-time backpropagation through fused reductions is plausible but not yet demonstrated (Tang et al., 24 Feb 2026).

Relative to prior work, RedFuser is presented as a generalization of hand-optimized online-softmax kernels such as FlashAttention and later attention-specialized implementations. FlashAttention introduced exact attention with IO-aware tiling and online softmax (Dao et al., 2022), and FlashAttention-2 improved parallelism and work partitioning (Dao, 2023). RedFuser’s novelty is not a new attention formula, but the formalization of cascaded reduction fusion for arbitrary patterns that satisfy the algebraic conditions, together with an automated legality test, a robust $K$0-state incremental computation framework, and a unified kernel-generation pipeline (Tang et al., 24 Feb 2026). In that sense, it shifts operator fusion from heuristic, pattern-specific implementations toward a principled transformation system for dependent reductions.

RedFuser is open source at https://github.com/alibaba/redfuser, and its intended workflow is to import models from PyTorch into TVM Relax, run the RedFuser pass to identify cascaded reduction subgraphs and derive fused or incremental forms, and then generate auto-tuned TileLang kernels for GPU execution (Tang et al., 24 Feb 2026).