triton.language.atomic_xchg

triton.language.atomic_xchg(pointer, val, mask=None, sem=None, scope=None)

Performs an atomic exchange at the memory location specified by pointer.

Return the data stored at pointer before the atomic operation.

Parameters:

  • pointer (Block of dtype=triton.PointerDType) – The memory locations to operate on

  • val (Block of dtype=pointer.dtype.element_ty) – The values with which to perform the atomic operation

  • sem (str) – Memory semantics to use (“ACQUIRE_RELEASE” (default), “ACQUIRE”, “RELEASE”, or “RELAXED”)

  • scope (str) – Scope of threads that observe synchronizing effect of the atomic operation (“GPU” (default), “CTA”, or “SYSTEM”)