triton.language.load

triton.language.load(pointer, mask=None, other=None, boundary_check=(), padding_option='', cache_modifier='', eviction_policy='', volatile=False)

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, and

   • boundary_check and padding_option must be empty.

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, and

   • boundary_check and padding_option must be empty.

3. If pointer is a block pointer defined by make_block_ptr, a tensor is loaded. In this case:

   • mask and other must be None, and

   • boundary_check and padding_option can be specified to control the behavior of out-of-bound access.

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] (must be None with block pointers)

• other (Block, optional) – if mask[idx] is false, return other[idx]

• boundary_check (tuple of ints, optional) – tuple of integers, indicating the dimensions which should do the boundary check

• padding_option – should be one of {“”, “zero”, “nan”}, the padding value to use while out of bounds. “” means an undefined value.

• cache_modifier (str, optional, should be one of {“”, “ca”, “cg”}, where “ca” stands for cache at all levels and “cg” stands for cache at global level (cache in L2 and below, not L1), 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