Skoči na vsebino

Obdelava slik1

Poudarjanje robov (ang. Edge enhancement) je slikovni filter, ki poveča kontrast robov. Rob je definiran kot znatna lokalna sprememba v intenziteti slike. Pri idealnem robu (ang. step) se sprememba intenzitete zgodi v koraku ene slikovne točke. Idealni robovi so v slikah zelo redki, še posebej, če slike prej gladimo, da iz njih odstranimo šum. Spremembe intenzitete se zato zgodijo v nekaj zaporednih slikovnih točkah, takemu robu pravimo nagib (ang. ramp). Filter za poudarjanje robov deluje tako, da na sliki prepozna ostre robove (na primer rob med motivom in ozadjem) in poveča kontrast slike na območju neposredno okoli roba. Zaradi tega je rob videti bolj definiran. Spodnja slika prikazuje primer poudarjanja robov:

Vhodna slika Izhodna slika
Vhodna slika Izhodna slika

Za postopek poudarjanja robov uporabimo Laplacovo jedro velikosti 3x3. Vsako slikovno točko in njeno okolico v izvorni sliki pomnožimo z Laplacovim jedrom, produkte seštejemo in vsoto uporabimo za vrednost slikovne točke v novi sliki s poudarjenimi robovi. Temu procesu pravimo konvolucija. Spodnja slika prikazuje en korak konvolucije z Laplacovim jedrom.

Poudarjanje robov

Na sliki je prikazano Lapacovo jedro, ki ga bomo uporabili za poudarjanje robov na sliki. Vidimo, kako poudarimo vrednost slikovne točke z vrednostjo 2, ki tvori rob s slikovno točko na svoji desni (vrednost 0). Novo vrednost, 6, dobimo tako, da jedro položimo na izvorno sliko, zmnožimo istoležne elemente ter seštejemo delne produkte.

Delo s slikami

Za branje slik z diska v pomnilnik ter zapisovanje nazaj na disk bomo uporabljali javno dostopno knjižnico STB. Knjižnica je dokaj obsežna, vendar bomo uporabili le funkciji stbi_load in stbi_write, ki sta definirani v zaglavjih stb_image.h in stb_image_write.h. Navodila za uporabo funkcij najdemo tukaj.

Slike v pomnilniku

Slike bomo najprej iz datotek na disku prebrali v pomnilnik. Slike so sestavljene iz množice slikovnih točk. Pri sivinskih slikah je običajno ena slikovna točka predstavljena z osmimi biti, ki določajo sivinski nivo (od črne do bele). Slikovna ravnina dimenzije Width x Height pa je v pomnilniku predstavljena kot vektor slikovnih točk, pri čemer do posamezne slikovne točke dostopamo na naslednji način: image[Y x Width + X]. Dostop do posamezne slikovne točke prikazuje spodnja slika.

Naslavljanje slikovnih točk - en kanal

Količine X (stolpec), Y (vrstica), Width (širina) in Height (višina) so izražene v slikovnih točkah.

Pri barvnih slikah je tipično vsaka slikovna točka predstavljena s 4 x 8 biti ali štirimi bajti. Prvi trije bajti določajo tri komponente barve (angl. RGB, Red, Green in Blue), medtem ko pri določenih zapisih (na primer png) zadnji bajt določa prosojnost. Pravimo tudi, da je slika sestavljena iz štirih kanalov. Vsaka slikovna točka je torej v pomnilniku predstavljena s štirimi vrednostmi, pri čemer je vsaka vrednost zapisana z enim bajtom. Štirikanalno slika dimenzij Width x Height v pomnilniku prikazuje spodnja slika.

Naslavljanje slikovnih točk - več kanalov

Do posamezne slikovne točke sedaj dostopamo na naslednji način: image[(Y x Width + X) x CPP + channel], pri čemer je CPP število kanalov na slikovno točko (ang. channels per pixel) in je lahko med 1 in 4, channel pa je indeks kanala z vrednostjo od 0 do 3.

