Compilers and program analysis: SSA, dataflow, and the optimizations that matter for ML

Why this subject matters now

For thirty years a compiler was something a systems programmer used and never thought about: type in C, get an object file, trust that -O2 did something reasonable. The optimizations that ran inside were mature by the late 1990s, and the interesting work had moved to the hardware. Machine learning reopened the field. A modern training or inference stack is a compiler stack: a PyTorch program is traced to a graph, the graph is lowered through several intermediate representations, fused into a handful of kernels, tiled to the cache and register hierarchy, and finally emitted as PTX or a hardware-specific binary. The people who build these systems, Google's XLA and MLIR teams, Meta's TorchInductor team, OpenAI's Triton team, are doing classical compiler construction with tensors as the primitive value instead of 64-bit integers.

The reason is economic in exactly the way it is for the kernels on the parallel computing page. A transformer's arithmetic is a small number of operation types repeated billions of times, and the gap between a naive lowering and a well-fused, well-tiled one is a factor of several in wall-clock time and therefore in dollars. In this repository, fusing a chain of elementwise operations that were memory-bound cut a measured kernel from 0.911 ms to 0.214 ms on an NVIDIA H100 80GB, a \(4.26\times\) speedup with no change to the arithmetic. That decision, whether to fuse and where to cut the graph into kernels, is made by a compiler pass whose correctness rests on a dependence analysis that is a hundred-year-old idea about affine inequalities. The practitioner is now expected to read an HLO or Inductor dump, understand why a fusion did or did not happen, and reason about the transformation in the vocabulary below.

This page is deliberately scoped to the analyses and transformations that recur, front end through back end, with the tensor-compiler payoff kept in view. It does not teach language design or type theory; those live on the companion programming languages page, which owns lambda calculus, Hindley-Milner inference, and operational semantics. It leans on the machine model developed in advanced systems architecture whenever a transformation is justified by the memory hierarchy. The anchors are the standard texts: Aho, Lam, Sethi, and Ullman's Dragon Book, Muchnick, Appel, and Cooper and Torczon, plus the primary papers for SSA, dataflow, register allocation, and the polyhedral model.

The front end, quickly

The front end turns text into a typed abstract syntax tree, and it is the part of a compiler with the cleanest theory and the least remaining research. Two layers do the work. Lexing partitions the character stream into tokens, and the token classes are exactly the languages a finite automaton can recognize: the regular languages. An identifier, a number, a keyword, an operator, each is described by a regular expression, the regular expressions are compiled to a single deterministic finite automaton by the Thompson construction followed by the subset construction, and the DFA runs in one pass, one table lookup per character. The reason lexers are generated rather than written is that this construction is mechanical and easy to get subtly wrong by hand; tools like lex, flex, and re2c emit the transition table from the regular expressions.

Parsing recognizes the nested structure that regular languages cannot express. The canonical proof is that \(\{a^n b^n : n \ge 0\}\) is not regular, shown by the pumping lemma: a DFA with \(k\) states reading \(a^{k}b^{k}\) must repeat a state inside the run of \(a\)'s, and pumping that cycle produces a string with unequal counts that the DFA still accepts, a contradiction. Balanced brackets and nested expressions have exactly this shape, so parsing needs a stack, which is a pushdown automaton, which recognizes the context-free languages. Grammars are written in Backus-Naur form and parsed by one of two families. Top-down LL(k) parsers predict the production from \(k\) tokens of lookahead and are what hand-written recursive-descent parsers implement; they cannot handle left recursion and must factor the grammar. Bottom-up LR(k) parsers, and the practical LALR(1) subset produced by yacc and bison, build the parse from the leaves using a shift-reduce automaton over item sets, and they accept a strictly larger class of grammars. Parser generators exist because constructing the LALR(1) automaton, computing FIRST and FOLLOW sets and detecting conflicts, is again mechanical and error-prone by hand. The output of the front end, for the rest of this page, is a control-flow graph of an intermediate representation, and everything interesting happens there.

Intermediate representation and SSA

Why an IR, and why SSA

An intermediate representation is a small, regular instruction set that is neither the source language nor the target machine. It exists so that the optimizer can be written once against a fixed vocabulary and then retargeted, and so that analyses do not have to cope with the irregularity of real source syntax. The dominant modern form is a three-address, register-based IR arranged as a control-flow graph: basic blocks of straight-line instructions, connected by edges wherever control can flow. LLVM IR, described by Lattner and Adve in 2004, is the reference example and the one most ML compilers ultimately reach.

The single most important normal form on top of that IR is static single assignment. In SSA form every variable is assigned exactly once, textually, and every use refers to exactly one definition. A variable reassigned three times in the source becomes three differently-named versions. This makes def-use chains explicit and free: to find the definition reaching a use you read its name, no analysis required. Constant propagation, value numbering, dead-code elimination, and register allocation all become dramatically simpler because the question "which definition does this use see" has a one-word answer. The difficulty is control-flow merges. If \(x\) is assigned in one arm of an if and differently in the other, the join block cannot name a single reaching definition. SSA resolves this with the phi-function: at a merge block, \(x_3 = \phi(x_1, x_2)\) selects \(x_1\) or \(x_2\) according to which predecessor edge was taken. The whole problem of constructing SSA reduces to one question: at which blocks, for which variables, must a phi-function be inserted. Placing a phi at every join for every variable is correct but wasteful; the minimal answer is computed from dominance.

Dominance

Fix a CFG with a unique entry node. A node \(d\) dominates a node \(n\), written \(d \operatorname{dom} n\), if every path from entry to \(n\) passes through \(d\). Dominance is reflexive (\(n\) dominates itself), transitive, and antisymmetric, so it is a partial order; restricted to the nodes reachable from entry it is in fact a tree. The immediate dominator \(\operatorname{idom}(n)\) is the unique dominator of \(n\) other than \(n\) itself that is dominated by every other dominator of \(n\); it is the parent of \(n\) in the dominator tree. Dominance is the answer to "what must have already run before control can be here," which is why it controls both phi placement and code motion.

Dominance is itself a dataflow problem, computed as a fixed point. Let \(D(n)\) be the set of dominators of \(n\). Then \(D(\text{entry}) = \{\text{entry}\}\) and, for every other node,

The equation says a node's dominators are itself plus everything that dominates all of its predecessors. Initializing every non-entry \(D(n)\) to the full node set and iterating this equation to convergence gives the dominator sets; on a reverse-postorder traversal it converges in a small number of passes. For the worked CFG below, initialize and iterate.

          B0            edges:  B0->B1
           |                    B1->B2   B1->B3
          B1 <-----+            B2->B4   B3->B4
         /    \     |            B4->B5   B4->B6
        B2    B3    |            B5->B1   (back edge)
         \    /     |            B6  exit
          B4        |
         /    \     |
        B5     B6   |
        |           |
        +-----------+

Iterating the dominance equation in reverse postorder converges in two passes (one productive pass and one confirming pass) to

Notice that B4 is dominated by B1 and not by B2 or B3, because control can reach B4 through either arm of the branch, so neither arm is on every path. The dominator tree therefore has B1 as the parent of B2, B3, and B4, and B4 as the parent of B5 and B6. In production, nobody iterates the set equation; the Lengauer-Tarjan algorithm computes the dominator tree in near-linear time, and the Cooper-Harvey-Kennedy iterative formulation is the one most compilers ship because it is simple and fast enough. The set equation is the definition; the algorithms are engineering.

