GPU Programming

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

  1. Diferentelor din ce in ce mai mare de putere de procesare dintre CPU si GPU in favoarea acestora din urma
  2. Aparitiei de API-uri care usureaza munca programatorilor pentru a folosi GPU
  3. Existentei GPU puternice in aproape fiecare computer general-purpose

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 descurca mai bine la rularea de software care nu executa calcule in mod intensiv. Un procesor de uz general are o cantitate mare din tranzistori dedicati cachingului si controlului fluxului. GPU-uri, in constrast, isi bazeaza puterea pe executia acelorasi instructiuni, in paralel, pe cat mai multe date, iar pentru obtinerea de performante mari trebuie evitat branchingul (if-uri).

Aplicatiile existente sunt in domenii precum:

  • prelucrari video si de imagini
  • simulari de fizica in jocuri
  • domeniile clasice pentru High Performance Computing:
    • finante
    • stiinte
      • industria petrolului
      • dinamica fluidelor
    • design electronic (VLSI)

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

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

Totusi, bariere semnificative in adoptarea GPU inca exista pentru programatorii care doresc sa utilizeze puterea de calcul ieftina a acestora. Placile video sunt construite pentru jocuri video: modelul de programare este neobisnuit, resursele sunt strans limitate si arhitectura pe care se bazeaza este secreta in mare parte.

Framework-uri pentru GPGPU

nVidia CUDA si ATI Stream

Implementarea nVidia pentru GPGPU (General-Purpose computation on Graphics Processing Units) se numeste CUDA (Compute Unified Device Architecture) si permite utilizarea limbajului C pentru programarea pe GPU-urile proprii. Deoarece una din zonele tinta pentru CUDA este High Performance Computing, in care limbajul Fortran este foarte popular, PGI ofera un compilator de Fortran ce permite generarea 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)
  • CUDA toolkit:
    • compilatorul cu suport CUDA (nvcc)
    • debugger (cuda-gdb)
    • profiler (cudaprof)
    • biblioteca FFT (Fast Fourier Transform - CUFFT)
    • biblioteca BLAS (Basic Linear Algebra - CUBLAS)
    • documentatie pentru cele de mai sus
  • CUDA SDK
    • exemple de folosire CUDA (atat pentru Linux cat si pentru Windows-Visual Studio)

Pentru placile video echipate cu GPU-uri de la ATI exista ATI Stream SDK.

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

Desi majoritatea placilor video de ultima ganeratie de pe piata suporta procesari pe GPU, amandoua companiile au dezvoltat si o serie de placi grafice special concepute ca GPGPU (General Purpose GPU): 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.

Alternative

ATI si NVIDIA nu au fost singurii care s-au hotarat sa intre in piata de procesoare masiv paralele:

  • IBM propune un BladeCenter bazat pe doua procesoare Cell. Comparativ cu GeForce 8800, Cell are mai putine unitati de calcul si latime de banda a memoriei mai mica, dar ofera o frecventa mai mare si 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 de GPU , 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 de 256KB, toate insotite de un cache general de 4MB. De asemenea, va avea avantajul major de a fi bazat pe arhitectura x86. Din cauza intarzierilor de dezvoltare insa, prima generatie de chipuri Larabee nu va fi accesibila consumatorilor, ci doar pentru High Performance Computing. Larrabee arata o scalabilitate excelenta.
  • ARM a intrat pe piata procesoarelor de inalta performanta prin quad-core-ul sau Cortex-A9, folosit in data centerul Calxeda. Estimarea consumului de energie este impresionanta: fiecare nod va consuma 5W, sau 1.25W per core.

Arhitectura NVIDIA CUDA

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

Comparatie a utilizarii tranzistorilor in CPU vs. GPU

GPU sunt potrivite pentru paralelismul de date (aceleasi instructiuni sunt executate in paralel pe mai multe unitati de procesare) intensiv computationale. Datorita faptului ca acelasi program este executat 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 prin calcule in locul unor cache-uri mari pentru date.

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

Detaliile de implementare ale arhitecturii chipurilor nVidia sunt secrete. In consecinta, modul in care functioneaza 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 de programare CUDA. A fost proiectata modular astfel incat sa permita cresterea numarului de elemente computationale in generatii succesive.

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

