5.7. CUDA C++ 内存模型

标准 C++ 所呈现的观点是，线程同步的成本均匀且较低。

CUDA C++ 则不同：线程之间的距离越远，线程同步的成本就越高。同一线程块内各线程之间的同步成本较低，但系统中运行于多个 GPU 和 CPU 上的任意线程之间的同步成本较高。

为应对不均匀且并非始终较低的线程同步成本，CUDA C++ 在cuda::命名空间中使用线程作用域扩展了标准 C++ 内存模型和并发机制，同时默认保留标准 C++ 的语法和语义。

5.7.1. 线程作用域

线程作用域指定了哪些类型的线程可以使用cuda::atomic或cuda::barrier等同步原语相互同步。

namespace cuda {

enum thread_scope {
  thread_scope_system,
  thread_scope_device,
  thread_scope_block,
  thread_scope_thread
};

}  // namespace cuda

5.7.1.1. 作用域关系

每个程序线程都通过一个或多个线程作用域关系与其他每个程序线程相关联：

系统中的每个线程都通过系统线程作用域与系统中的其他每个线程相关联：cuda::thread_scope_system。
每个 GPU 线程都通过设备线程作用域与同一 CUDA 设备中且位于同一内存同步域内的其他每个 GPU 线程相关联：cuda::thread_scope_device。
每个 GPU 线程都通过线程块线程作用域与同一 CUDA 线程块中的其他每个 GPU 线程相关联：cuda::thread_scope_block。
每个线程都通过线程线程作用域与自身相关联：cuda::thread_scope_thread。

5.7.2. 同步原语

命名空间std::和cuda::std::中的类型，与命名空间cuda::中以cuda::thread_scope_system作用域实例化的相应类型具有相同的行为。

5.7.3. 原子性

如果满足以下条件，则原子操作在其指定的作用域内具有原子性：

它指定的作用域不是cuda::thread_scope_system，或
作用域为cuda::thread_scope_system且：
- 它作用于系统分配的内存中的对象，并且pageableMemoryAccess为1[0,1]，或
- 它作用于托管内存中的对象，并且concurrentManagedAccess为1[1]，或
- 它作用于映射内存中的对象，并且hostNativeAtomicSupported为1，或
- 它是作用于映射内存[2] 中自然对齐且大小为1、2、4、8或16字节的对象的加载或存储操作，或
- 它作用于 GPU 内存中的对象，只有 GPU 线程访问该对象，且
  - 在每个进行访问的srcDev与对象所在的 GPU（dstDev）之间，cudaDeviceGetP2PAttribute(&val, cudaDevP2PAttrNativeAtomicSupported, srcDev, dstDev)为1[1]，或
  - 只有来自单个 GPU 的 GPU 线程并发访问它。

注意

[0] 如果PageableMemoryAccessUsesHostPagetables为0，则对内存映射文件或hugetlbfs分配的内存进行的原子操作不具有原子性。
[1] 在这些情况下，只有hostNativeAtomicSupported为1时，系统作用域的张量归约操作才在系统作用域内具有原子性。系统作用域的张量归约操作包括 PTX{multimem.}cp.reduce.async.bulk.tensor.relaxed.sys指令族及公开这些指令的 API，例如cuda::ptx::cp_reduce_async_bulk_tensor。
[2] 如果hostNativeAtomicSupported为0，则作用于系统分配的内存或映射内存中自然对齐的 16 字节宽对象的系统作用域原子加载或存储操作需要系统支持。NVIDIA 尚未知晓任何缺乏此项支持的系统，也没有可用于检测此类系统的 CUDA API 查询。

有关系统分配的内存、托管内存、映射内存、CPU 内存和 GPU 内存的更多信息，请参阅本指南的相关章节。

5.7.4. 数据竞争

将 ISO/IEC IS 14882（C++ 标准）的intro.races 第 21 段修改如下：

