OpenCL: materiały do laboratorium

Spis treści

1 Wstęp do OpenCL

OpenCL to zbiór standardów pomagających w programowaniu na heterogeniczne platformy obliczeniowe. Zwykłe biblioteki programowania współbieżnego (OpenMP, wątki w Javie i w pthreads) zakładają, że procesory są identyczne, a dostęp do każdej komórki współdzielonej pamięci zajmuje tyle samo czasu. Model ten w miarę dobrze odpowiadał klasycznym superkomputerom z dzieloną pamięcią z lat '80 i '90, czy współczesnym wielordzeniowym procesorom - ale już nie komputerom które oprócz zwykłego procesora mają wyspecjalizowane jednostki obliczeniowe, takie jak karty graficzne czy procesory Cell. By wydajnie wykorzystać te platformy, programista musi być świadomy pewnych ograniczeń sprzętu, takich jak grupowanie procesorów czy hierarchia pamięci - i ograniczenia te muszą być widoczne na poziomie języka programowania i API.

OpenCL definiuje abstrakcyjny model heterogenicznej platformy obliczeniowej i zależne od niego modele kodu źródłowego programu, wykonania i hierarchii pamięci. Kod źródłowy składa się z programu hosta (wykonywanego na zwykłym procesorze), który, używając API OpenCL, uruchamia na urządzeniach obliczeniowych kernele obliczeniowe (które pisze się w języku OpenCL).

OpenCL standaryzuje to, co producenci kart graficznych udostępniali przez swoje specyficzne narzędzia i języki programowania (np. nVidia i jej CUDA). Dzięki z jednej strony standaryzacji, a z drugiej odpowiednim sterownikom, program napisany w OpenCL powinien być przenośny nie tylko między kartami graficznymi różnych producentów, ale między urządzeniami obliczeniowymi różnych typów (np. karta graficzna i zwykły procesor).

1.1 Zrównoleglanie oparte o dane vs o zadania

Karty graficzne dysponują dużą liczbą (>100) procesorów. By w maksymalnym stopniu przyspieszyć program, jak najwięcej procesorów musi być jak najintensywniej wykorzystywanych. Dlatego zamiast dużych zadań typu:

for (i = 1; i < 100; i++) {
   output[i] = input[i] * input[i];
}

tworzymy wiele małych wątków, które równolegle wykonują iteracje pętli:

int i = get_global_id(0); # identyfikator wątku
if (i < 100) 
  output[i] = input[i] * input[i];

1.2 Model platformy obliczeniowej

Uruchom CLInfo

1.2.1 Host

Abstrakcja komputera do którego podłączone są urządzenia obliczeniowe.

1.2.2 Urządzenie obliczeniowe (Compute Device)

np. karta graficzna

1.2.3 Multiprocesor (Compute Unit, Streaming Multiprocessor, SM)

Zespół procesorów które mogą wykonywać obliczenia synchronizując się pomiędzy sobą. Procesory z jednej jednostki obliczeniowej zwykle współdzielą część pamięci; mają możliwość synchronizacji obliczeń; współdzielą cache pamięci urządzenia.

NVidia GeForce GTX 470 ma 14 multiprocesorów (CLInfo, max compute units).

1.2.4 Procesor (Processing Element, Streaming Processor, SP)

Pojedynczy procesor. Jednostka obliczeniowa może mieć procesory różnego typu, wyspecjalizowane w różnych obliczeniach.

NVidia GeForce GTX 470 w każdej jednostce obliczeniowej ma 32 Shader Processor (SP) i 4 SFU (Special Function Unit) (źródło: dokumentacja producenta).

1.3 Model programu

1.3.1 Program hosta

Zarządza urządzeniami obliczeniowymi i uruchamia kernele używając API OpenCL.

1.3.2 Kernel

Wykonywany na urządzeniu obliczeniowym. Tworzone w języku OpenCL (albo native).

  1. OpenCL a C99

    Więcej za chwilę. Na początek:

    • brak dostępu do wielu bibliotek (np. printf jest dopiero w OpenCL 1.2)
    • brak wskaźników do funkcji
    • dodane typy wektorowe i operacje na wektorach
    • dodane funkcje atomowe na int i float

1.4 Model wykonania, czyli abstrakcja wykonywanego programu

1.4.1 Wątek (work-item, thread)

Pojedyncza instancja wykonującego się kernela - wątek.

1.4.2 Grupa wątków (work-group, block)

