GPGPU Generalpurpose computing on graphics processing units CUDA

GPGPU - General-purpose computing on graphics processing units CUDA (Compute Unified Device Architecture) W. Bożejko

Plan • • • Wstęp Model programowania Model pamięci CUDA API Przykład – iloczyn skalarny

Wstęp

Tesla C 870 Produkt Tesla C 870 Obudowa ATX, 4. 38” x 12. 28” Ilość GPU Tesla 1 Dedykowana pamięć 1. 5 GB GDDR 3 Szczytowa wydajność Ponad 500 gigaflopów Precyzja obliczeń zmiennoprzecinkowych Interfejs pamięci Przepustowość pamięci Maksymalny pobór mocy Pojedyncza precyzja w standardzie IEEE 754 384 -bit 76. 8 GBps 170 W Interfejs systemowy PCI Express x 16 Dodatkowe źródła zasilania Liczba slotów Tak (2) 2 Chłodzenie wentylator • całkowity rozmiar pamięci globalnej 1, 61 GB • liczba multiprocesorów 16 • liczba rdzeni (procesorów) 128 • całkowity rozmiar pamięci stałej 65536 KB • całkowity rozmiar pamięci współdzielonej przypadającej na jeden blok 16384 KB • liczba rejestrów dostępna dla każdego bloku 8192 • częstotliwość zegara 1, 35 GHz

CUDA – model programowania • GPU jest widziane jako urządzenie obliczeniowe mogące wykonać część aplikacji która – musi być wykonana wielokrotnie – może być wyizolowana jako funkcja – działa niezależnie na różnych danych (model SIMD) • Taka funkcja może być skompilowana o wykonana na GPU

CUDA – model programowania • Blok wątków (Thread Block) – Wątki mogą kooperować • Mają szybką pamięć współdzieloną • Są zsynchronizowane • można je łatwo rozróżniać (mają Thread ID) – Blok może być 1, 2 lub 3 -wymiarową tablicą

CUDA – model programowania • Grid bloków wątków – Ograniczona ilość wątków w bloku – Pozwala wywołać większą liczbę wątków za pomocą jednego wywołania – Bloki są identyfikowane za pomocą block ID – Wymaga zmniejszenia kooperacji wątków – Bloki mogą być 1 lub 2 -wymiarowymi tablicami

CUDA – model programowania

CUDA – model pamięci

CUDA – model pamięci • Shared Memory – Wbudowana w chip • Znacznie szybsza niż pamięć lokalna i globalna • Tak szybka jak rejestry (jeśli nie ma konfliktów) • Dzielna na równej wielkości banki – Kolejne 32 -bitowe słowa są przypisane do kolejnych banków, – Każdy bank ma przepustowość (bandwidth) 32 bity na 1 cykl zegara

CUDA – model pamięci • Shared Memory

CUDA API • Rozszerzenie języka C – Kwalifikatory typu funkcji specyfikujące wykonanie na procesorze (host) lub na urządzeniu GPU – Kwalifikatory typu zmiennej specyfikujące rodzaj pamięci w GPU – Nowe składnia <<< mówiąca jak wykonać program na urządzeniu – Cztery wbudowane zmienne pamiętające rozmiary grid’a i bloku oraz numery bloku i wątku

CUDA API • Kwalifikatory typu funkcji __device__ • Wykonywane na GPU • Wywoływane tylko z GPU __global__ • Wykonywane na GPU • Wywoływane tylko z procesora głównego (host’a) __host__ • Wykonywane na host’cie, • Wywoływane tylko z procesora głównego (host’a)

CUDA API • Kwalifikatory typu zmiennych __device__ • Umieszone w pamięci globalnej • Widoczne przez cały czas działania programu • Dostępne dla wszystkich wątków w grid’zie oraz z hosta (poprzez runtime library) __constant__ (ewentulanie razem z __device__) • Umieszczone w pamięci stałej (constant memory space), • Widoczne przez cały czas działania programu • Dostępne dla wszystkich wątków w grid’zie oraz z hosta (poprzez runtime library) • __shared__ (ewentulanie razem z __device__) • Umieszczone w pamięci współdzielonej (shared memory) bloku danego wątku • Widoczne tak długo jak istnieje blok • Dostępne tylko dla wszystkich wątków w bloku

