Scuola Politecnica e delle Scienze di BaseCorso di Laurea in Ingegneria Informatica

Elaborato finale in Programmazione I

Modelli di programmazione per il calcolo paralleloParallel programming models

Anno Accademico 2015/16

candidatoMichele Caggianomatr. N46001601

"Ora non è il momento di pensarea quello che non hai. Pensa a quello che puoi fare con quello che hai." ­ Ernest Hemingway

Alla mia famiglia, ai miei compagni di viaggio, a chi mi ha supportato e a chi mi sopporta.

2

Indice

Indice....................................................................................................................................................3Introduzione..........................................................................................................................................4Capitolo 1: L'evoluzione del modello architetturale............................................................................5

1.1 Dal sequenziale al parallelo........................................................................................................51.2 Lo sviluppo tecnologico: il general purpose many core.............................................................6

Capitolo 2: I modelli per la programmazione parallela........................................................................82.1 Il modello pipeline......................................................................................................................82.2 Il modello master-slave..............................................................................................................9

Capitolo 3: Il mondo della GP-GPU..................................................................................................103.1 La tassonomia di Flynn.............................................................................................................103.2 L'architettura di una GPU.........................................................................................................113.3 Lo scontro generazionale: GP-GPU vs CPU............................................................................133.4 I modelli di programmazione su GP-GPU...............................................................................14

Capitolo 4: CUDA..............................................................................................................................154.1 La gestione dei thread in CUDA..............................................................................................154.2 L'organizzazione gerarchica della memoria in CUDA.............................................................164.3 La struttura di un codice CUDA...............................................................................................184.4 Implementazione di CUDA in C..............................................................................................194.5 Esempio di codice CUDA C.....................................................................................................21

4.5.1 provaCUDA.cu..................................................................................................................224.5.2 Analisi del codice..............................................................................................................264.5.3 Esecuzione su Nsight.........................................................................................................294.5.4 Profiling.............................................................................................................................294.5.5 Confronto con codice sequenziale di funzionalità equivalente.........................................32

Conclusioni.........................................................................................................................................35Sviluppi Futuri....................................................................................................................................37Bibliografia.........................................................................................................................................38

3

Introduzione

Questo elaborato è stato redatto al fine di presentare la natura architetturale ed il recente

sviluppo dei modelli di programmazione per il calcolo parallelo. Quest'ultima è una

tecnica informatica nata dall'evoluzione della classica programmazione sequenziale e

basata sull'esecuzione simultanea e concorrente del codice sorgente di uno o più

programmi, allo scopo di aumentare le prestazioni del calcolo attraverso l'ottimizzazione

dell'uso delle risorse a disposizione e del relativo tempo di esecuzione, pur necessitando di

una maggiore capacità del programmatore nel dividere e specificamente adattare su più

microprocessori o più core dello stesso processore le istruzioni da eseguire. Il processo

evolutivo a cui si fa riferimento nel primo capitolo di questa tesi è molto evidente

nell'emblematico passaggio dalla centralizzazione dell'elaborazione sulla CPU (“Central

Processing Unit”) alle tecniche di co-processing tra CPU e GPU (“Graphic Processing

Unit”) che permette netti aumenti delle prestazioni di computing grazie allo sfruttamento

della potenza di calcolo delle GPU. Questo lavoro di tesi, inoltre, presenta un'analisi

generale dei principali modelli di programmazione parallela definiti nel corso degli anni

all'interno del settore della ricerca informatica e volge particolare attenzione ad una delle

più recenti piattaforme di elaborazione in parallelo introdotte dalla NVIDIA Corporation e

denominata CUDA (“Compute Unified Device Architecture”) che sfrutta linguaggi di

programmazione ad alto livello ed aumenta esponenzialmente le performance di calcolo di

sistemi di svariata natura (scientifica, grafica, finanziaria, etc.). Nelle ultime pagine di

questo lavoro, infine, si evidenziano i vantaggi esecutivi effettivi e le potenzialità future

più importanti dell'applicazione del modello di programmazione parallela in confronto con

il tradizionale codice sequenziale attraverso un esempio di implementazione applicativa

basato sul linguaggio C.

4

Capitolo 1: L'evoluzione del modello architetturale

I modelli di programmazione si sono evoluti nel corso del tempo di pari passo con la

scienza informatica della quale fanno parte. Tali sviluppi si sono resi evidenti con le

mutazioni che i modelli architetturali, ovvero l'insieme di caratteristiche che determinano

le funzionalità dei vari componenti informatici, hanno subito. Tra tutti il processore è

sicuramente quello di maggiore interesse e, al fine di comprendere in maniera migliore

l'efficacia dei modelli, c'è bisogno di conoscere i passaggi evolutivi che hanno interessato

l'architettura di questo componente nevralgico nell'economia di un sistema informatico.

1.1 Dal sequenziale al parallelo

Fino ai primi anni 2000 il modello architetturale del processore è rimasto pressoché

invariato: si trattava, infatti, di un modello sequenziale nel quale un insieme di

componenti hardware era in grado di eseguire le istruzioni di un unico programma alla

volta. Tale modello architetturale si è evoluto grazie alla sempre maggiore densità di

componenti elettronici sul silicio di cui è composto, in concordanza con la legge di Moore,

formulata negli anni '60 da Gordon Moore e rispettata tuttora, che stabilisce che il numero

di componenti su una certa superficie di silicio raddoppia ogni due anni per via dei

miglioramenti delle tecnologie di produzione. Questi miglioramenti, però, non hanno

portato esclusivamente ad una costruzione di modelli architetturali che occupassero

porzioni di chip sempre maggiori, ma anche ad una tendenza alla miniaturizzazione dei

componenti e alla possibilità di lavorare a frequenze di clock sempre più elevate. Questo

andamento ha subito un brusco rallentamento proprio nei primi anni 2000 per due ragioni

fondamentali:

• non si è stati più in grado di trovare miglioramenti architetturali che portassero ad

un incremento delle prestazioni proporzionale all'area di silicio occupata per

introdurli;

• l'utilizzo di frequenze di clock sempre più elevate ha introdotto notevoli problemi

5

di consumo e di dissipazione termica.

A questo punto i produttori si sono trovati in una situazione che vedeva da una parte la

possibilità di realizzare componenti di calcolo sempre più complessi su un singolo chip e

dall'altra parte la difficoltà oggettiva a realizzarli all'interno del modello architetturale

sequenziale utilizzato fino a quel momento. Di fronte a questo criticismo si è delineata

l'idea di accantonare il modello sequenziale stesso: non si è cercato, quindi, di produrre

componenti che eseguissero le istruzioni di un singolo programma sempre più

velocemente, ma si sono sfruttate le potenzialità della tecnologia a disposizione per

produrre sullo stesso chip più core, ossia più componenti, che, operando in maniera

combinata e sfruttando un sottosistema di memoria comune, fossero in grado di eseguire

più istruzioni parallelamente.

1.2 Lo sviluppo tecnologico: il general purpose many core

Viene così segnato un nuovo solco particolarmente profondo nell'approccio al codice da

parte del programmatore: quest'ultimo, infatti, non tende più a concepire il proprio