Dominance frontiers, derived

The dominance frontier of a node \(d\) is the set of nodes where \(d\)'s dominance "just stops." Formally, \(n \in \operatorname{DF}(d)\) if \(d\) dominates a predecessor of \(n\) but does not strictly dominate \(n\) itself:

where \(\operatorname{sdom}\) is strict dominance (\(d \operatorname{dom} n\) and \(d \ne n\)). The intuition is precisely the phi-placement condition. If a value is defined in \(d\), that definition dominates everything in the subtree below \(d\), so within the subtree every use sees it and no phi is needed. But at a node \(n\) on the frontier, control can arrive both through \(d\) (carrying \(d\)'s definition) and through some other path (carrying a different one), so the definitions merge and \(n\) needs a phi. The dominance frontier is exactly the set of merge points a single definition can reach without dominating.

Cytron, Ferrante, Rosen, Wegman, and Zadeck gave the algorithm in 1991. Compute \(\operatorname{DF}\) bottom-up on the dominator tree. For each node \(b\) with two or more predecessors, walk from each predecessor \(p\) up the dominator tree, adding \(b\) to the frontier of every node visited, stopping just before \(\operatorname{idom}(b)\):

for each join node b (|pred(b)| >= 2):
    for each p in pred(b):
        runner = p
        while runner != idom(b):
            DF[runner] = DF[runner] union {b}
            runner = idom(runner)

Running this on the CFG above, the join nodes are B4 (predecessors B2, B3) and B1 (predecessors B0, B5, since B5 has the back edge). For B4: from B2, runner=B2, \(\operatorname{idom}(B4)=B1\), so add B4 to \(\operatorname{DF}(B2)\), move to B1, stop. From B3 symmetrically, add B4 to \(\operatorname{DF}(B3)\). For B1: from B5, runner walks B5, B4 up to \(\operatorname{idom}(B1)=B0\), adding B1 to the frontier of B5 and of B4; from B0 the loop does not run. The result, confirmed in Python:

The loop header B1 is in its own dominance frontier, which is the signature of a loop: a value defined inside the loop reaches the header again through the back edge, so the header needs a phi. This is why loop induction variables always appear as phi-nodes at the header in SSA.

Phi placement by iterated dominance frontier

A single definition of \(x\) in block \(d\) forces phi-nodes at every node in \(\operatorname{DF}(d)\). But a phi-node is itself a definition of \(x\), so it can force further phi-nodes at its dominance frontier. The correct set of phi locations for a variable defined in a set of blocks \(S\) is therefore the iterated dominance frontier \(\operatorname{DF}^+(S)\), the least fixed point of \(\operatorname{DF}\) starting from \(S\):

Suppose \(x\) is assigned in B0, B2, and B5 (an initialization, an update in one branch, and an update at the loop tail). Take the union of their frontiers: \(\operatorname{DF}(B0)=\{\}\), \(\operatorname{DF}(B2)=\{B4\}\), \(\operatorname{DF}(B5)=\{B1\}\), giving \(\{B1, B4\}\). Now iterate: B1 and B4 are themselves definitions of \(x\) (they hold phis), and \(\operatorname{DF}(B1)=\{B1\}\), \(\operatorname{DF}(B4)=\{B1\}\), which adds nothing new. The fixed point is \(\{B1, B4\}\): phi-nodes for \(x\) go at the loop header B1 and at the branch merge B4, and nowhere else. The Python computation agrees exactly. After placement, a single pre-order walk of the dominator tree renames variables into versions and wires each phi's operands to the versions live along each incoming edge, completing SSA construction.

Dataflow analysis as a lattice fixed point

The framework

Almost every classical analysis, liveness, reaching definitions, available expressions, constant propagation, is an instance of one framework, first stated in this generality by Kildall in 1973 and refined by Kam and Ullman in 1976 and 1977. The framework has four ingredients. A lattice \((L, \sqsubseteq)\) of dataflow facts, with a meet operation \(\sqcap\) that combines facts arriving from different paths. A transfer function \(f_B : L \to L\) for each basic block, describing how the block changes the facts. A direction, forward or backward. And a starting value at the boundary. The analysis computes, for each program point, the meet over the transfer functions along all paths reaching it, as a fixed point of the system of equations

for a forward analysis, and the mirror image (successors, IN and OUT swapped) for a backward one. The meet \(\sqcap\) is intersection or union depending on whether the analysis wants facts true on all paths or on some path; this is the choice that distinguishes "must" analyses from "may" analyses.

Three properties make this well-defined and computable. First, \(L\) must be a complete lattice, so the meet of any set of facts exists. Second, every transfer function must be monotone: if \(x \sqsubseteq y\) then \(f(x) \sqsubseteq f(y)\); more information in produces no less information out. Third, \(L\) must have finite height: no infinite strictly-descending chain \(x_0 \sqsupset x_1 \sqsupset \cdots\). Under these three conditions the iterative solver terminates, for a reason worth stating precisely.

Termination. Initialize every OUT to the lattice top \(\top\) (the most optimistic fact) and apply the equations repeatedly. Each application can only move a value down or leave it fixed, by monotonicity, so every program point traces a monotonically descending chain in \(L\). A descending chain in a finite-height lattice cannot descend more than \(h\) times, where \(h\) is the height. With \(n\) program points the whole system therefore stabilizes after at most \(n \cdot h\) value changes, and once no equation changes a value the system is at a fixed point. Because the iteration always decreases and is bounded below, it reaches the greatest fixed point consistent with the equations, which is the most precise sound solution the framework can express. This is the Knaster-Tarski theorem specialized to a finite lattice, and it is the entire reason dataflow analysis is guaranteed to work.

Reaching definitions, worked to fixpoint

A definition \(d\) (an assignment to a variable) reaches a program point if there is a path from \(d\) to that point along which the variable is not reassigned. It is a forward, may analysis: the fact set is a set of definitions, the meet is union (a definition reaches if it reaches along any path), and the block transfer function is

where \(\operatorname{gen}[B]\) is the definitions made in \(B\) and reaching its end, and \(\operatorname{kill}[B]\) is all other definitions of the same variables, which \(B\) overwrites. Consider this CFG, a branch that redefines the same variables on each arm:

B1: d1: a = ...        gen={d1}  kill={d3}
     |
B2: d2: b = ...        gen={d2}  kill={d4}
    /  \
   /    \
B3:d3   B4:d4          B3 gen={d3} kill={d1}
 a=..    b=..          B4 gen={d4} kill={d2}
   \    /
    \  /
B5: (use a, b)         gen={}   kill={}

edges: B1->B2, B2->B3, B2->B4, B3->B5, B4->B5

Here d1 and d3 both define \(a\), so each kills the other; d2 and d4 both define \(b\), likewise. Iterating the forward equations with meet = union, in block order, starting from empty sets:

The system reaches its fixed point in a single productive pass here (a second pass confirms nothing changes) because the CFG is acyclic and the blocks were visited in topological order. Read B3's OUT: d1 was killed (B3 redefines \(a\)), so only d2 and the new d3 survive. At the join B5, the union brings in all four definitions, but note that both a-definitions (d1 and d3) reach B5: along the B3 arm \(a\) is d3, along the B4 arm \(a\) is still d1. That ambiguity, two definitions of one variable reaching one point, is precisely the situation SSA eliminates with a phi at B5. Reaching definitions and phi placement are the same fact viewed two ways.

Liveness, worked to fixpoint

A variable is live at a point if its current value may be read before it is overwritten. It is a backward, may analysis: facts are sets of variables, meet is union over successors, and the transfer function is

where \(\operatorname{use}[B]\) is the variables read in \(B\) before any redefinition in \(B\) (upward-exposed uses) and \(\operatorname{def}[B]\) is the variables \(B\) assigns. Liveness is the analysis register allocation runs on: two variables that are never simultaneously live can share a register. Consider an accumulator loop, where \(c\) enters as a parameter:

B1: a = 0                use={}      def={a}
     |
B2: b = a + 1            (loop head)
    c = c + b            use={a,c}   def={a,b,c}
    a = b * 2
    if a < N goto B2 else B3
     |
B3: return c             use={c}     def={}

edges: B1->B2, B2->B2 (loop back), B2->B3

The use set of B2 is \(\{a, c\}\): \(a\) is read by b = a + 1 before B2 redefines \(a\), and \(c\) is read by c = c + b before B2 redefines \(c\); \(b\) is defined in B2 before its use, so it is not upward-exposed. Iterating the backward equations, processing blocks in reverse order B3, B2, B1, starting from empty:

Pass 1 processes B2 before its own loop-back fact has stabilized: at that point \(\operatorname{IN}[B2]\) has not yet been fed back into \(\operatorname{OUT}[B2]\), so \(\operatorname{OUT}[B2]=\{c\}\) (only the B3 successor contributes). Pass 2 picks up the loop edge: now \(\operatorname{OUT}[B2] = \operatorname{IN}[B2] \cup \operatorname{IN}[B3] = \{a,c\}\), because \(a\) computed in one iteration is read by the next. A third pass changes nothing, so the fixed point is reached in two productive passes. The back edge is exactly why more than one pass is needed: an acyclic CFG converges in a single reverse pass, a loop needs one extra pass per level of nesting to propagate a fact around the cycle. Reading the result, \(c\) is live throughout (it is the accumulator, correctly live on entry as a parameter), and \(a\) and \(b\) are live only across short ranges, so they could share physical storage.

The worklist algorithm

Re-evaluating every block on every pass, the "round-robin" solver above, is wasteful: most blocks' facts do not change on most passes. The worklist algorithm only re-evaluates a block when an input it depends on has changed. Seed a worklist with all blocks; pop a block, recompute its OUT (forward) or IN (backward), and if it changed, push the neighbors that read it. Terminate when the worklist is empty.

worklist = all blocks
while worklist not empty:
    B = worklist.pop()
    in  = meet over predecessors' out    (forward)
    out = transfer(B, in)
    if out changed:
        for S in successors(B):
            worklist.push(S)

It computes the same fixed point as round-robin, because it applies the same monotone equations until none changes; it just skips the applications that would not change anything. Termination is the same argument: each block's value only descends, the lattice has finite height, so only finitely many pushes cause a change and the queue drains. Visiting order affects speed but not the answer; reverse postorder for forward problems and postorder for backward ones minimize the number of re-evaluations. The runnable analyzer in the implementation section is a worklist solver, and its output on the reaching-definitions CFG above matches the table exactly.

MOP versus the fixpoint solution

There are two things one might mean by "the dataflow answer." The meet-over-all-paths (MOP) solution at a point is the meet, over every executable path \(P\) from entry to that point, of the composed transfer functions along \(P\) applied to the boundary fact:

The MOP is what we actually want: the meet of the facts that hold on each real path. But the number of paths through a CFG with loops is infinite, so MOP is not directly computable. The iterative solver computes the maximum fixed point (MFP) instead, which merges facts at every join before continuing rather than keeping paths separate. Kam and Ullman proved the relationship: MFP \(\sqsubseteq\) MOP always (the fixpoint is a sound, possibly conservative, approximation of the meet-over-paths), and MFP \(=\) MOP exactly when all transfer functions are distributive, meaning \(f(x \sqcap y) = f(x) \sqcap f(y)\). The classic bit-vector analyses, liveness, reaching definitions, available expressions, are all distributive, so for them the iterative solution is exactly the meet-over-paths, no precision is lost. Constant propagation is the standard example that is monotone but not distributive, so its iterative solution can be strictly weaker than MOP, which is the subtlety the next section confronts.

Constant propagation and SCCP

The lattice and the non-distributivity

Constant propagation asks which variables hold a known constant at each point. The per-variable lattice has three levels: \(\top\) meaning "not yet known to be anything" (optimistically assumed constant), a middle layer with one element per concrete constant \(c\), and \(\bot\) meaning "known to vary" (not a constant). The order is \(\top \sqsupset c \sqsupset \bot\), and the meet of two different constants is \(\bot\): if two paths deliver 3 and 5, the value is not constant. This lattice has height 2, so it satisfies the finite-height requirement, and the transfer functions (evaluate the right-hand side when all inputs are constant, otherwise \(\bot\)) are monotone.

It is not distributive. Take z = x + y where one path supplies \((x,y)=(1,4)\) and another \((x,y)=(4,1)\). On each path \(z = 5\), so the meet-over-paths says \(z\) is the constant 5. But the iterative solver meets the inputs first: \(x = 1 \sqcap 4 = \bot\), \(y = 4 \sqcap 1 = \bot\), and then \(z = \bot + \bot = \bot\), losing the constant. This is the price of merging at joins, and it is why plain iterative constant propagation misses correlated constants. It is still sound, it never claims a non-constant is constant, just conservative.

Sparse conditional constant propagation

Wegman and Zadeck's 1991 sparse conditional constant propagation (SCCP) is strictly more powerful than the naive analysis, not by fixing non-distributivity but by exploiting SSA and by tracking reachability. It runs on SSA form with two interacting worklists, one for CFG edges and one for SSA def-use edges, and it maintains a lattice value per SSA variable and a reachable flag per CFG edge. Two ideas give it its power. First, because it is on SSA, def-use edges are explicit and it only re-evaluates the instructions that actually read a changed value, which is the "sparse" part and makes it fast. Second, it evaluates branch conditions with the current lattice values, and if a condition is a known constant it marks only the taken edge reachable and never propagates facts along the other. This is the "conditional" part, and it is what lets SCCP fold away entire dead branches: if if (false) is discovered, the then-arm's assignments never contribute to any phi, so constants survive the merge that a path-insensitive analysis would have destroyed. SCCP simultaneously discovers constants and unreachable code, and each discovery feeds the other; it is the constant-propagation pass in production compilers for exactly this reason.

Loops and the transformations that matter

Natural loops from dominance

Loops are found from the dominator tree, not from source syntax, so that the optimizer handles goto-built loops and reducible control flow uniformly. A back edge is a CFG edge \(n \to h\) whose head \(h\) dominates its tail \(n\). The natural loop of that back edge is \(h\) together with every node that can reach \(n\) without passing through \(h\); \(h\) is the loop header, the single entry through which all iterations pass. In the first worked CFG, B5 → B1 is a back edge because B1 dominates B5, and the natural loop is \(\{B1, B2, B3, B4, B5\}\), header B1. Identifying the header matters because it is where loop-invariant code is hoisted to (into a preheader inserted just before it) and where induction-variable phis live. When every back edge's head dominates its tail the CFG is reducible, which real structured code always is, and the loops nest cleanly; irreducible CFGs (from unrestricted gotos) need node splitting first.

Loop-invariant code motion

An instruction inside a loop is loop-invariant if every operand is either defined outside the loop or is itself loop-invariant, and its value therefore does not change across iterations. LICM hoists such an instruction to the preheader so it runs once instead of once per iteration. Correctness has a subtlety: hoisting is only safe if the instruction dominates all loop exits (so it would have executed anyway on every path that leaves the loop) or is otherwise known to be side-effect-free and non-faulting, because hoisting a faulting operation, a division or a load, out of a loop that might execute zero times would introduce a fault the original never had. In SSA the invariance test is nearly free: an operation is invariant exactly when all its SSA operands are defined outside the loop or by already-hoisted instructions, which is a simple fixpoint over the loop body. The payoff on tensor code is large because address computations, base + i*stride with a loop-invariant base and stride, are the most common hoisting target and feed directly into the next transformation.

Induction variables and strength reduction

A basic induction variable is one incremented by a loop-invariant amount each iteration, the loop counter i += 1 being the archetype. A derived induction variable is an affine function of a basic one, \(j = c_1 \cdot i + c_2\) with loop-invariant \(c_1, c_2\). Induction-variable analysis finds these families, and strength reduction replaces the expensive recomputation \(j = c_1 \cdot i + c_2\) with a cheap running update \(j \mathrel{+}= c_1\) each iteration, turning a multiply into an add. This is the transformation that makes array-address arithmetic cheap: an access A[i] at byte address \(\text{base} + i \cdot \text{elt}\) becomes a pointer that advances by \(\text{elt}\) per iteration. Concretely, with j = 4*i + 7 and i stepping by 3, the derived variable starts at \(4\cdot 2 + 7 = 15\) (for \(i_0 = 2\)) and increments by \(4 \cdot 3 = 12\) each iteration, so the multiply-add per iteration collapses to a single add of 12. Trip-count analysis, counting iterations, follows from the same affine description: for for (i = 2; i < 100; i += 3) the count is \(\lceil (100 - 2)/3 \rceil = 33\) and the final value is \(2 + 3 \cdot 32 = 98\). Trip counts drive unrolling and vectorization decisions.

Tiling, fusion, and unrolling for tensors

Three loop transformations dominate tensor compilation. Unrolling replicates the loop body to expose instruction-level parallelism and amortize the loop overhead, at the cost of code size and register pressure. Fusion merges two loops with compatible iteration spaces into one, so a value produced by the first is consumed by the second while still in a register or cache rather than being written to and reread from memory; this is the graph-level decision that produced the measured \(4.26\times\) elementwise speedup on the H100. Tiling (also called blocking) splits a loop over a large index into an outer loop over tiles and an inner loop within a tile, so that the working set of the inner loop fits in a fast level of the memory hierarchy and is reused there instead of being streamed repeatedly from DRAM. Tiling is the single most important transformation for matrix and tensor code, and its benefit is a reuse computation worth doing exactly.

Take \(C = A B\) with \(A, B\) both \(N \times N\), \(N = 1024\), fp32 (4 bytes per element). The untiled triple loop, for each \((i,j)\) streaming a full row of \(A\) and column of \(B\), issues \(2N^3\) element loads of \(A\) and \(B\) combined: \(2 \cdot 1024^3 = 2.147 \times 10^9\) loads, which at 4 bytes each is 8.590 GB moved through the memory system for a problem whose inputs are only 8 MB. Every element of \(A\) is fetched \(N = 1024\) times. Tiling all three loops into \(T \times T\) blocks changes the reuse: an \(A\) tile is loaded once and reused across the \(N/T\) tiles of the \(j\) dimension, so each element of \(A\) is fetched \(N/T\) times instead of \(N\). The combined \(A\)-and-\(B\) load count drops to \(2 N^2 (N/T)\), a reduction of exactly \(T\times\):

The arithmetic intensity, useful flops per byte moved, rises in lockstep: the \(2N^3\) flops divided by the bytes moved climbs from 0.25 flop/byte untiled to 16 flop/byte at \(T = 64\). The footprint column is the constraint: three \(T \times T\) fp32 tiles must fit in the fast memory being targeted, and at \(T = 64\) that is 48 KB, exactly the shared-memory budget of an H100 streaming multiprocessor, which is not a coincidence: production GPU matmul tiles are sized to that budget. Larger tiles give more reuse but overflow the cache and stop helping, so the optimum is the largest tile whose footprint fits, precisely the reasoning behind the tiled kernels on the parallel computing page. Multi-level tiling (register tile inside cache tile inside a DRAM-blocked tile) applies the same computation at each level of the hierarchy.

Register allocation by graph coloring

The reduction to coloring

After optimization the IR still has unboundedly many virtual registers; the machine has \(k\) physical ones. Register allocation assigns each live value to a physical register or, failing that, to a memory slot (a spill). Chaitin's 1982 insight is that this is graph coloring. Build the interference graph: one node per value, an edge between two values that are simultaneously live and therefore cannot share a register. Liveness analysis, computed exactly as above, supplies the interferences: two values interfere if one is live at a point where the other is defined. A valid assignment of \(k\) registers is then a proper \(k\)-coloring of this graph, an assignment of one of \(k\) colors to each node so that adjacent nodes differ. Graph \(k\)-coloring is NP-complete in general, so compilers use a heuristic that is fast and usually optimal in practice.

Chaitin's algorithm and the k-coloring heuristic

The heuristic rests on one observation: a node with fewer than \(k\) neighbors can always be colored, whatever its neighbors get, because at most \(k-1\) colors are blocked and a \(k\)-th remains. So simplify: repeatedly remove a node of degree \(< k\) from the graph and push it on a stack, which lowers its neighbors' degrees and may expose more low-degree nodes. If the graph empties, pop the stack and give each node a color avoiding its already-colored neighbors; success is guaranteed by construction. If at some point every remaining node has degree \(\ge k\), pick a spill candidate (a high-degree, low-use node), optimistically push it anyway, and continue; on the way back, it might still find a free color (Briggs' improvement), and only if it genuinely cannot is it spilled to memory, its live range split by loads and stores, and the whole process repeated.

