Inductor: deciding what to fuse

Because the decisions interact, and greedy local choices are provably not optimal.

Consider a producer feeding two consumers. Fusing it into both means computing it twice — fine if it is cheap, wasteful if it is expensive. Fusing into neither means materialising it once. Fusing into one and not the other is sometimes best. That choice depends on the producer’s cost, the sizes, and what else is already in each consumer’s group — so it cannot be made by looking at any single edge.

Then the constraints couple the choices further. Iteration space: two ops fuse only if their output indices can be expressed in one loop nest, so a fusion may require a layout change, which itself costs. Registers: each op added to a group increases live values, and crossing the threshold spills, which is catastrophic and non-local — adding op $k$ can make ops $1..k{-}1$ slower. Reductions: fusing across one changes the loop structure entirely. Memory: materialising fewer intermediates lowers peak memory, which may be the binding constraint rather than speed.

So it is a partitioning problem over a DAG with non-linear costs, which is NP-hard in general. Inductor uses heuristics — fuse elementwise greedily, respect iteration-space compatibility, cap group size, estimate register pressure — and they are good on typical models and beatable on unusual ones.

The transferable point: when a compiler underperforms, the question is rarely “did it miss an optimisation” and usually “did its cost model mis-rank two legal choices”. Reading the generated code tells you which.

Because cuBLAS and cuTLASS represent an enormous amount of specialised effort that a general-purpose scheduler will not match, and matmul is the one op where the gap is decisive.

A production GEMM kernel encodes: tile sizes tuned per architecture and per shape; the three-level blocking from the hardware series; ldmatrix/mma tensor-core instruction selection; swizzled shared memory to avoid bank conflicts; software-pipelined cp.async stages; split-K for skinny problems; and a heuristic table mapping shapes to kernels built from exhaustive benchmarking. Every one of those is a decision Inductor’s general machinery would have to rediscover for each shape.

So Inductor does two more useful things instead. It fuses the epilogue — bias, activation, scaling, residual add — into the GEMM, which library calls alone cannot do and which is pure memory traffic in eager mode. And under max-autotune it benchmarks a Triton GEMM template against the library call and picks the winner per shape, which sometimes wins on unusual sizes the library’s heuristics handle poorly, especially skinny or oddly-aligned shapes.

The general principle is the pattern-matching one from Module 5: decompose to expose structure, then recognise the cases a human already solved better and defer to them. A compiler that insists on generating everything loses to a compiler that knows when to call a library — and knowing which is the actual engineering.

Rare but real, and the causes are enumerable.

1. Are you measuring steady state? Warm up past compilation, torch.cuda.synchronize() before timing. This explains most reports.
2. Is it recompiling? TORCH_LOGS=recompiles. Continuous recompilation is strictly worse than eager (Module 10).
3. Graph breaks in a loop? TORCH_LOGS=graph_breaks. Many small graphs pay compilation costs and keep eager’s overhead (Module 4).
4. Read the generated Triton. TORCH_LOGS=output_code. Look for a very large fused kernel — and check for register spilling, which Triton reports and which is usually the culprit when one kernel is unexpectedly slow.
5. Did layout changes cost more than fusion saved? Inductor may insert contiguous() to make a fusion legal; if the copy is bigger than the traffic saved, that fusion was a mistake.
6. Did a decomposition lose a good kernel? If a pattern match failed, a fused library op may have been replaced by a slower generated equivalent — attention is the classic case.

Then the escape hatches, in increasing order of effort: try mode="max-autotune"; mark the offending submodule with torch.compiler.disable and leave it eager; or write the kernel yourself in Triton and register it as a custom op.

The last of these is the point worth making about the whole design: because the output is Triton rather than an opaque binary, you can read it, understand the decision, and override it. That is a qualitatively better position than a black-box compiler.

It knows every buffer’s lifetime before anything runs, which turns allocation into planning.

Eager mode asks the caching allocator for each intermediate as it appears and returns it when the refcount drops (Module 8). The allocator is fast and good, but it is reacting to a stream of requests with no knowledge of the future. Inductor has the whole graph, so it can:

• Reuse buffers. Once an intermediate’s last consumer has run, its storage is available for a later tensor of compatible size — computed statically, so no allocator round trip at all.
• Not materialise fused intermediates. The largest saving is the tensors that never exist because they stayed in registers inside a fusion group.
• Order operations to lower peak memory. Among schedules that respect dataflow, some hold fewer tensors live simultaneously, and the peak is what determines whether you OOM.
• Choose recompute over storage, in concert with AOTAutograd’s partitioner (Module 5).

This is precisely the XLA-style compile-time planning that PyTorch’s dynamic allocator gave up in exchange for eager flexibility — and torch.compile recovers a large part of it within each graph, while the allocator still handles everything between graphs.

Which is the honest description of the whole system: it is not that dynamic allocation was wrong, it is that a graph is the scope in which planning is possible, and PyTorch spent years not having one.