1
00:00:01,000 --> 00:00:09,300
[Hal Turing] Alrighty! Thanks for tuning in! Hello AI world! I am your host, Hal Turing, and my co-host is Dr. Ada Shannon.

2
00:00:09,300 --> 00:00:48,825
[Hal Turing] Today's paper is RedFuser: An Automatic Operator Fusion Framework for Cascaded Reductions on AI Accelerators. Xinsheng Tang is the first author, with Yangcheng Li, Nan Wang and colleagues, all at Alibaba Cloud Computing, in February 2026. The question is whether a compiler can take a chain of dependent reductions, like softmax feeding a matrix multiply, derive a single fused loop from the math, and generate GPU kernels that beat general compilers and match hand-written ones. The code is open-sourced at github.com/alibaba/redfuser.

3
00:00:48,825 --> 00:01:04,900
[Dr. Ada Shannon] The stakes are that FlashAttention is a hand-derived trick, and every new attention variant needs a new derivation. If a compiler can do that derivation, kernel authors stop being the bottleneck. Whether this paper shows that is what we need to work out.

4
00:01:04,900 --> 00:01:09,700
[Hal Turing] Start at the bottom. What's a reduction, and why do GPUs care?

5
00:01:09,700 --> 00:01:31,150
[Dr. Ada Shannon] A reduction collapses an array into one value: a sum, a max. A GPU does it as a tree. Each thread block reduces its own chunk, then the partial results are combined. Harris's well-known CUDA reduction tutorial covers exactly this pattern. A matrix multiply is a grid of these. Every output element is a sum reduction along the K axis.

6
00:01:31,150 --> 00:01:35,500
[Hal Turing] So reductions are everywhere. What makes them hard to fuse?

7
00:01:35,500 --> 00:02:03,150
[Dr. Ada Shannon] First, what fusion is. Operator fusion merges consecutive operators into one kernel, so intermediates stay in registers or shared memory instead of making a round trip through global memory. For elementwise chains like bias, GELU, dropout, that's easy, because each output depends on one input. A reduction collapses an axis, so each output depends on a whole row. Compilers handle one reduction with elementwise neighbors fine. The trouble starts when one reduction needs the finished result of another.

8
00:02:03,150 --> 00:02:05,050
[Hal Turing] Like softmax.

9
00:02:05,050 --> 00:02:25,375
[Dr. Ada Shannon] Right. Safe softmax has three stages. First, reduce to find the row max m. Second, sum exp of x minus m to get the denominator. Third, divide. Each stage needs the finished answer from the one before. Attention stacks a GEMM, then softmax, then another GEMM. That chain is what the paper calls cascaded reductions.

10
00:02:25,375 --> 00:02:30,625
[Hal Turing] And the cost of running them separately is just the extra memory traffic?

11
00:02:30,625 --> 00:02:51,375
[Dr. Ada Shannon] Two costs. Each stage rereads the input from memory. And the stages can't overlap, because stage two can't start until stage one has its final max. The same shape shows up outside attention: MoE routing, which is softmax then top-k, and FP8 per-token quantization followed by a GEMM, where an abs-max reduction sets the scale.

12
00:02:51,375 --> 00:02:57,250
[Hal Turing] So why don't TVM, XLA or PyTorch's Dynamo just fuse these?

13
00:02:57,250 --> 00:03:17,449
[Dr. Ada Shannon] They fuse elementwise operations with a single reduction. Fusing across dependent reductions needs algebra, not loop merging. Some research compilers, Bolt, Chimera, MCFuser and SOUFFLE, touch attention, but only incidentally. The real solutions are hand-derived: FlashAttention, FlashDecoding, each specific to its pattern.

14
00:03:17,449 --> 00:03:22,024
[Hal Turing] Then show me the trick. I want a number I can hold in my head.

15
00:03:22,024 --> 00:03:53,250
[Dr. Ada Shannon] Take three scores: 1, 3, and 5, arriving in two tiles. First tile is 1 and 3. The running max is 3, and the denominator is e to the minus 2, about 0.14, plus 1. Call it 1.14. Now 5 arrives and the max jumps to 5. Every earlier term was computed relative to 3, so it's stale. But you don't recompute. Multiply the old denominator by e to the minus 2, the gap between old and new max, and add the new term, which is 1. You get about 1.15. Computing from scratch, e to the minus 4, plus e to the minus 2, plus 1, gives the same.

16
00:03:53,250 --> 00:03:59,550
[Hal Turing] So the correction is one multiplication by exp of old max minus new max.