V nadaljevanju se bomo omejili na štirikanalne slike (četudi so sivniske) in jih bomo s funkcijo stbi_load zapisali v pomnilnik. V primeru sivinskih slik (kanal Rn), bodo trije zadnji bajti slikovne točke neuporabljeni. Na prvi pogled je rešitev potratna, saj sivinske slike v pomnilniku zasedajo štirikrat več prostora, ampak bomo zaradi tega imeli isto kodo za obdelavo vseh vrst slik.

Ščepci za poudarjanje robov

Vsaka nit bo izračunala novo vrednost ene slikovne točke v sliki. V ta namen bo morala pri filtrih velikosti 3x3 dostopati še do osmih slikovnih točk v neposredni okolici slikovne točke, za katero računa novo vrednost. Pri izbranem Laplacovem jedru je dovolj, da dostopa samo do štirih sosednjih slikovnih točk, pri katerih ima jedro neničelne vrednosti.

V nadaljevanju predstavimo štiri ščepce - od najbolj neučinkovitega in hkrati najbolj preprostega do najbolj učinkovitega in najbolj kompleksnega.

Naiven filter za pudarjanje robov

Najbolj preprosto implementacijo ščepca sharpenGPU za uporabo filtra za poudarjanje robov prikazuje spodnja koda.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
//Simplest kernel (memory accesses not optimal)
__kernel void sharpenGPU(__global unsigned char *imageIn,
                         __global unsigned char *imageOut, int width,
                         int height, int cpp) {

  int y = get_global_id(0);
  int x = get_global_id(1);

  if (x < width && y < height) {
    // for each color channel
    for (int c = 0; c < cpp; c++) {
      unsigned char px01 =
          getIntensity(imageIn, y - 1, x, c, height, width, cpp);
      unsigned char px10 =
          getIntensity(imageIn, y, x - 1, c, height, width, cpp);
      unsigned char px11 = 
          getIntensity(imageIn, y, x, c, height, width, cpp);
      unsigned char px12 =
          getIntensity(imageIn, y, x + 1, c, height, width, cpp);
      unsigned char px21 =
          getIntensity(imageIn, y + 1, x, c, height, width, cpp);

      short pxOut = (5 * px11 - px01 - px10 - px12 - px21);
      if (pxOut > 255)
        pxOut = 255;
      if (pxOut < 0)
        pxOut = 0;
      imageOut[(y * width + x) * cpp + c] = (unsigned char)pxOut;
    }
  }
}

V zanki for se sprehodimo čez vse kanale in za vsak kanal posebej nit prebere vrednosti petih slikovnih točk iz slike: slikovne točke, za katero računa novo vrednost, ter njenih štirih sosedov (levo, desno, zgoraj in spodaj), kot to zahteva Laplacovo jedro. Za dostop do vrednosti posameznih slikovnih točk pripravimo funkcijo getIntensity:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
//Helper function to extract intensity value for a pixel form image
inline unsigned char getIntensity(
                        __global unsigned char *image, 
                        int y, int x,
                        int channel, int height, int width, 
                        int cpp) {
  if (x < 0 || x >= width)
    return 0;
  if (y < 0 || y >= height)
    return 0;
  return image[(y * width + x) * cpp + channel];
}

Tak ščepec je sicer preprost, ima pa eno hudo pomanjkljivost. Kot smo že večkrat zapisali, dostop do glavnega pomnilnika vedno opravijo vse niti iz snopa hkrati in je vedno dolžine 128 bajtov - štiri bajte za posamezno nit. V našem primeru vsaka nit v eni iteraciji potrebuje le enega od štirih bajtov, vendar pri branju prebere vse štiri, nepotrebne pa zavrže! Nato v naslednji iteraciji spet dostopa do istega segmenta v globalnem pomnilniku, čeprav je do tega segmenta dostopala že v prejšnji iteraciji! A kaj, ko je v prejšnji iteraciji podatke, ki jih takrat ni potrebovala, preposta zavrgla.

Dvodimenzionalni prostor NDRange nastavimo na naslednji način:

