Példaprogram - Pufferek, futtatás, eredmény

Memória objektumok

Nos, van már kernelünk, amit tudnánk futtatni, csak egy probléma van: nincsenek adataink, amin dolgozhatunk. No problemo, csinálunk 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árom dimenzió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:

cl_mem dev_x;
cl_mem dev_y;
dev_x = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(float) * VECTOR_SIZE, NULL, NULL);
dev_y = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float) * VECTOR_SIZE, NULL, NULL);

Ez a két függvényhívás létrehozza az adott kontextusban az x és y vektoroknak megfelelő puffereket. Az x vektort csak olvasni fogjuk, így a CL_MEM_READ_ONLY flaggel hozzuk létre, míg az y vektorhoz szükség van a CL_MEM_READ_WRITE 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 CL_MEM_COPY_HOST_PTR flaget. Puffert létrehozhatunk a host memóriájában is (CL_MEM_ALLOC_HOST_PTR), illetve felhasználhatunk már lefoglalt host memóriaterületet is (CL_MEM_USE_HOST_PTR).

A vektorokat a következő módon töltjük fel az eszköz memóriájába:

clEnqueueWriteBuffer(commands, dev_x, CL_TRUE, 0, sizeof(float) * VECTOR_SIZE, x, 0, NULL, NULL);
clEnqueueWriteBuffer(commands, dev_y, CL_TRUE, 0, sizeof(float) * VECTOR_SIZE, y, 0, NULL, NULL);

A clEnqueueWriteBuffer API függvény előjegyez egy pufferbe írást a megadott parancslistában. Itt konkrétan a host memóriában lévő x és y vektorokat másoljuk a dev_x és dev_y memória objektumokba. A függvény utolsó három paraméterét most nem használjuk, de azért elmondom, mire jók. Minden olyan függvény, ami a parancslistához ad elemeket, kér egy cl_event típusú pointert (utolsó paraméter). Az így visszaadott cl_event-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. A clEnqueueWriteBuffer hetedik argumentuma a lista mérete, a nyolcadik pedig maga a várólista (cl_event típusú tömb).

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.

clSetKernelArg(kernel, 0, sizeof(cl_mem), &dev_x);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &dev_x);
clSetKernelArg(kernel, 2, sizeof(float), &a);

Ezek az API 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 clSetKernelArg függvény második paramétere. A munkaegységek száma megegyezik a vektorok méretével. Azt, hogy ez hogyan van felosztva munkacsoportokra, 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.

const size_t work_items = VECTOR_SIZE;
clEnqueueNDRangeKernel(commands, kernel, 1, NULL, &work_items, NULL, 0, NULL, NULL);

Ezzel a függvénnyel előjegyeztük a kernelfuttatást, megadtuk az index tér dimenzióinak (harmadik paraméter), és a munkaegységeknek (ötödik paraméter) a számát. 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 clFinish függvénnyel biztosíthatjuk, ami csak akkor tér vissza, ha a paraméterként megadott parancslista minden parancsa befejeződött.

Az eredmények host memóriába olvasása:

clEnqueueReadBuffer(commands, dev_y, CL_TRUE, 0, sizeof(float) * VECTOR_SIZE, result, 0, NULL, NULL);

Amikor ez lefut, a dev_y memória objektum tartalma a host memóriában lévő result vektorba másolódik. A pufferekbe írásnál kimaradt, most megemlítem: ha a harmadik paraméter CL_TRUE, az olvasás/írás befejeztéig nem tért vissza a függvény, ha CL_FALSE, nem várja meg a műveletek befejeztét.

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!

clReleaseMemObject(dev_x);
clReleaseMemObject(dev_y);
clReleaseProgram(program);
clReleaseKernel(kernel);
clReleaseCommandQueue(commands);
clReleaseContext(context);

Végszó

Még mielőtt befejezem, egy megjegyzést engedjetek meg nekem, mely az Xcode használóknak szól: a Snow Leopard már csak az Intel platformot támogatja, így az OpenCL Framework nem is tartalmaz PPC binárisokat. Alapértelmezetten a C/C++ projectek Release konfigurációban fordítanának PPC-re is, de a hiányzó OpenCL binárisok miatt ez természetesen nem sikerül. A megoldás: a project Properties lapján a Build tabon a Valid Architectures listából ki kell venni a PPC architektúrákat.

Most már tényleg végeztünk :) Remélem érthető voltam, és talán néhányatoknak meghoztam a kedvét egy kis videokártya-programozáshoz. Természetesen minden kérdést/kritikát/javaslatot szívesen fogadok. Köszönöm, hogy elolvastátok az írást!

Kapcsolódó anyagok, források

BrookGPU - A Stanford University GPGPU megoldása
Sh - Egy metaprogramozási nyelv GPU-khoz
CUDA - Az nVidia GPGPU API-ja
Khronos OpenCL API Registry - A Khronos Group OpenCL-lel kapcsolatos anyagai
OpenCL Programming Guide for Mac OS X - Az Apple OpenCL doksija

Példaprogram forráskód: [link]
OpenCL Info Browser - első Cocoa progim, a gépen elérhető OpenCL platform/eszköz infókat mutatja