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?