3.4. 多 GPU 系统编程
多 GPU 编程通过利用多 GPU 系统提供的更高的总算术运算性能、更大的总内存容量和更高的总内存带宽，使应用程序能够处理超出单个 GPU 能力的问题规模，并达到超出单个 GPU 能力的性能水平。
CUDA 通过主机 API、驱动程序基础设施和配套的 GPU 硬件技术实现多 GPU 编程：
主机线程的 CUDA 上下文管理
系统中所有处理器的统一内存寻址
GPU 之间的对等批量内存传输
细粒度的对等 GPU 加载/存储内存访问
更高层次的抽象和配套系统软件，例如 CUDA 进程间通信、使用NCCL进行并行归约，以及通过NVSHMEM和 MPI 等 API，使用 NVLink 和/或 GPU-Direct RDMA 进行通信
在最基本的层面上，多 GPU 编程要求应用程序并发管理多个活动的 CUDA 上下文，将数据分发给各个 GPU，在 GPU 上启动核函数以完成其工作，并通信或收集结果，以便应用程序能够对这些结果采取操作。具体如何执行这些步骤，取决于如何将应用程序的算法、可用并行性和现有代码结构最有效地映射到合适的多 GPU 编程方法。一些最常见的多 GPU 编程方法包括：
由单个主机线程驱动多个 GPU
由多个主机线程各自驱动自己的 GPU
由多个单线程主机进程各自驱动自己的 GPU
多个主机进程包含多个线程，每个线程驱动自己的 GPU
通过 NVLink 连接的多节点簇，其中 GPU 由分布在簇内各节点上的多个操作系统实例中运行的线程和进程驱动
GPU 可以通过设备内存之间的内存传输和对等访问相互通信，涵盖上述每种多设备工作分配方法。通过查询并启用对等 GPU 内存访问，并利用 NVLink 实现设备之间的高带宽传输和更细粒度的加载/存储操作，可以支持高性能、低延迟的 GPU 通信。
CUDA 统一虚拟寻址允许同一主机进程内的多个 GPU 之间进行通信，只需极少的额外步骤即可查询并启用高性能对等内存访问和传输，例如通过 NVLink 进行访问和传输。
通过使用进程间通信（IPC）和虚拟内存管理（VMM）API，可以支持由不同主机进程管理的多个 GPU 之间的通信。进程间通信一节介绍了 IPC 的高层概念，并讨论了节点内 CUDA IPC API。高级虚拟内存管理（VMM）API 同时支持节点内和多节点 IPC，可在 Linux 和 Windows 操作系统上使用，并允许以单次分配为粒度控制内存缓冲区的 IPC 共享，如虚拟内存管理中所述。
对于 NVLink 互连网络（fabric）上的 GPU，计算互连传输采用了不同的方法：它不像上述方法那样通过映射的虚拟地址访问对等设备的内存，而是使用逻辑端点 ID 为一组资源命名，通过偏移量指定该组中的某个资源，并通过显式的异步互连网络（fabric）操作访问该资源。参见下文的计算互连传输。
CUDA 本身提供了在一组 GPU 内实现集合操作所需的 API，这些操作也可能包括主机，但 CUDA 本身并不提供高层多 GPU 集合 API。多 GPU 集合操作由NCCL和NVSHMEM等更高抽象层次的 CUDA 通信库提供。
3.4.1. 多设备上下文与执行管理
应用程序要使用多个 GPU，首先需要枚举可用的 GPU 设备，根据这些设备的硬件属性、CPU 亲和性以及与对等设备的连接情况，从可用设备中作出适当选择，并为应用程序将使用的每个设备创建 CUDA 上下文。
3.4.1.1. 设备枚举
以下代码示例展示了如何查询支持 CUDA 的设备数量、枚举每个设备并查询其属性。
int deviceCount;
cudaGetDeviceCount(&deviceCount);
int device;
for (device = 0; device < deviceCount; ++device) {
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, device);
printf("Device %d has compute capability %d.%d.\n",
device, deviceProp.major, deviceProp.minor);
}
3.4.1.2. 设备选择
主机线程可以随时通过调用cudaSetDevice()来设置其当前操作的设备。设备内存分配和核函数启动在当前设备上进行；流和事件在创建时与当前设置的设备关联。在主机线程调用cudaSetDevice()之前，当前设备默认为设备 0。
以下代码示例说明了设置当前设备如何影响后续的内存分配和核函数执行操作。
size_t size = 1024 * sizeof(float);
cudaSetDevice(0); // Set device 0 as current
float* p0;
cudaMalloc(&p0, size); // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1); // Set device 1 as current
float* p1;
cudaMalloc(&p1, size); // Allocate memory on device 1
MyKernel<<<1000, 128>>>(p1); // Launch kernel on device 1
3.4.1.3. 多设备流、事件和内存复制行为
如果向一个不与当前设备关联的流发出核函数启动请求，则启动将失败，如以下代码示例所示。
cudaSetDevice(0); // Set device 0 as current
cudaStream_t s0;
cudaStreamCreate(&s0); // Create stream s0 on device 0
MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 0 in s0
cudaSetDevice(1); // Set device 1 as current
cudaStream_t s1;
cudaStreamCreate(&s1); // Create stream s1 on device 1
MyKernel<<<100, 64, 0, s1>>>(); // Launch kernel on device 1 in s1
// This kernel launch will fail, since stream s0 is not associated to device 1:
MyKernel<<<100, 64, 0, s0>>>(); // Launch kernel on device 1 in s0
即使向一个不与当前设备关联的流发出内存复制请求，复制也会成功。
如果输入事件和输入流关联的是不同的设备，cudaEventRecord()将失败。
如果两个输入事件关联的是不同的设备，cudaEventElapsedTime()将失败。
即使输入事件关联的设备与当前设备不同，cudaEventSynchronize()和cudaEventQuery()也会成功。
即使输入流和输入事件关联的是不同的设备，cudaStreamWaitEvent()也会成功。因此，可以使用cudaStreamWaitEvent()使多个设备彼此同步。
每个设备都有自己的默认流，因此，向某个设备的默认流发出的命令，相对于向任何其他设备的默认流发出的命令，可能乱序执行或并发执行。
3.4.2. 多设备对等传输和内存访问
3.4.2.1. 对等内存传输
CUDA 可以在设备之间执行内存传输，并且在能够进行对等内存访问时，会利用专用复制引擎和 NVLink 硬件来最大限度地提高性能。
cudaMemcpy可以与复制类型cudaMemcpyDeviceToDevice或cudaMemcpyDefault一起使用。
否则，必须使用cudaMemcpyPeer()、cudaMemcpyPeerAsync()、cudaMemcpy3DPeer()或cudaMemcpy3DPeerAsync()执行复制，如以下代码示例所示。
cudaSetDevice(0); // Set device 0 as current
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size); // Allocate memory on device 0
cudaSetDevice(1); // Set device 1 as current
float* p1;
cudaMalloc(&p1, size); // Allocate memory on device 1
cudaSetDevice(0); // Set device 0 as current
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1); // Set device 1 as current
cudaMemcpyPeer(p1, 1, p0, 0, size); // Copy p0 to p1
MyKernel<<<1000, 128>>>(p1); // Launch kernel on device 1
两个不同设备的内存之间的复制（在隐式NULL流中）：
只有在先前向这两个设备中的任一设备发出的所有命令都已完成后才会开始，并且
只有在复制运行至完成后，在该复制之后向这两个设备中的任一设备发出的任何命令（参见异步执行）才能开始。
与流的正常行为一致，两个设备的内存之间的异步复制可以与另一个流中的复制或核函数重叠执行。
如果已在两个设备之间启用对等访问，例如按照对等内存访问中所述的方式启用，那么这两个设备之间的对等内存复制就不再需要通过主机暂存，因此速度更快。
3.4.2.2. 对等内存访问
根据系统属性，特别是 PCIe 和/或 NVLink 拓扑，设备能够对彼此的内存进行寻址（即，在一个设备上执行的核函数可以解引用指向另一个设备内存的指针）。如果cudaDeviceCanAccessPeer()对指定的设备返回 true，则这两个设备之间支持对等内存访问。
必须通过调用cudaDeviceEnablePeerAccess()来启用两个设备之间的对等内存访问，如以下代码示例所示。在未启用 NVSwitch 的系统上，每个设备在整个系统范围内最多可支持八个对等连接。
两个设备使用统一的虚拟地址空间（参见统一虚拟地址空间），因此可以在这两个设备上使用同一个指针对内存进行寻址，如下面的代码示例所示。
cudaSetDevice(0); // Set device 0 as current
float* p0;
size_t size = 1024 * sizeof(float);
cudaMalloc(&p0, size); // Allocate memory on device 0
MyKernel<<<1000, 128>>>(p0); // Launch kernel on device 0
cudaSetDevice(1); // Set device 1 as current
cudaDeviceEnablePeerAccess(0, 0); // Enable peer-to-peer access
// with device 0
// Launch kernel on device 1
// This kernel launch can access memory on device 0 at address p0
MyKernel<<<1000, 128>>>(p0);
注意
使用cudaDeviceEnablePeerAccess()启用对等内存访问会全局作用于对等设备上所有先前和后续的 GPU 内存分配。通过cudaDeviceEnablePeerAccess()启用对某个设备的对等访问，会增加该对等设备上设备内存分配操作的运行时开销，因为需要使分配的内存立即可供当前设备以及同样具有访问权限的任何其他对等设备访问，从而带来随对等设备数量增长的成倍开销。
与为所有设备内存分配启用对等内存访问相比，一种可扩展性更好的替代方法是使用 CUDA 虚拟内存管理 API，仅在需要时，在分配期间显式分配可供对等设备访问的内存区域。通过在内存分配期间显式请求对等可访问性，不可供对等设备访问的内存分配的运行时开销不会受到影响，同时可供对等设备访问的数据结构也能被正确限定作用域，从而改善软件调试并提高可靠性（参见 ref::virtual-memory-management）。
3.4.2.3. 对等内存一致性
对于分布在多个设备上的网格中并发执行的线程，必须使用同步操作来强制保证其内存访问的顺序和正确性。跨设备同步的线程在thread_scope_system同步作用域内操作。同样，内存操作属于thread_scope_system内存同步域。
当只有一个 GPU 访问对等设备内存中的某个对象时，CUDA原子函数可以对该对象执行读取-修改-写入操作。对等原子性的要求和限制在 CUDA 内存模型的原子性要求讨论中进行了说明。
3.4.2.4. 多设备托管内存
托管内存可以在支持对等访问的多 GPU 系统上使用。并发多设备托管内存访问的详细要求，以及用于 GPU 独占访问托管内存的 API，在多 GPU中进行了说明。
3.4.2.5. 主机 IOMMU 硬件、PCI 访问控制服务和虚拟机
具体而言，在 Linux 上，CUDA 和显示驱动程序不支持在启用 IOMMU 的裸机上进行 PCIe 对等内存传输。但是，CUDA 和显示驱动程序确实支持通过虚拟机直通使用 IOMMU。在裸机系统上运行 Linux 时，必须禁用 IOMMU，以防止设备内存发生静默损坏。相反，对于虚拟机，应启用 IOMMU，并使用 VFIO 驱动程序进行 PCIe 直通。
在 Windows 上，不存在上述 IOMMU 限制。
另请参见在 64 位平台上分配 DMA 缓冲区。
此外，可以在支持 IOMMU 的系统上启用 PCI 访问控制服务（ACS）。PCI ACS 功能会将所有 PCI 点对点流量重定向，经由 CPU 根联合体传输，这会降低总体对分带宽，从而可能造成显著的性能损失。
3.4.2.6. 计算互连传输
对等内存访问通过发起进程已预留并映射的虚拟地址访问对等设备上的内存分配。计算互连传输为大型 NVLink 互连网络（fabric）提供了一种替代方式：它使用逻辑端点 ID 标识对等端，并通过偏移量而不是映射地址来标识所访问的资源。它面向通信库和高级应用程序，详见计算互连传输。