Skoči na vsebino

Seštevanje vektorjev1

Začnimo s preprostim primerom seštevanja dveh vektorjev. Problem bomo rešili na dva načina. Najprej bomo začeli dokaj naivno in se osredotočili le na osnove programiranja GPE ter na korake na gostitelju, ki so potrebni, da naš prvi program zaženemo na GPE. V drugem načinu se bomo nato osredotočili na izboljšavo ščepca, ki se izvaja na GPE.

Spodnja koda prikazuje implementacijo funkcije v jeziku C za seštevanje dveh vektorjev, katerih elementi so realna števila, predstavljena z zapisom v plavajoči vejici z enojno natančnostjo (float).

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
// add the elements of two arrays
void VectorAdd(float *vecA,
               float *vecB,
               float *vecC,
               int iNumElements) {

    int iGID = 0;

    while (iGID < iNumElements) {
        vecC[iGID] = vecA[iGID] + vecB[iGID];
        iGID += 1;
    }
}
Argumenti funkcije so naslovi vseh treh vektorjev (vecA, vecB, vecC) ter število elementov v vektorjih (iNumElements). Vektorje seštevamo tako, da seštejemo vse istoležne elemente v vhodnih vektorjih vecA in vecB ter rezultat shranimo v istoležni element vektorja vecC. Ker imajo vektorji iNumElements elementov, bomo seštevanje vektorjev opravili v zanki, ki ser bo ponovila iNumElements krat.

V zanki posamezne elemente vektorja indeksiramo s celoštevilskim indeksom iGID. Ime smo izbrali namenoma (ang. Global Index), razlog za tako izbiro nam bo znan kmalu.

Naiven poskus seštevanja vektorjev

Funkcijo VectorAdd() bi sedaj radi izvajali na GPE. Spoznali smo, da so GPE idealne za izvajanje nalog, kjer imamo veliko podatkovnega paralelizma - v našem primeru je to še kako res, saj isto operacijo (seštevanje) izvajamo nad vsemi elementi vektorjev. Povrh pa so te operacije medseboj neodvisne (elemente vektorjev lahko seštejemo v poljubnem vrstnem redu) in jih lahko tako seštevamo vzporedno.

Program za napravo - ščepec

V prejšnjih poglavjih smo spoznali, da programe za GPE (ščepce) pišemo tako, da jih bodo lahko (hkrati) izvajale vse niti iz prostora NDRange (tega bomo seveda še definirali v nadaljevanju) - to so vse niti, ki jih bomo na GPE pognali. Zato moramo delo, ki ga opravi funkcija VectorAdd(), čim bolj eneakomerno razdeliti med niti. V našem primeru je to dokaj preprosta naloga, saj bomo za začetek delo razdelili tako, da bo vsaka nit seštela le istoležna elementa. Zato bomo na GPE pognali toliko niti, koliko je elementov v vektorjih - v našem primeru je to iNumElements niti.

Vseh iNumElements niti, ki jih bomo pognali, tvori prostor NDRange. Omenili smo že, da je ta prostor lahko eno-, dvo-, ali tri-dimenzionalen in da razsežnost prostora prilagajamo podatkom. Ker so naši vektorji enodimenzionalna polja, bomo prostor NDRange v našem primeru organizirali v eni dimenziji (X). Globalna velikost tega prostora bo kar iNumElements. Zato bo globalni indeks vsake niti, ki se bo izvajala na GPE, enolično določen iz tega prostora in bo pravzaprav ustrezal indeksu elementov, ki jih nit sešteva. Ščepec, ki ga bodo vse niti izvajale na GPE, je prikazan v spodnji kodi:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
__kernel void vecadd_naive (
                __global float* vecA,
                __global float* vecB,
                __global float* vecC,
                int iNumElemements){

    int iGID = get_global_id(0);

    if (iGID < iNumElemements) {
        vecC[iGID] = vecA[iGID] + vecB[iGID];
    }
}

