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.