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

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

1
DECLARAȚIE PRIVIND ORIGINALITATE ȘI RESPECTAREA
DREPTURILOR DE AUTOR

Prin prezenta declar că Lucrarea de licență cu titlul „ Utilizarea procesorului grafic (GPU) în aplicații
Java ” este scrisă de mine și nu a mai fost prezentată niciodată la o altă facultate sau instituție de
învățământ superior din țară sau străinătate. De asemenea, declar că toate sursele utilizate, inclusiv
cele preluate de pe Internet, sunt indicate în lucrare, cu respectarea regulilo r de evitare a plagiatului:

 toate fragmentele de text reproduse exact, chiar și în traducere proprie din altă limbă, sunt
scrise între ghilimele și dețin referința precisă a sursei; 

 reformularea în cuvinte proprii a textelor scrise de către alți autori deține referința precisă;

 codul sursă, imagini etc. preluate din proiecte open -source sau alte surse sunt utilizate cu
respectarea drepturilor de autor și dețin referințe precise;

 rezumarea ideilor altor autori precizează referința precisă la textul orig inal.

Iași, Absolvent: [anonimizat]

_________________________

2
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]

_________________________

3
Cuprins
Introducere ………………………….. ………………………….. ………………………….. ………………………….. ………………. 5
Contribuții ………………………….. ………………………….. ………………………….. ………………………….. ……………….. 7
1. Fundamentele teoretice ale programării CUDA ………………………….. ………………………….. …………… 8
1.1. Concepte de bază ale programării pe GPU ………………………….. ………………………….. ……………….. 8
1.2. Arhitectura CUDA ………………………….. ………………………….. ………………………….. ……………………. 9
1.3. Setarea mediului de lucru ………………………….. ………………………….. ………………………….. …………. 11
1.3.1. Microsoft Visual C++ ………………………….. ………………………….. ………………………….. ……….. 11
1.3.2. Nvidia CUDA Toolkit ………………………….. ………………………….. ………………………….. ………. 11
1.4. NVCC (Nvidia CUDA Compiler) ………………………….. ………………………….. ………………………….. 12
1.5. Noțiuni de bază ale programării GPU în CUDA ………………………….. ………………………….. ………. 13
1.5.1. Fire de execuție ( Threads ) și blocuri ………………………….. ………………………….. ……………….. 13
1.5.2. Kernel ………………………….. ………………………….. ………………………….. ………………………….. …14
1.6. Clasificarea memoriei ………………………….. ………………………….. ………………………….. ……………… 16
1.6.1. Registre ………………………….. ………………………….. ………………………….. ………………………….. 17
1.6.2. Memoria locală ………………………….. ………………………….. ………………………….. ………………… 17
1.6.3. Memoria partajată ………………………….. ………………………….. ………………………….. ……………. 17
1.6.4. Memoria globală ………………………….. ………………………….. ………………………….. ……………… 17
1.6.5. Memoria constantă ………………………….. ………………………….. ………………………….. …………… 18
1.7. Sincronizarea firelor de execuție ………………………….. ………………………….. ………………………….. ..18
1.7.1. Proce dura __syncthreads() ………………………….. ………………………….. ………………………….. …18
1.7.2. Pattern -uri de comunicare ………………………….. ………………………….. ………………………….. ….18
1.7.3. Funcții atomice ………………………….. ………………………….. ………………………….. ………………… 20
1.8. Runtime API vs. Driver API la nivel de host ………………………….. ………………………….. …………… 21
2. Utilizarea platformei CUDA în aplicații Java ………………………….. ………………………….. …………….. 22
2.1. Limbajul de programare Java ………………………….. ………………………….. ………………………….. ……. 22
2.2. Maven ………………………….. ………………………….. ………………………….. ………………………….. ……….. 22
2.3. JCuda ………………………….. ………………………….. ………………………….. ………………………….. ………… 22
2.4. Apache POI ………………………….. ………………………….. ………………………….. ………………………….. ..25
3. Implementare și benchmark -uri ………………………….. ………………………….. ………………………….. ……. 26
3.1. Structura pro iectului ………………………….. ………………………….. ………………………….. ………………… 26
3.2. Exemple de algoritmi pentru înțelegerea conceptelor de programare CUDA ………………………… 28
3.3. Scenarii Monte Carlo ………………………….. ………………………….. ………………………….. ………………. 32
3.3.1. Opțiuni de bursă de tip barieră ………………………….. ………………………….. ……………………….. 32
3.3.2. Imple mentare folosind JCuda ………………………….. ………………………….. …………………………. 34
3.3.3. Descrierea kernel -ului ………………………….. ………………………….. ………………………….. ………. 38