Vidimo, da je koda ščepca zelo podobna kodi funkcije VectorAdd(). Vsaka funkcija, ki jo želimo izvajati na GPE, mora biti deklarirana kot ščepec (__kernel). Ščepec smo poimenovali vecadd_naive(), njegovi argumenti pa so naslovi vektorjev ter število elementov v vektorjih. Naslovi vektorjev imajo tokrat specifikator __global, s katerim določimo, da se ti naslovi nanašajo na globalni pomnilnik na GPE. Vektorje, ki jih želimo seštevati na GPE, moramo namreč hraniti na GPE, saj računske enote na GPE ne morejo naslavljati glavnega pomnilnika na gostitelju. Ker so vektorji lahko precej veliki, je globalni pomnilnik edini ustrezen pomnilniški prostor za hranjenje vektorjev. Povrh pa vse računske enote lahko naslavljajo globalni pomnilnik in bodo zato elementi vektorjev dostopni vsem nitim, ne glede na to, na kateri računski enoti se bo posamezna nit izvajala (na to itak nimamo vpliva).

Vsaka nit najprej ugotovi svoj globalni indeks iGID v prostoru NDRange s funkcijo get_global_id(0). Argument 0 pri klicu funkcije določa, da zahtevamo indeks v dimenziji X. Nato vsaka nit sešteje istoležna elementa v vektorjih, katerih indeks je enak njenemu indeksu, vendar le, če je njen indeks manjši ali enak številu elementov v vektorjih (zato je v ščepcu if stavek). Na ta način zagotovimo dvoje:

  • podatki, s katerimi dela posamezna nit, so za vsako nit enolično določeni in ni možnosti, da bi morebiti dve niti skušali pisati v isti element vektorja vecC,
  • pomnilniške dostope združujemo v segmente, kjer vse niti v istem snopu dostopajo do zaporednih 8-, 16-, 24- ali 32 bitnih podatkov v pomnilniku (ang. memory coalescing).

Sedaj nam preostane le še, da si pogledamo, kako tak ščepec prevedemo za izbrano GPE napravo ter kako ga prenesemo na GPE in tam poženemo.

Podroben opis za OpenCL C programski jezik najdemo na spletni strani The OpenCL C Specification.

Program na gostitelju

Ko smo napisali ščepec, ga moramo najprej prevesti. Ščepce prevajamo za izbrano napravo, zato moramo najprej izbrati platformo,napravo v platformi, ustvariti kontekst in ukazno vrsto za vsako napravo kot smo to spoznali v prejšnjem poglavju. S tem delo programa na gostitelju še zdaleč ni končano. Ta mora sedaj opraviti še naslednje naloge:

  1. rezervirati prostor v glavnem pomnilniku in inicializirati podatke na gostitelju,
  2. rezervirati prostor v globalnem pomnilniku naprave, kjer bomo hranili naše vektorje,
  3. prebrati program za napravo iz datoteke v pomnilnik,
  4. prevesti program za napravo - program za napravo se bo prevajal med izvajanjem programa na gostitelju,
  5. iz prevedenega programa ustvariti ščepec, ki ga bomo zagnali na izbrani napravi,
  6. prenesti podatke iz gostitelja na napravo,
  7. nastaviti argumente za ščepec,
  8. nastaviti velikost in organizacijo prostora niti NDRange,
  9. zagnati ščepec na napravi ter
  10. prebrati podatke iz globalnega pomnilnika na napravi po končanem izvajanju ščepca.

Offloading to GPU

Zgornja slika prikazuje vse korake, ki jih mora izvesti program na gostitelju pred in po zaganjanju ščepcev na GPE naparavi.

Inicializacija podatkov na gostitelju

Program na gostitelju mora inicializirati vse podatke, ki so potrebni za računanje na napravi. Zavedati se moramo, da gostitelj dostopa lahko le do svojega glavnega pomnilnika in bodo zaenkrat vsi podatki, ki jih gostitelj inicializira, v glavnem pomnilniku. V našem primeru moramo rezervirati prostor za vse tri vektorje ter inicializirati vektorja, ki ju želimo sešteti:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
    cl_float* vecA_h;
    cl_float* vecB_h;
    cl_float* vecC_h;

    int iNumElements = 256*4; // works for vectors with up to 1024 elements

    // Allocate host arrays
    vecA_h = (void *)malloc(sizeof(cl_float) * iNumElements);
    vecB_h = (void *)malloc(sizeof(cl_float) * iNumElements);
    vecC_h = (void *)malloc(sizeof(cl_float) * iNumElements);
    // init arrays:
    for (int i = 0; i<iNumElements; i++ ) {
        vecA_h[i] = 1.0;
        vecB_h[i] = 1.0;
    }

Imenom vektorjev, ki jih hranimo na gostitelju smo dodali _h, da poudarimo, da so to podatki na gostitelju (ang. host). Vse elemente vhodnih vektorjev na gostitelju smo nastavili na vrednost 1.

