Gpu Programming

2015/05/08 14:50 1/15 GPU Programming

ASC Wiki - http://cs.curs.pub.ro/wiki/asc/

GPU ProgrammingIn ultimii ani folosirea GPU (Graphics Processing Unit) pentru anumite tipuri de procesari non-grafice aluat amploare. Acest lucru se datoreaza in principal:

Diferentelor din ce in ce mai mare de putere de procesare dintre CPU si GPU in favoarea acestora1.din urmaStandardizarea de API-uri care usureaza munca programatorilor pentru a folosi GPU-ul2.Existentei unitatilor GPU in aproape fiecare computer consumer fie integrat in soclul procesorului3.fie dedicat

Fig. 1:Cresterea uimitoare a puterii de procesare a GPU in ultimii ani

Observatie: Faptul ca un GPU are mai multi GFlops decat un CPU nu inseamna ca s-ar descurcamai bine la rularea de software care nu executa calcule in mod intensiv. Un procesor de uz generalare o cantitate mare din tranzistori dedicati cachingului si controlului fluxului. GPU-uri, in constrast, isibazeaza puterea pe executia acelorasi instructiuni, in paralel, pe cat mai multe date, iar pentruobtinerea de performante mari trebuie evitat branchingul (if-uri).

Aplicatiile existente sunt in domenii precum:

prelucrari video si de imaginilsimulari de fizica in jocurildomeniile clasice pentru High Performance Computing:l

finantem

Last update: 2015/05/08 14:36 asc:lab10:index http://cs.curs.pub.ro/wiki/asc/asc:lab10:index

http://cs.curs.pub.ro/wiki/asc/ Printed on 2015/05/08 14:50

stiintemindustria petroluluindinamica fluidelorn

design electronic (VLSI)m

Printre aplicatiile care folosesc GPU pentru calcule si sunt disponibile utilizatorilor obisnuiti, se numara:

Clientul Folding@Home - procesari legate de proteine ce se realizeaza si pe GPUlCoreAVC - Offloading pentru decodificarea video pe GPU, pentru un CPU mai putin utilizat (ex: 4%lpentru pentru FullHD)Jocurile care folosesc NVIDIA Physx : Gears of War, Age of Empires III, Assassin's Creed etclBadaboom Media Convertor - pentru unele placi video de la NVIDIAlAvivo Video Convertor - pentru unele placi video de la ATIlDXVA (Direct[X] Video Acceleration)lFiltre pentru GIMPlElcomsoft Password Recovery - password cracking atat pentru NVIDIA cat si pentru ATI.l

De asemenea, exista si niste demo-uri de la AMD.

Principalii producatori de core-uri IP (intellectual property) tip GPU sunt:

INTEL http://en.wikipedia.org/wiki/List_of_Intel_graphics_processing_units (integrat procesoare x86,lindustrie consumer/embedded/hpc)NVIDIA http://en.wikipedia.org/wiki/List_of_Nvidia_graphics_processing_units (dedicat pe magistralalAGP/PCIe sau integrat procesoare ARM, industrie consumer/embedded/hpc)AMD http://en.wikipedia.org/wiki/List_of_AMD_graphics_processing_units (dedicate pe magistralalAGP/PCIe sau integrat procesoare x86, industrie consumer/hpc)Imagination Technologies/PowerVR http://en.wikipedia.org/wiki/List_of_PowerVR_products (integratlprocesoare ARM, industrie consumer/embedded)Qualcomm/Adreno http://en.wikipedia.org/wiki/Adreno (integrat procesoare ARM, industrielconsumer/embedded )Vivante http://en.wikipedia.org/wiki/Vivante_Corporation (integrat procesoare ARM, industrielconsumer/embedded)