1
2
3
4
5
const size_t szLocalWorkSize[2] = {16, 16};
const size_t szGlobalWorkSize[2] = {
      ((width - 1) / szLocalWorkSize[1] + 1) * szLocalWorkSize[1], 
      ((height - 1) / szLocalWorkSize[0] + 1) * szLocalWorkSize[0]
      };

Niti se bodo izvajale v delovnih skupinah velikosti 16 x 16. Celotno število niti (globalna velikost prostora NDRange) je odvisno od velikosti vhodne slike in mora biti v obeh dimenzijah deljivo z velikostjo delovne skupine, v našem primeru 16. Ker dimenzije vhodnih slik niso nujno deljive s 16, pri določanju števila niti zaokrožimo velikost slike navzgor tako, da bo deljiva s 16.

Ščepec sharpenGPU se na GPE Nvidia Tesla K40 pri obdelavi slike velike 2580x1319 slikovnih točk izvede v 1.39 milisekundah: 

[patriciob@nsc-login 09-image-filter]$ srun prog sharpenGPU celada_in.png celada_out.png
Loaded image celada_in.png of size 2580x1319.
Program size = 5894 B 
Kernel Execution time is: 1.388 milliseconds

Filter za poudarjanje robov z upoštevanjem organizacije podatkov v pomnilniku

Bolj učinkovit ščepec dobimo, če upoštevamo mehanizem dostopanja do globalnega pomnilnika (ang. memory coalescing). Nit v snopu vedno dostopa do najmanj štirih bajtov, zato je prav, da to izkoristimo - z enim branjem preberimo vse štiri kanale! V ta namen uporabimo vektorske podatkovne tipe OpenCL. Vektorski podatkovni tipi nam omogočajo, da operacije izvajamo nad kratkimi vektorji dolžine 2, 4, 8 ali 16 elementov. Aritmetična operacija se izvede istočasno nad istoležnimi elementi vektorjev. Spodnja koda prikazuje ščepec sharpenGPU_vector, ki naenkrat prebere vse štiri kanale v vektor dolžine štiri ter operacijo, ki jo zahteva Laplacovo jedro, opravi nad vsemi elementi vektorja (vsemi kanali) naenkrat.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
__kernel void sharpenGPU_vector(__global uchar4 *imageIn,
                                __global uchar4 *imageOut, int width,
                                int height) {

  int y = get_global_id(1);
  int x = get_global_id(0);

  if (x < width && y < height) {
    short4 px01 = getIntensity4(imageIn, y - 1, x, height, width);
    short4 px10 = getIntensity4(imageIn, y, x - 1, height, width);
    short4 px11 = getIntensity4(imageIn, y, x, height, width);
    short4 px12 = getIntensity4(imageIn, y, x + 1, height, width);
    short4 px21 = getIntensity4(imageIn, y + 1, x, height, width);

    short4 pxOut = (short4)(5) * px11 - px01 - px10 - px12 - px21;

    imageOut[y * width + x] = convert_uchar4_sat(pxOut);
  }
}
Argumenti ščepca tipa uchar4 so vektorji, ki vsebujejo štiri 8-bitne vrednosti (tip uchar). Podobno je short4 vektor štirih 16-bitnih short podatkov. OpenCL funkcija convert_uchar4_sat (vrstica 17) pa eksplicitno pretvori vektor 16-bitnih elementov short4 v vektor 8-bitnih elementov uchar4 z uporabo saturacijske aritmetike. Za vektorsko branje posameznih slikovnih točk iz globalnega pomnilnika uporabimo funkcijo getIntensity4:

1
2
3
4
5
6
7
8
9
inline short4 getIntensity4(__global uchar4 *image, 
                        int y, int x, int height,
                        int width) {
  if (x < 0 || x >= width)
    return (short4)(0);
  if (y < 0 || y >= height)
    return (short4)(0);
  return convert_short4(image[y * width + x]);
}

OpenCL funkcija convert_short4 opravi eksplicitno pretvorbo iz vektorja uchar4 v vektor short4. Tokrat se ščepec sharpenGPU_vecotor na GPE Nvidia Tesla K40 pri obdelavi slike velike 2580x1319 slikovnih točk izvede v 0,66 milisekundah: 

