CUDA C编程（三十二）多GPU间细分计算

在 多 设 备 上 分 配 内 存
在从主机向设备分配计算任务之前，需要确定在当前中有多少可用的GPU:

int ngpus;
cudaGetDeviceCount(&ngpus);
printf("CUDA-capable devices: %i
",ngpus);

一旦GPU的数量已经确定，接下来就需要为多个设备声明主机内存、设备内存、流和事件。保存这些变量的一个简单方法是使用数组，声明如下：

float *d_A[NGPUS], *d_B[NGPUS], *d_C[NGPUS];
float *h_A[NGPUS], *h_B[NGPUS], *hostRef[NGPUS], *gpuRef[NGPUS];
cudaStream_t stream[NGPUS];

在向量加法的例子中，元素的总输入大小为16M，所有设备平分，给每个设备isize个元素：

int size = 1 << 24;
int iSize = size / ngpus;

设备上一个浮点向量的字节大小按如下方式进行计算：

size_t iBytes = iSize * sizeof(float);

现在可以分配主机和设备内存了，为每个设备创建CUDA流，代码如下：

for(int i = 0; i < ngpus; i++)
{
   cudaSetDevice(i);
   cudaMalloc((void **)&d_A[i], iBytes);
   cudaMalloc((void **)&d_B[i], iBytes);
   cudaMalloc((void **)&d_C[i], iBytes);
    
   cudaMallocHost((void **)&h_A[i], iBytes);
   cudaMallocHost((void **)&h_B[i], iBytes);
   cudaMallocHost((void **)&hostRef[i], iBytes);
   cudaMallocHost((void **)&gpuRef[i], iBytes);
    
   cudaStreamCreate(&stream[i]);
}

请注意，分配锁页(固定)主机内存是为了在设备和主机之间进行异步数据传输。同时，在分配任何内存或创建任何流之前，使用上面提到的cudasetDevice函数在每次循环迭代的开始，设置当前设备。

单 主 机 线 程 分 配 工 作
在设备间分配操作之前，需要为每个设备初始化主机数组的状态：

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]);
   
   kernel<<<grid, block ,0 , stream[i], iBytes, cudaMemcpyDeviceToHost, stream[i]);
    
   cudaMemcpyAsync(gpuRef[i], d_C[i], iBytes, cudaMemcpyDeviceToHost, stream[i]);
}
cudaDeviceSynchronize();

这个循环遍历多个GPU，为设备异步地复制输入数组。然后在相同的流中操作iSize个数据元素以便启动内核。最后，设备发出的异步拷贝命令，把结果从内核返回到主机。因为所有的函数都是异步的，所有控制会立即返回到主机线程。因此，当任务仍在当前设备上运行时，切换到下一个设备是安全的。
编 译 和 执 行
从Wrox.com中下载simpleMultiGPU.cu文件，其中含有完整的多GPU向量加法示例。用以下命令编译它：$ nvcc -O3 simpleMultiGPU.cu -o simpleMultiGPU,simpleMultiGPU函数的采样输出为：

通过将命令行选项设为1，尝试只用一个GPU运行它，代码如下：

尽管使用双倍数量的GPU时，运行时间并没有减少一半，但是仍然取得了显著的性能提升。使用nvprof可以获得每个设备行为的更多细节：$ nvprof --print-gpu-trace ./simpleMultiGPU,在有两个M2090GPU的系统上产生的输出总结如下：

在只有一个M2090GPU上运行，产生的输出如下所示：