Differences
This shows you the differences between two versions of the page.
|
asc:laboratoare:05 [2024/04/02 01:36] alexandru.bala [Ierarhia de memorie] |
asc:laboratoare:05 [2025/04/02 10:20] (current) alexandru.bala [Ierarhia de memorie] |
||
|---|---|---|---|
| Line 42: | Line 42: | ||
| <code sh> | <code sh> | ||
| /* marcam pentru compilator regValPi in register file */ | /* 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 */ | /* compilatorul cel mai probabil oricum incadreaza regVal2Pi ca registru */ | ||
| float regVal2Pi = 2 * 3.14f; | float regVal2Pi = 2 * 3.14f; | ||
| </code> | </code> | ||
| *Cea mai rapida forma de memorie de pe GPU | *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: | *Un kernel complex poate determina folosirea unui numar mare de registrii si astfel: | ||
| * limitarea executiei multor thread-uri simultan | * limitarea executiei multor thread-uri simultan | ||
| Line 55: | Line 55: | ||
| <code sh> | <code sh> | ||
| /* fiecare work item salveaza un element */ | /* fiecare work item salveaza un element */ | ||
| - | __local float lArray[lid] = data[gid]; | + | __local__ float lArray[lid] = data[gid]; |
| </code> | </code> | ||
| - | *Ca si in cazul registrilor, este accesibila doar de catre thread, durata de viata este aceeasi ca si a threadului | + | *In functie de implementarea hardware, 100GB/sec -> 2TB/sec |
| - | *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 |
| - | *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 |
| **Shared Memory** | **Shared Memory** | ||
| Line 66: | Line 66: | ||
| __shared__ int s[64]; | __shared__ int s[64]; | ||
| </code> | </code> | ||
| - | *accesibila tuturor threadurilor dintr-un bloc (warp/wavefront), durata de viata este aceeasi ca si a blocului | + | *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) | + | *Trebuie evitate conflictele de access (bank conflicts) |
| **Constant Memory** | **Constant Memory** | ||
| <code sh> | <code sh> | ||
| - | __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 -> 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** | **Global Memory** | ||
| <code sh> | <code sh> | ||
| - | __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 -> 500GB/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 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)** | **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 | Caracteristici GPU K40m (coada hpsl), via query device properties CUDA | ||
| Line 246: | Line 246: | ||
| 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. | 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> | <note important> | ||
| Line 253: | 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. | 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> | <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> | </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 ====== | ====== Aplicații ====== | ||
| Line 287: | Line 287: | ||
| </hidden> | </hidden> | ||
| - | * Responsabili laborator: Matei Barbu | + | * Responsabili laborator: Matei Barbu, Alexandru Bala |
| ==== Referinte ==== | ==== Referinte ==== | ||
| Line 293: | Line 293: | ||
| * [[https://booksite.elsevier.com/9780124077263/downloads/advance_contents_and_appendices/appendix_C.pdf|Graphics and Computing GPUs]] | * [[https://booksite.elsevier.com/9780124077263/downloads/advance_contents_and_appendices/appendix_C.pdf|Graphics and Computing GPUs]] | ||
| * Documentatie CUDA: | * 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/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/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]] | * [[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 | * 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]] | * [[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]] | ||
