[OpenCL] Wie is er al mee bezig?

[Iska]

Hey allemaal,

Eind december heeft AMD ATI Stream v2.0 uitgegeven met OpenCL (dat overigens als sinds eind 2008 bestaat). Tegelijkertijd werd er toen een hotfix uitgebracht voor FireStream en Radeon kaarten. Maar ondertussen zit de OpenCL ondersteuning standaard bij Catalyst.
OpenCL is ontwikkeld door een aantal bedrijven waaronder Apple en Khronos met als doel parallelle berekeningen op de CPU en de GPU te standaardiseren. Hierdoor kunnen bijvoorbeeld complexe wiskundige berekeningen voortaan op de grafische kaart worden uitgevoerd in plaats van de CPU, aangezien een grafische kaart toch al snel 20 tot 80 keer sneller is dan een CPU.

Nu is het qua tutorials en manuals nog helemaal niks rondom OpenCL, maar ik vroeg mij toch af of hier al mensen zijn die het al geprobeerd hebben en er wat werkends mee hebben kunnen maken.

Ik zou graag zelf willen aftrappen, maar ik ben nog niet veel verder dan de HelloCL
Ik ben er nu achter hoe OpenCL precies te werk gaat en dat er een werkende template zat bij de samples. Heb nu de GPU en de CPU dmv de taylorreeks van de arctan pi uit laten rekenen met het volgende resultaat:


C2D e8400@3,6ghz@1 core tegenover Radeon 4850@700mhz@60% load = factor 30 verschil!!!

[ Voor 22% gewijzigd door Iska op 12-01-2010 13:07 ]

[H!GHGuY]

Ik heb de API al eens zitten bekijken. Maar er echt iets mee gedaan is iets anders. Probleem is dat ik niet direct een toepassing kan bedenken waar ik iets aan heb...

[Haan]

H!GHGuY schreef op dinsdag 12 januari 2010 @ 12:42:
Probleem is dat ik niet direct een toepassing kan bedenken waar ik iets aan heb...

Dat was precies mijn gedachte, dat je sneller Pi kan berekenen is natuurlijk leuk voor benchmarks, maar nuttig, nee

Ik neem aan dat dit vooral voor wetenschappers interessant is, om complexe modellen of simulaties sneller te kunnen berekenen.

[Iska]

Ik studeer zelf Natuurkunde @ TU Delft dus ik kan zat toepassingen bedenken want soms word je echt gek van zo'n CPU...

[Bob]

Ik ga binnen dit en een maand werken met een GTX280, met als toepassing waarschijnlijk allerlei wavelet transformaties. OpenCL heb ik eens oppervlakkig bekeken, spreekt me wel aan omdat het open is. Maar ik zal nog moeten uitzoeken of de Tesla kaarten (waarmee ik uiteindelijk ook zou willen werken) daarmee volledig benut kunnen worden, want de mogelijkheden daarop zien er heeeel interessant uit (oa globale address space).

[ Voor 5% gewijzigd door Bob op 12-01-2010 13:15 ]

[Spinal]

De mensen van Dolphin (Gamecube/Wii emulator) gebruiken OpenCL, project is opensource dus dat zou je eens kunnen bekijken om te zien wat zij er mee doen.

[NC83]

Kun je niet hetzelfde doen wat DirectCompute toe laat dus bijvoorbeeld SSAO berekenen. Dat zijn de dingen waarvoor microsoft zegt dat DirectCompute geschickt is. Ik weet dat deze berekeningen ook in gewone shaders kunnen.

[Iska]

Een nieuwe update.. OpenCL zal binnenkort komen met standaard double drecision floating-points, maar de experimentele versie is al beschikbaar. Deze kan geactiveerd worden dmv de volgende regel:

code:

#pragma OPENCL EXTENSION cl_khr_fp64 : enable

Ik heb nu de volgende kernel

code:

#pragma OPENCL EXTENSION cl_khr_fp64 : enable

__kernel void MainKernel(__global double * output)
{
    uint xid = get_global_id(0);
    
    uint   iCount = xid * 100000;
    double dTemp  = 0.0;
    bool   bFlag  = false;
    while (iCount < (xid * 100000 + 100000))
    {
        if (bFlag)
        {
            dTemp += (double) -1 / (2 * iCount + 1);
            bFlag = false;
        }
        else
        {
            dTemp += (double) 1 / (2 * iCount + 1);
            bFlag = true;
        }
        
        iCount++;
    }
    output[xid] = dTemp;
}

