Utilizarea procesorului grafic (GPU) în aplicații Java propusă de Vlad-Adrian Iacobeț Sesiunea: Iulie, 2018 Coordonator științific Lector, dr…. [631611]

UNIVERSITATEA ALEXANDRU IOAN CUZA IAȘI
FACULTATEA DE INFORMATICĂ

LUCRARE DE LICENȚĂ
Utilizarea procesorului grafic (GPU) în
aplicații Java

propusă de
Vlad-Adrian Iacobeț

Sesiunea: Iulie, 2018

Coordonator științific
Lector, dr. Cristian Frăsinaru

UNIVERSITATEA ALEXANDRU IOAN CUZA IAȘI
FACULTATEA DE INFORMATICĂ

Utilizarea procesorului grafic (GPU) în
aplicații Java

Vlad -Adrian Iacobeț
Sesiunea: Iulie, 2018

Coordonator științific
Lector, dr. Cristian Frăsinaru

DECLARAȚIE DE CONSIMȚĂMÂNT

Prin prezenta declar că sunt de acord ca Lucrarea de licență cu titlul „Utilizarea procesorului grafic
(GPU) în aplicații Java ”, codul sursă al programelor și celelalte conținuturi (grafice, multimedia, date
de test etc.) care însoțesc această lucrare să fie utilizate în cadrul Facultății de Informatică.

De asemenea, sunt de acord ca Facultatea de Informatică de la Universitatea Alexandru Ioan Cuza Iași
să utilizeze, modifice, reproducă și să distribuie în scopuri necomerciale programele -calculator,
format executabil și sursă, realizate de mine în cadrul prezentei lucrări de licență.

Iași, Absolvent: [anonimizat]

_________________________

1
Cuprins
Introducere ………………………….. ………………………….. ………………………….. ………………………….. ………………. 3
Contribuții ………………………….. ………………………….. ………………………….. ………………………….. ……………….. 4
1. Fundamentele teoretice ale programării CUDA ………………………….. ………………………….. …………… 5
1.1. Concepte de bază ale programării pe GPU ………………………….. ………………………….. ……………….. 5
1.2. Arhitectura CUDA ………………………….. ………………………….. ………………………….. ……………………. 6
1.3. Setarea mediului de lucru ………………………….. ………………………….. ………………………….. …………… 7
1.3.1. Microsoft Visual C++ ………………………….. ………………………….. ………………………….. …………. 7
1.3.2. Nvidia CUDA Toolkit ………………………….. ………………………….. ………………………….. ………… 7
1.4. NVCC (Nvidia CUDA Compiler) ………………………….. ………………………….. ………………………….. .. 8
1.5. Noțiuni de bază ale programării GPU în CUDA ………………………….. ………………………….. ………… 9
1.5.1. Fire de execuție ( Threads ) și blocuri ………………………….. ………………………….. …………………. 9
1.5.2. Kernel ………………………….. ………………………….. ………………………….. ………………………….. …10
1.6. Clasifica rea memoriei ………………………….. ………………………….. ………………………….. ……………… 12
1.6.1. Regiștri ………………………….. ………………………….. ………………………….. ………………………….. .12
1.6.2. Memoria locală ………………………….. ………………………….. ………………………….. ………………… 12
1.6.3. Memoria partajată ………………………….. ………………………….. ………………………….. ……………. 12
1.6.4. Memoria globală ………………………….. ………………………….. ………………………….. ……………… 13
1.6.5. Memoria constantă ………………………….. ………………………….. ………………………….. …………… 13
1.7. Sincronizarea firelor de execuție ………………………….. ………………………….. ………………………….. ..13
1.7.1. Procedura __syncthreads() ………………………….. ………………………….. ………………………….. …13
1.7.2. Pattern -uri de comunicare ………………………….. ………………………….. ………………………….. ….14
1.7.3. Funcții atomice ………………………….. ………………………….. ………………………….. ………………… 15
1.8. Runtime API vs. Driver API la nivel de host ………………………….. ………………………….. …………… 16
2. Utilizarea platformei CUDA în aplicații Java ………………………….. ………………………….. …………….. 17
2.1. Limbajul de programare Java ………………………….. ………………………….. ………………………….. ……. 17
2.2. Maven ………………………….. ………………………….. ………………………….. ………………………….. ……….. 17
2.3. JCuda ………………………….. ………………………….. ………………………….. ………………………….. ………… 17
2.4. Apache POI ………………………….. ………………………….. ………………………….. ………………………….. ..19
3. Implementare și benchmark -uri ………………………….. ………………………….. ………………………….. ……. 20
3.1. Prezentarea proiectului ………………………….. ………………………….. ………………………….. …………….. 20
3.2. Exemple de algoritmi pentru înțelegerea conceptelor de prog ramare CUDA ………………………… 22
3.3. Utilizarea GPU -ului în cadrul algoritmilor Monte Carlo ………………………….. ……………………….. 25
3.3.1. Utilizarea algoritmului Monte Carlo în domeniul bursei ………………………….. ………………… 25
3.3.2. Implementare folosind JCuda ………………………….. ………………………….. …………………………. 26
3.3.3. Descrierea kernel -ului ………………………….. ………………………….. ………………………….. ………. 29

2
3.3.4. Clasa ReduceCudaExecutor ………………………….. ………………………….. ………………………….. .30
3.3.5. Benchmark -uri ………………………….. ………………………….. ………………………….. …………………. 32
Concluzie ………………………….. ………………………….. ………………………….. ………………………….. ………………… 34
Bibliografie ………………………….. ………………………….. ………………………….. ………………………….. ……………… 35

3
Introducere

Lucrarea de față are ca scop explorarea tehnicilor de programare a procesorului grafic
(GPU) prin intermediul platformei de programare paralelă Nvidia CUDA și integrarea acestei
platforme în aplicații dezvoltate folosind limbajul de programare Java. De asemenea, se vor expune
rezultatele de performanță (benchmarks ) obținute în urma implementării unor algoritmi într -o
manieră care se bazează pe puterea de calcul a GPU -ului, cât și concluziile car e pot fi trase pe baza
acestor rezultate.
În ultimii ani am asistat la o ascensiune a hardware -ului de calcul paralel, aces ta ajungând
sa fie folosit într -o arie cât mai vastă de domenii precum grafică, inteligență artificială, învățare
automată, algoritmi genetici , jocuri video , minarea de cripto monede și cam orice domeniu în care
o problemă poate fi divizată în probleme mai mici care sunt apoi rezolvate în paralel. Desigur,
această ascensiune a dus la crea rea unor platforme care să ofere un mediu de dezvoltare ce permite
utilizarea resurselor de calcul paralel într-un mod cât mai accesibil pentru a facilita crearea de
software car e poate beneficia de un spor de performanță considerabi l datorită unei abordări care se
bazează pe aceste concepte . Una dintre cele mai cunoscute platforme este CUDA , folosit ă pentru
a scrie software care să ruleze direct pe procesoarele grafice dezvoltate de Nvidia. Deși CUDA
limitează programatorul la hardware -ul produs de Nvidia am ales totuși această platformă în
detrimentul omologului său open source OpenCL deoarece această exclusivitate față de Nvidia
elimină problemele legate de compatibilitate între arhitecturi diferite de procesoare grafice. Practic,
este mult mai ușor de abordat programarea pe GPU în Nvidia CUDA decât în OpenCL.
Prin urmare, acea stă lucrare explorează avantajele, cât și dezavantajele unei astfel de
abordări de a scrie software și în special cum s -ar putea îmbunătății din punct de vedere al
performanței aplicațiile dezvoltate în Java dacă ar profita de o platformă precum CUDA. Astfel,
conținut ul lucrării este structurat pe trei mari părți. Prima parte are ca rol explicarea conceptelor de
calcul paralel și cum acestea pot fi abordate în cadrul platformei de programare CUDA . A doua
parte se axează pe modul în care programarea procesor ului grafic poate fi utilizată în Java prin
intermediul librăriei JCuda ce oferă o integrare cu mediul de dezvoltare CUDA printr -o
implementare de tip binding . A treia parte e xplorează diverși algoritmi implementați în Java
folosind librăria JCuda și descrie raportul de performanță din punct de vedere al CPU vs. GPU
obținut pentru acești algoritmi. O atenție deo sebită va fi asupra algoritmilor Monte Carlo și cum
implementarea unui algoritm de o asemenea natură poate beneficia de procesorul g rafic folos indu-
se un exemplu din domeniul bursei. Aceste rapoarte de performanță ne pot oferi o imagine mai
clară asupra avantajelor unei astfel de abordări dar și a dezavantajelor ce pot surveni.
Într-un final îmi doresc ca prin această lucrare să ofer o bază solidă de cunoștințe despre
programarea procesorului grafic, în special CUDA și JCuda dar să -mi expun ș i propriile concluzii
la care am ajuns încercând să observ sporul de performanță obținut în urma unei astfel de abordări
de programare .

4
Contribuții

Pentru a -mi atinge scopul de a descrie utilizarea procesorului grafic în aplicații Java cât și
avantajele și dezavantajele pe care le presupune, a fost necesar în p rimul rând studiul tehnologii
CUDA. Cu alte cuvinte, obiectivul meu a fost aprofundarea cunoștințelor despre programarea în
CUDA. După cum urmează, în prima parte a acestei lucrări am descris atât aspectele teoretice cât
și cele practice ale acestui mod de abordare a programării pentr u a oferi o imagine clară asupra
acestui subiect. Pentru a înțelege aceste aspecte a fost necesar consultarea unor anumite cărți cât și
consultarea diverselor resurse de pe internet precum articole sau documentația oficială pentru
CUDA. Toate aceste surse pot fi găsite în bibliografie.
După prezentarea tehnologiilor CUDA, a urmat înțelegerea modului prin care acestea pot
fi utilizate la nivelul platformei Java. Abia atunci când am înțeles toate aceste aspecte legate de
utilizarea procesorului grafic, am pu tut să încep să implementez algoritmi care să utilizeze
procesorul grafic. Pentru a oferi un exemplu clar de spor de performanță al procesorului grafic, am
ales să utilizez algoritmi Monte Carlo pentru a rezolva o problemă din domeniul bursei.
În capitolele care urmează va fi descrisă platforma CUDA cât și conceptele care stau la
baza acestei platforme, cum se poate implementa în Java și benc hmark -urile obținute pentru
anumite implementări folosind procesorul grafic.

