CUDA ( Compute Unified Device Architecture )

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

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

3. Wstęp

4. Tesla C870 • 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

5. 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

6. 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ą

7. 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

8. CUDA – model programowania

9. CUDA – model pamięci

10. 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

11. CUDA – model pamięci • Shared Memory

12. 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

13. 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)

14. 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

15. 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);

16. CUDA API • Konfiguracja wykonania gdzieDg, Db, Nssą: • Dg jest typudim3  wymiar i rozmiar grida Dg.x * Dg.y = ilość uruchamianych bloków; • Db jest typudim3 wymiar i rozmiar bloków Db.x * Db.y * Db.z = ilość wątków na blok; • Ns jest typusize_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.

17. CUDA API • Wbdowane zmienne • gridDimtypu dim3 wymiary grida. • blockIdxtypuuint3 number bloku w grid’zie • blockDimtypudim3 wymiary bloku • threadIdxis of type uint3 numer wątku w bloku

18. 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

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

20. Przykład – iloczyn skalarny int main(int argc, char *argv[]){ CUT_CHECK_DEVICE(); … h_A = (float *)malloc(DATA_SZ); … cudaMalloc((void **)&d_A, DATA_SZ); … cudaMemcpy(d_A, h_A, DATA_SZ, cudaMemcpyHostToDevice); … ProdGPU<<<BLOCK_N, THREAD_N>>>(d_C, d_A, d_B); … cudaMemcpy(h_C_GPU, d_C, RESULT_SZ, cudaMemcpyDeviceToHost); … CUDA_SAFE_CALL( cudaFree(d_A) ); free(h_A); … CUT_EXIT(argc, argv); } Program dla host’a

21. Przykład – iloczyn skalarny __global__ void ProdGPU(float *d_C, float *d_A, float *d_B){ __shared__ float t[THREAD_N]; __shared__ float r[SLICE_N]; const int I = threadIdx.x; for(int vec_n=blockIdx.x; vec_n<VECTOR_N; vec_n+=gridDim.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(); } 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]; } } Funkcja dla GPU (Kernel Function) • Parametry: • d_C: wskaźnik do wyniku (tj. tablicy) • d_A, d_B wskaźniki do danych (tablic) • Tablice lokalne: • t[]: wynkki8 pojedynczego wątku • r[]: używane do dodawania wyników segmentów • I: numer (Id) wątku w bloku

22. Przykład – iloczyn skalarny __global__ void ProdGPU(float *d_C, float *d_A, float *d_B){ __shared__ float t[THREAD_N]; __shared__ float r[SLICE_N]; const int I = threadIdx.x; for(int vec_n=blockIdx.x; vec_n<VECTOR_N; vec_n+=gridDim.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(); } 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]; } } Funkcja dla GPU • Uruchamiane dla każdej pary wektorów wejściowych • Zostanie uruchomione tylko raz, ponieważ: Grid dimension == number of vectors  vector number = block Id

23. Przykład – iloczyn skalarny __global__ void ProdGPU(float *d_C, float *d_A, float *d_B){ __shared__ float t[THREAD_N]; __shared__ float r[SLICE_N]; const int I = threadIdx.x; for(int vec_n=blockIdx.x; vec_n<VECTOR_N; vec_n+=gridDim.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(); } 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]; } } Funkcja dla GPU • Uruchamiane dla każdego segmentu wektorów wejściowych • Każdy wątek wylicza jeden iloczyn i zapamiętuje go

24. Przykład – iloczyn skalarny Funkcja dla GPU • Wyliczenie wyniku częściowego dla segmentu • Zapamiętanie wyniku częsciowego __global__ void ProdGPU(float *d_C, float *d_A, float *d_B){ __shared__ float t[THREAD_N]; __shared__ float r[SLICE_N]; const int I = threadIdx.x; for(int vec_n=blockIdx.x; vec_n<VECTOR_N; vec_n+=gridDim.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(); } 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]; } } t[0] += t[128] t[1] += t[129] t[0] += t[64] t[2] += t[130] t[1] += t[65] … t[0] += t[1] … … … t[64]+= t[127] t[127]+= t[255]

25. Przykład – iloczyn skalarny __global__ void ProdGPU(float *d_C, float *d_A, float *d_B){ __shared__ float t[THREAD_N]; __shared__ float r[SLICE_N]; const int I = threadIdx.x; for(int vec_n=blockIdx.x; vec_n<VECTOR_N; vec_n+=gridDim.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(); } 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]; } } Funkcja dla GPU • Dodanie wyników dla wszystkich segmentów • Zapisanie wyniku w pamięci