triton.experimental.gluon.language.amd.cdna5.tdm.async_load

triton.experimental.gluon.language.amd.cdna5.tdm.async_load(src: tensor_descriptor, offsets: List[ttgl.constexpr | ttgl.tensor] = None, dest: shared_memory_descriptor = None, pred=None, mbarrier: shared_memory_descriptor = None, warp_used_hint=None, cache_modifier='', _semantic=None) → None

Load a block of tensor specified in tensor descriptor from global memory to shared memory asynchronously.

This operation expects a prior update_tensor_descriptor() to position the descriptor for the load offsets and bounds. Doing this explicitly lets the developer control when and where the descriptor update happens, which has performance implications for downstream code generation. For convenience, offsets (and optionally pred) may be passed here, in which case an update_tensor_descriptor() is emitted immediately before the load.

Parameters:

  • src (tensor_descriptor) – the source tensor descriptor.

  • offsets (List[int], optional) – if given, the offsets from the base pointer used to position the descriptor before the load.

  • dest (shared_memory_descriptor) – the shared memory destination to store the loaded data.

  • pred (bool, optional) – if given, predicate to enable or disable the load.

  • mbarrier (shared_memory_descriptor, optional) – The barrier object to signal “arrive” on.

  • warp_used_hint (int, optional) – Bitmask selecting the active warp subset for descriptor layout (bit n => warp n). Doesn’t affect the data in dest, only the work split. The number of active warps must be a power of two, and the active warps must follow a regular bit pattern for efficient lowering. Examples: 0b00001111 (warps 0..3), 0b11110000 (warps 4..7), 0b01010101 (warps 0,2,4,6). Omit / None = all warps participate; explicit 0 and other invalid hints are rejected by the verifier.

  • cache_modifier (str, optional) – Cache behavior.