CUDA (Compute Unified Device Architecture)

1. CUDA (Compute Unified Device Architecture)

2. Wzrost mocy obliczeniowej GPU i CPU Wniosek: ze względu na olbrzymią moc obliczeniową karty graficznej coraz częściej są wykorzystywane w celu wykonywania coraz większej klasy algorytmów równoległych nie związanych z grafiką

3. Podstawowe różnice między architekturą GPU i CPU Przykładowe GPU - GeForce 8800 GTX:16 multiprocesorówKażdy multiprocesor zawiera: 8 rdzeni 16 KB szybkiej pamięci 8192 rejestry 8 KB cache dla stałych + 8KB cache dla tekstur

4. Architektura GPU (G80 i późniejsze) Każdy multiprocesor działa w trybie SIMD: w każdym cyklu każdy rdzeń (procesor) wykonuje tą samą instrukcję, ale na różnych danych

5. Czym jest CUDA? Język programowania oraz trzypoziomowe API Tylko dla kart NVIDII Tylko pod Windows XP i Linux Rozwiązanie ATI-AMD: Close To Metal – assembler

6. Co CUDA wnosi do programowania GPGPU? • Wygodny, podobny do C język programowania równoległego, który nie wymaga znajomości grafiki komputerowej • Operacje na intach, floatach i doublach • Swobodę zapisów do pamięci GPU

7. Co CUDA wnosi do programowania GPGPU? • Komunikacja wątków poprzez szybką pamięć dzieloną

8. Podstawowe pojęcia • Host i Device • Kernel • Strumieńmożna zsynchronizować CPU z danym strumieniem(CPU czeka na zakończenie wykonywania wszystkich operacji w strumieniu) cudaStreamSynchronize()można zsynchronizować CPU ze wszystkimi strumieniamicudaThreadSynchronize()

9. Podstawowe pojęcia 4. Wątki, Bloki, Siatka bloków Uruchomienie kernela wymaga podania: liczby i organizacji wątków w bloku (struktura 1D, 2D lub 3D) liczby i organizacji bloków w siatce bloków (struktura 1D lub 2D)Wątki/bloki aktywne – wykonywane w tym samym czasie na przemian lub równolegle jeden blok może się wykonywać tylko na jednym multiprocesorze maksymalnie 512 wątków na blokmaksymalnie 768 aktywnych wątków na multiprocesormaksymalnie 8 aktywnych bloków na multiprocesor pamięć i rejestry multiprocesora dzielone pomiędzy wszystkie aktywne wątki

10. Wątki, Bloki, Siatki Bloków

11. Pamięć • rejestry • pamięć globalna • pamięć lokalna • pamięć stałych • pamięć tekstur • pamięć dzielona

12. Pamięć - rejestry Rejestry- 8192 rejestry na multiprocesor dzielone pomiędzy wszystkie aktywne wątki- użycie mniejszej liczby rejestrów na wątek pozwala jednocześnie obsługiwać więcej wątków na multiprocesorze - większa liczba wątków na multiprocesorze pozwala ukryć opóźnienia związane z czekaniem na dane z pamięci globalnej- kompilator często generuje nie optymalny kod, który niepotrzebnie zużywa dodatkowe rejestry- można wymusić przy kompilacji ograniczenie liczby używanych rejestrów

13. Pamięć - globalna Pamięć globalna- brak cache- koszt dostępu 400-600 cykli- dostęp do 4/8/16 bajtów struktury wyrównanej do 4/8/16 bajtów w jednej operacji- half-warp (16 kolejnych wątków bloku) adresujący spójny i wyrównany blok 64/128/256 bajtów kilkukrotnie szybszy od dostępu do bloku niespójnego

14. Przykłady:szybki dostęp sekwencyjny

15. Przykłady:wolny dostęp niesekwencyjny

16. Pamięć - lokalna Pamięć lokalna - blok przestrzeni pamięci globalnej- brak kontroli ze strony programisty- Kompilator automatycznie zapisuje w niej zmienne lokalne. Ponieważ wszystkie wątki bloku wykonują ten sam program i mają te same zmienne lokalne, to dostępy do tej pamięci zawsze są szybkimi dostępami sekwencyjnymi.

17. Pamięć - stałych • Pamięć stałych- blok przestrzeni pamięci globalnej • tylko do odczytu- 64 KB pamięci- 8 KB cache stałych dla każdego multiprocesora- jeśli wszystkie 16 kolejnych wątków bloku czyta ten sam adres i wartość jest w cache’u to dostęp tak szybki jak do rejestrów

18. Pamięć - tekstur Pamięć tekstur- blok przestrzeni pamięci globalnej - tylko do odczytu- 8KB cache tekstur dla każdego multiprocesora- specjalne tryby adresowania (normalizacja, zawijanie współrzędnych)- możliwa interpolacja (dwu-)(trzy-)liniowa - ograniczone formaty danych

19. Pamięć - tekstur Filtrowanie liniowe tekstur

