Show page

Differences

This shows you the differences between two versions of the page.

--- asc:laboratoare:05 [2022/04/04 15:26]
emil.slusanschi [Exercitii]
+++ asc:laboratoare:05 [2024/04/08 13:48] (current)
emil.slusanschi [Referinte]
@@ Line 1: / Line 1: @@
-====== Laboratorul 05 - Tehnici de Optimizare de Cod – Inmultirea Matricelor ======
+====== Laboratorul 05 - Arhitectura GPU NVIDIA CUDA ======
-===== Obiective =====
+Pentru o înțelegere profundă a arhitecturii CUDA vă recomandăm să citiți din resursele bilbiografie "Graphics and Computing GPUs". În continuare o să discutăm mai aplicat despre implementări ale acestei arhitecturi.
-In acest laborator vom exemplifica o serie de optimizari de cod pe una dintre cele mai simple, si in acelasi timp utilizate probleme, si anume, inmultirea matricelor.
+Arhitectura NVIDIA FERMI [[https://www.nvidia.com/content/PDF/fermi_white_papers/NVIDIA_Fermi_Compute_Architecture_Whitepaper.pdf|aici]], Tesla 2070, coada executie fep8.grid.pub.ro -> dp
-==== De ce inmultirea matricelor? ====
+Arhitectura NVIDIA KEPLER [[https://www.nvidia.com/content/PDF/kepler/NVIDIA-Kepler-GK110-Architecture-Whitepaper.pdf|aici]], Tesla K40M, coada executie fep8.grid.pub.ro -> hpsl
-Este o operatie fundamentala si elementara in algebra liniara ce serveste la rezolvarea unui numar extrem de mare de probleme, cum ar fi: rezolvarea sistemelor liniare de ecuatii in majoritatea domeniilor stiintifice si economice (operatiile cu matrice sunt practic prezente pretudindeni); calcule si operatii cu grafuri; inversari de matrice. Problema inmultirii matricelor este in mod cert cea mai bine studiata problema in HPC (High Performance Computing), ea beneficiind de o multitudine de algoritmi inteligenti si implementari performante pe toate arhitecturile existente astazi. Pentru a simplifica lucrurile, in acest laborator ne vom ocupa doar de inmultirea matricelor patratice.
+Prima arhitectura NVIDIA complet programabila a fost G80 (ex. [[http://www.nvidia.com/page/8800_tech_briefs.html|Geforce 8800]], lansat in anul 2006). Cu aceasta arhitectura s-a trecut de la unitati hardware fixe vertex/pixel la cele de unified shader care puteau procesa atat vertex/pixel cat si geometry. Evolutia arhitecturilor GPU de la NVIDIA este detaliata [[http://s08.idav.ucdavis.edu/luebke-nvidia-gpu-architecture.pdf|aici]].
-===== Cel mai simplu algoritm =====
+Implementarea NVIDIA pentru GPGPU se numeste CUDA (Compute Unified Device Architecture) si permite utilizarea limbajului C pentru programarea pe GPU-urile proprii. Lista de GPU-uri ce suporta API-ul CUDA sau OpenCL se regaseste pe site-ul oficial [[https://www.geforce.com/hardware/technology/cuda/supported-gpus|aici]] sau pe wiki [[https://en.wikipedia.org/wiki/List_of_Nvidia_graphics_processing_units|aici]]. Fiecare noua arhitectura are un codename (ex Fermi, Pascal) si este reprezentata de un "compute capability" (list [[https://developer.nvidia.com/cuda-gpus|aici]]). Cu cat arhitectura este mai noua, cu atat sunt suportate mai multe facilitati din API-urile CUDA si OpenCL.
-Intuitiv, cel mai simplu algoritm, urmeaza formularea matematica:
+Unitatea GPU este potrivita pentru paralelismul de date SIMD (Single Instruction Multiple Data), astfel aceleasi instructiuni sunt executate in paralel pe mai multe unitati de procesare. 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.
-{{:asc:lab5:cij.jpg|}}
-Matricele A = [aij], i,j=1,...,N  si B = [bij], i,j=1,...,N sunt salvate ca vectori bidimensionali de marime N x N. Matricea rezultat C = A x B = [cij], i,j=1,...,N, avand fireste aceeasi dimensiune.
-{{:asc:lab5:axb_c.jpg|}}
+Motivul discrepantei intre performanta paralela 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.
-Cum este si de asteptat, similar cu majoritatea operatiilor din algebra liniara, formula de mai sus se transforma in urmatorul program extrem de simplu:
+La GPU-urile NVIDIA, un Streaming Processor (SP) este un microprocesor cu executie secventiala, ce contine un pipeline, unitati aritmetico-logice (ALU) si 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.
-<code cpp>
+SP impreuna cu Special Function Units (SFU) sunt incapsulate intr-un Streaming Multiprocessor (SM/SMX). Fiecare SFU contine 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.
-int i,j,k;
-double a[N][N], b[N][N], c[N][N];
-// initializarea matricelor a si b
-for (i=0;i<N;i++){
-   for (j=0;j<N;j++){
-      c[i][j] = 0.0;
-      for (k=0;k<N;k++){
-	c[i][j] += a[i][k] * b[k][j];
-      }
-   }
-}
-</code>
-<note important>
+Pe langa acestea, exista si un cache (de dimensiuni reduse) pentru instructiuni, unul pentru date precum si memorie shared, partajata de SP-uri. Urmatorul nivel de incapsulare este Texture / Processor Cluster (TPC). Acesta contine SM-uri, 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.
-**Cat de bun este acest algoritm?**
+{{:asc:lab11:cuda-arch.png?direct&720|}}
-Algoritmul este bun pentru ca:
+Filosofia din spatele arhitecturii este permiterea rularii unui numar foarte mare de threaduri. Acest lucru este facut posibil prin paralelismul existent la nivel hardware.
-  *	Se poate specifica in doar cateva linii;
-  *	Este o mapare directa a formulei de calcul pentru Cij (din algebra liniara); este usor de inteles si de urmarit de catre oricine poseda cunostinte minime de matematica;
-  *	In sfarsit, in mod sigur nu contine bug-uri datorita simplitatii extreme pe care o manifesta algoritmul!
-Algoritmul este prost pentru ca:
+Documentatia NVIDIA recomanda rularea unui numar cat mai mare threaduri pentru a executa un task. Arhitectura CUDA de exemplu suporta zeci de mii de threaduri, numarul acestora fiind mult mai mare decat unitatile fizice existente pe chip. Acest lucru se datoreaza faptului ca un numar mare de threaduri poate masca latenta accesului la memorie.
-  *	Are performante extrem de reduse!
-De aceea ne vom ocupa in acest laborator de optimizarea acestei operatii din punctul de vedere al performantei.
+Urmarind acelasi model modular din punct de vedere al arhitecturii, threadurile sunt incapsulate in blocuri (thread blocks / warps), 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 programarea pentru probleme ce utilizeaza structuri de date cu mai multe dimensiuni. Se poate observa ca thread-urile dintr-un thread block trebuie sa execute cat mai multe instructiuni identice spre a nu irosi resurse.
-</note>
+{{:asc:lab11:thread.blocks.jpg?direct&360|{{thread.blocks.jpg|''Structura threadurilor in blocuri''}}
-===== Optimizarea algoritmului de inmultire a doua matrice =====
+Threadurile dintr-un bloc pot coopera prin partajarea de date prin intermediul memoriei shared si prin sincronizarea executiei. Functia de bariera functioneaza doar pentru threadurile dintr-un bloc. Sincronizarea nu este posibila la alt nivel (intre blocuri/grila etc.). Mai multe explicatii se regasesc in [[https://ocw.cs.pub.ro/courses/asc/laboratoare/06|Laboratorul 6]].
-==== Detectarea constantelor din bucle ====
+===== Ierarhia de memorie =====
-Prima optimizare, consta in a observa ca c[i][j] este o constanta in cadrul ciclului interior k. Totusi, pentru un compilator acest fapt nu este neaparat evident deoarece c[i][j] este o referinta in cadrul unui vector. Astfel, o prima optimizare va arata asa:
+Intelegerea ierharhiei de memorie este esentiala in programarea eficienta a unitatii GPU. Capacitatea mare de executie in paralel a unui GPU necesita ascunderea latentei de acces catre memoria principala (fie VRAM pentru dGPU sau RAM pentru iGPU).
-<code cpp>
+{{:asc:lab11:mem.hierarchy.png?direct|Ierarhia memoriei in CUDA}}
-for (i=0;i<N;i++){
-   for (j=0;j<N;j++){
+**Register File**
-      register double suma = 0.0;
+<code sh>
-      for (k=0;k<N;k++) {
+/* marcam pentru compilator regValPi in register file */
-         suma += a[i][k] * b[k][j];
+__private float regValPi = 3.14f;
-      }
+/* compilatorul cel mai probabil oricum incadreaza regVal2Pi ca registru */
-      c[i][j] = suma;
+float regVal2Pi = 2 * 3.14f;
-   }
-}
 </code>
+  *Cea mai rapida forma de memorie de pe GPU
+  *Accesibila doar de catre thread, iar durata de viata este aceeasi ca si a threadului
+  *Un kernel complex poate determina folosirea unui numar mare de registrii si astfel:
+    * limitarea executiei multor thread-uri simultan
+    * register spill, atunci cand valorile registrilor sunt salvate in memoria globala
-In acest mod, compilatorul va putea avea grija ca variabila suma sa fie tinut intr-un registru, permitand astfel o utilizare optima a acestei resurse. Astfel utilizarea keyword-ului "register" este util de folosit ca hint pentru compilator, atunci cand socotiti ca acest lucru este util.
+**Local Memory**
+<code sh>
+/* fiecare work item salveaza un element */
+__local float lArray[lid] = data[gid];
+</code>
+  *In functie de implementarea hardware, 100GB/sec -> 2TB/sec
+  *Pentru GPU o memorie rapida, actioneaza ca un cache L1/alt register file, la CPU de regula este doar o portiune din RAM
+  *Ca si in cazul registrilor, este accesibila doar de catre thread, iar durata de viata este aceeasi ca si a threadului
-==== Accesul la vectori ====
+**Shared Memory**
+<code sh>
-Un alt aspect care necesita resurse din plin, este utilizarea si accesul variabilelor de tip vectorial. De fiecare data cand programul face o referinta la un obiect de tipul X[i][j][k] compilatorul trebuie sa genereze expresii aritmetice complexe, pentru a calcula aceasta adresa, in cadrul vectorului muldimensional X. De exemplu, iata cum arata un vector bidimensional in limbajul C (salvat row-major):
+/* elementele sunt salvate la nivel de bloc */
+__shared__ int s[64];
-{{:asc:lab5:aij.jpg|}}
-Astfel, pentru N = 6, M = 4: a[2][3] = a[0][0] + 2*6 + 3 = a[0][0] + 15
-In limbaje de programare ca FORTRAN-ul, formula este inversata, deoarece aceste limbaje salvează vectorii în format column-major:
-a[i][j] = a[0][0] + j*M + i
-Oricare ar fi asezarea vectorilor in memorie, accesele la vectori sunt scumpe din punctul de vedere al performantelor. Noi vom considera de aici inainte o asezare row-major, ca in limbajul C. Conform acestei formule, pentru vectori bidimensionali (matrice), fiecare acces presupune doua adunari si o inmultire (de numere intregi). Evident, pentru vectori cu mai multe dimensiuni, aceste costuri cresc considerabil. Astfel, in momentul in care compilatorul intalneste instructiunea:
-suma += a[i][k] * b[k][j]
-se vor efectua implicit, suplimentar inmultirii si adunarii in virgula mobila implicata de codul de mai sus, patru adunari si doua inmultiri in numere intregi pentru a calcula adresele necesare din vectorii a si b. Se intampla astfel destul de frecvent ca procesorul sa nu aiba date disponibile pentru a lucra in continuu, din cauza faptului ca overhead-ul pentru calculul adreselor este semnificativ.
-Astfel, un mod de a spori viteza programului este renuntarea la accesele vectoriale prin derefentiere utilizand in acest scop pointeri. De exemplu:
-<code cpp>
-for (j=0;j<N;j++)
-   a[i][j] = 2;         // 2*N adunari si N inmultiri
 </code>
+  *Accesibila tuturor threadurilor dintr-un bloc (warp/wavefront), iar durata de viata este aceeasi ca si a blocului
+  *Trebuie evitate conflictele de access (bank conflicts)
-se va inlocui cu:
+**Constant Memory**
+<code sh>
-<code cpp>
+__const float pi = 3.14f
-double *ptr=&(a[i][0]); // 2 adunari si o inmultire
-for (j=0;j<N;j++) {
-   *ptr = 2;
-   ptr++;               // N adunari in numere intregi
-}
 </code>
+  * In functie de implementarea hardware, 100GB/sec -> 1TB/sec
+  * In general performanta foarte buna, (cache L1/L2, zona dedicata),
+  * Are durata de viata a aplicatiei kernel
-In mod similar se procedeaza si pentru cazul in care indexul incrementat este cel al liniilor si nu cel al coloanelor. In ambele cazuri, practic se va calcula "de mana" adresa in cadrul vectorului, exact in modul in care ar face-o compilatorul limbajului folosit. Totusi, rezolvarea noastra este mai rapida, deoarece ea tine cont de pozitia in care ne aflam in cadrul vectorului, lucru destul de complicat de facut automat. De exemplu, pentru a trece la urmatoarea coloana, e suficient sa adunam N pointer-ului, fata de recalcularea pornind de la @(a[0][0]) ce necesita doua inmultiri si o adunare in intregi. Evident, facilitatile oferite de limbaje ca C-ul, ne vin in ajutor: astfel incrementarile de pointeri de tip char * vor face incrementarea cu un byte, in vreme ce pentru int * se va face cu patru bytes. Ca urmare a aspectelor prezentate mai sus, iata forma optimizata in care ajunge algoritmul nostru:
+**Global Memory**
+<code sh>
-<code cpp>
+__kernel void process(__global float* data){ ... }
-for(i = 0; i < N; i++){
-  double *orig_pa = &a[i][0];
-  for(j = 0; j < N; j++){
-    double *pa = orig_pa;
-    double *pb = &b[0][j];
-    register double suma = 0;
-    for(k = 0; k < N; k++){
-      suma += *pa * *pb;
-      pa++;
-      pb += N;
-    }
-    c[i][j] = suma;
-  }
-}
 </code>
+  * In functie de implementarea hardware, 30GB/sec -> 500GB/sec
+  * Video RAM (VRAM), de regula cu o capacitate intre 1GB si 12GB in functie de placa video
+  * Memorie dedicata specializata doar pentru placile grafice discrete (GPU-urile integrate in CPU folosesc RAM)
+  * In general latime mare de banda (256-512 biti) si chipuri de memorii de mare viteza (GDDR5)
-<note tip>Atentie! Codul de mai sus va da rezultate corecte doar daca matricile sunt declarate global sau pe stivă pentru că în felul acesta sunt stocate continuu în memorie (și are sens pb += N). Dacă alocați dinamic, atunci folosiți matrici liniarizate și adaptați acest cod pentru cazul lor.</note>
+**Host Memory (RAM)**
+  * In general, 4GB/sec -> 30GB/sec
+  * Pentru acces din kernel trebuie transfer/mapare explicita RAM->VRAM pe partea de host/CPU
+  * Memoria RAM accesibila direct de CPU si indirect de GPU via DMA si magistrala PCIe
+  * Viteza de transfer (throughput/latenta) este limitata de magistrala PCIe cat si de memoria RAM
-<note important>
+Caracteristici GPU K40m (coada hpsl), via query device properties CUDA
-Din primele doua optimizari se pot desprinde cateva concluzii. Prima ar fi ca optimizarea unui cod (din punct de vedere al performantelor), presupune utilizarea a cat mai putine constructii complexe (high-level), puse la dispozitie de limbajul folosit. Aceasta concluzie poate suna extrem de ciudat pentru cineva care porneste de la ideea ca facilitatile limbajelor de programare sunt acolo pentru a fi folosite. Da, este adevarat acest lucru, insa atunci cand vrei performanta, trebuie sa stii ce constructii sa eviti! Astfel, apare concluzia a doua: vectorii sunt concepte mai abstracte decat pointerii (ca implementare), asadar, utilizati pointeri cand vreti viteza. Viteza crescuta insa, va fi obtinuta cu pretul unui cod mult mai dificil de urmarit si de inteles, mai rau, mult mai greu de debug-at. Un cod complex si performant, de multe ori poate contine bug-uri extrem de subtile si greu de depistat. Asadar, e util sa stii exact ceea ce faci cand incepi sa faci astfel de optimizari!
-</note>
+<code sh>
+Device 0: "Tesla K40m"
-----
+  CUDA Driver Version / Runtime Version          9.1 / 9.1
+  CUDA Capability Major/Minor version number:    3.5
-==== Activitate practica - Optimizare constantelor si al accesului la vectori ====
+  Total amount of global memory:                 11441 MBytes (11996954624 bytes)
+  (15) Multiprocessors, (192) CUDA Cores/MP:     2880 CUDA Cores
-Intrebarea este acum: aduc ceva imbunatatiri optimizarile 1 si 2? Pentru a afla raspunsul la aceasta intrebare, va invitam sa implementati problema, cu optimizarile sugerate, si sa observati singuri ce se intampla.
+  GPU Max Clock rate:                            745 MHz (0.75 GHz)
+  Memory Clock rate:                             3004 Mhz
-----
+  Memory Bus Width:                              384-bit
+  L2 Cache Size:                                 1572864 bytes
-==== Optimizarea pentru accesul la memorie ====
+  Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
+  Maximum Layered 1D Texture Size, (num) layers  1D=(16384), 2048 layers
-Dupa cum ar trebui sa va fie destul de evident pana acum, din experienta voastra de programatori, memoria este in general cel mai problematic bottleneck. Optimizarile prezentate mai sus reduc timpul de executie intr-o oarecare masura, insa ele nu schimba in nici un fel modul in care memoria este accesata in cadrul algoritmului. Cu alte cuvinte, aceleasi locatii de memorie sunt accesate in aceeasi ordine, indiferent daca am operat sau nu optimizarile prezentate. O intrebare interesanta ar fi acum: ce se intampla, daca am schimba ordinea in care se executa buclele? S-ar obtine performante diferite?
+  Maximum Layered 2D Texture Size, (num) layers  2D=(16384, 16384), 2048 layers
-Pentru problema noastra, care contine trei bucle, exista asadar sase secvente posibile, si anume: i-j-k, i-k-j, j-i-k, j-k-i, k-i-j, si k-j-i. Fiecare dintre aceste secvente corespunde unui tip diferit de acces la memorie pentru matricele considerate. Deoarece bucla interioara este cea mai des executata, ne vom concentra acum atentia un pic asupra ei. Operatia executata acolo ramane:
+  Total amount of constant memory:               65536 bytes
+  Total amount of shared memory per block:       49152 bytes
-c[i][j] += a[i][k] * b[k][j]
+  Total number of registers available per block: 65536
+  Warp size:                                     32
-Pentru fiecare dintre cele trei matrice, a, b si c, fiecare element poate fi accesat in trei moduri diferite, si anume:
+  Maximum number of threads per multiprocessor:  2048
-  * Constant: accesul nu depinde de indexul buclei interioare
+  Maximum number of threads per block:           1024
-  * Secvential: aceesul la memorie este contiguu (adica in celule succesive de memorie)
+  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
-  * Nesecvential: accesul la memorie nu este contiguu (celulele de memorie logic succesive, sunt de fapt adresate cu pauze de dimensiune N)
+  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
+  Maximum memory pitch:                          2147483647 bytes
-Astfel, pentru cele sase configuratii, se obtine:
+  Texture alignment:                             512 bytes
+  Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
-^ Loop order ^ c[i][j] += ^ a[i][k] ^ * b[k][j] ^
+  Run time limit on kernels:                     No
-| i-j-k: | Constant   | Secvential | Nesecvential |
+  Integrated GPU sharing Host Memory:            No
-| i-k-j: | Secvential | Constant | Secvential |
+  Support host page-locked memory mapping:       Yes
-| j-i-k: | Constant | Secvential | Nesecvential |
+  Alignment requirement for Surfaces:            Yes
-| j-k-i: | Nesecvential | Nesecvential | Constant |
+  Device has ECC support:                        Enabled
-| k-i-j: | Secvential | Constant | Secvential |
+  Device supports Unified Addressing (UVA):      Yes
-| k-j-i: | Nesecvential | Nesecvential | Constant |
+  Device PCI Domain ID / Bus ID / location ID:   0 / 8 / 0
-Care sunt totusi, comparativ, performantele celor trei moduri de acces? In mod clar, accesul constant este mai bun decat cel secvential – aceste constante in cadrul unor bucle, sunt in general puse in registri, ducand la imbunatatirea performantelor  algoritmului, dupa cum s-a aratat in optimizarea 1. Accesul secvential la randul sau, este mai bun decat cel nesecvential, in principal pentru ca utilizeaza considerabil mai bine cache-ul.
-<note important>
-Luand in considerare aceste observatii, putem concluziona ca:
-  * Configuratiile k-i-j si i-k-j ar trebui sa aiba cele mai bune performante
-  * Configuratiile i-j-k si j-i-k ar trebui sa fie mai proaste decat primele, si
-  * Configuratiile j-k-i si k-j-i ar trebui sa fie cele mai proaste!
-</note>
-==== Activitate practica - Ordinea buclelor ====
-----
-Efectiv, care este adevarul? Construiti singuri aceste scenarii si analizati aceasta problema!
-----
-Pentru a studia mai in detaliu problema, sa analizam un pic configuratia i-j-k (desi nu este cea mai buna configuratie, cum vedem de mai sus):
-<code cpp>
-for (i=0;i<N;i++){
-   for (j=0;j<N;j++){
-      sum=0;
-      for (k=0;k<N;k++)
-         sum+=a[i][k]*b[k][j];
-      c[i][j] = sum;
-   }
-}
 </code>
-Cate cache-miss-uri sunt generate in acest algoritm, cu aceasta secventa de acces la memorie? In mod evident, aceasta nu este o intrebare usoara. De exemplu: daca fiecare matrice ar fi de doua ori mai mare decat cache-ul, ar avea loc multe incarcari si eliberari de linii, ducand astfel la o formula complicata. Astfel, cel mai simplu aproximam, si consideram ca dimensiunea matricei este mult mai mare decat cea a Cache-ului. Astfel, fie C, numarul de elemente din matrice ce intra in Cache.
+Caracteristici GPU M2070 (coada dp), via query device properties CUDA
-Astfel, considerand algoritmul de mai sus (fara optimizarea pentru constante):
-<code cpp>
+<code sh>
-for (i=0;i<N;i++){
+Device 0: "Tesla M2070"
-   // Citeste linia i pt a in Cache (Ra)
+  CUDA Driver Version / Runtime Version          9.1 / 9.1
-   // Scrie linia i a lui c in Memorie (Wc)
+  CUDA Capability Major/Minor version number:    2.0
-   for (j=0;j<N;j++){
+  Total amount of global memory:                 5302 MBytes (5559156736 bytes)
-      // Citeste coloana j a lui b in Cache (Rb)
+  (14) Multiprocessors, ( 32) CUDA Cores/MP:     448 CUDA Cores
-      for (k=0;k<N;k++){
+  GPU Max Clock rate:                            1147 MHz (1.15 GHz)
-         c[i][j] += a[i][k] * b[k][j];
+  Memory Clock rate:                             1566 Mhz
-      }
+  Memory Bus Width:                              384-bit
-   }
+  L2 Cache Size:                                 786432 bytes
-}
+  Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65535), 3D=(2048, 2048, 2048)
+  Maximum Layered 1D Texture Size, (num) layers  1D=(16384), 2048 layers
+  Maximum Layered 2D Texture Size, (num) layers  2D=(16384, 16384), 2048 layers
+  Total amount of constant memory:               65536 bytes
+  Total amount of shared memory per block:       49152 bytes
+  Total number of registers available per block: 32768
+  Warp size:                                     32
+  Maximum number of threads per multiprocessor:  1536
+  Maximum number of threads per block:           1024
+  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
+  Max dimension size of a grid size    (x,y,z): (65535, 65535, 65535)
+  Maximum memory pitch:                          2147483647 bytes
+  Texture alignment:                             512 bytes
+  Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
+  Run time limit on kernels:                     No
+  Integrated GPU sharing Host Memory:            No
+  Support host page-locked memory mapping:       Yes
+  Alignment requirement for Surfaces:            Yes
+  Device has ECC support:                        Enabled
+  Device supports Unified Addressing (UVA):      Yes
+  Device PCI Domain ID / Bus ID / location ID:   0 / 20 / 0
 </code>
-Astfel, daca L este dimensiunea unei linii de Cache: pentru (Ra) obtinem aproximativ N*(N/L) cache-miss-uri, pentru (Wc) la fel, iar pentru (Rb) un dezastruos N*N*N! Acest lucru se intampla deoarece, desi accesul la b este secvential intr-o coloana, matricea este salvata in memorie utilizand row-major! Concluzia este descurajatoare: 2N<sup>2</sup>/L + N<sup>3</sup> -> N<sup>3</sup> cache-miss-uri! Se adauga la acest aspect si cele 2N<sup>3</sup> operatii aritmetice, si se ajunge la raportul: operatii aritmetice / operatii cu memoria -> 2. Acest lucru este extrem de rau, deoarece noi stim de la (curs) si de la alte materii, ca arhitecturile calculatoarelor NU sunt echilibrate, si ca operatiile aritmetice sunt de ordine de marime mai rapide decat operatiile cu memoria. De aceea, memoria ramane in continuare bottleneck-ul pentru aceasta implementare a inmultirii de matrice. Pentru a obtine performante mai bune, este necesara obtinerea unui raport considerabil mai mare.
+Caracteristici GPU P100 (coada xl), via query device properties CUDA
-Cum se face insa, ca pentru N<sup>2</sup> elemente intr-o matrice, ajungem la N<sup>3</sup> cache-miss-uri? Pai am stabilit ca acest lucru se datoreaza accesului ineficient al lui b, deoarece se incearca incarcarea coloana cu coloana a matricei!
+<code sh>
+Device 1: "Tesla P100-PCIE-16GB"
-Concluzia acestei analize este ca nu putem spune, doar dupa numarul de operatii efectuate si dimensiunea datelor folosite, daca un algoritm va suferi sau nu din cauza unui bottleneck la memorie.
+  CUDA Driver Version / Runtime Version          12.2 / 11.4
+  CUDA Capability Major/Minor version number:    6.0
-Solutia este: utilizarea mai ingenioasa a cache-ului.
+  Total amount of global memory:                 16276 MBytes (17066885120 bytes)
+  (056) Multiprocessors, (064) CUDA Cores/MP:    3584 CUDA Cores
-Acest lucru se poate realiza prin reorganizarea operatiilor din cadrul inmultirii de matrice pentru a obtine mai multe cache-hit-uri. Faptul ca adunarea si inmultirea sunt atat operatii asociative, cat si comutative face posibila aceasta reordonare a operatiilor. Acesta este un subiect de cercetare asupra caruia si-au indreptat atentia numerosi cercetatori de-a lungul timpului, generand o multitudine de algoritmi si de teoreme matematice care sa ii sustina. In orice caz, daca vom considera r = raportul intre operatiile aritmetice si operatiile  la memorie (cu cache-miss-uri), este evident ca se doreste un r maxim, pentru a elimina bottleneck-ul de la memorie. S-a aratat ca orice reorganizare a acestui algoritm este limitata la r = O(sqrt(C)), unde C este dimensiunea Cache-ului (in numar de elemente ce intra in Cache). Acest lucru arata ca r nu scaleaza cu dimensiunea matricei N, indiferent de impartirea intuitiva a lui 2N<sup>3</sup> la N<sup>2</sup>...
+  GPU Max Clock rate:                            1329 MHz (1.33 GHz)
+  Memory Clock rate:                             715 Mhz
-==== Solutia: “Blocked Matrix Multiplication” ====
+  Memory Bus Width:                              4096-bit
+  L2 Cache Size:                                 4194304 bytes
-Pentru a rezolva problema accesului in b pentru coloane intregi, se va trece la accesarea unui subset a unei coloane in b, sau a mai multor coloane la un moment dat. Pentru o mai buna intelegere, urmariti desenele de mai jos:
+  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
+  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
-{{:asc:lab5:c_axb_1.jpg|}}
+  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
+  Total amount of constant memory:               65536 bytes
-Ideea de baza este refolosirea cat mai buna a elementelor aflate in cache (pentru matricea b). Astfel odata cu calculul lui c[i][j], de ce nu am calcula si c[i][j+1], daca tot se afla in cache si coloana j+1. Acest lucru presupune insa reordonarea operatiilor astfel: calculeaza primii b termeni pentru c[i][j], calculeaza primii b termeni pentru c[i][j+1], calculeaza urmatorii b termeni pentru c[i][j], calculeaza urmatorii b termeni pentru c[i][j+1], etc.
+  Total amount of shared memory per block:       49152 bytes
+  Total shared memory per multiprocessor:        65536 bytes
-{{:asc:lab5:c_axb_2.jpg|}}
+  Total number of registers available per block: 65536
+  Warp size:                                     32
-In acest mod, de ce nu am calcula o intreaga sectiune de linie din c, folosind aceste reordonari de operatii?
+  Maximum number of threads per multiprocessor:  2048
+  Maximum number of threads per block:           1024
-Ce s-ar intampla daca am incerca sa calculam o intreaga linie din c?
+  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
+  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
-{{:asc:lab5:c_axb_3.jpg|}}
+  Maximum memory pitch:                          2147483647 bytes
+  Texture alignment:                             512 bytes
-Ar insemna ca trebuie sa incarcam toate coloanele lui b in memorie (cache), lucru pe care am incercat sa il evitam aici! Astfel, se vor refolosi doar acele blocuri din b ce au fost deja incarcate. De aici nu ne mai ramane decat sa utilizam intreaga linie de cache din b, si obtinem ideea de baza a algoritmului “Blocked Matrix Multiplication”:
+  Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
+  Run time limit on kernels:                     No
-{{:asc:lab5:c_axb_4.jpg|}}
+  Integrated GPU sharing Host Memory:            No
+  Support host page-locked memory mapping:       Yes
-Operatiile trebuie reordonate astfel: calculeaza primii b termeni pentru c[i][j] din blocul C, calculeaza urmatorii b termeni pentru c[i][j] din blocul C, ..., calculeaza ultimii b termeni pentru c[i][j] din blocul C. Generalizand:
+  Alignment requirement for Surfaces:            Yes
+  Device has ECC support:                        Enabled
-{{:asc:lab5:bmm.jpg|}}
+  Device supports Unified Addressing (UVA):      Yes
+  Device supports Managed Memory:                Yes
-Pentru a calcula blocul C<sub>22</sub> folosim formula:
+  Device supports Compute Preemption:            Yes
+  Supports Cooperative Kernel Launch:            Yes
-C<sub>22</sub> %%=%% A<sub>21</sub>B<sub>12</sub> + A<sub>22</sub>B<sub>22</sub> + A<sub>23</sub>B<sub>32</sub> + A<sub>24</sub>B<sub>42</sub>
+  Supports MultiDevice Co-op Kernel Launch:      Yes
+  Device PCI Domain ID / Bus ID / location ID:   0 / 142 / 0
-ce presupune patru inmultiri si patru adunari de matrice. Ideea este ca fiecare inmultire opereaza pe un block suficient de mic ca dimensiune astfel incat sa intre in Cache!
-Versiunea inmultirii de matrice utilizand metoda bloc si ordonarea i-j-k devine:
-<code cpp>
-for (i=0;i<N/b;i++){
-   for (j=0;j<N/b;j++){
-      for (k=0;k<N/b;k++){
-         C[i][j] += A[i][k]*B[k][j]
-      }
-   }
-}
 </code>
-unde:
+Caracteristici GPU A100 (coada ucsx), via query device properties CUDA
-  * b este dimensiunea blocului (presupunem ca b divide N)
-  * C[i][j] este un bloc al matricei C pe linia i si coloana j
-  * "+=" inseamna adunare de matrice
-  * si "*" inseamna inmultire de matrice
-Ce se intampla cu Cache-miss-urile acum?
+<code sh>
-<code cpp>
+Device 0: "NVIDIA A100-PCIE-40GB"
-for (i=0;i<N/b;i++){
+  CUDA Driver Version / Runtime Version          12.4 / 11.4
-   for (j=0;j<N/b;j++){
+  CUDA Capability Major/Minor version number:    8.0
-      // Scrie blocul C[i][j] al lui c in Memorie (Wc)
+  Total amount of global memory:                 40326 MBytes (42285268992 bytes)
-      for (k=0;k<N/b;k++){
+  (108) Multiprocessors, (064) CUDA Cores/MP:    6912 CUDA Cores
-         // Citeste blocul A[i][k] pt a in Cache (Ra)
+  GPU Max Clock rate:                            1410 MHz (1.41 GHz)
-         // Citeste blocul B[k][j] pt b in Cache (Rb)
+  Memory Clock rate:                             1215 Mhz
-         C[i][j] += A[i][k] * B[k][j];
+  Memory Bus Width:                              5120-bit
-      }
+  L2 Cache Size:                                 41943040 bytes
-   }
+  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
-}
+  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
+  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
+  Total amount of constant memory:               65536 bytes
+  Total amount of shared memory per block:       49152 bytes
+  Total shared memory per multiprocessor:        167936 bytes
+  Total number of registers available per block: 65536
+  Warp size:                                     32
+  Maximum number of threads per multiprocessor:  2048
+  Maximum number of threads per block:           1024
+  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
+  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
+  Maximum memory pitch:                          2147483647 bytes
+  Texture alignment:                             512 bytes
+  Concurrent copy and kernel execution:          Yes with 3 copy engine(s)
+  Run time limit on kernels:                     No
+  Integrated GPU sharing Host Memory:            No
+  Support host page-locked memory mapping:       Yes
+  Alignment requirement for Surfaces:            Yes
+  Device has ECC support:                        Enabled
+  Device supports Unified Addressing (UVA):      Yes
+  Device supports Managed Memory:                Yes
+  Device supports Compute Preemption:            Yes
+  Supports Cooperative Kernel Launch:            Yes
+  Supports MultiDevice Co-op Kernel Launch:      Yes
+  Device PCI Domain ID / Bus ID / location ID:   0 / 49 / 0
 </code>
-Pentru (Wc) avem acum (N/b)*(N/b)*b*b Cache-miss-uri, in vreme ce pentru (Ra) si (Rb) avem (N/b)*(N/b)*(N/b)*b*b, astfel ducand la N<sup>2</sup> + 2N<sup>3</sup>/b -> 2N<sup>3</sup>/b Cache-miss-uri pentru intregul algoritm. Combinand acest calcul cu faptul ca avem 2N<sup>3</sup> operatii aritmetice, rezulta un raport r = 2N<sup>3</sup>/b / 2N<sup>3</sup> -> b. Dupa cum am stabilit, r trebuie sa fie maxim (mai mare oricum decat 2-ul obtinut in varianta anterioara). Daca mergem pana la cazul extrem, il vom face pe b = N, dar asta nu este viabil, pentru ca atunci suntem din nou la cazul fara blocuri, de la care tocmai venim...
+===== Optimizarea accesului la memorie =====
-Astfel, acest algoritm functioneaza doar daca blocurile intra in Cache. Acest lucru inseamna ca trei blocuri diferite, de dimensiune b*b, trebuie sa intre in Cache, pentru toate cele trei matrice (a, b si c). Daca C este dimensiunea Cache-ului in elemente de matrice, atunci trebuie sa fie 3b<sup>2</sup> ≤ C sau b ≤ √(C / 3) . Astfel, in cel mai bun caz, r-ul trebuie sa fie si el  √(C / 3).
+Modul cum accesam memoria influenteaza foarte mult performanta sistemului. Cum putem avea arhitecturi foarte diferite din punctul de vedere al ierarhiei de memorie este important de inteles ca nu putem dezvolta un program care sa ruleze optim in toate cazurile. Un program CUDA este portabil caci poate fi usor rulat pe diferite arhitecturi NVIDIA CUDA, insa de cele mai multe ori trebuie ajustat in functie de arhitectura pentru o performanta optima.
-Putem astfel spune, pentru diverse procesoare, cunoscand rata de operatii aritmetice la cache-miss-uri r, care este dimensiunea necesara a Cache-ului, pentru a rula acest algoritm, astfel incat procesorul sa NU astepte niciodata memoria:
+In general pentru arhitecturile de tip GPU, memoria shared este impartita in module de SRAM identice, denumite 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.
-/* PLEASE UPDATE THIS!
+<note important>
-^ Procesor ^ Dimensiune Cache (KB) ^
+Conflictele de access la bancuri de memorie (cache) pot reduce semnificativ performanta.
-|Ultra 2i   | 14.8 |
+</note>
-|Ultra 3    | 4.7 |
-|Pentium 3  | 0.9 |
-|Pentium 3M | 2.4 |
-|Power 3    | 1.8 |
-|Power 4    | 5.4 |
-|Itanium 1  | 31.1 |
-|Itanium 2  | 0.7 |
-*/
-==== Activitate practica - BMM & Optimizare pentru Cache ====
-De aceea incercati sa experimentati cele prezentate in acest laborator, in C. Pentru cei interesati, incercati completarea tabelului de mai sus cu dimensiunea Cache-ului pentru procesoarele voastre personale. Acest lucru presupune evident, si o documentare asupra caracteristicilor sistemului propriu (determinarea r-ului, a dimensiunii Cache-ului etc.).
-<code cpp>
+Cand are loc un bank conflict, hardware-ul serializeaza operatiile cu memoria (warp/wavefront serialization), si face astfel toate threadurile sa astepte pana cand operatiile de memorie sunt efectuate. In unele cazuri, 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.
-#include <stdio.h>
+Spre exemplu daca linia de cache este alcatuita din 16 bancuri de memorie, avem urmatoarele situatii care impacteaza performanta accesului la cache (in comentarii apare echivalentul OpenCL, intrucat memoria shared din CUDA are ca echivalent memoria locala in OpenCL):
-#include <stdlib.h>
-#include <sys/time.h>
-void BMMultiply(int n, double** a, double** b, double** c)
+<code sh>
-{
+__global__ void func(...) {  // __kernel void func(...)
-    int bi=0;
+...
-    int bj=0;
+   __shared__ int *array; // __local int *array;
-    int bk=0;
+   x = array[threadIdx.x]; // x = array[get_local_id(0)]; => performanta 100%, 0 bank conflicts
-    int i=0;
+   x = array[threadIdx.x + 1]; // x = array[get_local_id(0) + 1]; => performanta 100%, 0 bank conflicts
-    int j=0;
+   x = array[threadIdx.x * 4]; // x = array[get_local_id(0) * 4]; => performanta 25%, 4 bank conflicts
-    int k=0;
+   x = array[threadIdx.x * 16]; // x = array[get_local_id(0) * 16]; => performanta 6%,  16 bank conflicts
-    // TODO: set block dimension blockSize
+...
-    int blockSize=100;
-    for(bi=0; bi<n; bi+=blockSize)
-        for(bj=0; bj<n; bj+=blockSize)
-            for(bk=0; bk<n; bk+=blockSize)
-                for(i=0; i<blockSize; i++)
-                    for(j=0; j<blockSize; j++)
-                        for(k=0; k<blockSize; k++)
-                            c[bi+i][bj+j] += a[bi+i][bk+k]*b[bk+k][bj+j];
-}
-int main(void)
-{
-    int n;
-    double** A;
-    double** B;
-    double** C;
-    int numreps = 10;
-    int i=0;
-    int j=0;
-    struct timeval tv1, tv2;
-    struct timezone tz;
-    double elapsed;
-    // TODO: set matrix dimension n
-    n = 500;
-    // allocate memory for the matrices
-    // TODO: allocate matrices A, B & C
-    ///////////////////// Matrix A //////////////////////////
-    // TODO ...
-    ///////////////////// Matrix B //////////////////////////
-    // TODO ...
-    ///////////////////// Matrix C //////////////////////////
-    // TODO ...
-    // Initialize matrices A & B
-    for(i=0; i<n; i++)
-    {
-        for(j=0; j<n; j++)
-        {
-            A[i][j] = 1;
-            B[i][j] = 2;
-        }
-    }
-    //multiply matrices
-    printf("Multiply matrices %d times...\n", numreps);
-    for (i=0; i<numreps; i++)
-    {
-        gettimeofday(&tv1, &tz);
-        BMMultiply(n,A,B,C);
-        gettimeofday(&tv2, &tz);
-        elapsed += (double) (tv2.tv_sec-tv1.tv_sec) + (double) (tv2.tv_usec-tv1.tv_usec) * 1.e-6;
-    }
-    printf("Time = %lf \n",elapsed);
-    //deallocate memory for matrices A, B & C
-    // TODO ...
-    return 0;
 }
 </code>
-===== In loc de concluzie =====
-Intelegerea reala a comportamentului unei aplicatii (algoritm), din punctul de vedere al utilizarii cache-ului (si al performantelor in general), este o chestiune complexa, ce necesita multa rabdare si cunostinte diverse. Deseori, aproximatii utile pot fi folosite pentru a imbunatati unele aspecte ale implementarii curente. Utilizarea blocurilor este intalnita deseori in algoritmi si aplicatii ce necesita performante crescute.
+In cazul arhitecturilor de tip CPU, memoria shared este doar o regiune din RAM. Optimizarile pentru a tine datele critice in memoria shared pentru GPU nu ar prezenta deci aceleasi imbunatatiri de performanta.
+====== Aplicații ======
-===== Exercitii =====
+Urmăriți instrucțiunile de pe [[https://gitlab.cs.pub.ro/asc/asc-public/-/tree/master/labs/cuda/arch|GitLab]]
+<note important>
+Recomandăm sa va delogati mereu de pe serverele din cluster dupa terminarea sesiunii, utilizand comanda ''exit''
-**Task 0**  - Rulați ''task01'', ''task02'' si ''task03''ca exemple pentru [[#activitate_practica_-_optimizare_constantelor_si_al_accesului_la_vectori | Optimizarea constantelor si al accesului la vectori]] folosind matrici liniarizate.
+Alternativ, daca ati uitat sesiuni deschise, puteti verifica acest lucru de pe fep8.grid.pub.ro, utilizand comanda ''squeue''. In cazul in care identificati astfel de sesiuni "agatate", le puteti sterge (si va rugam sa faceti asta), utilizand comanda ''scancel  ID'' unde ID-ul il identificati din comanda anterioara ''squeue''. Puteți folosi mai precis ''squeue -u username'' (username de pe fep8.grid.pub.ro) pentru a vedea doar sesiunile care vă interesează.
-**Task 1**  -  Implemenati  [[#activitate_practica_-_ordinea_buclelor | Ordonarea buclelor]] folosind matrici liniarizate in ''task11.c'' si ''task12.c''. Rulati si observati diferentele.
+Daca nu veti face aceasta delogare, veti putea ajunge in situatia in care sa nu va mai puteti loga pe nodurile din cluster.
+</note>
-**Task 2**  - Implementati  [[#activitate_practica_-_bmm_optimizare_pentru_cache | Optimizari pentru Cache]] folosind matrici liniarizate.
-    * ''task21.c'' Implementati BMM
-    * ''task22.c'' Implementati BMM alaturi de  [[#activitate_practica_-_optimizare_constantelor_si_al_accesului_la_vectori | Optimizarea constantelor si al accesului la vectori]].
-**Task 3**  - In fisierul ''task3.c'' implementati [[#activitate_practica_-_ordinea_buclelor | Ordonarea buclelor]] i-k-j folosind [[#activitate_practica_-_optimizare_constantelor_si_al_accesului_la_vectori | Optimizarea constantelor si al accesului la vectori]]
-**Task 4**  -  (Bonus) In general, nu recomandam alocarea matricelor ca vectori de vectori. Ca bonus va sugeram sa realizati un test unde se face acest tip de alocare si se verifica "performantele" obtinute.
-<note important>
-Toate rularile din acest laborator trebuiesc facute cu optiunea -O0 (O zero), adica ''gcc -O0 -o binary source-file.c''
-Motivul principal este ca optimizarile de compilator ar putea sa ascunda optimizarile pe care le veti incerca (si reusi) voi la laborator.
-</note>
-Si pentru acest laborator se pot utiliza sistemele din cluster prin intermediul fep.grid.pub.ro:
-  - ''ssh -Y username@fep8.grid.pub.ro'' (natural puneti utilizatorul vostru in loc de username)
-  - ''%%srun --x11 -p nehalem --pty /bin/bash%%'' - va permite conectarea pe coada Nehalem cu 14 servere
-  - ''%%singularity run docker://gitlab.cs.pub.ro:5050/asc/asc-public/c-labs:1.3.1 /bin/bash%%'' - accesăm imaginea de docker în cadrul
 ===== Resurse =====
-  * Responsabilul acestui laborator: [[emil.slusanschi@cs.pub.ro|Emil Slușanschi]]
-  * <html><a class="media mediafile mf_pdf" href=":asc:lab5:index?do=export_pdf">PDF laborator</a></html>
-{{:asc:lab5:lab5_skl.tar.gz|Schelet Laborator 5}}
 <hidden>
-{{:asc:lab5:sol:lab5_sol.tar.gz|Soluție Laborator 5}}
+{{:asc:lab8:sol:lab8_sol.tar.gz|Soluție Laborator 8}}
 </hidden>
-==== Discutii interesante ====
+  * Responsabili laborator: Matei Barbu, Alexandru Bala
-  * [[http://stackoverflow.com/questions/11227809/why-is-processing-a-sorted-array-faster-than-an-unsorted-array | De ce este mai rapida procesarea unui vector ordonat?]]
-  * {{:asc:lab5:what_every_programmer_should_know_about_memory_by_ulrich_drepper_.pdf|What every programmer should know about memory.pdf}}
-==== Valgrind ====
+==== Referinte ====
-  * http://valgrind.org/docs/manual/cg-manual.html
+  * Bibliografie
+    * [[https://booksite.elsevier.com/9780124077263/downloads/advance_contents_and_appendices/appendix_C.pdf|Graphics and Computing GPUs]]
+  * Documentatie CUDA:
+    * [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html|CUDA C Programming]]
+    * [[https://docs.nvidia.com/cuda/pdf/CUDA_Compiler_Driver_NVCC.pdf| CUDA NVCC compiler]]
+    * [[https://docs.nvidia.com/cuda/profiler-users-guide/index.html| CUDA Visual Profiler]]
+    * [[https://docs.nvidia.com/cuda/cuda-toolkit-release-notes/index.html|CUDA Dev Toolkit]]
+    * [[https://developer.nvidia.com/cuda-gpus|CUDA GPUs]]
+  * Acceleratoare xl (NVidia P100)
+    * [[https://www.nvidia.com/en-us/data-center/tesla-p100/|NVIDIA Pascal P100]]
+  * Advanced CUDA
+    * [[https://developer.download.nvidia.com/CUDA/training/StreamsAndConcurrencyWebinar.pdf|CUDA Streams 1]]
+    * [[https://devblogs.nvidia.com/gpu-pro-tip-cuda-7-streams-simplify-concurrency/|CUDA Streams 2]]
+    * [[https://devblogs.nvidia.com/introduction-cuda-dynamic-parallelism/|CUDA Dynamic Parallelism]]
+    * [[http://www-personal.umich.edu/~smeyer/cuda/grid.pdf | CUDA Thread Basics]]
+    * [[https://devblogs.nvidia.com/even-easier-introduction-cuda/ | An Even Easier Introduction to CUDA]]
-==== Referințe ====
-  * [[https://github.com/metallurgix | Exemple inmultire matrice]]