Calculatoare și tehnologia informației [603872]
UNIVERSITATEA TEHNICĂ „Gheorghe Asachi” din IAȘI
FACULTATEA DE AUTOMATICĂ ȘI CALCULATOARE
DOMENIUL: Calculatoare și tehnologia informației
SPECIALIZAREA:Tehnologia informației
Benchmark GPU
Student: [anonimizat], 2017
Cuprins
Introducere…………………………………………………………………………………………………………………….. 1
Capitolul 1. Fundamente teoretice. Arhitectura GPU …………………………………………………………… 3
1.1. Cum fucționează un GPU ……………………………………………………………………………………… 3
1.2. Arhitectura CUDA ……………………………………………………………………………………………….. 4
1.3. General purpose computing on GPU (GPGPU) ……………………………………………………….. 5
Capitolul 2. Descrierea aplicației ………………………………………………………………………………………. 7
2.1. Mod de utilizare…………………………………………………………………………………………………… 7
2.2. Analiza aplicațiilor existente pe piață ……………………………………………………………………… 8
Capitolul 3. Implementarea aplicației ………………………………………………………………………………. 10
3.1. Proiectarea software a aplicației ……………………………………………………………………………. 10
3.2. Sincronizarea view-urilor și creearea interfeței în Qt ………………………………………………. 13
3.3. Elemente de prelucrare ale imaginilor …………………………………………………………………… 15
3.3.1. Scena 1………………………………………………………………………………………………………. 15
3.3.2. Scena 2………………………………………………………………………………………………………. 20
3.4. Lățimea de bandă. Teoretic și Practic …………………………………………………………………….. 24
3.5. Gflop………………………………………………………………………………………………………………… 25
3.6. Interfata Web……………………………………………………………………………………………………… 25
Capitolul 4. Evaluare …………………………………………………………………………………………………….. 26
4.1. Testarea aplicației ……………………………………………………………………………………………….. 26
4.2. Rezultate……………………………………………………………………………………………………………. 26
4.3. Probleme întâmpinate și soluții de rezolvare ………………………………………………………….. 28
4.3.1. Creearea backend-ului aplicației ……………………………………………………………………. 28
4.3.2. Calcularea corectă a lățimii de bandă ……………………………………………………………… 29
Capitolul 5. Concluzii ……………………………………………………………………………………………………. 30
5.1. Gradul în care s-a realizat tema propusă ………………………………………………………………… 30
5.2. Dezvoltare ulterioară …………………………………………………………………………………………… 30
Bibliografie………………………………………………………………………………………………………………….. 31
Anexe………………………………………………………………………………………………………………………….. 32
Benchmark GPU
Adrian Aioanei
Rezumat
Programarea pe unitățile de procesare grafică (GPGPU) a devenit un subiect fierbinte
pentru oamenii de știință datorită cantităților mari de date de procesat. În special în cadrul
comunității High Performance Computing (HPC), aceste procesoare grafice masiv paralele au
primit multă atenție și cu un motiv întemeiat. Benchmarking-ul problemelor obișnuite rezolvate
pe cele mai noi unități de procesare grafică de la NVIDIA și alți furnizori arată o accelerare de
până la 70 de ori mai mare decât cea a unei implementări pe un procesor obișnuit.
GPU-ul oferă rezultate rapide și are o putere mare de procesare însă aici intervin și unele
limitări, iar aceste limitări constau în procesul te transfer al datelor către GPU și dinspre el. Aici
intervin două metrici importante, latența și lățimea de bandă. Chiar și pentru transferurile de mici
dimensiuni latența poate avea impact negativ asupra perofrmanțelor, lățimea de bandă
influențând performnțele pentru transferurile de dimensiuni mai mari.
Lucrarea de față vine spre întâmpinarea utilizatorilor care doresc să verifice capabilitățile
de care dispun dispozitivele lor GPU. În acest fel proiectul urmărește testarea unor puncte
importante ce influențează în mod direct performanțele dispozitivului precum capacitățile de
randare, lățimea de bandă, numărul de operații realizate pe secundă, toate acestea printr-o
interfață intuitivă, ușor de utilizat și susținută de o platformă online ce are drept scop
centralizarea tuturor rezultatelor opținute.
Introducere
Introducere
Aplicația își propune testarea și interogarea performanțelor plăcilor grafice dezvoltate de
compania Nvidia pentru a sublinia performanțele superioare pe care le oferă plăcile video acuale
în defavoarea CPU-urilor. Dezvoltarea plăcilor video a fost întotdeauna îndreptată spre
procesarea cantităților mari de date cât mai rapid și mai eficient posibil. Acest fapt a dus la o
creștere exponențiala în ceea ce privește puterea de procesare a GPU-urilor spre deosebire de
CPU-uri unde se pare că a ajuns într-un impas datorită constrângerilor fizice. Acest lucru
împreună cu posibilitatea rulării propriului cod pe GPU a dus la apariția conceptului de GPGPU
(General-Purpose computing on GPU), concept care a făcut posibilă utilizarea performanțelor
plăcilor video în domenii diverse.
Companiile care dezvoltă plăci grafice precum AMD/ATI și NVIDIA au dezvoltat
procesoare grafice capabile de procesare masivă în paralel care dispun de arhitecturi hardware
special dezvoltate în ideea procesării paralele. Utilizarea procesoarelor grafice pentru a executa
algoritmi în paralel diminuează considerabil timpul de execuție în comparație cu implementarea
clasica, secvențiala pe CPU.
Tehnologia este într-o continuă dezvoltare iar pe măsură ce știința avansează apar
aplicații tot mai complexe care tind să ofere un răspuns utilizatorului în timp real, însă de cele
mai multe ori nu este atât de simplu. Oamenii de știință s-au confruntat continuu cu nevoia de a
procesa cantități imense de date într-un timp cât mai scurt posibil, iar pentru rezolvarea acestei
probleme plăcile video vin cu o rezolvare.
În momentul de față programarea paralelă pe GPU este întâlnită în domenii diverse
începând de la cercetare, simulări fizice, procesare de semnale, rețele neuronale pana la sisteme
predicție unde viteza de procesare în cazul unor aplicații a fost crescută de 100 de ori față de
cazul general unde era folosit un procesor obișnuit.
Datorită dezvoltării spectaculoase și a promisiunilor fondate programarea și paralelizarea
aplicațiilor pe GPU câștigă teren continuu fiind atrase de performanțe sporite, o comunitate
puternică și produse performante, însă momentan cel mai frecvent plăcile video sunt folosite în
cercetare acolo unde volumul de date este impresionant și necesită capabilități hardware de care
un procesor (CPU) nu dispune.
Aplicația de față este un benchmark, termen larg utilizat în multiple domenii precum
informatica, marketing-ul, financiar și în principal este un proces de măsurare a performanțelor și
a celor mai bune practici specifice fiecărui domeniu în parte. Astfel termenul de benchmark
poate fi utilizat diferit în funcție de domeniul asupra căruia face referire. De exemplu în
domeniul economic termenul de benchmark reprezintă un proces continuu de căutare a celor mai
bune practici care să conducă organizația către un nivel superior de performanță. Ideea de bază a
acestui concept în informatică, este aceea de a analiza performanțele unui device anume cu
privire la capabilitățile pe care le posedă.
NVIDA în anul 2000 a lansat prima placă ce putea fi programată GeForce 3. Acum
programatorul are posibilitatea ca fiecare pixel să fie prelucrat de un „program” care să-i adauge
textură și chiar fiecare vertex poate fi procesat la rândul său de un program(shader). În momentul
de față există două alternative care merită atenția pentru programarea plăcilor video. Una este
CUDA, iar cealaltă este OpenCL care oferă și portabilitate pentru dispozitivele creeate de
AMD/ATI, cealaltă alternativă CUDA fiind disponibilă doar plăcile NVIDIA.
Sunt utilizate o varietate de benchmark-uri pentru a testa diferite aspecte ale plăcilor
video cât mai corect posibil. Dacă placa video este una de consum este o idee bună să fie testată
cu scene intensive 3D pentru jocuri. Aceste benchmark-uri sunt greu de reprodus pe dispozitive
1
Adrian Aioanei
diferite însă oferă o idee asupra capabilităților. De asemenea aici s-a descoperit și o problemă în
ceea ce privește frame-ul care este randat și frame-ul care este prezentat utilizatorului, acestă
problemă fiind rezolvată cu ultimele update-uri de la Nvidia.
Un alt tip de benchmark este numit sintetic benchmark. Are multe în comun cu modul în
care se face testarea scenelor 3D pentru jocuri însă în plus oferă o reprezentare foarte exactă a
testelor în fiecare ciclu. Acest lucru permite compararea diferitelor unități de procesare grafică cu
o mare precizie și permite evaluarea capabilităților unde procesorul grafic excelează.
Al treilea tip de benchmark face referire la puterea de calcul. Aceste benchmark-uri arată
adevărata putere de prelucrare a plăcilor video și sunt utilizate în aplicații științifice. Un astfel de
exemplu ar fi interfețele GEMM și FFT. GEMM măsoară performanța plăcii video în
multiplicarea matricilor de dimensiuni marei pe când FFT este Fast Fourier Transformations.
Ambele sunt task-uri paralelizate care pot dovedi potențialul uriaș al device-ului.
O primă problemă care apare în realizarea benchmark-urilor reprezintă posibilitatea de a
reproduce exact același test,în special a scenelor intense din jocuri ce necesită randare 3D pe
device-uri diferite pentru a avea un punct de comparație comun. De cele mai multe ori acest
lucru nu este posibil deoarece în testarea unui procesor video nu intră doar capabilitățile proprii
ci și restul componentelor hardware ce compun sistemul.
2
Capitolul 1. Fundamente teoretice. Arhitectura GPU
Capitolul 1. Fundamente teoretice. Arhitectura GPU
1.1. Cum fucționează un GPU
GPU-ul a fost inventat de NVIDIA în anul 1999. Inițial unitățile de procesare grafică au
fost dispozitive cu funcție fixă, ceea ce înseamnă că au fost proiectate să proceseze anumite
stadii din pipeline precum vertex shaders sau pixel shaders, însă au evoluat în procesoare
programabile din ce în ce mai flexibile.
Un GPU sau o unitate de procesare grafica este folosită în principal pentru aplicații sau
task-uri grafice. Este un procesor capabil să creeze efecte de lumină, umbră, și transformă
obiectele din contexul curent de fiecare dată când o scenă este redesenată. Acestea sunt sarcini de
ordin matematic care ar pune destul de multă presiune asupra CPU-ului, iar ridicarea acestor
sarcini eliberează cicluri procesor care pot fi utilizate pentru alte task-uri urgente.
Procesoarele moderne cum sunt și GPU-urile folosesc mii de fire de execuție pentru a
conduce sistemul către o lățime de bandă ridicată și latență cât mai mică, iar o dată cu
maturizarea device-urilor, GPU-urile au devenit cipuri programabile manycore în jurul unui șir
de procesoare paralele. Device-ul este construit dintr-un număr de Streaming Multiprocessors
(SM) fiecare dintre ele fiind capabile să suporte mii de thread-uri concurente, mai exact până la
2048 pe arhitecturile moderne.
Managementul thread-urilor incluzând crearea, programarea, sincronizarea este realizat
în totalitate de către SM. Pentru a gestiona eficient numărul mare de thred-uri fiecare SM
folosește arhitectura SIMT (Single Instruction, Multiple Thread) .
Pentru a executa intrucțiuni pe GPU se lansează funcția care se dorește a fi executată (de
obicei aceasta se numește kernel) sub forma unui grid de blocuri de thread-uri care este trimis la
device. O dată ajuns kernel-ul la GPU acesta trimite către SM (Streaming Processor) unul sau
mai multe blocuri. Fiecare SM își crează o coada de warp-uri (un warp conține de obicei 32 de
thred-uri) care este executat pe un clock. Dacă execuția warp-ului nu s-a încheiat într-un clock
acesta este adăugat din nou în coadă. Acest lucru se repetă până se termină de executat blocul
curent. Atunci când s-a terminat execuția este solicitat un nou bloc de la GPU.
Un alt aspect important al GPU-ului este pipeline-ul. Termenul de pipeline este des
întâlnit în grafica computerizată și descrie pașii pe care trebuie să îi urmeze un sistem grafic
pentru a crea o scenă 2D sau 3D, mai concret pipeline-ul este procesul de transformare a
modelului 2D sau 3D în ceea ce se afișează ca rezultat final pe display.
Deoarece pașii necesari pentru această operație depind în mare măsură de software-ul și
hardware-ul folosit și de caracteristicile de afișare dorite, nu există un pipeline universal potrivit
pentru toate cazurile. Cu toate acestea, au fost create API-uri grafice, cum ar fi Direct3D și
OpenGL, pentru a unifica pașii asemănători și pentru a controla pipeline-ul unui accelerator
hardware dat. Practic, aceste interfețe API abordează hardware-ul de bază și îl împiedică pe
programator să scrie un cod greu pentru a manipula acceleratoarele de hardware grafice
(NVIDIA / AMD / Intel etc.).
Randarea propriu zisă începe cu una sau mai multe coordonate ce servesc ca intrare
pentru programul de procesarea a coordonatelor(vertex shader) iar în total pipeline-ul cuprinde 4
etape importante precum : Vertex Shader (VS), Geometry Shader, Rasterization și Fragment
Shader (FS). Geometry Shader folosește output-ul primit de la Vertex Shader și operează la nivel
geometric (linie, triunghi). Acest shrader este programabil și poate fi modificat astfel încât sa
adauge sau să elimine puncte dintr-o primitivă. Dacă etapele inițiale se ocupau de coordonate și
geometrie acum apare nevoia randării pixelilor iar pentru acesta trebuie rasterizate primitivele
inițiale. Acest lucru se realizează împărțind primitiva inițială în fragmente individuale care în
3
Adrian Aioanei
cele din urmă v-or fi colorate de Fragment Shader, iar ulterior trasformat în pixeli într-un buffer.
Ultima etapă este cea de Fragment Shader care preia fragmentele generate în timpul procesului
de rasterizare unde se aplică un algoritm creeat de programator pentru a crea valorile finale
pentru culoare și adâncime.
1.2. Arhitectura CUDA
CUDA (Compute Unified Device Architecture) este o arhitectură software și hardware
pentru calculul paralel al datelor dezvoltată de către compania americană NVIDIA. CUDA este
utilizată atât în serile de procesoare grafice destinate utilizatorilor obișnuiți cât și în cele
profesionale. O serie de interfețe de calcul din arhitectura CUDA sunt similare cu cele ale
principalilor competitori: OpenCL de la Khronos Group și DirectCompute de la Microsoft.
Dezvoltatorii pot accesa prin intermediul CUDA setul de instrucțiuni și memoria
elementelor de calcul paralel din procesoarele grafice. Utilizând CUDA, cele mai recente
procesoare grafice NVIDIA pot realiza calcule specifice microprocesoarelor. Totuși, spre
deosebire de acestea, arhitectura procesoarelor video este concepută pentru execuția simultană a
numeroase fire, cu o viteză scăzută și nu a unui singur fir dar foarte rapid. Această tehnică de
rezolvare a problemelor de uz general cu ajutorul procesoarelor video este cunoscută ca GPGPU.
În industria jocurilor pe calculator, pe lângă generarea graficii, procesoarele video mai
realizează și calculele pentru interacțiunea fizică dintre obiecte (fum, foc, fluide). Un exemplu în
acest sens este tehnologia PhysX. CUDA mai este utilizată și în domeniile bioinformaticii,
criptografiei precum și în alte arii ale științei și tehnologiei.
CUDA pune la dispoziție atât un API de nivel low cât și unul de nivel înalt. Primul SDK
CUDA a fost făcut public în data de 15 februarie 2007, având versiuni pentru Microsoft
Windows și Linux. Versiunea 2.0 oferea suport și pentru Mac OS X. Toate seriile de procesoare
NVIDIA (GeForce, Quadro și Tesla) începând cu G8X sunt compatibile CUDA.
CUDA oferă o serie de avantaje față de API-urile tradiționale de prelucrare a datelor cu
ajutorul procesoarelor video:
4
Ilustrație 1: GPU Pipeline
Capitolul 1. Fundamente teoretice. Arhitectura GPU
•Citiri nesecvențiale – se pot face citiri din locații de memorie arbitrare
•Memorie partajată – CUDA pune la dispoziție o regiune cu memorie partajată de mare
viteză (până la 48KB per Multi-Procesor) care poate fi împărțită între firele de execuție.
•Descărcări rapide și recitiri spre și dinspre procesorul video
•Suport complet pentru operațiile cu întregi și pe bit.
Cu toate argumentele mai sus mentionate am putea afirma superioritatea pe care o deține
CUDA însă aici există și anumite limitări:
•CUDA (având capacitatea de calcul 1.x) suportă un subset al limbajului de programare C
și o serie de extensii simple. Lipsesc pointerii la funcții și recursivitatea. Un process ce
rulează va utiliza mai multe spații de memorie disjuncte.
•Pentru formatul în virgulă mobilă cu dublă precizie există niște abateri de la standardul
IEEE 754
•Spre deosebire de OpenCL, procesoarele compatibile CUDA sunt fabricate doar de
NVIDIA (începând cu seria GeForce 8, Quadro și Tesla)
•Pentru a se obține cea mai bună performanță, thread-urile ar trebui să ruleze cel puțin în
grupuri de câte 32, atunci când numărul total de thread-uri este de ordinul miilor.
Ramificarea execuției codului programului nu afectează semnificativ performanța dacă
fiecare din cele 32 de thread-uri urmează aceeași cale de execuție.
•Uneori codul valid C/C++ poate împiedica procesul de compilare datorită tehnicilor de
optimizare pe care compilatorul le aplică dacă este obligat să utilizeze resurse limitate.
1.3. General purpose computing on GPU (GPGPU)
General-purpose computing on graphics processing units (GPGPU, de asemenea referit
ca GPGP și mai puțin ca GP²) este tehnica de utilizare a unui GPU, care de obicei manevrează
calculul doar pentru grafica pe calculator, pentru a efectua calcul în aplicații tratate de obicei de
microprocesor. Acest lucru este posibil prin adăugarea de etape de programare și aritmetică de
mare precizie la pipeline-urile de randare, ceea ce permite dezvoltatorilor software să utilizeze
procesarea în flux asupra datelor non-grafice.
Plăcile grafice DirectX 9 suportau doar tipuri de culori întregi sau paletate. Diferite
formate sunt disponibile, fiecare conținând un element roșu, un element verde și un element
albastru. Câteodată o valoare alpha suplimentară este adăugată, pentru a fi utilizată pentru
transparență. Formatele obișnuite sunt :
•8 biți per pixel
•16 biți per pixel – de obicei alocați ca cinci biți pentru roșu, șase biți pentru verde, și cinci
biți pentru albastru.
•24 biți per pixel – opt biți pentru fiecare din roșu, verde, și albastru
•32 biți per pixel – opt biți pentru fiecare din roșu, verde, albastru, și alpha
Pentru grafica de început cu funcție fixă sau cu programabilitate limitată (adică până la și
incluzând DirectX 8.1-GPU) a fost suficientă deoarece aceasta este de asemenea reprezentarea
utilizată în afișare. Această reprezentare are totuși anumite limitări. Având putere de procesare
grafică suficientă, chiar și programatorii grafici doresc să utilizeze formate mai bune, cum ar fi
formatele de date în virgulă mobilă, pentru a obține efecte cum ar fi imagistica de gamă dinamică
extinsă. Multe aplicații GPGPU necesită precizie în virgulă mobilă, care vine cu plăcile grafice
în conformitate cu specificațiile Direct9 X.
Implementările în virgulă mobilă pe GPU-urile NVidia sunt în mare majoritate conform
IEEE; totuși, acest lucru nu se întâmplă pentru toți vânzătorii. Aceasta are implicații pentru
corectitudine care este considerată importantă pentru anumite aplicații științifice. În timp ce
5
Adrian Aioanei
valorile virgulei mobile pe 64 de biți (mobilă în dublă precizie) sunt disponibile de obicei pe
microprocesoare, acestea nu sunt mereu suportate pe GPU-uri; anumite arhitecturi GPU sacrifică
conformitatea cu IEEE în timp ce altele duc lipsă de precizie dublă. Au existat eforturi de a imita
valorile virgulei mobile în dublă precizie pe GPU-uri, totuși compromisul de viteză neagă orice
beneficiu de a scăpa de calcul pe GPU de la bun început.
Majoritatea operațiilor de pe GPU funcționează în mod vectorial: o singură operație poate
fi executată cu până la patru valori simultan. De exemplu, dacă o culoare <R1, G1, B1> trebuie
să fie modulată de culoarea <R2, G2, B2>, GPU poate determina culoarea rezultată <R1*R2,
G1*G2, B1*B2> printr-o singură operație. Această funcționalitate este utilă în grafică deoarece
aproape fiecare tip de date de bază este un vector (de 2,3, sau 4 dimensiuni). Exemplele includ
noduri, culori, vectori normali, și coordonate de textură. Multe aplicații pot gestiona acest lucru
într-un mod util, și datorită performanței sporite, instrucțiunile vector (SIMD) au fost mult timp
disponibile pe microprocesoare.
GPU-urile pot procesa doar noduri și fragmente independente, dar poate procesa multe
dintre ele în paralel. Acest lucru este eficient în mod deosebit atunci când programatorul dorește
să proceseze multe noduri și fragmente în același mod. În acest sens, GPU-urile sunt procesoare
în flux – procesoare care pot opera în paralel prin rularea unui singur nucleu în același timp pe
mai multe înregistrări dintr-un flux.
Un flux este un set simplu de înregistrări care necesită calcul similar. Fluxurile furnizează
paralelism de date. Nucleele (în engleză kernel) sunt funcțiile care sunt aplicate fiecărui element
din flux. În GPU-uri, nodurile și fragmentele sunt elementele din fluxuri, shader vertex și
fragment sunt nucleele pe care sunt rulate acestea. Deoarece GPU-urile procesează elemente
independent nu există posibilitatea de a avea date partajate sau statice. Pentru fiecare element se
poate citi de la intrare, se poate executa operații pe el, și se poate afișa la ieșire. Este permis să se
aibă intrări și ieșiri multiple, dar niciodată o parte de memorie care poate fi și citită și scrisă în
același timp.
Deoarece CPU-ul și GPU-ul sunt utilizate într-o gamă largă de aplicații, s-a recunoscut că
ambele unități de procesare (PU) au caracteristicile și punctele lor forte și, prin urmare,
colaborarea CPU-GPU este inevitabilă pentru realizarea computerelor de înaltă performanță în
contexul programării GPGPU și nu numai. Acest lucru a motivat o cantitate semnificativă de
cercetări pe tehnici de calcul eterogene, împreună cu proiectarea cipurilor fuzionate CPU-GPU și
a supercomputerelor eterogene petascale.
Ca și idee principala GPU-urile au evoluat constant iar o dată cu apariția conceputului
GPGPU puntem afirma ca programarea paralelă a urcat o nouă treaptă unde aturile cele mai
importante ce o definesc sunt :
•precizia
•performanța
•utilizarea performanțelor crescute oferite de GPU-urilor utilizând un limbaj de
programare confortabil
6
Capitolul 2. Descrierea aplicației
Capitolul 2. Descrierea aplicației
2.1. Mod de utilizare
Proiectul a fost dezvoltată urmărind un scenariu simplu care să necesite un număr cât mai
mic de pași din partea utilizatorilor pentru a nu crea confuzie în utilizarea interfeței, astfel
aplicația dispune de două view-uri principale și două secundare.
Primul view al aplicației oferă utilizatorului două opțiuni simple și anume aceea de a
lansa în execuție testele sau de a închide aplicația. O dată ce butonul „Start Benchmark” a fost
apăsat utilizatorului îi este deschis un nou view unde îi este indicat stadiul în care se află testele
ce sunt rulate.
O dată ce procesul a fost completat cu succes view-ul curent este închis și este creeat un
nou view care prezintă utilizatorului rezultatele finale obținute centralizate în pagini. Astfel
prima pagina prezintă detalii și specificații generale despre dispozitiv, a doua pagina prezintă
detalii strict referitoare la memorie, ce-a de-a treia rezultatele pe care le-a obținut dispozitivul în
urma testelor, iar cel de-al patrulea prezintă termeni și condiții.
Dacă răspunsul utilizatorului la întrebarea referitoare la publicarea datelor online este
pozitiv se creează o conexiune la baza de date ce se află local unde sunt înserate în cele 4 tabele
memory_performance, core_performance, general_information și practical_results datele
obținute. Pe lângă acestea mai există o tabelă unde sunt salvate feedback-urile primite de la
utilizatori cu privire la aplicație. În plus față de interfața desktop, platforma web afișează și date
referitoare la sistemul pe care s-au realizat testele.
7
Ilustrație 2: User Interface
Ilustrație 3: Loading bar
Adrian Aioanei
2.2. Analiza aplicațiilor existente pe piață
O dată cu dezvoltarea plăcilor video și prezentarea de către dezvoltatori a configurațiilor
tot mai impresionante, a apărut și nevoie de testare practică a specificațiilor teoretice oferite de
dezvoltator în vederea obținerii unei imagini clare asupra dispozitivului.
Pentru testarea acestor performanțe au apărut soft-uri specializate care supun
dispozitivele la teste specifice precum randarea, lățimea de bandă, latența și multe altele. În
momentul de față există câteva soft-uri care par să fie mai agreate și utilizate de publicul larg.
Trei dintre acestea sunt : CUDA-Z, NovaBench, 3DMark. Aceste nu sunt singurele însă sunt cele
mai utilizate de către publicul larg deoarece au o interfață intuitivă și sunt ușor de utilizat. Mai
jos este prezentat un scurt preview al acestor aplicații.
PerformanceTest este un soft oferit de către PassMark Software ce este o firmă cu
vechime în Sydney, California fiind înființată în anul 1998. Firma oferă o gamă largă de soluții
software începând cu soluții pentru statusul și disponibilitatea rețelelor wireless din jur până la
soluții ce investighează performanțele ecranelor LCD. PerformanceTest nu se adresează doar
procesoarelor grafice, ea fiind capabilă să ruleze teste de performanța pentru CPU, disk și
memoria RAM. Aplicația dispune de o interfața modernă însă nu este foarte bine organizată
acesta fiind un prim dezavantaj al aplicației. Un alt dezavantaj îl consideră modul în care pot fi
rulate testele. Nu există nici o opțiune care să realizeze un test complet asupra unui anumit
device. Utilizatorul trebuie să își selecteze personal fiecare test pe care vrea să-l ruleze (Image
filters, Image rendering,GPU compute) iar o dată cu apariția noului rezultat, fereastra anterioară
este ștearsă neavând posibilitatea vizualizării rezultatelor cumulate. Un ultim impediment este
prețul, fiind o aplicație care necesită licență.
A două aplicație, NovaBench este o aplicație free ce este disponibila pe Windows și
MacOS. Aplicația este una desktop însă conține și o interfață web ce permite utilizatorilor ce
folosesc aplicația să-și vizualizeze rezultatele cât și nota asignată dispozitivului testat. Chiar dacă
NovaBench vine cu un element de noutate, cu aceasta interfață web, testele pe care le realizează
sunt destul de minimale deoarece într-un sigur test se încearcă un mix de operații ceea ce
conduce de cele mai multe ori la rezultate parțial eronate. Comunitatea este aproape inexistentă
iar suportul asupra aplicației este inexistent.
Ultima aplicație prezentată este CUDA-Z. Aceasta este o aplicație multiplatformă ce
poate rula pe 3 dintre cele mai importante sisteme de operare: Windows, MacOS și Linux.
Aplicația fiind disponibilă pe toate cele trei sisteme de operare uneori are un aspect
8
Ilustrație 4: Benchmark-uri existente
Capitolul 2. Descrierea aplicației
inestetic. Un alt incovenient al aplicației îl consideră gama de dispozitive pe care le poate testa.
Soft-ul este capabil să testeze device-uri ce aparțin doar companiei Nvidia deoarece la baza
aplicației stă un model de programare paralelă CUDA ce este suportat doar de arhitectura
plăcilor video dezvoltate de Nvidia. Aplicația oferă informații minimale referitoare strict la GPU
precum: lățimea de bandă, dimensiunea memoriei, versiunea de CUDA și dll și alte câteva
informații de interes general. Rezultatele afișate de aplicație pot fi salvate în orice moment într-
un format html sau txt.
9
Adrian Aioanei
Capitolul 3. Implementarea aplicației
3.1. Proiectarea software a aplicației
Aplicația a fost dezvoltată utilizând două API-uri (CUDA și OpenGL), o extensie pentru
OpenGL (GLEW), un framework pentru interfață (Qt), C++ ca și limbaj de programare și două
compilatoare (Visual C++ 14.0 și NVCC) pentru genrarea executabilului.
Backend-ul aplicației este separat de interfață astfel fiind realizată o modularizare mult
mai bună. După cum se poate obseva și în Ilustrația 5 aplicația lansează un prim view care oferă
posibilitatea utilizatorului să lanseze testele sau să renunțe la execuția aplicației. Execuția
urmărește 3 pași principali. Primul pas este cel de testare propriu zisă a device-ului și de rulare a
testelor în vederea stabilirii performanțelor, apoi este creeat un fișier xml care este populat cu
toate datele obținute în urma execuției, iar ultimul pas constă într-o conexiune la o bază de date
ce este localizată pe localhost în vederea centralizării datelor dacă utilizatorul este de acord cu
acest lucru.
Fișierul xml creeat în urma execuției este procesat și afișat utilizatorului în interfața
desktop, unde este generat și un link corespunzător rezultatelor sale pentru platforma online unde
pe lângă rezultatele proprii utilizatorii au posibilitatea să urmărească rezultatele obținute de toți
utilizatorii ce au testat aplicația. Platforma online aduce în plus pe lângă interfața desktop câteva
informații despre specificațiile hardware disponibile în sistemul utilizatorului pentru a crea o
imagine mai clară despre modul în care diferite componente hardware pot influența rezultatele
obținute.
Un prim pas pe care aplicația îl realizează atunci când este lansată în execuție este acela
de a colecta specificații și informații despre dispozitivul pe care se execută testele deoarece
arhitectura CUDA nu este disponibilă pe toate device-urile ci doar pe dispozitivele ce au fost
lansate după arhitectura Tesla aceasta fiind și ea de asemenea compatibilă. Totodată este necesară
și determinarea capacității de calcul a dispozitivului (compute capability) deoarece anumite teste
necesită o capacitate de calcul ridicată (>=3.0) în vederea obținerii unor performanțe sporite în
ceea ce privește transferul de memorie între host și device.
Aplicația dispune de două scene în OpenGL care au ca obiectiv testarea puterii de randare
a device-ului și sunt contorizate numărul de FPS-uri (Frames per secound) pe care reușește să-l
redea dispozitivul. Acesta este un test important deoarece FPS-ul este o unitate de măsură a
numărului maxim de repetiții pe care îl poate realiza o placă video în vederea creării percepției
de mișcare. Cele două scene realizează diferite operații asupra pixelilor astfel încât prima scenă
realizează un număr mai ridicat de operații, asta rezultând într-un număr mai scăzut de FPS-uri
(aproximativ ~60 pe o placă video GTX 750M), iar cea de-a doua realizează un număr mai
ridicat, ~70 FPS-uri pe aceeași placă mai sus menționată.
O dată cu lansarea CUDA 6 cei de la Nvidia au oferit posibilitatea de a crește lățimea de
bandă prin implementarea funcțiilor cudaMallocHost sau cudaHostAlloc. Acest lucru a fost
posibil deoarece în mod normal device-ul nu poate accesa în mod direct date din memoria host-
ului. În acest mod de fiecare dată când se dorea transferul dinspre host către device, driver-ul
CUDA era nevoit să aloce un vector “pinned“ apoi să copie datele din memorie în vectorul nou
alocat ulterior realizându-se un nou transfer dinspre vectorul alocat către memoria device-ului.
Aceste operații fiind consumatoare de timp s-au creeat aceste funcții care nu mai implică
alocarea și transferul datelor în host ci se realizează un transfer direct între datele din host și
memoria GPU.
Ultimul pas al aplicației îl constituie conexiunea remote la baza de date MySQL. Fiecare
utilizator realizează o conexiune proprie la baza de date unde îi sunt publicate rezultatele
10
Capitolul 3. Implementarea aplicației
obținute în urma testelor.
Mai jos este prezentat blocul functional al aplicației:
Ilustrație 5: Proiectarea aplicației
Aplicația a fost dezvoltată utilizând și respectând principiile POO pe cât posibil, pentru a
putea realiza o modularizare cât mai corectă și pentru a împedica nevoia terminării task-urilor
repetitive pentru a ajunge la un rezultat final.
De asemenea aplicația a fost gândită în așa fel încât pe device-urile mai actuale, ce oferă
putere de calcul suplimentară unele teste să fie tratate în mod diferit spere deosebire de
arhitecturile ce dispun de un număr limitat de f eature-uri. Un astfel de exemplu ar putea fi clasa
ce calculează Gflop (floating point operation per secound).Gflop reprezintă o unitate de măsură a
puterii de calcul a unui sistem, măsurând numărul maxim de operații în virgulă mobilă ce sunt
executate pe secundă. Aici se verifcă dacă device-ul utilizat are o arhitectură Fermi sau mai sus
atunci numărul de blocuri de thread-uri este 32 altfel pentru versiunile mai vechi se lansează doar
16 blocuri de thread-uri.
template <int BLOCK_SIZE> __global__ void
matrixMulCUDA(float *C, float *A, float *B, int wA, int wB) //function defintion
if (block_size == 16)
multiplyMatrix<16> << <grid, threads >> >(d_C, d_A, d_B, dimsA.x, dimsB.x);
else
multiplyMatrix<32> << < grid, threads >> >(d_C, d_A, d_B, dimsA.x, dimsB.x);
Acesta este doar un exemplu, însă aproximativ fiecare test lucrează în mod diferit cu
device-ul pe care îl testează pentru a putea profita la maxim de capabilitățile oferite de
procesoarele puternice și de le feri de eventuale crash-uri pe cele ce dispun de versiuni sau
feature-uri depășite. Aplicația dispune și de un test status care poate fi PASS sau FAIL în cazul în
11
Adrian Aioanei
care aplicația este testată pe dispozitive ce nu aparțin Nvidia.
După cum am menționat aplicația dispune și de o interfață web iar structura de tabele din
spatele platformei arată în felul următor:
Interfața web este una minimală fiind dezvoltată doar pentru a oferi informații
suplimentare utilizatorilor cu privire la sistemul pe care s-au realizat testele și totodată pentru
compararea rezultatelor între diferitele device-uri.
Astfel dispune de o primă pagină ce oferă informații despre produs, cui se adresează,
capabilitățile hardware de care are nevoie aplicația și un link de download. Ce-a de-a doua
pagină prezintă prezintă sub forma unui tabele unde sunt extrase cele mai importante detalii
valorile obținute, fiecare înregistrare fiindu-i asociat un ID care este un link ce conduce spre o
pagină unde sunt afișate toate valorile obținute. Ultima pagină este una de feedback unde
utilizatorii își pot exprima părerile și pot anunța anumite bug-uri.
12
Ilustrație 6: Tabele din baza de date
Capitolul 3. Implementarea aplicației
3.2. Sincronizarea view-urilor și creearea interfeței în Qt
Interfața aplicației utilizează framework-ul Qt, care este cea mai bună soluție posibilă
atunci când limbajul utilizat este C++. Este un framework stabil, utilizat la scară globală care
pune la dispoziție programatorului un IDE. Când discutăm despre interfață sunt disponibile două
variante mai utilizate, una este utilizarea Qt Widget Application iar cea de-a doua Qt Quick
Application. Prima variantă este mai ușor de utilizat fiind asemănătoare cu ideea de C# în ceea
ce privește widget-urile (drag and drop). A doua variantă, Qt Quick Application utilizează QML
pentru interfață. Acesta este un limbaj de programare conform documentației Qt și permite
definirea obiectelor într-o sintaxă asemănătoare JSON și totodată oferă suport pentru JavaScript
și expresii combinate cu legături dinamice de proprietăți.
În acest proiect Qt nu este folosit la adevăratele sale capabilități deoarece principalele
operații sunt realizate în executabilul extern, în Qt existând doar 4 clase, care arată în felul
următor:
13
Ilustrație 7: Class View
Adrian Aioanei
Clasa model, moștenește clasa QAbstractListModel care este o clasă ce permite definirea
unei liste de modele iar în cazul de față este utilizată pentru generarea automată a câmpurilor
citite din fișierul xml, astfel nu mai este nevoie să fie înserat un nou label în cazul în care se
dorește adăugarea unei noi valori interfeței, acest lucru fiind suplinit și realizat automat de clasa
model.
Clasa Launcher se ocupă de administrarea primelor două view-uri și tot de aici este lansat
în execuție procesul extern ce formează backend-ul aplicației. Moștenește clasa QObject care
este clasa de bază a tuturor obiectelor din Qt deoarece aici se realizează anumite conexiuni între
semnale și sloturi.
Clasa AddData este cea care se ocupă de crearea view-ului final cu toate datele
obținute.Această clasa de asemenea moștenește QObject din același motiv mai sus menționat
(vezi Ilustrația 8). Clasa cupride 4 metode : readFromXml, insertDataToLabels, insertDataToDb
și checkError. Prima metodă stochează datele din xml într-o listă ce urmează a fi afișată, a doua
metodă se ocupă de popularea câmpurilor pentru interfața desktop, iar ultima se ocupă de
comunicarea cu baza de date. Funcțiile pentru înserare ar fi putut fi tratate unitar (aceiași metodă
să facă înserarea în baza de date cât și în interfața web) însă am optat pentru această metodă
deoarece utilizatorul are obțiunea a-și salva sau nu datele, astfel dacă este primit un răspuns
pozitiv (semnalul qmlSigYes) atunci se execută ambele metode, în caz contrar (qmlSigNo) este
excutată doar metoda insertDataToLabels().
Clasa OperatingSystemDetails este utilizată pentru obținerea de informații privitoare la
sistemul pe care s-a executat benchmark-ul. Rezultatele obținute de această clasă sunt afișate
doar în interfața web.
Mai jos este prezentat modul în care lucrează sloturile și semnalele în aplicație.
14
Ilustrație 8: View sync
Capitolul 3. Implementarea aplicației
3.3. Elemente de prelucrare ale imaginilor
Aplicația dispune de două scene realizate în OpenGl.
3.3.1. Scena 1
Testele ce privesc capabiliăți de randare sunt apreciate deoarece sunt creeate să utilizeze
cât mai bine posibil capabilitatile hardware disponibile, acest lucru reflectând performențele
dispozitivului testat.
Există mai multe metode de randare și fiecare metodă poate fi descrisă printr-o secvență
de pași care mai este numită și „commun volum rendering pipeline” și conține următorii pași:
1.Obținerea datelor
2.Procesarea datelor
3.Clasificarea datelor
4.Proiectarea datelor în planul imaginii
5.Afișarea pixelilor în planul imaginii
1. Obtinerea datelor
În acest pas datele sunt măsurate,calculate sau generate. Metodele de vizualizare
presupun ca aceste date din prima etapa să fie procesate corect și totodată se utilizază și scalarea
pentru vizualizare. În cazul de față datele de intrare se citesc dintr-un fișier raw.
2. Procesarea datelor
Datele de intrare nu pot fi imediat trimise către randare, ele ar trebui să fie procesate
corespunzător. Aceasta lucru poate include ajustarea luminozității,a constrastului, egalizarea
histogramei sau eliminarea zgomotului.
3. Clasificarea datelor
Clasificarea datelor este un proces de mapare a valorilor volumului de intrare catre
valorile utilizate pentru vizualizare. Pentru metode care afișează datele direct acest lucru de cele
mai multe ori implică in special setarea culorii și opacității. Propagarea luminii prin volum poate
fi precompusă și în timpul acestei etape. Clasificarea datelor este cel mai important pas din
"pipeline reandering".
4. Proiectarea datelor în planul imaginii
Acesta este un pas important în cadrul procesului de randare, unde în general se decide ce
voxeli și cum afectează pixelii rezultați planul imaginii virtuale.
5. Afișarea pixelilor în planul imaginii
Ultimul pas este afișarea rezultatului etapelor anterioare. Este de obicei similară sau chiar
aceeași pentru toate metodele și depinde în principal de arhitectura pe care algoritmul este
implementat.
Pentru creearea scenei s-a folosit un volum obținut de la adresa
(http://www9.informatik.uni-erlangen.de/External/vollib/ ) care este o pagină ce pune la
dispoziție volume strânse de la diferite persoane ce au fost în mare parte obținute utilizând un
computed tomography scan . Volumul este de dimensiune 32x32x32 (conține 32768 valori)
fiecare fiind reperezentată pe 8 biți. V olumul a fost obținut de la AVS, SUA și reprezintă
simularea densității unui electron Buckminster-Fullerene.
15
Adrian Aioanei
Mai jos este prezentat un preview al scenei ulterior urmând să fie descrise metodele ce
compun clasa:
16
Illustration 9: Scena 1
Capitolul 3. Implementarea aplicației
Clasa Scene1 conține urmtoarele metode:
•static void numberOfFps()
•static void randare()
•static void showResults()
•static void idle()
•static void update(int, int)
•static void clean()
•void bufferGLInit()
•void initiateOpenGL();
•void readDataFromFile( char*)
•uchar* readFile(const char*, size_t)
•extern ”C” void initCuda(const uchar*,cudaExtent) //Cuda call
•extern ”C” renderKernel(dim3, dim3, uint*, uint, uint, float) //Cuda call
•extern ”C” setTextureFilterMode( bool) //Cuda call
Scena are definită o serie de parametri pentru randarea pe GPU precum dimensiunea
blocului de thread-uri, dimensiunea gridului, care este setată în mod dinamic ca fiind lățimea și
înălțimea imaginii împărțită la dimensiunea blocului de thread-uri pe x și y. De asemenea este
definit un obiect de tip cudaExtent care este folosit să rețină dimensiunile volumului citit (width,
height, depth), un PBO (pixel buffer object) de tip Gluint și o structură cudaGraphicsResource
pentru transferul obiectului către device.
Prima metoda numberOfFps() îi poate fi intuit ușor rolul. Metoda calculează și setează
bara de titlu a scenei cu valoarea FPS obținută.
Cea de-a doua metodă în principal se ocupă de invocarea unei funcții kernel și anume
renderKernel(). Pe lângă acest lucru metoda utilizează cudaGraphicsMapResources ce primește
3 parametri, un număr întreg ce reprezintă cantitatea de resurse ce urmează să fie accesate de
CUDA, un pointer la un obiect de tip cudaGraphicsResource unde datele v-or fi mapate pentru
accesul CUDA și un obiect stream care este utilizat pentru sincornizare. După ce resursele au
fost mapate este nevoie de adresa prin intermediul căreia CUDA poate obține acces asupra
datelor, iar acest lucru este realizat de cudaGraphicsResourceGetMappedPointer ce primește 3
parametri. Primul returnează un pointer prin intermediul căruia se realizeaza accesul asupra
resurselor, al doilea parametru returnează dimensiunea buffer-ului accesat, iar ultimul parametru
este un pointer către adresa resursei mapate. Este important ca maparea oricărei resurse să fie
demapată deoarece, remaparea unei resurse fără demparea celei existente rezultă în undefined
behavior.
Următoarea metodă idle() este utilizată pentru incrementarea unei variabile globale ce
reprezintă dimensiunea pe axa z. După ce sunt setate toate funcțiile corespunzătoare în main este
apelată metoda glutMainLoop() ce este un loop infinit, acest lucru împiedicând executabilul să
randeze următoarea scenă. Datorită acestui fapt de fiecare dată când metoda idle() este apelată se
incrementează un index, iar cănd acesta ajunge la o valoare prestabilită se părăsește contextul
curent și se execută următoarele instrucțiuni.
Metoda update() este o metodă simplă ce primește doi parametri și se ocupă de setarea
perspectivei pentru obiectele desenate și de asemenea setează poziția camerei.
După cum am menționat anterior dacă resursele mapate în randare() nu sunt demapate
acest lucru condeuce către undefined behavior. Metoda clean() este creeată special pentru
prevenirea acestei probleme. Astfel metoda utilizează cudaGraphicsUnmapResources și
cudaGraphicsUnregisterResource care au ca rol eliberearea resurselor din contextul curent.
Metoda bufferGLInit se ocupă de inițializarea buffer-ului.Prima dată se utilizează
17
Adrian Aioanei
glGenBuffers pentru a genera un obiect (parametru de tip int), iar ca al doilea parametru este o
variabilă globală trimisă prin referință. Mai departe se face un bind urmat de glBufferData care
creează și inițializează buffer-ul. Ultima operație ce se realizează în funcția curentă constă în
utilizarea funcției cudaGraphicsGLRegisterBuffer ce primește trei parametri. Primul parametru
este un pointer către obiectul CUDA ce reține valorile la momentul actual, următorul este
obiectul creeat anterior cu glGenBuffers(), iar ultimul parametru este
cudaGraphicsMapFlagsWriteDiscard ce face ca datele din această resursă să nu fie păstrate, v-or
fi rescrise și toate valorile inițiale storate în această resursă nu v-or fi șterse.
Următoarea metodă initiateOpenGL, inițializează funcțiile de callback pentru GLUT. Se
setează dimensiunile ferestrei și titlul. Ca și funcții de callback pentru glutDisplayFunc() este
setată funcția showResults(), pentru glutReshapeFunc() funcția update(), iar pentru
glutIdleFunc() funcția idle(). Aici este verificată și versiunea minimă de openGL necesară.
ShowResults este metoda din care sunt apelate metoda de randare și numberOfFPS. Mai
departe este specificată poziția curentă de rasterizare ca fiind în colțul stânga jos, urmând a fi
apelată metoda glDrawPixels care citește pixeli din memorie și îi scrie în frame buffer relativ la
poziția curentă de rasterizare. Tot în acestă metodă este pornit timer-ul la începutul metodei și
este oprit la final apoi apleată metoda de calcul pentru FPS pentru obținerea rezultatelor.
ReadDataFromFile primește ca parametru path-ul către fișierul din care se face citirea.
Numărul de valori citit este returnat de dimensiunile oferite de un obiect cudaExtent care
precizează width, height and depth. Apoi într-un pointer se rețin valorile returnate de metoda
ReadFile. Tot în această metoda este aplelată funcția initCuda cu cei doi parametri despre care v-
om discuta ulterior.
Metodele ce se ocupă de rularea datelor pe GPU se află într-un fișier separat cu
extensia .cu ce va fi compilat separat de către NVCC (compilatorul celor de la Nvidia). Astfel în
fișierul cuda avem două variabile globale, o variabilă ce creează un obiect de tip textură ce
primește doi parametri. Primul paramtru specifică tipul de date iar al doilea faptul că datele v-or
fi normalizate. Pentru optimizare se creează un pointer către un obiect de tip cudaArray deoarece
acest obiect conține optimizări pentru legăturile cu texturi.
Prima metodă din fișierul cuda este InitCuda care este apletată din metoda
ReadDataFromFile și primește ca parametru datele citite și un obiect cudaExtent ce reprezintă
dimensiunea ce o ocupă datele citite. Inițial metoda creează un volum în memoria device-ului
utilizând valorile citite din fișier, dimensiunea volumului și un canal (cudaChannelFormatDesc)
care are ca rol descrierea formatului unui element din textura curentă. După ce volumul a fost
creeat datele sunt copiate într-un vector 3D utilizând structura cudaMemcpy3DParms, unde sunt
setate câmpurile pentru sursă, destinație, direcția de transfer (cudaMemcpyHostToDevice) și
dimensiunea volumului. În final sunt setate detaliile referitoare la textură, astfel textura este
setată ca fiind normalizată și returnează date în virgulă mobilă (float) ceea ce este foarte
important pentru cel de-al doilea parametru ce specifică modul de filtrare. Modul de filtrare este
setat ca cudaFilterModeLinear ce implică o interpolare liniară ale celor mai apropiate coordonate
din textura de intrare. Numărul de valori pentru interpolare este ales în funcție de dimensiunea
setului de date, astfel pentru 1D avem interpolare între două valori, pentru 2D avem interpolare
între 4 valori, iar în cazul de față fiind un volum interpolarea se face între 8 valori. Ultimul
parametru ce trebuie setat pentru textură este modul de adresare (cudaAddressModeWrap) care
este setat pentru toate cele trei dimensiuni.
18 text.normalized = true; // access with normalized texture coordinates
text.filterMode = cudaFilterModeLinear; // linear interpolation
text.addressMode[0] = cudaAddressModeWrap ; // wrap texture coordinates
text.addressMode[1] = cudaAddressModeWrap ;
text.addressMode[2] = cudaAddressModeWrap ;
Capitolul 3. Implementarea aplicației
RenderKernel este doar o metoda intermediară ce este invocată pentru rularea kernel-
ului. Metoda primește șase parametri (gridSize,blockSize, output, width, height, depth) care sunt
utilizați pentru instanțierea metodei rander.
Este important de menționat modul în care CUDA realizează management-ul thread-
urilor lansate. Deoarece toate firele dintr-un grid execută aceleși kernel, managemant-ul lor se
bazează pe un sistem de coordonate care are ca rol identificarea unei resurse în mod unic. Aceste
thread-uri sunt organizate într-o ierarhie pe trei niveluri numite grid, block și thread, iar
coordonatele pentru aceste entități sunt asignate de către CUDA la runtime (programatorul nu
poate influența în nici un fel modul de indexare al thread-urilor), ulterior având posibilitatea de a
fi accesate valorile indicate de aceste thread-uri în funcția kernel. Când un thread execută funcția
kernel se compun coordonatele corespunzătoare oferite de threadId și blockId astfel
identificându-se în mod unic valoarea dorită. Deasupra thread-urilor și al block-urilor se află
grid-ul care este organizat ca o matrice bidimesională de blocuri. Numărul de block-uri pentru
fiecare dimensiune este specificat de primul parametru dat la lansarea kernel-ului. La nivelul
inferior toate block-urile unui grid sunt organizate într-un vector tridimensional de thred-uri, iar
fiecare threadId se compune din coordonata pe x care este threadId.x, coordoante pe y care este
threadId.y și coordonata pe z care este threadId.z.
Managementul thread-urilor incluzând creearea, programarea, sincronizarea este realizat
în totalitate de către SM. Pentru a gestiona eficient numărul mare de thread-uri fiecare SM
folosește arhitectura SIMT (Single Instruction, Multiple Thread) .
Pentru a executa intrucțiuni pe GPU se lansează funcția care se dorește a fi executată (de
obicei aceasta se numește kernel) sub forma unui grid de blocuri de threduri care este trimis la
device. O dată ajuns kernel-ul la GPU acesta trimite căte SM (Streaming Processor) unul sau mai
multe blocuri. Fiecare SM își creeaza o coada de warp-uri (un warp conține de obicei 32 de
thred-uri) care este executat pe un clock. Dacă execuția warp-ului nu s-a încheiat într-un clock
acesta este adăugat din nou în coadă. Acest lucru se repetă până se termină de executat blocul
curent. Atunci când s-a terminat execuția este solicitat un nou bloc de la GPU.
Mai jos este prezentată o imagine ce sugerează modul de organizare al thread-urilor:
19
Adrian Aioanei
Ultima metodă utilizată este render ce primește 4 parametri, un pointer unsigned int
pentru output, lățimea imaginii, înălțimea și depth. Prima dată kernel-ul identifică resursele
necesare utilizând operațiile descrise mai sus pentru identificarea coordonatelor curente:
Amintim faptul că în contextul curent a fost creeată și textura normalizată iar pentru a fi
posibil citirea din obiectul tex este nevoie de o normalizare pentru valorile reținute în x și y
pentru a putea citi din textură de la coordonatele corecte. După acest pas se folosește funcția
tex3D ce primește ca parametri obiectul ce conține textura și coordonatele de la care se face
citirea, iar datele returnate sunt reținute într-un voxel. După ce acești pași au fost îndepliniți se
face un test ce implică ca x să fie mai mic decât lățimea imaginii și y să fie mai mic decât
înălțimea imaginii. Dacă ambele condiții sunt respectate se compune valoarea finală a piexelui.
Ca și o optimizare pentru a reduce numărul de biți implicați în calcul se poate folosi __umul24
care realizează o înmulțire pe 24 de biți, însă s-a dovedit că această abordare ar renta doar pentru
dispozitivele mai vechi care au o arhitectură orientată în acest sens, deoarece dispozitivele mai
actuale pierd mai mult timp să emuleze cei 24 de biți pe 32 de biți decât ar dura propriu zis
înmulțirea.
3.3.2. Scena 2
Scena doi iși propune realizarea unui filtru Gaussian. De cele mai multe ori acest filtru
este utilizat pentru a reduce zgomotul din imagini, rezultatul acestui filtru fiind unul de netezire.
Acest filtru setează valoarea curentă ca fiind o medie a vecinilor. Numărul vecinilor luați în
considerare la fiecare interație este specificat printr-o variabilă numită radius. Filtrul este
modelat de următoarea relație matematică:
f(x,y)∗k(x,y)=∑i=−w/2w/2
∑j=−w/2w/2
f(x+i,y+i)
De cele mai multe ori valorile măștior sunt normalizate astfel încât:
∑i=−w/2w/2
∑j=−w/2w/2
k(i,j)=1
Avantaje ale unui filtru Gaussian:
•muchiile și liniile în direcții diferite sunt tratate la fel, este circular simetric
•este posibil ca matricea nucleu să fie separată într-un produs de 2 vectori (orizontal și
vertical).
h3x3=
[121
242
121] Sau h3x3=
[1
2
1]∗[121]
Algoritmul care realizeaz ă acest lucru exprimat în mod empiric este:
•Pentru fiecare pixel (i,j):
◦Se aplică formula operația de convoluție pentru matricea de pixeli a imaginii și
vectorul orizontal al matricii nucelu descompuse : (1 2 1).
20 uint x = ( blockIdx .x * blockDim .x) + threadIdx .x;
uint y = ( blockIdx .y * blockDim .y) + threadIdx .y;
Capitolul 3. Implementarea aplicației
•Pentru fiecare pixel (i,j):
◦Se aplică formula de convoluție pentru matricea de pixeli și vectorul vertical al
matricii filtru descompuse : (1 2 1 ) ^ T.
Apoi este necesară o matrice intermediară în care se v-or reține rezultatele obținute după
aplicarea primei etape de convoluție și asupra căreia se va aplica a doua operație de convoluție
(multiplicarea cu vectorul de valori verticale).
Clasa conține următorele metode :
static float fps_cout_complex;
static void varySigma();
static void computeFPS();
static void display();
static void timerEvent(int value);
static void reshape(int x, int y);
void initCuda(bool useRGBA);
static void cleanup();
void initGLResources();
void initGL();
int runSample();
void loadImageData();
bool checkCUDAProfile(int dev, int min_runtime, int min_compute);
O parte din metode au fost detaliate anterior așa ca de acesta data v-om urmări principiile
de baza care stau la creeare acestei scene.
Pentru valoarea fiecărui pixel se realizează doar două adunări și o înmulțire, iar pentru
îmbunătațirea performanțelor de citire se folosește tex2D.
Este setat un radius, iar fiecare thread ce operează la nivel de coloană adună (2 * r + 1)
pixeli, iar media lor este stocată pentru pixelul curent (i,j).
21
∑i=0(w)
∑j=0(2∗r+1)
ai,j
2∗r+1
Adrian Aioanei
În urma prelucrării inițiale valorile primilor pixeli v-or fi reprezentați ca în imaginea de mai jos:
Fiecare thread operează pe un rând și ajustează valorile pixelilor pe întreg rândul. De aici
se desprinde o regula generală de prelucrare a pixelilor.
•(i – (r + 1), j)
•(i + (r + 1), j)
Mai jos este prezentată procesarea pixelului 12, a cărui valoare este formată din suma
reprezentată de pixelii de la (12- (r + 1), j) la (12 + (r + 1), j).
22
Illustration 10
Illustration 11
Capitolul 3. Implementarea aplicației
Ulterior operația se realizează și pe coloană așa cum este exemplificat în imaginea de mai
jos:
După ce au fost realizate toate aceste operații rezultatul opținut este exemplificat în
imaginea următoare:
23
Illustration 12: Scena 2
Adrian Aioanei
3.4. Lățimea de bandă. Teoretic și Practic
Unul dintre principalele lucruri pe care trebuie luate în considerare atunci când se
testează o placă video este lățimea de bandă. Latimea de banda este de fapt viteza RAM și se
măsoară în gigaocteți pe secundă (GB/s sau MB/s). Cu cât lățime de bandă este mai mare, cu
atât placa dispune de performanțe sporite. O placă video cu o lățime de bandă mai mare a
memoriei poate desena mai rapid și poate desena imagini de calitate superioară. De asemenea,
trebuie luată în considerare viteza de desenare a unității de procesare grafică. O problemă sau
chiar un dezavantaj îl poate constitui o placă video cu un GPU foarte rapid și o lățime de bandă
limitată de memorie. Lățimea de bandă limitată conduce spre blocaje și forțează GPU-ul să
petreacă mult timp fără să facă nimic în timp ce așteaptă o memorie RAM lentă.
Valoarea teoretică a lățimea de bandă este dată de următoarea formulă:
BWTheoretical=MemoryClockRate∗106+((MemoryBuss)/8)∗2/109
Metodele ce implementează testarea lățimii de bandă au următoarea definiție:
bool testDeviceToHostTransfer (unsigned int memSize, memoryMode memMode,
int& mem, int& bands);
bool testHostToDeviceTransfer (unsigned int memSize, memoryMode memMode,
int& mem, int& bands);
bool testDeviceToDeviceTransfer (unsigned int memSize, memoryMode memMode,
int& mem, int& bands);
Metodele sunt dezvoltate modular pentru a face posibilă o dezvoltare ulterioară. Astefel
că fiecare metodă primește ca parametru dimeniunea memoriei care urmează a fi transferată,tipul
memoriei, iar prin referința sunt reținute valoarea setată pentru dimensiunea memoriei transferate
și lățimea de bandă opținută în urma transferului.
Fiecare metodă primește ca parametru un enum memoryMode care specifică tipul de
transfer al datelor ce urmează a fi testate. Aici dispunem de două variante : Pageable Data
Transfer și Pinned Data Transfer. Alocaările pe host (GPU) pot fi accesate în mod implicit, însă
GPU-ul nu poate accesa date în mod direct de la o memorie gazdă. Așadar atunci când este
invocat un transfer dinspre host către device, dirver-ul CUDA trebuie să aloce o zonă de
memorie temporară „ pinned”, iar apoi să se inițializeze transferul dintre zona temporară alocată
de driver-ul CUDA și device (GPU). Acest proces necesitând operații suplimentare și consumând
timp s-a optat pentru modelul de transfer pinned care nu mai implică alocarea zonei de memorie
suplimentare, transferul realizându-se direct.
Metoda verifică parametrul primit dacă este PINNED sau PAGEABELE și acționează în
consecință. Dacă parametru este PINNED se folosește cudaMallocHost, iar dacă este
PAGEABLE se folsește o funcția de alocare malloc pentru alocarea spațiului de memorie pe
host. Apoi urmează inițializarea datelor cu niște variabile aleatorii, urmând a se realiza transferul
dintre host și device utilizând cudaMemcpyAsync respectiv cudaMemcpy.
Pentru o mai bună precizie în ceea ce privește rezultatele opținute se realizează un set de
repetări ale testului. Astfel timer-ul se pornește o dată, apoi se realizează 100 de interații cu
detaliile mai sus prezentate, după care este oprit.
Valoarea finală este calculată după cum urmează:
//calculate bandwidth in MB/s
bandwidthInMBs = (( float)(1 << 10) * memTrans * ( float)MEMCOPY_ITERATIONS ) /
(elapsedTime * (float)(1 << 20));
24
Capitolul 3. Implementarea aplicației
3.5. Gflop
În informatică (știința calculatoarelor), FLOPS (scris și flops sau flop/s) este un acronim
ce provine de la expresia engleză floating point operations per second (tradus: operații în virgulă
mobilă pe secundă). FLOPS reprezintă o unitate de măsură a puterii (= a vitezei) de calcul a unui
calculator sau sistem de calcul, măsurând numărul maxim de operații în virgulă mobilă, (de
regulă adunări și înmulțiri), ce sunt executate pe secundă. Unitatea FLOPS își găsește folosul
mai ales în domeniul calculelor științifice, la care calculul în virgulă mobilă este frecvent folosit.
FLOPS nu este unitate SI, dar poate fi apreciat ca o unitate de măsură având mărimea 1/s.
Clasa ce implementează măsurarea numărului de flop/s realizează inmulțirea a două
matrici în paralel. Clasa dispune de două metode principale care realizează aproximativ toată
funcționalitatea. Prima metodă matrixMultiply primește ca parametru block size care este un
parametru dependent de device și cele două matrici ce urmează a fi înmulțite. Metoda
inițializează matricile A și B cu valori aleatoare și alocă memorie corespunzătoare matricilor de
pe device. După aceste operații se pornește un timer iar înmulțirea matricelor se realizează de
300 de ori pentru o măsurare mai exactă.
În cadrul metodei matrixMul apelată din matrixMultiply se realizează înmulțirea propriu
zisă. Funcția primște ca parametri matricea rezultat, cele două matrici ce urmează a fi înmulțite și
dimensiunea lor (widthA, widthB). Matricea este împărțită în submatrici și sunt reținute într-o
zonă de memorie shared. Ulterior se realizează înmulțirea propriu zisă unde fiecare thread se
ocupă de un element, scrierea rezultatelor în matricea finală se realizează tot de fiecare thread în
parte. Metoda returnează un tuplu ce conține valoarea obținută pentru Gflop, timul necesar
execuției, numărul de operații realizate și dimensiunea grupului de lucru.
3.6. Interfata Web
Interfața web a fost gândită în așa fel încât să ofere utilizatorului un mod de comparare a
performanțelor obținute raportat fiind la alți utilizatori. Interfața web comunică cu aplicația
desktop printr-un driver și dispune de o pagină unde este prezentat produsul și utilizatorii au
posibilitatea să descarce aplicația, o pagină unde este prezentat un preview al rezultatelor
obținute cu un id care face legătura către o pagină unde sunt disponibile rezultatele complete, iar
ultima pagină constă într-un formular ce are ca scop primirea feedback-ului de la utilizatori în
vederea îmbunătățirii aplicației.
Datorită faptului că în testarea performanțelor unui GPU intră și factori externi și
rezultatele nu depind stric numai de capabilitățile GPU, în interfața web s-au inclus o serile de
valori care ar putea influența obținera de rezultate diferite în urma testării aceleași plăci video dar
pe sisteme diferite. Pentru a preveni această confuzie o dată cu transmiterea datelor obținute în
urma testelor sunt transmise și date referitoare la tipul de procesor, arhitectura acestuia, numărul
de core-uri de care dispune CPU-ul, dimensiunea RAM, ș.a.m.d.
Am putea afirma faptul că acestă interfață are rol dublu. Pe lângă aspectele mai sus
menționate aplicația web ar putea juca rol și de urmărire a bug-urilor și a eventualelor
incompatibilități. Acest lucru fiind posibil deoarece atunci când aplicația este livrată
utilizatorului valorile câmpurilor din fișierul xml sunt setate pe NULL, acestea urmând a fi
modificate o dată ce testele returnează valorile. Astfel valorile ce sunt introduse în baza de date
cu valori de NULL indică existența unor probleme ce trebuie investigate.
25
Adrian Aioanei
Capitolul 4. Evaluare
4.1. Testarea aplicației
Aplicația a fost testată pe următoarele sisteme de operare :
•Windows 7
•Windows 8
•Windows 8.1
•Windows 10
Device-urile pe care a fost testată sunt următoarele:
•GeForce GTX 660 Ti
•GeForce GTX 860M
•GeForce GTS 450
•GeForce GTX 950M
•GeForce GT 750M
•GeForce GT 740M
Ca aplicația să ruleze are nevoie de:
•Dispozitiv Nvidia ce dispune de compute camability >= 3.0
•Microsoft visual C++ redistributable 12
4.2. Rezultate
Primul grafic prezintă valorile FPS obținute pentru cele două scene:
26 GTX 660 TiGTX 860MGTS 450GTX 950MGT 750MGT 740M0102030405060708090100
89
7582
7072
6478
7075
6165
51
Scena 1 (FPS)
Scena 2 (FPS)
Capitolul 4. Evaluare
Diagrama de mai jos prezintă rezultatele obținute în urma testării memoriei de tip
pinned . Este de la sine înțeles că transferul cel mai rapid v-a fi realizat în memorie device-ului.
Urmează prezentarea valorilor obținute pentru memoria de tip pagable care este cu mult
inferioară valorilor obținute pentru memoria pinned. După cum se poate observa și din grafic în
unele situații memoria de tip pinned are un randament cu 30% mai ridicat lund în considerare
transferul host-device și device-host.
27GTX 660 TiGTX 860MGTS 450TX 950MGT 750MGT 740M020000400006000080000100000120000
Host to Device
Device to Host
Device to Dev ice
660 Ti860M450950M750M740M020000400006000080000100000120000
Host to Device
Device to Host
Device to Device
Adrian Aioanei
Următorul grafic prezintă valorile obținute pentru GFLOP (giga floting point operation
per secound).
4.3. Probleme întâmpinate și soluții de rezolvare
4.3.1. Creearea backend-ului aplicației
O prima problemă care a intervenit a constat în găsirea unei modalitati de integrarea a
tuturor tehnologiilor mai sus enunțate. Dacă partea de openGL și interfață grafică nu au constituit
neapărat o provocare în integrare, problema reală a intervenit o dată cu integrarea API-ului
CUDA pentru procesarea GPU. Din punct de vedere tehnic existau 3 opțiuni valide:
•Importarea tuturor componentelor în interfață
•Creearea unui dll care să fie importat în interfață
•Lansarea unui nou proces extern din interfață
Chiar daca cei de la Nvidia oferă suport pentru integrarea CUDA în Visual Studio, acest
lucru nu este disponibil și pentru Qt (framework-ul pentru interfață). Sunt disponibile anumite
soluții pe Linux însă având în vedere că aplicația este destinată utilizatoruilor Windows deoarece
nu sunt multe persoane interesate de performanțele GPU care să folosescă Linux am renunțat la
această idee. Există variante și pentru Windows însă configurarea mediului ridică anumite
probleme suplimentare.
A doua opținue a constat în creearea unui dll în Visual Studio care să realizeze operațiile
necesare și să exporte funcțiile ce urmează a fi apelate în Qt. La acestă variantă am renunțat
datorită debug-ului greu de urmărit.
Ultima variantă și totodată cea aleasă de mine constă în creearea backend-ului separat de
interfață. Astfel am creeat un executabil separat cu VS/OpenGL/CUDA care este lansat în
execuție din Qt, tot acolo realizându-se și sincronizarea. Interfața așteaptă până ce aplicația și-a
încheiat execuția și a populat xml-ul, apoi afișează datele finale.
28 GTX 660 TiGTX 860MGTS 450GTX 950MGT 750MGT 740M050100150200250
223
118131
115
6561GFlop
Capitolul 4. Evaluare
4.3.2. Calcularea corectă a lățimii de bandă
Pentru a putea afirma că lățimea de bandă este calculată corect, conform datelor oferite
de Nvidia valoarea obținută trebuie să se încadreze între 70 – 85% din lățimea de bandă teoretică
dată de următoarea formulă:
BWTheoretical=MemoryClockRate∗106+((MemoryBuss)/8)∗2/109
Aplicând formula mai sus specificată pe un device GTX 750M cu următoarele
specificații :
•Memory Clock Rate = 1000 MHZ
•Memory Bus = 128 bits
Aplicând formula de mai sus pe specificațiile deținute de GTX 750M obținem o lățime de
bandă de 64.16 GB/sec (aici trebuie luat în considerare tipul de memorie, astfel pentru DDR
Memory Clock Rate va fi înmulțit cu 2 deoarece memoria DDR transferă 2 biți/clock). Datorită
faptului că array-ul transferat nu era suficient de mare pentru a satura întreaga lățime de bandă,
rezultatul obținut în urma testării era de aproximativ 30 GB/s, această valoare fiind aproximativ
50% din lățimea de bandă teoretică. Modificând dimensiunea datelor de transfer și tipul acestora
s-a obținut o lățime de bandă pentru dispozitivul menționat de 47.41 GB/s (Illustration 3) care
este aproximativ 73.9% ceea ce înseamnă că se încadrează în specificațiile menționate. O altă
problemă ce conducea spre rezultate eronate era metoda de calcul. Cum am specificat anterior
clock-ul trebuie înmulțit cu timpul de memorie, în cazul lui 750M acesta fiind DRAM, valoarea
frecvenței trebuie multiplicată de două ori. Aceaste variații între valori pot să apară din motive
diverse precum: ECC pornit sau oprit, tipul de memorie utilizat, frecvența GPU.
29
Adrian Aioanei
Capitolul 5. Concluzii
5.1. Gradul în care s-a realizat tema propusă
Gradul propus de dezvoltare al aplicației a fost atins, însă pe măsură ce am lucrat la
proiect au apărut ideii și soluții noi de implementare și îmbunătățire care datorită timpului scurt
nu au putut fi implementate complet.
Aplicația urmărește testarea celor mai importante capabilități de care ar trebui să dispună
un GPU pentru obținera de performanțe crescute fie vorba de jocuri, procesare de imagini sau
procesarea masivă de date.
Această aplicație este destinată gamei de mijloc de utilizatori. Este destinată acelor
utilizatori care care doresc să-și interogheze capabilitățile hardware într-un mod predefinit fără a
avea momentan posiblilitatea de a-și personaliza aceste teste.
Consider că aplicația în momentul de față se află într-un stadiu matur, stabil fiind
capabilă să ofere utilizatorului final date concrete referitoare la capabilitățile device-ului, fiind
testată pe șase dispozitive diferite și 4 sisteme de operare. Aplicația este momentan disponibilă
doar pentru Windows însă este posibil și un build pentru Linux și MacOS.
5.2. Dezvoltare ulterioară
Aplicația a fost dezvoltată modular pentru a putea fi îmbunătățită fără a aduce modificări
substanțiale codului. Astfel intr-o următoare dezvoltare mi-aș dori să implementez un nou view
care să ofere posibilitatea utilizatorului de a-și personaliza benchmark-ul. Rudimentar noul view
ar trebui să ofere următoarele posibilități:
30
Bibliografie
Bibliografie
[1] Claus Jespersen “ Monte Carlo Evaluation of Financial Option using a GPU ”, Aarhus
University Denmark, 2015
[2] Enhua Wu, Youquan Liu “ Emerging Tehnology about GPGPU ”, University of Macau, 2008
[3] Vaslily Volkov, “ Understanding Latency Hiding on GPUs ”, University of Calofornia at
Berkeley, 2016
[4] Yusuke Fuji, Takuya Azumiy, Nobuhiko Nishioy, Shinpei Katoz and Masato Edahiroz “ Data
Transfer Matter for GPU Computing ”
[5] H.G. Schmidt “ Benchmarking Benchmarks ” ,University Rotterdam, doctoral dissertation
2012
[6] Rune Johan Hovland “ Latency and Bandwidth Impact on GPU-systems ”, Trondheim,
Norway, December 17, 2008
[7] Ondrej Mosnácek “ Key derivation functions and their GPU implementation ”, Brno, Spring
2015
[8] David Luebke, Greg Humphreys, “ How GPUs Work” February 2007
[9] Henry Wong, Misel-Myrto Papadopoulou, Maryam Sadooghi-Alvandi, and Andreas
Moshovos, “Demystifying GPU Microarchitecture through Microbenchmarking ”
[10] Michael Romero, “ Volume Ray Casting Techniques and Applications using General
Purpose Computations on Graphics Processing Units ”
[11] Elizabeth G. Reid “ Design and Evaluation of a Benchmark for Main Memory Transaction
Processing Systems ”, MIT June 2009
[12] Jukius Sandgren, “ Transfer Time Reduction of Data Transfers between CPU and GPU ”,
July 2013
[13] Desislava Ivanova, “ Performance Evaluation and Benchmarking of Modern GPU
Architectures”, September 2015
[14] CHRISTIAN BIENIA, “ BENCHMARKING MODERN MULTIPROCESSORS ”, JANUARY
2011
31
Adrian Aioanei
Anexe.
Anexa 1.
32
Copyright Notice
© Licențiada.org respectă drepturile de proprietate intelectuală și așteaptă ca toți utilizatorii să facă același lucru. Dacă consideri că un conținut de pe site încalcă drepturile tale de autor, te rugăm să trimiți o notificare DMCA.
Acest articol: Calculatoare și tehnologia informației [603872] (ID: 603872)
Dacă considerați că acest conținut vă încalcă drepturile de autor, vă rugăm să depuneți o cerere pe pagina noastră Copyright Takedown.
