triton.experimental.gluon.language.load

triton.experimental.gluon.language.load(pointer, mask=None, other=None, *, cache_modifier='', eviction_policy='', volatile=False, _semantic=None)

Return a tensor of data whose values are loaded from memory at location defined by pointer:

  1. If pointer is a single element pointer, a scalar is be loaded. In this case:

    • mask and other must also be scalars,

    • other is implicitly typecast to pointer.dtype.element_ty.

  2. If pointer is an N-dimensional tensor of pointers, an N-dimensional tensor is loaded. In this case:

    • mask and other are implicitly broadcast to pointer.shape,

    • other is implicitly typecast to pointer.dtype.element_ty.

Parameters:

  • pointer (triton.PointerType, or block of dtype=triton.PointerType) – Pointer to the data to be loaded

  • mask (Block of triton.int1, optional) – if mask[idx] is false, do not load the data at address pointer[idx]

  • other (Block, optional) – if mask[idx] is false, return other[idx]. If other is None, the masked-out value is undefined.

  • cache_modifier (str, optional, should be one of {“”, “.ca”, “.cg”, “.cv”}, where “.ca” stands for cache at all levels, “.cg” stands for cache at global level (cache in L2 and below, not L1), and “.cv” means don’t cache and fetch again. see cache operator for more details.) – changes cache option in NVIDIA PTX

  • eviction_policy (str, optional) – changes eviction policy in NVIDIA PTX

  • volatile (bool, optional) – changes volatile option in NVIDIA PTX