Oktatási Segédlet. OpenCL

Miskolci Egyetem Gépészmérnöki és Informatikai Kar

Oktatási Segédlet

OpenCL Új párhuzamos szoftverfejlesztési lehetőségek

Dr. Nehéz Károly

Miskolc, 2012.02.29

Informatikai Rendszerek Építése

Oktatási segédlet

Tartalomjegyzék

Tartalomjegyzék Tartalom Bevezetés ................................................................................................................................................ 4 2. Grafikus processzorok ......................................................................................................................... 5 2.1. A GPU-k előnyei ................................................................................................................................ 5 2.2. CPU vagy GPU ................................................................................................................................... 9 3. Általános célú GPU programozás ...................................................................................................... 10 3.1. CUDA Eszköztár .............................................................................................................................. 10 4. OpenCL .............................................................................................................................................. 10 4.1 Platform modell ............................................................................................................................... 11 4.2 A végrehajtási modell ...................................................................................................................... 12 4.3 Memória modell .............................................................................................................................. 13 4.4 Programozási Modell....................................................................................................................... 13 5. Beépített típusok ............................................................................................................................... 14 6. Értékadások ....................................................................................................................................... 17 6.1. Alprogramok, modulok................................................................................................................... 20 6.2 Nyomkövetés, debuggolás .............................................................................................................. 21 7. Párhuzamosság .................................................................................................................................. 21 7.2 Kommunikáció ................................................................................................................................. 21 7.3 Szinkronizáció .................................................................................................................................. 22 7.4 Kapcsolat az OpenGL-el ................................................................................................................... 22 8. Egyszerű példa kernel ........................................................................................................................ 24 9. Az Aparapi java OpenCL technológia................................................................................................. 25 9.1 Fordítás és futtatás .......................................................................................................................... 26

2


Oktatási segédlet

9.3 Mandelbrot halmaz vizualizátor példa ............................................................................................ 27 10. Aparapi gyors referencia ................................................................................................................. 34 11. Egy prototípus marás szimulátor aparapival ................................................................................... 43 11.1 A prototípus alkalmazás bemutatása ............................................................................................ 44 11.2 A marás szimulációban használt kernel ........................................................................................ 45 11.3 A masírozó kockák algoritmusa ..................................................................................................... 48 12. Referenciák ...................................................................................................................................... 50 13. Összefoglalás ................................................................................................................................... 51

3


Oktatási segédlet

Bevezetés Napjaink grafikus kártyái már nem csak megjelenítésre használhatóak. A rajtuk található feladat specifikus processzor, a GPU képes fix és lebegőpontos számítások elvégzésére és a megjelenítési követelmények miatt erősen párhuzamos architektúrájú. A grafikai algoritmusok kifinomultsága miatt, ezeket a hardver eszközöket programozhatóvá tették. Mindez ahhoz vezetett, hogy a grafikus processzorokat általános számítási műveletek végrehajtására is használni kezdték. Ez a folyamat már körülbelül 10 éve elkezdődődött, először általános célú tárterületeket lehetett használni a grafikus kártyák saját memóriaterületein, majd később jöttek a programozható „shader”-ek. Pixel és vertex shaderek, amelyek a korábban beégetett algoritmusokat tették módosíthatóvá. A vertex shaderek alkalmazásával a geometriát lehetett rugalmasan módosítani, majd a pixel shaderekkel a megjelenítést. Egy megjelenítésre szánt buffer tartalmát a korábbi eljárások pixelenként tudták feltölteni és ezek a képpontok elvileg egymástól független számítás eredményeképpen kapták meg a „színüket” (nagyfokú párhuzamosság), ezért már korán természetesen jelentkezett az igény ezen feladatok a párhuzamos végrehajtására. Pixel shaderek esetén már 4-8 párhuzamos pixelszámítás végrehajtásra alkalmas eszközök jelentek meg (2002 körül). A nagyfokú párhuzamossága miatt a videokártyák számítási teljesítménye jóval meghaladja a CPU-k számítási teljesítményét, így általános célú használata

4


Oktatási segédlet

gyorsan elterjedt. Később ezt megkönnyítendő különféle fejlesztőkörnyezeteket és függvénykönyvtárakat kezdtek fejleszteni az általános számítási feladatok implementálására. Jelenleg bizonyos feladatokat a grafikus processzorok képesek akár 150-szer gyorsabban is elvégezni a hagyományos processzorokhoz képest, a nagy probléma kisebb feladatokra osztásával, majd az eredmények későbbi egyesítésével. Jelenleg az NVIDIA Tesla típusú kártyák már 512 maggal rendelkeznek és több kártya is egyidejűleg használható egy alaplapon, ennek eredményeképpen

az

Alkalmazott

Informatika

Tanszéken

hamarosan

beszerzésre kerül egy ~ 2000 magos számítógép 4 TESLA kártyával, amely alkalmas modern termelésinformatikai feladatok megoldására.

2. Grafikus processzorok A GPU architektúrájának köszönhetően erősen párhuzamos műveletvégzésre alkalmas, emellett nagy és gyorsan elérhető gyorsítót tárat is tartalmaz. Eredetileg ezek a processzorokat arra fejlesztették ki, hogy kétdimenziós képet állítsanak elő 3D modellekből és textúrákból a lehető leggyorsabban. Így a belső felépítésük és utasítás készletük ennek a megvalósítására lett megtervezve. Emiatt a specializáltság miatt nagy teljesítményt lehet elérni.

2.1. A GPU-k előnyei

5


Oktatási segédlet

 Nagymértékű párhuzamos lebegőpontos számítás A

jelenlegi

GPUk

több

mint

400

lebegőpontos

számításra

képes

processzormagot tartalmaznak, így számítási teljesítményük többszörösen meghaladja a CPUk számítási teljesítményét. A számítási teljesítmény alakulását az 1. ábra mutatja [1].

1. ábra. A CPU és GPU lebegőpontos számításai sebességeinek alakulása

6


Oktatási segédlet

A jelenlegi leggyorsabb Tesla GPU dupla pontosságú lebegőpontos számítási sebessége 515 Gflop [2]. Ezt a számítási sebbegéset a jelenlegi CPU-k meg sem közelítik. Nagy memória sávszélesség A GPUk a belső memóriájukat gyorsabban tudják elérni, mint a CPU a rendszermemóriát. A GPUk belső memóriája 128MB-tól akár 6GB-ig is terjedhet és a sávszélessége akár 144 GB/sec is lehet. A GPUk memória sávszélességének növekedését a 2. ábra mutatja [1]. Multi-GPUs számítás Egy számítógépen belül akár több grafikus kártyát is lehet alkalmazni, így a rendszer számítási teljesítménye még inkább megnövelhető.

7


Oktatási segédlet

2. ábra. A CPU és GPU memóriasávszélességének növekedése