Rezervacija prostora v globalnem pomnilniku naprave

Sedaj moramo rezervirati prostor v globalnem pomnilniku naprave, kjer bomo hranili vse vektorje, ki nastopajo pri računanju. Za rezervacijo prostora v globalnem pomnilniku naprave uporabljamo funkcijo clCreateBuffer(). Ta rezervira prostor velikosti datasize. Pri rezervaciji prostora v globalnem pomnilniku navedemo še, ali bo naprava iz rezerviranega prostora samo brala (CL_MEM_READ_ONLY), vanj samo pisala (CL_MEM_WRITE_ONLY) ali pa bo imela bralne in pisalne dostope (CL_MEM_READ_WRITE). Funkcija clCreateBuffer() vrne rokovalnik (tipa cl_mem) na pomnilniški objekt (rezervirani prostor) v globalnem pomnilniku naprave. Ne pozabimo, da gostitelj (CPE) nima neposrednega dostopa do tega pomnilnika zato nikoli ne poskusite uporabiti tega rokovalnika kot kazalca in ga dereferencirati na gostitelju! 

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
    //***************************************************
    // Create device buffers
    //***************************************************
    cl_mem vecA_d; // Input array on the device
    cl_mem vecB_d; // Input array on the device
    cl_mem vecC_d; // Output array on the device

    // Size of data:
    size_t datasize = sizeof(cl_float) * iNumElements;

    // Use clCreateBuffer() to create a buffer object (d_A)
    // that will contain the data from the host array A
    vecA_d = clCreateBuffer(
                             context,
                             CL_MEM_READ_ONLY,
                             datasize,
                             NULL,
                             &status);
    clerr_chk(status);

    // Use clCreateBuffer() to create a buffer object (d_B)
    // that will contain the data from the host array B
    vecB_d = clCreateBuffer(
                             context,
                             CL_MEM_READ_ONLY,
                             datasize,
                             NULL,
                             &status);
    clerr_chk(status);

    // Use clCreateBuffer() to create a buffer object (d_C)
    // with enough space to hold the output data
    vecC_d = clCreateBuffer(
                             context,
                             CL_MEM_WRITE_ONLY,
                             datasize,
                             NULL,
                             &status);
    clerr_chk(status);

Branje programa za napravo - ščepca

Programe za napravo pišemo v datotekah s končnico .cl. V našem primeru bo datoteka kernel.cl vsebovala le ščepec vecadd_naive(). V splošnem bi datoteka kernel.cl lahko vsebovala poljubno število ščepcev in funkcij, ki jih ti ščepci kličejo. Program moramo prevesti za izbrano napravo, zato to počnemo takrat, ko ima gostitelj že vse informacije o platformi in napravah. Pred samim prevajanjem, moramo program iz datoteke kernel.cl prebrati v pomnilnik gostitelja in nato iz njega ustavriti programski objekt, primeren za prevajanje. To storimo s spodnjo kodo.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
    //***************************************************
    // Create a program object for the context
    //***************************************************

    FILE* programHandle;            // File that contains kernel functions
    size_t programSize;
    char *programBuffer;
    cl_program cpProgram;
    // 6 a: Read the OpenCL kernel from the source file and
    //      get the size of the kernel source
    programHandle = fopen("kernel.cl", "r");
    fseek(programHandle, 0, SEEK_END);
    programSize = ftell(programHandle);
    rewind(programHandle);

    printf("Program size = %lu B \n", programSize);

    // 6 b: read the kernel source into the buffer programBuffer
    //      add null-termination-required by clCreateProgramWithSource
    programBuffer = (char*) malloc(programSize + 1);

    programBuffer[programSize] = '\0'; // add null-termination
    fread(programBuffer, sizeof(char), programSize, programHandle);
    fclose(programHandle);

    // 6 c: Create the program from the source
    //
    cpProgram = clCreateProgramWithSource(
                                          context,
                                          1,
                                          (const char **)&programBuffer,
                                          &programSize,
                                          &status);
    clerr_chk(status);
    free(programBuffer);

Najprej odpremo datoteko kernel.cl in iz nje preberemo kodo ter jo kot zaporedje znakov shranimo v znakovni niz programBuffer. Šele nato s funkcijo clCreateProgramWithSource() preberemo znakovni niz programBuffer in iz njega ustvarimo ustrezen programski objekt cpProgram, primeren za prevajanje.