programma come una serie di istruzioni eseguite sequenzialmente, ma deve pensare ad

esso come ad un insieme di attività da coordinare esplicitamente ed in grado di essere

eseguite in modo autonomo. Questa metodologia di programmazione diventa obbligatoria

nel momento in cui il progresso tecnologico di produzione dei chip porta all'eliminazione

dall'area di silicio di tutti quei componenti che introducevano miglioramenti minimi a

6

Figura 1.1: confronto tra modello sequenziale e paralello

fronte dello spazio occupato, al fine di far posto ad un numero maggiore di core sulla

superficie a disposizione. Si determina così la nascita dei cosiddetti general purpose

many core, che dispongono di decine o addirittura centinaia di core ottimizzati per

l'esecuzione parallela di più flussi di controllo. È importante notare che, eseguendo un

programma implementato appositamente per un processore single core su uno dei core di

un processore many core, si otterrebbero prestazioni leggermente inferiori. L'avanzamento

tecnologico non si ferma, però, solo a queste innovazioni architetturali, ma include anche

la diffusione di nuove componenti che si adattano perfettamente alle esigenze della

programmazione parallela. Queste novità sono:

• le GPU (Graphics Processing Unit), ovvero acceleratori in grado di processare

molto velocemente tutte le operazioni tipiche della grafica computazionale, che

utilizzano migliaia di componenti capaci di eseguire ciascuno semplici operazioni

su un singolo pixel dell'immagine;

• le FPGA (Field Programmable Gate Array), ovvero dispositivi programmabili via

software formati da migliaia di porte logiche che permettono l'implementazione di

funzioni logiche anche abbastanza complesse. La loro programmazione è realizzata

attraverso linguaggi descrittivi come VHDL o Verilog, i quali richiedono una

profonda conoscenza del dispositivo stesso.

Entrambe queste tecnologie erano inizialmente messe a disposizione come componenti

che interagivano con la CPU attraverso il bus del processore, ma successivamente ci si è

resi conto che, soprattutto in riferimento alle GPUs, potessero essere più efficaci se

programmate ad-hoc per eseguire particolari tipi di computazioni, anche non

esclusivamente grafiche. Nacquero così le cosiddette GP-GPU (General Purpose-GPU),

che tutt'oggi sono molto impiegate nei sistemi informatici più all'avanguardia.

7

Capitolo 2: I modelli per la programmazione parallela

Fin dagli anni '90 ci si è resi conto che molte applicazioni parallele in realtà sfruttavano

modelli computazionali profondamente simili tra loro. Si è reputato utile, quindi, definire

alcuni pattern, ossia schemi prestabiliti per la risoluzione di algoritmi specifici, che

organizzassero le modalità di creazione ed interazione delle varie azioni svolte

parallelamente e che ora andiamo brevemente ad analizzare.

2.1 Il modello pipeline

Il modello pipeline è basato su stadi che si susseguono e che hanno il compito di calcolare

risultati parziali che si avvicinano sempre più al risultato finale dell'operazione da

effettuare e che siano posti in ingresso allo stadio immediatamente successivo. In termini

di programmazione parallela, l'idea è di assegnare ciascuno stadio ad una unità di

elaborazione e provvedere ad implementare un meccanismo per il quale ogni stadio possa

interagire con lo stadio successivo attraverso un invio di dati. Questo modello è

particolarmente utilizzato nelle applicazioni di calcolo matematico e nel video processing

o comunque risulta molto valido se gli stadi della pipeline sono tutti computazionalmente

equivalenti, altrimenti stadi meno onerosi potrebbero dare vita ad effetti “a collo di

bottiglia”.

8

Figura 2.1: operazioni di un modello pipeline

2.2 Il modello master­slave

Nel modello master-slave un'attività, che prende il nome di master, distribuisce

l'esecuzione dei vari task ad un insieme di più attività, ciascuna denominata slave, che

restituiranno il risultato dell'elaborazione al master. Questo pattern è particolarmente

utilizzato nell'ambito di applicazioni che prevedono una serie di elaborazioni che possono

essere eseguite indipendentemente l'una dall'altra e in maniera parallela, come per esempio

il calcolo della moltiplicazione tra due matrici.

9

Figura 2.2: operazioni di un modello master-slave

Capitolo 3: Il mondo della GP-GPU

La GPU è particolarmente efficace nel calcolo di tutte quelle operazioni che richiedono

l'esecuzione di una singola funzione su tutti gli elementi di un vettore o di una matrice,

grazie all'elevata potenza di calcolo dovuta alla presenza di migliaia di componenti su un

unico chip. Ben presto questa unità ha abbandonato la semplice elaborazione grafica per

essere destinata ad ambiti più generali, evolvendosi in GP-GPU, ovvero un nuovo settore

della ricerca informatica che ha come scopo l'utilizzo del processore della scheda grafica

in modo da accelerare notevolmente l'esecuzione di calcoli data parallel, ovvero calcoli il

cui risultato dipende da un insieme di risultati parziali ottenuti attraverso computazioni

indipendenti eseguite in maniera altamente parallela su partizioni del dato in input. Prima

di analizzare nel dettaglio l'architettura di una GPU, può essere sicuramente utile chiarire

la distinzione tra le varie tipologie di architetture dei calcolatori esistenti. Tale

classificazione prende il nome di tassonomia di Flynn.

3.1 La tassonomia di Flynn

Nel 1966 Michael J. Flynn ha classificato i sistemi di calcolo a seconda della molteplicità

del flusso di istruzioni e del flusso dei dati che possono gestire. Nella tassonomia di Flynn

sono presenti quattro diverse categorie:

• SISD (Single Instruction, Single Data): in un sistema SISD (e.g., macchina di Von

Neumann) un singolo stream di istruzioni processa un singolo flusso di dati senza

alcun tipo di parallelismo;

• SIMD (Single Instruction, Multiple Data): un sistema SIMD (e.g., supercomputer

vettoriali) prevede che un singolo stream di istruzioni venga trasmesso

contemporaneamente a più unità di elaborazione, ognuna caratterizzata da un

proprio flusso di dati. Tale sistema è molto utilizzato per applicazioni con

parallelismo a grana fine, ossia gestito in maniera del tutto trasparente al

programmatore, o con semplici comunicazioni tra processi;

10

• MISD (Multiple Instruction, Single Data): un sistema MISD è basato su più flussi

di istruzioni che lavorano contemporaneamente su un singolo stream di dati. Tale

sistema non ha ancora trovato un'implementazione pratica;

• MIMD (Multiple Instruction, Multiple Data): in un sistema MIMD (e.g., cluster di

computer) più flussi di istruzioni vengono eseguiti contemporaneamente su più

flussi di dati. L'architettura della maggior parte dei più moderni sistemi paralleli

richiama questo schema e si è soliti suddividere questa categoria in sotto-categorie

discriminando l'architettura della memoria a disposizione.

3.2 L'architettura di una GPU

Ma perché la GPU ha sostituito la CPU come processore di riferimento nell'ambito della

programmazione parallela? La risposta a questa domanda è racchiusa all'interno

