W ramach naszej witryny stosujemy pliki cookies. Kontynuując przeglądanie strony, wyrażasz zgodę na używanie przez nas plików cookies. Dowiedz się więcej tutaj
X

Księgarnia PWN

   

   

   

   

.

.

   

   

   

   

Bezpłatny newsletter


OpenCL – dodawanie wektorów

OpenCL – dodawanie wektorów

Proponujemy podrozdział książki Marka Sawerwaina „OpenCL. Akceleracja GPU w praktyce”. Autor w publikacji skupił się na przedstawieniu licznych przykładów, które ułatwią pracę programistom.

W prezentowanej przez nas części autor omawia dodawanie wektorów w języku C.

1.1. Dodawanie wektorów

W tym przykładzie będziemy obliczać sumę dwóch ciągów liczb, czyli wektorów. Z punktu widzenia OpenCL wektory będą reprezentowane przez jednowymiarowe tablice. W języku C czy C++ operację dodawania dwóch tablic możemy zapisać np. za pomocą jednej pętli for, co przedstawia się to następująco:

int i;

float a[N],b[N],c[N];

for( i=0 ; i<N ; i++ ) {

    c[i] = a[i] + b[i];}

Poszczególne elementy tablicy o nazwie c stanowią rezultat sumy elementów tablic a oraz b. Podstawowym wymaganiem jest, aby wszystkie trzy tablice miały takie same wymiary. Zamiana powyższego kodu na równoważny kod w OpenCL nie jest trudna, a więcej troski będzie wymagać przygotowanie danych i obsługa jądra obliczeniowego.

Operacja dodawania wektorów charakteryzuje się łatwością podania programu równoległego, co bardzo sprzyja implementacji równoległej. Poszczególne elementy tablicy mogą być przetwarzane w sposób niezależny. Dysponując wieloma jednostkami obliczeniowymi, można dodawać różne fragmenty wektorów, nie martwiąc się o spójność danych. Obliczanie sumy np. siódmego elementu zapiszemy jako

a[7] = b[7] + c[7];

obliczenie tego częściowego wyniku nie ma żadnego wpływu na obliczenia pozostałych elementów wektora. Na rysunku 1.1 przedstawiono schemat obliczania sumy wektora. Każdy element jest dodawany niezależnie od pozostałych.

 

wektory

Rys. 1.1. Ogólna idea przetwarzania równoległego wektora za pomocą wielu jednostek obliczeniowych

 

1.1.1. Jądro obliczeniowe

Jądro obliczeniowe OpenCL w dwóch wersjach zostało przedstawione na listingu 1.1. Przekazujemy trzy lub cztery argumenty. W argumencie pierwszym o nazwie a podajemy wskaźnik do wektora wynikowego, a dwa pozostałe parametry o nazwach b oraz c stanowią wektory wejściowe. Ostatni argument o nazwie N określa wielkość wektorów.

W pierwszym argumencie nie pojawia się słowo const, bowiem jest to wektor, do którego będziemy kierować obliczone elementy sumy. Słowo const pojawia się w dwóch pozostałych wektorach, bowiem z tych zmiennych tylko odczytujemy dane.

Nie występuje także pętla for, bowiem jej rolę przejęły poszczególne jednostki obliczeniowe. Dlatego w pierwszej kolejności odczytujemy indeks, jakim będziemy się zajmować:

int i = get_global_id(0);

Wartości funkcji get_global_id są związane z siatką obliczeniową. W naszym przypadku będzie ona miała kształt i wielkość wektora, krótko mówiąc, będzie to liczba z zakresu od zera do N-1, gdzie N to liczba elementów, co oznacza iż get_global_id będzie zwracać numery indeksów dla naszych tablic reprezentujących wektory. Wykorzystujemy tu ideę pokazaną na rys. 4.3. Proces sumowania możemy zrealizować na dwa sposoby: można przydzielać poszczególne jednostki przetwarzające oddzielnie. Inaczej mówiąc, jeden indeks/element tablicy jest przetwarzany tylko przez jeden rdzeń obliczeniowy. Nie jest to efektywne rozwiązanie, bowiem lepiej, aby pojedyncza jednostka obliczeniowa przetwarzała więcej elementów wektora. W naszym przypadku na razie pozostaniemy przy prostszym rozwiązaniu, tzn. będziemy obliczać tylko sumę jednego elementu bezpośrednio wskazywanego przez indeks z siatki obliczeniowej.

Jeśli wielkość siatki jest równa wielkości przetwarzanego wektora, to można usunąć instrukcję warunkową odpowiedzialną za sprawdzanie, czy został przekroczony zakres rozmiaru wektora. W takim przypadku do obowiązku programisty należy precyzyjne określenie rozmiaru siatki obliczeniowej na równy co do liczby elementów zawartych w wektorze. Z drugiej strony, obecność instrukcji if pozwala na zastosowanie większej siatki obliczeniowej o wielkości dopasowanej do liczby dostępnych rdzeni obliczeniowych, co może się przyczynić do większej wydajności końcowych obliczeń. Należy jednak podkreślić raz jeszcze, że obecność instrukcji if zabezpiecza przed przekroczeniem zakresu przydzielonej pamięci, jeśli rozmiar siatki byłby większy niż przetwarzany wektor.

Listing 1.1. Suma dwóch wektorów w OpenCL (wersja pierwsza sum1 pozbawiona instrukcji warunkowej oraz wersja sum2 z instrukcją if sprawdzającą zakres danych)

__kernel void sum1( __global const float *a, __global const float *b, __global float *c ) {

    int gid = get_global_id(0);

    c[gid] = a[gid] + b[gid];}

