Differences
This shows you the differences between two versions of the page.
|
asc:laboratoare:06 [2026/02/23 18:47] giorgiana.vlasceanu |
asc:laboratoare:06 [2026/04/08 08:14] (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. | 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 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. | + | 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, blocurilor si grilei pot fi uni/bi/tri-dimensionale. 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''}} | {{:asc:lab11:thread.blocks.jpg?direct&360|{{thread.blocks.jpg|''Structura threadurilor in blocuri''}} | ||
| Line 41: | Line 41: | ||
| **Register File** | **Register File** | ||
| <code sh> | <code sh> | ||
| - | /* marcam pentru compilator regValPi in register file */ | + | /* compilatorul incadreaza regVal2Pi ca registru in CUDA */ |
| - | __private__ float regValPi = 3.14f; | + | |
| - | /* compilatorul cel mai probabil oricum incadreaza regVal2Pi ca registru */ | + | |
| float regVal2Pi = 2 * 3.14f; | float regVal2Pi = 2 * 3.14f; | ||
| + | /* marcam pentru compilator regValPi in register file in OpenCL */ | ||
| + | __private__ float regValPi = 3.14f; | ||
| </code> | </code> | ||
| - | *Cea mai rapida forma de memorie de pe GPU | + | *Cea mai rapida forma de memorie de pe GPU, care lucreaza la aceeasi frecventa cu core-urile CUDA (poate atinge viteze mai mari de 300 TB/s) |
| *Accesibila doar de catre thread, iar 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: | *Un kernel complex poate determina folosirea unui numar mare de registrii si astfel: | ||
| Line 54: | Line 54: | ||
| **Local Memory** | **Local Memory** | ||
| <code sh> | <code sh> | ||
| - | /* fiecare work item salveaza un element */ | + | /* fiecare work item salveaza un element in OpenCL */ |
| __local__ float lArray[lid] = data[gid]; | __local__ float lArray[lid] = data[gid]; | ||
| </code> | </code> | ||
| - | *In functie de implementarea hardware, 100GB/sec -> 2TB/sec | + | *In functie de implementarea hardware, 100GB/sec -> 300TB/sec |
| *Pentru GPU o memorie rapida, actioneaza ca un cache L1/alt register file, la CPU de regula este doar o portiune din RAM | *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 | *Ca si in cazul registrilor, este accesibila doar de catre thread, iar durata de viata este aceeasi ca si a threadului | ||
| Line 66: | Line 66: | ||
| __shared__ int s[64]; | __shared__ int s[64]; | ||
| </code> | </code> | ||
| - | *Accesibila tuturor threadurilor dintr-un bloc (warp/wavefront), iar durata de viata este aceeasi ca si a blocului | + | *In functie de implementarea hardware, 100GB/sec -> 200TB/sec |
| + | *Accesibila tuturor threadurilor dintr-un bloc (care conține mai multe warp-uri/wavefront-uri), iar durata de viata este aceeasi ca si a blocului | ||
| *Trebuie evitate conflictele de access (bank conflicts) | *Trebuie evitate conflictele de access (bank conflicts) | ||
| Line 73: | Line 74: | ||
| __const__ float pi = 3.14f | __const__ float pi = 3.14f | ||
| </code> | </code> | ||
| - | * In functie de implementarea hardware, 100GB/sec -> 1TB/sec | + | * In functie de implementarea hardware, 100GB/sec -> 100TB/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 | ||
| Line 81: | Line 82: | ||
| __kernel__ void process(__global__ float* data){ ... } | __kernel__ void process(__global__ float* data){ ... } | ||
| </code> | </code> | ||
| - | * In functie de implementarea hardware, 30GB/sec -> 500GB/sec | + | * In functie de implementarea hardware, 30GB/sec -> 8TB/sec |
| - | * Video RAM (VRAM), de regula cu o capacitate intre 1GB si 12GB in functie de placa video | + | * Video RAM (VRAM), de regula cu o capacitate intre 4GB (entry-level) si 288GB (Nvidia B300) 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 (GDDR7) | * In general latime mare de banda (256-512 biti) si chipuri de memorii de mare viteza (GDDR7) | ||
| **Host Memory (RAM)** | **Host Memory (RAM)** | ||
| - | * In general, 4GB/sec -> 30GB/sec | + | * In general, 4GB/sec -> 400GB/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 | ||
| + | |||
| + | ^ Memory Type ^ Scope ^ Lifetime ^ Location ^ | ||
| + | | Global | Grid | Application | Device (VRAM) | | ||
| + | | Constant | Grid | Application | Device (VRAM, cached) | | ||
| + | | Shared | Block | Kernel | SM (on-chip SRAM) | | ||
| + | | Local CUDA* | Thread | Kernel | Device (VRAM) | | ||
| + | | Register | Thread | Kernel | SM (Register File) | | ||
| + | |||
| + | *A nu se confunda cu Local din OpenCL. | ||
| Caracteristici GPU K40m (coada hpsl), via query device properties CUDA | Caracteristici GPU K40m (coada hpsl), via query device properties CUDA | ||
| Line 253: | Line 263: | ||
| 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. | 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): | + | 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,luandu-se in considerare ca asignarea valorilor in bancuri este ciclica, deci de la o adresa incolo se revine la primul banc): |
| <code sh> | <code sh> | ||
