cuda - 如何在 CUDA 中实现原子负载

Question

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

Answer 1

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

虽然有些评论说只使用普通读取，因为它不能撕裂，但这与原子加载不同。原子不仅仅是撕裂:

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

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

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

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

放置良好的内存栅栏通过强制指令重新排序不会影响您所需的内存排序并且使之前的写入对其他线程可见来防止此类问题。 __threadfence()和 friend 也被覆盖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 原子的 NVIDIA 开发人员交谈，看起来我们应该开始看到 CUDA 中对原子的更好支持，并且 PTX 已经包含 load/store instructions with acquire/release memory ordering语义——但目前没有办法访问它们而不求助于内联 PTX。他们希望在今年的某个时候添加它们。一旦这些到位，一个完整的 std::atomic实现应该不会太落后。