Grupa wykonuje się na jednej jednostce obliczeniowej. Wątki wewnątrz grupy mogą wykonywać się równolegle i synchronizować się między sobą (ale nie ma synchronizacji między grupami). Ilość wątków z grupy faktycznie wykonujących się równolegle zależy od sprzętu: typowo jest to tzw warp, czyli 32 wątki (CLInfo, .

Na GPU wątki te wykonują na raz tę samą instrukcję — więc jeśli wątki jednego warpu chcą wykonać różne instrukcje (np. warunek w instrukcji warunkowej dla części wątków jest prawdziwy, a dla części - fałszywy), grupy wykonają się sekwencyjnie.

Karta graficzna przełącza grupy wątków (a dokładnie: pół-warpy).

Ilość wątków wewnątrz grupy jest ograniczona (GTX470: 1024). Również ograniczona jest liczba wątków wykonywanych na multiprocesorze (GTX470: 1024).

Ilość grup które mogą wykonywać się równolegle na multiprocesorze jest ograniczona (np. GTX470: 8).

1.4.3 Indeksowanie

OpenCL Runtime Environment definiuje jedno-, dwu-, lub trójwymiarową przestrzeń indeksów (NDRange) pozwalających odróżniać poszczególne work-item i work-group. Jest to wielowymiarowy i (do pewnego stopnia) konfigurowalny system identyfikacji wątków (Thread ID). Wielowymiarowe indeksy pozwalają numerować wątki w zależności od dziedziny problemu (np. przy działaniach na macierzach - wiersz i kolumna).

Załóżmy indeksowanie dwuwymiarowe, wtedy gx = wx * Lx + Ix + ox gy = wy * Ly + Iy + oy

gdzie:

g
to współrzędna globalna work-item
w
to współrzędna work-group do której należy work-item
L
to wielkość work-group (do jednej work-group należy Lx * Ly work-item)
I
to lokalna (wewnątrzgrupowa) współrzędna work-item
o
to przesunięcie (określane przed uruchomieniem kernela)

1.4.4 Kontekst

urządzenia
kernele
obiekty programu
źródła i pliki wykonywalne kerneli; pliki wykonywalne budowane w czasie wykonywania programu hosta, dostosowane do konkretnego urządzenia
obiekty pamięci
fragmenty pamięci widoczne urządzeniom OpenCL

1.4.5 Kolejki wykonywania

Kernele wykonywane są asychronicznie poprzez przesyłanie do kolejki związanej z danym urządzeniem.

Polecenia:

  • wykonanie kernela
  • operacje na pamięci (np. transfer danych między hostem a urządzeniem)
  • synchronizacja

Polecenia w kolejce mogą być wykonywane w kolejności wysłania (in-order) albo tak szybko, jak możliwe (out-of-order).

1.5 Model pamięci

1.5.1 Typy obiektów

  • bufory
  • obrazy

1.5.2 Abstrakcje rodzajów pamięci

Transfery do i z pamięci są zwykle wąskim gardłem obliczeń, dlatego kluczowe jest odpowiednie użycie odpowiednich typów pamięci.

pamięć hosta
niedostępna z poziomu kernela;
pamięć globalna (urządzenia, device)
dostępna r/w dla wszystkich wątków wszystkich grup; zwykle cache'owana; duża latencja (ok. 200 cykli - a wykonanie typowej instrukcji arytmetycznej zajmuje 4 cykle); sporej wielkości (GTX 470 – 1280MB)
pamięć stałych (constant)
dostępna do czytania dla wszystkich wątków wszystkich grup; zwykle cache'owana; optymalizowana do odczytu tej samej lokalizacji przez wszystkie wątki; niewielka (64 KB)
pamięć lokalna (local, shared)
dostępna r/w dla wszystkich wątków pojedynczej grupy; ``blisko'' jednostki obliczeniowej na której wykonują się wątki; szybka, ale niewielka (np. GTX 470 – 48 KB)
pamięć prywatna
dostępna dla pojedynczego wątku. Szybka ale mała (i wielkość nie jest znana ) - w sprzęcie odpowiada rejestrom procesora. Po przepełnieniu rejestrów dramatycznie zwalnia (bo jest symulowana przez pamięć globalną).

Interakcje host-device przez kopiowanie obszarów pamięci lub przez mapowanie (więcej o tym później)

1.5.3 Spójność pamięci

  1. brak spójności pomiędzy grupami

    Relaxed consistency: kolejność load/store może być różna dla różnych wątków.

  2. lokalna spójność wewnątrz grupy

    Pamięć lokalna (i globalna) jest spójna po instrukcji synchronizacji grupy (np. barrier dla grupy).

2 Pierwszy program w OpenCL

Źródło programu w first.cpp.

2.1 Informacje techniczne

OpenCL Toolkit jest dystrybuowane przez AMD, Intela i nVidię. Do uruchomienia przykładów potrzebujesz bibliotek (libOpenCL.so) i kompatybilnych sterowników urządzeń (karty graficznej lub procesora). Wersja sterowników zainstalowana w systemie: nvidia-smi. Linki:

nvidia1.mimuw.edu.pl ma zainstalowane dwa drivery do OpenCL: nVidia do uruchamiania OpenCL na GPU i AMD do uruchamiania OpenCL na CPU (sterowniki Intela wymagają nowszego procesora).

Do kompilacji i uruchomienia programu potrzebujesz biblioteki libOpenCL.so i plików nagłówkowych <CL/opencl.h>.

Kompilacja na nvidia1.mimuw.edu.pl:

g++ -lOpenCL first.cpp

Kompilacja w laboratorium 3043 / 3045:

g++ -L /usr/lib64/nvidia -I /opt/cuda/include/ -lOpenCL  first.cpp

2.2 Kod hosta

2.2.1 dostęp do platformy:

clGetPlatformIDs(1, &cpPlatform, NULL)

Platforma reprezentuje model systemu OpenCL i umożliwia dostęp do poszczególnych urządzeń.

2.2.2 dostęp do urządzenia obliczeniowego:

clGetDeviceIDs(cpPlatform, devType, 1, &device_id, NULL);

2.2.3 stwórz kontekst obliczeniowy:

context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);