5
1. Fundamentele teoretice ale programării CUDA

1.1. Concepte de bază ale programării pe GPU

Procesorul grafic (GPU) a fost inițial conceput pentru generarea de imagini grafice pentru
a fi redate pe un monitor. Termenul „GPU” a fost popularizat de către Nvidia în 1999 când au
lansat mode lul de placă video GeForce 256 pe care l -au comercializat sub mottoul „ primul GPU
din lume”1. Începând cu anii 2000, procesoarele grafice au cunoscut o evoluție puternică în cea ce
privește puterea de calcul paralel ajungând în ziua de azi să rivalizeze cu procesoarele (CPU) în
anumite aspecte computaționale care nu au legătură cu domeniul graficii 3D, domeniu pentru care
au fost inițial concepute. Deși procesoarele moderne sunt și ele capabile de calcul paralel datorită
tehnologii multi -core, performanța lor din punct de vedere al calculu i paralel nu se compară cu cea
prezentă pe procesorul grafic.
Programarea pe procesorul grafic diferă de programarea secvențială deoarece performanța
asociată rezolvării unei probleme se obține prin divizarea problemei într -un număr cât mai mare de
probleme de dimensiuni mai mic i și rezolvarea lor simultană în timp ce programarea secvențială
dorește rezolvarea problemelor pe rând și într -un timp cât mai scurt. Doi termeni foarte des întâlniți
în programarea pe GPU sunt:
– Host care reprezintă CPU -ul și memoria sa asociată (host memory)
– Device care reprezintă GPU -ul și memoria sa asociată (device memory)
La baza programării pe GPU stă conceptul de kernel care reprezintă o secvență de cod care
se execută pe device la cererea host -ului. Ne putem imagina kernel -ul ca un set de instrucțiuni ce
se execută în interiorul unei structuri repetitive de tip for doar că diferența e ste că acest set în loc
să fie acces at pe rând pentru fiecare iterație, se execută simultan pentru toate iterațiile for -ului.
Conform taxonomiei lui Flynn2, care clasifică în patru categorii arhitecturile de calcul paralel,
conceptul de kernel poate fi asociat cu „Single Instruction, Multiple Data”. Acest tip de arhitectură
„dispune de elemente multiple de procesare care toate efectuează în paralel aceeași operație sau
instrucțiune, dar pe date diferite” [1].
În general, modelul de programare pentru o aplicație care folosește procesorul grafic este
constituit din următorii pași:
1. Host -ul inițializează device -ul și alocă spațiul de memorie necesar pentru procesarea pe
device (la nivelul host -ului se stabilește cât spațiu de memorie este necesar să fie alo cat pe
device)
2. Host -ul inițializează transferul de date de input de la host la device
3. Se rulează kernel -ul pentru datele de input și se obțin datele de output
4. Host -ul inițializează transferul de date de output de la device la host
Foarte important de observat din acești pași că într -o aplicație care folosește procesorul
grafic doar anumite porțiuni de cod ajung să fie rulate pe GPU, motiv pentru care este foarte
important pentru programator să știe cum să -și separe partea aplicației destinată pentru GP U față
de restul aplicației.
Maximizarea performanței în programarea procesorului grafic este exprimată prin raportul:

𝑐𝑜𝑚𝑝𝑙𝑒𝑥𝑖𝑡𝑎𝑡𝑒 𝑚𝑎𝑡𝑒𝑚𝑎𝑡𝑖𝑐 ă
𝑚𝑒𝑚𝑜𝑟𝑖𝑒 𝑢𝑡𝑖𝑙𝑖𝑧𝑎𝑡 ă

1 https://en.wikipedia.org/wiki/GeForce_256
2 https://ro.wikipedia.org/wiki/Taxonomia_lui_Flynn

6

Utilizarea procesorului grafic se justifică cu cât avem o complexitate matematică mai mare
per fir de execuție și memoria utilizată per fir de execuție este mai mică. Un alt procedeu pentru
maximizarea performanței este utilizarea corectă a diverselor tipuri de memorie prezente pe GPU,
tipuri de memorie care dife ră în funcție de latență și viteză de scriere și citire.

1.2. Arhitectura CUDA

CUDA ( Compute Unified Device Architecture ) este o platformă de calcul paralel
proprietară dezvoltat ă de cătr e compania americană Nvidia, prima versiune fiind lansată în 2007.
Scopul acestei platforme de programare este de a abstractiza programarea proces orului grafic într –
un device care poate fi programat prin intermediul unui API. Astfel, programatorul nu trebuie să
cunoască detalii arhitecturale de nivel jos legate de proceso rul grafic, detalii ce pot să difere de la
model la model. Un alt avantaj mare este compatibilitatea asigurată de această platformă care
permite ca un software scris pentru un model de procesor grafic mai vechi să ruleze și pe ultimele
modele fără a fi nevoie de modificări ale codului sursă.
CUDA oferă programatorului un nivel abstract de interacțiune cu procesorul grafic printr –
un API scris pentru limbajul de programare C3 ce permite scrierea de kernel într -un limbaj familiar
fără a fi nevoie de cunoștințe despre primitive grafice pentru a folosi puterea d e calcul a
procesorului grafic. Astfel, se poate utiliza procesorul grafic la o gamă largă de aplicații și nu doar
din domeniul graficii.

Figura 1 .1. Arhitectura platformei CUDA

La nivel ha rdware, platforma CUDA conține procesoare grafice compatibile. În general,
orice procesor grafic produs de Nvidia începând cu anul 2008 este compatibil cu CUDA. Fiecare
GPU Nvidia are o versiune de Compute Capabilit y. Această versiune „identifică capabilitățile
hardware oferite de GPU și este folosit de către aplicații la runtime pentru a stabili care
capabilități și/sau instrucțiuni sunt disponibile pe respectivul GPU. Versiunea de Compute
Capability nu trebuie conf undată cu versiunea de CUDA (de ex. CUDA 7.5, CUDA 8.0 etc) care

3 C este un limbaj de programare procedural dezvoltat de Dennis Ritchie la Bell Labs și lansat în 1972

7
reprezintă versiuni ale platformei software CUDA. Platforma software CUDA este folosită de către
programatori pentru a crea aplicații care pot rula pe o gamă largă de arhitecturi de GPU,
incluzând arhitectu ri care nu au fost lansate încă.” [2] De exemplu procesoarele grafice care au
versiunea de Compute Capability sub 2.0 pot avea maxim 512 fire de execuție p er bloc, în timp ce
procesoarele grafice care au de la 2.0 în sus pot avea maxim 1024 fire de execuție per bloc.

La nivelul software, platforma CUDA conține driver -ele pentru integrarea cu sistemul de
operare și două API -uri: unul de tip „low level” , Drive r API, și unul de tip „ high level ”, Runtime
API.
CUDA Driver API este un API care are la bază componente care permit scrierea de kernel
în cod C precum suport pentru ierarhia de fire de execuție, utilizarea diverselor tipuri de memorie
disponibilă în CUDA și alte proceduri utilizate pentru a profita de puterea de calcu l oferită de
procesorul grafic. Pe lângă partea de kernel, oferă și un API implementat tot în C pentru
interacțiunea cu device -ul de pe partea de host cum ar fi transferul de date între host și device sau
apelarea pentru execuție a kernel -ului. Acest API e ste însă unul de nivel jos deoarece pe partea de
host sunt prezenți pași în plus care nu sunt prezenți în Runtime API. Cu toate acestea Driver API,
poate fi integrat cu alte limbaje de programare prin implementări de tip binding cum ar fi JCuda, o
librărie care permite aplicațiilor Java să utilizeze platforma CUDA.
CUDA Runtime API este un API disponibil tot în C și este co nstruit ca un API de tip „ high
level ” având la bază to tuși Driver API. Partea de API care se ocupă de kernel rămâne neschimbată,
în schimb pe partea de host este prezent un API mult mai facil de utilizat. În plus Runtime API
transformă CUDA într -un limbaj de programare de sine stătător deoarece este implementat ca o
extensie a limbajului de programare C. Dezavantajul e că Runtime API nu poate fi folosit în alte
limbaje de programare decât într -un mod limitat care nu permite apeluri de kernel .

1.3. Setarea mediului de lucru

Pentru a seta mediul de lucru necesar dezvoltării și rulării de aplicații CUDA pe sistemul
de operare Windows este necesar un procesor grafic compatibil4, Microsoft Visual C++ și Nvidia
CUDA Toolkit5.

1.3.1. Microsoft Visual C++
Microsoft Visual C++ este un mediu integrat de dezvoltare (IDE) dezvoltat de Microsoft
pentru limbajele de programare C și C++. Oferă o serie de unelte pentru dezvoltarea și testarea de
cod C și C++. Inițial era dezvoltat ca produs separat dar în prezent face parte din suita Microsoft
Visual Studio. Instalarea s a este necesară deoarece compilatorul său C integrat este folosit în cadrul
platformei CUDA.

1.3.2. Nvidia CUDA Toolkit
Nvidia CUDA Toolkit reprezintă o serie de unelte care asigură un mediu de dezvoltare
pentru a plicații care folosesc puterea de calcul oferită de procesoarele grafice compatibile cu
CUDA. Printre acestea se numără driver -ul pentru GPU, compilatorul NVCC, librării pentru

4 https://developer.nvidia.com/cuda -gpus
5 https://developer.nvidia.com/cuda -toolkit

8
dezvoltarea de aplicații și unel te pentru testare și optimizare precum Nvidia Nsight, un plu gin
pentru Visual Studio ce permite folosirea sa ca IDE pentru CUDA.

1.4. NVCC (Nvidia CUDA Compiler)

NVCC (Nvidia CUDA Compiler) este un compilator cu licență comercială proprietară
dezvoltat de Nvidia. Este folosit pentru a compila fișiere de tip .cu ce conțin cod sursă CUDA.
Pentru a verifica versiunea de NVCC instalată se poate utiliza următoarea comandă: nvcc –version.

Figura 1.2.

Să presupunem următorul cod CUDA de tip „hello world” salvat în fișierul hello_world.cu ,
vizibil și în tabelul 1.1. În figura 1.3 avem compilarea fișierului la linia de comandă și executarea
sa.

#include <stdio.h>

__global__ void mykernel() {
printf("Hello world" );
}

int main() {
mykernel<<<1,1>>>();
return 0;
}
Tabelul 1.1. Exemplu de „hello world” pentru CUDA

Figura 1.3. Compilarea unui fișier cu cod sursă CUDA

