4.4 Technologia CUDA
4.4.4 Założenia programistyczne
4.4.3 Schemat przetwarzania w technologii CUDA
Na rys. 15 przedstawiony został schemat przetwarzania w technologii CUDA. Obejmuje ona trzy elementy: procesor główny komputera (CPU), pamięć operacyjną oraz kartę graficzną (GPU) wraz z wyodrębnioną pamięcią własną. Na początku działania każdego programu w technologii CUDA dane z pamięci operacyjnej komputera są kopiowane do karty pamięci karty graficznej (1). Następnie procesor główny komputera wysyła instrukcję do procesora graficznego (2), zgodnie z którymi następuje równoległe przetwarzanie danych (3). Po uzyskaniu wyników obliczeń są one kopiowane do pamięci operacyjnej komputera (4), skąd są dostępne dla użytkownika.
Rys. 15 Schemat przetwarzania w technologii CUDA.
4.4.4 Założenia programistyczne
Programy napisane z wykorzystaniem technologii CUDA są kompilowane przez NVCC (kompilator NVidia). W rzeczywistości kod źródłowy jest przez NVCC rozdzielany na dwie części. Część przeznaczona dla urządzenia, jakim jest karta graficzna lub specjalizowana jednostka obliczeniowa, kompilowana jest przez kompilator NVidia, natomiast część kodu przeznaczona do wykonania przez procesor główny komputera hosta (CPU) kompilowana jest przez przeznaczony do tego kompilator standardowo używany w systemie, np. gcc.
W języku C/C++ deklaracje funkcji przeznaczonych do wykonania w urządzeniu poprzedzone są słowem kluczowym __global__. Funkcje te choć wykonywane w urządzeniu, są wywoływane z poziomu hosta. Noszą one nazwę kerneli. Kernel wywoływany jest w kodzie programu z użyciem potrójnych nawiasów ostrych.
Zarządzanie pamięcią
Pamięć w urządzeniu, podobnie jak w hoście musi być alokowana, zwalniana i inicjowana. Służą do tego funkcje cudaMalloc(), cudaFree() oraz cudaMemCopy() będące odpowiednikami funkcji Malloc(), Free() i MemCopy() w języku C. Wskaźniki do pamięci urządzenia mogą być przekazywane do i z hosta, jednak nie można ich z poziomu jego kodu
46
zmieniać. Podobnie jest ze wskaźnikami do pamięci komputera, które można przekazywać do i z urządzenia, jednak nie mogą być zmieniane z poziomu kodu karty graficznej czy specjalizowanej jednostki obliczeniowej.
Obliczenia równoległe
Na przykładzie dodawania dwóch wektorów o długości N pokazany zostanie sposób w jaki implementowane są obliczenia równoległe w technologii CUDA. Gdyby dodać te dwa wektory w sposób tradycyjny, funkcja dodająca do siebie po jednym elemencie każdego wektora musiałaby zostać wywołana N razy. W obliczeniach równoległych, wywołanych zostanie natomiast N instancji takiej funkcji działających jednocześnie, a więc jak wspomniano wyżej wywołanych zostanie N kerneli. Każde wywołanie kernela rozpoczyna działanie nowego bloku. Ponadto zawsze określa się ile pojedynczych wątków przypada na jeden blok. Bloki z kolei grupowane są w większe struktury zwane z angielska gridami. Każdy wywołany kernel jest indeksowany. Jego indeks można określić za pomocą indeksu wątku, indeksu bloku oraz indeksu gridu. Istnieją różne sposoby indeksowania, które wybiera się odpowiednio do implementowanego zadania obliczeniowego. W poniższym przykładzie dla jasności pominięty został indeks gridu, który nie jest wymagany. Kernel dodający do siebie po jednym elemencie z dwóch wektorów będzie miał postać:
__global__ void add(int *a, int *b, int *c) { c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x]; }
Oznaczenie blockIdx oznacza numer bloku, a .x numer wątku w tym bloku. Wywołanie dodawania dwóch wektorów z użyciem powyższej funkcji odbywa się następująco:
add<<< N, 1 >>>();
Pierwszy parametr N to liczba bloków, które mają zostać wywołane, a drugi to liczba wątków przypadających na jeden blok. W przypadku rozważanej funkcji dodającej dwa wektory, wywołanych zostanie N bloków funkcji, a więc tyle ile wynosi długość dodawanych wektorów. Bloki wykonywane są równocześnie, a każdy z nich doda do siebie elementy wektorów o indeksie blockIdx.x. Indeks ten zależy od numeru bloku i numeru wątku. Ponieważ w add<<< N, 1 >>>() wywołane zostało N bloków po jednym wątku na blok, indeks N wątków doda wszystkie wartości wektorów. W ten sposób zaimplementowane zostało równoległe dodawanie wektorów.
Cały kod programu łączącego kod wykonywany w hoście i w urządzeniu przedstawiony został poniżej.
47 #define N 512
#include <stdlib.h>
__global__ void add(int *a, int *b, int *c) {
c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x]; }
int main(void) {
int *a, *b, *c; // Wskaźniki do pamięci w hoście
int *d_a, *d_b, *d_c; // Wskaźniki do pamięci w urządzeniu
int size = N * sizeof(int);
// Alokowanie pamięci dla zmiennych w urządzeniu cudaMalloc((void**)&d_a, size);
cudaMalloc((void**)&d_b, size); cudaMalloc((void**)&d_c, size);
// Alokowanie pamięci dla zmiennych w hoście a = (int*)malloc(size);
b = (int*)malloc(size); c = (int*)malloc(size);
// Inicjalizacja zmiennych for(int i = 0 ; i< N ; i++){ a[i]= rand();
b[i]= rand(); c[i]=0; }
// Kopiowanie wartości zmiennych do pamięci urządzenia cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice); cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);
// Wywołanie kernela (N bloków po jednym wątku) add<<<N,1>>>(d_a, d_b, d_c);
// Kopiowanie wyników z urządzenia do hosta cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);
48 // Zwolnienie pamięci
free(a); free(b); free(c);
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
return 0; }
Długość dodawanych wektorów wynosi N = 512. Na początku programu alokowana jest pamięć dla dwóch wektorów (tablica zmiennych typu int), które zostaną do siebie dodane oraz dla wektora, który stanowił będzie wynik tej operacji. Alokacja odbywa się zarówno dla pamięci w hoście jak i w urządzeniu. Następnie pamięć ta jest inicjalizowana – każdemu elementowi wektorów przypisywane są losowe wartości. Warto nadmienić, że polecenia alokujące pamięć jej nie inicjalizują, dlatego wektor wyjściowy zostaje zainicjalizowany poprzez przypisanie każdemu jego elementowi wartości zero. W tym konkretnym przypadku nie jest to konieczne, jednak przy prostej modyfikacji programu mogłoby być to źródłem błędów i dlatego jest to dobra praktyka programistyczna. Po zainicjalizowaniu zmiennych, ich wartości są kopiowane do pamięci urządzenia. Następnie wywoływany jest kernel add<<< N, 1 >>>() w 512 blokach, po jednym wątku na blok. Po zakończeniu działania wszystkich 512 funkcji działających równolegle, pamięć zawierająca końcowy wynik, a więc tablicę c jest kopiowana z pamięci urządzenia do pamięci hosta. Na koniec działania programu, pamięć zarówno w urządzeniu jak i hoście jest zwalniana.
49