Multi-gpu CUDA Thrust

Question

У меня есть код Cuda C++, который использует Thrust, который в настоящее время работает правильно на одном графическом процессоре. Теперь я хотел бы изменить его для multi-gpu. У меня есть функция хоста, которая включает в себя ряд вызовов Thrust, которые сортируют, копируют, вычисляют различия и т. Д. На массивах устройств. Я хочу использовать каждый графический процессор для выполнения этой последовательности вызовов Thrust на своем собственном (независимом) наборе массивов одновременно. Я прочитал, что функции Thrust, возвращающие значения, являются синхронными, но я могу использовать OpenMP для того, чтобы каждый поток хоста вызывал функцию (с вызовами 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); 

} `

Я думаю, что также необходима структура "частицы [0]", чтобы быть сохранены на GPU 0, частица [1] на GPU 1 и т. Д., И я думаю, это невозможно. Опцией может быть использование «switch» с отдельным кодом для каждого случая 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)