vs dit voor de CPU

code:

    dPi = 0.0;
    int iCount  = 0;
    int iCount0 = 0;
    while (iCount0 < 256)
        {
        iCount = iCount0 * 100000;
        while (iCount < (iCount0 * 100000 + 100000))
        {
            dPi += pow((-1.0), iCount) / (2 * iCount + 1);

            iCount++;
        }

        iCount0++;
    }

met het volgende als resultaat (taylor serie met n=256.000.000):

Dit is dus een 4850@700::1020 tegen een E8400@3,6::400

Het nadeel is wel dat deze experimentele versie alleen nog maar +-*/ ondersteund

[Zoijar]

Niet om fanboy te spelen, maar dit draaide natuurlijk al een tijd op nvidia. OpenCL zelf ga ik wel op overstappen, maar op het moment is het nog CUDA bij mij. Dat is wel bijna hetzelfde.

Grappigs dat ik wavelets zag; dat wordt waarschijnlijk ook mijn volgende toepassing.

[Iska]

Je hebt helemaal gelijk, ik vind het eigenlijk ook belachelijk dat ATI nu pas met iets als OpenCL komt. Maar ik ben aangewezen op mijn ATI's dus dan moet je toch wat

[oeLangOetan]

Jasper91 schreef op donderdag 14 januari 2010 @ 21:34:
Je hebt helemaal gelijk, ik vind het eigenlijk ook belachelijk dat ATI nu pas met iets als OpenCL komt. Maar ik ben aangewezen op mijn ATI's dus dan moet je toch wat

Brook+ ? en die kan trouwens de local share van de hd4xxx kaarten gebruiken in tegenstelling tot OpenCL.

[Nick The Heazk]

Jasper91 schreef op donderdag 14 januari 2010 @ 20:13:
Een nieuwe update.. OpenCL zal binnenkort komen met standaard double drecision floating-points, maar de experimentele versie is al beschikbaar. Deze kan geactiveerd worden dmv de volgende regel:

code:

...

Ik heb nu de volgende kernel

code:

#pragma OPENCL EXTENSION cl_khr_fp64 : enable

__kernel void MainKernel(__global double * output)
{
    uint xid = get_global_id(0);
    
    uint   iCount = xid * 100000;
    double dTemp  = 0.0;
    bool   bFlag  = false;
    while (iCount < (xid * 100000 + 100000))
    {
        if (bFlag)
        {
            dTemp += (double) -1 / (2 * iCount + 1);
            bFlag = false;
        }
        else
        {
            dTemp += (double) 1 / (2 * iCount + 1);
            bFlag = true;
        }
        
        iCount++;
    }
    output[xid] = dTemp;
}

Toch even een inhoudelijke opmerking op je codefragment. Ik ben niet erg vertrouwd met de architectuur van GPUs, maar ik betwijfel dat deze uitgerust worden met geavanceerde branch-prediction mechanismen. Ik denk dan ook dat je stukje code heel wat sneller zou lopen als je die if-test elimineert en een nieuwe integer variabele "factor" introduceerd. Deze factor wissel je aan het einde van de lus steeds van teken, door te vermenigvuldigen met -1 (waarvan ik vermoed dat dit op een GPU heel wat efficienter zal zijn dan op branch prediction te steunen). Het zou interessant zijn indien je deze hypothese zou kunnen verifiëren .

Merk overigens op dat je algoritme niet numeriek stabiel is. Je begint best te sommeren bij de hoge-orde termen. De latere iteraties zijn op deze manier namelijk vrij nutteloos (als in: dragen niet meer bij tot de som).

Zelf ben ik geïnteresseerd in deze technologie, voornamelijk omdat het aansluit bij mijn eindwerk. De numerieke technieken die ik in dit kader bestudeer lenen zich zeer goed tot parallelle verwerking (voornamelijk matrix-matrix en matrix-vector bewerkingen op relatief kleine matrices met dimensies kleiner dan pakweg 100x100). Het is niet uitgesloten dat ik later op het jaar nog wat probeer te experimenteren met OpenCL of CUDA.

