Show page

Differences

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

--- asc:laboratoare:05 [2024/04/02 11:14]
alexandru.bala [Ierarhia de memorie]
+++ asc:laboratoare:05 [2026/02/23 18:48] (current)
giorgiana.vlasceanu
@@ Line 1: / Line 1: @@
-====== Laboratorul 05 - Arhitectura GPU NVIDIA CUDA ======
+====== Laboratorul 05 - Arhitecturi de tip GPGPU ======
-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.
+===== Introducere =====
-Arhitectura NVIDIA FERMI [[https://www.nvidia.com/content/PDF/fermi_white_papers/NVIDIA_Fermi_Compute_Architecture_Whitepaper.pdf|aici]], Tesla 2070, coada executie fep8.grid.pub.ro -> dp
+==== Motivație ====
-Arhitectura NVIDIA KEPLER [[https://www.nvidia.com/content/PDF/kepler/NVIDIA-Kepler-GK110-Architecture-Whitepaper.pdf|aici]], Tesla K40M, coada executie fep8.grid.pub.ro -> hpsl
+Succesorul chipurilor de prelucrare grafică din jocurile aracade, procesoarul grafic, acronim **GPU** (**G**raphics **P**rocessing **U**nit), este un circuit electronic, specializat, în crearea și manipularea imaginilor trimise către un afișaj electronic (e.g. monitor).
-Prima arhitectura NVIDIA complet programabila a fost G80 (ex. [[http://www.nvidia.com/page/8800_tech_briefs.html|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 [[http://s08.idav.ucdavis.edu/luebke-nvidia-gpu-architecture.pdf|aici]].
+Utilitatea lui s-a extins ulterior către probleme "embarrassingly parallel", iar astăzi sunt folosite la antrenarea rețelelor neurale și minarea de criptomonede. Vorbim aici despre întrebuințarea unui **GPGPU** (**G**eneral **P**urpose GPU), un procesor grafic cu o flexibilitate ridicată de programare, capabil de a rezolva și probleme generale.
-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 [[https://www.geforce.com/hardware/technology/cuda/supported-gpus|aici]] sau pe wiki [[https://en.wikipedia.org/wiki/List_of_Nvidia_graphics_processing_units|aici]]. Fiecare noua arhitectura are un codename (ex Fermi, Pascal) si este reprezentata de un "compute capability" (list [[https://developer.nvidia.com/cuda-gpus|aici]]). Cu cat arhitectura este mai noua, cu atat sunt suportate mai multe facilitati din API-urile CUDA si OpenCL.
+==== Teorie ====
-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.
+În execuție, o arhitectură de tip GPU folosește paradigma **SIMD** (**S**ingle **I**nstruction **M**ultiple **D**ata, vezi taxonomia lui Flynn). Acesta presupune:
+  * schimb rapid de context între thread-uri,
+  * planificarea în grupuri de thread-uri,
+  * și orientare către prelucrari masive de 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.
+Deci, unitățile de tip GPU sunt potrivite pentru paralelismul de date, adică pentru un flux intensiv computațional, cu puține decizii de control.
-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.
+<note important>
+Nu orice algoritm paralel rulează optim pe o arhitectură GPGPU. În principiu, probleme de tip SIMD sau MIMD se pretează rulării pe GPU-uri.
+</note>
-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.
+De obicei, termenul de GPGPU apare atunci când unitatea GPU este folosită ca și //coprocesor matematic//. Astăzi, majoritatea unităților de tip GPU sunt și GPGPU.
+Ampla folosire a acestora se datorează:
+  * diferențelor de putere de procesare brută dintre CPU și GPU (instrucțiuni/secundă)
+  * standardizarea de API-uri care ușurează munca programatorilor
+  * răspândirea aplicațiilor ce pot beneficia de pe urma paralelismului de tip SIMD
+  * cererii pe piața unităților computaționale destinate:
+    * atât consumatorilor (PC, Smartphone, TV, etc.),
+    * cât și mediilor industriale (Automotive, HPC etc).
-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.
+Principalii producatori de core-uri IP (intellectual property) tip GPU sunt:
+  * Intel http://en.wikipedia.org/wiki/List_of_Intel_graphics_processing_units
+  * Nvidia http://en.wikipedia.org/wiki/List_of_Nvidia_graphics_processing_units
+  * Amd http://en.wikipedia.org/wiki/List_of_AMD_graphics_processing_units
+  * Apple https://en.wikipedia.org/wiki/Apple_silicon
+  * Imagination http://en.wikipedia.org/wiki/List_of_PowerVR_products
+  * Qualcomm http://en.wikipedia.org/wiki/Adreno
+  * Vivante http://en.wikipedia.org/wiki/Vivante_Corporation
-{{:asc:lab11:cuda-arch.png?direct&720|}}
+Dacă un IP de GPU este integrat pe aceeași //pastilă de siliciu// a unui SoC (**S**ystem-**o**n-a-**C**hip), spunem că este un GPU integrat. Exemple de SoC-uri cu IP de GPU integrat includ procesoarele x86 Intel/AMD, cât și majoritatea SoC-urilor pentru dispozitive mobile bazate pe arhitectura ARM (ex. Qualcomm Snapdragon). Un GPU integrat împarte ierarhia de memorie cu alte IP-uri (ex. controllere PCIe/USB/SATA/ETH).
-Filosofia din spatele arhitecturii este permiterea rularii unui numar foarte mare de threaduri. Acest lucru este facut posibil prin paralelismul existent la nivel hardware.
+De altfel, un GPU dedicat (discrete GPU) presupune valorificarea unei unui spațiu de memorie, mapat peste **VRAM** (**V**ideo **R**andom-**A**ccess **M**emory), cât și o magistrală PCIe/AGP8x/USB pentru comunicarea cu sistemul. Exemple de GPU-uri dedicate sunt seriile de plăci grafice Geforce (Nvidia) și Radeon (AMD).
-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.
+{{:asc:lab10:dgpu_igpu.png?direct&750|}}
-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.
+==== Aplicații ====
-{{:asc:lab11:thread.blocks.jpg?direct&360|{{thread.blocks.jpg|''Structura threadurilor in blocuri''}}
+Exemple de folosire de GPGPU-uri: prelucrări video, audio și de imagini, simulări ale fenomenelor fizice, finanțe, criptografie, design electronic (VLSI), mașini autonome.
-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]].
+  * https://www.nvidia.com/en-us/self-driving-cars/partners/bmw/
+  * https://blogs.nvidia.com/blog/2018/09/18/audi-unveils-e-tron-electric-suv/
-===== Ierarhia de memorie =====
+Rețele neurale - antrenare vs. inferență.
+  * https://www.forbes.com/sites/forbestechcouncil/2017/12/01/for-machine-learning-its-all-about-gpus/
+  * https://www.quora.com/Why-are-GPUs-well-suited-to-deep-learning
-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).
+Criptomonede - mining via hashing.
+  * https://coincentral.com/best-gpu-for-mining-2018-edition/
-{{:asc:lab11:mem.hierarchy.png?direct|Ierarhia memoriei in CUDA}}
+SmartTV, Smartphone - accelerare video, recunoaștere facială/audio.
-**Register File**
+Simulări fizice - NVIDIA Physx, Folding@Homel
-<code sh>
+  * https://blogs.nvidia.com/blog/2018/11/13/weather-predicted-sc18-gpu-hpc-jensen-huang/
-/* 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;
-</code>
-  *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**
+Prelucrări multimedia - filtre imagini GIMP/Photoshop.
-<code sh>
-/* fiecare work item salveaza un element */
+Alte domenii - arhivare (WinZip), encriptare.
-__local float lArray[lid] = data[gid];
+{{:asc:lab10:gpu-use.png?direct&700|}}
+===== Programarea GPGPU =====
+În cadrul unui sistem ce conține un GPU, procesorul general (CPU) coordonează execuția și este numit "HOST"; pe când unitatea care efectuează calculele (GPU) este numită "DEVICE".
+HOST-ul controlează toate schimbarile de stare în cadrul unui GPU, alocările/transferurile de memorie și evenimentele ce țin de sistemul de operare.
+O unitate GPU conține un procesor de comandă ("command processor") care citește comenzile scrise de către HOST (CPU) în anumite zone de memorie mapate spre access atât către unitatea GPU, cât și către CPU.
+În cazul GPU-urilor dedicate, o prelucrare de date necesită în prealabil un transfer din RAM către VRAM. Acest transfer se face printr-o magistrală (PCIe, AGP, USB…). Viteza de transfer RAM-VRAM via magistrală este inferioară vitezei de acces la RAM sau la VRAM. O potențială optimizare în cadrul acestui transfer ar fi intercalarea cu procesarea.
+În cazul GPU-urilor integrate transferul RAM<->"VRAM" presupune o mapare de memorie, de multe ori translatată printr-o operație de tip zero-copy.
+Programarea unui GPU se face printr-un API (Application Programming Interface). Cele mai cunoscute API-uri orientate către folosirea unui GPU ca coprocesor matematic sunt: CUDA, OpenCL, DirectCompute, OpenACC și Vulkan.
+Dezvoltarea de cod pentru laboratoarele de GPU se va face folosind CUDA.
+==== De ce CUDA? ====
+CUDA este un API introdus în 2006 de Nvidia pentru GPU-urile sale. În prezent CUDA este standardul //de facto// pentru folosirea unităților GPU în industrie și cercetare. Aceasta se datorează faptului că este o platformă stabilă cu multe facilități. Dacă o nouă versiune de CUDA introduce noi funcționalități, dar arhitectura nu le suportă, acestea sunt dezactivate.
+În mare toate GPU-urile oferite de Nvidia sunt suportate, diferența fiind la facilitățile suportate. Singura limitare, majoră, a platformei CUDA este că suportă **numai** unități de procesare de tip GPU de la Nvidia.
+Un standard alternativ la CUDA este **OpenCL**, suportat de Khronos și implementat de majoritatea producătorilor de GPU (inclusiv Nvidia ca o extensie la CUDA). OpenCL suferă de următoarele lipsuri:
+* suportul este fragmentat
+* standardul este mult mai restrictiv (decât CUDA)
+* mai complicat de scris programe (decât CUDA)
+===== Arhitectura Nvidia CUDA =====
+CUDA (**C**ompute **U**nified **D**evice **A**rchitecture) permite utilizarea limbajului C pentru programarea pe GPU-urile Nvidia cât și extensii pentru alte limbaje (exp. Python). Deoarece una din zonele țintă pentru CUDA este HPC (**H**igh **P**erformance **C**omputing), în care limbajul Fortran este foarte popular, PGI ofera un compilator de Fortran ce permite generarea de cod și pentru GPU-urile Nvidia. Există binding-uri pentru Java (jCuda), Python (PyCUDA) și .NET (CUDA.NET).
+{{:asc:lab7:cuda-software.png?800|}}
+Unitatea de bază în cadrul arhitecturii CUDA este numită **SM** (**S**treaming **M**ultiprocessor). Ea conține în funcție de generație un număr variabil de CUDA Cores sau **SP** (**S**tream Processors) - de regulă între 8SP și 128SP. Unitatea de bază în scheduling este denumită "warp" și este alcatuită din 32 de thread-uri. Vom aborda mai amănunțit arhitectura CUDA în laboratorul următor.
+==== Compute capability ====
+Versiunea de [[https://docs.nvidia.com/deploy/cuda-compatibility/index.html
+|"compute capability"]] a unui SM, are formatul X.Y, unde X este versiunea majoră, pe când Y este versiunea minoră. Partea majoră identifică generația din care face parte arhitectura.
+Partea minoră identifică diferențe incrementale în arhitectură și posibile noi funcționalități.
+Știind versiunea majoră și cea minoră cunoaștem facilitățile hardware oferite de către arhitectură.
+O listă a GPU-urilor NVIDIA și versiunile lor majore/minore se regăsește [[https://developer.nvidia.com/cuda-gpus|aici]].
+===== Programarea in CUDA =====
+CUDA extinde limbajul C prin faptul că permite unui programator să definească funcții C, denumite //kernels//, care urmează a fi execute în paralel de N thread-uri CUDA. Scopul este de a abstractiza arhitectura GPU astfel încat partea de scheduling cât și gestiunea resurselor se face de catre stiva software CUDA împreună cu suportul hardware. Figura de mai jos denotă distribuirea thread-urilor către două arhitecturi partiționate diferit.
+Un kernel se definește folosind specificatorul ''__global__'' iar execuția sa se face printr-o configurație de execuție folosind <nowiki> <<<...>>> </nowiki>. Configurația de execuție denotă numarul de blocuri și numărul de thread-uri dintr-un block. Fiecare thread astfel poate fi identificat unic prin ''blockIdx'' și ''threadIdx''.
+{{:asc:lab7:cuda-scalability.png?640|}}
+Mai jos avem definit un kernel, ''vector_add'', care are ca argumente pointeri de tip ''float'', respectiv ''size_t''. Acesta calculează $ f(x) = 2x + 1/(x + 1) $, pentru fiecare element din vector. Numărul total de thread-uri este dimensiunea vectorului.
+<code C>
+__global__ void vector_add(const float *a, float *b, const size_t n) {
+   // Compute the global element index this thread should process
+   unsigned int i = threadIdx.x + blockDim.x * blockIdx.x;
+   // Avoid accessing out of bounds elements
+   if (i < n) {
+      b[i] = 2.0 * a[i] + 1.0 / (a[i] + 1.0);
+   }
+}
 </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
-**Shared Memory**
+Configurația de execuție denotă maparea între date și instrucțiuni. În funcția de kernel, se definește setul de instrucțiuni ce se va executa repetat pe date. Mai jos ''vector_add'' este lansat în execuție cu N thread-uri (''blocks_no'' x ''block_size'') organizate câte ''block_size'' thread-uri per bloc.
-<code sh>
-/* elementele sunt salvate la nivel de bloc */
+<code C>
-__shared__ int s[64];
+// Launch the kernel
+vector_add<<<blocks_no, block_size>>>(device_array_a, device_array_b, num_elements);
 </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**
+==== HelloWorld CUDA ====
-<code sh>
-__const float pi = 3.14f
+<code C>
+#include <stdio.h>
+__global__ void kernel_example(int value) {
+   /**
+   * This is a kernel; a kernel is a piece of code that
+   * will be executed by each thread from each block in
+   * the GPU device.
+   */
+   printf("[GPU] Hello from the GPU!\n");
+   printf("[GPU] The value is %d\n", value);
+   printf("[GPU] blockDim=%d, blockId=%d, threadIdx=%d\n",blockDim.x, blockIdx.x, threadIdx.x);
+}
+int main(void) {
+   /**
+   * Here, we declare and/or initialize different values or we
+   * can call different functions (as in every C/C++ program);
+   * In our case, here we also initialize the buffers, copy
+   * local data to the device buffers, etc (you'll see more about
+   * this in the following exercises).
+   */
+   int nDevices;
+   printf("[HOST] Hello from the host!\n");
+   /**
+   * Get the number of compute-capable devices. See more info
+   * about this function in the Cuda Toolkit Documentation.
+   */
+   cudaGetDeviceCount(&nDevices);
+   printf("[HOST] You have %d CUDA-capable GPU(s)\n", nDevices);
+   /**
+   * Launching the above kernel with a single block, each block
+   * with a single thread. The syncrhonize and the checking functions
+   * assures that everything works as expected.
+   */
+   kernel_example<<<1,1>>>(25);
+   cudaDeviceSynchronize();
+   /**
+   * Here we can also deallocate the allocated memory for the device
+   */
+   return 0;
+}
 </code>
-  * 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**
+===== Aplicatie compute CUDA =====
-<code sh>
-__kernel void process(__global float* data){ ... }
+O aplicatie CUDA are ca scop executia de cod pe GPU-uri NVIDIA CUDA.
+In cadrul laboratoarelor partea de CPU (host) va fi folosita exclusiv pentru managementul executiei partii de GPU (device).
+Aplicatiile vor viza executia folosind un singur GPU NVIDIA CUDA.
+==== 0. Definire functie kernel ====
+In codul prezentat mai jos, functia vector_add este marcata cu "__global__" si va fi compilata de catre [[https://docs.nvidia.com/cuda/pdf/CUDA_Compiler_Driver_NVCC.pdf| CUDA NVCC compiler]] pentru GPU-ul de pe sistem (in cazul cozii xl va fi NVIDIA Pascal P100).
+<code C>
+/**
+ * This kernel computes the function f(x) = 2x + 1/(x + 1) for each
+ * element in the given array.
+ */
+__global__ void vector_add(const float *a, float *b, const size_t n) {
+  // Compute the global element index this thread should process
+  unsigned int i = threadIdx.x + blockDim.x * blockIdx.x;
+  // Avoid accessing out of bounds elements
+  if (i < n) {
+    b[i] = 2.0 * a[i] + 1.0 / (a[i] + 1.0);
+  }
+}
 </code>
-  * 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)**
+==== 1. Definire zone de memorie host si device ====
-  * 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
+Din punct de vedere hardware, partea de host (CPU) are ca memorie principala RAM (chip-uri memorie instalate pe placa de baza via slot-uri memorie) iar partea de device (GPU) are VRAM (chip-uri de memorie prezente pe placa video). Cand vorbim de memoria host (CPU) ne referim la RAM, iar in cazul memoriei device (GPU) la VRAM.
-<code sh>
+La versiunile mai recente de CUDA, folosind limbajul C/C++, un pointer face referire la spatiul virtual care este unificat pentru host (CPU) si device (GPU). Adresele virtuale insa sunt translatate catre adrese fizice ce rezida ori in memoria RAM (CPU) ori in memoria VRAM (GPU). Astfel este important cum alocam memoria (fie cu malloc pentru CPU sau cudaMalloc pentru GPU) si respectiv sa facem cu atentie transferurile de memorie intre zonele virtuale definite (de la CPU la GPU si respectiv de la GPU la CPU).
-Device 0: "Tesla K40m"
-  CUDA Driver Version / Runtime Version          9.1 / 9.1
+<code C>
-  CUDA Capability Major/Minor version number:    3.5
+// Declare variable to represent ~1M float values and
-  Total amount of global memory:                 11441 MBytes (11996954624 bytes)
+// computes the amount of bytes necessary to store them
-  (15) Multiprocessors, (192) CUDA Cores/MP:     2880 CUDA Cores
+const int num_elements = 1 << 16;
-  GPU Max Clock rate:                            745 MHz (0.75 GHz)
+const int num_bytes = num_elements * sizeof(float);
-  Memory Clock rate:                             3004 Mhz
-  Memory Bus Width:                              384-bit
+// Declaring the 'host arrays': a host array is the classical
-  L2 Cache Size:                                 1572864 bytes
+// array (static or dynamically allocated) we worked before.
-  Maximum Texture Dimension Size (x,y,z)         1D=(65536), 2D=(65536, 65536), 3D=(4096, 4096, 4096)
+float *host_array_a = 0;
-  Maximum Layered 1D Texture Size, (num) layers  1D=(16384), 2048 layers
+float *host_array_b = 0;
-  Maximum Layered 2D Texture Size, (num) layers  2D=(16384, 16384), 2048 layers
-  Total amount of constant memory:               65536 bytes
+// Declaring the 'device array': this array is the equivalent
-  Total amount of shared memory per block:       49152 bytes
+// of classical array from C, but specially designed for the GPU
-  Total number of registers available per block: 65536
+// devices; we declare it in the same manner, but the allocation
-  Warp size:                                     32
+// process is going to be different
-  Maximum number of threads per multiprocessor:  2048
+float *device_array_a = 0;
-  Maximum number of threads per block:           1024
+float *device_array_b = 0;
-  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
 </code>
-Caracteristici GPU M2070 (coada dp), via query device properties CUDA
+==== 2. Alocare memorie host (CPU) ====
-<code sh>
+Functia malloc va intoarce o adresa virtuala ce va avea corespondent o adresa fizica din RAM.
-Device 0: "Tesla M2070"
-  CUDA Driver Version / Runtime Version          9.1 / 9.1
+<code C>
-  CUDA Capability Major/Minor version number:    2.0
+// Allocating the host array
-  Total amount of global memory:                 5302 MBytes (5559156736 bytes)
+host_array_a = (float *) malloc(num_bytes);
-  (14) Multiprocessors, ( 32) CUDA Cores/MP:     448 CUDA Cores
+host_array_b = (float *) malloc(num_bytes);
-  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
 </code>
-Caracteristici GPU P100 (coada xl), via query device properties CUDA
+==== 3. Alocare memorie device (GPU) ====
-<code sh>
+Functia cudaMalloc va intoarce o adresa virtuala ce va avea corespondent o adresa fizica din VRAM.
-Device 1: "Tesla P100-PCIE-16GB"
-  CUDA Driver Version / Runtime Version          12.2 / 11.4
+<code C>
-  CUDA Capability Major/Minor version number:    6.0
+// Allocating the device's array; notice that we use a special
-  Total amount of global memory:                 16276 MBytes (17066885120 bytes)
+// function named cudaMalloc that takes the reference of the
-  (056) Multiprocessors, (064) CUDA Cores/MP:    3584 CUDA Cores
+// pointer declared above and the number of bytes.
-  GPU Max Clock rate:                            1329 MHz (1.33 GHz)
+cudaMalloc((void **) &device_array_a, num_bytes);
-  Memory Clock rate:                             715 Mhz
+cudaMalloc((void **) &device_array_b, num_bytes);
-  Memory Bus Width:                              4096-bit
-  L2 Cache Size:                                 4194304 bytes
+// If any memory allocation failed, report an error message
-  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
+if (host_array_a == 0 || host_array_b == 0|| device_array_a == 0 || device_array_b == 0) {
-  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
+  printf("[HOST] Couldn't allocate memory\n");
-  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
+  return 1;
-  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
+==== 4. Initializare memorie host (CPU) si copiere pe device (GPU) ====
-<code sh>
+{{:asc:lab10:cpu_to_gpu.png?720|}}
-Device 0: "NVIDIA A100-PCIE-40GB"
-  CUDA Driver Version / Runtime Version          12.4 / 11.4
+<code C>
-  CUDA Capability Major/Minor version number:    8.0
+// Initialize the host array by populating it with float values
-  Total amount of global memory:                 40326 MBytes (42285268992 bytes)
+for (int i = 0; i < num_elements; ++i) {
-  (108) Multiprocessors, (064) CUDA Cores/MP:    6912 CUDA Cores
+  host_array_a[i] = (float) i;
-  GPU Max Clock rate:                            1410 MHz (1.41 GHz)
+}
-  Memory Clock rate:                             1215 Mhz
-  Memory Bus Width:                              5120-bit
+// Copying the host array to the device memory space; notice the
-  L2 Cache Size:                                 41943040 bytes
+// parameters of the cudaMemcpy function; the function default
-  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
+// signature is cudaMemcpy(dest, src, bytes, flag) where
-  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
+// the flag specifies the transfer type.
-  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
+//
-  Total amount of constant memory:               65536 bytes
+// host -> device: cudaMemcpyHostToDevice
-  Total amount of shared memory per block:       49152 bytes
+// device -> host: cudaMemcpyDeviceToHost
-  Total shared memory per multiprocessor:        167936 bytes
+// device -> device: cudaMemcpyDeviceToDevice
-  Total number of registers available per block: 65536
+cudaMemcpy(device_array_a, host_array_a, num_bytes, cudaMemcpyHostToDevice);
-  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 =====
+==== 5. Executie kernel ====
-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.
+{{:asc:lab10:exec_gpu.png?720|}}
-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.
+<code C>
+// Compute the parameters necessary to run the kernel: the number
+// of blocks and the number of threads per block; also, deal with
+// a possible partial final block
+const size_t block_size = 256;
+size_t blocks_no = num_elements / block_size;
-<note important>
+if (num_elements % block_size)
-Conflictele de access la bancuri de memorie (cache) pot reduce semnificativ performanta.
+  ++blocks_no;
-</note>
-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.
+// Launch the kernel
-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):
+vector_add<<<blocks_no, block_size>>>(device_array_a, device_array_b, num_elements);
+cudaDeviceSynchronize();
+</code>
-<code sh>
-__global__ void func(...) {  // __kernel void func(...)
+==== 6. Copiere date inapoi de la device (GPU) catre host (CPU) ====
-...
-   __shared__ int *array; // __local int *array;
+<code C>
-   x = array[threadIdx.x]; // x = array[get_local_id(0)]; => performanta 100%, 0 bank conflicts
+// Copy the result back to the host memory space
-   x = array[threadIdx.x + 1]; // x = array[get_local_id(0) + 1]; => performanta 100%, 0 bank conflicts
+cudaMemcpy(host_array_b, device_array_b, num_bytes, cudaMemcpyDeviceToHost);
-   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
+// Print out the first 10 results
-...
+for (int i = 0; i < 10; ++i) {
+  printf("Result %d: 2 * %1.1f + 1.0/(%1.1f + 1.0)= %1.3f\n",
+  i, host_array_a[i], host_array_a[i], host_array_b[i]);
 }
 </code>
-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.
+{{:asc:lab10:gpu_to_cpu.png?720|}}
-====== Aplicații ======
+==== 7. Cleanup ====
-Urmăriți instrucțiunile de pe [[https://gitlab.cs.pub.ro/asc/asc-public/-/tree/master/labs/cuda/arch|GitLab]]
+<code C>
+// Deallocate memory
+free(host_array_a);
+free(host_array_b);
+cudaFree(device_array_a);
+cudaFree(device_array_b);
+</code>
+===== Aplicații =====
+{{:asc:lab7:cuda-nvcc.png?480|}}
 <note important>
-Recomandăm sa va delogati mereu de pe serverele din cluster dupa terminarea sesiunii, utilizand comanda ''exit''
+Pentru editarea fișierelor pe cluster, recomandam sa va montați sistemul de pe fep8 pe mașină locală. Pașii sunt detaliați mai jos - multumiri lui Radu Millo pentru redactare.
-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ă.
+Tutorial chei ssh: https://www.ssh.com/academy/ssh/keygen
-Daca nu veti face aceasta delogare, veti putea ajunge in situatia in care sa nu va mai puteti loga pe nodurile din cluster.
+Pe fep8, din directorul vostru home (''%%~%%''), dați comanda:
+<code sh>
+mkdir asc
+cd asc
+pwd # acesta va returna <cale_absolută_director_asc_fep8>
+</code>
+Tutorial montare filesystem din fep pe local - comenzi date pe local:
+<code sh>
+mkdir asc # puteți da comanda din directorul vostru home
+cd asc
+pwd # acesta va returna <cale_absolută_director_asc_local>
+sudo chown -R <user> asc # numele user-ului vostru de pe local
+decomentăm linia 'user_allow_other' din /etc/fuse.conf
+sshfs -o allow_other <user.moodle>@fep8.grid.pub.ro:<cale_absolută_director_asc_fep8> <cale_absolută_director_asc_local>
+</code>
 </note>
+Urmăriți instrucțiunile de pe GitLab [[https://gitlab.cs.pub.ro/asc/asc-public/-/tree/master/labs/cuda/intro|GitLab]].
 ===== Resurse =====
 <hidden>
-{{:asc:lab8:sol:lab8_sol.tar.gz|Soluție Laborator 8}}
+{{:asc:lab7:sol:lab7_sol.tar.gz|Soluție Laborator 7}}
 </hidden>
-  * Responsabili laborator: Matei Barbu
+{{:asc:lab7:asc_lab7.pdf|Enunt Laborator 4}}
+  * Responsabili laborator: Costin Carabaș, Tudor Calafeteanu, Grigore Lupescu, Mihnea Mitroi, Irinel Gul, Alex 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]]
-    * [[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]]