Wstęp do OpenCL

Przez większość użytkowników komputerów karta graficzna kojarzona jest jedynie jako element, który służy do wyświetlania grafiki. Jednak każdy kto lepiej orientuje się w nowoczesnych technologiach wie, że obecnie nie jest to jedyne zastosowanie tych układów (w dodatku bardzo uproszczone, bo jak wiadomo pod pojęciem „wyświetlanie grafiki” kryją się wszystkie operacje związane z przetworzeniem danych na piksele ekranu). Chociaż z GPGPU można było już korzystać na przełomie XX i XXI wieku (rok 1999) do jednak dopiero kiedy Nvidia stworzyła w 2006 roku technologię CUDA procesory graficzne coraz częściej zaczęły być wykorzystywane jako bardzo wydajne jednostki obliczeniowe ogólnego przeznaczenia. Zaś w roku 2009 powstała wersja 1.0 standardu OpenCL, który pozwolił na użycie tego samego kodu do obliczeń zarówno na kartach firmy Nvidia, jak i AMD (dawniej ATI), w przeciwieństwie do CUDA, które ograniczone było do kart tylko pierwszego z producentów. I właśnie o pierwszych krokach z OpenCL będzie ten wpis.

1. Wstęp

Na początku warto zaznaczyć, że chociaż na GPU mogą być teoretycznie wykonywane dowolne obliczenia to jednak dopiero przy algorytmach dających możliwość obliczeń równoległych ta technologia pokazuje swoją przewagę. Wynika to z budowy samych układów, które posiadają bardzo dużą liczbę rdzeni, jednak każdy z nich jest w pewien sposób mniej wydajny niż zwykły rdzeń CPU (po więcej informacji odsyłam np. do Wikipedii).

Dosyć istotną informacją jest to, że kernele OpenCL mogą być wykonywane zarówna na karcie graficznej jak i na procesorze głównym, co ma tą zaletę, że pozwala stosować w większości ten sam kod również na maszynach nie mających układów bezpośrednio wspierających tą technologię.

W poniższym tekście nie będę opisywał wszystkich kroków związanych np. z dodawaniem ścieżek do plików w IDE, a jedynie krótkie informacje co należy dodać. Zakładam, że zainteresowani potrafią posługiwać się używanym IDE lub ostatecznie potrafią korzystać z wyszukiwarki Google. Zakładam również, że czytelnik zna język C++.

Poniższy tekst oparłem głównie o tutorial wprowadzający ze strony AMD, jednak znalazłem w nim kilka nieścisłości i braków, z którymi musiałem się uporać dlatego myślę, że również tutaj każdy znajdzie coś przydatnego.

2. Czego potrzebujemy

Na początku musimy zaopatrzyć się w SDK producenta naszej karty, które zawierać będzie pliki potrzebne do budowania programów z wykorzystaniem OpenCL. W moim przypadku jest to SDK dostarczone przez AMD. Poniżej podaję linki dla AMD i Nvidii:

AMD: http://developer.amd.com/tools-and-sdks/heterogeneous-computing/amd-accelerated-parallel-processing-app-sdk/downloads/

Nvidia: https://developer.nvidia.com/cuda-downloads

Jak widać w przypadku Nvidii OpenCL obsługiwany jest przez pakiet CUDA. Dodatkowo taka uwaga, że, przynajmniej w przypadku AMD (nie wiem jak jest z SDK Nvidii), nawet jeśli chcemy budować program dla platformy x86, to mając system 64 bitowy ściągamy SDK przygotowane właśnie pod 64 bity ponieważ inaczej instalator nie uruchomi się.

Mając już SDK wypadałoby się teraz zaopatrzyć w jakieś IDE, no chyba, że ktoś bardzo lubi kompilację z poziomu konsoli i pisanie kodu w osobnym edytorze bez np. podpowiadania składni. Ponieważ pracuję pod systemem Windows 7 jako IDE wybrałem Visual C++ 2010 Express, z którego korzystam już od dłuższego czasu.

Kiedy posiadamy już odpowiednie narzędzia czas zacząć działać.

3. Pierwszy program

Pierwszy program jaki napiszemy korzystając z OpenCL nie będzie jeszcze wykonywał żadnych obliczeń. Na początku wyświetlimy jedynie dostępne urządzenia, które będą mogły służyć jako jednostki obliczeniowe.