Prevajanje programa za izbrano GPE napravo

Prevajanje prebranega programa, ki ga hranimo v objektu cpProgram, prevedemo s funkcjo clBuildProgram(), kot to prikazuje spodnja koda.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
    //***************************************************
    // Build the program
    //***************************************************

    status = clBuildProgram(
                           cpProgram,
                           0,
                           NULL,
                           NULL,
                           NULL,
                           NULL);
    clerr_chk(status);

Zavedati se moramo, da se šele sedaj, torej v času izvajanja programa na gostitelju, prevede program za napravo, ki smo ga prebrali iz datoteke kernel.cl. Kakršnekoli napake v kodi v kernel.cl se bodo pokazale šele sedaj, med prevajanjem, in ne takrat, ko prevajamo program za gostitelja. Zato moramo v primeru napak med prevajanjem (pri funkciji clBuildProgram()), le te shraniti v podatkovni niz v glavnem pomnilniku ter jih po potrebi prikazati. To lahko storimo s funkcijo clGetProgramBuildInfo():

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
    if (status != CL_SUCCESS)
    {
        size_t len;

        printf("Error: Failed to build program executable!\n");
        // Firstly, get the length of the error message:
        status = clGetProgramBuildInfo(cpProgram,
                              devices[0],
                              CL_PROGRAM_BUILD_LOG,
                              0,
                              NULL,
                              &len);
        clerr_chk(status);

        // allocate enough memory to store the error message:
        char* err_buffer = (char*) malloc(len * sizeof(char));

        // Secondly, copy the error message into buffer
        status = clGetProgramBuildInfo(cpProgram,
                              devices[0],
                              CL_PROGRAM_BUILD_LOG,
                              len * sizeof(char),
                              err_buffer,
                              NULL);
        clerr_chk(status);
        printf("%s\n", err_buffer);
        free(err_buffer);
        exit(1);
    }

Funkcijo clGetProgramBuildInfo() kličemo dvakrat. Prvič, da ugotovimo dolžino sporočila o napaki, in drugič, da preberemo celotno sporočilo o napaki. Spodaj je prikazan primetr izpisa napake v primeru, ko smo v ščepcu uporabili nedeklarirano ime za vektor B:

Error: Failed to build program executable!
<kernel>:9:35: error: use of undeclared identifier 'vecBB'; did you mean 'vecB'?
        vecC[myID] = vecA[myID] + vecBB[myID];
                                  ^~~~~
                                  vecB
<kernel>:2:33: note: 'vecB' declared here
                __global float* vecB,
                                ^

Ustvarjanje ščepca iz prevedenega programa

V splošnem bi lahko naš prevedeni program vseboval več ščepcev. Zato sedaj iz programa ustvarimo le ščepec vecadd_naive(), ki ga želimo zagnati na napravi. To storimo s funkcijo clCreateKernel().

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
    //***************************************************
    // Create and compile the kernel
    //***************************************************

    cl_kernel ckKernel;
    // Create the kernel
    ckKernel = clCreateKernel(
                              cpProgram,
                              "vecadd_naive",
                              &status);
    clerr_chk(status);

Prenos podatkov na napravo

Vektorja, ki smo ju inicializirali v prvem koraku, prenesemo iz glavnega pomnilnika gostitelja v prej ustvarjena pomnilniška objekta v globalnem pomnilniku naprave. Prenos podatkov pravzaprav sprožimo s pisanjem ustreznega ukaza v ukazno vrsto cmdQueue naprave. To storimo s funkcijo clEnqueueWriteBuffer().

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
    //***************************************************
    // Write host data to device buffers
    //***************************************************
    // Use clEnqueueWriteBuffer() to write input array A to
    // the device buffer bufferA
    status = clEnqueueWriteBuffer(
                                  cmdQueue,
                                  vecA_d,
                                  CL_FALSE,
                                  0,
                                  datasize,
                                  vecA_h,
                                  0,
                                  NULL,
                                  NULL);
    clerr_chk(status);

    // Use clEnqueueWriteBuffer() to write input array B to
    // the device buffer bufferB
    status = clEnqueueWriteBuffer(
                                  cmdQueue,
                                  vecB_d,
                                  CL_FALSE,
                                  0,
                                  datasize,
                                  vecB_h,
                                  0,
                                  NULL,
                                  NULL);
    clerr_chk(status);