dell'architettura della GPU stessa: l'unità di elaborazione grafica, infatti, al contrario della

CPU, presenta un numero molto elevato di processori relativamente semplici che basano

la propria efficienza sull'esecuzione parallela di moltissimi thread, anziché sull'esecuzione

ottimizzata di un unico processo. Tale architettura è basata sulla struttura definita SIMT

(Single Instruction, Multiple Thread) che non è altro che un'estensione del concetto

precedentemente discusso di SIMD. Il funzionamento di una GPU è legato all'esecuzione

fisica in parallelo di blocchi di thread, chiamati warp, da parte di più SM (Streaming

Multiprocessor). Se ciascun SM deve eseguire più blocchi di thread, un componente

interno chiamato warp scheduler li partiziona in singoli warp di grandezza fissa in base

all'ID dei thread, iniziando dal thread con ID 0 e proseguendo consecutivamente. Ciascuno

11

Figura 3.1: tassonomia di Flynn

di questi SM racchiude al suo interno tre tipologie di unità di esecuzione:

• otto processori scalari (SP) che eseguono operazioni aritmetico-logiche in virgola

mobile e virgola fissa;

• due unità funzionali speciali (SFU) che eseguono varie funzioni matematiche oltre

alla moltiplicazione in virgola mobile;

• un'unità di doppia precisione (DPU) che esegue operazioni aritmetiche su operandi

in virgola mobile di 64 bit.

Gruppi di SM appartengono a Thread Processing Clusters (TPC) che contengono anche

altre risorse hardware tipiche delle applicazioni grafiche e in genere non visibili al

programmatore.

La GPU è costituita dall'insieme dei TPC, dalla rete di interconnessione e dai controllori

della memoria esterna.

12

Figura 3.2: visione di dettaglio dell'architettura di una GPU

Figura 3.3: visione globale dell'architettura di una GPU

3.3 Lo scontro generazionale: GP­GPU vs CPU

Dopo aver analizzato il processore grafico nel dettaglio a livello architetturale passiamo a

confrontarlo con la CPU a livello di performance. Facendo riferimento al grafico

sottostante, possiamo facilmente intuire che, in termini di Gflops/s, ossia il numero di

operazioni in virgola mobile eseguite al secondo, il divario tra CPUs e GPUs è in costante

crescita, sebbene entrambe le tecnologie siano segnate da un processo evolutivo in

costante ascesa.

Il divario tra le diverse architetture e, di conseguenza, tra le performance delle due

tipologie di processori è ancora più netto se andiamo a confrontare due modelli di

processore pressoché coetanei come l'Intel Core 2 Extreme QX9650 (CPU, 2009) e la

NVIDIA GeForce GTX 280 (GPU, 2008). La scheda sottostante ci mette in risalto come la

differenza architetturale basata sulla maggiore quantità ma minore complessità dei

componenti presenti sulla GPU sia decisiva anche a livello prestazionale. Fattori come il

minor numero di context switch, la minore latenza o i livelli di memoria cache a

disposizione hanno decretato la preferenza riguardo l'uso delle GPU come processori a

tutto tondo nell'ambito di una programmazione fortemente parallela.

13

Figura 3.4: confronto delle performance tra GPUs e CPUs

3.4 I modelli di programmazione su GP­GPU

Anche se, come abbiamo constatato nei paragrafi precedenti, i processori grafici sono

nettamente più efficienti delle CPU, l'approccio odierno è comunque un approccio ibrido e

congiunto tra le due unità di elaborazione principali. L'interazione tra le due componenti è

possibile attraverso lo sfruttamento di un'area di memoria condivisa che permette di

scambiare esclusivamente i riferimenti agli oggetti e non intere copie degli stessi, evitando

la latenza dovuta ai trasferimenti. Tutto questo è possibile attraverso la Application

Programming Interface (API) messa a disposizione dal modello di programmazione. Le

piattaforme più recenti ed efficienti su GP-GPU sono OpenCL e CUDA: il primo è un

framework multi-piattaforma basato sul linguaggio C o C++ portato a compimento dal

consorzio no-profit Khronos Group; il secondo è una piattaforma ideata e sviluppata

completamente da NVIDIA Corporation e destinata all'uso esclusivo su schede grafiche

prodotte dall'azienda statunitense stessa. Oltre ad offrire un nuovo modello architetturale

orientato all'alto tasso di parallelismo, CUDA permette di utilizzare una API e diversi

linguaggi di programmazione ad alto livello, per sfruttare il processore grafico a tutto

tondo.

14

Figura 3.5: confronto tra modelli specifici di GPU e CPU

Capitolo 4: CUDA

Nel novembre 2006 la NVIDIA Corporation ha presentato CUDA, una piattaforma di

calcolo parallelo general purpose e modello di programmazione che permette di

demandare la risoluzione di problemi computazionali di natura altamente parallela al

processore posto all'interno della GPU NVIDIA. Tale piattaforma fornisce una serie di

strumenti per facilitarne la programmazione, tra i quali un supporto software, scaricabile

all'interno del CUDA Toolkit dal sito web ufficiale della casa produttrice statunitense, che

comprende un ambiente di sviluppo denominato Nsight, il quale permette anche di

analizzare dettagliatamente il proprio programma attraverso un visual profiler, e un

compilatore apposito (nvcc) che permette di programmare sia direttamente a livello di

architettura, sia con linguaggi di programmazione ad alto livello. CUDA introduce un set

di istruzioni assembly a più alto livello, denominato Parallel Thread Execution (PTX), che

esula anche dallo specifico dispositivo utilizzato e semplifica il processo compilativo. In

questo capitolo, infine, verrà valutato un esempio di codice in CUDA C, ossia

un'estensione del classico linguaggio C con una API studiata appositamente da NVIDIA,

che verrà analizzato nei suoi aspetti più significativi e confrontato con un codice C

sequenziale.

4.1 La gestione dei thread in CUDA

Sebbene la piattaforma CUDA sia esclusivamente destinata a prodotti NVIDIA, le

architetture delle varie generazioni di schede grafiche cambiano costantemente, soprattutto

per quanto riguarda il numero di microprocessori operanti, rispettando l'andamento

pronosticato dalla legge di Moore. Per fronteggiare questo processo evolutivo senza troppi

sforzi per il compilatore, CUDA prevede la suddivisione del programma multithread in

una serie di blocchi indipendenti fra loro (blocks) che risiedono tutti sullo stesso

multiprocessore ma ognuno eseguito da un singolo SM e dotato di una porzione di

memoria propria. L'insieme di blocks è organizzato logicamente in griglie (grids) e

15

ciascun blocco contiene un numero limitato di thread a seconda delle dimensioni del

multiprocessore, ma comunque non superiore a 1024 thread.

Nei capitoli precedenti si è già discusso approfonditamente dell'architettura di una GPU

generica: in questo caso, ogni multiprocessore di una GPU CUDA-Enabled gestisce warp

composti da 32 thread che iniziano fisicamente l'esecuzione in parallelo. L'intero contesto

di esecuzione (program counters, registri, etc.) di ogni warp è mantenuto in memoria per

