triton.experimental.gluon.language.store

triton.experimental.gluon.language.store(pointer, value, mask=None, *, cache_modifier='', eviction_policy='', _semantic=None)

Store a tensor of data into memory locations defined by pointer.

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

    • mask must also be scalar.

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

    • mask is implicitly broadcast to pointer.shape.

value is implicitly broadcast to pointer.shape and typecast to pointer.dtype.element_ty.

Parameters:

  • pointer (triton.PointerType, or block of dtype=triton.PointerType) – The memory location where the elements of value are stored

  • value (Block) – The tensor of elements to be stored

  • mask (Block of triton.int1, optional) – If mask[idx] is false, do not store value[idx] at pointer[idx]

  • cache_modifier (str, optional, should be one of {“”, “.wb”, “.cg”, “.cs”, “.wt”}, where “.wb” stands for cache write-back all coherent levels, “.cg” stands for cache global, “.cs” stands for cache streaming, “.wt” stands for cache write-through, see cache operator for more details.) – changes cache option in NVIDIA PTX

  • eviction_policy (str, optional, should be one of {"", "evict_first", "evict_last"}) – changes eviction policy in NVIDIA PTX

This function can also be called as a member function on tensor, as x.store(...) instead of store(x, ...).