În timpul procesului de compilare, NVCC separă codul sursă pentru host de codul sursă
pentru device . Codul sursă pentru h ost este apoi compilat de Viusal C++ în timp ce codul sursă
pentru device este co mpilat în fișier de t ip .ptx care este folosit mai departe de către GPU , unde are
un proces de compilare de tip JIT6.
Fișierele de tip .ptx conțin cod sursă PTX (Parallel Thread Execution) care reprezintă un
limbaj de nivel jos.

6 JIT – Just In Time

9

Figura 1.4. Pașii compilării unui fișier de tip .cu

1.5. Noțiuni de bază ale programării GPU în CUDA

În acest capitol voi descrie elementele care alcătuiesc ierarhia de procesare paralelă în
CUDA, cum se implementează un kernel și voi oferi un exemplu de abordare a unui algoritm
simplu din perspectiva programării CUDA.

1.5.1. Fire de execuție ( Thread s) și blocuri
Noțiunea de fir de execuție stă la baza programării CUDA și reprezintă cea mai mică
unitate de procesare la nivel de GPU. Suportul pentru fire de execuție multiple este prezent în multe
limbaje de programare moderne dar în CUDA reprezintă un concept de bază deoarece creșterea de
performanță din punct de vedere al programării paralele constă în mărirea numărului de fire de
execuție și nu îmbunătățirea puterii de calcul a firului de execuție individual.
Noțiunea de bloc reprezintă un grup de maxim 1024 de fire de execuție (512 pe procesoarele
grafice mai vechi) care se execută simultan, comunică între ele și pot avea resurse comune. Această
grupare a firelor de execuție pe blocuri este foarte importantă în cadrul programării CUDA
deoarece sincronizarea firelor de execuție este una dintre cele mai des întâlnite probleme. Este
foarte important de precizat și faptul că mai multe blocuri care se execută în paralel au obligatoriu
număr egal de fire de execuție.
Mai multe blocuri care rulează în paralel constituie un grid. Atât grid -urile cât și blocurile
pot fi unidimensionale, bidimensionale sau tridimensionale.
Indexarea reprezintă un alt concept important din cadrul programării CUDA. Având în
vedere că programarea pe GPU constă în rularea unui număr foarte mare de fire de execuție, este
foarte important să putem identifica firele de execuție într-un mod facil . Pentru aceas ta CUDA
oferă 3 structuri la nivel de API :
– threadIdx – reprezintă indexul thread -ului în cadrul blocului
– blockIdx – reprezintă indexul blocului în cadrul grid -ului
– block Dim – reprezintă numărul de thread -uri dintr -un bloc
– gridDim – reprezintă numărul de blocuri dintr -un grid
Fiecare dintre aceste structuri are 3 câmpuri (x, y și z), unde fiecare câmp reprezintă o
dimensiune a blocului. De exe mplu, dacă avem un număr de blocuri unidimensionale ce se execută
în paralel și vrem să aflăm indexul unui fir de execuție în cadrul grid -ului și nu la nivel de bloc
putem utiliza formula:
Index = threadIdx.x + blockIdx.x * blockDim.x
Să presupunem că avem 3 blocuri unidimensionale a câte 6 fire de execuție și ca input avem
un tablou unidimensional de 16 elemente, e lementul de pe ultima poziție va fi procesat de către

10
firul de execuție cu indexul 15, adică al patrulea fir de execuție din al treilea bloc deoarece 3 + (2
* 6) = 15 .
Aceste noțiuni stau la baza creării de ierarhii de fire de execuție în CUDA și sunt foarte
importante deoarece astfel se pot gestiona un număr foarte mare de fire de execuție. În figura 1.5
putem observa o astfel de ierarhie constituită dintr -un grid bidim ensional alcătuit din blocuri
bidimensionale. Aflarea indexului unui fir de execuție în cadrul grid -ului din figura 1.5 se poate
face astfel:
Index = threadIdx.x + threadIdx.y * blockDim.x + (blockIdx.x + blockIdx.y * gridDim.x) *
(blockDim.x * blockDim.y )

Figura 1.5 . Ierarhie constituită dintr -un grid bidimensional alcătuit din blocuri bidimensionale 7

1.5.2. Kernel
Kernel -ul reprezintă o procedură care se execută pe device și este apelată de către host.
Kernel -ul se declară prin cuvântul cheie __global__ la începutul semnăturii. Pentru a apela un
kernel se folosește următoarea sintaxă: nume kernel <<<N, M>>>(list ă parametrii) unde N
reprezintă numărul de blocuri folosite și M numărul de fir e de execuție prezente per bloc, adică
dimensiunea grid-ului.

Pentru a exemplifica mai bine concepte le de bază prezentate mai sus vom considera
următorul algoritm simplu: fie un tablou unidimensional cu n elemente unde fiecare element este
parcurs și ridicat la pătrat. Dacă conform abordării clasice a p rogramării secvențiale se parcurge pe
rând fiecare element din tablou, din perspectiva programării paralele se vor accesa fiecare element
simultan. Un kernel care să îndeplinească această cerință poate fi observat în tabelul 1.2 .

7 https://en.wikipedia.org/wiki/File:Block -thread.svg

11

__global__ void computeSquare( float * array, int size) {

int index = threadIdx.x + blockIdx.x * blockDim.x;
if (index < size)
array[index] = array[index] * array[index];

}

Tabe lul 1.2 . Exemplu de kernel CUDA

În exemplul de cod de mai sus avem o procedură numită computeSquare care primește ca
parametrii un tablou unidimensional cu variabile de tip float și numărul de elemente din tablou.
Cuvântul cheie __global__ atenționează compilatorul că procedura aceasta e ste un kernel și prin
urmare va rula pe GPU. Având în vedere că se va folosi câte un fir de execuție pentru fiecare
element din tablou, este important să știm indexul firului de execuție care rulează iar pentru asta
avem variabila index care reprezintă suma dintre indexul firului de execuție în cadrul blocului și
produsul dintre indexul blocului și numărul tota l de fire de execuție per bloc.
Condiția (index < size) este importantă și o putem explica astfel: dacă avem 16 elemente în
tablou dar 3 blocuri a câte 6 fire de execuție atunci se vor rula 18 fire de execuție dar ultimele două
fire de execuție nu au ce element să acceseze. Astfel, pentru a evita potențiale erori continuarea
firului de execuție este limitată de această condiție.
În cont inuare se v a exemplifica cum această proce dură este apelată de către host.

int main() {

const int size = 16;
float host_array[size];
float *device_array;
int block_size = 6;

for (int i = 0; i < size; i++) {
host_array[i] = ( float)rand();
printf("%f\n", host_array[i]);
}

cudaMalloc(( void **)&device_array, size * sizeof(float));
cudaMemcpy(device_array, host_array, size * sizeof(float),
cudaMemcpyHostToDevice );

computeSquare << <3, 6 >> >(device_array, size);

cudaMemcpy(host_array, device_array, size * sizeof(float),
cudaMemcpyDeviceToHost );
cudaFree(device_array);

printf("\nAfter computeSquare \n");
for (int i = 0; i < size; i++) {
printf("%f\n", host_array[i]);
}
return 1;

}

Tabelul 1 .3. Exemplu de cod C care apelează un kernel de CUDA

12
Codul din tabelul 1.3 . rulează pe host și este implementat folosind CUDA Runtime API .
Inițial se populează tabloul unidimensional cu 16 elemente de tip float. Apoi se alocă spațiu în
memoria procesorului grafic prin procedura cudaMalloc și se transferă cele 16 elemente din
memoria RAM în memoria GPU -ului prin procedura cudaMemcpy . Linia de cod
computeSquare <<< 3, 6>>>(device_array, size) apelează kernel -ul folosind un grid alcătuit din 3
blocuri a câte 6 fire de execuție. În final elementele modificate sunt copiate înapoi în memoria
RAM iar memoria din GPU este eliberată prin procedura cudaFree . Din punct de vedere al
performanței putem observa că atunci când vrem să folosim GPU -ul trebuie luat în considerare și
timpul de transfer a datelor din memoria host -ului în memoria device -ului și invers.

1.6. Clasificarea memoriei

Procesoarele grafice compatibile cu CUDA au tipuri multiple de memorie: registre, locală,
partajată, globală, constantă . Utilizarea tipului de memorie corectă este foarte important din punct
de vedere al performanței deoarece fiecare tip diferă în funcție de rol, latență și locație.

Figura 1.6 . Tipuri de memorie GPU, viteza de scriere/citire și latență [3]

1.6.1. Regi ștri
Regiștri sunt un tip de memorie specifică firului de execuție fiind accesată doar de către
firul de execuție asociat. Durata sa de viață depinde de durat a de viață a firului de execuție adică
imediat ce acesta și -a terminat execuția, memoria este ștearsă. Deoarece dimensiunea acestui ti p
de memorie este una foarte mică, singurele variabile care ajung să fie stocate sunt variabilele din
interiorul kernel -ului inițializate la compilare . Cu toate acestea, acest tip de memorie are avantajul
de a fi cea mai rapidă din punct de vedere al latenței și vitezei de scriere și citire.

1.6.2. Memoria locală
Memoria locală reprezintă o extensie a memoriei registrului. Din punct de vedere al
accesibilității și duratei de viață este ase mănătoare cu un registru dar este folosită pentru a stoca
variabile inițializate la runtime sau este folosită în cazul în care memoria registrului și -a atins limita.
Din punct de vedere al latenței este asemănătoare cu memoria globală însă începând cu
procesoarele grafice care au nivelul de Compute Capability minim 2.0, această memorie este de tip
cache și este apr oape la fel de rapidă ca memoria registru lui datorită u nei latențe foarte scăzute .
Prin urmare, la procesoarele grafice curente putem considera memoria locală în aceeași categorie
cu memoria registrului la capitolul latenț ă.

1.6.3. Memoria partajată
Memoria partajată este memoria specifică unui bloc și prin urmare poate fi accesată de orice
fir de execuție din respectivul bloc. Durata de viață depinde de durata de viață a blocului și imediat

13
ce acesta își termină execuția, memoria este golită. Din punct de vedere al la tenței și vitezei de
scriere și citire este mult mai rapidă decât alte tipuri de memorie dar nu la fel ca memoria oferită
de regiștri. Acest tip de memorie este foarte important deoarece este folosit la procedee de
sincronizare a firelor de execuți e. În un ele cazuri este de preferat ca datele din memoria globală să
fie transferate și grupate în memoria partajată pentru ca apoi operațiile de procesare a datelor de
către firele de execuție să se facă mai rapid.
La niv el de API, variabilele marcate cu __shared__ la declarare su nt stocate în memori a
partajată precum în figura 1.7 .

