CUDA C编程（三十三）多GPU上的点对点通信

  本篇主要介绍两个GPU之间的数据传输。将测试以下3种情况：

• 两个GPU之间的单向内存复制；
• 两个GPU之间的双向内存复制；
• 内核中对等设备内存的访问。

实 现 点 对 点 访 问

首先，必须对所有设备启用双向点对点访问，如以下代码所示：

inline void enableP2P(int ngpus){
   for(int i = 0; i < ngpus; i++){
      cudaSetDevice(i);
      for(int j = 0; j < ngpus; j++){
         if(i == j)   continue;
          
         int peer_access_available = 0;
         cudaDeviceCanAccessPeer(&peer_access_available, i, j);
         
         if(peer_access_available){
            cudaDeviceEnablePeerAccess(j,0);
            printf("> GPU%d enabled direct access to GPU%d\n", i, j);
         }
         else{
            printf("(%d, %d)\n", i, j);
         }
      }
   }
}

  函数enableP2P遍历所有设备对(i, j),如果支持点对点访问，则使用cudaDeviceEnablePeerAccess函数启用双向点对点访问。

点 对 点 的 内 存 复 制

  启用点对点访问后，可以在两个设备之间直接复制数据。如果不支持点对点访问，该例子输出不能启用点对点访问的设备ID(不能启用的最有可能的原因是因为它们没有连接到同一个PCIe根节点上)，并且没有错误继续运行了。然而，回想一下可知，如果在两个GPU之间不支持点对点访问，那么这两个设备之间的点对点内存复制将通过主机内存中转，从而会降低其性能。性能降低对应用程序的影响程度取决于内核进行时间计算和执行对等传输需要的时间。如果有足够的时间来进行计算，那么可以隐藏点对点复制的延时，该延时主要是通过主机内存使用设备计算进行重叠的。

  启用点对点访问后，下面的代码在两个设备间执行ping-pong同步内存复制，次数为100次。如果点对点访问在所有设备上都被成功启用了，那么将直接通过PCIe总线进行数据传输而不用与主机交互。

cudaEventRecorf(start, 0);
for(int i = 0; i < 100; i++){
   if(i % 2 == 0){
      cudaMemcpy(d_src[1], d_src[0], iBytes, cudaMemcpyDeviceToDevice);
   }
   else{
      cudaMemcpy(d_src[0], d_src[1], iBytes, cudaMemcpyDeviceToDevice);
   }
}

  请注意，在内存复制之前没有设备转换，因为跨设备的内存复制不需要显式地设置当前设备。如果在内存复制前指定了设备，也不会影响它的行为。

  如需衡量设备之间数据传输的性能，需要把启动和停止事件记录在统一设备上，并将ping-pong内存复制包含在内。然后，用cudaEventElapseTime计算两个事件之间消耗的事件。

cudaSetDevice(0);
cudaEventRecord(start, 0);
for(int i = 0; i < 100; i++){
   ...
}
cudaSetDevice(0);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);

float elapsed_time_ms;
cudaEventElapsedTime(&elapsed_time_ms, start, stop);

然后通过ping-pong测试所获得的带宽按照下面所示的代码进行估计：

elapsed_time_ms /= 100.0f;
printf("Ping-pong unidirectional cudaMemcpy:\t\t %8.2f ms", elapsed_time_ms);
printf("performance: %8.2f GB/s\n",(float)iBytes / (elapsed_time_ms * 1e6f)));

从Wrox.com中可以下载包含这个例子的文件simpleP2P_PingPong.cu。编译和运行如下所示：

$ nvcc -O3 simpleP2P_PingPong.cu -o simplePingPong
$ ./simplePingPong

simpleP2P_PingPong的输出如下所示：

因为PCIe总线支持任何两个端点的全双工通信，所以也可以使用异步复制函数来进行双向的且点对点的内存复制：

cudaEventRecord(start, 0);
for(int i = 0; i < 100; i++){
   cudaMemcpyAsync(d_src[1], d_src[0], iBytes, cudaMemcpyDeviceToDevice, stream[0]);
   cudaMemcpyAsync(d_src[0], d_src[1], iBytes, cudaMemcpyDeviceToDevice, stream[1])
}

双向内存复制的测试在同一个文件中可以被实现。下面是一个示例输出：

注意，因为PCIe总线是一次在两个方向上使用的，所以获得的带宽增加了一倍。如果通过在simpleP2P_PingPong.cu中移除对enableP2P的调用来禁用点对点访问，那么无论是单向还是双向的例子都会不带任何错误的运行，但由于通过主机内存中转传输，所以测得的带宽将会下降。

统 一 虚 拟 寻 址 的 点 对 点 内 存 访 问
  统一虚拟寻址(UAV),是将CPU系统内存和设备的全局内存映射到一个单一的虚拟地址空间，如下图所示。所有由cudaHostAlloc分配的主机内存和由cudaMalloc分配的设备内存留在这个统一的地址空间内。内存地址所驻留的设备可以根据地址本身确定。

  将点对点CUDA API与UVA相结合，可以实现对任何设备内存的透明访问。不必手动管理单独的内存缓冲区，也不必从主机内存种进行显式地复制。底层系统能使我们避免显式地执行这些操作，从而简化了代码。请注意，过于依赖UAV进行对等访问对性能将产生负面的影响，如跨PCIe总线的许多小规模的传输会明显地有过大的消耗。下面的代码演示了如何检查设备是否支持统一寻址：

int deviceId = 0;
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, deviceId);
printf("GPU%d: %s unified addressing\n",deviceId, prop.unifiedAddressing ? "supports" : "does not support");

  为了使用UVA，应用程序必须在设备的计算能力为2.0及以上的64位架构上进行编译，并且CUDA版本为4.0或以上。如果同时启用点对点访问和UVA,那么在一个设备上执行的核函数，可以解除另一个设备上存储的指针。可以使用以下简单的核函数（该函数将输入数组扩展了2倍，并将结果存储在输出数组中），来测试GPU的直接点对点内存访问：

__global__ void iKernel(float *src, float *dst){
   const int idx = blockIdx.x * blockDim.x + threadIdx.x;
   dst[idx] = src[idx] * 2.0f;
}

  以下代码将设备0设置为当前设备，并有一个核函数使用指针d_src[1]从设备1中读取全局内存，同时通过全局内存的指针d_rcv[0]将结果写入当前设备中。

cudaSetDevice(0);
iKernel<<<grid, block>>>(d_rcv[0], d_src[1]);

  以下代码将设备1设置为当前设备，并有一个核函数使用指针d_src[0]从设备0中读取全局内存，同时通过全局内存的指针d_rcv[1]将结果写入当前设备中。

cudaSetDevice(1);
iKernel<<<grid, block>>>(d_rcv[1], d_src[0]);

  这些代码包含在simpleP2P_PingPong.cu文件中，以下输出表明这些核函数运行成功：

  如果GPU没有连接到相同的PCIe根节点上，或点对点访问被禁止，那么将会出现以下的错误信息：