| Zobrazit předchozí téma :: Zobrazit následující téma |
| Autor |
Zpráva |
pcmaster
Založen: 28. 07. 2007
Příspěvky: 1827
|
 Zaslal: 10. únor 2010, 10:01:43 Předmět: OpenCL na AMD |
 |
|
Zdar,
v diplomke budem pisat tak trochu aj ray-tracer v OpenCL a k dispozicii mam len ATI HW - Radeon HD 4670. Ako vsetci tak nejako vieme, s ATI drivermi to nie je nijaka slava. Hned ako sa mi podarilo napisat male hello world pokusy s OpenCL vo Visual Studio 2008 C++, chcel som skusit aj OpenGL interop, ale tam to nejako nefacha. Co je divnejsie, neskompiluje mi ani takyto kernel (bez OpenCL from OpenGL kontextu, uplne normalne) (cela sprava z clGetProgramBuildInfo "Link Failed"), takze mam podozrenie, ze je nieco zle:
| kód: |
__kernel void whatever(__write_only image2d_t, const int w, const int h) {
.. absolutne cokolvek, i nic!
}
|
 
Mam podozrenie, ze by to mohlo suvisiet s tym 2. screenshotom (Image 2D size 0x0). To ze je tam Athlon si nevsimajte, Radeon vyzera podobne (akurat mi da len 128 MB globalne, 16 kB lokalne, 128x128x128 work sizes a 8 compute units) - zabudol som to prepnut v tom liste  Velmi mile od ATI je, ze v SDK SAMPLES nie je nic na image2d_t kernely a niekde som sa docital, ze to dokonca zatial mooozno ani nepodporuju
Cely kod sem davat zatial nejdem (vsak tam dokopy nic nie je), rad by som bol, ak by sa nasiel niekto, kto ma tiez ATI HW a urobil aspon screenshot z GPU caps. To plati aj pre stastnejsich nVidia computistov
Este spomeniem, ze mam Windows 2003 R2 SP2 x64 a 9.12 drivery (ale skusal som aj najnovsie 10.1 aj 9.11 a rozne "specialne" beta a fixy pre SDK v 2.0) a vsade je to podobne. Ze sa neda nainstalovat CCC z dovodu rozjebaneho .NETu radsej nejdem ani spominat 
_________________
Off-topic flame-war addict since the very beginning. Registered since Oct. 2003!
Interproductum fimi omne est. |
|
| Návrat nahoru |
|
 |