Ahitectura generala GPUProcesorul grafic are o arhitectura de tip SIMD si un spatiu propriu de memorie (GPU dedicat VRAM,GPU integrat parte din RAM). In cadrul unui sistem procesorul general care coordoneaza executiaeste numit HOST (CPU) pe cand unitatea care efectueaza calculele este numit TARGET (GPU). Orice felde procesare necesita in prealabil un transfer din spatiul de memorie de la CPU catre spatiul dememorie de la GPU. In cazul unui procesor grafic dedicat acest transfer se face printr-o magistrala(PCIe, AGP, USB). Viteza de transfer RAM-VRAM via magistral este inferioara vitezei RAM/VRAM siprin urmare necesita eficientizata (exemplu facuta in paralel cu calculele).

2015/05/08 14:50 3/15 GPU Programming


Alternative la arhitecturi tip GPUIBM propune un BladeCenter bazat pe doua procesoare Cell. Comparativ cu GeForce 8800, Cell arelmai putine unitati de calcul si latime de banda a memoriei mai mica, dar ofera o frecventa mai maresi memorie mai multa pentru unitatile de calcul (256KB versus 8-16KB pentru GeForce 8)Intel este implicat prin proiectul Larrabee, care este un chip masiv paralel destinat sa aiba rolul delGPU , dar si de coprocesor. Se vorbeste despre 16 pana la 24 de nuclee cu unitati SSE de 512biti(sau 16 operatii pe 32 de biti per ciclu per core). Fiecare core are un cache L1 de 32KB si L2 de256KB, toate insotite de un cache general de 4MB. De asemenea, va avea avantajul major de a fibazat pe arhitectura x86. Din cauza intarzierilor de dezvoltare insa, prima generatie de chipuriLarabee nu va fi accesibila consumatorilor, ci doar pentru High Performance Computing. Larrabeearata o scalabilitate excelenta.

Totusi, bariere semnificative in adoptarea GPU inca exista pentru programatorii care doresc sautilizeze puterea de calcul a acestora. GPU-urile sunt construite pentru calcul paralel (indeosebimultimedia), modelul de programare este diferit de cel conventional iar optimizarile sunt de regulastrans legata de arhitectura target pe care va rula aplicatia.

Framework/API pentru GPGPU - NVIDIA CUDA

Implementarea nVidia pentru GPGPU (General-Purpose computation on Graphics Processing Units) senumeste CUDA (Compute Unified Device Architecture) si permite utilizarea limbajului C pentruprogramarea pe GPU-urile proprii. Deoarece una din zonele tinta pentru CUDA este High PerformanceComputing, in care limbajul Fortran este foarte popular, PGI ofera un compilator de Fortran ce permitegenerarea de cod si pentru GPU-urile Nvidia. Exista binding-uri pana si pentru Java (jCuda), Python (PyCUDA) sau .NET (CUDA.NET).

Pentru folosirea CUDA, Nvidia pune la dispozitie:

driver cu suport CUDA (Windows, Linux)l



CUDA toolkit:lcompilatorul cu suport CUDA (nvcc)mdebugger (cuda-gdb)mprofiler (cudaprof)mbiblioteca FFT (Fast Fourier Transform - CUFFT)mbiblioteca BLAS (Basic Linear Algebra - CUBLAS)mdocumentatie pentru cele de mai susm

CUDA SDKlexemple de folosire CUDA (atat pentru Linux cat si pentru Windows-Visual Studio)m

Pentru eliminarea segmentarii pe diferite produse si frameworkuri si ingreunarea astfel a dezvoltariide software care sa profite de capabilitatile de procesare oferite de procesoare tip GPU, CELL sausimilare, a fost adoptat standardul OpenCL. Conceput de acelasi grup de lucru care, printre altele, seocupa de standardele OpenGL, OpenCL permite programarea pe platforme eterogene. Exista dejasuport pentru acesta din partea NVIDIA, ATI si IBM. Mac OS X 10.6 foloseste deja OpenCL pentruaccelerarea interfetei grafice.

