cuda - 多 GPU CUDA 推力

Question

我有一个使用 Thrust 目前在单个 GPU 上正常工作的 Cuda C++ 代码。我现在想为多 GPU 修改它。我有一个主机函数，其中包括许多对设备数组进行排序、复制、计算差异等的推力调用。我想使用每个 GPU 同时在它自己的(独立的)数组集上运行这个 Thrust 调用序列。我读过返回值的 Thrust 函数是同步的，但我可以使用 OpenMP 让每个主机线程调用一个在单独 GPU 上运行的函数(使用 Thrust 调用)吗？

例如(在浏览器中编码):

#pragma omp parallel for 
for (int dev=0; dev<Ndev; dev++){
   cudaSetDevice(dev);
   runthrustfunctions(dev);
}

void runthrustfunctions(int dev){
  /*lots of Thrust functions running on device arrays stored on corresponding GPU*/
 //for example this is just a few of the lines"

 thrust::device_ptr<double> pos_ptr = thrust::device_pointer_cast(particle[dev].pos);
 thrust::device_ptr<int> list_ptr = thrust::device_pointer_cast(particle[dev].list);
 thrust::sequence(list_ptr,list_ptr+length);
 thrust::sort_by_key(pos_ptr, pos_ptr+length,list_ptr);
 thrust::device_vector<double> temp(length);
 thrust::gather(list_ptr,list_ptr+length,pos_ptr,temp.begin());   
 thrust::copy(temp.begin(), temp.end(), pos_ptr);

}`

我认为我还需要将结构“particle[0]”存储在 GPU 0 上，将particle[1] 存储在 GPU 1 上等，我猜这是不可能的。一个选项可能是为每个 GPU 案例使用带有单独代码的“开关”。

我想知道这是正确的方法还是有更好的方法？谢谢

Answer 1

是的，您可以将推力和 OpenMP 结合使用。

这是一个完整的工作示例和结果:

$ cat t340.cu
#include <omp.h>
#include <stdio.h> 
#include <stdlib.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <thrust/copy.h>
#include <time.h>
#include <sys/time.h>

#define DSIZE 200000000

using namespace std;

int main(int argc, char *argv[])
{
    timeval t1, t2;
    int num_gpus = 0;   // number of CUDA GPUs

    printf("%s Starting...\n\n", argv[0]);

    // determine the number of CUDA capable GPUs
    cudaGetDeviceCount(&num_gpus);

    if (num_gpus < 1)
    {
        printf("no CUDA capable devices were detected\n");
        return 1;
    }

    // display CPU and GPU configuration
    printf("number of host CPUs:\t%d\n", omp_get_num_procs());
    printf("number of CUDA devices:\t%d\n", num_gpus);

    for (int i = 0; i < num_gpus; i++)
    {
        cudaDeviceProp dprop;
        cudaGetDeviceProperties(&dprop, i);
        printf("   %d: %s\n", i, dprop.name);
    }

    printf("initialize data\n");


    // initialize data
    typedef thrust::device_vector<int> dvec;
    typedef dvec *p_dvec;
    std::vector<p_dvec> dvecs;

    for(unsigned int i = 0; i < num_gpus; i++) {
      cudaSetDevice(i);
      p_dvec temp = new dvec(DSIZE);
      dvecs.push_back(temp);
      }

    thrust::host_vector<int> data(DSIZE);
    thrust::generate(data.begin(), data.end(), rand);

    // copy data
    for (unsigned int i = 0; i < num_gpus; i++) {
      cudaSetDevice(i);
      thrust::copy(data.begin(), data.end(), (*(dvecs[i])).begin());
      }

    printf("start sort\n");
    gettimeofday(&t1,NULL);

    // run as many CPU threads as there are CUDA devices
    omp_set_num_threads(num_gpus);  // create as many CPU threads as there are CUDA devices
    #pragma omp parallel
    {
        unsigned int cpu_thread_id = omp_get_thread_num();
        cudaSetDevice(cpu_thread_id);
        thrust::sort((*(dvecs[cpu_thread_id])).begin(), (*(dvecs[cpu_thread_id])).end());
        cudaDeviceSynchronize();
    }
    gettimeofday(&t2,NULL);
    printf("finished\n");
    unsigned long et = ((t2.tv_sec * 1000000)+t2.tv_usec) - ((t1.tv_sec * 1000000) + t1.tv_usec);
    if (cudaSuccess != cudaGetLastError())
        printf("%s\n", cudaGetErrorString(cudaGetLastError()));
    printf("sort time = %fs\n", (float)et/(float)(1000000));
    // check results
    thrust::host_vector<int> result(DSIZE);
    thrust::sort(data.begin(), data.end());
    for (int i = 0; i < num_gpus; i++)
    {
        cudaSetDevice(i);
        thrust::copy((*(dvecs[i])).begin(), (*(dvecs[i])).end(), result.begin());
        for (int j = 0; j < DSIZE; j++)
          if (data[j] != result[j]) { printf("mismatch on device %d at index %d, host: %d, device: %d\n", i, j, data[j], result[j]); return 1;}
    }
    printf("Success\n");
    return 0;

}
$ nvcc -Xcompiler -fopenmp -O3 -arch=sm_20 -o t340 t340.cu -lgomp
$ CUDA_VISIBLE_DEVICES="0" ./t340
./t340 Starting...

number of host CPUs:    12
number of CUDA devices: 1
   0: Tesla M2050
initialize data
start sort
finished
sort time = 0.398922s
Success
$ ./t340
./t340 Starting...

number of host CPUs:    12
number of CUDA devices: 4
   0: Tesla M2050
   1: Tesla M2070
   2: Tesla M2050
   3: Tesla M2070
initialize data
start sort
finished
sort time = 0.460058s
Success
$

我们可以看到，当我将程序限制为使用单个设备时，排序操作大约需要 0.4 秒。然后，当我允许它使用所有 4 台设备(在所有 4 台设备上重复相同的排序)时，整个操作只需要 0.46 秒，即使我们做了 4 倍的工作。

对于这种特殊情况，我碰巧使用 CUDA 5.0 和推力 v1.7 和 gcc 4.4.6 (RHEL 6.2)