triton.experimental.gluon.language.amd.cdna4.async_copy.global_load_to_shared

triton.experimental.gluon.language.amd.cdna4.async_copy.global_load_to_shared(dest, ptr, mask=None, other=None, cache_modifier='', _semantic=None)

AMD global load to shared operation.

This operation loads data directly from global memory to shared memory without going through registers. It happens asynchronously; call wait_group() before accessing dest. The operation still completes in order with load(), store(), buffer_load(), and buffer_store() on CDNA4, so interleaving with them will hurt performance.

Compared to buffer_load_to_shared(), it requires a tensor pointer which supports 64-bit indexing range for each thread in a block, which gives more flexibility, but at the cost of higher register pressure and no hardware out-of-bound masking support. Prefer buffer_load_to_shared() when possible for better performance.

The underlying hardware instruction uses separate registers for global memory address for each thread but the same register for local memory address for the whole warp. Therefore, while using this operation the following conditions must be met or lowering to LLVM will fail:

  • For the ptr layout, size per thread * bits per element must be 128 or 32. To get ideal performance, it is recommended to use 128 bits per element.

  • Writes to dest must be coalesced.

  • If dest is swizzled, it can only be swizzled within a warp boundary.

Parameters:

  • dest (shared_memory_descriptor) – Destination shared memory descriptor.

  • ptr (pointer tensor) – Tensor of pointers to global memory to load from.

  • mask (tensor, optional) – Mask tensor for predicated loads. Defaults to None.

  • other (tensor or scalar, optional) – Tensor or scalar providing default values for masked elements. Defaults to None.

  • cache_modifier (str) – Cache modifier specifier. Defaults to “”.