KÉPFELDOLGOZÓ ALGORITMUSOK FEJLESZTÉSE GRAFIKUS HARDVER KÖRNYEZETBEN

Budapesti Műszaki és Gazdaságtudományi Egyetem Villamosmérnöki és Informatikai Kar Irányítástechnika és Informatika Tanszék

KÉPFELDOLGOZÓ ALGORITMUSOK FEJLESZTÉSE GRAFIKUS HARDVER KÖRNYEZETBEN

Takács Gábor

Konzulens: Vajda Ferenc PhD, adjunktus

1


TARTALOMJEGYZÉK: A kutatási projekt ismertetése .................................................................................................... 3 A téma ismertetése ................................................................................................................. 3 Bevezetés.................................................................................................................................... 4 Elmélet ....................................................................................................................................... 4 Számítógépes látás ................................................................................................................. 4 Szegmentálás .......................................................................................................................... 5 GPU (Graphics Processing Unit) ........................................................................................... 5 CUDA (Compute Unified Device Architecture).................................................................... 6 Megvalósított algoritmusok........................................................................................................ 8 CPU-n futó kód ...................................................................................................................... 8 Szürkeárnyalatosítás............................................................................................................. 10 Küszöbözés........................................................................................................................... 11 Konvolúció ........................................................................................................................... 12 Alakzatok keresése ............................................................................................................... 14 Jövőbeli tervek ......................................................................................................................... 15 Felhasznált irodalom ................................................................................................................ 15

2


A kutatási projekt ismertetése A számítási teljesítmény növekedésének és a magas integráltságú eszközök költségeinek csökkenésének

köszönhetően

egyre

növekszik

az

igény

valósidejű

képfeldolgozó

rendszerekre. A képfeldolgozás eszköztára igen széles, és gyakorlatilag elképzelhetetlen, hogy azok, akik különleges algoritmusokat kutatnak, mindig ismerjék az adott feladathoz használandó legmegfelelőbb platformot. A tanszék 3D érzékelés és Mobilrobotika kutatócsoportjában elindult kutatási projekt keretében azt vizsgáljuk, hogy egy adott feladat esetén melyik eszközt esetleg eszközkészletet célszerű alkalmazni. A projekt részét képezi a képfeldolgozás során alkalmazott robusztus szegmentálási algoritmusok kialakítása. A feladat egy régóta kutatott részterület, amelynek valósidejű megoldása egyelőre várat magára.

A téma ismertetése A téma keretében meglévő szegmentálási algoritmusok GPU-n történő implementálása és továbbfejlesztése a cél, ami a következő lépéseket tartalmazza:  A szegmentálás és hozzá kapcsolódó képfeldolgozási részterület alapos áttekintése, illetve az egyes algoritmusok grafikus kártyán történő implementálás lehetőségeinek vizsgálata. Képfeldolgozó könyvtárrendszerek megismerése.  A szegmentálás egyes algoritmusainak implementálása Direct3D illetve CUDA környezetben.  Az algoritmusok továbbfejlesztési, gyorsítási lehetőségeinek vizsgálata, illetve a szegmentálás háromdimenziós kiterjesztése videoszekvenciákra mozgáskövetéssel.

3


Bevezetés A képfeldolgozás egyik leginkább kutatott részfeladata a szegmentálás, melyre számos jó algoritmus létezik, de ezek viszonylag lassúak. Manapság viszont egyre nagyobb az igény a feladatok valósidejű végrehajtására. Az algoritmusok gyorsítására több lehetőség is létezik. Ilyen például a CPU-k sebességének növelése, de belátható, hogy rövid időn belül a processzorok sebessége nem fog olyan mértékben fejlődni, hogy valós időben le tudja futtatni ezeket az algoritmusok, ezért ez nem jó megoldás. Jó megoldás viszont, a szemléletváltás, vagyis az algoritmusok párhuzamosítása. Adódik a lehetőség: a GPU architektúrájának kihasználása. Ezért ebben a dokumentációban röviden áttekintjük a grafikus kártyák felépítését, megismerkedünk a CUDA környezettel, valamint néhány algoritmust is tanulmányozunk.

