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

triton.experimental.gluon.language.amd.cdna5.tdm.async_scatter(desc: tensor_descriptor, dst_row_indices: ttgl.tensor, src: shared_memory_descriptor, mbarrier: shared_memory_descriptor = None, _semantic=None) → None

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

This operation uses TDM scatter mode to write data to non-contiguous rows in global memory. Unlike async_store(), which writes to contiguous rows, scatter allows writing to arbitrary rows specified by the dst_row_indices tensor.

The dtype of dst_row_indices determines the index size:

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

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

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

The column offset is carried by the descriptor: position it beforehand with update_tensor_descriptor() with add_offsets=[0, col].

Parameters:

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

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

• src (shared_memory_descriptor) – the shared memory source containing data to scatter. Must be 2D.

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