有关多GPU编程

加速计算 CUDA
1

或许大家在写程序的时候，会遇到多GPU编程的问题，这里我简单的介绍一下方法：

  1. 主线程用cudaSetDevice()函数切换不同的设备
  2. linux下使用pthread实现多GPU

以上两种方法都可以实现多GPU的编程，但是区别还是很大的！比如使用第一种方法，有两块GPU的情况下，我们把任务分在不同的GPU上，执行的时候需要：



cudaSetDevice(0);

lanuch_kernel1<<<...>>>(...);

cudaSetDevice(1);

lanuch_kernel2<<<...>>>(...);

cudaSetDevice(0);

cudaMemcpy(...);

cudaSetDevice(1);

cudaMemcpy(...);

这样虽然在内核执行的时候是并行的，但是在内存拷贝的时候是串行的。因为cudaMemcpy是同步拷贝函数，CPU执行到这里必须等待GPU完成所有工作后才能进行内存拷贝。

接下来换种方法，在linux下使用pthread来实现多GPU并行。


pthread_t *thread = (pthread_t *)malloc(sizoef(pthread_t)*devCnt);

for(int i=0; i<devCnt; i++)

pthread_create(&thread[i], NULL, lanuch_kernel, (void*)param);
 //注意这里的lanuch_kernel并非CUDA内核函数，而是一个线程调用的函数，参数将通过param传递到函数内

for(int i=0; i<devCnt; i++)

pthread_join(thread[i], NULL);

在当今多核CPU遍布天下的时代，每个线程会占用一个核心来执行，那么此时内核会并行执行，而内存拷贝也是并行的，消耗的时间会比方法1要少！（pthread的用法请各位自行查阅资料~）

以上是经过我实际测试的，大家可以参考！

除此以外还有OPENMP和MPI的方式，也可以在单节点使用多GPU（OPENMP没试过，我只用MPI）。

2

接下来补充一个MPI的多GPU方式，假设一个节点上有N个设备：

MPI_Comm_rank(MPI_COMM_WORLD, &rank);
MPI_Comm_size(MPI_COMM_WORLD, &size);

int dev = rank % size;

//在这里指定设备
cudaSetDevice(dev);
...

通过上述方式，那么一个节点上便可以使用多个MPI进程，每个MPI进程将对应一个计算设备。

3

我想问个问题，就是cudaSetDevice这个切换函数本身开销如何？
cudaMemcpy这个本来是可以用流异步函数cudaMemcpyasync的，无需等待。
多线程内存访问时需要加锁等很繁琐，如果cudaSetDevice这个函数本身开销可以忽略的话是不是单线程更好啊？

4

如果用stream的话，我想应该也是没问题的，至于开销如何我倒没试过。。。
cudaSetDevice的开销非常小，多线程的话如果每个线程分配区域不冲突的话，应该也不需要互斥访问的！