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

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

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

参数:

pointer (Block of dtype=triton.PointerDType) – 要操作的内存位置
  • val (Block of 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, ...)