有关多GPU编程

或许大家在写程序的时候,会遇到多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)。

接下来补充一个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进程将对应一个计算设备。

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

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