A GPU-k architektúráját folyamatosan fejlesztik és nem csak a számítási teljesítmény növelése céljából, hanem a könnyebb programozhatóság érdekében is. A kezdeti szigorúan feladat-párhuzamos működést támogató architektúrát

8


Oktatási segédlet

leváltotta az adat-párhuzamos működési architektúra [3]. Napjainkban a grafikus

3. ábra. Egységesített GPU architektúra, amely 8 szálprocesszort tartalmaz egyenként 16 streaming multiprocesszorral.

kártyákat gyártó AMD és NVIDIA egységesített shader architektúrájú GPU-kat fejleszt. Az alkalmazások ez által ugyanolyan architektúrájú, de különböző hardverkörnyezetben futtathatóak.

2.2. CPU vagy GPU

9


Oktatási segédlet

A CPU-t és a GPU-t más céllal fejlesztik, ebből kifolyólag más-más feladatokat képesek hatékonyan végrehajtani. A CPU a számítógép központi egységeként sok folyamat és szál kezelését végzi hatékonyan. A szálak közötti kontextus váltásokra és az egyes szálak szekvenciális végrehajtására tervezték. Ezzel szemben a GPU-t egyetlen folyamat szálainak párhuzamos végrehajtásra tervezték. Emiatt jól lehet alkalmazni olyan esetekben, ahol nagy adathalmazon kell műveteket végrehajtani és az egyes adatok között minimális a függőség.

3. Általános célú GPU programozás A GPU-k programozása szinte tetszőleges programozási nyelven történhet, csak a megfelelő függvénykönyvtárakat kell használni. Több ilyen függvénykönyvtár használata is elterjedt.

3.1. CUDA Eszköztár A Compute Unified Device Architecture (CUDA), az nVidia párhuzamos programozási eszköztára, melynek segítségével csak az nVidia által gyártott grafikus kártyák programozhatóak. Ez az eszköztár könnyen használható GPUra optimalizált

matematikai

függvénykönyvtárakat

tartalmaz

C

és

C++

programozási nyelvhez.

4. OpenCL Az Open Computing Language, egy olyan szabvány, melynek segítségével párhuzamosan programozhatók heterogén rendszerek. A heterogenitás ebben az

10


Oktatási segédlet

esetben azt jelenti, hogy ugyanaz a kód képes módosítás nélkül futni bármelyik gyártó GPU-ján vagy CPUn. Hasonló használhatóságot biztosít, mint a CUDA eszköztár,

de

annak

függvénykönyvtárat

is

megkötési

nélkül.

definiál,

melynek

A

szabvány

segítségével

nyelvet

és

megvalósított

alkalmazásaink hordozhatóak lesznek az OpenCL képes hardvereken és az operációs rendszerek között, kezdve akár a mobiltelefonoktól egészen a szuperszámítógépekig. A nyelv a C99 ismert szabványára épül, de azt bizonyos részeken természetesen leszűkítve, máshol pedig kiegészítve alkalmazza. A szabvány létrehozását eredetileg az Apple kezdeményezte, és végül az ipar más szereplői dolgozták ki mint az AMD, az Intel, az IBM és az nVidia.

4.1 Platform modell Az OpenCL úgynevezett Platform Modellje egy gazda eszközből és a hozzá kapcsolódó OpenCL eszközökből állnak. Az OpenCL eszközök számoló egységekből állnak, melyek további végrehajtó egységekre bonthatóak. Végeredményben az OpenCL nyelven megírt alkalmazások ezeken a végrehajtó egységeken futnak. Érdekes módon az OpenCL eszköz például egy GPU vagy akár egy CPU is lehet, végrehajtó egységei pedig azok magja, illetve stream processzorai. Tehát az OpenCL program nem feltétlenül egy a videokártyákon futó alkalmazást jelent a gyakorlatban, olyan számítógépeken is futtatható, amelyik nem is rendelkezik videokártyával. Ez a tulajdonsága segíti elő azt, hogy szervergépeken is futtathatóvá teszi az alkalmazásokat.

11


Oktatási segédlet

Egy OpenCL programnak, a gazda eszközön futó része egy parancs listán keresztül vezérli a párhuzamos számításokat végző OpenCL eszközök működését.

4.2 A végrehajtási modell Az OpenCL programok végrehajtása alapvetően két fő részre bontható: a gazda program a gazdaeszközön, amíg az úgynevezett kernelek az OpenCL eszközökön

futnak.

A

kernelek

környezetének

beállítását

és

azok

végrehajtásának ütemezését a gazdaprogram végzi. Fogalmazhatunk úgy is, hogy a kernel kódja az a kód, amelyet OpenCL nyelven kell fejleszteni, a gazdakód pedig lehet bármilyen más hagyományos programnyelven íródott kód is. A kerneleket úgynevezett munkatereken hajtjuk végre. A szabvány egy-, kétilletve háromdimenziós munkatereket definiál. A végrehajtás során a munkaterület összes elemén lefut egy kernel. A kernel egy aktuális példányát munkaegységnek

hívjuk.

Minden

munkaegység

egy

egyedi

globális

azonosítóval (GlobalID-val) rendelkezik a munkatéren belül. A munkaegységek munkacsoportokba szerveződnek, ami egy durvább felosztása a munkatérnek. A munkacsoportok is rendelkeznek egy egyedi Work-group ID-val. Azaz azt mondhatjuk, hogy a kernelek futás közben azonosítani tudják, hogy melyik munkacsoportba illetve munkatérbe tartoznak. Ez azért lényeges, mert az alapkoncepció szerint a minden mag ugyanazt a kernelt futtatja adott esetben több ezer, vagy akár százezer példányban is!

12


Oktatási segédlet

4.3 Memória modell A munkaegységek négy egymástól elkülönülő különböző memóriaterületet érhetnek el: 

Globális memória – Minden munkacsoport minden egyes munkaegysége írhatja, olvashatja ezt a memóriát. Az OpenCL eszköz képességeinek megfelelően gyorsító tárazhatja az adott műveleteket.



Lokális memória – Egy munkacsoport helyi memóriájának tekinthető. Az itt allokált változók megosztottak a munkacsoport összes munkaegysége között.



Konstans memória – Olyan adatokat tartalmaz, mely nem változik a kernel lefutása során. A gazdaeszköz inicializálja és tölti fel adatokkal.



Privát memória – Egy munkaegység privát memóriája. Az itt meghatározott változók láthatatlanok a többi munkaegység részéről. Természetesen ennek a memóriának a mérete a legkisebb.

4.4 Programozási Modell Az OpenCL végrehajtási modellje támogatja mind az adatközpontú, mind a feladatközpontú párhuzamosságot. Elsődlegesen mégis inkább az adatközpontú párhuzamosság preferált a szabvány részéről. Adatközpontú párhuzamosság esetén egy számítási műveletsorozatot adunk meg, amit nagy mennyiségű adatra hajtunk végre egymástól függetlenül. Egy szigorú modell esetén a számítások és az adatok között 1-1 kapcsolat áll fenn.