Work a concrete interference graph over values \(\{a, b, c, d, e\}\) with edges \(a\!-\!b\), \(a\!-\!c\), \(a\!-\!e\), \(b\!-\!c\), \(b\!-\!d\), \(c\!-\!d\), \(c\!-\!e\), \(d\!-\!e\). Degrees are \(a{:}3\), \(b{:}3\), \(c{:}4\), \(d{:}3\), \(e{:}3\). Try \(k = 3\).

interference graph (degrees):    a(3)  b(3)  c(4)  d(3)  e(3)

        a --- b
        |\   /|
        | \ / |
        |  c  |        c is adjacent to a,b,d,e
        | / \ |
        |/   \|
        e --- d

No node has degree \(< 3\) initially (the minimum degree is 3), so Chaitin's original algorithm would declare a spill immediately. Briggs' optimistic variant instead picks the highest-degree node, \(c\) (degree 4), and pushes it as a potential spill without committing. Removing \(c\) drops every other node to degree 2, all now trivially colorable: push \(a, b, d, e\) in turn. The push order is \([c, a, b, d, e]\) (computed in Python). Now pop and color, each node taking the lowest color not used by an already-colored neighbor: \(e \to 0\), \(d \to 1\), \(b \to 0\), \(a \to 1\), and finally \(c\), whose neighbors \(a, b, d, e\) hold colors \(\{1, 0, 1, 0\} = \{0, 1\}\), leaving color 2 free, so \(c \to 2\). The coloring \(\{a{:}1, b{:}0, c{:}2, d{:}1, e{:}0\}\) uses 3 colors and is legal on every edge, verified in Python. The teaching point is exactly the Chaitin-versus-Briggs distinction: the pessimistic algorithm spilled \(c\) needlessly, the optimistic one found that \(c\) fit after its neighbors were constrained. This graph is 3-colorable, and only optimistic coloring finds it without a spill.