Tworzymy więc nowy projekt konsolowy w Visual Studiu. Dodajemy w ustawieniach projektu ścieżkę do plików nagłówkowych biblioteki OpenCL (w przypadku AMD jest to folder include znajdujący się w folderze z SDK). Dodajemy również w opcjach linekra ścieżkę do folderu z plikami lib (jeśli kompilujemy pod platformę 32 bitową to wybieramy folder z bibliotekami x86) oraz dodajemy do linkera plik OpenCL.lib. Teraz mając skonfigurowany projekt możemy zająć się kodem.

Na samym początku definiujemy, że nie będziemy używali wektora ze standardowej biblioteki C++, zamiast niego OpenCL będzie korzystał z własnego wektora.

Teraz musimy dodać potrzebne includy

Użyłem tutaj pliku cl.hpp zamiast cl.h ponieważ uznałem, że od razu będę korzystał z wersji obiektowej przygotowanej dla C++. Czas przejść do funkcji main. Tworzymy zmienną err typu cl_int (odpowiednik zwykłego inta zawarty w bibliotece OpenCL) oraz wektor przechowujący typ cl::Platform.

Za pomocą funkcji get() z klasy cl::Platform pobieramy dostępne platformy i wyświetlamy ich ilość

Kiedy już pobraliśmy dostępne platformy czas wyświetlić jakieś informacje na ich temat. Będziemy to robić w pętli for iterującej od do platformList.size() -1. Na początek informacja o producencie platformy. Pobieramy ją za pomocą funkcji getInfo() dostępnej w klasie cl::Platform

Oprócz CL_PLATFORM_VENDOR mogliśmy też użyć np. CL_PLATFORM_VERSION lub CL_PLATFORM_NAME, które zwróciły by odpowiednio wersję i nazwę platformy. Warto w tym miejscu powiedzieć, że platformą jest po prostu zainstalowany sterownik OpenCL.

Następnie tworzymy właściwości kontekstu, w których podajemy dla jakiej platformy będziemy go tworzyli

 I samo utworzenie kontekstu

Parametr CL_DEVICE_TYPE_ALL określa jakiego typu urządzenia chcemy obsługiwać, w tym wypadku wybrane będzie zarówno GPU jak i CPU. Jednak w kolejnych przykładach będę korzystał jedynie z GPU, a więc w tym miejscu znajdzie się wartość CL_DEVICE_TYPE_GPU.

Mając już utworzony kontekst możemy pobrać informacje o wszystkich urządzeniach jakie obsługuje. A służy do tego szablonowa funkcja getInfo<>() klasy cl::Context, która w tym przypadku jako nazwę typu przyjmuje CL_CONTEXT_DEVICES

Teraz zostało już tylko wypisanie informacji o dostępnych urządzeniach. Ja w przykładzie pobieram tylko nazwę każdego z nich jednak możliwe jest również pobranie np. wersji lub producenta, tak samo jak w przypadku platformy

Na koniec tylko standardowe zatrzymanie programu (jeśli programujesz pod linuxem nie masz problemu ze znikającą konsolą po wykonaniu się programu dlatego możesz pominąć oczekiwanie na wciśnięcie klawisza) i zakończenie funkcji main()

 Cały kod tego przykładu prezentuje się następująco:

 A wynik jego wykonania na moim komputerze wygląda tak:

p1

 

Jak można zauważyć, dostępną mam jedną platformę (AMD APP SDK), która posiada 2 urządzenia. Tajemniczy Pitcairn to po prostu nazwa rdzenia użytego w moim Radeonie HD7870 (nie wiem czemu wyświetla tą nazwę zamiast nazwy karty, ale to nie istotne), drugie urządzenie to procesor Intela i5-2550K.

4. Pora coś policzyć

Poprzedni program był tylko wstępem, który pozwolił sprawdzić czy wszystko jest dobrze skonfigurowane. Teraz przejdę do bardziej praktycznej części. W końcu wykonamy jakieś obliczenia na GPU. W tym wypadku będzie to proste mnożenie elementów dwóch tablic i przypisywanie wyniku do trzeciej. Niby niewiele, ale jest to wystarczający przykład aby poznać ideę obliczeń z wykorzystaniem GPU.

Nie rozpisując się za bardzo przejdźmy do samego programu. Standardowo zaczynamy od dołączenia potrzebnych plików nagłówkowych. Jest ich trochę więcej niż w poprzednim przykładzie.

 Teraz funkcja, którą wziąłem z tutoriala AMD. Pozwala ona wypisać w konsoli wiadomość i numer błędu. Zastępuje if-a, którego by trzeba użyć do sprawdzania błędów, przez co kod jest trochę czytelniejszy. Jednak i tak najlepszym rozwiązaniem byłoby użycie wyjątków, ale nie jest to istotą tego tekstu.

