uso básico multi-GPU

¿Cómo puedo usar dos dispositivos para mejorar, por ejemplo, el rendimiento del siguiente código (sum 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  #include  #include  #include  #include  #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<<>>( 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

    Desde que se lanzó CUDA 4.0, los cálculos multi-GPU del tipo sobre el que está preguntando 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 asignación de memoria de su código de 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) ); } 

    (descargo de responsabilidad: escrito en el navegador, nunca comstackdo, nunca probado, 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); } 

    (descargo de responsabilidad: escrito en el navegador, nunca comstackdo, nunca probado, uso bajo su propio riesgo).

    La sección de inicio de kernel de tu código podría verse más o menos así:

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

    (descargo de responsabilidad: escrito en el navegador, nunca comstackdo, nunca probado, uso bajo su propio riesgo).

    Tenga en cuenta que he agregado un argumento adicional a su llamada al kernel, ya que cada instancia del kernel puede ser llamada con un número diferente de elementos de la matriz para procesar. Te dejo a ti para resolver las modificaciones requeridas. Pero, una vez más, la idea básica es la misma: use cudaSetDevice para seleccionar una GPU determinada, luego ejecute kernels en ella de la manera habitual, y cada núcleo obtiene 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.