Show page

Differences

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

--- asc:laboratoare:07 [2023/04/28 10:46]
rares.folea [Programarea in CUDA]
+++ asc:laboratoare:07 [2026/02/23 18:47] (current)
giorgiana.vlasceanu
@@ Line 1: / Line 1: @@
-====== Laboratorul 07 - Arhitecturi de tip GPGPU ======
+====== Laboratorul 07 - Advanced CUDA ======
-===== Intro =====
+===== Spatiu unificat memorie =====
-Procesorul grafic (GPU - graphics processing unit) reprezinta un circuit electronic specializat in crearea si manipularea imaginilor trimise catre o unitate de display (e.g. monitor). Termenul GPGPU (general purpose graphics processing unit) denota un procesor grafic cu o flexibilitate ridicata de programare, capabil de a rezolva si probleme generale. In executie, o arhitectura de tip GPU foloseste paradigma SIMD (single instruction multiple data, taxonomia Flynn), ceea ce presupune, schimb rapid de context intre thread-uri, planificarea in grupuri de thread-uri si orientare catre prelucrari masive de date. Procesorul grafic dispune si de un spatiu propriu de memorie (GPU dedicat -> VRAM, GPU integrat -> RAM).
+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.
-Unitatile tip GPU sunt potrivite pentru paralelismul de date, intensiv computationale. Datorita faptului ca aceleasi instructiuni sunt executate pentru fiecare element, nu sunt necesare mecanisme complexe pentru controlul fluxului. Ierarhia de memorie este simplificata comparativ cu cea a unui procesor x86/ARM. Deoarece calculele sunt intensive computational, latenta accesului la memorie poate fi ascunsa prin paralelism (massive multithreading, SIMT sau Single Instruction Multiple Threads) in locul folosirii extensive a memoriei cache.
+{{:asc:lab9:nv-unified.png?640|NVIDIA Unified Memory}}
-<note important>
+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.
-Nu orice algoritm paralel ruleaza optim pe o arhitectura GPGPU. De principiu probleme de tip SIMD sau MIMD se preteaza rularii pe GPU-uri.
-</note>
-In multe cazuri, termenul de GPGPU apare atunci cand unitatea GPU este folosita ca si coprocesor matematic. In ziua de azi, majoritatea unitatilor de tip GPU sunt si GPGPU. In ultimii ani folosirea unitatilor GPGPU a luat amploare. Acest lucru se datoreaza:
+<code C>
-  * diferentelor de putere de procesare bruta dintre CPU si GPU in favoarea acestora din urma
+#include <iostream>
-  * standardizarea de API-uri care usureaza munca programatorilor pentru a folosi GPU-ul
+#include <math.h>
-  * raspandirea aplicatiilor ce pot beneficia de pe urma paralelismului tip SIMD
-  * regasirea unitatilor GPU atat in unitatile computationale consumer (PC, Smartphone, TV etc) cat si cele industriale (Automotive, HPC etc).
+// 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>
-Principalii producatori de core-uri IP (intellectual property) tip GPU sunt:
+===== Operatii atomice CUDA =====
-  * 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
-Daca un IP de GPU este integrat pe aceeasi pastila de siliciu a unui SoC (system on chip), acesta se numeste GPU integrat (integrated GPU). Exemple de SoC-uri cu IP de GPU integrat includ procesoarele x86 Intel si Amd cat si majoritatea SoC-urilor pentru dispozitive mobile bazate pe arhitectura ARM (ex. Qualcomm Snapdragon). Un GPU integrat imparte mare parte din ierarhia de memorie cu alte IP-uri (ex core-uri ARM/x86, controller PCIe/USB/SATA/ETH). Pe de alta parte un GPU dedicat (discrete GPU) presupunea integrarea IP-ului de GPU pe o placa cu memorie dedicata (VRAM) cat si o magistrala PCIe/AGP8x/USB pentru comunicare cu sistemul. Exemple de GPU-uri dedicate sunt seriile de placi grafice Geforce (Nvidia) si Radeon (Amd).
+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
-{{:asc:lab10:dgpu_igpu.png?direct&750|}}
+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]]
-==== Aplicatii arhitecturi GPGPU ====
+<note important>
+A se consulta cu atentie documentatia CUDA inainte de folosirea unei operatii atomice (legat de contextul in care se aplica, cum opereaza, limitari etc).
+</note>
-Exemple de domenii ce folosesc procesare GPGPU: prelucrari video si de imagini, simulari de fizica, finante, dinamica fluidelor, criptografie, design electronic (VLSI). Exemple de aplicatii pentru GPGPU:
+In codul de mai jos se lanseaza un kernel concurrentRW in configuratie numBlocks=8 fiecare cu cate 10 thread-uri.
-Automotive - self driving cars (BMW, Continental etc)
-  * 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/
-Inteligenta artificala - antrenare retele neurale, inferenta
-  * 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
-Criptomonede - mining via hashing
-  * https://coincentral.com/best-gpu-for-mining-2018-edition/
-SmartTV, Smartphone - accelerare video, recunoastere faciala/audio
-Simulari fizice - NVIDIA Physx, Folding@Homel
-  * https://blogs.nvidia.com/blog/2018/11/13/weather-predicted-sc18-gpu-hpc-jensen-huang/
-Prelucrari multimedia - filtre imagini GIMP/Photoshop
-Alte domenii - arhivare (WinZip), encriptare
-{{:asc:lab10:gpu-use.png?direct&700|}}
+<code C>
+#include <iostream>
-===== Programarea GPGPU =====
+#define NUM_ELEM        8
+#define NUM_THREADS     10
-In cadrul unui sistem ce contine o unitate IP de tip GPU, procesorul general care coordoneaza executia este numit "HOST" (CPU) pe cand unitatea care efectueaza calculele este numita "DEVICE" (GPU). O unitate GPU contine un procesor de comanda ("command processor") care citeste comenzile scrise de catre HOST (CPU) in anumite zone din RAM mapate spre access atat catre unitatea GPU cat si catre unitatea CPU. Toate schimbarile de stare in cadrul unui GPU, alocarile/transferurile de memorie si evenimentele ce tin de sistemul de operare sunt controlate de catre CPU (HOST).
+using namespace std;
-In general, o prelucrare de date folosind unitatea GPU, necesita in prealabil un transfer din spatiul de memorie de la CPU catre spatiul de memorie de la GPU. In cazul unui procesor grafic dedicat acest transfer se face printr-o magistrala (PCIe, AGP, USB…). Viteza de transfer RAM-VRAM via magistrala este inferioara vitezei RAM sau VRAM. O potentiala optimizare in transferul RAM<->VRAM ar fi intercalarea cu procesarea. In cazul unui procesor integrate transferul RAM<->VRAM presupune o mapare de memorie, de multe ori translatata printr-o operatie de tip zero copy.
+__global__ void concurrentRW(int *data) {
+...
+}
-Programarea unui GPU se face printr-un API (Application Programming Interface). Cele mai cunoscute API-uri orientate catre folosirea unui GPU ca coprocesor matematic sunt: Cuda, OpenCL, DirectCompute, OpenACC, Vulkan. Dezvoltarea de cod pentru laboratoarele de GPU se va face folosind Cuda.
+int main(int argc, char *argv[]) {
+    int* data = NULL;
+    bool errorsDetected = false;
-//De ce CUDA ?//
+    cudaMallocManaged(&data, NUM_ELEM * sizeof(unsigned long long int));
+    if (data == 0) {
+        cout << "[HOST] Couldn't allocate memory\n";
+        return 1;
+    }
-CUDA este un API introdus in 2006 de catre NVIDIA pentru GPU-urile sale. In prezent CUDA este standardul de facto pentru folosirea unitatilor GPU in industrie si cercetare. Aceasta se datoreaza faptului ca este o platforma stabila cu multe facilitati. O noua versiune de CUDA introduce noi functiontionalitati dar acestea uneori necesita versiuni recente ale arhitecturilor fiind dezactivate daca nu exista suport hardware. O versiune noua de CUDA extinde versiunea mai veche - de exemplu versiunea CUDA 9.0 reprezinta in mare o extensie/update asupra versiunii CUDA 8.0. In mare toate GPU-urile oferite de NVIDIA sunt suportate, diferenta fiind la facilitatile suportate. Singura limitare majora a platformei CUDA este ca suporta numai unitati de procesare de tip GPU de la NVIDIA.
+    // init all elements to 0
+    cudaMemset(data, 0, NUM_ELEM * sizeof(unsigned long long int));
-Un standard alternativ la CUDA este OpenCL, suportata de Khronos ca standard si implementata de majoritatea producatorilor de GPU (inclusiv NVIDIA ca o extensie la CUDA). Problema majora la OpenCL este ca suportul este fragmentat si standardul este mult mai restrictiv decat CUDA si mai complicat de scris programe.
+    // launch kernel writes
+    concurrentRW<<<NUM_ELEM, NUM_THREADS>>>(data);
+    cudaDeviceSynchronize();
+    if (cudaSuccess != cudaGetLastError()) {
+        return 1;
+    }
-===== Arhitectura NVIDIA CUDA =====
+    for(int i = 0; i < NUM_ELEM; i++) {
+        cout << i << ". " << data[i] << endl;
+        if(data[i] != (NUM_THREADS * (NUM_THREADS - 1) / 2)) {
+            errorsDetected = true;
+        }
+    }
-Implementarea NVIDIA pentru GPGPU se numeste CUDA (Compute Unified Device Architecture) si permite utilizarea limbajului C pentru programarea pe GPU-urile proprii cat si extensii pentru alte limbaje (ex Python). Deoarece una din zonele tinta pentru CUDA este HPC (High Performance Computing), in care limbajul Fortran este foarte popular, PGI ofera un compilator de Fortran ce permite generarea de cod si pentru GPU-urile Nvidia. Exista binding-uri pana si pentru Java (jCuda), Python (PyCUDA) sau .NET (CUDA.NET). Framework-ul/arhitectura CUDA expune si API-ul de OpenCL prin intermediul caruia vom interactiona cu GPGPU-ul Nvidia Tesla disponibil pe dp.
+    if(errorsDetected) {
+        cout << "Errors detected" << endl;
-{{:asc:lab7:cuda-software.png?800|}}
+    } else {
+        cout << "OK" << endl;
-Arhitectura CUDA (toate GPU-urile, seriile Geforce (consumer), Tesla (HPC), Jetson (automotive)).
+    }
-Driver cu suport Windows, Linux, ce suporta atat CUDA API cat si OpenCL API.
-Framework/toolkit compilator cu suport CUDA/OpenCL API (nvcc), debugger/profiler (CUDA API only)
-Numeroase biblioteci si exemple CUDA/OpenCL API
-Unitatea de baza in cadrul arhitecturii CUDA este numita SM (Streaming Multiprocessor). Ea contine in functie de generatie un numar variabil de Cuda Cores sau SP (Stream Processors) - de regula intre 8SP si 128SP. Unitatea de baza in scheduling este denumita "warp" si alcatuita din 32 de thread-uri. Vom aborda mai amanuntit arhitectura CUDA in laboratorul urmator. Ultima versiune de CUDA 8.0 suport OpenCL 1.2.
-===== Compute capability =====
-Versiunea de "compute capability" a unui SM (Streaming Multiprocessor), in cadrul arhitecturii CUDA, este reprezentat de un format X.Y, unde X este versiunea majora pe cand Y este versiunea minora. Partea majora identifica generatia din care face parte arhitectura. Astfel revizia 7 denota arhitectura Volta, 6 este pentru arhitectura Pascal, 5 pentru arhitectura Maxwell, 3 pentru arhitectura Kepler, 2 pentru Fermi iar 1 pentru Tesla. Partea minora identifica diferente incrementale in arhitectura si posibile noi functionalitati. Stiind versiunea majora si cea minora cunoastem facilitatile hardware oferite de catre arhitectura. GPU-urile care au aceasi versiune suporta aceleasi capabilitati.
-O lista a GPU-urile NVIDIA si versiunile lor majore/minore se regaseste [[https://developer.nvidia.com/cuda-gpus|aici]].
-In cadrul cozii hpsl se regasesc GPU-uri [[https://www.nvidia.com/content/pdf/kepler/tesla-k40-active-board-spec-bd-06949-001_v03.pdf|Tesla K40M]], iar in cadrul cozii dp GPU-uri [[https://www.nvidia.com/docs/IO/43395/NV_DS_Tesla_C2050_C2070_jul10_lores.pdf|Tesla 2070]].
-===== Programarea in CUDA =====
-CUDA extinde limbajul C prin faptul ca permite unui programator sa defineasca functii C, denumite kernels, care urmeaza a fi execute de N ori in paralel de N thread-uri CUDA. Scopul este de a abstractiza arhitectura GPU astfel incat partea de scheduling cat si gestiunea resurselor se face de catre stack-ul software CUDA impreuna cu suportul hardware. Figura de mai jos denota distribuirea thread-urilor catre 2 arhitecturi partitionate diferit.
-Un kernel se defineste folosind specificatorul __global__ iar executia sa se face printr-o configuratia de executie folosind <nowiki> <<<...>>> </nowiki>. Configuratia de executie denota numarul de blocks si numarul de thread-uri dintr-un block. Fiecare thread astfel poate fi identificat unic prin blockIdx si threadIdx.
-{{:asc:lab7:cuda-scalability.png?640|}}
-Mai jos avem definit un kernel, vector_add, care are ca argumente pointers de tip float respectiv size_t. Acesta calculeaza f(x) = 2x + 1/(x + 1), pentru fiecare elemente din vector. Numarul 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
+    return 0;
-   if (i < n) {
-      b[i] = 2.0 * a[i] + 1.0 / (a[i] + 1.0);
-   }
 }
 </code>
-Configuratia de executie denota maparea intre date si instructiuni. In functia de kernel, se defineste setul de instructiuni ce se va executa repetat pe date. Mai jos vector_add este lansat in executie cu N thread-uri (blocks_no x block_size) organizate cate block_size thread-uri per block.
+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>
-// Launch the kernel
+__global__ void concurrentRW(int *data) {
-vector_add<<<blocks_no, block_size>>>(device_array_a, device_array_b, num_elements);
+    // NUM_THREADS try to read and write at same location
+    data[blockIdx.x] = data[blockIdx.x] + threadIdx.x;
+}
 </code>
-===== Aplicatie HelloWorld CUDA =====
+Exemplu rezultat:
+<code sh>
-<code C>
+. 9
-#include <stdio.h>
+. 9
+. 9
-__global__ void kernel_example(int value) {
+. 9
-   /**
+. 9
-   * This is a kernel; a kernel is a piece of code that
+. 9
-   * will be executed by each thread from each block in
+. 9
-   * the GPU device.
+. 9
-   */
+Errors detected
-   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>
-===== Aplicatie compute CUDA =====
+Corect ar fi folosirea functiei atomicAdd pentru a serializa accesul.
-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).
-Aplicatiilor 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 hpsl va fi NVIDIA Tesla K40M).
 <code C>
-/**
+__global__ void concurrentRW(int *data) {
- * This kernel computes the function f(x) = 2x + 1/(x + 1) for each
+    // NUM_THREADS try to read and write at same location
- * element in the given array.
+    atomicAdd(&data[blockIdx.x], threadIdx.x);
- */
-__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>
-==== 1. Definire zone de memorie host si device ====
+Rezultatul rularii este:
+<code sh>
-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.
+. 45
+. 45
-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).
+. 45
+. 45
-<code C>
+. 45
-// Declare variable to represent ~1M float values and
+. 45
-// computes the amount of bytes necessary to store them
+. 45
-const int num_elements = 1 << 16;
+. 45
-const int num_bytes = num_elements * sizeof(float);
+OK
-// Declaring the 'host arrays': a host array is the classical
-// array (static or dynamically allocated) we worked before.
-float *host_array_a = 0;
-float *host_array_b = 0;
-// Declaring the 'device array': this array is the equivalent
-// of classical array from C, but specially designed for the GPU
-// devices; we declare it in the same manner, but the allocation
-// process is going to be different
-float *device_array_a = 0;
-float *device_array_b = 0;
 </code>
-==== 2. Alocare memorie host (CPU) ====
+==== Operatii atomice system wide ====
-Functia malloc va intoarce o adresa virtuala ce va avea corespondent o adresa fizica din RAM.
+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>
-// Allocating the host array
+__global__ void mykernel(int *addr) {
-host_array_a = (float *) malloc(num_bytes);
+  atomicAdd_system(addr, 10);       // only available on devices with compute capability 6.x
-host_array_b = (float *) malloc(num_bytes);
-</code>
-==== 3. Alocare memorie device (GPU) ====
-Functia cudaMalloc va intoarce o adresa virtuala ce va avea corespondent o adresa fizica din VRAM.
-<code C>
-// Allocating the device's array; notice that we use a special
-// function named cudaMalloc that takes the reference of the
-// pointer declared above and the number of bytes.
-cudaMalloc((void **) &device_array_a, num_bytes);
-cudaMalloc((void **) &device_array_b, num_bytes);
-// If any memory allocation failed, report an error message
-if (host_array_a == 0 || host_array_b == 0|| device_array_a == 0 || device_array_b == 0) {
-  printf("[HOST] Couldn't allocate memory\n");
-  return 1;
 }
-</code>
-==== 4. Initializare memorie host (CPU) si copiere pe device (GPU) ====
+void foo() {
+  int *addr;
+  cudaMallocManaged(&addr, 4);
+  *addr = 0;
-{{:asc:lab10:cpu_to_gpu.png?720|}}
+   mykernel<<<...>>>(addr);
+   __sync_fetch_and_add(addr, 10);  // CPU atomic operation
-<code C>
-// Initialize the host array by populating it with float values
-for (int i = 0; i < num_elements; ++i) {
-  host_array_a[i] = (float) i;
 }
-// Copying the host array to the device memory space; notice the
-// parameters of the cudaMemcpy function; the function default
-// signature is cudaMemcpy(dest, src, bytes, flag) where
-// the flag specifies the transfer type.
-//
-// host -> device: cudaMemcpyHostToDevice
-// device -> host: cudaMemcpyDeviceToHost
-// device -> device: cudaMemcpyDeviceToDevice
-cudaMemcpy(device_array_a, host_array_a, num_bytes, cudaMemcpyHostToDevice);
 </code>
-==== 5. Executie kernel ====
+===== Operatii asincrone CUDA =====
-{{:asc:lab10:exec_gpu.png?720|}}
+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.
-<code C>
+==== Executie asincrona Host si Device ====
-// 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;
-if (num_elements % block_size)
+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.
-  ++blocks_no;
+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)
-// Launch the kernel
+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).
-vector_add<<<blocks_no, block_size>>>(device_array_a, device_array_b, num_elements);
-cudaDeviceSynchronize();
-</code>
+==== Fluxuri nonimplicite ====
-==== 6. Copiere date inapoi de la device (GPU) catre host (CPU) ====
+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>
-// Copy the result back to the host memory space
+cudaStream_t stream1;
-cudaMemcpy(host_array_b, device_array_b, num_bytes, cudaMemcpyDeviceToHost);
+cudaError_t result;
+result = cudaStreamCreate(&stream1);
+result = cudaStreamDestroy(stream1);
-// 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>
-{{:asc:lab10:gpu_to_cpu.png?720|}}
+Odata creat un astfel de flux, el poate fi utilizat in procesul de copiere a memoriei host -> device astfel:
-==== 7. Cleanup ====
 <code C>
-// Deallocate memory
+// num_bytes = N * sizeof (type_a);
-free(host_array_a);
+result = cudaMemcpyAsync(d_a, a, num_bytes, cudaMemcpyHostToDevice, stream1);
-free(host_array_b);
-cudaFree(device_array_a);
-cudaFree(device_array_b);
 </code>
-===== Compilare si executie =====
+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.
-Desi pentru un programator partile de host/CPU respectiv device/GPU pot fi in acelasi fisier *.cu, compilatorul CUDA (nvcc) le separa facand o compilare diferita pentru partea de host/CPU respectiv device/GPU. Figura de mai jos denota acest aspect.
+<code C>
+increment<<<1,N,0,stream1>>>(d_a);
-{{:asc:lab7:cuda-nvcc.png?480|}}
-Intrati pe frontend-ul ''fep8.grid.pub.ro'' folosind contul de pe curs.upb.ro. Executati comanda
-<code sh>
-srun --pty -p hpsl /bin/bash
 </code>
-pentru a accesa una din statiile cu GPU-uri. Cozile ce au unitati GPU NVIDIA Tesla sunt {{:asc:lab10:hpsl-wn0x.pdf|hpsl}} si {{:asc:lab10:dp-wn0x.pdf|dp}}.
+==== Executie asincrona programe kernel ====
-<code sh>
+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.
-[@fep8 ~]$ srun --pty -p hpsl /bin/bash
-[@hpsl-wn01 ~]$ nvidia-smi       # NVIDIA System Management Interface program
-+-----------------------------------------------------------------------------+
-| NVIDIA-SMI 470.82.01    Driver Version: 470.82.01    CUDA Version: 11.4     |
-|-------------------------------+----------------------+----------------------+
-| GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
-| Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
-|                               |                      |               MIG M. |
-|===============================+======================+======================|
-|   0  Tesla K40m          Off  | 00000000:08:00.0 Off |                    0 |
-| N/A   40C    P0    63W / 235W |      0MiB / 11441MiB |      0%      Default |
-|                               |                      |                  N/A |
-+-------------------------------+----------------------+----------------------+
-|   1  Tesla K40m          Off  | 00000000:24:00.0 Off |                    0 |
-| N/A   44C    P0    64W / 235W |      0MiB / 11441MiB |      0%      Default |
-|                               |                      |                  N/A |
-+-------------------------------+----------------------+----------------------+
-|   2  Tesla K40m          Off  | 00000000:27:00.0 Off |                    0 |
-| N/A   40C    P0    62W / 235W |      0MiB / 11441MiB |     52%      Default |
-|                               |                      |                  N/A |
-+-------------------------------+----------------------+----------------------+
-</code>
+==== Executie si transfer date asincron  ====
-Pentru laboratoarele de GP-GPU Computing vom folosi CUDA 9.1 [[https://developer.download.nvidia.com/compute/cuda/9.1/Prod/docs/sidebar/CUDA_Toolkit_Release_Notes.pdf|aici]].
+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.
-SDK-ul CUDA de la NVidia include atat implementarea de CUDA API cat si cea de OpenCL API. In cadrul laboratoarelor vom programa numai folosind CUDA. Verificam mai jos ca scheletul laboratorului compileaza.
+{{:asc:lab9:cuda_async.png?900|}}
-<code sh>
+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.
-[@hpsl-wn01 ~]$ wget -O lab7_skl.tar.gz https://ocw.cs.pub.ro/courses/_media/asc/lab7/lab7_skl.tar.gz
---2020-03-22 18:52:14--  http://ocw.cs.pub.ro/courses/_media/asc/lab7/lab7_skl.tar.gz
-Resolving ocw.cs.pub.ro (ocw.cs.pub.ro)... 141.85.227.65
-Connecting to ocw.cs.pub.ro (ocw.cs.pub.ro)|141.85.227.65|:80... connected.
-HTTP request sent, awaiting response... 302 Found
-Location: https://ocw.cs.pub.ro/courses/_media/asc/lab7/lab7_skl.tar.gz [following]
---2020-03-22 18:52:14--  https://ocw.cs.pub.ro/courses/_media/asc/lab7/lab7_skl.tar.gz
-Connecting to ocw.cs.pub.ro (ocw.cs.pub.ro)|141.85.227.65|:443... connected.
-HTTP request sent, awaiting response... 200 OK
-Length: 4884 (4.8K) [application/octet-stream]
-Saving to: 'lab7_skl.tar.gz'
-%[==================================================================================================>] 4,884       --.-K/s   in 0s
+{{:asc:lab9:cuda_async_2.png?900|}}
--03-22 18:52:14 (11.1 MB/s) - 'lab7_skl.tar.gz' saved [4884/4884]
+===== Dynamic Paralellism  =====
-[@hpsl-wn01 ~]$ tar -xvzf lab7_skl.tar.gz
+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.
-task1/
-task1/task1.cu
-task1/Makefile
-task1/Makefile_Cluster
-task2/
-task2/task2.cu
-task2/Makefile
-task2/Makefile_Cluster
-...
-[@hpsl-wn01 ~]$ cd task1/
-</code>
-<code sh>
+{{:asc:lab9:dynamic-paralellism.png?560|Fluid simulation}}
-[@hpsl-wn01 ~]$ apptainer run --nv docker://gitlab.cs.pub.ro:5050/asc/asc-public/cuda-labs:1.9.1
-INFO:    Using cached SIF image
-Apptainer>
-</code>
-În cadrul imaginei CUDA avem compilatorul nvcc.
+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]].
-<code Makefile>
+===== Exercitii =====
-COMPILER=nvcc
-LIBS=-lm
-%: %.cu
-	$(COMPILER) $^ -o $@ $(LIBS)
-clean:
+Urmăriți instrucțiunile de pe [[https://gitlab.cs.pub.ro/asc/asc-public/-/tree/master/labs/cuda/advanced|GitLab]].
-	rm -rf task0
-</code>
-Exemplu de compilare si rulare interactiva pe coada hpsl
-<code sh>
-Apptainer> make task1
-nvcc task1.cu -o task1 -lm
-Apptainer> ./task1
-[HOST] Hello from the host!
-[HOST] You have 3 CUDA-capable GPU(s)
-Apptainer> make clean
-rm -rf task1
-Apptainer> cd ../task2
-Apptainer> make task2
-nvcc task2.cu -o task2 -lm
-Apptainer> ./task2
-Result 0: 2 * 0.0 + 1.0/(0.0 + 1.0)= 0.000
-Result 1: 2 * 1.0 + 1.0/(1.0 + 1.0)= 0.000
-Result 2: 2 * 2.0 + 1.0/(2.0 + 1.0)= 0.000
-Result 3: 2 * 3.0 + 1.0/(3.0 + 1.0)= 0.000
-Result 4: 2 * 4.0 + 1.0/(4.0 + 1.0)= 0.000
-Result 5: 2 * 5.0 + 1.0/(5.0 + 1.0)= 0.000
-Result 6: 2 * 6.0 + 1.0/(6.0 + 1.0)= 0.000
-Result 7: 2 * 7.0 + 1.0/(7.0 + 1.0)= 0.000
-Result 8: 2 * 8.0 + 1.0/(8.0 + 1.0)= 0.000
-Result 9: 2 * 9.0 + 1.0/(9.0 + 1.0)= 0.000
-Apptainer>
-</code>
-Exemplu executie program CUDA folosind rularea ne-interactiva pe coada de executia hpsl (apelul trebuie facut de pe fep8.grid.pub.ro):
-<code sh>
-[fep8 ~]$ cd task2/
-[fep8 task2]$ cat ../utils/batch_run.sh
-#!/bin/bash
-apptainer exec --nv $CONTAINER_IMAGE  \
-./$TASK
-[fep8 task2]$ sbatch --time 01:00:00 -p hpsl --export=TASK=task2,CONTAINER_IMAGE=docker://gitlab.cs.pub.ro:5050/asc/asc-public/cuda-labs:1.9.1 ../utils/batch_run.sh
-Submitted batch job 1816
-[fep8 task2]$ squeue | grep 1816
-      hpsl batch_ru stefan_d  R       0:07      1 hpsl-wn01
-[fep8 task2]$ cat slurm-1816.out
-INFO:    Using cached SIF image
-Result 0: 2 * 0.0 + 1.0/(0.0 + 1.0)= 1.000
-Result 1: 2 * 1.0 + 1.0/(1.0 + 1.0)= 2.500
-Result 2: 2 * 2.0 + 1.0/(2.0 + 1.0)= 4.333
-Result 3: 2 * 3.0 + 1.0/(3.0 + 1.0)= 6.250
-Result 4: 2 * 4.0 + 1.0/(4.0 + 1.0)= 8.200
-Result 5: 2 * 5.0 + 1.0/(5.0 + 1.0)= 10.167
-Result 6: 2 * 6.0 + 1.0/(6.0 + 1.0)= 12.143
-Result 7: 2 * 7.0 + 1.0/(7.0 + 1.0)= 14.125
-Result 8: 2 * 8.0 + 1.0/(8.0 + 1.0)= 16.111
-Result 9: 2 * 9.0 + 1.0/(9.0 + 1.0)= 18.100
-</code>
-Puteți folosi Makefile_Cluster:
-<code sh>
-[fep8 ~]$ cd task2/
-[fep8 task2]$ make -f Makefile_Cluster clean
-rm -rf task2
-rm -rf slurm-*
-[fep8 task2]$ make -f Makefile_Cluster task2
-sbatch --time 01:00:00 -p hpsl --export=TASK=task2,CONTAINER_IMAGE=docker://gitlab.cs.pub.ro:5050/asc/asc-public/cuda-labs:1.9.1 ../utils/batch_build.sh | ../utils/batch_wait.sh
-INFO:    Using cached SIF image
-nvcc task2.cu -o task2 -lm
-[fep8 task2]$ make -f Makefile_Cluster run_task1
-sbatch --time 01:00:00 -p hpsl --export=TASK=task2,CONTAINER_IMAGE=docker://gitlab.cs.pub.ro:5050/asc/asc-public/cuda-labs:1.9.1 ../utils/batch_run.sh | ../utils/batch_wait.sh
-INFO:    Using cached SIF image
-Result 0: 2 * 0.0 + 1.0/(0.0 + 1.0)= 1.000
-Result 1: 2 * 1.0 + 1.0/(1.0 + 1.0)= 2.500
-Result 2: 2 * 2.0 + 1.0/(2.0 + 1.0)= 4.333
-Result 3: 2 * 3.0 + 1.0/(3.0 + 1.0)= 6.250
-Result 4: 2 * 4.0 + 1.0/(4.0 + 1.0)= 8.200
-Result 5: 2 * 5.0 + 1.0/(5.0 + 1.0)= 10.167
-Result 6: 2 * 6.0 + 1.0/(6.0 + 1.0)= 12.143
-Result 7: 2 * 7.0 + 1.0/(7.0 + 1.0)= 14.125
-Result 8: 2 * 8.0 + 1.0/(8.0 + 1.0)= 16.111
-Result 9: 2 * 9.0 + 1.0/(9.0 + 1.0)= 18.100
-</code>
 <note important>
 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ă.
+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ă. (Sau ''squeue --me'').
 Daca nu veti face aceasta delogare, veti putea ajunge in situatia in care sa nu va mai puteti loga pe nodurile din cluster.
-</note>
-<note important>
-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.
-Tutorial chei ssh: https://www.ssh.com/academy/ssh/keygen
-Tutorial montare filesystem din fep pe local - comenzi date pe local:
-<code sh>
-mkdir /mnt/asc
-sudo chown -R <user> /mnt/asc
-decomentam linia 'user_allow_other' din /etc/fuse.conf
-sshfs -o allow_other <user.moodle>@fep8.grid.pub.ro:/ /mnt/asc
-</code>
-</note>
-===== Exercitii =====
-Pentru inceput:
-  - Logati-va pe ''fep8.grid.pub.ro'' folosind contul de pe ''curs.upb.ro''
-  - Executați comanda:
-<code>wget -O lab7_skl.tar.gz http://ocw.cs.pub.ro/courses/_media/asc/lab7/lab7_skl.tar.gz</code>
-  - Dezarhivati folosind comanda ''tar -xzvf lab7_skl.tar.gz''
-<note tip>
-Debug aplicatii CUDA [[https://docs.nvidia.com/cuda/cuda-gdb/index.html#introduction|aici]]
-</note>
-Modificarile se vor face acolo unde este necesar in ''task_<i>.cu'' unde ''<i>'' este numarul taskului. Urmariti indicatiile ''TODO'' din cod. De asemenea, va recomandam sa folositi documentatia oficiala CUDA Toolkit Documentation de la adresa: https://docs.nvidia.com/cuda/. Aici veti gasi informatii despre majoritatea functiilor de care aveti nevoie (folositi functia search).
-**Task 1**  - Rulați task1 ca exemplu pentru a verifica funcționalitatea CUDA pe GPU
-**Task 2**  - Rulați task2 ca exemplu pentru efectuarea unor operații pe GPU
-**Task 4**  - Efectuați adunarea a doi vectori folosind CUDA în task4.cu
-    * Sugestia este de a face intai taskul 4 si apoi taskul 3 pentru ca are sens dpdv logic - e mai usor de inteles ce se intampla.
-**Task 3**  - Urmăriți TODO--uri din taks3.cu
-    * Listați informații despre device-urile existente și selectați primul device
-    * Completați și rulați kernelul kernel_parity_id
-    * Completați și rulați kernelul kernel_block_id; explicați rezultatul
-    * Completați și rulați kernelul kernel_thread_id; explicați rezultatul
-**Task 5**  - Urmăriți instrucțiunile din task5.cu pentru a realiza interschimbarea a doi vectori
-<note important>
-Recomandăm sa folosiți pentru compilarea și rularea task-urilor ''sbatch'' sau Makefile_Cluster
 </note>
 ===== Resurse =====
-{{:asc:lab7:lab7_skl.tar.gz|Schelet Laborator 7}}
 <hidden>
-{{:asc:lab7:sol:lab7_sol.tar.gz|Soluție Laborator 7}}
+{{:asc:lab9:sol:lab9_sol.tar.gz|Soluție Laborator 6}}
 </hidden>
-{{:asc:lab7:asc_lab7.pdf|Enunt Laborator 7}}
+/* {{:asc:lab9:sol:lab9_sol.zip|Solutie Laborator 6}} */
-  * Responsabili laborator: Grigore Lupescu, Ștefan-Dan Ciocîrlan, Costin Carabaș
+/* {{:asc:lab6:asc_lab9.pdf|Enunt Laborator 6}} */
-==== Referinte ====
+  * Responsabili laborator: Matei Barbu, Alexandru 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/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, dp-wn04)
-    * [[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://devblogs.nvidia.com/introduction-cuda-dynamic-parallelism/|CUDA Dynamic Parallelism]]
-  * Extra
     * [[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]]
-    * [[https://www.nvidia.com/docs/IO/116711/sc11-cuda-c-basics.pdf | Supercomputing 2011 Tutorial ]]
+    * [[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://developer.nvidia.com/blog/how-overlap-data-transfers-cuda-cc/|How to Overlap Data Transfers in CUDA C/C++]]