Elmélet Számítógépes látás Három szintet különböztetünk meg: 1) Alacsonyszintű képfeldolgozás Alapvető, a képhez szervesen hozzátartozó tulajdonságok meghatározása (élek, szín, textúra, ...). Ezen feldolgozás eredménye általában egy újabb kép, amely az általunk keresett tulajdonságokat tartalmazza (pl. élkép). 2) Középső szint A képen látható objektumok tulajdonságainak meghatározása. A kiindulási adatokat az alacsony szintű feldolgozásból kapjuk. A feldolgozás eredményeként a kép tartalmának egy kezdetleges, szimbolikus leírását kapjuk, amely főleg a képen látható alakzatok jellemzőit foglalja magába (felületek, kontúrvonalak helyzete, mozgás, sztereó...). 3) Felső szint Képi tartalom értelmezése (objektumok felismerése, képi esemény leírása, …)

4


Szegmentálás A képszegmentálás a képfeldolgozás egyik legfontosabb alapproblémája, mely a hasonló tulajdonságú pixelek homogén régiókba történő csoportosításával foglalkozik. A képszegmentálást felületleírásra, alakfelismerésre, képi adatbázis indexelésére, keresésre és még számos más területen alkalmazzák. Máig, az egyik leginkább kutatott részterület, ugyanis a meglévő algoritmusok lassúak, illetve nem létezik abszolút értelemben jó vagy rossz algoritmus, mindig az adott probléma analízise során derül az ki, hogy melyik eljárástól várhatjuk a legjobb eredményt. K-means algoritmus, mely minden egyes pixelt ahhoz a csoporthoz sorol, amelyiknek a középpontja a legközelebb esik az adott pixelhez.

GPU (Graphics Processing Unit) A GPU a jobb gyorsítótár kihasználáson túl erőteljes párhuzamosításra képes, architektúrájának köszönhetően. Számos fix funkció van benne hardveresen megvalósítva, melyek sokkal gyorsabb működésre képesek, mint a szoftveres társaik. Ilyen például a logaritmus, hatványozás, raszterizálás, stb. A GPU a CPU-val ellentétben, nem foglalkozik spekulációval, azaz nem tesz sejtést arra, hogy később milyen adatokra lesz szükség, és nem foglalkozik azok cache-be töltésével, inkább a feldolgozásra fektet nagyobb hangsúlyt.

GT200-as architektúra

5


A GT200-as architektúrára épülő GPU-k tíz darab párhuzamosan elhelyezett Texture Processor Cluster-t (TPC) tartalmaznak, melyekben további három-három darab Streaming Multiprocessor (SM) található, és ezek mindegyike egyenként nyolc darab Streaming Processor-t (SP) tartalmaz. Ez tehát azt jelenti, hogy egy GT200-as architektúrára épülő videokártya összesen kettőszáznegyven darab feldolgozóegységet tartalmaz. Ezen okokból használható a GPU párhuzamos algoritmusok futtatására.

CUDA (Compute Unified Device Architecture) Az Nvidia által, grafikus kártyák programozásához fejlesztett egységes eszközökre épülő számítási architektúra, melyet általános számítási célok megvalósítására terveztek. SIMD architektúráról van szó, minden eszköz egyidejűleg ugyanazt a feladatot hajtja végre, különböző adatokon. A GeForce 8000-esnél nagyobb sorozatszámú videokártyák képesek futtatni a CUDA-t.

Feldolgozási folyamatábra

6


A

feldolgozási

folyamatábrán

megtekinthetjük, hogy CUDA alkalmazásával mi a programvégrehajtás menete. A jobboldali képen pedig a CUDA futási modelljét láthatjuk, ahol a rácsban egyidejűleg egy kernel fut, a szekvenciális program egy párhuzamosítható részlete található itt. A blokk a rács maximum 2D-s felosztása, párhuzamos, de független programelemeket tartalmaz, melyek futhatnak egyszerre, vagy egymás után. A szál a blokk maximum

