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
Az elsővel tulajdonképpen már megismerkedtünk: a host/device kapcsolatról van szó.
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);
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
OK, már majdnem kész vagyunk, már csak futtatnunk kell a kernelt. Egyszerűen hangzik, de azért ez nem csak annyiból áll, hogy meghívunk egy függvényt, aztán fut. Be kell állítani a kernel-argumentumokat, meg kell határozni a munkacsoportok és munkaegységek számát, és elő kell jegyezni a kernelt a parancslistában.
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
Hát igen, nem valami meggyőző. Nagy elemszámnál már a Cloo átveszi a vezetést, de nem sokkal. Mi lehet ennek az oka?! Először is a .NET nem is olyan lassú, mint hinnénk, másodszor pedig az OpenCL API-t nem tisztán értük el, hanem a Cloo wrapperen keresztül, így egy kis overhead mindig is lesz. Ez főleg akkor megy a teljesítmény rovására, ha a kernel program olyan apró és egyszerű, mint a mi VectorAdd példánk. Jóval nagyobb és összetettebb feladatok esetén már jóval nagyobb különbség is elérhető a CPU és a GPU között. Remélem érthető volt ez kis példa és talán meghozza a kedvet egy kis videokártya-programozáshoz. A teljes forrás letölthető:
innen
[iPon] A Khronos Group szoftvermérnökei 18 hónappal az OpenCL 1.1-es változatának hivatalos bejelentése után végre elkészültek az OpenCL 1.2-es verziójával. A nyílt, mindenféle díjaktól mentes szabvány tökéletesen működik a különböző platformok között, így széleskörű támogatást nyújt ahhoz, hogy a mai modern processzorokat – CPU-kat, GPU-kat – hatékonyan vethessék be a fejlesztők párhuzamos számítási feladatokra.
Bejegyzések
