Architetture Avanzate, Appunti di Elementi di Informatica

Scarica Architetture Avanzate e più Appunti in PDF di Elementi di Informatica solo su Docsity!

Architetture Avanzate 2016

Fuser Alessandro

Sistemi Embedded

Università degli Studi di

Verona

Prof. Nicola Bombieri

Indice

• 1 Introduzione
  • 1.1 Tassonomia di Flynn
  • 1.2 Definizioni
  • 1.3 Shared Memory vs Distributed Memory
  • 1.4 Modi di Programmare parallelo
• 2 Pipeline
  • 2.1 MIPS Pipeline
  • 2.2 Pipeline Hazards
  • 2.3 Operazioni Multi Ciclo
• 3 Instruction Level Parallelism
  • 3.1 Tecniche basi di compilazione
  • 3.2 Predizione su branch
  • 3.3 Dynamic Scheduling
  • 3.4 Hardware-based Speculation
• 4 GP-GPUs
  • 4.1 CUDA C
  • 4.2 CUDA Threads
  • 4.3 GPU Merory Model
  • 4.4 Performance Measurement
  • 4.5 Data transfer and streams
• 5 OpenCL
• 6 OpenACC
• 7 Memory Hierarchy
  • 7.1 Block Placement
  • 7.2 Block Search
  • 7.3 Block Replacement
  • 7.4 Write Strategy
  • 7.5 Cache Performance
  • 7.6 Memoria Virtuale

• Task: set di istruzioni eseguite da un processore;
• Pipelining: dividere un task in step eseguiti da tipi diversi di elementi;
• Shared Memory: tutti i processori hanno accesso diretto ad una memo- ria fisica comune oppure possono indirizzarsi alla stessa locazione virtuale di memoria;
• Distributed Memory: accesso alla memoria fisica non comune ma co- municazione tra macchine tramite network;
• Granularity: presente nei computer paralleli, è il rapporto tra i tempi di computazione e di comunicazione (Coarse se intervallo grande, Fine altrimenti);
• SpeedUp: rapporto tra tempi di esecuzione seriale ed esecuzione paral- lela;
• Parallel Overhead: tempo richiesto per coordinare i task paralleli; com- prende il tempo di avvio, di sincronizzazione, di comunicazione, di software eseguito e di terminazione;
• Scalability: abilità di un sistema parallelo di avere un incremento di speedup proporzionale all’aggiunta di processori;

1.3 Shared Memory vs Distributed Memory

La Shared Memory può essere di due tipi diversi: UMA se ho processori identici e lo stesso tempo di accesso alla memoria; NUMA se ho la connessio- ne di 2+ Symmetric Multi Processors e il tempo di accesso non è lo stesso. I vantaggi della Shared Memory sono una facile programmazione alla memoria ed una condivisione dati tra task veloce ed uniforme, ma presenta anche svataggi tra i quali assenza di scalabilità, costo elevato ed il programmatore è responsa- bile della sincronizzazione. Nella Distributed Memory invece è richiesta una connessione di rete per connettere i vari sistemi processore-memorie, ma ogni processore ha la propri memoria ed opera indipendentemente, per cui ho vantaggi in termini di scala- bilità, interferenze e problemi di coerenza della cahe, ma con una più difficile mappatura dei dati e un modello NUMA.

1.4 Modi di Programmare parallelo

Esistono 3 modi principali per trasformare un lavoro sequenziale in un lavoro parallelo:

• Threads Model: Il programma viene schedulato ed acquisisce le risorse necessarie, poi crea un certo numero di thread. Ogni thread possiede dati locali ma condivide le sue risorse col programma che l’ha creato, comunicando con gli altri thread attraverso la memoria globale. Mentre il programma deve restare fino a che non termina, i thread possono essere creati e rimossi in qualsiasi momento;

• Message Passing Model: ho un insieme di task che usano la memoria locale per la computazione e lo scambio di dati avviene tramite invio e ricezione di messaggi, per cui il trasferimento dati richiede operazioni cooperative;
• Data Parallel Model: il lavoro viene focalizzato sull’eseguire operazioni di un data set organizzato in strutture comuni, dove ogni task lavora su parti differenti della stessa struttura dati, eseguendo la stessa operazione nella loro partizione;

2.2 Pipeline Hazards

Gli Hazards sono situazioni che prevengono l’esecuzione della prossima istru- zione nello stream di esecuzione nel designato ciclo di clock e possono essere di 3 tipi:

• Structural: due istruzioni differenti usano la stessa risorsa nello stesso ciclo. La soluzione può essere lo stallo oppure il replicare una risorsa. Tipi hazard di questo tipo sono sulla memoria e sui floating point;
• Data: due instruzioni differenti devono accedere alla stessa locazione di un dato. Di 3 tipi: RAW, l’istruzione successiva legge un dato che deve ancora finire di essere computato; WAR, un dato viene scritto da un’istruzione successiva ma deve an- cora essere letto da quella precedente; WAW, un dato viene scritto da un’istruzione successiva ma poi viene sovrascritto da un’istruzione precedente; La soluzione può essere nuovamente lo stallo oppure si può utilizzare il Forwarding, ossia rendere disponibile un dato all’ALU di un’istruzione successiva che lo necessita tramite collegamento logico. Un’altra soluzione software può essere la schedulazione dell’istruzione, cosi da poter scambia- re tra loro istruzioni nel flusso di esecuzione per rendere minimo il numero di hazard.
• Control: esecuzione di branch o istruzioni di salto.SI può risolvere at- traverso l’anticipazione del calcolo della condizione del branch nella fase Decode oppure aggiungere un processore di calcolo. Un’altra soluzione è l’eliminazione del branch in 3 possibili modi: Branch Not Taken(33%), dove prevedo che la condizione non si veri- fica e quindi eseguo le istruzioni successive. Se viene invece preso, sosti- tuisco le operazioni correnti con NOP e metto uno stallo, e poi eseguo le operazioni dentro il branch; Branch Taken,(66%) dove prevedo stavolta che la condizione si veri- fiche, se non viene preso fare come prima; Delayes Branch, dove eseguo le istruzioni nel BDS (numero di cicli richiesto per risolvere branch) fregandomene del risultato della condizione, prendendole da prima il branch, da dopo il branch oppure dal target del branch.

Lo SpeedUp della Pipeline è ora P ipelineDeath/(1+BranchF requency∗BranchP enalty). Ma allora cosa rende così difficile l’implementazione della Pipeline? Le inter- ruzioni, poichè un’istruzione completa il N cicli di clock ma questo è il minimo intervallo per lanciare un’interruzione, per cui bisogna avere la capacità di can- cellare un’istruzione mentre sta eseguendo. Queste interruzioni possono essere in più stadi della pipeline (IF, ID, EX, MEM), perciò l pipeline deve poter essere spenta in modo sicuro ed il PC deve essere salvato per ripartire da quel punto. Perciò quando si presenta un’eccezione:

• Forzo la trap nella pipeline come prossima IF;

• Flasho tutte e istruzioni che seguono l’eccezione (evita cambiamento sta- to);
• Completa le istruzioni precedenti se possibile;
• Salva il valore del PC
• Gestisce istruzione

Un’eccezione viene chiamata precise se lascia la macchina in uno stato consi- stente con il modello sequenziale di esecuzione; per essere tale tutte le istruzioni prima della fault sono completate, tutte quelle dopo non sono completate e la fault esegue correttamente.

2.3 Operazioni Multi Ciclo

Poichè è impossibile pretendere che tutte le operazioni completino in un ciclo di clock, si usa una pipeline modificata dove il ciclo EX viene ripetuto tante volte quante ne sono necessarie per completare l’operazione. Le unità funzionali separate sono quattro: integer unit, FP (Floating Point) e Integer multiplier, FP adder, FP e Integer divider. I possibili problemi sono che operazioni senza pipeline aumentano la possibi- lità di hazard strutturali, differenti latenze possono richiedere registri multi- pli su cui scrivere, le istruzioni finiscano in ordine diverso da quello previsto.

3.2 Predizione su branch

La predizione statica sul branch è fatta a tempo di compilazione, ma dicendo che un branch verrà preso abbiamo uno sbaglio che varia dal 59% al 9% (troppo). Quello dinamico (decisioni cambiano durante il tempo) viene fatto dall’hardware e si basa sul passato: sia F una funzione che esprime il risultato di una predizione su branch, siano x1..xn parametri che influenzano F, allora F(x1..xn) è preso se > 0.5. Più semplicemente, posso usare una tabella che salva la storia dei branch: durante IF accedo alla BHT(Branch History Table) per predire il branch e durante ID verifico se è un branch (implementata come una cache), di solito usa 2-bit per la predizione