4
3.3.4. Clasa ReduceCudaExecutor ………………………….. ………………………….. ………………………….. .39
3.3.5. Benchmark -uri ………………………….. ………………………….. ………………………….. …………………. 41
3.4. Generarea rapoartelor de performanță ………………………….. ………………………….. …………………….. 43
Concluzie ………………………….. ………………………….. ………………………….. ………………………….. ………………… 45
Bibliografie ………………………….. ………………………….. ………………………….. ………………………….. ……………… 46

5
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 ase menea, 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,
acesta 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
crearea 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 deoare ce 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, această 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 procesorului 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 metodei de simulare Monte Carlo și cum implementarea unui algoritm
de o asemenea natură poate beneficia de procesorul g rafic folosindu -se un exemplu
real 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.

6
Într-un final îmi doresc ca prin această lucrare să ofer o bază solidă d e
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 .

7
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. După cum urmează, în prima parte a lucrării vor fi
descrise atâ t conceptele teoretice ale programării bazate pe procesorul grafic cât și
modul practic în care acestea sunt utilizate în CUDA. Pentru a înțelege aceste aspecte
a fost necesar consultarea unor anumite cărți cât și consultarea diverselor resurse de
pe inter net 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 putut să încep să
implementez algoritmi care să utilizeze procesorul grafic. Inițial acești algoritmi au
fost de o natură simplă doar pentru a obține o mai bună perspectivă asupra
programării CUDA . Implementarea scenariilor Monte Carlo a fost făcută pentru a
ilustra un exemplu care poate beneficia considerabil de sporul de performanță al
procesorului grafic.
Î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 becnhmark –
urile obținute pentru anumite implementări folosind procesorul grafic.

8
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 modelul 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 -core2, 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 memo ry)
– 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țiu ni ce se execută în interiorul unei structuri repetitive
de tip for doar că diferența e ste că acest set în loc să fie accesat pe rând pentru fiecare
iterație, se execută simultan pentru toate iterațiile for -ului. Conform taxonomiei lui
Flynn3, 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 alocat pe device)
2. Host -ul inițializează transferul de date de input de la host la device

1 https://en.wikipedia.org/wiki/GeForce_256
2 https://en.wikipedia.org/wiki/Multi -core_processor
3 https://ro.wikipedia.org/wiki/Taxonomia_lui_Flynn

9
3. Se rulează kernel -ul pen tru 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 GPU față de restul aplicației.
Maximizarea performanței în programarea procesorului grafic este exprimată
prin raport ul:

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

Utilizarea procesorului grafic se justifică cu cât avem o complexitate
matematică mai mare pe 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
diferă î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ă Nvidia4, prima versiune
fiind lansată în 2007. Scopul acestei platforme de programare este de a abstractiza
programarea proces orului grafic într -un device c are poate fi programat prin
intermediul unui API. Astfel, programatorul nu trebuie să cunoască detalii
arhitecturale de nivel jos legate de procesorul 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ă programato rului un nivel abstract de interacțiune cu procesorul
grafic printr -un API scris pentru limbajul de programare C5 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 calcu l a procesorului grafic. Astfel, se poate utiliza
procesorul grafic la o gamă largă de aplicații și nu doar din domeniul graficii.

4 https://en.wikipedia.org/wiki/Nvidia
5 C este un limbaj de programare procedural dezv oltat de Dennis Ritchie la Bell Labs și lansat în 1972

10

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 Capability .
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
confundată cu versiunea de CUDA (de ex. CUDA 7.5, CUDA 8.0 etc) care reprez intă
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 ce au versiunea de Compute Capability sub 2.0 pot
avea maxim 512 fire de execuție per 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 nivel jos, Driver API, și unul de nivel
superior, 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 putea utiliza puterea de calcu l oferită de procesorul grafic. Pe lângă
partea de ke rnel, 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 este î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, fiind un API ne nivel jos poate fi integrat cu alte limbaje de

11
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 construit ca un
API de nivel înalt 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 al te limbaje de
programare . În capitolul Runtim e API vs Driver API la nivel de host vor fi detaliate
diferențele dinte cele două API -uri.

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 Windows6 este necesar un procesor grafic compatibil7,
Microsoft Visual C++ și Nvidia CUDA Toolkit8.

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 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.

