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, optional) – Specifies the memory semantics for the operation. Acceptable values are “acquire”, “release”, “acq_rel” (stands for “ACQUIRE_RELEASE”), and “relaxed”. If not provided, the function defaults to using “acq_rel” semantics.

• scope (str, optional) – Defines the scope of threads that observe the synchronizing effect of the atomic operation. Acceptable values are “gpu” (default), “cta” (cooperative thread array, thread block), or “sys” (stands for “SYSTEM”). The default value is “gpu”.

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