3.3 Dynamic Scheduling

L’HW riarrangia l’esecuzione delle istruzioni per ridurre gli stalli mantenendo il comportamento delle eccezioni e del flusso di dati. Permette alla CPU di tollerare ritardi impredicibili come mancanza di cache e di compilare codice che era pensato con una sola pipeline di funzionare efficaciemente in pipeline diffe- renti, ma non si può cambiare l’ordine delle istruzioni ordinate e non si sfrutta lo scheduling run-time -> dynamic scheduling! Nelle pipeline tradizionali, un hazard manda in stallo istruzioni che non sono affatte da tale hazard, per cui si cerca di completare questa operazione; facendo cosi però c’è il rischio di introdurre hazard di tipo WAR e WAW e problemi di imprecisioni degli errori. Distinguiamo l’inizio dell’esecuzione di un’istruzione dalla sua fine, per cui tra questi due tempi essa è in esecuzione. Assumiamo inoltre che la pipeline permet- te l’esecuzione multipla di più istruzioni in contemporanea e che il processore ha più unità funzionali -> dividiamo la fase ID in 2 stadi: Issue che decodifica istruzioni e controlla per hazard strutturali e Read Operands, che aspetta fino a che non ho più data hazard. poi legge gli operandi; uno stage IF precede lo stage Issue e può portare entrambi nell’IR oppure in una coda di istruzioni pendenti. Viene usato l’approccio di Tomasulo: evita WAR e WAW rinominando i regi- stri ed evita RAW eseguendo istruzioni solo quando gli operandi sono disponibili. Introduce le stazioni di prenotazione,che mandano gli operandi delle istru- zioni ad operare nell’Issue, eliminando il bisogno di prendere gli operandi dai registri (posso avere più stazioni di prenotazioni che registri reali). Ognuna ha 7 campi: Op (operazione da eseguire), Qj Qk (stazioni che producono l’ope- rando corrispondente), Vj Vk (valori operandi), A (info per calcolo memoria di store/load), Busy, Qi (numero di stazioni che contengono l’operazione il cui risultato deve essere salvato). Compie 3 passi base:

• Issue -> prende la prossima istruzione dalla testa della coda, se c’è una stazione vuota corrispondente manda l’istruzione alla stazione, se non c’è lancia uno structural hazard e tiene l’istruzione
• Execute -> se qualche operando no è disponibile monitora il CDB (Com- mon Data Bus, dove sono salvati i risultati da FP o LS) e quando divnta disponibile viene piazzato nella stazione che lo sta aspettando, quando

tutti gli operandi sono disponibili l’istruzione viene eseguita dalla corri- spondente unità funzionale

• Write -> quando il risultato è disponibile lo scrive nel CDB e poi dal CDB ai registri e alle stazioni

I vantaggi quindi di questo approccio sono che se un singolo risultato atteso da più istruzioni questo viene propagato e l’esecuzione riprende subito e che WAR e WAW sono risolvi semplicemente rinominando i registri

3.4 Hardware-based Speculation

Sfruttare più parallelismo richiede che superiamo il limite delle dipendenze di controllo. Un’idea è quella di permettere alle istruzioni di eseguire non in or- dine ma forzarle a registrarsi in ordine e di prevenire ogni azione irrimediabile, separando il processo di esecuzione completata da quello di istruzione registrata usando un buffer di riordine (ROB): questo tiene i risultati di un’istruzione tra i due tempi.

l’host e chiamabile solo dall’host). Per compilare un programma CUDA si prende il codice scritto in C con le esten- sioni del CUDA, si usa il compilatore NVCC che compila il codice diversamente a seconda che sia Host o Device e poi il tutto viene rimesso assieme.

4.2 CUDA Threads

L’hardware è libero di assegnare i blocchi ad ogni processore quando vuole, men- tre un kernel scala attraverso qualsiasi numero di processori paralleli. Ogni bloc- co può eseguire in qualsiasi ordine relativamente agli altri blocchi. Ogni blocco viene eseguiro come un Warp a 32 thread. I thread sono assegnati a Streaming Multiprocessors (MP) in blocchi granulari e possono essere eseguiti concor- rentemente. I blocchi di thread sono partizionati in warp (che mantengono lo ID consistenti) e questo partizionamento è lo stesso sempre. Thread all’interno di un singolo warp possono prendere strade differenti e queste diverse esecuzioni sono serializzate nelle GPU correnti, quindi è da evitare -> if (threadIdx.x / WARPSIZE > 2) e non if (threadIdx.x > 2).