20. Pamięć - dzielona Pamięć dzielona- 16 KB dla każdego multiprocesora - wątki jednego bloku mogą wymieniać informacje poprzez pamięć dzieloną i mechanizm synchronizacji (synchronizacja w obrębie kernela)- wątki różnych bloków mogą wymieniać informacje tylko poprzez pamięć globalną i mechanizm synchronizacji z CPU (synchronizacja tylko między kernelami)- jeśli kilka wątków half-warpa odwołuje się do tego samego 32-bitowego słowa, to wykonywany jest tylko jeden dostęp do pamięci (jeden odczyt i rozesłanie informacji lub jeden zapis zakończony sukcesem dla pewnego z wątków) - 16 banków = 16 wątków half-warpa; indeks banku = (addr/4)%16 (kolejne 32-bitowe słowa pamięci dzielonej należą do kolejnych banków; co 16-te słowo 32-bitowe należy do tego samego banku) - koszt dostępu do pamięci dla half-warpa to liczba konfliktów banków

21. Przykład:brak konfliktówjedna operacja dostępu do pamięci

22. Przykład:brak konfliktówjedna operacja dostępu do pamięci

23. Przykład po lewej:2 operacje dostępu do pamięciPrzykład po prawej: 8 operacji dostępu do pamięci

24. Wydajność operacji Dostępne (wybrane) operacje:4 cykle:+, * dla floatów+ dla intówmin, max, (int), (float), &, |, ^, ~, <, ==, <=, min, max16 cykli:1/x, 1/sqrtf(x), logf(x)* dla intów32 cykle:sinf, cosf, sincosf, tanf, expf36 cykli:/ dla floatówdrogie:operacje na doublach/ i % dla intówoperacje atomowedostępy do pamięci 400-600 cykli

25. Kod w języku CUDA Kod dzieli się na wykonywany na CPU i wykonywany na GPU. __global__ - funkcje wykonywane na GPU, wywoływane z CPU__device__ - funkcje wykonywane na GPU, wywoływane z GPU (nie ma rekurencji, domyślnie zawsze są inline)__host__ - funkcje wykonywane na CPU, wywoływane przez CPU (równoważne braku oznaczenia)

26. Kod w języku CUDA • Kod funkcji uruchamianych na GPU jest kodem C z ograniczeniami:- można korzystać tylko ze standardowej biblioteki matematycznej CUDy • funkcje __global__ zawsze są typu void • wbudowane typy- specjalne zmienne tylko do odczytu: blockIdx, threadIdx, blockDim, gridDim- oznaczenia zmiennych __device__, __constant__, __shared__- specjalna instrukcja __syncthreads()- wywołania funkcji __global__

27. Kompilator języka CUDA Kompilator nvcc:- możliwość skompilowania do kodu PTX (assembler GPU), który można potem podejrzeć- możliwość utworzenia od razu kodu wykonywalnego- możliwość kompilacji kodu uruchamianego na GPU do pliku binarnego (cubin), który można potem załadować i uruchamiać z dowolnego programu w C/C++ (podobnie jak biblioteki dll)

28. 1 #define ACCUM_N 1024 2 3 __global__ void scalarProdGPU( 4 float *d_C, 5 float *d_A, 6 float *d_B, 7 int vectorN, 8 int elementN 9 ){ 10 __shared__ float accumResult[ACCUM_N]; 11 12 for(int vec = blockIdx.x; vec < vectorN; vec += gridDim.x){ 13 int vectorBase = vec * elementN; 14 int vectorEnd = vectorBase + elementN; 15 16 for(int iAccum = threadIdx.x; iAccum < ACCUM_N; iAccum += blockDim.x){ 17 float sum = 0; 18 19 for(int pos = vectorBase + iAccum; pos < vectorEnd; pos += ACCUM_N) 20 sum += d_A[pos] * d_B[pos]; 21 22 accumResult[iAccum] = sum; 23 } 24 25 for(int stride = ACCUM_N / 2; stride > 0; stride >>= 1){ 26 __syncthreads(); 27 for(int iAccum = threadIdx.x; iAccum < stride; iAccum += blockDim.x) 28 accumResult[iAccum] += accumResult[stride + iAccum]; 29 } 30 31 if(threadIdx.x == 0) d_C[vec] = accumResult[0]; 32 } 33 }

