4

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

4

2 回答 2

4

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

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

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

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

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

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

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

将所有这些放在一起，在 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 已经包含具有获取/释放内存排序语义的加载/存储指令——但没有办法当前访问它们而无需使用内联 PTX。他们希望在今年的某个时候添加它们。一旦这些都到位，一个完整的std::atomic实施应该不会落后。

于 2019-02-28T11:42:33.293 回答
2

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

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

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

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

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

于 2015-09-02T01:41:07.233 回答