l'intero ciclo di vita del warp stesso, ottimizzando il numero e il costo dei context switch.

Ad ogni iterazione il warp scheduler seleziona il warp che ha thread pronti ad eseguire la

prossima istruzione tra i cosiddetti active threads. In particolare, ogni microprocessore ha

un set di registri a 32 bit da dividere tra i warp e una porzione di memoria condivisa

(shared memory) tra i blocchi di thread alla quale questi ultimi possono accedere venendo

sincronizzati esplicitamente dal programmatore.

4.2 L'organizzazione gerarchica della memoria in CUDA

L'elevata efficienza dell'architettura CUDA, però, non è dovuta esclusivamente alla

gestione ottimizzata dei thread. Le performance di un codice CUDA sono ulteriormente

migliorate dall'organizzazione gerarchica della memoria a disposizione, che fornisce una

notevole flessibilità di utilizzo e permette di ottimizzare la velocità di accesso ai dati ed il

16

Figura 4.1: esempio di gestione bidimensionale dei thread in CUDA

transfer rate tra memoria e processore, contribuendo anche a minimizzare l'effetto

bottleneck nei trasferimenti di dati. Su un'architettura CUDA, infatti, sono disponibili le

seguenti aree di memoria:

• local memory: memoria a rapido accesso locale ai singoli thread;

• registers: insieme di registri a rapidissimo accesso senza latenza locali ai singoli

thread;

• shared memory: memoria a rapido accesso con visibilità globale ai thread di un

blocco e locale a ciascun blocco e con durata pari. Viene principalmente utilizzata

per operazioni di lettura/scrittura comuni a più thread appartenenti ad un singolo

blocco e la sua dimensione è diversa a seconda della versione di compute capability

del dispositivo che esegue il programma, passando, infatti, dai 16 kB montati su

ogni core dei chip versione 1.x ai 96 kB equipaggiati da quelli versione 5.2;

• global memory: memoria con visibilità globale a tutti i thread dell'applicazione e

all'host, il quale ne gestisce allocazione e de-allocazione;

• constant memory e texture memory: memorie a sola lettura accessibili da tutti i

thread dell'applicazione.

17

Figura 4.2: interazione tra memorie eprocessori in CUDA

L'unica porzione di memoria che può essere acceduta e modificata sia dall'host che dal

device è la global memory. Questa viene utilizzata svariate volte all'interno di un

programma ed è, quindi, necessario minimizzare il bandwidth. Le performance ottimali si

ottengono quando la GPU, nell'esecuzione della funzione kernel, appronta il cosiddetto

coalescing degli accessi alla memoria globale: gli accessi da parte di un thread di un warp

nella memoria globale sono agglomerati dal device nel numero minore possibile di

transazioni, permettendo una esecuzione più efficiente e fluida. Le tecniche ed i requisiti

per gestire il coalescing variano a seconda della compute capability dell'architettura in

uso: per le GPUs precedenti la compute capability 2.0, infatti, le transazioni sono unite per

un numero di thread pari alla metà di un warp anziché pari alla totalità dell'insieme.

4.3 La struttura di un codice CUDA

Il modello CUDA prevede l'utilizzo della CPU, identificata come host, coadiuvato dalla

GPU, identificata come device. Un codice CUDA alterna porzioni di codice sequenziale,

eseguite dall'host, a porzioni di codice parallelo, eseguite dal device. La sequenza di

istruzioni eseguite dalla CPU segue fondamentalmente tre passi:

I. caricamento dei dati nella GPU;

II. chiamata della funzione kernel;

III. recupero dei risultati dell'elaborazione della GPU.

La funzione kernel rappresenta il modo di implementare l'esecuzione parallela di una

porzione di codice da parte introdotto da CUDA: questa, infatti, permette di eseguire una

funzione più volte in parallelo su più thread, al contrario della singola esecuzione a cui fa

riferimento una tradizionale funzione C. La GPU esegue un solo kernel alla volta. Per

operare sui singoli thread o sui singoli blocchi, CUDA mette a disposizione degli

identificatori, la cui natura è approfondita nel paragrafo successivo. Ogni kernel è

suddiviso in grids di blocks, ma è l'hardware a gestire in maniera totalmente autonoma la

distribuzione dei blocchi su ogni multiprocessore. Così facendo, CUDA si conferma come

una piattaforma efficiente, altamente scalare e per lo più trasparente al programmatore per

quanto concerne i passaggi generalmente più delicati di un'applicazione parallela.

18

4.4 Implementazione di CUDA in C

CUDA C estende il linguaggio ad alto livello C e permette, ovviamente, la gestione di

porzioni di codice destinate alla CPU congiuntamente a porzioni di codice destinate al

coprocessore GPU. L'estensione CUDA C introduce alcuni qualificatori del tipo della

funzione che vanno a completare la firma tradizionale del linguaggio C, quali:

• __device__: si inserisce se un metodo deve essere eseguito e richiamato dalla GPU;

• __host__: si inserisce se un metodo deve essere eseguito e richiamato dalla CPU;

• __global__: si inserisce per identificare un kernel, metodo chiamato dall'host e

demandato alla GPU. Deve avere come tipo di ritorno un parametro void.

Per quanto concerne i tipi e le variabili disponibili, sono stati introdotte diverse novità,

quali:

• __device__: qualificatore di tipo che definisce una variabile che risiede nella

memoria globale, che ha ciclo di vita pari all'applicazione;

• __constant__: qualificatore di tipo che definisce una variabile che risiede nella

constant memory, che ha ciclo di vita pari all'applicazione;

19

Figura 4.3: interazione tra host e device in un codice CUDA

• __shared__: qualificatore di tipo che definisce una variabile che risiede nella

memoria condivisa di un blocco di thread;

• dim3: struttura che identifica una tripla di interi basato su un uint3 (intero senza

segno tridimensionale), il cui valore di default è 1, usato per specificare le

dimensioni di una griglia o un blocco;

• gridDim: variabile di tipo dim3 che contiene le dimensioni di una griglia;

• blockIdx: variabile di tipo uint3 che contiene l'identificativo del singolo blocco

all'interno della griglia;

• blockDim: variabile di tipo dim3 che contiene le dimensioni del blocco;

• threadIdx: variabile di tipo uint3 che contiene l'identificativo del singolo thread

all'interno del blocco e permette l'accesso alle singole dimensioni tramite i suffissi

.x, .y e .z;

• warpSize: variabile di tipo int che contiene il numero di thread all'interno del

blocco;

• cudaError_t: tipo di errore introdotto da CUDA C.

Inoltre CUDA C estende i singoli tipi primitivi del tradizionale C includendo anche tipi

multidimensionali fino a tre dimensioni (e.g., uint3).

La sintassi di invocazione dell'esecuzione di un kernel è totalmente nuova ed è la

seguente:

__global__ void KernelFuncion(int a, int b, int c){

…}

int main(){

…// Invocazione di un kernel su un unico blocco// di N*N*1 threadsint numBlocks = 1;dim3 threadsPerBlock(N, N);KernelFunction<<<numBlocks, threadsPerBlock>>>(a, b, c);…

}

20

