In questo tutorial vedremo come sfruttare l'architettura Compute Unified Device Architecture (CUDA) e il linguaggio C per CUDA di NVIDIA per realizzare programmi paralleli di computazione generale in grado di girare sui chip grafici.
CUDA è, insieme, un modello architetturale, una Application Programming Interface (API) per sfruttare tale architettura e una serie di estensioni (C for CUDA) al linguaggio C per descrivere un'applicazione parallela in grado di girare sulle GPU che adottano quel modello.
Il compilatore di CUDA è basato sulla piattaforma Open64, un progetto inizialmente sviluppato da SGI e integrato con una serie di contributi dai gruppi di ricerca nell'ambito dei compilatori, che l'hanno reso uno dei più interessanti ambienti per la valutazione di nuove idee nei compilatori. Di fatto, data la licenza GPL 2.0, anche NVIDIA è tenuta a rilasciare i sorgenti del suo compilatore al pubblico, cosa che difatti avviene in un apposito sito FTP, benché la cosa non sia ai più nota. Ciò porta anche alla considerazione che, effettivamente, CUDA non avrebbe vincoli al supporto di piattaforme non di NVIDIA. D'altra parte, però, con gli standard oggi emergenti, OpenCL e DirectX Compute, i concorrenti di NVIDIA sembrano meno che mai interessati a sfruttare CUDA. È, però, innegabile che al momento CUDA sia la piattaforma più sviluppata e matura per poter progettare i propri programmi di GPGPU computing, disponibile ormai da molto tempo per lo scaricamento gratuito da parte di tutti gli interessanti, e con una grande mole di chip grafici prodotti in grado di supportarlo.
Nello specifico, qualsiasi GPU NVIDIA a partire dalla Generazione GeForce 8 (G80) può da subito eseguire applicazioni scritte in CUDA. Il discorso vale ovviamente per le schede con chip grafici di generazione equivalente della famiglia Quadro (dedicate principalmente al video e alla grafica industriale) e della famiglia Tesla (specificatamente indirizzate al GPGPU computing, non avendo neanche l'uscita video). Vedremo nel prosieguo del tutorial cosa cambia nelle caratteristiche supportate dalle varie versioni dell'hardware fino all'attuale serie di schede basata su GT200 (GeForce GTX).
Le motivazioni sull'uso di una ambiente di GPGPU per eseguire certe classi di applicazione sono motivate dall'immagine seguente.
Come si può vedere, la potenza computazionale di queste architetture sta crescendo in maniera considerevole, se comparata con i normali processori (Central Processing Unit - CPU), perchè nate con uno specifico orientamento, quello delle grafica, dove centinaia di migliaia di operazioni vengono eseguite in parallelo senza necessità di dover spendere transistor per gestire la logica di controllo. La maggioranza delle operazioni, infatti, è estremamente ripetiva e le condizioni per cui delle parti di codice vengono eseguite o meno sono molto rare. A grandi linee, infatti, possiamo dire che nelle CPU standard, con un numero di core che oggigiorno arriva generalmente ad 8 o poco più, la maggior parte dei transistor è spesa per gestire il controllo (con le cache e la logica di salto), mentre nel caso delle GPU i core sono molto più semplici (effettivamente solo delle unità in virgola mobile molto veloci) e quasi tutta l'area è spesa per implementarli.
L'installazione
Andando per gradi, il primo passo da fare per poter scrivere software in CUDA e quello di installare il toolkit,contenente il compilatore, e, a partire dalla versione 2.3, debugger e visual profiler. Per scrivere programmi non è necessario avere a disposizione una GeForce, e neanche per poterli eseguire in modalità di emulazione, benchè con prestazioni molto limitate: viene eseguito un thread per volta dalla CPU di sistema, e considerando che le applicazioni CUDA possono avere svariate migliaia di thread, è evidente che non si raggiungeranno velocità elevate. É invece ovviamente necessaria una scheda GeForce almeno della serie 8 per eseguire i programmi nativamente.
Tutto il software da installare relativo a cuda si trova sulle pagine della CUDA Zone del sito NVIDIA: http://www.nvidia.com/object/cuda_home.html
Da queste pagine è possibile scaricare, oltre al toolkit, il driver e il kit di sviluppo con gli esempi di codice. Per quanto riguarda i driver, generalmente l'ultima versione ufficiale disponibile nella sezione per le schede grafiche o per le schede della famiglia Tesla, supporta l'equivalente ultima versione di CUDA. Consigliamo comunque di controllare che le versioni siano adeguatamente allineate, e che la versione di driver installata sul proprio sistema sia successiva rispetto a quella fornita in CUDA Zone per la versione di CUDA installata. Il toolkit è invece la parte fondamentale da installare, e al momento in cui scriviamo è già alla versione 2.3, che prenderemo come riferimento per tutto il tutorial. CUDA si evolve molto velocemente ed è possibile che già per la fine dell'anno si arrivi alla versione 3.0. Infine, il kit di sviluppo (SDK) è un pacchetto contenente un elevato numero di esempi di codice, i template per sviluppare le prime applicazioni, ed una serie di librerie di utilità liberamente utilizzabili nei propri programmi. L'installazione è caldamente consigliata per avere un'ottima base di partenza per i proprio lavori.
Sono adeguatamente supportati sia Windows, che Linux. Anche MacOS X è un'ottima piattaforma per lo sviluppo CUDA. In generale, per applicazioni di supercomputing puro, con tanti numeri da macinare su un numero elevato di GPU, è preferito l'uso di Linux, data la sua grande diffusione sui sistemi di supercalcolo. Per applicazioni più casalinghe, comunque, nulla impedisce di adottare le versioni Windows o Mac OS X, dal punto di vista delle caratteristiche pienamente equivalenti. Noi come riferimento prenderemo comunque la versione Linux. Attenzione nella scelta della versione: per ambienti a 64 bit, consigliamo di installare appunto il pacchetto a 64 bit, benchè anche quello a 32 bit possa funzionare senza problemi, per trarre i benefici dal maggiore spazio di indirizzamento offerti dal sistema. Attenzione infine alle distribuzioni: in generale, Debian, OpenSuse, Fedora e Ubuntu sono ben supportate, con appositi pacchetti per ciascuna versione. Ci sono però attualmente delle incompatibilità con le versioni di GCC superiori alla 4.3 (la 4.4 di Fedora 11 per esempio), che richiedono alcuni aggiustamenti per permettere a CUDA di utilizzare la versione di compatibilità del compilatore di sistema normalmente inserita nelle varie distribuzioni.
Scaricati i pacchetti, l'ordine di installazione consigliato è: driver (eventualmente), toolkit e infine SDK. Il driver e il toolkit vanno evidentemente installati con i diritti di root, mentre raccomandiamo invece di installare l'SDK nella propria home con i diritti del proprio utente. Per ambienti simili ad Ubunto che sfruttano sudo, installare con diritti di root significa anteporre ai comandi la particella sudo, appunto. Per altre distribuzioni, basta passare in modalita superuser con il comando su.
Potrebbe essere necessario verificare l'installazione dei pacchetti relativi al g++, automake e GLUT (librerie OpenGL), queste ultime in particolare per permettere di funzionare agli esempi che sfruttano l'interfacciamento tra CUDA e le librerie 3D. Una volta installato il tutto, dovrebbe bastare inserire i path ai binari e alle librerie di cuda. Per farlo, una soluzione molto semplice è quella di editare il file .bashrc nella propria home aggiungendo sulle ultime righe i comandi che seguono, se i pacchetti sono stati intallati nelle directory di default:
per ambienti a 32 bit:
per ambienti a 64 bit:
per ambienti a 64 bit, nei quali si voglia compilare anche in modalità compatibilità a 32 bit:
A questo punto, basta uscire dalla shell, rientrarci e, per verificare il corretto funzionamento di tutto quanto, entrare nella cartella dell'SDK (di default NVIDIA_GPU_Computing_SDK) e da qua in quella “C” che contiene i sorgenti e le librerie comuni a tutti gli esempi, e dare il comando make. Se tutto funzionza, gli esempi verranno compilati e saranno pronti per essere lanciati nella directory C/bin/linux/release e C/bin/linux/emurelease dell'SDK, rispettivamente per le versioni native e le versioni emulate. Anche per il kit di sviluppo consideriamo la versione 2.3, che ha avuto cambiamenti abbastanza significativi nell'organizzazione delle directory poiché integra adesso sia gli esempi in C per CUDA che in OpenCL. Daremo una breve idea nel prosieguo del tutorial delle differenze principali tra CUDA e OpenCL.
Le basi del modello architetturale di CUDA
CUDA lavora, concettualmente, sul modello architetturale riportato nella figura seguente. La GPU accede ad una memoria locale ed è connessa al sistema, e quindi alla CPU, tramite un bus (normalmente, il PCI Express). Il chip grafico, nel modello di CUDA, è costituito da una serie di multiprocessori, denonimati Streaming MultiProcessor. Il numero di multiprocessori dipende dalle caratteristiche specifiche e dalla classe di prestazioni di ciascuna GPU. Ciascun multiprocessore è a sua volta formato da 8 Stream Processors. In questo caso, invece, il numero è fisso, indipendente dalla fascia di mercato del chip grafico. Ognuno di questi processori può eseguire una operazione matematica fondamentale (addizione, moltiplicazione, sottrazione, ecc) su interi o su numeri in virgola mobile in singola precisione (32 bit). In ciascun multiprocessore ci sono anche due unità per funzioni speciali (che eseguono operazioni trascendenti come seno, coseno, inverso ecc.) e, solo per i chip basati su architettura GT200, una singola unità in virgola mobile a doppia precisione (64 bit). Già da questo fattore, è evidente che le prestazioni massime a 64 bit saranno un ottavo rispetto a quelle a 32 bit, supposto di poter sfruttare appieno le capacità computazionali dell'architettura. In un multiprocessore è anche presente una shared memory, accessibile da tutti gli streaming processor, delle cache per le istruzioni e per i dati e, infine, una unità di decodifica delle istruzioni. Un particolare che accenniamo adesso, ma diverrà chiaro probabilmente solo nel prosieguo del corso, è che c'è una sola unità per 8 processori. Siamo in una situazione di tipo Single Instruction, Multiple Data (SIMD), dove un'istruzione viene eseguita per una serie di dati diversi. NVIDIA preferisce chiamare questa situazione SIMT, Single Instruction Multiple Thread, perchè di fatto nel modello di CUDA vengono eseguite le stesse istruzioni da thread diversi. Ancora di più, questa unità decodifica 1 istruzione ogni 4 cicli di clock dei processori. Ciascun processore, invece, di picco può lanciare una istruzione per ogni ciclo di clock. Questo significa che ad ogni istruzione decodificata corrispondono 32 esecuzioni di quella stessa istruzione (8 moltiplicato per 4).
Un'applicazione CUDA è composta da parti seriali, normalmente eseguite dalla CPU di sistema, o host, e da parti parallele, denominate kernel, che vengono invece eseguite dalla GPU, o meglio ancora nei termini usati da NVIDIA, dal device.
Un kernel, nei termini usati da CUDA, è definito come una grid (griglia), e può a sua volta essere decomposto in blocchi, che vengono assegnati, sequenzialmente, ai vari multiprocessori, e rappresentano un parallelismo a grana grossa. All'interno dei blocchi, c'è l'unità di computazione fondamentale, il thread, ad una granularità di parallelismo molto fine. Un thread può appartenere ad un solo blocco, ed è identificato da un indice univoco per tutto il kernel. Per comodità, c'è la possibilità di utilizzare indici bidimensionali per i blocchi e tridimensionali per i thread.
I kernel (le griglie) sono eseguite sequenzialmente tra loro. I blocchi e i thread, invece, sono eseguiti logicamente in parallelo. Il numero di thread fisici in esecuzione in parallelo dipende dalla loro organizzazione in blocchi e dalle loro richieste in termini di risorse rispetto alle risorse disponibili nel device. I blocchi sono pensati per garantire la scalabilità: supponendo di avere un'architettura con solo 2 multiprocessori e un'altra con 4, e un'applicazione decomposta in 8 blocchi, essa potrà essere eseguita su entrambe, ovviamente con tempi (e livelli di parallelismo diversi). Nel caso si renda successivamente disponibile un'architettura con 8 multiprocessori, l'applicazione automaticamente si adatterà ad essa e potenzialmente scalerà ancora con le prestazioni.
La caratteristica fondamentale di CUDA, che rende il modello di programmazione sostanzialmente differente da altri modelli paralleli, normalmente usati dalle CPU, è che per essere efficiente richiede migliaia di thread. Ciò è reso possibile dalla struttura tipica delle architetture grafiche, che impiegano thread molto leggeri e permettono di creare e cambiare i contesti di esecuzione in maniera molto rapida ed efficiente (nessun ciclo di clock).
Dal punto di vista dell'architettura di memoria, a questo stadio ci soffermiamo solo sulle caratteristiche base. La GPU (o device) può accedere alla sua memoria privata, o global memory. L'host ha la possibilità di spostare i dati in questa memoria, primitive. Ciascun multiprocessore può poi accedere alla propria shared memory, che invece non può in alcun modo essere gestita dall'host. I dati allocati in shared memory hanno una visibilità limitata al singolo blocco di thread CUDA. Esistono altre classi di memoria (memoria per le costanti, local memory, memoria per le texture) che vedremo in dettaglio solo successivamente. Di base, per fare funzionare bene un programma CUDA, generalmente basta avere conoscenza di global memory e shared memory.
Si comprende abbastanza facilmente che il modello di CUDA si adatta a specifiche classi di applicazione. In particolare, le caratteristiche principali di queste applicazioni risiedono nella presenza di molte operazioni matematiche (grande intensità aritmetica), elevato grado di parallelismo (le stesse operazioni vengono ripetute per una grande quantità di dati), elavata richiesta di banda di memoria e condizioni di controllo molto limitate. Fortunatamente, possiedono queste caratteristiche algoritmi che appartengono ai campi più disparati: da applicazioni per gestione di oleodotti e individuazione di petrolio o gas, a crittografia, da chimica all'analisi di immagini e audio. Gli algoritmi di codifica e decodifica audio e video, analisi del traffico di rete fino alla gestione di database si sposano anche molto bene con il modello di CUDA.
Il primo programma
Per iniziare a sperimentare con CUDA vi consigliamo di creare una apposita cartella e creare al suo interno un file. È possibile utilizzare qualsiasi tipo di editor.
Il concetto base di CUDA, e di qualsiasi applicazione di tipo GPGPU, è che si sta eseguendo un “offloading” di una parte di computazione ad un coprocessore, rispetto al processore (la CPU) del sistema. Tale coprocessore, che viene chiamato device, accede ad una memoria privata (la global memory) nella quale devono essere mossi i dati prima di poter eseguire la computazione. La CPU e la memoria di sistema vengono identificati come host. L'host, prima di far iniziare una computazione al device, muove i dati nella global memory. Al termine della computazione, i risultati sono copiati dalla memoria del device alla memoria dell'host.
Per eseguire queste operazioni, sono necessarie delle apposite primitive per allocare la memoria sul device e eseguire le copia dall'host al device. La memoria sull'host è allocata nelle modalità standard del C, con variabili automatiche o, nel caso di allocazione dinamica, con la classica malloc.
Per eseguire l'allocazione della memoria sul device, CUDA offre la seguente primitiva:
Dove va specificato il nome del puntatore all'area di memoria che si sta allocando (attenzione al puntatore di puntatore), e la grandezza in bytes dell'area di memoria da allocare.
Per eseguire le operazioni di copia invece si usa la primitiva:
Dove si specificano puntatore ad area di memoria sorgente, puntatore ad area di memoria destinazione, la grandezza in byte dei dati da copiare, e la direzione dell'operazione di copia. La direzione può essere uno dei seguenti tipi:
La prima copia i dati dall'host al device, la seconda dal device all'host (per esempio, per leggere i risultati) e la terza i dati in due aree di memoria diverse sul Device.
Per liberare l'area di memoria usata, si usa invece la primitiva:
nella quale basta specificare il puntatore all'area di memoria da liberare.
Esiste anche una primitiva:
che serve ad inizializzare ad un valore voluto (value) un'area di memoria di grandezza (in numero di elementi) count.
Con queste primitive, è possibile realizzare il primissimo programma, che esegue una copia dei dati dall'host al device, una copia di dati sul device, e poi riporta tali dati sull'host. Apriamo il nostro editor e chiamiamo tale file cudamemcpy.cu. I file sono tratti da uno dei training disponibili sul sito di NVIDIA, per consistenza con la documentazione prodotta dalla società stessa.
// include di sistema
#include <stdio.h>
#include <assert.h>
// Semplice funzione per verificare gli errori in cuda
void checkCUDAError(const char *msg);
//dichiariamo il main
int main( int argc, char** argv)
{
//puntatore all'area di memoria sull'host
float *h_a;
//dimensioni di tale area di memoria
int dimA;
//variabile contatore
int n;
// puntatori alle aree di memoria sul device
float *d_a, *d_b;
// allocazione ed inizializzazione della memoria sull'host
// definizione delle dimensioni
dimA = 8;
//esecuzione della malloc
h_a = (float *) malloc(dimA*sizeof(float));
//inizializzazione
for (n=0; n<dimA; n++)
{
h_a[n] = (float) n;
}
// Allocazione della memoria sul device
size_t memSize = dimA*sizeof(float);
cudaMalloc( (void**)&d_a, memSize );
cudaMalloc( (void**)&d_b, memSize );
// Memcopy dall'host al device
cudaMemcpy( d_a, h_a, memSize, cudaMemcpyHostToDevice );
// Memcopy da device a device
cudaMemcpy( d_b, d_a, memSize, cudaMemcpyDeviceToDevice );
// Cancellazione dei dati nell'area di memoria sull'host
for (n=0; n<dimA; n++)
{
h_a[n] = 0.f;
}
// Copia dal device all'host
cudaMemcpy( h_a, d_b, memSize, cudaMemcpyDeviceToHost );
// Check for any CUDA errors
checkCUDAError("cudaMemcpy calls");
// verify the data on the host is correct
for (n=0; n<dimA; n++)
{
assert(h_a[n] == (float) n);
}
// Liberazione della memoria sul device
cudaFree( d_b );
cudaFree( d_a );
// Verifica di eventuali errori in CUDA
checkCUDAError("cudaFree");
// liberazione della memoria dell'host
free(h_a);
// Se il programma arriva fin qua, tutto corretto!
printf("Corretto!\n");
return 0;
}
void checkCUDAError(const char *msg)
{
cudaError_t err = cudaGetLastError();
if( cudaSuccess != err)
{
fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString( err) );
exit(-1);
}
}
Per compilare l'applicazione, basta lanciare nvcc con il nome del file. NVCC offre una serie di opzioni che possono risultare molto utili nello sviluppo di un programma:
- nvcc <nomefile>.cu [-o <executable>
Comando standard per compilare il codice in modalità rilascio, con eventuale nome dell'eseguibile specificato dall'utente
- nvcc -g <nomefile>.cu
Compila il file in modalità di debug, rendendo utilizzabile sia a gdb per il codice sull'host, sia a cuda-gdb per il codice che gira sul device
- nvcc -deviceemu <filename>.cu
Compila il codice in modalità emulazione
Tutto il codice gira sulla CPU
- nvcc -deviceemu <filename>.cu
Compila il codice in modalità emulazione, ma con i simboli di debug
Permette di debuggare il codice con gdb facendolo girare completamente sulla CPU