Desi majoritatea placilor video de ultima ganeratie de pe piata suporta procesari pe GPU, amandouacompaniile au dezvoltat si o serie de placi grafice special concepute ca GPGPU (General PurposeGPU): ATI are FireStream si NVidia are TESLA.NVidia ofera si sisteme complete dedidate acestui fel de procesari prin seria Quadro.

Placile NVIDIA Tesla sunt similare placilor desktop, dar le lipsesc conexiunile de output video. Imagini:NVIDIA, ATI.

Arhitectura NVIDIA

Motivul discrepantei intre performanta in virgula mobila dintre CPU si GPU este faptul ca GPU suntspecializate pentru procesare masiv paralela si intensiva computational (descrierea perfecta ataskurilor de randare grafica) si construite in asa fel incat majoritatea tranzistorilor de pe chip seocupa de procesarea datelor in loc de cachingul datelor si controlul fluxului executiei.

GPU sunt potrivite pentru paralelismul de date (aceleasi instructiuni sunt executate in paralel pe maimulte unitati de procesare) intensiv computationale. Datorita faptului ca acelasi program esteexecutat pentru fiecare element de date, sunt necesare mai putine elemente pentru controlul fluxului.Si deoarece calculele sunt intensive computational, latenta accesului la memorie poate fi ascunsa princalcule in locul unor cache-uri mari pentru date.

Procesarea data-parallel mapeaza elemente de date pe theaduri de procesare paralele.

2015/05/08 14:50 5/15 GPU Programming


Detaliile de implementare ale arhitecturii chipurilor nVidia sunt secrete. In consecinta, modul in carefunctioneaza unitatile aritmetice si benchmark-uri oficiale ale performantelor nu sunt cunoscute.

Arhitectura GT 200

GT 200 este o extensie a arhitecturii G80. G80 este prima arhitectura care a suportat modelul deprogramare CUDA. A fost proiectata modular astfel incat sa permita cresterea numarului de elementecomputationale in generatii succesive.

Un Streaming Processor (SP) este un microprocesor cu executie secventiala, ce contine un pipeline, 2unitati aritmetico-logice (ALU) si o unitate de calcul in virgula mobila (FPU). Nu are un cache, fiind bundoar la executia multor operatii matematice. Un singur SP nu are performante remarcabile, insa princresterea numarului de unitati, se pot rula algoritmi ce se preteaza paralelizarii masive.

8 SP impreauna cu 2 Special Function Units (SFU) sunt incapsulate intr-un Streaming Multiprocessor.Fiecare SFU contine 4 unitati pentru inmultire in virgula mobila, utilizate pentru operatiitranscendente (sin, cos) si interpolare. MT se ocupa cu trimiterea instructiunilor pentru executie la SPsi SFU.

Pe langa acestea, exista si un cache (de dimensiuni reduse) pentru instructiuni, unul pentru dateprecum si memorie shared de 16KB, partajata de cele 8 SP. Urmatorul nivel de incapsulare esteTexture / Processor Cluster (TPC). Acesta contine 3 SM, logica de control si un bloc de handling pentrutexturi. Acest bloc se ocupa de modul de adresare al texturilor, logica de filtrare a acestora precum siun cache pentru texturi. Pentru GF 200, un TPC contine deci 24 SP si 6 SFU.



Ultimul nivel de incapsulare este Streaming Processor Array (SPA), ce contine 10 TPC. Se ajunge astfella 240 SP si 60 SFU.

Arhitectura GF 100 (Fermi)

Ultima iteratie a acestei arhitecturi este GF 100, nume de cod Fermi. Printre noutati, a fost crescutnumarul de SP la 512, memoria shared are 64 KB (fata de 16 KB la GT 200) si a fost introdus unmecanism de caching.

2015/05/08 14:50 7/15 GPU Programming


La GF100, Texture Processor Cluster a disparut, fiind inlocuit de Graphics Processing Cluster, cecontine SM-uri mult mai mari, cu 32 SP fiecare. Un GPC contine 4 SM si sunt 4 GPC in total.