17
00:03:59,550 --> 00:04:23,000
[Dr. Ada Shannon] Yes. That's online softmax, from Milakov and Gimelshein at NVIDIA, 2018, in their paper on the online normalizer calculation for softmax. One pass keeps a running max and a running denominator. Rabe and Staats at Google, 2021, in Self-attention Does Not Need O of n squared Memory, extended it to chunked, numerically stable attention. FlashAttention, by Tri Dao and colleagues at Stanford in 2022, pushed it into a fused GPU kernel.

18
00:04:23,000 --> 00:04:27,575
[Hal Turing] What about the second reduction, the one that multiplies by V?

19
00:04:27,575 --> 00:04:49,350
[Dr. Ada Shannon] Same move. Keep a running output accumulator, the sum of exp of x minus m times v. When the max changes, multiply the accumulator by that same correction factor. Divide by the denominator once, at the end. One pass over K and V replaces three. It's exact math, not an approximation, apart from floating-point reordering. And the cost is still quadratic in sequence length. It saves memory traffic, not FLOPs.

20
00:04:49,350 --> 00:04:54,176
[Hal Turing] Does this connect to the split-across-blocks version, FlashDecoding?

21
00:04:54,176 --> 00:05:08,526
[Dr. Ada Shannon] Yes. The corrected state merges. Two partial states, each with its own max, combine with the same rescaling. So the update works across tiles in one block, and also across blocks. Keep that in mind for the next stretch.

22
00:05:08,526 --> 00:05:15,926
[Hal Turing] One thing I noticed. You named two prior sources for this trick, and the RedFuser paper doesn't cite either.

23
00:05:15,926 --> 00:05:26,476
[Dr. Ada Shannon] Correct, as far as I can tell. The claim to watch isn't the rescaling idea. It's whether a compiler can derive it for patterns nobody has hand-derived yet.

24
00:05:26,476 --> 00:05:32,351
[Hal Turing] So the bar is a compiler that does what FlashAttention's authors did by hand.

25
00:05:32,351 --> 00:05:51,326
[Dr. Ada Shannon] Right. Existing compilers mostly choose a schedule for a fixed algorithm. Ansor, from Zheng and colleagues at Berkeley in 2020, searches tiling and thread mapping, but it can't discover that softmax admits a one-pass form. RedFuser works one level up. It rewrites the algorithm first, then schedules it.

26
00:05:51,326 --> 00:05:58,751
[Hal Turing] So the compiler rewrites the algorithm. What does it need to be true about the math before it's allowed to do that?

27
00:05:58,751 --> 00:06:33,076
[Dr. Ada Shannon] Three conditions on each reduction. First, the per-element function has to split into two parts: one that depends only on the input, and one that depends only on the earlier results. Call them G and H, joined by an operator the paper writes as ⊗. Second, that operator has to be a commutative monoid, meaning order and grouping don't matter and there's an identity element. Third, the reduction operator has to distribute over it. Softmax passes. The second stage sums exp of x minus the max. That's exp of x, times exp of minus the max. The max enters as one multiplicative factor, so it comes out of the sum.

28
00:06:33,076 --> 00:06:37,801
[Hal Turing] And pulling it out is what lets you stop waiting for the final max.

29
00:06:37,801 --> 00:06:57,426
[Dr. Ada Shannon] Right. Each tile computes its partial sum using only its own local max. When a better max shows up, you multiply the partial sum by the ratio of the old factor to the new one. That's the inverse step. If the factor isn't invertible, say it can hit zero, the appendix substitutes the identity element so the correction stays valid.

30
00:06:57,426 --> 00:07:02,576
[Hal Turing] So that's the running rescale, written for any operator that qualifies.

31
00:07:02,576 --> 00:07:24,426
[Dr. Ada Shannon] Yes. That's the incremental form, equations 15 and 16. Keep one running state per row. When a tile arrives, rescale the state for the changed earlier result, then add the new tile's term. On-chip storage drops from the length of the sequence to a constant. For attention the update is exactly the FlashAttention one, and it sits at the block level of the reduction tree. The price is extra correction arithmetic on every tile.

32
00:07:24,426 --> 00:07:30,626
[Hal Turing] Softmax is the example everyone knows. What does a less familiar case look like?

33
00:07:30,626 --> 00:08:00,751
[Dr. Ada Shannon] FP8 per-token quantization followed by a GEMM. The first reduction is the absolute max of the token. The second is the GEMM sum of values scaled by the FP8 maximum over that abs-max. The abs-max again enters as a single multiplicative factor. So when the max grows, you multiply the running accumulator by the old abs-max over the new one. One caveat: the case study is written as real-number algebra. The main text doesn't say how FP8 rounding is handled, so we don't know what that does to accuracy.