Now try \(k = 2\). The same simplify loop finds no degree-\(<2\) node and pushes \(c\) then \(e\) as potential spills; on the way back \(c\) cannot be colored (its neighbors already use both colors 0 and 1), so \(c\) is a real spill. With two registers this graph genuinely cannot be colored, because it contains a triangle (\(a, b, c\) are mutually adjacent) and a triangle needs 3 colors; the algorithm correctly reports that no 2-coloring exists and spills. The chromatic number of an interference graph is a hard lower bound on registers, and any allocator confronted with fewer must spill; the art is spilling the cheapest value, typically the one used least inside the hottest loop.

Instruction scheduling

The last back-end analysis reorders instructions to keep a pipelined, multiple-issue machine busy without changing results. Two instructions may be swapped only if they are not dependent: a true (read-after-write) dependence, an anti (write-after-read) dependence, or an output (write-after-write) dependence forces order. These dependences form a directed acyclic graph within a basic block, edges weighted by the latency of the producing instruction, and scheduling is choosing a topological order that minimizes total cycles subject to the machine's issue width and functional-unit counts. Optimal scheduling of a general DAG is NP-hard, so compilers use list scheduling: maintain the set of ready instructions (all predecessors scheduled and their latencies elapsed), and each cycle greedily issue the ready instruction of highest priority, where priority is usually the length of the longest dependence path to the end of the block (the "critical path"), breaking ties toward instructions that unblock the most successors. List scheduling is not optimal but is within a small constant of optimal in practice and runs in near-linear time.

