A program ebben az esetben is a main függvényben indul.

A korábbiakhoz képest ez a program már komplexebbnek tűnhet. Ennek fő oka, hogy számos előkészítő lépést tartalmaz, amik igazából nem bonyolultak, csak időigényel a leírásuk.

A textúrák használatához először is még a globális scope-ban definiálnunk kell a textúra leírókat.

texture<float4, cudaTextureType2D, cudaReadModeElementType> tex_in;
texture<float4, cudaTextureType2D, cudaReadModeElementType> tex_out;

Erre azért van szükség, mert a CUDA-ban egyszerűen így van definiálva a textúrák előkészítése. A textúrázó egy generikus objektum, aminek három paraméter kell megadnunk.

• Az bemeneti képen a pixelek intenzitások típusát;
• A textúra típusát (1-, 2-, vagy 3-D);
• Illetve az olvasás közben megadott típus konverziót.

Következőnek nézzük meg a main függvényt!

Az első rész, ami korábbi példáktól eltérő lehet, az az adatok előkészítése. Ezt több lépésben kell megoldanunk, miközben két új függvényt használunk.

    unsigned char* image;
    int h, w;

    readRGBImageFromFile("Lena.png", image, w, h);

    //  tömbök a GPU-nál
    unsigned char* dev_byte_ptr;
    float* dev_src;
    float* dev_dst;

    //  grafikus memóriaterületek foglalása
    cudaMalloc((void**)&dev_byte_ptr, 4 * h * w * sizeof(unsigned char));
    cudaMalloc((void**)&dev_src, 4 * h * w * sizeof(float));
    cudaMalloc((void**)&dev_dst, 4 * h * w * sizeof(float));

    //  kiinduló kép másolás a grafikus memóriába
    cudaMemcpy(dev_byte_ptr, image, 4 * h * w * sizeof(unsigned char), cudaMemcpyHostToDevice);

    //  kép konvertálás float formátumra
    convByte2FloatGPU(dev_src, dev_byte_ptr, 4*h*w);

Az adatok előkészítésétben először beolvassuk a képet a readRGBImageFromFile() függvénnyel.

A képen viszont az adatok byte-os formában vannak amit float-á kell konvertálni. Itt jön a képbe a GPU, mivel ha belegondolunk a konverziót el tudjuk végezni a GPU-n is. A byte-os pixel intenzitásokat átmásoljuk a GPU-ra, majd a convByte2FloatGPU(...) függvényt használva átkonvertáljuk azokat float típusúvá. Erről a függvényről még később részletesebben írunk.

Az adatok előkészítése után a következő lépés a textúrázó előkészítése. Ez viszonylag egyszerűen megy, csak létre kell hoznunk egy csatornaleírót, ami a textúrázó paramétereit adja meg, és hozzá kell kötnünk a globális scope-ban definiált textúrázókhoz a egfelelő memóriaterületeket.

    cudaChannelFormatDesc chDesc = cudaCreateChannelDesc();
    cudaBindTexture2D(NULL, tex_in, dev_src, chDesc, w, h, 4*w*sizeof(float));
    cudaBindTexture2D(NULL, tex_out, dev_dst, chDesc, w, h, 4*w*sizeof(float));

Majd később látjuk, hogy miért is, de ezen a ponton két textúrázót készítünk. Egyet a bemeneti képhez, a másikat pedig a kimeneti képhez kötjük hozzá.

Következőnek jöhet maga aszámítás.

    for(i=0; i<ITERATION_COUNT; i++) {
        if(i%2) {
            //  ping
            //  páratlan cilkusban a 2. textúrából olvasunk és az eredményt az 1.-be írjuk
            smoothKernel<<<grid, block>>>((float4*)dev_src, w, h, 1);
        } else {
            //  pong
            //  páros cilkusban az 1. textúrából olvasunk és az eredményt a 2.-ba írjuk
            smoothKernel<<<grid, block>>>((float4*)dev_dst, w, h, 0);
        }
    }

A számítás iteratívan meg végbe többször, amit a smoothKernel<<<grid, block>>>(...) függvény segítségével végzünk. Itt észre kell venni, hogy minden konvolúciónak két memóriája van egy kimenet és egy bemenet. A következő iterációban viszont az előző interció kimenetén kell tovább dolgozni, és kelleni kell egy újabb kimenet, ami lehet az előző iteráció bemeneti memóriája.

