Példaprogram - Előkészületek

Kernel függvény

OK, most, hogy átrágtuk magunkat a száraz tényeken, következzen egy konkrét feladat megvalósítása OpenCL-ben. A mellékelt forráskódot érdemes nyitva tartani olvasás közben, illetve természetesen az OpenCL specifikáció is hasznos olvasmány. Ez utóbbit azért is célszerű olvsani, mert a használt API függvények argumentumait nem mutattam be részletesen, csupán leírtam, hogy adott paraméterezéssel mit csinál az adott függvény.

Példának egy egyszerű dolgot választottam: a BLAS csomag SAXPY függvényét fogjuk megvalósítani. A SAXPY fog két vektort, az elsőt megszorozza egy skalárral, majd az eredményt hozzáadja a második vektorhoz (y = a * x + y). 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 i = 0; i < n; ++i)
{
y[i] = a * x[i] + y[i];
}

A ciklus bejárja a vektorok elemeit, és elvégzi rajtuk a megfelelő műveletet. Szépen sorban, egymás után. Egy i-edik elem kiszámolásához csak a két vektor i-edik elemeire, és a konstansra 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 SAXPY-t 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, fogjunk is neki!

Első lépésként írjuk meg a kernelt!

__kernel void saxpy(__constant float * x, __global float * y, __private const float a)
{
const uint i = get_global_id(0);
y[i] += a * x[i];
}

Ennyi az egész. A __kernel kulcsszó - meglepő módon :) - azt jelzi, hogy az adott függvény egy kernel függvény.(FIGYELEM! A következő rész át lett írva, és a forráskód is javítva lett. Az ok a 32-es hsz-ben található. Köszönet karics-nak, hogy felhívta a figyelmemet a hibára) Egy kernel argumentumai a __private névtérben kell legyenek, és alapértelmezetten oda is kerülnek, ezért az "a" változó elől akár el is hagyhatnánk a __private kulcsszót. Ha egy argumentum mutató (a mi kernelünk esetében ez ugye igaz x-re és y-ra is), megmondhatjuk, hogy a __global, __local és __constant névterek melyikébe mutasson. Az x a __constant névtérbe mutat, ugyanis az x vektor értékei nem fognak változni a program futása során (ez CUDA architektúrán gyorsabb, mintha a __global címtérbe mutatna, ugyanis a konstans memória cache-elt). Az y vektorba tesszük a számolás eredményét, ezért az a __global névtérbe mutat.

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 második 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 a következőképpen szerezhetjük meg:

cl_uint num_platforms;
cl_platform_id * platforms;
clGetPlatformIDs(0, NULL, &num_platforms);
platforms = malloc(sizeof(cl_platform_id) * num_platforms);
clGetPlatformIDs(num_platforms, platforms, NULL);

Ez a kódrészlet az első clGetPlatforms hívással kideríti az elérhető OpenCL platformok számát, lefoglal egy megfelelő méretű tömböt a platform ID-knek, majd a második hívással lekárdezi őket. A mellékelt példaprogramban az első elérhető platformot használjuk.

Eszközök

Ha megvan a platform azonosítónk (platform_id), lekérdezhetjük az elérhető eszközök listáját:

cl_uint num_devices;
cl_device_id * devices;
clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_ALL, 0, NULL, &num_devices);
devices = malloc(sizeof(cl_device_id) * num_devices);
clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_ALL, num_devices, devices, NULL);

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. A clGetDeviceIDs második paramétere meghatározza, hogy milyen típusú eszközökkel szeretnénk foglalkozni. Most minden elérhető eszközt felsoroltattunk, de lehetne pl. csak a GPU-kat (CL_DEVICE_TYPE_GPU), vagy csak a CPU-kat (CL_DEVICE_TYPE_CPU). Egy megjegyzés: ha a platform id (a clGetDeviceIDs első paramétere) NULL, akkor a specifikáció szerint a függvény viselkedése implementációfüggő. Az Apple példaprogramjaiban NULL-t használ. Bár erről külön nem írnak sehol, gondolom azért, mert úgyis csak egy platform van, annak az ID-jét használja.

Környezet

Most, hogy megvan az eszköz, létre kell hozni egy környezetet a számoláshoz. Erre való a clCreateContext API függvény:

cl_context context;
context = clCreateContext(0, 1, &device_id, NULL, NULL, &err);

Ez a függvény létrehoz egy környezetet a harmadik paraméterben megadott azonosítójú eszközhöz (vagy eszközökhöz, ugyanis többet is megadhatunk - a második paraméter közli a függvénnyel ezek számát). Az err változóban egy hibakódot kapunk vissza, melynek értéke sikeres végrehajtás esetén CL_SUCCESS. Ezt érdemes ellenőrizni egy programban. Környezetet nem csak ily módon hozhatunk létre, lehet pl. minden, adott típusú eszközhöz környezetet kreálni. GOTO OpenCL doksi :)

Parancslista

OK, van környezetünk, most kell hozzá egy parancslista, ahová majd a végrehajtandó kerneleket pakoljuk.

cl_command_queue commands;
commands = clCreateCommandQueue(context, device_id, 0, &err);

Ez a hívás a device_id 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 ö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. Ja, é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 az írásban nem térek ki, a példaprogi forrásában viszont ott van, ki lehet lesni :) )!

cl_program program;
program = clCreateProgramWithSource(context, 1, &kernel_source, NULL, &err);

OK, van egy programunk, de ez még csak a forráskódót tartalmazza, tehát le kell fordítanunk:

err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);

Ha a függvény visszatérési értéke nem CL_SUCCESS, valami hiba történt a fordításkor. Erről bővebb információt a clGetProgramBuildInfo függvény szolgáltat, ha elkérjük tőle a build logot:

if (err != CL_SUCCESS)
{
size_t size;
clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &size);
char * log = malloc(size);
clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, size, log, NULL);
printf("%s\n", log);
exit(1);
}

Kernel objektum

Miután sikeresen felépítettük a programot, kinyerhetjük belőle a kernel objektumot.

cl_kernel kernel;
kernel = clClreateKernel(program, "saxpy", &err);

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á.

A cikk még nem ért véget, kérlek, lapozz!