Differences
This shows you the differences between two versions of the page.
|
asc:laboratoare:06 [2026/04/08 06:51] alexandru.bala [Optimizarea accesului la memorie] |
asc:laboratoare:06 [2026/04/08 08:14] (current) alexandru.bala [Ierarhia de memorie] |
||
|---|---|---|---|
| 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 | ||
| Line 96: | Line 97: | ||
| | Constant | Grid | Application | Device (VRAM, cached) | | | Constant | Grid | Application | Device (VRAM, cached) | | ||
| | Shared | Block | Kernel | SM (on-chip SRAM) | | | Shared | Block | Kernel | SM (on-chip SRAM) | | ||
| - | | Local | Thread | Kernel | Device (VRAM) | | + | | Local CUDA* | Thread | Kernel | Device (VRAM) | |
| | Register | Thread | Kernel | SM (Register File) | | | 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 | ||