13


Oktatási segédlet

Ezzel ellentétben az OpenCL egy jóval megengedőbb modellt definiál, ahol ez a kapcsolat nem követelmény. Az OpenCL egy hierarchikus végrehajtási modellt kínál a programozóknak. Két lehetőségünk van a hierarchia kialakítására a végrehajtásban. Explicit módon – mi adjuk meg mind a munkaegységek számát, és azt is, hogyan képeznek munkacsoportokat vagy Implicit módon azaz – mi csak a munkaegységek számát kötjük meg, a munkacsoport kezelését az OpenCL végrehajtó környezetre bízzuk. A feladat központú párhuzamosságot az OpenCL úgy teszi lehetővé, hogy megengedi a különböző kernelek különböző munkatereken való egyidejű végrehajtását. Ebben a modellben a felhasználó többféle módon érhet el párhuzamosságot:  a keretrendszer által implementált vektor típusok használata  több feladat egy időben való végrehajtása  olyan natív kernelek hívása, melyek az OpenCL-től különböző párhuzamosságot valósítanak meg.

5. Beépített típusok A következő táblázat tartalmazza a név bool

meghatározás Feltételes típus. Értéke true vagy false. Az igaz érték

14


Oktatási segédlet

megfelel az egész 1-nek, a hamis az egész 0-nak char

8 bites, előjeles egész

unsigned char, 8 bites, előjeletlen egész uchar short

16 bites, előjeles egész

unsigned short, 16 bites, előjeletlen egész ushort int unsigned int, uint long

32 bites, előjeles egész 32 bites, előjeletlen egész 64 bites, előjeles egész

unsigned long, 64 bites, előjeletlen egész ulong float

Egyszeres pontosságú lebegőpontos szám. A műveleteknek meg kell felelnie a IEEE-754 szabványnak.

half

Fél (16 bit) pontosságú lebegőpontos szám. A műveleteknek meg kell felelnie a IEEE-754-2008 szabványnak.

size_t

Előjel nélküli egész, a sizeof operátor eredményének típusa

Lehetőségünk van vektorokat is képezni a char, unsigned char, short, unsigned short, integer, unsigned integer, long, unsigned long, float típusokból, a következő szintaktika alapján: charn, ucharn, shortn, ushortn, intn, uintn, longn, ulongn, floatn, ahol n 2, 4, 8 és 16 közül kerülhet ki.

Vektor literálok segítségével vektorokat hozhatunk létre skalárok, más vektorok vagy ezek keverékének listájával. A literál egy zárójelezett vektortípusnévből és egy zárójelezett értéklistából áll. Ezek felhasználhatók inicializálásra vagy konstansként egy kifejezésben. Egy literál értéknek kell lennie vagy pontosan

15


Oktatási segédlet

annyinak, amekkora a létrehozandó vektor mérete. Egy érték esetén az minden vektorkomponensnek értékül adódik. Például: float4 vf = (float4)(1.0f, 1.0f, 1.0f, 1.0f); // egyszerűsített megadás uint4 vu = (uint4)(1); // eredmény (1, 1, 1, 1). float4 vf = (float4)((float2)(1.0f, 2.0f),(float2)(3.0f, 4.0f)); float4 vf = (float4)(1.0f, (float2)(1.0f, 1.0f), 1.0f); // a követező hibás megadás, mert csak 2 értéket tartalamaz float4 vf = (float4)(1.0f, 1.0f);

Kiegészítő típusok a grafikai műveletek fejlesztésének segítéséhez:

Típus

Leírás

image2d_t

2D-s képobjektum. Grafikus hardveren annak 2D-s textúráit reprezentálja.

image3d_t

3D-s képobjektum. Grafikus hardveren annak 3D-s textúráit reprezentálja.

Mintavételező képobjektumok kezelésére. Megadhatunk különböző sampler_t szűrési módszereket (legközelebbi, lineáris). Grafikus hardver esetén annak mintavételező egységeit reprezentálja.

Változó

deklarációknál

kulcsszavakkal

a

__global,

meghatározhatjuk

a

__local, változó

__private, helyét

a

__constant memóriában.

Alapértelmezetten minden változó __private minősítőt kap. Különböző típusú memóriára

hivatkozó

pointerek

nem

adhatók

értékül

egymásnak.

minősítőkkel ekvivalensek a global, local, private, constant kulcsszavak is.

16

A


Oktatási segédlet

A kernelek nem írhatnak és olvashatnak egyszerre egy képobjektumot. A __read_only (read_only), __write_only (write_only) minősítőkkel állíthatjuk be a hozzáférés módját. Az alapértelmezett a __read_only kulcsszó. Továbbá van lehetőség a tömb, rekord, unió típus létrehozására. A szintaxis megfelel a szokásos C99 szabványnak.

typedef struct { float2 a[3]; int4 b[2]; } mytype; foo_t my_info[10]; union{ float a; uint b; int i} myinion;

A C99 szabvánnyal ellentétben a következő felsorolás tartalamazza, hogy mi az amit nem enged az OpenCL:  rekurzió nem támogatott,  bitmezők nincsenek támogatva,  függvénypointerek nem hozhatók létre,  változó hosszúságú dinamikus tömbök nem használhatóak,

6. Értékadások A C-ben megszokott módon lehet értékadásokat programozni, a legnagyobb eltérés a vektorok esetén van. A 2, 3, 4 elemű vektorok komponenseit elérhetjük rendre a megszokott x, y, z, w tagokkal. Egy vektor nem létező elemére való hivatkozás hibát okoz.

17


Oktatási segédlet

float2 position; position.x = 4.1f; // a következő értékadás hibát okoz, mert 2 elemnél nincs z index! position.z = 3.2f;

A bal oldalon lehetőséget kapunk a több index egyidejű megadására is. float2 position; position.xy = (float2)(4.1f,3.33f); float4 pos; pos.xyz = (float3)(4.1f,3.33f,-2.0f);

Megengedett a komponensek keverése, jobb oldali kifejezés esetén azok többszörözése is. float4 pos = (float4)(4.0f, 3.0f, 2.0f, 1.0f); // keverés, többszörözés float4 rev = pos.wzyx; // rev = (1.0f, 2.0f, 3.0f, 4.0f) float4 dup = pos.xxyy; // dup = (1.0f, 1.0f, 2.0f, 2.0f) pos.xw = (float2)(5.0f, 6.0f); // pos = (5.0f, 2.0f, 3.0f, 6.0f) pos.wx = (float2)(7.0f, 8.0f); // pos = (8.0f, 2.0f, 3.0f, 7.0f) pos.xyz = (float3)(3.0f, 5.0f, 9.0f); // pos = (3.0f, 5.0f, 9.0f, 4.0f) // ez hibás lesz, mert kétszeres x komponens pos.xx = (float2)(3.0f, 4.0f); // hibás - float2 és float4 típusok között értékadás nem lehetséges pos.xy = (float4)(1.0f, 2.0f, 3.0f, 4.0f); float4 a, b, c, d; float16 x; x = (float16)(a, b, c, d); x = (float16)(a.xxxx, b.xyz, c.xyz, d.xyz, a.yzw);