2.2.4 stwórz kolejkę zadań:

commands = clCreateCommandQueue(context, device_id, 0, &err);

Ponieważ OpenCL jest pomyślany jako środowisko asynchroniczne, zadania (takie jak kernele czy polecenia transferu pamięci) będziemy wysyłać do urządzenia za pośrednictwem kolejki. Kolejka jest interfejsem do programu szeregującego (schedulera) działającego na urządzeniu i decydującego o tym które zadania zacząć wykonywać.

2.2.5 stwórz obiekt programu (używając napisu KernelSource):

program = clCreateProgramWithSource(context, 1, (const char **) &KernelSource, NULL, &err);

Kernel inicjowany jest z napisu (char*). Dla uproszczenia kodu, w first.cpp kernel zdefiniowany jest jako stała typu char *. W innych programach kernel będzie czytany z niezależnego pliku. Zauważ, że program tworzony jest dla konkretnego kontekstu (parametr context).

2.2.6 zbuduj program:

clBuildProgram(program, 0, NULL, NULL, NULL, NULL);

Program budowany jest dla określonego urządzenia (skojarzonego z kontekstem z którym tworzyliśmy program w poprzednim kroku). Jeśli wartością urządzenia jest NULL, program budowany jest dla wszystkich urządzeń.

2.2.7 stwórz kernel z programu:

kernel = clCreateKernel(program, "square", &err);

Program może zawierać wiele kerneli - ta funkcja tworzy obiekt kernela z wybranej funkcji (tu square) skompilowanego programu. Obiekt kernela będziemy mogli wysłać do wykonania (oczywiście po ustaleniu parametrów).

2.2.8 zmienne: podział na dane w pamięci hosta i w pamięci urządzenia

float* data = new float[DATA_SIZE];    // host: parametry
float* results = new float[DATA_SIZE]; // host: rezultaty

cl_mem input = clCreateBuffer(context,  CL_MEM_READ_ONLY,  
sizeof(float) * count, NULL, NULL); // bufor wejściowy w pamięci urządzenia
cl_mem output = clCreateBuffer(context, CL_MEM_WRITE_ONLY, 
sizeof(float) * count, NULL, NULL); // bufor wyjściowy w pamięci urządzenia

2.2.9 polecenie transferu danych z pamięci hosta do pamięci urządzenia

Kernele OpenCL nie mają bezpośredniego dostępu do pamięci hosta, więc wszystkie dane (tu float* data) trzeba przepisać z pamięci hosta (float* data) do buforu w pamięci urządzenia (cl_mem input).

clEnqueueWriteBuffer(commands, input, 
                         CL_TRUE, 0, sizeof(float) * count, 
                         data, 0, NULL, NULL);

2.2.10 argumenty dla kernela

Ponieważ kernel będziemy wykonywać poprzez wysłanie do kolejki, trzeba ustawić jego parametry za pomocą specjalnych funkcji (a nie - po prostu wołając funkcję z odpowiednimi argumentami).

