Differences

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

Link to this comparison view

asc:laboratoare:08 [2022/03/30 22:27]
stefan_dan.ciocirlan [NVIDIA Visual Profiler]
asc:laboratoare:08 [2024/02/29 13:09] (current)
giorgiana.vlasceanu created
Line 1: Line 1:
-====== Laboratorul 08 - Arhitectura GPU NVIDIA CUDA ======+====== Laboratorul 08 - Arhitecturi de Microprocesoare si Sisteme de Calcul ​======
  
-Arhitectura NVIDIA FERMI [[https://​www.nvidia.com/​content/​PDF/​fermi_white_papers/​NVIDIA_Fermi_Compute_Architecture_Whitepaper.pdf|aici]],​ Tesla 2070, coada executie fep.grid.pub.ro -> dp+===== Obiective =====
  
-Arhitectura NVIDIA KEPLER [[https://www.nvidia.com/​content/​PDF/​kepler/​NVIDIA-Kepler-GK110-Architecture-Whitepaper.pdf|aici]],​ Tesla K40M, coada executie fep.grid.pub.ro ​-> hpsl+Doua componente esentiale ale structurii unui sistem de calcul sunt reprezentate de catre procesor si placa de baza. Astfel, in acest laborator vom vorbi despre: 
 +  * Diferitele abordari ale structurii unui procesor 
 +  * Ce probleme si ce imbunatatiri ale performantei au motivat aparitia acestor structuri 
 +  * Arhitecturi CISC/RISC bazate pe procesoare Intel/AMD inclusiv modele standard de placi de baza 
 +  * Comparatii intre abordarile curente utilizate in implementarea procesoarelor moderne si a placilor de baza 
 +  * Solutii de interconectare a procesoarelor AMD si Intel Hypertransport si QuickPath Interconnect
  
-Prima arhitectura NVIDIA complet programabila a fost G80 (ex. [[http://​www.nvidia.com/​page/​8800_tech_briefs.html|Geforce 8800]], lansat in anul 2006). Cu aceasta arhitectura s-a trecut de la unitati hardware fixe vertex/​pixel la cele de unified shader care puteau procesa atat vertex/​pixel cat si geometry. Evolutia arhitecturilor GPU de la NVIDIA este detaliata [[http://​s08.idav.ucdavis.edu/​luebke-nvidia-gpu-architecture.pdf|aici]].+===== Instruction level parallelism =====
  
-Implementarea NVIDIA pentru GPGPU se numeste CUDA (Compute Unified Device Architecture) si permite utilizarea limbajului C pentru programarea pe GPU-urile proprii. Lista de GPU-uri ce suporta API-ul CUDA sau OpenCL se regaseste pe site-ul oficial [[https://​www.geforce.com/​hardware/​technology/​cuda/​supported-gpus|aici]] sau pe wiki [[https://​en.wikipedia.org/​wiki/​List_of_Nvidia_graphics_processing_units|aici]]Fiecare noua arhitectura are un codename (ex FermiPascal) si este reprezentata ​de un "​compute capability"​ (list [[https://​developer.nvidia.com/​cuda-gpus|aici]]). Cu cat arhitectura ​este mai noua, cu atat sunt suportate ​mai multe facilitati din API-urile CUDA si OpenCL.+Masina Turing executa cate o instructiune la un moment datCand programatorul scrie un programii este foarte simplu sa considere ca programul sau va fi executat in acest mod. Pe de alta parte, o masina care executa cate o instructiune ​este mai lenta decat una care executa mai multe instructiuni in paralel. Pentru a cumula avantajele celor doua abordariar trebui ca programatorul sa poata inca scrie cod ca pentru o masina seriala, iar procesorul sa execute acest cod cu un nivel de paralelism cat mai ridicat. Cum este posibil asa ceva? Cineva trebuie sa faca trecerea dintre perspectiva seriala a programatorului si perspectiva paralela pe care ne-ar placea sa o aiba procesorul. Acest cineva poate fi ori un compilator, ori un hardware specializat aflat in structura procesorului.
  
-Unitatea GPU este potrivita pentru paralelismul ​de date SIMD (Single ​Instruction ​Multiple Data), astfel aceleasi ​instructiuni ​sunt executate ​in paralel ​pe mai multe unitati de procesareDatorita faptului ca acelasi program este executat pentru fiecare element de datesunt necesare mai putine elemente pentru controlul fluxului. Si deoarece calculele sunt intensive computational,​ latenta accesului la memorie ​poate fi ascunsa prin calcule ​in locul unor cache-uri mari pentru ​date.+Aceste considerente influenteaza structura procesorului si au dus la aparitia conceptului ​de Instruction ​Level Paralelism (ILP). Astfelprocesorul ia instructiuni ​dintr-un singur flux de control, le decodifica si executa ​in paralel. ​De exempluun procesor cu ILP poate sa scrie simultan rezultatele a doua instructiuni ​in registre, sa faca operatii aritmetice ​pentru ​alte trei, sa citeasca operanzii pentru alte doua, sa decodifice alte patru si sa ia (fetch) din fluxul de intrare inca patru instructiuni
  
-Motivul discrepantei intre performanta paralela dintre CPU si GPU este faptul ca GPU sunt specializate pentru procesare masiv paralela si intensiva computational (descrierea perfecta a taskurilor ​de randare grafica) si construite in asa fel incat majoritatea tranzistorilor de pe chip se ocupa de procesarea datelor in loc de cachingul datelor si controlul fluxului executiei. ​+<note important>​ Cateva implementari ​de ILP includ:
  
-La GPU-urile NVIDIAun Streaming Processor (SP) este un microprocesor cu executie secventialace contine un pipelineunitati aritmetico-logice (ALU) si de calcul in virgula mobila ​(FPU). Nu are un cache, fiind bun doar la executia multor operatii matematiceUn singur SP nu are performante remarcabile, insa prin cresterea numarului ​de unitati, se pot rula algoritmi ​ce se preteaza paralelizarii masive.+  * Pipeline: In acelasi ciclu de ceasprocesorul scrie rezultatul unei instructiuni in registreexecuta operatia aritmetica a instructiunii urmatoare, si citeste operanzii instructiunii ​de dupa instructiunea urmatoare ​(la doua instructiuni dupa prima) . 
 +  * VLIW (Very Long Instruction Word): Lanseaza mai multe instructiuni in acelasi ciclu de ceasCompilatorul trebuie sa se asigure ca nu exista dependente de date intre acestea. La procesoarele superscalarenumarul de unitati de executie este transparent pentru setul de instructiuni. VLIW este insa constient de numarul ​de unitati ​de executie. 
 +  * Superscalar:​ Lanseaza mai multe instructiuni in acelasi ciclu de ceas. Dependenta de date este insa verificata de hardware aditional. Daca nu pot fi lansate in paralel, se va executa cate o instructiune secvential (neavand suport din partea compilatorului,​ exista si aceasta posibilitate). 
 +  * Planificare-dinamica:​ Instructiunile sunt reordonate in timp ce sunt executate. In modul acesta, poate sa gaseasca usor instructiuni care nu au dependenta de date intre ele, pentru a fi executate simultan. 
 +</​note>​
  
-SP impreuna cu Special Function Units (SFU) sunt incapsulate intr-un Streaming Multiprocessor (SM/SMX)Fiecare SFU contine unitati pentru inmultire ​in virgula mobilautilizate ​pentru ​operatii transcendente (sincos) si interpolare. MT se ocupa cu trimiterea instructiunilor pentru executie la SP si SFU.+ILP-ul mareste asadar performanta procesoruluiDar de ce nu executam toate instructiunile deodata ​in paralel? Acesta ar fi de fapt modul cel mai rapid de a executa un program. Acest lucru nu se intampla deoarece ILP-ul are si anumite limitarirespectiv:​ 
 +  * Dependenta de date: Daca rezultatul instructiunii A este operand ​pentru ​instructiunea Batunci evident B nu poate fi executata inainte ca A sa se fi terminat. 
 +  * Numar limitat de unitati functionale:​ Daca avem 5 sumatoare in procesor, nu putem executa mai mult de 5 sume simultan. 
 +  * Numar limitat de instructiuni lansate: Daca unitatea de lansare de instructiuni poate lansa maxim 5 instructiuni simultan, un program ​cu 500 de instructiuni va avea nevoie de 100 de operatii ale acestei unitati.  
 +  * Numar limitat de registre.
  
-Pe langa acestea, exista si un cache (de dimensiuni reduse) pentru instructiuni,​ unul pentru date precum si memorie shared, partajata de SP-uri. Urmatorul nivel de incapsulare este Texture / Processor Cluster (TPC). Acesta contine SM-uri, logica de control si un bloc de handling pentru texturi. Acest bloc se ocupa de modul de adresare al texturilor, logica de filtrare a acestora precum si un cache pentru texturi.+Mai multe detalii despre Explicitly Parallel Instruction Computing pot fi gasite aici [[:​asc:​extra:​epic?​480 | EPIC]].
  
-{{:​asc:​lab11:​cuda-arch.png?​direct&​720|}}+===== Comparatie CISC vsRISC =====
  
-Filosofia din spatele arhitecturii este permiterea rularii unui numar foarte mare de threaduri. Acest lucru este facut posibil ​prin paralelismul existent ​la nivel hardware.+Cand a aparut CISC, ideea era sa se aduca in hardware stilul ​de programare specific unui limbaj care sa se aproprie (pe cat e posibil la nivelul ​hardware) de un limbaj cat mai inaltAstfel, instructiunile complexe au acelasi efect ca micile secvente de instructiuni simple. Implemenarea acestor instructiuni complexe in hardware insemana insa:
  
-Documentatia NVIDIA recomanda rularea unui numar cat mai mare threaduri ​pentru a executa un task. Arhitectura CUDA de exemplu suporta zeci de mii de threaduri, numarul acestora fiind mult mai mare decat unitatile fizice existente pe chip. Acest lucru se datoreaza faptului ca un numar mare de threaduri poate masca latenta accesului la memorie.+<note important>​ 
 +  * Hardware complex 
 +  * Locul ocupat de hardul pentru instructiunile complexe ar fi putut fi utilizat ​pentru a avea mai multe unitati ​de executie (si deci grad de paralelism ​mai mare
 +  * Secvente ​de microcod, care sunt lente comparativ cu restul procesorului 
 +</​note>​
  
-Urmarind acelasi model modular din punct de vedere al arhitecturiithreadurile ​sunt incapsulate in blocuri (thread blocks / warps), iar blocurile in grile (thread grid). Fiecare thread este identificat prin indexul threadului in bloc, indexul blocului in grila si indexul grilei. Indexurile threadurilor si ale blocurilor pot fi uni/​bi/​tri-dimensionale,​ iar  indexul grilei poate fi uni sau bi-dimensional. Acest tip de impartire are rolul de a usura programare pentru probleme ce utilizeaza structuri ​de date cu mai multe dimensiuniSe poate observa ca thread-urile dintr-un thread block trebuie sa execute cat mai multe instructiuni identice spre a nu irosi resurse.+Datorita setului redus de instructiuni de asamblarecompilatoarele optimizate pentru RISC sunt capabile sa organizeze mai eficient fluxul ​de instructiuni ​de asamblare. Pe de alta parte insa, compilatoarele optimizate pentru RISC necesita ​mai mult timp de compilare decat cele pentru CISCAceasta deoarece ​trebuie sa se ocupe si de managementul benzii de asamblare, anticiparea ramificatiilor (branch prediction) sau reorganizarea codului.
  
-{{:​asc:​lab11:​thread.blocks.jpg?​direct&​360|{{thread.blocks.jpg|''​Structura threadurilor in blocuri''​}}+Ca principiu, o arhitectura RISC are mai multe registre generale, in timp ce CISC are mai multe registre specialePractic toate procesoarele moderne imprumuta atat caracteristici CISC cat si RISC.
  
-Threadurile dintr-un bloc pot coopera prin partajarea ​de date prin intermediul memoriei shared ​si prin sincronizarea executiei. Functia de bariera functioneaza doar pentru threadurile dintr-un bloc. Sincronizarea nu este posibila ​la alt nivel (intre blocuri/​grila etc.). Mai multe explicatii se regasesc in [[http://​cs.curs.pub.ro/​wiki/​asc/​asc:​lab9:​index|Laboratorul 9]].+Exista trei tipuri ​de categorii de instructiuni CISC, si anume:  
 +  * Aritmetico-logice 
 +  * De control secvential 
 +  * De acces la memorie
  
-===== Ierarhia ​de memorie ​=====+Formatul instructiunilor RISC are o lungime fixa, cu lungimea unei instructiuni in general egala cu lungimea cuvantului ​de memorie; in cazul CISC, lungimea unei instructiuni variaza in functie de formatul instructiunii. RISC are un numar mic de moduri de adresare, spre deosebire de CISC, care are un numar mare de moduri de adresare (utilizate mai rar).
  
-Intelegerea ierharhiei ​de memorie este esentiala in programarea eficienta a unitatii GPUCapacitatea mare de executie in paralel ​unui GPU necesita ascunderea latentei de acces catre memoria principala ​(fie VRAM pentru dGPU sau RAM pentru iGPU).+Setul de instructiuni RISC este orientat pe registre (peste 32 de registre). Pentru ca accesul la memorie este mult mai lent decat lucrul cu registrele, RISC incurajeaza lucrul cu acestiaFace acest lucru prin cresterea numarului ​de registre si prin limitarea explicita ​acceselor la memorie. In general instructiunile au 2 operanzi ​(registresi un registru destinatie.
  
-{{:asc:​lab11:​mem.hierarchy.png?​direct|Ierarhia memoriei in CUDA}}+<note important>​ 
 +In cadrul arhitecturilor RISC exista o limitare explicita, si anumesingurul mod de acces la memorie este prin load si store. Aceasta se deosebeste fundamental de CISC care are instructiuni cu operanzi locatii de memorieTotusi, desi RISC impune aceasta disciplina de lucru cu memoria, doar 20-25% din codul unui program e reprezentat de operatii de tip load sau store 
 +</​note>​
  
-**Register File** +Fie urmatorul exemplu de cod in C: 
-<​code ​sh+<​code ​C
-/* marcam pentru compilator regValPi in register file */ +#include <stdio.h>
-__private float regValPi = 3.14f; +
-/* compilatorul cel mai probabil oricum incadreaza regVal2Pi ca registru */ +
-float regVal2Pi = 2 * 3.14f; +
-</code>+
  
-  *Cea mai rapida forma de memorie de pe GPU +#define N 100
-  *Accesibila doar de catre thread, durata de viata este aceeasi ca si a threadului +
-  *Un kernel complex poate determina folosirea unui numar mare de registrii si astfel: +
-    * limitarea executiei multor thread-uri simultan +
-    * register spill, atunci cand valorile registrilor sunt salvate in memoria globala+
  
-**Local Memory** +int main() {
-<code sh> +
-/* fiecare work item salveaza un element */ +
-__local float lArray[lid] = data[gid];​ +
-</​code>​+
  
-  *in functie de implementarea hardware100GB/sec -> 2TB/sec +    int a[N]i; 
-  ​*pentru GPU o memorie rapida, actioneaza ca un cache L1/alt register file, la CPU de regula este doar o portiune din RAM +    for(i=0;​i<​N;​i++) 
-  *accesibila tuturor threadurilor dintr-un bloc (warp/​wavefront), durata de viata este aceeasi ca si blocului +        ​a[i] = 2*i; 
-  ​*trebuie evitate conflictele de access (bank conflicts)+    ​return 0; 
 +}
  
-**Constant Memory** 
-<code sh> 
-__const float pi = 3.14f 
 </​code>​ </​code>​
-  * in functie de implementarea hardware, 100GB/sec -> 1TB/sec 
-  * in general performanta foarte buna, (cache L1/L2, zona dedicata), 
-  * are durata de viata a aplicatiei kernel 
  
-**Global Memory** +Rezultatul compilarii codului pana la nivel de asamblare, pe arhitectura CISC si RISC este urmatorul:​ 
-<​code ​sh+<spoiler Arata codul> 
-__kernel void process(__global float* data)... }+  ​RISC: 
 +<​code ​asm
 +root@raspberrypi:​~#​ cat risc.s 
 + .arch armv6 
 + .eabi_attribute 28, 1 
 + .eabi_attribute 20, 1 
 + .eabi_attribute 21, 1 
 + .eabi_attribute 23, 3 
 + .eabi_attribute 24, 1 
 + .eabi_attribute 25, 1 
 + .eabi_attribute 26, 2 
 + .eabi_attribute 30, 6 
 + .eabi_attribute 34, 1 
 + .eabi_attribute 18, 4 
 + .file "​risc.c"​ 
 + .text 
 + .align 2 
 + .global main 
 + .arch armv6 
 + .syntax unified 
 + .arm 
 + .fpu vfp 
 + .type main,​ %function 
 +main: 
 + @ args = 0, pretend = 0, frame = 408 
 + @ frame_needed = 1, uses_anonymous_args = 0 
 + @ link register save eliminated. 
 + str fp, [sp, #-4]! 
 + add fp, sp, #0 
 + sub sp, sp, #412 
 + mov r3, #0 
 + str r3, [fp, #-8] 
 + b .L2 
 +.L3: 
 + ldr r3, [fp, #-8] 
 + lsl r2, r3, #1 
 + ldr r3, [fp, #-8] 
 + lsl r3, r3, #2 
 + sub r1, fp, #4 
 + add r3, r1, r3 
 + str r2, [r3, #-404] 
 + ldr r3, [fp, #-8] 
 + add r3, r3, #1 
 + str r3, [fp, #-8] 
 +.L2: 
 + ldr r3, [fp, #-8] 
 + cmp r3, #99 
 + ble .L3 
 + mov r3, #0 
 + mov r0, r3 
 + add sp, fp, #0 
 + @ sp needed 
 + ldr fp, [sp], #4 
 + bx lr 
 + .size main,​ .-main 
 + .ident "​GCC: ​(Raspbian 8.3.0-6+rpi18.3.0"​ 
 + .section .note.GNU-stack,"",​%progbits
 </​code>​ </​code>​
-  * in functie de implementarea hardware, 30GB/sec -> 500GB/sec 
-  * Video RAM (VRAM), de regula cu o capacitate intre 1GB si 12GB in functie de placa video 
-  * memorie dedicata specializata doar pentru placile grafice discrete (GPU-urile integrate in CPU folosesc RAM) 
-  * in general latime mare de banda (256-512 biti) si chipuri de memorii de mare viteza (GDDR5) 
  
-**Host Memory ​(RAM)** +  ​CISC: 
-  * in general4GB/​sec ​-> 30GB/sec +<code asm> 
-  * pentru acces din kernel trebuie transfer/​mapare explicita RAM->VRAM pe partea de host/CPU +[razvan.dobre@fep7-1 ~]$ cat cisc.s 
-  * memoria RAM accesibila direct de CPU si indirect de GPU via DMA si magistrala PCIe + .file "​cisc.c"​ 
-  * viteza de transfer ​(throughput/​latentaeste limitata de magistrala PCIe cat si de memoria RAM+ .text 
 + .globl main 
 + .type main,​ @function 
 +main: 
 +.LFB0: 
 + .cfi_startproc 
 + pushq %rbp 
 + .cfi_def_cfa_offset 16 
 + .cfi_offset 6, -16 
 + movq %rsp, %rbp 
 + .cfi_def_cfa_register 6 
 + subq $296, %rsp 
 + movl $0, -4(%rbp
 + jmp .L2 
 +.L3: 
 + movl -4(%rbp)%eax 
 + leal (%rax,​%rax),​ %edx 
 + movl -4(%rbp), %eax 
 + cltq 
 + movl %edx, ​-416(%rbp,​%rax,​4) 
 + addl $1, -4(%rbp) 
 +.L2: 
 + cmpl $99, -4(%rbp) 
 + jle .L3 
 + movl $0, %eax 
 + leave 
 + .cfi_def_cfa 7, 8 
 + ret 
 + .cfi_endproc 
 +.LFE0: 
 + .size main,​ .-main 
 + .ident "​GCC:​ (GNU) 4.8.5 20150623 (Red Hat 4.8.5-44)"​ 
 + .section .note.GNU-stack,"",​@progbits 
 +</​code>​ 
 +</​spoiler>​ 
 +===== Arhitectura Intel =====
  
-Caracteristici GPU K40m (coada hpsl), via query device properties CUDA+O schema clasica pentru un sistem CISC este prezentata in Figura 1. Aici se poate distinge usor in partea de sus Procesorul, legat de restul sistemului prin Front-Side-Bus ​(FSBde 400/​533/​800MHz catre North Bridge (i.e. 82865PE MCH). Pe North Bridge se afla controllerul de memoriesi ca atare si memoriile sunt conectate direct aici prin canale intre 2.1GB/s si 3.2GB/s. De asemenea pe North Bridge se conecteaza atat placa grafica (AGP 8x/4x) cat si interfata de retea de mare viteza Gigabit Ethernet. 
 +  
 +{{ :​asc:​lab4:​865pe.jpg?​640 |Figura 1. Schema bloc a Chipsetului Intel® 865PE }}
  
-<code sh> +La randul sau North Bridge-ul este conectat printr-o legatura de 266MB/s catre South Bridge (i.e82801EB ICH5 82801ER ICH5R)Dupa cum se poate vedeaNorth Bridge-ul impreauna cu South Bridge-ul formeaza impreauna ceea ce se numeste Chipsetul Intel® 865PEUrmarind in continuare schema din Figura 1se observa ca pe South Bridge se conecteaza o multitudine de componente periferice cu o viteza si rata de transfer de date considerabil mai scazuta decat elementele conectate pe North Bridgecum ar fiAC97 (placa audio), porturi ATA si Serial ATAporturi USBsistem de management si de control al consumuluietc.
-Device 0: "Tesla K40m"​ +
-  CUDA Driver Version ​Runtime Version ​         9.1 / 9.+
-  CUDA Capability Major/Minor version number: ​   3.+
-  Total amount of global memory: ​                11441 MBytes (11996954624 bytes) +
-  (15) Multiprocessors(192) CUDA Cores/​MP: ​    2880 CUDA Cores +
-  GPU Max Clock rate:                            745 MHz (0.75 GHz) +
-  Memory Clock rate:                             3004 Mhz +
-  Memory Bus Width: ​                             384-bit +
-  L2 Cache Size:                                 ​1572864 bytes +
-  Maximum Texture Dimension Size (x,y,z)         ​1D=(65536),​ 2D=(65536, 65536), 3D=(4096, 4096, 4096) +
-  Maximum Layered 1D Texture Size, (num) layers ​ 1D=(16384), 2048 layers +
-  Maximum Layered 2D Texture Size, (num) layers ​ 2D=(16384, 16384), 2048 layers +
-  Total amount of constant memory: ​              65536 bytes +
-  Total amount of shared memory per block: ​      49152 bytes +
-  Total number of registers available per block: 65536 +
-  Warp size:                                     32 +
-  Maximum number of threads per multiprocessor: ​ 2048 +
-  Maximum number of threads per block          1024 +
-  Max dimension size of a thread block (x,y,z): (1024102464) +
-  Max dimension size of a grid size    (x,y,z): (2147483647,​ 65535, 65535) +
-  Maximum memory pitch: ​                         2147483647 bytes +
-  Texture alignment: ​                            512 bytes +
-  Concurrent copy and kernel execution: ​         Yes with 2 copy engine(s) +
-  Run time limit on kernels: ​                    No +
-  Integrated GPU sharing Host Memory: ​           No +
-  Support host page-locked memory mapping: ​      Yes +
-  Alignment requirement for Surfaces: ​           Yes +
-  Device has ECC support: ​                       Enabled +
-  Device supports Unified Addressing (UVA): ​     Yes +
-  Device PCI Domain ID / Bus ID / location ID:   0 / 8 / 0 +
-</​code>​+
  
-Caracteristici GPU M2070 (coada dp)via query device properties CUDA+Din aceasta schema se poate usor deduce ca punctul vulnerabil al acestor sisteme il constituie integrarea controllerului de memorie pe North Bridgesi in special legaturile de marime limitata intre South si North Bridge, precum si intre North Bridge si procesor. In mod evident, dimensionarea acestora este un compromis de design al sistemelor Intel, menit sa deserveasca majoritatea sistemelor hardware bazate pe acest Chipset, si a aplicatiilor ce ruleaza pe ele. 
  
-<code sh> +Pentru a atinge insa performante mai inaltemai ales in contextul aparitiei sistemelor multi-procesor si multi-coreeste insa nevoie de imbunatatiri ale acestei abordaricum se poate observa la Chipsetul Intel 7300 din Figura 2. 
-Device 0: "Tesla M2070"​ +  
-  CUDA Driver Version / Runtime Version ​         9.1 / 9.1 +{{ :asc:lab4:intel7300chipset.jpg?​640 |Figura 2. Schema bloc Chipsetului Intel® 7300}}
-  CUDA Capability Major/Minor version number: ​   2.0 +
-  Total amount of global memory: ​                5302 MBytes (5559156736 bytes) +
-  (14) Multiprocessors( 32) CUDA Cores/​MP: ​    448 CUDA Cores +
-  GPU Max Clock rate:                            1147 MHz (1.15 GHz) +
-  Memory Clock rate:                             1566 Mhz +
-  Memory Bus Width: ​                             384-bit +
-  L2 Cache Size:                                 ​786432 bytes +
-  Maximum Texture Dimension Size (x,y,z)         ​1D=(65536),​ 2D=(65536, 65535), 3D=(2048, 2048, 2048) +
-  ​Maximum Layered 1D Texture Size, (num) layers ​ 1D=(16384), 2048 layers +
-  Maximum Layered 2D Texture Size, (num) layers ​ 2D=(16384, 16384), 2048 layers +
-  Total amount of constant memory              65536 bytes +
-  Total amount of shared memory per block      49152 bytes +
-  Total number of registers available per block32768 +
-  Warp size:                                     32 +
-  Maximum number of threads per multiprocessor: ​ 1536 +
-  Maximum number of threads per block: ​          ​1024 +
-  Max dimension size of thread block (x,y,z): (1024, 1024, 64) +
-  Max dimension size of a grid size    (x,y,z): (65535, 65535, 65535) +
-  Maximum memory pitch: ​                         2147483647 bytes +
-  Texture alignment: ​                            512 bytes +
-  Concurrent copy and kernel execution: ​         Yes with 2 copy engine(s) +
-  Run time limit on kernels: ​                    No +
-  Integrated GPU sharing Host Memory: ​           No +
-  Support host page-locked memory mapping: ​      Yes +
-  Alignment requirement for Surfaces: ​           Yes +
-  Device has ECC support: ​                       Enabled +
-  Device supports Unified Addressing (UVA): ​     Yes +
-  Device PCI Domain ID / Bus ID / location ID:   0 / 20 / 0 +
-</​code>​+
  
-===== Optimizare accesului la memorie ​=====+Aici se poate observa conectarea a patru procesoare catre North Bridge (MCH) prin canale distincte FSB, de 1066MHz fiecare, menite sa asigure o alimentare eficienta cu date a acestora. Controllerul de memorie ​ramane pe North Bridge insa memoria se conecteaza prin patru canale de 8GB/s. Apar pe North Bridge conexiuni multiple PCI-Express (PCI-E), placa de retea devine doar "una dintre acestea";,​ iar largimea de banda este crescuta pentru fiecare dintre acest componente fata de versiunile anterioare de Chipseturi. ​
  
-Modul cum accesam memoria influenteaza foarte mult performanta sistemuluiCum putem avea arhitecturi foarte diferite din punctul ​de vedere al ierarhiei de memorie este important de inteles ​ca nu putem dezvolta un program care sa ruleze optim in toate cazurileUn program CUDA este portabil caci poate fi usor rulat pe diferite arhitecturi NVIDIA CUDAinsa de cele mai multe ori trebuie ajustat in functie ​de arhitectura ​pentru ​o performanta optima.+Pentru simplitate, si conexiunea intre North Bridge si South Bridge (631xESB I/O Controller) este realizata prin conexiuni PCI Express 2x sau chiar 4xEste sugestiv ​de asemenea faptul ​ca South Bridge-ul poarta acum numele de "I/O Controller"​ si se consfinteste astfel si prin nume direct rolul South Bridge-uluiPorturile ATA sunt inlocuite acum de SATA si PATA, apar din nou numeroase porturi PCI-E si PCI-X, USB, de gestiune a consumului, a biosului sau placi aditionale ​de retea pentru ​management.
  
-In general pentru arhitecturile ​de tip GPUmemoria locala este impartita in module de SRAM identicedenumite bancuri ​de memorie (memory banks)Fiecare banc contine o valoare succesiva de 32 biti (de exempluun int sau un float), astfel incat accesele consecutive intr-un array provenite ​de la threaduri consecutive sa fie foarte rapidBank conflicts au loc atunci cand se fac cereri multiple asupra datelor aflate in acelasi banc de memorie.+Din prezentarea celor doua chipseturi 865PE si 7300 se poate vedea evolutia arhitecturala din jurul procesoarelor CISC de la Intelcu o modularitate si o flexibilitate considerabil mai mare a chipsetului 7300care ofera un potential ​de performanta mult crescut pentru sistemele ce il utilizeazaAtentieatat procesoarele Intel considerate pentru Chipset-urile prezentate, cat si cele AMD din sectiunile urmatoare, sunt din familii dedicate sistemelor ​de inalta performanta si cele mai puternice din clasa lorDiscutia prezentata este insa relevanta, la o scara corespunzator mai scazuta, si pentru celelalte sisteme si procesoare oferite ​de cele doua mari firme.
  
-<​note ​important+Anexa: 
-Conflictele ​de access la bancuri ​de memorie ​(cachepot reduce semnificativ performanta.+Magistrale uzuale si largimea lor de banda maxima  
 + 
 +^ Bus ^ Max Bandwidth ^ 
 +|PCI|132 MB/s| 
 +|AGP 8X|2,100 MB/s| 
 +|PCI Express 1x|250 [500]* MB/s| 
 +|PCI Express 2x|500 [1000]* MB/s| 
 +|PCI Express 4x|1000 [2000]* MB/s| 
 +|PCI Express 8x|2000 [4000]* MB/s| 
 +|PCI Express 16x|4000 [8000]* MB/s| 
 +|PCI Express 32x|8000 [16000]* MB/s| 
 +|IDE (ATA 100)|100 MB/s| 
 +|IDE (ATA 133)|133 MB/s| 
 +|SATA|150 MB/s| 
 +|Gigabit Ethernet|125 MB/s| 
 +|IEEE 1394B [Firewire]|100 MB/s| 
 + 
 +<​note ​tip
 +PCI Express este o magistrala seriala (datele pot circula simultan in ambele directii). In tabelul ​de mai sus cele doua valori pentru largimea ​de banda corespund largimii de banda intr-o singura directie respectiv in ambele directii ​(combinat).
 </​note>​ </​note>​
  
-Cand are loc un bank conflict, hardware-ul serializeaza operatiile cu memoria (warp/​wavefront serialization),​ si face astfel toate threadurile sa astepte pana cand operatiile de memorie sunt efectuate. In unele cazuri, daca toate threadurile citesc aceeasi adresa de memorie shared, este invocat automat un mecanism de broadcast iar serializarea este evitata. Mecanismul de broadcast este foarte eficient si se recomanda folosirea sa de oricate ori este posibil.  +===== Arhitectura AMD Hammer =====
-Spre exemplu daca linia de cache este alcatuita din 16 bancuri de memorie. Avem urmatoarele situatii care impacteaza performanta accesului la cache.+
  
-<code sh> +Din familia Hammer, sau AMD64, face parte cel mai putermic procesor de la AMD, si anume OpteronOpteronul este echivalentul familiilor Intel Itanium si Intel Xeon, destinat serverelor si sitemelor de inalta performantaOpteron este un procesor out-of-order si in interiorul unitatii de executie ordinea instructiunilor este schimbata, pentru a maximiza eficientaPentru utilizatorul extern insa, instructiunile par a se executa in aceeasi ordine in care au fost lansateDe asemeneael este 3-way superscalaradica poate decodaexecuta si incheia trei instructiuni x86 la fiecare ciclu masina. Desi poate lucra in paralel la 3 instructiuniaceasta nu insemna neaparat ca cele 3 instructiuni sunt procesate in intregime pe acea perioda de ceasOpteronul a fost creat pentru a putea lucra in sisteme multiprocesor si fiind primul care a oferit o scalabilitate sporita, el a acaparat la vremea respectiva o portiune semnificativa din piata comerciala de servere
-__kernel func(...) {  +
-...  +
-   ​__local int *array; +
-   x = array[get_local_id(0)]; ​   // performanta 100%0 bank conflicts +
-   x = array[get_local_id(0)+1]; ​ // performanta 100%0 bank conflicts +
-   x = array[get_local_id(0)*4]; ​ // performanta 25%4 bank conflicts +
-   x = array[get_local_id(0)*16];​ // performanta 6% 16 bank conflicts +
-... +
-+
-</​code>​+
  
-In cazul arhitecturilor ​de tip CPU, memoria locala ​este doar regiune din RAMOptimizarile pentru ​tine datele critice in memoria locala pentru GPU nu ar prezenta deci aceleasi imbunatatiri ​de performanta.+Printre cele mai importante imbunatatiri arhitecturale cu care vine Hammer (generatia 8) fata de Athlon si AthlonXP (generatia 7) se numara: doua stagii in plus la pipeline, algoritmi imbunatatiti de predictie a ramificatiilor,​ suport pentru SSE2 (Streaming Multimedia Instructions),​ controller de memorie integrat in CPU si extensie completa pentru setul de instructiuni pe 64 biti pentru x86. Toate procesoarele Hammer au viteze de maxim 2.5GHz. Pe de alta parteIntel s-a concentrat mai mult pe cresterea vitezei, fara a se preocupa excesiv de paralelism, acest trend a fost insa oprit din 2007-2008, cand si Intel, si AMD au trecut la producerea de sisteme multi-core. Acest lucru este important in lumea serverelor, unde disiparea caldurii ​este o problema principalaIntel insa venit si cu o noua "​arma"​ si anume tehnologia de integrare bazata pe dielectrici High-K, ce a dus la o scadere drastica a consumului, si a permis de asemenea performante considerabile la frecvente ​de executie scazute.
  
-===== Debug aplicatii CUDA =====+{{ :​asc:​lab4:​amdquadcorephenom.jpg?​640 |Figura 3. Arhitectura Procesorului Quad-Core Phenom}}
  
-Cele mai des intalnite probleme sunt cele de acces invalid ​la memorieNu de putine ori vom observa ca aceste accese invalide pot crea efecte secundare sau erori ce apar/sunt semnalate abia ulterior. +Dupa cum am spus, Hammer pastreaza compatibilitatea cu 16 si 32 biti (fata de Intel, care renunta complet ​la x86 si trece la IA-64)Pentru a putea face acest lucru, Hammer are doua moduri ​de operare: Legacy- si Long-ModeLong Mode este subdivizat si el in Compatibility mode si 64bit Mode. Legacy mode este destinat exclusiv sistemelor ​de operare pe 16 si 32 biti. Compatibility mode este destinat sistemelor de operare pe 64 de biti, dar care ruleaza programe scrise pentru 32 biti. Astfel, desi programul ​in sine nu beneficiaza ​de facilitatile 64 bitimanagamentul resurselor, facut de sistemul de operare pe 64 de biti beneficiaza de totate avantajele date de rularea pe 64 biti.
-Sa luam de exemplu cazul in care 1 thread acceseaza 1 element ​de date si sa dam in executie mai multe thread-uri decat elemente ​de memorie alocate (8x16 ~ 128 thread-uri vs 100 elemente)dat fiind MAGNITUDE=1.+
  
-<code C> +{{ :​asc:​lab4:​amdquadcorephenomcache.jpg?640 |Figura 4Arhitectura Cache-ului la Phenom}}
-..+
-#define MAGNITUDE ​      (1) +
-#define NUM_BLOCKS ​     8 * MAGNITUDE +
-#define NUM_THREADS ​    16 +
-#define NUM_ELEM ​       100 * MAGNITUDE+
  
-__global__ void kernel_compute(int* data) { +{{ :​asc:​lab4:​corex64.jpg?480 |Figura 5Schema Bloc a unui Core AMD x64}}
-    int idx = blockIdx.x * blockDim.x + threadIdx.x; +
-    // invalid access +
-    data[idx] = 1111 * idx; +
-}+
  
-int main(int argc, char *argv[]) { +==== Controller-ul de memorie integrat in procesor la AMD ====
-    int* data NULL;+
  
-    HANDLE_ERROR( cudaMalloc(&​data1 * sizeof(int)) );+In generalun procesor foloseste doua cipuri de pe placa de baza pentru a accesa memoria si perifericele. Aceste sunt numite North Bridge si South Bridge, dupa cum au fost descrise in sectiunile anterioare. Se observa ca North Bridge-ul joaca un rol esential, el facand legatura cu memoria. De acea, de la generatia Hammer, AMD a integrat in cipul procesorului North Bridge-ul. In felul acesta se obtine o latenta de acces la memorie redusa cu cel putin 20%. Acest controller are o legatura de 128 biti cu memoria. In plus, el functioneaza dupa ceasul procesorului,​ acest lucru marind inca o data viteazapermite de asemenea ca marirea frecventei de ceas a procesorului sa imbunatateasca si performantele controller-ului.
  
-    // launch kernel +Controller-ul de memorie are grija si de coerenta cacheului. Cipul integrat se concentreaza acum doar pe comunicarea cu memoria. Alte functionalitati ale North Bridge-uluicum erau de exemplu comunicarea cu AGP sau Placa de Retea, au fost mutate pe un cip extern.
-    kernel_compute<<<​NUM_BLOCKSNUM_THREADS>>>​(data);​ +
-    HANDLE_ERROR( cudaDeviceSynchronize() );+
  
-    return 0; +{{ :​asc:​lab4:​opteronmem.jpg?​640 |Figura 6. Controller de memorie integrat de la procesoarele Opteron}}
-} +
-</​code>​+
  
-Daca rulam programul vom observa ca nu intoarce nici o eroare. Deoarece sunt putine accese invalide HW-ul nu semnaleaza vreo problema.+==== Hypertransport ====
  
-CUDA insa ofera aplicatii care sa analizeze si sa detecteze accese invalide cu precizie ridicataDaca rulam de exemplu [[https://​docs.nvidia.com/​cuda/​cuda-memcheck/​index.html|cuda-memcheck]] vom vedea instant ca avem accese invalide ​la memorie.+Hypertransport este tehnologie pentru I/O dezvoltata initial de AMDEa este o alternativa la sistemele actuale ​de busFoloseste legaturi duble, punct la punct, pentru a lega componentele intre eleEste, in termeni de retele, echivalentul unei legaturi full-duplex punct la punct fata de o topologie buss.
  
-<code sh+<note tip
-[stefan_dan.ciocirlan@fep8 ~]$ srun --pty -p hpsl /bin/bash +O astfel de lagatura poate avea intre 2 si 64 biti, si poate opera la viteze de 400Mhz-2.6GHzDatele sunt impachetate ​ si trimise folosind un protocol care  prevede trimiterea de pachete multiplu de bytescu marimi intre 4 si 64 bytes
-[stefan_dan.ciocirlan@hpsl-wn01 ~]$ singularity run --nv docker://​nvidia/​cudagl:​9.1-devel-centos7 +</note>
-INFO:    Using cached SIF image +
-Singularity>​ cd skl/​task0/​ +
-Singularity>​ cuda-memcheck ./task0 +
-========= CUDA-MEMCHECK +
-========= Invalid __global__ write of size 4 +
-========= ​    at 0x00000050 in kernel_compute(int*) +
-========= ​    by thread (15,0,0) in block (225,0,0) +
-========= ​    ​Address 0x13040e387c is out of bounds +
-========= ​    Saved host backtrace up to driver entry point at kernel launch time +
-========= ​    Host Frame:/.singularity.d/​libs/​libcuda.so.1 [0x20d50a] +
-========= ​    Host Frame:​./​task0 [0x1a1fb] +
-========= ​    Host Frame:​./​task0 [0x3744e] +
-========= ​    Host Frame:​./​task0 [0x35aa] +
-========= ​    Host Frame:​./​task0 [0x34a5] +
-========= ​    Host Frame:​./​task0 [0x34bf] +
-========= ​    Host Frame:​./​task0 [0x33f8] +
-========= ​    Host Frame:/​lib64/​libc.so.6 (__libc_start_main + 0xf5) [0x22555] +
-========= ​    Host Frame:​./​task0 [0x31e9] +
-========= +
-... +
-========= +
-========= ERROR SUMMARY: 3841 errors +
-</code>+
  
-Daca insa avem multe accese invalide (de exe '#​define MAGNITUDE (1024 * 1024)'​) o sa vedem ca API-ul arunca erori la executia ​de kernel.+<note tip> 
 +Hypertransport e compatibil cu PCI, de aceea a fost usor de introdus. El poate lucra in doua moduri: coerent si non-coerent. Modul coerent ​ e folosit pentru comunicatiile interprocesor. Modul non-coerent e optimizat pentru comunicatiile I/O. 
 +</​note>​
  
-<code sh> +==== Integrarea ​in Arhitectura Hammer ====
-Singularity>​ ./task0  +
-an illegal memory access was encountered ​in task0.cu at line 33 +
-</​code>​+
  
-In acest caz eroare semnalata apare la cudaDeviceSynchronize() desi problema ​este la kernel.+E folosit pentru a lega controller-ul de memorie integrat ​(fostul NorthBridgede memorie. In mod similar ​este folosit in sistemele multiprocesor pentru comunicarea interprocesor,​ folosind modul coerent.
  
-Folosind insa [[http://​developer.download.nvidia.com/​GTC/​PDF/​1062_Satoor.pdf|cuda-gdb]] putem gasi rapid ca problema este la executia de kernelatunci cand se acceseaza zone de memorie nealocate.+AMDOpteron are 3 legaturi HypertransportSeria 100 are 3 legaturi non-coerentedeoarece, fiind destianta monoprocesoareleor,​ nu are nevoie ​de comunicatie interprocesor. Seria 200 are 2 linii non-coerente si una coerenta, pentru unica legatura dintre cele doua procesoare (seria 200 e pentru dual-procesor). Si seria 300 are toate cele 3 legaturi coerente.
  
-<code sh> +{{ :asc:lab4:hypertransport.jpg?640 |Figura7 Hypertransport intern si extern}}
-Singularity>​ cuda-gdb task0 +
-NVIDIA (R) CUDA Debugger +
-9.1 release +
-... +
-(cuda-gdb) run +
-Starting program/​export/​home/​acs/​prof/​stefan_dan.ciocirlan/​skl/​task0/​task0  +
-[Thread debugging using libthread_db enabled] +
-Using host libthread_db library "/​lib64/​libthread_db.so.1"​. +
-warningCannot parse .gnu_debugdata section; LZMA support was disabled at compile time +
-warningCannot parse .gnu_debugdata section; LZMA support was disabled at compile time +
-warning: Cannot parse .gnu_debugdata section; LZMA support was disabled at compile time +
-warning: Cannot parse .gnu_debugdata section; LZMA support was disabled at compile time +
-[New Thread 0x7ffff3f7b700 (LWP 475929)] +
-[New Thread 0x7ffff377a700 (LWP 475930)]+
  
-CUDA Exception: Device Illegal Address +<note tip> 
-The exception was triggered ​in device 0.+Liniile sunt de 16 biti, bidirectionale cu frecvente intre 200 Mhz si 800Mhz, de aici rezultand o viteza de 6.4Gbytes/​sec (3.2Gbytes/​sec  ​in fiecare directie)Cum Opteron are 3 astfel de lagaturi, poate comunica deci 19.2 Bytes/​sec. 
 +</​note>​
  
-Thread 1 "​task0"​ received signal CUDA_EXCEPTION_10,​ Device Illegal Address. +===== Cipuri ce conecteaza prin Hypertransport core-uri AMD =====
-[Switching focus to CUDA kernel 0, grid 1, block (16563,​0,​0),​ thread (0,0,0), device 0, sm 0, warp 6, lane 0] +
-0x0000000000684558 in kernel_compute(int*)<<<​(8388608,​1,​1),​(16,​1,​1)>>>​ () +
-</​code>​+
  
-<note important>​ +  *AMD8151 Hypertransport AGP Tunnel: Controller grafic AGP3.0 . Este practic ce mai ramas din NorthBridge dupa integrarea controller-ului 
-Pentru ​folosi cuda-gdb asupra surselor nu uitați să adăugați ''​-g''​ la compilare în Makefile +  *AMD8131 Hypertransport PCI-X Tunnel: Are rol de buss cu PCI-X 
-</note> ​+  ​*AMD8111 Hypertransport I/O Hub: Are functionalitate standard de SouthBridge,​ incluzand controller PCI, BIOS,​USB,​hard disk, retea si audio.
  
-===== Analiza de performanta in aplicatiile CUDA =====+{{ :​asc:​lab4:​htchips.jpg?​480 | Figura 8. Cipuri ce asigura interconectarea prin Hypertransport a core-ului}}
  
-In aceast sectiune vom explora cateva metode pentru a evalua performantele programelor CUDA.+===== Raspunsul Intel la Hypertransport - QuickPath Interconnect =====
  
-==== Timing via executie kernel (host/CPU) ====+In mod evident Intel nu putea ramane indiferent avantajelor oferite de sistemul de interconectare oferit de catre HyperTransport. Si astfel putem vedea in figura urmatoare cum a aparut QuickPath Interconnect:​
  
-Putem masura timpul de executie al diverselor operatii (executie kernel, transfer date etc), cand acestea sunt blocanteAstfel obtinem timpi de executie al operatiilor,​ asa cum sunt percepute din perspectiva host/CPUAceasta metoda nu este foarte precisa deoarece in timpul de executie sunt incluse si toate operatiile de control CPU<->GPU.+{{ :​asc:​lab4:​fsb_evolution.jpg?640 |Figura 9Evolutia Front-side-Bus-ului in sistemele Intel}}
  
-Mai jos avem un exemplu ​de folosire a functiei cudaDeviceSynchronize pentru a forta o blocare pe partea ​de host/CPU pana cand toate operatiile ​pe partea ​de GPU au fost executate. +Acest sistem ​de interconectare este practic identic functional cu HyperTransport. Intre timp, procesoarele Intel au incorporat si ele controller-ul ​de memorie, si astfel cei doi mari competitori ​pe piata procesoarelor ​de uz general sunt pregatiti in aceeasi masura pentru sisteme multi-core cu multe procesoarememorie multa si aplicatii multi-pthreading pe scara larga:
-<code C> +
-cudaMemcpy(d_xx, N*sizeof(float),​ cudaMemcpyHostToDevice);​ +
-cudaMemcpy(d_y,​ y, N*sizeof(float),​ cudaMemcpyHostToDevice);​+
  
-t1 = myCPUTimer();​ +{{ :​asc:​lab4:​intel-qpi-1.jpg?640 |Figura 10. Arhitectura Intel QuickPath Interconnect }}
-saxpy<<<​(N+255)/​256,​ 256>>>​(N,​ 2.0, d_x, d_y); +
-cudaDeviceSynchronize();​ +
-t2 = myCPUTimer();​+
  
-cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);​ +Sistemul QuickPath ofera o rata de transfer de pana la 25.6GB/s pentru fiecare port (pereche de linii). Legatura poate fi utilizata atat pentru interconectarea controller-ului I/O cu CPU-urilesau a CPU-urilor intre ele. Astfel atat sisteme mono- cat si multi-procesor pot fi realizate cu usurinta cu acest tip de interconectare. Mai multe detalii pot fi vazute in figura urmatoare:
-</​code>​+
  
-==== Timing via CUDA events (device/​GPU) ====+{{ :​asc:​lab4:​intel-qpi-2.jpg?​640 |Figura 11. Exemple de utilizare a QPI pentru sisteme mono si multi-procesor}}
  
-O varianta ​mai buna decat operatiile blocante sunt CUDA events. Acestea au suport hardware la GPU si ofera timpi de executie din perspectiva device/GPU. Mai jos avem un exemplu folosind CUDA events.+In mod similar cu sistemul Hypertransport,​ si interconectarea QuickPath poate fi utilizata pentru a lega sisteme HPC cu mai multe procesoare ​si siteme ​de intrare iesire, dupa cum se poate observa in urmatoarea figura:
  
-<code C> +{{ :​asc:​lab4:​qpi-cpu-mem-io.gif?​480 |Figura 12. Conectarea Utilizand QPI a mai multor procesoare memorii si sisteme I/O}}
-cudaEvent_t start, stop; +
-cudaEventCreate(&​start);​ +
-cudaEventCreate(&​stop);​+
  
-cudaMemcpy(d_x,​ x, N*sizeof(float),​ cudaMemcpyHostToDevice);​ +===== Sisteme Multiprocesor Intel si AMD =====
-cudaMemcpy(d_y,​ y, N*sizeof(float),​ cudaMemcpyHostToDevice);​+
  
-cudaEventRecord(start)+Familia Hammer a fost creata pentru a putea oferi un multiprocesor scalabil, eficient din punctul de vedere al pretului raportat la numarul de procesoare. AMD a mai avut o tentativa in trecut de a crea procesoare pentru sisteme multiprocesor,​ cu AthlonMP. Desi acesta nu a fost o reusita de piata, datorita lui AMD am putut studia problemele aparute in astfel de sisteme. La Athlon MP, memoria ​(care era partajataera botleneck-ul principal. Fiind memorie partajata, toate procesoarele imparteau FSB (Front Side Bus). Cu alte cuvinte viteza cu care procesoarele puteau teoretic accesa memoria era mult mai mare decat viteza cu care putea fi aceasta accesata. Solutia de la Hammer ar fi fost sa ofere fiecarui CPU propria sa conexiune la North Bridgedar acest lucru ar fi fost foarte scump. Solutia relativ ieftina si care nu are nici penalizari de performanta a fost includerea controller-ului de memorie in procesorceea ce s-a si facutAstfelfiecare procesor are propria sa legatura de 128 biti cu memoriaavand pana la 5.3 Gbytes/sec.
-saxpy<<<​(N+255)/256256>>>​(N2.0fd_xd_y); +
-cudaEventRecord(stop);​+
  
-cudaMemcpy(y,​ d_y, N*sizeof(float),​ cudaMemcpyDeviceToHost);​+{{ :​asc:​lab4:​amdvsmp.jpg?​640 |Figura 13. Sistem AMD vs sisteme Muti Procesor Clasice}}
  
-cudaEventSynchronize(stop);​ +In plusdatorita hypertransportfiecare procesor poate accesa memoria celorlate procesoare la viteze de 3.2Gbytes/​sec. Datorita acestui fapt, implementarea unui sistem dual-procesor e la fel de "​usoara"​ ca a unuia cu 8 procesoare, deoarece partile componenete sunt scalabile prin utilizarea Hypertransport. AMD numeste acesta abordare "​glueless multiprocessing",​ deoarece procesoarele sunt legate slab prin Hypertransport. De fapt, e diferenta dintre o cuplare puternica gen circuit-switched versus o cuplare slaba, gen packet-switched,​ cum se intampla in cazul de fata. Figura de mai jos face o comparatie intre arhitecturile de la Intel (Xeonsi AMD (Athlon/Opteron).
-float milliseconds = 0; +
-cudaEventElapsedTime(&​millisecondsstartstop)+
-</code>+
  
 +{{ :​asc:​lab4:​multiprocs.gif?​640 |Figura 14. Comparatie intre sisteme multiprocesor Intel si AMD}}
  
-===== Exercitii =====+Se observa ca cele doua procesoare Athlon (a) impart ​ acelasi controller de memorie. Desi aceasta abordare nu are repercusiuni asupra performantei,​ sistemul nu e scalabil, adica pentru un sistem cu 3 procesoare ar trebui creat un controller separat. Este practic un sistem puternic cuplat (circuit-switched). Cum am mai mentionat, sistemele puternic cuplate sunt greu de scalat. In sistemul (b) cu Intel Xeon, procesoarele impart FSB-ul care, dupa cum am arata mai sus, duce la un botleneck semnificativ. In final, la (c)(d)(e) avem sisteme cu Hammer Opteron. Acestea , fiind slab cuplate prin HyperTransort sunt usor de scalat la 2, 4 sau 8 procesoare. In plus, fiecare are propria sa legatura la memorie, neaparand botleneckuri,​ ca in cazul Intel Xeon.
  
-  - logati-va pe ''​fep.grid.pub.ro''​ folosind contul de pe ''​cs.curs.pub.ro''​ +===== Comparatie intre servere AMD si Intel =====
-  - executati comanda ''​wget https://​ocw.cs.pub.ro/​courses/​_media/​asc/​lab8/​lab8_skl.tar.gz -O lab8_skl.tar.gz''​ +
-  - dezarhivati folosind comanda ''​tar -xzvf lab8_skl.tar.gz''​+
  
-<note tip> +Generatia Phenom de microprocesoare de la AMD este competitorul direct al Intel Xeon si Intel ItaniumCele doua arhitecturi , Intel versus AMD sunt fundamental diferite, dar ofera performante comparabile,​ in functie de domeniul de aplicatie ales pentru comparatie.
-Debug aplicatii CUDA [[https://​docs.nvidia.com/​cuda/​cuda-gdb/​index.html#​introduction|aici]] +
-</​note>​+
  
-Modificarile se vor face in task1.cu si task2.cu - urmariti indicatiile TODO din cod.+O prima mare diferenta intre cele doua arhitecturi este modul in care cele doua abordeaza compatibilitatea cu 32 de bitiAstfel, de la Hammer incoace, AMD a ales sa extinda setul actual de instructiuni x86 pentru 32biti ​cu instructiuni pentru 64, in timp ce Intel a renuntat complet la setul x86, trecand in mod radical la IA64Compatibilitatea la AMD este asigurata automat, noul set de instructiuni fiind doar o extensie e acelui vechi. La Intel, compatibilitatea ​cu 32 biti se face prin emularea vechiului set. Fiind o emulare, exista penalizari de performanta. Pe de alta parte insa, Intel a reusit sa scape in acest fel de complicatii de arhitectura inutile intr-o lume numai de 64 biti. Ambele tipuri mari de arhitecturi de procesoare, Intel si AMD, au trecut in ultimele generatii de procesoare la emularea setului de instructiuni x86 catre micro-operatii specifice fiecarei generatii de procesoare.
  
-**Task 0**  ​Rulați task0 ca exemplu pentru debug (vedeți text laborator)+O alta diferenta este legatura dintre performanta,​ viteza ceasului si gradul de paralelism oferit. Astfel, la inceput Intel s-a concentrat pe marirea frecventei de ceas, si mai putin pe efectuarea de mai multe operatii in paralel. AMD ofera frecvente mai mici, dar a pus mult accent pe paralelism. Astfel performanta e oferita de AMD la frecvente mult mai mici decat Intel. Un avantaj al acestui fapt este ca AMD elimina astfel problemele de disipare a caldurii. Aceste aspecte sunt esentiale mai ales in domeniul serverelor. In ultimii ani Intel merge pe aceeasi cale, si se axeaza in principal pe cresterea numarului de core-uri si nu a frecventei de procesare, cu aceleasi avantaje mentionate anterior.
  
-**Task 1**  - Deschideți fișierul task1.cu și urmăriți instrucțiunile pentru a măsura performanța maximă a unitații GPU, înregistrând numărul de GFLOPS +{{ :​asc:​lab4:​amdvsintel_2procs.jpg?800 |Figura 15. Comparatie intre servere multiprocesor Intel si AMD cu doua procesoare}} 
-    * Masurati timpul petrecut in kernel HintFolositi evenimente CUDA. +{{ :asc:​lab4:​amdvsintel_4procs.jpg?800 |Figura 16Comparatie intre servere multiprocesor Intel si AMD cu patru procesoare}}
-    * Realizati profiling pentru functiile implementate folosind utilitarul nvprof.+
  
 +O alta direrenta ​ e numarul de registre generale; acest aspect are efecte imediate asupra costului de productie. Intel are 128 de registre pentru numere intregi si 128 pentru numere in virgula mobila, in timp ce AMD are numai 16 registre generale. AMD a decis aceasta abordare in urma constatarii ca 80% din cod foloseste maxim 16 registre. Aceasta abordare i-a permis sa reduca semnificativ costurile.
  
-**Task 2**  - Urmăriți TODO-uri din cadrul ​fișierului task2.cu +In sfarsit, o serie de exemple de utilizare ale sistemului de interconectare QuickPath pot fi observate in urmatoarea figuraAici procesoare din generatia Xeon modelele Bloomfield si Nehalem - sunt legate de controller-ele placilor de baza si de sistemele de intrare - iesire prin legaturi Intel QuickPathAcelasi tip de interconectare este utilizat si pentru ​legatura mai multor chipseturi intre ele.
-    * Completati functia matrix_multiply_simple care va realiza inmultirea a 2 matrice primite ca parametru. +
-    * Completati functia matrix_multiply care va realiza o inmultire optimizata a 2 matrice, folosind Blocked Matrix Multiplication. Hint: Se va folosi directiva __shared__ pentru a aloca memorie partajata intre thread-uri. Pentru sincronizarea thread-urilor se foloseste functia __syncthreads. +
-    * Masurati timpul petrecut in kernel Hint: Folositi evenimente CUDA. +
-    * Realizati profiling ​pentru ​functiile implementate folosind utilitarul nvprof.+
  
 +{{ :​asc:​lab4:​qpi-vs-procs.gif?​800 |Figura 17. Exemple de utilizare a QPI impreuna cu arhitecturi de procesoare Intel (Bloomfield / Nehalem)}}
  
 +Dupa cum se poate observa, scalabilitatea sistemelor de calcul de-a lungul vremii este dependenta puternic de solutiile de interconectare a elementelor de procesare. Astfel, de la magistrale simple, la magistrale multiple, la solutii complexe si scalabile cum sunt Hypertransport si QuickPath Interconnect,​ toate s-au dezvoltat in paralel cu modelele de procesoare si chipseturi pe care le interconecteaza.
  
 +===== Exerciții =====
 +
 +Se recomanda utilizarea fep.grid.pub.ro sau a sistemelor din cluster pentru rezolvarea acestui laborator:
 +  - ''​ssh -Y username@fep8.grid.pub.ro''​ (natural puneti utilizatorul vostru in loc de username)
 +  - ''​%%srun --x11 -p nehalem --pty /​bin/​bash%%''​ - va permite conectarea pe coada Nehalem cu 14 servere
 +  - ''​%%apptainer run docker://​gitlab.cs.pub.ro:​5050/​asc/​asc-public/​c-labs:​1.3.1 /​bin/​bash%%''​ - accesăm imaginea de docker în cadrul căreia avem permisiunile necesare realizării laboratorului
 +
 +Alte comenzi utile: ​
 +  - ''​sinfo''​ - va arata cozile existente
 +  - ''​squeue''​ - va arăta informațiile despre joburile în execuție
 +  - ''​sbatch -p nehalem ./​script.sh''​ - va lansa în execuție un task pe una dintre cozi
 +  - ''​%%srun --x11 -p hpsl --pty /​bin/​bash%%''​ - va permite conectarea pe coada HP-SL cu 3 servere
 +  - ''​%%srun --x11 -p dp --pty /​bin/​bash%%''​ - va permite conectarea pe coada DP cu 3 servere
 +  - ''​srun -p nehalem --pty -w nehalem-wn17 /​bin/​bash''​ permite conectarea pe un anumit nod din cadrul unei cozi (17 in acest caz), si nu doar conectarea pe orice nod dintr-o coada
  
 <note important>​ <note important>​
-Recomandăm sa va delogati mereu de pe serverele din cluster dupa terminarea sesiunii, utilizand comanda ''​exit''​+Dacă folosiți WSL2 urmați indicațiile ​de [[:​asc:​res:​wsl-x11]] 
 +</​note>​
  
-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ă. 
  
-Daca nu veti face aceasta delogareveti putea ajunge in situatia in care sa nu va mai puteti loga pe nodurile ​din cluster+**Task 0**  - Alocați un vector de elemente ''​struct particle''​ în următoarele moduri: global, pe stivă și dinamic. 
-</note+    * Porniți de la ''​task01.c''​ (pentru alocare globală), ''​task02.c'' ​ (pentru alocare pe stivă) și, respectiv, ''​task03.c''​ (pentru alocare dinamică). 
 +    * Pentru compilare utilizați fișierul Makefile și comanda ''​make task0<​n>''​. 
 +    * Compilați și rulați executabilele generate pentru un număr ​din ce în ce mai mare de elemente alocate în vector - porniți de la 1,000,000
 +    * Câte elemente pot fi alocate maxim prin fiecare metodă? Utilizati comanda ''​size ​<executabil>''​ pentru a vedea ''​bss''​ unde se afla stocate variabilele globale. 
 +    * Pentru a creste dimensiunea stack-ului utilizat de sistem puteti folosi ''​ulimit -s unlimited''​. Pentru vizualizare si verificare a limitelor sistemului puteti incerca ''​ulimit -a''​. 
 +    * Pentru a face verificarile codului mai rapide, parcurgeti vectorul cu verificarea vitezei din 5M in 5M - nu pentru fiecare valoare in parte. Exercitiul este despre alocare.
  
 +**Task 1**  - Alocați dinamic o matrice de elemente ''​struct particle'',​ in mod liniarizat. Populați aleator matricea astfel încât liniile pare să conțină particule care au componentele vitezei pozitive, iar liniile impare să conțină particule care au componentele vitezei negative. Scalați apoi vitezele tuturor particulelor cu 0.5, ignorând structura matricii, prin folosirea unui cast. Introduceti o verificare pentru alocarea dinamica vs. alocarea "​clasica"​ pe linii. Porniti de la ''​task1.c''​
  
-===== Resurse =====+**Task 2**  - Studiați alinierea variabilelor în C in fisierul ​ ''​task2.c''​ 
 +    * Afișati adresele variabilelor declarate în schelet. Ce observați despre aceste adrese? Ce legatură există între acestea și dimensiunea tipului? 
 +    * Calculați dimensiunea structurilor a si b din scheletul de program. Afișati dimensiunea acestora folosind ''​sizeof''​. Explicați cele observate. 
 +    * Studiați exemplul de aliniere manuală la un multiplu de 16 sau 32 bytes. In cazul in care lucrati pe o arhitectura de 32 de biti, explicati rezultatele observate. 
 +    * Studiati efectul optimizarilor de compilator -O2 sau -O3 asupra alinierii datelor. ​
  
-{{:asc:lab8:lab8_skl.tar.gz|Schelet Laborator 8}}+**Task 3**  - Determinați dimensiunea cache-urilor L1 și L2 folosindu-vă de metoda prezentată în curs ({{asc:lab4:asc_-_03_-_ierarhia_de_memorie_si_optimizarea_pentru_memorie.pdf slide-urile 67-77}}). Utilizati variabile char/int8_t pentru acest task. 
 +    * Folosiți makefile-ul pus la dispoziție pentru a genera graficele. 
 +    * Porniti de la ''​task3.c''​
  
-/* {{:​asc:​lab8:​sol:​lab8_sol.tar.gz|Solutie Laborator 8}} */ 
  
-{{:asc:lab8:asc_lab8.pdf|Enunt Laborator ​8}}+**Task 4**  - Determinați dimensiunea unei linii de cache folosindu-vă de metoda prezentată în curs ({{asc:lab4:asc_-_03_-_ierarhia_de_memorie_si_optimizarea_pentru_memorie.pdf | slide-urile 67-77}}). 
 +    * Folosiți makefile-ul pus la dispoziție pentru a genera graficele. Utilizati variabile char/int8_t pentru acest task. 
 +    * Porniti de la ''​task4.c''​ 
 +    * Generați grafice pentru mai multe dimensiuni ale vectorului parcurs astfel încât să depășească mărimea cache-ului L1, L2, respectiv, L3. 
 +    * Mecanismele hardware avansate implementate în arhitecturile de procesoare actuale generează comportamente complexe care nu corespund nepărat modelului simplu prezentat în curs. Aceste mecanisme pot chiar masca dimensiunea reală a linei de cache, fiind astfel necesară testarea cu diferite valori ale vectorului parcurs pentru a putea trage o concluzie informată. 
 +    * Studiati si explicati de ce codurile (identice) si makefile-urile (diferite) de la taskurile 3 si 4 duc la graficele obtinute in cadrul laboratorului.  
 +    * Combinati intr-un singur grafic "​relevant"​ rezultatele prezentate in taskurile 3 si 4. 
 + 
 + 
 +/* 
 +  - Studiați și documentați cum poate fi utilizată dimensiunea determinată mai sus în contextul rezolvării unor probleme simple de genul: înmulțirea unui vector cu un scalar sau înmulțirea a doi vectori (fragment de cod relevant pe slide-ul 77). Dimensiunea acestor vectori va varia între limitele următoare:​ 
 +      * Mai mici decât cache-ul sistemului (trebuie determinat de dumneavoastră din specificațiile sistemului). 
 +      * Mai mari decât cache-ul și mai mici decât memoria principala a sistemului de calcul din laborator (aceasta din urmă o determinați dumneavoastră din specificațiile sistemului). 
 +      * Mai mare decât memoria principala a sistemului. 
 +*/ 
 + 
 +==== Tips ==== 
 +Pentru verificarea rezultatelor obtinute pentru dimensiunea cache-ului puteti folosi si rezultatele obtinute cu: 
 +  * <code bash> 
 +[student@localhost ~]$ getconf -a 
 +[student@localhost ~]$ cat /​sys/​devices/​system/​cpu/​cpu1/​cache/​index0/​coherency_line_size 
 +[student@localhost ~]$ cat /​proc/​cpuinfo 
 +</​code>​ 
 +  * <code bash> 
 +[student@localhost ~]$ sudo dmidecode | less  
 +</​code>​ 
 + 
 +Cautati dupa "​Cache"​ si veti afla informatiile despre memoria cache instalata 
 + 
 +  * <code bash> 
 +[student@localhost ~]$ valgrind --tool=cachegrind ./​program_test 
 +</​code>​ 
 + 
 +Acesta va simula rularea programului si va afisa informatii legate de accesul la L1, L2 cache si miss rate. 
 + 
 +Un tutorial interesant si util pentru utilizarea gdb poate fi gasit in pagina echipei de SO, aici: http://​ocw.cs.pub.ro/​courses/​so/​laboratoare/​resurse/​gdb 
 + 
 +===== Resurse ===== 
 +  * Responsabilii acestui laborator: [[emil.slusanschi@cs.pub.ro|Emil Slușanschi,​ Cosmin Samoila]] 
 +  * <​html><​a class="​media mediafile mf_pdf"​ href=":​asc:​lab4:​index?​do=export_pdf">​PDF laborator</​a></​html>​ 
 +  * {{asc:​lab4:​lab4_skl.tar.gz|Schelet laborator}} 
 +  <​hidden> ​ * {{:​asc:​lab4:​sol:​lab4_sol.tar.gz|Soluție ​Laborator ​4}}   </​hidden>​
  
-  * Responsabili laborator: Andreea Birhala, Roxana Balasoiu, Ovidiu Dancila, Mihai Volmer, Grigore Lupescu 
  
-==== Referinte ​====+==== Referințe ​====
  
-  ​* Documentatie CUDA: +  * [[http://en.wikipedia.org/wiki/Data_structure_alignment#​Typical_alignment_of_C_structs_on_x86 ​Alinierea structurilor in pe x86]] 
-    ​* [[https://docs.nvidia.com/pdf/CUDA_C_Programming_Guide.pdf|CUDA Programming]] +  {{:asc:​lab4:​amd_cpu_roadmap.pdf|}} 
-    [[https://docs.nvidia.com/​cuda/​pdf/​CUDA_Compiler_Driver_NVCC.pdf| CUDA NVCC compiler]] +  * {{:​asc:​lab4:​amd-quad-core_opteron_2proc_server-ws_comparison.pdf|}} 
-    [[https://​docs.nvidia.com/​cuda/​profiler-users-guide/​index.htmlCUDA Visual Profiler]] +  {{:asc:lab4:amd-4proc_server_comparison.pdf|}} 
-    [[https://developer.download.nvidia.com/​compute/​cuda/​9.1/​Prod/​docs/​sidebar/​CUDA_Toolkit_Release_Notes.pdf|CUDA 9.1 Toolkit]] +  {{:asc:lab4:amd_designing_for_n_cores.pdf|}} 
-    [[https://developer.nvidia.com/​cuda-gpus|CUDA GPUs]] +  * {{:​asc:​lab4:​introduction_to_amd64_2005.pdf|}} 
-  * Acceleratoare hpsl (hpsl-wn01, hpsl-wn02, hpsl-wn03) +  {{:asc:lab4:​coleamd.pdf|}} 
-    [[http://​international.download.nvidia.com/tesla/pdf/​tesla-k40-passive-board-spec.pdf|NVIDIA Tesla K40M]] +  * {{:​asc:​lab4:​multi-core_codingphenom.pdf|}} 
-    [[https://en.wikipedia.org/​wiki/​Nvidia_Tesla|NVIDIA Tesla]] +  {{:asc:​lab4:​opteronoverview.pdf|}} 
-  * Acceleratoare dp (dp-wn01, dp-wn02, dp-wn03) +  * {{:​asc:​lab4:​318082.pdf|}} 
-    [[https://​www.nvidia.com/​docs/​IO/​43395/​NV_DS_Tesla_C2050_C2070_jul10_lores.pdf|NVIDIA Tesla C2070]] +  * {{:​asc:​lab4:​opteron_mpf.pdf|}} 
-    [[http://​www.nvidia.com/​docs/​io/​43395/​nv_ds_tesla_c2050_c2070_apr10_final_lores.pdf|NVIDIA Tesla 2050/2070]] +  * {{:​asc:​lab4:​opteron_data_sheet.pdf|}} 
-    [[https://​cseweb.ucsd.edu/​classes/​fa12/​cse141/​pdf/​09/​GPU_Gahagan_FA12.pdf|NVIDIA CUDA Fermi/​Tesla]] +  {{:asc:lab4:xeon-5600-brief.pdf|}} 
-  ​* Advanced CUDA +  * {{:​asc:​lab4:​xeon-5600-vol-1-datasheet.pdf|}} 
-    ​* [[https://devblogs.nvidia.com/gpu-pro-tip-cuda-7-streams-simplify-concurrency/|CUDA Streams]] +  {{:asc:​lab4:​xeon-5600-vol-2-datasheet.pdf|}} 
-    * [[https://devblogs.nvidia.com/introduction-cuda-dynamic-parallelism/|CUDA Dynamic Parallelism]]+  {{:asc:lab4:​nv_ds_tesla_c2050_c2070.pdf|}} 
 +  {{:asc:lab4:​tesla-kseries-overview-lr.pdf|}} 
 +  * [[http://blog.regehr.org/archives/898Sequential Consistency vs TSO]]   
 +  * [[http://www.catb.org/esr/​structure-packing/ | The Lost Art of C Structure Packing]]
  
asc/laboratoare/08.1648668420.txt.gz · Last modified: 2022/03/30 22:27 by stefan_dan.ciocirlan
CC Attribution-Share Alike 3.0 Unported
www.chimeric.de Valid CSS Driven by DokuWiki do yourself a favour and use a real browser - get firefox!! Recent changes RSS feed Valid XHTML 1.0