3 / 11 Kernels (CUDA / Triton / Pallas)

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?

A

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

B

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

C

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

D

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

E

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