CUDA C编程（三十一）从一个GPU到多个GPU

  在应用程序中添加对多GPU的支持，其最常见的原因是以下几个方面：

• 问题域的大小：现有的数据集太大，单GPU内存大小与其不相符合；
• 如果单GPU适合处理单任务，那么可以通过使用多GPU并发地处理多任务来增加应用程序地吞吐量。

  在多GPU系统中，允许分摊跨GPU的服务器节点的功率消耗，具体方式是为给定的功率消耗单元提供更多的性能，同时提高吞吐量。当使用多GPU运行应用程序时，需要正确设计GPU间的通信。GPU间数据传输的效率取决于GPU是如何连接在一个节点上并跨集群的。在多GPU系统里有两种连接方式：

• 多GPU通过单个节点连接到PCIe总线上；
• 多GPU连接到集群中的网络交换机上。

  这些连接拓扑结构不是互斥的。下图展示了一个集群的简化拓扑结构，它其中有两个计算节点。GPU0和GPU1通过PCIe总线连接到node0上。同样，GPU2和GPU3在通过PCIe总线连接到node1上。两个节点(node0 和 node1)通过Infiniband交换机互相连接。

  每个节点可能包括以下内容中的一个或多个：通过CPU插槽和主机芯片连接的多个CPU，主机DRAM，本地存储设备，网络主机卡适配器(HCA),板载网络和USB端口，以及连接多个GPU的PCIe交换机。系统可能有一个PCIe根节点和多个PCIe交换机，这些PCIe交换机连接到根节点上，并在一个树结构中连接GPU。因为PCIe链路是双工的，所以可以使用CUDA API在PCIe链路之间映射一条路径，以避免总线竞争，同时也可以在GPU间共享数据。

  为了设计一个利用多GPU的程序，需要跨设备分配工作负载。根据应用程序，这种分配会导致两种常见的GPU间通信模式:

• 问题分区之间没有必要进行数据交换，因此在各GPU间没有数据共享；
• 问题分区之间有部分数据交换，在各GPU间需要冗余数据存储。

  第一种模式是最基本的情况：每个问题分区可以在不同的GPU上独立运行。要处理这些情况，只需了解如何在多个设备中传输数据及调用内核。在第二种情况下，GPU之间的数据交换是必需的，必须考虑数据如何在设备之间实现最优移动。总之，要避免通过主机内存中转数据(即数据复制到主机，只能将它复制到另一个GPU上)，重要的是要注意有多少数据被传输了和发生了多少次传输。

在 多 GPU 上 执 行
  CUDA 4.0中增加的功能使CUDA程序员能更容易地使用多GPU。CUDA运行时API支持在多GPU系统中管理设备和执行内核的多种方式。单个主机线程可以管理多个设备。一般来说，第一步是确定系统内可用的使能CUDA设备的数量是，使用如下函数可获得：cudaError_t cudaGetDeviceCount(int *count);该函数返回计算能力为1.0或更高的设备数量。下面的代码说明了如何确定使能CUDA设备的数量，对其进行遍历，并查询性能。

int ngpus;
cudaGetDeviceCount(&ngpus);

for(int i = 0; i < ngpus; i++)
{
   cudaDeviceProp devProp;
   cudaGetDeviceProperties(&devProp,i);
   printf("Device %d has compute capability %d.%d. \n"i, devProp.major,devProp.minor);
}

  在利用与多GPU一起工作的CUDA应用程序时，必须显式地指定哪个GPU是当前所有CUDA运算的目标。使用以下函数设置当前设备：cudaError_t cudaSetDevice(int id);该函数将具有标识符id的设备设置为当前设备。该函数不会与其他设备同步，因此是一个低开销的调用。使用此函数，可以在任何时间从任何主机线程中选择任何设备。有效的设备标识符是从0到ngpus-1。如果在首个CUDA API调用发生之前，没有显式地调用cudaSetDevice函数，那么当前设备会被自动设置设备0。一旦选定了当前设备，所有的CUDA运算将会被应用到那个设备上：

• 任何从主线程中分配来的设备内存将完全地常驻于那个设备上；
• 任何由CUDA运行时函数分配的主机内存都会有与该设备相关的生存时间；
• 任何由主机线程创建的流或事件都会与该设备相关；
• 任何由主机线程启动的内核都会在该设备上执行。

  可以在以下情况中同时使用多GPU：

• 在一个节点的单CPU线程上；
• 在一个节点的多CPU线程上；
• 在一个节点的多CPU进程上；
• 在一个节点的多CPU进程上。

  下面的代码准确展示了如何执行内核和单一的主机线程中进行内存拷贝，使用循环遍历设备：

