Show page

Differences

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

--- asc:laboratoare:05 [2024/03/31 23:54]
matei.barbu1905
+++ asc:laboratoare:05 [2025/04/02 10:20] (current)
alexandru.bala [Ierarhia de memorie]
@@ Line 27: / Line 27: @@
 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.
-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 programare 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.
+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.
 {{:asc:lab11:thread.blocks.jpg?direct&360|{{thread.blocks.jpg|''Structura threadurilor in blocuri''}}
-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/09|Laboratorul 9]].
+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]].
 ===== Ierarhia de memorie =====
@@ Line 42: / Line 42: @@
 <code sh>
 /* marcam pentru compilator regValPi in register file */
-__private float regValPi = 3.14f;
+__private__ float regValPi = 3.14f;
 /* compilatorul cel mai probabil oricum incadreaza regVal2Pi ca registru */
 float regVal2Pi = 2 * 3.14f;
 </code>
   *Cea mai rapida forma de memorie de pe GPU
-  *Accesibila doar de catre thread, durata de viata este aceeasi ca si a threadului
+  *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
@@ Line 56: / Line 55: @@
 <code sh>
 /* fiecare work item salveaza un element */
-__local float lArray[lid] = data[gid];
+__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
-  *in functie de implementarea hardware, 100GB/sec -> 2TB/sec
+**Shared Memory**
-  *pentru GPU o memorie rapida, actioneaza ca un cache L1/alt register file, la CPU de regula este doar o portiune din RAM
+<code sh>
-  *accesibila tuturor threadurilor dintr-un bloc (warp/wavefront), durata de viata este aceeasi ca si a blocului
+/* elementele sunt salvate la nivel de bloc */
-  *trebuie evitate conflictele de access (bank conflicts)
+__shared__ int s[64];
+</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)
 **Constant Memory**
 <code sh>
-__const float pi = 3.14f
+__const__ float pi = 3.14f
 </code>
-  * in functie de implementarea hardware, 100GB/sec -> 1TB/sec
+  * In functie de implementarea hardware, 100GB/sec -> 1TB/sec
-  * in general performanta foarte buna, (cache L1/L2, zona dedicata),
+  * In general performanta foarte buna, (cache L1/L2, zona dedicata),
-  * are durata de viata a aplicatiei kernel
+  * Are durata de viata a aplicatiei kernel
 **Global Memory**
 <code sh>
-__kernel void process(__global float* data){ ... }
+__kernel__ void process(__global__ float* data){ ... }
 </code>
-  * in functie de implementarea hardware, 30GB/sec -> 500GB/sec
+  * 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)
+  * 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)
+  * In general latime mare de banda (256-512 biti) si chipuri de memorii de mare viteza (GDDR7)
 **Host Memory (RAM)**
-  * in general, 4GB/sec -> 30GB/sec
+  * In general, 4GB/sec -> 30GB/sec
-  * pentru acces din kernel trebuie transfer/mapare explicita RAM->VRAM pe partea de host/CPU
+  * 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
+  * 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
+  * Viteza de transfer (throughput/latenta) este limitata de magistrala PCIe cat si de memoria RAM
 Caracteristici GPU K40m (coada hpsl), via query device properties CUDA
@@ Line 157: / Line 162: @@
 </code>
-===== Optimizare accesului la memorie =====
+Caracteristici GPU P100 (coada xl), via query device properties CUDA
+<code sh>
+Device 1: "Tesla P100-PCIE-16GB"
+  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>
+Caracteristici GPU A100 (coada ucsx), via query device properties CUDA
+<code sh>
+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>
+===== Optimizarea accesului la memorie =====
 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.
-In general pentru arhitecturile de tip GPU, memoria locala 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.
+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.
 <note important>
@@ Line 168: / Line 253: @@
 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.
+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):
 <code sh>
-__kernel func(...) {
+__global__ void func(...) {  // __kernel void func(...)
 ...
-   __local int *array;
+   __shared__ int *array; // __local int *array;
-   x = array[get_local_id(0)];    // performanta 100%, 0 bank conflicts
+   x = array[threadIdx.x]; // x = array[get_local_id(0)]; => performanta 100%, 0 bank conflicts
-   x = array[get_local_id(0)+1];  // performanta 100%, 0 bank conflicts
+   x = array[threadIdx.x + 1]; // x = array[get_local_id(0) + 1]; => performanta 100%, 0 bank conflicts
-   x = array[get_local_id(0)*4];  // performanta 25%, 4 bank conflicts
+   x = array[threadIdx.x * 4]; // x = array[get_local_id(0) * 4]; => performanta 25%, 4 bank conflicts
-   x = array[get_local_id(0)*16]; // performanta 6%,  16 bank conflicts
+   x = array[threadIdx.x * 16]; // x = array[get_local_id(0) * 16]; => performanta 6%,  16 bank conflicts
 ...
 }
 </code>
-In cazul arhitecturilor de tip CPU, memoria locala este doar o regiune din RAM. Optimizarile pentru a tine datele critice in memoria locala pentru GPU nu ar prezenta deci aceleasi imbunatatiri de performanta.
+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 ======
@@ Line 202: / Line 287: @@
 </hidden>
-  * Responsabili laborator: Barbu Matei
+  * Responsabili laborator: Matei Barbu, Alexandru Bala
 ==== Referinte ====
+  * Bibliografie
+    * [[https://booksite.elsevier.com/9780124077263/downloads/advance_contents_and_appendices/appendix_C.pdf|Graphics and Computing GPUs]]
   * Documentatie CUDA:
-    * [[https://docs.nvidia.com/pdf/CUDA_C_Programming_Guide.pdf|CUDA C Programming]]
+    * [[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://developer.download.nvidia.com/compute/cuda/9.1/Prod/docs/sidebar/CUDA_Toolkit_Release_Notes.pdf|CUDA 9.1 Toolkit]]
+    * [[https://docs.nvidia.com/cuda/cuda-toolkit-release-notes/index.html|CUDA Dev Toolkit]]
     * [[https://developer.nvidia.com/cuda-gpus|CUDA GPUs]]
-  * Acceleratoare hpsl (hpsl-wn01, hpsl-wn02, hpsl-wn03)
+  * Acceleratoare xl (NVidia P100)
-    * [[http://international.download.nvidia.com/tesla/pdf/tesla-k40-passive-board-spec.pdf|NVIDIA Tesla K40M]]
+    * [[https://www.nvidia.com/en-us/data-center/tesla-p100/|NVIDIA Pascal P100]]
-    * [[https://en.wikipedia.org/wiki/Nvidia_Tesla|NVIDIA Tesla]]
-  * Acceleratoare dp (dp-wn01, dp-wn02, dp-wn03)
-    * [[https://www.nvidia.com/docs/IO/43395/NV_DS_Tesla_C2050_C2070_jul10_lores.pdf|NVIDIA Tesla C2070]]
-    * [[http://www.nvidia.com/docs/io/43395/nv_ds_tesla_c2050_c2070_apr10_final_lores.pdf|NVIDIA Tesla 2050/2070]]
-    * [[https://cseweb.ucsd.edu/classes/fa12/cse141/pdf/09/GPU_Gahagan_FA12.pdf|NVIDIA CUDA Fermi/Tesla]]
   * Advanced CUDA
-    * [[https://devblogs.nvidia.com/gpu-pro-tip-cuda-7-streams-simplify-concurrency/|CUDA Streams]]
+    * [[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]]