clSetKernelArg(kernel, 0, sizeof(cl_mem), &input);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &output);
clSetKernelArg(kernel, 2, sizeof(unsigned int), &count);

2.2.11 polecenie uruchomienia kernela

OpenCL działa asynchronicznie, więc uruchomienie kernela to tak na prawdę wysłanie gotowego do wykonania kernela do kolejki commands związanej z urządzeniem obliczeniowym.

err = clEnqueueNDRangeKernel(commands, kernel, 
                           1, NULL, &global, &local, 
                           0, NULL, NULL);

Liczbę tworzonych work-items określają parametry global (typ: const size_t *global_work_size) oraz local (typ: const size_t *local_work_size). Oba parametry są tablicami indeksowanymi wymiarem. global to całkowita liczba work-item które zostaną stworzone. Stworzone work-item zostaną podzielone na grupy zawierające local work-item. Uwaga: dla każdego wymiaru d, global[d] musi być podzielne przez local[d].

Zwykle liczba wątków które możemy stworzyć jest bardzo duża (232), natomiast ograniczona jest maksymalna liczba wątków wewnątrz jednej work-group (np. 1024 na lepszych kartach).

2.2.12 bariera: czekaj na koniec obliczeń

Bariera to kolejny element związany z asynchronicznością OpenCL. Kod hosta musi być powiadomiony kiedy skończy się wykonanie kernela. Czekamy na zakończenie wszystkich poleceń z kolejki commands - a jedynym poleceniem było tam wykonanie kernela.

clFinish(commands);

2.2.13 polecenie transferu wyników do pamięci hosta

Ponieważ host również nie ma bezpośredniego dostępu do pamięci urządzenia, wynik (bufor output) musi być jawnie przepisany do pamięci hosta.

err = clEnqueueReadBuffer( commands, output,
                         CL_TRUE, 0, sizeof(float) * count,
                         results, 0, NULL, NULL );

2.2.14 końcowe porządki

Niszczymy wszystkie zainicjowane obiekty OpenCL.

clReleaseMemObject(input);
clReleaseMemObject(output);
clReleaseProgram(program);
clReleaseKernel(kernel);
clReleaseCommandQueue(commands);
clReleaseContext(context);

2.3 Kod kernela

Kod kernela jest bardzo prosty: i-ty work-item oblicza kwadrat i-tego elementu tablicy input i zapisuje wynik do tablicy output.

__kernel void square( __global float* input, __global float* output,
                      const unsigned int count) { 
  int i = get_global_id(0); 
  if (i < count) 
    output[i] = input[i] * input[i]; 
}

