triton.language.atomic_cas

triton.language.atomic_cas(pointer, cmp, val, sem=None, scope=None)

Performs an atomic compare-and-swap 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

  • cmp (Block of dtype=pointer.dtype.element_ty) – The values expected to be found in the atomic object

  • 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”)