Az interációban tehát folyamatosan cserégetjük a bemenetet és a kimenetet, és a konvolúciós kernelt éis kétféleképpen használjuk. Az egyik iterációban az első memóriából olvasul az első textúrázón keresztül, és követlenül a második memóriába írunk. A másik iterációban pedig fordítva a második memóriából olvasunk a második textúrázón keresztül, és közvetlenül az első memóriába írunk.

Ez nem okoz gondot, mert - bár a textúrázó csak olvasható, rajta keresztül nem tudunk írni a memóriába - a textúrázó alatti memória közvetlenül még mindig írható.

A kernel ennek megfelelően megkapja a cél memória pointerét, a kép szélességét, és magasságát, illetve egy 0 vagy 1-es paramétert attól függően, hogy melyik textúrázóból olvas.

Végül a számítás után az eredményben a pixelek típusát visszaalakítjuk byte-á; az eredményt pedig visszamásoljuk a központi memóriába és kiírjuk a lemezre.

    convFloat2ByteGPU(dev_byte_ptr, dev_src, 4*w*h);

    cudaMemcpy(image, dev_byte_ptr, 4*w*h*sizeof(unsigned char), cudaMemcpyDeviceToHost);
    
    writeRGBImageToFile("out.png", image, w, h);

Következőnek nézzük meg a kernel felépítését.

Ahogy említttük a kernel négy paramétert vár: a kimenetre mutató pointert, a kép szélességét, és magasságát, illetve egy 0/1 válzotót ami meghatározza, hogy melyik textúrából olvasunk.

A kernel a korábban megszokottnak megfelelően a kiszámított képpont koordináták meghatározásával kezd.

    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

Ezek után jöhet a számítás. Ehhez egy dupla for ciklus összeadja a képpont 3*3-as környezetében a pixel intenzitásokat, majd majd osztunk a pixelek számával.

    //  pixel értékeinek előkészítése
    pixel.x = 0;
    pixel.y = 0;
    pixel.z = 0;
    pixel.w = 0;

    //  bejárjuk a 3x3-as szomszédságot
    for(i=-1; i<2; i++) {
        for(j=-1; j<2; j++) {
            //  páros lépésben a egyik textúrából olvasunk, páratlanban a másikból
            //  a kilógó pixelekkel nem foglalkozunk, a textúrázó lekezeli.
            if(odd) {
                temp = tex2D(tex_out, x+i, y+j);
            } else {
                temp = tex2D(tex_in, x+i, y+j);
            }
            pixel.x += temp.x;
            pixel.y += temp.y;
            pixel.z += temp.z;
            pixel.w += temp.w;
        }
    }

    //  osztás a maszk méretével
    pixel.x /= 9;
    pixel.y /= 9;
    pixel.z /= 9;
    pixel.w /= 9;

Ahogy a kódrészben is írjuk, a képről kilógó pixelekkel nem kell foglalkozni azokat a textúrázó automatikusan kezeli.

Ugyanígy nem kell foglalkozni a rgb érteékek beolvasásával sem, azokat is automatikusan együtt kezeli a textúrázó.

végül pedig minden bonyolítás nélkül kiírjuk az eredményet a cél tömb megfelelő helyére. A pixeleket most egyszerre írjuk ki minden csatornájukkal együtt, így az indexelés is egyszerűbb egy kicsit.

Van még két függvény, amiről eddig nem értekeztünk részletesebben. Az egyik byte-os tömböt konvertál float-ossá, a másik pedig fordítva.

Ezek a függvények külön lettek implementálva, hogy a fő kódunk tisztább maradjon. Kiszervezésük már csak azzért is logikus, mert tipikusan sokszor használt műveletek, így ha kell ki tudjuk tenni őket külön fileba, és egy nagyobb projekt bármely műveletéből használhatjuk őket.

A byte -> float konverzió egyszerű feladat. Ebben a vektor összeadáshoz hasonló struktúrát írunk, ahol a pixeleket egy lineáris tömb-ként dolgozzuk fel. Ha belegondolunk itt nincs is szükség rá, hogy kihasználjuk a kép struktúráját, vagy az rgb színcsatornák meglétét. Elég ha úgy gondolunk az adatra, mint egy csimó byte értékre amiket át kell konvertálni float-á.

Ehhez kiszámítunk egy lináris indítási struktúrát, ami hasonlít a belső szorzatos példaprogramunkra, és elindítjuk fele a emgfelelő kernelt.

__host__ void convByte2FloatGPU(float* dev_dst, unsigned char* dev_src, int count) {
    //  csak kiszámítjuk az indított rács strukturáját, és elindítuk a megfelelő kernelt.
    dim3 grid;
    dim3 block;
    grid.x = min(DIM_GRID_LIN, (count + DIM_BLOCK_LIN - 1)/DIM_BLOCK_LIN);
    block.x = DIM_BLOCK_LIN;

    byte2FloatKernel<<<grid, block>>>(dev_dst, dev_src, count);

    return;
}