Memoria shared a crescut de 4 ori, la 64 KB. Din acesti 64 KB, o parte sunt rezervati cache-ului L1. Potfi folosite configuratii cache-memorie shared 16KB:48KB sau invers, 48KB:16KB.

Exista si un L2 cache de 768KB, partajat de toate GPC.

Chipurile bazate pe arhitectura Fermi fac parte din seria GeForce 400. La data scrierii acestui text nuexista un chip care sa aiba implementata arhitectura Fermi la capacitatea sa maxima de 512 SP. Celmai mare numar de SP active, 480, este utilizat de chipul GTX 480:

Chip GTX 480 GTX 470 GTX 295 GTX 285Arhitectura GF 100 GF 100 GT 200 GT 200Stream Processors 480 448 2 x 240 240



Chip GTX 480 GTX 470 GTX 295 GTX 285Transistor Count 3B 3B 2 x 1.4B 1.4B

n prezent (2012) cele mai performante plci grafice de la Nvidia sunt bazate pe arhitectura Kepler:familia GeForce 600. Nvidia GeForce GTX 690 aparut n aprilie 2012 este dual-GPU i conine 2 x1536 SP-uri si 2128 Texture units.

Modelul de threading CUDA

Filosofia din spatele arhitecturii este permiterea rularii unui numar foarte mare de threaduri. Acestlucru este facut posibil prin paralelismul existent la nivel hardware.

Documentatia nVidia recomanda rularea unui numar cat mai mare threaduri pentru a executa un task.Arhitectura GT 200 suporta 30720 de threaduri active, iar Fermi 24576 (numarul a fost scazut pentrueficienta). Deci numarul este mult mai mare decat unitatile fizice existente pe chip. Acest lucru sedatoreaza faptului ca un numar mare de threaduri poate masca latenta accesului la memorie.

De asemenea, exista si un model de programare numit modelul Volkov. Acesta implica rularea unuinumar de threaduri egal cu numarul de theaduri suportate fizic de catre arhitectura. Rezultatul esteun model de programare SIMD. Pentru anumite probleme, folosind acest model, se obtine un spor deperformata de 5-20%.

Urmarind acelasi model modular ca si arhitectura, threadurile sunt incapsulate in blocuri (threadblocks), iar blocurile in grile (thread grid). Fiecare thread este identificat prin indexul threadului inbloc, indexul blocului in grila si indexul grilei. Indexurile threadurilor si ale blocurilor pot fiuni/bi/tri-dimensionale, iar indexul grilei poate fi uni sau bi-dimensional. Acest tip de impartire arerolul de a usura programare pentru probleme ce utilizeaza structuri de date cu mai multe dimensiuni.Pentru GT 200, numarul maxim de threaduri dintr-un bloc este de 512.

2015/05/08 14:50 9/15 GPU Programming


Threadurile dintr-un bloc pot coopera prin partajarea de date prin intermediul memoriei shared si prinsincronizarea executiei. Functia __syncthreads() are rolul de bariera pentru threadurile dintr-un bloc.Sincronizarea nu este posibila la alt nivel (intre blocuri/grila etc.).



Ierarhia de memorie

Registrilcea mai rapida forma de memorie de pe multi-procesormsunt accesibili doar de catre threadmdurata de viata este aceeasi ca si a threaduluim

Memoria partajata (shared memory)l

2015/05/08 14:50 11/15 GPU Programming


poate fi la fel de rapida ca registrii atunci cand nu exista conflicte de acces intre threadurimaccesibila tuturor threadurilor dintr-un blocmdurata de viata este aceeasi ca si a bloculuim

Memoria globalalpoate fi de 150 de ori mai lenta decat registrii sau memoria partajata, mai ales pentru citiri saumscrieri divergente (uncoalesced)accesibila atat de catre host cat si de devicemare durata de viata a aplicatieim

