2. Az OpenCL API eszközei
7.17. példa - build.c
int i= get_global_id(0);
b[i]= sqrt(a[i]);
}
A példaprogram elsődleges célja a program és kernel objektumok használatának, kapcsolatának demonstrálása, a kernel-kód ismeretlen elemeire később térünk vissza.
7.17. példa - build.c
#include <stdio.h>
#include <string.h>
#include <error.h>
#include <kernelio.h>
#include <CL/opencl.h>
#define ARRAY_SIZE 20
#define MAX_DEVICES 2
#define MAX_PLATFORMS 2
int main(int argc, char** argv) {
cl_int err;
cl_context context;
cl_command_queue queue;
cl_mem memobjInput, memobjOutput;
float input[ARRAY_SIZE];
int i;
cl_event event;
cl_uint numDevices;
cl_device_id devices[MAX_DEVICES];
cl_uint numPlatforms;
cl_platform_id platforms[MAX_PLATFORMS];
cl_context_properties properties[MAX_PLATFORMS*2]= {0};
cl_uint size;
cl_program program;
cl_kernel kernel;
char* kernelSource;
unsigned int* kernelLength;
for ( i= 0; i < ARRAY_SIZE; ++i ) input[i]= i;
err= clGetPlatformIDs(MAX_PLATFORMS, platforms, &numPlatforms);
ERROR(err, "clGetPlatformIDs") for ( i= 0; i < numPlatforms; ++i ) {
properties[i*2]= (cl_context_properties)CL_CONTEXT_PLATFORM;
properties[i*2 + 1]= (cl_platform_id)(platforms[i]);
}
properties[i*2]= 0;
context= clCreateContextFromType(properties, CL_DEVICE_TYPE_ALL, NULL, NULL, &err);
ERROR(err, "clCreateContextFromType")
err= clGetContextInfo(context, CL_CONTEXT_DEVICES, MAX_DEVICES*sizeof(cl_device_id), devices, &size);
ERROR(err, "clGetContextInfo")
err= clGetContextInfo(context, CL_CONTEXT_NUM_DEVICES, sizeof(cl_uint), &numDevices,
&size);
ERROR(err, "clGetContextInfo")
queue= clCreateCommandQueue(context, devices[0], NULL, &err);
ERROR(err, "clCreateCommandQueue")
memobjInput= clCreateBuffer(context, 0, ARRAY_SIZE * sizeof(float), NULL, &err);
ERROR(err, "clCreateBuffer")
memobjOutput= clCreateBuffer(context, 0, ARRAY_SIZE * sizeof(float), NULL, &err);
ERROR(err, "clCreateBuffer")
readSourceProgram("sqrtKernel.k", &kernelSource, &kernelLength);
program= clCreateProgramWithSource( context, 1, &kernelSource, NULL, &err);
ERROR(err, "clCreateProgramWithSource")
err= clBuildProgram(program, numDevices, devices, NULL, NULL, NULL);
ERROR(err, "clBuildProgram")
kernel= clCreateKernel(program, "sqrtKernel", &err);
ERROR(err, "clCreateKernel")
err= clEnqueueWriteBuffer(queue, memobjInput, 1, 0, sizeof(float)*ARRAY_SIZE, input, NULL, NULL, &event);
ERROR(err, "clEnqueueWriteBuffer")
err= clSetKernelArg(kernel, 0, sizeof(memobjInput), &memobjInput);
ERROR(err, "clSetKernelArg")
err= clSetKernelArg(kernel, 1, sizeof(memobjOutput), &memobjOutput);
ERROR(err, "clSetKernelArg") size_t global_work_offset= 0;
size_t global_work_size= ARRAY_SIZE;
size_t local_work_size= 2;
err= clEnqueueNDRangeKernel(queue, kernel, 1, &global_work_offset, &global_work_size, NULL, 0, NULL, &event);
ERROR(err, "clEnqueueNDRangeKernel") err= clWaitForEvents(1, &event);
ERROR(err, "clWaitForEvents")
err= clEnqueueReadBuffer(queue, memobjOutput, 1, 0, sizeof(float)*ARRAY_SIZE, input, NULL, NULL, &event);
ERROR(err, "clEnqueueReadBuffer") err= clWaitForEvents(1, &event);
ERROR(err, "clWaitForEvents");
for ( i= 0; i < ARRAY_SIZE; ++i ) printf("%f ", input[i]);
printf("\n");
clReleaseMemObject(memobjInput);
clReleaseMemObject(memobjOutput);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseCommandQueue(queue);
clReleaseContext(context);
free(kernelSource);
return 0;
}
Az OpenCL C nyelvvel a következő szakaszban foglalkozunk, egyelőre fogadjuk el azt, hogy a megadott kernel gyököt von a memobjInput memória objektumba feltöltött egész számokból és az eredményt a memobjOutput memóriaterületre írja. Az OpenCL környezet beállítása és a kernel forráskódjának beolvasása után a clCreateProgramWithSource függvény segítéségével hozzuk létre a program objektumot, majd a futtatni kívánt kernel-t kiemeljük kernel objektum szintre. Paramétereként beállítjuk a megfelelő memóriaterületeket, majd a kernelek futtatásának kezdeményezésére a clEnqueueNDRangeKernel függvényt használjuk, paramétereként megadva az indextartományt, amelyet használni szeretnénk. A fenti példában a globális indextartományt csak egy dimenzióban specifikáltuk, azaz a kernel-függvény a 0, ..., 19 indexekre fog lefutni. Mivel a példában használt paraméterezéssel a clEnqueueNDRangeKernel függvény nem blokkolja a főprogram futását, gondoskodnunk kell róla, hogy ne fejeződjön be a főprogram (gazdaprogram), mielőtt amíg a kernel függvény le nem fut az indextartomány minden elemére. Erre használjuk a clWaitForEvents függvényt.
Az eredményeket a clEnqueueReadBuffer függvény segítségével az input tömbbe olvassuk, s elemeit a kimenetre írjuk.
A program futásának kimenete az elvárásoknak megfelelően az egész számok négyzetgyökeit tartalmazza 20-ig.
user@home> ./programkernel
0.000000 1.000000 1.414214 1.732051 2.000000 2.236068 2.449490 2.645751 2.828427 3.000000 3.162278 3.316625 3.464102 3.605551 3.741657 3.872983 4.000000 4.123106 4.242640 4.358900
2.4.3. Az OpenCL C programozási nyelv
Az OpenCL C programozási nyelv az ANSI C99 szabvány kiterjesztésének tekinthető. Alkalmasan definiált makrókkal egy OpenCL C nyelvű program egy szabványos ANSI C99 fordítóval lefordítható. Az ANSI C99 nyelvben definiált nyelvi eszközök az OpenCL C nyelvben is elérhetőek és kevés kivételtől eltekintve azonos működésűek. Az OpenCL C nyelv azonban számos olyan kiegészítést tartalmaz az ANSI C99 nyelvvel összehasonlítva, melyeket nem áll módunkban részletesen tárgyalni, ezért csak felsorolás szintjén ismertetjük a leglényegesebbeket, hogy az olvasónak benyomása legyen az elérhető eszközökről:
• bool típus, azaz explicit módon megjelenik a logikai típus;
• half típus fél pontosságú lebegőpontos ábrázolásra;
• minden atomi típushoz 2, 4, 8, 16 elemű vektor típus (például float2, int8; és ezeknek megfelelő aritmetikai operátorok;
• image* típusok képek optimalizált reprezentálására és kezelésére;
• a megjelenő új típusokhoz kapcsolódó optimalizált sztenderd könyvtári függvények tömkelege.
• címtér minősítő, hozzáférést minősítő és függvény minősítő alapszavak.
A fenti elemek közül elsősorban a minősítő alapszavakkal és néhány beépített függvénnyel ismerkedünk meg, amelyek használata elengedhetetlen OpenCL C nyelvű kernelek írásához.
Minősítők. Az OpenCL memória modelljének áttekintésénél négy különböző memóriaterületről beszéltünk:
globális, konstans, lokális és privát memóriákról. Koncepcionális szinten a fő különbség az egyes memóriaterületek méretében és a gazdaprogram, illetve a munkaelemek számára biztosított hozzáférésben van.
A gyakorlatban ezek a memóriaterületek az OpenCL eszköz és implementáció függvényében a hardver
különböző részén helyezkedhetnek el, így a hozzáférés sebességében is lényeges különbség mutatkozhat.
Jellemzően a konstans memória elérése jóval gyorsabb, mint a globális memóriáé, azonban ettől is gyorsabb a lokális vagy privát memóriák elérése. A különböző memóriaterületekhez tartozó címtereket négy minősítő segítségével azonosíthatjuk: __global, __local, __constant és __private. Ezen minősítő alapszavakat a kernelek formális paraméterlistáján, illetve deklarációs utasításokban használhatjuk, megadva velük a programozási eszköz létrehozásának helyét, azaz szintaktikailag a típusmódosítókhoz hasonlóan használhatjuk őket.
• A __global minősítővel létrehozott mutatók a globális memória címteréhez tartozó címet vehetnek fel értékül. Megjelenhetnek formális paraméter listán és lokális változó deklarációjában egyaránt.
• A __constant minősítésű mutatók a konstans memória címteréhez tartozó címet vehetnek fel értékül. A program szintű nevesített konstansokat a konstans memóriában __constant minősítővel kell deklarálni, s azok fordítási időben, konstansokkal inicializálhatók. Futásidőben történő módosításuk futási hibát okoz.
• A __local minősítésű eszközök a lokális memóriában jönnek létre, vagy mutatók esetén a lokális memória címteréből vehetnek fel értékeket. A lokális memóriában létrehozott eszközökhöz a munkacsoportokhoz tartozó munkaelemek osztott hozzáféréssel rendelkeznek, azaz közös (shared) változókat valósíthatunk meg velük. __local minősítésű változókat csak kernel-függvényekben deklarálhatunk, s azok minden munkacsoporthoz csak egyszer jönnek létre. __local minősítésű változóknak nem adhatunk értéket deklarációs utasításban, csak külön, önálló értékadó utasításban.
• Alapértelmezetten __private minősítésűek a kernel-függvények minősítés nélküli lokális változói, és csak __private minősítésűek lehetnek nem-kernel függvények lokális változói, illetve a kernel- és nem-kernel-függvények formális paraméterei. A __private minősítésű eszközök mind a privát memóriában jönnek létre, a minősítés nélkül deklarált mutatók a privát memóriába mutathatnak. A privát memória jellemzően igen kicsiny, implementációtól függően csupán néhány regisztert foglal magában. Ennek megfelelően elérésük igen gyors, azonban nagyon takarékosan kell bánni a privát memória foglalásával. Ha túl sok programozási eszközt helyezünk el a privát memóriában (megtelnek a regiszterek), akkor sem futási, sem fordítási hibát nem kapunk, azonban a regiszterekben el nem férő eszközök a globális memóriában kerülnek tárolásra, ami a leglassabb eléréssel rendelkezik az OpenCL eszköz különböző memóriaterületei közül. Egy ilyen jelenség nagyban lassíthatja programunk végrehajtását.
A címtér minősítők mellett megjelenik a kernel függvényeket minősítő __kernel alapszó. A kernel-függvények szintaktikájukat tekintve egyszerű C függvények, amelyek az OpenCL-eszközön történő végrehajtás belépési pontjaiként szolgálhatnak.
A kernel és címtér minősítők használata egyszerű, jelentésük könnyen értelmezhető, azonban néhány fontos szabályt be kell tartanunk:
1. A kernel függvények visszatérési értékének típusa mindig void kell, hogy legyen.
2. Kernel-függvények nem kaphatnak mutatóra mutató mutatót paraméterként, azonban egy kernel függvényen belül már használhatunk és nem-kernel függvényeknek át is adhatunk paraméterként mutatóra mutató típusú változókat.
3. Kernel függvények nem rendelkezhetnek bool, half, size_t, ptrdiff_t, intptr_t és uintptr_t típusú formális paraméterekkel, vagy olyan összetett paraméterrel, amelynek valamely mezője ezen típusok közül való.
4. Rekurziót nem valósíthatunk meg sem kernel, sem nem-kernel függvényekkel.
5. Nem használhatunk függvény mutatókat.
6. A kernel függvények minden mutató típusú paramétere a __global, __constant vagy __local minősítőkkel kell, hogy rendelkezzen.
7. A lokális memóriában létrejött változókat nem inicializálhatjuk a deklarációs utasításban.
8. __global, __constant vagy __local minősítésű nem-mutató formális paramétert nem használhatunk.
9. __global és __constant minősítésű nem-mutató típusú lokális változókat nem hozhatunk létre.
10. Nem végezhetünk implicit konverziót, illetve értékadást különböző minősítésű mutató típusú változók között.
11. Program szintű változót csak nevesített konstansként, __constant minősítővel hozhatunk létre.
A minősítők használatára vonatkozó szabályokat röviden a következő módon foglalhatjuk össze: címtér minősítők mindegyikét használhatjuk mutatókkal, azonban lokális változót csak __local vagy __private minősítővel deklarálhatunk, program szintű változót pedig csak a __constant minősítővel hozhatunk létre.
Kernel-függvények visszatérési értéke mindig void, minden mutató paraméterük __global, __constant vagy __local minősítésű kell legyen, nem használhatunk a formális paraméterlistán kétszeres indirekciót, speciális típusú változókat és függvénymutatókat. Egy mutató nem kaphat értékül más címtér minősítésű mutatót.
Könyvtári függvények. Az OpenCL C könyvtári függvényeinek az OpenCL C nyelvű kódokban történő használatához nem kell include direktívákat megadnunk. Az OpenCL C szabvány számos könyvtári függvényt specifikál, amelyek listája és részletes leírása megtalálható az OpenCL specifikációban. A következő táblázatban elsősorban az indexek meghatározására szolgáló függvényeket tekintjük át.
Függvény Leírás
uint get_work_dim(); Az clEnqueueNDRangeKernel függvény hívásakor definiált indextartomány dimenzióinak számát kérdezhetjük le vele.
uint get_global_size(uint dimindx); Visszatérési értéke a dimindx dimenzióban specifikált munkaelemek száma.
uint get_global_id(uint dimindx); Visszatérési értéke a munkaelem indexe a dimindx dimenzióban.
uint get_local_size(uint dimindx); Visszatérési értéke a munkacsoportok mérete dimindx dimenzióban.
uint get_local_id(uint dimindx); Visszatérési értéke a munkaelem lokális indexe a dimindx dimenzióban.
uint get_num_groups(uint dimindx); A dimindx dimenzió munkacsoportjainak számát kérdezhetjük le vele.
uint get_group_id(uint dimindx); A munkaelem munkacsoportjának dimindx dimenzióbeli indexét kérdezhetjük le.
uint get_global_offset(uint dimindx); Visszatérési értéke a globális indextartomány eltolása azaz a kezdőindex a dimindx dimenzióban.
A matematikai függvények zöme az ANSI C math.h header-jében definiált függvényekkel azonos néven és paraméterezéssel érhető el, de mellettük számos további matematikai függvény érhető el, melyek használata erősen javasolt.
Bár a printf függvényt az OpenCL 1.2 definiálja, működése implementáció függő, az implementáció által specifikált output stream-re kerül a printf függvény kimenete. Paramétereit tekintve megegyezik az ANSI C99 printf függvénnyel, visszatérési értéke azonban sikeres végrehajtás esetén 0, egyébként -1. Egy kernel függvény által meghívott printf függvények kimenetei a kernel végrehajtáshoz rendelt esemény befejezésekor kerülnek a kimenetre. A formátum sztring használatának fő különbsége, hogy az s konverziós konstans csak sztring literálok kiírásakor használható.
A függvények egy fontos csoportját képezik a szinkronizációs függvények, mint explicit szinkronizációs eszközök az kernel kódban. Ezek közül a barrier függvénnyel ismerkedünk meg, ezt ugyanis használni fogjuk az illesztett szűrés párhuzamosítása során. A függvény specifikációja:
Függvény Leírás
void barrier(cl_mem_fence_flags flags); A függvény az egy munkacsoporton belüli munkaelemek szinkronizációját szolgálja: egy munkacsoport munkaelemeinek futása megáll amíg a munkacsoport minden munkaeleme meg nem hívja a
Függvény Leírás
függvényt. Lehetséges paraméterei a
CLK_LOCAL_MEM_FENCE, CLK_GLOBAL_MEM_FENCE konstansok és ezek logikai ,,vagy'' művelettel képzett kombinációja.
A barrier függvény CLK_LOCAL_MEM_FENCE paraméter esetén a szinkronizációs ponton túl további funkcióval rendelkezik: megfelelő módon frissíti a lokális memória tartalmát, így az egyes munkaelemek a barrier pont után ugyanazon memóriaképet látják. CLK_GLOBAL_MEM_FENCE használata esetén hasonlóan jár el a barrier függvény a globális memóriával: az egyes munkaelemek globális memóriában végrehajtott módosításai alapján frissíti azt, így az egyes munkaelemek ugyanazon memóriaképet látják a szinkronizációs pont után.
A kernel függvényre már két példát is láthattunk a korábbiakban, egyrészt az üres kernelt (emptyKernel), másrészt a globális memóriába feltöltött számokból négyzetgyököt vonó sqrtKernel-t. A továbbiakban alkalmazás nélkül mutatunk be néhány egyszerű kernelt, amelyek klasszikus, jól párhuzamosítható feladatok megoldására használhatók. Megjegyezzük, hogy naiv, könnyen értelmezhető implementációkat mutatunk be, az OpenCL architektúrát alaposabban ismerő olvasó számos optimalizálási lehetőséget fedezhet fel bennük.
Az első példaprogram a gyökvonás egy variánsa, segítségével azt szemléltetjük, hogy több, és nem feltétlenül kernel függvény is szerepelhet egy OpenCL C nyelvű kódban. Működése a korábban bemutatott sqrtKernel függvényével azonos: lekérdezi a munkaelem (kernel-végrehajtás) globális indexét az első dimenzióban (feltételezve, hogy csak egydimenziós indextartományt használunk), majd gyököt von az index által specifikált adatelemből és a megoldást elhelyezi a második paraméterként kapott tömb megfelelő helyén.
7.18. példa - sqrtKernel.k
float gyok(float a) {
return sqrt(a);
}
__kernel void sqrtKernel(__global float* a, __global float* b) {
int i= get_global_id(0);
b[i]= gyok(a[i]);
}
A második példa a klasszikus, brute-force prímtesztet valósítja meg: a munkaelem indexét sorra osztjuk az egész számokkal annak gyökéig és ha maradék nélküli oszthatóságot tapasztalunk, az eredményeket tartalmazó b tömb i. elemét logikai hamisra állítjuk.