3D-s

felosztása,

szekvenciális

programrészlet fut rajta. 32 db párhuzamos szál

Futási modell

fonatot alkot, mely egyszerre futtat egy SIMD utasítást. A

CUDA

memóriamodelljén

látszik, hogy minden egyes szálnak saját

regiszterei

statikus

RAM-ok,

vannak,

melyek

tehát

gyorsak,

továbbá mindegyikhez tartozik lokális memória is, de ezek dinamikus RAMok, így lassúak. Minden blokkban van osztott

memória,

melyen keresztül

kommunikálnak a szálak, ez is statikus RAM, tehát gyors. A globális, a konstans, és a textúra memóriát minden szál és rács eléri, ezek dinamikus RAM-ok, de utóbbi kettő ki van

Memóriamodell

egészítve egyirányú gyorsítótárral, így azok viszonylag gyorsak, ellentétben a globális memóriával. CUDA-ra algoritmust áttenni könnyű, mert C nyelven lehet programozni, nem szükséges hozzá párhuzamos programozási nyelv ismerete. Jelentős gyorsítást elérni nagyon nehéz, elengedhetetlen hozzá a GPU felépítésének alapos ismerete, a jó algoritmusválasztás, továbbá a szűk keresztmetszetek feloldása. 7


Megvalósított algoritmusok Az algoritmusok megvalósításához Microsoft Visual Studio 2008-at és CUDA Toolkitet (ingyenesen elérhető az Nvidia weboldalán) használtam, futtatásuk pedig egy GeForce 9600 GT videokártyán történt. Az algoritmusok azonos részekből állnak: egy, a CPU-n futó kódból és egy GPU-n futó kernelből. A CPU-n futó kód szinte ugyanaz minden esetben, a textúra beállításánál találunk csak lényegi különbséget. A processzoron futó kód fájlból olvassa be a memóriába a feldolgozandó képeket, majd átmásolja azokat az eszköz memóriába, paraméterezi a textúrát és meghívja a GPU-n futó kernelt. Itt a textúra blokkjain a feldolgozóegységek párhuzamosan végzik el a szükséges műveleteket, például konvolúció során a mátrix elemeivel való szorzást. A kernel lefutása után a CPU visszamásolja az eredményül kapott képeket a memóriába, majd fájlba menti azokat, így kapjuk meg mind a három esetben az eredményképeket.