[ Voor 11% gewijzigd door Nick The Heazk op 14-01-2010 23:23 ]

[Iska]

Nick The Heazk schreef op donderdag 14 januari 2010 @ 23:18:
[...]


Toch even een inhoudelijke opmerking op je codefragment. Ik ben niet erg vertrouwd met de architectuur van GPUs, maar ik betwijfel dat deze uitgerust worden met geavanceerde branch-prediction mechanismen. Ik denk dan ook dat je stukje code heel wat sneller zou lopen als je die if-test elimineert en een nieuwe integer variabele "factor" introduceerd. Deze factor wissel je aan het einde van de lus steeds van teken, door te vermenigvuldigen met -1 (waarvan ik vermoed dat dit op een GPU heel wat efficienter zal zijn dan op branch prediction te steunen). Het zou interessant zijn indien je deze hypothese zou kunnen verifiëren .

Merk overigens op dat je algoritme niet numeriek stabiel is. Je begint best te sommeren bij de hoge-orde termen. De latere iteraties zijn op deze manier namelijk vrij nutteloos (als in: dragen niet meer bij tot de som).

Zelf ben ik geïnteresseerd in deze technologie, voornamelijk omdat het aansluit bij mijn eindwerk. De numerieke technieken die ik in dit kader bestudeer lenen zich zeer goed tot parallelle verwerking (voornamelijk matrix-matrix en matrix-vector bewerkingen op relatief kleine matrices met dimensies kleiner dan pakweg 100x100). Het is niet uitgesloten dat ik later op het jaar nog wat probeer te experimenteren met OpenCL of CUDA.

Ik vond jouw hypothese zeer logisch klinken dus ik heb het geprobeerd met een integer. Het probleem was echter dat als ik een integer gebruikte de doubles niet meer werkten (ook niet na conversie), dus toen ben ik maar overgestapt naar een double als sign:

code:

#pragma OPENCL EXTENSION cl_khr_fp64 : enable

__kernel void MainKernel(__global double * output)
{
    uint xid = get_global_id(0);
    
    uint   iCount = xid * 100000;
    double dSign  = 1;
    double dTemp  = 0.0;
    bool   bFlag  = false;
    while (iCount < (xid * 100000 + 100000))
    {
        dTemp += (double) dSign / (2 * iCount + 1);
        bFlag = false;
        
        dSign = dSign * (-1.0);
        
        iCount++;
    }
    output[xid] = dTemp;
}

Met dit resultaat:


Toch een stuk langzamer...

Verder is het inderdaad instabiel, maar als ik echt pi wilde weten met 10^8 decimalen ging ik wel naar de TU in Delft, dit is dus alleen om het verschil tussen de CPU en de GPU te meten

[Nick The Heazk]

Kun je de berekening ook eens een paar keren herhalen (pakweg 500 keer) en de opgemeten timings uitmiddelen? Een enkele meting is natuurlijk niet betrouwbaar (op de CPU bijvoorbeeld is er een verschil van 10% merkbaar, op de GPU bijna 25%). De code voor de CPU is dezelfde gebleven neem ik aan?

Voor het opmeten van de tijd gebruik je best de wall-clock time; start de timer aan het begin van de loop en stop hem na de 500 iteraties. Je CPU wordt in het geval van de GPU code niet gebruikt natuurlijk (je process is idle), waardoor de clock cycles die je uitleest ook niet worden aangepast.

Het zou kunnen, al acht ik dat onwaarschijnlijk, dat je compiler erin slaagt om de while lus met if-tests om te vormen in twee while-lussen. In de eerste lus tel je alle even bijdragen op en in de tweede lus worden alle oneven termen er vanaf getrokken. Ik vermoed dat zulk een codefragment nog iets performanter kan zijn dan de versie met if-testen in de lus (wederom eliminatie van branch prediction en tov. de andere versie een floating-point bewerking minder).

[ Voor 16% gewijzigd door Nick The Heazk op 14-01-2010 23:52 ]

[NC83]

Jasper91 schreef op donderdag 14 januari 2010 @ 21:34:
Je hebt helemaal gelijk, ik vind het eigenlijk ook belachelijk dat ATI nu pas met iets als OpenCL komt. Maar ik ben aangewezen op mijn ATI's dus dan moet je toch wat

