HLO

An HloModule is the top of the tree: it holds one entry computation, the function actually called when the program runs, plus however many other computations get called from inside it. Each HloComputation is in turn a list of HloInstructions connected in something close to SSA form, every value defined exactly once and consumed by name wherever it is used.

Every one of those instructions carries its own shape and its own layout from the moment it is created, and both survive every pass that touches it. Some instructions carry more than a shape: a reduction needs to know how to combine two elements, so it carries an entire sub-computation of its own, a small nested function that XLAThe compiler: brilliant at fusing along dataflow edges, structurally unable to change your algorithm. That gap is why kernels exist.taught in /l/xla → calls once per reduction step. Those show up in a dump as separate %region blocks sitting right next to the instruction that owns them.

Optimization happens on HLO, nowhere else.

Every HLO instruction also carries a metadata field linking it back to the Python that produced it: an op_name describing the operation's place in the traced program, a source_file, a source_line. None of the two hundred passes that run between here and codegen strip this out. It survives fusionSeveral ops compiled into one kernel so intermediates stay in fast memory instead of round-tripping through HBM. XLA’s central optimization, with an exact limit.taught in /l/xla →, it survives layout assignment, it survives scheduling.

That survival is not incidental. It is the entire reason reading a dump is a viable debugging technique at all. Pick any instruction out of an optimized module, no matter how many passes have rewritten the graph around it, and its metadata still points at the exact line of Python that caused it to exist. Losing that thread anywhere along the pipeline would make every later chapter's dump-reading exercises useless.

Below is one actual reduction sub-computation from an optimized attention program, captured whole. Two parameters come in, a maximum gets taken, and the metadata on that maximum instruction still names reduce_max, still names attend.py, still names line 12, exactly as if no pass had ever touched the module.

%region_0.13 (Arg_0.14: f32[], Arg_1.15: f32[]) -> f32[] {
  %Arg_0.14 = f32[] parameter(0), metadata={op_name="jit(attend)/jit(main)/reduce_max"}
  %Arg_1.15 = f32[] parameter(1), metadata={op_name="jit(attend)/jit(main)/reduce_max"}
  ROOT %maximum.16 = f32[] maximum(f32[] %Arg_0.14, f32[] %Arg_1.15), metadata={op_name="jit(attend)/jit(main)/reduce_max" source_file="attend.py" source_line=12}
}

Everything above sets up the load-bearing point: later passes do not re-derive facts about the graph from scratch every time they need one. Two static analyses sit underneath the whole pass pipeline, and they answer different questions. Dataflow analysis asks what values can reach a given point in the program; fusionSeveral ops compiled into one kernel so intermediates stay in fast memory instead of round-tripping through HBM. XLA’s central optimization, with an exact limit.taught in /l/xla → depends on this, because deciding whether one instruction can merge into another requires knowing exactly which values flow between them.

Alias analysis asks a narrower and stricter question: which positions in the program are required to share the exact same physical buffer. Buffer assignment depends on this one, because it is the pass that actually hands out memory, and it cannot hand two required-aliased positions two different addresses without producing a wrong program. Fusion cares about value flow; buffer assignment cares about physical sharing. They sound similar and they answer genuinely different questions, which is exactly why XLAThe compiler: brilliant at fusing along dataflow edges, structurally unable to change your algorithm. That gap is why kernels exist.taught in /l/xla → keeps them as two separate analyses instead of one.