Z racji tego, że ostatnio na uczelni zderzyłem się z tematem współbieżnych programów opartych o procesor karty graficznej to postanowiłem podzielić się z Wami tymi rewalacjami. Postaram się tutaj po krótce opisać podstawowe elementy i koncepcje, którymi rządzi się takie podejście do programowania. Od razu zawęże, że mowa tu będzie tylko o programowaniu kart graficznych opartych o architekturze CUDA (Compute Unified Device Architecture).
Zacznijmy od scharakteryzowania głównej różnicy. Mamy tutaj dwa procesory. Procesor CPU zwany w terminologii hostem oraz procesor karty graficznej (GPU) zwany urządzeniem. Ogólna zasada programowania sprowadza się do:
- uruchomienia programu na procesorze CPU,
- skopiowaniu potrzebnych danych do pamięci karty graficznej (tak GPU posiada własną pamięć L2, DRAM),
- wykonaniu kodu współbieżnego (z podziałem na bloki/wątki) przez procesor karty graficznej,
- skopiowaniu wyników z urządzenia (GPU) do pamięci operacyjnej RAM i dalsze kroki np. prezentacja wyników.
Podział kodu na część, która ma być wykonywana przez GPU określamy za pomocą słowa kluczowego __global__ przy deklaracji funkcji. Taka funkcja zostanie wywołana przez procesor CPU, i wykonana przez procesor karty graficznej. Podczas wywołania w nawiasach ostrych <<< >>> określamy na ilu blokach i ilu wątkach ma się uruchomić nasz kod.
__global__ void mykernel(void) {
}
int main(void) {
mykernel<<<1,1>>>();
printf("Hello World!\n");
return 0;
}Pomówmy teraz o alokacji pamięci, czyli o wymianie danych między procesorami. W przypadku zarządzania pamięcią GPU dostajemy funkcje cudaMalloc(), cudaFree(), cudaMemcpy() będące odpowiednikiem z języka C do zarządzania pamięcią operacyjną – malloc(), free(), memcpy(). Przy czym należy pamiętać, że wskaźniki operujące na pamięci urządzenia mogą być przetworzone do i z kodu na hoście. Nie można natomiast wykonywać dereferencji (odwołania się do wartości danej zmiennej) przez kod na hoście. Analogicznie dla wskaźników operujących na pamięci hosta.
int main(void) {
int a, b, c; // host copies of a, b, c
int *d_a, *d_b, *d_c; // device copies of a, b, c
int size = sizeof(int);
// Allocate space for device copies of a, b, c
cudaMalloc((void **)&d_a, size);
cudaMalloc((void **)&d_b, size);
cudaMalloc((void **)&d_c, size);
// Setup input values
a = 2;
b = 7;
// Copy inputs to device
cudaMemcpy(d_a, &a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, &b, size, cudaMemcpyHostToDevice);
// Launch add() kernel on GPU
add<<<1,1>>>(d_a, d_b, d_c);
// Copy result back to host
cudaMemcpy(&c, d_c, size, cudaMemcpyDeviceToHost);
// Cleanup
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
return 0;
}Do powyższego kodu można dopisać funkcje dodawania, która wykona się równolegle w ramach bloków. Koncepcja bloków, choć podobna do wątków ma zasadniczą różnice – bloki są od siebie odseparowane tj. nie mogą wpływać wzajemnie na swoją pamięć. W architekturze CUDA, bloki stanowią dodatkową abstrakcje tj. grupują wątki, które z kolei posiadają wspólną pamięć. Zmienna blockIdx.x dostarcza identyfikator konkretnego wątku. Stąd przykładowa definicja metody sumującej macierz do poprzedniego kodu mogła by wyglądać w następujący sposób.
__global__ void add(int *a, int *b, int *c) {
c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}Dla osób, które złapały bakcyla do programowania kart graficznych odsyłam po więcej informacji do strony producenta https://www.nvidia.com/docs/IO/116711/sc11-cuda-c-basics.pdf