A példaprogramunk egy egyszerű számláló növelést old meg.

A keret úgy indul, hogy lefoglalunk egy float változót a GPU memóriában, aminek értéket adunk.

int* x = new int[1];
    int* dev_x;

    int res_cpu;
    int* res_gpu = new int[1];

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

    *x = 4;

    cudaMemcpy(dev_x, x, sizeof(int), cudaMemcpyHostToDevice);

Ezek után elindítunk egy kernelt, ami 4 szálon kezdi meg a számítást.

addKernel<<<2, 2>>>(dev_x);

Eddig még akár rendben is lehetne a program. Nézzük viszont meg a kernel.

__global__ void addKernel(int *x)
{
    x[0]++;
    return;
}

A kernelben megadott kód csak a változó értékét növeli egyel. Látható viszont, hogy - az eddig megszokottól eltérően - minden szál ugynazt az értéket módosítja. Ez vezet a konfliktushoz.

Ha kiszámítjuk, hogy mit csinál a programunk, akkor rájöhetünk, hogy egy 4-es értéket növelünk 4-szer, vagyis az eredmény elvileg 8 lenne. A program a legtöbb esetben viszont 5-öt fog visszaadni, bár még ez sem garantált.

Ha jobban megnézzük mi történik, akkor az alábbi folyamatot láthatjuk.

1. Minden szál elindít egy olvasást a globális memóriából. Ezek a műveletek egyszerre indulnak, tehát minden szál a 4-es értéket ovassa.
2. A szálak a 4-es értéket növelik 1-el tehát 5-öt kapnak.
3. Az eredményt elkezdjük visszírni a memóriába. Mivel minden szál az 5-ös értéket próbálja kiírni ezért valamelyik 5-ös szám kerül a memóriába. Azt nem tudhatjuk, hogy melyik, de ez most nem is probléma.

A práhuzamos 1-el növelés eredménye tehát 5.

A helyzet viszont sokkal bonyolultabb is lehet, ha több szálat indítunk. Azt tudjuk, hogy a GPU-n több szálat idíthatunk, mint ahány multiprocesszor van. Ilyenkor a GPU ütemezéje tölti be, és terminálja a szálakat. A szálak tehát nem teljesen párhuzamosan futnak, hanem vannak közöttük amelyek valamilyn módon párhuzamosan, és vannak olyanok amik egymás után futnak. Ha ez történik, akkor lehetnek szálak, amik csak az után indulnak el, hogy más szálak már végeztek a számítással, tehát némely szál 5-ös, 6-os, stb. értékeket olvashat attól függően, hogy hány szál felyezte be előttük a kiírást. Ez a sorrendiség pedig teljesen véletlenszerű, a GPU-től és az ötemezőjétől függ, így kiszámíthatatlan eredményre vezet.

Szerencsére az Nvidia mérnökei legalább annyit garantálnak, hogy az eredmény bitjei nem keverednek össze. A végső eredmény valamelyik szál kimenete lesz.

A következő fejeleztben megnézzük, hogy hogyan lehet kezelni a versenyhelyzeteket.

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

#include <stdio.h>

#define BLOCK_SIZE 16
#define GRID_SIZE  16

__global__ void addKernel(int *x)
{
    x[0]++;
    return;
}

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

    int* x = new int[1];
    int* dev_x;

    int res_cpu;
    int* res_gpu = new int[1];

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

    *x = 4;

    //  Számítás a GPU-n
    cudaMemcpy(dev_x, x, sizeof(int), cudaMemcpyHostToDevice);

    addKernel<<<2, 2>>>(dev_x);

    cudaMemcpy(res_gpu, dev_x, sizeof(int), cudaMemcpyDeviceToHost);

    //   Számítás a CPU-n
    res_cpu = x[0];
    for(i=0; i<BLOCK_SIZE*GRID_SIZE; i++) {
        res_cpu++;
    }

    printf("Result on GPU: %d\n", *res_gpu);
    printf("Result on CPU: %d\n", res_cpu);
    
    return 0;
}