Opencl - képmanipulálás

Opencl - képmanipulálás
2013-08-26T14:24:28+02:00
2013-08-28T22:21:34+02:00
2022-11-29T16:00:41+01:00
l0rika
Sziasztok

Hobbi szinten elkezdtem foglalkozni az Opencl-el, azonban vannak részek amelyek némi magyarázatra szorulnak.

Képfeldolgozás céljából szeretném megtanulni, ezért első programomban egy képet szeretnék módosítani. Jelenleg a kép egy 4x4-es mátrix, amit én adtam meg példaként. Kezdésnek csak annyit akarok, hogy a kép minden pixeléhez adjon hozzá 5-öt és tegye bele egy bufferbe.

A kernel-em:


const sampler_t s_nearest = CLK_NORMALIZED_COORDS_FALSE| CLK_ADDRESS_NONE | CLK_FILTER_NEAREST; __kernel void image_test(__read_only image2d_t im, __global float4* out) { int2 pos = {get_global_id(0), get_global_id(1)}; int i = get_global_id(2) ; out[i] = read_imagef(im, s_nearest, pos ) + 5; }

A kérdésem az lenne, hogy hol adom meg, hogy a kernel hány példányba hajtsa végre? Ugyanakkor hol adom meg, hogy a kernelek között mi legyen a különbség. Például, hogy mindegyik példányban a pos.x és pos.y változzon, hasonlóan az i változó is növekedjen.
Mutasd a teljes hozzászólást!
"Esetleg te hogyan írnád meg a kernelt?"

Mivel én meg már eleve öreg vagyok ahhoz, hogy a nyers opencl api-val kínlodjak, ezert a leheto legegyszerubben oldanam meg: a sajat ocl wrapper-emben, haha.

var code:=' __kernel void main(__global float* buf) { buf[get_global_id(0)]+=5.0; }'; const WorkCount:=16; var dev:=cl.devices[0]; //select device var kernel:=dev.NewKernel(code); //compile kernel procedure DumpBuf; //matrix kiiratasa begin for var i:=0 to WorkCount-1 do begin write(format('%6.2f',buf.floats[i]),(i and 3)=3 ? #13#10 : ''); end;end; var buf:=dev.NewBuffer('rw',WorkCount*4{in bytes}); //initialize test data for var i:=0 to WorkCount-1 do buf.floats[i]:=i*0.3; writeln('Kiindulo matrix'); DumpBuf; with kernel.run(WorkCount,buf)do begin waitfor;free;end; //futtatas, 1D kernel writeln('Eredmeny matrix'); DumpBuf; //eredmeny kiirasa buf.free; //free kernel.free; --Output------------------------------------------------------ Kiindulo matrix 0.00 0.30 0.60 0.90 1.20 1.50 1.80 2.10 2.40 2.70 3.00 3.30 3.60 3.90 4.20 4.50 Eredmeny matrix 5.00 5.30 5.60 5.90 6.20 6.50 6.80 7.10 7.40 7.70 8.00 8.30 8.60 8.90 9.20 9.50

Amit egyszerusitettem a feladat miatt:
- image helyett sima clCreateBuffer()
- 2 parameter helyett 1 RW parameter. A kernel felulirja az eredetit.
- csak 1D domain (0..15 range)
- a kernel egyszeru, mint a faék, azaz nem bonyolultabb, mint a feladat.
- az image-kel viszont sztem raersz kesobb is foglalkozni, mivel azt letrehozni nem olyan egyszeru, mint a buffert (esetunkben az 1D buffert). Nameg ehhez a matrix-skalar osszeadashoz, nem hiszem hogy szukseg lenne hardveres anizotrop textura-mintavetelre. :D

Egy fontos dolog az lenne, hogy amikor te csinalod ezt az ocl api-ban, akkor minden cl() hivas utan ellenorizd a result-ot. Altalaban mindig szol, ha baja van.

