Tesi di Laurea Triennale - ingegneria-informatica.unina.it · Due sono le scuole di pensiero: la...

Universita degli Studi di Federico II

FACOLTA DI INGEGNERIA

Corso di Laurea Triennale in Ingegneria Informatica

Tesi di Laurea Triennale

Acceleratori hardware:architetture hardware delle odierne GPU e relative

ottimizzazioni software

Candidato:

Maria Luisa FarinaMatricola N46001592

Relatore:

Alessandro Cilardo

Anno Accademico 2017-2018

Indice

1 Introduzione 21.1 Cos'è una GPU . . . . . . . . . . . . . . . . . . . . . . . . . . 21.2 GPGPU . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 31.3 Acceleratori Hardware . . . . . . . . . . . . . . . . . . . . . . 41.4 Intelligenza Arti�ciale . . . . . . . . . . . . . . . . . . . . . . 5

2 NVIDIA 62.1 Architetture . . . . . . . . . . . . . . . . . . . . . . . . . . . . 6

2.1.1 Tesla . . . . . . . . . . . . . . . . . . . . . . . . . . . . 62.1.2 Fermi . . . . . . . . . . . . . . . . . . . . . . . . . . . 72.1.3 Kepler . . . . . . . . . . . . . . . . . . . . . . . . . . . 82.1.4 Maxwell . . . . . . . . . . . . . . . . . . . . . . . . . . 92.1.5 Pascal . . . . . . . . . . . . . . . . . . . . . . . . . . . 102.1.6 Volta . . . . . . . . . . . . . . . . . . . . . . . . . . . . 122.1.7 Turing . . . . . . . . . . . . . . . . . . . . . . . . . . . 14

2.2 CUDA . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 152.2.1 Compilazione ed esecuzione . . . . . . . . . . . . . . . 152.2.2 Funzionamento . . . . . . . . . . . . . . . . . . . . . . 162.2.3 Modello di Memoria . . . . . . . . . . . . . . . . . . . 172.2.4 Esempio di Programmazione CUDA . . . . . . . . . . . 19

3 OPENCL 243.1 Cos'è e come funziona . . . . . . . . . . . . . . . . . . . . . . 24

3.1.1 Modello di Memoria . . . . . . . . . . . . . . . . . . . 253.1.2 Portabilità . . . . . . . . . . . . . . . . . . . . . . . . . 27

4 Conclusioni 284.1 Non solo NVIDIA . . . . . . . . . . . . . . . . . . . . . . . . . 28

1

Capitolo 1

Introduzione

1.1 Cos'è una GPU

GPU è l' acronimo di Graphic Processing Unit. Come suggerisce ilnome una GPU non è altro che un'unità di elaborazione in grado di pro-cessare (per quanto riguarda la gra�ca tale elaborazione prende il nome direndering1) immagini gra�che lavorando in sinergia con la CPU (CentralProcessing Unit, conosciuto semplicemente come processore), dando luogoad un esempio di architetture di calcolo eterogenee.Basandosi sulla GPU è nato il cosiddetto GPGPU( General-purposecomputing on graphics processing units). Di questa de�nizione fannoparte tutti quegli scenari in cui la sola CPU non riesce a garantire risultatiottimali in quanto deve confrontarsi con un'enorme mole di dati da eleborare,nella maggioranza dei casi in maniera parallela per questioni di e�cacia masopratutto di e�cienza. Ed è per questo che l'architettura intrinsecamenteparallela delle GPU si presta molto bene a questi scopi.Generalmente la parte sequenziale della applicazione è gestita dalla CPUmentre la porzione di codice che richiede una maggiore potenza compu-tazionale è svolta dalla GPU. Su questo tipo di soluzione si basa CUDA(Compute Uni�ed Device Architecture) proprietà di casa NVIDIA,una delle leader delle settore delle schede gra�che.

1Rendering: ricostruzione di un'immagine gra�ca a partire dalla sua rappresentazionematematica

2

CAPITOLO 1. INTRODUZIONE 3

Figura 1.1: CUDA

1.2 GPGPU

Utilizzare una GPU per elaborazioni poi denominate proprio GPGPU per ilruolo centrale e cruciale dell'unità di elaborazione gra�ca, è stato il modopiù vantaggioso per a�rontare quelle situazioni in cui è richiesta un'elevatapotenza di calcolo per determinate applicazioni. Si pensi, ad esempio, aprogrammi di risoluzione di funzioni matematiche, oppure in ambito medico,a macchinari che consentono di sfruttare modelli ed immagini sempre piùso�sticati di una particolare situazione clinica, oppure a tutti quei programmidi simulazione utilizzati per la ricerca in ambiente chimico-�sico.I principali vantaggi risultano:

• Vantaggi prestazionali: rispetto ad una CPU i tempi di elaborazionepossono essere abbattuti del 100% previa riscrittura del codice volta aparallelizzarne la struttura, laddove possibile, per renderlo più fruibileall'architettura della GPU.

• Costi e tecnologie: se dal lato economico abbiamo un'equivalenzatra il costo di una CPU e di una GPU della stessa fascia, dal latodel progresso tecnologico abbiamo invece un'enorme disparità sia perquanto riguarda la frequenza di rilascio di aggiornamenti architetturali,


sia per quanto riguarda l'aumento delle prestazioni, che nel caso delleGPU arriva a raddoppiare le performance della generazione precedente.

• Consumo: nonostante le GPU consumino molto in termini energetici èanche vero che producono risultati eccellenti in frazioni di tempo moltoridotte. Di conseguenza il trade-o� tra consumo/prestazioni rappresen-ta, ad oggi, un aspetto da tenere in considerazione al �ne di raggiungereprestazioni ancora più vantaggiose.