nou
Založen: 28. 07. 2007
Příspěvky: 1050
|
| Zaslal: 10. únor 2010, 10:24:24 Předmět: |
|
|
http://nopaste.ceske-hry.cz/222887
je to SDL/GLX takze treba vymenit glx funkcie za wgl ekvivalenty. potom aj CL_GLX_DISPLAY_KHR za CL_WGL_HDC_KHR
a AMD teraz nepodporuje textury v OpenCL takze sa neda pouzit image2d.
za dalsie CL/GL interoperabilita je v stadiu preview takze to nemusi uplne dobre behat.
_________________
Najjednoduchšie chyby sa najtažšie hľadajú. |
|
| Návrat nahoru |
|
|
pcmaster
Založen: 28. 07. 2007
Příspěvky: 1827
|
| Zaslal: 10. únor 2010, 18:49:05 Předmět: |
|
|
Dik nou, aj som sa bal, ze to bude tak. Podla tvojho kodu mi to beha, to je ok.
Teraz ma uz len zaujima, ako najlepsie dostat obrazok, ktory spocitam v OpenCL (cize napriklad 800x600x4 floaty ci ubyte) do OpenGL textury. Ako na to z toho ARRAY_BUFFERu, asi tazko, ze? Tak, aby som to nemusel kopirovat na host (cpu). EDIT: zeby glCopyTexImage2D?
_________________
Off-topic flame-war addict since the very beginning. Registered since Oct. 2003!
Interproductum fimi omne est. |
|
| Návrat nahoru |
|
|
nou
Založen: 28. 07. 2007
Příspěvky: 1050
|
| Zaslal: 10. únor 2010, 19:21:54 Předmět: |
|
|
no teoreticky by mohlo ist toto. mal by sa k tomu vyjadrit Eosie ze ako co najrychlejsie spravt z bezneho buffera texturu.
| kód: |
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, CLbuffer);
glBindTexture2D(GL_TEXTURE_2D, textura);
glTexImage2D(..... , NULL);//ak je data NULL a je nieco bindnute do GL_PIXEL_UNPACK_BUFFER tak sa cita z neho => format bude musiet byt GL_FLOAT
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); |
ak ti ide o to vykreslit vysledny buffer tak by slo aj pouzit metodu 1 pixel - 1 bod. ale to by asi zabilo vykon este viac nez to kopirovat cez hosta. aj ked v jednom projekte to idem pouzit ale tam to bude 1 bod - 5x5 pixlov.
alebo pockat kym ATI zacne podporovat textury. podla vyjadreni na AMD fore je to pre nich dolezita feature takze v dalsom vydani sa toho asi dockame. otazne je ci to daju na 4xxx radu. takze otazka ma tvoja 4670 nejake extensions? pretoze moja 5850 podporuje napr. atomicke operacie.
EDIT: chybicka posledny bind ma byt 0
_________________
Najjednoduchšie chyby sa najtažšie hľadajú.
Naposledy upravil nou dne 11. únor 2010, 16:13:34, celkově upraveno 1 krát |
|
| Návrat nahoru |
|
|
pcmaster
Založen: 28. 07. 2007
Příspěvky: 1827
|
| Zaslal: 10. únor 2010, 19:29:09 Předmět: |
|
|
4670 podla GPU Caps vieweru nepodporuje nijake extensions Co je celkom fail, kedze lacnejsie a starsie karty od nVidie podporuju toho kopec (afaik) Skusim to s tou texturou a dam vediet.
EDIT: Facha to a celkom dobre. Urcite nebude hrdlo pri vyrobe tej textury z bufferu na GPU bez hosta
Tak, dalsia vec:
| kód: |
| clGetKernelWorkGroupInfo(hKernel, hDevice, CL_KERNEL_WORK_GROUP_SIZE, sizeof(cl_uint), &localWorkGroupSize, NULL) |
podla helpu vracia jedine cislo (sice 128), no ale do clEnqueueNDRangeKernel() potrebujem 2 (3) cisla ako local work size. Alebo to nema suvis? Ako uhadnem spravne local work group sizes pre 2D a 3D problemy? Ked dam privela (napriklad 16x16), clEnqueueNDRangeKernel vrati chybu (aj ked podporujem udajne 128x128x128). Alebo si to zle interpretujem?
A hned este dalej, budem potrebovat vyrobit aj buffer z textury, teda opacne. To bude?
| kód: |
glBindBuffer(GL_PIXEL_PACK_BUFFER);
glBindTexture(GL_TEXTURE_2D, tex);
glGetTexImage(..., NULL); |
_________________
Off-topic flame-war addict since the very beginning. Registered since Oct. 2003!
Interproductum fimi omne est. |
|
| Návrat nahoru |
|
|
Marek
Založen: 28. 07. 2007
Příspěvky: 1782
Bydliště: Velká Morava
|
| Zaslal: 11. únor 2010, 23:54:30 Předmět: |
|
|
Ten přesun do bufferu z textury a naopak přes ty PACK/UNPACK buffery je asi nejrychlejší způsob.
Je možno přímo přetypovat buffer na texturu (v 3D API nevím, jestli to jde, ale hw to umí), nicméně pro efektivní čtení ta textura musí mít hardwarově-specifické uspořádání pixelů v paměti a ta kopie mezi bufferem a texturou se už postará o přeuspořádání dat pro texturu. (to je vlastní jediný důvod, proč to API je tak navržené)
_________________
AMD Open Source Graphics Driver Developer |
|
| Návrat nahoru |
|
|
nou
Založen: 28. 07. 2007
Příspěvky: 1050
|
| Zaslal: 12. únor 2010, 09:18:58 Předmět: |
|
|
CL_DEVICE_MAX_WORK_ITEM_SIZES co ti vrati 128x128x128 je teoreticke maxmimum ktore je dalej obmedzene s CL_DEVICE_MAX_WORK_GROUP_SIZE ktore je na tvojej karte taktiez 128. to znamena ze v jednej workgroup moze byt maximalne 128 prvkov a maximalne dimenzie skupiny su 128x128x128. lenze sucin rozmerov skupiny musi byt mensi alebo rovny velkosti skupiny. teda mozne skupiny su 128x1x1 64x2 8x8x1 16x8x1 8x4x4.
na skompilovany kernel musis zavolat CL_KERNEL_WORK_GROUP_SIZE pretoze podla toho kolko pouzivas registrov sa ti moze stat ti celkova velkost skupiny skoci len 64 alebo 32.
NDRange global_offset ignoruj to musi byt NULL, globa_work_size davas pocet prvkov ktore chces spracovat. musi byt delitelny s jednotlivymi rozmermi v local_work_size. sucin vsetkych rozmerov local_work_size musi byt <= CL_KERNEL_WORK_GROUP_SIZE <= CL_DEVICE_MAX_WORK_GROUP_SIZE
_________________
Najjednoduchšie chyby sa najtažšie hľadajú. |
|
| Návrat nahoru |
|
|
pcmaster
Založen: 28. 07. 2007
Příspěvky: 1827
|
| Zaslal: 12. únor 2010, 10:35:57 Předmět: |
|
|
Dik, vyborne, takto si to uz aj ja interpretujem. Az na to prve CL_DEVICE_MAX_WORK_ITEM_SIZES, to ma nenapadlo zistovat, asi som to v tej 'prehladnej' prirucke prehliadol 
Teraz idem este zistovat, preco moj "raytracer" (to je velmi silne slovo ) chce fachat len do asi 200x200 px vystupneho obrazku a s vyssim rozlisenim hadze Integer division by zero exception, ktoru VS chyti... To be continued! 
_________________
Off-topic flame-war addict since the very beginning. Registered since Oct. 2003!
Interproductum fimi omne est. |
|
| Návrat nahoru |
|
|
|
Hlavní strana
FAQ
Test
Hledat
Seznam uživatelů
Uživatelské skupiny
Registrace
Ke Stažení
Nastavení
Přihlášení
Calendar