Arhitectura unui Single Processing (SP) core

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 operatii transcendente (sin, cos) si interpolare. MT se ocupa cu trimiterea instructiunilor pentru executie la SP si SFU.

Pe langa acestea, exista si un cache (de dimensiuni reduse) pentru instructiuni, unul pentru date precum si memorie shared de 16KB, partajata de cele 8 SP.

Arhitectura unui Streaming Multiprocessor (SM) pentru GT 200

Urmatorul nivel de incapsulare este Texture / Processor Cluster (TPC). Acesta contine 3 SM, logica de control si un bloc de handling pentru texturi. Acest bloc se ocupa de modul de adresare al texturilor, logica de filtrare a acestora precum si un cache pentru texturi. Pentru GF 200, un TPC contine deci 24 SP si 6 SFU.

Arhitectura unui Texture/Processor Cluster (TPC) pentru GT 200

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

Nucleul de procesare al chipului GT 200: Streaming Processor Array (SPA)

Arhitectura GF 100 (Fermi)

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

Fermi - 3 miliarde de tranzistori intr-un singur chip

La GF100, Texture Processor Cluster a disparut, fiind inlocuit de Graphics Processing Cluster, ce contine 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. Pot fi folosite configuratii cache-memorie shared 16KB:48KB sau invers, 48KB:16KB.

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

Arhitectura unui Streaming Multiprocessor (SM) pentru GF 100

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

Chip GTX 480 GTX 470 GTX 295 GTX 285
Arhitectura GF 100 GF 100 GT 200 GT 200
Stream Processors 480 448 2 x 240 240
Transistor Count 3B 3B 2 x 1.4B 1.4B

În prezent (2012) cele mai performante plăci grafice de la Nvidia sunt bazate pe arhitectura Kepler: familia GeForce 600. Nvidia GeForce GTX 690 aparută în aprilie 2012 este dual-GPU și conține 2 x 1536 SP-uri si 2×128 Texture units.

Modelul de threading CUDA

Filosofia din spatele arhitecturii este permiterea rularii unui numar foarte mare de threaduri. Acest lucru 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 pentru eficienta). Deci numarul este mult mai mare decat unitatile fizice exitente pe chip. Acest lucru se datoreaza 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 unui numar de threaduri egal cu numarul de theaduri suportate fizic de catre arhitectura. Rezultatul este un model de programare SIMD. Pentru anumite probleme, folosind acest model, se obtine un spor de performata de 5-20%.

Urmarind acelasi model modular ca si arhitectura, threadurile sunt incapsulate in blocuri (thread blocks), iar blocurile in grile (thread grid). Fiecare thread este identificat prin indexul threadului in bloc, indexul blocului in grila si indexul grilei. Indexurile threadurilor si ale blocurilor pot fi uni/bi/tri-dimensionale, iar indexul grilei poate fi uni sau bi-dimensional. Acest tip de impartire are rolul 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.

''Structura threadurilor in blocuri''

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

''Modul de indexare al threadurilor intr-un bloc'' ''Utilizarea indexului''

Tipuri de memorie

Ierarhia memoriei in CUDA

  • Registri
    • cea mai rapida forma de memorie de pe multi-procesor
    • sunt accesibili doar de catre thread
    • durata de viata este aceeasi ca si a threadului
  • Memoria partajata (shared memory)
    • poate fi la fel de rapida ca registrii atunci cand nu exista conflicte de acces intre threaduri
    • accesibila tuturor threadurilor dintr-un bloc
    • durata de viata este aceeasi ca si a blocului
  • Memoria globala
    • poate fi de 150 de ori mai lenta decat registrii sau memoria partajata, mai ales pentru citiri sau scrieri divergente (uncoalesced)
    • accesibila atat de catre host cat si de device
    • are durata de viata a aplicatiei
  • Memoria locala
    • o posibila capcana; se afla de fapt in memoria globala si are deci aceleasi penalitati ca si aceasta
    • in momentul in care compilatorul nu mai are loc in registri (si din alte motive…), va stoca datele in memoria locala in loc sa o faca in registri. Se examineaza codul .PTX pentru a determina acest lucru. Profilerul de asemenea va indica un numar mare de „local store” si „local load”.
    • este accesibila doar de catre thread
    • are durata de viata a threadului
  • Memoria constanta
    • se afla in memoria globala
    • read-only din device context
    • cached
  • Memoria dedicata texturilor
    • se afla in memoria globala
    • read-only
    • mod special de acces la date