Nastavljanje argumentov ščepca

Da bi pognali izbrani ščepec vecadd_naive(), mu moramo nastaviti še argumente. Za to uporabimo funkcijo clSetKernelArg().

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
    //***************************************************
    // Set the kernel arguments
    //***************************************************

    // Set the Argument values
    status = clSetKernelArg(ckKernel,
                           0,
                           sizeof(cl_mem),
                           (void*)&vecA_d);
    status |= clSetKernelArg(ckKernel,
                            1,
                            sizeof(cl_mem),
                            (void*)&vecB_d);
    status |= clSetKernelArg(ckKernel,
                            2,
                            sizeof(cl_mem),
                            (void*)&vecC_d);
    status |= clSetKernelArg(ckKernel,
                            3,
                            sizeof(cl_int),
                            (void*)&iNumElements);
    clerr_chk(status);

Drugi argument funkcije clSetKernelArg() označuje vrstni red argumenta pri klicu ščepca.

Nastavitev in organizacija prostora NDRange

Smo tik pred tem, da zaženemo niti, ki bodo izbrani ščepec izvajale na napravi GPE. Določiti moramo še koliko niti bomo pognali na napravi GPE, kako bodo te niti organizirane v delovne skupine ter kako bodo niti in delovne sklupine organizirane v prostoru NDRange. Temu pravimo nastavitev prostora NDRange. Zaenkrat bomo delali v enodimenzionalnem prostoru NDRange, zato njegovo velikost in organizacijo določimo z naslednjima spremenljivkama szLocalWorkSize in szGlobalWorkSize:

1
2
3
    // set and log Global and Local work size dimensions
    const size_t szLocalWorkSize = 128;
    const size_t szGlobalWorkSize = iNumElements;

V zgornji kodi smo določili, da želimo imeti 128 niti v delovni skupini in da želimo pognati toliko niti, kolikor so dolgi vektorji (iNumElements).

Zagon ščepca na izbrani GPE napravi

Ščepec na izbrani napravi zaženemo tako, da v ukazno vrsto cmdQueue zapišemo ukaz za poganjanje ščepca ter opis prostora NDRange. Oboje storimo s funkcijo clEnqueueNDRangeKernel().

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
    //***************************************************
    // Enqueue the kernel for execution
    //***************************************************

    // Launch kernel
    status = clEnqueueNDRangeKernel(
                                   cmdQueue,
                                   ckKernel,
                                   1,
                                   NULL,
                                   &szGlobalWorkSize,
                                   &szLocalWorkSize,
                                   0,
                                   NULL,
                                   NULL);
    clerr_chk(status);

V zgornji kodi smo rekli, da želimo zagnati ščepec ckKernel, da ima prostor NDRange razsežnost 1, da je njegova globalna velikost szGlobalWorkSize in da so niti organizirane v delovne skupine velikosti szLocalWorkSize.

Na tem mestu je potrebno opozoriti, da ko v ukazno vrsto ustavimo nek ukaz (na primer za prenos podatkov ali zaganjanje ščepca), nimamo nobenega vpliva na to, kdaj se bo ukaz dejansko izvedel. Funkciji clEnqueueWriteBuffer() in clEnqueueNDRangeKernel() sta neblokirajoči in se zaključita takoj. Tako pravzaprav ne vemo, kdaj bodo podatki prenešeni, kdaj se ščepec začne in konča izvajati na napravi. Vrsta nam zagotavlja le vrstni red izvajanja ukazov v njej. Tako smo lahko prepiričani, da se ščepec zagotovo ne bo začel izvajati, preden se podatki ne prenesejo iz gostitelja v napravo.

Branje podatkov iz GPE naprave

Po končanem izvajanju ščepca prenesemo vektor z vsotami iz globalnega pomnilnika naprave v glavni pomnilnik gostitelja.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
    //***************************************************
    // Read the output buffer back to the host
    //***************************************************
    // Synchronous/blocking read of results
    status = clEnqueueReadBuffer(
                                cmdQueue,
                                vecC_d,
                                CL_TRUE,
                                0,
                                datasize,
                                vecC_h,
                                0,
                                NULL,
                                NULL);
    clerr_chk(status);


    // Block until all previously queued OpenCL commands in a command-queue
    // are issued to the associated device and have completed
    clFinish(cmdQueue);
