A most következő írás egy kezdő lecke azok számára, akik valamilyen okból kifolyólag menedzselt környezetből szeretnék használni az OpenCL API-t. Az írás sghc_toma OpenCL kedvcsináló című cikkén alapszik, ezért az azzal való egyezés nem a véletlen műve. Ezer köszönet érte sghc_toma-nak. Kezdjünk is hozzá.
Az OpenCL számolás nagy vonalakban
Az OpenCL eszköz (compute device) az a hardver, amin a párhuzamos feldolgozás történik. Ezekben a hardverekben egy, vagy több számolási egység (compute unit) van, melyek egy, vagy több feldolgozó egységet (processing element) tartalmaznak – tulajdonképpen ezek hajtják végre az utasításokat. Például egy videokártya minden stream processzora, és egy CPU minden magja is egy feldolgozó egység, az nVidia videokártyák egy multiprocesszora pedig egy számolási egység.
A videokártyán futó programrészlet neve kernel kód vagy függvény. Az egyes feldolgozó egységek a kernel kód egy példányát futtatják, mindegyikük más-más adatokon. A kerneleket OpenCL-C nyelven (ISO C99 alapokon) lehet írni, és minden eszközre külön le kell fordítani őket. Az OpenCL eszközre írt kódokat természetesen nem kell egy darab függvénybe sűríteni, a kernel hívhat más függvényeket, illetve lehet több kernel is. A kernelek, a kiegészítő függvények, illetve a kernel által használt konstansok együtt egy programot alkotnak.
A kernelek egy úgynevezett kontextusban (context), környezetben futnak, mely magába foglalja a használható eszközöket, az általuk elérhető memória objektumokat, illetve a kernelek futtatásának ütemezését végző parancslistákat (command queue). A programot, amely létrehozza a kontextusokat, előjegyzi a kernelek futtatását, host programnak nevezzük, az őt futtató hardvert pedig a host eszköznek. A videokártya szempontjából a host (gazda) a CPU. A host kódból lehet memóriát lefoglalni a kártyán, feltölteni adattal, és a kernel kód futása után a host kódból lehet elérni a számolás eredményét. A kernel kódból ez a memória hozzáférhető. A kernel kódot úgy kell elkészíteni, hogy az a feladatnak csak egy részét végzi el; hogy melyik részét, azt beépített paraméterek adják meg. A feladat felosztását mi határozzuk meg, azaz kijelöljük a GPU számára a paraméter tartományokat, amelyekben végre kell hajtani a kernelt, és a GPU saját erőforrásait figyelembe véve maga beosztja, hogy melyik rész mikor és melyik multiprocesszoron fut le.
A fentieket egy egyszerű példán át könnyedén megérthetjük. Tegyük fel, hogy egy mátrixszorzást szeretnénk elvégezni az OpenCL-el. Ehhez a kernel kódban általánosan, sor és oszlop paraméterre leírjuk azt, hogy a szorzatmátrix adott sorában és oszlopában, hogy áll elő az eredmény. A kernel számára kijelöljük, hogy a sor paraméter egytől a sorok számáig, az oszlop paraméter egytől az oszlopok számáig kell, hogy terjedjen. A GPU tehát kap egy feladatot: sor × oszlop-szor végrehajtani a kernelt. A feladat végrehajtását optimálisan a GPU határozza meg magának.
A platform modell
Az OpenCL működése leírható négy modell segítségével:
- a platform modell
- a futtatási modell
- a memória modell
- a programozási modell
A futtatási modell
Mint már említettük, az OpenCL eszköz egy feldolgozó egysége a kernel egy példányát futtatja. Egy kernel-példányt munkaegységnek (work-item) nevezünk. Mikor elindítunk egy számolást, meg kell adnunk, hogy összesen hány munkaegységre lesz szükségünk. A munkaegységek összessége az index tér (index space), mely lehet 1, 2, vagy 3 dimenziós. A munkaegységek munkacsoportokba (work-group) szervezhetők. Ez azért fontos, mert az egy munkacsoportba tartozó munkaegységek között lehetséges szinkronizáció, és mindegyikük hozzáfér a csoport lokális memóriájához (erről később bővebben), míg ez a különböző csoportba tartozó egységekről ez nem mondható el.
Minden munkaegységnek van egy úgynevezett globális azonosítója (global ID), mely egyértelműen meghatározza annak helyét az index térben. Hasonlóan, minden munkacsoportnak is van egy azonosítója (work-group ID). A munkaegységeknek ezen felül van egy helyi azonosítója (local ID), mely a munkacsoporton belüli helyét határozza meg. A fentiekből következik, hogy a munkaegység pozíciója az index térben meghatározható a csoport azonosító és a helyi azonosító kombinációjával. Az index tér dimenzióinak maximális száma, az egyes dimenziókban a maximális méret, illetve egy munkacsoport maximális mérete eszközönként eltérő lehet, ezt figyelembe kell venni programozás közben! Az OpenCL API természetesen lehetőséget nyújt ezen adatok lekérdezésére.
A memória modell
A munkaegységek/kernelek által hozzáférhető memória négy típusra van osztva. A globális memóriához (global memory) az index tér minden egyes munkaegysége hozzáfér, azt írni és olvasni is tudják. Eszköztől függően a globális memória írása/olvasása lehet cache-elt. A konstans memória (constant memory) a globális memória egy olyan része, melynek tartalma nem változik kernelfuttatás közben, azt csak a host módosíthatja.
Minden munkacsoport rendelkezik egy lokális memória (local memory) területtel, melyet minden munkaegység a csoportban képes írni/olvasni. A privát memória (private memory) minden egyes munkaegységnek a sajátja, ő írni/olvasni tudja, de más munkaegység nem fér hozzá. A CUDA-t ismerőknek: érdemes odafigyelni az elnevezésekre, mert bár a két környezet felépítése hasonló, a terminológia nagyon nem. Például a CUDA-féle lokális memória az, amit itt privátnak neveznek, és ami itt lokális memória, az a CUDA-ban megosztott (shared memory).
A programozási modell
OpenCL-t használva kétféle párhuzamosítást érhetünk el: data, illetve task parallel módon programozhatunk. Az első az OpenCL fő profilja, ez jelenti azt, hogy sok kernel példány csinálja ugyanazt az index tér más-más elemein. Ha több, más feladatot végző kernelünk van, betehetjük őket egy parancslistába, és az OpenCL megtesz minden tőle telhetőt, hogy ezek optimálisan használják ki a hardvert. Ez utóbbi a task parallel módszer, hiszen egymástól független folyamatok futnak párhuzamosan.
OpenCL a gyakorlatban
Most hogy átrágtuk magunkat a száraz tényeken, következzen egy konkrét feladat megvalósítása OpenCL-ben. Először is szükségünk van egy wrapperre amivel elérhetjük az OpenCL API függvényeit .NET környezetből. Erre jelenleg az egyik legígéretesebb projekt a Cloo és az OpenCL.NET. Jelen bejegyzésben a Cloo-val fogunk megismerkedni mivel ez objektum orientált szemlélettel készült. Töltsük le a Cloo legfrissebb változatát majd indítsuk el a Visual Studio-t. Hozzunk létre egy konzolos alkalmazást majd a referenciákhoz adjuk hozzá a Cloo.dll szerelvényt. Ezután már csak a szükséges namespace-ket kell elhelyeznünk:
using System;
using System.Collections.Generic;
using System.Collections.ObjectModel;
using System.Runtime.InteropServices;
using System.Diagnostics;
using Cloo;
Első programunk rendkívül egyszerű lesz, összeadunk két vektort és az eredményt eltároljuk egy harmadikban. Persze ezt a műveletet elvégezzük, mondjuk úgy 10 milliószor. Ha nem párhuzamos megvalósításban gondolkodnánk, akkor ezt valószínűleg egy ciklussal oldanánk meg, valahogy így.
for (int index = 0; index < n; index++) c[index]= a[index] + b[index];
A számoláshoz a vektorok i-edik elemeire van szükség - tehát az egyes számolások teljesen függetlenek egymástól, így a probléma szinte felkínálja magát párhuzamosításra. A terv az, hogy írunk egy kernelt, ami elvégzi a fenti műveletet egy vektor egy elemére, majd egy (n elemű vektorokkal számolva) n elemből álló index térre rászabadítjuk ezt a kernelt. Egyszerűen hangzik, mert az is.
A kernel
Első lépés a kernel kód vagy függvény elkészítése:
private static string code = @" kernel void VectorAdd(
global read_only float* a,
global read_only float* b,
global write_only float* c )
{ int index = get_global_id(0);
c[index] = a[index] + b[index]; }";
Ennyi az egész. A kernel kulcsszó - meglepő módon - azt jelzi, hogy az adott függvény egy kernel függvény. Egy kernel argumentumai a private névtérben kell legyenek, és alapértelmezetten oda is kerülnek. Ha egy argumentum mutató, megmondhatjuk, hogy a global, local és constant névterek melyikébe mutasson. A get_global_id(0) függvény az adott munkaegység globális azonosítójának első koordinátáját adja vissza. Erre azért van szükség, hogy tudjuk, hogy egy adott kernelpéldány a vektor hányadik elemének számolásáért felel. Mivel vektorokon dolgozunk, az index terünk egydimenziós, így valóban csak az első koordinátára van szükségünk. Ezután a kernel már csak elvégzi a megfelelő műveleteket a vektorok megfelelő elemein, és az eredményt visszaírja a harmadik vektor megfelelő elemébe.
Platformok
A számoláshoz szükségünk van egy OpenCL eszközre. Ahhoz, hogy létrehozhassunk egy eszközt, szükségünk lesz egy platform ID-re. Az elérhető platformok listáját a következőképpen szerezhetjük meg:
int platformNo = ComputePlatform.Platforms.Count;
ComputePlatform platform = ComputePlatform.Platforms[0];
Ez a kódrészlet az első ComputePlatform.Platforms.Count hívással kideríti az elérhető OpenCL platformok számát. A mellékelt példaprogramban az első elérhető platformot használjuk.
Eszközök
Ha megvan a platformazonosítónk (platform), lekérdezhetjük az elérhető eszközök listáját:
ReadOnlyCollection<ComputeDevice> devices = platform.Devices;
Ez nagyon hasonlít az előző kódrészlethez, a különbség annyi, hogy itt eszközökről szerzünk listát.
Környezet
Most, hogy megvan az eszköz, létre kell hozni egy környezetet a számoláshoz. Ehhez első lépésben készítünk egy környezet tulajdonság listát a platform felhasználásával:
ComputeContextPropertyList properties =
new ComputeContextPropertyList(platform);
Majd a tényleges környezet létrehozása a ComputeContext osztály segítségével történik:
ComputeContext context =
new ComputeContext(devices, properties, null, IntPtr.Zero);
Az első paraméterben az elérhető eszközöket adjuk át, persze minden adott eszközhöz külön is kreálható környezet. Ilyenkor a ComputeDeviceTypes felsorolás tagjaiból választhatunk: Default, CPU, GPU, Accelerator, All. A következő a környezet tulajdonságait tartalmazó lista. Majd ezután, ha szükségesnek látjuk létrehozhatunk call-back függvényt amivel a hibakódokról értesülhetünk, de mi most ezt nem alkalmazzuk, ezért null.
Parancslista
Most már van környezetünk, amihez már csak hozzá kell adni egy parancslistát, ahová majd a végrehajtandó kerneleket pakoljuk.
ComputeCommandQueue commands =
new ComputeCommandQueue(context, context.Devices[0], ComputeCommandQueueFlags.None);
Ez a hívás a context.Devices[0] azonosítójú eszközhöz hoz létre egy parancslistát. Több parancslistát is létrehozhatnánk, de a kitűzött feladat elég egyszerű, egy lista is elég. Több lista esetén, amennyiben az egyes parancsok használnak közös objektumokat, figyelni kell a szinkronizálásra; erről bővebben olvashattok a OpenCL specifikációban. Még egy dolog a parancslistákkal kapcsolatban: több eszköz esetén mindegyiknek saját listára van szüksége!
Foglaljuk akkor most össze, mink van eddig. Kiválasztottunk egy OpenCL platformot, és erről a platformról egy eszközt, amin számolni fogunk. Ehhez az eszközhöz készítettünk egy környezetet, melyhez létrehoztunk egy parancslistát és írtunk egy kernel függvényt is. Most már csak pár dolgot kell végrehajtanunk: kellenek memória objektumok, amikben átadjuk a kernelnek a két vektort, illetve visszakapjuk az eredményt. A kernelfuttatáshoz szükségünk lesz egy kernel objektumra, amit csak egy, az adott eszközhöz felépített program objektumból nyerhetünk ki. Folytassuk tehát a munkát, készítsük el a program objektumot.
Program objektum
Programot készíthetünk forráskódból, illetve binárisból is. A gyorsabb inicializálás érdekében célszerű az első futtatáskor lefordítani a forrást, majd a kapott binárist (több OpenCL eszköz esetén az eltérő gépi kód miatt binárisokat) elraktározni, s következő alkalommal abból készíteni a program objektumot. Lássuk, hogy forrásból hogyan készítünk programot (a binárisból készítésre most nem térünk ki):
ComputeProgram program =
new ComputeProgram(context, new string[] { code });
Most már van egy programunk, de ez még csak a forráskódot tartalmazza, tehát le kell fordítanunk:
program.Build(null, null, null, IntPtr.Zero);
A fordítás állapotáról a program.GetBuildStatus(devices[0]) tulajdonságon keresztül kaphatunk információt. Ha minden rendben volt a ComputeProgramBuildStatus.Success enumot kapjuk vissza.
Kernel objektum
Miután sikeresen felépítettük a programot, kinyerhetjük belőle a kernel objektumot:
ComputeKernel kernel = program.CreateKernel("VectorAdd");
Mint látható, a kernel objektum gyártás elég egyszerű művelet, csupán egy lefordított programra, és a kernel függvény nevére van szükségünk hozzá.
Memória objektumok
Nos, van már kernelünk, amit tudnánk futtatni, csak egy probléma van: nincsenek adataink, amin dolgozhatunk. Ezért most létre kell hozni memória objektumokat, s feltöltjük őket a vektorainkkal. OpenCL-ben az eszköz memóriáját kétféle objektumon keresztül érhetjük el: puffer és kép (image) objektumok. A pufferekbe bármilyen adatot tehetünk, míg a kép objektumok kettő, vagy háromdimenziós képek tárolására alkalmasak. A kép objektumokkal most nem foglalkozom részletesebben, kitűzött célunk eléréséhez nincs is szükség rájuk, a vektorokat pufferekben tároljuk:
ComputeBuffer<float> a = new ComputeBuffer<float>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.CopyHostPointer, arrA);
ComputeBuffer<float> b = new ComputeBuffer<float>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.CopyHostPointer, arrB);
ComputeBuffer<float> c = new ComputeBuffer<float>(context, ComputeMemoryFlags.WriteOnly, arrC.Length);
ComputeBuffer<float> b = new ComputeBuffer<float>(context, ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.CopyHostPointer, arrB);
ComputeBuffer<float> c = new ComputeBuffer<float>(context, ComputeMemoryFlags.WriteOnly, arrC.Length);
Ez a három függvényhívás létrehozza az adott kontextusban az a, b és c vektoroknak megfelelő puffereket. Az a és b vektort csak olvasni fogjuk, így a ComputeMemoryFlags.ReadOnly flaggel hozzuk létre, míg a c vektorhoz szükség van a ComputeMemoryFlags.WriteOnly flagre, hiszen abba írjuk majd a végeredményt. Lehetőség van a puffer készítésekor megadni egy host pointert (negyedik paraméter), és az adatot, amire mutat, rögtön felmásolni a pufferbe. Ehhez szükséges megadni a ComputeMemoryFlags.CopyHostPointer flaget. Puffert létrehozhatunk a host memóriájában is (ComputeMemoryFlags.AllocateHostPointer), illetve felhasználhatunk már lefoglalt host memóriaterületet is (UseHostPointer).
Kernel futtatás
kernel.SetMemoryArgument(0, a);
kernel.SetMemoryArgument(1, b);
kernel.SetMemoryArgument(2, c);
Ezek a függvény hívások állítják be a kernel argumentumokat. Az egyes argumentumokat a kernel függvény fejlécében elfoglalt helyük alapján azonosítjuk, ez a szám a SetMemoryArgument függvény első paramétere.
ICollection<ComputeEvent> events =
new Collection<ComputeEvent>();
commands.Execute(kernel, null, new long[] { count }, null, events);
Most már csak meg kell hívnunk a commands parancslista Execute() metódusát megfelelően felparaméterezve. Az első paraméter a kernel objektum, a második a GlobalWorkOffset amit általában null. A GlobalWorkSize vagyis a munkaegységek száma megegyezik a vektorok méretével. Azt, hogy ez hogyan van felosztva munkacsoportokra (LocalWorkSize), most teljesen mindegy (az egyes vektor elemek teljesen függetlenek egymástól, a munkaegységeknek nincs szüksége közös memóriára), ezért az OpenCL-re bízzuk.
Minden olyan függvény, ami a parancslistához ad elemeket, kér egy ComputeEvent típusú lista pointert (utolsó paraméter). Az így visszaadott ComputeEvent-tel lehet lekérdezni a parancs aktuális állapotát (pl. végzett-e?), illetve más parancsok várólistájához lehet adni a parancsot. A várólistában szereplő összes parancsnak le kell futnia, mielőtt az adott parancs lefuthatna. Nos, ha minden jól ment, ezen hívás után már fut is a kernelünk.
Eredmény olvasása
A számolás eredménye persze az eszköz memóriájában van, tehát ki kell olvasnunk onnan. Ehhez persze biztosnak kell lennünk benne, hogy a számolás befejeződött. Ezt a commands.Finish() függvénnyel biztosíthatjuk, ami csak akkor tér vissza, ha a paraméterként megadott parancslista minden parancsa befejeződött. Ezután még figyelembe kell vennünk azt is, hogy a host memóriájában menedzselt az az objektum amibe bele kívánjuk másolni a számítás eredményét és így meg kell védeni a GC-től.
A menedzselt objektumok nem-menedzselt kódból történő elérésének egyik módja, hogy használjuk a GCHandle struktúrát. A struktúra legfontosabb tulajdonsága, hogy akkor sem engedi a menedzselt objektumot kisöpörni a memóriából, ha már csak egy nem-menedzselt objektum hivatkozik rá (ellenkező esetben a memóriaterület felszabadul). Azt is megakadályozza, hogy az adott cím egy más szegmensre tolódjon a memórián belül. Az azonban lényeges, hogy a nem-menedzselt kódnak rendelkeznie kell a megfelelő jogosultságokkal ehhez, melyet a SecurityPermission osztály segítségével adhatunk meg.
GCHandle arrCHandle = GCHandle.Alloc(arrC, GCHandleType.Pinned);
Az eredmények host memóriába olvasása:
commands.Read(c, false, 0, count,
arrCHandle.AddrOfPinnedObject(), events);
commands.Finish();
Amikor ez lefut, a c memória objektum tartalma a host memóriában lévő arrC vektorba másolódik. Az arrCHandle AddrOfPinnedObject() metódusa kinyeri a szükséges címet. A pufferekbe írásnál kimaradt, most megemlítem: ha a második paraméter true, az olvasás/írás befejeztéig nem tért vissza a függvény, ha false, nem várja meg a műveletek befejeztét. Végül:
arrCHandle.Free();
Ezuán már a számítás eredménye elérhető az arrC tömbben és azt tehetünk vele amit csak akarunk. Végeztünk a számolással, de még nem fejeződött be a dolgunk. Illik eltakarítani magunk után, szabadítsuk fel az OpenCL erőforrásokat:
a.Dispose();
b.Dispose();
c.Dispose();
kernel.Dispose();
program.Dispose();
context.Dispose();
commands.Dispose();
Végszó
Elsőre kicsit bonyolultnak tűnhet a dolog, de higgyük el, hogy megéri a kínlódás mivel így új távlatok nyílhatnak meg előttünk. De még mielőtt elbúcsúznánk érdemes egy kis teljesítmény tesztet végrehajtani a klasszikus .NET for ciklus és az OpenCL/Cloo (GeForce 8800 GTS) páros között. Az eredmények a lenti ábrán láthatóak:
.NET for vs. Cloo
Bejegyzések
0 megjegyzés :
Megjegyzés küldése