Programowanie równoległe i rozproszone – opracowanie lab nr 2

GPGPU (ang. General-Purpose computing on Graphics Processor Units – obliczenia ogólnego przeznaczenia na układach GPU) - technika, dzięki której GPU, zwykle zajmujący się tylko obliczeniami związanymi z grafiką komputerową, umożliwia wykonywanie obliczeń ogólnego przeznaczenia tak jak CPU. Dzięki temu wiele obliczeń, głównie obliczenia równoległe, można przeprowadzić znacznie szybciej.

Procesor jest przede wszystkim jednostką wykonującą i podejmującą decyzje, zgodnie z zaleceniami oprogramowania.

Procesory wideo są znacznie lepsze w wykonywaniu rutynowych zadań. GPU mają duże ilości ALU w porównaniu do procesorów. W wyniku tego, mogą one wykonywać duże ilości zadań matematycznych w odróżnieniu od CPU (źródło: http://infocoin.pl/tech/teoria/cpu-vs-gpu/).

CPU składa się z kilku rdzeni zoptymalizowanych pod sekwencyjne szeregowe działanie, podczas gdy GPU ma w znacznym stopniu równoległą architekturę składają się z tysięcy małych, bardziej efektywnych rdzeni stworzonych do obsługi wielokrotnych zadań równocześnie (za http://www.nvidia.com/object/what-is-gpu-computing.html).

Aktualnie istnieją dwa kompleksowe środowiska do programowania procesorów GPU (ATI Stream, nVidia CUDA). Równolegle powstaje OpenCL — otwarty standard do obliczeń równoległych.

CUDA (ang. Compute Unified Device Architecture) - opracowana przez firmę Nvidia uniwersalna architektura procesorów wielordzeniowych (głównie kart graficznych) umożliwiająca wykorzystanie ich mocy obliczeniowej do rozwiązywania ogólnych problemów numerycznych w sposób wydajniejszy niż w tradycyjnych, sekwencyjnych procesorach ogólnego zastosowania.

CUDA posiada oparte na języku programowania C środowisko programistyczne wysokiego poziomu, w którego skład wchodzą m.in. specjalny kompilator (nvcc), debugger (cuda-gdb), profiler oraz interfejs programowania aplikacji. Istnieją biblioteki do Python, Fortran, Java, C# oraz Matlab. Wspiera Windows, Linux oraz Mac Os X (od wersji 2.0).

Zalety i wady:

• Zalety:

• Język programowania oparty na językach C/C++, w tym pełna obsługa szablonów C++.

• Model pamięci procesora ściśle odpowiadający architekturze sprzętowej, co umożliwia świadome, efektywne wykorzystywanie dostępnych zasobów GPU, w tym pamięci współdzielonej.

• Kod uruchamiany na GPU może odczytywać i zapisywać dane z dowolnego adresu w pamięci GPU.

• Pełna kompatybilność programów – napisany dziś program wykonywalny ma w przyszłości działać bez żadnych zmian na coraz wydajniejszych procesorach graficznych posiadających coraz większą liczbę rdzeni, rejestrów, pamięci operacyjnej i innych zasobów.

• Dostępność na wszystkich kartach firmy NVIDIA począwszy od serii GeForce 8 w tym Quadro oraz Tesla.

• Obsługa procesorów ARM (od wersji 6.5).

• Wady:

• CUDA korzysta z podzbioru języka C++. Nie można na przykład definiować statycznych zmiennych wewnątrz funkcji, a funkcje mogą mieć tylko stałą liczbę parametrów.

• Dla liczb zmiennoprzecinkowych podwójnej precyzji (dostępnych w nowszych procesorach) istnieją pewne odstępstwa od standardu w zakresie zaokrąglania liczb.

• Przepustowość i opóźnienia magistrali PCI-Express łączącej CPU i GPU mogą być wąskim gardłem w przypadku przesyłania dużej ilości danych.

• CUDA jest dostępna jedynie dla kart graficznych produkowanych przez firmę NVIDIA

Jądra (ang. kernels) i wątki w CUDA

Współbieżne bloki aplikacji są wykonywane na urządzeniu jako jądra (kernels). Tylko jedno jądro może się wykonywać w danej chwili. Wiele wątków (tysiące) wykonuje to samo jądro. Różnice pomiędzy wątkami CUDA a CPU: - wątki CUDA są „bardzo lekkie” (bardzo krótki czas tworzenia, natychmiastowe przełączanie), - CUDA używa tysiące wątków do osiągnięcia wydajności, wielordzeniowe CPU może obsługiwać tylko kilka wątków.

Słowniczek: - urządzenie = GPU, - host = CPU, - kernel – funkcja uruchamiana na urządzeniu.

Jądro CUDA jest wykonywane przez tablicę wątków. Każdy wątek uruchamia ten sam kod. Każdy wątek posiada ID używany do obliczania adresów pamięci oraz podejmowania decyzji.

Siatka (grid) bloków z wątkami

Jądro uruchamia siatkę bloków z wątkami. Wątki w obrębie bloku korzystają ze wspólnego obszaru pamięci oraz mogą być synchronizowane. Wątki w różnych blokach nie mogą współpracować.

Adres bloku w siatce może być dwuwymiarowy (jest strukturą z polami x oraz y). Adres wątku w bloku może być dwuwymiarowy (jest strukturą z polami x, y). Na powyższym rysunku adres jest jednowymiarowy — pola y, z są równe 1. Zatem jeżeli wątek operuje na wektorze 0, ..., N*M to numer elementu wektora może być obliczony przy pomocy wyrażenia – numer_wątku_w_bloku+rozmiar_bloku*numer_bloku.

Skalowalność

Sprzęt może zaszeregować dowolny blok na dowolnym procesorze. Brak założeń o kolejności wykonywania bloków.

Model pamięci w architekturze CUDA

• 
  Global memory (odczyt i zapis) – wolna i niecache'owalna, host może odczytywać i zapisywać do tej pamięci

• Texture memory (tylko odczyt) – pamięć cache zooptymalizowana do wzorców z dostępem dwuwymiarowym

• Constant memory – przechowywane są tam stałe oraz argumenty jądra, wolna ale cache'owalna

• Shared memory – mała, szybka, współdzielona przez blok

• Local memory – używana na wszystko, co się nie zmieściło w rejestrach wątku, część pamięci globalnej więc wolna i niecache'owalna

• Registers – szybka, każdy wątek posiada swoje

CUDA SDK

• Pliki źródłowe kompilowane za pomocą nvcc mogą zawierać kod hosta jak również kod urządzenia

• nvcc rozdziela kod dla urządzenia oraz kod dla hosta

• kod dla urządzenia kompiluje do postaci assemblerowej (PTX kod) i/lub binarnej postaci (obiekt cubin),

• modyfikuje kod hosta zamieniając znaczniki <<<...>>> występujące przy wywoływaniu jądra na potrzebne funkcje środowiska CUDA ładujące i uruchamiające każde skompilowane jądro z kodu PTX i/lub obiektu cubin.

• Zmodyfikowany plik hosta jest traktowany jako plik w języku C i jest zostawiany do kompilacji za pomocą innego narzędzia lub do skompilowania podczas ostatniego etapu kompilacji nvcc.

• Aplikacja następnie może:

• zlinkować do skompilowanego pliku hosta (najpopularniejszy przypadek)

• lub zignorować zmodyfikowany plik hosta i użyć API sterownika CUDA w celu załadowania i wykonania kodu PTX lub obiektu cubin.

Wbudowane typy wektorów

• int1, float2, double3, char4, … - typy wektorów zbudowane na podstawie typów całkowitych i zmiennoprzecinkowych. Możliwość stworzenia maksymalnie czterowymiarowego wektora, dostęp do składowych poprzez zmienne x, y, z i w. Generowanie: „make_<type name>”.

  int2 make_int2(int x, int y);

• dim3 – wektor liczb całkowitych bazujący na typie uint3, który używany jest to określania wymiarów. Jeśli, któryś z elementów pozostawimy niezdefiniowany to domyślnie przyjmie on wartość 1.

dim3 grid( 512 ); // 512 x 1 x 1 dim3 block( 1024, 1024 ); // 1024 x 1024 x 1

Wbudowane zmienne

Wbudowane zmienne określają wymiary siatki oraz blok wątków i indeks bloku oraz wątku w bloku. Zmienne te dostępne są tylko w obrębie funkcji wykonywanych w urządzeniu (GPU).

• gridDim – zmienna typu dim3, zawiera wymiary siatki.

• blockIdx – zmienna typu uint3, zawiera indeks bloku w siatce.

• blockDim – zmienna typu dim3, zawiera wymiary bloku.

• threadIdx – zmienna typu uint3, zawiera indeks wątku w bloku.

int ThreadIDX= threadIdx.x+blockDim.x*blockIdx.x;

Wykonywanie kodu na GPU

Jądra są funkcjami języka C z pewnymi ograniczeniami:

• Nie mogą odwoływać się do pamięci CPU.

• Muszą zwracać void jako wynik.

• Nie mogą przyjmować zmiennej liczby argumentów.

• Nie mogą być rekurencyjne.

• Nie mogą deklarować zmiennych statycznych.

Parametry są kopiowane automatycznie z CPU na GPU. Kod jąder umieszcza się w plikach

.cu.

Deklarowane z kwalifikatorem __global__ oznaczającym funkcję wywoływaną

z CPU i wykonującą się na GPU. Inne kwalifikatory to:

• __device__ oznacza funkcję wywoływaną z GPU i wykonującą się na GPU.

• __host__ oznacza funkcję wywoływaną z CPU w wykonującą się na CPU (domyślnie)

• __device__ i __host__ mogą być łączone.

• __global__ i __host__ nie mogą być łączone.

Przykład deklaracji:

__global__ void doSomethingOnDevice(...) { …. }

Wywołanie jądra

Wywoływanie jądra odbywa się w podobny sposób do wywołania zwykłej funkcji w C, z tym że dodatkowo musimy zadeklarować wymiar siatki, wymiar każdego bloku wątków i opcjonalnie liczbę bajtów w przestrzeni wspólnej, która jest dynamicznie alokowana przez blok dodatkowo do już przydzielonej statycznie pamięci (domyślnie 0) używając nawiasów <<< … , …[, …] >>>.

int blocksize=256; dim3 dimBlock (blocksize); dim3 dimGrid( ceil( N / (float)blocksize) ); inc_gpu<<<dimGrid, dimBlock>>>(d_a, N);

Kwalifikatory typów zmiennych w CUDA

Deklaracja zmiennej	Pamięć	Zasięg	Czas życia taki jak
__device__ __local__ int LocalVar;	lokalna	wątek	wątek
__device__ __shared__ int SharedVar;	współdzielona	blok	blok
__device__ int GlobalVar;	globalna	siatka bloków	aplikacja
__device__ __constant__ int ConstantVar;	stała	siatka bloków	aplikacja

Opcja __device__ jest opcjonalna, kiedy używana jest z __local__, __shared__ lub __constant__.

Zmienne bez żadnego kwalifikatora zapisywane są w rejestrach wątku (skalary oraz wbudowane typy wektorów). Te dane, które nie zmieszczą się w rejestrach lub tablice zapisywane są do pamięci lokalnej.

Kiedy deklarujemy zmienną w pamięci współdzielonej jako zewnętrzną tablice w następujący sposób:

extern __shared__ float shared[];

rozmiar tablicy określany jest podczas uruchamiania jądra (parametr trzeci wywołania jądra). Wszystkie zmienne zadeklarowane w ten sposób mają początek w tym samym adresie w pamięci.

__global__ void kernel(…) { … extern __shared__ float sData[]; … } int main(void) { … smBytes = blockSize*sizeof(float); kernel<<<nBlocks, blockSize, smBytes>>>(…); … }

Wskaźniki mogą wskazywać tylko na pamięć za alokowaną lub zadeklarowaną w pamięci globalnej GPU.

Gdzie deklarować zmienne?

Alokacja pamięci GPU

CPU zarządza pamięcią GPU:

• cudaMalloc(void ** pointer, size_t nbytes)

• cudaMemset(void * pointer, int value, size_t count)

• cudaFree(void* pointer)

int n = 1024; int nbytes = 1024*sizeof(int); // Wskaźnik na pamięć GPU; int * d_a = 0; // cudaMalloc( (void**)&d_a, nbytes ); // Próba odwołania się do pamięci wskazywanej przez d_a // w kodzie CPU prowadzi do błędu ochrony pamięci cudaMemset( d_a, 0, nbytes); cudaFree(d_a);

Kopiowanie danych

cudaMemcpy(void *dst, void *src, size_t nbytes,enum cudaMemcpyKind direction);

„Direction” określa lokalizację (urządzenie lub host) zmiennych src lub dst.

Blokuje wątek CPU – wznawia wątek jak wszystkie dane zostaną przekopiowane.

Nie rozpoczyna kopiowania, dopóki wcześniejsze wywołania CUDA nie zostały skończone.

• enum cudaMemcpyKind

• cudaMemcpyHostToDevice

• cudaMemcpyDeviceToHost

• cudaMemcpyDeviceToDevice

  Funkcja cudaMemcpyAsync() nie blokuje kodu CPU

Synchronizacja CPU z GPU

• Wywołania jąder są asynchroniczne

• Sterowanie powraca od razu do CPU bez oczekiwania na zakończenie wątków jądra.

• Jądro jest kolejkowane i wykona się po zakończeniu wszystkich poprzednich wywołań CUDA.

• Funkcja cudaMemcpy jest synchroniczna

• Kopiowanie rozpocznie się po wykonaniu poprzednich wywołań CUDA.

• Sterowanie powraca do CPU po wykonaniu kopii.

• Kopie pomiędzy pamięciami GPU są asynchroniczne.

• Funkcja cudaThreadSynchronize() wstrzymuje kod CPU do momentu wykonania wszystkich kolejkowanych wywołań CUDA.

  Synchronizacja wątków w bloku

• CUDA umożliwia współpracę wątków w ramach bloku. Do tego celu służy funkcja void _syncthreads().

• Synchronizacja barierowa: każdy wątek w bloku oczekuje, aż wszystkie inne wątki osiągną punkt wywołania funkcji _syncthreads().

• Bariera dotyczy wyłącznie wątków w bloku, nie w siatce. Tylko wątki w obrębie danego bloku mogą korzystać z pamięci współdzielonej.

• Stosowana w celu uniknięcia wyścigów przy dostępie do pamięci wspólnej.

• Jeżeli występuje w kodzie warunkowym np. instrukcja if, to musimy zadbać aby wszystkie wątki ją wywołały.

Atomowe operacje w kodzie GPU

Operacje atomowe na danych typu int/unsigned w pamięci globalnej.

• add, sub, min, max

• and, or, xor

• inc, dec

atomicAdd()

Kody błędów CUDA

• Wszystkie funkcje API, z wyjątkiem wywołań jąder zwracają kody błędów.

• Wyrażenie typu cudaError_t

• Funkcja cudaError_t cudaGetLastError(void) zwraca kod ostatniego błędu

• char* cudaGetErrorString(cudaError_t code) zwraca ciąg znaków zakończony zerem zawierający opis błędu w języku angielskim

printf(“%s\n”, cudaGetErrorString( cudaGetLastError() ) );

Przykładowe biblioteki CUDA

• CUBLAS — podzbiór bibliotek BLAS, operacje ma macierzach i wektorach

• CUFFT — szybka transformata Fourier'a (jednowymiarowa i dwuwymiarowa)

Do zadania:

Filtracja obrazów — filtracja obrazów jest zaliczana do metod cyfrowego przetwarzania sygnałów. Filtracja jest operacją matematyczną na pikselach obrazu źródłowego, w której wyniku uzyskiwany jest nowy, przekształcony obraz. Filtrację określa się jako przekształcenie kontekstowe, gdyż dla wyznaczenia nowej wartości piksela obrazu docelowego potrzebna jest informacja z wielu pikseli obrazu źródłowego.

Filtracja stosowana jest przeważnie jako metoda wydobycia z oryginalnego obrazu szeregu informacji w celu ich dalszej obróbki. Informacjami takimi mogą być: położenie krawędzi, pozycje rogów obiektów itp. Innym zastosowaniem filtracji jest usuwanie szumów (filtr medianowy i inne) lub rozmycie obrazu (filtry uśredniające, Gaussa). Filtrację można przeprowadzać zarówno w dziedzinie przestrzennej, jak i częstotliwościowej. Filtracje w dziedzinie przestrzennej uzyskuje się wykorzystując operacje splotu. W dziedzinie częstotliwości odpowiednikiem splotu jest operacja mnożenia transformat obrazu i filtru.

Filtry dolnoprzepustowe. Działanie tego typu filtrów opiera się na usuwaniu elementów obrazu o wysokiej częstotliwości i przepuszczaniu elementów o niskiej częstotliwości. Ponieważ większość szumów występujących w obrazach zawiera się w wysokich częstotliwościach, filtry te przeważnie wykorzystuje się właśnie do eliminacji zakłóceń.

Filtry górnoprzepustowe. Filtry te działają w sposób odwrotny do filtrów dolnoprzepustowych, tłumią one nisko częstotliwościowe elementy obrazu, wzmacniają natomiast elementy o wysokich częstotliwościach (szczegóły). Wynikiem działania tego typu filtrów jest podkreślenie, uwypuklenie elementów obrazu o dużej częstotliwości poprzez zwiększenie ich jasności, koloru itp. Dla obrazu jako całości efektem jest zazwyczaj zwiększenie kontrastu poprzez podkreślenie ostrych krawędzi obiektów.

OpenCV – biblioteka funkcji wykorzystywanych podczas obróbki obrazu, oparta na otwartym kodzie i zapoczątkowana przez Intela. Biblioteka ta jest wieloplatformowa, można z niej korzystać w Mac OS X, Windows, jak i Linux. Autorzy jej skupiają się na przetwarzaniu obrazu w czasie rzeczywistym.

Pytania prowadzącego:

• jakie mamy rodzaje pamięci w CUDA?

• co to jest i jak działa Kernel?

• dlaczego w obecnych czasach zauważamy wzrost zainteresowania obliczeniami w CUDA?

• rozwiń skrót CUDA

• jakie zasięgi mają pamięci?

• co potrzebujesz żeby programować w CUDA?

• jakie mamy kompilatory?