Scheduling and register allocation are in tension: aggressive reordering to fill pipeline slots lengthens live ranges and raises register pressure, which can force spills that cost more than the stalls the scheduling avoided. This phase-ordering conflict, schedule-then-allocate versus allocate-then-schedule, has no universally right answer, and production compilers run a prepass schedule, then allocation, then a postpass schedule to repair damage from spill code. For tensor code the scheduling that matters most is software pipelining of the innermost loop: overlapping the memory loads of iteration \(i+1\) with the arithmetic of iteration \(i\), which is exactly what a good GPU matmul does to hide HBM latency behind tensor-core math, and what Triton's pipeliner inserts automatically.

The polyhedral model

The loop transformations above, tiling, fusion, interchange, skewing, are ad-hoc when done one at a time. The polyhedral model, developed by Feautrier in the late 1980s and early 1990s and made practical by tools like Pluto (Bondhugula and colleagues, 2008), gives them a single algebraic foundation. It represents a loop nest not as syntax but as geometry. The set of iterations of a nest is its iteration domain, a polyhedron carved out of \(\mathbb{Z}^d\) by the affine loop bounds: a nest for i in 0..N, for j in 0..i is the integer points of \(\{(i, j) : 0 \le i < N,\; 0 \le j \le i\}\), a triangle. Each array access is an affine map from the iteration vector to a memory location. A dependence between two iterations, one writes a location a later one reads, is itself described by an affine relation, and the whole dependence structure is a union of polyhedra computable by integer linear programming.

A program transformation is then a schedule: an affine map assigning each iteration a logical execution time. Loop interchange, skewing, tiling, and fusion are all just different affine schedules over the same domain, and the model can search the space of legal schedules, those that preserve every dependence, meaning no iteration is moved before one it depends on, by solving a linear program. The payoff is that the model can find a tiling and fusion automatically, together, optimizing for locality and parallelism at once, rather than relying on a fixed pattern-matcher. This is why the polyhedral model underlies the most aggressive automatic parallelizers and appears inside ML compilers wherever the loop structure is affine, which for dense tensor operators it almost always is. Its limitation is equally clear: it requires affine bounds and affine accesses, so data-dependent control flow and indirect indexing (gather/scatter, sparse formats) fall outside it, which is why production stacks use it for the dense core and fall back to pattern-based rewrites elsewhere.

All of it, applied to tensor programs

A tensor compiler is the whole pipeline above with a tensor as the primitive value. The front end is a tracer that turns an eager PyTorch or JAX program into a graph; the IR is a graph of tensor operations; the analyses are shape and layout inference, dependence, and fusibility; the transformations are fusion, tiling, and layout assignment; the back end is code generation to PTX, LLVM IR, or a hardware ISA. Four systems, from three groups, define the current landscape.

XLA and HLO (Google) compile JAX and TensorFlow graphs. HLO (High-Level Optimizer IR) is the tensor IR; XLA runs algebraic simplification, common-subexpression elimination, layout assignment, and, centrally, operator fusion, deciding which elementwise and reduction operators to merge into a single kernel so intermediates never touch HBM. Its fusion heuristics are the graph-level analogue of the classical loop fusion above. TorchInductor (Meta), the default backend of torch.compile introduced with PyTorch 2, lowers captured FX graphs to a loop-level IR, performs fusion and tiling, and generates Triton for GPUs and C++/OpenMP for CPUs; Ansel and colleagues described it in 2024. Triton (OpenAI), from Tillet, Kung, and Cox in 2019, is a block-level programming language and compiler: the programmer writes a kernel over tiles and Triton's compiler does the register allocation, instruction scheduling, software pipelining, and shared-memory management, all the back-end passes above, automatically. MLIR (Google, led by Lattner and colleagues, 2021) is not a compiler but the infrastructure the others increasingly sit on: a framework for defining multiple IRs ("dialects") at different abstraction levels, from a high-level tensor dialect down through affine and vector dialects to LLVM, with reusable passes that lower progressively between them. The whole XLA and Triton stacks are being rebuilt on MLIR dialects.

The unifying observation is that nothing in these systems is new theory. Operator fusion is loop fusion; tiling a matmul in HLO is the same reuse computation done in the tiling table above; Triton's pipeliner is software pipelining; layout assignment is a dataflow analysis over a lattice of layouts; the affine dialect in MLIR is the polyhedral model. The novelty is entirely in the primitive value and the scale, and the measured payoff is real: the same fusion that classical compilers apply to scalar loops cut a memory-bound elementwise chain from 0.911 ms to 0.214 ms on an H100 in this repository, and on a matmul the tiling arithmetic above is what carries the kernel from tens to hundreds of measured TFLOPS.

