OpenCL: materiały do laboratorium

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

np. karta graficzna

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).

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).

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

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

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

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).

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 g_x = w_x * L_x + I_x + o_x g_y = w_y * L_y + I_y + o_y

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 L_x * L_y work-item)

I

to lokalna (wewnątrzgrupowa) współrzędna work-item

o

to przesunięcie (określane przed uruchomieniem kernela)

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

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).

• bufory
• obrazy

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)

clGetPlatformIDs(1, &cpPlatform, NULL)

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

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ć.

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).

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ń.

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).

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);

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);

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 (2³²), natomiast ograniczona jest maksymalna liczba wątków wewnątrz jednej work-group (np. 1024 na lepszych kartach).

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);

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 );

Niszczymy wszystkie zainicjowane obiekty OpenCL.

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

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.

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)

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.

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.

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

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

Przykład w matrixmul/Matrixmul_basic.cpp.

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.

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.

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

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.

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).

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?

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.

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