Async Copy Fence

In a Pallas TPU kernel using manual memory management (an HBM-resident Ref with memory_space=pltpu.ANY), you issue copy = pltpu.make_async_copy(hbm_ref, vmem_ref, sem) to stage a tile into VMEM. What must happen before the kernel body reads from vmem_ref?

You must call pl.program_id(0) again to force a re-sync of the grid iteration

Issuing make_async_copy is synchronous by default, so no wait is needed

You must call copy.wait() to block on the semaphore until the DMA completes

Nothing extra; the Mosaic compiler auto-inserts a barrier before any read of vmem_ref

You must call pltpu.repeat on vmem_ref to materialize the copy