2.4 Zadania

  • uruchom kernel na GPU i na CPU (uruchamiając program hosta z jednym parametrem)
  • porównaj czas działania (polecenie time)
  • co może być powodem wolniejszego działania programu na GPU niż na CPU?
  • sprawdź przepustowość pamięci uruchamiając /usr/local/cuda/samples/1_Utilities/bandwdithTest/bandwidthTest (na nvidia1/2 w: /usr/local/cuda/samples/bin/x86_64/linux/release)
  • dodaj instrukcję printf("thread %d", i") do kodu kernela. Uruchom program na GPU i na CPU. (printf działa w OpenCL 1.2; sterownik nvidia jest dla wersji 1.1).

2.5 Kod template

Będziemy budować aplikacje w oparciu o kod Template ( template/ )

  • skompiluj kod
  • uruchom go; jeśli coś nie działa, zmień kontekst z CL_DEVICE_TYPE_CPU na CL_DEVICE_TYPE_GPU
  • porównaj ten kod z pierwszym programem

3 Dostosowanie programu do platformy obliczeniowej

OpenCL API definiuje funkcje którymi można uzyskać informacje o parametrach platformy i poszczególnych urządzeń. Dzięki temu można np. dostosować wielkość tablic do wielkości pamięci dostępnej na urządzeniu.

Przykłady użycia znajdują się w info.cpp.

Ważniejsze funkcje:

3.1 Informacje o platformie:

cl_int clGetPlatformInfo (cl_platform_id platform,
       cl_platform_info param_name,
       size_t param_value_size,
       void * param_value,
       size_t * param_value_size_ret)

Wartości dla param_name znajdziesz w http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetPlatformInfo.html

3.2 Informacje o urządzeniu:

cl_int clGetDeviceInfo (cl_device_id device,
       cl_device_info param_name,
       size_t param_value_size,
       void * param_value,
       size_t * param_value_size_ret)

Wartości dla param_name znajdziesz w http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/clGetDeviceInfo.html

Najważniejsze:

CL_DEVICE_LOCAL_MEM_SIZE
wielkość pamięci lokalnej
CL_DEVICE_GLOBAL_MEM_SIZE
wielkość pamięci globalnej
CL_DEVICE_MAX_COMPUTE_UNITS
liczba jednostek obliczeniowych
CL_DEVICE_MAX_WORK_GROUP_SIZE
maksymalna liczba work-items wewnątrz work-group
CL_DEVICE_MAX_WORK_ITEM_SIZES
maksymalna liczba work-items

Dodaj do info.cpp wyświetlenie liczby work-items na urządzeniu.

3.3 Informacja o wielkości (liczbie) zwróconych elementów:

standardowo przez ustawienie parametrów na 0 / NULL np. clGetPlatformIDs(0, NULL, &num_platforms) zwraca w num_platforms liczbę platform.

4 Programowanie kerneli

Kernele OpenCL pisane są w języku OpenCL, opartym o C99.

4.1 Nie ma wielu bibliotek

printf nie działa

4.2 Kwalifikatory

  • oznaczenie kodu funkcji kernela przez kernel (lub __kernel)
  • oznaczenie pamięci w której przechowywane są zmienne (bądź parametry)
    • global
    • constant
    • local
    • private

4.3 Wektorowe typy danych

Przykład użycia w vectors/vectorsq3.cl.

Uwaga: karty graficzne na razie nie optymalizują obliczeń na wektorach (CLInfo, Preferred vector width = 1).

Wektorowe typy danych rozszerzają zwykłe typy ( char, int, long ), np. intn, gdzie n to 2, 3, 4, 8, 16: int8 to wektor 8 liczb całkowitych.

Typom wektorowym odpowiadają typy w OpenCL API (np. cl_int8), które można wykorzystywać w programie hosta.

Inicjacja wartości przez równoczesne przypisanie: uint4 v = (uint4) (3, 5, 7, 9)

Można uzyskać dostęp do poszczególnych elementów wektora przez kropkę i numer elementu:

  • v.y = 9 przypisanie do drugiego elementu wektora (możliwe indeksy: x y z w)
  • v.s0 = 2 przypisanie do pierwszego elementu wektora - dzięki indeksowaniu można zaadresować do 16 wartości wektora.

Na wektorach można wykonywać działania matematyczne i wywoływać funkcje, np:

uint4 v1 = (uint4) (3, 5, 7, 9);
uint4 v2 = (uint4) (1, 2, 4, 8);

uint4 v3 = v1 + v2; // (4, 7, 11, 17)
uint4 v4 = 2 * v1; // (6, 10, 14, 18)

4.4 Funkcje obliczeniowe

OpenCL ma zdefiniowanych wiele wysoko-wydajnych funkcji operujących także na wektorowych typach danych. Kompletną listę znajdziesz w http://www.khronos.org/files/opencl-quick-reference-card.pdf

4.4.1 funkcje matematyczne ( sin, pow, exp, … )

4.4.2 funkcje geometryczne ( dot, distance, length, normalize)

4.4.3 funkcje porównań ( isequal, isnan, isinf, isgreater, …)

4.5 Funkcje identyfikujące work-item

Uruchom dims/dims.cpp z lid-kernel.cl ; w dims.cpp modyfikuj wartość localThreads[0] określającą ile work-item wykonuje się w ramach jednej grupy.

uint get_work_dim()
liczba wymiarów użytych w wywołaniu kernela;
site_t get_global_size(uint D)
ogólna liczba work-item w wymiarze D;
size_t get_global_id(uint D)
współrzędna globalna work-item;
size_t get_local_size(uint D)
liczba work-item w danej work-group;
size_t get_local_id(uint D)
współrzędna lokalna work-item (w ramach jednej work-group);
size_t get_num_groups(uint D)
liczba work-group;
size_t get_group_id*uint D)
współrzędna aktualnej work-group;

4.6 Funkcje synchronizacyjne

Synchronizacja możliwa jest tylko między work-item wewnątrz jednej work-group. Nie działa synchronizacja pomiędzy work-group.

void barrier (cl_mem_fence_flags flags)

flags: CLK_LOCAL_MEM_FENCE CLK_GLOBAL_MEM_FENCE

Bariera wstrzymuje wykonanie work-item aż wszystkie work-item wykonają rozkaz bariery. Równocześnie po rozkazie bariery gwarantowana jest spójność pamięci - wszystkie przypisania zainicjowane przed barierą zostaną zakończone.