18


Oktatási segédlet

// hibás – a.xxxxxx nem érvényes vektor típus x = (float16)(a.xxxxxxx, b.xyz, c.xyz, d.xyz);

További lehetőség a vektor komponenseinek indexel való elérése. A lehetséges indexek a következők:

VEKTOR

INDEX

2 elemű 0, 1 3 elemű 0, 1, 2 4 elemű 0, 1, 2, 3 8 elemű 0, 1, 2, 3, 4, 5, 6, 7 16 elemű 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, a, A, b, B, c, C, d, D, e, E, f, F

Az index jelölést egy s vagy S betű előzi meg, de a jelölés nem keverhető az .xyzw szelektorokkal. Például: float8 f; // f.s0 - f első eleme // f.s7 - f 8. eleme float16 x; // x.sa vagy x.sA - f 11. eleme // x.Sf vagy x.SF - f 16. eleme

A vektor műveletek komponens szinten működnek, a skalárok automatikusan konvertálódnak a megfelelő méretű vektorra. Például: float4 v, u, w; float f;

19


Oktatási segédlet

v = u + f; // ekvivalens v.x = u.x + f; v.y = u.y + f; v.z = u.z + f; v.w = u.w + f; w = v + u; // ekvivalens w.x = v.x + u.x; w.y = v.y + u.y; w.z = v.z + u.z; w.w = v.w + u.w;

