C# Fine Matching + CUDA
2012-06-05T10:44:21+02:00
2012-06-21T08:13:38+02:00
2022-08-06T07:15:30+02:00
thezollie
Sziasztok!

Segítségeteket szeretném kérni egy nem hétköznapi problémában.
Képfeldolgozásról lenne szó.
Fine Matching algoritmust írok CUDA vga-ra (konkrétan: GeForce 210 videokártyára).
Algoritmus lényege annyi lenne, hogy 2 képünk van, a nagyobb képben keresem a kisebb képet.
Ezt meg is valósítottam CPU-ra és GPU-ra is, csak annyi velük a probléma, hogy rettenetesen lassúak.
Főleg a GPU-s rész érdekel engem.

Működő programom röviden mit csinál:
- a képet, melyet keresünk a nagyképen, bemásolom a GPU memóriájába
- egy dupla for ciklussal a nagyképből kivágok vagyis klónozok (Clone()) pixelenként egy kisképnyi területet (Rectangle()), majd szépen egyesével továbbküldöm az összehasonlítást végző CUDA-s algoritmusomnak ami pedig histogram-ban tárolja el, hogy melyik területen mennyi egyezőség volt megtalálható. Pl: histogram[47]=1225
Az indexből pedig ki tudom számolni az x, y koordinátát, majd onnan indulva Rectangle() segítségével be tudom keretezni a területet, ahol egyezőség van (kiskép szélességű és magasságú terület).
(Memóriába másolgatás végett lassú alapból az algoritmus)

Ez eddig tökéletesen és lassan működik.
A gyorsításhoz szükséges nekem, hogy a nagykép a GPU memóriájában legyen, persze a kiskép is.

Nah most CUDA alatt szeretnék vagyis próbálok olyan algoritmust írni, mely ugyanazt teszi mint előző, csak minden a GPU-n futna.
Tehát:
- 2 kép a GPU memóriájában
- el kell osztanom a blokkokat és szálakat optimálisan, képmérettől függetlenül.
- byte-osan tárolom a képeket 24 bites formátumban (B,G,R), így pixelenként mehetek rajtuk végig:

int b = large[largeTID] - small[smallTID]; int g = large[largeTID + 1] - small[smallTID + 1]; int r = large[largeTID + 2] - small[smallTID + 2]; if (-treshold < b && b < treshold && -treshold < g && g < treshold && -treshold < r && r < treshold) thread.atomicAdd(ref histogram[histogramY + histogramX], 1);

Ezt a "szerkezetet meghagynám, mert eddig bevált"

A hisztogram mérete kötött (az algoritmust karakterfelismerésnél is használom):
(nagykép szélessége - kiskép szélessége)*(nagykép magassága - kiskép magassága)
Ez azért ennyi, mert:
pl:
150*100 - nagykép
100*50 - kiskép
kezdem nagykép (0,0)-nál, akkor 100*50-es területet néz meg, mely a kisképem mérete
(1,0)-nál is ekkora területet vizsgál, de (101,50)-ig megy, ezért:
(50,50)-nél tovább nem érdemes menni, mert lelógna a kiskép.

GPU-n sokféle megoldásra gondoltam, de az indexeléssel meggyűlt a problémám sajnos.
Úgy gondoltam, hogy a blockIdx-be a nagyképet indexelném. Offset=nagykép pixelszáma lenne, és ez által a kisképet pedig a nagykép utána indexelném: KISKEP_TID=kiskép pixelje + offset
Sajnos nem oldotta meg a problémát


Bízok benne, hogy érthetően fogalmaztam. :)
Köszönöm szépen előre is a segítséget!
Üdv!
Zollie
Mutasd a teljes hozzászólást!
Szia!

Átfutottam a kódodat (most kevés időm van).
A legnagyobb hiba az az, hogy a histogramot globálisan számolod.
Így teljesen mást jelent: valami olyasmit, hogy a kép eltérése globálisan a nagytól, mert minden egyes pozicióban ugyanazt a kinullázatlan kódot növeled (amihez kívülről hozzá sem férsz mindaddig, amíg nem végzett az egész számítás).
(Arról nem is beszélve, hogy ez így sajnos nem túl párhuzamos az atomicAdd miatt).

