Kernel fusion

01 · The elementwise-chain tax

Every intermediate tensor is a full trip through DRAM

Consider the most common shape in a transformer's feed-forward path — a chain of elementwise ops: y = gelu(x) * scale + residual. Eager mode runs three kernels. The GELU reads x and writes a temporary. The multiply reads the temporary and writes another. The add reads that plus the residual and writes y. Count it in tensor-sized units: 2 + 2 + 3 = 7 tensor-trips through DRAM for math that, done in registers, needs 3 — read x, read residual, write y.

Nothing about the arithmetic is hard — GELU is a handful of FLOPs. The cost is that each intermediate tensor exists: 100M FP32 elements is 400 MB, and day 10's bandwidth charges ~0.8 ms for every 400 MB round trip at 504 GB/s. Fusion — doing the whole chain between one load and one store — deletes the round trips and, with them, the intermediates' VRAM footprint too (the same accounting as the VRAM budgeting playbook on this site).

02 · Why frameworks fuse

torch.compile is a fusion engine with a compiler attached

This is the load-bearing reason torch.compile and XLA speed up real models: they trace the op graph, find runs of elementwise (and reduction-shaped) ops, and emit one kernel per run — automatically doing what you're about to do by hand in Triton. The wins concentrate exactly where the calculator says: long chains of cheap elementwise math between expensive memory moves. When a compiled model's speedup disappoints, it's usually because its hot path was already fused — or because the bottleneck was reductions and GEMMs all along, which fusion alone can't restructure.

03 · Fusion in Triton, and its limits

Between tl.load and tl.store, registers are free real estate

Writing the fused version is almost anticlimactic — the chain just happens between the load and the store, in registers (day 09), without ever being named in memory:

@triton.jit
def fused_chain(x_ptr, r_ptr, y_ptr, n, scale, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    mask = offs < n

    x = tl.load(x_ptr + offs, mask=mask)
    t = tl.math.gelu(x) * scale   # lives in registers
    r = tl.load(r_ptr + offs, mask=mask)
    tl.store(y_ptr + offs, t + r, mask=mask)

The limits are structural, not syntactic. Fusing ops that need different loop structures — an elementwise op into a GEMM's inner loop, or two reductions over different axes — requires the tiling to be compatible, which is exactly the engineering inside FlashAttention (day 24) and inside fused attention kernels generally. Fusion also couples scheduling: one mega-kernel can't start until all its inputs are ready, so over-fusing can cost concurrency that separate kernels would have exploited. Day 11's discipline applies: fuse where the roofline says memory-bound and the calculator says round trips.

Next: day 24 — the Reading Real-World Kernels phase. Anatomy of FlashAttention: the kernel that made long context possible, built from ideas you've been collecting for weeks.