如何在CUDA中进行原子加载

Question

如何在CUDA中进行原子加载

cudagpu-atomics

5

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

- kirill

所以你想要一个阻塞式加载？那听起来你需要自己编写互斥锁。 - talonmies

更具体地说，我想要类似于C++中的原子加载和存储的东西。http://en.cppreference.com/w/cpp/atomic/atomic/load - kirill

我真的不理解这个问题。每个线程最多加载128位数量的适当负载在“原子”意义上是指负载的任何部分都不会被“干扰”（负载或）存储修改。存储本身也保证是原子性的。原子函数的目的是提供一个不间断的RMW设施。 - Robert Crovella

2个回答

2

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

有两个准备替代方案，它们各有优缺点：

1. 像您建议的那样使用无操作原子读修改写。我以前提供了一个类似的答案。保证原子性和内存一致性，但需要支付不必要的写入成本。

2. 实际上，最接近原子加载的第二种方法可能是将变量标记为volatile，尽管严格来说语义完全不同。该语言不保证加载的原子性（例如，理论上可能会出现破碎的读取），但您保证获得最新值。但是，根据@Robert Crovella在注释中指出的，在最多32字节的正确对齐交易中，无法出现破碎的读取，这确实使它们成为原子操作。

方案2有点不正规，我不建议使用，但目前它是唯一的不需要写入操作的替代方案1。理想的解决方案是在语言中添加一种直接表达原子加载的方法。

- user703016

1

我不确定volatile限定符是否可以帮助加载的原子性。我认为它只强制生成的PTX加载操作具有.cv后缀，并认为缓存中的现有值已过时。它是否还会使线程看不到加载操作被撕裂？ - Farzad

@Farzad volatile 实际上对原子性没有帮助，因此 OP 如果想要这个保证，应该使用一个无操作的 RMW。对于任何小于或等于本机字长的内容，都不会发生撕裂写入或读取，因此 32 位类型不会出现这种情况。对于 64 位，则有可能。我不建议使用 volatile，但 OP 说他们不想为额外的原子写入付费。我会进行编辑。 - user703016

1

一个64位类型的正确对齐加载不会被“撕裂”或部分修改，也不会被“干扰”的写入所影响。我认为这个问题很愚蠢。所有内存事务都是针对L2缓存执行的。L2缓存仅提供32字节的高速缓存行。没有其他可能的交易。一个正确对齐的64位类型将始终落入单个L2高速缓存行中，并且该高速缓存行的服务不可能由某些在冗余写入之前的数据（这些数据将被冗余写入修改）和相同冗余写入之后的一些数据组成。 - Robert Crovella

@Robert 就 编程语言 而言，理论上是允许的（这就是我所说的“理论上”）。目前在CUDA中还没有一种表达“原子加载这个64位类型”的方法。 - user703016

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Allison Vacanti · Accepted Answer

除了在其他答案中建议使用volatile之外，适当地使用__threadfence也是必要的，以获得具有安全内存顺序的原子加载。

虽然一些评论说只需使用普通读取即可，因为它不会破坏数据的完整性，但这与原子加载不同。原子操作不仅仅是数据完整性：

普通读取可能会重用已经在寄存器中的先前读取数据，因此可能不反映由其他SM使用所需的内存顺序进行的更改。例如， int *flag = ...; while (*flag) { ... } 可能仅在循环的每次迭代中读取

 flag 一次并重复使用该值，如果你正在等待另一个线程更改标志的值，那么你永远无法观察到更改。  volatile 修饰符确保每次访问实际从内存中读取该值。 有关详细信息，请参见 CUDA文档中对volatile的描述 。
此外，您还需要使用内存障碍来在调用线程中强制执行正确的内存顺序。 如果没有内存障碍，则在C++11中，您会得到“松散”的语义，当使用原子通信时，这可能是不安全的。
例如，假设您的代码（非原子性地）将一些大型数据写入内存，然后使用普通写入来设置原子标志以指示已写入数据。 指令可以被重新排序、硬件缓存线可能在设置标志之前未被刷新等等。 结果是这些操作不能保证以任何顺序执行，并且其他线程可能无法按照您期望的顺序观察这些事件：在写入标志之前发生写入。
同时，如果读取线程也在使用普通读取来检查标志是否有条件地加载数据，则在硬件级别上会出现竞争。 乱序和/或推测性执行可能在标志读取完成之前加载数据。 然后使用推测性加载的数据，这可能无效，因为它是在标志读取之前加载的。
通过正确放置内存障碍，可以通过强制执行指令重排来避免这些问题，使之不影响所需的内存顺序，并且可以使先前的写入对其他线程可见。  __threadfence()等相关信息也在 CUDA文档中进行了介绍。
综上所述，在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已经包含了带有获取/释放内存顺序的读取/存储指令语义 - 但目前没有办法访问它们而不使用内联PTX。他们希望在今年的某个时候添加它们。一旦这些位置就位，完整的std::atomic实现就不会太远了。