Show page

Differences

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

--- asc:laboratoare:06 [2025/04/09 09:11]
alexandru.bala [Fluxuri nonimplicite]
+++ asc:laboratoare:06 [2026/02/23 18:47] (current)
giorgiana.vlasceanu
@@ Line 1: / Line 1: @@
-====== Laboratorul 06 - Advanced CUDA ======
+====== Laboratorul 06 - Arhitectura GPU NVIDIA CUDA ======
-===== Spatiu unificat memorie =====
+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.
-De la [[http://developer.download.nvidia.com/compute/cuda/6_0/rel/docs/CUDA_Toolkit_Release_Notes.pdf| CUDA 6.0]], NVIDIA a schimbat semnificativ modelul de programare prin facilitarea comunicarii unitatii CPU (host) cu unitatea GPU (device), in mod transparent prin acelasi set de adrese de memorie virtuale. Astfel exista posibilitatea ca prin acelasi pointer de memorie sa se scrie date atat de catre CPU cat si de catre GPU. Evident transferurile de memorie au loc intre spatii diferite de adresare (ex RAM vs VRAM), dar acest lucru se intampla transparent la nivel de aplicatie CUDA / pentru programator.
+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
-{{:asc:lab9:nv-unified.png?640|NVIDIA Unified Memory}}
+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
-Mai jos avem un exemplu de folosire a memoriei unificate. Singura diferenta fata de alocarea pe CPU/HOST este ca memoria trebuie alocata cu cudaMallocManaged si dealocata cu cudaFree.
+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]].
-<code C>
+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.
-#include <iostream>
-#include <math.h>
-// CUDA kernel to add elements of two arrays
-__global__
-void add(int n, float *x, float *y)
-{
-  int index = blockIdx.x * blockDim.x + threadIdx.x;
-  int stride = blockDim.x * gridDim.x;
-  for (int i = index; i < n; i += stride)
-    y[i] = x[i] + y[i];
-}
-int main(void)
-{
-  int N = 1<<20;
-  float *x, *y;
-  // Allocate Unified Memory -- accessible from CPU or GPU
-  cudaMallocManaged(&x, N*sizeof(float));
-  cudaMallocManaged(&y, N*sizeof(float));
-  // initialize x and y arrays on the host
-  for (int i = 0; i < N; i++) {
-    x[i] = 1.0f;
-    y[i] = 2.0f;
-  }
-  // Launch kernel on 1M elements on the GPU
-  int blockSize = 256;
-  int numBlocks = (N + blockSize - 1) / blockSize;
-  add<<<numBlocks, blockSize>>>(N, x, y);
-  // Wait for GPU to finish before accessing on host
-  cudaDeviceSynchronize();
-  // Check for errors (all values should be 3.0f)
-  float maxError = 0.0f;
-  for (int i = 0; i < N; i++)
-    maxError = fmax(maxError, fabs(y[i]-3.0f));
-  std::cout << "Max error: " << maxError << std::endl;
-  // Free memory
-  cudaFree(x);
-  cudaFree(y);
-  return 0;
-}
-</code>
-===== Operatii atomice CUDA =====
+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.
-CUDA ofera acces la multiple operatii atomice tip citire-modificare-scriere. Acestea presupun serializarea accesului in contextul mai multor thread-uri. Functiile sunt limitate la anumite tipuri de date:
+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.
-  - int
-  - unsigned int
-  - unsigned long long int
-  - float
-  - double
-Exemple de functii atomice:
+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.
-  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicadd|atomicAdd]]
-  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicsub|atomicSub]]
-  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicexch|atomicExch]]
-  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicmin|atomicMin]]
-  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicmax|atomicMax]]
-  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicinc|atomicInc]]
-  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicdec|atomicDec]]
-  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicand|atomicAnd]]
-  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicor|atomicOr]]
-<note important>
+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.
-A se consulta cu atentie documentatia CUDA inainte de folosirea unei operatii atomice (legat de contextul in care se aplica, cum opereaza, limitari etc).
-</note>
-In codul de mai jos se lanseaza un kernel concurrentRW in configuratie numBlocks=8 fiecare cu cate 10 thread-uri.
+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.
-<code C>
+{{:asc:lab11:cuda-arch.png?direct&720|}}
-#include <iostream>
-#define NUM_ELEM        8
+Filosofia din spatele arhitecturii este permiterea rularii unui numar foarte mare de threaduri. Acest lucru este facut posibil prin paralelismul existent la nivel hardware.
-#define NUM_THREADS     10
-using namespace std;
+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.
-__global__ void concurrentRW(int *data) {
+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.
-...
-}
-int main(int argc, char *argv[]) {
+{{:asc:lab11:thread.blocks.jpg?direct&360|{{thread.blocks.jpg|''Structura threadurilor in blocuri''}}
-    int* data = NULL;
-    bool errorsDetected = false;
-    cudaMallocManaged(&data, NUM_ELEM * sizeof(unsigned long long int));
+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]].
-    if (data == 0) {
-        cout << "[HOST] Couldn't allocate memory\n";
-        return 1;
-    }
-    // init all elements to 0
+===== Ierarhia de memorie =====
-    cudaMemset(data, 0, NUM_ELEM * sizeof(unsigned long long int));
-    // launch kernel writes
+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).
-    concurrentRW<<<NUM_ELEM, NUM_THREADS>>>(data);
-    cudaDeviceSynchronize();
-    if (cudaSuccess != cudaGetLastError()) {
-        return 1;
-    }
-    for(int i = 0; i < NUM_ELEM; i++) {
+{{:asc:lab11:mem.hierarchy.png?direct|Ierarhia memoriei in CUDA}}
-        cout << i << ". " << data[i] << endl;
-        if(data[i] != (NUM_THREADS * (NUM_THREADS - 1) / 2)) {
-            errorsDetected = true;
-        }
-    }
-    if(errorsDetected) {
+**Register File**
-        cout << "Errors detected" << endl;
+<code sh>
-    } else {
+/* marcam pentru compilator regValPi in register file */
-        cout << "OK" << endl;
+__private__ float regValPi = 3.14f;
-    }
+/* compilatorul cel mai probabil oricum incadreaza regVal2Pi ca registru */
+float regVal2Pi = 2 * 3.14f;
-    return 0;
-}
 </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