4.7 Kopiowanie i pre-fetch danych

event_t async_work_group_copy(__local T* dst, const __global T* src, size_t num_elements, event_t event)

zwraca event, na którym poniższa funkcja blokuje się do zakończenia kopiowania:

void wait_group_events(int num_events, event_t* event_list)

4.8 Funkcje atomowe

4.8.1 sposób działania

T atomic_add(global T* p, T* val)
atomic_cmpxchg(global T* p, T cmp, T val)
jeśli (*p == cmp), to *p = val

Działają niestety tylko dla liczb całkowitych.

4.8.2 typy funkcji: add, sub, xchg, inc, dec, min, max, and, or, xor

4.9 Funkcje operujące na obrazach

5 Obiekty pamięci

memobjects/memobj.cpp

5.1 użycie w kernelu pamięci lokalnej

Ilość alokowanej pamięci musi być znana w momencie kompilacji kernela; dlatego nie można napisać:

uint temp_size = get_local_size(0);
__local float temp[temp_size];

Wskaźnik do pamięci lokalnej najwygodniej jest przekazać jako parametr kernela (typu __local float* temp). Parametr ten w kodzie hosta inicjuje się przez:

clSetKernelArg(kernel, INDEKS_PARAMETRU, sizeof(cl_float) * localThreads[0], NULL);

5.2 Bufory i subbufory

Bufor to jednowymiarowy obszar pamięci dostępnej dla kerneli OpenCL.

5.2.1 tworzenie bufora

cl_mem clCreateBuffer(  
       cl_context context, // bufor zawsze dla jakiegoś kontekstu
       cl_mem_flags flags,
       size_t size, // wielkość bufora
       void *host_ptr,  // wskaźnik na pamięć hosta
       cl_int *errcode_ret) // miejsce na błędy

Ważniejsze flagi:

  • CL_MEM_READ_WRITE, CL_MEM_WRITE_ONLY, CL_MEM_READ_ONLY
  • CL_MEM_COPY_HOST_PTR skopiuj zawartość pamięci hosta spod host_ptr
  • CL_MEM_USE_HOST_PTR używaj pamięci hosta (czyli bufor może być adresowany z kodu hosta przez host_ptr)

5.2.2 asynchroniczne operacje kopiowania na buforach

Transfery między pamięcią hosta i urządzenia mogą być kolejkowane w podobny sposób co polecenia wykonania kerneli.

cl_int clEnqueueWriteBuffer (   cl_command_queue command_queue,
    cl_mem buffer,
    cl_bool blocking_write, // jeśli CL_TRUE, to operacja blokująca
    size_t offset,
    size_t cb, // ile bajtów skopiować
    const void *ptr, // skąd kopiować 
    cl_uint num_events_in_wait_list,
    const cl_event *event_wait_list, // eventy na które czekam przed rozpoczęciem operacji
    cl_event *event // event generowany po zakończeniu operacji
    )

Analogicznie: clEnqueueReadBuffer

Zastąp wywołanie clCreateBuffer z kopiowaniem zawartości pamięci (flaga CL_MEM_COPY_HOST_PTR) przez clCreateBuffer a następnie clEnqueueWriteBuffer.

5.2.3 mapowanie buforów

Pozwala mapować bufor OpenCL w przestrzeni adresowej hosta. Dzięki temu program hosta może manipulować taką zamapowaną pamięcią nie wiedząc o tym, że jest ona zarządzana przez OpenCL.

void * clEnqueueMapBuffer (     cl_command_queue command_queue,
    cl_mem buffer,
    cl_bool blocking_map,
    cl_map_flags map_flags, // CL_MAP_READ albo CL_MAP_WRITE
    size_t offset,
    size_t cb,
    cl_uint num_events_in_wait_list,
    const cl_event *event_wait_list,
    cl_event *event,
    cl_int *errcode_ret)

Zwolnienie przez clEnqueueUnmapMemObject.

Zastąp wywołanie clEnqueueReadBuffer przez clEnqueueMapBuffer.

6 Tuning programów openCL

6.1 Profilling źródeł

Chcesz przyspieszyć obszerny program który nie korzysta z karty graficznej?

Zacznij od znalezienia celu - funkcji której wykonanie zajmuje większość czasu. W tym celu zrób profilling oryginalnego programu.

