Ha a programunkban szeretnénk kihasználni a GPU lehetőségeit is, akkor GPU kódot is írnunk kell. A CUDA forráskódok általában keverten tartalmazzák a CPU, és GPU kódokat.

Az alábbi példaprogram GPU-n végzi el két szám összeadását.

#include <cuda.h>
#include <stdio.h>

__global__ void add( int a, int b, int* c )
{
    *c = a + b;

    return;
}

int main(int argc, char** argv)
{
    int c;
    int* dev_c;

    cudaMalloc((void**)&dev_c, sizeof(int) );

    add<<<1,1>>>(1, 2, dev_c);

    cudaMemcpy(&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost);

    printf("a + b = %d\n", c);

    cudaFree(dev_c);

    return 0;
}

A programban számos részlet, és úgy vezérlési forma szerepel, amelyeket a következőkben mutatunk be. A teljes forráskódfile Megtalálható itt.

A CUDA három fontos kiegészítést tesz a C++ szintaxisához.

Egyrészt bejön három új jelölő a függvények megkülönböztetésére.

• __host__ előtaggal jelöljük a CPU-n hutó normál kódot;
• __global__ előtagot kap a GPU-n futó kód, amelyet a CPU kódból meghívhatunk;
• __device__ jelölővel illetjük a függvényeket, amik a GPU-n futnak, és amit csak a GPU-ról hívhatunk.
• Ha valaminek nem adunk jelölőt, akkor automatikusan __host__ előtagot feltételez a fordító.

A fenti példában volt két függvény

__global__ void add( int a, int b, int* c ) {
    ...
}

int main(int argc, char** argv)
{
    ...
}

Az első add(...) függvény ezek szerint a GPU-n fut. A main függvény előtag nélküli, tehát a CPU-n indul el.

A programokkal kapcsolatba alapszabály, hogy a GPU nem képes kezelni a program teljes futását. A GPU kódjaink általában a CPUn futtató keretprogramban futnak. A keret végzi a program koordinálását, az adatok előksézítését, és amikor olyan feladathoz ér, amit a GPU fel tud gyorsítani, akkor azt átadja a GPU számára.

A GPU-n futó számítást úgynevezett kernel függvények megadásával tudjuk megvalósítani. Ezek a __global__ előtaggal rendelkező függvények.

Amikor elindítjuk őket a kerlnelek több szálon kezdik meg a számítást. Ezt haszn áljuk ki a párhuzamos adatfeldolgozásra.

Ügyelni kell viszont, hogy a párhuzamos szálakra vonatkoznak korlátozások:

• A szálakban korlátozott a maximális utasítások száma. Egy szálban számítási képessőgtől függően csak néhány millió utasítás lehet.
• A szálak csak korlátozottan kommunikálhatnak egymással. Vannak szálak, amelyek úgynevezett blokkokban csoportokban dolgoznak, de alapvetően fel kell tennünk, hogy a szálak egymástól függetlenül futnak, és a futásuk sorrendjére nincs ráhatásunk. A száímtást ezért általában úgy kell megoldanunk, hogy a kernelből indított szálak teljesen függetlenül képesek legyenek működni.
• Figyelni kell az adatok kiírásakor. A szálak párhuzamossága miatt nem tudni (nem ddefiniált) mi történik, ha két szál ugyanazt a memőriaterületet próbálja írni. Ezért ezt általában kerüljük is. Van mód az ütemezés áthidalására, de az költségekkel jár (Erről részletesebben az Atomikus műveletek fejezetben beszélünk).

A fenti megfontolások miatt a GPU számára legjobban megfelelő a nagyban párhuzamosítható számítások programozása úgynevezett Gather (begyűjtéses) megközelítésben. Ez azt jelenti, hogy a probléma kimenetét osztjuk fel darabokra. Például egy kép feldolgozásakor a kimenti kép pixeleit állítjuk elő. Ilyenkor a kimenet minden eleméhez (minden eredmény pixelhez) hozzárendelünk egy szálat. A szál indítása után begyűjti az eredmény pixel kiszámításához szükséges adatokat, elvégzi a számításokat, és kiírja az eredményt. Mindezt ideális esetben az összes szál teljesen függetlenül végzi.

A kernelek indítására van egy új direktívánk. Az ilyen függvényeket a CPU kódből indítjuk a függvény paraméterlistája előtt viszont beékelünk egy második paraméterlistát a <<<X, Y>>> hármas relációs jellel adott zárójelek közé. Ez fogja megadni, hogy az adott kernelt hány szálon fogjuk elindítani.

Esetünkben az elábbi kód egy szálat indít.

    add<<<1,1>>>(1, 2, dev_c);

A GPU programozásánál külön kell foglalkoznunk a memória kezelésével is. A GPU-nak ugyanis saját memóriája (grafikus memória) van, és a kernelek csak ebből a memóriából képesek adatokat olvasni, és csak ide képesek adatokat írni.

A GRafikus memória kezelésére a CUDA függvénykönyvtárban külön függvények vannak. A memória fogalalására a cudaMalloc függvényt használhatjuk. Formája:

• void cudaMalloc(void** devPtr, size_t size): lefoglal számunkra a GPU-n egy memóriaterületet. Paraméterei:
  • devPtr: Pointer a lefoglalt memória címének;
  • size: A lefoglalni kívánt memóriaterület mérete.

A cuda Malloc parancs segítségével általában minden foglalunk egy területet a GPU-n minden szükséges adatnak:

cudaMalloc((void**)&dev_c, sizeof(int) );

Ha le van foglalva a memória, akkor át kell tenni a feldolgozni kívánt adatokat a GPU-ra.el kell végezni a számítást, majd az eredményeket vissza kell másolni a központi memóriába.

A memória másolásához használhatjuk a cudaMemcpy() függvényt.

• cudaMemcpy( void* dst, const void* src, size_t count, enum cudaMemcpyKind kind): adatot másol két memóriaterület között. Paraméterei:
  
  • dst: A cél terület első byte-jára mutató pointer;
  • src: A forrás terület első byte-jára mutató pointer;
  • count: Másoln kívánt bype-ok száma;
  • kind: Másolás iránya. Ez a paraméter meghatrozza, hogy a központi, és a grafikus memória között milyen irányba másolunk. Értéke lehet:
    • cudaMemcpyHostToHost,
    • cudaMemcpyHostToDevice,
    • cudaMemcpyDeviceToHost,
    • cudaMemcpyDeviceToDevice.

A fenti kódrészletben a bemeneti adatokat nem kell átmásolni a kernelnek. A keretrendszer a függvények paramétereit automatikusan átteszi a grafikus memóriába. A visszatérési értéket viszont a kernel csak a GPU RAM-ba tudja menteni így azt utólag kell visszamásolni.

cudaMemcpy(&c, dev_c, sizeof(int), cudaMemcpyDeviceToHost);

Az adatkezelés végére már csak a takarítás marad. A lefoglalt memóriaterületek felszabadításáról ugyanis nekünk kell gondoskodni a cudaFree() függvénnyel.

• cudaFree(void* ptr): Felszabadít egy megadott memóriaterületet a grafikus memóriában. Paramétere:
  
  • ptr: A felszabadítandó memóriaterületre mutató pointer.