A kernel hasonlóan egyszerű. Meglehetősen hasonlít a belsőszorzatot számító kernelhez, csak nem kell bele a kód végi részösszegek kezelése, és a fő művelet is egy egyszerű érték másolás.

__global__ void byte2FloatKernel(float* dst, unsigned char* src, int count) {
    //  kiszámítjuk az első konvertálandó elem pozícióját.
    int tid = blockDim.x * blockIdx.x + threadIdx.x;

    //  majd ciklikusan
    while(tid < count) {
        //  konvertálunk (csak egy másolás)
        dst[tid] = src[tid];
        //  és lépünk a többi szál által konvertált értékeken túlra
        tid += blockDim.x * gridDim.x;
    }

    return;
}

A másik irányú konverzióhoz nem kell sok magyarázat. Ugyanúgy működik, mint az előző kódrészlet.

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <cuda.h>

#include <iostream>
#include <cstdio>
#include <ctime>
#include <cmath>
#include <algorithm>

#include "ImageIO.h"

#define DIM_BLOCK_LIN   256
#define DIM_GRID_LIN    1024
#define DIM_BLOCK_2D    16

#define ITERATION_COUNT 100

using namespace std;

texture<float4, cudaTextureType2D, cudaReadModeElementType> tex_in;
texture<float4, cudaTextureType2D, cudaReadModeElementType> tex_out;

//------------------------------------------------------------------------------
// Kernel Definíciók

// byte <-> float konverziókra
__global__ void byte2FloatKernel(float* dst, unsigned char* src, int count);
__global__ void float2ByteKernel(unsigned char* dst, float* src, int count);

// Kép simításra
__global__ void smoothKernel(float4* out, int w, int h, int odd);

//------------------------------------------------------------------------------
// CPU-n futó segédfüggvények
__host__ void convByte2FloatGPU(float* dev_dst, unsigned char* dev_src, int count);
__host__ void convFloat2ByteGPU(unsigned char* dev_dst, float* dev_src, int count);



//------------------------------------------------------------------------------
// main
int main()
{
    int i;
    unsigned int time;

    unsigned char* image;
    int h, w;

    readRGBImageFromFile("Lena.png", image, w, h);

    //  egy kics időmérés
    time = clock();

    //  tömbök a GPU-nál
    unsigned char* dev_byte_ptr;
    float* dev_src;
    float* dev_dst;

    //  grafikus memóriaterületek foglalása
    cudaMalloc((void**)&dev_byte_ptr, 4 * h * w * sizeof(unsigned char));
    cudaMalloc((void**)&dev_src, 4 * h * w * sizeof(float));
    cudaMalloc((void**)&dev_dst, 4 * h * w * sizeof(float));

    //  kiinduló kép másolás a grafikus memóriába
    cudaMemcpy(dev_byte_ptr, image, 4 * h * w * sizeof(unsigned char), cudaMemcpyHostToDevice);

    //  csak kiírjuk az időt
    printf("Kep beolvasva: %d s\n", clock()-time);

    //  kép konvertálás float formátumra
    convByte2FloatGPU(dev_src, dev_byte_ptr, 4*h*w);

    //  textúrák előkészítése
    cudaChannelFormatDesc chDesc = cudaCreateChannelDesc();
    cudaBindTexture2D(NULL, tex_in, dev_src, chDesc, w, h, 4*w*sizeof(float));
    cudaBindTexture2D(NULL, tex_out, dev_dst, chDesc, w, h, 4*w*sizeof(float));

    //  még egy idő kiírás
    printf("Kep konvertalva es elokeszitve: %d s\n", clock()-time);

    //  blokkméret előkészítése
    dim3 grid((w + DIM_BLOCK_2D - 1)/DIM_BLOCK_2D, (h + DIM_BLOCK_2D - 1)/DIM_BLOCK_2D);
    dim3 block(DIM_BLOCK_2D, DIM_BLOCK_2D);

    //  iteratív simítás.
    for(i=0; i<ITERATION_COUNT; i++) {
        if(i%2) {
            //  ping
            //  páratlan cilkusban a 2. textúrából olvasunk és az eredményt az 1.-be írjuk
            smoothKernel<<<grid, block>>>((float4*)dev_src, w, h, 1);
        } else {
            //  pong
            //  páros cilkusban az 1. textúrából olvasunk és az eredményt a 2.-ba írjuk
            smoothKernel<<<grid, block>>>((float4*)dev_dst, w, h, 0);
        }
    }

    //  Bevárjuk, hogy a kernelek lefussanak
    //  csak az időmérés miatt kell. a CPU futása aszinkrom megy elindítja a kerneleket
    //    és megy tovább amíg nem kell az általuk szolgáltatott adat.
    cudaThreadSynchronize();

    printf("Iterativ simitas kesz: %d s\n", clock()-time);

    //  ha kész van az eredmény, akkor viszakonvertáljuk byte-tömbbé
    convFloat2ByteGPU(dev_byte_ptr, dev_src, 4*w*h);

    //  és visszamásoljuk a központi memóriába
    cudaMemcpy(image, dev_byte_ptr, 4*w*h*sizeof(unsigned char), cudaMemcpyDeviceToHost);
    
    //  utolsó időmérés
    printf("Kep visszalakitva: %d s\n", clock()-time);

    writeRGBImageToFile("out.png", image, w, h);

    //  program vége, memóriafelszabadítás
    cudaFree(dev_byte_ptr);
    cudaFree(dev_src);
    cudaFree(dev_dst);

    return 0;
}

