triton.experimental.gluon.language.amd.cdna5.tdm.async_gather

triton.experimental.gluon.language.amd.cdna5.tdm.async_gather(desc: tensor_descriptor, src_row_indices: ttgl.tensor, dst: shared_memory_descriptor, mbarrier: shared_memory_descriptor = None, _semantic=None) → None

Gather data from non-contiguous rows in global memory to shared memory asynchronously.

This operation uses TDM gather mode to read data from non-contiguous rows in global memory. Unlike async_load(), which reads from contiguous rows, gather allows reading from arbitrary rows specified by the src_row_indices tensor.

The dtype of src_row_indices determines the index size:

• int16: up to 16 rows can be gathered per TDM instruction.

• int32: up to 8 rows can be gathered per TDM instruction.

If more rows are needed, multiple TDM instructions will be automatically issued.

The column offset and predicate are carried by the descriptor: position it beforehand with update_tensor_descriptor(), passing add_offsets=[0, col] and pred=p.

Parameters:

• desc (tensor_descriptor) – The source tensor descriptor. Must be 2D and positioned through update_tensor_descriptor() for the column and predicate.

• src_row_indices (tensor) – 1D tensor of row indices (int16 or int32) in the source tensor.

• dst (shared_memory_descriptor) – the shared memory destination to store gathered data. Must be 2D.

• mbarrier (shared_memory_descriptor, optional) – The barrier object to signal “arrive” on.