Show page

Differences

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

--- asc:laboratoare:07 [2024/04/15 00:13]
mihnea.mitroi
+++ asc:laboratoare:07 [2026/02/23 18:47] (current)
giorgiana.vlasceanu
@@ Line 1: / Line 1: @@
-====== Laboratorul 07 - CUDA ======
+====== Laboratorul 07 - Advanced CUDA ======
-===== Teorie =====
+===== Spatiu unificat memorie =====
-Scopul acestui laborator este aprofundarea și îmbinarea noțiunilor pe care le-ați învățat deja la APD și în laburile trecute (poate și la IA).
+De la [[http://developer.download.nvidia.com/compute/cuda/6_0/rel/docs/CUDA_Toolkit_Release_Notes.pdf| CUDA 6.0]], NVIDIA a schimbat semnificativ modelul de programare prin facilitarea comunicarii unitatii CPU (host) cu unitatea GPU (device), in mod transparent prin acelasi set de adrese de memorie virtuale. Astfel exista posibilitatea ca prin acelasi pointer de memorie sa se scrie date atat de catre CPU cat si de catre GPU. Evident transferurile de memorie au loc intre spatii diferite de adresare (ex RAM vs VRAM), dar acest lucru se intampla transparent la nivel de aplicatie CUDA / pentru programator.
-Vă încurajăm să profitați de acest laborator ca să vă clarificați posibile nelămuriri din laburile trecute. Sau, de ce nu, să vă șlefuiți abilitatea de a optimiza programe CUDA dacă terminați mai devreme. :P
+{{:asc:lab9:nv-unified.png?640|NVIDIA Unified Memory}}
-==== Odd-Even Transposition Sort ====
+Mai jos avem un exemplu de folosire a memoriei unificate. Singura diferenta fata de alocarea pe CPU/HOST este ca memoria trebuie alocata cu cudaMallocManaged si dealocata cu cudaFree.
-Bubble Sort este un algoritm secvențial în care se parcurge șirul de sortat element cu element, comparând elementul curent cu vecinul din dreapta. Dacă numărul din dreapta este mai mic, se realizează o interschimbare între elementul curent și cel din dreapta sa.
+<code C>
+#include <iostream>
+#include <math.h>
+// CUDA kernel to add elements of two arrays
+__global__
+void add(int n, float *x, float *y)
+{
+  int index = blockIdx.x * blockDim.x + threadIdx.x;
+  int stride = blockDim.x * gridDim.x;
+  for (int i = index; i < n; i += stride)
+    y[i] = x[i] + y[i];
+}
+int main(void)
+{
+  int N = 1<<20;
+  float *x, *y;
+  // Allocate Unified Memory -- accessible from CPU or GPU
+  cudaMallocManaged(&x, N*sizeof(float));
+  cudaMallocManaged(&y, N*sizeof(float));
+  // initialize x and y arrays on the host
+  for (int i = 0; i < N; i++) {
+    x[i] = 1.0f;
+    y[i] = 2.0f;
+  }
+  // Launch kernel on 1M elements on the GPU
+  int blockSize = 256;
+  int numBlocks = (N + blockSize - 1) / blockSize;
+  add<<<numBlocks, blockSize>>>(N, x, y);
+  // Wait for GPU to finish before accessing on host
+  cudaDeviceSynchronize();
+  // Check for errors (all values should be 3.0f)
+  float maxError = 0.0f;
+  for (int i = 0; i < N; i++)
+    maxError = fmax(maxError, fabs(y[i]-3.0f));
+  std::cout << "Max error: " << maxError << std::endl;
+  // Free memory
+  cudaFree(x);
+  cudaFree(y);
+  return 0;
+}
+</code>
-Operațiile pe elemente adiacente nu se pot realiza simultan, pentru că se poate ajunge la un race condition. Din acest motiv, un mod de a paraleliza Bubble Sort este Odd-Even Transposition Sort, un algoritm in 2 două faze. În faza pară, elementele de pe poziții pare din șirul de sortat sunt comparate (și eventual interschimbate) cu vecinii din dreapta. După ce se termină faza pară (adică după ce toate elementele pare au fost procesate), urmează faza impară, în care elementele impare sunt analizate și comparate cu vecinii din dreapta.
+===== Operatii atomice CUDA =====
-{{:asc:lab7:oets.png?200 | OETS}}
+CUDA ofera acces la multiple operatii atomice tip citire-modificare-scriere. Acestea presupun serializarea accesului in contextul mai multor thread-uri. Functiile sunt limitate la anumite tipuri de date:
+  - int
+  - unsigned int
+  - unsigned long long int
+  - float
+  - double
-==== Merge Sort ====
+Exemple de functii atomice:
+  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicadd|atomicAdd]]
+  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicsub|atomicSub]]
+  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicexch|atomicExch]]
+  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicmin|atomicMin]]
+  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicmax|atomicMax]]
+  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicinc|atomicInc]]
+  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicdec|atomicDec]]
+  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicand|atomicAnd]]
+  - [[https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomicor|atomicOr]]
-Merge sort (sau sortarea prin interclasare) este un algoritm de sortare de tip divide et impera care presupune următorii pași generali:
+<note important>
-  * se împarte șirul de N elemente de sortat în N șiruri de lungime 1
+A se consulta cu atentie documentatia CUDA inainte de folosirea unei operatii atomice (legat de contextul in care se aplica, cum opereaza, limitari etc).
-  * se aplica operația de interclasare ("merge") între câte două astfel de șiruri de lungime 1, rezultând N/2 șiruri sortate de lungime 2
+</note>
-  * se repetă pașii de mai sus realizând interclasări între șiruri din ce în ce mai mari, până se ajunge la un șir sortat de N elemente.
-Pentru a paraleliza acest algoritm, putem observa că operațiile de interclasare de la fiecare pas se pot realiza în paralel. Totuși, operațiile de "merge" de la fiecare pas trebuie terminate în totalitate înainte de a trece la următorul pas, deci avem nevoie de o barieră (sau un mecanism similar) după fiecare pas de interclasare. Se poate observa că gradul de paralelism de la un pas de interclasări este din ce în ce mai mic pe măsură ce avansăm în algoritm, pentru că numărul de operații de "merge" de la fiecare pas scade.
+In codul de mai jos se lanseaza un kernel concurrentRW in configuratie numBlocks=8 fiecare cu cate 10 thread-uri.
-{{:asc:lab7:mergesort.png?300 | Mergesort}}
+<code C>
+#include <iostream>
-==== Gaussian Blur ====
+#define NUM_ELEM        8
+#define NUM_THREADS     10
-În prelucrarea imaginilor, blurarea gaussiană este rezultatul estomparii unei imagini printr-o funcție gaussiană.
+using namespace std;
-Este un efect larg utilizat în software-ul grafic, în mod tipic pentru a reduce zgomotul imaginii și pentru a reduce detaliile. Efectul vizual al acestei tehnici de estompare este un blur neted care seamănă cu cel obținut privind imaginea printr-un ecran translucid. Estomparea gaussiană este de asemenea folosită ca o etapă de prelucrare în algoritmi de viziune artificială pentru a îmbunătăți structurile imaginii la diferite scări.
+__global__ void concurrentRW(int *data) {
+...
+}
-Aplicarea unei estompări gaussiene asupra unei imagini este echivalentă cu convoluția imaginii cu o funcție gaussiană. Transformarea (funcția) de aplicat fiecărui pixel din imagine este urmatoarea:
+int main(int argc, char *argv[]) {
-\[ G(x) = \frac{1}{\sqrt{2\pi \sigma^2}} e^{-\frac{x^2}{2\sigma^2}} \]
+    int* data = NULL;
+    bool errorsDetected = false;
-{{:asc:lab7:gaussian-blur.png?400 | Gaussian Blur}}
+    cudaMallocManaged(&data, NUM_ELEM * sizeof(unsigned long long int));
+    if (data == 0) {
+        cout << "[HOST] Couldn't allocate memory\n";
+        return 1;
+    }
+    // init all elements to 0
+    cudaMemset(data, 0, NUM_ELEM * sizeof(unsigned long long int));
+    // launch kernel writes
+    concurrentRW<<<NUM_ELEM, NUM_THREADS>>>(data);
+    cudaDeviceSynchronize();
+    if (cudaSuccess != cudaGetLastError()) {
+        return 1;
+    }
+    for(int i = 0; i < NUM_ELEM; i++) {
+        cout << i << ". " << data[i] << endl;
+        if(data[i] != (NUM_THREADS * (NUM_THREADS - 1) / 2)) {
+            errorsDetected = true;
+        }
+    }
+    if(errorsDetected) {
+        cout << "Errors detected" << endl;
+    } else {
+        cout << "OK" << endl;
+    }
+    return 0;
+}
+</code>
+Functia concurrentRW citeste valoarea de la adresa data[blockIdx.x], o incrementeaza cu threadIdx.x si apoi o scrie.
+In acest caz avem 10 thread-uri care fac operatii citire/scriere la aceeasi adresa, deci un comportament nedefinit.
+<code C>
+__global__ void concurrentRW(int *data) {
+    // NUM_THREADS try to read and write at same location
+    data[blockIdx.x] = data[blockIdx.x] + threadIdx.x;
+}
+</code>
+Exemplu rezultat:
+<code sh>
+. 9
+. 9
+. 9
+. 9
+. 9
+. 9
+. 9
+. 9
+Errors detected
+</code>
+Corect ar fi folosirea functiei atomicAdd pentru a serializa accesul.
+<code C>
+__global__ void concurrentRW(int *data) {
+    // NUM_THREADS try to read and write at same location
+    atomicAdd(&data[blockIdx.x], threadIdx.x);
+}
+</code>
+Rezultatul rularii este:
+<code sh>
+. 45
+. 45
+. 45
+. 45
+. 45
+. 45
+. 45
+. 45
+OK
+</code>
+==== Operatii atomice system wide ====
+Unitatile GPU ce au Compute capability 6.x permit largirea scopului operatiilor atomice. De exemplu atomicAdd_system garanteaza ca operatia este atomica cand atat thread-urile de pe unitatea GPU cat si cele de pe unitatea CPU incearca sa acceseze datele. Mai jos avem un exemplu de folosire al functiei atomicAdd_system.
+<code C>
+__global__ void mykernel(int *addr) {
+  atomicAdd_system(addr, 10);       // only available on devices with compute capability 6.x
+}
+void foo() {
+  int *addr;
+  cudaMallocManaged(&addr, 4);
+  *addr = 0;
+   mykernel<<<...>>>(addr);
+   __sync_fetch_and_add(addr, 10);  // CPU atomic operation
+}
+</code>
+===== Operatii asincrone CUDA =====
+In CUDA, urmatoarele operatii sunt definite ca fiind independente si pot fi executate concurent:
+  - Calcule pe unitatea host
+  - Calcule pe unitatea device
+  - Transfer memorie host -> device
+  - Transfer memorie device -> host
+  - Transfer memorie device -> device
+Nivelul de concurenta o sa depinda si de capabilitatea unitatilor GPU (compute capability).
+In continuare vom explora mai multe scenarii de executie concurenta a operatiilor descrise.
+==== Executie asincrona Host si Device ====
+Folosind apeluri asincrone, operatiile de executie catre device sunt puse in coada avand controlul intors catre host instant. Astfel unitatea host poate continua executia fara sa fie blocata in asteptarea altor task-uri.
+Urmatoarele operatii sunt asincrone relativ la host:
+  - Lansari de kernel
+  - Copieri in cadrul spatiului de memorie al unui device
+  - Copiere memorie host -> device, avand < 64 KB
+  - Copiere memorie host -> device, avand functii cu sufix Async
+  - Functii memorie set (setare / initializare de memorie la o valoare)
+Pentru a face debug unor scenarii de executie asincrona se poate dezactiva complet executia asincrona setand variabila de mediu CUDA_LAUNCH_BLOCKING la 1. Executia de kernels este sincrona cand se ruleaza cu un profiler (Nsight, Visual Profiler).
+==== Fluxuri nonimplicite ====
+Pentru a folosi cudaMemcpyAsync, este necesar lucrul cu fluxuri nonimplicite (non-default streams), care, in C/C++ pot fi declarate, create si distruse in partea de cod de pe host (CPU) in urmatorul fel:
+<code C>
+cudaStream_t stream1;
+cudaError_t result;
+result = cudaStreamCreate(&stream1);
+result = cudaStreamDestroy(stream1);
+</code>
+Odata creat un astfel de flux, el poate fi utilizat in procesul de copiere a memoriei host -> device astfel:
+<code C>
+// num_bytes = N * sizeof (type_a);
+result = cudaMemcpyAsync(d_a, a, num_bytes, cudaMemcpyHostToDevice, stream1);
+</code>
+Pentru a emite un kernel către un flux nonimplicit, specificăm identificatorul fluxului ca al patrulea parametru de configurare a execuției. Se observă și un al treilea parametru de configurare a execuției, care este folosit pentru a aloca memorie partajată (shared memory) device-ului (GPU-ului), utilizându-se 0 dacă nu se dorește acest aspect.
+<code C>
+increment<<<1,N,0,stream1>>>(d_a);
+</code>
+==== Executie asincrona programe kernel ====
+Arhitecturile cu compute capability 2.x sau mai nou, pot executa in paralel instante de kernel diferite. Aceste unitati de executie o sa aibe proprietate concurrentKernels setata la 1 (se face query la device properties inainte). Numarul maxim de lansari asincrone de kernele diferite este dependent de arhitectura (se verifica in functie de compute capability). Singura restrictie este ca programele kernel sa fie in acelasi context.
+==== Executie si transfer date asincron  ====
+Anumite device-uri pot executa un transfer asincron memorie alaturi de o executie de kernel. Acest lucru este dependent de compute capability si se poate verifica in device property asyncEngineCount.
+{{:asc:lab9:cuda_async.png?900|}}
+De asemenea, se pot face transferuri de memorie intra-device simultan cu executia de kernel cand atat device property concurrentKernels, cat si asyncEngineCount sunt 1.
+{{:asc:lab9:cuda_async_2.png?900|}}
+===== Dynamic Paralellism  =====
+Paralelismul dinamic consta in posibilitatea de a lansa programe kernel din thread-urile ce ruleaza pe device/GPU. In alte cuvinte, unitatea GPU poate sa isi  atribuie noi task-uri/thread-uri fara interventia unitatii host/CPU. Aceasta manifestare este utila in problemele unde maparea threaduri<->date nu este simpla/triviala. De exemplu, in situatia unde unele thread-uri ar avea prea putin de lucru, iar altele prea mult (imaginea de mai jos, simulare fluide) - o situatia debalansata computational.
+{{:asc:lab9:dynamic-paralellism.png?560|Fluid simulation}}
+Cerintele pentru paralelism dinamic sunt CUDA 5.0 ca Toolkit si respectiv Compute Capability 3.5. O lista cu GPU-uri NVIDIA si Compute Capability se regaseste [[https://developer.nvidia.com/cuda-gpus|aici]].
 ===== Exercitii =====
-Urmăriți instrucțiunile de pe [[https://gitlab.cs.pub.ro/asc/asc-public/-/tree/master/labs/cuda/practice|GitLab]].
-<note important>
+Urmăriți instrucțiunile de pe [[https://gitlab.cs.pub.ro/asc/asc-public/-/tree/master/labs/cuda/advanced|GitLab]].
-Puteți găsi materiale ajutătoare în cadrul [[https://mobylab.docs.crescdi.pub.ro/docs/parallelAndDistributed/laboratory3|laboratorului 3 de la APD]] și în videoclipul [[https://www.youtube.com/watch?v=KuXjwB4LzSA&t=512s|Gaussian Blur]].
-</note>
 <note important>
@@ Line 51: / Line 280: @@
 Daca nu veti face aceasta delogare, veti putea ajunge in situatia in care sa nu va mai puteti loga pe nodurile din cluster.
 </note>
-* Responsabili laborator: Matei Barbu, Mihnea Mitroi
+===== Resurse =====
+<hidden>
+{{:asc:lab9:sol:lab9_sol.tar.gz|Soluție Laborator 6}}
+</hidden>
+/* {{:asc:lab9:sol:lab9_sol.zip|Solutie Laborator 6}} */
+/* {{:asc:lab6:asc_lab9.pdf|Enunt Laborator 6}} */
+  * Responsabili laborator: Matei Barbu, Alexandru Bala
 ==== Referinte ====
@@ Line 73: / Line 313: @@
     * [[https://devblogs.nvidia.com/introduction-cuda-dynamic-parallelism/|CUDA Dynamic Parallelism]]
     * [[https://developer.nvidia.com/blog/how-overlap-data-transfers-cuda-cc/|How to Overlap Data Transfers in CUDA C/C++]]