1.3 Acceleratori Hardware

La GPU è quindi un valido esempio di acceleratore hardware. Con questotermine s'intende una componente hardware che riesce a migliorare le pre-stazioni generale del calcolatore andando a potenziare speci�che operazioni(ad esempio calcoli in virgola mobile) del PC. Un ampio uso delle GPU vienefatto per migliorare l'esperienza di gioco nei videogames, che ad oggi ormaisono completamente basati su gra�ca 3D.Ma perchè le GPU risultano essere tanto e�caci ed e�cienti in questo am-bito?La risposta va ricercata nella struttura di un generico componente 3D. Sem-pli�cando possiamo dire che un oggetto 3D è un insieme di tanti triangolitridimensionali, ognuno dei quali presenta tre vertici, che nel piano x,y,z han-no delle speci�che coordinate. Tali coordinate devono essere processate al �nedi voler applicare una qualsivoglia trasformazione (ad esempio: traslazione,scaling.. ) all'oggetto. Ad ogni trasformazione corrisponde una speci�ca 3Dscaling matrix, ovvero una matrice standard che moltiplicata per la matri-ce contenente le coordinate dei vertici, dà luogo proprio al comportamentodesiderato sull'oggetto 3D. Dunque è importante, per una corretta �uiditàdi gioco, una componente hardware che sia in grado di svolgere operazioniquali la moltiplicazione tra matrici, calcoli in virgola mobile, moltiplicazionitra vettori etc. in maniera rapida. Tale componente hardware è proprio laGPU, che per costruzione è in grado di svolgere questo tipo di calcoli rag-giungendo velocità e performance che una CPU non fornisce perchè, ad oggi,la sua tecnologia è più orientata per l'elaborazione di istruzioni.


1.4 Intelligenza Arti�ciale

Ostico ma a�ascinante.L' AI (Arti�cial Intelligence) è un argomento che ad oggi suscita grandeinteresse e grandi apprensioni. Un giorno i robot davvero conquisteranno ilmondo?Due sono le scuole di pensiero: la prima, di cui fa parte Ray Kurzweil, capotecnologo di Google, sostiene che l'AI sarà in grado di dar luogo alla cosid-detta Singolarità tecnologica de�nita come il radicale cambiamento dellasocietà umana nel momento in cui una superintelligenza arti�ciale innescheràl'inizio di una crescita tecnologica senza precendenti, il tutto in termini po-sitivi. La seconda scuola di pensiero, Future of Life Institute di Boston,tra i cui sostentiori troviamo scienzati del calibro di Stephen Hawking e im-prenditori come Bill Gates e Elon Musk, si impegnano a lanciare l'oppostoavvertimento: l'intelligenza arti�cale è un "rischio esistenziale" con un altopotenziale di danneggiamento per la specie umana. Mettendo da parte lequestioni etiche e morali, le GPU si inseriscono in questo scenario come ipilastri dell'AI, in particolar modo, come i componenti principali che hannodato luogo al deep learing, processo coaudivante la machine learing, chesfruttando le reti neurali 2 permette l'apprendimento automatizzato. Il deeplearing è diventato possibile grazie alla convergenza di tre fattori: l'aumentodella capacità di calcolo dei microprocessori, la disponibilità di grandi basi didati su cui fare allenare le macchine arti�cialmente intelligenti ed ultimo, manon per importanza, la crescente so�sticazione anche dal lato software conla creazione di tecniche ad algoritmi sempre più complessi e funzionali. Diquesti tre fattori il primo è però vincolato alla Legge di Moore (la potenzadi calcolo raddoppia ogni due anni) che sta per superare i suoi limiti �sici neichip di silicio. Proprio per questo da circa tre anni le GPU sono diventatel'unità di elaborazione prediletta dall'AI, in quanto, come già detto, operanoin parallelo su grandi moli di dati e sono più e�cienti.

2Con "rete neurale" s'intende un sistema di network arti�ciale che s'inspira di fatto allastruttura biologica del cervello umano, astraendone il funzionamento attraverso relazionimatematiche/logiche

Capitolo 2

NVIDIA

2.1 Architetture

Nvidia è la casa produttrice leader del settore delle schede gra�che. Nel1999 viene lanciata sul mercato la prima GPU, GeForce256, de�nita comeun processore su singolo chip capace di operare trasformazioni su 10 milionidi poligoni per secondo. La tecnologia da quel lontano 1999 è andata moltoavanti. Le attuali GPU superano di gran lunga quei record con circa 7 mi-liardi di poligoni elaborati al secondo.

In ordine cronologico, le sette micro-architetture 1 GPU sono:

2.1.1 Tesla

Tesla è stata la primissima microarchitettura lanciata sul mercato da NVI-DIA.Rispetto all'architetture odierne ogni SM, presentava solamente 8 ShaderProcessor, o CUDA core, e 2 SFU, unità dedicate al calcolo delle funzionitrascendentali. Ogni SM poteva eseguire �no a 24 operazioni per clock, aseconda del tipo di operazione.

