如何在 CUDA 中进行原子负载

Question

我的问题是如何在 CUDA 中实现原子负载。原子交换可以模拟原子存储。可以以类似的方式非昂贵地模拟原子负载吗？ 我可以使用带有 0 的原子添加来原子地加载内容，但我认为它很昂贵，因为它执行原子读取-修改-写入，而不仅仅是读取。

Answer 1

除了按照其他答案的建议使用volatile 之外，还需要适当地使用__threadfence 以获得具有安全内存排序的原子负载。

虽然一些 cmets 说只使用普通读取，因为它不会撕裂，但这与原子负载不同。原子不仅仅是撕裂：

正常读取可能会重用已在寄存器中的先前加载，因此可能无法反映其他 SM 以所需内存顺序所做的更改。例如，int *flag = ...; while (*flag) { ... } 可能只读取一次flag，并在循环的每次迭代中重复使用该值。如果您正在等待另一个线程更改标志的值，您将永远不会观察到更改。 volatile 修饰符确保在每次访问时实际从内存中读取该值。请参阅CUDA documentation on volatile 了解更多信息。

此外，您需要使用内存栅栏在调用线程中强制执行正确的内存排序。如果没有栅栏，您将获得 C++11 用语中的“宽松”语义，这在使用原子进行通信时可能是不安全的。

例如，假设您的代码（非原子）将一些大数据写入内存，然后使用普通写入设置原子标志以指示数据已被写入。指令可能会被重新排序，硬件缓存线可能不会在设置标志之前被刷新等等。结果是这些操作不能保证以任何顺序执行，其他线程可能不会按照你期望的顺序观察这些事件: 允许在写入保护数据之前写入标志。

同时，如果读取线程在有条件地加载数据之前也使用正常读取来检查标志，则会在硬件级别出现竞争。乱序和/或推测执行可能会在标志读取​​完成之前加载数据。然后使用推测加载的数据，这可能是无效的，因为它是在读取标志之前加载的。

放置良好的内存栅栏通过强制指令重新排序不会影响您所需的内存顺序并且使之前的写入对其他线程可见来防止此类问题。 __threadfence()和朋友们也被in the CUDA docs覆盖了。

将所有这些放在一起，在 CUDA 中编写您自己的原子加载方法如下所示：

// addr must be aligned properly.
__device__ unsigned int atomicLoad(const unsigned int *addr)
{
  const volatile unsigned int *vaddr = addr; // volatile to bypass cache
  __threadfence(); // for seq_cst loads. Remove for acquire semantics.
  const unsigned int value = *vaddr;
  // fence to ensure that dependent reads are correctly ordered
  __threadfence(); 
  return value; 
}

// addr must be aligned properly.
__device__ void atomicStore(unsigned int *addr, unsigned int value)
{
  volatile unsigned int *vaddr = addr; // volatile to bypass cache
  // fence to ensure that previous non-atomic stores are visible to other threads
  __threadfence(); 
  *vaddr = value;
}

对于其他非撕裂加载/存储大小，可以类似地编写。

通过与一些从事 CUDA atomics 工作的 NVIDIA 开发人员的交谈，看起来我们应该开始看到 CUDA 中对 atomics 的更好支持，并且 PTX 已经包含 load/store instructions with acquire/release memory ordering 语义——但目前没有办法访问它们诉诸内联 PTX。他们希望在今年的某个时候添加它们。一旦这些都到位，完整的std::atomic 实施应该不会落后。

Answer 2

据我所知，目前无法在 CUDA 中请求原子负载，这将是一个很棒的功能。

有两种准替代方案，各有优缺点：

1. 按照您的建议使用无操作原子读取-修改-写入。我过去提供过similar answer。保证原子性和内存一致性，但您需要为不必要的写入付出代价。

2. 实际上，第二接近原子负载的事情可能是标记变量volatile，尽管严格来说语义完全不同。该语言不保证负载的原子性（例如，您理论上可能会被撕毁），但保证您会获得最新的价值。但是在实践中，正如@Robert Crovella 在 cmets 中所指出的那样，对于最多 32 个字节的正确对齐事务，不可能获得撕裂读取，这确实使它们成为原子。

解决方案 2 有点老套，我不推荐它，但它是目前唯一替代 1 的无写替代方案。理想的解决方案是添加一种直接用语言表达原子负载的方法。