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?