Obliczanie CU Multi-GPU

Question

Jestem początkującym programistą multi-gpu i mam kilka pytań na temat obliczeń wielogigowych. Na przykład, weźmy przykład dot-produktu. Używam wątku procesora, który tworzy 2 duże tablice A [N] i B [N]. Ze względu na rozmiar tych tablic muszę podzielić obliczenia ich iloczynu punktowego na 2 GPU, zarówno Tesla M2050 (zdolność obliczeniowa 2.0). Problem polega na tym, że muszę kilka razy obliczyć te produkty-punkty w pętli do-do-kontrolowanej przez mój wątek CPU. Każdy produkt-kropka wymaga wyniku poprzedniego. Czytałem o tworzeniu 2 różnych wątków, które kontrolują 2 różne procesory graficzne oddzielnie (jak to opisano na przykładzie cuda), ale nie mam pojęcia, jak synchronizować i wymieniać między nimi dane. Czy jest inna alternatywa? Naprawdę doceniam każdy rodzaj pomocy/przykładu. Z góry dziękuję!

Answer 1

Przed CUDA 4.0, programowanie wielu GPU wymagało wielowątkowego programowania CPU. Może to być trudne, zwłaszcza gdy zachodzi potrzeba synchronizacji i/lub komunikacji między wątkami lub procesorami graficznymi. A jeśli cały twój paralelizm jest w twoim kodzie GPU, to posiadanie wielu wątków procesora może dodać do złożoności twojego oprogramowania bez poprawy wydajności poza to, co robi GPU.

Tak więc, począwszy od wersji CUDA 4.0, można z łatwością programować wiele GPU z programu z pojedynczym gwintem. Here are some slides I presented last year about this.

Programowanie wielu GPU może być tak proste, jak to:

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

dla konkretnego przykładu iloczyn skalarny, można użyć thrust::inner_product jako punkt wyjścia. Zrobiłbym to dla prototypowania. Ale zobacz moje komentarze na końcu o wąskich gardłach przepustowości.

Ponieważ nie podano wystarczająco dużo szczegółów na temat zewnętrznej pętli, która uruchamia wiele produktów, nie próbowałem nic z tym zrobić.

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

(muszę przyznać, że indeksowanie powyżej nie jest poprawne, jeśli n nie jest nawet wielokrotnością numDevs, ale zostawię ustalenie, że jako ćwiczenie dla czytelnika. :)

Jest to proste i jest świetnym początkiem. Zacznij działać, a następnie zoptymalizuj.

Po uruchomieniu, jeśli wszystko, co robisz na urządzeniach, to produkty dot, okaże się, że jesteś związany przepustowością - głównie przez PCI-e, a także nie uzyskasz współbieżności między urządzeniami, ponieważ :: inner_product jest synchroniczny ze względu na odczyt z powrotem, aby zwrócić wynik ..Więc możesz użyć cudaMemcpyAsync (konstruktor device_vector użyje cudaMemcpy). Ale łatwiejszym i bardziej efektywnym podejściem byłoby użycie "zerowej kopii" - bezpośredni dostęp do pamięci hosta (omówionej również w przedstawionej powyżej prezentacji dotyczącej programowania multi-GPU). Ponieważ wszystko, co robisz, czyta każdą wartość raz i dodaje ją do sumy (ponowne użycie równoległe dzieje się w kopii z dzieloną pamięcią), możesz równie dobrze przeczytać ją bezpośrednio z hosta, zamiast kopiować ją z hosta do urządzenia, a następnie odczytać to z pamięci urządzenia w jądrze. Ponadto, chciałbyś asynchroniczne uruchamianie jądra na każdym GPU, aby zapewnić maksymalną współbieżność.

Można zrobić coś takiego:

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

Aby utworzyć kilka wątków, można użyć OpenMP lub pthreads. Aby zrobić to, o czym mówisz, wydaje się, że musisz wykonać i uruchomić dwa wątki (omp równoległy lub pthread_create), każdy z nich wykona swoją część obliczeń i zapisze wynik pośredni w oddzielnych zmiennych procesowych (Przypomnijmy, zmienne globalne są automatycznie dzielone między wątki procesu, więc oryginalny wątek będzie mógł zobaczyć zmiany dokonane przez dwa zarodkowane wątki). Aby uzyskać oryginalne wątki oczekiwania na ukończenie innych, zsynchronizuj (używając globalnej bariery lub operacji łączenia nici) i połącz wyniki w oryginalnym wątku po ukończeniu dwóch odrodzonych wątków (jeśli dzielisz tablice na pół i obliczając produkt kropki poprzez pomnożenie odpowiednich elementów i wykonanie globalnej redukcji sumy na połówkach, powinno być konieczne tylko dodanie dwóch pośrednich wyników z dwóch zarodkowanych wątków).

Można również użyć MPI lub widełek, w takim przypadku komunikacja może być wykonana w sposób podobny do programowania sieci ... rury/gniazda lub komunikacja i synchronizacja przez (blokowanie) wysyła i odbiera.