CUDA API • Konfiguracja wykonania – Musi być sprecyzowana dla kazdego wywołania funkcji typu __global__ – Definiuje rozmiary grid’a i bloków – Umieszczana pomiędzy nazwą funkcji a listą argumentów: funkcja: __global__ void Func(float* parameter); musi być wywołana tak: Func<<< Dg, Db, Ns >>>(parameter);

CUDA API • Konfiguracja wykonania gdzie Dg, Db, Ns są: – Dg jest typu dim 3 wymiar i rozmiar grida Dg. x * Dg. y = ilość uruchamianych bloków; – Db jest typu dim 3 wymiar i rozmiar bloków Db. x * Db. y * Db. z = ilość wątków na blok; – Ns jest typu size_t ilość bajtów w pamięci współdzielonej (shared memory) która jest dynamiczne alokowana dodatkowo do pamięci alokowanej statycznie • Ns jest opcjonalne; domyślnie 0.

CUDA API • Wbdowane zmienne – grid. Dim typu dim 3 wymiary grida. – block. Idx typu uint 3 number bloku w grid’zie – block. Dim typu dim 3 wymiary bloku – thread. Idx is of type uint 3 numer wątku w bloku

Przykład – iloczyn skalarny • Policzyć iloczyn skalarny – 32 par wektorów – Kożdy po 4096 elementów • Efektywna organizacja obliczeń: – grid składający się z 32 bloków – z 256 wątkami na blok • Otrzymamy 4096/265 = 16 segmentów na wektor

Przykład – iloczyn skalarny • Dane będą trzymane w GPU jako dwie tablice; wynik umieszczony zostanie w tablicy Vector A 0 Vector A 1 Vector B 0 Vector B 1 Vector AN-1 Vector BN-1 Results 0 to N-1 segment 0 • Każdy iloczyn par wektórw An, Bn będzie obliczany w segmentach, dodawanych do wyniku … … segment 1 … segment S-1 Vector A 0 Vector B 0 Partial results 0 to S-1 Results 0 Results 1

Przykład – iloczyn skalarny Program dla host’a int main(int argc, char *argv[]){ CUT_CHECK_DEVICE(); … h_A = (float *)malloc(DATA_SZ); … cuda. Malloc((void **)&d_A, DATA_SZ); … cuda. Memcpy(d_A, h_A, DATA_SZ, cuda. Memcpy. Host. To. Device); … Prod. GPU<<<BLOCK_N, THREAD_N>>>(d_C, d_A, d_B); … cuda. Memcpy(h_C_GPU, d_C, RESULT_SZ, cuda. Memcpy. Device. To. Host); … CUDA_SAFE_CALL( cuda. Free(d_A) free(h_A); … CUT_EXIT(argc, argv); } );

Przykład – iloczyn skalarny Funkcja dla GPU (Kernel Function) __global__ void Prod. GPU(float *d_C, float *d_A, float *d_B){ __shared__ float t[THREAD_N]; __shared__ float r[SLICE_N]; const int I = thread. Idx. x; • Parametry: – d_C: wskaźnik do wyniku (tj. tablicy) – d_A, d_B wskaźniki do danych (tablic) for(int vec_n=block. Idx. x; vec_n<VECTOR_N; vec_n+=grid. Dim. x){ int base = ELEMENT_N * vec_n; for(int slice = 0; slice < SLICE_N; slice++, base += THREAD_N){ t[I] = d_A[base + I] * d_B[base + I]; __syncthreads(); for(int stride = THREAD_N / 2; stride > 0; stride /= 2){ if(I < stride) t[I] += t[stride + I]; __syncthreads(); } • Tablice lokalne: – t[]: wynkki 8 pojedynczego wątku – r[]: używane do dodawania wyników segmentów • I: numer (Id) wątku w bloku if(I == 0) r[slice] = t[0]; } for(int stride = SLICE_N / 2; stride > 0; stride /= 2){ if(I < stride) r[I] += r[stride + I]; __syncthreads(); } if(I == 0) d_C[vec_n] = r[0]; } }

