4.15. 内存同步域
4.15.1. 内存栅栏干扰
某些 CUDA 应用程序可能会出现性能下降，因为内存栅栏或刷新操作等待的事务多于 CUDA 内存一致性模型所要求的事务。
__managed__ int x = 0;
__device__ cuda::atomic<int, cuda::thread_scope_device> a(0);
__managed__ cuda::atomic<int, cuda::thread_scope_system> b(0);
|
||
线程 1（SM） x = 1;
a = 1;
|
线程 2（SM） while (a != 1) ;
assert(x == 1);
b = 1;
|
线程 3（CPU） while (b != 1) ;
assert(x == 1);
|
考虑上面的示例。CUDA 内存一致性模型保证断言条件为真，因此，线程 1 对x的写入必须先于线程 2 对b的写入对线程 3 可见。
对a执行释放和获取所提供的内存顺序，仅足以使x对线程 2 可见，而不足以使其对线程 3 可见，因为这是设备作用域的操作。因此，对b执行释放和获取所提供的系统作用域顺序，不仅需要确保线程 2 自身发出的写入对线程 3 可见，还需要确保来自其他线程且对线程 2 可见的写入也对线程 3 可见。这称为累积性。由于 GPU 在执行时无法知道哪些写入在源代码层面已被保证可见，哪些写入只是由于偶然的时序而可见，因此它必须以保守方式涵盖较大范围内正在进行的内存操作。
这有时会导致干扰：由于 GPU 在等待源代码层面并不要求它等待的内存操作，栅栏或刷新操作所花费的时间可能会超过必要的时间。
请注意，栅栏可以像示例中那样，以内建函数或原子操作的形式显式出现在代码中，也可以隐式出现，以实现任务边界处的与之同步关系。
一个常见的示例是，某个核函数正在本地 GPU 内存中执行计算，而一个并行核函数（例如来自 NCCL 的核函数）正在与对等设备通信。完成时，本地核函数会隐式刷新其写入，以满足与下游工作之间的任何与之同步关系。这可能会不必要地全部或部分等待通信核函数发出的较慢的 nvlink 或 PCIe 写入。
4.15.2. 使用域隔离流量
从计算能力为 9.0（Hopper 架构）的 GPU 和 CUDA 12.0 开始，内存同步域功能提供了一种减轻此类干扰的方法。借助代码提供的显式协助，GPU 可以缩小栅栏操作所涵盖的范围。每次核函数启动都会被赋予一个域 ID。写入和栅栏都标记有该 ID，栅栏只会对与其域匹配的写入施加顺序约束。在上述计算与通信并发的示例中，可以将通信核函数放置在不同的域中。
使用域时，代码必须遵守以下规则：同一 GPU 上不同域之间的排序或同步需要系统作用域的栅栏操作。在一个域内，设备作用域的栅栏操作仍然足够。这对于累积性是必要的，因为一个核函数的写入不会被另一个域中的核函数发出的栅栏所涵盖。实质上，通过确保跨域流量提前刷新到系统作用域，即可满足累积性。
请注意，这会修改thread_scope_device的定义。但是，由于核函数将默认为域 0，如下文所述，因此仍保持向后兼容性。
4.15.3. 在 CUDA 中使用域
可以通过新的启动属性cudaLaunchAttributeMemSyncDomain和cudaLaunchAttributeMemSyncDomainMap访问域。前者在逻辑域cudaLaunchMemSyncDomainDefault和cudaLaunchMemSyncDomainRemote之间进行选择，后者提供从逻辑域到物理域的映射。远程域旨在用于执行远程内存访问的核函数，以将其内存流量与本地核函数的内存流量隔离。但是请注意，选择某个特定域并不影响核函数可以合法执行哪些内存访问。
可以通过设备属性cudaDevAttrMemSyncDomainCount查询域的数量。计算能力为 9.0（Hopper）的设备有 4 个域。为了便于编写可移植代码，所有设备上都可以使用域功能，对于计算能力低于 9.0 的设备，CUDA 将报告域的数量为 1。
逻辑域的存在便于应用程序组合。在软件栈较低层级进行的单次核函数启动（例如来自 NCCL 的启动）可以选择一个具有语义意义的逻辑域，而无需关心周围的应用程序架构。较高层级可以使用映射来调控逻辑域。如果未设置逻辑域，其默认值为默认域；默认映射将默认域映射到 0，将远程域映射到 1（在具有多于 1 个域的 GPU 上）。在 CUDA 12.0 及更高版本中，某些特定库可能会将启动标记为远程域；例如，NCCL 2.16 会这样做。这些机制共同为常见应用程序提供了一种开箱即用的有益使用模式，无需更改其他组件、框架或应用程序层面的代码。另一种使用模式可能是对并行流进行划分，例如在使用 NVSHMEM 或核函数类型没有明确区分的应用程序中。流 A 可以将两个逻辑域都映射到物理域 0，流 B 则映射到 1，依此类推。
// Example of launching a kernel with the remote logical domain
cudaLaunchAttribute domainAttr;
domainAttr.id = cudaLaunchAttrMemSyncDomain;
domainAttr.val = cudaLaunchMemSyncDomainRemote;
cudaLaunchConfig_t config;
// Fill out other config fields
config.attrs = &domainAttr;
config.numAttrs = 1;
cudaLaunchKernelEx(&config, myKernel, kernelArg1, kernelArg2...);
// Example of setting a mapping for a stream
// (This mapping is the default for streams starting on compute capability 9.0 (Hopper) or later if not
// explicitly set, and provided for illustration)
cudaLaunchAttributeValue mapAttr;
mapAttr.memSyncDomainMap.default_ = 0;
mapAttr.memSyncDomainMap.remote = 1;
cudaStreamSetAttribute(stream, cudaLaunchAttributeMemSyncDomainMap, &mapAttr);
// Example of mapping different streams to different physical domains, ignoring
// logical domain settings
cudaLaunchAttributeValue mapAttr;
mapAttr.memSyncDomainMap.default_ = 0;
mapAttr.memSyncDomainMap.remote = 0;
cudaStreamSetAttribute(streamA, cudaLaunchAttributeMemSyncDomainMap, &mapAttr);
mapAttr.memSyncDomainMap.default_ = 1;
mapAttr.memSyncDomainMap.remote = 1;
cudaStreamSetAttribute(streamB, cudaLaunchAttributeMemSyncDomainMap, &mapAttr);
与其他启动属性一样，这些属性以统一的方式提供给 CUDA 流、使用cudaLaunchKernelEx的单次启动以及 CUDA 图中的核函数节点。典型用法是在流级别设置映射，并在启动级别设置逻辑域（或在流的一段使用区间前后设置），如上所述。
在流捕获期间，这两个属性都会被复制到图节点。图从节点本身获取这两个属性，这实质上是间接指定物理域的一种方式。在图所启动到的流上设置的域相关属性，不会用于该图的执行。