"de nem szeretném az egészet kiolvasni, mert addig öreg leszek." -> No gain without pain.
Mutasd a teljes hozzászólást!

  • Hi,

    Itt adod meg: clEnqueueNDRangeKernel()
    Ezen belul a work_dim, global_work_offset, global/local work_size parameterekkel.

    A peldadban minden workitem, ami a 3. dimenziojat tekintve ugyanott van, az ugyanabba az out-be fog beleirni. Tehat out-ben VALAMELYIK pixel+5 fog szerepelni.

    Ha image-be akarsz visszairni, akkor -> write_imagef is van.

    A 3. domain dimenzio (get_gid(2)) meg eleg erdekesen van hasznalva ebben a peldaban.

    Az osszeadasnal meg nem dob hibat az opencl fordito, hogy vektort akarsz skalarral osszeadni?
    Mutasd a teljes hozzászólást!
  • Hali,

    A clEnqueueNDRangeKernel() függvényt a következőképpen adtam meg. A
    work_items
    azért 16 mert a képem mérete 4x4-es.

    const size_t work_items = 16; ret = clEnqueueNDRangeKernel(command_queue, kernel, 2, &work_items, NULL, NULL, 0, NULL, NULL);


    Továbbá a,

    read_imagef(im, s_nearest, pos )

    eredménye nem egy szám? Az egyik pixel értéke. A pos tartalmazza a koordinátáját a pixelnek. Az nem akar változni a kernel különböző példányain belül. Mindig 0,0 ha nem adok kezdeti feltételt. Na de nem tapasztaltam olyan problémát, hogy az vektor lenne.
    Mutasd a teljes hozzászólást!
  • Azt mondod neki, hogy 2 dimenzios legyen a domain. De csak az elso dimenzio meretet hatarozod meg (16), ekkor a masodik dimenzio memoriaszemet lesz, es biztosan behal.

    const size_t work_items[2]={4,4}

    Az image az tuti, float4-et ad vissza, mert az komplett texturamintavetelezest csinal. Akkor gondolom az ocl kezeli a vektor+skalar osszeadast.

    Van egy rakas tutorial amugy, valassz egyet es csinald vegig! Azt pl. jó lenne érteni, hogy mi az az NDRange. Az ilyeneket abrakkal szepen el szoktak magyarazni.
    Mutasd a teljes hozzászólást!
  • Olvastam egy pár leírást, de az vagy csupán elméleti szöveg (kódok nélkül) volt, vagy csak kódleírás volt. Amit találtam, hogy volt benne kódrészlet elmagyarázva, az nem képekkel foglalkozott. Most megint nekiestem az elméleti résznek. Mostmár tisztában vagyok vele, hogy mi a work-grop, a work-item és az ilyen finomságok.

    Most az a kérdésem lenne, hogy a host-on belül hol adom meg a fenti példával élve, hogy az i, a pos.x és a pos.y változó hogyan változzon kernel-példányonként?
    Pl: Van egy 4x4-es képem. Az azt jelenti, hogy a kernel 16x kell majd végrehajtódjon. Azt hol adom meg, hogy az elsőben az i is a pos.x is és a pos.y is nulla legyen, a második esetben az i=1, a pos.x=0 és a pos.y=1, a harmadiknál az i=2, a pos.x=0 és a pos.y=2 és így tovább? Gondolom, hogy ezt a host-ban kell valahol megadni.
    Valószínűleg ez is benne van a könyvben amit olvasok (385 oldal), de nem szeretném az egészet kiolvasni, mert addig öreg leszek.
    Mutasd a teljes hozzászólást!
  • Esetleg te hogyan írnád meg a kernelt?
    Mutasd a teljes hozzászólást!
  • "Esetleg te hogyan írnád meg a kernelt?"

    Mivel én meg már eleve öreg vagyok ahhoz, hogy a nyers opencl api-val kínlodjak, ezert a leheto legegyszerubben oldanam meg: a sajat ocl wrapper-emben, haha.

    var code:=' __kernel void main(__global float* buf) { buf[get_global_id(0)]+=5.0; }'; const WorkCount:=16; var dev:=cl.devices[0]; //select device var kernel:=dev.NewKernel(code); //compile kernel procedure DumpBuf; //matrix kiiratasa begin for var i:=0 to WorkCount-1 do begin write(format('%6.2f',buf.floats[i]),(i and 3)=3 ? #13#10 : ''); end;end; var buf:=dev.NewBuffer('rw',WorkCount*4{in bytes}); //initialize test data for var i:=0 to WorkCount-1 do buf.floats[i]:=i*0.3; writeln('Kiindulo matrix'); DumpBuf; with kernel.run(WorkCount,buf)do begin waitfor;free;end; //futtatas, 1D kernel writeln('Eredmeny matrix'); DumpBuf; //eredmeny kiirasa buf.free; //free kernel.free; --Output------------------------------------------------------ Kiindulo matrix 0.00 0.30 0.60 0.90 1.20 1.50 1.80 2.10 2.40 2.70 3.00 3.30 3.60 3.90 4.20 4.50 Eredmeny matrix 5.00 5.30 5.60 5.90 6.20 6.50 6.80 7.10 7.40 7.70 8.00 8.30 8.60 8.90 9.20 9.50

    Amit egyszerusitettem a feladat miatt:
    - image helyett sima clCreateBuffer()
    - 2 parameter helyett 1 RW parameter. A kernel felulirja az eredetit.
    - csak 1D domain (0..15 range)
    - a kernel egyszeru, mint a faék, azaz nem bonyolultabb, mint a feladat.
    - az image-kel viszont sztem raersz kesobb is foglalkozni, mivel azt letrehozni nem olyan egyszeru, mint a buffert (esetunkben az 1D buffert). Nameg ehhez a matrix-skalar osszeadashoz, nem hiszem hogy szukseg lenne hardveres anizotrop textura-mintavetelre. :D

    Egy fontos dolog az lenne, hogy amikor te csinalod ezt az ocl api-ban, akkor minden cl() hivas utan ellenorizd a result-ot. Altalaban mindig szol, ha baja van.

    "de nem szeretném az egészet kiolvasni, mert addig öreg leszek." -> No gain without pain.
    Mutasd a teljes hozzászólást!
  • Hi,

    Hát igen...te "valamivel" jobban vágod ezeket.

    A bufferek tényleg egyszerűbbek, azonban, mint írtam képfeldolgozás céljából lenne szükségem erre, ezért jó lenne megtanulni az image-eket...úgy olvastam vhol, hogy így gyorsabb.

    Esetleg nem tudsz egy olyan leírást az ocl-ről amely párhuzamosan tárgyalja az elméletet a gyakorlattal? Amit jelenleg olvasok az túl nyers.

    Ha idegesít a sok kérdés, akkor szólj, vagy ne is válaszolj.
    Mutasd a teljes hozzászólást!
  • De ha egyszer van egy konyved, ami reszletesen targyalja az egeszet az a legjobb, ha vegigolvasod, aztan feluletesen legalabb tudod, hogy mi van.
    Amugy meg -> google opencl tutorial
    Rogton a 3. talalat egy video eloadas, ha nem gond a szakmai angol.

    En Delphizek, szoval a ocl.dll header sem volt egyszeru dolog, de mire azt osszeraktam, legalabb egy kicsit meg is ertettem...

    "Amit jelenleg olvasok az túl nyers."
    Márpedig egyet csinalj vegig :D

    Az opencl image az joval tagabb fogalom, mint egy altalanossagban vett bitmap image. Annak a textura mintavetelezesnel van jelentosege.

    image: A jelen helyzetben, amikor nearest texturamintavetelezes van, az image csak 4 byte-ot (rgba) olvas be annyi ido alatt, mint amennyi ido alatt a buffer 16-ot. (Nem a nearest mintavetelezesre lett optimizalva, na :D)
    Ha anizotrop texture filteringet akarsz, akkor persze erdemesebb a hardveres image-t hasznalni, de amugy meg...
    Mutasd a teljes hozzászólást!
Tetszett amit olvastál? Szeretnél a jövőben is értesülni a hasonló érdekességekről?
abcd