//------------------------------------------------------------------------------
// Kernel megvalósítások

// a simító kernel
__global__ void smoothKernel(float4* out, int w, int h, int odd) {
    //  koordináták kiszámítása
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    float4 pixel;
    float4 temp;
    int i, j;

    //  ha a szél "lelóg" a képről akkor leálítjuk
    if(x>=w || y>=h)
        return;

    //  pixel értékeinek előkészítése
    pixel.x = 0;
    pixel.y = 0;
    pixel.z = 0;
    pixel.w = 0;

    //  bejárjuk a 3x3-as szomszédságot
    for(i=-1; i<2; i++) {
        for(j=-1; j<2; j++) {
            //  páros lépésben a egyik textúrából olvasunk, páratlanban a másikból
            //  a kilógó pixelekkel nem foglalkozunk. a textúrázó lekezeli.
            if(odd) {
                temp = tex2D(tex_out, x+i, y+j);
            } else {
                temp = tex2D(tex_in, x+i, y+j);
            }
            pixel.x += temp.x;
            pixel.y += temp.y;
            pixel.z += temp.z;
            pixel.w += temp.w;
        }
    }

    //  osztás a maszk méretével
    pixel.x /= 9;
    pixel.y /= 9;
    pixel.z /= 9;
    pixel.w /= 9;

    //  eredmény kiírása
    out[y*w+x] = pixel;

    return;
}

//  byte -> float konverzió
__global__ void byte2FloatKernel(float* dst, unsigned char* src, int count) {
    //  kiszámítjuk az első konvertálandó elem pozícióját.
    int tid = blockDim.x * blockIdx.x + threadIdx.x;

    //  majd ciklikusan
    while(tid < count) {
        //  konvertálunk (csak egy másolás)
        dst[tid] = src[tid];
        //  és lépünk a többi szál által konvertált értékeken túlra
        tid += blockDim.x * gridDim.x;
    }

    return;
}

//  float -> byte konverzió
__global__ void float2ByteKernel(unsigned char* dst, float* src, int count) {
    int tid = blockDim.x * blockIdx.x + threadIdx.x;

    while(tid < count) {
        dst[tid] = src[tid];
        tid += blockDim.x * gridDim.x;
    }

    return;
}

//------------------------------------------------------------------------------
// CPU segédfüggvény megvalósítások

// keret a byte->float konverzióhoz
__host__ void convByte2FloatGPU(float* dev_dst, unsigned char* dev_src, int count) {
    //  csak kiszámítjuk az indított rács strukturáját, és elindítuk a megfelelő kernelt.
    dim3 grid;
    dim3 block;
    grid.x = min(DIM_GRID_LIN, (count + DIM_BLOCK_LIN - 1)/DIM_BLOCK_LIN);
    block.x = DIM_BLOCK_LIN;

    byte2FloatKernel<<<grid, block>>>(dev_dst, dev_src, count);

    return;
}

// keret a byte->float konverzióhoz
__host__ void convFloat2ByteGPU(unsigned char* dev_dst, float* dev_src, int count) {
    dim3 grid;
    dim3 block;
    grid.x = min(DIM_GRID_LIN, (count + DIM_BLOCK_LIN - 1)/DIM_BLOCK_LIN);
    block.x = DIM_BLOCK_LIN;

    float2ByteKernel<<<grid, block>>>(dev_dst, dev_src, count);

    return;
}