Performanta diferitelor tipuri de memorie

Memoria locala si memoria globala nu sunt cached, deci orice acces la acestea genereaza un acces explicit 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 la memorie 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 la memoria locala si reutilizarea datelor din memoria shared. De asemenea, latenta memoriei globale poate fi ascunsa destul de mult prin specificarea unui numar mare de blocuri pentru executie si folosirea variabilelor de tip registru, __shared__ si __constant__.

Cum accesul la memoria shared este mult mai rapid decat accesul la memoria globala, tehnica principala de optimizare este evitarea de bank conflicts (warp serialization). Insa, in ciuda vitezei mari a memoriei shared, imbunatatiri majore ale CUBLAS si CUFFT au fost aduse prin evitarea memoriei shared 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 un float), astfel incat accesele consecutive intr-un array provenite de la threaduri consecutive sa fie foarte rapid. Bank conflicts au loc atunci cand se fac cereri multiple asupra datelor aflate in acelasi banc 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 sunt satisfacute. Insa, daca toate threadurile citesc aceeasi adresa de memorie shared, este invocat automat un mecanism de broadcast iar serializarea este evitata. Mecanismul de broadcast este foarte eficient 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 Shared Memory

Programare in CUDA

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

  1. pregatirea mediului de dezvoltare: instalare de driver, toolkit si SDK
  2. scrierea si testarea codului pe CPU (host)
  3. adaptarea codului pentru modelul CUDA prin introducerea de kernel-uri si prin copierea datelor intre host si device
  4. testare cu diferite configuratii de blocuri si numar de threaduri per bloc
  5. profiling si optimizare

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

Kernel CUDA

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

  1. // Kernel definition
  2. __global__ void VecAdd(float* A, float* B, float* C)
  3. {
  4. ...
  5. }
  6.  
  7. int main()
  8. {
  9. ...
  10. // Kernel invocation
  11. VecAdd<<<1, N>>>(A, B, C);
  12. }

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 de 25-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 fara branching divergent pentru threadurile din acelasi bloc. Modul de executie al instructiunilor de branching al GPU este una din principalele diferente fata de CPU. Un posibil model pentru intelegerea constructiei 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

Activitate practica

  1. In loc de Hello world:
    1. Intrati pe masina din cluster cu capabilitati pentru GPU Computing:
      1. fiecare student intra pe fep.grid.pub.ro cu ssh -X folosind contul de pe cs
      2. fiecare student ruleaza comanda: apprun.sh xterm fs-dual.q
      3. acum fiecare student va avea o consola/terminal pe masina dual-wn21.grid.pub.ro pe care este montanta placa video
    2. Copiati aplicatia deviceQuery din SDK (/opt/NVIDIA_GPU_Computing_SDK/C/bin/linux/release/) in Home-ul vostru si observati caracteristicile placii video, in special compute capability.
  2. Exemple simple - in sectiunea Exemple de pe pagina Dezvoltarea de cod pentru CUDA, primele doua exemple:
    1. pentru “Dublarea valorilor dintr-un array” aveti codul complet, trebuie doar compilat (hint: folositi nvcc) si rulat
    2. pentru “Adunare de matrice” aveti doar sectiunea principala a codului (main si kernel) - completati codul, compilati si rulati programul
  3. Inmultirea de matrice - in sectiunea Exemple de pe pagina Dezvoltarea de cod pentru CUDA ultimele doua exemple:
    1. pentru varianta fara shared memory, completati din arhiva matmul.tar fisierul matmul_noshared.cu
    2. pentru varianta cu shared memory, completati din arhiva matmul.tar fisierul matmul_shared.cu
    3. ganditi-va pentru cele doua exemple: ce ar trebui modificat ca dimensiunea matricei sa nu trebuiasca sa fie multiplu de block size? ce ar trebui modificat pentru ca matricele sa nu trebuiasca sa fie patratice?

Resurse

asc/lab12/index.txt · Last modified: 2013/02/07 12:41 (external edit)
CC Attribution-Share Alike 3.0 Unported
www.chimeric.de Valid CSS Driven by DokuWiki do yourself a favour and use a real browser - get firefox!! Recent changes RSS feed Valid XHTML 1.0