Script.tma.wait_group

Script.tma.wait_group

Script.tma.wait_group(n, read=False)[source]

Wait for TMA async copy commit groups to complete.

Blocks until at most n commit groups remain pending. Use n=0 to wait for all committed groups.

When read=False (default), waits for all operations to complete, including writes being visible to the executing thread.

When read=True, only waits for reads from source locations to complete. This is useful when the source shared memory needs to be reused, but there is no subsequent instruction that reads the destination global memory. If subsequent instructions need to read the global memory written by TMA, use the default read=False.

Parameters:

  • n (int) – The number of groups to allow to be on-the-fly. It should be an integer larger or equal to 0.

  • read (bool) – If True, only wait for reads to complete (not writes). Default is False.

Return type:

None

Notes

  • Thread group: Can be executed by any sized thread group.

  • Hardware: Requires compute capability 9.0+ (sm_90).

  • PTX: cp.async.bulk.wait_group or cp.async.bulk.wait_group.read

Contents