Vendar kako sedaj vemo, kdaj se bodo podatkli dejansko prepisali in kdaj smemo dostopati do njih? To storimo tako, da pokličemo funkcijo clFinish(), s katero blokiramo izvajanje glavnega programa, dokler se ukazna vrsta cmdQueue ne sprazni.

Brisanje pomnilnika na gostitelju

Na koncu programa moramo še sprostiti ves rezervirani pomnilniški prostor na gostitelju.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
    //*************************************************
    // Cleanup
    //*************************************************

    if (srcA_h) free(vecA_h);
    if (srcB_h) free(vecA_h);
    if (srcC_h) free(vecA_h);


    if (platforms) free(platforms);
    if (devices) free(devices);

    if(ckKernel) clReleaseKernel(ckKernel);
    if(cpProgram) clReleaseProgram(cpProgram);
    if(cmdQueue) clReleaseCommandQueue(cmdQueue);
    if(context) clReleaseContext(context);

    if(srcA_d) clReleaseMemObject(srcA_d);
    if(srcB_d) clReleaseMemObject(srcB_d);
    if(srcC_d) clReleaseMemObject(srcC_d);

Celotno kodo iz tega poglavja najdete v mapi 02-vector-add-short tukaj.

Seštevanje poljubno dolgih vektorjev

Ščepec vecadd_naive(), ki smo ga uporabili za seštevanje vektorjev ima eno pomanjklivost - vsaka nit, ki bo ščepec vecadd_naive() izvedla, bo izračunala le eno vsoto. Spomnimo se, da je število niti, ki tvorijo delovno skupino omejeno in da je prav tako omejeno število niti, ki se hkrati izvajajo na eni računski enoti. Pri napravi Tesla K40 je največje število niti, ki se hkrati izvajajo na eni računski enoti, 2048. Ker ima naprava Tesla K40 le 15 računskih enot, je največje število niti, ki se lahko naenkrat izvajajo 30.720. Če poženemo več kot toliko niti, bo moral gonilnik serializirati njihovo izvajanje na GPE. Zato je veliko bolje, da maksimalno število niti, ki jih zaženemo, določimo iz tega, koliko niti bomo dali v eno delovno skupino, koliko skupin se lahko hkrati izvaja na eni računski enoti in koliko je računskih enot na GPE. V tem primeru bomo verjetno imeli manjše število niti kot je dolžina vektorjev. Problem rešimo s spodnjim ščepcem.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
__kernel void VectorAdd_arbitrary(__global float* vecA,
                __global float* vecB,
                __global float* vecC,
                int iNumElements) {

    // get index into global data array
    int iGID = get_global_id(0);
    int iGS = get_global_size(0);

    while (iGID < iNumElements) {
        //add the vector elements
        vecC[iGID] = vecA[iGID] + vecB[iGID];
        iGID = iGID + iGS;
    }
}

Sedaj bo vsaka nit najprej ugotovila svoj globalni index iGID ter globalno velikost prostora NDRange iGS (to je pravzaprav število vseh aktivnih niti). Nato bo istoležne elemente seštevala v zanki while tako, da bo v eni iteraciji seštela dva istoležna elementa in nato povečala iGID za toliko, kolikor je vseh aktivnih niti. Tako se bo premaknila na naslednja dva istoležna elementa in ju seštela. Zanka se zaključi, ko indeks iGID preseže velikost vektorjev.

Zagotovo se sprašujete, zakaj ena nit ne bi seštela sosednji par istoležnih elementov ali 4 sosednje pare istoležnih elementov? Razlog tiči v načinu dostopa do pomnilnika! Spomnimo se, da se do pomnilnika dostopa po segmentih. Da bi zagotovili usklajen dostop do pomnilnika, morata dve sosednji niti v snopu dotopati do dveh sosednjih pomnilniških besed.

Program na gostitelju tokrat ostane enak - le drugi ščepec je treba naložiti. Lahko se tudi sami poskusite poigrati z različnimi velikostmi vektorjev ter z različnimi nastavitvami prostora NDRange.

Celotno kodo iz tega poglavja najdete v mapi 04-vector-add-arb tukaj.

  1. © Patricio Bulić, Univerza v Ljubljani, Fakulteta za računalništvo in informatiko. Gradivo je objavljeno pod licenco Creative Commons Priznanje avtorstva-Nekomercialno-Deljenje pod enakimi pogoji 4.0 Mednarodna. ↩