triton.language.atomic_xchg

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

在 pointer 指定的内存位置执行原子交换。

返回原子操作之前存储在 pointer 的数据。

参数:

• pointer (Block 的 dtype=triton.PointerDType) – 要操作的内存位置

• val (Block 的 dtype=pointer.dtype.element_ty) – 用于执行原子操作的值

• sem (str, 可选) – 指定操作的内存语义。可接受的值为 "acquire"（获取），"release"（释放），"acq_rel"（代表 "ACQUIRE_RELEASE"）和 "relaxed"（宽松）。如果未提供，函数默认为使用 "acq_rel" 语义。

• scope (str, 可选) – 定义观察原子操作同步效果的线程范围。可接受的值为 "gpu"（默认），"cta"（协同线程数组，线程块），或 "sys"（代表 "SYSTEM"）。默认值为 "gpu"。

此函数也可以作为 tensor 的成员函数调用，例如 x.atomic_xchg(...) 而不是 atomic_xchg(x, ...)。