Kolejną przydatną funkcją będzie funkcja losująca wartości w tablicy. Ona również została napisana jedynie w celu zwiększenia przejrzystości kodu

 Jeśli jesteśmy już przy wypełnianiu tablicy to warto jeszcze dodać sobie dwie stałe, które będą określały wielkość wykorzystywanych tablic i maksymalną wartość jaką będą mogły mieć ich elementy

W końcu przyszedł czas na funkcję main(). Początek wygląda podobnie jak poprzednio. Jest tutaj pobranie wszystkich dostępnych platform, wypisanie informacji na temat pierwszej z nich, a następnie utworzenie kontekstu na jej podstawie

 Kiedy udało się utworzyć prawidłowo kontekst czas na utworzenie buforów, które posłużą do wymiany danych pomiędzy naszym programem, a kernelem

 Najpierw utworzyliśmy standardowe tablice intów. Dwie z nich, które posłużą jako dane wejściowe zostały wypełnione losowymi wartościami. Ostatnia posłuży do zapisania wartości zwróconych przez kernel. W konstruktorze bufora podajemy kontekst, flagę oznaczającą czy będą to dane tylko do odczytu (CL_MEM_READ_ONLY) czy zapisu (CL_MEM_READ_ONLY) oraz rozmiar bufora (Tutaj uwaga: podczas prób doszedłem do tego, że jeśli korzystamy z urządzeń GPU to wielkością bufora musi być sizeof(int) * arraySize, zaś jeśli korzystamy z CPU to wielkość bufora ustawiamy jedynie na arraySize. Niestety nie udało mi się jak na razie znaleźć informacji dlaczego tak się dzieje).

Teraz pobieramy informację o dostępnych urządzeniach

 W tym momencie następuje otwarcie pliku .cl z kodem naszego kernela, który podam na końcu i wczytanie go do zmiennej typu std::string.

 Tworzymy zmienną przechowującą źródła kernela i przekazujemy jej wczytany przed chwilą ciąg znaków

 Teraz utworzymy zmienną reprezentującą program OpenCLa i zbudujemy go

Stwórzmy teraz kolejkę poleceń, i za jej pomocą wypełnijmy bufory wejściowe wartościami znajdującymi się w tablicach A i B

Pierwszym parametrem jest stworzony wcześniej bufor, który chcemy wypełnić, następnie wartość CL_TRUE oznacza, że czynność będzie blokująca, a więc program nie przejdzie dalej dopuki nie skończy tej operacji, później jest offset, rozmiar bufora i na końcu tablica, z której pobierzemy dane.

W takim razie nadeszła w końcu pora na utworzenie zmiennej reprezentującej kernel czyli funkcję zapisaną w naszym pliku .cl. Jednym z parametrów jest nazwa tej funkcji, którą mamy zamiar użyć

 Kiedy utworzyliśmy już kernel czas przekazać mu parametry z jakimi zostanie uruchomiony

Pierwszy parametr w funkcji setArg() mówi, który argument mamy zamiar zmieniać, gdzie oznacza pierwszy parametr, zaś 3 ( w tym wypadku) ostatni.

Czas na wykonanie w końcu obliczeń. Najpierw tworzymy zmienną cl::Event, która pozwoli nam poczekać na zakończenie wykonywania się kernela korzystając do tego z funkcji wait(), która zablokuje program dopóki nie zostanie wywołane zdarzenie. Sam kernel uruchamiamy funkcją enqueueNDRangeKernel(). Ważniejsze parametry to kernel, który chcemy wykonać, ilość danych (zauważ, że teraz już nie mnożymy ilości elementów przez rozmiar typu) określająca ile razy kernel zostanie uruchomiony, ilość wątków przypadających na jedną jednostkę obliczeniową, tworzą one grupę, wewnątrz której mogą się synchronizować, jeśli wszystkie wątki wykonują tą samą operację to robią to równolegle, jeśli część z nich, np. z powodu użycia instrukcji if wykonuje inne operacje to wszystkie wątki w grupie wykonają się sekwencyjnie, a ostatni parametr to wskaźnik na naszą zmienną event, który pozwoli wywołać zdarzenie po skończeniu operacji.

Teraz zostało tylko poczekanie na zakończenie obliczeń przez GPU i pobranie wyników z bufora

 Na sam koniec wypisujemy wyniki w konsoli, czyścimy pamięć i wychodzimy z funkcji main()

Została jeszcze kwestia pliku multiply.cl, w końcu jest on bardzo ważnym elementem całości i to w nim zapisane mamy operacje, jakie ma wykonać nasze GPU. Prezentuje się on następująco