如果程序的执行包含两个可能并发且相互冲突的动作，其中至少有一个在包含执行另一操作的线程的作用域内不是原子操作，且二者互不先发生于对方，则该执行包含数据竞争，但下文所述的信号处理程序的特殊情况除外。任何此类数据竞争都会导致未定义行为。[…]

将 ISO/IEC IS 14882（C++ 标准）的thread.barrier.class 第 4 段修改如下

4. 对barrier除析构函数以外的成员函数的并发调用不会引入数据竞争，如同它们是原子操作一样。[…]

将 ISO/IEC IS 14882（C++ 标准）的thread.latch.class 第 2 段修改如下：

2. 并发调用latch的成员函数（析构函数除外）不会引入数据竞争，如同这些调用是原子操作一样。[…]

将 ISO/IEC IS 14882（C++ 标准）的thread.sema.cnt 第 3 段修改如下：

3. 并发调用counting_semaphore的成员函数（析构函数除外）不会引入数据竞争，如同这些调用是原子操作一样。

将 ISO/IEC IS 14882（C++ 标准）的thread.stoptoken.intro 第 5 段修改如下：

调用函数request_stop、stop_requested和stop_possible不会引入数据竞争，如同这些调用是原子操作一样。[…]

将 ISO/IEC IS 14882（C++ 标准）的atomics.fences 第 2 至第 4 段修改如下：

如果存在均作用于某个原子对象 M 的原子操作 X 和 Y，使得 A 按顺序先于 X，X 修改 M，Y 按顺序先于 B，并且 Y 读取 X 写入的值，或者读取假设 X 是释放操作时以其为首的假想释放序列中任何副作用写入的值，且每个操作（A、B、X 和 Y）指定的作用域都包含执行其他各操作的线程，则释放栅栏 A 与获取栅栏 B 同步。

如果存在原子操作 X，使得 A 按顺序先于 X，X 修改 M，并且 B 读取 X 写入的值，或者读取假设 X 是释放操作时以其为首的假想释放序列中任何副作用写入的值，且每个操作（A、B 和 X）指定的作用域都包含执行其他各操作的线程，则释放栅栏 A 与在原子对象 M 上执行获取操作的原子操作 B 同步。

如果存在某个作用于 M 的原子操作 X，使得 X 按顺序先于 B，并且读取 A 写入的值，或者读取以 A 为首的释放序列中任何副作用写入的值，且每个操作（A、B 和 X）指定的作用域都包含执行其他各操作的线程，则在原子对象 M 上执行释放操作的原子操作 A 与获取栅栏 B 同步。

5.7.5. 示例：消息传递

以下示例通过标志f，将线程块0中的一个线程存储到x变量的消息传递给线程块1中的一个线程：

初始状态

int x = 0, f = 0;

线程块 0 的线程 0

x = 42;
cuda::atomic_ref<int, cuda::thread_scope_device> flag(f);
flag.store(1, memory_order_release);

线程块 1 的线程 0

cuda::atomic_ref<int, cuda::thread_scope_device> flag(f);
while(flag.load(memory_order_acquire) != 1);
assert(x == 42);

在前一示例的以下变体中，两个线程在没有同步的情况下并发访问f对象，这会导致数据竞争，并表现出未定义行为：

初始状态

int x = 0, f = 0;

线程块 0 的线程 0

x = 42;
cuda::atomic_ref<int, cuda::thread_scope_block> flag(f);
flag.store(1, memory_order_release); // UB: data race

线程块 1 的线程 0

cuda::atomic_ref<int, cuda::thread_scope_device> flag(f);
while(flag.load(memory_order_acquire) != 1); // UB: data race
assert(x == 42);

虽然对f的内存操作（存储和加载）是原子的，但存储操作的作用域是“线程块作用域”。由于该存储由线程块 0 的线程 0 执行，因此它仅包含线程块 0 中的所有其他线程。然而，执行加载的线程位于线程块 1 中，也就是说，它不在由线程块 0 中执行的存储操作所包含的作用域内，导致该存储和加载不是“原子的”，从而引入数据竞争。

更多示例请参见PTX 内存一致性模型试金石测试。