6 https://en.wikipedia.org/wiki/Microsoft_Windows
7 https://developer.nvidia.com/ cuda -gpus
8 https://developer.nvidia.com/cuda -toolkit

12
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:

#include <stdio.h>
__global__ void mykernel() {
printf("Hello world");
}
int main() {
mykernel<<<1,1>>>();
return 0;
}

În figura 1.3 avem compila rea fișierul ui la linia de comandă și executarea sa :

Figura 1.3.

Î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 compilat în fișier de t ip .ptx care este folosit
mai departe la compilare de tip JIT9 pentru a obți ne cod ul care este executat pe GPU .

9 JIT – Just In Time

13
Fișierele de tip .ptx conțin cod sursă PTX (Parallel Thread Execution) care
reprezintă un limbaj de nivel jos.

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 fir e 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

14
– 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, elementul de pe ultima
poziție va fi procesat de către 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
bidimensional 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 10

1.5.2. Kernel
Kernel -ul reprezintă o procedură care se execută pe device și este apelată de
către host. K ernel -ul se declară prin cuvântul cheie __global__ la începutul

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

15
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. Foarte
important de specificat că un kernel poate fi apelat doar de pe host.

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

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

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

}

Tabelul 1.1.

Î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 este 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 impor tant
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,
pentr u a evita potențiale erori continuarea firului de execuție este limitată de această
condiție.
În continuare vom vedea 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]);
}

16

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

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 .2.

Codul din tabelul 1.2. 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 vr em 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]

17
1.6.1. Registre
Registrele 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, această 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 asemă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 a u 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ă unei 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ță și viteză de scriere și citire.

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 vi ață depinde de
durata de viață a blocului și imediat ce acesta își termină execuția, memoria este
golită. Din punct de vedere al latenței și vitezei de scriere și citire este mult mai
rapidă decât alte tipuri de memorie dar nu la fel ca registrul. Acest ti p de memorie
este foarte important deoarece este folosit la procedee de sincronizare a firelor de
execuți e. În unele cazuri este de preferat ca datele din memoria globală să fie
transferate și grupate în memoria partajată pentru ca apoi operațiile de proce sare 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 memoria partajată precum în figura 1.9.

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 la când aplicația își încheie execuția.

18
Dezavantajul acestui tip de memorie este însă latența foarte mare. La nivel de API,
variabilele marcate cu __device__ la declarare sunt stocate în memoria globală
precum în figura 1.10.

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 singură
diferență: este read -only. Pentru procesoarele 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 independent unul față de
celălalt.

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

19
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 u n 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 reprezintă mereu o problemă de
sincronizare .

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.

20

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 suma
unor numere implementată din perspectiva programă rii paralele.

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 alte 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 parametrii: 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 ;
– 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 ;

21
– 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 es te
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
librării pentru scrierea de kernel cât și librării care să permită utilizarea device -ului
de către host pentru a facilita integrarea CUDA în aplicații C/C++. Cu toate acestea,
librăriile destinate pentru a fi folosite pe partea de host s -au dovedit a fi de tip low –
level, motiv pentru care s -a dovedit mai greu de învățat. Ca răspuns, Nvidia a lansat
mai târz iu CUDA Runtime API pentru aplicații C/C++. Acest nou API păstrează
toate librăriile prezente în Driver API dar oferă noi librării intenționate a fi folosite
la nivel de host, librării 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
host și mai exact felul în care host -ul apelează ker nel-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ții resursele
procesorului grafic în 3 contexte diferite într -o singură aplicație și fieca re 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ă. Mai 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.

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

În acest 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, puternic tipizat, conceput
de către James Gosling la Sun Microsystems (acum filială Oracle) la începutul
anilor ʼ90, fiind lansat în 1995. Cele mai multe aplicații distribuite sunt scrise în
Java, iar noile evoluții tehnologice permit utilizarea sa și pe dispozitive mobile gen
telefon, agenda electronică , palmtop etc. În felul acesta se creează o platformă
unică, la nivelul programatorului, deasupra unui mediu eterogen extrem de
diversificat. Limbajul împrumută o mare parte din sintaxă de la C și C++, dar are
un model al obiectelor mai simplu și prezintă mai puține facilități de nivel jos. Un
program Java compilat, corect scris, poate fi rulat fără modificări pe orice platformă
care e instalată o mașină virtuală Java (engleză Java Virtual Machine, prescurtat
JVM). Acest nivel de portabilitate (inexistent p entru limbaje mai vechi cum ar fi C)
este posibil deoarece sursele Java sunt compilate într -un format standard numit cod
de octeți (engleză byte -code) care este intermediar între codul mașină (dependent de
tipul calculatorului) și codul sursă.” [4]

