cp_async_bulk_tensor_2d_gather4

cp_async_bulk_tensor_2d_gather4[dst_type: AnyType, mbr_type: AnyType, /, *, cta_group: Int = 1, eviction_policy: CacheEviction = CacheEviction.EVICT_NORMAL](dst_mem: UnsafePointer[dst_type, address_space=AddressSpace.SHARED], tma_descriptor: UnsafePointer[NoneType], mem_bar: UnsafePointer[mbr_type, address_space=AddressSpace.SHARED], col_idx: Int32, row0: Int32, row1: Int32, row2: Int32, row3: Int32)

Initiates an asynchronous gather4 copy of 4 non-contiguous rows from a 2D tensor in global memory into shared memory using TMA.

This is a Blackwell (SM100+) instruction that loads 4 rows at arbitrary row indices from a 2D tensor map, placing them contiguously in shared memory. The tensor map must be created with box dim1=1 (one row per tile). Each row is a full tile along the column dimension.

PTX: cp.async.bulk.tensor.2d.shared::cta.global.tile::gather4 .mbarrier::complete_tx::bytes

Constraints:

Requires SM100 (Blackwell) or newer GPU architecture.

Parameters:

• ​dst_type (AnyType): The data type of the destination memory.
• ​mbr_type (AnyType): The data type of the memory barrier.
• ​cta_group (Int): The CTA group for the copy operation. Must be 1 or 2.
• ​eviction_policy (CacheEviction): Cache eviction policy. Defaults to EVICT_NORMAL.

Args:

• ​dst_mem (UnsafePointer[dst_type, address_space=AddressSpace.SHARED]): Pointer to the destination in shared memory. Must be 128-byte aligned.
• ​tma_descriptor (UnsafePointer[NoneType]): Pointer to the TMA descriptor (created with cuTensorMapEncodeTiled and box dim1=1).
• ​mem_bar (UnsafePointer[mbr_type, address_space=AddressSpace.SHARED]): Pointer to a shared memory barrier for synchronization.
• ​col_idx (Int32): Column offset in the source tensor (typically 0 for full-row loads).
• ​row0 (Int32): Row index of the first row to gather.
• ​row1 (Int32): Row index of the second row to gather.
• ​row2 (Int32): Row index of the third row to gather.
• ​row3 (Int32): Row index of the fourth row to gather.