Читать книгу Przetwarzanie obrazów grafiki 2D - Группа авторов - Страница 2

Pierwszy przykład polega na nałożeniu gradientu na istniejący obraz. Na podstawie tego przykładu pokażemy podstawowe elementy OpenCL, jakie wykorzystuje się w przetwarzaniu obrazów.

Pierwszy problem, jaki pojawia się nawet w najprostszym zadaniu, to wybór formatu przechowywania informacji o obrazie. W tym przykładzie będziemy stosować format RGBA1, czyli trójkę kolorów, tj. czerwony, zielony, niebieski oraz kanał alfa. Wymaga to naturalnie przygotowania odpowiednich buforów danych, dla których trzeba określić format obrazu. W tym celu stosowany jest typ cl_image_format

cl_image_format image_format;

Typ cl_image_format ma dwa pola. W przykładzie poszczególne kanały kolorów oraz kanał alfa są opisywane przez liczby całkowite oraz układ kanałów – wspominany już układ RGBA:

image_format.image_channel_data_type = CL_UNSIGNED_INT8;

image_format.image_channel_order = CL_RGBA;

Ponieważ mamy bufor wejściowy oraz wyjściowy, to potrzebne są dwa obiekty do buforów:

cl_mem input_image;

cl_mem output_image;

Utworzenie obiektów reprezentujących obrazy to zadanie funkcji clCreateImage2D, gdzie z punktu widzenia przetwarzanego obrazu najważniejsze parametry to format oraz szerokość (parametr WIDTH) i wysokość (HEIGHT) obrazu. W przypadku obrazu wejściowego:

input_image = clCreateImage2D(cl_compute_context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &image_format, WIDTH, HEIGHT, 0, (void*)bytes, &cli_err_num);

dodatkowo podana została zmienna bytes, z której kopiowane są dane do bufora wejściowego. Należy podkreślić, iż dane są kopiowane ze zmiennej umieszczonej w pamięci gospodarza do pamięci urządzenia obliczeniowego.

Zmienna reprezentująca bufor wyjściowy jest tworzona w następujący sposób:

output_image = clCreateImage2D( cl_compute_context, CL_MEM_WRITE_ONLY, &image_format, WIDTH, HEIGHT, 0, NULL, &cli_err_num );

Poszczególne etapy przygotowawcze do uruchomienia jądra obliczeniowego są podobne jak w innych przykładach. Jądro obliczeniowe z listingu 1 używa dwóch parametrów, w których należy podać obiekty reprezentujące bufor wejściowy oraz wyjściowy:

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

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

Po ustaleniu parametrów procedury obliczeniowej należy określić postać siatki obliczeniowej. Ponieważ przetwarzany będzie obraz dwuwymiarowy, to deklaracja zmiennej opisującej siatkę jest następująca:

size_t globalWorkSize[2];

Będzie ona zawierać wymiary tworzonego obrazu, dlatego umieszczamy tu informacje o szerokości i wysokości:

globalWorkSize[0]=HEIGHT;

globalWorkSize[1]=WIDTH;

W tym momencie można uruchomić jądro obliczeniowe za pomocą clEnqueueNDRangeKernel w następujący sposób:

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

Sposób, w jaki realizujemy nałożenie gradientu na istniejący obraz, został zawarty w kodzie procedury obliczeniowej pokazanej w listingu 1, jednak opis zostanie przedstawiony później, bowiem po wykonaniu obliczeń przez jądro obliczeniowe, należy odebrać dane. Wymaga to zdefiniowania dwóch dodatkowych zmiennych, w których podaje się współrzędne źródłowe (wskazujące na początek obrazu) origin oraz współrzędne docelowe region:

size_t origin[3], region[3];

Wpisanie wartości określających początek obrazu w przypadku źródła przedstawia się następująco:

origin[0] = 0; origin[1] = 0; origin[2] = 0;

Tradycyjnie początek jest umieszczony we współrzędnych (0, 0), ale mamy też trzeci parametr, który jest stosowany dla obrazów trójwymiarowych. W przykładzie zajmujemy się obrazem dwuwymiarowym, więc trzeci parametr przyjmuje wartość zero.

Dla zmiennej region dwa pierwsze parametry to lewy dolny wierzchołek obrazu, trzeci parametr znajduje zastosowanie w przypadku obrazu trójwymiarowego. Jednak jeśli region odnosi się do obrazu dwuwymiarowego, standard OpenCL wymaga, aby była wstawiona tam wartość jeden:

region[0] = WIDTH; region[1] = HEIGHT; region[2] = 1;