Conseguentemente alle modifiche architetturali approntate dalla piattaforma CUDA, l'API

CUDA C include anche nuovi metodi principalmente volti alla gestione della memoria e

alla sincronizzazione, quali:

• cudaMalloc(void** pointer, size_t nbytes): alloca un oggetto all'interno della

memoria globale. Riceve come parametri di input l'indirizzo del puntatore

all'oggetto allocato e la dimensione dello spazio da allocare;

• cudaFree(void* pointer): libera lo spazio di memoria precedentemente allocato.

Riceve in ingresso il puntatore all'oggetto da rimuovere;

• cudaGetErrorString(cudaError_t err): ritorna i dettagli dell'errore passatogli in

ingresso;

• cudaMemcpy(void* dst, void* src, size_t nbytes, cudaMemcpyKind direction):

gestisce il trasferimento dall'host al device e viceversa. Riceve in ingresso il

puntatore alla destinazione, il puntatore alla sorgente, il numero di bytes da

trasferire e un tipo di trasferimento. Quest'ultimo è solitamente una costante

simbolica che può assumere un valore tra cudaMemcpyHostToDevice,

cudaMemcpyDeviceToHost oppure cudaMemcpyDeviceToDevice;

• cudaThreadSynchronize(): blocca l'esecuzione fino al momento in cui il device ha

completato tutti i task richiesti. Se uno dei task fallisce, restituisce un errore;

• __syncthreads(): utilizzato per coordinare la comunicazione tra thread dello stesso

blocco, funziona come una barriera alla quale tutti i thread all'interno del blocco

devono attendere fintantoché ognuno sia autorizzato a procedere.

4.5 Esempio di codice CUDA C

Il modo migliore per evidenziare le capacità del linguaggio CUDA C è sicuramente

attraverso l'analisi di un esempio di programma che fornisca una panoramica completa ed

intuitiva sulle funzionalità della piattaforma: il seguente codice, ottenuto a partire da

quello riportato come esempio nel paragrafo 3.2.3 della CUDA C Programming Guide di

NVIDIA Corporation, calcola il prodotto tra due matrici A e B, rispettivamente di

dimensioni 80x144 e 144x96, entrambe contenenti numeri interi compresi tra 0 e 5, e ne

21

riporta il risultato in una terza matrice C di dimensioni 80x96. Il listato in CUDA C

implementato di seguito prevede l'utilizzo della memoria condivisa e stampa a video la

durata dell'esecuzione della funzione kernel.

4.5.1 provaCUDA.cu

#include <stdio.h>#include <stdlib.h>#include <cuda.h>

// Grandezza del singolo blocco di thread#define BLOCK_SIZE 16

cudaEvent_t start, stop;

typedef struct {// Numero di righe e colonneint width;int height; // Passo che coinciderà con il numero di colonne per// un rapido accesso agli elementi ordinati nell'arrayint stride;// Puntatore all'array contenente gli elementiint* elements;

} Matrix;

// Ritorna un elemento di una matrice in // posizione(row, col)__device__ int GetElement(const Matrix A, int row, int col) {

return (A.elements[row * A.stride + col]); }

// Set di un elemento in una matrice in posizione (row, col)__device__ void SetElement(Matrix A, int row, int col,

int value) {

A.elements[row * A.stride + col] = value; }

// Ritorna la sotto-matrice Asub di dimensioni// BLOCK_SIZExBLOCK_SIZE di A che è posizionata a col // sotto-matrici di distanza a destra e a row sotto-matrici// di distanza in basso rispetto al primo elemento in alto// a sinistra__device__ Matrix GetSubMatrix(const Matrix A, int row,

int col) {

Matrix Asub;

22

Asub.width = BLOCK_SIZE; Asub.height = BLOCK_SIZE; Asub.stride = A.stride; Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row

+ BLOCK_SIZE * col]; return Asub;

}

// Implementazione funzione kernel __global__ void MatMulKernel(const Matrix A, const Matrix B,

Matrix C){

// Assegnazione riga e colonna su cui costruire la// sotto-matrice corrispondente al blocco di thread int blockRow = blockIdx.y; int blockCol = blockIdx.x;

// Ogni blocco opera su una sotto-matrice Csub di CMatrix Csub = GetSubMatrix(C, blockRow, blockCol); // Ogni thread calcola un elemento di Csub// salvandolo in Cvalueint Cvalue = 0; // Assegnazione riga e colonna corrispondente al threadint row = threadIdx.y; int col = threadIdx.x;

// Scorri tutte le sotto-matrici di A e B necessarie per// calcolare Csub. Moltiplica ogni coppia di // sotto-matrici e somma il risultatofor (int m = 0; m < (A.width / BLOCK_SIZE); ++m) {

// Estrazione della sotto-matrice Asub of A Matrix Asub = GetSubMatrix(A, blockRow, m);

// Estrazione della sotto-matrice Bsub of B Matrix Bsub = GetSubMatrix(B, m, blockCol);

// Assegna spazio in shared memory per Asub e Bsub__shared__ int As[BLOCK_SIZE][BLOCK_SIZE]; __shared__ int Bs[BLOCK_SIZE][BLOCK_SIZE];

// Carica Asub e Bsub da device a shared memory// Ogni thread carica un singolo elementoAs[row][col] = GetElement(Asub, row, col); Bs[row][col] = GetElement(Bsub, row, col); // Sincronizzazione delle operazioni per essere // sicuri che le sotto-matrici siano state caricate// prima delle operazioni __syncthreads();

// Moltiplica Asub e Bsubfor (int e = 0; e < BLOCK_SIZE; ++e)

Cvalue += As[row][e] * Bs[e][col];

23

// Sincronizzazione necessaria per essere certi che// le operazioni precedenti siano state completate // prima di caricare le prossime sotto-matrici__syncthreads();

}

// Scrivi Csub sulla device memory// Ogni thread scrive un elementoSetElement(Csub, row, col, Cvalue);

}

// Le dimensioni delle matrici sono multiple di BLOCK_SIZEfloat MatMul(const Matrix A, const Matrix B, Matrix C) {

// Creazione checkpoint per il calcolo del// tempo di esecuzione del kernelcudaEventCreate(&start);cudaEventCreate(&stop);

// Definizione e caricamento delle matrici A e B// in device memoryMatrix d_A; d_A.width = d_A.stride = A.width;d_A.height = A.height; size_t size = A.width * A.height * sizeof(int); cudaMalloc(&d_A.elements, size); cudaMemcpy(d_A.elements, A.elements, size,

cudaMemcpyHostToDevice);

Matrix d_B; d_B.width = d_B.stride = B.width;d_B.height = B.height; size = B.width * B.height * sizeof(int);cudaMalloc(&d_B.elements, size); cudaMemcpy(d_B.elements, B.elements, size,

cudaMemcpyHostToDevice);

// Definizione e allocazione della matrice C in device// memory Matrix d_C; d_C.width = d_C.stride = C.width;d_C.height = C.height; size = C.width * C.height * sizeof(int); cudaMalloc(&d_C.elements, size);

// Invocazione funzione kerneldim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); dim3 dimGrid((B.width / dimBlock.x),

