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

Mojo function

async_copy

def async_copy[dtype: DType, //, size: Int, *, fill: Optional[Scalar[dtype]] = None, bypass_L1_16B: Bool = True, l2_prefetch: Optional[Int] = None, eviction_policy: CacheEviction = CacheEviction.EVICT_NORMAL](src: Pointer[Scalar[dtype], address_space=AddressSpace.GLOBAL], dst: Pointer[Scalar[dtype], address_space=AddressSpace.SHARED], src_size: Int32 = SIMD(size), predicate: Bool = False)

Asynchronously copies data from global memory to shared memory.

This function provides a high-performance asynchronous memory copy operation with configurable caching behavior, prefetching, and fill values. It maps directly to the PTX cp.async instruction on NVIDIA GPUs.

Constraints:

• Fill value only supported for dtypes <= 32 bits.
• Size must be 4, 8, or 16 bytes.
• Cannot enable both L2 prefetch and L1 bypass.
• L2 prefetch size must be 64, 128, or 256 bytes.

Parameters:

• ​dtype (DType): The data type to copy (e.g. float32, int32).
• ​size (Int): Number of bytes to copy (must be 4, 8, or 16).
• ​fill (Optional[Scalar[dtype]]): Optional fill value for uncopied bytes when src_size < size.
• ​bypass_L1_16B (Bool): If True, bypasses L1 cache for 16-byte copies.
• ​l2_prefetch (Optional[Int]): Optional L2 prefetch size (64, 128, or 256 bytes).
• ​eviction_policy (CacheEviction): Cache eviction policy for the copy operation.

Args:

• ​src (Pointer[Scalar[dtype], address_space=AddressSpace.GLOBAL]): Source pointer in global memory.
• ​dst (Pointer[Scalar[dtype], address_space=AddressSpace.SHARED]): Destination pointer in shared memory.
• ​src_size (Int32): Actual bytes to copy from src (remaining bytes use fill value).
• ​predicate (Bool): Optional predicate to conditionally execute the copy.