A GPU-n indított práhuzamos kódokban aszálak struktúráját a függvény és paraméterlista között megadott <<<rácsméret, blokkméret>>> jelölés adja meg.

Mind a két paraméter dim3 tílusú, tehát a rács és blokk is lehet 3-dimenziós.

A korábbi többszálú példána van egy nagy problémája. Nevezetesen, hogy rácsonként csak egy szál van elindítva. Ez azért problémás, mert a multiprocesszorokban az ALU-k csoportokban egyszerre ugyanazt a műveletet hajtják végre egy blokk szálait futtatva. A csoportok mérete a jelenlegi GPU-kon 32 ALU, ami práhuzamosan fut. Tehát ha egy blokkban csak 1 szál van, akkor a 32 ALU-ból egy dolgozik, a többi pedig tétlenül áll.

A lenti példaprogramban ezt oldjuk meg. A teljes példaprogram forráskódja Letölthető innen.

A vektorok összeadásának módosított teljes kódja itt lesz megadva. A kontextur kedvéért a teljes kód egyben lentebb megtalálható.

Az első feladatunk, hogy a kernel hívásakor módosítanunk kell a szálstruktúrát, hogy az egy megfelelő rácsban (grid) megfelelő blokkokkal (block) működjön.

Ehhez először is meg kell adnunk a blokkok méretét. Ez általában egy konstans szám, amit több tényező figyelembevételével kell kialakítanunk.

• A blokkok méretének 32 egész többszörösének kell lennie. Erre azért van szükség, mert a GPU-ban a szálakat 32-es csoportokban futtatjuk úgy, hogy 32 szálító egység egyszerre ugyanazt végzi. Ha kevesebb szálat definiálunk, akkor üresen maradt kihasználatlan szálaink lesznek.
• A blokkok méretének minél nagyobbnak kell lennie, de nem túl nagynak. Ez most még kicsit ködös megfogalmazás, de a magyarázot csak a későbbi anyagrészben a memóriakezelénél kapunk. Irányelvnek tipikusan 256-512 körüli számot érdemes megadni. Ha túl nagy számot adunk meg a kernel hívás hibával eláll.
• Figyelni kell, hogy legyen elég blokk. Ez a szempont is a blokkonkénti szálak számával függ össze, csak fordítva mint az előző irányelv. Úgy érdemes méretezni a blokkokat, hogy a rácsban sok blokk legyen (vagyis jó a kis blokk). Ennek magyarázata, hogy a korábban megadott GPU architektúrával függ össze. A GPU-ban az ALU-k multiprocesszorokban vannak, és minden multiprocesszor kap blokkokat végrehajtásra. Ha kecé a blokk, akkor némely ALU feladat nélkül maradhat.

Esetünkben a fenti feladatok mindegyikét nem tudjuk teljesíteni, mert az első példaprogramunk csak kevés adattal dolgozik. Innen is látszik, hogy a GPU-t akkor érdemes használni, ha nagy számú párhuzamos számítás van, különben nem használjuk ki a teljes számítási kapacitást.

De a példa okáért most megfelelő a 32-es blokkok használata is.

#define BLOCKDIM    32

A következő feladatunk, hogy ki kell számítanunk a rács méretét.

Fontos még, hogy annyi blokkot kell használnunk, hogy a feladatok a blokkok között szétosztva lefedjék az összes feldolgozandó lemet. Ehhel lesztjuk az elemek számát a blokkok szálainak számával, és az eredményt felfelé kerekítjük. Egész aritmetikával az alábbi módszer célravezető:

addKernel << <(N + BLOCKDIM - 1) / (BLOCKDIM), BLOCKDIM >> >(dev_a, dev_b, dev_c);

Itt meg kell jegyeznünk, hogy ha 100 vektor elemet dolgounk fel 32-es blokkmérettel, akkor lesznek szálak amik nem kapnak feladatot. Ez előfordul a programok írásakor, de a mivel sok a szempont nem tudunk mindennek megfelelni. Általános szempont, hogy a GPU kihasználása minél nagyobb legyen (tehát kevés legyen az üresen indított szál).

