uso básico multi-GPU

Question

¿Cómo puedo usar dos dispositivos para mejorar, por ejemplo, el rendimiento del siguiente código (suma de vectores)? ¿Es posible usar más dispositivos "al mismo tiempo"? En caso afirmativo, ¿cómo puedo gestionar las asignaciones de los vectores en la memoria global de los diferentes dispositivos?

#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; 
} 

Gracias de antemano. Michele

Answer 1

Desde que se lanzó CUDA 4.0, los cálculos multi-GPU del tipo sobre el que pregunta son relativamente fáciles. Antes de eso, habría necesitado utilizar una aplicación de host de subprocesos múltiples con un subproceso de host por GPU y algún tipo de sistema de comunicación entre subprocesos con el fin de utilizar GPU mutliple dentro de la misma aplicación de host.

Ahora es posible hacer algo como esto para la parte de la asignación de memoria de su código 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: escrito en el navegador, no compilado, nunca se probó, su uso bajo su propio riesgo).

La idea básica aquí es que use cudaSetDevice para seleccionar entre dispositivos cuando está realizando operaciones en un dispositivo. Entonces, en el fragmento de arriba, he asumido dos GPU y memoria asignada en cada [(N/2) dobles en el primer dispositivo y N- (N/2) en el segundo].

La transferencia de datos desde el host al dispositivo podría ser tan simple como:

// 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: escrito en el navegador, no compilado, nunca se probó, el uso a su propio riesgo).

El núcleo sección de su código de lanzamiento a continuación, podría ser algo como:

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: escrito en el navegador, no compilado, no probado, el uso bajo su propio riesgo).

Tenga en cuenta que he agregado un argumento adicional a su llamada al kernel, porque cada instancia del kernel puede ser llamada con un número diferente de elementos de matriz para procesar. Te dejo a ti para resolver las modificaciones requeridas. Pero, de nuevo, la idea básica es la misma: use cudaSetDevice para seleccionar una GPU determinada, luego ejecute kernels en ella de la manera habitual, con cada kernel obteniendo sus propios argumentos únicos.

Debería poder unir estas partes para producir una aplicación simple multi-GPU. Hay muchas otras características que se pueden usar en versiones recientes de CUDA y hardware para ayudar a múltiples aplicaciones GPU (como el direccionamiento unificado, las instalaciones peer-to-peer son más), pero esto debería ser suficiente para comenzar. También hay una aplicación muLti-GPU simple en el SDK de CUDA que puede consultar para obtener más ideas.