Memoria localalo posibila capcana; se afla de fapt in memoria globala si are deci aceleasi penalitati ca si aceastamin momentul in care compilatorul nu mai are loc in registri (si din alte motive), va stoca datelemin memoria locala in loc sa o faca in registri. Se examineaza codul .PTX pentru a determina acestlucru. Profilerul de asemenea va indica un numar mare de local store si local load.este accesibila doar de catre threadmare durata de viata a threaduluim

Memoria constantalse afla in memoria globalamread-only din device contextmcachedm

Memoria dedicata texturilorlse afla in memoria globalamread-onlymmod special de acces la datem

Performanta diferitelor tipuri de memorie

Memoria locala si memoria globala nu sunt cached, deci orice acces la acestea genereaza un accesexplicit la memorie. Este important deci de stiut care este costul pentru un astfel de acces.

Fiecare multiprocesor are nevoie de 4 cicluri de ceas pentru a transmite o instructiune de acces lamemorie unui warp. Accesul la memoria locala sau globala dureaza inca 400-600 de cicluri de ceas.

Cu o diferenta de 100-150x in timpul de acces, este evident ca trebuie redus la minim accesul lamemoria locala si reutilizarea datelor din memoria shared. De asemenea, latenta memoriei globalepoate fi ascunsa destul de mult prin specificarea unui numar mare de blocuri pentru executie sifolosirea variabilelor de tip registru, __shared__ si __constant__.

Cum accesul la memoria shared este mult mai rapid decat accesul la memoria globala, tehnicaprincipala de optimizare este evitarea de bank conflicts (warp serialization). Insa, in ciuda vitezei maria memoriei shared, imbunatatiri majore ale CUBLAS si CUFFT au fost aduse prin evitarea memorieishared in favoarea registrilor deci este recomandata folosirea acestora oricand este posibil.

Memoria shared este impartita in module de memorie identice, denumire bancuri de memorie(memory banks). Fiecare banc contine o valoare succesiva de 32 biti (de exemplu, un int sau unfloat), astfel incat accesele consecutive intr-un array provenite de la threaduri consecutive sa fiefoarte rapid. Bank conflicts au loc atunci cand se fac cereri multiple asupra datelor aflate in acelasibanc de memorie. Cand acest lucru are loc, hardware-ul serializeaza operatiile cu memoria (warp



serialization), fortand toate threadurile sa astepte pana cand toate cererile de memorie suntsatisfacute. Insa, daca toate threadurile citesc aceeasi adresa de memorie shared, este invocatautomat un mecanism de broadcast iar serializarea este evitata. Mecanismul de broadcast este foarteeficient si se recomanda folosirea sa de oricate ori este posibil.

Pentru mai multe detalii si exemple, consultati in CUDA Programming Guide, sectiunea 5.3.2.3 SharedMemory

Programare in CUDA

Dezvoltarea de cod pentru CUDA este un proces complex, desfasurat in mai multe etape:

pregatirea mediului de dezvoltare: instalare de driver, toolkit si SDK1.scrierea si testarea codului pe CPU (host)2.adaptarea codului pentru modelul CUDA prin introducerea de kernel-uri si prin copierea datelor3.intre host si devicetestare cu diferite configuratii de blocuri si numar de threaduri per bloc4.profiling si optimizare5.

Pentru mai multe detalii cu privire la acesti pasi, accesati pagina Dezvoltarea de cod CUDA

Kernel CUDA

2015/05/08 14:50 13/15 GPU Programming


Principala caracteristica a extensiei CUDA, fata de C, este ca permite utilizatorului definirea de functiispeciale, care vor rula pe GPU, denumite kernel. Identificarea sintactica a unui kernel se face pe bazanotatiei __global__ la declaratie si a notatiei cu tripla paranteza unghiulara > la apel,pentru a specifica numarul de threaduri care vor executa kernelul.

// Kernel definition__global__ void VecAdd(float* A, float* B, float* C) { ... } int main() { ... // Kernel invocation VecAdd(A, B, C); }