6.1. Alprogramok, modulok Modulok létrehozása a nyelvben nem támogatott. Függvények használata C-ben megszokott módon történik. A __kernel módosítószóval jelölt függvények a program elsődleges belépési pontjai. Egy forrásállományban több belépési pont is lehet. A gazda program dönti el, hogy melyiket kívánja futtatni, viszont a visszatérési értéke csak void lehet. Függvényre mutató pointerek használatára nincs lehetőség. A C nyelvben megszokott módon hozhatunk létre függvényeket. int sample_fuction(float4 par1, float2 *par2) { // ... return 0; }

20


Oktatási segédlet

Paraméterátadásnál, az egyszerű változók érték szerint, a tömbök pedig cím szerint kerülnek átadásra.

6.2 Nyomkövetés, debuggolás A kód debuggolása és profilozása is nehézkes lehet a különböző architektúrák miatt. Az egyes gyártók készítenek eszközöket, hogy saját termékeiken lehetővé tegyék a hibakeresést. Az egyik legnagyobb tudású eszköz az nVidia Parallel Nsight, amely nVidia GPU-kon könnyíti meg a fejlesztést.

7. Párhuzamosság Nyelvi szinten a párhuzamosság támogatása a vektor típusokon való műveletek hatékony elvégzésével ki is merül. Ez a módszer kihasználják a grafikus processzorok hatékony utasításkészletét, illetve a CPU-k

újabb

SSE

(párhuzamos vektorkezelő) lehetőségeit. További nyelvi megoldások hiányának oka, hogy az OpenCL programok alapvetően párhuzamos futásra készülnek, melyek vezérléséért az OpenCL API-n keresztül a gazdaprogram a felelős. [4]

7.2 Kommunikáció A teljesítményre való tekintettel a szálak a kommunikációhoz osztott memóriát használnak. [4] Erre két szinten van lehetőség: 

a globális memória minden szál által használt memória terület, míg

21




Oktatási segédlet

a lokális memóriát az egyes munkacsoportok használhatják.

7.3 Szinkronizáció Szinkronizációt nyelvi szinten a munkacsoportok között, illetve API szinten a parancslistában elhelyezett parancsok között tudunk kiváltani. A kódban kiadott barrier parancs segítségével egy szinkronizációs pontot hozhatunk létre. A munkacsoporton belüli szálak futása felfüggesztődik addig, míg a munkacsoport összes tagja el nem éri ezt parancsot. Abban az esetben, ha egy szál nem jut el a barrier utasításig, akkor a többi szál blokkolódik. Ezért veszélyes a parancs használata feltételes szerkezetekben illetve ciklusokban. [4]

7.4 Kapcsolat az OpenGL-el Az OpenCL-t napjainkban már számos GPU és CPU támogatja. Mivel a videókártya az egyik legelterjedtebb végrehajtási platformja az OpenCL programoknak, ezért természetes igény, hogy a szabvány együttműködést biztosítson a különböző grafikus API-kal. Támogatja mind az DirectX mind az OpenGL-lel való együttműködést. Azonban ez utóbbi, mivel a két szabványt fejlesztő szervezet ugyanaz, nagyobb támogatást élvez a szabvány részéről.

22


Oktatási segédlet

Az együttműködés során az optimálisabb működés érdekében közös memóriaterületek használatára van lehetőség, illetve a két API működésének szinkronizációjára is lehetőség van. [4] A következő erőforrások megosztására van lehetőségünk: 

CL Buffer Objects - OpenGL Buffer objektumokból hozhatjuk létre a clCreateFromGLBuffer függvény segítségével.



CL Image Objects - OpenGL textúrákat és Render targeteket használhatunk

OpenCL

Image

clCreateFromGLTexture2D,

erőforrásként.

A

clCreateFromGLTexture2D,

clCreateFromGLTexture2D függvény segítségével hozhatjuk létre a megosztott objektumot Az együttműködés biztosításához az OpenCL kontextust egy már inicializált OpenGL kontextus segítségével kell létrehozni. Ezt a inicializálás során megfélően beállított paraméterek segítségével tehetjük meg. Például Windows operációs rendszer alatt a következő módon: [4]

cl_context_properties props[] = { CL_GL_CONTEXT_KHR, (cl_context_properties)wglGetCurrentContext(), CL_WGL_HDC_KHR, (cl_context_properties)wglGetCurrentDC(), CL_CONTEXT_PLATFORM, (cl_context_properties)platforms[platformID], 0 }; clContext = clCreateContext(props,deviceCount,devices,NULL,NULL,&errCode) ;

23


Oktatási segédlet

Az OpenCL-nek jeleznie kell igényét, illetve lemondását a közös erőforrásokról a clEnqueueAcquireGLObjects, clEnqueueReleaseGLObjects függvények segítségével. Ezek a megosztási lehetőségek implementációtól függően rendelkezésre állhatnak x86 architektúrákon is. Mind az AMD, mind az Intel támogatja ezt a lehetőséget. Azonban ennek használata esetén számolni kell a memória műveletek miatt fellépő teljesítmény csökkenéssel. Ez a negatívum azonban már valószínűleg nem jelentkezik az idén bevezetett új generációs processzor architektúrák esetén, mivel az egy lapkára integrált grafikus és x86 magok közös memóriával, sőt gyorsító tárral is rendelkeznek. [4]

8. Egyszerű példa kernel A példaként bemutatott kernel két integer típusú tömb összegzését végzi.

__kernel void vectorSum (__global int* a, __global int* b, __global int* c) { int i = get_global_id(0); c[i] = a[i] + b[i]; }

Bemenő paraméterek a globális memóriában lesznek tárolva, három darab bemenő paramétert látunk. A get_global_id(0) hívás megadja azt, hogy a szál hányadik a sorban, természetesen ez alapján fog megtörténni a tömbindex

24


Oktatási segédlet

megadása. A kernelt annyi példányban érdemes elindítani, amennyi a tömbelemek száma.

9. Az Aparapi java OpenCL technológia Az aparapi segítségével java fejlesztők könnyen fejleszthetnek adatpárhuzamos alkalmazásokat OpenCL alapon. Az aparapi technológia képes arra, hogy a java byte kódot futás időben OpenCL kóddá alakítsa és a gazdakód megírásának idejét is nagyban lerövidíti. A kód futási idejének növekedését nagyon könnyű érzékelni és mérni, mert egy kapcsoló segítségével ugyanaz a kód a CPU-n fog lefutni java thread pool alkalmazásával. Természetesen az OpenCL közvetlen használata nagyobb rugalmasságot biztosít, a párhuzamos fejlesztések tanulmányozására az aparapi nagyon jó lehetőséget biztosít. Vegyük példaként a következő egyszerű kódot: final float inA[] = .... // A tömb final float inB[] = .... // B tömb ahol (inA.length==inB.length) final float result = new float[inA.length]; for (int i=0; i<array.length; i++){ result[i]=intA[i]+inB[i]; }

25


Oktatási segédlet

A két tömb méretének meg kell egyeznie, az eredménytömb a két tömb egyes elemeinek az összegét fogja tartalmazni. Hogyan néz ki ennek az aparai refaktorált változata? Kernel kernel = new Kernel(){ @Override public void run(){ int i = getGlobalId(); result[i]=intA[i]+inB[i]; } }; Range range = Range.create(result.length); kernel.execute(range);

A kernel osztálynak tartalmaznia kell egy run() metódust, aminek a tartalma nagyon hasonlít a korábban bemutatott OpenCL kód logikájához. A Range osztály segítségével meg lehet adni, hogy mennyi szál induljon el párhuzamosan, majd a kernel.execute() hívással lehet elindítani a folyamatot.

9.1 Fordítás és futtatás Aparapi technológia alkalmazásához két feltétel szükséges:  aparapi.jar fájlnak benne kell lenni a class path-ban fordításkor  a legenerált class fájlnak kell tartalmazni debug információt javac –g –cp ${APARAPI_DIR}/aparapi.jar Squares.java

Futtatáskor a következő parancsot adhatjuk ki: java –cp ${APARAPI_DIR}/aparapi.jar;. Squares

26


Oktatási segédlet

Ebben az esetben a futás a java thread pool segítségével a CPU-n fog végrehajtódni, ahhoz hogy a GPU-n hajtódjon végre, meg kell adni a JNI shared lib-eket is. A –Djava.library.path megadásával. java –Djava.library.path=${APARAPI_DIR} –cp ${APARAPI_DIR}/aparapi.jar;. Squares

Az aparapi detektálja a JNI könyvtárak meglétét. A végrehajás után egy feltétel beiktatásával megbizonyosodhatunk arról, hogyan történt a futtatás: Kernel kernel = new Kernel(){ @Override public void run(){ int i = getGlobalId(); out[i]=in[i]*in[i]; } }; kernel.execute(in.length); if (!kernel.getExecutionMode().equals(Kernel.EXECUTION_MODE.GPU)){ System.out.println(”Kernel did not execute on the GPU!”); }

9.3 Mandelbrot halmaz vizualizátor példa A következőkben bemutatjuk hogy lehet mandelbrot halmaz vizualizációt segíteni aparapi segítségével, a kód tartalmazza az értelmezéshez szükséges magyarázatokat. /* Copyright (c) 2010-2011, Advanced Micro Devices, Inc. All rights reserved.

27


Oktatási segédlet

Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met: Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer. Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution. Neither the name of the copyright holder nor the names of its contributors may be used to endorse or promote products derived from this software without specific prior written permission. THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. If you use the software (in whole or in part), you shall adhere to all applicable U.S., European, and other export laws, including but not limited to the U.S. Export Administration Regulations ("EAR"), (15 C.F.R. Sections 730 through 774), and E.U. Council Regulation (EC) No 1334/2000 of 22 June 2000. Further, pursuant to Section 740.6 of the EAR, you hereby certify that, except pursuant to a license granted by the United States Department of Commerce Bureau of Industry and Security or as otherwise permitted pursuant to a License Exception under the U.S. Export Administration Regulations ("EAR"), you will not (1) export, re-export or release to a national of a country in Country Groups D:1, E:1 or E:2 any restricted technology, software, or source code you receive hereunder, or (2) export to Country Groups D:1, E:1 or E:2 the direct product of such technology or software, if such foreign produced direct product is subject to national security controls as identified on the Commerce Control List (currently found in Supplement 1 to Part 774 of EAR). For the most current Country Group listings, or for additional information about the EAR or your obligations under those regulations, please refer to the U.S. Bureau of Industry and Security's website at http://www.bis.doc.gov/. */

28


Oktatási segédlet

package com.amd.aparapi.sample.mandel; import import import import import import import import import import

java.awt.Color; java.awt.Dimension; java.awt.Graphics; java.awt.Point; java.awt.event.MouseAdapter; java.awt.event.MouseEvent; java.awt.event.WindowAdapter; java.awt.event.WindowEvent; java.awt.image.BufferedImage; java.awt.image.DataBufferInt;

import javax.swing.JComponent; import javax.swing.JFrame; import com.amd.aparapi.Kernel; public class Main { public static class MandelKernel extends Kernel{ /** RGB buffer, ami tartalamazni a fogja a megjelenítendő képet. */ final private int rgb[]; /** kép szélessége. */ final private int width; /** kép magassága */ final private int height; /** Paletta 0..maxIterations. */ final private int pallette[]; /** Maximum interációk száma */ final private int maxIterations; /** scale, offsetX, offsetY változók arra szolgálnak, hogy a képre kattintáskor a nagyítást a megfelelő irányba szolgáltassák*/ private float scale = .0f; private float offsetx = .0f; private float offsety = .0f;

29


Oktatási segédlet

/** * Initialize the Kernel. * * @param _width image width * @param _height image height * @param _rgb image RGB buffer * @param _pallette image palette */ public MandelKernel(int _width, int _height, int[] _rgb, int[] _pallette) { width = _width; height = _height; rgb = _rgb; pallette = _pallette; maxIterations = pallette.length - 1; } @Override public void run() { /** meghatározza, hogy melyik pixelre történik a számítás, a getGlobalId() függvény visszaadja, melyik szálban vagyunk éppen, ennek lehetséges értékei: 0…width*height */ int gid = getGlobalId(); /** A következő 2 utasítás visszaszámolja a globalID alapján, hogy pontosan melyik x és y koordinátát kell kiszámolni. Továbbá figylelmbe veszi az eltolást és az aktuális nagyítási faktort is. */ float x = (((gid % width * scale) - ((scale / 2) * width)) / width) + offsetx; float y = (((gid / height * scale) - ((scale / 2) * height)) / height) + offsety; int count = 0; float zx = x; float zy = y; float new_zx = 0f; // addig iterálunk, ameddig elérjuk a maximum iterációk számát, vagy már nem kell tovább számolni while (count < maxIterations && zx * zx + zy * zy < 8) { new_zx = zx * zx - zy * zy + x; zy = 2 * zx * zy + y;

30


Oktatási segédlet

zx = new_zx; count++; } // a paletta értéke az adott iterációhoz rgb[gid] = pallette[count]; } public void setScaleAndOffset(float _scale, float _offsetx, float _offsety) { offsetx = _offsetx; offsety = _offsety; scale = _scale; } } /** ez tartalmazza azt apontot, amire a felhasználó rákattintott */ public static volatile Point to = null; @SuppressWarnings("serial") public static void main(String[] _args) { JFrame frame = new JFrame("MandelBrot"); /** Width of Mandelbrot view. */ final int width = 768; /** Height of Mandelbrot view. */ final int height = 768; /** Maximum iterations for Mandelbrot. */ final int maxIterations = 256; /** Palette which maps iteration values to RGB values. */ final int pallette[] = new int[maxIterations + 1]; //Initialize palette values for (int i = 0; i < maxIterations; i++) { float h = i / (float) maxIterations; float b = 1.0f - h * h; pallette[i] = Color.HSBtoRGB(h, 1f, b); } /** Image for Mandelbrot view. */

31


Oktatási segédlet

final BufferedImage image = new BufferedImage(width, height, BufferedImage.TYPE_INT_RGB); final BufferedImage offscreen = new BufferedImage(width, height, BufferedImage.TYPE_INT_RGB); // Draw Mandelbrot image JComponent viewer = new JComponent(){ @Override public void paintComponent(Graphics g) { g.drawImage(image, 0, 0, width, height, this); } }; // Set the size of JComponent which displays Mandelbrot image viewer.setPreferredSize(new Dimension(width, height)); final Object doorBell = new Object(); // Mouse listener which reads the user clicked zoom-in point on the Mandelbrot view viewer.addMouseListener(new MouseAdapter(){ @Override public void mouseClicked(MouseEvent e) { to = e.getPoint(); synchronized (doorBell) { doorBell.notify(); } } }); // Swing housework to create the frame frame.getContentPane().add(viewer); frame.pack(); frame.setLocationRelativeTo(null); frame.setVisible(true); // Extract the underlying RGB buffer from the image. // Pass this to the kernel so it operates directly on the RGB buffer of the image final int[] rgb = ((DataBufferInt) offscreen.getRaster().getDataBuffer()).getData(); final int[] imageRgb = ((DataBufferInt) image.getRaster().getDataBuffer()).getData(); // Create a Kernel passing the size, RGB buffer and the palette. final MandelKernel kernel = new MandelKernel(width, height, rgb, pallette);

32


Oktatási segédlet

float defaultScale = 3f; // Set the default scale and offset, execute the kernel and force a repaint of the viewer. kernel.setScaleAndOffset(defaultScale, -1f, 0f); kernel.execute(width * height); System.arraycopy(rgb, 0, imageRgb, 0, rgb.length); viewer.repaint(); // Report target execution mode: GPU or JTP (Java Thread Pool). System.out.println("Execution mode=" + kernel.getExecutionMode()); // Window listener to dispose Kernel resources on user exit. frame.addWindowListener(new WindowAdapter(){ public void windowClosing(WindowEvent _windowEvent) { kernel.dispose(); System.exit(0); } }); // Wait until the user selects a zoom-in point on the Mandelbrot view. while (true) { // Wait for the user to click somewhere while (to == null) { synchronized (doorBell) { try { doorBell.wait(); } catch (InterruptedException ie) { ie.getStackTrace(); } } } float float float float float

x = -1f; y = 0f; scale = defaultScale; tox = (float) (to.x - width / 2) / width * scale; toy = (float) (to.y - height / 2) / height * scale;

// This is how many frames we will display as we zoom in and out. int frames = 128;

33


Oktatási segédlet

long startMillis = System.currentTimeMillis(); for (int sign = -1; sign < 2; sign += 2) { for (int i = 0; i < frames - 4; i++) { scale = scale + sign * defaultScale / frames; x = x - sign * (tox / frames); y = y - sign * (toy / frames); // Set the scale and offset, execute the kernel and force a repaint of the viewer. kernel.setScaleAndOffset(scale, x, y); kernel.execute(width * height); System.arraycopy(rgb, 0, imageRgb, 0, rgb.length); viewer.repaint(); } } long elapsedMillis = System.currentTimeMillis() startMillis; System.out.println("FPS = " + frames * 1000 / elapsedMillis); // Reset zoom-in point. to = null; } } }

A példa alapján látható, hogy mennyire hatékony tud lenni az aparapi alkalmazása, mert a programozási nyelv java marad a továbbiakban is, de ki tudjuk használni az OpenCL ben rejlő előnyöket is.

10. Aparapi gyors referencia Kernel létrehozása a com.amd.aparapi.Kernel osztályból származó saját osztály létrehozásával.

34


Oktatási segédlet

A Kernel.run() és az általa elérhető függvények olvashatnak és írhatnak tömböket vagy final jelölésű egyszerű változókat is.

class MyKernel extends Kernel { int data[] = new int[1024]; final float pi = 3.1415f;

@Override public void run() { // data[] és pi elérése } }; Névtelen osztály segítségével is kompakt módon lehet kerneleket létrehozni. Ebben az esetben is a kernel képes elérni azokat a final-al megadott változókat, amelyeket a run() függvénye is elérhet.

final int data[] = new int[1024]; final float pi = 3.1415f; Kernel kernel = new Kernel() { @Override public void run() { // data[] és pi elérése } };

35


Oktatási segédlet

A kernel elindítása előre megadott szál segítségével. Ekkor a Kernel.execute(int ) hívással lehet megadni a szálak számát. Minden szál indulásakor megkapja, hogy hányadik a sorban.

final int data[] = new int[1024]; Kernel kernel = new Kernel() { @Override public void run() { Data[getGlobalId()] = getGlobalId(); } }; kernel.execute(data.length); Az alapértelmezett futási mód beállítása a következő módon lehetséges:

Kernel kernel = new Kernel() {…} kernel.setExecutionMode(Kernel.EXECUTION_MODE.JTP); kernel.execute(data.length); A lefuttatott kernel hiba esetén automatikusan java thread pool módban futtatja az alkalmazást. Hogyan lehet ezt utólag lekérdezni? Kernel kernel = new Kernel() {…} kernel.setExecutionMode(Kernel.EXECUTION_MODE.JTP); kernel.execute(data.length); if(!kernel.getExecutionMode().equals(Kernel.EXECUTION_MODE.GPU)) { // nem a GPU futtatta a kernelt

36


Oktatási segédlet

} Az aparapi futásidejű viselkedését módosító konstansok és változók:

Mi történik amikor meghívóvódik a kernel.execute metódus?

37


Oktatási segédlet

Explicit buffer menedzsmentre szükségünk lehet, ha nagyméretű tömböket másolunk, mivel lehet, hogy a kernel elindítása előtt még nincs teljesen átmásolva a buffer tartalma a videokártya memóriába.

final int hugeArray[] =new int[HUGE]; final boolean[] done = new boolean[] {false}; Kernel kernel = new Kernel(){ // reads and writes hugeArray sets done[0] to true when complete. }; while (!done[0])}

38


Oktatási segédlet

kernel.execute(range); } Viszont hatékonyabb megoldás, ha explicite megadjuk, hogy mely bufferekre van szükségünk.

final int hugeArray[] =new int[HUGE]; final boolean[] done = new boolean[false]; Kernel kernel = new Kernel(){ // reads and writes hugeArray }; kernel.setExplicit(true); // we take control of all transfers kernel.put(hugeArray); kernel.put(done); while (!done[0])} kernel.execute(HUGE); kernel.get(done); } kernel.get(hugeArray); Egyszerűbb megoldást is kínál a rendszer

final int hugeArray[] =new int[HUGE]; final boolean[] done = new boolean[false]; Kernel kernel = new Kernel(){ // reads and writes hugeArray

39


Oktatási segédlet

}; kernel.setExplicit(true); // we take control of all transfers kernel.put(hugeArray).put(done); while (!done[0])} kernel.execute(range).get(done); } kernel.get(hugeArray); A végrehajtás finomítása létrejöhet úgy, ha az egyes kernelek meg tudják határozni, hogy melyik munkacsoportba tartoznak

A következő táblázat bemutatja az egyes függvények értelmezését. Függvény neve getGlobalId()

Értelmezése 0.. intervallumban megadja az adott kernel globális szálazonosítóját

getGroupSize()

A munkacsoportok mérete. Minden kernel ugyanazt az értéket kapja természetesen.

getGroupId()

Az

40

aktuális

kernel

melyik


Oktatási segédlet

munkacsoportban dolgozik. getLocalId()

Az aktuális kernel a munkacsoportján belül hányadik.

getGlobalSize()

Hány kernel indult el összesen.

getPassId()

A rendszer képes arra, hogy a kernelt többször visszadja,

indítja

el.

A

hogy

függvény hanyadik

munkamenetben vagyunk. Fontos megjegyezni, hogy használhatjuk az OpenCL beépített matematikai függvényeit is a következő 2 táblázat megmutatja melyek ezek.

41


Oktatási segédlet

42


Oktatási segédlet

11. Egy prototípus marás szimulátor aparapival A megmunkálás szimuláció problémájának matematikai leírása a következő [6]. Legyen M vonalszakaszok halmaza a térben, amely leírja a szerszámmozgást. Legyen f függvény a szerszám és alakjuk szerint ujjmaró (1) vagy gömbvégű maró (2). f ( x, y) : r , x, y  D( f ) : {( x, y) | x 2  y 2  r 2 } (1)

f ( x, y) :  r 2  x 2  y 2 , x, y  D( f ) : {( x, y) | x 2  y 2  r 2 } (2)

43


Oktatási segédlet

Mindkét esetben az origót tekinthetjük a programozott pontnak. Az x és y értékei leírják a szerszámmozgást a nyersdarab vízszintes síkjában, z pedig a magasságot jelenti. Ha a munkadarabot egy téglatestnek tekintjük xmin, xmax, ymin, ymax, zmin, zmax méreteivel megadva, akkor az anyagleválasztás matematikailag a következő implicit formulával írható le (3). F ( x, y) : min{ z max, min{ z | z : s( x' , y' )  f ( x  x' , y  y' ), ( x  x' , y  y' )  D( f ), z  z min , ( x' , y' , s( x' , y' ))  M }} (3)

Az

F(x,y)

függvény

kiszámítása

időigényes

feladat

lehet,

főleg

ha

szabadformájú felületek esetén ahol a szerszámhelyzetek száma több tízezer is lehet. Az irodalomban egzakt és analitikus megközelítések léteznek F(x,y) meghatározására. Jelen cikkben a OpenCL technológiát alkalmazunk GPU segítségével.

11.1 A prototípus alkalmazás bemutatása Az irodalomban implicit függvények megjelenítésére számos módszer ismert: az u.n. masírozó kockák algoritmusa, valamint duális kontúrozás. A prototípusban a masírozó kockák algoritmusát alkalmaztuk, amelynek hátrány az élek automatikus letörése. A szimuláció lépései a következőek: 1.) szerszám, munkadarab definíció implicit felületekkel; 2.) szerszámhelyzetek előállítása; 2.) a (3) egyenlet kiszámítása az egyes szerszámhelyzetekre OpenCL segítségével, párhuzamosan;

44


Oktatási segédlet

Gömbvégű szerszám pályaszimulációja ~1600 vonalszakasszal 2500ms számítási idővel (azaz: 1.5 ms szerszámhelyzetenként) A 1. ábrán a marási folyamat vizualizálását látjuk. Összehasonlítva a GPU alapú megoldással körülbelül 20 szoros sebességnövekedést tudunk elérni. Szimuláció környezet mindössze egy laptop AMD Radeon HD 5700-as videokártyával, ami 4 GPU maggal rendelkezik. Az implementáció Aparapi OpenCL technológiát alkalmaz [7]. A forráskód letölthető: http://users.iit.uni-miskolc.hu/~simon11/conferences/FMTU2012/millingsimulator.zip

11.2 A marás szimulációban használt kernel package org.ait.millingsimulator; import com.amd.aparapi.Kernel;

45


Oktatási segédlet

public class MillSimKernel extends Kernel { private int mSize; private int mSizePow; private float onePerMsize; final private float ToolX = -0.6f; final private float ToolY = 0.7f; final private float ToolZ = 0.3f; final private float ToolR = 0.05f; final private float[] valueVector; public MillSimKernel() { valueVector = null; mSizePow = 0; } public MillSimKernel(int _mSize, float[] _valueVector) { mSize = _mSize; mSizePow = mSize * mSize; onePerMsize = 1.0f / (float) (mSize); valueVector = _valueVector; } @Override public void run() { int gid = getGlobalId(); int k = gid / (mSizePow); int j = (gid - (k * mSizePow)) / mSize; int i = gid - (k * mSizePow) - (j * mSize); float x = (float) (2 * i) * onePerMsize - 1.0f; float y = (float) (2 * j) * onePerMsize - 1.0f; float z = (float) (2 * k) * onePerMsize - 1.0f; valueVector[gid] = workpiece(x, y, z); } public float workpiece(float x, float y, float z) { float val = 0.0f; // F1 kivonása F2 ből // F1 - F2 // // max( -f1(x,y,z) , f2(x,y,z) ) // return val; float toolX = ToolX;

46


Oktatási segédlet

float toolY = ToolY; float toolZ = ToolZ; val = max(-tool(x, y, z, toolX, toolY, toolZ), max3(abs(x) - 0.7f, abs(y) - 0.7f, abs(z) 0.35f)); for (int k = 0; k < 23; k++) { // ide kell a vizszintes mozgás... toolX = -0.6f; toolZ = 0.3f; for (int i = 0; i < 70; i++) { toolX += 0.02f; toolZ -= 0.003f * 8f * sin(toolX) * sin(toolY); // fuggolegesen // emeli if (val > 0.1f) { return val; // ha már nincs anyag akkor anyag nem lesz később sem } val = max(-tool(x, y, z, toolX, toolY, toolZ), val); } toolY -= 0.05f; // elmozditja sikban } return val; } float tool(float x, float y, float z, float tx, float ty, float tz) { float val; val = pow2(x - tx) + pow2(y - ty) + pow2(z - tz) - ToolR; return val; } float pow2(float x) { return x * x; } private float max3(float x, float y, float z) { if (x > y) {

47


Oktatási segédlet

if (x > z) { return x; } else { return z; } } else { if (y > z) { return y; } else { return z; } } } }

A run() metódusban használt érdekes megoldás az, hogy a kernelek száma megegyezik a téglatest 3 oldalhosszának szorzatával. Kérdésként merül fel, hogyan lehet megállapítani a kernel által megkapott golbalId ből, hogy melyik térbeli pontot számoljuk? Erre ad választ a következő képlet: int k = gid / (mSizePow); int j = (gid - (k * mSizePow)) / mSize; int i = gid - (k * mSizePow) - (j * mSize);

Először a z iránynak megfelelő k éréket, majd a, j és i értékeket határozzuk meg.

11.3 A masírozó kockák algoritmusa Érdekes probléma, hogyan lehet egy térbeli adathalmazt megjeleníteni akkor, ha a térbeli pontokban adottak a sűrűségértékek. Jelen esetben ezt úgy kell elképzelni, hogy egy adott pont a térben tartalmazza, hogy ott van-e a anyag vagy nincs. Hogyan lehet hatékonyan megállapítani a felület határát? Erre szolgál a masírozó kockák algoritmusa. Ennek lényege az, hogy egy 3 szoros ciklussal végigmegyünk a kockánként az adott térfogaton és a kocka 8

48


Oktatási segédlet

csúcspontjában megállapítjuk, hogy van-e anyag vagy nincs. Ha mind a 8 pont azonos értéket tartalmaz, akkor vagy a test belsejében, vagy kívül vagyunk. Összesen a 8 csúcspont miatt 256 variáció lehet. A következő ábra azt mutatja, milyen felület lehet a határfelület a köztes esetekben.

A 256 esetet le lehet egyszerüsíteni a fenti 16 elemi esetre. Az első mutatja azt az esetet, amikor nem kell felületet rajzolni, mert vagy bent vagy kint vagyunk. A második kép azt mutatja, hogy a 7 pont benne van a testben, de a 8. kívül esik. Ilyenkor kell egy sarkot rejzolni egyetlen háomszöggel. A második sor első eleme egy olyan esetet vizsgál, amikor 4 csúcs benne van és 4 csúcs nincs benne a testben, ilyenkor 2 háromszöggel lapot készítünk. Az algoritmus ki van egészítve egy olyan finomítással, hogy a csúcspontokban az implicit függvények értéke van eltárolva, azaz nem true/false értékek, ez

49


Oktatási segédlet

alapján kiszámolható, hogy a háromszögek pontosan hová helyezendőek a kockán belül, ezzel valóságosabban lehet a felületet megjeleníteni.

12. Referenciák [1] NVIDIA Corporation: NVIDIA CUDA compute unified device architecture programming guide, Verzió 4.0, 5/6/2011, http://developer.nvidia.com/cuda, [utoljára megtekintve 2012. January 6.]

[2]

http://www.nvidia.com/object/personal-supercomputing.html,

[utoljára

megtekintve 2011 November 2.]

[3] Owens, J., Houston, M., Luebke, D., Green, S., Stone, J., Phillips, J.: GPU Computing, Proceedings of the IEEE 96(5), 2008, 879–899 oldal. [4] Az OpenCL programozási nyelv, http://nyelvek.inf.elte.hu/leirasok/OpenCL/index.php, 2012

[5] Aparapi quick reference guide. http://aparapi.googlecode.com/svn/trunk/QuickReference.pdf, 2012 [6] H. Müller, F. Albarsmann, A. Zabel, Efficient raster-based simulation and visualization of 3-axis Milling of Free-formd Shapes. Research report No. 667/1998, Fachbereich Informatik Universitat Dortmund, Germany, 1998

50


Oktatási segédlet

[7] Az aparapi hivatalos weboldala, http://code.google.com/p/aparapi/, 2012

13. Összefoglalás Ebben a segédletben betekintést nyerhetünk az OpenCL programozási technológia alapvető tulajdonságaival, példákon keresztül bemutatásra kerülnek a nyelv alaptulajdonságai. Java fejlesztőknek szánt aparapi technológia bemutatásával és egy saját fejlesztésű marás szimulátor prototípus alkalmazás bemutatásával a hallgatók példák alapján is tanulmányozhatják a technológia lényegét. Jelen oktatási segédlet a TÁMOP-4.2.1.B-10/2/KONV-2010-0001 jelű projekt részeként az Európai Unió támogatásával, az Európai Szociális Alap társfinanszírozásával valósult meg.

51

Oktatási Segédlet. OpenCL

Recommend Documents