多GPU的CUDA执行

GPU间的并发：

在设备分配操作之前，我们首先获取服务器GPU的个数与编号，后为每个设备初始化主机数组的状态：分配代码需要的存储空间，计算时的中间变量等。

for(int i = 0; i < ngpus; i++){
    cudaSetDevice(i);
    initialData(h_A[i], iSize);
    initialData(h_B[i], iSize);
}

在所有的资源都被分配和初始化后，可以使用一个循环在多个设备之间分配数据和计算。

for(int i = 0; i < ngpus; i++){
    cudaSetDevice(i);
    
    cudaMemcpyAsync(d_A[i], h_A[i], iBytes, cudaMemcpyHostToDevice, stream[i]);
    cudaMemcpyAsync(d_B[i], h_B[i], iBytes, cudaMemcpyHostToDevice, stream[i]);

    iKernal<< <grid, block, 0, stream[i]> >> (d_A[i], d_B[i], d_C[i], iSize);
    
    cudaMemcpyAsync(gpuRef[i], d_C[i], iBytes, cudaMemcpyDeviceToHost, stream[i]);
}
cudaDeviceSynchronize();

这个循环会遍历多个GPU，为设备异步的复制输入数组，然后在相同的流中操作传输的元素以启动内核进行数据处理，数据处理结束后设备发送异步拷贝指令将结果从设备返回到主机。因为所有的函数都是异步的，因此控制全会立即返回到主机线程，因此，当任务仍在当前设备上执行时，切换到下一个设备时安全的。

GPU间的通信：

对于连接到同一条PCIE总线上的GPU，可以通过CUDA P2P API的支持实现GPU之间的点对点通信，该通信模式分为以下两种：

1. 点对点访问：在CUDA内核和GPU之间加载和存储地址，用于GPU交互处理
2. 点对点传输：在GPU之间复制数据，用于GPU并发

在使用点对点传输（显存复制）时，首先需要启动点对点之间的对等访问请求：

cudaError_t cudaDeviceCanAccessPeer(int* canAccessPeer, int device, int peerDevice);

收到支持对等访问的返回值后，我们便可以显示的异步复制设备上的数据，该步骤不经过主机端，且数据的传输会自动沿着PCIE的最短路径进行执行：

cudaError_t cudaMemcpyPeerAsync(void* dst, int dstDev, void* src,
                 int srcDev, size_t nBytes, cudaStream_t stream);

点对点的访问使用结束后不会自动关闭，需要被显示的禁用：

cudaError_t cudaDviceDisablePeerAccess(int peerDevice);