Figura 1.7 . Exemplu de cod de stocare a datelor în memoria partajată

1.6.4. Memoria globală
Memoria globală, așa cum sugerează și numele, este cea mai accesibilă memorie și poate fi
scrisă și citită de orice fir de execuție. Host -ul are și el acces la memoria globală și poate executa
operațiuni de citire și scriere. Durata de viață este pe tot parcursul aplicației și se golește doar când
aplicația își încheie execuția. Dezavantajul acestui tip de memorie este însă latența foarte mare. La
nivel de API, variabilele marcate cu __device__ la declarare sunt stocate în memor ia globală
precum în figura 1.8 . De asemenea, parametrii de kernel sunt stocați tot în memoria globală.

Figura 1. 8. Exemplu de cod de stocare a datelor în memoria globală

1.6.5. Memoria constantă
Memoria constantă este asemănătoare cu memoria globală cu o s ingură diferență: din
perspectiva device -ului poate fi doar citită și este scrisă doar de către host . Pentru procesoar ele
grafice care au Compute Capability minim 2.0 această memorie este de tip cache și poate oferi
astfel o latență mai mică față de memoria globală. Memoria constantă maximă care poate fi utilizată
este de 64 KB.

1.7. Sincronizarea firelor de execuție

1.7.1. Procedura __syncthreads()
Procedura __syncthreads() este oferită de API -ul CUDA la nivel de kernel pentru a
sincroniza firele de execuție la nivel de bloc . Rolul său este de a stabili o barieră pentru toate firele
de execuție dintr -un bloc, mai exact atunci când este apelată firul de execuție se oprește și așteaptă
ca toate celelalte fire de execuție din bloc să ajungă în același punct. Doar atunci când toate firele
de execuție au ajuns în același punct, activitatea lor poate continua mai departe indep endent unul
față de celălalt.

14

Figura 1.9 . Exemplificare vizuală a procedurii __syncthreads()

1.7.2. Pattern -uri de comunicare
Sincronizarea firelor de execuție este necesară atunci când acestea au de comunicat între
ele pentru a realiza operații de citire sau scriere asupra memoriei. O implementare greșită poate
duce la rezultate eronate sau la fenomenul de deadlock după cum este descris în exemplul următor:
„firul de execuție 0 preia resursa 0, în timp ce firul de execuție 1 preia resursa 1. Firul de execuție
0 încearcă să preia resursa 1, în timp ce firul de execuție 1 încearcă să preia resursa 0. Cum
resursele nu sunt valabile, ambele fire de execuție așteaptă ca resursele să devină valabile. Cum
nici un fir de execuție nu va elibera resursa deja deținută, toate firele de execuție vor aștepta până
la nesfârșit.” [3]
În cele ce urmează voi descrie câteva pattern -uri de comunicare între fire de execuție,
pattern -uri pe care le -am folosit în cadr ul proiectului. Aceste pattern -uri nu se limitează numai la
programarea pe GPU și sunt general valabile în cadrul programării paralele.
Map este un pattern de tip one-to-one în care fiecare fir de execuție citește dintr -o zonă de
memorie și scrie în altă zonă de memorie fără a exista conflicte între firele de execuție. Kernel -ul
prezentat ca exemplu în capitolul 1.5. este de tipul map.

Figura 1.10 . Pattern -ul map

Gather este un pattern de tip many -to-one în care fiecare fir de execuție citește din mai
multe zone de memorie și scrie într -o singură zonă, zonă în care doar acel fir de execuție scrie.
Două sau mai multe fire de execuție ajung astfel să citească din zone comune de memorie dar
aceasta nu reprez intă mereu o problemă de sincronizare .

15

Figura 1.11 . Pattern -ul gather

Scatter este un pattern de tip one-to-many în care rezultatul obținut de un fir de execuție
este scris în zone multiple de memorie. Un astfel de pattern necesită o implementare care să țină
cont de sincronizarea firelor de execuție.

Figura 1.12 . Pattern -ul scatter

Reduce este un pattern de tip all-to-one în care firele de execuție citesc toate datele din
memoria de input și le procesează pentru a obține un rezultat final care este salvat într -o singură
zonă de memorie. Un exemplu de acest pattern este un algoritm care aplică o operație binară pe un
șir de numere dat ca input, obținându -se un singur număr în final.

1.7.3. Funcții atomice
Funcțiile atomice sunt oferite de API-ul CUDA la nivel de device pentru a veni în ajutorul
programatorului la capitolul fire de execuție care accesează zone comune de memorie . Acestea
sunt descrise conform documentației CUDA oficial e ca „o funcție atomică realizează o operațiune
de tipul citește -modifică -scrie pe o resursă de 32 -bit sau 64 -bit ce se găsește în memoria globală
sau partajată. De exemplu atomicAdd() citește un numă r de la o anumită adresă din memoria
globală sau partajată, face suma cu un alt număr, și scrie rezultatul înapoi la aceeași adresă.
Operația este atomică în sensul că este garantat că se va executa fără intervenții din partea altor
fire de execuție. În al te cuvinte, nici un alt fir de execuție nu poate accesa această adresă de
memorie până când operația nu s -a terminat” . [2] Funcțiile atomice au în general doi parametri :
adresa de memorie la care se găsește valoarea care trebuie modificată și valoarea cu care se
modifică. Funcțiile atomice sunt:
– atomicAdd folosit pentru operații de sumă ;
– atomicSub folosit pentru operații de scădere ;
– atomicExch folosit pentru a inter schimba două valori ;
– atomicMin folosit pentru a compara două valori; valoarea mai mică este stocată înapoi
în memorie ;

16
– atomicMax folosit pentru a compara două valori; valoarea mai mare este stocată înapoi
în memorie ;
– atomicInc folosit pentru a incr ementa valoarea stocată la adresa de memorie cu 1 dacă
e mai mică decât al doilea parametru, altfel se salvează 0 la respectiva adresă de
memorie ;
– atomicDec folosit pentru a decrementa valoarea stocată la adresa de memorie cu 1 dacă
e diferită de 0 și este mai mică sau egală cu valoarea dată ca al doilea parametru, altfel
se salvează valoarea dată ca al doilea parametru la respectiva adresă de memorie ;
– atomicCAS singura funcție atomică cu trei parametrii: adresa de memorie, valoarea cu
care se verifică dacă valoarea prezentă curent în memorie este egală și valoarea care se
salvează în memorie dacă condiția de egalitate este adevărată ;

1.8. Runtime API vs . Driver API la nivel de host

CUDA Driver API a fost inițial singurul API pentru CUDA și a introdus atât biblioteci
pentru s crierea de kernel cât și biblioteci care să permită utilizarea device -ului de către host pentru
a facilita integrarea CUDA în aplicații C/C ++. Cu toate acestea, bibliotecile destinate pentru a fi
folosite pe partea de ho st sunt de tip „ low level ” și necesită pași în plus pentru apelarea de kernel –
uri. Ca răspuns, Nvidia a lansat mai târziu CUDA Runtime API pentru aplicații C/C++. Acest nou
API păstrează toate bibliotecile prezente în Driver API dar oferă noi biblioteci intenționate a fi
folosite la nivel de host, biblioteci care au o abordare mult mai „high level ”. Cu alte cuvinte, codul
sursă a uni kernel poate fi utilizat atât pe Driver API cât și pe Runtime API fără a fi nevoie de nici
o modificare de cod și poate fi folosit chiar și în alte limbaje de programare pentru care există
librării de tip binding. Diferențele apar la partea de ho st și mai exact felul în care host -ul apelează
kernel-ul pentru a rula pe device. Dezavantajul la Runtime API este că momentan este disponibil
doar pentru limbajele de programare C și C++ .
După cum spuneam, diferențele sunt la nivelul host -ului. Principala diferență constă în felul
în care se face management -ul contextului aplicației. Atunci când o aplicație folosește procesorul
grafic, la nivelul host -ului se creează un context pentru device iar asta presupune rezervarea
resurselor oferite de procesorul gr afic către acea aplicație. Dacă aplicația se oprește, resursele nu
mai sunt rezervate și devin disponibile către alte aplicații în funcție de nevoie. Runtime API
realizează inițializarea contextului automat și practic limitează programatorul să folosească un
singur context per aplicație și astfel se elimină o bună parte din codul asociat utilizării device -ului
de către host. În Device API pe de altă parte, contextul este inițializat manual, motiv pentru care la
nivelul aplicației se pot inițializa contexte multiple care presupune că o aplicație își poate crea
procese multiple, unde fiecare proces are propriul context care folosește device -ul. De exemplu s –
ar putea împărți resursele procesorului grafic în 3 contexte diferite într -o singură aplicație și fiecar e
context ar fi responsabil cu a rezolva o problemă diferită. Dar în general o astfel de abordare este
de evitat deoarece management -ul memoriei device -ului ar fi foarte greu de realizat în acest caz,
motiv pentru care o serie de erori ar putea să apară. M ai mult, pe device nu pot fi rulate mai multe
kernel -uri în paralel, motiv pentru care nu se justifică utilizarea mai multor contexte într -o aplicație.

17
2. Utilizarea platformei CUDA în aplicații Java

În acest e capitole voi prezenta tehnologiile Java utilizat e pentru a fi posibil integrarea
platformei CUDA în proiecte Java.

2.1. Limbajul de programare Java

Java este un limbaj de programare orientat obiect lansat în 1995 de către Sun Microsystems
(acum filială Oracle) . Acest limbaj este compilat în byte code care ru lează pe JVM (Java Virtual
Machine) indiferent de arhitectura calculatorului sau sistemul de operare. Sintaxa limbajului este
inspirată din limbajul C++ dar spre deosebire de acesta oferă puține facil ități de tip „ low level ”. În
prezent, Java este unul dintre cele mai populare limbaje de programare fiind prezent într -o arie
largă de aplicații precum aplicații desktop, web, mobile, embedded .

2.2. Maven

Maven8 este un sistem de build și gestionat dependențele în materie de librării pentru
limbajul de programare Java. Constă într -un fișier de co nfigurare XML denumit pom.xml unde
sunt descrise toate dependențele folosite de proiect, cât și plugin -urile folosite la realizat build -ul
proiectului. Maven este foarte util deoarece face integrarea anumitor librării mult mai facilă și
asigură independență față de IDE.