-Functia concurrentRW citeste valoarea de la adresa data[blockIdx.x], o incrementeaza cu threadIdx.x si apoi o scrie.
+**Local Memory**
-In acest caz avem 10 thread-uri care fac operatii citire/scriere la aceeasi adresa, deci un comportament nedefinit.
+<code sh>
+/* fiecare work item salveaza un element */
-<code C>
+__local__ float lArray[lid] = data[gid];
-__global__ void concurrentRW(int *data) {
-    // NUM_THREADS try to read and write at same location
-    data[blockIdx.x] = data[blockIdx.x] + threadIdx.x;
-}
 </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
-Exemplu rezultat:
+**Shared Memory**
 <code sh>
-. 9
+/* elementele sunt salvate la nivel de bloc */
-. 9
+__shared__ int s[64];
-. 9
-. 9
-. 9
-. 9
-. 9
-. 9
-Errors detected
 </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)
-Corect ar fi folosirea functiei atomicAdd pentru a serializa accesul.
+**Constant Memory**
+<code sh>
-<code C>
+__const__ float pi = 3.14f
-__global__ void concurrentRW(int *data) {
-    // NUM_THREADS try to read and write at same location
-    atomicAdd(&data[blockIdx.x], threadIdx.x);
-}
 </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
-Rezultatul rularii este:
+**Global Memory**
 <code sh>
-. 45
+__kernel__ void process(__global__ float* data){ ... }
-. 45
-. 45
-. 45
-. 45
-. 45
-. 45
-. 45
-OK
 </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 (GDDR7)