Implementation

The first block is a complete, runnable worklist dataflow solver, instantiated for reaching definitions on the branch CFG worked above. It builds the predecessor map, iterates the monotone equations off a worklist until no OUT set changes, and prints the fixpoint. Running it reproduces the table exactly: B5 receives \(\{d1, d2, d3, d4\}\), the two definitions of \(a\) both reaching the join. The JAX tab shows the same solver written functionally with frozensets, to make the point that a dataflow solver is pure and has no framework dependency.

from collections import deque

# CFG: B1->B2, B2->{B3,B4}, B3->B5, B4->B5  (a branch that redefines a,b)
succ = {"B1": ["B2"], "B2": ["B3", "B4"], "B3": ["B5"], "B4": ["B5"], "B5": []}
pred = {b: [] for b in succ}
for u, vs in succ.items():
    for v in vs:
        pred[v].append(u)

# d1:a  d2:b  d3:a  d4:b  ; defs of the same variable kill each other
gen  = {"B1": {"d1"}, "B2": {"d2"}, "B3": {"d3"}, "B4": {"d4"}, "B5": set()}
kill = {"B1": {"d3"}, "B2": {"d4"}, "B3": {"d1"}, "B4": {"d2"}, "B5": set()}

def reaching_definitions(succ, pred, gen, kill):
    IN  = {b: set() for b in succ}
    OUT = {b: set() for b in succ}
    work = deque(succ)                       # seed with all blocks
    while work:
        b = work.popleft()
        IN[b] = set()
        for p in pred[b]:                    # meet = union (may analysis)
            IN[b] |= OUT[p]
        new_out = gen[b] | (IN[b] - kill[b]) # transfer function
        if new_out != OUT[b]:
            OUT[b] = new_out
            for s in succ[b]:                # only re-examine dependents
                work.append(s)
    return IN, OUT

IN, OUT = reaching_definitions(succ, pred, gen, kill)
for b in sorted(succ):
    print(b, "IN=", sorted(IN[b]), "OUT=", sorted(OUT[b]))
# B5 IN= ['d1', 'd2', 'd3', 'd4']  OUT= ['d1', 'd2', 'd3', 'd4']
# two definitions of 'a' (d1, d3) reach B5 -> SSA would insert a phi there

# The same fixpoint written functionally; no jax arrays are needed because a
# dataflow lattice is a set lattice, but the structure mirrors a jax fixpoint.
from functools import reduce

succ = {"B1": ["B2"], "B2": ["B3", "B4"], "B3": ["B5"], "B4": ["B5"], "B5": []}
pred = {b: [p for p in succ if b in succ[p]] for b in succ}
gen  = {"B1": {"d1"}, "B2": {"d2"}, "B3": {"d3"}, "B4": {"d4"}, "B5": set()}
kill = {"B1": {"d3"}, "B2": {"d4"}, "B3": {"d1"}, "B4": {"d2"}, "B5": set()}

def step(OUT):                               # one round-robin sweep
    new = {}
    for b in succ:
        inb = reduce(lambda a, p: a | OUT[p], pred[b], frozenset())
        new[b] = frozenset(gen[b]) | (inb - frozenset(kill[b]))
    return new

def fixpoint(f, x):                           # iterate to a fixed point
    while True:
        y = f(x)
        if y == x:
            return x
        x = y

OUT = fixpoint(step, {b: frozenset() for b in succ})
for b in sorted(succ):
    print(b, "OUT=", sorted(OUT[b]))
# B5 OUT= ['d1', 'd2', 'd3', 'd4']  -- identical fixpoint, reached monotonically

The second block computes dominators and dominance frontiers on the seven-node CFG from the SSA section, then places phi-nodes by iterated dominance frontier. It is the code that produced every dominance number quoted above; run it and it prints \(\operatorname{DF}(B1) = \{B1\}\) and phi-nodes at \(\{B1, B4\}\) for a variable defined in \(\{B0, B2, B5\}\).

succ = {0: [1], 1: [2, 3], 2: [4], 3: [4], 4: [5, 6], 5: [1], 6: []}
nodes = sorted(succ)
pred = {n: [] for n in nodes}
for u in nodes:
    for v in succ[u]:
        pred[v].append(u)

def dominators(nodes, pred, entry=0):
    alln = set(nodes)
    Dom = {n: set(alln) for n in nodes}
    Dom[entry] = {entry}
    changed = True
    while changed:                            # fixpoint of the dom equation
        changed = False
        for n in nodes:
            if n == entry:
                continue
            inter = set(alln)
            for p in pred[n]:
                inter &= Dom[p]
            new = {n} | inter
            if new != Dom[n]:
                Dom[n], changed = new, True
    return Dom

Dom = dominators(nodes, pred)
idom = {n: max(Dom[n] - {n}, key=lambda d: len(Dom[d]))
        for n in nodes if n != 0}

def dominance_frontiers(nodes, pred, idom):
    DF = {n: set() for n in nodes}
    for b in nodes:
        if len(pred[b]) < 2:
            continue
        for p in pred[b]:                     # Cytron et al. 1991
            runner = p
            while runner != idom.get(b, 0):
                DF[runner].add(b)
                if runner == 0:
                    break
                runner = idom[runner]
    return DF

DF = dominance_frontiers(nodes, pred, idom)

def place_phis(defsites, DF):                 # iterated dominance frontier
    phi, work = set(), set(defsites)
    while work:
        n = work.pop()
        for y in DF[n]:
            if y not in phi:
                phi.add(y)
                if y not in defsites:
                    work.add(y)
    return phi

print("DF(1) =", sorted(DF[1]))               # [1]  (loop header in own DF)
print("phis  =", sorted(place_phis({0, 2, 5}, DF)))  # [1, 4]

// Iterative dominators + Cytron dominance frontiers, LLVM-flavored but
// self-contained. Nodes are ints; entry is 0.
#include <vector>
#include <set>
#include <map>
using namespace std;

map<int, set<int>> dominators(const vector<vector<int>>& pred, int n) {
    set<int> all;
    for (int i = 0; i < n; ++i) all.insert(i);
    map<int, set<int>> Dom;
    for (int i = 0; i < n; ++i) Dom[i] = all;
    Dom[0] = {0};
    bool changed = true;
    while (changed) {                       // fixpoint of D(n)={n} U cap D(preds)
        changed = false;
        for (int nd = 1; nd < n; ++nd) {
            set<int> inter = all;
            for (int p : pred[nd]) {
                set<int> tmp;
                for (int x : inter)
                    if (Dom[p].count(x)) tmp.insert(x);
                inter = tmp;
            }
            set<int> nw = inter; nw.insert(nd);
            if (nw != Dom[nd]) { Dom[nd] = nw; changed = true; }
        }
    }
    return Dom;
}

// idom(n) = strict dominator with the largest dominator set (closest above n).
map<int,int> immediate_doms(map<int,set<int>>& Dom, int n) {
    map<int,int> idom;
    for (int nd = 1; nd < n; ++nd) {
        int best = -1; size_t bestsz = 0;
        for (int d : Dom[nd])
            if (d != nd && Dom[d].size() > bestsz) { best = d; bestsz = Dom[d].size(); }
        idom[nd] = best;
    }
    return idom;                            // walk idom up to build DF (Cytron)
}

