Computación multi-GPU Cuda

Question

soy un novato en la programación multi-gpu y tengo algunas preguntas sobre la computación multi-gpu. Por ejemplo, tomemos el ejemplo del producto punto. Estoy ejecutando un subproceso de CPU que crea 2 matrices grandes A [N] y B [N]. Debido al tamaño de estas matrices, necesito dividir el cálculo de su producto de puntos en 2 GPU, ambas Tesla M2050 (capacidad de cálculo 2.0). El problema es que necesito calcular estos productos de puntos varias veces dentro de un do-loop controlado por mi hilo de CPU. Cada producto de punto requiere el resultado del anterior. He leído sobre la creación de 2 subprocesos diferentes que controlan las 2 GPU diferentes por separado (como se describe en cuda por ejemplo), pero no tengo idea de cómo sincronizar e intercambiar datos entre ellos. ¿Hay otra alternativa? Realmente agradecería cualquier tipo de ayuda/ejemplo. ¡Gracias de antemano!

Answer 1

Antes de CUDA 4.0, la programación multi-GPU requería programación de CPU con múltiples subprocesos. Esto puede ser un desafío especialmente cuando necesita sincronizar y/o comunicarse entre los subprocesos o las GPU. Y si todo su paralelismo está en su código GPU, entonces tener múltiples subprocesos de CPU puede aumentar la complejidad de su software sin mejorar el rendimiento más allá de lo que hace la GPU.

Por lo tanto, a partir de CUDA 4.0, puede programar fácilmente múltiples GPU desde un programa de host de subproceso único. Here are some slides I presented last year about this.

Programación de múltiples GPU puede ser tan simple como esto:

int numDevs = 0; 
cudaGetNumDevices(&numDevs); 
... 
for (int d = 0; d < numDevs; d++) { 
    cudaSetDevice(d); 
    kernel<<<blocks, threads>>>(args); 
} 

Para su ejemplo específico de productos punto, se puede usar thrust::inner_product como punto de partida. Yo haría eso para prototipos. Pero vea mis comentarios al final sobre los cuellos de botella de ancho de banda.

Dado que no proporcionó suficientes detalles sobre su bucle externo que ejecuta los productos dot varias veces, no intenté hacer nada con eso.

// assume the deviceIDs of the two 2050s are dev0 and dev1. 
// assume that the whole vector for the dot product is on the host in h_data 
// assume that n is the number of elements in h_vecA and h_vecB. 

int numDevs = 0; 
cudaGetNumDevices(&numDevs); 
... 
float result = 0.f; 
for (int d = 0; d < numDevs; d++) { 
    cudaSetDevice(d); 
    device_vector<float> vecA(h_vecA + d*(n/d), h_vecA + (d+1)*(n/d)-1); 
    device_vector<float> vecB(h_vecB + d*(n/d), h_vecB + (d+1)*(n/d)-1); 
    result += thrust::inner_product(vecA.begin(), vecA.end(), vecB.begin(), 0.f); 
} 

(Tengo que reconocer que la indexación anterior no es correcta si n no es un múltiplo par del numDevs, pero voy a dejar que la fijación como ejercicio para el lector. :)

Esto es simple y es un gran comienzo. Haz que funcione primero, luego optimiza.

Una vez que lo tengas funcionando, si todo lo que estás haciendo en los dispositivos es productos dot, verás que estás vinculado al ancho de banda, principalmente por PCI-e, y tampoco obtendrás concurrencia entre los dispositivos porque empujar :: inner_product es sincrónico debido a la lectura para devolver el resultado.Así que podrías usar cudaMemcpyAsync (el constructor device_vector usará cudaMemcpy). Pero el enfoque más fácil y probablemente más eficiente sería usar "copia cero": acceder directamente a la memoria del host (también se discute en la presentación de programación multi-gpu vinculada anteriormente). Como todo lo que hace es leer cada valor una vez y agregarlo a la suma (la reutilización paralela ocurre en una copia de memoria compartida), también puede leerlo directamente desde el host en lugar de copiarlo de host a dispositivo, y luego leer desde la memoria del dispositivo en el kernel. Además, desearía lanzar de forma asíncrona el kernel en cada GPU, para garantizar la máxima concurrencia.

se podría hacer algo como esto:

int bytes = sizeof(float) * n; 
cudaHostAlloc(h_vecA, bytes, cudaHostAllocMapped | cudaHostAllocPortable); 
cudaHostAlloc(h_vecB, bytes, cudaHostAllocMapped | cudaHostAllocPortable); 
cudaHostAlloc(results, numDevs * sizeof(float), cudaHostAllocMapped | cudaHostAllocPortable); 
// ... then fill your input arrays h_vecA and h_vecB 


for (int d = 0; d < numDevs; d++) { 
    cudaSetDevice(d); 
    cudaEventCreate(event[d])); 
    cudaHostGetDevicePointer(&dptrsA[d], h_vecA, 0); 
    cudaHostGetDevicePointer(&dptrsB[d], h_vecB, 0); 
    cudaHostGetDevicePointer(&dresults[d], results, 0); 
} 

... 

for (int d = 0; d < numDevs; d++) { 
    cudaSetDevice(d); 
    int first = d * (n/d); 
    int last = (d+1)*(n/d)-1; 
    my_inner_product<<<grid, block>>>(&dresults[d], 
             vecA+first, 
             vecA+last, 
             vecB+first, 0.f); 
    cudaEventRecord(event[d], 0); 
} 

// wait for all devices 
float total = 0.0f; 
for (int d = 0; d < devs; d++) { 
    cudaEventSynchronize(event[d]); 
    total += results[numDevs]; 
} 

Answer 2

Para crear varios hilos, puede usar OpenMP o pthreads. Para hacer lo que estás diciendo, parece que necesitarías crear y lanzar dos subprocesos (omp parallel section, o pthread_create), haz que cada uno haga su parte del cálculo y almacene su resultado intermedio en variables separadas del proceso por separado (recuerde, las variables globales se comparten automáticamente entre los hilos de un proceso, por lo que el hilo original podrá ver los cambios realizados por los dos hilos generados). Para que los subprocesos originales esperen a que los demás completen, sincronice (usando una operación global de barrera o unión de subprocesos) y combine los resultados en el subproceso original una vez que los dos subprocesos generados estén completos (si está dividiendo los arreglos por la mitad y calculando el producto escalar multiplicando los elementos correspondientes y realizando una reducción de suma global en las mitades, solo debería ser necesario agregar los dos resultados intermedios de los dos hilos generados).

También puede usar MPI o una horquilla, en cuyo caso la comunicación podría realizarse de forma similar a la programación de red ... tuberías/enchufes o comunicación y sincronización a través de (bloqueo) envía y recibe.