-==== Operatii atomice system wide ====
+**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
-Unitatile GPU ce au Compute capability 6.x permit largirea scopului operatiilor atomice. De exemplu atomicAdd_system garanteaza ca operatia este atomica cand atat thread-urile de pe unitatea GPU cat si cele de pe unitatea CPU incearca sa acceseze datele. Mai jos avem un exemplu de folosire al functiei atomicAdd_system.
+Caracteristici GPU K40m (coada hpsl), via query device properties CUDA
-<code C>
+<code sh>
-__global__ void mykernel(int *addr) {
+Device 0: "Tesla K40m"
-  atomicAdd_system(addr, 10);       // only available on devices with compute capability 6.x
+  CUDA Driver Version / Runtime Version          9.1 / 9.1
-}
+  CUDA Capability Major/Minor version number:    3.5
+  Total amount of global memory:                 11441 MBytes (11996954624 bytes)
-void foo() {
+  (15) Multiprocessors, (192) CUDA Cores/MP:     2880 CUDA Cores
-  int *addr;
+  GPU Max Clock rate:                            745 MHz (0.75 GHz)
-  cudaMallocManaged(&addr, 4);
+  Memory Clock rate:                             3004 Mhz
-  *addr = 0;
+  Memory Bus Width:                              384-bit
+  L2 Cache Size:                                 1572864 bytes
-   mykernel<<<...>>>(addr);
+  Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
-   __sync_fetch_and_add(addr, 10);  // CPU atomic operation
+  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: 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 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 / 8 / 0
 </code>
-===== Operatii asincrone CUDA =====
+Caracteristici GPU M2070 (coada dp), via query device properties CUDA
-In CUDA, urmatoarele operatii sunt definite ca fiind independente si pot fi executate concurent:
-  - Calcule pe unitatea host
-  - Calcule pe unitatea device
-  - Transfer memorie host -> device
-  - Transfer memorie device -> host
-  - Transfer memorie device -> device
-Nivelul de concurenta o sa depinda si de capabilitatea unitatilor GPU (compute capability).
-In continuare vom explora mai multe scenarii de executie concurenta a operatiilor descrise.
-==== Executie asincrona Host si Device ====
-Folosind apeluri asincrone, operatiile de executie catre device sunt puse in coada avand controlul intors catre host instant. Astfel unitatea host poate continua executia fara sa fie blocata in asteptarea altor task-uri.
-Urmatoarele operatii sunt asincrone relativ la host:
-  - Lansari de kernel
-  - Copieri in cadrul spatiului de memorie al unui device
-  - Copiere memorie host -> device, avand < 64 KB
-  - Copiere memorie host -> device, avand functii cu sufix Async
-  - Functii memorie set (setare / initializare de memorie la o valoare)
-Pentru a face debug unor scenarii de executie asincrona se poate dezactiva complet executia asincrona setand variabila de mediu CUDA_LAUNCH_BLOCKING la 1. Executia de kernels este sincrona cand se ruleaza cu un profiler (Nsight, Visual Profiler).
-==== Fluxuri nonimplicite ====
-Pentru a folosi cudaMemcpyAsync, este necesar lucrul cu fluxuri nonimplicite (non-default streams), care, in C/C++ pot fi declarate, create si distruse in partea de cod de pe host (CPU) in urmatorul fel:
-<code C>
-cudaStream_t stream1;
-cudaError_t result;
-result = cudaStreamCreate(&stream1);
-result = cudaStreamDestroy(stream1);
+<code sh>
+Device 0: "Tesla M2070"
+  CUDA Driver Version / Runtime Version          9.1 / 9.1
+  CUDA Capability Major/Minor version number:    2.0
+  Total amount of global memory:                 5302 MBytes (5559156736 bytes)
+  (14) Multiprocessors, ( 32) CUDA Cores/MP:     448 CUDA Cores
+  GPU Max Clock rate:                            1147 MHz (1.15 GHz)
+  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>
-Odata creat un astfel de flux, el poate fi utilizat in procesul de copiere a memoriei host -> device astfel:
+Caracteristici GPU P100 (coada xl), via query device properties CUDA
-<code C>
+<code sh>
-// num_bytes = N * sizeof (type_a);
+Device 1: "Tesla P100-PCIE-16GB"
-result = cudaMemcpyAsync(d_a, a, num_bytes, cudaMemcpyHostToDevice, stream1);
+  CUDA Driver Version / Runtime Version          12.2 / 11.4
+  CUDA Capability Major/Minor version number:    6.0
+  Total amount of global memory:                 16276 MBytes (17066885120 bytes)
+  (056) Multiprocessors, (064) CUDA Cores/MP:    3584 CUDA Cores
+  GPU Max Clock rate:                            1329 MHz (1.33 GHz)
+  Memory Clock rate:                             715 Mhz
+  Memory Bus Width:                              4096-bit
+  L2 Cache Size:                                 4194304 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:        65536 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 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 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 / 142 / 0
 </code>
-Pentru a emite un kernel către un flux nonimplicit, specificăm identificatorul fluxului ca al patrulea parametru de configurare a execuției. Se observă și un al treilea parametru de configurare a execuției, care este folosit pentru a aloca memorie partajată (shared memory) device-ului (GPU-ului), utilizându-se 0 dacă nu se dorește acest aspect.
+Caracteristici GPU A100 (coada ucsx), via query device properties CUDA
-<code C>
+<code sh>
-increment<<<1,N,0,stream1>>>(d_a);
+Device 0: "NVIDIA A100-PCIE-40GB"
+  CUDA Driver Version / Runtime Version          12.4 / 11.4
+  CUDA Capability Major/Minor version number:    8.0
+  Total amount of global memory:                 40326 MBytes (42285268992 bytes)
+  (108) Multiprocessors, (064) CUDA Cores/MP:    6912 CUDA Cores
+  GPU Max Clock rate:                            1410 MHz (1.41 GHz)
+  Memory Clock rate:                             1215 Mhz
+  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>
-==== Executie asincrona programe kernel ====
+===== Optimizarea accesului la memorie =====
-Arhitecturile cu compute capability 2.x sau mai nou, pot executa in paralel instante de kernel diferite. Aceste unitati de executie o sa aibe proprietate concurrentKernels setata la 1 (se face query la device properties inainte). Numarul maxim de lansari asincrone de kernele diferite este dependent de arhitectura (se verifica in functie de compute capability). Singura restrictie este ca programele kernel sa fie in acelasi context.
+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.
-==== Executie si transfer date asincron  ====
+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.
-Anumite device-uri pot executa un transfer asincron memorie alaturi de o executie de kernel. Acest lucru este dependent de compute capability si se poate verifica in device property asyncEngineCount.
+<note important>
+Conflictele de access la bancuri de memorie (cache) pot reduce semnificativ performanta.
+</note>
-{{:asc:lab9:cuda_async.png?900|}}
+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.
+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):
-De asemenea, se pot face transferuri de memorie intra-device simultan cu executia de kernel cand atat device property concurrentKernels, cat si asyncEngineCount sunt 1.
+<code sh>
+__global__ void func(...) {  // __kernel void func(...)
+...
+   __shared__ int *array; // __local int *array;
+   x = array[threadIdx.x]; // x = array[get_local_id(0)]; => performanta 100%, 0 bank conflicts
+   x = array[threadIdx.x + 1]; // x = array[get_local_id(0) + 1]; => performanta 100%, 0 bank conflicts
+   x = array[threadIdx.x * 4]; // x = array[get_local_id(0) * 4]; => performanta 25%, 4 bank conflicts
+   x = array[threadIdx.x * 16]; // x = array[get_local_id(0) * 16]; => performanta 6%,  16 bank conflicts
+...
+}
+</code>
-{{:asc:lab9:cuda_async_2.png?900|}}
+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.
-===== Dynamic Paralellism  =====
+====== Aplicații ======
-Paralelismul dinamic consta in posibilitatea de a lansa programe kernel din thread-urile ce ruleaza pe device/GPU. In alte cuvinte, unitatea GPU poate sa isi  atribuie noi task-uri/thread-uri fara interventia unitatii host/CPU. Aceasta manifestare este utila in problemele unde maparea threaduri<->date nu este simpla/triviala. De exemplu, in situatia unde unele thread-uri ar avea prea putin de lucru, iar altele prea mult (imaginea de mai jos, simulare fluide) - o situatia debalansata computational.
+Urmăriți instrucțiunile de pe [[https://gitlab.cs.pub.ro/asc/asc-public/-/tree/master/labs/cuda/arch|GitLab]]
-{{:asc:lab9:dynamic-paralellism.png?560|Fluid simulation}}
-Cerintele pentru paralelism dinamic sunt CUDA 5.0 ca Toolkit si respectiv Compute Capability 3.5. O lista cu GPU-uri NVIDIA si Compute Capability se regaseste [[https://developer.nvidia.com/cuda-gpus|aici]].
-===== Exercitii =====
-Urmăriți instrucțiunile de pe [[https://gitlab.cs.pub.ro/asc/asc-public/-/tree/master/labs/cuda/advanced|GitLab]].
 <note important>
 Recomandăm sa va delogati mereu de pe serverele din cluster dupa terminarea sesiunii, utilizand comanda ''exit''
-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ă. (Sau ''squeue --me'').
+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ă.
 Daca nu veti face aceasta delogare, veti putea ajunge in situatia in care sa nu va mai puteti loga pe nodurile din cluster.
 </note>
 ===== Resurse =====
 <hidden>
-{{:asc:lab9:sol:lab9_sol.tar.gz|Soluție Laborator 6}}
+{{:asc:lab8:sol:lab8_sol.tar.gz|Soluție Laborator 8}}
 </hidden>
-/* {{:asc:lab9:sol:lab9_sol.zip|Solutie Laborator 6}} */
-/* {{:asc:lab6:asc_lab9.pdf|Enunt Laborator 6}} */
   * Responsabili laborator: Matei Barbu, Alexandru Bala
@@ Line 307: / Line 301: @@
     * [[https://www.nvidia.com/en-us/data-center/tesla-p100/|NVIDIA Pascal P100]]
   * Advanced CUDA
-    * [[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]]
     * [[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]]
-    * [[https://developer.nvidia.com/blog/how-overlap-data-transfers-cuda-cc/|How to Overlap Data Transfers in CUDA C/C++]]
+    * [[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]]