For the complete documentation index, see llms.txt. Markdown versions of all pages are available by appending .md to any URL (e.g. /max/get-started.md).

Mojo function

async_copy_with_bound_check

def async_copy_with_bound_check[dtype: DType, src_layout: TensorLayout, dst_layout: TensorLayout, src_element_width: Int, dst_element_width: Int, //, thread_layout: Layout[thread_layout.shape_types, thread_layout.stride_types], swizzle_mode: TensorMapSwizzle](src: TileTensor[dtype, src_layout, ImmutAnyOrigin, Storage=PointerStorage[element_width=src_element_width], linear_idx_type=src.linear_idx_type], dst: TileTensor[dtype, dst_layout, MutAnyOrigin, Storage=PointerStorage[element_width=dst_element_width], address_space=AddressSpace.SHARED, linear_idx_type=dst.linear_idx_type])

Helper function for cp.async with boundary checking.

This method performs element-wise async copies with per-element boundary checking. Out-of-bounds accesses are automatically zero-filled, ensuring safe operation near matrix edges.

The method also handles shared memory swizzling to avoid bank conflicts and maximize memory bandwidth utilization.

Parameters:

dtype (DType): Element type of the source and destination tiles (inferred).
src_layout (TensorLayout): Static layout of the source tile in global memory (inferred).
dst_layout (TensorLayout): Static layout of the destination tile in shared memory (inferred).
thread_layout (Layout[thread_layout.shape_types, thread_layout.stride_types]): Thread mapping that partitions the source and destination tiles across threads.
swizzle_mode (TensorMapSwizzle): Shared memory swizzle pattern applied to avoid bank conflicts.

Args:

src (TileTensor[dtype, src_layout, ImmutAnyOrigin, Storage=PointerStorage[element_width=src_element_width], linear_idx_type=src.linear_idx_type]): Source tensor fragment in global memory.
dst (TileTensor[dtype, dst_layout, MutAnyOrigin, Storage=PointerStorage[element_width=dst_element_width], address_space=AddressSpace.SHARED, linear_idx_type=dst.linear_idx_type]): Destination tensor fragment in shared memory.