Végül pedgi módosítanunk kell a kernelen is, hogy az kezelni tudja az új szálstruktúrát. Ez az alábbi módon történik:

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

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

    return;
}

Az új kernelünkben először le kell kérdeznünk, hogy a szál hol található a szálstruktúrában. Ehhez használhatjuk a blockIdx, és threadIdx beépített változókat, amik megadják a szál blokkjának és indexének helyét. Ezek az értékek dim3 típusúak, így (x,y,z) koordinátákban egy 3D rácsban tudunk pozíciót lekérdezni, de nálunk csak az "x" koordináta van használatban.

Szükségünk van még a blokk méretére, amit megkapunk a blockDim.x beépített változóból. Vagyis kiszámíthatjuk, hogy a blokkunk hanyadik szálnál kezdődik, és hogy azon belül hanyadik szálat dolgozzuk fel éppen.

Ez után lellenőrizzük, hogy az aktuálisan futó szához tartozik-e adat. Emlékezzünk vissza, hogy lehet, hogy több szálat indítunk, mint amennyi adatunk van tehát e nélkül túlindexelhetjük a memóriát.

Vlgül pedig elvégezzük a számítást.

A vektorok összeadásának teljes kódja a lenti példaprogram.

A részletes magyrázat lentebb található.

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

#include <iostream>
#include <ctime>

using namespace std;

#define N           100
#define BLOCKDIM    10

__host__ void addOnGPU(int* a, int* b, int* c);
__host__ void addOnCPU(int* a, int* b, int* c);

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

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

    int* a = new int[N];
    int* b = new int[N];
    int* c = new int[N];

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

    addOnGPU(a, b, c);
    addOnCPU(a, b, c);

    return 0;
}

__host__ void addOnGPU(int* a, int* b, int* c)
{
    int* dev_a;
    int* dev_b;
    int* dev_c;

    int start_gpu_full = clock();

    cudaMalloc((void**)&dev_a, N * sizeof(int));
    cout << cudaGetErrorString(cudaGetLastError()) << endl;
    cudaMalloc((void**)&dev_b, N * sizeof(int));
    cout << cudaGetErrorString(cudaGetLastError()) << endl;
    cudaMalloc((void**)&dev_c, N * sizeof(int));
    cout << cudaGetErrorString(cudaGetLastError()) << endl;

    cudaMemcpy(dev_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
    cout << cudaGetErrorString(cudaGetLastError()) << endl;
    cudaMemcpy(dev_b, b, N * sizeof(int), cudaMemcpyHostToDevice);
    cout << cudaGetErrorString(cudaGetLastError()) << endl;

    cudaDeviceSynchronize();

    int start_kernel = clock();

    addKernel << <(N + BLOCKDIM - 1) / (BLOCKDIM), BLOCKDIM >> >(dev_a, dev_b, dev_c);
    cout << cudaGetErrorString(cudaGetLastError()) << endl;

    cudaDeviceSynchronize();

    cout << "Kernel Excution time: " << clock() - start_kernel << endl;

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

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

    cout << "Full GPU time: " << clock() - start_gpu_full << endl;

    int correct = 1;
    for (int i = 0; i<N; i++)
    {
        if (a[i] + b[i] != c[i])
        {
            correct = 0;
            break;
        }
    }

    if (correct) {
        cout << "GPU: A szamitas helyes. :)" << endl;
    }
    else {
        cout << "GPU: A szamitas helytelen. :(" << endl;
    }

}

__host__ void addOnCPU(int* a, int* b, int* c)
{
    int start_cpu = clock();

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

    cout << "Full CPU time: " << clock() - start_cpu << endl;

    int correct = 1;
    for (int i = 0; i<N; i++)
    {
        if (a[i] + b[i] != c[i])
        {
            correct = 0;
            break;
        }
    }

    if (correct) {
        cout << "CPU: A szamitas helyes. :)" << endl;
    }
    else {
        cout << "CPU: A szamitas helytelen. :(" << endl;
    }

    return;
}


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

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

    return;
}