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

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

AMD buffer load to shared operation.

Buffer load is similar to global_load_to_shared(), but it accesses global memory through a scalar base pointer and a tensor of 32-bit offsets rather than a tensor of pointers. 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 global_load_to_shared(), it has better performance and also supports hardware out-of-bound masking. But it strictly requires a 32-bit offset instead of a 64-bit tensor pointer.

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 offsets 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 to scalar) – Global memory scalar base pointer to load from.

  • offsets (tensor) – Offsets tensor for the load operation.

  • 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 “”.