2.3. JCuda

JCuda este o librărie pentru Java care oferă o implementare de tip binding pentru a utiliza
platforma CUDA. Implementarea de tip binding este o metodă de a permite unui limbaj de
programare să utilizeze o librărie sau un API scris în alt limbaj de programare . Pentru platforma
Java putem implementa soluții de tip binding utilizând JNI (Java Native Interface), un framework
care permite codului Java care rulează pe JVM (Java Virtual Machine) să apeleze metode
implementate într -un alt limbaj de programare precum C, C ++ sau Assembly.
Integrarea librării JCuda presupunea inițial pe lâng ă adăugarea fișierelor .jar în proiect și
includerea anumitor fișiere .dll dar în prezent, începând cu versiunea 8.0, JCu da este disponibil și
ca dependență de tip Maven și astfel integrarea este mult mai facil ă și nu mai este nevoie de
repectivele fișiere .dll.
JCuda permite rularea de kernel -uri CUDA în aplicații Java utilizând CUDA Driver API.
Totuși, deoarece compilatorul NVCC nu poate fi utilizat la nivelul platformei Java, intervine un
pas în plus și anume kernel -ul rulat de către JCuda trebuie să fi e scris în cod PTX, acesta fiind ușor
de rulat pe GPU datorită compilării de tip JIT. Astfel, fișierele .cu care conțin cod sursă pentru
kernel trebuie compilate în fișiere .ptx înainte de a fi uti lizate de către JCuda. D e exemplu , kernel –
ul prezent în t abelul 1.2 . din capitolul 1.5 care primește ca input un tablou unidimensional și
mărimea tabloului și ridică la pătrat fiecare element din tablou. Acesta este salvat în fișierul

8 https://en.wikipedia.org/wiki/Apache_Maven

18
kernel_computeSquare.cu și la începutul fișierului se adaugă extern “C” precum în tabelul 2.1.
Pentru a rula acest kernel în Java , va trebui compilat în fișier .ptx utilizând comanda nvcc -ptx
nume_fișier.cu –o nume_fișier.ptx precum în figura 2.1.

extern "C"

__global__ void computeSquare( float * a, int size) {

int index = threadIdx.x + blockIdx.x * blockDim.x;
if (index < size)
a[index] = a[index] * a[index];

}

Tabelul 2.1 . Conținutul f ișierul ui kernel_computeSquare.cu

Figura 2.1 . Compilarea unui kernel din fișier .cu în fișier .ptx

.version 5.0
.target sm_20
.address_size 64

// .globl computeSquare

.visible .entry computeSquare(
.param .u64 computeSquare_param_0,
.param .u32 computeSquare_param_1
)
{
.reg .pred %p<2>;
.reg .f32 %f<3>;
.reg .b32 %r<6>;
.reg .b64 %rd<5>;

ld.param.u64 %rd1, [computeSquare_param_0];
ld.param.u32 %r2, [computeSquare_param_1];
mov.u32 %r3, %tid.x;
mov.u32 %r4, %ntid.x;
mov.u32 %r5, %ctaid.x;
mad.lo.s32 %r1, %r4, %r5, %r3;
setp.ge.s32 %p1, %r1, %r2;
@%p1 bra BB0_2;

cvta.to.global.u64 %rd2, %rd1;
mul.wide.s32 %rd3, %r1, 4;
add.s64 %rd4, %rd2, %rd3;
ld.global.f32 %f1, [%rd4];
mul.f32 %f2, %f1, %f1;
st.global.f32 [%rd4], %f2;

BB0_2:
ret;
}

Tabelul 2.2 . Conținutul fișierului kernel_computeSuqre.ptx

19
Codul din tabelul 2.2 . poate fi apelat de către o aplicație Java prin intermediul metodelo r
oferite de clasa JCudaDriver din librăria JCuda. Această clasă oferă și metode pentru a aloca spațiu
de memorie pe procesorul grafic și de a realiza transferul de date între ho st și device sau inv ers. În
tabelul 2.3 . avem un fragment de cod Java în care este utilizată clasa JCudaDriver pentru a rula
kernel -ul găsit în kernel_computeSquare.ptx pe device pentru un tablou unidimensional salvat în
hostArray.

CUmodule module = getCUmodule( "src/main/resources/kernel_computeSquare.ptx" );
CUfunction ptxKernelFunction = getCUfunction(module, "computeSquare" );

CUdeviceptr deviceArray = new CUdeviceptr();
JCudaDriver. cuMemAlloc (deviceArray, hostArray. length * Float. SIZE);
JCudaDriver. cuMemcpyHto D(deviceArray, Pointer. to(hostArray),
hostArray. length * Float. SIZE);
Pointer kernelParameters = Pointer. to(
Pointer. to(deviceArray),
Pointer. to(new int[] {size})
);
JCudaDriver. cuLaunchKernel (ptxKernelFunction,
3, 1, 1,
6, 1, 1,
0, null,
kernelParameters, null);
JCudaDriver. cuCtxSynchronize ();
JCudaDriver. cuMemcpyDtoH (Pointer. to(hostArray), deviceArray,
hostArray. length * Float. SIZE);
JCudaDriver. cuMemFree (deviceArray);

Tabelul 2.3 . Fragment de cod Java care utilizează clasa JCudaDriver

2.4. Apache POI

Apache POI [4] este o librărie Java de tip open -source care oferă un API pentru generarea
de fișiere Microsoft Office. Am inclus această librărie în proiect pentru a salva statistici despre
timpii de execuție ai CPU -ului și ai GPU -ului în funcție de inputul primit. Aceste statistici sunt
salvate în fișiere Excel de tip .xlsx . Fișierele generate pot fi utilizate mai departe pentru generarea
de benchmark -uri, foarte util e pentru a înțelege felul de abordare a programării utilizând procesorul
grafic.

20
3. Implementare ș i benchmark -uri

În acest capitol voi prezenta algoritmii pe care i -am implementat folosind JCuda pentru a
observa diferențele de performanță între CPU și GPU. Pentru fiecare algoritm voi expune
benchmark -urile obținute. Hardware -ul pe care s -au realizat benchmark -urile îl reprezintă un
procesor Intel i5 dual -core cu 2.5 GHz, 8 GB memorie RAM și un procesor grafic Nvidia 940MX.

3.1. Prezentarea proiectului

Proiectul este implementat ca proiect de tip Maven din două motive. În primul rând este
asigurată independența față de IDE, pentru a seta proiectul fiind necesar doar descărcarea sa și
importarea ca proiect Maven în IDE -ul folosit. Al doilea motiv este implementarea librării JCuda
ca dependențe de Maven deoarece astfel nu mai este nevoie de adăugat în proiect fișierel e .dll
necesare librării JCuda.

Figura 3.1. Struc tura proiectului

Fișierul pom.xml conține toate dependențele JCuda, setările pentru build și alte setări la
nivel de configurare a proiectului. Folder -ul src/main/java conține codul sursă Java iar folder -ul
src/main/resources conține toate fișierele de tip .cu și .ptx folo site în proiect. În pachetul
cudaExecutors se găsesc clasele care utilizează procesorul grafic. Foarte importantă este clasa
AbstractCudaExecutor deoarece toate clasele care se folosesc de procesorul grafic extind această
clasă abstractă.

