Show page

Differences

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

--- asc:laboratoare:04 [2024/03/30 21:02]
alexandru.bala [Aplicații]
+++ asc:laboratoare:04 [2026/02/23 18:48] (current)
giorgiana.vlasceanu
@@ Line 1: / Line 1: @@
-====== Laboratorul 04 - Arhitecturi de tip GPGPU ======
+====== Laboratorul 4 - Analiza Performantei Programelor ======
-===== Introducere =====
+The purpose of this lab is the familiarization with the field of application profiling & debugging, through the means of dedicated tools for spotting performance bottlenecks.
-==== Motivație ====
+We will focus on several open source and commercial tools, as follows:
+  * valgrind / kcachegrind (on own systems & hpslcluster queue)
+  * perf (on your own systems)
-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).
+===== 1. Valgrind / KCachegrind=====
-Utilitatea lui s-a extins ulterior către probleme "embarrassingly parallel", iar în astăzi sunt folosite la antrenarea 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.
+Valgrind is a tool used for memory debugging, memory leak detection and profiling. It is also a generic framework for creating dynamic analysis tools, such as memory checkers [1].
-==== Teorie ====
+Valgrind is in essence a virtual machine using just-in-time compilation techniques, including dynamic recompilation. It is important to keep in mind that nothing from the original program ever gets run directly on the host processor. Instead, it will translate the input program into a simpler form called Intermediate Representation (IR), which is processor neutral. After this transformation, a tool [2] is called to do whatever transformation of the IR it needs and the resulting IR is then translated back into machine code and ran on the host processor.
-Î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:
+The tools available in Valgrind are:
-  * schimb rapid de context între thread-uri,
+  * **memcheck**. Used to detect memory-management problems and it is aimed at C and C++ programs. All memory reads and writes are checked, and all calls to malloc/new/free/delete are intercepted. Therefore it can detect memory leaks, access to invalid memory, weird initialization values, overflows, etc. Memcheck runs programs about 10-30x slower than normal;
-  * planificarea în grupuri de thread-uri,
+  * **cachegrind**. Used to profile CPU cache. It performs detailed simulation of the I1, D1 and L2 caches in order to pinpoint the sources of cache misses. It identifies the number of cache misses, memory references and instructions executed for each line of source code. Cache grind runs programs about 20-100x slower than normal;
-  * și orientare către prelucrari masive de date.
+  * **callgrind**. It is an extension to cachegrind and provides all the information that the latter offers, plus extra information regarding call graphs. In order to view the results, a visualization tool called KCachegrind [3] can be used;
+  * **massif**. It is a heap profiler and it performs detailed heap profiling by taking regular snapshots of a program's heap and produces a graph showing heap usage over time, including information about which parts of the program are responsible for the most memory allocations. Massif runs programs about 20x slower than normal;
+  * **helgrind** and drd. These tools are thread debuggers which find data races in multithreaded programs. They look for memory locations which are accessed by more than one (POSIX) pthread, but for which no consistently used (pthread_mutex_) lock can be found;
+  * other 3rd party tools can be found here [4].
-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.
-<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>
-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.
+===== 2. Analysis of the Tachyon raytracing engine =====
-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).
-Principalii producatori de core-uri IP (intellectual property) tip GPU sunt:
+In this section, we will focus on analyzing a software application. We will analyze both a serial and a parallel implementation. The application is called “tachyon” and you can find the source code attached to this lab.
-  * 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
-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 (exp. controllere PCIe/USB/SATA/ETH).
+On your own system, before compilation, you must install the X11 dev tools and create a set of symlinks. For Ubuntu 64 bit, we must do the following:
+  * install dependencies <code bash> sudo apt-get install libx11-dev </code>
+  * create the symlinks:
+    * <code bash> sudo mkdir /usr/lib64 </code>
+    * <code bash> sudo ln -s /usr/lib/x86_64-linux-gnu/libX11.so /usr/lib64/libX11.so </code>
+    * <code bash> sudo ln -s /usr/lib/x86_64-linux-gnu/libXext.so /usr/lib64/libXext.so </code>
-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 comunicare cu sistemul. Exemple de GPU-uri dedicate sunt seriile de plăci grafice Geforce (Nvidia) și Radeon (AMD).
+To compile it, you must extract the archive {{asc:lab6:tachyon_vtune_amp_xe.tgz| Tachyon}} to local disk and run make. You can test the compilation by running in the same directory: <code bash>./tachyon_find_hotspots dat/balls.dat </code>
-{{:asc:lab10:dgpu_igpu.png?direct&750|}}
+You should see a window like the one below:
-==== Aplicații ====
+{{ :asc:lab6:image09.png?300 |}}
-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.
+To run and observe the functionalities of this tool, follow the following sequence of steps and instructions:
+<code sh>
-  * https://www.nvidia.com/en-us/self-driving-cars/partners/bmw/
+wget -O tachyon_vtune_amp_xe.tgz http://ocw.cs.pub.ro/courses/_media/asc/lab6/tachyon_vtune_amp_xe.tgz
-  * https://blogs.nvidia.com/blog/2018/09/18/audi-unveils-e-tron-electric-suv/
+gunzip tachyon_vtune_amp_xe.tgz
+tar -xvf tachyon_vtune_amp_xe.tar
-Rețele neurale - antrenare vs. inferență.
+cd tachyon
-  * https://www.forbes.com/sites/forbestechcouncil/2017/12/01/for-machine-learning-its-all-about-gpus/
+make
-  * 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, recunoaștere facială/audio.
-Simulări fizice - NVIDIA Physx, Folding@Homel
-  * https://blogs.nvidia.com/blog/2018/11/13/weather-predicted-sc18-gpu-hpc-jensen-huang/
-Prelucrări multimedia - filtre imagini GIMP/Photoshop.
-Alte domenii - arhivare (WinZip), encriptare.
-{{: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-urile 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>
-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.
+=== Using Valgrind on the Tachyon Code ===
-<code C>
+. Make sure you have Valgrind and KCachegrind installed on the system (or login on the hp-sl.q queue) and the application in the initial state, without any modifications on your system
-// Launch the kernel
+<code bash>
-vector_add<<<blocks_no, block_size>>>(device_array_a, device_array_b, num_elements);
+sudo apt-get update
+sudo apt-get install valgrind kcachegrind
 </code>
+. We will use the tool //callgrind //to get information from the running application. Run the following command line:
-==== HelloWorld CUDA ====
+<code>
+valgrind --tool=callgrind --collect-jumps=yes --dump-instr=yes --collect-systime=yes -- ./tachyon_find_hotspots dat/balls.dat
-<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>
+. Open the profile in KCachegrind and click on the Calee Map tab. Also, make sure that the buttons //% Relative//, //Cycle detection// and //Relative to parent// are selected. You should see something like this:
+{{ :asc:lab6:image20.png?800p |}}
+From this image, we can see that valgrind measured that about 78% of the total time was spent in the initialize_2D_buffer function. Double click the square containing the function name, then select the “Source code” tab and you will see the problematic code.
+{{ :asc:lab6:image05.png |}}
-===== Aplicatie compute CUDA =====
-O aplicatie CUDA are ca scop executia de cod pe GPU-uri NVIDIA CUDA.
+===== 3. Perf =====
-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 ====
+Perf is a performance analysis tool, available in the Linux kernel since version 2.6.31 [5]. The userspace control application is accessed from the command line and provides a number of subcommands. Unlike Valgrind, perf is capable of statistical profiling of both the entire system (kernel and userspace) and per process PID basis. It supports hardware performance counters, tracepoints, software performance counters (e.g. hrtimer), and dynamic probes (for example, kprobes or uprobes).
-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).
+Perf is used with several subcommands:
+  * **stat**: measure total event count for a single program or for the whole system for a specified time period;
+  * **top**: top-like dynamic view of hottest functions;
+  * **record**: measure and save sampling data for single program;
+  * **report**: analyze file generated by perf record;
+  * **annotate**: annotate sources or assembly;
+  * **sched**: tracing/measuring of scheduler actions and latencies;
+  * **list**: list available events.
-<code C>
+. Make sure you have perf installed on the system and the application in the initial state, without any modifications. **You can only run perf as root. You can only do this on your system.**
-/**
+. Run the following command line:
- * This kernel computes the function f(x) = 2x + 1/(x + 1) for each
+<code bash>
- * element in the given array.
+perf record -a -g -- ./tachyon_find_hotspots
- */
-__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>
+For other perf parameters, you can read this [[http://www.brendangregg.com/perf.html|link]]
-==== 1. Definire zone de memorie host si device ====
+. Run the following command line to view the collected results:
+<code>
-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.
+perf report
-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).
-<code C>
-// Declare variable to represent ~1M float values and
-// computes the amount of bytes necessary to store them
-const int num_elements = 1 << 16;
-const int num_bytes = num_elements * sizeof(float);
-// 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>
+You should see a screen like the following:
+{{ :asc:lab6:image19.png |}}
+From this image you can see that perf will display the symbol for the function that takes the most amount of CPU time in red. In our case it’s the //_Z20initialize_2D_bufferPjS_//, which translates in the C source code into the same function as with VTune and Valgrind.
+<note tip>Hint: To find out the demangled name, use the c++filt command: <code> c++filt _Z20initialize_2D_bufferPjS_</code></note>
-==== 2. Alocare memorie host (CPU) ====
-Functia malloc va intoarce o adresa virtuala ce va avea corespondent o adresa fizica din RAM.
+===== Exercitii =====
-<code C>
-// Allocating the host array
-host_array_a = (float *) malloc(num_bytes);
-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>
+Pentru acest laborator se pot utiliza sistemele din cluster prin intermediul fep.grid.pub.ro:
-// Allocating the device's array; notice that we use a special
+  - ''ssh -Y username@fep.grid.pub.ro'' (natural puneti utilizatorul vostru in loc de username)
-// function named cudaMalloc that takes the reference of the
+  - ''wget https://ocw.cs.pub.ro/courses/_media/asc/lab6/lab6_skl.tar.gz -O lab6_skl.tar.gz'' - Downloadati arhiva laboratorului
-// pointer declared above and the number of bytes.
+  - ''tar -xzvf lab6_skl.tar.gz'' - dezarhivati arhiva downloadata mai sus
-cudaMalloc((void **) &device_array_a, num_bytes);
+  - ''%%srun --time 00:05:00 --x11 -p haswell --pty /bin/bash%%'' - conectati-va pe coada Haswell cu 14 servere
-cudaMalloc((void **) &device_array_b, num_bytes);
+  - ''%%apptainer run docker://gitlab.cs.pub.ro:5050/asc/asc-public/c-labs:1.3.1 /bin/bash%%'' - accesati imaginea de docker in cadrul căreia avem permisiunile necesare realizării laboratorului
-// If any memory allocation failed, report an error message
+**Task 0**  - Folosit valgrind pentru ''task0.c'' urmarind TODO-uri pentru teste.
-if (host_array_a == 0 || host_array_b == 0|| device_array_a == 0 || device_array_b == 0) {
+    * ''make task0'' pentru versiunea seriala
-  printf("[HOST] Couldn't allocate memory\n");
+    * ''make openmp_task0'' pentru versiunea paralelizata
-  return 1;
+    * <code sh>
-}
+[@fep ]$ srun --time 00:05:00 --x11 -p haswell --pty /bin/bash
+[@haswell-wnxx ~]$ apptainer run docker://gitlab.cs.pub.ro:5050/asc/asc-public/c-labs:1.3.1 /bin/bash
+Apptainer> make task0
+Apptainer> valgrind --tool=callgrind -v --dump-every-bb=10000000 ./task0
+Apptainer> kcachegrind
+Apptainer> make clean
+Apptainer> make openmp_task0
+Apptainer> valgrind --tool=callgrind -v --dump-every-bb=10000000 ./task0
+Apptainer> kcachegrind
 </code>
-==== 4. Initializare memorie host (CPU) si copiere pe device (GPU) ====
+**Task 1**  - Analizati aplicatia Tachyon.
+    * Rulati scriptul ''task1.sh'' pentru a descarca si compila Tachyon. Atentie, trebuie sa rulati pe fep scriptul task1.sh pentru ca doar aici aveti disponibil libX11.so de care aveti nevoie pentru compilare. Pe masinile proprii aveti de asemenea nevoie de libX11.so asa incat e bine sa il intalati inainte de rularea acestui task.
-{{:asc:lab10:cpu_to_gpu.png?720|}}
+    * Varianta seriala ''tachyon_find_hotspots''
+    * Varianta paralelizata ''tachyon_analyze_locks''
-<code C>
+    * <code sh>
-// Initialize the host array by populating it with float values
+[@fep ~]$ ./task1.sh
-for (int i = 0; i < num_elements; ++i) {
+[@fep ]$ srun --time 00:05:00 --x11 -p haswell --pty /bin/bash
-  host_array_a[i] = (float) i;
+[@haswell-wn01 ~]$ apptainer run docker://gitlab.cs.pub.ro:5050/asc/asc-public/c-labs:1.3.1 /bin/bash
-}
+Apptainer> cd tachyon
+Apptainer> valgrind --tool=callgrind --collect-jumps=yes --dump-instr=yes --collect-systime=yes -- ./tachyon_find_hotspots dat/balls.dat
-// Copying the host array to the device memory space; notice the
+Apptainer> valgrind --tool=callgrind --collect-jumps=yes --dump-instr=yes --collect-systime=yes -- ./tachyon_analyze_locks dat/balls.dat
-// 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>
+    * Analizati cu perf
-==== 5. Executie kernel ====
-{{:asc:lab10:exec_gpu.png?720|}}
-<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;
-if (num_elements % block_size)
-  ++blocks_no;
-// Launch the kernel
+**Task 2**  - Folositi tool-ul cachegrind din valgrind pentru a analiza codul care realizeaza inmultirea de matrice folosind diferite reordonari ale buclelor.
-vector_add<<<blocks_no, block_size>>>(device_array_a, device_array_b, num_elements);
+    * Compilati si rulati ''task2.c''
-cudaDeviceSynchronize();
+    * Notati observatiile voastre legate de numarul de I refs, D refs, D1 misses, branches si mispredicts.
+    * <code sh>
+[@fep ]$ srun --time 00:05:00 --x11 -p haswell --pty /bin/bash
+[@haswell-wn01 ~]$ apptainer run docker://gitlab.cs.pub.ro:5050/asc/asc-public/c-labs:1.3.1 /bin/bash
+Apptainer> make task2
+Apptainer> valgrind --tool=cachegrind --branch-sim=yes ./task2 1
+Apptainer> valgrind --tool=cachegrind --branch-sim=yes ./task2 2
+Apptainer> valgrind --tool=cachegrind --branch-sim=yes ./task2 3
 </code>
-==== 6. Copiere date inapoi de la device (GPU) catre host (CPU) ====
+<note important>
+Recomandăm sa va delogati mereu de pe serverele din cluster dupa terminarea sesiunii, utilizand comanda ''exit''
-<code C>
+Alternativ, daca ati uitat sesiuni deschise, puteti verifica acest lucru de pe fep.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 fep.grid.pub.ro) pentru a vedea doar sesiunile care vă interesează.
-// Copy the result back to the host memory space
-cudaMemcpy(host_array_b, device_array_b, num_bytes, cudaMemcpyDeviceToHost);
-// Print out the first 10 results
+Daca nu veti face aceasta delogare, veti putea ajunge in situatia in care sa nu va mai puteti loga pe nodurile din cluster.
-for (int i = 0; i < 10; ++i) {
+</note>
-  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|}}
-==== 7. Cleanup ====
-<code C>
+===== Reference =====
-// Deallocate memory
+  * http://valgrind.org/
-free(host_array_a);
+  * http://valgrind.org/info/tools.html
-free(host_array_b);
+  * http://kcachegrind.sourceforge.net/html/Usage.html
-cudaFree(device_array_a);
+  * http://valgrind.org/downloads/variants.html
-cudaFree(device_array_b);
+  * https://perf.wiki.kernel.org/index.php/Main_Page
-</code>
+  * https://software.intel.com/en-us/intel-parallel-studio-xe
+  * http://www.brendangregg.com/perf.html
-===== Aplicații =====
+  * https://www.oracle.com/tools/developerstudio/downloads/developer-studio-jsp.html
-{{:asc:lab7:cuda-nvcc.png?480|}}
-<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>
-Urmăriți instrucțiunile de pe GitLab [[https://gitlab.cs.pub.ro/asc/asc-public/-/tree/master/labs/cuda/intro|GitLab]].
-===== Resurse =====
+==== Resources ====
+  * Responsabilii acestui laborator: [[emil.slusanschi@cs.pub.ro|Emil Slusanschi]], Alexandru Patrascu si Octavian Moraru.
+  * <html><a class="media mediafile mf_pdf" href=":asc:lab6:index?do=export_pdf">PDF laborator</a></html>
+  * {{asc:lab6:tachyon_vtune_amp_xe.tgz|Aplicatie laborator (tachyon)}}
+  * {{asc:lab6:lab6_skl.tar.gz|Schelet laborator}}
 <hidden>
-{{:asc:lab7:sol:lab7_sol.tar.gz|Soluție Laborator 7}}
+  * {{:asc:lab6:sol:lab6_sol.tar.gz|Soluție Laborator 6}}
 </hidden>
-{{:asc:lab7:asc_lab7.pdf|Enunt Laborator 4}}
-  * Responsabili laborator: Costin Carabaș, Tudor Calafeteanu, Grigore Lupescu, Mihnea Mitroi, Irinel Gul, Alex Bala
-==== Referinte ====
-  * 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://docs.nvidia.com/cuda/cuda-toolkit-release-notes/index.html|CUDA Dev Toolkit]]
-    * [[https://developer.nvidia.com/cuda-gpus|CUDA GPUs]]
-  * Acceleratoare xl (NVidia P100)
-    * [[https://www.nvidia.com/en-us/data-center/tesla-p100/|NVIDIA Pascal P100]]
-  * Advanced CUDA
-    * [[https://developer.download.nvidia.com/CUDA/training/StreamsAndConcurrencyWebinar.pdf|CUDA Streams]]
-  * 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]]