(A.height / dimBlock.y));cudaEventRecord(start, 0);MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);

24

cudaThreadSynchronize();cudaEventRecord(stop, 0);cudaEventSynchronize(stop);

// Lettura della matrice C dalla device memory cudaMemcpy(C.elements, d_C.elements, size,

cudaMemcpyDeviceToHost);

// De-allocazione della device memory cudaFree(d_A.elements); cudaFree(d_B.elements); cudaFree(d_C.elements);

float el_time = 0.0;cudaEventElapsedTime(&el_time, start, stop);

cudaEventDestroy(start);cudaEventDestroy(stop);

return el_time;}

int main(int argc, char* argv[]){Matrix A, B, C;int a1, a2, b1, b2;float duration;srand(time(NULL));

// Numero di righe di Aa1 = 80;// Numero di colonne di Aa2 = 144;// Numero di righe di Bb1 = a2;// Numero di colonne di Bb2 = 96;

A.height = a1;A.width = a2;// Alloca spazio necessario nella memoria dell'hostA.elements = (int*)malloc(A.width * A.height

* sizeof(int));

B.height = b1;B.width = b2;// Alloca spazio necessario nella memoria dell'hostB.elements = (int*)malloc(B.width * B.height

* sizeof(int));

C.height = A.height;C.width = B.width;// Alloca spazio necessario nella memoria dell'host

25

C.elements = (int*)malloc(C.width * C.height * sizeof(int));

// Riempimento matrici A e Bfor(int i = 0; i < A.height; i++)

for(int j = 0; j < A.width; j++)A.elements[i * A.width + j] = rand() %6);

for(int i = 0; i < B.height; i++)for(int j = 0; j < B.width; j++)

B.elements[i * B.width + j] = rand() %6;

// Lancio funzione di prodottoduration = MatMul(A, B, C);

// Stampa a video le matricifor(int i = 0; i < A.height; i++){

for(int j = 0; j < A.width; j++)printf("%f ", A.elements[i * A.width + j]);

printf("\n");}printf("\n");

for(int i = 0; i < B.height; i++){for(int j = 0; j < B.width; j++)

printf("%f ", B.elements[i * B.width + j]);printf("\n");

}printf("\n");

for(int i = 0; i < C.height; i++){for(int j = 0; j < C.width; j++)

printf("%f ", C.elements[i * C.width + j]);printf("\n");

}printf("\n");

printf("\nDurata calcolo: %f millisecondi", duration);

return 0;}

4.5.2 Analisi del codice

Come detto all'inizio del paragrafo, il codice precedentemente riportato calcola il prodotto

di due matrici A e B, salvandone il risultato su una terza matrice C. Le matrici trattate sono

variabili il cui tipo Matrix è definito all'interno del codice stesso come una struttura dati

contenente le seguenti variabili:

26

• width, un intero che rappresenta il numero di colonne della matrice;

• height, un intero che rappresenta il numero di righe della matrice;

• stride, un intero che rappresenta il passo con cui scorrere l'array in cui sono salvati

gli elementi della matrice stessa per minimizzare il bandwidth;

• elements, un puntatore agli elementi contenuti nella matrice che sono salvati in un

array monodimensionale e disposti per numero di riga crescente.

Il filo conduttore di questo modello di programmazione parallela in CUDA C è quello di

implementare il prodotto matriciale delegando il calcolo di ogni elemento di una sotto-

matrice quadrata Csub della matrice C ad un singolo thread di un blocco. Csub è ottenuta dal

prodotto di due matrici rettangolari: una sotto-matrice di A di dimensioni (A.width,

BLOCK_SIZE), che ha lo stesso numero di righe di Csub, e una sotto-matrice di B di

dimensioni (BLOCK_SIZE, B.height), che ha lo stesso numero di colonne di Csub. Per

adattarsi alle dimensioni del blocco di thread del device, le due sotto-matrici vengono

suddivise in tante matrici quadrate di dimensione BLOCK_SIZE quante necessario

affinché non ci siano elementi in comune tra loro e Csub viene calcolata come la somma dei

prodotti delle matrici quadrate costruite. Ognuno di questi prodotti è ottenuto caricando le

matrici quadrate utilizzate dalla memoria globale a quella condivisa, utilizzando ogni

thread per calcolare un singolo elemento della matrice prodotto. Ciascun thread somma il

proprio risultato a quello degli altri e lo salva in una variabile di appoggio comune

(Cvalue) che contiene l'elemento da calcolare. Al termine dell'esecuzione di tutti i thread,

il risultato definitivo viene scritto in memoria globale nella posizione che rappresenta

all'interno della matrice Csub.

Il codice è ottimizzato per l'utilizzo della memoria condivisa da parte di host e device, al

fine di non dover accedere in memoria un numero eccessivo di volte rispetto a quante sono

realmente necessarie. Idealmente, il contesto migliore su cui operare sarebbe quello di

accedere una sola volta in memoria, avendo precedentemente caricato entrambe le intere

matrici A e B necessarie per il calcolo. Se, però, operassimo su matrici di dimensioni

cospicue, questo metodo comporterebbe un calo di efficienza nell'uso della shared

memory dato che quest'ultima è di grandezza limitata e potrebbe non riuscire ad ospitare la

27

totalità dei dati. Il codice mostrato precedentemente implementa una modalità di

sfruttamento di questa risorsa che rappresenta un giusto compromesso tra praticità d'uso e

gestione dello spazio disponibile: attraverso la suddivisione in blocchi, infatti, si sfrutta la

velocità di accesso alla memoria condivisa e si ottimizza l'uso della memoria globale, dal

momento che la matrice A viene letta esclusivamente (B.width / BLOCK_SIZE) volte e la

matrice B solo (A.height / BLOCK_SIZE) volte.

L'utilizzo della shared memory permette di minimizzare il bandwidth, sfruttando anche la

tecnica di coalescing della memoria globale attraverso un accesso regolato dalla variabile

membro stride della struttura dati Matrix. L'applicazione di questa tecnica per il

trasferimento di più dati dalla memoria globale nel numero minimo di transazioni permette

di sfruttare una memoria, quale è quella condivisa, con un bandwidth molto alto, una

latenza molto bassa e che non soffre la gestione di vettori multidimensionali. Gli elementi

delle sotto-matrici, infatti, sono prelevati dalla memoria globale, nella quale risiedono

come array monodimensionali ordinati per righe, e immagazzinati nella shared memory,

dove tutti i thread in esecuzione possono estrapolarne i dati necessari alle proprie

operazioni con trasferimenti nettamente più veloci.

28Figura 4.4: prodotto matriciale con memoria condivisa

4.5.3 Esecuzione su Nsight

Non resta che implementare il listato precedentemente riportato all'interno di Nsight,

l'ambiente di programmazione incluso nel CUDA Toolkit che permette di interagire con la

piattaforma CUDA. Per l'esecuzione di questo codice si è sfruttata una scheda video

NVIDIA GeForce GT 330M, dotata di architettura Tesla 2.0, con compute capability 1.2

montata su una macchina fornita di sistema operativo Linux Ubuntu 14.04 LTS a 64 bit.

