utilizzo di base multi-GPU

Question

Come posso utilizzare due dispositivi per migliorare ad esempio le prestazioni del seguente codice (somma di vettori)? È possibile utilizzare più dispositivi "allo stesso tempo"? Se sì, come posso gestire le allocazioni dei vettori sulla memoria globale dei diversi dispositivi?

#include <stdio.h> 
#include <stdlib.h> 
#include <math.h> 
#include <time.h> 
#include <cuda.h> 

#define NB 32 
#define NT 500 
#define N NB*NT 

__global__ void add(double *a, double *b, double *c); 

//=========================================== 
__global__ void add(double *a, double *b, double *c){ 

    int tid = threadIdx.x + blockIdx.x * blockDim.x; 

    while(tid < N){ 
     c[tid] = a[tid] + b[tid]; 
     tid += blockDim.x * gridDim.x; 
    } 

} 

//============================================ 
//BEGIN 
//=========================================== 
int main(void) { 

    double *a, *b, *c; 
    double *dev_a, *dev_b, *dev_c; 

    // allocate the memory on the CPU 
    a=(double *)malloc(N*sizeof(double)); 
    b=(double *)malloc(N*sizeof(double)); 
    c=(double *)malloc(N*sizeof(double)); 

    // allocate the memory on the GPU 
    cudaMalloc((void**)&dev_a, N * sizeof(double)); 
    cudaMalloc((void**)&dev_b, N * sizeof(double)); 
    cudaMalloc((void**)&dev_c, N * sizeof(double)); 

    // fill the arrays 'a' and 'b' on the CPU 
    for (int i=0; i<N; i++) { 
     a[i] = (double)i; 
     b[i] = (double)i*2; 
    } 

    // copy the arrays 'a' and 'b' to the GPU 
    cudaMemcpy(dev_a, a, N * sizeof(double), cudaMemcpyHostToDevice); 
    cudaMemcpy(dev_b, b, N * sizeof(double), cudaMemcpyHostToDevice); 

    for(int i=0;i<10000;++i) 
     add<<<NB,NT>>>(dev_a, dev_b, dev_c); 

    // copy the array 'c' back from the GPU to the CPU 
    cudaMemcpy(c, dev_c, N * sizeof(double), cudaMemcpyDeviceToHost); 

    // display the results 
    // for (int i=0; i<N; i++) { 
    //  printf("%g + %g = %g\n", a[i], b[i], c[i]); 
    // } 
    printf("\nGPU done\n"); 

    // free the memory allocated on the GPU 
    cudaFree(dev_a); 
    cudaFree(dev_b); 
    cudaFree(dev_c); 
    // free the memory allocated on the CPU 
    free(a); 
    free(b); 
    free(c); 

    return 0; 
} 

Grazie in anticipo. Michele

Answer 1

Da quando è stato rilasciato CUDA 4.0, i calcoli a più GPU del tipo di cui si sta parlando sono relativamente facili. Prima di ciò, sarebbe necessario utilizzare un'applicazione host multi-thread con un thread host per GPU e una sorta di sistema di comunicazione inter-thread per utilizzare GPU multiple all'interno della stessa applicazione host.

Ora è possibile fare qualcosa di simile per la parte allocazione di memoria del codice host:

double *dev_a[2], *dev_b[2], *dev_c[2]; 
const int Ns[2] = {N/2, N-(N/2)}; 

// allocate the memory on the GPUs 
for(int dev=0; dev<2; dev++) { 
    cudaSetDevice(dev); 
    cudaMalloc((void**)&dev_a[dev], Ns[dev] * sizeof(double)); 
    cudaMalloc((void**)&dev_b[dev], Ns[dev] * sizeof(double)); 
    cudaMalloc((void**)&dev_c[dev], Ns[dev] * sizeof(double)); 
} 

(disclaimer: scritto in del browser, mai compilato, mai testato, l'uso a proprio rischio).

L'idea di base qui è che si utilizza cudaSetDevice per selezionare tra i dispositivi quando si eseguono operazioni su un dispositivo. Quindi nel frammento sopra, ho assunto due GPU e ho allocato memoria su ogni [(N/2) doppio sul primo dispositivo e N- (N/2) sul secondo].

Il trasferimento dei dati dall'host al dispositivo potrebbe essere semplice come:

// copy the arrays 'a' and 'b' to the GPUs 
for(int dev=0,pos=0; dev<2; pos+=Ns[dev], dev++) { 
    cudaSetDevice(dev); 
    cudaMemcpy(dev_a[dev], a+pos, Ns[dev] * sizeof(double), cudaMemcpyHostToDevice); 
    cudaMemcpy(dev_b[dev], b+pos, Ns[dev] * sizeof(double), cudaMemcpyHostToDevice); 
} 

(disclaimer: scritto in del browser, mai compilato, mai testato, l'uso a proprio rischio).

Il kernel lanciare sezione del vostro codice potrebbe poi essere simile:

for(int i=0;i<10000;++i) { 
    for(int dev=0; dev<2; dev++) { 
     cudaSetDevice(dev); 
     add<<<NB,NT>>>(dev_a[dev], dev_b[dev], dev_c[dev], Ns[dev]); 
    } 
} 

(disclaimer: scritto in del browser, mai compilato, mai testato, l'uso a proprio rischio).

Si noti che ho aggiunto un argomento aggiuntivo alla chiamata del kernel, poiché ogni istanza del kernel può essere chiamata con un numero diverso di elementi di matrice da elaborare. Lo lascerò a te per elaborare le modifiche richieste. Ma, ancora una volta, l'idea di base è la stessa: usare cudaSetDevice per selezionare una data GPU, quindi eseguire i kernel su di essa nel modo normale, con ogni kernel che ottiene i propri argomenti univoci.

Dovreste essere in grado di mettere insieme queste parti per produrre una semplice applicazione multi-GPU. Ci sono molte altre funzionalità che possono essere utilizzate nelle recenti versioni di CUDA e hardware per supportare più applicazioni GPU (come l'indirizzamento unificato, le funzionalità peer-to-peer sono più), ma questo dovrebbe essere sufficiente per iniziare. C'è anche una semplice applicazione muLti-GPU nell'SDK CUDA che puoi guardare per altre idee.