是否有与 CUDA 的 atomicCAS 等效的 DPC++/SYCL？答案

【问题标题】：Is there a DPC++/SYCL equivalent of CUDA's atomicCAS?是否有与 CUDA 的 atomicCAS 等效的 DPC++/SYCL？
【发布时间】：2022-06-17 22:46:33
【问题描述】：

据我了解，CUDA的atomicCAS有如下定义（这是四个之一）

int atomicCAS(int* address, int compare, int val);

它会将全局共享内存中位于address（在文档old 中命名）的值与compare 进行原子比较，如果相等，则将该值分配给val，否则什么也不做。在这两种情况下都返回old。

查看 SYCL API，我只能找到 compare_exchange_strong，不幸的是，它并没有做我想要的，因为它使用与上面相同的命名，它将 old 与 compare 进行比较，如果不成功，改变compare（通过引用传递）。

【问题讨论】：

是的，C++ std::atomic's compare_exchange_strong 通过引用更新 expected，但它旨在成为局部变量。该输出 arg 的更新是非原子的。这只是暴露相同底层原始操作（原子 CAS）的不同方式。更改代码以使用它纯粹是声明局部变量的问题。如果比较为真，那么compare 已经等于old 的值，所以不需要更新它。如果这就是您认为它不同或代码有问题的原因，请 IDK。

【解决方案1】：

正如 Peter Cordes 在评论中指出的那样，sycl::compare_exchange_strong 是正确的选择。从 SYCL 2020 rev. 4compare_exchange_strong的描述：

以原子方式将此atomic_ref 引用的对象的值与expected 的值进行比较。如果值相等，将引用对象的值替换为desired的值；否则将引用对象的原始值分配给expected。

所以，

int old = compare;
ref.compare_exchange_strong(old, val);

就更新ref而言，相当于

old = atomicCAS(address, compare, val);

【讨论】：