Az előző ismereteink alapján a CUDA-ban a GPU kód párhuzaomsan több szálon futtatott programot jelent. A korábbi kód viszont még csak egyetlen szálon végzett egy összeadást. Nézzünk meg egy olyan kódrészletet, amely már több szálon fut.

A feladata mit elvégzünk ebben az esetben egy vektor összeadás. Tegyük fel, hogy adott két vektor a és b a feladat hogy számítsuk ki a két vektor

c = a + b

összegét. Mivel a három vektor mérete megegyezik a feladatot elvégezhetjük elemenként úgy, hogy minden elempár összeadására egy-egy szálat indítunk, ahogy a következő ábrán is látható.

A kernel indításakor tehát annyi elemet szálat kell indítani, ahány elempár van.

A megoldás teljes kódja Letölthető innen.

A fenti feladat megoldás itt látható. A részlete magyarázat lentebb olvasható.

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

#define N   10

__global__ void addKernel(int* a, int* b, int* c);

int main(int argc, char** argv)
{
    int i;

    int a[N];
    int b[N];
    int c[N];

    int* dev_a;
    int* dev_b;
    int* dev_c;

    cudaMalloc((void**)&dev_a, N * sizeof(int));
    cudaMalloc((void**)&dev_b, N * sizeof(int));
    cudaMalloc((void**)&dev_c, N * sizeof(int));

    for(i=0; i<N; i++)
    {
        a[i] = -i;
        b[i] = i * i;
    }

    cudaMemcpy(dev_a, a, N*sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, N*sizeof(int), cudaMemcpyHostToDevice);

    addKernel<<<N, 1>>>(dev_a, dev_b, dev_c);

    cudaMemcpy(c, dev_c, N*sizeof(int), cudaMemcpyDeviceToHost);

    for(i=0; i<N; i++)
    {
        printf("%d + %d = %d\n", a[i], b[i], c[i]);
    }

    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

    return 0;
}

__global__ void addKernel(int* a, int* b, int* c)
{
    int tid = blockIdx.x;

    if(tid < N)
        c[tid] = a[tid] + b[tid];

    return;
}

A korábbi példaprogramhoz képest az egyik változás a kódban, hogy a bemeneti adatokat (vagyis a vektorokat) egy tömbben kell kezelni. Ennek megoldás már nem lehetséges egyszerű paraméter-ként a kernel fejlécében, tömböket kell foglalnunk és előkészíteni a grafikus memóriában.

Az első lépés a tömbök foglalása. Minden CPU-n szereplő adatnak le kell foglalni a GPU-n található megfelelőjét.

    int* dev_a;
    int* dev_b;
    int* dev_c;

    cudaMalloc((void**)&dev_a, N * sizeof(int));
    cudaMalloc((void**)&dev_b, N * sizeof(int));
    cudaMalloc((void**)&dev_c, N * sizeof(int));

A tömbök foglalása után az adatokat át kell másolni a grafikus memóriába.

    cudaMemcpy(dev_a, a, N*sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, N*sizeof(int), cudaMemcpyHostToDevice);

Ekkor következik a számítás maga. Végül pedig a program végén a foglalt területek felszabadítása a felat.

    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

A többszál programozás másik lépése, hogy a kernelt több példányban - több szálon - kell elindítani.

Ennek legyegyszeűbb módja, hogy a kernel hívásban az első paraméter megadásával több számítási blokkot indíthatunk, amelyek párhuzamosan futnak.

    addKernel<<<N, 1>>>(dev_a, dev_b, dev_c);

FONTOS! Majd később látjuk, hogy ez ellenjavattoll módszer, mert nem túl hatékony, de első nekifutásnak megfelel.