Ooit van CTM(Close To Metal) gehoord dat was een soort gelijke techniek als CUDA alleen dan van ATi, dus ze hadden het wel al alleen was het niet zo bekend en was de taal niet prettig om mee te werken. Alle Open standaards zijn nu eenmaal trager dan commerciele en dat zal ook wel zo blijven.

[Maurits van Baerle]

Toevallig was ik een paar dagen geleden begonnen met dingen te verzamelen om een topic over OpenCL te openen. Dan kan ik net zo goed alles wat ik verzameld had hier neerzetten. Ik had eigenlijk het plan om 'Het grote OpenCL topic' te beginnen met een forse openingspost maar dat was er dus nog niet van gekomen.

Als iemand meer tijd (en vooral ook meer verstand van OpenCL) heeft dan ik, voel je vrij om alles hier onder te gebruiken voor een eventuele mooie openingspost. Ik heb het tenslotte ook maar op het web gevonden.



Benchmarks

• Een paar aardige benchmark resultaten op InsanelyMac
• SiSoftware Sandra 2010 ondersteunt OpenCL benchmarking

Tutorials, demo's en SDK's

• De eerste plek om te kijken als je aan de slag wil met OpenCL is natuurlijk bij Khronos zelf
• AMD's Introductory Tutorial to OpenCL
• MacResearch heeft een aardige tutorial, ook voor niet-wetenschappelijke toepassingen.
• ATI Stream Software Development Kit (SDK) v2.0
• NVIDIA's OpenCL Developer Center
• NVIDIA OpenCL Best Practices Guide Version 1.0
• Aangezien OpenCL eigenlijk een soort doorontwikkeling is op CUDA is het porten van oude CUDA code naar OpenCL niet heel erg complex. AMD over het porten van CUDA code naar OpenCL.
• HardOCP over AMDs developer demo videos: Geeking Out with AMD and OpenCL

[ Voor 13% gewijzigd door Maurits van Baerle op 20-01-2010 23:20 ]

[Juup]

Jasper91 schreef op donderdag 14 januari 2010 @ 20:13:

code:

#pragma OPENCL EXTENSION cl_khr_fp64 : enable

__kernel void MainKernel(__global double * output)
{
    uint xid = get_global_id(0);
    
    uint   iCount = xid * 100000;
    double dTemp  = 0.0;
    bool   bFlag  = false;
    while (iCount < (xid * 100000 + 100000))
    {
        if (bFlag)
        {
            dTemp += (double) -1 / (2 * iCount + 1);
            bFlag = false;
        }
        else
        {
            dTemp += (double) 1 / (2 * iCount + 1);
            bFlag = true;
        }
        
        iCount++;
    }
    output[xid] = dTemp;
}

Jasper91 schreef op donderdag 14 januari 2010 @ 23:35:
Ik vond jouw hypothese zeer logisch klinken dus ik heb het geprobeerd met een integer.
...
Toch een stuk langzamer...

Of je doet

C:

    while (iCount < (xid * 100000 + 100000))
    {
        dTemp += (double) 1 / (2 * iCount + 1) - (1 / (2 * iCount + 2));
        iCount+=2;
    }

Ofzoiets

Edit: waarom eigenlijk die iCount+1?

[ Voor 14% gewijzigd door Juup op 20-01-2010 23:19 ]

[Verwijderd]

Juup schreef op woensdag 20 januari 2010 @ 23:16:

Edit: waarom eigenlijk die iCount+1?

Ik denk om division by zero te voorkomen.

C++:

     uint xid = get_global_id(0);
    
     uint   iCount = xid * 100000;

Als xid toevallig als global_id 0 terugkrijgt, is iCount ook 0. Zonder de +1 zou je dan dit krijgen:

C++:

(double) -1 / (2 * iCount); // == (double) -1 / (2 * 0) == division by zero!

Met de +1 geeft dit gewoon een division door één(2* 0 + 1).

[The Flying Dutchman]

Jasper91 schreef op donderdag 14 januari 2010 @ 21:34:
Je hebt helemaal gelijk, ik vind het eigenlijk ook belachelijk dat ATI nu pas met iets als OpenCL komt. Maar ik ben aangewezen op mijn ATI's dus dan moet je toch wat