29. 50 int main(int argc, char **argv){ 51 float *h_A, *h_B, *h_C; 52 float *d_A, *d_B, *d_C; 53 54 CUT_DEVICE_INIT(argc, argv); 55 56 h_A = (float *)malloc(DATA_SZ); 57 h_B = (float *)malloc(DATA_SZ); 58 h_C = (float *)malloc(RESULT_SZ); 59 cudaMalloc((void **)&d_A, DATA_SZ); 60 cudaMalloc((void **)&d_B, DATA_SZ); 61 cudaMalloc((void **)&d_C, RESULT_SZ); 62 63 // Tutaj powinno być wypełnienie h_A i h_B danymi 64 65 cudaMemcpy(d_A, h_A, DATA_SZ, cudaMemcpyHostToDevice); 66 cudaMemcpy(d_B, h_B, DATA_SZ, cudaMemcpyHostToDevice); 67 68 cudaThreadSynchronize(); 69 scalarProdGPU<<<128, 256>>>(d_C, d_A, d_B, 70 VECTOR_N, ELEMENT_N); 71 cudaThreadSynchronize(); 72 73 cudaMemcpy(h_C, d_C, RESULT_SZ, cudaMemcpyDeviceToHost); 74 75 cudaFree(d_C); 76 cudaFree(d_B); 77 cudaFree(d_A); 78 free(h_C); 79 free(h_B); 80 free(h_A); 81 82 CUT_EXIT(argc, argv); 83 } 34 #include <stdio.h> 35 #include <stdlib.h> 36 #include <time.h> 37 #include <string.h> 38 #include <cutil.h> 39 40 const int VECTOR_N 41 = 256; 42 const int ELEMENT_N 43 = 4096; 44 const int DATA_N 45 = VECTOR_N * ELEMENT_N; 46 const int DATA_SZ 47 = DATA_N*sizeof(float); 48 const int RESULT_SZ = 49 = VECTOR_N*sizeof(float);

30. Korzystanie z wielu kart graficznych - Możliwość korzystania z wielu kart graficznych w jednym komputerze (przy wyłączonym SLI)- Wsparcie dla obliczeń na klastrach z wieloma kartami graficznymi.- Zmiana rozdzielczości może spowodować przerwanie działania programów i utratę danych w pamięci karty graficznej (jeśli brakuje pamięci do zainicjalizowania ekranu w nowej rozdzielczości). - Tesla – karty graficzne NVIDII bez wyjścia video (!)

31. Komunikacja Host-Device - PCIe 2.0 16x ma przepustowość 8 GB/s- CUDA 1.1 wprowadza (ponoć) full-duplex (jednoczesny zapis do karty i odczyt z karty)- Operacje przesyłania danych asynchroniczne z punktu widzenia CPU (poprzez DMA).- Z punktu widzenia GPU, jeśli operacja kopiowania dotyczy i-tego strumienia, to bloki wątków z i-tego strumienia nie mogą się wykonywać, ale mogą dla każdego innego strumienia

32. Debugowanie Możliwe skompilowanie programu w trybie emulacji GPU na CPU. Jest to emulacja, a nie symulacja zatem pewne błędy nie zostaną zauważone:- różnica w obliczeniach zmiennoprzecinkowych - wątki w trybie emulacji na jednordzeniowym komputerze nie działają równolegleMakro __DEVICE_EMULATION__Ponieważ w trybie emulacji wszystkie funkcje są wykonywane na CPU, to w funkcjach __device__ i __global__ można korzystać ze wszystkich funkcji C/C++ (np. printf) oraz odwoływać się do zmiennych zarówno związanych z CPU jak i GPU w programie

33. Współpraca z OpenGL / DirectX OpenGLMożna czytać i zapisywać pamięć obiektów OpenGL:- pixel buffer- vertex buffer GLuint bufferObj; cudaGLRegisterBufferObject(bufferObj); GLuint bufferObj; float* devPtr; cudaGLMapBufferObject((void**)&devPtr, bufferObj); DirectXMożna czytać i zapisywać pamięć obiektów DirectX:- LPDIRECT3DVERTEXBUFFER9 - LPDIRECT3DSURFACE9

34. Biblioteki CUBLAS (Basic Linear Algebra Subprograms) - operacje na dowolnych rozmiarów wektorach i macierzach w pojedynczej lub podwójnej precyzji o elementach rzeczywistych lub zespolonych - BLAS1 – operacje wektor-wektor - norma Euklidesowa - iloczyn skalarny - obroty Givensa - BLAS2 – operacje macierz-wektor - * - rozwiązanie Ax=b - BLAS3 – operacje macierz-macierz - * - rozwiązanie AX=B CUFFT (Fast Fourier Transform)- szybka transformata Fouriera w 1D, 2D lub 3D, dla liczb rzeczywistych lub zespolonychCUDPP (Data Parallel Primitives)- segmentowane sumy progresywne

35. Zastosowania Przykłady: 1. Zastąpienie vertex i geometry shaderów2. Przetwarzanie obrazów - histogramy - filtry 3. Haszowanie i szyfrowanie 4. Kompresja MP35. Sortowanie 6. Systemy cząsteczkowe, N-body system7. Symulacja cieczy i gazów8. Detekcja kolizji i symulacja fizyki 9. k najbliższych cząstek10. Smith-Waterman - podobieństwo między łańcuchami białek

39. LCP Algorithms for Collision Detection

40. Bibliografiahttp://www.nvidia.com/object/cuda_home.htmlhttp://www.gpgpu.org/Bibliografiahttp://www.nvidia.com/object/cuda_home.htmlhttp://www.gpgpu.org/