A program kerete a korábbiakhoz nagyon hasonlít, így most csak a kernellel foglalkozunk.

A közzös memória definiálása a kód elején van __shared__ előtaggal. Az így megadott változókat a fordító közösként kezeli, és a szálak közsen használhatják.

    __shared__ float shared[BLOCKDIM][BLOCKDIM];

Ebben az esetében pont akkora közös memóriát használtunk, amekkora a blokk.

A következő lépés a közös memória kinullázása. Mivel annyi szál van ahány elem a meóriában ezt megtehetjük párhuzamosan úgy, hogy minden szál egy-egy memória elemet nulláz ki.

    shared[threadIdx.x][threadIdx.y] = 0;

    __syncthreads();

A nullázást egy szinkronizáció követi, hogy a további számítás előtt minden elem 0-s értékre álljon be.

Ezek után következhet a számítás. Ebbe a kódban egy szinuszos fehér foltot rajzounk minden 32x32-es négyzetbe, de bármi mást is rajzolhatnánk.

    shared[threadIdx.x][threadIdx.y] = 255 * sinf(threadIdx.x * PI / blockDim.x) * sinf(threadIdx.y * PI / blockDim.y);

A kódnak ezena pontján kellene szinkronizálni. A programban szerepel is a megfelelő függvényhívás, de azt kikommenteztük.

Végül pedig az eredmények kiírása következik a globális memóriába.

    result[tid + 1] = shared[threadIdx.y][threadIdx.x];
    result[tid + 0] = shared[threadIdx.y][threadIdx.x];
    result[tid + 2] = shared[threadIdx.y][threadIdx.x];
    result[tid + 3] = 255;

Itt észre kell venni, hogy minden szál egy másik szál áltel beírt értéket olvas a közös memóriából. Ezért a szinkronizáció hiányában lehet olyan szál, ami az eredeti 0-s értéket olvassa, és lesz olyan, ami az új számítás eredményét. A kimenet pedig a lenti képen látható mintázat.

A programot ki lehet próbálni úgy is, hogy visszatesszük a megfelelő helyen kikommentezett szinkronizációt.

A program érdekes pontja még, hogy a kiírásban nem volt ellenőrzés, hogy a kimeneti képen létezik-e az elérni kívánt képont. Ahogy korábban is itt is lehet, hogy nagyobb a blokstruktúra mint maga a kép, és vannak túllógó szálak. Ezeket viszont a kernel elején kezeltük. A túllógó szálak visszatértek, így a futásuk leállt.

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

#include "ImageIO.h"

#include <iostream>
#include <ctime>

using namespace std;

#define N           512
#define BLOCKDIM    32

#define PI 3.1415926535897932

__global__ void kernel(unsigned char* result, int w, int h);

int main(int argc, char** argv)
{
    unsigned char* result = new unsigned char[N*N * 4];
    unsigned char* dev_result;

    cudaMalloc((void**)&dev_result, N*N * 4 * sizeof(unsigned char));

    dim3 blockDim = dim3(BLOCKDIM, BLOCKDIM, 1);
    dim3 gridDim = dim3((N + BLOCKDIM - 1) / BLOCKDIM, (N + BLOCKDIM - 1) / BLOCKDIM, 1);

    kernel << <gridDim, blockDim >> > (dev_result, N, N);

    cudaMemcpy(result, dev_result, N*N * 4 * sizeof(unsigned char), cudaMemcpyDeviceToHost);

    writeRGBImageToFile("image.png", result, N, N);

    return 0;
}

__global__ void kernel(unsigned char* result, int w, int h)
{
    int x = blockDim.x *blockIdx.x + threadIdx.x;
    int y = blockDim.y *blockIdx.y + threadIdx.y;

    int tid = (y*w + x) * 4;

    __shared__ float shared[BLOCKDIM][BLOCKDIM];

    if (x >= w || y >= h)
        return;

    shared[threadIdx.x][threadIdx.y] = 0;

    __syncthreads();

    shared[threadIdx.x][threadIdx.y] = 255 * sinf(threadIdx.x * PI / blockDim.x) * sinf(threadIdx.y * PI / blockDim.y);

    //__syncthreads();

    result[tid + 1] = shared[threadIdx.y][threadIdx.x];
    result[tid + 0] = shared[threadIdx.y][threadIdx.x];
    result[tid + 2] = shared[threadIdx.y][threadIdx.x];
    result[tid + 3] = 255;

    return;
}