Ostatecznie odczyt danych z przetworzonego obrazu do zmiennej bytes (umieszczonej w pamięci gospodarza) zrealizujemy w następujący sposób:

cli_err_num = clEnqueueReadImage(cl_compute_command_queue, output_image, CL_TRUE, origin, region, 0, 0, (void*)bytes, 0, NULL, NULL);

Wartość CL_TRUE umieszczona w trzecim parametrze oznacza, że proces kopiowania będzie tzw. procesem blokowanym, inaczej mówiąc, funkcja clEnqueueReadImage nie zwróci sterowania do następnej instrukcji do czasu zakończenia procesu kopiowania danych. Naturalnie można to zmienić za pomocą odpowiednio zbudowanej procedury do obsługi zdarzeń.

Po omówieniu niezbędnych czynności, jakie należy wykonać, aby uruchomić jądro obliczeniowe, możemy omówić samo jądro obliczeniowe. Funkcja imageKernel przyjmuje dwa argumenty: src_image – reprezentuje obraz wejściowy oraz dst_image – obraz wyjściowy.

Bardzo istotną rolę odgrywa deklaracja tzw. samplera (obiekt sampler). Obiekt ten opisuje sposób, w jaki odczytywane są piksele z obrazu wejściowego. Listing 1 definiuje sampler jako stałą. Poszczególne flagi, jakie zostały użyte, mają następujące znaczenie: CLK_NORMALIZED_COORDS_FALSE oznacza, że współrzędne poszczególnych pikseli obrazu nie będą znormalizowane, co pozwala zapisać wartości współrzędnych np. od 0 do 768. Normalizacja oznacza, że będą stosowane współrzędne od zera do jedności.

Druga zastosowana flaga CLK_ADDRESS_CLAMP oznacza, iż wszystkie piksele odczytywane spoza zakresu obrazu będą posiadały wartości pikseli znajdujących się na brzegu oryginalnego obrazu.

Ostatnia, trzecia flaga dotyczy sposobu odczytywania pikseli. Dla wartości CLK_FILTER_NEAREST odczyt pikseli polega na stosowaniu tzw. filtrowania bezpośredniego. Łatwo to wytłumaczyć, posługując się np. jedną linią obrazu. Jeśli odczytuje się piksel np. numer dwa, to można zwrócić informację o nim, podobnie w przypadku piksela numer trzy. Jednakże, gdyby obraz jednej linii został powiększony dwukrotnie, to np. piksel numer dwa zostanie powielony. Podobnie będzie z pikselem numer trzy. Jeśli obraz nie jest powiekszany o całkowitą wielokrotność, to trzeba wybrać piksel, który zostanie powielony (numer dwa czy też numer trzy). Niewątpliwie pogorszy to jakość obrazu. Jeszcze gorszy przypadek to zmniejszenie obrazu, trzeba bowiem część pikseli usunąć. Dlatego lepszym rozwiązaniem jest stosowanie filtru liniowego CLK_FILTER_LINEAR, który poprawia znacznie jakość odczytywanych pikseli przez uśrednienie wartości piksela za pomocą pikseli sąsiednich.

Jednak w naszym przypadku wystarczy odczyt wartości piksela sąsiedniego, dodatkowo nie są stosowane współrzędne znormalizowane, więc poszczególne piksele są adresowane bezpośrednio.

Po definicji samplera jądro wykonuje operację odczytania piksela z obrazu źródłowego:

uint4 pixel = read_imageui(src_image, sampler, coord);

Następnie współrzędne zapisane w zmiennej coord są odczytywane przez globalne identyfikatory uzyskane za pomocą siatki globalnej:

int2 coord = (int2)(get_global_id(0), get_global_id(1));

Oczywiście czynność budowy zmiennej coord wykonuje się przed odczytem wartości piksela. Kolejna linia jądra obliczeniowego do zmiennej pixel2 wpisuje nową wartość. Polega to na tym, iż wartości współrzędnych x oraz y stanowić będą składowe kolorów czerwonego (składowa R) oraz zielonego (składowa G) pozostałe dwie składowe zostają ustalone na zero:

pixel2 = (uint4)(coord.x, coord.y, 0, 0);

Teraz dodajemy do siebie dwa piksele:

pixel = pixel + pixel2;

Może się okazać, iż wartości dwóch pierwszych składowych przekroczą wielkość 255, dlatego dodajemy dwie instrukcje warunkowe sprawdzające zakres wartości w dwóch pierwszych składowych piksela, co widać na listingu 1.