Per interagire con la GPU è stato scaricato ed installato il CUDA Toolkit versione 6.5,

release massima per l'architettura a disposizione, che include il compilatore nvcc release

5.5 e l'estensione Nsight per l'ambiente di sviluppo Eclipse.

L'output ottenuto ha evidenziato che l'esecuzione della sola porzione di codice che include

la funzione kernel, la quale rappresenta il fulcro del calcolo parallelo adottato

nell'esempio, ha impiegato circa 6 millisecondi.

4.5.4 Profiling

CUDA mette a disposizione il software NVIDIA Visual Profiler (nvvp) che fornisce gli

strumenti di profiling per visualizzare ed analizzare le performance dell'applicazione,

senza distinzione riguardo il linguaggio di programmazione adottato. È possibile utilizzare

anche un applicativo, incluso sempre nel toolkit, per l'analisi direttamente da terminale

29

Figura 4.5: output dell'esecuzione di provaCUDA.cu

(nvprof), previo precedente spostamento all'interno della directory contenente il binario

dell'applicazione da analizzare. Sebbene l'esempio riportato sia abbastanza basilare e

presenti una sola funzione kernel che comporta un utilizzo della GPU per meno del 20%

del tempo complessivo di esecuzione del programma, è possibile notare alcune

ottimizzazioni nel codice evidenti dai dati estrapolati attraverso il profiler.

Lanciando l'applicazione da terminale attraverso la riga di comando nvprof ./myApp

si può ottenere un breve riepilogo delle funzioni kernel e delle operazioni sulla memoria

del device.

Una vista più dettagliata sull'esecuzione del kernel e sulle operazioni in memoria possono

essere reperite lanciando la riga di comando nvprof --print-gpu-trace

./myApp.

I dati più interessanti e significativi nella fase di profiling dell'applicazione sono

sicuramente gli events e le metrics. Un event è un'attività, un'azione o un'occorrenza

numerabile sul device che corrisponde ad un singolo valore accumulato durante

l'esecuzione del kernel. Una metric è una caratteristica dell'applicazione che è calcolata a

partire da uno o più valori di events del kernel. Per tracciare events e metrics relative ad

una determinata applicazione possiamo lanciare da terminare i comandi nvprof

--events all ./myApp e nvprof --metrics all ./myApp.

30

Figura 4.6: nvprof ./provaCUDA

Figura 4.7: nvprof --print-gpu-trace ./provaCUDA

Gli events più significativi sono quelli legati alle operazioni di lettura/scrittura in memoria

globale. Dalla figura sovrastante è desumibile, attraverso l'osservazione dei valori relativi

al gld_incoherent e gst_incoherent che stabiliscono il numero di operazioni

rispettivamente di load e store in global memory e che risultano in questo caso nulli, che il

kernel implementato nell'applicazione di esempio assicura il coalescing della totalità delle

operazioni in memoria. Questo è sicuramente un vantaggio ed un dato indicativo della

bontà del codice parallelo implementato, permettendo di ottimizzare l'uso delle risorse che

la GPU offre.

Per quanto riguarda le metrics ottenute dall'esecuzione del kernel, si può definire più che

soddisfacente il throughput raggiunto per le operazioni in memoria in relazione al device a

disposizione. Notevole è sicuramente la percentuale di branch_efficiency

raggiunta: tutti i salti all'interno del codice sono non divergenti.

31

Figura 4.8: nvprof --events all ./provaCUDA

Figura 4.9: nvprof --metrics all ./provaCUDA

4.5.5 Confronto con codice sequenziale di funzionalità equivalente

Ma quanto è remunerativa l'applicazione di un modello di programmazione parallelo

rispetto al tradizionale modello sequenziale? Ovviamente non può essere data una risposta

di carattere generale, poiché i vantaggi dell'esecuzione parallela rispetto a quella

sequenziale non sono assoluti. È possibile, per esempio, valutare se la scelta di far

intervenire la GPU nell'esecuzione del programma riportato nel paragrafo precedente,

assegnandole mansioni di calcolo parallelo, sia stata una scelta efficiente. L'indice più

attendibile e sicuramente immediatamente più indicativo è quello relativo al tempo di

esecuzione del programma: se, eseguendo il codice parallelo e quello sequenziale in

successione, si ottiene un tempo di esecuzione minore per l'implementazione con CUDA,

allora si può certamente identificare la scelta del modello parallelo come la più corretta per

tale applicazione, almeno dal punto di vista temporale.

Un esempio di codice sequenziale in C da me appositamente implementato per effettuare il

confronto può essere il seguente (matrixMul.c):

#include <stdio.h> #include <stdlib.h> #include <time.h>

int main() {

srand(time(NULL)); int a1, a2, b1, b2;

int** A; int** B; int** C;

int i, j, k; float TempoInMillisecondi; double TempoInSecondi; clock_t c_start, c_end;

// Numero di righe di A a1 = 80; // Numero di colonne di A a2 = 144; // Numero di righe di B b1 = a2; // Numero di colonne di B b2 = 96;

// Alloca spazio necessario nella memoria A = (int**)malloc(a1 * sizeof(int*));

32

for(i = 0; i < a1; i++) A[i] = (int*)malloc(a2 * sizeof(int));

B = (int**)malloc(b1 * sizeof(int*)); for(i = 0; i < b1; i++)

B[i] = (int*)malloc(b2 * sizeof(int));

C = (int**)malloc(a1 * sizeof(int*)); for(i = 0; i < a1; i++)

C[i] = (int*)malloc(b2 * sizeof(int));

// Riempimento matrici A e B for(i = 0; i < a1; i++)

for(j = 0; j < a2; j++) A[i][j] = rand() %6;

for(i = 0; i < b1; i++) for(j = 0; j < b2; j++)

B[i][j] = rand() %6;

c_start = clock();

// Calcolo prodotto matriciale for(i = 0; i < a1; i++){

for(j = 0; j < b2; j++){ C[i][j] = 0; for(k = 0; k < a2; k++){

C[i][j] += A[i][k] * B[k][j]; }

} }

c_end = clock();

// Stampa a video le matrici for(i = 0; i < a1; i++){

for(j = 0; j < a2; j++) printf("%d ", A[i][j]);

printf("\n"); } printf("\n");

for(i = 0; i < b1; i++){ for(j = 0; j < b2; j++)

printf("%d ", B[i][j]); printf("\n");

} printf("\n");

for(i = 0; i < a1; i++){ for(j = 0; j < b2; j++)

printf("%d ", C[i][j]);

33

printf("\n"); } printf("\n"); TempoInSecondi = (c_end - c_start) / (double)

CLOCKS_PER_SEC; TempoInMillisecondi = TempoInSecondi * 1000; printf("\nDurata calcolo: %f millisecondi\n", TempoInMillisecondi);

// De-allocazione delle matrici for(i = 0; i < a1; i++) free(C[i]); free(C);

for(i = 0; i < b1; i++) free(B[i]); free(B);

for(i = 0; i < a1; i++) free(A[i]); free(A);

return 0; }