1progetto di un processore e non di un calcolatore nella sua interezza (a quest'ultimoci si riferisce come architettura). La microarchitettura de�nisce l'implementazione di unISA, nonchè la disposizione di unità funzionali come, ad esempio, cache, pipeline,unitàSIMD etc

6

CAPITOLO 2. NVIDIA 7

2.1.2 Fermi

Fermi, in onore del �sico italiano Enrico Fermi, è la seconda microarchitet-tura di NVIDIA, subentrata a Tesla.Conta ben 3 miliardi di transistor per un totale di 512 CUDA cores (an-che chiamati Stream Processors), disposti a gruppi di 32 in 16 distintiStreaming Multiprocessors. Rispetto all'architettura precedente si riuscìad ottenere un miglioramento di circa 8 volte superiore nello speedup deicalcoli in double precision, che per l'epoca rappresentava un notevole passoavanti per lo sviluppo di applicazioni orientate all'elaborazioni scienti�che.Oltre a questo, altre importanti novità furono: l'introduzione dell' ECC (Error Correction Code) per consentire il rilevamento, la correzione e laprotezione delle caches, dei registri e del resto dei componenti di memoria daerrori dovuti a cause esogene, e il completo supporto per il linguaggio C++grazie all'uni�ed address space che consentì l'uso agevole delle classichefunzioni C++ di allocazione/deallocazione della memoria.All'interno di ogni singolo cuda core troviamo una FPU (Floating-PointUnit) e una ALU (Arithmetic logic unit) ognuna con propri pipelinedatapath. 2

Importante ruolo riveste il Wrap Scheduler e la Instruction DispatchUnit entrambi doppi in Fermi. Questi due componenti si occupano di sche-dulare e quindi di eseguire due wrap concorrentemente ed indipendentementein maniera lock-step. Qualora le risorse per l'esecuzione di un wrap non fos-sero disponibili l'SM può schedulare e quindi mandare in esecuzione un altrowrap senza perdita di e�cienza dal momento che l'overhead di un contexswitch fra thread è pari a zero.Per quanto riguarda le memorie, Fermi presentava due livelli di cache L1 eL2, la prima posizionata all'interno di ogni singolo SM è con�gurabile se-condo la disposizione di 48 KB di shared memory e 16 KB di L1 cache oviceversa a seconda delle necessità e L2 cache, che si rapportava direttamen-te con la DRAM (Global memory) e utilizzata per tagliare i tempi e i costidei trasferimenti di memoria con l'host.

2struttura che consente un aumento di throughtput, consentendo l'esecuzione di piùistruzioni contemporaneamente.


Figura 2.1: microarchitettura Fermi

2.1.3 Kepler

La microarchitettura Kepler, in memoria di Johannes Kepler, scienziato te-desco, fu la successiva soluzione hardware lanciata sul mercato dopo Fermi.Fu la prima microarchitettura a concentrarsi in particolar modo al migliorarel'e�cienza energetica. Con Kepler furono introdotte numerose novità:

• SMX: nuovi Streaming Multiprocessors, caratterizzati, rispetto ai pre-cedenti SM, da 192 cores ( rispetto ai 32 di Fermi), da 32 SFU (SpecialFunction Unit, rispetto alle 4 di Fermi e altre migliorie hardware. Trai vari perfezionamenti spicca la presenza di 4 unità wrap schedulingcapaci di smistare istruzioni e schedulare esecuzioni su 8 dispatch uni-ts, aumentando il complessivo throughtput. Al posto dei due wraps (64threads) che potevano essere schedulati ed eseguiti su Fermi contempor-taneamente e concorrentemente in modalità SIMD, il numero di wrapstrattabili sale a 4, con la capacità di poter a�dare a coppie di wrapsdue istruzioni diverse.

• Hyper-Q: in Fermi si riscontrò un problema di e�cienza che com-portava la concentrazione del carico di lavoro solo su parte dei cuda


cores, mentre altri restavano senza lavoro. Questo comporta una per-dita di e�cienza e soprattutto un tempo di elaborazione più elevato.Per questo su Kepler sono aumentate le linee di comunicazione con laCPU, 32 simultanee e concorrenti in totale, rispetto all'unica presen-te in Fermi, che consente ai singoli core della CPU di poter inviaremaggiore workload alla singola GPU, inducendo di fatto un aumentodelle prestazione e una drastica diminuizione dei tempi morti, nonchèun quantitativo di lavoro disponibile per la maggior parte dei cores.L'Hyper-Q quindi aggira in questo modo il problema delle dipenden-ze e dell'accodamento di lavoro sull'unica linea di comunicazione la-sciando Kepler con prestazioni 32 volte superiori rispetto l'antecententearchitettura.

• Dynamic-Parallelism: questa funzionalità è esprimibile come "laGPU crea al volo lavoro per se stessa". L'espressione vuole indicareche la GPU è in grado di sbrigare da sola eventuali innesti e salti pre-senti all'interno di kernel, adattandosi quindi alle dipendenze intrisecheche presentano i dati, senza rimandare i dati alla CPU. In questo modosi limitano i contatti con la CPU, di conseguenza i tempi di attesa edi elaborazione, e la CPU riceve solamente i risultati �nali richiesti dalkernel.In questo modo la GPU crea lavoro per se stessa, essendo indipendentedi lanciare altri �ow di lavoro dal kernel di partenza.

2.1.4 Maxwell

Maxwell, dal �sico e matematico scozzese James Clerk Maxwell, è la microar-chitettura che subentò a Kepler e consta di due generazione.Riassumento le novità introdotte per entrambe le generazioni:

• SMM: Streaming Multiprocessors (e.g. 16 sulla GX980 con 128 cu-da cores ciascuno, organizzati in 4 blocchi da 32 cores ognuno con unproprio wrap scheduler e shared memory) volti a migliorare ulterior-mente l'e�cienza energetica, attraverso la riorganizzazione di un nuovodatapath e l'introduzione di istruction bu�ers, che rispetto a Keplerche ne era sprovvista, ha aumentato molto le performance diminuendoi tempi di attesa per le istruzioni.


• Instruction Scheduling: Nonostante la diminuzione del numero dicore presenti rispetto a Kepler, la loro distribuzione su chip, in blocchida 32, rende possibile un migliore è più facile smistamento delle istru-zioni mantenendo lo stesso numero di istruzioni per clock di Kepler maal contempo riducendo tempi di latenza.Infatti, ogni wrap scheduler rilascia un wrap su un gruppo di coresesattamente pari alla lunghezza del wrap in esecuzione (32 threads su32 cores).

• Aumento Shared Memory: Quantitativo disponibile per la sharedmemory privata per thread è, in Maxwell, di 48KB �ssi e di 16 diL1 cache. La novità introdotta è che è stato possibile combinare lefunzionalità della L1 e delle texture caches in una singola unità. Nellaseconda generazione, anche il quantitativo di cache L2 fu incrementato.

• Compressione Dati Delta: Dal momento che texture, informazionisul lighting, bump, frame etc.. sono salvate nel frame bu�er la velocitàcon cui si comunicava con questo componente era sì vitale (essendopoi stata risolta con un'ampli�cazione di bandwitdh 3per l'accesso allamemoria) ma non era l'unica componente che determinava un aumentodelle performance. La compressione Delta di dati permette di calcolareil delta confrontando l'ultimo dato immagazzinato nel frame bu�er (edè qui che entra in gioco la compressione, dati uguali di un'immagine,ad esempio il colore rosa, venivano compressi) con l'informazione cheinvece stava per essere scritta sul frame bu�er. Questo espediente riuscìa ridurre del 30% il tra�co verso la memoria rispetto a Kepler.

2.1.5 Pascal

Pascal, in ricordo di Blaise Pascal, è l'architettura successiva a Maxwell. Ri-lasciata nell'aprile del 2016, nell'era in cui l'AI sta impossessandosi sempredi più dell'attenzione tecnologica, Pascal nasce con l'intento di accelerare inparticolar modo le applicazioni deep learing tassello chiave dell' arti�calintelligence.L'architettura è stata realizzando sfruttando la tecnologia a 16nm FinFET,

3quantità di informazioni massime che riescono a passare su un percorso hardware


Figura 2.2: microarchitettura Maxwell

il che ha comportanto una completa rivoluzione rispetto all'architetture pre-cedenti.A seconda del modello, si può avere una GPU con numero di SM (�no a 60)aventi numero di cores interni variabili (nel top di gamma GP100, la con�-gurazione conta 64 cores per SM, mentre per le versioni per desktop si arrivaa 128 cores per SM).Tra le novità introdotte troviamo delle unità addizionali per il calcolo indouble-precision, un aumento del quantitativo di shared memory, L1 cache edi register �le bu�er. Ovviamente a seconda del modello si avranno diversecon�gurazioni hardware.Da un punto di vista prestazioni si incontrano le seguenti innovazioni:

• Memoria di tipo 3D: cambia la posizione dei chip di memoria che,invece di risiedere a distanza di centimetri dalla GPU, verranno impilatiin maggiore prossimità del processore gra�co in maniera tale da ridurreil percorso che i bit devo attraversare per ogni load/store dalla memo-ria con il risultato di avere maggiore e�cienza energetica e maggioriprestazioni. Il che comporterà una triplicazione della banda passanterispetto alla soluzione architetturale precedente.


• Calcoli a precisione mista: grazie ai processori migliorati, al quan-titativo di memoria disponibile maggiore, alla maggiore larghezza dibanda e ad un'ampli�cazione delle istruzioni eseguibili in un singologiro di clock, i calcoli a precisione mista cioè calcoli in virgola mobilea 16 bit sono eseguiti con una precisione doppia rispetto ai calcoli invirgola mobile a 32 bit.

• NVLink: si tratta di un collegamento Host-GPU che sostituisce difatto il precedente PCI bus, garantendo un passaggio di dati tra i duesistemi, di 12 volte più veloce rispetto a Maxwell.

Figura 2.3: microarchitettura Pascal

2.1.6 Volta

Volta, in onore di Alessandro Volta, è l'architettura che vede come principaleinnovazione l'introduzione dei Tensor Core.Queste unità sono state introdotte per migliorare i tempi e le prestazionidel deep learing tipico delle nuove tecnologie di intelligenza arti�ciale. Sulla


scheda gra�ca V100, se ne contano ben 640. Ciascuno di essi opera, molti-plicando, due matrici 4x4 FP16 (�oating point a half precision, ovvero cheoccupa due byte di memoria), e poi aggiunge una terza matrice sempre di�oating point, a 16 o 32 bit, utilizzando operazioni FMA4. Il risultato ot-tenuto può essere sia a 32 bit che ridotto a 16 bit. Ogni Tensor core puòsvolgere 64 operazioni FMA il che permette, ad ogni sessione di deep learingdi utilizzare i risultati scegliendo quale tipo di precisione è più opportuna peri dati e per i percorsi intrapresi durante l'allenamento della macchina.Le fasi di training e di inferenza 5 dei sistemi AI subiscono un drastico taglioin materia di tempo di assimilazione e un incremento di prestazioni di circa3 volte superiore rispetto alla generazione precedente.

Figura 2.4: microarchitettura Volta, per uso professionale

4operazione di moltiplicazione-addizione in virgola mobile eseguita in un unico passocon un solo tentativo di approssimazione

5capacità basata su algoritmi per simulare come una mente umana prende decisioni siain senso deduttivo che induttivo


2.1.7 Turing

Ultima architettura rilasciata da NVIDIA nel 2018, che eredita molte caratte-ristiche della precedente architettura Volta. Per quanto riguarda le speci�chetecniche la documentazione è ancora scarseggiante ma da anticipazioni da-te dal CEO NVIDIA nell'ultima presentazione in cui non si speci�ca unaparticolare scheda gra�ca, sono quelle presenti in �gura.

Figura 2.5: microarchitettura Turing

Principale protagonista è l'intoduzione diRT (Ray Tracing), particolariunità nate con l'intento di migliorare ulteriormente il rendering ibrido (chemixa il ray tracing6 con la rasterizzazione) al �ne di ottenere una simulazio-ne della realtà d'interesse dell'elaborazione quanto più fedele possibile allarealtà �sica.Le nuove schede gra�che saranno quindi i pilastri su cui si poggerà, secondoJen-Hsun Huang, CEO NVIDIA, l'intelligenza arti�ciale e i suoi meccanismidi deep learing.Ritroviamo nella con�gurazione Turing anche i Tensor core, introdotti in Vol-

6tradotto "tracciamento di raggi", indica un modo di progettazione usato in computergra�ca, che modella le scene 3D in base alla ri�essione della luce della telecamera suglioggetti della scena.


ta.Miglioramenti sono stati e�ettuati anche sul fronte Streaming processor, an-dando a separare le unità integer da quelle �oating point. In questo modo siottiene, stando a quanto riporta l'azienda di Santa Clara, un miglioramentonelle prestazioni FMA(Fused Multiply Add).Ritroviamo anche in Turing, l'NVLink, ovvero il bus che sostituisce l'ormaiobsoleto PCI bus, che permette il collegamento di due schede in parallelosu un collegamento ad altà velocità con annessa possibilità di condividererisorse tra le due GPU.L'ultima novità è il cosiddetto VirtualLink, un collegamento che consen-tirà di connettersi a futuri visori di realtà aumentata con un incremento di�uidità.

2.2 CUDA

CUDA (Compute Uni�ed Device Architecture) nasce in casa NVI-DIA nel 2006 e de�nisce un'architettura di elaborazione in parallelo nonchéun modello di programmazione che permette e�ettivamente di programmarele strutture hardware intensamente parallele tipiche delle GPU. IntroducendoCUDA le ultime GPU NVIDIA hanno subito una grande rivoluzione passan-do da sistemi chiusi, cioè programmabili solamente da NVIDIA, a sistemiaperti, programmabili da terzi.L'introduzione di CUDA è nata dalla necessità di facilitare il general-purposecomputing on graphics processing units, che vede la GPU come ruolodi co-processore della CPU per l'accelerazione di applicazioni che richiedonograndi capacità computazionali.CUDA non è un linguaggio del tutto nuovo, ma rappresenta un'estensionedel C e del C++, fornendo delle funzionalità che facilitano lo sviluppo dielaborazioni parallele.

2.2.1 Compilazione ed esecuzione

A�nchè un programma, e quindi il suo codice, venga eseguito dalla GPU,questo ha bisogno di essere compilato. La compilazione avviene grazie alnvcc che è appunto il compilatore CUDA. Il codice scritto in CUDA vie-ne quindi tradotto dal nvcc in PTX (Parallel Thread Execution) che èun linguaggio pseudo-assembly tipico dell'ambiente CUDA. Successivamen-


te il PTX viene tradotto in codice binario dal compilatore contenuto neldriver della scheda gra�ca interessata per poi essere eseguito dall'hardwaresottostante.

2.2.2 Funzionamento

Concetto centrale della programmazione in ambiente CUDA è il kernel, cheè assimilabile ad una funzione del linguaggio C. Il kernel contiene il codiceche verrà poi eseguito dai threads, che sono l'unità base di esecuzione. Talithreads vengono raggruppati in blocks, e il numero di thread appartenentiad un blocco è de�nito al momento del lancio del kernel da parte del pro-grammatore ed è soggetto comunque alle speci�che hardware della schedagra�ca che si utilizza. I blocks sono poi raggruppati in grids, ovvero in gri-glie. Quando lanciamo in esecuzione un kernel stiamo di fatti lanciando unagriglia, dal momento che andiamo a speci�care sia il numero di blocchi inte-ressati dall'elaborazione e sia il numero di thread appartenenti a tali blocchi.Il "lancio" di un kernel avviene attraverso la notazione <<< e >>>Ad esempio, la con�gurazione:esempioKernel<<<100, 128>>>(...)prevede 100 blocchi di 128 thread ciascuno per un totale di 12800 threads.Ogni kernel viene eseguito dagli SM (Streaming Multiprocessor) pre-senti sulla scheda gra�ca. Precisando, un kernel viene eseguito dal numerodi blocchi de�niti dalla con�gurazione di lancio e dai thread facente partedi tali blocchi. Quindi all'esecuzione di un kernel avremo che i blocchi dithread verranno eseguiti su un SM, ma lo streaming multiprocessor suddi-viderà ulteriormente il blocco di thread in un'altra unità funzionale, dettawrap, formata da 32 threads, al �ne di garantire una corretta condivisionedelle proprie risorse interne perchè se un wrap necessita di un accesso in me-moria, quell'eventuale tempo di attesa, viene recuperato dall'SM mandandoin esecuzione un altro wrap. Ogni wrap viene eseguito in modalità SIMD(Single instruction Multiple data), cioè i 32 thread eseguono la stessaistruzione su dati diversi, tutti nello stesso momento. Può succedere che peròsi veri�chi una condizione che comporti un branching all'interno del wrap.Tale situazione viene a�rontata eseguendo ogni rami�cazione del codice inmaniera seriale andando a disattivare i thread che in quel momento non sitrovano sul percorso in esecuzione. Una volta eseguiti tutti i percorsi natidal branching i thread riconvergono nello stesso percorso di esecuzione.


2.2.3 Modello di Memoria

Ogni GPU basata su CUDA ha a disposizione diversi tipi di memoria chesi di�erenziano in base all'uso, allo spazio di indirizzamento, allo scope esoprattutto alla velocità che in�uisce quindi sul tempo di latenza.

La memoria in CUDA è suddivisa:

• Globale: accessibile in R/W da tutti i thread e dall'host (CPU), tempidi latenza elevati ma ampio spazio di indirizzamento.

• Costante: accessibile in sola lettura, da tutti thread e dalla CPU,anch'essa presenta tempi di latenza elevati ma comunque più velocirispetto alla memoria globale perchè è cached e dato che è in solalettura è garantito che i dati non verranno mai invaldati dai thread ameno che non vengano esplicitamente modi�cati dalla CPU. Rispettoalla memoria globale, lo spazio di indirizzamento è nettamente inferiore.

• Locale: accessibile in R/W da un singolo thread, ha tempi di latenzapari a quelli della memoria globale, dunque è molto lenta. Viene utiliz-zata nei casi di register-spilling, cioè quando �nisce lo spazio fruibiledei registri e di conseguenza serve altro spazio per altre variabili ne-cessarie ai thread ad esempio. Oppure quando sono coinvolte grandistrutture di dati che consumano molto spazio nei registri.

• Shared: accessibile in R/W, da tutti thread appartenenti allo stessoblocco. Ha una velocità cento volte superiore alla memoria globale,ma di contro ha dimensioni ridotte. Può essere modi�cata dal kernel eper evitare eventuali situazioni di pericolo e di invalidità l'accesso deveessere sincronizzato.

• Registri: accessibili in R/W, da un solo thread, caratterizzati dallavelocità più elevata rispetto a tutte le altre memorie. Le variabili de�-nite in un kernel sono di default disposte all'interno dei registri a menoche queste siano troppe. In quel caso le rimanenti variabili verrannotrasferite, come già detto, nella memoria locale. Dal momento che loscope è per un singolo thread gli accessi ai registri non necessitano disincronizzazione.

Una considerazione importante da fare è collegata alla compute capa-bility 7 dei dispositivi. Per schede gra�che con compute capability di 1.x,

7indica le speci�che e le caratteristiche disponibili su una data GPU


le memorie possono essere "on-chip" e sono quelle cache, ovvero memoriemolto veloci e quindi con bassa latenza, oppure possono essere posizionate"o�-chip" e per questi tipi di memorie i tempi di accesso aumentano in ma-niera considerevole. Tra le memorie on-chip troviamo la shared memory ela memoria constante, mentre le memorie o�-chip sono la globale, la localee la parte uncached della memoria costante. Su dispostivi che vantano unacompute capability di 2.x o superiore ci sono due addizionali memorie: unacache L1 installata all'interno di ogni SM e un cache L2 condivisa tra tuttigli SM. Dato che la memoria globale è distante rispetto ai multiprocessori,sono di supporto la cache L2 e la cache L1. Se i thread vogliono accederead un dato presente nella shared memory oppure dalla cache L1, i tempi dilatenza sono estremamente ridotti. Ma se un thread deve leggere dalla globalmemory deve controllare prima che:

• Il valore è presente nella cache L1. Se lo è, ci si ferma qui.

• Se non è presente nella L1, si controlla se è presente nella L2, ma sel'esito è negativo allora bisogna per forza accedere alla memoria globalee quindi l'accesso diventa molto lento.

Per questo motivo l'introduzione delle due memorie cache ha favoritonotevoli aumenti di performance.

Sincronizzazione

Parlare di sincronizzazione quando si ha che fare con programmazione pa-rallela potrebbe sembrare un controsenso, perchè in linea di massima unasincronizzazione prevede poi una scelta di chi e cosa far eseguire, il che puòportare poi ad una serializzazione del codice e quindi alla snaturazione dellaconcetto di parallelismo. Ma in alcuni casi questa è indispensabile al �nedi evitare casi di inconsistenza di dati e conseguenti malfunzionamenti, tra iquali troviamo i bank con�icts, che avvengono quando un halfwrap (cioèuna collezione di 16 threads) tenta di leggere/scrivere dati dal/su lo stessobanco 8 di memoria. Questo tipo di problema può essere risolto facendo usodi request patterns, particolari modalità standard di richiesta dati che re-golamentano quindi gli accessi in memoria. Un altro modo per sincronizzare

8La memoria condivisa è suddivisa in memorynbank che a loro volta sono suddivisi inword, ovvero in sezioni di 4 byte ciascuna.


Figura 2.6: Modello di Memoria CUDA

gli accessi in memoria, e quindi consentire o meno manipolazioni di dati,consiste con l'uso di syncthreads() che ha la funzione di barriera a livellodi blocco. Nessun thread del blocco può sorpassare la condizione de�nita dasyncthreads() prima che tutti i thread del blocco non hanno anch'essi rag-giunto quel punto. Un ulteriore modo consiste, invece, nel blocking. Ruolocentrale ha la shared memory in cui vengono caricate porzione della memo-ria globale. Una volta che i thread hanno eseguito il loro compito sui datipresenti in shared memory, i risultati vengono spostati in memoria globaleed eventualmente verrà copiata in shared memory un'altra porzione dellamemoria globale per essere oggetto di altri calcoli.

2.2.4 Esempio di Programmazione CUDA

Di seguito riporto la programmazione di una semplice elaborazione in paral-lelo, sfruttando la scheda gra�ca Nvidia GeForce 940M ( compute capability5.0 basata su microarchitettura Maxwell ).L'obiettivo del programma è quello di e�ettuare un'operazione di moltiplica-zione tra vettori di lunghezza arbitraria de�nita dal programmatore eseguen-dola in parallelo, ovvero ogni thread, in base al proprio id, "sceglie" i datisu cui lavorare contemporaneamente agli altri threads in modalità SIMD, in


maniera tale da ottenere un taglio di tempo e di costi rispetto ad un codiceseriale. Il codice è stato sviluppato in Visual Studio 2017, con plug-in CUDA9.0, sfruttando il toolkit Visual Studio 2015.

Figura 2.7: librerie utilizzate e de�nizione del kernel.


Figura 2.8: creazione spazio di memoria e de�nizione valori vettori sull'Host(CPU)

Figura 2.9: creazione spazio di memoria su GPU, tramite estensione fornitada CUDA


Figura 2.10: copia dei valori de�niti lato Host sullo spazio di memoriaprecedentemente allocato sulla GPU

Figura 2.11: con�gurazione di lancio del kernel da lato host, dal momentoche il kernel è stato de�nito come global


Figura 2.12: copia dei dati elaborati dai threads dalla GPU alla CPU, tramitecudaMemcpy

Figura 2.13: stampa a video dell'elaborazione eseguita

Capitolo 3

OPENCL

3.1 Cos'è e come funziona

OPENCL (Open Computing Language), è un framework 1 che è utiliz-zato per scrivere codice portabile, quindi fruibile da più architetture comeCPU, GPU e DSP (Digital Signal Processor). Consente notevoli aumento diprestazioni per applicazioni di calcolo intensivo e di conseguenza, è utilizzatosu ampia scala per accelerare la computazione parallela.Così come in CUDA, anche in OPENCL l'unità base di codice eseguibile èchiamata kernel, scritti in OPENCL C. La codi�ca del kernel consente losviluppo di applicazioni sia orientate al task parallelism, ovvero al paralle-lismo/distrubuzione di compiti su diversi processori, e al data-parallelism,concetto precedentemente trattato.I kernel, insieme ad altre funzioni, costituiscono un programma. L'esecu-zione del programma e quindi dei relativi kernel può avvenire in due modi: oin ordine o fuori ordine. Spiegandolo meglio, a�nchè i kernel siano eseguitidevono essere prima accodati (cioè disposti in una command queue, unadiversa per ogni host e ogni device) da parte dell'host (CPU) per mandarelavoro al device (ad esempio una GPU). Il lavoro da fare è accodato in ordi-ne di rilascio da parte dell'host ma sul device può essere eseguito in ordineoppure fuori ordine, a seconda della disponibilità di risorse e di altri fatto-ri che concorrono, ovviamente, ad un'ottimizzazione dell'esperienza d'uso.Questa decisione è presa run-time dal device stesso. L'insieme dei device,delle command queues e delle loro memorie è detto context, che ha anche

1una struttura logica di "guida" e di supporto per lo sviluppo di applicazioni software

24

CAPITOLO 3. OPENCL 25

l'onere di abilitare la condivisione di memoria tra i device d'interesse. InCUDA abbiamo che i kernel sono eseguiti da threads, mentre in OPENCLla terminologia cambia: ad eseguire i kernel sono i cosiddetti work-items,cioè singoli elementi di esecuzione tra loro indipendenti.L'intero spazio del problema è de�nito da un dominio globale n-dimensionale(con N =1,2,3) che de�nisce il numero totale di work-item che eseguono inparallelo. La suddivisione del dominio globale del problema in domini lo-cali, di spazio inferiore, de�nisce invece, la dimensione dei work-groups,cioè una collezione di work-items (un po' come accade in CUDA con i blocksdi threads). Facendo riferimento all'architettura hardware in �gura, ogniwork-item è eseguito da una compute unit (in rosso, come un singolo coredi un processore), ogni work-group è eseguito da un processing element(in blu, un processore/microprocessore), ed in�ne un kernel è eseguito dall'OPENCL Compute Device (CPU,GPU.. etc).

3.1.1 Modello di Memoria

Il modello di memoria utilizzato in ambiente OPENCL, molto simile a quelloCUDA, è descritto in �gura:

• Memoria Privata: riservata ed accessibile ad un singolo work-item.

• Memoria Locale: condivisa in un work-group (consente quindi lacollaborazione/sincronizzazione dei work-item appartenenti allo stessowork-group)

• Memoria Globale/Costante: non è sincronizzabile, ma è visibile datutti i work-group.


• Memoria Host: presente sull'host. Utilizzata per scambiare informa-zioni e comunicare con la memoria globale del device.

Una considerazione importante da fare è che l'uso della memoria è espli-cito, ovvero lo spostamento dei dati dall'host alla memoria globale accessibiledal device e viceversa è a carico del programmatore.

Figura 3.1: Modello di Memoria OPENCL

Sincronizzazione

La sincronizzazione in OPENCL è uno strumento che è utilizzato solamenteall'interno di work-groups, ciò signi�ca che tra i work-item che appertengonoallo stesso gruppo di lavoro si può sincronizzare il loro operato attraversomeccansimi di barriere (intesi come condizione di sincronizzazione) e dimemory fances che sono anch'esse barriere che però regolamentano gli ac-cessi alle memorie. Due work-items che invece non appartengono allo stessogruppo, sebbene condividano lo stesso dominio globale, non possono esseresincronizzati in alcun modo. La sincronizzazione può essere regolamentata


anche da eventi qualora sia presente un grado di collaborazione tra kernel.Come mostrato in �gura:

Figura 3.2: scenario in cui il kernel2 deve usare i dati modi�cati dal kernel1,serve una sincronizzazione

3.1.2 Portabilità

La caratteristica principale che contraddistingue OPENCL da CUDA è ap-punto la portabilità. Mentre CUDA rende programmabili solamente dispo-stivi CUDA-Enabled, comunque di casa NVIDIA, OPENCL permette diutilizzare lo stesso codice per diversi device.

Capitolo 4

Conclusioni

4.1 Non solo NVIDIA

Nvidia è ad oggi l'azienda leader del settore schede gra�che, ma questo nonsigni�ca che sia l'unica.AMD(Advanced Micro Devices), conosciuta anche come storica anta-gonista della Intel nel campo CPU, starebbe progettando una nuova GPUa 7 nanometri 1 in arrivo entro la �ne del 2018, mentre NVIDIA invece èimpegnata nel lancio della nuova microarchitettura Turing basata su unatecnologia a 12 nanometri. Nanometri a parte, anche Intel, sebbene abbiaperso terreno nello scenario GPU è in procinto di rientrare gloriosamentenel mercato della graphic computing nel 2020, avendo annunciato una nuo-va GPU discreta che, secondo indiscrezioni, sarebbe la protagonista di unascheda video a singolo slot con un connettore PCIe e una ventola a turbinaper dissipare il calore.Se NVIDIA verrà spodestata dalla Intel e le sue turbine o da AMD con i suoi7 nanometri, sarà il tempo a dircelo.

1con 7 nanometri s'indica la dimensione del gate di un transistor, non la sua dimensionetotale

28

Bibliogra�a

[1] Cuda Programming Guide, link: https://developer.download.nvidia.com/compute/DevZone/docs/html/C/doc/CUDA_C_Programming_Guide.pdf.

[2] Sito NVIDIA, link: https://www.nvidia.com/it-it/.

[3] NVIDIA forum, link: https://devtalk.nvidia.com/

[4] Cervello. Manuale dell'utente. Guida sempli�cata alla macchina piùcomplessa del mondo, Marco Magrini.

[5] Nvidia Pascal: i punti chiave dell'architettura che sostituirà Maxwell,Manolo De Agostini.

[6] Turing, Nvidia svela l'architettura delle nuove GPU, Manolo De Agostini.

[7] Conferenza: GTC 2012 Keynote (Part 03): TheNVIDIA Kepler GPU Architecture, link youtu-be: https://www.youtube.com/watch?v=TxtZwW2Lf-w&index=23&list=PLQgupYgwgFT6Fcj70Ph8v82i6vLQXvrYN&t=0s

[8] Conferenza: Nvidia GTC 2009: GT300 "Fermi" ar-chitecture unveiled by Jen Hsun Huang, link youtube:https://www.youtube.com/watch?v=fYuH2Kl_b98&index=20&list=PLQgupYgwgFT6Fcj70Ph8v82i6vLQXvrYN&t=493s

[9] Conferenza: NVIDIA Keynote 2014 � Max-well, GeForce GTX 980 & 970, link youtube:https://www.youtube.com/watch?v=LUdP8_NQmgk&index=21&list=PLQgupYgwgFT6Fcj70Ph8v82i6vLQXvrYN&t=546s

29

BIBLIOGRAFIA 30

[10] Conferenza: GTX 1080 Pascal Architecture Explai-ned, link youtube: https://www.youtube.com/watch?v=-jgb1Uizz7o&list=PLQgupYgwgFT6Fcj70Ph8v82i6vLQXvrYN&index=14&t=458s

[11] Conferenza: SIGGRAPH 2018 - NVIDIA CEO Jen-sen Huang - Reinventing Computer Graphics, link youtube:https://www.youtube.com/watch?v=jY28N0kv7Pk&list=PLQgupYgwgFT6Fcj70Ph8v82i6vLQXvrYN&index=13&t=1518s

[12] CUDA Part A: GPU Architecture Overview and CU-DA Basics; Peter Messmer (NVIDIA), link youtube:https://www.youtube.com/playlist?list=PLQgupYgwgFT6Fcj70Ph8v82i6vLQXvrYN

[13] WhitePaper Volta, link: https://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf

[14] WhitePaper Fermi, link: https://www.nvidia.com/content/PDF/fermi_white_papers/NVIDIA_Fermi_Compute_Architecture_Whitepaper.pdf

[15] WhitePaper Maxwell,link: https://international.download.nvidia.com/geforce-com/international/pdfs/GeForce_GTX_980_Whitepaper_FINAL.PDF

[16] OpenCL wikipedia, link: https://it.wikipedia.org/wiki/OpenCL

[17] OpenCL Technical Overview, link youtube:https://www.youtube.com/watch?v=aKtpZuokeEk&list=PLVk9nlso0x0K0pRMzEj0-kbidFBxohTTk

Tesi di Laurea Triennale - ingegneria-informatica.unina.it · Due sono le scuole di pensiero: la...

Documents

Transcript of Tesi di Laurea Triennale - ingegneria-informatica.unina.it · Due sono le scuole di pensiero: la...