2.2. Maven

Maven11 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 p rogramare . Pentru platforma Java putem implementa soluții de tip binding

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

23
utilizând JNI12 (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 programa re 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,
JCuda este disponibil ca repository de Maven și astfel integrarea este mult mai facilă
și nu mai este nevoie de acele 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 Jav a, intervine un pas în plus și anume kernel -ul rulat de către JCuda
trebuie să fie 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 .p tx înainte de a fi utilizate de către JCuda. Să luăm de exemplu kernel -ul
prezent în figura 1.6 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 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
)
{

12 https://en.wikipedia.org/wiki/Java_Native_Interface

24
.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

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 invers. Î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. cuMemcpyHtoD (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. cuCtxSynch ronize();
JCudaDriver. cuMemcpyDtoH (Pointer. to(hostArray), deviceArray,
hostArray. length * Float. SIZE);
JCudaDriver. cuMemFree (deviceArray);

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

25
2.4. Apache POI

Apache POI [5] 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
genera 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 utile
pentru a înțelege felul de abordare a programării utilizând procesorul grafic.

26
3. Implementare ș i benchmark -uri

În acest capitol voi prezenta algoritmii pe care i -am implementat folosind
JCuda pentru a obse rva diferențele de performanță între CPU și GPU. Am ales să
încep cu algoritmi simpli și după primele benchmark -uri obținute am decis să trec la
implementarea unor algoritmi de tip Monte Carlo folosiți în domeniul bursei. 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. Structura 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.

Figur a 3.1. Structura proiectului

27
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și erele de tip .cu și .ptx
folosite î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ă.

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 v oid 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 getCUmodule ș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.

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

Primul algoritm pe care l -am folosit 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 punct 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ă valoarea 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.

29
Inițial se calculează pentru fiecare fir de execuție index -ul său în cadrul grid –
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ă, adi că 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 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 memoria 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, e lement 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 10 24 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 n umere pozitive generate aleatoriu.

30

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

Î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 trag e astfe,
concluzia că GPU -ul este mai rapid în acest caz dacă ne referim strict la 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 datelor
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.
Un alt algoritm simplu din punct de vedere al programării secvențiale
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 p recedent 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:

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;

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

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

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
element ul se găsește pe la finalul tabloului.

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

32
3.3. Scenarii Monte Carlo

Scenariile Monte Carlo reprezintă o categorie de algoritmi care utilizează
multe date generate aleatoriu pentru a obține, cu o anumită rată de eroare, o estimare
a rezultatului unei probleme. Numele „Monte Carlo” provine de la subdiviziunea
Monte Carlo din orașul -stat Monaco fiind o zonă celebră pentru caz inourile sale.
Acest tip de algoritmi este util pentru rezolvarea de probleme care implică un
grad foarte ridicat de incertitudine. Mai exact nu se obține un rezultat clar dar în
schimb se oferă rezultatele a câtor mai multe scenarii posibile, scenarii care depind
de variabile aleatoare , care pot fi folosite mai departe pentru a realiza o analiză a
riscului și eventual se poate estima un rezultat final.
Algoritmii de tip Monte Carlo au în general următorii patru pași:
1. Se stabilește domeniul problemei
2. Se generează date de input aleatoriu pe baz a unei distribuții de probabilitate
3. Pentru fiecare set de date de input generat la pasul 2, se realizează un calcul
determinist pentru a se obține un rezultat
4. Rezultatele finale sunt agregate

Scenariile Monte Carlo se pliază foarte bine pe o implementare bazată pe
procesorul grafic . În primul rând, estimarea rezultatului depinde de un număr cât mai
mare de scenarii și prin urmare rularea scenariilor în paralel poate reduce
semnificativ timpul de execuție . În al doilea rând în funcție de felul în care se
realizează agregarea, se poate utiliza o abordare din perspectiva programării paralele.
În al treilea rând, CUDA oferă metode foarte eficiente pentru generare de numere în
mod aleatoriu iar aceste numere sunt salvate direct în memoria GPU -ului și astfel
timpul de transfer al datelor este redus considerabil deoarece transferul de la host la
device va conține doar câteva date de input comune pentru fiecare scenariu iar restul
datelor de input se generează direct în memoria globală. Latența pentru scrierea
datel or în memoria globală a GPU -ului este mai mică când este accesată de către
GPU decât când este accesată de către CPU.
Prin urmare, pentru a implementa scenarii Monte Carlo putem utiliza
procesorul grafic la pașii 2, 3 și 4.

3.3.1. Opțiuni de bursă de tip barieră
Pentru a exemplifica implementarea scenariilor Monte Carlo, am ales un
exemplu din domeniul financiar . O opțiune reprezintă un contract încheiat între un
cumpărător și un vânzător care dă dreptul cumpărătorului de a cumpăra sau vinde un
activ anume dacă sunt îndeplinite anumite condiții. Activul poate fi de exemplu o
proprieta te sau o acțiune la bursă.

33
În cazul opțiunilor de tip barieră, valoarea opțiunii la scadență, numită payoff ,
depinde de două valori: valoarea de strike și valoarea ba rierei. În cazul valorii de
strike , payoff -ul reprezintă diferența dintre valoarea finală a opțiunii și valoarea de
strike doar dacă valoarea finală a opțiunii este mai mare decât valoarea de strike . În
caz contrar, valoarea de payoff este 0 .
Bariera reprezintă o limită de care dacă se trece opțiunea devine inactivă iar
valoarea de payoff va fi 0 sau poate fi o limită care trebuie trecută pentru ca opțiunea
să devină activă, altfel valoarea de payoff va fi 0 . Opțiunile de tip barieră pot fi
împărțite în 4 categorii [6]:
 „Down and In” atunci când valoarea opțiunii trebuie să treacă în partea de
jos a limitei stabilite de barieră pentru a deveni activă
 „Up and In” atunci când valoarea opțiunii trebuie să treacă în p artea de sus
a limitei stabilite de barieră pentru a deveni activă
 „Down and Out” atunci când valoarea opțiunii devine inactivă dacă trece
în partea de jos a limitei stabilite de barieră
 „Up and Out” atunci când valoarea opțiunii devine inactivă dacă trece în
partea de sus a limitei stabilite de barieră
În afară de valoarea inițială, valoarea de strike și valoarea barierei, alte
variabile legate de opțiunile de tip barieră sunt:
– Timpul până la at ingerea nivelului de maturitate, adică atunci când
opțiunea ajunge la scadență
– Volatilitatea anuală sau gradul de variație a unei opțiuni
– Randamentul anual al opțiunii sau nivelul câștigului anual adus de o
investiție exprimat sub formă de procent
– Rata do bânzii fără risc
Folosind metoda Monte Carlo putem obține estimări pentru opțiunile de tip
barieră. Practic, pentru o astfel de opțiune se generează un număr mare de simulări,
unde o simulare constă în evoluția valorii opțiunii pe baza unor numere care
reprezintă cu c ât se modifică valoarea opțiunii pentru o anumită zi , numere care sunt
generate aleatoriu.

Figura 3.5. Simulări Mon te Carlo multiple pentru opțiune de tip barieră „Down and out ”

34

După cum se observă și în figura 3.5 . fiecare simulare are același punct de
pornire dar datorită datelor generate aleatoriu evoluează în mod diferit. Este evident
că rata de eroare este una foarte mare dar cu cât obținem rezultatele a cât mai multe
simulări putem să stabilim o zonă în care s -ar încadra rezultatul.
Pentru exemplu meu am ales să fac o medie aritmetică pentru toate rezultatele
obținute, această medie reprezentând rezultatul estimat.
Prețul curent al unei opțiuni pentru o zi anume se calculează cu formula [7]:
𝑃𝑛= 𝑃𝑛−1+ 𝜇 ∗𝑃𝑛−1∗∆𝑡+ 𝜎∗ 𝑃𝑛−1∗𝑅𝑎𝑛𝑑

– 𝑃𝑛 reprezintă prețul care trebuie calculat și 𝑃𝑛−1 reprezintă ultimul preț
calculat pentru o zi anterioară.
– 𝜇 reprezintă randamentul anual
– ∆𝑡 reprezintă diferența de timp între două prețuri calculate
– 𝜎 reprezintă volatilitatea anuală
– 𝑅𝑎𝑛𝑑 este un număr generat aleatoriu conform unei distribuții de tip Gauss
cu media 0 și varianța ∆𝑡

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ă AbstractCudaExecu tor și
suprascriu metoda execute() , metodă în care se află toată logica pentru algoritmul de
tip Monte Carlo care utilizează procesorul grafic pentru generarea scenariilor.
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 GPU. 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.

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

public MonteCarloDownAndOutBarrier(BarrierOptionInputData inputDa ta, int
nrOfSimulations) {
this.inputData = inputData;

35
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 fiecare scenariu care se
generează, variabila de tip int monteCarloTotalSimulations care reprezintă numărul
total de simulări care se vor genera și rest ul atributelor sunt folosite pentru a stoca
informații legate de execuția CPU -ului și a GPU -ului. Fiecare clasă are un singur
constructor care preia ca parametrii un obiect de tip BarrierOptionInputData și
numărul de simulări.
Clasa BarrierOptionInputData conține atributele asociate unei bariere, cât și
metodele se tip setter și getter . Aceste atribute sunt prețul inițial al opțiunii, valoarea
de strike , valoarea barierei, timpul de atingere a maturității (data scadenței) exprimat
în ani, volatilitatea, randamentul anual și rata dobânzii . Toate aceste atribute sunt de
tip float.

final int numberOfDaysUntilMaturity = ( int) Math.ceil(365 *
inputData .getMaturityTimeInYears());

final int DEVICE_MEMORY_SIZE = monteCarloTotalSimulations *
numberOfDaysUntilMaturity;

float timeDifference = inputData .getMaturityTimeInYears()/numberOfDaysUntilMaturity;
float standardDeviation = ( float) Math.sqrt(timeDifference);

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

// Enable exceptions and omit all subsequent error checks
JCuda.setExceptionsEnabled (true);
JCurand. setExceptionsEn abled(true);

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

//allocate device memory
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() care obține numerele generate aleatoriu în
memoria globală a device -ului

În cele care urmează voi descrie implementarea metodei execute() pentru clasa
MonteCarloDownAndOutBarrier , implementările pentru celelalte 3 clase fiind

36
foarte 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 fișierului de tip .ptx care conține codul sursă
a kernel -ului în cod PTX și respectiv numele funcției de kernel. Apelul metodei
setExceptionsEnab led 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 pe ntru 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 mem orie este alocat prin metoda statică cuMemAlloc având o mărime
exprimată în bytes. Constanta Sizeof.FLOAT este egală cu 4.
Variabila deviceInputData este un obiect de tip Pointer care reprezintă adresa
pentru locația din memoria globală unde sunt stocate datele generate aleatoriu de
metoda getCUDAGeneratedInput .

private Pointer getCUDAGeneratedInput( 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,
standardDeviation);

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 devic e-ului și de
a returna un obiect de tip Pointer care conține adresa de memorie către datele
generate. Clasa JCurand reprezintă o implementare de tip binding pentru librăria
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ă

37
curandGenerateNormal pentru a genera datele conform unei distribuți i 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 mereu 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 de mai jos . Inițial se stabilesc parametrii kernel -ului prin
variabila kernelParameters . Parametrii kernel -ului sunt variabilele devicePrices și
deviceInputData , atributele din obiectul inputData folosite pentru a calcula fiecare
simulare, numărul de zile până la scadență, varianța și numărul total de simulări
Monte Carlo.
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 simulări care
trebuie generate. Metoda st atică 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 eliberează memoria
de pe device care a fost utilizat ă pentru rularea kernel -ului.

gpuTimeStart = System. nanoTime ();

Pointer kernelParameters = Pointer. to(
Pointer. to(devicePrices),
Pointer. to(new float []{inputData .getMaturityTimeInYears()}),
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 .getAnnua lDrift()}),
Pointer. to(new float []{inputData .getInterestRate()}),
Pointer. to(new float []{timeDifference}),
Pointer. to(deviceInputData),
Pointer. to(new int[]{numberOfDaysUntilMaturity}),
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);
JCudaDriver. cuCtxSynchronize ();

float[] simulationPrices;
ReduceCudaExecutor reduceCudaExecutor =
new ReduceCudaExecutor( monteCarloTotalSimulations ,
devicePrices);
reduceCudaExecutor.execute();
simulationPrices = reduceCudaExecutor.getResult();
double averagePrice = 0.0;
for (float price : simulationPrices) {

38
averagePrice += price;
}
averagePrice = averagePrice/ monteCarloTotalSimulations ;
gpuTimeEnd = System. nanoTime();

// GPU Cleanup
JCudaDriver. cuMemFree (devicePrices);
JCuda.cudaFree (deviceInputData);

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, deși diferențele dintre ele sunt foarte mici:
– kernel_montecarlo_downandin
– kernel_montecarlo_downandout
– kernel_montecarlo_upandin
– kernel_montecarlo_upandout

extern "C"

__global__ void mc_kernel(
float * d_prices,
float years_to_maturity_time,
float strike,
float barrier_value,
float initial_price,
float volatility,
float annual_drift,
float interest_rate,
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 + volatility *
s_curr * 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();
d_prices[simulation_idx] = exp( -interest_rate *
years_to_maturity_time) * payoff;

39
}
}

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 în așa fel să folosească câte un fir de execuție
pentru fiecare simulare. 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ă index ul simulării este mai mic
decât numărul total de simulă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 maturitate 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ă termina rea structurii de tip do -while, se calculează valoarea
finală doar în cazul în care prețul curent din ultima zi este mai mare decât valoarea
de strike . Altfel rezultatul simulării 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
Până acum am văzut cum procesorul grafic a fost utilizat la pașii 2 și 3 ai unui
algoritm de tip Monte Carlo. Am folosit procesorul grafic pentru a genera datele
aleatorii și pentru a calcula rezultatele scenariilor. 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 acest e numere trebuie transferate în memoria procesorului grafic, în cazul
nostru numerele sunt deja în memoria globală deci prin urmare putem obține un spor
de performanță printr -o astfel de implementare.