Dall'output riportato nella figura sottostante è facile verificare che il tempo di esecuzione

della porzione di codice sequenziale relativa al solo calcolo del prodotto matriciale ha

impiegato circa 16 millisecondi di tempo, ossia più del doppio del tempo di esecuzione

relativo alla stessa funzionalità implementata attraverso il modello parallelo precedente.

34

Figura 4.10: output dell'esecuzione di matrixMul.c in Nsight

Conclusioni

Così come la programmazione orientata agli oggetti ha fatto con il metodo non strutturale

in voga fino agli anni '80, si può senza dubbio affermare che la programmazione parallela

sta rivoluzionando il modo di concepire e pensare le applicazioni. Questo processo

innovativo, a mio avviso, è di portata maggiore rispetto alle novità annunciate come

rivoluzionarie per il settore informatico negli anni passati, dato che la programmazione

parallela apre ad innumerevoli nuove opportunità di ricerca nei più svariati ambiti di

interesse, non esclusivamente scientifici. Il GPU-Computing, in particolare, è il fulcro di

svariati progetti che hanno segnato un deciso progresso nei rispettivi settori: dal progetto

COSMO che studia e prevede i cambiamenti climatici globali al progetto AMBER14 che

simula i movimenti molecolari delle biomolecole, dall'applicazione Smilart Platform che

riconosce in tempo reale i volti delle persone in qualsiasi scenario a quella PowerGrid che

applica la trasformata discreta di Fourier nell'ambito del medical imaging. Inoltre, se

all'inizio l'approccio parallelo veniva accantonato nei programmi a causa della gestione

troppo elaborata e dettagliata degli strumenti necessari a renderlo efficiente, ora sta

prendendo sempre più piede tra gli addetti ai lavori grazie ai notevoli sviluppi

architetturali dei processori, di cui si è trattato approfonditamente nel presente elaborato.

In questo senso, la piattaforma CUDA rappresenta il primo esempio di applicazione del

modello parallelo a tutti gli aspetti della computazione, dall'architettura del processore

all'esecuzione del codice, senza impegnare il programmatore con complesse operazioni di

sincronizzazione o di gestione della memoria.

La parte nevralgica e, dal mio punto di vista, più interessante e coinvolgente di questo

lavoro di tesi è stata quella di analizzare in prima persona le grandi potenzialità di

applicazione di CUDA, permettendomi di allargare le mie conoscenze grazie ad una prima

completa esperienza di programmazione parallela. Partendo da un codice di esempio della

CUDA C Programming Guide, ho sviluppato ed analizzato un'intera applicazione che

racchiudesse tutte le caratteristiche per cui la piattaforma CUDA è considerata altamente

35

innovativa. I vantaggi dell'esecuzione parallela si sono rivelati palesemente evidenti nel

confronto con il rispettivo modello sequenziale riportato nell'ultimo capitolo di questo

elaborato. L'esecuzione del primo codice, infatti, ha richiesto meno della metà del tempo

di esecuzione delle stesse funzionalità eseguite in maniera sequenziale dal secondo codice.

Inoltre, bisogna ricordare che i benefici della programmazione parallela, sia in termini di

tempo di esecuzione che di sfruttamento delle risorse messe a disposizione

dall'architettura, sono tanto maggiori quanto è più grande la mole di dati indipendenti da

trattare.

Alla luce della trattazione complessiva emersa da questa tesi, la piattaforma CUDA e i

modelli di programmazione per il calcolo parallelo in generale possono essere considerati,

a ragion veduta, il presente ed il futuro della tecnologia applicata ai sistemi informatici di

tutto il mondo, dato che il progresso tecnologico richiede sempre maggiore capacità

computazionale in tempo reale e sempre maggiore velocità di esecuzione in tutti gli ambiti

di applicazione.

36

Sviluppi Futuri

Nonostante le tante innovazioni già operative ed analizzate all'interno di questo elaborato,

i modelli di programmazione per il calcolo parallelo sono destinati a svilupparsi e a

sviluppare il settore informatico in maniera ancora più decisa e rivoluzionaria già nel

prossimo futuro. In particolare, la NVIDIA Corporation sta introducendo nuove migliorie

alla propria piattaforma CUDA con il lancio di nuove generazioni di supporti video. Dopo

la Tesla, infatti, l'ultima novità architetturale dell'azienda statunitense, che porta il nome

del fisico italiano Fermi, presenta migliorie sotto l'aspetto delle operazioni in virgola

mobile a doppia precisione, del supporto ECC, della shared memory e della velocità dei

context switch e delle operazioni atomiche. Tutto questo grazie a 512 CUDA core divisi in

16 SM, porzioni di memorie condivise configurabili, un livello di cache aggiuntivo e un

doppio warp scheduler. Il supporto Error Correcting Code (ECC) assicura, inoltre,

l'integrità dei dati in memoria, e risulta utile soprattutto durante le esecuzioni di sistemi

che trattano dati sensibili, come in ambito medico. La tecnologia ECC rileva e corregge

errori su singoli bit dovuti a naturali radiazioni accidentali prima che danneggino il

sistema e procurino errori indesiderati e indelebili.

La sempre maggiore affidabilità ed efficienza dei sistemi paralleli consente di segnalare

questi modelli come i più avveniristici ed interessanti per il futuro del settore informatico.

37

Figura 19: evoluzione dell'architettura di una GPU dal 2006 ad oggi

Bibliografia

[1] M. Danelutto, Programmazione parallela: evoluzione e nuove sfide, Mondo Digitale, Aprile2015

[2] NVIDIA Corporation, CUDA C Programming Guide, PG-02829-001_v7.5, Settembre 2015[3] T. G. Mattson, B. A. Sanders, B. L. Massingill, Patterns for parallel programming, Addison

Wesley, 2013[4] NVIDIA Corporation, http://www.nvidia.it/object/gpu-computing-it.html, 7/6/2016[5] NVIDIA Corporation, http://www.nvidia.it/content/EMEAI/images/tesla/tesla-server-

gpus/nvidia-gpu-application-catalog-eu.pdf, 10/6/2016[6] NVIDIA Corporation,

http://www.nvidia.it/content/PDF/fermi_white_papers/NVIDIA_Fermi_Compute_Architecture_Whitepaper.pdf, 10/6/2016

[7] NVIDIA Corporation, http://www.nvidia.it/object/cuda-parallel-computing-it.html, 9/6/2016[8] NVIDIA Corporation, https://devblogs.nvidia.com/parallelforall/how-access-global-memory-

efficiently-cuda-c-kernels/, 13/6/2016[9] NVIDIA GeForce GT330M, http://www.geforce.com/hardware/notebook-gpus/geforce-gt-

330m/specifications, 13/6/2016[10] NVIDIA CUDA Toolkit 6.5,

http://developer.download.nvidia.com/compute/cuda/6_5/rel/docs/CUDA_Toolkit_Release_Notes.pdf, 13/6/2016

[11] CUDA Pro Tip: nvprof is Your Handy Universal GPU Profiler, https://devblogs.nvidia.com/parallelforall/cuda-pro-tip-nvprof-your-handy-universal-gpu-profiler/, 18/7/2016