Miközben neked (ha jól gondolom) arra van szükséged, hogy ha a kis kép mondjuk a 100:100 pozicióban van, akkor hogy néz ki a histogram.
Mutasd a teljes hozzászólást!

  • Ugyan nem az eredeti problémádra fogok válaszolni, de remélem segít, amit mondok:
    Lineáris képfeldolgozás esetén, ha egy kisebb képet keresel egy nagyobban, és a konvolúciós megoldást választod, akkor sokkal jobb, ha Fourier térben csinálod, mert ott a konvolúció a maga O(n^2)-es igénye helyett csak egy szorzás van, így megúszod O(n*log(n))-el, és ezt még egy jó kis GPU se tudja túlszárnyalni. A GPU-t akkor érdemes bevetni, ha már minden polinomiális shortcutot megoldottál, vagy ha a problémád NP-teljes, de jól párhuzamosítható.

    Ajánlom a "The Scientist and Engineer's Guide to
    Digital Signal Processing" könyvet. A te problémád a 24. fejezetben van. (FFT Convolution, 416. oldaltól)
    link:
    http://www.analog.com/static/imported-files/tech_docs/dsp_book_Ch24...
    Ennél is jobb, ha a képekből piramist készítesz, és nem teljes egyezőséget keresel, hanem először a kis képeken dolgozol. Ha elmélyedsz benne, akkor olyan megoldást is le tudsz programozni, ami nemlineáris torzítások esetén is megtalálja a képet. (Például a kép el van forgatva, nyújtva van, stb.)

    Ehhez olvass el egy SIFT leírást. (Találsz a neten CUDA free kódot is.)
    Mutasd a teljes hozzászólást!
  • Szia!

    Köszönöm szépen a tippeket!
    Mindenképpen GPU-n szeretném a problémát megoldani, ráadásul párhuzamosítással.
    GPU-n a legegyszerűbb módszert kell erre alkalmazni, nem szabad elbonyolítani, mert úgy a leggyorsabb.
    Szépen pixelenként rá kell illeszteni a kisképet a nagyképre, majd ebből hisztoigramot csinálok, melyben lineárisan benne van, hogy melyik nagykép-pixelnél mennyi egyező pixel található.
    Paddingra ügyelve csinálom természetesen.
    A cpu-n ha kivágok a nagyképből egy kisképnyi képet és átadom az összehasonlító algoritmusnak, ez működik, viszont, ha GPU-n van a nagykép és azt próbálom beindexelni, az sajnos nem működik és ennek okát keresem. Ad vissza értékeket a histogram, de nem helyeseket.

    Az algoritmusnál a párhuzamosítás működik, legolcsóbb CUDA-s VGA-nál 1 sec alatt meg is találja az egyezőséget, izmosabbnál ez olyan 70 millisecet jelent.

    Csak tényleg nem a helyes értékeket dobja vissza, megeshet hogy az AtomicAdd-al vannak problémák, párhuzamosításnál nem a legmegfelelőbben működik.

    Mindenesetre hasznos amit küldtél, köszönöm szépen!

    Zollie
    Mutasd a teljes hozzászólást!
  • Szia!

    Pár kérdés:
    Melyik memóriába akarod írni az eredményt?
    Milyen compute capability-je van a videókártyádnak?
    Mekkora maximálisan az input kép (a nagy)?
    Honnan jön a histogramX és a histogramY értéke?
    Miért adod össze őket?
    Nem így kéne?
    histogram[histogramY * widthStep + histrogramX]

    Nem lenne jobb az egészet histogram nélkül kiszámolni?
    Minden egyes x,y-nál kiszámolod a kis kép és a nagy kép szeletének a távolságát. Ez lehet a sum of absolute differences vagy a least square distance is.
    Utána ezt az egy értéket tárolod el egy globális tömbben:
    distances[y * widthStep + x] = distance

    Majd pedig egy második pass-ban kiválasztod ebből a tömbből maximumot. (Az indexet egészértékesen elosztod widthStep-el, ekkor megkapod az y-t, majd index % (y * widthStep) és megvan az x is.


    Én így csinálnám az indexelést (ha nem túl nagy a nagy kép):
    int x = threadIdx.x + blockIdx.x * blockDim.x; int y = threadIdx.y + blockIdx.y * blockDim.y;


    Szerk:
    Most átfutottam az első hozzászólásodat.
    Ez a folyamatos kisebb kép kivágása a nagy képből nagyon-nagyon-nagyon lassú!!
    GPU-nál messze a legszűkebb keresztmetszet a memória! Minél kisebb kell, hogy legyen a memóriaforgalom. Egyszer be kell másolni a nagy képet és utána direktben azzal dolgozni (ahogy fent vázoltam).
    Egy régi GPU-n sem szabad, hogy 50-60 ms-nál tovább tartson (főleg, ha jól optimalizálod a dolgokat, pl kisképet konstans memóriába rakod stb).
    Mutasd a teljes hozzászólást!
  • Szia!

    Végre egy hozzáértő ember :)
    Compute Capability: 1.2 GeForce 210-es kártyám van.

    A kódot privát üzenetben mellékelem mindjárt.

    Itt vannak a dimenziók:
    dim3 blockDim = new dim3(small_Bitmap.Width);//maximum 512 lehet dim3 gridDim = new dim3(histogram_width * small_Bitmap.Height, histogram_height); //Max: (X = 65535, Y = 65535)

    blockDim egyértelmű: kiskép szélessége, mely max 512 pixel lehet
    gridDim:
    X dimenzio: historgram szélessége * kiskép magassága
    Y dimenzio: hisztogram magassága

    Tehát:
    Van 2 képem: kicsi és nagy, és van a hisztogram, mely mérete: nagykép szélessége - kiskép szélessége, nagykép magassága - kiskép magassága
    Azért van rá szükség, mert létre fog jönni egy padding rész.
    Képen lerajzoltam hogy kell elképzelni.
    A keresés csak a hisztogram koordinátáin belül történik , mivel a kiskép kinyúlik majd belőle és az algoritmus a padding részekre is lefut így.

    A hisztogrammos AtomicAdd-os megoldás működött szépen, ha CPU-ról croppoltam a nagyképből egy kisképnyit.
    De nagyon lassú volt úgy.
    Ezért döntöttem úgy hogy GPU-n tárolom a nagyképet is és inkább indexelek.

    Amúgy 6 index kell:
    Kiksép x, y ;
    Histogram x, y;

    Nagykép x, y ez megegyezik az előzőek megfelelőjének összegével.
    (Vektor összeadás)

    Mindjárt küldöm privát a kódot :)
    Köszi a segítséget!
    Mutasd a teljes hozzászólást!
    Csatolt állomány
  • Szia!

    Átfutottam a kódodat (most kevés időm van).
    A legnagyobb hiba az az, hogy a histogramot globálisan számolod.
    Így teljesen mást jelent: valami olyasmit, hogy a kép eltérése globálisan a nagytól, mert minden egyes pozicióban ugyanazt a kinullázatlan kódot növeled (amihez kívülről hozzá sem férsz mindaddig, amíg nem végzett az egész számítás).
    (Arról nem is beszélve, hogy ez így sajnos nem túl párhuzamos az atomicAdd miatt).

    Miközben neked (ha jól gondolom) arra van szükséged, hogy ha a kis kép mondjuk a 100:100 pozicióban van, akkor hogy néz ki a histogram.
    Mutasd a teljes hozzászólást!
  • Szia!

    Köszi szépen a tippet, már én is gyanakodtam rá, nagy eséllyel az a ludas :/
    Megpróbálom orvosolni a hibát.
    Köszönettel!

    Zollie
    Mutasd a teljes hozzászólást!
abcd