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 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;
        }
"

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 przenieść 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 (!).

in=np.zeros((300,200,2))
in[:,:,0]=np.arange((300*200)).reshape((300,200))
in[:,:,1]=np.arange((300*200)).reshape((300,200)) - 100
in=in.astype(np.float32) #zmieniamy dane na pojedynczą precyzję
out = np.zeros((300,200)).astype(np.float32)
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= cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=in)
output= cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=out)

Jesteśmy już gotowi do wywołania naszego kernela:

event=mod.Sum(queue,(np.int32(300),np.int32(200),1),(1,1,1),input,output,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ść spowrotem 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 z wartości bezwzględnych elementów obu tablic i zwracać największy wspólny dzielnik elementów tych tablic.