34
00:08:00,751 --> 00:08:06,526
[Hal Turing] How does the compiler know a reduction qualifies? It can't just try everything.

35
00:08:06,526 --> 00:08:33,751
[Dr. Ada Shannon] It doesn't. The algorithm, called ACRF, is a restricted search. It assumes real numbers and only sum, product, max and min as reduction operators. Table 1 then fixes the matching ⊗. It picks a reference point and tests an identity: F at x and d, times F at the reference, must equal F at x and the reference d, times F at the reference x and d. If that holds, G and H fall out. SymPy does the symbolic work. So it's an algebraic test on a narrow family, not general program analysis.

36
00:08:33,751 --> 00:08:36,801
[Hal Turing] Then what turns that into a kernel?

37
00:08:36,801 --> 00:09:04,251
[Dr. Ada Shannon] A lowering chain. The framework lowers a Relax graph to TensorIR, extracts the math from the syntax tree, applies the fusion, and converts to tile operations. It emits TileLang, a tile-level GPU language. TileLang supplies software pipelining, warp specialization, and TMA and WGMMA instructions on Hopper, or cp.async and MMA on Ampere. Tile size, threads, pipeline depth and segment count are autotuned. I'm flagging that now because it matters when we assess the speedups.

38
00:09:04,251 --> 00:09:06,426
[Hal Turing] You mentioned segments.

39
00:09:06,426 --> 00:09:27,326
[Dr. Ada Shannon] There are two strategies. Single-Segment has one thread block stream the whole reduction incrementally. Multi-Segment splits the sequence across blocks and merges their partials with the fused expression. Decode has one query row, so a single block would leave most of the GPU idle. Splitting the key sequence recovers parallelism. That's the FlashDecoding trade.

40
00:09:27,326 --> 00:09:31,201
[Hal Turing] Okay, results. Start with the hand-written kernels.

41
00:09:31,201 --> 00:09:51,401
[Dr. Ada Shannon] Roughly a tie, and that's the strongest evidence here. On an A10, generated attention averages about 9 percent faster than FlashAttention-2, and wins in several configurations. On an H800, the MLA kernel is within 2 percent of DeepSeek's FlashMLA. No variance is reported, so I'd call it parity, not a win.

42
00:09:51,401 --> 00:09:54,051
[Hal Turing] And against other compilers?

43
00:09:54,051 --> 00:10:22,501
[Dr. Ada Shannon] Much bigger, and it depends on the baseline. Against PyTorch Dynamo with default Inductor, speedups run from 1.7x on MoE routing to 3.4x on Quant plus GEMM. Against TVM version 0.21, they run from 2.6x on the LLaMA attention to 12.1x on Quant plus GEMM. The authors turned off TVM's CUTLASS and FlashInfer backends to isolate its own fusion. That's a stated choice, and it also widens the gap. I'll come back to whether that's fair.

44
00:10:22,501 --> 00:10:28,001
[Hal Turing] Does the extra arithmetic from incremental mode ever pay for itself?

45
00:10:28,001 --> 00:10:51,076
[Dr. Ada Shannon] On BERT-base attention, yes. The non-incremental version only fits when each block handles short sequences. At matched parallelism it's actually faster. But the best configuration needs longer per-block sequences, and only incremental mode can reach it. That setting is up to 1.25x faster than the baseline configuration in the sweep. It's one workload on one GPU, and the comparison is within their own sweep, not against another system.

46
00:10:51,076 --> 00:10:54,601
[Hal Turing] What about fusing at different levels of the tree?

47
00:10:54,601 --> 00:11:15,376
[Dr. Ada Shannon] On safe softmax from 1K to 8K elements, every fusion level beats the unfused version. Block-level is fastest and thread-level is slowest. They also run variance and moment of inertia, plus an appendix on A100, H800 and an AMD MI308X, but those are reported only against PyTorch Eager.

48
00:11:15,376 --> 00:11:22,651
[Hal Turing] Those are encouraging results. Before we treat them as settled, which of the compiler numbers would you trust?

49
00:11:22,651 --> 00:11:53,751
[Dr. Ada Shannon] The Dynamo ones. The TVM ratios come from a default pipeline with CUTLASS and FlashInfer switched off. The authors say that isolates TVM's own fusion, but it also leaves TVM at its weakest. Dynamo is at default settings too, so there's no max-autotune mode and no hand-written Triton or TileLang kernels for these patterns. I'd read the 1.7x to 3.4x range as the more meaningful signal. The abstract's '2x to 5x' undersells the body, which runs from 1.7x to 12.1x.

