Expand description
The spatial IR as a PERSISTENT CUDA kernel (sm_80+).
A site is a resident thread block that loops over ITS work items — the same static placement every other
target runs, so no atomic work counter appears. A FIFO buffer is a shared-memory ring whose pages are
cp.async commit groups: Reserve names the next slot, a load is every thread’s share of the page in
16-byte cp.async copies, Push commits them as one group, Wait waits until nothing but the prefetched
stage is still in flight and then crosses a block barrier (a thread’s wait completes only its own copies;
the barrier publishes everyone’s), and Pop crosses a block barrier so no thread refills a slot another
still reads. The ring has no run-time flow control, so the lowering proves that no ring is filled past its
depth before it prints (ring_occupancy). A tile matmul is WMMA fragments held per warp across the whole
reduction. Every statement is the image of one choreography op, which is what the census asserts.