Przykład – iloczyn skalarny Funkcja dla GPU __global__ void Prod. GPU(float *d_C, float *d_A, float *d_B){ __shared__ float t[THREAD_N]; __shared__ float r[SLICE_N]; const int I = thread. Idx. x; • Uruchamiane dla każdej pary wektorów wejściowych for(int vec_n=block. Idx. x; vec_n<VECTOR_N; vec_n+=grid. Dim. x){ int base = ELEMENT_N * vec_n; for(int slice = 0; slice < SLICE_N; slice++, base += THREAD_N){ • Zostanie uruchomione tylko raz, ponieważ: t[I] = d_A[base + I] * d_B[base + I]; __syncthreads(); for(int stride = THREAD_N / 2; stride > 0; stride /= 2){ Grid dimension == number of vectors if(I < stride) t[I] += t[stride + I]; __syncthreads(); vector number = block Id } if(I == 0) r[slice] = t[0]; } for(int stride = SLICE_N / 2; stride > 0; stride /= 2){ if(I < stride) r[I] += r[stride + I]; __syncthreads(); } if(I == 0) d_C[vec_n] = r[0]; } }

Przykład – iloczyn skalarny Funkcja dla GPU __global__ void Prod. GPU(float *d_C, float *d_A, float *d_B){ __shared__ float t[THREAD_N]; __shared__ float r[SLICE_N]; const int I = thread. Idx. x; • Uruchamiane dla każdego segmentu wektorów wejściowych for(int vec_n=block. Idx. x; vec_n<VECTOR_N; vec_n+=grid. Dim. x){ int base = ELEMENT_N * vec_n; for(int slice = 0; slice < SLICE_N; slice++, base += THREAD_N){ t[I] = d_A[base + I] * d_B[base + I]; __syncthreads(); • Każdy wątek wylicza jeden iloczyn i zapamiętuje go for(int stride = THREAD_N / 2; stride > 0; stride /= 2){ if(I < stride) t[I] += t[stride + I]; __syncthreads(); } if(I == 0) r[slice] = t[0]; } for(int stride = SLICE_N / 2; stride > 0; stride /= 2){ if(I < stride) r[I] += r[stride + I]; __syncthreads(); } if(I == 0) d_C[vec_n] = r[0]; } }

Przykład – iloczyn skalarny Funkcja dla GPU __global__ void Prod. GPU(float *d_C, float *d_A, float *d_B){ __shared__ float t[THREAD_N]; __shared__ float r[SLICE_N]; const int I = thread. Idx. x; • Wyliczenie wyniku częściowego dla segmentu for(int vec_n=block. Idx. x; vec_n<VECTOR_N; vec_n+=grid. Dim. x){ int base = ELEMENT_N * vec_n; for(int slice = 0; slice < SLICE_N; slice++, base += THREAD_N){ t[I] = d_A[base + I] * d_B[base + I]; __syncthreads(); t[0] += t[128] t[1] += t[129] t[0] += t[64] t[2] += t[130] t[1] += t[65] for(int stride = THREAD_N / 2; stride > 0; stride /= 2){ … … … t[64]+= t[127] … if(I < stride) t[I] += t[stride + I]; __syncthreads(); t[0] += t[1] } if(I == 0) r[slice] = t[0]; t[127]+= t[255] } • Zapamiętanie wyniku częsciowego for(int stride = SLICE_N / 2; stride > 0; stride /= 2){ if(I < stride) r[I] += r[stride + I]; __syncthreads(); } if(I == 0) d_C[vec_n] = r[0]; } }

Przykład – iloczyn skalarny Funkcja dla GPU __global__ void Prod. GPU(float *d_C, float *d_A, float *d_B){ __shared__ float t[THREAD_N]; __shared__ float r[SLICE_N]; const int I = thread. Idx. x; • Dodanie wyników dla wszystkich segmentów for(int vec_n=block. Idx. x; vec_n<VECTOR_N; vec_n+=grid. Dim. x){ int base = ELEMENT_N * vec_n; for(int slice = 0; slice < SLICE_N; slice++, base += THREAD_N){ • Zapisanie wyniku w pamięci t[I] = d_A[base + I] * d_B[base + I]; __syncthreads(); for(int stride = THREAD_N / 2; stride > 0; stride /= 2){ if(I < stride) t[I] += t[stride + I]; __syncthreads(); } if(I == 0) r[slice] = t[0]; } for(int stride = SLICE_N / 2; stride > 0; stride /= 2){ if(I < stride) r[I] += r[stride + I]; __syncthreads(); } if(I == 0) d_C[vec_n] = r[0]; } }