USG/GPU: Różnice pomiędzy wersjami

Z Brain-wiki
(Utworzono nową stronę "=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....")
 
Linia 1: Linia 1:
 
=Wstęp do obliczeń równoległych na 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.  
 
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.
+
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.
 
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ę:
 
Na sam początek importujemy bibliotekę:
Linia 26: Linia 26:
 
</source>
 
</source>
  
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.
+
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
 
Na początek sprawdzimy jakie procesory mamy do dyspozycji
 
<source lang=python>
 
<source lang=python>
Linia 37: Linia 37:
 
Wybierzemy jedno z urządzeń i skompilujemy nasz kernel do obliczeń na tym urządzeniu
 
Wybierzemy jedno z urządzeń i skompilujemy nasz kernel do obliczeń na tym urządzeniu
 
<source lang=python>
 
<source lang=python>
my_device = [platform[0].get_devices()[0]] #ta czesc kodu musi byc dostosowana do sprzetu na ktorym sa prowadzone cwiczenia
+
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)
 
ctx = cl.Context(my_gpu_device)
 
queue=cl.CommandQueue(ctx, device=my_gpu_device[0])  
 
queue=cl.CommandQueue(ctx, device=my_gpu_device[0])  
 
mod = cl.Program(ctx,kernels).build()
 
mod = cl.Program(ctx,kernels).build()
 
</source>
 
</source>
Musimy jeszcze przygotować dane. Na początek zdefiniujemy przykładowe dane - pamiętaj przy tym, że obliczenia wykonywać możemy na liczbach pojedynczej precyzji (!).
+
Musimy jeszcze przygotować dane. Na początek zdefiniujemy przykładowe dane - pamiętaj przy tym, że obliczenia wykonywać możemy na liczbach pojedynczej precyzji.
  
 
<source lang=python>
 
<source lang=python>
 
in=np.zeros((300,200,2))
 
in=np.zeros((300,200,2))
in[:,:,0]=np.arange((300*200)).reshape((300,200))
+
in[:,:,0] = np.arange((300*200)).reshape((300,200))
in[:,:,1]=np.arange((300*200)).reshape((300,200)) - 100
+
in[:,:,1] = np.arange((300*200)).reshape((300,200)) - 100
in=in.astype(np.float32) #zmieniamy dane na pojedynczą precyzję
+
in = in.astype(np.float32) # zmieniamy dane na pojedynczą precyzję
 
out = np.zeros((300,200)).astype(np.float32)
 
out = np.zeros((300,200)).astype(np.float32)
 
add = np.float32(add)
 
add = np.float32(add)
Linia 54: Linia 54:
 
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.
 
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.
 
<source lang=python>
 
<source lang=python>
input= cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=in)
+
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)
+
output = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=out)
 
</source>
 
</source>
 
Jesteśmy już gotowi do wywołania naszego kernela:
 
Jesteśmy już gotowi do wywołania naszego kernela:
 
<source lang=python>
 
<source lang=python>
event=mod.Sum(queue,(np.int32(300),np.int32(200),1),(1,1,1),input,output,add)
+
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)
 
# 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()     
 
event.wait()     
 
</source>
 
</source>
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).
+
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).
 
<source lang=python>
 
<source lang=python>
event=cl.enqueue_copy(queue,out,output)
+
event = cl.enqueue_copy(queue, out, output)
 
event.wait()
 
event.wait()
 
</source>
 
</source>
 
Proszę sprawdzić dane wynikowe i porównać je z analogicznymi danymi otrzymanymi z obliczeń liniowych:
 
Proszę sprawdzić dane wynikowe i porównać je z analogicznymi danymi otrzymanymi z obliczeń liniowych:
 
<source lang=python>
 
<source lang=python>
out_liniowy = in[:,:,0]+in[:,:,1]
+
out_liniowy = in[:,:,0] + in[:,:,1]
 
</source>
 
</source>
  
 
===Zadanie===
 
===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.
+
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.

Wersja z 19:42, 8 maj 2016

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

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.

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