The last block illustrates the fusion decision a tensor compiler makes, at the level a Triton kernel exposes. The unfused version applies three elementwise operations as three PyTorch calls, materializing two intermediate tensors in HBM; the fused Triton kernel does all three inside one program, reading the input once and writing the output once, so the intermediates never leave registers. This is precisely the loop fusion of the theory section, and it is the transformation that produced the measured \(4.26\times\) elementwise speedup on the H100 in this repository.

import triton
import triton.language as tl

@triton.jit
def fused_gelu_scale(x_ptr, out_ptr, scale, n, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)     # this block's element indices
    mask = offs < n
    x = tl.load(x_ptr + offs, mask=mask)         # one HBM read
    # three elementwise ops fused: all stay in registers, no intermediate in HBM
    y = x * scale                                # op 1: scale
    g = 0.5 * y * (1.0 + tl.math.erf(y * 0.70710678))  # op 2: gelu
    z = g * g                                    # op 3: square
    tl.store(out_ptr + offs, z, mask=mask)       # one HBM write
    # unfused this touches HBM 6 times (3 reads + 3 writes); fused, twice.

def run(x, scale):
    out = torch.empty_like(x)
    n = x.numel()
    grid = (triton.cdiv(n, 1024),)
    fused_gelu_scale[grid](x, out, scale, n, BLOCK=1024)
    return out

import torch
import torch.nn.functional as F

# Unfused: each op is a separate kernel that reads and writes all of HBM.
def run_unfused(x, scale):
    y = x * scale          # kernel 1: read x, write y   (y materialized in HBM)
    g = F.gelu(y)          # kernel 2: read y, write g   (g materialized in HBM)
    z = g * g              # kernel 3: read g, write z
    return z
    # 3 reads + 3 writes of an N-element tensor. For a memory-bound chain the
    # runtime is (bytes moved)/(bandwidth), so removing the 4 intermediate
    # touches is the entire speedup -- flops are unchanged.

# torch.compile fuses this automatically via TorchInductor, emitting one Triton
# kernel equivalent to the hand-written fused version:
run_fused = torch.compile(run_unfused)
# measured on an H100 80GB in this repo: 0.911 ms unfused -> 0.214 ms fused (4.26x)

How it is done in practice

The gap between the framework and a shipping compiler is mostly engineering discipline around the same ideas. Dominators are computed by Lengauer-Tarjan or the Cooper-Harvey-Kennedy iterative method, not by the set equation, because the set equation is quadratic in space. SSA is constructed once early and maintained incrementally through every subsequent pass, because reconstructing it is expensive; LLVM keeps values in SSA from the first optimization to just before register allocation, where it leaves SSA via a phi-elimination pass that inserts copies on the incoming edges. Dataflow analyses are run with bit-vectors, one bit per fact, so the meet is a machine AND or OR over words and a whole block's transfer is a couple of instructions; this is why the classical analyses are effectively free and run many times during a compilation.

Register allocation in production is more than textbook Chaitin. LLVM's default allocator is not graph coloring at all but a greedy linear-scan-derived allocator with live-range splitting, chosen because it is faster to run and its results are close; graph coloring survives in GCC and in JITs where compile time is less critical. Spilling is guided by loop-nesting depth: a value used inside a triply nested loop is spilled only as a last resort because each spill costs a memory access per iteration, multiplied by the trip counts. Instruction scheduling is often folded into the same phase to manage the pressure-versus-latency tension described above.

In the tensor world the same engineering shows up at graph scale. TorchInductor's fusion is a scheduling pass over a dependency graph of loop-level operations, with heuristics for when a fusion helps (memory-bound producers and consumers) versus hurts (fusing past a reduction, or into a kernel already register-starved). XLA's layout assignment is a dataflow analysis choosing tensor layouts to minimize transposes, a direct analogue of register-class assignment. Triton's compiler runs the full back end, an autotuner searches tile sizes and pipeline depths, and the winning configuration is cached, which is tiling and scheduling with an empirical cost model substituting for a static one. Across all of them the measured lesson is the one the tiling table makes precise: the transformation that moves a kernel from memory-bound to compute-bound is worth more than any local micro-optimization, and it is a compiler pass, not a hand edit.

The current research frontier

Four threads are active. The first is MLIR as a universal substrate: the Google-led effort (Lattner and colleagues, 2021) to express every level of a tensor compiler as a dialect and reuse lowering passes between them, now adopted well beyond Google, including by parts of the Triton and IREE stacks. The bet is that the historical fragmentation of ML compilers, each reinventing SSA, dataflow, and tiling, collapses onto one infrastructure. The second is autotuning and learned cost models: rather than a static machine model, systems like TVM (from the Washington group, Chen and colleagues) and Triton's autotuner search the transformation space and fit a cost model to measurements, trading compile time for kernel quality, an approach that is winning on hardware too irregular to model analytically.

The third is the polyhedral model meeting machine learning: extending affine scheduling to the fusion and tiling decisions of full networks, and coping with the non-affine parts (dynamic shapes, data-dependent control) that classical polyhedral compilation excludes; the tension between the model's optimality guarantees and its restricted applicability is the open problem. The fourth is verified and translation-validated optimization: as compilers grow, so does the risk that an optimization is subtly wrong, and lines of work on translation validation (checking after the fact that a specific compilation preserved semantics, as in the Alive tooling for LLVM peephole optimizations from the Utah and collaborators) and on fully verified compilers (CompCert, from Leroy and INRIA, a C compiler proved correct in Coq) are increasingly relevant as tensor compilers make ever more aggressive, ever less obviously-correct rewrites. The unifying question across all four is how to get the aggressiveness of an autotuned, fusing, tiling compiler with the correctness guarantees of the classical framework, whose soundness the lattice theory above actually establishes.

Open source to read

• llvm/llvm-project — the reference production compiler. Read llvm/lib/Analysis/ for the dataflow and dominator implementations and llvm/lib/Transforms/Scalar/ for SSA-based passes; SROA.cpp and SCCP.cpp are the concrete versions of the algorithms on this page.
• llvm/mlir — the multi-level IR infrastructure. Start with the Affine and Linalg dialects to see the polyhedral model and tensor tiling as first-class passes.
• openai/triton — the block-level kernel compiler. Read python/triton/language/ for the programming model and the lib/ passes for the automatic register allocation, scheduling, and pipelining.
• pytorch/pytorch — torch/_inductor/ is TorchInductor. Open scheduler.py for the fusion decisions and codegen/triton.py for the lowering to Triton.
• openxla/xla — the XLA compiler and HLO IR. The xla/service/ directory holds the fusion, layout-assignment, and algebraic-simplification passes.
• halide/Halide — the language that separated algorithm from schedule and inspired much of the tensor-compiler wave; the clearest place to internalize tiling and fusion as explicit, composable schedule directives.
• google/jax — the tracing front end and its lowering to XLA; jax/_src/interpreters/ shows how an eager program becomes a graph an optimizing compiler can consume.

Common misconceptions