4.3 GPU Merory Model

Ogni thread può leggere/scrivere: registri per-thread (1 ciclo), shared memory per-blocco (5 cicli), global memory per-grid (500 cicli). Le variabili senza qua- lificatore (shared o constant) risiedono in un registro, eccetto che per array di thread che stanno nella memoria globale. Le variabili vanno dichiarate in base alla possibilità dell’host di accedervi, per cui è globale se può ed è all’esterno di ogni funzione, è shared/local/automatica se non può ed è all’interno del kernel. Un modo vantaggioso di migliorare le performance è di partizionare dati in sot- toinsiemi che possono stare nella shared memory e gestire ognuno di questi con un blocco di thread caricandolo dalla global memory sulla shared memory e poi ripassarlo (buono quando i thread hanno lo stesso tempo di accesso). Per esempio con una lunghezza di 16, posso avere 1616 = 256 thread, che vuol dire 2256=512 load da global memory per 256216=8192 operazioni +/; quindi nella shared memory da 16KB, ogni blocco usa 2256*4B=2KB di shared me- mory, ciò vuole 8 potenziali blocchi attivi contemporaneamente. Per la sincronizzazione si usa __syncthreads(), che è una sorta di barriera dove tutti i blocchi dello stesso thread devono arrivare prima di poter prosegui- re. Si può interrogare il device chiedendone le proprietà tramite: cudaGetDevice- Count(&devcount) per il numero di device nel sistema, cudaGetDevicePro(&devprop), devprop.maxThreadsPerBlock, devprop.sharedMemoryPerBlock. La struttura tipica di un programma CuDA è dunque questa:

4.4 Performance Measurement

Per l’utente ha importanza il tempo di esecuzione (latenza), mentre per un ge- store interessa principalmente il throughput, ossia l’ammontare complessivo di lavoro svolto in un dato tempo (bandwith). Il tempo di esecuzione è il reciproco della prestazione, quindi T empoesecuzioney T empoesecuzionex =^

P restazionex P restazioney =^ 1 +^

n 100 Si deve favorire il caso più frequente a discapito del più raro e, per la legge di Amdahl, il miglioramento di prestazione che può essere ottenuto usando alcune modalità di esecuzione più veloci è limitato dalla frazione di tempo nella quale tali modalità possono venire impiegate. Lo speedup fornisce informazioni su quanto più velocemente un lavoro verrà eseguito usando la macchina con la mi- glioria rispetto alla macchina originale: speedup = tempoesecuzionenomiglioriatempoesecuzionesimigliora. Il tempo è la misura delle prestazioni di un computer, dove il tempo di risposta rappresenta la latenza per il completamento di un lavoro includendo accessi al disco, memoria e I/O, mentre il tempo di CPU corrisponde a cicli CPU per un programma * tempo CPU sistema, il CPI è il numero di cicli di clock diviso il numero di istruzioni. Il tempo di CPU dipende dal ciclo di clock (tecnologia HW ed organizzazione), dal ciclo di clock per istruzione (organizzazione ed ar- chitettura set istruzioni) e dal numero di istruzioni (architettura set istruzioni e tecnologia compilatori).

4.5 Data transfer and streams

cudaMemcpy serializza il traferimento dei dati e quindi anche la computazio- ne della GPU poichè usa una sola direzione di svolgimento. Ma alcuni device CUDA supportano il device overlap, ossia l’esecuzione simultanea di kernel mentre è in corso una processo di copiatura tra device e host. CUDA supporta esecuzione parallela di kernel e cudaMemcpy con Streams, ossia code di operazioni (lanco di kernel e cudaMemcpy), dove operazioni (task)

Capitolo 5

OpenCL

OpenCl è un acronimo per "Open Computin Language", che sta ad indicare la sua natura aperta e senza royalty, permette la programmazione parallela tra più piattaforme e può usare tutte le risorse di computazione (CPU, GPU, .. ). Consiste di 3 componenti principali:

• Platform Model: CPU, GPU e acceleratori come devices, dove ogni device contiene una o più compute units, le quali contengono una o più SIMD processing elements; un host è connesso ad uno o più devices (men- tre in CUDA no); i devices sono contenuti nei context ed ognuno necessita la propria coda di lavoro, cosi come i trasferimenti sono associati con una coda di comandi;
• Execution Model: un programma viene diviso in kernels, che eseguono uno o più devices, e host, che definisce il contesto per il kernel e gestisce la loro esecuzione; ogni coda li lavoro può eseguire in ordine o meno, ma comunque ci deve essere una sincronizzazione esplicita con il ritorno di un evento;
• Memory Model: private memory (registers), local memory, global e host memory; i dati vengono trasferiti dall’host->global->local e indietro; le variabili possono essere global, private, local o constant;

Rispetto a CUDA non supporta certe funzioni hardware specifiche a favore della portabilità, permette la compilazione run-time ma non la texture memory. Di simile hanno i modelli e la sintassi.

Capitolo 7

Memory Hierarchy

La cache è il nome dato al primo livello di gerarchia della memoria incontrato quando l’indirizzo lascia il processore. Posso avere cache hit e cache miss; il tempo richiesto per il cache miss dipende dalla latenza (tempo di trovare la prima parola) e dalla bandwidth (tempo di trovare il blocco) della memoria.

7.1 Block Placement

Posso avere 3 modalità di piazzamento del blocco nella cache:

• Direct Mapping: ogni blocco può essere mappato solo in un posto secondo la formula (block address) MOD (n. of block in cache)
• Fully Associative: un blocco può essere piazzato ovunque nella cache
• Set Associative: un blocco può essere piazzato in un set ristretto di posti nella cache secondo la mappatura (block address) MOD (n. of sets in cache)

7.2 Block Search

In generale, un indirizzo è partizionato in t bits per il TAG, s bits per il SET INDEX e b bits per l’OFFSET, che vuol dire che la cache consiste in 2 s^ classi equivalenti, ogni set consiste in t entries e che ogni entry consiste in 2 b^ bytes. Il dato comunque viene validato attraverso un valid bit finale.

7.3 Block Replacement

Con la Direct Mapping, viene verificato solo un blocco per hit e quindi posso sostituire solo quello; con le altre posso usare politiche Random, LRU o FIFO.

7.4 Write Strategy

Due possibili scelte di scrittura: write through con la quale l’informazione viene scritta sia in cache che in memoria per avere consistenza, oppure write back con la quale l’informazione viene scritta solo nella cache e modifico il blocco scritto in memoria solo quando nella cache viene rimpiazzato. Con un write miss, posso decidere di usare due strategie: write allocate per cui write miss corrisponde a read miss, oppure no-write allocate che non carica in cache ma modifica il blocco nel livello più basso

7.5 Cache Performance

Il tempo di accesso alla memoria medio AMAT = Hit time + (Miss Rate x Miss Penalty), dove Hit Time è il tempo di accedere alla cache e trovare il dato, la Miss Rate è la possibilità di fare miss e il Miss Penalty è quanto mi costa il Miss. Per migliorare le performance si possono seguire 3 strade: ridurre il miss rate, ridurre il miss penalty oppure ridurre il time hit. I miss che si possono avere sono: Compulsory (all’inizio quando ancora non ho niente nella cache), Capacity (oltre tot informazioni non posso salvarle), Conflict (un blocco può essere rimosso e poi rimesso).

• Dimensioni blocchi maggiori: riduce miss rate ma potrebbe aumentare i miss di conflitto e di capacità se la cache è piccola, porta vantaggi di locazione dello spazio, potrebbe aumentare la AMAT
• Cache maggiore: riduce miss rate ma ho un tempo di hit più lungo con costi maggiori
• Associatività maggiore: riduce miss di conflitto, può incrementare il tempo di hit e AMAT non lineare nel numero di way
• Cache Multilivello: riduce i miss di penalità con primo livello veloce e piccolo, il secondo grande; il secondo livello può includere quello che c’è nel primo livello (inclusion) oppure non aver nessun dato in comune (exclusion)
• Priorità alla lettura: riduce miss di penalità con una write-through cache
• Evitare traduzione indirizzi: indirizzo la cache con VA, cosi ho tradu- zioni solo per i miss; ho problemi però di protezione (risolvo aggiungendo un campo di verifica), aliasing (risolvo usando gli ultimi n bit (page co- loring)) e flusso di cache (risolvo aggiungendo un tag identificatore del processo)

7.6 Memoria Virtuale

La memoria virtuale è quella vista dal programma, mentre quella fisica è quella vista dalla memoria, con necessità di traduzione da una all’altra fatta per ogni accesso in memoria tramite paging o segmentatio.