21
public abstract class AbstractCudaExecutor {

private static final CUdevice GPU;
private static final CUcontext context;

static {
System.load("C:\\Program Files \\NVIDIA GPU Computing
Toolkit\\CUDA\\v8.0\\bin\\curand64_80.dll" );
JCudaDriver. cuInit(0);
GPU = new CUdevice();
JCudaDriver. cuDeviceGet (GPU, 0);
context = new CUcontext();
JCudaDriver. cuCtxCreate (context, 0, GPU);
System.out.println( "\nGPU context created successfully \n");
}

protected CUmodule getCUmodule(String ptxFileName) {
CUmodule module = new CUmodule();
cuModuleLoad (module, ptxFileName);
return module;
}

protected CUfunction getCUfunction(CUmodule module, String functionName) {
CUfunction function = new CUfunction();
cuModuleGetFunction (function, module, functionName);
return function;
}

public abstract void execute() throws Exception;

}

Tabelul 3.1 . Clasa AbstractCudaExecutor

În primul rând, clasa AbstractCudaExecutor se ocupă de inițializarea device -ului și de creat
contextul , acestea fiind realizate la prima instanțiere a unei clase care extinde această clasă.
Contextul creat este unul singur și rămâne activ pe durata aplicației. Metodele getCUmod ule și
getCUfunction sunt pentru a încărca fișiere de tip .ptx care conțin cod sursă pentru kernel. În final
metoda abstractă execute este destinată pentru a fi folosită ca metoda în care se lansează kernel -ul
pentru a rula pe device.

22
3.2. Exemple de algoritmi pentru înțelegerea conceptelor de
programare CUDA

Unul dinte primii algoritm i pe care i -am implementat pentru a observa performanța CPU
vs. GPU a fost un algoritm de găsire a elementului cu valoarea maximă dintr -un tablou
unidimensional. Din pun ct de vedere al programării secvențiale , implementarea este una simplă:
se parcurg toate elementele și se compară fiecare cu elementul maxim găsit până în prezent.

extern "C"

__global__ void find_maximum_kernel( float *array, float *max, int *mutex, unsigned
int n) {
unsigned int index = threadIdx.x + blockIdx.x*blockDim.x;
unsigned int stride = gridDim.x*blockDim.x;
__shared__ float cache[1024];

float temp = -1.0;
for (int i = index; i < n; i = i + stride) {
temp = fmaxf(temp, array[i]);
}

cache[threadIdx.x] = temp;
__syncthreads();

unsigned int i = blockDim.x / 2;
while (i != 0){
if (threadIdx.x < i){
cache[threadIdx.x] = fmaxf(cache[threadIdx.x],
cache[threadIdx.x + i]);
}

__syncthreads();
i /= 2;
}

if (threadIdx.x == 0){
while (atomicCAS(mutex, 0, 1) != 0); //lock
*max = fmaxf(*max, cache[0]);
atomicExch(mutex, 0); //unlock
}
}

Tabelul 3.2 . Codul sursă a kernel -ului pentru găsirea maximului într -un tablou unidimensional

Kernel -ul din tabelul 3.2 reprezintă logica care se execută pe GPU pentru fiecare fir de
execuție lansat. Inputul kernel -ului este alcătuit din tabloul unidimensional pentru care trebuie găsit
elementul cu valoarea maximă, pointer către adresa de memorie unde va fi salvată valoare a
maximă, pointer către adresa de memorie folosit pentru valoarea de mutex folosită la sincronizarea
firelor de execuție și ultimul parametru îl reprezintă dimensiunea tabloului.
Inițial se calculează pentru fiecare fir de execuție index -ul său în cadrul g rid-ului iar apoi
se calculează valoarea stride reprezentând numărul total de fire de execuție din grid. Fiecare fir de
execuție parcurge apoi tabloul din „ stride în stride ” elemente și salvează în variabila temp valoarea
maximă găsită, adică dacă grid -ul este format din 1024 de blocuri a câte 1024 de fire de execuție
atunci valoarea variabilei stride este 1 048 576. Prin urmare fiecare fir de execuție sare din 1 048
576 elemente în 1 048 576 elemente atunci când parcurge tabloul unidimensional. Elementul
maxim găsit de către fiecare fir de execuție este salvat în memoria partajată la care acces. Memoria

23
partajată pentru fiecare bloc conține un tablou unidimensional de 1024 elemente deoarece kernel –
ul este rulat cu blocuri a câte 1024 de fire de execuție, unde fiecare element reprezintă maximul
găsit de fiecare fir de execuție din bloc. Procedura __syncthreads() asigură că nici un fir de execuție
din cadrul unui bloc nu -și continuă execuția până când fiecare fir de execuție din bloc a terminat
de scris în m emoria partajată.
După ce memoria partajată a fost scrisă, firele de execuție încep să o utilizeze pentru a găsi
elementul maxim pentru fiecare bloc, element care ajunge să fie salvat pe poziția 0 a fiecărui cache .
Astfel, se obține elementul maxim pentru fiecare bloc. În final, firele de execuție care au index -ul
0 în cadrul blocului stabilesc valoarea maximă a tabloului pe baza operațiilor atomic.

Figura 3.2 . Logica kernel -ului pentru găsit elementul maxim

Acest kernel este apelat în clasa FindMaxCudaExecutor unde se realizează și transferul
parametrilor de input de la host către device. Dimensiune grid -ului cu care este apelat kernel -ul
este de 1024 de blocuri a câte 1024 de fire de execuție, adică 1048576 fire de execuție în total.
Pentru a obține un benchmark pe baza acestui algoritm am folosit tablouri unidimensionale
cu dimensiuni cuprinse între 1000000 de elemente și 50000000 de elemente cu numere pozitive
generate aleatoriu.

Figura 3.3 . Benchmark pentru găsirea elementului maxim într-un tablou unidimensional

24

În figura 3.3 . avem raportul de performanță pentru CPU, GPU și GPU la care se adaugă
timpul de transfer al datelor de la host la device. Putem observa că inițial pentru un număr mai mic
de elemente, diferențele sunt minuscule. Pe parcurs ce crește numărul de elemente, timpul de
execuție pentru CPU și GPU cresc constant deși timpul pentru GPU rămâne mai mic decât cel
pentru CPU. Putem trage ast fel concluzia că GPU -ul este mai rapid în acest caz dacă ne referim
strict l a rularea kernel -ului. Pe de altă parte diferența dintre timpul de execuție reprezentat de
transferul datelor către device la care se adaugă și timpul de execuție a kernel -ului și timpul de
execuție al CPU -ului crește foarte mult pe parcurs ce numărul date lor crește. Prin urmare deși
procesorul grafic ajunge mai repede la rezultat decât CPU -ul, transferul de date pe care -l presupune
face nerentabilă o implementare pe bază de GPU.

extern "C"

__global__ void find_element_kernel( float *array, float *element, int *elementIndex,
unsigned int n) {
unsigned long index = threadIdx.x + blockIdx.x*blockDim.x;
unsigned long stride = gridDim.x*blockDim.x;

while (index < n && *elementIndex > index){
if (array[index] == *element) {
atomicMin(elementIn dex, index);
}
index += stride;
}
}

Tabelul 3.3 . Codul sursă al kernel -ului pentru găsirea primei apariții a unui element într –
un tablou

Un al t algoritm implementat pentru a rula pe GPU a fost căutarea unui element într -un
tablou unidimensional și aflarea indexului primei apariții. În mare parte acest algoritm este
asemănător cu cel precedent din punct de vedere al timpului de execuție dar avantajează mai mult
procesorul decât primul algoritm în funcție de unde se găsește elementul. Kernel -ul pentru acest
algoritm este în tabelul 3.3.

Figura 3.4. Benchmark pentru găsirea primei apariții a unui element într -un tablou

25
Benchmark -ul obținut scoate în evidență ac eeași problemă ca în benchmarl -ul precedent.
Deși procesorul grafic are potențial de a ajunge mai rapid la rezultat, transferul de date de la host
la device este prea costisitor. Mai mult, se poate vedea că diferențele mari între GPU și CPU sunt
mari doar în anumite cazuri în care elementul se găsește pe la finalul tabloului.

3.3. Utilizarea GPU -ului în cadrul algoritmilor Monte Carlo

Algoritmii Monte Carlo reprezintă o categorie de algoritmi probabiliști care estimează
rezultatul unei probleme cu o probabilitate de eroare strict mai mare decât 0. Acest tip de algoritm
este ideal atunci când pentru o problemă dată nu dispunem de o soluț ie pe baza unui algoritm de
tip determinist. Denumirea de „Monte Carlo” provine de l a subdiviziunea orașului -stat Monaco
renumită pentru cazinourile sale.
Motivul pentru care acest tip de algoritm poate beneficia de puterea de calcul paralel oferită
de GPU constă în faptul că probabilitatea de eroare scade cu fiecare rulare a algoritmu lui. Evident
că aceste rulări multiple pentru un algoritm de acest tip duc la creșterea timpului de execuție dar
această creștere se datoreaz ă rulării î n mod secvențial . Abordarea din perspectiva programării
paralele ne permite să micșorăm această rată de eroare dar fără să mărim timpul de execuție
deoarece aceste rulări se vor face în paralel iar la final putem agrega rezultatele pentru a obține o
estimare mai bună a re zultatului.

3.3.1. Utilizarea algoritmului Monte Carlo în domeniul bursei
Pentru a exemplifica implementarea scenariilor Monte Carlo, am ales un exemplu din
domeniul financiar , mai exact din domeniul opțiunilor de bursă .
Să considerăm următoarea problemă: un cumpărător achiziționează de la un vânzător o
opțiune pentru bursă (din eng. stock option ) pe baza unor speculații că în viitor prețul opțiunii va
crește. La achiziționare se încheie u n contract cu următoarele reguli:

