Mojo struct
ScatterGather
@register_passable(trivial)
struct ScatterGather
Implemented traits
AnyType,
Copyable,
ImplicitlyCopyable,
Movable,
UnknownDestructibility
Aliases
__copyinit__is_trivial
alias __copyinit__is_trivial = True
__del__is_trivial
alias __del__is_trivial = True
__moveinit__is_trivial
alias __moveinit__is_trivial = True
Methods
load_tile
static load_tile[dtype: DType, tile_layout: Layout, desc_layout: Layout, dst_layout: Layout, //, cluster_size: Int, use_partitioned_multicast: Bool](tma_op: TMATensorTile[dtype, tile_layout, desc_layout], dst: LayoutTensor[dtype, dst_layout, origin, address_space=AddressSpace(3), element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=128], ref [3] mem_barrier: SharedMemBarrier, rank: UInt, coords: Tuple[UInt, UInt], multicast_mask: UInt16)
static load_tile[dtype: DType, src_layout: Layout, dst_layout: Layout, //, thread_layout: Layout, swizzle_mode: TensorMapSwizzle, vector_size: Int](src: LayoutTensor[dtype, src_layout, MutableAnyOrigin, element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], dst: LayoutTensor[dtype, dst_layout, MutableAnyOrigin, address_space=AddressSpace(3), element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], tile_idx_m: Int, tile_idx_n: Int)
async_copy_with_bound_check
static async_copy_with_bound_check[dtype: DType, src_layout: Layout, dst_layout: Layout, //, thread_layout: Layout, swizzle_mode: TensorMapSwizzle](src: LayoutTensor[dtype, src_layout, MutableAnyOrigin, element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment], dst: LayoutTensor[dtype, dst_layout, MutableAnyOrigin, address_space=AddressSpace(3), element_layout=element_layout, layout_int_type=layout_int_type, linear_idx_type=linear_idx_type, masked=masked, alignment=alignment])
Helper function for cp.async with bound checking.
Was this page helpful?
Thank you! We'll create more content like this.
Thank you for helping us improve!