Przyspieszenie programu jako całości będzie ograniczone przez prawo Amdahla (czyli przez obliczenia wykonywane przez fragmenty niezmodyfikowane) oraz transfery pamięci między hostem i GPU. Kod wykonywany na GPU po pewnych optymalizacjach jest tak szybki, że czas wykonywania można pominąć.

Załóżmy na przykład, że przepisałeś na GPU fukcję która zajmowała 90% czasu oryginalnych obliczeń. Teoretyczną granicą przyspieszenia jest 10 razy; ale trzeba jeszcze brać pod uwagę transfery do i z pamięci. 10-krotne przyspieszenie to nie jest rewelacyjny wynik przy obliczeniach na GPU; zwykle celem jest przyspieszenie 50-80 krotne - ale pamiętaj, że dla większości programów każde przyspieszenie ma znaczenie.

Problem pojawia się, gdy wykonanie żadnej z funkcji nie zabiera większości czasu obliczeń. By uniknąć kosztów transferu pamięci, trzeba przepisać na GPU całą część obliczeniową aplikacji - albo, jeśli funkcje są w miarę niezależne, zrównoleglać obliczenia jednej funkcji z transferami pamięci dla innej funkcji.

6.2 Ocena czasu wykonania

6.2.1 całego programu - time

Bierze pod uwagę część niezrównolegloną, transfery pamięci i kernel.

6.2.2 kernelu

(albo dowolnej innej operacji OpenCL wysłanej do kolejki zadań)

Przykład w matrixmul/Matrixmul_basic.cpp.

  1. Przy tworzeniu kolejki zadań dodaj do properties flagę CL_QUEUE_PROFILING_ENABLE.
    commandQueue = clCreateCommandQueue(
                       context, 
                       devices[0], 
                       CL_QUEUE_PROFILING_ENABLE, // enables to profile the code
                       &status);
    
  2. Przy wysyłaniu kernela do kolejki ustaw niepuste zdarzenie (tu events[0]):
    status = clEnqueueNDRangeKernel(
                 commandQueue,
                 kernel,
                 2,
                 NULL,
                 globalThreads,
                 localThreads,
                 0,
                 NULL,
                 &events[0]); // not-null event for tracing execution
    
  3. Po zakończeniu kernela posłuż się funkcją clGetEventProfilingInfo do odczytania czasu startu i zakończenia obliczeń.
    clGetEventProfilingInfo (   cl_event event,
        cl_profiling_info param_name, // CL_PROFILING_COMMAND_START lub CL_PROFILING_COMMAND_END
        size_t param_value_size, // sizeof(cl_ulong)
        void *param_value, // parametrem powinno być cl_ulong
        size_t *param_value_size_ret) // może być NULL
    

    Czas zwracany w param_value mierzony jest w nanosekundach. Podaje on wartość wewnętrznego zegara urządzenia w momencie zajścia zdarzenia event - param_name. Np. przy param_name = CL_PROFILING_COMMAND_START, funkcja zwróci w param_value wartość wewnętrznego zegara urządzenia z chwili, kiedy wysłany do kolejki kernel z argumentem event zaczął się wykonywać.

7 Zadania

7.1 Jak to rozwiązać

Prosimy o rozwiązanie zadania w kilku etapach.

7.1.1 Przygotuj kod hosta i podstawową (naiwną) wersję kodu kernela.

W kodzie kernela korzystaj z wielowątkowości (czyli ogranicz użycie pętli), ale nie staraj się używać instrukcji wektorowych czy optymalizować dostępy do pamięci.

7.1.2 Zmierz czas wykonania naiwnego kernala.

Wykorzystaj zarówno polecenie time jak i profilowanie openCL.

Zmierz czas wykonania dla różnych wielkości parametrów wejściowych.

Porównaj czas wykonania kernela na GPU i na CPU.

7.1.3 Optymalizuj kernel

Pamiętaj o:

  • optymalizacji wielkości work-group
  • instrukcjach wektorowych
  • szybkich wersjach funkcji matematycznych
  • użyciu pamięci lokalnej, a nie globalnej (co zwykle wymaga gruntownej zmiany algorytmu)

Po każdej optymalizacji mierz czas wykonania i zapisuj wyniki.

Pamiętaj, że w programowaniu na GPU wyzwaniem jest danie karcie wystarczająco dużo obliczeń.

Gdzie szukać pomysłów co optymalizować:

  • oclReduction z nVidia SDK

7.2 przykład optymalizacji: mnożenie macierzy

Pliki w katalogu matrixmul.

7.2.1 algorytm naiwny

