CUDA中自旋锁的正确实现

Question

很多资源都提供了 CUDA 中自旋锁的实现：

• https://devtalk.nvidia.com/default/topic/1014009/try-to-use-lock-and-unlock-in-cuda/
• Cuda Mutex, why deadlock?
• How to implement Critical Section in cuda?
• Implementing a critical section in CUDA
• https://wlandau.github.io/gpu/lectures/cudac-atomics/cudac-atomics.pdf。

它们遵循相同的模式：

1. LOCK：等待lock的值从0到1的原子变化
2. 做一些关键操作
3. UNLOCK：通过将锁的值设置为0来释放锁

假设我们没有 warp-divergence，或者换句话说，我们不使用锁来进行 interwarp 同步。

实施步骤 1 的正确方法是什么？

一些答案​​建议使用atomicCAS，而其他atomicExch。两者是等价的吗？

while (0 != (atomicCAS(&lock, 0, 1))) {}
while (atomicExch(&lock, 1) != 0) {}

实施步骤 3 的正确方法是什么？

几乎所有消息来源都建议为此使用atomicExch：

atomicExch(&lock, 0);

一位用户提出了一个替代方案 (Implementing a critical section in CUDA)，这也很有意义，但对他不起作用（因此可能会导致 CUDA 中的未定义行为）：

lock = 0;

似乎对于 CPU 上的一般自旋锁，这样做是有效的：https://stackoverflow.com/a/7007893/8044236。为什么我们不能在 CUDA 中使用它？

我们是否必须在第 2 步中使用内存栅栏和 volatile 说明符进行内存访问？

关于原子的 CUDA 文档 (https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomic-functions) 说他们不保证排序约束：

原子函数不充当内存栅栏，也不意味着内存操作的同步或排序约束

这是否意味着我们必须在临界区 (2) 的末尾使用内存栅栏来确保临界区 (2) 内的更改在解锁 (3) 之前对其他线程可见？

CUDA 是否保证其他线程会看到在步骤 (1) 和 (3) 中具有原子操作的线程所做的更改？

这不适用于内存栅栏 (https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#memory-fence-functions)：

内存栅栏函数只影响线程内存操作的顺序；它们不确保这些内存操作对其他线程可见（就像 __syncthreads() 对块内的线程所做的那样（请参阅同步函数））。

所以原子操作可能也不是这样？如果是，CUDA 中的所有自旋锁实现都依赖于 UB。

我们如何在存在 warp 的情况下实现可靠​​的自旋锁？

现在，如果我们对上述所有问题都有答案，那么让我们移除我们没有翘曲发散的假设。这种情况下是否可以实现自旋锁？

主要问题（死锁）在https://wlandau.github.io/gpu/lectures/cudac-atomics/cudac-atomics.pdf 的幻灯片 30 中表示：

是否是在步骤 (1) 中将 while 循环替换为 if 并将所有 3 个步骤包含在单个 while 循环中的唯一选项，例如，在 Thread/warp local lock in cuda 或 CUDA, mutex and atomicCAS() 中？

Answer 1

实施第 1 步的正确方法是什么？一些答案建议使用 atomicCAS 而其他 atomicExch。两者是等价的吗？

不，它们不是，只有atomicCas 是正确的。该代码的重点是检查锁的状态从解锁到由给定线程锁定的变化是否有效。 atomicExch 版本不会这样做，因为它不会在执行分配之前检查初始状态是否已解锁。

似乎对于 CPU 上的一般自旋锁，这样做是有效的：https://stackoverflow.com/a/7007893/8044236。为什么我们不能在 CUDA 中使用它？

如果您阅读该答案的 cmets，您会发现它在 CPU 上也无效。

我们是否必须在第 2 步中对内存访问使用内存栅栏和 volatile 说明符？

这完全取决于您的最终用途以及您首先使用关键部分的原因。显然，如果您希望给定线程对全局可见内存的操作是全局可见的，则需要栅栏或原子事务来执行此操作，并且需要确保编译器不会将值缓存在寄存器中。

[原文如此] CUDA 是否保证其他线程将永远看到具有原子操作的线程在步骤 (1) 和 (3) 中所做的更改？

是的，但仅适用于原子执行的其他操作。原子操作意味着在给定地址上执行的所有内存事务的序列化，并且当线程执行操作时它们返回地址的先前状态。

我们如何在存在 warp 的情况下实现可靠​​的自旋锁？

翘曲散度无关紧要。当来自同一个warp的多个线程尝试获取锁时，原子内存操作的序列化意味着warp发散