[patriciob@nsc-login 09-image-filter]$ srun prog sharpenGPU_vector celada_in.png celada_out.png
Loaded image celada_in.png of size 2580x1319.
Program size = 5894 B 
Kernel Execution time is: 0.662 milliseconds

Vidimo, da je pohitritev zaradi upoštevanja dostopv do globalnega pomnilnika znatna in zato moramo vedno paziti, kako niti dostopajo do podatkov v globalnem pomnilniku.

Filter za poudarjanje robov z uporabo lokalnega pomnilnika

V prejšnjih ščepcih vsaka nit prebere pet slikovnih točk. Ker je vsaka slikovna točka v sliki soseda štirim slikovnim točkam (levo, desno, zgoraj in spodaj), bomo vsako slikovno točko brali iz glavnega pomnilnika štirikrat. 256 niti v delovni skupini velikosti 16x16 si praktično deli sosednje slikovne točke. Zato je bolje, če iz slike v globalnem pomnilniku preberemo podsliko ustrezne velikosti v lokalni pomnilnik. Tako bodo imele niti v delovni skupini vse potrebne slikovne točke v lokalnem pomnilniku. Dostop do podatkov o slikovnih točkah bo tako nekaj stokrat hitrejši, pri tem pa bomo vsako slikovno točko samo enkrat brali iz globalnega pomnilnika. Pravzaprav moramo v lokalni pomnilnik prebrati 18x18 slikovnih točk, kajti niti na robovih delovne skupine potrebujejo še dodaten rob slikovnih točk.

Ščepec sharpenGPU_local_vector za poudarjanje robov z uporabo lokalnega pomnilnika predstavlja spodnja koda.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
__kernel void sharpenGPU_local_vector_linear(__global uchar4 *imageIn,
                                      __global uchar4 *imageOut, int width,
                                      int height, int cpp) {
  int gy = get_global_id(1);
  int gx = get_global_id(0);

  int ly = get_local_id(1);
  int lx = get_local_id(0);

  int lsy=get_local_size(1);
  int lsx=get_local_size(0);


  // reserve local memory
  __local short4 imgChunk[CHUNK_SIZE * CHUNK_SIZE];

  // load an image chunk to local memory
  // linearize local indices and move block through the image chunk
  for (int lli=ly*lsx+lx;lli<CHUNK_SIZE*CHUNK_SIZE;lli+=lsx*lsy){

    int groupx=get_group_id(0)*lsx;   // global X offset for the work-group
    int groupy=get_group_id(1)*lsy;   // global Y offset for the work-group


    // Now calculate 2D global indices back from local linear ones:
    int gyi=groupy+(lli/CHUNK_SIZE)-1;    // ROW: global Y offset + row in the workgroup
    int gxi=groupx+lli%CHUNK_SIZE-1;      // COLUMN: global X offset + column in the workgroup

    // read pixel flom global to local memory
    imgChunk[lli]=getIntensity4(imageIn, gyi, gxi,
                                   height, width, cpp);
  }

  // wait for threads to finish
  barrier(CLK_LOCAL_MEM_FENCE);

  if (gx < width && gy < height) {
    int pxi = (ly + 1) * CHUNK_SIZE + (lx + 1);
    short4 px01 = imgChunk[pxi - CHUNK_SIZE];
    short4 px10 = imgChunk[pxi - 1]; 
    short4 px11 = imgChunk[pxi]; 
    short4 px12 = imgChunk[pxi + 1]; 
    short4 px21 = imgChunk[pxi + CHUNK_SIZE];

    short4 pxOut = (short4)(5) * px11 - px01 - px10 - px12 - px21;
    imageOut[gy * width + gx] = convert_uchar4_sat(pxOut);
  }
}

Najprej niti prenesejo košček slike (chunk) iz globalnega pomnilnika v lokalni pomnilnik (vrstice 19-32).