Executia unui kernel, solicitata de catre CPU, se face in felul urmator:

Mai multe detalii in privinta kernelurilor se gasesc in sectiunea Dezvoltarea de cod pentru CUDA

Speedup

Speedup mediu de 10x, daca un kernel poate ocupa suficiente threaduri. Se poate obtine speedup de25-400x daca accesul la date si fluxul executiei se muleaza pe modelul GPU si kernelul este optimizat.

In general, codul de kernel trebuie sa fie foarte CPU intensive, cu acces rar la memoria globala si farabranching divergent pentru threadurile din acelasi bloc. Modul de executie al instructiunilor debranching al GPU este una din principalele diferente fata de CPU. Un posibil model pentru intelegereaconstructiei unui kernel este vederea GPU-ului ca o masina SIMD multithread.

Exemple de software implementat cu success pe GPU si statistici in privinta speedup-ului obtinut pot



fi consultate pe CUDA Community Showcase

ExercitiiIn loc de Hello world:1.

Intrati pe frontend-ul fep.grid.pub.ro folosind contul de pe cs.curs.pub.ro. Executai1.comanda qlogin -q ibm-dp.q pentru a accesa una din staiile cu GPU-uri.

Exemple simple - in sectiunea Exemple de pe pagina Dezvoltarea de cod pentru CUDA, primele2.doua exemple:

(2p) pentru Dublarea valorilor dintr-un array aveti codul complet, trebuie doar compilat (hint:1.folositi nvcc) si rulat(4p) pentru Adunare de matrice aveti doar sectiunea principala a codului (main si kernel) -2.completati codul, compilati si rulati programul

Inmultirea de matrice - in sectiunea Exemple de pe pagina Dezvoltarea de cod pentru CUDA3.ultimele doua exemple:

(4p) pentru varianta fara shared memory, completati din arhiva matmul.tar fisierul1.matmul_noshared.cu(2p) pentru varianta cu shared memory, completati din arhiva matmul.tar fisierul2.matmul_shared.cuganditi-va pentru cele doua exemple: ce ar trebui modificat ca dimensiunea matricei sa nu3.trebuiasca sa fie multiplu de block size? ce ar trebui modificat pentru ca matricele sa nutrebuiasca sa fie patratice?

ResurseResponsabili laborator: Alexandra Vieru, Silvia Stegaru, Alexandru Olteanu, Grigore LupesculPDF laboratorl

Referinteworkshop CUDAlnVidia CUDAl

documentatie oficiala nVidia:mCuda Zone pe site-ul nVidianDownloads si documentatie nVidia pentru CUDA 3.2nCuda Programming Guide v3.2n

alte surse:mCurs programare pe procesoare masiv paralele si abstractnArhitectura KeplernKepler architecture whitepapernSerialul dedicat CUDA de la Dr Dobbsn

Referinte aditionale GP-GPUllaborator_4_the_democratization_of_parallel_computing.pdfm

2015/05/08 14:50 15/15 GPU Programming


laborator_4_accelerator.pdfmlaborator_4_a_survey_of_gpgpu.pdfmlaborator_4_the_graphics_card_as_a_stream_computer.pdfmlaborator_4_supercomputing_with_streams.pdfmlaborator_4_rapidmind_platform.pdfm

From:http://cs.curs.pub.ro/wiki/asc/ - ASC Wiki

Permanent link:http://cs.curs.pub.ro/wiki/asc/asc:lab10:index

Last update: 2015/05/08 14:36

GPU ProgrammingAhitectura generala GPUAlternative la arhitecturi tip GPUFramework/API pentru GPGPU - NVIDIA CUDAArhitectura NVIDIAArhitectura GT 200Arhitectura GF 100 (Fermi)Modelul de threading CUDAIerarhia de memoriePerformanta diferitelor tipuri de memorie

Programare in CUDAKernel CUDASpeedup

ExercitiiResurseReferinte

Gpu Programming

Documents

Transcript of Gpu Programming