50
00:11:53,751 --> 00:12:04,951
[Hal Turing] The backend is TileLang, which brings pipelining, warp specialization and Hopper's async copy instructions. How much of the gain is fusion and how much is TileLang?

51
00:12:04,951 --> 00:12:26,026
[Dr. Ada Shannon] The paper can't say. The experiment that would answer it is Inductor or TVM emitting through the same TileLang backend without fusion, against RedFuser on a naive backend. Neither is run. The one partial control is the safe-softmax study, where unfused versus fused isolates fusion. But that's small softmax kernels, not the end-to-end workloads.

52
00:12:26,026 --> 00:12:32,626
[Hal Turing] And the tuning? RedFuser searches tile size, threads, pipeline depth and segment count.

53
00:12:32,626 --> 00:12:41,226
[Dr. Ada Shannon] Whether the baselines got comparable tuning effort isn't stated, and tuning time isn't reported. I'd mark both as unknown.

54
00:12:41,226 --> 00:12:49,326
[Hal Turing] The hand-written comparison was FlashAttention-2 on an A10. Does that support 'matches hand-written kernels'?

55
00:12:49,326 --> 00:13:16,426
[Dr. Ada Shannon] It supports matching FlashAttention-2 on Ampere, plus FlashMLA on Hopper for decode shapes, where there's one query. The Hopper reference for dense attention is FlashAttention-3, from Jay Shah and colleagues in 2024, which uses Hopper's asynchrony and adds FP8 attention. It isn't compared. FlashInfer, from Ye and colleagues in 2025, is a programmable attention template system, an alternative route to similar coverage. Here it appears only as a TVM backend that was turned off.

56
00:13:16,426 --> 00:13:21,826
[Hal Turing] What about the claim that FlashAttention just falls out as a special case?

57
00:13:21,826 --> 00:13:41,526
[Dr. Ada Shannon] Of the algebra, yes, and the paper itself says FlashAttention and FlashDecoding emerge as special cases. That's why the defensible contribution is generalization and automation, to routing and quantization cascades, not the rescaling idea, which the two earlier papers already gave and RedFuser doesn't cite.

58
00:13:41,526 --> 00:13:48,451
[Hal Turing] Then the fixed-point check worries me. Passing at one point doesn't prove the function factors everywhere.

59
00:13:48,451 --> 00:14:06,526
[Dr. Ada Shannon] Fair concern. The paper doesn't spell out what guards against an expression that coincides at the test point. The operator set is also narrow: sum, product, max, min. Top-k doesn't obviously fit a scalar running state, so I can't tell from the text exactly what is fused in the routing benchmark.

60
00:14:06,526 --> 00:14:13,526
[Hal Turing] And the FP8 case. Does the fused kernel round to FP8 the way a real quantizer would?

61
00:14:13,526 --> 00:14:31,126
[Dr. Ada Shannon] The main text doesn't say, and that's the biggest gap. The case study is real-number algebra. There are no accuracy or error results in the main text for the rescaled kernels or for Quant plus GEMM. MoE routing and Quant plus GEMM also have no hand-written baseline.

62
00:14:31,126 --> 00:14:34,176
[Hal Turing] Anything else the authors concede?

63
00:14:34,176 --> 00:14:54,026
[Dr. Ada Shannon] Yes. Fusion adds compute and register pressure and can be suboptimal, and they plan a cost model for when to fuse. Their main experiments are on two NVIDIA GPUs, an A10 and an H800. The AMD result is in the appendix, against Eager only. So 'AI accelerators' is still mostly a GPU story.

64
00:14:54,026 --> 00:14:56,276
[Hal Turing] So who should care today?

65
00:14:56,276 --> 00:15:16,926
[Dr. Ada Shannon] Kernel authors facing a new attention variant or routing pattern. It's a plausible way to get a fused kernel without a hand derivation, and decode is a good fit because parallelism is scarce there. Treat it as promising, not proven against tuned alternatives. The missing pieces are the ablation, the cost model, relaxed algebraic conditions, and more backends.

66
00:15:16,926 --> 00:15:19,151
[Hal Turing] So what's established?

67
00:15:19,151 --> 00:15:34,851
[Dr. Ada Shannon] The math generalizes the online-softmax trick, and the automated pipeline lands near FlashAttention-2 and FlashMLA on those tests. The boundary is that the large compiler speedups can't yet be attributed to the fusion method itself.

68
00:15:34,851 --> 00:15:44,576
[Hal Turing] That's the paper, then: a sound idea that is partly known and partly new, with the strongest evidence in the parity results. Thanks for listening.