"SSA is just renaming variables." Renaming is the easy half. The content of SSA is the phi placement, and placing phis minimally requires computing dominance frontiers and their iterated closure. A naive "phi at every join for every variable" is correct but bloats the IR by an order of magnitude and slows every later pass; the Cytron algorithm exists precisely to place the few phis that are actually needed.

"Dataflow analysis needs the program to halt or to be executed." It is a purely static fixed-point computation over a lattice, and it terminates by the finite-height-plus-monotonicity argument regardless of whether the analyzed program terminates. It never runs the program; it computes facts true of all runs at once, which is exactly why it can reason about infinite loops.

"The iterative solver gives the exact meet-over-all-paths answer." Only when the transfer functions are distributive, which the bit-vector analyses (liveness, reaching definitions, available expressions) happen to be. Constant propagation is monotone but not distributive, so its fixpoint solution can be strictly less precise than the meet-over-paths, which is the entire motivation for the smarter SCCP algorithm.

"More registers always means no spilling." Spilling is forced by the chromatic number of the interference graph, not by a shortage relative to some average. A single point where \(k+1\) values are simultaneously live forces a spill on a \(k\)-register machine no matter how few values the rest of the function uses, because that point contains a \((k{+}1)\)-clique. The remedy is live-range splitting to break the clique, not more registers.

"Fusion is always a win." Fusion removes intermediate memory traffic, which helps only while the operations are memory-bound. Fusing into a kernel that is already compute-bound adds register pressure and can force spills or cut occupancy; fusing across a reduction forces a grid-wide synchronization or recomputation. The measured \(4.26\times\) came from a chain that lived far on the memory-bound side; the same fusion applied to a saturated matmul would gain nothing. Fusion is a decision the cost model makes, not a universal good.

"Tiling helps because it reduces the number of operations." Tiling changes no arithmetic at all; the flop count is identical. It reduces the number of times each datum is fetched from slow memory, raising arithmetic intensity from 0.25 to tens of flop/byte in the worked example, which is a memory-traffic win, not a compute win. Confusing the two leads to wrong predictions everywhere, because it is the traffic reduction, capped by cache size, that sets the achievable speedup.

"The polyhedral model can optimize any loop nest." It requires affine loop bounds and affine array subscripts. Data-dependent bounds, indirect indexing (gather/scatter), and sparse formats fall outside it, which is why production tensor compilers use it for the dense affine core and fall back to pattern-based rewriting for everything else. Its strength (a complete algebraic search over legal schedules) and its limitation (only affine programs) are two sides of the same restriction.

"A tensor compiler is a fundamentally new kind of compiler." It is the classical pipeline with a tensor as the primitive value. Operator fusion is loop fusion, layout assignment is a dataflow analysis, the affine dialect is the polyhedral model, Triton's pipeliner is software pipelining, and its allocator is graph coloring or linear scan. The novelty is the value type and the scale, not the theory, which is exactly why the framework on this page transfers directly.

Self-check

References

1. Aho, A., Lam, M., Sethi, R., and Ullman, J. Compilers: Principles, Techniques, and Tools (the Dragon Book), 2nd ed., Addison-Wesley, 2006.
2. Muchnick, S. Advanced Compiler Design and Implementation, Morgan Kaufmann, 1997.
3. Appel, A. Modern Compiler Implementation in ML, Cambridge University Press, 1998.
4. Cooper, K. and Torczon, L. Engineering a Compiler, 2nd ed., Morgan Kaufmann, 2011.
5. Cytron, R., Ferrante, J., Rosen, B., Wegman, M., and Zadeck, F. K. "Efficiently Computing Static Single Assignment Form and the Control Dependence Graph," ACM TOPLAS 13(4), 1991. doi:10.1145/115372.115320
6. Kildall, G. "A Unified Approach to Global Program Optimization," POPL, 1973. doi:10.1145/512927.512945
7. Kam, J. and Ullman, J. "Global Data Flow Analysis and Iterative Algorithms," JACM 23(1), 1976. doi:10.1145/321921.321938
8. Kam, J. and Ullman, J. "Monotone Data Flow Analysis Frameworks," Acta Informatica 7, 1977. doi:10.1007/BF00290339
9. Wegman, M. and Zadeck, F. K. "Constant Propagation with Conditional Branches," ACM TOPLAS 13(2), 1991. doi:10.1145/103135.103136
10. Chaitin, G. "Register Allocation and Spilling via Graph Coloring," SIGPLAN Symposium on Compiler Construction, 1982. doi:10.1145/800230.806984
11. Briggs, P., Cooper, K., and Torczon, L. "Improvements to Graph Coloring Register Allocation," ACM TOPLAS 16(3), 1994. doi:10.1145/177492.177575
12. Lengauer, T. and Tarjan, R. "A Fast Algorithm for Finding Dominators in a Flowgraph," ACM TOPLAS 1(1), 1979. doi:10.1145/357062.357071
13. Cooper, K., Harvey, T., and Kennedy, K. "A Simple, Fast Dominance Algorithm," Rice University technical report, 2001. cs.rice.edu/~keith/EMBED/dom.pdf
14. Lattner, C. and Adve, V. "LLVM: A Compilation Framework for Lifelong Program Analysis & Transformation," CGO, 2004. doi:10.1109/CGO.2004.1281665
15. Feautrier, P. "Some Efficient Solutions to the Affine Scheduling Problem, Part I & II," International Journal of Parallel Programming 21, 1992. doi:10.1007/BF01407835
16. Bondhugula, U., Hartono, A., Ramanujam, J., and Sadayappan, P. "A Practical Automatic Polyhedral Parallelizer and Locality Optimizer" (Pluto), PLDI, 2008. doi:10.1145/1375581.1375595
17. Lattner, C., Amini, M., Bondhugula, U., Cohen, A., Davis, A., Pienaar, J., Riddle, R., Shpeisman, T., Vasilache, N., and Zinenko, O. "MLIR: Scaling Compiler Infrastructure for Domain Specific Computation," CGO, 2021. arXiv:2002.11054
18. Tillet, P., Kung, H. T., and Cox, D. "Triton: An Intermediate Language and Compiler for Tiled Neural Network Computations," MAPL, 2019. doi:10.1145/3315508.3329973
19. Ansel, J. et al. "PyTorch 2: Faster Machine Learning Through Dynamic Python Bytecode Transformation and Graph Compilation" (TorchInductor), ASPLOS, 2024. doi:10.1145/3620665.3640366
20. Chen, T. et al. "TVM: An Automated End-to-End Optimizing Compiler for Deep Learning," OSDI, 2018. arXiv:1802.04799
21. Ragan-Kelley, J., Barnes, C., Adams, A., Paris, S., Durand, F., and Amarasinghe, S. "Halide: A Language and Compiler for Optimizing Parallelism, Locality, and Recomputation in Image Processing Pipelines," PLDI, 2013. doi:10.1145/2491956.2462176
22. Leroy, X. "Formal Verification of a Realistic Compiler" (CompCert), CACM 52(7), 2009. doi:10.1145/1538788.1538814
23. Lopes, N., Menendez, D., Nagarakatte, S., and Regehr, J. "Provably Correct Peephole Optimizations with Alive," PLDI, 2015. doi:10.1145/2737924.2737965