Voor Ati is de OpenCL driver ook al maandenlang beschikbaar hoor. Ik denk sinds september of oktober ofzo (voor de AMD cpu's was het zelfs nog eerder).

Zelf wil ik binnenkort snel ook eens aan de slag met OpenCL, vooral in combinatie met OpenGL om bijvoorbeeld transparante objecten perfect te kunnen renderen (daarvoor moet je alle triangles in de objecten sorteren, of misschien valt er met OpenCL iets te bedenken hoe volgorde onafhankelijk zou kunnen). Ook denk ik dat er leuke mogelijkheden zijn op het gebied van berekenen van schaduwen en dergelijk... verder lijkt het me leuk om weer eens met fluid simulations aan de slag te gaan (eerder wel met OpenGL en GLSL gedaan).

[ Voor 42% gewijzigd door The Flying Dutchman op 21-01-2010 09:40 ]

[Zoijar]

The Flying Dutchman schreef op donderdag 21 januari 2010 @ 09:35:
Zelf wil ik binnenkort snel ook eens aan de slag met OpenCL, vooral in combinatie met OpenGL om bijvoorbeeld transparante objecten perfect te kunnen renderen (daarvoor moet je alle triangles in de objecten sorteren, of misschien valt er met OpenCL iets te bedenken hoe volgorde onafhankelijk zou kunnen). Ook denk ik dat er leuke mogelijkheden zijn op het gebied van berekenen van schaduwen en dergelijk...

Daar heb je geen opencl voor nodig. Sterker nog, waarschijnlijk zijn shaders gewoon efficienter. Ik weet niet hoe dat in opencl is geregeld, maar in cuda bijvoorbeeld was de interoperabiliteit tussen cuda en opengl vrij traag; je moet dan steeds buffers heen en weer kopieren. Zie ook: http://portal.acm.org/cit...Type=Proceedings&title=GH

[The Flying Dutchman]

Zoijar schreef op donderdag 21 januari 2010 @ 09:56:
[...]

Daar heb je geen opencl voor nodig. Sterker nog, waarschijnlijk zijn shaders gewoon efficienter. Ik weet niet hoe dat in opencl is geregeld, maar in cuda bijvoorbeeld was de interoperabiliteit tussen cuda en opengl vrij traag; je moet dan steeds buffers heen en weer kopieren. Zie ook: http://portal.acm.org/cit...Type=Proceedings&title=GH

Ik weet dat er een voorbeeld applicatie was om order independent transparent rendering te doen met OpenCL. Ik weet ook dat schaduwen en transparantie met shaders kunnen. Echter OpenGL is nou niet echt de juiste tool voor berekeningen (en daar gaat het om bij transparantie en ook bij schaduwen). Zover ik weet is er wel interoperabiliteit tussen OpenGL en OpenCL. Men wil juist berekeningen uit OpenGL houden, dat is het domein van OpenCL. Hoe goed switchen tussen die twee werkt weet ik niet, maar ik geloof dat ze wel buffers kunnen sharen.

Het zal in eerste instantie echter vooral exploratief zijn wat ik met OpenCL ga doen. Ik wil kijken wat de mogelijkheden ermee zijn en wat nieuwe dingen proberen op gebied van schaduwen bijvoorbeeld (ik heb hier wat vage ideeën over, dingen die in OpenGL heel moeilijk zijn, dus ik ben benieuwd of OpenCL betere mogelijkheden hiervoor biedt).

Even nagezocht, sharing is inderdaad snel zoals ik zei.

quote uit nvidia slide serie (slide 11):

Can efficiently share resources with OpenGL:

• Textures, Buffer Objects and Renderbuffers
• Data is shared, not copied
• OpenCL objects are created from OpenGL objects.

[ Voor 12% gewijzigd door The Flying Dutchman op 21-01-2010 10:10 ]

[.oisyn]

The Flying Dutchman schreef op donderdag 21 januari 2010 @ 10:02:
Echter OpenGL is nou niet echt de juiste tool voor berekeningen (en daar gaat het om bij transparantie en ook bij schaduwen)

Dit vind ik een onzinstatement. Ten eerste kun je het ook omdraaien: OpenCL is niet echt de juiste tool voor rendering (en daar gaat het om bij transparatie en ook bij schaduwen). Ten tweede is er niets mis met berekeningen in shaders, daar zijn ze nou net voor bedoeld. Anders gebruikten we de fixed function pipeline nog.

[The Flying Dutchman]

.oisyn schreef op donderdag 21 januari 2010 @ 10:10:
[...]

Dit vind ik een onzinstatement. Ten eerste kun je het ook omdraaien: OpenCL is niet echt de juiste tool voor rendering (en daar gaat het om bij transparatie en ook bij schaduwen). Ten tweede is er niets mis met berekeningen in shaders, daar zijn ze nou net voor bedoeld. Anders gebruikten we de fixed function pipeline nog.

Laat ik het dan iets genuanceerder zeggen:

OpenGL is voornamelijk voor graphics.
OpenCL is voornamelijk voor berekening.

Bij een aantal dingen die komen kijken bij het renderen van graphics komen veel berekeningen kijken. Dan lijkt het mij heel interessant om te kijken wat OpenCL daar kan betekenen voor de efficientie, het implementatie gemak en wellicht nieuwe methoden om mooiere graphics neer te zetten. Ik verwacht ook dat Direct Compute veel gebruikt zal worden om bepaalde rekentaken in het rendering proces over te nemen. Je moet de combinatie OpenGL-OpenCL dan ook echt zien als Direct3D-DirectCompute.

Maar nogmaals, ik ben nog niet met OpenCL bezig geweest. Ik wil er voornamelijk graag eens naar gaan kijken om de mogelijkheden te ontdekken. Het kan best zijn dat OpenCL minder geschikt is dan ik denk. Ik denk echter dat het nieuwe mogelijkheden biedt. Zeg nou zelf, het is niet echt gemakkelijk om berekeningen te doen in shaders, dat limiteerd de implementatie van vernieuwende technieken. Als dit met OpenCL gemakkelijker is, dan zijn er misschien leuke dingen mee te doen in combinatie met OpenGL.

edit:
Om nog even in te gaan op transparantie en schaduwen en berekeningen in shaders: ik denk dat er voornamelijk een taak voor OpenCL weggelegd is als het gaat om het sorteren van objecten en misschien ook wel het genereren van shadow volumes etc. Het renderen zelf is uiteraard een taak van OpenGL. Het zijn vooral de berekeningen vóór het renderen waar ik denk dat OpenCL van toepassing zou kunnen zijn.

Filmpje gevonden van order independent transparency. Blijkbaar is het met DirectX 11 gedaan (meende me te herinneren dat het OpenCL was... helaas...). Ik ben niet zo bekend met DirectX, dus ik weet niet of dit een pure Direct3D implementatie is, of dat er ook DirectCompute voor nodig is. Echter dit zijn het soort render toepassingen waar ik aan denk waarbij OpenCL geschikt zou kunnen zijn: de stap waarin de pixels gesorteerd worden.
link

[ Voor 23% gewijzigd door The Flying Dutchman op 21-01-2010 10:28 ]

[.oisyn]

The Flying Dutchman schreef op donderdag 21 januari 2010 @ 10:18:
[...]


Laat ik het dan iets genuanceerder zeggen:

OpenGL is voornamelijk voor graphics.
OpenCL is voornamelijk voor berekening.

Laat ik het dan weer nog iets verder nuanceren:
OpenGL is voornamelijk voor graphics, met berekeningen op vertex- en pixelniveau.
OpenCL is voornamelijk voor algemene berekening.

Ik zeg niet dat OpenCL niet kan helpen, het is alleen niet direct handig om berekeningen die je op vertex- en pixelniveau doet in OpenCL te implementeren (en schaduwen zijn bijv. aan de ene kant een depth render met een doodsimpele shader en aan de andere kant een sampling van die shadowmap). Zoals ik het zie is OpenGL een specialisatie van OpenCL op het gebied van rendering. En the right tool for the job is over het algemeen de meest specialistische variant in het gebied van die job.

[Zoijar]

OpenCL is vooral handig als je veel scatter en gather operaties nodig hebt, of een random access data patroon. Physics simulaties doen het goed in opencl en ook transformaties zoals FFTs ed, of sparse linear algebra (ray tracing ...). Dat soort dingen. Een paar data punten in (texture samples), een fixed aantal data punten uit (pixels) is ideaal voor OpenGL.

[The Flying Dutchman]

.oisyn schreef op donderdag 21 januari 2010 @ 10:28:
[...]

Laat ik het dan weer nog iets verder nuanceren:
OpenGL is voornamelijk voor graphics, met berekeningen op vertex- en pixelniveau.
OpenCL is voornamelijk voor algemene berekening.

Ik zeg niet dat OpenCL niet kan helpen, het is alleen niet direct handig om berekeningen die je op vertex- en pixelniveau doet in OpenCL te implementeren (en schaduwen zijn bijv. aan de ene kant een depth render met een doodsimpele shader en aan de andere kant een sampling van die shadowmap). Zoals ik het zie is OpenGL een specialisatie van OpenCL op het gebied van rendering. En the right tool for the job is over het algemeen de meest specialistische variant in het gebied van die job.

Ik denk dat we het dan aardig met elkaar eens zijn. Wat betreft schaduwen: ik was niet direct aan het overwegen om traditionele methoden voor het berekenen van schaduwen met OpenCL te doen. Wat me vooral leuk lijkt is om eens te kijken of er met behulp van OpenCL mogelijkheden zijn om in schaduwen ook umbra's/penumbra's te implementeren (but then again: wellicht kun je met geometry shaders net zo goed een umbra volume genereren oid, een soort extended shadow volumes dus).

[Maurits van Baerle]

De nieuwe ATI Catalyst 10.1 heeft een nieuwe OpenCL driver, kennelijk...

[ Voor 16% gewijzigd door Maurits van Baerle op 27-01-2010 22:29 ]

[PrisonerOfPain]

offtopic:
Overigens is het resultaat dat je bij Pi neerzet fout. De laatste 3 gerapporteerde cijfers horen niet 145 te zijn maar 535

@Zoijar, de huidige technieken voor GPU raytracing zijn vziw nog steeds bandwidth-limited. Parallel tracen word gedaan dmv packet tracing.

[ Voor 32% gewijzigd door PrisonerOfPain op 27-01-2010 23:15 ]

[Maurits van Baerle]

Twee nieuwtjes:

ATI heeft een update van de Stream SDK uitgebracht, versie 2.01.

What’s New in v2.01
• Update release for ATI Stream SDK v2.0.
• Support for Red Hat® Enterprise Linux® 5.3.
• Support for debugging OpenCL™ with GDB on x86 CPUs under Linux® (see application note for more details).
• Preview: Support for OpenCL™ / Microsoft® DirectX® 9 interoperability. Please see this knowledge base article for more information about this preview feature.
• Additional OpenCL™ samples:
• BoxFilter
• FFT
• GaussianNoise (under cpp_cl)
• URNG
• Stream KernelAnalyzer with OpenCL™ support (available for download separately from Stream KernelAnalyzer Product Page).
• Various OpenCL™ compiler and runtime fixes and enhancements (see developer release notes for more details).
• Support for ATI Radeon™ HD 5670 GPU and ATI Radeon™ HD 5570 GPU.

Verder heeft GraphicRemedy de gDEBugger for OpenCL aangekondigd die later dit jaar uit moet komen. Er is wel al een beta program waar je je voor aan kunt melden.

Graphic Remedy is proud to announce the upcoming release of gDEBugger for OpenCL™. This new product will bring gDEBugger's advanced Debugging, Profiling and Memory Analysis abilities to the OpenCL developer's world, helping OpenCL developers find bugs, optimize parallel computing application performance and memory consumption.

[Maurits van Baerle]

Swan, een (GPL2) tool om bestaande CUDA code te converteren naar OpenCL.

What is it?

Swan is a small tool that aids the reversible conversion of existing CUDA codebases to OpenCL. It does several useful things:

• Translates CUDA kernel source-code to OpenCL.
• Provides a common API that abstracts both CUDA and OpenCL runtimes.
• Preserves the convenience of the CUDA <<< grid, block >>> kernel launch syntax by generating C source-code for kernel entry-point functions.

It can also be usefully used for compiling and managing kernels written directly for OpenCL.

Why might you want it?

Possible uses include:

• Evaluating OpenCL performance of an existing CUDA code.
• Maintaining a dual-target OpenCL and CUDA code.
• Reducing dependence on NVCC when compiling host code.
• Support multiple CUDA compute capabilities in a single binary
• A runtime library for managing OpenCL kernels for new development

http://www.multiscalelab.org/swan

[Maurits van Baerle]

Khronos heeft deze week OpenCL 1.1 gelanceerd.

The Khronos™ Group today announced the ratification and public release of the OpenCL™ 1.1 specification, the latest version of the open, royalty-free standard for cross-platform, parallel programming of modern processors. OpenCL 1.1 provides enhanced performance and functionality for parallel programming in a backwards compatible specification that is the result of cooperation between industry-leading companies.

Today Khronos also announced the release of a C++ wrapper API for use with OpenCL, and the immediate availability of OpenCL 1.1 conformance tests. The OpenCL 1.1 specifications, online reference pages and reference cards are available at www.khronos.org/opencl.

Khronos Drives Momentum of Parallel Computing Standard with Release of OpenCL 1.1 Specification

Ars Technica:

OpenCL 1.1 enhances performance with backward compatibility
The Khronos Group, the organization responsible for maintaining a number of open standards such as OpenGL, WebGL, and OpenMAX, has announced the release of OpenCL 1.1, an update to the cross-platform standard for making it easier to target GPUs and other specialized hardware for parallel processing. The latest update includes a number of performance improvements and added instructions while maintaining full backward compatibility for OpenCL 1.0. The release also establishes an 18-month cadence for updates to the standard.

The original OpenCL 1.0 specification was released just six months after Apple announced that it planned to spearhead the creation of the Open Compute Language standard and include support for it system-wide in Snow Leopard. The 1.1 update was timed to give those working on OpenCL 1.0 compatibility time to deliver while still being responsive to developer feedback. "Any real standard is a living thing—it has to grow and respond to feedback and developer requirements," Neil Trevett, OpenCL working group chairman, Khronos Group president, and NVIDIA vice president of mobile content, told Ars. "We think 18 months is good balance."

The updated version adds additional built-in data types and C functions for increased flexibility and improved compatibility between OpenCL and OpenGL pipelines, especially useful for many scientific visualization and gaming tasks. Support for coordinating commands from multiple hosts makes it possible to spread workloads to resources available in computing clusters. And while OpenCL is C-based, OpenCL 1.1 includes a new C++ wrapper API.

Even with these additions, though, developers won't have to worry about existing code breaking. "Any 1.0-compliant program will continue to run with strict backwards compatibility," Trevett said. "So, there's extra goodness with little disruption."

The OpenCL model makes it possible for developers to poll a system for available OpenCL-compliant resources, and allows developers to create small processing functions called kernels to send to those resources. Operating at a low level, OpenCL provides a generalized model for managing memory and running instructions on CPUs, GPUs, DSP chips, and more. "OpenCL deliberately stays focused at the silicon/software interface—we let others add value at higher levels in the software stack," Trevett told Ars. "It really isn't magic, per se, but there is a lot of value in the industry agreeing on a cross-platform standard."

That wide, cross-platform support will also offer benefits to mobile computing as well. Imagination Technologies recently announced built-in support for OpenCL in the latest revision of its PowerVR SGX graphics cores—the ones used on many smartphones and portable devices, including Apple's iPhone and iPad. So improvements developed on the desktop will be instantly portable to the mobile space as well.

Though Apple was the first to offer OpenCL support on a systemwide basis, NVIDIA supports Windows 7 and Linux with necessary API and compiler support on those platforms. AMD is also offering support, and Intel is working on integrating support for its range of silicon (though support for its IGPs has yet to appear, much to Apple's chagrin). An impressive list of industry players are contributing to OpenCL development, including AMD, Apple, ARM, Blizzard Activision, Broadcom, Codeplay, Electronic Arts, Ericsson, Freescale, Graphic Remedy, IBM, Imagination Technologies, Intel, Kestrel Institute, Los Alamos National Laboratory, Movidia, Nokia, NVIDIA, Petapath, Presagis, Qualcomm, Renesas, S3 Graphics, Seaweed Systems, Sony, STMicroelectronics, Symbian, and Texas Instruments.

Trevett told Ars that NVIDIA will release OpenCL 1.1 drivers today, so ISVs will be able to try out OpenCL 1.1 code on NVIDIA hardware immediately. Khronos Group has full specifications, online reference pages, downloadable reference cards, and complete conformance tests available immediately as well.

OpenCL 1.1 enhances performance with backward compatibility