__kernel void sum2( __global const float* a, __global const float *b, __global float *c, const int N) {

    int gid = get_global_id(0);

    if(gid < N)

        c[gid] = a[gid] + b[gid];}

Przed uruchomieniem obliczeń należy przygotować dane i przesłać je do urządzenia obliczeniowego. Deklaracja zmiennych po stronie gospodarza może zostać zrealizowana w postaci statycznych tablic:

float host_in_a[N];

float host_in_b[N];

float host_out_c[N];

gdzie: N – stała oznaczająca wielkość wektora np.:

const int N = 16;

Niezbędne jest też obliczenie rozmiaru danych, jakie będą kopiowane pomiędzy pamięcią urządzenia a pamięcią gospodarza. Ponieważ mamy do czynienia z prostą tablicą, to wystarczy przemnożyć liczbę elementów przez wielkość pojedynczego elementu:

size_t dataBytes = sizeof(float)*N;

Należy też przygotować zmienne tablicowe po stronie hosta. W naszym przykładzie tablica host_in_a zostanie wypełniona elementami w kolejności rosnącej, a tablica host_in_b w kolejności malejącej. W zmiennej host_out_c umieszczone zostaną zera, a właściwe dane zostaną umieszczone w tej zmiennej dopiero po zakończeniu obliczeń przez jądro obliczeniowe. Zamieniając te objaśnienia na odpowiednią pętlę for, otrzymujemy następujące linie kodu:

for(i=0;i<N;i++) {

    host_in_a[i]=static_cast<float>(i);

    host_in_b[i]=static_cast<float>(N-i);

    host_out_c[i]=0.0f;}

Kolejne czynności są identyczne jak omówione w poprzednim rozdziale, dlatego nie będziemy ich w tym miejscu powtarzać, możemy od razu przejść do utworzenia odpowiednich buforów w pamięci urządzenia:

device_in_a = clCreateBuffer( cl_compute_context, CL_MEM_READ_ONLY, dataBytes, NULL, &cli_err_num);

device_in_b = clCreateBuffer( cl_compute_context, CL_MEM_READ_ONLY, dataBytes, NULL, &cli_err_num);

device_out_c = clCreateBuffer( cl_compute_context, CL_MEM_WRITE_ONLY, dataBytes, NULL, &cli_err_num);

Dwa pierwsze bufory reprezentują wektory a oraz b, a trzeci bufor – tylko do zapisu – to bufor wektora c, gdzie zapiszemy wynik obliczeń. Deklaracje są identyczne jak w poprzednim przykładzie.

Bufory dla wektorów a oraz b wymagają przesłania danych z pamięci gospodarza, co zrealizujemy w sposób podobny jak w poprzednim przykładzie:

cli_err_num = clEnqueueWriteBuffer( cl_compute_command_queue, device_in_a, CL_TRUE, 0, dataBytes, &host_in_a[0], 0, NULL, NULL);

cli_err_num = clEnqueueWriteBuffer( cl_compute_command_queue, device_in_b, CL_TRUE, 0, dataBytes, &host_in_b[0], 0, NULL, NULL);

Do zmiennej c nie warto przesyłać danych, bowiem po wywołaniu jądra obliczeniowego będziemy odczytywać dane z urządzenia obliczeniowego. Kolejnym zadaniem jest utworzenie obiektu jądra oraz kompilacja, ale te czynności są podobne do poprzednio omówionych, więc możemy przejść do ustalenia parametrów jądra. W zależności od tego, jaką wersję będziemy implementować, mamy trzy lub cztery parametry. W obu przypadkach trzy pierwsze argumenty określamy w podobny sposób:

cli_err_num  = clSetKernelArg( kernel, 0, sizeof(cl_mem), &device_in_a );

cli_err_num |= clSetKernelArg( kernel, 1, sizeof(cl_mem), &device_in_b );

cli_err_num |= clSetKernelArg( kernel, 2, sizeof(cl_mem), &device_out_c );

Różnica tkwi w ostatnim parametrze. Brakujący czwarty parametr dla drugiej wersji jądra, gdzie podajemy dodatkowy parametr N określający wielkość tablicy, jest definiowany w następujący sposób:

cli_err_num |= clSetKernelArg( kernel, 3, sizeof(int), &N );

Mimo obecności dodatkowego argumentu dla jądra, jeśli wykorzystamy tylko globalną siatkę obliczeniową, to wywołanie jądra ma w obu przypadkach identyczną postać:

globalWorkSize = N;

cli_err_num = clEnqueueNDRangeKernel( cl_compute_command_queue, kernel, 1, NULL,  &globalWorkSize, NULL, 0, NULL, NULL );

Do zmiennej globalWorkSize przepisywany jest wymiar wektora i siatka obliczeniowa ma kształt 1 x N. Jednakże wyznaczony wektor, będący sumą dwóch innych wektorów, znajduje się nadal w urządzeniu obliczeniowym, toteż należy go przesłać do pamięci hosta. Proces odczytania danych będzie podobny do podanego w poprzednim przykładzie w rozdz. 4.2:

cli_err_num = clEnqueueReadBuffer( cl_compute_command_queue, device_out_c, CL_TRUE, 0, dataBytes, &host_out_c[0], 0, NULL, NULL);

Pozostało już tylko wyświetlić dane, które zostały przeniesione z pamięci urządzenia OpenCL do pamięci gospodarza. Najłatwiej wykorzystać do tego printf oraz pętlę for:

printf("Obliczone wartości:\n");

for(i=0;i<N;i++)

    printf("%f ", host_out_c[i]);

printf("\n")

Polecamy

Partnerzy