USG/GPU

Wstęp do obliczeń równoległych na GPU

W tej części ćwiczeń stworzymy pierwszy program wykorzystujący bibliotekę OpenCL do obliczeń na kartach graficznych. Posłużymy się pythonowym wrapperem do OpenCL - biblioteką PyOpencl. Same kernele będziemy pisać w języku C, ale wszystkie dodatkowe procedury będziemy mogli pisać przy użyciu standardowych poleceń Pythona. Zaczniemy od stworzenia prostego programu sumującego dwie tablice liczb zmiennoprzecinkowych i następnie mnożącego tę sumę przez pewien ustalony mnożnik. Na sam początek importujemy bibliotekę:

import pyopencl as cl

Kod OpenCL przechowywać będziemy w postaci zmiennej tekstowej. Stworzymy prosty kernel, który dostaje wejściu dwie tablice o wymiarach odpowiednio [math]N\times M\times 2[/math] i [math]N\times M[/math] oraz zmienną z mnożnikiem. Kernel wykorzystywać będzie [math]N\cdot M[/math] jednostek obliczeniowych. Pierwsza tablica zawiera nasze dane wejściowe (ostatni wymiar koduje indeks tablicy wejściowej); druga tablica będzie tablicą w której przechowywać będziemy wynik obliczeń. Przykładowy kernel zdefiniowany jest poniżej:

kernels="
        __kernel void Sum(__global float *in, __global float *out, const float add)
        {
        //zaczynamy od zadeklarowania zmiennych
        const int n = get_global_id(0); //indeks w pierwszym wymiarze danej jednostki roboczej
        const int m = get_global_id(1); //indeks w drugim wymiarze danej jednostki roboczej
        const int M = get_global_size(1); 
        const int nm = n*M + m; // indeks w tablicy wyjsciowej odpowiadajacy wartości o współrzędnych (n,m)
        __private int index1; // indeks w zmiennej wejsciowej odpowiadajacy wartości o współrzędnych (n,m,0) 
        __private int index2; // indeks w zmiennej wejsciowej odpowiadajacy wartości o współrzędnych (n,m,1) 
        index1 = 0 + m * 2 + n * M  * 2;
        index2 = 1 + m * 2 + n * M  * 2;
        out[nm] = (in[index1]+in[index2]) * add;
        }
"

Kernel ten dodaje do każdej komórki tablicy wejściowej wartość zmiennej add. W celu wywołania kernela będziemy musieli przygotować środowisko OpenCL, skompilować kod OpenCL (za pomocą odpowiednich funkcji w Pythonie) oraz zarezerwować miejsce w pamięci i przesłać dane do bufora pamięci na GPU. Na początek sprawdzimy jakie procesory mamy do dyspozycji

platform = cl.get_platforms()

Proszę wypisać zmienną platform. Zależnie od sprzętu i oprogramowania do dyspozycji będziemy mieli jedną lub więcej platform (gdy np. mamy proces Intela i kartę graficzną AMD) wraz z określoną liczbą urządzeń na każdej platformie. Dostępne urządzenia na pierwszej platformie możemy podejrzeć wywołując:

print(platform[0].get_devices())

Wybierzemy jedno z urządzeń i skompilujemy nasz kernel do obliczeń na tym urządzeniu

my_device = [platform[0].get_devices()[0]] # ta czesc kodu musi byc dostosowana do sprzetu na ktorym sa prowadzone cwiczenia
ctx = cl.Context(my_gpu_device)
queue = cl.CommandQueue(ctx, device=my_gpu_device[0]) 
mod = cl.Program(ctx,kernels).build()

Musimy jeszcze przygotować dane. Na początek zdefiniujemy przykładowe dane - pamiętaj przy tym, że obliczenia wykonywać możemy na liczbach pojedynczej precyzji.

inp = np.zeros((300,200,2))
inp[:,:,0] = np.arange((300*200)).reshape((300,200))
inp[:,:,1] = np.arange((300*200)).reshape((300,200)) - 100
inp = inp.astype(np.float32) # zmieniamy dane na pojedynczą precyzję
out = np.zeros((300,200)).astype(np.float32)
add=3.14
add = np.float32(add)

Następnie zarezerwujemy pamięć, przenosząc od razu dane wejściowe. Poniższe zmienne stanowią rodzaj wskaźnika do interesujących nas miejsc w pamięci urządzenia.

input_buffer = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=inp)
output_buffer = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=out)

Ponieważ zmienna add występuje jako stała, nie musimy wykonywać dodatkowych operacji rezerwowania pamięci i możemy podać ją bezpośrednio przy wywoływaniu kernela. Jesteśmy już do tego gotowi:

event = mod.Sum(queue, (np.int32(300), np.int32(200), 1), (1,1,1), input_buffer, output_buffer, add)
# jak widać nasz kernel możemy wywołać jako metodę zmiennej mod. przyjmuje ona na wejściu potok (?) obliczeniowy, całkowite wymiary grup roboczych, wymiary ???? oraz wskaźniki do buforów)
event.wait()

Po wywołaniu kernela pozostaje nam tylko przenieść z powrotem dane wynikowe do pamięci hosta (tak aby były one dostępne do dalszych, standardowych obliczeń z poziomu Pythona).

event = cl.enqueue_copy(queue, out, output)
event.wait()

Proszę sprawdzić dane wynikowe i porównać je z analogicznymi danymi otrzymanymi z obliczeń liniowych:

out_liniowy = in[:,:,0] + in[:,:,1]

Zadanie

Proszę stworzyć kernel, który będzie dostawać dwie tablice liczb zmiennoprzecinkowych, a następnie liczyć sufit (ang. ceil) z wartości bezwzględnych elementów obu tablic i zwracać największy wspólny dzielnik elementów tych tablic. Porównać działanie (wyniki) programu z poniższą liniową implementacją:

import numpy as np
import fractions
gcd = np.frompyfunc(fractions.gcd, 2, 1)

def funkcja(a,b):
    a=np.ceil(np.abs(a))
    b=np.ceil(np.abs(b))
    return gcd(a,b)

Profilowanie

Główną przyczyną naszego zainteresowania obliczeniami na GPU jest potencjalny spadek czasu wykonywania wymaganych obliczeń. Z tego punktu widzenia istotna jest możliwość dokładnego określenia czasu obliczeń i porównania algorytmów liniowych z równoległymi. Praktyczny czas trwania obliczeń jest zależny od implementacji sprzętowej, parametrów komputera itp. Nie jest to pojęcie tożsame ze złożonością obliczeniową.
Do mierzenia czasu obliczeń liniowych wykorzystamy standardową bibliotekę time. Aby wykonać pojedynczy pomiar czasu działania funkcji y(x) należy wywołać kod poniższej postaci

from time import clock
start = clock() # wskazanie zegara w danej chwili
y(x)
end = clock()
t=end-start # zmierzony czas

Oczywiście, pomiar taki jest obarczony błędem pomiarowym. Z tego względu rozsądne szacunki należy wykonywać w oparciu o dużą liczbę pomiarów.
Biblioteka time niestety nie współpracuje dobrze z obliczeniami równoległymi - zawyżając wartość efektywnego czasu obliczeń. Z tego względu, OpenCL dostarcza własnych narzędzi do profilowania. Aby możliwe było profilowanie danej kolejki należy deklarując użyć odpowiedniego parametru:

queue=cl.CommandQueue(ctx, device=my_gpu_device[0], properties=cl.command_queue_properties.PROFILING_ENABLE)

Aby zmierzyć czas wykonywania naszego kernela Sum z wcześniejszej części tekstu, wykonujemy

event = mod.Sum(queue, (np.int32(300), np.int32(200), 1), (1,1,1), input, output, add)
event.wait()
t = (event.profile.end-event.profile.start)*1e-9 # zmierzony czas

Zwróćmy uwagę na to, że wynik dotyczy wyłącznie samych obliczeń - nie uwzględnia on czasu potrzebnego na przeniesienie danych z pamięci hosta do pamięci urządzenia i z pamięci urządzenia do hosta. Czas komunikacji między urządzeniami jest względnie długi - sprawia to, że obliczenia równoległe są opłacalne dopiero przy odpowiednio dużych danych wejściowych. W oparciu o poniższy kod możemy zmierzyć, ile czasu trwa skopiowanie zawartości zmiennej in do bufora input w pamięci urządzenia:

start_event=cl.enqueue_marker(queue)
input = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=in)
t = (event.profile.end-start_event.profile.start)*1e-9 # zmierzony czas

Analogicznie zmierzyć możemy czas potrzebny na przeniesienie wyników z pamięci urządzenia do pamięci hosta.

Zadanie

1. Dla kilku wybranych rozmiarów tablicy danych wejściowych zmierzyć (liczba powtórzeń >10000) czas potrzebny na:
   • przeniesienie danych z hosta do urządzenia;
   • wykonanie kernela Sum;
   • przeniesienie danych z urządzenia do hosta.

   Czy czasy te rosną liniowo z rozmiarem tablicy?

2. Zmierzyć (liczba powtórzeń >10000) czas trwania obliczeń funkcji np.add(x,y) dla kilku wybranych rozmiarów tablic danych wejściowych (rozmiary tożsame z rozmiarami z podpunktu pierwszego).

   Czy wielkość ta zmienia się liniowo z rozmiarem tablicy?

3. Na podstawie powyższego ocenić: (dla konkretnego komputera) od jakiego rozmiaru tablicy zrównoleglenie obliczeń zaczyna być opłacalne?