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:
pointer could be a single element pointer, then a scalar will be loaded - mask and other must be scalar too - other is implicitly typecast to pointer.dtype.element_ty - boundary_check and padding_option must be empty
pointer could be element-wise tensor of pointers, in which case: - mask and other are implicitly broadcast to pointer.shape - other is implicitly typecast to pointer.dtype.element_ty - boundary_check and padding_option must be empty
pointer could be a block pointer defined by make_block_ptr, in which case: - mask and other must be None - 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”}, do padding while out of bound
cache_modifier (str, optional) – 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