Show page

Differences

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

--- asc:laboratoare:04 [2025/03/26 11:00]
alexandru.bala [Compute capability]
+++ asc:laboratoare:04 [2026/03/26 14:24] (current)
emil.slusanschi [Reference]
@@ 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 the following open source tools:
+  * ''valgrind'' (''callgrind'', ''cachegrind'') with ''callgrind_annotate'' and ''cg_annotate'' for results visualization (on ''haswell'' partition)
+  * ''perf'' (on your own systems only, as hardware counters are not available inside Apptainer containers)
-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 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.
+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:
+<note tip>
-  * schimb rapid de context între thread-uri,
+Valgrind makes programs run 10-100x slower than normal, depending on the tool used. This is expected behavior.
-  * planificarea în grupuri de thread-uri,
-  * și orientare către prelucrari masive de date.
-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.
+The tools available in Valgrind are:
-Ampla folosire a acestora se datorează:
+  * **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;
-  * diferențelor de putere de procesare brută dintre CPU și GPU (instrucțiuni/secundă)
+  * **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;
-  * standardizarea de API-uri care ușurează munca programatorilor
+  * **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;
-  * răspândirea aplicațiilor ce pot beneficia de pe urma paralelismului de tip SIMD
+  * **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;
-  * cererii pe piața unităților computaționale destinate:
+  * **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;
-    * atât consumatorilor (PC, Smartphone, TV, etc.),
+  * other 3rd party tools can be found here [4].
-    * cât și mediilor industriale (Automotive, HPC etc).
-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
-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).
-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).
+===== 2. Analysis of the Tachyon raytracing engine =====
-{{:asc:lab10:dgpu_igpu.png?direct&750|}}
+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.
-==== Aplicații ====
+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>
-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 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>
-  * https://www.nvidia.com/en-us/self-driving-cars/partners/bmw/
+You should see a window like the one below:
-  * https://blogs.nvidia.com/blog/2018/09/18/audi-unveils-e-tron-electric-suv/
-Rețele neurale - antrenare vs. inferență.
+{{ :asc:lab6:image09.png?300 |}}
-  * 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.
+To run and observe the functionalities of this tool, follow the following sequence of steps and instructions:
-  * https://coincentral.com/best-gpu-for-mining-2018-edition/
+<code sh>
+wget -O tachyon_vtune_amp_xe.tgz http://ocw.cs.pub.ro/courses/_media/asc/lab6/tachyon_vtune_amp_xe.tgz
-SmartTV, Smartphone - accelerare video, recunoaștere facială/audio.
+gunzip tachyon_vtune_amp_xe.tgz
+tar -xvf tachyon_vtune_amp_xe.tar
-Simulări fizice - NVIDIA Physx, Folding@Homel
+cd tachyon
-  * https://blogs.nvidia.com/blog/2018/11/13/weather-predicted-sc18-gpu-hpc-jensen-huang/
+make
-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-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>
-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) ====
+Pentru acest laborator se pot utiliza sistemele din cluster prin intermediul fep.grid.pub.ro:
+  - Vă conectați pe ''fep'', conform [[https://ocw.cs.pub.ro/courses/asc/laboratoare/01|Laboratorului 1]] (Vezi ''Ce este FEP-ul?)''.
+  - Downloadați arhiva laboratorului: ''wget https://ocw.cs.pub.ro/courses/_media/asc/laboratoare/lab4_skl.tar.gz -O lab4_skl.tar.gz''.
+  - Dezarhivați arhiva downloadată mai sus: ''tar -xzvf lab4_skl.tar.gz''.
+  - Intrați în mod interactiv pe Haswell: ''%%srun --pty bash%%''.
+  - Accesați imaginea de docker în cadrul căreia veți face laboratorul:''%%apptainer run docker://gitlab.cs.pub.ro:5050/asc/asc-public/c-labs:1.3.1 /bin/bash%%''
-Functia cudaMalloc va intoarce o adresa virtuala ce va avea corespondent o adresa fizica din VRAM.
+**Task 0**  - Folositi Callgrind pentru ''task0.c'', urmărind TODO-urile din cod.
+<code sh>
-<code C>
+# Versiunea serială
-// Allocating the device's array; notice that we use a special
+Apptainer> make task0
-// function named cudaMalloc that takes the reference of the
+Apptainer> valgrind --tool=callgrind -v --dump-every-bb=10000000 ./task0
-// pointer declared above and the number of bytes.
+Apptainer> callgrind_annotate callgrind.out.<pid>
-cudaMalloc((void **) &device_array_a, num_bytes);
-cudaMalloc((void **) &device_array_b, num_bytes);
-// If any memory allocation failed, report an error message
+# Versiunea paralelizată
-if (host_array_a == 0 || host_array_b == 0|| device_array_a == 0 || device_array_b == 0) {
+Apptainer> make clean
-  printf("[HOST] Couldn't allocate memory\n");
+Apptainer> make openmp_task0
-  return 1;
+Apptainer> valgrind --tool=callgrind -v --dump-every-bb=10000000 ./task0
-}
+Apptainer> callgrind_annotate callgrind.out.<pid>
 </code>
-==== 4. Initializare memorie host (CPU) si copiere pe device (GPU) ====
+Notați și explicați următoarele observații:
+  * Care este funcția cu cel mai mare număr de instrucțiuni (coloana ''Ir'')?
+  * Cum diferă profiling-ul serial de cel paralel? Ce funcții noi apar în versiunea paralelizată?
+  * De ce ordinea primelor 100 de prime diferă între versiunea serială și cea paralelă?
-{{:asc:lab10:cpu_to_gpu.png?720|}}
+**Task 1**  - Analizati aplicatia Tachyon (ray tracer).
-<code C>
+<code sh>
-// Initialize the host array by populating it with float values
+Apptainer> ./task1.sh
-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 ====
-{{:asc:lab10:exec_gpu.png?720|}}
-<code C>
+# Versiunea serială
-// Compute the parameters necessary to run the kernel: the number
+Apptainer> valgrind --tool=callgrind --collect-jumps=yes --dump-instr=yes --collect-systime=yes -- ./tachyon_find_hotspots dat/balls.dat
-// of blocks and the number of threads per block; also, deal with
+Apptainer> callgrind_annotate callgrind.out.<pid>
-// a possible partial final block
-const size_t block_size = 256;
-size_t blocks_no = num_elements / block_size;
-if (num_elements % block_size)
+# Versiunea paralelizată
-  ++blocks_no;
+Apptainer> valgrind --tool=callgrind --collect-jumps=yes --dump-instr=yes --collect-systime=yes -- ./tachyon_analyze_locks dat/balls.dat
+Apptainer> callgrind_annotate callgrind.out.<pid>
-// Launch the kernel
+# Perf (doar software counters în container)
-vector_add<<<blocks_no, block_size>>>(device_array_a, device_array_b, num_elements);
+Apptainer> perf stat ./tachyon_find_hotspots dat/balls.dat
-cudaDeviceSynchronize();
+Apptainer> perf stat ./tachyon_analyze_locks dat/balls.dat
 </code>
+Notați și explicați următoarele observații:
+  * Care sunt funcțiile cu cel mai mare număr de instrucțiuni în versiunea serială?
+  * Ce funcție consumă cel mai mult din versiunea serială și de ce? (hint: ''initialize_2D_buffer'')
+  * Cum diferă profiling-ul serial de cel paralel? Ce funcții noi apar?
+  * Unde este plasat mutex-ul în ''analyze_locks''? Este necesar acolo?
-==== 6. Copiere date inapoi de la device (GPU) catre host (CPU) ====
+**Task 2** - Analizați înmulțirea de matrice cu diferite ordonări ale buclelor folosind Cachegrind.
+    * Compilati si rulati ''task2.c''
+    * Notati observatiile voastre legate de numarul de I refs, D refs, D1 misses, branches si mispredicts.
-<code C>
+<code sh>
-// Copy the result back to the host memory space
+Apptainer> make task2
-cudaMemcpy(host_array_b, device_array_b, num_bytes, cudaMemcpyDeviceToHost);
+Apptainer> valgrind --tool=cachegrind --cache-sim=yes ./task2 1
+Apptainer> valgrind --tool=cachegrind --cache-sim=yes ./task2 2
-// Print out the first 10 results
+Apptainer> valgrind --tool=cachegrind --cache-sim=yes ./task2 3
-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|}}
+Completați tabelul cu rezultatele obținute:
-==== 7. Cleanup ====
+^ Metric        ^ Mode 1 (ijk) ^ Mode 2 (ikj) ^ Mode 3 (jki) ^
+| I refs        |              |              |              |
+| D refs        |              |              |              |
+| D1 miss rate  |              |              |              |
+| LLd miss rate |              |              |              |
-<code C>
+Notați și explicați următoarele observații:
-// Deallocate memory
+  * Comparați valorile ''I refs'' și ''D refs'' între cele 3 ordonări. Ce observați?
-free(host_array_a);
+  * Comparați valorile ''D1 miss rate'' între cele 3 ordonări. Care mod este mai eficient din punct de vedere al cache-ului și de ce?
-free(host_array_b);
+  * Comparați valorile ''LLd miss rate'' între cele 3 ordonări. Ce observați?
-cudaFree(device_array_a);
+  * Cum influențează ordinea buclelor accesul la memorie în C?
-cudaFree(device_array_b);
-</code>
-===== Aplicații =====
-{{: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
+===== Reference =====
+  * http://valgrind.org/
-Pe fep8, din directorul vostru home (''%%~%%''), dați comanda:
+  * http://valgrind.org/info/tools.html
-<code sh>
+  * http://kcachegrind.sourceforge.net/html/Usage.html
-mkdir asc
+  * http://valgrind.org/downloads/variants.html
-cd asc
+  * https://perf.wiki.kernel.org/index.php/Main_Page
-pwd # acesta va returna <cale_absolută_director_asc_fep8>
+  * https://software.intel.com/en-us/intel-parallel-studio-xe
-</code>
+  * http://www.brendangregg.com/perf.html
+  * https://www.oracle.com/tools/developerstudio/downloads/developer-studio-jsp.html
-Tutorial montare filesystem din fep pe local - comenzi date pe local:
+  * https://guides.upb.ro/docs/grid/apptainer-usage
-<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 =====
+==== 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]]
-    * [[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]]