the path · chapter 07 of 15 · part i, the descent

Mosaic → LLO

The TPU backend: where your kernel becomes tiled vector operations, and why you read it but never write it.

mastery work · this chapter0/1
manual items are your word; auto items complete from your streaks, labs, and can-you ticks · stored in your browser only
the layer

What happens below pallas_call

Pallas does not execute your kernel; it lowers it to Mosaic, an MLIR dialect for TPU kernels. Mosaic maps your block operations onto the chip's real units: matmuls onto the MXU's systolic array, elementwise work onto the VPU's lanes, and everything onto the (8, 128) register tiling. Below Mosaic sits LLO, the TPU's near-assembly, where operations become VLIW bundles the sequencer issues. You will never write either.

You will, however, read Mosaic's complaints. When a block shape violates the lattice, when a dynamic slice defeats vectorization, when VMEM overflows: the error surfaces from this layer, in this layer's vocabulary. The reason this track teaches the machine (stage 0) before the kernel language (stage 1) is exactly so those errors read as information rather than noise.

One design fact worth knowing: this layer is why TPU kernel scheduling is tractable at all. TPU grid steps are a sequential pipeline, not thousands of independent thread blocks, so questions like "will this transfer hide under this compute" have answers you can reason about from constants. The whole roofline discipline of stage 0 works because the hardware below is this predictable.

later on the path Stage 1's LAB·1.4 provokes and reads this layer's limits (chapter 11). Keep descending; the path arrives there in order.