Zacznij od programu hosta Matrixmul_basic.cpp i kernela matrixmul_kernel.cl

  1. Zanotuj czas działania programu dla różnych wielkości macierzy (wartości rowA, colA, itp w initializeHost).
  2. Spróbuj uruchomić program na CPU i porównać czas działania.
  3. Uruchom program na macierzy z liczbami zmiennoprzecinkowymi; czy pojawiły się błędy? Z czego one wynikają?
  4. Jakie widzisz możliwości przyspieszenia kernela?
  5. Spróbuj uruchomić program dla dużych macierzy. Dlaczego zgłaszane są błędy? Zmodyfikuj program tak, by można było go uruchomić dla bardzo dużych macierzy.

7.2.2 optymalizacja wielkości grup wątków

Algorytm naiwny korzysta z grup zawierających tylko 1 wątek (zmienna localThreads). Optymalna wielkość grup zależy od:

  • maksymalnej liczby grup w multiprocesorze i na karcie
  • maksymalnej ilości wątków w multiprocesorze i na karcie
  • zużyciu rejestrów przez wątek

Zadania:

  1. Jaki jest górny limit wielkości grupy?
  2. Stopniowo zwiększaj grupy i notuj czas działania.
  3. Sprawdź czas działania dla wielkości grup niepodzielnych przez wielkość warp-u (32 wątki).

7.2.3 drobne optymalizacje

Zmień kernel tak, by mnożenia kolejnych elementów modyfikowały zmienną prywatną wątku; dopiero po zakończeniu wszystkich mnożeń wartość zapisywana jest do macierzy wynikowej.

Czy program działa szybciej? Dlaczego?

7.2.4 algorytm 3d

Zmień program hosta na Matrixmul_3d.cpp; program ten korzysta z kernela matrixmul_kernel3d.cl.

  1. Co autor chciał zoptymalizować?
  2. Czy program daje poprawne rezultaty?
  3. Zmień program tak, by był poprawny.
  4. Porównaj czas działania z poprzednim podejściem dla macierzy różnych rozmiarów.

7.2.5 optymalizacja dostępu do pamięci przez kafelkowanie

Do obliczenia sąsiednich elementów macierzy C używane są pokrywające się elementy macierz A i B, np. zarówno C[1,1] jak i C[1,2] wymagają m.in. A[1,1]. Zamiast dwukrotnie sprowadzać A[1,1] z pamięci globalnej, można A[1,1] "na chwilę" sprowadzić do pamięci local - a następnie skorzystać z niego do obliczenia C[1,1], C[1,2] - i następnych elementów pierwszego wiersza i pierwszej kolumny macierzy C.

Algorytm:

  1. Podziel macierze A i B na kafelki wielkości group_size[0] x group_size[1].
  2. Wątki współdzielą kafelki w pamięci dzielonej.
  3. Dla każdego kolejnego kafelka:
    1. Każdy z wątków grupy ładuje 1 element A i 1 element B do kafelków
    2. Każdy z wątków grupy oblicza iloczyn wiersza i kolumny kafelka

7.3 redukcja wektora

Zaimplementuj algorytm redukcji wektora w czasie O(logn) ( opisywany m.in. w artykule ``Data-parallel algorithms'',WD Hills, GL Steele, http://dx.doi.org/10.1145/7902.7903 ).

8 Źródła i bibliografia

http://www.khronos.org/developers/library/overview/opencl_overview.pdf
wstęp do openCL
http://www.codeproject.com/KB/showcase/Portable-Parallelism.aspx
niezły wstęp
http://www.codeproject.com/KB/showcase/Memory-Spaces.aspx
jak korzystać z pamięci
http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/
OpenCL Reference
http://www.khronos.org/files/opencl-quick-reference-card.pdf
Reference Card
https://developer.nvidia.com/nsight-eclipse-edition
debugger nvidia
http://developer.amd.com/tools-and-sdks/heterogeneous-computing/codexl/
debugger do kart AMD
http://documen.tician.de/pyopencl/
language bindings dla pythona
OpenCL Programming Guide
Aaftab Munshi; Benedict Gaster; Timothy G. Mattson; James Fung, Addison-Wesley 2011
Programming Massively Prallel Processors
David B. Kirk; Wen-mei W. Hwu; Morgran Kaufmann 2010

Data: $Date: 2014-03-20 18:19:31 +0100 (Thu, 20 Mar 2014) $

Autor: Krzysztof Rządca

Created: 2014-03-26 Wed 17:34

Emacs 24.3.50.2 (Org mode 8.2.5h)

Validate