Laboratorul 05 - Arhitectura GPU NVIDIA CUDA
- Ierarhia de memorie
- Optimizarea accesului la memorie
Aplicații
- Resurse
  - Referinte

Laboratorul 05 - Arhitectura GPU NVIDIA CUDA

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.

Arhitectura NVIDIA FERMI aici, Tesla 2070, coada executie fep8.grid.pub.ro → dp

Arhitectura NVIDIA KEPLER aici, Tesla K40M, coada executie fep8.grid.pub.ro → hpsl

Prima arhitectura NVIDIA complet programabila a fost G80 (ex. 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 aici.

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 aici sau pe wiki aici. Fiecare noua arhitectura are un codename (ex Fermi, Pascal) si este reprezentata de un “compute capability” (list aici). Cu cat arhitectura este mai noua, cu atat sunt suportate mai multe facilitati din API-urile CUDA si OpenCL.

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.

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.

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.

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.

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.

Filosofia din spatele arhitecturii este permiterea rularii unui numar foarte mare de threaduri. Acest lucru este facut posibil prin paralelismul existent la nivel hardware.

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.

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 Laboratorul 6.

Ierarhia de memorie

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).

Register File

/* marcam pentru compilator regValPi in register file */
__private float regValPi = 3.14f;
/* compilatorul cel mai probabil oricum incadreaza regVal2Pi ca registru */
float regVal2Pi = 2 * 3.14f;

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

Local Memory

/* fiecare work item salveaza un element */
__local float lArray[lid] = data[gid];

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

Shared Memory

/* elementele sunt salvate la nivel de bloc */
__shared__ int s[64];

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

__const float pi = 3.14f

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

Global Memory

__kernel void process(__global float* data){ ... }

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 (GDDR5)

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

Caracteristici GPU K40m (coada hpsl), via query device properties CUDA

Device 0: "Tesla K40m"
  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)
  (15) Multiprocessors, (192) CUDA Cores/MP:     2880 CUDA Cores
  GPU Max Clock rate:                            745 MHz (0.75 GHz)
  Memory Clock rate:                             3004 Mhz
  Memory Bus Width:                              384-bit
  L2 Cache Size:                                 1572864 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
  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

Caracteristici GPU M2070 (coada dp), via query device properties CUDA

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

Caracteristici GPU P100 (coada xl), via query device properties CUDA

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

Caracteristici GPU A100 (coada ucsx), via query device properties CUDA

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

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 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.

Conflictele de access la bancuri de memorie (cache) pot reduce semnificativ performanta.

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):

__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
...
}

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

Urmăriți instrucțiunile de pe GitLab

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ă.

Daca nu veti face aceasta delogare, veti putea ajunge in situatia in care sa nu va mai puteti loga pe nodurile din cluster.

Resurse

Responsabili laborator: Matei Barbu, Alexandru Bala

Referinte

Bibliografie
- Graphics and Computing GPUs
Documentatie CUDA:
Acceleratoare xl (NVidia P100)
- NVIDIA Pascal P100
Advanced CUDA

Table of Contents