Predpostavimo, da je CHUNK_SIZE = 4 in da je velikost delovnih skupin 2x2 (get_local_size() nam v obeh dimenzijah vrne 2). To pomeni, da 4 niti morajo iz globalnega pomnilnika prenesti 16 elementov v imgChunk v lokalnem pomnilniku. Bodite pozorni, da je imgChunk deklariran kot 1D vektor (vrstica 12 v ščepcu sharpenGPU_local_vector ) in ne kot 2D polje! Niti podatke iz globalnega pomnilnika v lokalni pomnilnik prenašajo, kot prikazuje spodnja slika:

Branje dela slike v lokalni pomnilnik

Da bomo imeli dostope do zaporednih besed v globalnem pomnilniku, za en blok niti vedno izračunamo vse linearizirane indekse (lli) elementov v imgChunk, ki jih posamezne niti v bloku prenašajo. Te linearizirane indekse uporabimo tudi za določanje elementov iz globalnega pomnilnika, ki jih posamezna nit prenaša (slika na desni), s sledečo kodo: 

25
26
27
// Now calculate 2D global indices back from local linear ones:
int gyi=groupy+(lli/CHUNK_SIZE)-1; // ROW: global Y offset + row in the workgroup
int gxi=groupx+lli%CHUNK_SIZE-1;   // COLUMN: global X offset + column in the workgroup

kjer groupx in groupy določata globalni odmik delovne skupine v globalnem pomnilniku.

Pred računanjem izhodne slikovne točke se vse niti počakajo pri prepreki. Nato vsaka nit izračuna novo vrednost svoje slikovne točke, pri čemer podsliko bere iz lokalnega pomnilnika (imgChunk):

34
35
36
37
38
39
40
41
42
43
44
45
46
47
  // wait for threads to finish
  barrier(CLK_LOCAL_MEM_FENCE);

  if (gx < width && gy < height) {
    int pxi = (ly + 1) * CHUNK_SIZE + (lx + 1);
    short4 px01 = imgChunk[pxi - CHUNK_SIZE];
    short4 px10 = imgChunk[pxi - 1]; 
    short4 px11 = imgChunk[pxi]; 
    short4 px12 = imgChunk[pxi + 1]; 
    short4 px21 = imgChunk[pxi + CHUNK_SIZE];

    short4 pxOut = (short4)(5) * px11 - px01 - px10 - px12 - px21;
    imageOut[gy * width + gx] = convert_uchar4_sat(pxOut);
  }

Tokrat se ščepec sharpenGPU_local_vector na GPE Nvidia Tesla K40 pri obdelavi slike velike 2580x1319 slikovnih točk izvede v 0.56 milisekundah:

[patriciob@nsc-login 09-image-filter]$ srun prog sharpenGPU_local_vector_linear celada_in.png celada_out.png
Loaded image celada_in.png of size 2580x1319.
Program size = 7000 B 
Kernel Execution time is: 0.557 milliseconds

Vidimo, da nam uporaba lokalnega pomnilnika prinese dodatno pohitritev izvajanja, saj se posamezen slikovni element tokrat samo enkrat prebere iz globalnega pomnilnika (razen slikovnih elementov na robovih, ki se berejo dvakrat). V primerjavi z naivno različico ščepca pa je pohitritev 2,5 kratna! Slednja ugotovitev nas pravzaprav ne sme presenentiti, saj vemo, da je ozko grlo v računalnikih pravzprav pomnilnik DRAM in z upoštevanjem organizacije podatkov v pomnilniku DRAM, načinom dostopa do pomnilnika DRAM ter predpomnjenjem (saj je lokalni pomnilnik programsko krmiljen predpomnilnik) lahko bistveno vplivamo na čas izvajanja programov.

Celotno kodo iz tega poglavja najdete v mapi 09-image-filter tukaj.

  1. © Patricio Bulić in Davor Sluga, Univerza v Ljubljani, Fakulteta za računalništvo in informatiko. Gradivo je objavljeno pod licenco Creative Commons Priznanje avtorstva-Nekomercialno-Deljenje pod enakimi pogoji 4.0 Mednarodna. ↩