for(int i = 0; i < ngpus; i++)
{
   cudaSetDevice(i);
   kernel<<<grid,block>>>(...);
   cudaMemcpyAsync(...);
}

  因为循环中的内核启动和数据传输是异步的，因此在每次调用操作后控制将很快返回到主机线程。但是，即使内核或由当前线程发出的传输仍然在当前设备上执行时，也可以安全地转变设备，因为cudaSetDevice函数不会导致主机同步。总之，想要在单一节点内获取GPU的数量和他们的性能，可以使用下述函数：

cudaError_t cudaGetDeviceCount(int *count);
cudaError_t cudaGetDeviceProperties(struct cudaDeviceProp *prop,int device);

  然后使用下述函数设置当前设备：cudaError_t cudaSetDevice(int device);一旦设置好了当前设备，所有CUDA操作都会在那个设备的上下文发出。

点 对 点 通 信
  在计算能力为2.0或以上的设备中，在64位应用程序上执行的内核，可以直接访问任何GPU的全局内存，这些GPU连接到同一个PCIe根节点上。如果想这样操作，必须使用CUDA点对点(P2P)API来实现设备间的直接通信。点对点通信需要CUDA 4.0或更高版本，相应的GPU驱动器，以及一个具有两个或两个以上连接到同一个PCIe根节点上的Fermi或Kepler GPU系统。有两个由CUDA P2P API支持的模式，它们允许GPU之间直接通信：

• 点对点访问：在CUDA内核和GPU间直接加载和存储地址；
• 点对点传输：在GPU间直接复制数据。

  在一个系统内，如果两个GPU连接到不同的PCIe根节点上，那么不允许直接进行点对点访问，并且CUDA P2P API将会通知你。仍然可以使用CUDA P2P API在这些设备之间进行点对点传输，但是驱动器将通过主机内存透明地传输数据，而不是通过PCIe总线直接传输数据。

启用点对点访问
  点对点访问允许各GPU连接到同一个PCIe根节点上，使其直接引用存储在其他GPU设备内存上的数据。对于透明的内核，引用的数据将通过PCIe总线传输到请求的线程上。因为并不是所有的GPU都支持点对点访问，所以需要使用下述函数显式地检查设备是否支持P2P:cudaError_t cudaDeviceCanAccessPeer(int* canAccessPeer,int device,int peerDevice)；如果设备device能够直接访问对等设备peerDevice的全局内存，那么函数变量canAccessPeer返回值为整型1,否则返回0。在两个设备间，必须用以下函数显式地启用点对点内存访问：cudaError_t cudaDeviceEnablePeerAccess(int peerDevice,unsigned int flag);这个函数允许从当前设备到peerDevice进行点对点访问。flag参数被保留以备将来使用，目前必须将其设置为0。一旦成功，该对等设备的内存将立即由当前设备进行访问。这个函数授权的访问是单向的，即这个函数允许从当前设备到peerDevice的访问，但不允许从peerDevice到当前设备的访问。如果希望对等设备能直接访问当前设备的内存，则需要另一个方向单独的匹配调用。点对点访问保持启用状态，知道它被以下函数显式地禁用：cudaError_t cudaDeviceDisablePeerAccess(int peerDevice)；32位应用程序不支持点对点访问。

点对点内存复制
  两个设备之间启用对等访问之后，使用下面的函数，可以异步地复制设备上的数据：cudaError_t cudaMemcpyPeerAsync(void* dst, int dstDev, void* src, int srcDev, size_t nBytes, cudaStream_t stream);这个函数将数据从设备的srcDev设备内存传输到设备dstDev地设备内存中。函数cudaMemcpyPeerAsync对于主机和所有其他设备来说是异步地。如果srcDev和dstDev共享相同的PCIe根节点，那么数据传输是沿着PCIe最短路径执行的，不需要通过主机内存中专。

多 GPU 间 的 同 步
  在多GPU应用程序上可以使用和单GPU应用程序相同的同步函数，但是必须指定适合的当前设备。多GPU应用程序中使用流和事件的典型工作流程如下所示：

1. &选择这个应用程序将使用的GPU集；
2. 为每个设备创建流和事件；
3. 为每个设备分配设备资源(如设备内存)；
4. 通过流在每个GPU上启动任务(例如，数据传输或内核执行)；
5. 使用流和事件来查询和等待任务完成；
6. 清空所有设备的资源。

  只有与该流相关联的设备是当前设备时，在流中才能启动内核。只有与该流相关联的设备是当前设备时，才可以在流中记录事件。任何时间都可以在任何流中进行内存拷贝，无论该流与什么设备相关联或当前设备是什么、即使流或事件与当前设备不相关，也可以查询或同步它们。