Multi-GPU Computazione Cuda

Question

Sono un principiante nella programmazione multi-GPU e ho alcune domande sul multi-GPU computing. Ad esempio, prendiamo l'esempio del punto-prodotto. Sto eseguendo un thread della CPU che crea 2 array di grandi dimensioni A [N] e B [N]. A causa delle dimensioni di questi array, ho bisogno di dividere il calcolo del loro prodotto dot in 2 GPU, entrambi Tesla M2050 (capacità di calcolo 2.0). Il problema è che ho bisogno di calcolare questi dot-product più volte all'interno di un do-loop controllato dal mio thread della CPU. Ogni punto-prodotto richiede il risultato del precedente. Ho letto sulla creazione di 2 thread diversi che controllano separatamente 2 GPU diverse (come descritto su cuda dall'esempio) ma non ho idea di come sincronizzare e scambiare i dati tra di loro. C'è un'altra alternativa? Apprezzerei molto ogni tipo di aiuto/esempio. Grazie in anticipo!

Answer 1

Prima di CUDA 4.0, la programmazione multi-GPU richiedeva la programmazione della CPU multi-thread. Questo può essere difficile soprattutto quando è necessario sincronizzare e/o comunicare tra i thread o le GPU. E se tutto il tuo parallelismo è nel tuo codice GPU, avere più thread della CPU può aumentare la complessità del tuo software senza migliorare le prestazioni oltre a ciò che fa la GPU.

Quindi, a partire da CUDA 4.0, è possibile programmare facilmente più GPU da un programma host a thread singolo. Here are some slides I presented last year about this.

di programmazione più GPU può essere semplice come questo:

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

Per esempio specifico di prodotto scalare, è possibile utilizzare thrust::inner_product come punto di partenza. Lo farei per la prototipazione. Ma vedi i miei commenti alla fine sui colli di bottiglia della larghezza di banda.

Poiché non hai fornito abbastanza dettagli sul tuo ciclo esterno che esegue i prodotti punto più volte, non ho provato a fare nulla con quello.

// 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); 
} 

(ammetto che l'indicizzazione di cui sopra non è corretto se n non è un multiplo pari numDevs, ma lascio che fissa come un esercizio per il lettore. :)

Questo è semplice ed è un ottimo inizio. Scaricalo prima, quindi ottimizza.

Una volta che hai funzionato, se tutto ciò che stai facendo sui dispositivi è dot-products, scoprirai di essere legato alla larghezza di banda, principalmente da PCI-e, e inoltre non otterrai la concorrenza tra i dispositivi perché la spinta :: inner_product è sincrono a causa della lettura indietro per restituire il risultato ..Quindi è possibile utilizzare cudaMemcpyAsync (il costruttore device_vector utilizzerà cudaMemcpy). Ma l'approccio più semplice e probabilmente più efficiente sarebbe quello di usare "zero copy" - accedere direttamente alla memoria dell'host (discusso anche nella presentazione di programmazione multi-gpu collegata sopra). Dato che tutto ciò che stai facendo è leggere ogni valore una volta e aggiungerlo alla somma (il riutilizzo parallelo avviene in una copia di memoria condivisa), potresti leggerlo direttamente dall'host piuttosto che copiarlo da host a dispositivo e quindi leggere dalla memoria del dispositivo nel kernel. Inoltre, vorrai i lanci asincroni del kernel su ogni GPU, per garantire la massima concorrenza.

Si potrebbe fare qualcosa di simile:

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

Per creare più thread, è possibile utilizzare OpenMP o pthreads. Per fare ciò di cui parli, sembra che tu abbia bisogno di creare e lanciare due thread (omp parallel section, o pthread_create), ognuno faccia la sua parte del calcolo e memorizzi il suo risultato intermedio in variabili separate del processo-wIDE (Ricordate, le variabili globali vengono automaticamente condivise tra i thread di un processo, quindi il thread originale sarà in grado di vedere le modifiche apportate dai due thread generati). Per ottenere i thread originali in attesa che gli altri si completino, sincronizzare (usando un'operazione di global barrier o di join del thread) e combinare i risultati nel thread originale dopo che i due thread spawn sono completi (se si dividono gli array a metà e calcolando il prodotto punto moltiplicando gli elementi corrispondenti ed eseguendo una riduzione globale delle somme sulle metà, dovrebbe essere solo necessario aggiungere i due risultati intermedi dai due thread generati).

È possibile utilizzare anche MPI o fork, nel qual caso la comunicazione può essere eseguita in modo simile alla programmazione di rete ... tubi/prese o comunicazione e sincronizzazione tramite (blocco) invia e riceve.