Mamy tutaj tylko jedną funkcję multiply(), która jako parametry przyjmuje trzy wskaźniki oraz jedną wartość. __global oznacza, że jest to wskaźnik na bufor za alokowany w globalnej puli pamięci dostępnej dla urządzenia. Funkcja get_global_id(0) pobiera unikalne id uruchomionego kernela, dzięki czemu możemy określić, do którego elementu tablicy będziemy się odwoływać, ponieważ jeden uruchomiony kernel wykonuje tutaj operacje na jednym elemencie. Jeśli wszystko zrobiliśmy tak jak trzeba to w naszym przypadku get_global_id(0) powinno zawsze dać wartości od 0 do arraySize – 1 jednak dobrym nawykiem jest sprawdzanie tego przy pomocy instrukcji if.

A tak prezentuje się cały kod naszego programu

 Wynik jego wykonania:

p2

 

5. Zakończenie

Tym oto sposobem opisałem podstawy budowania programów wykorzystujących do obliczeń GPU i bibliotekę OpenCL. Mam nadzieję, że ta treść przyda się komuś. Jeśli znalazłeś błąd to śmiało pisz o nim w komentarzu lub na maila contact@zajacmarek.com . Jeśli chciałbyś gdzieś wykorzystać fragmenty tego wpisu to również proszę o wcześniejszy kontakt. Kopiowanie części lub całości bez mojej zgody jest zabronione.

Dla osób, które wolą wersję off-line zamieszczam też wpis w formie pdf-a: Wstęp do OpenCL

5 przemyśleń nt. „Wstęp do OpenCL

    1. Jeśli tylko znajdzie się czas to oczywiście postaram się zamieszczać kolejne wpisy, bo tematyka dość popularna i ciekawa. Jednak w tym momencie nie ma żadnych konkretów na ich temat.

  1. Fajnie, że napisałeś samouczek do OpenCL, zwłaszcza że nie musiałeś tego robić. GPU są bezkonkurencyjne – tanie, wydajne, dokładniejsze od CPU ( błąd kumulacyjny ), o dużej i kilkanaście razy bardziej wydajnej pamięci. Podrzucę ciekawostkę odnośnie różnego typu urządzeń w moim posiadaniu, które przetestowałem, aby ślepo nie  powtarzać „że są szybkie, bo są szybkie i jak tak słyszałem, że są szybkie bo są szybkie”:) :

    Celeron G3900 (2 rdzenie ~150zł) – 6GFLOPs,
    i7-7100 (4 rdzenie ~1100zł) – 8GFLOPs,
    Allwinner H3 Orange Pi One (4rdzenie ~5zł) 0.2GFLOPs,
    GT1030 (6144 rdzeni ~300zł ) 1200GFLOPs,
    GTX780 (24576 rdzeni ~800) 3700GFLOPs.

    Ja tam do niedawna myliłem wydajność obliczeniową z przepływnością ( praktyczną przepustowością ) pamięci i tak:

    powyższy Celeron 2 kanałowa pamięć DDR3 1333MHz ~5.5GBps,
    powyższy i7 z 1 kanałową pamięcią DDR4 2133MHz ~7GBps,
    powyższy H3 z 1 kanałową pamięcią DDr3 800MHz(?) ~0.7GBps
    GT1030 30GBps,
    GTx780 95GBps.

    Powodzenia we wdrażaniu oprogramowania od początku na GPU ( przyśpieszanie pojedynczych funkcji jest bez sensu – nawet jak 80%’ową zredukujesz do 0 sekund to masz marne przyśpieszenie rzędu 5x! ).
    Post Scriptum: mało rdzyniuf i mało ramu, to popularne błędy amatorów – 2GB RAM i pamięć globalna GPU 1GB  to aż za dużo!

    1. Wydajność procesorów Intela została zmierzona na pojedynczym rdzeniu, tak więc:

      Intel Celeron G3900: ~12GFLOPs,

      Intel i7 7100 ~32GFLOPs.

      Wydajność dowolnego procesora ściśle zależy od typu problemu, oraz umiejętności użycia dostępnego sprzętu przez programistę. Warto zwrócić uwagę, że przy aplikacji wielordzeniowej pasmo pamięci RAM jest dzielone między rdzenie i rzeczywista przepływność pamięci jest ponad x – razy mniejsza.

Dodaj komentarz

Twój adres email nie zostanie opublikowany. Pola, których wypełnienie jest wymagane, są oznaczone symbolem *