【CUDA编程】系列博客参考NVIDIA官方文档“CUDA C++ Programming Guide（v12.6）”。
本文为原创文章，未经本人允许，禁止转载。转载请注明出处。
1.Memory Fence Interference
先介绍两个概念：内存屏障（memory fence）和内存刷新（memory flush）。
内存屏障的主要目的是控制内存访问的顺序，确保不同线程或处理器之间的内存操作按照预期顺序发生。在并发编程中，不同线程或处理器对内存的访问可能会因为编译器优化、处理器指令重排、缓存行为等原因，导致内存操作的顺序与程序代码中的顺序不一致。在单线程程序中，内存操作通常按照程序的顺序执行，因此不会存在顺序问题。但在并发编程中，由于不同线程或处理器同时访问共享内存，可能会出现顺序错乱的情况。例如：
1
2
3
4
5
6
7
8
// Thread 1
x = 1; // 写操作
flag = 1; // 写操作
// Thread 2
if (flag == 1) { // 读操作
assert(x == 1); // 期望 x 已经被写入 1
}
在线程1中，x的值被设置为1，然后flag被设置为1。理论上，线程2检测到flag==1时，x的值应该已经是1，但如果没有内存屏障，处理器或编译器可能会重排操作顺序，导致线程2读取到的x还是旧值。
内存刷新操作通常是指将某个缓存或临时存储中的数据写回到主存，以确保数据的修改对其他处理器、线程或设备是可见的。这种操作在并发编程和多处理器系统中非常重要，因为现代计算系统通常会使用缓存来提升性能，而内存刷新操作可以确保缓存中更新的数据被传播到共享的内存区域，使其他处理单元能够访问到最新的数据。
鉴于CUDA的内存一致性模型的要求，一些CUDA应用程序可能会因为内存屏障或内存刷新而导致性能下降。
__managed__声明了一个位于统一内存（Unified Memory）中的变量，该变量可以在host和device之间共享使用，而无需手动管理内存拷贝。
考虑以上例子，如果程序按照顺序执行，那么在a被赋值为1之前，x肯定已经是1了（见线程1），因此在线程2中，如果a不等于1，则while会一直循环，当a等于1时，跳出while循环，走到assert，此时x肯定是等于1的，所以判定成功，程序继续执行，然后b被赋值为1。同理，在线程3中，如果b不等于1，则while会一直循环，当b等于1时，跳出循环，走到assert，因为程序是顺序执行的缘故，如果有b等于1，则此时x肯定已经也是1了，所以assert判定成功，程序可以继续执行。
但在多线程程序中，上述顺序可能会被打乱，比如a已经被赋值为1了，但此时x还未被赋值为1，如果是这种情况，程序就有可能会出现问题。此时就需要内存屏障和内存刷新来保证上述程序的正常运行。
2.Isolating Traffic with Domains
从Hopper架构和CUDA 12.0开始，内存同步域（memory synchronization domains）功能提供了一种方法用于缓解内存同步操作（比如内存屏障和内存刷新）所带来的开销。GPU可以通过一个屏障操作（fence operation）减少同步覆盖的范围。每个kernel启动时会被赋予一个域ID（domain ID）。一个域ID对应一个域，一个域内可以有多个kernel。在传统的内存同步模型中，内存屏障操作往往会影响所有相关内存，导致同步范围过于宽泛，从而增加了额外的开销。而通过引入域的概念，CUDA可以将同步操作的影响限制在一个特定的kernel集合中，减少不必要的同步操作。
使用域时，必须遵守以下规则：在同一GPU上，不同域之间的排序或同步需要系统范围的屏障（system-scope fencing）。而在同一域内，device范围的屏障（device-scope fencing）就足够了。
注意，在引入内存同步域后，thread_scope_device只同步所在域内的所有线程。kernel默认会被分配到域0。
3.Using Domains in CUDA
首先介绍下逻辑域和物理域的概念：
- 逻辑域（logical domain）：逻辑域是一种在高层应用程序或库中定义的抽象概念，用于区分kernel启动时的不同同步环境。逻辑域帮助开发者不必直接考虑硬件的同步实现细节，而是可以在代码中使用更高层次的抽象来定义同步策略。逻辑域的引入使得应用的逻辑与硬件同步解耦。开发者可以指定一个逻辑域而不必知道具体的硬件架构，这使得代码更加模块化和可移植。逻辑域又被分为默认域和远程域（remote domain）：
- 默认域：默认域是CUDA中逻辑域的默认设置，如果开发者没有显式指定逻辑域，系统会将kernel启动映射到默认域，其内存访问和同步需求仅涉及本地的内存（也就是在同一device内）。在默认域中，kernel的同步行为不会涉及跨device的远程访问，而是集中在device范围内的同步。这种设置适合于那些需要在同一个device内进行多个线程块协作的计算任务。默认的逻辑域会映射到物理域0。
- 远程域：专门用于执行涉及远程内存访问的任务。这些远程内存访问通常包括跨device的数据交换，比如GPU与其他GPU之间的通信。远程域的主要作用是将远程内存流量与本地内存操作隔离，使得远程访问的内存操作不会影响到同一个device内其他kernel的内存同步。这对于大规模并行计算任务中，多GPU之间进行数据通信的场景非常有用。通常，远程域被映射到物理域1（在具有多个物理域的GPU上），从而将远程内存访问的同步行为与默认kernel操作的同步行为隔离开来。
- 物理域（physical domain）：物理域是实际硬件上的一个同步范围，用于定义内存同步的物理隔离。它指的是GPU硬件中划分的不同同步区域，这些区域的内存同步可以独立进行，以减少整体的同步开销。物理域的数量在不同架构上可能不同，例如在Hopper架构的GPU中，有4个物理域，这意味着GPU可以划分为4个独立的同步区域。在CUDA中，每个逻辑域都可以映射到一个物理域，逻辑域和物理域之间的映射关系是可以自定义的，通过启动属性来控制。多个逻辑域可以映射到同一个物理域。
域的指定可以通过设置kernel启动属性来实现。cudaLaunchAttributeMemSyncDomain用于设置逻辑域，其中cudaLaunchMemSyncDomainDefault用于设置默认域，cudaLaunchMemSyncDomainRemote用于设置远程域。cudaLaunchAttributeMemSyncDomainMap用于设置逻辑域到物理域的映射。
可以通过cudaDevAttrMemSyncDomainCount查询device中域的数量。Hopper架构有4个域，再之前的架构都是1个域。
下面是一个使用远程域启动kernel的例子：
1
2
3
4
5
6
7
8
9
// 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...);
下面的例子是为CUDA stream设置逻辑域到物理域的映射：
1
2
3
4
5
6
7
// Example of setting a mapping for a stream
// (This mapping is the default for streams starting on Hopper if not
// explicitly set, and provided for illustration)
cudaLaunchAttributeValue mapAttr;
mapAttr.memSyncDomainMap.default_ = 0; //默认域映射到物理域0
mapAttr.memSyncDomainMap.remote = 1; //远程域映射到物理域1
cudaStreamSetAttribute(stream, cudaLaunchAttributeMemSyncDomainMap, &mapAttr);
下面的例子是将不同的CUDA stream映射到不同的物理域，忽略逻辑域的默认设置：
1
2
3
4
5
6
7
8
9
// 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);