public class ReduceCudaExecutor extends AbstractCudaEx ecutor {

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

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

40

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);
CUfunction reduceFunction = getCUfunction(reduceModule,
FUNCTION_NAME_REDUCE);

CUdeviceptr devicePricesAfterReduce = new CUdeviceptr();
JCudaDriver. cuMemAlloc (devicePricesAfterReduce, Sizeof. FLOAT * GRID_SIZE);
Pointer kern elParameters = 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 ca re 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 calcula t î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[];

41
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();
}

//only thread 0 writes the result for this block back to global memory
if (threadId == 0) {
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 realiza benchmark -ul pentru scenarii Monte Carlo am ales să măsor
timpul de execuție începând cu 10000 de scenarii pentru estimarea rezultatului unei
opțiuni de tip barieră și m-am oprit la un milion de scenarii. După un milion de
scenarii, încep să apară probleme legate de memorie insuficientă pe procesorul
grafic.
Benchmark -ul din figura 3.6 . ne confirmă superioritatea procesorului grafic
pentr u o astfel de problemă. Dacă timpul de execuție al CPU -ului crește constant cu
cât crește numărul de simulări, timpul de execuție al procesorului grafic aproape că
nu se modifică . Deși cele două secunde de diferență pentru un milion de scenarii nu

42
înseamnă așa mult, această diferență este totuși importantă atunci când avem de
calculat opțiuni multiple. Pentru a calcula opțiuni multiple, kernel -ul pentru rulat
scenariile Monte Carlo se execută secvențial pentru fiecare opțiune dar totuși d upă
cum se observă în figura 3.7 . diferența de timp între GPU și CPU este enormă.

Figura 3.6 . Benchmark pentru estimarea unei opțiuni utilizând scenarii Monte Carlo

Figura 3.7 . Benchmark pentru estimarea unui număr de opțiuni de tip barieră ; pentru fiecare
opțiune se utilizează un milion de scenarii Monte Carlo

43
3.4. Generarea rapoartelor de performanță

Pentru a putea genera bench mark -urile am utilizat fișiere E xcel de tip .xlsx î n
care pe prima coloană se găsește dimensiunea input -ului iar pe următoarele col oane
timpii de execuție pentru GPU, respectiv C PU. Aceste fișiere sunt generate de către
clasa SpreadsheetGenerator , clasă implementată folosind librăria Apache POI
pentru a genera fișierele.

public class SpreadsheetGenerator {

private String location ;
private List<Integer> inputSizeList ;
private List<Float> gpuResults ;
private List<Float> gpuNoDataTransferResults ;
private List<Float> cpuResults ;

public SpreadsheetGenerator(String loca tion, List<Integer> inputSi zeList,
List<Float> gpuResults, List<Float> cpuResults) {
this.location = location;
this.inputSizeList = inputSizeList;
this.gpuResults = gpuResults;
this.cpuResults = cpuResults;
}

public SpreadsheetGenerator(String location, List<Integer> inputSizeList ,
List<Float> gpuResults, List<Float> cpuResults, List<Float>
gpuNoDataTransferResults) {
this.location = location;
this.inputSizeList = inputSizeList;
this.gpuResults = gpuResults;
this.cpuResults = cpuResults;
this.gpuNoDataTransferResults = gpuNoDataTransferResults;
}

public void generateSpreadsheet() {

Workbook wb = new XSSFWorkbook();
Sheet sheet = wb.createSheet( "results");

Row firstRow = sheet.createRow( 0);
firstRow.createCell( 0).setCellValue( "Elemente procesate" );
firstRow.createCell( 1).setCellValue( "Timp execuție GPU + timp pentru
transfer de date" );
firstRow.createCell( 2).setCellValue( "Timp execuție GPU " );
firstRow.createCell( 3).setCellValue( "Timp execuție CPU" );

for (int i = 0; i < inputSizeList .size(); i++) {
Row row = sheet.createRow(i + 1);

row.createCell( 0).setCellValue( inputSizeList .get(i));
Cell gpuResultCell = row.createCell( 1);
gpuResultCell.setCellValue( gpuResults .get(i));
Cell gpuNoDataTransferResultCell = row.createCell( 2);

gpuNoDataTransferResultCell.setCellValue( gpuNoDataTransferResults .get(i));
Cell cpuResultCell = row.createCell( 3);
cpuResultCell.setCellValue( cpuResults .get(i));

}

try (OutputStream fileOut = new FileOutputStream( this.location )) {

44
wb.write(fileOut);
} catch (IOException e) {
e.printStackTrace();
}

}

public void generateSpreadsheetMonteCarlo() {

Workbook wb = new XSSFWorkbook();
Sheet sheet = wb.createSheet( "results" );

Row firstRow = sheet .createRow( 0);
firstRow.createCell( 0).setCellValue( "Număr scenarii" );
firstRow.createCell( 1).setCellValue( "Timp execuție GPU " );
firstRow.createCell( 2).setCellValue( "Timp execuție CPU" );

for (int i = 0; i < inputSizeList .size(); i++) {
Row row = sheet.createRow(i + 1);

row.createCell( 0).setCellValue( inputSizeList .get(i));
Cell gpuResultCell = row.createCell( 1);
gpuResultCell.setCellValue( gpuResults .get(i));
Cell cpuResultCell = row.createCell( 2);
cpuResultCell.setCellValue( cpuResults .get(i));

}

try (OutputStream fileOut = new FileOutputStream( this.location )) {
wb.write(fileOut);
} catch (IOException e) {
e.printStackTrace();
}
}
}
Tabelul 3.11. Clasa SpreadsheetGenerator

45
Concluzie

La finalul acestei lucrări consider că am reușit să ofer o introducere
cuprinzătoare a acestui domeniu exotic al programării 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 cât timp o problemă poate fi divizată în
probleme de dimensiune mai mic ă care pot fi rezolvat simultan în paralel. Cu toate
acestea însă memoria procesorului grafic are un rol important per total asupra
performanței. În primul rând complexitatea problemei executate de fiecare fir de
execuție în parte trebuie să justifice tran sferul de date către memoria procesorului
grafic deoarece de multe ori acest transfer face o implementare care utilizează
procesorul grafic să fie nerentabilă. În al doilea rând, programatorului îi revine
sarcina de a știi cum să jongleze cu diferitele tip uri 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 ar putea să beneficieze enorm de paralelismul
oferit d e procesoarele grafice pentru a procesa cantități masive de date.

46
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] „Java” https://ro.wikipedia.org/wiki/Java_(limbaj_de_programare)
[5] „Apache POI” https://poi.apache.org/
[6] E. Derman și I. Kani, „The Ins and Outs of Barrier Option Part 1” Preluată de pe:
http://www.emanuelderman.com/media/insoutbarriers1.pdf
[7] D. Emanuel și I. Kani, „The Ins and Outs of Barrier Options Part 2” Preluată de pe:
http://www.emanuelderman.com/media/insoutbarriers2.pdf
[8] „JCuda” [Interactiv]. Available: http://www.jcuda.org/