 Cumpărătorul are voie să vândă opț iunea la data scadenței , care este estimată la un
an
 Pe parcursul evoluției prețului până la data scadenței, dacă se întâmplă ca prețul
opțiunii să scadă sub limita minimă stabilită la încheierea contractului, atunci
opțiunea se anulează și profitul cumpăr ătorului va fi 0.

Acest tip de opțiune se numește opțiune de tip barieră „ Down and Out” , unde barier a
reprezintă o limită minimă pentru prețul opțiunii . Dacă această limită este atinsă, atunci opțiunea
este anulată iar profitul este 0. Dacă limita nu se atinge, atunci la data scadenței se calculează
profitul ca diferența dintre prețul opțiunii la data scadenței și pr ețul de la care opțiunea poate fi
considerată profitabilă numit și strike . De obicei prețul de strike este egal cu prețul achiziției
opțiunii.
Opțiunile de tip barieră sunt în total de patru feluri:
– Down and Out descries anterior
– Up and Out atunci când o opțiune este anulată dacă atinge o limită maximă de preț
– Down and In atunci când o opțiune devine valabilă dacă atinge o limită minimă de preț
– Up and In atunci când o opțiune devine valabilă dacă atinge o limită maximă de preț

26
Toate patru tipurile de opțiuni diferă doar din perspectiva condiției impuse de barieră, motiv
pentru care implementările lor sunt destul de asemănătoare.
Din pe rspectiva algoritmului Monte Carlo, obiectivul este de a calcula profitul la data
scadenței imediat după ce opțiunea a fost achiziționată. Ca date de input avem:
– prețul inițial la care opțiunea a fost achiziționată
– prețul minim de la care opțiunea poate fi considerată profitabilă ( strike) ; în multe cazuri
acesta este egal cu prețul inițial al opțiunii
– valoarea barierei care reprezintă un preț care odată atins anulează opțiunea în cazul
opțiunilor de tip „Out” sau care face opțiunea valabilă în cazul opțiunilor de tip „In”
– Timpul până l a data scadenței (pentru exemplul prezentat în această lucrare s -a utilizat
timpul până la scadență egal cu un an)
– Randamentul anual sau profitul anual asociat opțiunii
– Volatilitatea sau procentul cu care variază prețul opțiunii

Pentru a calcula prețul se utilizează formula:

𝑃𝑛= 𝑃𝑛−1+𝑟𝑎𝑛𝑑𝑎𝑚𝑒𝑛𝑡𝑢𝑙 𝑎𝑛𝑢𝑎𝑙 ∗ 𝑃𝑛−1∗ ∆𝑡+𝑅𝑎𝑛𝑑 𝑛−1

𝑃𝑛 reprezintă prețul care trebuie calculat și 𝑃𝑛−1 ultimul preț calculat . ∆𝑡 reprezintă
diferența de ti mp între două prețuri calculate. 𝑃0 reprezintă prețul de achiziție iar în cazul
exemplului prezent în această lucrare , 𝑃364 reprezintă prețul la data scadenț ei.
𝑅𝑎𝑛𝑑 𝑛−1 este un numărul generat aleatoriu pe baza unei distribuției de tip Gauss cu media
0 și deviația standard calculată pe baza volatilității. Aceste numere generate reprez intă cu cât crește
sau scade noul preț calculat . Deviația standard este calculată pe mulțimea de numere întregi din
intervalul:

[(−1)∗𝑣𝑜𝑙𝑎𝑡𝑖𝑙𝑖𝑡𝑎𝑡𝑒 ∗𝑃0
2,𝑣𝑜𝑙𝑎𝑡 𝑖𝑙𝑖𝑡𝑎𝑡𝑒 ∗𝑃0
2 ]

După ce fiecare rulare a algoritmului a produs un rezultat se va face o medie aritmetică
asupra acestor rezultate, media reprezentând rezultatul final cu o probabilitate de eroare mai mică
decât a unei rulări individuale. Fiecare rulare a algoritmului se numește și simulare Monte Carlo
deoarece reprezintă practic o simulare a evoluției prețului opțiunii până la data scadenței.

3.3.2. Implementare folosind JCuda
Pachetul cudaExecutors.montecarlo conține implementările pentru toate cele patru tipuri
de opțiuni și anume clasele :
– MonteCarloDownAndInBarrier
– MonteCarloDownAndOutBarrier
– MonteCarloUpAndInBarrier
– MonteCarloUpAndOutBarrier

Toate cele patru clase extind clasa abstractă AbstractCudaExecutor și suprascriu metoda
execute() , care conține codul care utilizează procesorul grafic . Fiecare dintre aceste patru clase
conține și o metodă numită executeCPU() care oferă o implementare clasică care nu folosește nici
un kernel pentru G PU. Acea stă metodă este folosită pentru a putea la final compara diferența dintre
timpul de execuție al CPU -ului cu cel al GPU -ului.

27
private BarrierOptionInputData inputData ;
private int monteCarloTotalSimulations ;
private long gpuTimeStart ;
private long gpuTimeEnd ;
private long cpuTimeStart ;
private long cpuTimeEnd ;
private float gpuTime;
private float cpuTime;
private float gpuMonteCarloResult ;
private float cpuMonteCarloResult ;

public MonteCarloDownAndOutBarrier(BarrierOptionInputData inpu tData, int
nrOfSimulations) {
this.inputData = inputData;
this.monteCarloTotalSimulations = nrOfSimulations;
}

Tabelul 3.4 . Atributele clasei MonteCarloDownAndOutBarrier și constructorul său

Fiecare dintre aceste patru clase are atributele ca re se pot observa în tabelul 3.4. și anume
o un obiect de tipul BarrierOptionInputData care conține toate atributele asociate unei opțiuni,
date comune pentru fi ecare execuție a algoritmului Monte Carlo.
Clasa BarrierOptionInputData conține atributele asociate unei bariere, cât și metodele se
tip setter și getter . Aceste atribute sunt prețu l inițial al opțiunii, prețul de strike , valoarea barierei,
timpul de atingere a maturității (data scadenței) exprimat în ani, randamentul anual și volatilitatea .
Această clasă c onține și metoda care calculează deviația standard pentru opțiunea în cauză.

final int NUMBER_OF_PRICE_COMPUTATIONS = 365;

final int DEVICE_MEMORY_SIZE =
monteCarloTotalSimulations * NUMBER_OF_PRICE_COMPUTATIONS;

float timeDifference = inputData .getMaturityTimeInYears()/
NUMBER_OF_PRICE_COMPUTATIONS;
float standardDeviation = inputData .computeStandardDeviation();

final String PTX_FILE = "src/main/resources/kern el_montecarlo_downandout.ptx" ;
final String FUNCTION_NAME = "mc_kernel" ;

JCuda.setExceptionsEnabled (true);
JCurand. setExceptionsEnabled (true);

CUmodule module = getCUmodule(PTX_FILE);
CUfunction function = getCUfu nction(module, FUNCTION_NAME);

CUdeviceptr devicePrices = new CUdeviceptr();
JCudaDriver. cuMemAlloc (devicePrices, monteCarloTotalSimulations * Sizeof. FLOAT);

Pointer deviceInputData = getCUDAGeneratedInput(DEVICE_MEMORY_SIZE,
standardDeviation);

Tabelul 3.5. Partea de cod din metoda execute()

În cele care urmează voi descrie implementarea metodei execute() pentru clasa
MonteCarloDownAndOutBarrier , implementările pentru celelalte 3 clase fiind foa rte
asemănătoare. În tabelul 3.5 . este partea de cod care reprezintă setarea device -ului și obținerea
datelor generate aleatoriu. Variabilele finale PTX_FILE și FUNCTION_NAME reprezintă locația

28
fișierului de tip .ptx care conține codul sursă a kernel -ului în cod PTX și respectiv numele funcției
de kernel. Ap elul metodei setExceptionsEnabled cu parametrul true atât pentru JCuda cât și pentru
JCurand permit aruncarea de excepții în caz că intervine o situație neprevăzută la nivel de procesor
grafic care poate compromite rezultatul obținut. Variabilele module și function sunt inițializate
pentru fi folosite la lansarea kernel -ului. În final, declarăm și inițializăm variabila devicePrices care
reprezintă o abstracțiune a pointer -ului către adresa de memorie globală a GPU -ului unde se vor
stoca rezultatele fiecărei simulări. Acest spațiu de memorie este alocat prin metoda statică
cuMemAlloc având o mărime exprimată în bytes. Constanta Sizeof.FLOAT este egală cu 4.

private Pointer getCUDAGen eratedInput( int deviceMemorySize, float standardDeviation)
{
Pointer deviceInputData = new Pointer();
JCuda.cudaMalloc (deviceInputData, deviceMemorySize * Sizeof. FLOAT);

//Generate random values
curandGenerator generator = new curandGenerator();
JCurand. curandCreateGenerator (generator, CURAND_RNG_PSEUDO_MTGP32 );
JCurand. curandSetPseudoRandomGeneratorSeed (generator, 1234L);
JCurand. curandGenerateNormal (generator, deviceInputData, deviceMemorySize, 0.0f,
standardDeviat ion);

JCurand. curandDestroyGenerator (generator);

return deviceInputData;
}

Tabelul 3.6 . Metoda getCUDAGeneratedInput

Metoda privată getCUDAGeneratedInput se găsește tot în interiorul clasei
MonteCarloDownAndOutBarrier și are rolul de a genera datele folosite pentru rularea simulărilor
Monte Carlo, de a le stoca în memoria globală a device -ului. Clasa JCurand reprezintă o
implementar e de tip binding pentru biblioteca CURAND din CUDA, librărie folosită pentru a
genera numere aleatoriu. Generatorul utilizat în această metodă este unul de tip pseudo random
care conform documentației oficiale Nvidia este un algoritm care face parte din „familia Mersenne
Twister de algoritmi pentru generare pseu do random ” [2]. Fiind un algoritm pseudo random, este
folosită metoda curandSetPseudoRandomGeneratorSeed pentru a stabili valoarea de seed pentru
generatorul de numere. În final se apelează metoda statică curandGenerateNormal pentru a genera
datele conform unei distribuții de tip Gauss (distribuție norm ală). Această metodă primește ca
parametrii generatorul folosit, pointer -ul către adresa de memorie la care vor fi salvate datele,
media distribuții care în cazul nostru va fi m ereu 0 și deviația standard a distribuției. La final,
generatorul este distrus pentru a elibera din spațiul de memorie al procesorului grafic și se
returnează obiectul de tip Pointer, deviceInputData .

Revenind la metoda execute() avem ilustrat partea de cod pentru lan sarea kernel -ului în
tabelul 3.7 . Inițial se stabilesc parametrii kernel -ului prin variabila kernelParameters .
Pentru dimensiunea grid -ului se utilizează blocuri de câte 1024 de fire de execuție și
numărul de blocuri este calculat î n funcție de numărul de rulări ale algoritmului . Metoda statică
cuLauchKernel din clasa JCudaDriver lansează kernel -ul pentru a rula pe device folosind
parametrii setați. După execuția kernel -ului, se calculează media aritmetică a acestor rezultate. În
final se elib erează memoria de pe device care a fost utilizată pentru rularea kernel -ului.

29

Pointer kernelParameters = Pointer. to(
Pointer. to(devicePrices),
Pointer. to(new float []{inputData .getStrikePrice()}),
Pointer. to(new float []{inputData .getBarrierPrice()}),
Pointer. to(new float []{inputData .getInitialPrice()}),
Pointer. to(new float []{inputData .getVolatility()}),
Pointer. to(new float []{inputData .getAnnualDrift()}),
Pointer. to(new float []{timeDifference}),
Pointer. to(deviceInputData),
Pointer. to(new int[]{NUMBER_OF_PRICE_COMPUTATIONS }),
Pointer. to(new int[]{monteCarloTotalSimulations })
);

final int BLOCK_SIZE = 1024;
final int GRID_SIZE = ( int)Math.ceil((double) monteCarloTotalSimulations /
BLOCK_SIZE);
JCudaDriver. cuLaunchKernel (function,
GRID_SIZE, 1, 1,
BLOCK_SIZE, 1, 1,
0, null,
kernelParameters, null);

Tabelul 3.7 . Partea de cod a metodei execute() care lansează rularea kernel -ului

3.3.3. Descrierea kernel -ului
În folder -ul resources găsim atât fișierele de tip .cu pentru kernel cât și cele de tip .ptx. Deși
la rularea aplicației Java, sunt utilizate fișierele de tip .ptx, acestea sunt obținute compilând fișierele
de tip .cu. Pe lângă acestea, codul din fișierele .cu poate fi înțeles ușor fiind vorba de cod C. Pentru
fiecare tip de opțiune avem un kernel separat .

extern "C"

__global__ void mc_kernel(
float * d_prices,
float strike,
float barrier_value,
float initial_price,
float volatility,
float annual_drift,
float dt,
float * d_input_data,
unsigned int numberOfDaysUntilMaturity,
unsigned int MONTE_CARLO_TOTAL_SIMULATIONS) {

int simulation_idx = threadIdx.x + blockIdx.x * blockDim.x;
int input_data_idx = threadIdx.x + blockIdx. x * blockDim.x;

float s_curr = initial_price;
if (simulation_idx < MONTE_CARLO_TOTAL_SIMULATIONS) {
int day = 0;
do {
s_curr = s_curr + annual_drift * s_curr * dt +
d_input_data[input_data_idx];
input_data_idx++;
day++;
}
while (day < numberOfDaysUntilMaturity && s_curr > barrier_value);
double payoff = (s_curr > strike ? s_curr – strike : 0.0);
__syncthreads();

30
d_prices[simulation_idx] = payoff;
}
}

Tabelul 3.8 . Kernel -ul pentru opțiunea de tip „ down and out ”

În continuare voi descrie kernel -ul pentru opțiunea de tip barieră de tip „ down and out ”.
Kernel -ul este conceput încât fiecare fir de execuție reprezintă o rulare a algoritmului Monte Carlo .
Variabila de tip int simulation_idx reprezintă indexul simulării iar variabila input_data_index
reprezintă indexul de la care firul de execuție începe să parcurgă tabloul unidimensional
d_input_data care conține numerele generate aleatoriu. Ambele se calculează în același mod și
reprezintă practic indexul firului de execuție în cadrul grid -ului.
Fiecare fir de execuție verifică la început dacă indexul simulării este mai mic decât numărul
total de sim ulări deoarece putem avea un număr de fire de execuție mai mare decât numărul de
simulări.
Dacă condiția este îndeplinită începe logica efectivă a simulării. Mai exact, se calculează
prețul curent pentru fiecare zi până la atingerea nivelului de maturitat e cât timp prețul curent este
mai mare decât prețul impus de barieră. Dacă prețul scade sub barieră nu se mai justifică să
continuăm deoarece rezultatul fina l va fi 0, opțiunea fiind anulată. După terminarea structurii de
tip do -while, se calculează valoar ea finală doar în cazul în care prețul curent din ultima zi este mai
mare decât valoarea de strike , altfel rezultatul este 0.
Înainte de a se scrie rezultatele în d_prices se așteaptă ca fiecare fir de execuție să termine
de ca lculat .

3.3.4. Clasa ReduceCudaExecutor
Clasa ReduceCudaExecutor este folosită la partea de agregare a da telor și mai exact
reprezintă o implementare a operației de adunare folosind procesorul grafic. Deși suma chiar și a
unei mulțimi mari de numere nu s -ar justifica pentru folosirea procesorului grafic având în vedere
că toate aceste numere trebuie transferate în memoria proc esorului grafic, în cazul acestei
implementări numerele sunt deja în memoria globală deci prin urmare putem obține un spor de
performanță printr -o astfe l de implementare.

public class ReduceCudaExecutor extends AbstractCudaExecutor {

private int numberOfElements ;
private CUdeviceptr deviceData ;
private float [] result;

public ReduceCudaExecutor( int numberOfElements, CUdeviceptr deviceData) {
this.numberOfElements = numberOfElements;
this.deviceData = deviceData;
}

public float [] getResult() {
return result;
}

@Override
public void execute() {
final String PTX_FILE_REDUCE = "src/main/resources/reduce.ptx" ;
final String FUNCTION_NAME_REDUCE = "global_reduce_kernel" ;
final int BLOCK_SIZE = 1024;
final int GRID_SIZE = ( int)Math.ceil((double)numberOfElements / BLOCK_SIZE);

CUmodule reduceModule = getCUmodule(PTX_FILE_REDUCE);

31
CUfunction reduceFunction = getCUfunction(reduceModule,
FUNCTION_NAME_REDUCE);

CUdeviceptr devicePricesAfterReduce = new CUdeviceptr();
JCudaDriver. cuMemAlloc (devicePric esAfterReduce, Sizeof. FLOAT * GRID_SIZE);
Pointer kernelParameters = Pointer. to(Pointer. to(deviceData ),
Pointer. to(devicePricesAfterReduce),
Pointer. to(new int[]{numberOfElements }));
JCudaDriver. cuLaunchKernel (reduceFunction ,
GRID_SIZE, 1, 1,
BLOCK_SIZE, 1, 1,
Sizeof.FLOAT * 1024, null,
kernelParameters, null);
JCudaDriver. cuCtxSynchronize ();
result = new float [GRID_SIZE];
Pointer resultPointer = Pointer. to(result);
JCudaDriver. cuMemcpyDtoH (resultPointer, devicePricesAfterReduce, GRID_SIZE *
Sizeof.FLOAT);

JCuda.cudaFree (kernelParameters);
JCuda.cudaFree (resultPointer);
JCudaDriver. cuMemFree(devicePricesAfterReduce);
}
}

Tabelul 3 .9. Clasa ReduceCudaExecutor

La fel ca celelalte clase din proiect care utilizează procesorul grafic și această clasă extinde
AbstractCudaExecutor . Atributele clasei sunt numărul de elemente care vor fi adunate, datele din
memoria globală a device -ului care vor fi adunate și rezultatul sumei.
Metoda execute() folosește kernel -ul reduce.ptx care conține procedura
global_reduce_kernel utilizată pentru a rula efectua calculele folosind procesorul grafic. Acest
kernel este lansat cu un număr de blocuri calculat în funcție de numărul de elem ente, fiecare bloc
având 1024 fire de execuție. În plus, tot în parametrii de lansare ai kernel -ului avem alocarea
dimensiunii pentru memoria partajată, mai exact 4096 de bytes pentru fiecare bloc.

extern "C"

__global__ void global_reduce_kernel( float * d_in, float * d_out, unsigned int
inputSize) {

extern __shared__ float sdata[];

int index = threadIdx.x + blockDim.x * blockIdx.x;
int threadId = threadIdx.x;

if (index < inputSize)
sdata[threadId] = d_in[index];
else
sdata[threadId] = 0;
__syncthreads();

for (unsigned int i = blockDim.x / 2; i > 0; i >>= 1) {
if (threadId < i) {
sdata[threadId] += sdata[threadId + i];
}
__syncthreads();
}

if (threadId == 0) {

32
d_out[blockIdx.x] = sdata[0];
}
}
Tabelul 3.10. Kernel -ul utilizat de clasa ReduceCudaExecutor

Fiecare fir de execuție care rulează pe baza kernel -ului din tabelul 3.10. citește un element
din input și îl stochează în memoria partajată. După ce fiecare fir de execuție a scris în memoria
partajată, firele de execuție care se găsesc în prima jumătate a blocului de care aparțin însumează
elementul la care au acces cu elementul oglindit din a doua jumătate a blocului. Acest proces se
repetă până când pe prima poziție a memoriei partajate a fiecărui bloc avem rezultatul sumei
elementelor din blocul respe ctiv. Evident că la fiecare iterație se așteaptă ca toate firele de execuție
să-și termine activitatea pentru a nu apărea conflicte. La final doar primul fir de execuție din cadrul
fiecărui bloc scrie rezultatul pentru blocul respectiv în memoria globală.
De exemplu pentru a însuma un milion de numere se utilizează 977 de blocuri a câte 1024
de fire de execuție. La finalul execuției kernel -ului vom avea 977 de numere care însumate
reprezintă rezultatul adunării pentru cele un milion de numere. Avantajul e c ă reducerea de la un
milion la doar 977 de numere se face într -o singură execuție.

3.3.5. Benchmark -uri
Pentru a real iza benchmark -ul pentru rulările multiple ale algoritmului de tip Monte Carlo
am ales să măsor timpul de execuție începând cu 10000 de rulări pentru estimarea rezultatului unei
opțiuni de tip barieră și m-am oprit la un milion de rulări . Pentru mai mult de un milion de rulări ,
am început să experimentez probleme legate de memorie insuficientă pe procesorul grafic. Timpii
de execuție din cadrul benchmark -ului nu includ și generarea de variabile aleatoare.
Benchmark -ul din figura 3.5 . ne confirmă superioritatea procesorului grafic pentru o astfel
de problemă. Dacă timpul de execuție al CPU -ului ajunge la aproape 4 secunde pentru un milioan
de rulări , timpul de execuție al procesorului grafic rămâne sub jumătate de secundă indiferent de
numărul de rulări .

Figura 3.5 . Benchmark pentru estimarea profitului unei opțiuni utilizând scenarii generate de
rulări multiple ale algoritmului Monte Carlo

33
Pentru a calcula opțiuni multiple, k ernel -ul responsabil de Monte Carlo se execută
secvențial pentru fiecare opțiune dar totuși d upă cum se observă în figura 3.6 . diferența de timp
între GPU și CPU este enormă.

Figura 3.6 . Benchmark pentru estimarea unui număr de opțiuni de tip barieră ; pentru fiecare
opțiune se utilizează un milion de rulări ale algoritmului de tip Monte Carlo

34
Concluzie

La finalul acestei lucrări consider că am reușit să -mi îndeplinesc obiectivul de a aprofunda
acest domeniu aparte al programăr ii CUDA și cum poate fi folosit în dezvoltarea de aplicații Java.
Algoritmii implementați cât și benchmark -urile obținute de pe urma lor ne pot oferi o imagine
asupra avantajelor procesorului grafic cât și a punctelor sale slabe . După cum am observat,
procesorul grafic poate aduce un spor de performanță considerabil în anumite situații în care un
volum larg de date poate fi procesat în paralel . Cu toate acestea însă memoria procesorului grafic
are un rol important per total asupr a performanței. În primul rând complexitatea problemei
executate de fiecare fir de execuție în parte trebuie să justifice transferul de date către memoria
procesorului grafic deoarece de multe ori acest transfer face o implementare care utilizează
procesor ul grafic să fie nerentabilă. În al doilea rând, programator ului îi revine sarcina de a ști cum
să jongleze cu diferitele tipuri de memorie prezente pe procesorul grafic în funcție de problema pe
care încearcă să o rezolve.
Putem concluziona că utilizarea procesorului grafic a ajuns să însemne mai mult decât
domeniul graficii și că în viitor ne putem aștepta la utilizarea într -o arie cât mai vastă de domenii.
Aș vedea ca potențiale direcții de dezvoltare a acestei lucrări explorarea din perspectiva CUDA a
unor domenii precum Big Data sau Învățarea Automată. Astfel de domenii a r putea să beneficieze
enorm de paralelismul oferit de procesoarele grafice pentru a procesa cantități masive de date.

35
Bibliografie

[1] „Single instruction, multiple data” https://ro.wikipedia.org/wiki/SIMD
[2] „CUDA C Programming Guide” https://docs.nvidia.com/cuda/cuda -c-programming -guide/index.html
[3] S. Cook, CUDA Programming: A Developer’s Guide to Parallel Computing with GPUs, Morgan
Kaufmann Publishers
[4] „Apache POI” https://poi.apache.org/
[5] E. Derman și I. Kani, „The Ins and Outs of Barrier Options Part 2” Preluată de pe:
http://www.emanuelderman.com/media/insoutbarriers2.pdf
[6] E. Derman și I. Kani, „The Ins and Outs of Barrier Option Part 1” Preluată de pe:
http://www.emanuelderman.com/media/inso utbarriers1.pdf
[7] „JCuda” http://www.jcuda.org/
[8] C. Frăsinaru, Curs practic de Java. Preluată de pe:
https://profs.info.uaic.ro/~acf/java/Cristian_Frasinaru -Curs_practic_de_Java.pdf