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
cp_async_bulk_global_shared_cta
def cp_async_bulk_global_shared_cta[dst_type: AnyType, src_type: AnyType, /, *, eviction_policy: CacheEviction = CacheEviction.EVICT_NORMAL](dst_mem: Pointer[dst_type, address_space=dst_mem.address_space], src_mem: Pointer[src_type, address_space=AddressSpace.SHARED], size: Int32)
Initiates an asynchronous bulk copy from shared CTA memory to global memory.
Performs a non-blocking copy of size bytes from shared memory to global
memory using the cp.async.bulk PTX instruction with the .bulk_group
completion mechanism. Use cp_async_bulk_commit_group and
cp_async_bulk_wait_group from max.gpu.sync to synchronize.
Both dst_mem and src_mem must be 16-byte aligned, and size must be a
multiple of 16. Requires sm_100 or higher.
Parameters:
- dst_type (
AnyType): The element type of the destination global memory. - src_type (
AnyType): The element type of the source shared memory. - eviction_policy (
CacheEviction): Cache eviction policy for the L2 cache. Defaults toEVICT_NORMAL.
Args:
- dst_mem (
Pointer[dst_type, address_space=dst_mem.address_space]): Destination pointer in global or generic memory (16-byte aligned). - src_mem (
Pointer[src_type, address_space=AddressSpace.SHARED]): Source pointer in shared CTA memory (16-byte aligned). - size (
Int32): Number of bytes to copy (must be a multiple of 16).