CPU-n futó kód int main(int argc, char **argv) { // kép megnyitása float* h_data = NULL; unsigned int width, height; char* image_path = cutFindFilePath(image_filename, argv[0]); cutLoadPGMf(image_path, &h_data, &width, &height); unsigned int size = width * height * sizeof( float ); float* d_data = NULL; cudaMalloc( (void**) &d_data, size); // tömb lefoglalása és az adatok feltöltése cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc( 8, 8, 8, 8, cudaChannelFormatKindUnsigned); cudaArray* cu_array; // textúramemóriába foglaljuk cudaMallocArray( &cu_array, &channelDesc, width, height ); cudaMemcpyToArray( cu_array, 0, 0, h_data, size, cudaMemcpyHostToDevice);

8


// textúra paraméterek beállítása tex.addressMode[0] = cudaAddressModeWrap; tex.addressMode[1] = cudaAddressModeWrap; tex.filterMode = cudaFilterModePoint; // ne legyen átlagolás tex.normalized = false; // a textúra koordináták legyenek valósak cudaBindTextureToArray( tex, cu_array, channelDesc ); dim3 dimBlock( 8, 8, 1 ); dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1 ); // GPU-n futó kernel meghívása kernel<<< dimGrid, dimBlock, 0 >>>( d_data, width, height); // eredmény visszaírása a host memóriába float* h_odata = (float*) malloc( size); cudaMemcpy( h_odata, d_data, size, cudaMemcpyDeviceToHost); // eredmény írása fájlba char output_filename[1024]; strcpy(output_filename, image_path); strcpy(output_filename + strlen(image_path) - 4, "_out.pgm"); cutSavePGMf( output_filename, h_odata, width, height); cudaFreeArray( cu_array ); } Itt a GPU-n futó kerneleket nem részletezem, ugyanis azok az egyes feladatoknál különböznek, részletezésüket lásd később.

9


Szürkeárnyalatosítás Megvalósítása során, a korábban bemutatott kód fut a CPU-n, mely a most részletezésre kerülő GPU-n futó kernelt hívja meg. __global__ void kernel( float* result, int width, int height ) { // textúrakoordináták kiszámítása unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; // szükséges pixelek olvasása a textúrából uchar4 value = tex2D( tex, ( float )x, ( float )y ); // szürkeárnyalatosítás unsigned char pixel = value.x*.3 + value.y*.59 + value.z*.11; // eredmény előállítása result[ y * width + x ] = pixel; }

A szürkeárnyalatosítás eredménye

10


Küszöbözés Megvalósítása során, a korábban bemutatott kód fut a CPU-n, mely a most részletezésre kerülő GPU-n futó kernelt hívja meg. __global__ void kernel( float* result, int width, int height ) { // textúrakoordináták kiszámítása unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; // szükséges pixelek olvasása a textúrából uchar4 value = tex2D( tex, ( float )x, ( float )y ); // küszöbözés if ( value < thresholdValue ) value = 0; else value = 1; // eredmény előállítása result[ y * width + x ] = value.x, value.x, value.x, 1; }

Küszöbözés eredménye

11


Konvolúció Megvalósítása során, a korábban bemutatott kód fut a CPU-n, mely a most részletezésre kerülő GPU-n futó kernelt hívja meg. __global__ void kernel( float* result, int width, int height ) { // textúrakoordináták kiszámítása unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; // szükséges pixelek olvasása a textúrából float pixel0 = tex2D( tex, ( float )x-1, ( float )y-1 ); float pixel1 = tex2D( tex, ( float )x, ( float )y-1 ); . . . float pixel8 = tex2D( tex, ( float )x+1, ( float )y+1 ); // konvolúciós mátrix elemeinek megadása const float m00 = 2.0; const float m01 = -1.0/4.0; const float m10 = -1.0/4.0; const float m11 = 0.0; // az egyes pixelértékek kiszámítása float value0 = pixel0*m11; . . . float value8 = pixel8*m11; // eredmény előállítása result[ y * width + x ] = value0 + value1 + value2 + … + value7 + value8; }

12


Konvolúció eredménye

13


Alakzatok keresése Ebben a feladatban a kép pixeleinek színértékei alapján szegmentáltam a képet, hogy megtaláljam rajta a négy kört. A megvalósítás menete hasonló a korábbiakban bemutatottakhoz, a GPU-n futó kernel a pixelek színértékeit hasonlítja össze a 3D színtérben kijelölt színcsatornánként meghatározott alsó és felső korlát értékeivel, ha a két határ között van, akkor meghagyja a pixelt, ha nem akkor a pixelt fehér színűre állítja. __global__ void kernel( uchar4* result, int width, int height ) { // textúrakoordináták kiszámítása unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; // szükséges pixelek olvasása a textúrából uchar4 value = tex2D( tex, ( float )x, ( float )y ); if (27
Körök keresésének eredménye

14


Jövőbeli tervek Következő félévben tervezem újabb szegmentálási algoritmusok megismerését, kipróbálását, illetve párhuzamosíthatóságuknak tesztelését. További célom az elkészült algoritmusok gyorsítása, optimalizálása, végső fázisban pedig, kamerakép valósidejű feldolgozása.

Felhasznált irodalom Vajda Ferenc PhD, adjunktus – Képfeldolgozás GPU segítségével előadásanyag Vajda Ferenc PhD, adjunktus – Képfeldolgozás CUDA segítségével előadásanyag Nvidia CUDA Reference Manual Nvidia CUDA Programing Guide

15

KÉPFELDOLGOZÓ ALGORITMUSOK FEJLESZTÉSE GRAFIKUS HARDVER KÖRNYEZETBEN

Recommend Documents